54 lines · plain
1// RUN: echo "GPU binary would be here" > %t2 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 %s6// 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 %s9 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 %s14// 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 %s18 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 %s22// 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 %s26 27#include "Inputs/cuda.h"28 29extern "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 835 36// Device side kernel names37// CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"38 39// DEV: define {{.*}}@ckernel{{.*}}!dbg40// DEV: store {{.*}}!dbg41// DEV: ret {{.*}}!dbg42 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{{.*}}!dbg46// CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[SYM]]47// CHECK-NOT: ret {{.*}}!dbg48 49// CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg50// O0: call void @[[CSTUB]]{{.*}}!dbg51void hostfunc(int *a) {52 ckernel<<<1, 1>>>(a);53}54