1 // RUN: echo "GPU binary would be here" > %t
2
3 // RUN: %clang_cc1 -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 -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 -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 -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 -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 -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
ckernel(int * a)29 extern "C" __global__ void ckernel(int *a) {
30 *a = 1;
31 }
32
33 // Kernel symbol for launching kernel.
34 // CHECK: @[[SYM:ckernel]] = constant ptr @__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
hostfunc(int * a)51 void hostfunc(int *a) {
52 ckernel<<<1, 1>>>(a);
53 }
54