[LLVMbugs] [Bug 21685] New: NVPTX emits incorrect PTX for weak_odr (and potentially other linkage types)

[bugzilla-daemon at llvm.org]

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: unassignedbugs at nondot.org
          Reporter: wujingyue at gmail.com
                CC: justin.holewinski at gmail.com, llvmbugs at cs.uiuc.edu
    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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20141127/28fa8c79/attachment.html>