[llvm-bugs] [Bug 27276] New: global functions in anonymous namespace cause runtime cuda error

[via llvm-bugs]

https://llvm.org/bugs/show_bug.cgi?id=27276

            Bug ID: 27276
           Summary: global functions in anonymous namespace cause runtime
                    cuda error
           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 16185
  --> https://llvm.org/bugs/attachment.cgi?id=16185&action=edit
Reproducer

When defining a global function in anonymous namespace, it causes a "invalid
device function" error if the function is called from the same compilation
unit. 

i.e. in simplified form this doesn't work: 
namespace {
__global__ void foo() {}

void bar() {
   foo<<<1,1>>>();
}
}

Attached is a reproducer. The output with NVCC is:
////
[crtrott at apollo test_kernel]$ ./a.out 
Error PreKernel: no error
Error PostKernel: no error
Error PostSync: no error
Arch: 350
////

The output with clang is:
////
[crtrott at apollo test_kernel]$ ./a.out 
Error PreKernel: no error
Error PostKernel: invalid device function
Error PostSync: no error
Arch: 9
////

Compile with ./build_clang after adjusting the paths in the build script. 
For a workaround (i.e. remove the namespace) compile with 
./build_clang -DWORKAROUND

-- 
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/20160408/b9ecbe78/attachment-0001.html>