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

    <tr>
        <th>Summary</th>
        <td>
            C++/Cuda __global__ kernels pre-declared in header results into nv-link error in separable compilation (gpu-rdc)
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

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

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

<pre>
    When enabling compilation with relocatable device code flag (--gpu-rdc), with cuda code with header and implementation files, I ran into a linking error:
```
nvlink error: Function A() not declared __global__ in all source files 
nvlink fatal: merge_elf failed
``` 

The error can be reproduced with the following set-up:
```
// bug.cuh
__global__ void foo();

// bug.cu
__global__ void foo() {}

// main.cu
#include "bug.cuh"

int main() { foo<<<1, 1, 1>>>(); }

>clang++ bug.cu main.cu --offload-arch=sm_89 -fgpu-rdc --offload-new-driver -lcudart
```

@jhuber6's insight:
Solution is to have .extern .entry instead of .extern .func
  ```
  >clang++ -x cuda main.cu --offload-arch=sm_89 --offload-device-only -o - -S -fgpu-rdc
  
 .version 8.5
  .target sm_89
  .address_size 64
  
  .extern .func _Z3foov
  ()
  ;
  .global .align 8 .u64 __clang_gpu_used_external[1] = {_Z3foov};
  ```


Workaround, enable -foffload-lto, which . 
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJyEVE2P4ygQ_TXkUsLCOB_OwYekM5H2PCuNtBcLQ9lmmkDEh7t7f_3KJumke6VdCSUyVL2qB6-eCEEPFrEhmyPZnFYixdH5ZkB3cTqsOqc-ml8jWkArOqPtANJdrtqIqJ2FNx1H8GicFFF0BkHhpCWCdAqhN2IAwmtKh2uiXknC94S_5CSZlMhhy-eIQqEHYRXoy9XgBW3MJXptMMxpf4AXFrSNDgQYbV_nZtB750l1IOxAtuy22MFOc8DnKZyTlQvagfCa8D1YF0GhNMKjgrYdjOuEaVvQFoQxEFzyEnNteOD1Igoz413QD9ii6aEX2qB6Lj_HE3b4c8RcH6Sw0CF4vHqnkkSVKccRoXfGuLeZSMBI0_XfTAg_E36GLg2FTCNhh6dmJ6cV9M5lTqQ65spfUv4rA8juSHanL1kXoW1OI7zSVpqkEAjn9wY4z-HaxiX2gbTgVi95lfOL3X6qH7d1bxMeRasf0gg7EH4k_Hhr-d4DUOr63jihqPByJNUpXNp6D7S_6ekpwOIbVV5P6IGaWVs-fr9IdiBr9ntMHfot4bsA2gY9jDHf-U9n0qIQHSA6GMWEUOB7RG-hQBv9xxwfUShw_eOkT1YSdgB4LgXwjRd9z3r_P2Kf23mMqLPmA6gDCvTng3auwA5QTOjD3HNdbJbNIgo_YIQFLu8IpTyG0Ab9N8J2_Zn7lQK0f1W9c1M-zu-UeRxzcFYQFMLowUINRdquoW0Xku1wTW0KqNoMKQzZHEuyOQGpTrMy7uC70x3v-8Owwy_nX4V3yapZNIvZIND-fiEmusU6Ri1HKGClmkrtq71YYVPuqopt2Z7tVmNTdltV9dtyy2u-Y1IhK7tSqs1u322k7OqVbjjj65KzsqzWNd8VKNR-zevtdsdELypG1gwvQpvCmOlSOD-sdAgJm7KseclXRnRowuKWnFt8g-V0HovNaeWbOYl2aQhkzYwOMTxgoo4Gm5csCcLPL7MinmbzFb1FE-DqkX5ak7Z3a_QYkokhG6Cd6MPg5qCAV-GXO3u2Z8Lrh_WukjfNGOM1zIpfpn3QcUxdId2F8PPc6O2PXr37jTISfl7oBcLPN_5Tw_8JAAD__4k_7ao">