98 lines · plain
1// REQUIRES: x86-registered-target, amdgpu-registered-target, lld2 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.o6 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.o10 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.o14 15// RUN: llvm-nm %t.1.o | FileCheck -check-prefix=OBJ1 %s16// 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 %s21// 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.o26 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 %s31// 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" {{.*}}-unbundle36// LD-R: "{{.*[/\\]}}lld" -flavor gnu -m elf64_amdgpu37// LD-R: "{{.*[/\\]}}clang-offload-bundler"38// LD-R: "{{.*[/\\]}}clang{{.*}}" -target x86_64-unknown-linux-gnu39// LD-R: "{{.*[/\\]}}ld.lld" {{.*}} -r40 41// RUN: llvm-nm %t.lib.o | FileCheck -check-prefix=OBJ %s42// 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 %s53// LINK-O-NOT: Found undefined HIP {{.*}}symbol54 55// Generate a static lib with %t.1.o and %t.2.o then link with %t.main.o56 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 %s61// 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" {{.*}}-unbundle66// STATIC: "{{.*[/\\]}}lld" -flavor gnu -m elf64_amdgpu67// STATIC: "{{.*[/\\]}}clang-offload-bundler"68// STATIC: "{{.*[/\\]}}clang{{.*}}" -target x86_64-unknown-linux-gnu69// 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 %s75// LINK-A-NOT: Found undefined HIP {{.*}}symbol76 77#include "hip.h"78 79#ifdef LIB80__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}89void run() {90 kern<<<1,1>>>();91}92#else93extern void run();94int main() {95 run();96}97#endif98