https://llvm.org/bugs/show_bug.cgi?id=27270
Bug ID: 27270
Summary: Macro protected __device__ function and identical host
function collide in global
Product: clang
Version: trunk
Hardware: PC
OS: Linux
Status: NEW
Severity: normal
Priority: P
Component: CUDA
Assignee: unassignedclangb...@nondot.org
Reporter: crtr...@sandia.gov
CC: llvm-bugs@lists.llvm.org
Classification: Unclassified
Created attachment 16183
--> https://llvm.org/bugs/attachment.cgi?id=16183&action=edit
Reproducer
If I have a __device__ function and a __host__ function with only one of them
visible at the time through protection of __CUDACC__ && __CUDA_ARCH__ the
compile will stop on global functions because it thinks I want to call the
__host__ function. My current suspicion is that during the host compilation
pass it still sees the global function, and it sees the __host__ function and
thinks I call the latter from the former. Here is the simplified pattern:
#ifdef __CUDA_ARCH__
__device__ inline void foo() {}
#else
inline void foo() {}
#endif
__global__ bar() {
foo();
}
In the attached reproducer this produces the following error with clang (while
it works with nvcc):
/////////////
main.cpp:21:43: error: no matching function for call to 'c_func'
a[blockIdx.x*blockDim.x+threadIdx.x] = c_func();
^~~~~~
main.cpp:11:12: note: candidate function not viable: call to __host__ function
from __global__ function
inline int c_func() {return 1;}
^
1 error generated.
////////////
The attached reproducer has two workaround solutions in it. To use them build
with
./build_clang -DWORKAROUND=1
./build_clang -DWORKAROUND=2
otherwise just use
./build_clang
to get the error. You will need to change the paths in the build script though.
--
You are receiving this mail because:
You are on the CC list for the bug.
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
