[llvm-bugs] [Bug 27270] New: Macro protected __device__ function and identical host function collide in global

[via llvm-bugs]

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: unassignedclangbugs at nondot.org
          Reporter: crtrott at sandia.gov
                CC: llvm-bugs at 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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20160407/0c3898e3/attachment-0001.html>