<html>
<head>
<base href="http://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 --- - NVPTX emits incorrect PTX for weak_odr (and potentially other linkage types)"
href="http://llvm.org/bugs/show_bug.cgi?id=21685">21685</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>NVPTX emits incorrect PTX for weak_odr (and potentially other linkage types)
</td>
</tr>
<tr>
<th>Product</th>
<td>libraries
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>All
</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>Backend: PTX
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>wujingyue@gmail.com
</td>
</tr>
<tr>
<th>CC</th>
<td>justin.holewinski@gmail.com, llvmbugs@cs.uiuc.edu
</td>
</tr>
<tr>
<th>Classification</th>
<td>Unclassified
</td>
</tr></table>
<p>
<div>
<pre>Given the following IR
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-unknown-unknown"
define weak_odr void @foo() {
entry:
ret void
}
!nvvm.annotations = !{!0}
!0 = metadata !{void ()* @foo, metadata !"kernel", i32 1}
The NVPTX backend emits
/
// Generated by LLVM NVPTX Back-End
//
.version 3.2
.target sm_35
.address_size 64
.weak foo
// @foo
.weak .entry foo(
)
{
// BB#0: // %entry
ret;
}
I found a similar issue was discussed in
<a href="http://lists.cs.uiuc.edu/pipermail/llvmdev/2012-June/050610.html">http://lists.cs.uiuc.edu/pipermail/llvmdev/2012-June/050610.html</a>. While it's
possible for us to workaround, it prevents explicit template instantiation
which is something we need in longer term. For instance,
template <typename T>
__global__ void foo(T x) {
bar<T>(x);
}
template __global__ void foo<float>(float x);
I looked at the code of AsmPrinter. It looks like we can solve this issue at
least for GlobalVariables by inheriting function EmitGlobalVariable. A lot of
logic there (such as thread local variable) is unnecessary/non-existent for
NVPTX anyway.
Jingyue</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>