http://llvm.org/bugs/show_bug.cgi?id=21685
Bug ID: 21685
Summary: NVPTX emits incorrect PTX for weak_odr (and
potentially other linkage types)
Product: libraries
Version: trunk
Hardware: PC
OS: All
Status: NEW
Severity: normal
Priority: P
Component: Backend: PTX
Assignee: [email protected]
Reporter: [email protected]
CC: [email protected], [email protected]
Classification: Unclassified
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
http://lists.cs.uiuc.edu/pipermail/llvmdev/2012-June/050610.html. 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
--
You are receiving this mail because:
You are on the CC list for the bug.
_______________________________________________
LLVMbugs mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/llvmbugs
