clang/test/Driver/hip-partial-link.hip

   1 // REQUIRES: x86-registered-target, amdgpu-registered-target, lld, system-linux
   2
   3 // RUN: %clang -x hip --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
   4 // RUN:   --offload-arch=gfx906 -c -nostdinc -nogpuinc -nohipwrapperinc \
   5 // RUN:   -nogpulib -fgpu-rdc -I%S/Inputs %s -o %t.1.o
   6
   7 // RUN: %clang -x hip --target=x86_64-unknown-linux-gnu -DLIB --no-offload-new-driver \
   8 // RUN:   --offload-arch=gfx906 -c -nostdinc -nogpuinc -nohipwrapperinc \
   9 // RUN:   -nogpulib -fgpu-rdc -I%S/Inputs %s -o %t.2.o
  10
  11 // RUN: %clang -x hip --target=x86_64-unknown-linux-gnu -DMAIN --no-offload-new-driver \
  12 // RUN:   --offload-arch=gfx906 -c -nostdinc -nogpuinc -nohipwrapperinc \
  13 // RUN:   -nogpulib -fgpu-rdc -I%S/Inputs %s -o %t.main.o
  14
  15 // RUN: llvm-nm  %t.1.o | FileCheck -check-prefix=OBJ1 %s
  16 // OBJ1:  B __hip_cuid_[[ID:[0-9a-f]+]]
  17 // OBJ1:  U __hip_fatbin_[[ID]]
  18 // OBJ1:  U __hip_gpubin_handle_[[ID]]
  19
  20 // RUN: llvm-nm  %t.2.o | FileCheck -check-prefix=OBJ2 %s
  21 // OBJ2:  B __hip_cuid_[[ID:[0-9a-f]+]]
  22 // OBJ2:  U __hip_fatbin_[[ID]]
  23 // OBJ2:  U __hip_gpubin_handle_[[ID]]
  24
  25 // Link %t.1.o and %t.2.o by -r and then link with %t.main.o
  26
  27 // RUN: %clang -v --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
  28 // RUN:   --hip-link -fgpu-rdc --offload-arch=gfx906 \
  29 // RUN:   -r -fuse-ld=lld -nostdlib %t.1.o %t.2.o -o %t.lib.o \
  30 // RUN:   2>&1 | FileCheck -check-prefix=LD-R %s
  31 // LD-R: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID1:[0-9a-f]+]]
  32 // LD-R: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID2:[0-9a-f]+]]
  33 // LD-R: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID1]]
  34 // LD-R: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID2]]
  35 // LD-R: "{{.*}}/clang-offload-bundler" {{.*}}-unbundle
  36 // LD-R: "{{.*}}/lld" -flavor gnu -m elf64_amdgpu
  37 // LD-R: "{{.*}}/clang-offload-bundler"
  38 // LD-R: "{{.*}}/clang{{.*}}" -target x86_64-unknown-linux-gnu
  39 // LD-R: "{{.*}}/ld.lld" {{.*}} -r
  40
  41 // RUN: llvm-nm  %t.lib.o | FileCheck -check-prefix=OBJ %s
  42 // OBJ:  B __hip_cuid_[[ID1:[0-9a-f]+]]
  43 // OBJ:  B __hip_cuid_[[ID2:[0-9a-f]+]]
  44 // OBJ:  R __hip_fatbin_[[ID1]]
  45 // OBJ:  R __hip_fatbin_[[ID2]]
  46 // OBJ:  D __hip_gpubin_handle_[[ID1]]
  47 // OBJ:  D __hip_gpubin_handle_[[ID2]]
  48
  49 // RUN: %clang -v --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
  50 // RUN:   --hip-link -fgpu-rdc --offload-arch=gfx906 \
  51 // RUN:   -fuse-ld=lld -nostdlib -r %t.main.o %t.lib.o -o %t.final.o \
  52 // RUN:   2>&1 | FileCheck -check-prefix=LINK-O %s
  53 // LINK-O-NOT: Found undefined HIP {{.*}}symbol
  54
  55 // Generate a static lib with %t.1.o and %t.2.o then link with %t.main.o
  56
  57 // RUN: %clang -v --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
  58 // RUN:   --hip-link -fgpu-rdc --offload-arch=gfx906 \
  59 // RUN:   --emit-static-lib -fuse-ld=lld -nostdlib %t.1.o %t.2.o -o %t.a \
  60 // RUN:   2>&1 | FileCheck -check-prefix=STATIC %s
  61 // STATIC: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID1:[0-9a-f]+]]
  62 // STATIC: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID2:[0-9a-f]+]]
  63 // STATIC: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID1]]
  64 // STATIC: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID2]]
  65 // STATIC: "{{.*}}/clang-offload-bundler" {{.*}}-unbundle
  66 // STATIC: "{{.*}}/lld" -flavor gnu -m elf64_amdgpu
  67 // STATIC: "{{.*}}/clang-offload-bundler"
  68 // STATIC: "{{.*}}/clang{{.*}}" -target x86_64-unknown-linux-gnu
  69 // STATIC: "{{.*}}/llvm-ar"
  70
  71 // RUN: %clang -v --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
  72 // RUN:   --hip-link -no-hip-rt -fgpu-rdc --offload-arch=gfx906 \
  73 // RUN:   -fuse-ld=lld -nostdlib -r %t.main.o %t.a -o %t.final.o \
  74 // RUN:   2>&1 | FileCheck -check-prefix=LINK-A %s
  75 // LINK-A-NOT: Found undefined HIP {{.*}}symbol
  76
  77 #include "hip.h"
  78
  79 #ifdef LIB
  80 __device__ int x;
  81 __device__ void libfun() {
  82   x = 1;
  83 }
  84 #elif !defined(MAIN)
  85 __device__ void libfun();
  86 __global__ void kern() {
  87   libfun();
  88 }
  89 void run() {
  90   kern<<<1,1>>>();
  91 }
  92 #else
  93 extern void run();
  94 int main() {
  95   run();
  96 }
  97 #endif