<html>
<head>
<base href="https://llvm.org/bugs/" />
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW --- - Cuda device constant doesn't work"
href="https://llvm.org/bugs/show_bug.cgi?id=26340">26340</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Cuda device constant doesn't work
</td>
</tr>
<tr>
<th>Product</th>
<td>clang
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>Linux
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>normal
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>-New Bugs
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedclangbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>crtrott@sandia.gov
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr>
<tr>
<th>Classification</th>
<td>Unclassified
</td>
</tr></table>
<p>
<div>
<pre>Created <span class=""><a href="attachment.cgi?id=15729" name="attach_15729" title="Reproduction code without extern symbols">attachment 15729</a> <a href="attachment.cgi?id=15729&action=edit" title="Reproduction code without extern symbols">[details]</a></span>
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.</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>