1 // RUN: echo "GPU binary would be here" > %t 2 3 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -O0 \ 4 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 5 // RUN: -o - -x hip | FileCheck -check-prefixes=CHECK,O0 %s 6 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \ 7 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 8 // RUN: -o - -x hip -fcuda-is-device | FileCheck -check-prefix=DEV %s 9 10 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -O0 \ 11 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 12 // RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \ 13 // RUN: | FileCheck -check-prefixes=CHECK,O0 %s 14 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \ 15 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 16 // RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \ 17 // RUN: -fcuda-is-device | FileCheck -check-prefix=DEV %s 18 19 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -O3 \ 20 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 21 // RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 | FileCheck %s 22 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -emit-llvm %s -O3 \ 23 // RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \ 24 // RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \ 25 // RUN: -fcuda-is-device | FileCheck -check-prefix=DEV %s 26 27 #include "Inputs/cuda.h" 28 29 extern "C" __global__ void ckernel(int *a) { 30 *a = 1; 31 } 32 33 // Kernel symbol for launching kernel. 34 // CHECK: @[[SYM:ckernel]] = constant void (i32*)* @__device_stub__ckernel, align 8 35 36 // Device side kernel names 37 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00" 38 39 // DEV: define {{.*}}@ckernel{{.*}}!dbg 40 // DEV: store {{.*}}!dbg 41 // DEV: ret {{.*}}!dbg 42 43 // Make sure there is no !dbg between function attributes and '{' 44 // CHECK: define{{.*}} void @[[CSTUB:__device_stub__ckernel]]{{.*}} #{{[0-9]+}} { 45 // CHECK-NOT: call {{.*}}@hipLaunchByPtr{{.*}}!dbg 46 // CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[SYM]] 47 // CHECK-NOT: ret {{.*}}!dbg 48 49 // CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg 50 // O0: call void @[[CSTUB]]{{.*}}!dbg 51 void hostfunc(int *a) { 52 ckernel<<<1, 1>>>(a); 53 } 54