[llvm-bugs] [Bug 50316] New: [NVPTX] Invalid PTX when initializing struct from shared-memory pointer

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=50316

            Bug ID: 50316
           Summary: [NVPTX] Invalid PTX when initializing struct from
                    shared-memory pointer
           Product: libraries
           Version: 11.0
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: Backend: PTX
          Assignee: unassignedbugs at nondot.org
          Reporter: bugs-llvm at fabian-knorr.info
                CC: llvm-bugs at lists.llvm.org

Aggregate-initializing a struct from an immediate pointer-to-shared-memory
produces invalid PTX:

    struct wrap { int *mem; };

    __global__ void kernel() {
        __shared__ int mem;
        wrap ptr{&mem};
    }

Compiled with 

    clang++ -c bug.cu --cuda-gpu-arch=sm_75 -std=gnu++17 -U__FLOAT128__
-U__SIZEOF_FLOAT128__

results in

        // .globl       _Z6kernelv
    // _ZZ6kernelvE3mem has been demoted
    .global .align 8 .u64 __const_$__Z6kernelv_$_ptr[1] =
{generic(_ZZ6kernelvE3mem)};

    .visible .entry _Z6kernelv()
    {
        .local .align 8 .b8     __local_depot0[8];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .b64       %rd<4>;
        // demoted variable
        .shared .align 4 .u32 _ZZ6kernelvE3mem;
        mov.u64         %SPL, __local_depot0;
        cvta.local.u64  %SP, %SPL;
        mov.u64         %rd1, __const_$__Z6kernelv_$_ptr;
        cvta.global.u64         %rd2, %rd1;
        ld.u64  %rd3, [%rd2];
        st.u64  [%SP+0], %rd3;
        ret;
    }

which ptxas rejects:

    ptxas /tmp/bug-26fed7.s, line 11; fatal   : Invalid initial value
expression
    ptxas fatal   : Ptx assembly aborted due to errors

This issue does not appear with -O1/-O2/-O3. In the unoptimized case it can be
worked around by storing the pointer in a temporary first or by adding an
explicit constructor to the struct.

    * Clang versions tested: 11.1.0, 10.0.1
    * CUDA versions tested: 11.3, 10.2
    * Systems tested: x86_64 host and sm_61 / sm_75 GPUs

-- 
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/20210512/990a4226/attachment.html>