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

[prev in list] [next in list] [prev in thread] [next in thread] 

List:       llvm-bugs
Subject:    [LLVMbugs] [Bug 21685] New: NVPTX emits incorrect PTX for weak_odr (and potentially other linkage ty
From:       bugzilla-daemon () llvm ! org
Date:       2014-11-27 18:56:27
Message-ID: bug-21685-206 () http ! llvm ! org/bugs/
[Download RAW message or body]

--1417114588.F46db0.18102
Date: Thu, 27 Nov 2014 12:56:28 -0600
MIME-Version: 1.0
Content-Type: text/plain; charset="UTF-8"

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@nondot.org
          Reporter: wujingyue@gmail.com
                CC: justin.holewinski@gmail.com, llvmbugs@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.

--1417114588.F46db0.18102
Date: Thu, 27 Nov 2014 12:56:28 -0600
MIME-Version: 1.0
Content-Type: text/html; charset="UTF-8"

<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&#64;nondot.org
          </td>
        </tr>

        <tr>
          <th>Reporter</th>
          <td>wujingyue&#64;gmail.com
          </td>
        </tr>

        <tr>
          <th>CC</th>
          <td>justin.holewinski&#64;gmail.com, llvmbugs&#64;cs.uiuc.edu
          </td>
        </tr>

        <tr>
          <th>Classification</th>
          <td>Unclassified
          </td>
        </tr></table>
      <p>
        <div>
        <pre>Given the following IR

target datalayout = &quot;e-i64:64-v16:16-v32:32-n16:32:64&quot;
target triple = &quot;nvptx64-unknown-unknown&quot;

define weak_odr void &#64;foo() {
entry:
  ret void
}

!nvvm.annotations = !{!0}

!0 = metadata !{void ()* &#64;foo, metadata !&quot;kernel&quot;, i32 1}

The NVPTX backend emits

/
// Generated by LLVM NVPTX Back-End
//

.version 3.2
.target sm_35
.address_size 64

        .weak   foo
                                        // &#64;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 &lt;typename T&gt;
__global__ void foo(T x) {
  bar&lt;T&gt;(x);
}

template __global__ void foo&lt;float&gt;(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>

--1417114588.F46db0.18102--

_______________________________________________
LLVMbugs mailing list
LLVMbugs@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmbugs

[prev in list] [next in list] [prev in thread] [next in thread] 