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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] Creating relocatable object (-r) from rdc objects (-fgpu-rdc) fails with lld error undefined protected symbol referenced by __clang_gpu_used_external
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            lld
      </td>
    </tr>

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

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

<pre>
    Hi @yxsamliu, I'm using clang -r to lower bitcode in objects that I compiled -fgpu-rdc to AMDGPU device code before link-time so that I can link without -fgpu-rdc, as I described in #77018. I am hitting lld error undefined protected symbol due to an extern device variable. I reference the variable in host code, but not device code, so why does that stop me from lowering the bitcode? I found your [commit](https://github.com/llvm/llvm-project/commit/0424b5115cffad73a0f6e68affed603a7ed9a692) which created the __clang_gpu_used_external symbol that appears in my error message, and I read your phabricator [discussion](https://reviews.llvm.org/D123441), but I still don't understand why you have to "Emit a global array containing all external kernels or device variables used by host functions and mark it as used..." (see your CodeGenModule.cpp [comment](https://github.com/llvm/llvm-project/blob/main/clang/lib/CodeGen/CodeGenModule.cpp#L924-L925)). Here's a tiny reproducer:

```
$ make
clang -O2 -fgpu-rdc --offload-arch=gfx942 -x hip -c -o alpha1.o alpha1.c
clang -O2 -fgpu-rdc --offload-arch=gfx942 -x hip -c -o alpha2.o alpha2.c
clang -r --hip-link -O2 -fgpu-rdc --offload-arch=gfx942 -o alpha_no_bitcode.o alpha1.o alpha2.o
lld: error: undefined protected symbol: globalfoo
>>> referenced by /var/tmp/pozulp1/alpha1-gfx942-637988.out.lto.o:(__clang_gpu_used_external)
clang: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)
make: *** [makefile:12: alpha_no_bitcode.o] Error 1
$ cat alpha1.c
#include <hip/hip_runtime.h>
#if defined(__HIP_DEVICE_COMPILE__)
extern __device__ int globalfoo;
#else
extern int globalfoo;
#endif
__device__ int alpha_add_impl(int x);
__global__ void alpha_kernel(int *x) { *x = alpha_add_impl(*x); }
void alpha_add(int *x) {
 alpha_kernel<<<1,1>>>(x);
    globalfoo++;
}
$ cat alpha2.c
__device__ int alpha_add_impl(int x) { return x + 1; }
$ cat makefile
COMPILER=/path/to/clang
COMPILE_FLAGS=-O2 -fgpu-rdc --offload-arch=gfx942 -x hip
RELOCATABLE_FLAGS=-r --hip-link -no-hip-rt -O2 -fgpu-rdc --offload-arch=gfx942

all: alpha_no_bitcode.o

.c.o:
    $(COMPILER) $(COMPILE_FLAGS) -c -o $@ $<

alpha_no_bitcode.o: alpha1.o alpha2.o
    $(COMPILER) $(RELOCATABLE_FLAGS) -o $@ $^

clean:
    rm -f *.o
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJysVluP4yoS_jXkpWQrxk5iP-Qh15lIPZrR2curhQHHbGOwAKc7--tXYMdJprdXfVZHQgnmUpevqj6KWCvOivM1WmzRYj8jvWu0WXf6373sZpVm1_V3ASibX98taaXoEd7BCeFVC70V6gxUEnWGyIDTIPUbN1AJRzXjIBTo6l-cOguuIQ5OQHXbCckZRPW56yPDqL-1-bH_9usfwPhFUA7hasVrbThIoV4jJ1oOVk8yiArr8CZco3t3l-UtIxZOwLilRlSceRMQTlereZLHcALSQiOc82ZLyYAbow30ivFaKM6gM9px6jgDe20rLYH13BtIFPB3x4262XghRpBKci_T8JobrigH19x3vOZGWxfc8YZVvQOl3aOXftlqeGuuwDQfQbJOd9ByqI1uBzy9tV70CCtKj3CCWveKwVX3BtBiS3XbCocWe4TzxrnOonSD8BHh41m4pq9iqluEj1Jebn9RZ7QPDcLH8TI-zjOcVYskWdC6JmyVknm95Muc1DVny3lKVpwVZFlghAt4awRtgBpOPFzevLIMmVCeu77sLWflABmRNzCDe6TrODHWw9NexwC03FpyDnAQxQKkZPSta0hlBCVOBz-ZsLS3Vmj133w1_CL4m429f7E2Z4SP-wSnWZYgXNxicALrhJTAtEJ45UL0jXVesQ_EVffQkEsIO8L40AoHBM5SV0QCMYZcgWrliFA-LERKmNx85UZxaUGb39PEggcEquuQEnWvqBNa2eBuS8wreC3DqTiOEcaAcG45H0DYaca_cfVDs17ymHbdLeRc_b8xr6SuED62RCifAT5w_pDwi6O6--yuGOH0pcBZ9FLgRcC0iOE7NxzhlQUCTqgrGN4ZzXrKjTdovkfz2-9yPo7hE2fQklc-fI0s8hM_UEMU6bqWmrCIGNqgdH-u34sMQ_QOjeggohBpILJrSBJPE_pXyMPxNHmWZyCKGtFFgYC-KH0UVSpdjiUc_242jvWgRkqG0s1QF37yOTf53SEtaz1eRulhGHdOCkmH8PFCDMJH13YIHwdqTxA-DkZEg6HRMl0VeR7r3sXS6ViHfMo_LWsf_zs0T1aTlp2pGkDyeerTvCaB-D1pA38XAzFC4hO9txyii684n_JCXTQlvj4mDSFN0g0gvBmGLwC_WAvpNxIctH5AGS32cAgUk9xzjnoWesoVhFOhqOwZB5TuGuFBakRXml75tyduPKrTyRrGkARwvp9-lfvDP0-7Q7n7-ePX6eVQlpPd46NRlgMflCUI5R6ilm4nsVxa_nTp85OKiXr4-k3wAABhrBRtJxHO_eK7t-Z2vSwHkWUJFy3YeGMgrvE8wht_BdBqG-aA0v1HyeMplG4BrfaD8AeJhLGP4oZTzzrT3TAShHfJlL8I509mA8ADFnjrx4TITf1TbKe6_SpEwV_DXW8UvAPCW0ienLtJn7IurI4h_wOle19YxDW-yvREqY-HyuPL5tvfULr_M6Q0SPjj8PJzt_n7Zvso5ZmKlA4fxn2NlR55mUj5Sfk8HIrpQAhTPBDOEM4nADyCjyujnbgYadVvZvPwl-6etX9QezPmIz3-L8UfQfLKnzQvDo-aqeREPflkWohqn7KT77cna8bWKSvSgsz4OlkluMiWRZrPmjVd0hVN6oJWCzKvF0ld13nNWF3UeJViWs3EGs9xNs8SnOTpIs3jNC0IpkWdcYZXhK1QNuctEXJqW2bC2p6v83yxSGaSVFza0Jxj7F8HjH2XbtbhQa_6s0XZXArr7m3PzAknQz8_NNaeBXe-UfMdi-HSs2toUYfm3DNwZDxYoeX0SXPr2v3OvbkuAofbgcK_1D4_P0OfviSz3sj1n25hAkoW4WMA6j8BAAD__0jvB2E">