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

    <tr>
        <th>Summary</th>
        <td>
            [HIP] Link error when compiling the template kernel function code with -fgpu-rdc
        </td>
    </tr>

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

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

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

<pre>
    **Overview**
---
Here is a code that contains calling the template kernel function. When I compile the code with `-fgpu-rdc` option, there will be undefined error in the final link stage.
`lld: errcor: undefined protected symbol: void kernel1<(ENUM)11111>()`

**Error**
---
I tested it with LLVM 17.0.6 in the ROCM environment. When compiling the example code in the following steps, an error will be reported in the link.
```sh
clang++ -c -fgpu-rdc -x hip -o call.o call.cu --offload-arch=gfx906 -include hip/hip_runtime.h
clang++ -c -fgpu-rdc -x hip -o kernel.o kernel.cu --offload-arch=gfx906 -include hip/hip_runtime.h
clang++ -c -fgpu-rdc -x hip -o main.o main.cu --offload-arch=gfx906 -include hip/hip_runtime.h 
clang++ -o test call.o kernel.o  main.o --hip-link -fgpu-rdc
```

```sh
[xxx]$ clang++ -o test call.o kernel.o main.o --hip-link -fgpu-rdc
lld: error: undefined protected symbol: void kernel1<(ENUM)11111>()
>>> referenced by lto.tmp:(__clang_gpu_used_external)

lld: error: undefined protected symbol: void kernel1<(ENUM)11155>()
>>> referenced by lto.tmp:(__clang_gpu_used_external)

lld: error: undefined protected symbol: void kernel1<(ENUM)11144>()
>>> referenced by lto.tmp:(__clang_gpu_used_external)

lld: error: undefined protected symbol: void kernel1<(ENUM)11133>()
>>> referenced by lto.tmp:(__clang_gpu_used_external)
clang++: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)
[xxx]$ clang --version
clang version 17.0.6
Target: x86_64-unknown-linux-gnu
Thread model: posix
```

**Code example**
---
```C++
template<ENUM METHOD>
void callkernel() {
    dim3 grid(1,1);
    dim3 block(1,1);

    if ((METHOD == ENUM_55)) {

    } else if ((METHOD == ENUM_66) || (METHOD == ENUM_77)){
        kernel1<METHOD><<<grid, block>>>();
    }

 switch(METHOD){
        case ENUM_11: {
           break;
 }
        case ENUM_22: {
            kernel1<METHOD><<<grid, block>>>();
            break;
        }
        case ENUM_33: {
 break;
        }
        case ENUM_44: {
 break;
        }
        default:{
            break;
 }
      }
}
```
The above is some code fragments. I know that my template function may not be standard enough. However, when compiling normally and does not involve `__global__` function and `-fgpu-rdc` option, the compiler can help me optimize and eliminate useless template function symbols. This code usage will not report errors. Only when the `__global__` function is called and the `-fgpu-rdc` option is used, the optimization of the compiler does not seem to take effect. Finally, all the template function symbols are retained, and then an error is reported. There is no such problem in native nvcc.

I have an idea. When processing `__clang_gpu_used_external` structure, can it be optimized to eliminate useless template function symbols, thereby supporting my code usage.

The complete example is here, please help me analyze the reason for me.
Thanks.
[template-issue-test.zip](https://github.com/llvm/llvm-project/files/14477350/template-issue-test.zip)

</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzUWN1u2zoSfhr6ZiBBpmTZvvBFasdogHazWGR3Lw1aGkncUKRAUv7p0y9ISnbcOjmnBz0XNYRKlYYz33xDzk-YMbyWiCsy-0RmmwnrbaP0quzrvkN-RDnZq_K8IvSB0IfnA-oDx2P4H0k2JHmIoig8fEaNwA0wKFSJYBtmoVDSMi4NFEwILmuwDYLFthPMIryiliig6mVhuZIx_LdBCU9QqLbjAr2w13XktgGSJ1FVd32ky4LkCajOLSJ07eS0ExIC9gi9LLHiEktArZUGLr2iiksmQHD5CsayGuOAmuSJECVJH5x0obR7umrotLJYWCzBnNu9Eu7rQfFygD4l6ZrQxeM__v2V0OXU_Uj6SOiC0CXJk8FC-NdT9ugQ3WXvCSwaZ4jb4O6XL__5CtN5nMT56MK_ntdfAeWBayVblHYgLNA1sosn1nZiIG70XQmhjk7CWOyM44zJgZ6RNo2d0h5AWOOYunIULtOEF4Vgsib0E6GfICrgEhaITtDwDiLlAx4Pt6KHKFJVJRQrI6aLhqSbujotkxwiLgvRl-iWEbpteLfTvbS8xfjP2wrBiC8Pf7e9lnEZD7e_aAvuGFN-C4zMXZwazUVRw7vIb-DrObiNz82G-z5oZPbpdDqR2YbQDP6M4T-0ez04v_LcBLDpY7hAY4UaZYEl7M8grIpt25H0gdDFbufd2NVdv-sNljs8WdSSiauaX4xzNvs9cGbZ74EzTf8enG-29w1S1pZ1IcNmLlTbMllCxbjAMiRdPHEbMucUCF30BiE6gFVg0CXTgypYKDrL984URNEBtXFCVyQwvBryefjywnSN1qE6LfJdnkW9fJXq6NH1p6iW_SDXaGQltKpEz2enDD99dPB9eVk7J4ZacLfgXNauB57827E2k3TtogRfH18-P29cSPxnH0uXJ0I8Q-SAzIfVAAAlb1OoNS8JXUwJXU8dWen3Anuhitd7Elc5XoHXvwgYgKQbkm7AwdrNZm7NW9PXdWS-ARQGP1SQ52H1mszXcF9kPg82brxzv-tevrKTrsMVHF8P_o2bedjhb1kg880NcHPktmguSO7ZLZjBAM2lywf4QQAA9hrZ69XQxcqPOih9T8evcfAdRMPvA2AuJ7wF9tMKXPL7CwpKrFgv3IG8S8rHzF7DeXm4PZsvDQLbq4NvkY1qhwat0qx2rZyJ4Qnc8Q99c3u-dsljewwtO4NU1jVrxjJZMl0CStXXTQyf1REPqF1kjrc9oVS6ZUKcweW6UqHxOlwuEwd0TfVuVwu1Z2K3c131xZoT_7DlHtt0DQWT0KDooEUv0vJv6Nej4C2XzoneoEBj7ngVCoWJ4aXhJpDSG1YP_bzDGjrTkMRNDM9SnIOTDsT7DvAwdWDpoQyyd9xxgq6KjG4NHvhED6q6dfVCoEFsXWGw7BUBqwoLG8PWzRji7PtrIW6Hne8dBqZd0-3Go2B7QCmvrTk3l67c0TPMV1KB6YvG1dq9wNb165JZfkCQh6KI32aVJ2jYwYUCeIlsGBc6rQo0xu0NT957tTRPwFjdF7bX6PC5KHO_-8YYl46An4jxZVbbn8H0nfPMoWjPb8J-g_9loF6gvQ423IDT4ZR1At2pH_cek0ycv4W5USMzSkKlNLSj0peGyVcTX0r3CDTixvQYuSY4_sY7X84XjbWd8f3GltBtzW3T7-NCtYRuhTiMt6jT6n9YWEK3FRdoCN1Os2w-T2cJodv3DIztw6RcpeUyXbIJrqbzZDlL0zyZTppVRtlimc6XaU6zIksKmmSsSqZJlk-z5bxYTPiKJjRL0iRLcppn0zilNEnnS0Q2X7LZMidZgi3jInYoY6XricewWqTz-XQi2B6F8TM_pRKP4D8SSslsM9Er79m-rw3JEsGNNVctllvh_1jw-emfZLaBL66VGibJH4fR90b9N2P95UROei1WP826B-5o9479PwAA__8etiXZ">