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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] Creating relocatable object (-r) from rdc objects (-fgpu-rdc) fails with lld error attempted static link of dynamic object in /opt/rocm-5.7.1/lib
        </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>
    Hey @arsenm and @jdoerfert, how do I generate a relocatable object (-r) for the amdgpu target? I am linking a large code containing a few millions of lines of C++ with an optional library dependency containing about 300,000 lines of C++. The library requires relocatable device code (-fgpu-rdc) because it has many kernels which reference device functions defined in separate translation units. The large code does not. A driver for the library links in 30 minutes. The large code takes 2 minutes to link without the optional library and over 8 hours with the library (the lld process is still running after 8 hours). I don't want to use rdc to link the large code, but I have to because of the optional library: if even a single object needs rdc, then the link needs it too. Perhaps an intermediate step between compiling the library and linking the large code, in which I generate a relocatable object (-r) from the rdc-compiled library, would allow me to link the large code without rdc even when I'm using the optional library.

# x86+LTO (good)

Consider using LTO to target x86, which works as expected. During compilation, clang -flto emits LLVM IR, which lld uses to perform link time optimizations like cross translation unit inlining. Here is an example:
```
$ cat main.c
int add1(int y);
int main(int argc, char **argv) { return add1(argc); }
$ cat add.c
int add1(int y) { return y + 1; }
$ cat build.sh
dirs="normal_build lto_build relocatable_build"
for dir in $(echo $dirs); do rm -rf $dir; mkdir $dir; done

# build separate compilation executable
dir=normal_build
clang -O2 -c add.c -o $dir/add.o                          # add.o  contains object code
clang -O2 -c main.c -o $dir/main.o                        # main.o contains object code
clang -O2 $dir/add.o $dir/main.o -o $dir/foo              # linker sees object code

# build lto executable
dir=lto_build
clang -flto -O2 -c add.c -o $dir/add.o                    # add.o  contains llvm IR
clang -flto -O2 -c main.c -o $dir/main.o                  # main.o contains llvm IR
clang -flto -O2 $dir/add.o $dir/main.o -o $dir/foo # linker sees llvm IR

# build lto executable but with an intermediate step between compiling and
# linking which creates relocatable uber.o
dir=relocatable_build 
clang    -flto -O2 -c add.c -o $dir/add.o # add.o  contains llvm IR
clang    -flto -O2 -c main.c -o $dir/main.o               # main.o contains llvm IR
clang -r -flto -O2 $dir/add.o $dir/main.o -o $dir/uber.o  # uber.o contains object code
clang -O2 $dir/uber.o -o $dir/foo                # linker sees object code
```
Building and then disassembling the executables shows that add1, which is referenced and defined in separate translation units, is inlined for the two LTO builds but not for the separate compilation build, as expected:
```
$ sh < build.sh 
$ llvm-objdump --disassemble-symbols=main */foo 
lto_build/foo:  file format elf64-x86-64

Disassembly of section .text:

0000000000400540 <main>:
  400540: 8d 47 01                      leal    0x1(%rdi), %eax
  400543: c3                            retq
  400544: 66 2e 0f 1f 84 00 00 00 00 00 nopw    %cs:(%rax,%rax)
  40054e: 66 90 nop
  
normal_build/foo:   file format elf64-x86-64

Disassembly of section .text: 

0000000000400550 <main>:
  400550: e9 eb ff ff ff                jmp 0x400540 <add1>
 400555: 66 2e 0f 1f 84 00 00 00 00 00 nopw    %cs:(%rax,%rax)
  40055f: 90 nop 

relocatable_build/foo:  file format elf64-x86-64

Disassembly of section .text:

0000000000400570 <main>:
  400570: 8d 47 01                      leal    0x1(%rdi), %eax
  400573: c3                            retq
  400574: 66 2e 0f 1f 84 00 00 00 00 00 nopw    %cs:(%rax,%rax)
  40057e: 66 90 nop
```
The difference in the two LTO builds is that one had -flto on the link line and the other didn't. The one which included an intermediate step between compiling and linking to create a relocatable object did *not* need -flto on the link line because I gave the linker object code, not LLVM IR.

# amdgpu+rdc (bad)

Now consider my use case. I'm building with rocm 6.0.0, the latest rocm clang distribution installed on my system, and I am targeting the amd mi250x. I modified my x86+LTO code to use hip with rdc:
```
$ cat main.c
#include <hip/hip_runtime.h>
__device__ int add1(int y);
__global__ void mykernel(int *y) { *y = add1(*y); } 
int main(int argc, char **argv) {
 mykernel<<<1,1>>>(&argc);
    return argc;
}
$ cat add.c
__device__ int add1(int y) { return y + 1; }
$ cat build.sh
dirs="rdc_build relocatable_build"
for dir in $(echo $dirs); do rm -rf $dir; mkdir $dir; done

dir=rdc_build
hipcc -O2 -fgpu-rdc --offload-arch=gfx90a -x hip -c add.c -o $dir/add.o         # add.o contains llvm IR
hipcc -O2 -fgpu-rdc --offload-arch=gfx90a -x hip -c main.c -o $dir/main.o       # main.o contains llvm IR
hipcc -O2 -fgpu-rdc --offload-arch=gfx90a $dir/add.o $dir/main.o -o $dir/foo    # linker sees llvm IR

dir=relocatable_build
hipcc    -O2 -fgpu-rdc --offload-arch=gfx90a -x hip -c add.c -o $dir/add.o          # add.o contains llvm IR
hipcc    -O2 -fgpu-rdc --offload-arch=gfx90a -x hip -c main.c -o $dir/main.o        # main.o contains llvm IR
hipcc -r -O2 -fgpu-rdc --offload-arch=gfx90a $dir/add.o $dir/main.o -o $dir/uber.o  # uber.o contains object code 
hipcc    -O2           --offload-arch=gfx90a -o $dir/uber.o                         # linker sees object code
```
The second-to-last line, which uses -r to make the relocatable object, fails with `ld.lld: error: attempted static link of dynamic object` and references shared libraries in /opt/rocm:
```
$ sh < build.sh
clang: warning: argument unused during compilation: '--rtlib=compiler-rt' [-Wunused-command-line-argument]
clang: warning: argument unused during compilation: '-unwindlib=libgcc' [-Wunused-command-line-argument]
ld.lld: error: attempted static link of dynamic object /opt/rocm-6.0.0/lib/libamdhip64.so
ld.lld: error: attempted static link of dynamic object /opt/rocm-6.0.0/lib/libhsa-runtime64.so
ld.lld: error: attempted static link of dynamic object /opt/rocm-6.0.0/lib/libamd_comgr.so
ld.lld: error: attempted static link of dynamic object /opt/rocm-6.0.0/lib/libamdhip64.so
clang: error: linker command failed with exit code 1 (use -v to see invocation)
ld.lld: error: undefined symbol: main
>>> referenced by /lib/../lib64/crt1.o:(_start)
clang: error: linker command failed with exit code 1 (use -v to see invocation)
```
Ignore the last 3 lines above, which are due to my attempt to link using the non-existent object file uber.o.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzEWVtv6zby_zTMy0CCTNux_ZCHxG7-DdD-Wyy6u48BLY4snlCkSlKx00-_GFLyJXFynLOnZw3DUXiZGc79Rwnv1cYg3rDpHZuurkQXautuWvtXp9urtZUvNz_jC7BJIZxH04Awkv77Ii26Cl1gfAm13YK08AAbNOhEQBDgUNtSBLHWCHb9BcsAjM8zx_gCKusg1AiikZu2gyDcBgMb38MDiAa0Mk_KbECApgkoraQfE4QyabzCLTRKa2WNB1vRDowPS8bvGL-DrQo1CAO2DcoaoUGrtRPuBSS2aCSa8uWE4tp2AcZFwfiyKIo39HL4o8Y9DYd_dsqhPzmixGdV9rLSOatN22VOlnTcNZai8wgqQC08NMK8wBM6g9rDtlZlDQ4rdGjKPZ2qM2WIx5NYKYMSlAGPrYjaDU4YrwUtgM6o4HsBD-qSFj0YG3K4BenUM7q90odjkJo9kR0X0CjTBXxLJogn9MCHeQg2bov6JZURvTc6Jg-xxHEOte2cT9Y4Zs34PP6rJbTOlug9KA8-KK3BdSbZpAoHEowvcngAaQ3jswBbYQLJQkp1styLFU6kJ89cdwEeoBbPSIsGQ9jqrORsfAuqAnxGAwK8MpuD7xpE6SEadEmbTX8g89RPKZLI5vA7ulq0nrxPmYCuQanIZj5gC2sMW0QDpW1apemYx2ohxQ3O__YoyvTOcnmYOdtEQk6WWWKJcn9YvoSt7bQEobXdQoPvqHFvbNJ01M2Wjv_A-KyBzg_CvlZmzooVK277Xz6G3fya8btf_viNBNxYKxlfHK9ZWuOVRNfTpIXB9rkhbV72Ctha9-RBeMBdi2VAmcOqc7QpHTIGBi0vtTAbyCodLGCjgodffvnXr_DwjwMt8sHOJ9du0VXWNb0SVJMO1ai_RApFrZ4QSme9fxOBoIyOuSSHn9EhubMwgDvRtBrZeFDDddF_e61MoBQBGqFMXqYxZQIIKUeMz-nxhZQ0vjvM0dp-TrhN9MayFg4Yv2X8VrjNM1meze7AYeicGYilxUQL2Gx1yl9I-QH7Y2IUuncwOktl3Sktc1-nUamcZ-MV49xY1wj9GKdBB9s_HTluGmGcp62UqKRy5PCMTxifY1lbeow00xmkBddA5qp-nMaaJ9p1-F9ag6-dMPHeJ9IjfwHcYdlFefYnYOPVsfRpvHeq3zhkZVIdZIN4jN_TiIV3PyREv6QvQH6I2xjob1kk7zjhEYfeZUIs-hWXsHgl-Wsmx3wra9-yomhBBx7xDJvXmo-ReFbRe884kS-G7ud1fU7LWj83FPrvkf-Ens9p-EPyn9Txa7We0P5YpbHiDe3PJQVIGHmgOFSflBtLhyK86nO6NbrcHhvuTSDDsQ4ALjPihRZ7Te5So11qMfdtRktaSWz650-FXr_nw1i7JNpO68sdWaO3cepZpPLCe2zW-97j4DkefG23HkIthhowFEnlDx2qjNQu6kpj1-JTaUS5b0DD1sbqHp3FR4c1NuynzybnvkIsj6v-R2XV18DGy31NgsMMWT2z6y-ya1rIsoNGMPMvzdpqKltk4lRSUzzG3YcMFUepW4RKaSTJGxEAdXU9yXbz6-x6chyqqz2LF-o8PcbOHvKAu3A4Q_wt9p9JUUwnBR0iVvzxT_uVAGmO-M8lTGZQjM4XAo1C099iR_Wc8amTiuonXwLjUxS7E4JjIliO36sq9HEY_jzZM6E919fAEYoKRhXMJ1AUJ19j221y3mnp6RBJErFjfDk8LE6IYk90ETcPU-nvSTk-mOF72AHeN8T0A0NMoyFwAbiGquq_rz5fmhaK3cGkMbrGP_VkIpXp36DJaUVEkxpPTnem9_oxHj37QJGz7-3Rs2_w6Nnf4dGzcx79KmkR8JaqGq4BlDmXKlWfna1BqIXsa5U9wqKUaYd0DzbUSK20jKg5gXva2id1U-pOxnx-aaNwQKe2bw7OI1CpJGVPYwPjtxEfvyfqgMgfYBMxej-J7qS48WWsED14e4Ms0zUS43eEUBmfr8VrbPn_dksFOeHL5iVeHJTCY97D2PVQKWPn5GzZwHVe5EUP9kFTHxTSRCrfUvng1LqLjq-MD0ITuLaGyPsXH7CJ5crIdKmVcOxQdUUjoVF8WuxyeIDGSlUplLT1gJLTBUy65KhV24smy8-gScbHvZkp7mrVMn5fq_bRdYbwbV7vs9DjY7p5enyEjyDo4-NG27XQj4_wbBVJnC6y-rWM3-4hIz0DG68GWv1cjx3hGyBtH1N7nuNl-lKnEvNp-hKr6wPeHSIR9oiYpobxj8Dwxzr5b3Gxk-X_DAr3nfsgQRqsVVuWqbEebi8hy2xVaStkJlxZs_FqU-0WhYBsF33yAlR2aOzPN97fxvVrff9X-_3PsP08Qr4AwL0Dno6lI7Dzna1xoTm-gfNXkdilJnHf3yoXwjM4o4LD510VnGH0zuezAO6PCIdKa2QWbKaFD7FqHoBZvLnMHJWJRjyl-vm2HtP6Sijd38Sz60LLXGsZG1fnrKMHEQI2bUAJPoigylSkbQXyxYhGlQOt6yIWtT0eJNgo3P5eWaFP2eretoHxeyqZnwBqRwiZhNoKZ1R6FG7TNWgCdKbzKEG-vfGlHp7PsswFrdZsvOqvvF3mAuMzYNO77N9pd1baphFGZqTObCDNpqvvwb8zW2VkEkGr9aYsP8f9241zqvasb2HuSZT4KxpZq_Z6knv7d7Oqvcj6LuOH8BONfCxts3E_htepGvf-sufUR3lv6Bh9KFP44U716WZEzSq1d9kzRbBH6vufKXjj-4vFu-fozHD_kq4saCx2USmshk7o-NJmTS1Kf4Q8T4_XE8bvSxdGuU0w5tEHQbGy-AHHepULHjbGuh4AUJ4b9y9hxdo-HyU84RBkF_vi5mWw5_7F1eFtlLEmw53ygSK2t2pEtilD51fyZiwX44W4wpvRrJhMivmUT6_qGz6Tci6qSvKRmMsRFliOxawcFfPxghc4ulI3vOCTYlRMOJ_MJ9N8LsbzRblAWQhZlKM5mxTYCKVzKm65dZsr5X2HN7NZMZpfabFG7ePbds51avWmqyt3E--l1t3Gs0mhlQ_-sD-ooOML-ttfV__3-z_ZdAVLQl902K--_6MSmib8mxfTR0VBa5nMfHGQvM7y2TSf5aPkWVed0zd1CG3Cx_eM329UqLt1XtqGlujn4U_WOpuK1H3Uk2f8PqrqPwEAAP__xhQWmA">