[llvm-bugs] [Bug 26340] New: Cuda device constant doesn't work

[via llvm-bugs]

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

            Bug ID: 26340
           Summary: Cuda device constant doesn't work
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: -New Bugs
          Assignee: unassignedclangbugs at nondot.org
          Reporter: crtrott at sandia.gov
                CC: llvm-bugs at lists.llvm.org
    Classification: Unclassified

Created attachment 15729
  --> https://llvm.org/bugs/attachment.cgi?id=15729&action=edit
Reproduction code without extern symbols

It looks like cuda __device__ __symbols__ don't work. 
Effectively it looks like the host side functions such as cudaGetSymbolAddress
or cudaMemcpyToSymbol have an issue at link time, where it says something with
the same name as the cuda symbol is an undefined reference. The compilation it
self gets through. If you cheat the linker by adding another .o file with a
global host symbol of the same name the linking works, but at runtime the
aforementioned functions produce an error of invalid device symbol. 

I attach a reproduction example with build scripts for NVCC (which works) and
Clang. In the build script for clang you have to comment in the line with the
workaround c.cpp file.

I also attach a similar example, where a constant symbol is declared extern.
This requires compilation with -rdc true for NVCC. It fails right now for
similar reasons I assume as the first example, but I'd like to get that working
as well. I'll attach that later.

-- 
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/20160127/a32c8186/attachment.html>