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: ) 61