clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp

   1 // RUN: %clang_cc1 %s -o - -O0 -emit-llvm                                     \
   2 // RUN:            -triple spir64-unknown-unknown                             \
   3 // RUN:            -aux-triple x86_64-unknown-linux-gnu                       \
   4 // RUN:            -fsycl-is-device                                           \
   5 // RUN:            -finclude-default-header                                   \
   6 // RUN:            -debug-info-kind=limited -gno-column-info                  \
   7 // RUN:   | FileCheck %s
   8 //
   9 // In spir functions, validate the llvm.dbg.declare intrinsics created for
  10 // parameters and locals refer to the stack allocation in the alloca address
  11 // space.
  12 //
  13
  14 #define KERNEL __attribute__((sycl_kernel))
  15
  16 template <typename KernelName, typename KernelType>
  17 KERNEL void parallel_for(const KernelType &KernelFunc) {
  18   KernelFunc();
  19 }
  20
  21 void my_kernel(int my_param) {
  22   int my_local = 0;
  23   my_local = my_param;
  24 }
  25
  26 int my_host() {
  27   parallel_for<class K>([=]() { my_kernel(42); });
  28   return 0;
  29 }
  30
  31 // CHECK:      define {{.*}}spir_func void @_Z9my_kerneli(
  32 // CHECK-SAME    i32 %my_param
  33 // CHECK-SAME:   !dbg [[MY_KERNEL:![0-9]+]]
  34 // CHECK-SAME: {
  35 // CHECK:        %my_param.addr = alloca i32, align 4
  36 // CHECK:        %my_local = alloca i32, align 4
  37 // CHECK:        call void @llvm.dbg.declare(
  38 // CHECK-SAME:     metadata i32* %my_param.addr,
  39 // CHECK-SAME:     metadata [[MY_PARAM:![0-9]+]],
  40 // CHECK-SAME:     metadata !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
  41 // CHECK-SAME:     )
  42 // CHECK:        call void @llvm.dbg.declare(
  43 // CHECK-SAME:     metadata i32* %my_local,
  44 // CHECK-SAME:     metadata [[MY_LOCAL:![0-9]+]],
  45 // CHECK-SAME:     metadata !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
  46 // CHECK-SAME:     )
  47 // CHECK:      }
  48
  49 // CHECK:      [[MY_KERNEL]] = distinct !DISubprogram(
  50 // CHECK-SAME:   name: "my_kernel"
  51 // CHECK-SAME:   )
  52 // CHECK:      [[MY_PARAM]] = !DILocalVariable(
  53 // CHECK-SAME:   name: "my_param"
  54 // CHECK-SAME:   arg: 1
  55 // CHECK-SAME:   scope: [[MY_KERNEL]]
  56 // CHECK-SAME:   )
  57 // CHECK:      [[MY_LOCAL]] = !DILocalVariable(
  58 // CHECK-SAME:   name: "my_local"
  59 // CHECK-SAME:   scope: [[MY_KERNEL]]
  60 // CHECK-SAME:   )