<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class="">
Hello all,
<div class=""><br class="">
</div>
<div class="">I believe I have found a bug in overload resolution for cuda code. I am currently awaiting permission to post to the bug tracker. </div>
<div class=""><br class="">
</div>
<div class="">The following code doesn’t compile with newer versions of clang:</div>
<div class=""><br class="">
</div>
<div class="">
<div class="">template<class T></div>
<div class="">__device__ __host__ int foo(T *x) {</div>
<div class=""> return 1;</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">__device__ int foo(int *x) {</div>
<div class=""> return 2;</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">__host__ int foo(long *x) {</div>
<div class=""> return 3;</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">__device__ __host__ int bar() {</div>
<div class=""> auto long_val = 1l;</div>
<div class=""> return foo(&long_val);</div>
<div class="">}</div>
</div>
<div class=""><br class="">
</div>
<div class=""><br class="">
</div>
<div class="">clang++ -O2 -g -x cuda --cuda-gpu-arch=sm_61 -std=c++14 -o main -c main.cpp give me:</div>
<div class="">
<div class="">error: reference to __host__ function 'foo' in __host__ __device__ function</div>
<div class=""> return foo(&long_val);</div>
<div class=""> ^</div>
<div class="">main.cpp:10:14: note: 'foo' declared here</div>
<div class="">__host__ int foo(long *x) {</div>
</div>
<div class=""><br class="">
</div>
<div class="">I believe that the issue is at <a href="https://github.com/llvm/llvm-project/blob/8224c5047e9cef2db4b0e31427cdf90a2568a341/clang/lib/Sema/SemaOverload.cpp#L9860" class="">https://github.com/llvm/llvm-project/blob/8224c5047e9cef2db4b0e31427cdf90a2568a341/clang/lib/Sema/SemaOverload.cpp#L9860</a></div>
<div class=""><br class="">
</div>
<div class="">It’s possible that IdentifyCUDAPreference will return CFP_HostDevice for valid overloads, but this code doesn’t erase the wrong side candidates in that case. Then because the wrong side candidate is an exact match, minus its host device attributes,
clang picks it as the best overload. </div>
<div class=""><br class="">
</div>
<div class="">If I rewrite those lines as:</div>
<div class=""><br class="">
</div>
<div class="">
<div class=""> bool ContainsSameSideCandidate =</div>
<div class=""> llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {</div>
<div class=""> // Check viable function only.</div>
<div class=""> if (Cand->Viable && Cand->Function) {</div>
<div class=""> auto MatchType = S.IdentifyCUDAPreference(Caller, Cand->Function);</div>
<div class=""> return MatchType == Sema::CFP_HostDevice ||</div>
<div class=""> MatchType == Sema::CFP_SameSide;</div>
<div class=""> }</div>
<div class=""> return false;</div>
<div class=""> });</div>
</div>
<div class=""><br class="">
</div>
<div class="">My code compiles again. I can submit a bug report once I am approved, but I figured I would post here in the mean time. </div>
<div class=""><br class="">
</div>
<div class="">-Drew </div>
</body>
</html>