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 // Device side kernel names
34 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
35
36 // DEV: define {{.*}}@ckernel{{.*}}!dbg
37 // DEV: store {{.*}}!dbg
38 // DEV: ret {{.*}}!dbg
39
40 // Make sure there is no !dbg between function attributes and '{'
41 // CHECK: define void @[[CSTUB:__device_stub__ckernel]]{{.*}} #{{[0-9]+}} {
42 // CHECK-NOT: call {{.*}}@hipLaunchByPtr{{.*}}!dbg
43 // CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[CSTUB]]
44 // CHECK-NOT: ret {{.*}}!dbg
45
46 // CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg
47 // O0: call void @[[CSTUB]]{{.*}}!dbg
hostfunc(int * a)48 void hostfunc(int *a) {
49 ckernel<<<1, 1>>>(a);
50 }
51