<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/54934>54934</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            Compiling empty CUDA kernel with RDC enabled fails in PTXAS
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            cuda
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
            yxsamliu
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          jhuber6
      </td>
    </tr>
</table>

<pre>
    The following CUDA kernel is failing to compile when RDC-mode is enabled because of PTXAS.

```
template <typename T>                    
__global__ void EmptyKernel(void) { }    
                
void foo() {                        
  EmptyKernel<void><<<1, 1>>>();    
} 
```

This function will create the following LLVM-IR
```
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

$_Z11EmptyKernelIvEvv = comdat any

@gpu.used.external = appending global [1 x ptr] [ptr @_Z11EmptyKernelIvEvv]
@llvm.compiler.used = appending global [1 x ptr] [ptr @gpu.used.external], section "llvm.metadata"

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define weak_odr dso_local void @_Z11EmptyKernelIvEvv() #0 comdat {
entry:
  ret void
}

attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_70" "target-features"="+ptx75,+sm_70" }

!llvm.module.flags = !{!0, !1, !2, !3}
!nvvm.annotations = !{!4}

!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 11, i32 5]}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 4, !"nvvm-reflect-ftz", i32 0}
!3 = !{i32 7, !"frame-pointer", i32 2}
!4 = !{ptr @_Z11EmptyKernelIvEvv, !"kernel", i32 1}
```
Which makes the following PTX
```
//
// Generated by LLVM NVPTX Back-End
//

.version 6.0
.target sm_70
.address_size 64

        // .weak        _Z11EmptyKernelIvEvv            // -- Begin function _Z11EmptyKernelIvEvv
.weak .entry _Z11EmptyKernelIvEvv
()
;
.global .align 8 .u64 gpu.used.external[1] = {_Z11EmptyKernelIvEvv};
                                        // @_Z11EmptyKernelIvEvv
.weak .entry _Z11EmptyKernelIvEvv()
{


// %bb.0:                               // %entry
        ret;
                                        // -- End function
}
```
Which fails when compiled with PTXAS
```
$ ptxas "-m64" "-O3" "--gpu-name" "sm_70" "--output-file" "empty.cubin" "empty.s" "-c"
ptxas empty.s, line 13; fatal   : Parsing error near '.used': syntax error
ptxas fatal   : Ptx assembly aborted due to errors
```

This seems to have started with 0424b5115cffa but is probably not the root cause.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJylV1tzmzgU_jX4RQMDAt8e_ODY6U6n3d1Om83u7ItHgLDVCImRBLH763skgY0Tp91LhoAQ53w693Ocy_K0ejhQVEnO5TMTe7T5Y7tGT1QJyhHTqCKM220jUSHrhnGKng9UoM_bTVjLkloaKkjOaYlyWpBWUyQr9Onhr_WXKIi3Qbzu77O4v9yroXXDiaEoSDfm1ABETdFDkN6jG3-eZbfbc5kTvtuhTrIS3deNOX1wkgZ4YbcCvETB_A7-txe222AOoZISOAeuN_4GlPFx6cYdl97Dyl9JgDcocTvucrBBendBsELdNIW_PxystVtRGCYFemaco0JRayFz5Z-PHx9_Dd9_vm1UovbUoJIYwslJtgasC4diTEM2y4J0PctCloBo6RruYZfM7GoWdimGRYpD4XbcG9BjfAVrFGs4HSBF15gj4ImOlYyERVuSM0N_x9nu7yQZ2e19d991DgBiCaRERJyuOLJ437QRxFAZ0aMBHsIdOWkgQkqrvg8BFEzvEnREjVHBdGvfYIWA_daBQHGG57yroz6QlTvo3-G_Es-Cg-c19X4DE7gjamqIdcNLk0BAvBt8vDZGabA0GEN0FEwsDKpbbRol94pqjYRkApKPwkLRolXarloBYVAi2RghBfWwJa0s2TMlTztZKlRqueOyAD1cmL9llz70cRoP_oA08IggizqBbEPwK_C_C_khlMdaEVCE5a2h2oO5CIGE-h96WUNWCmpC2AAxmNoaMgVTY8Ih27H9XjMRcronPOzA-lKFz6w0hzNh3JMJGRoFDgbvhjUZERjV0p5GG1I8hSCg8Uh5W1VUhZp9o2fyRU_rsyEsmvb8Sde7eXz9uYLUbUHbM02AIYaO8ylECywvHNemDHDi40eWLadRxcle9xmXWOfgJLbhBs-kf-L-mV6QcCI6wCBCSENsqL2AyF4fGo8oWIrRgAoiftl-QI9UacBxamxsOmBIDqDrk8NyeIHsampzYiRM8gI7uWA_FweidoOZe_5szIxfMGcXZqtkqGjFwWdhZb6NIOIxRPoCYn6BeBViPT8e82cj_h-VmTPqU9-TznDJBe66YP95YMUB1eQJUue6zEMDvd0s8Dt7jdboFyqogkYBDfjk2gP67RHY0Z2N6XtR3mB096jzXkWzqAeP-krvo9NvkbK0SeuchKApjOMmXvYiRLb0wOvNej_upZ48DNEd3TNx6Xc3DeoFsNAocgXpB2R9wx2qbM_bV_OIcLYXaIGidpahGzX8LnGR7OvWzUPAgQPqW4PCq8HBK_tWuPxD7cZ6DdX5Onn9KXia5-DH9GfiXeh9jR88CRX-v2oI7oQwOzvzZY-4FfJ2rNR-kuy7cQkzjzn4wfGNyM-gHx-JbTI4rN184pa_p8MqBNeGdpIcyvqoLochzENNC5UCDuv3qDV2VLQ5E1c7euApzv3bn3z-vEGugUHVhY5e2XnLmgOM_4lATkH-UqWkQoISqBd47uINnpZCn4QhR08wxr5CMUdEtKZ1zk-I5FLZ7C5basdwx6h_NklqSmttyQ-kowj6m4NwJo4znOXTJJkWVUUQ9G07w0Pvy4k9DXqGK0VKwsLN89GkXKXlMl2SiWGG09XGecxpae1x9ZPBnQA_Ds4_CrynIdedZyet4quDMY0dfHz47IGjze1QBi-2-fUP246_QmmHV6Z1a1vpu2m2TLPJYZUvllm6iEtoXPM0m6XJIiHTCi-yxazA83k84SSnXK8gs8GDw2w63U7YCscwGmTJNJnHSZZEy2WygJ3ZrMqmWZyWkK20BpEj14al2k_UykmTt3ttB0imjb58BCdBZaG0P-l01KTmrO1PI605SLX6CvpRNZs4NVZOh-_xydYi">