clang/test/CodeGenCUDA/kernel-dbg-info.cu

   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
  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 }