In spir functions, llvm.dbg.declare intrinsics created

for parameters and locals need to refer to the stack
allocation in the alloca address space.

Author: zahiraam

clang/lib/CodeGen/CGDecl.cpp

@@ -1447,6 +1447,7 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
 
   if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
     address = OpenMPLocalAddr;
+    AllocaAddr = OpenMPLocalAddr;
   } else if (Ty->isConstantSizeType()) {
     // If this value is an array or struct with a statically determinable
     // constant initializer, there are optimizations we can do.
@@ -1492,6 +1493,7 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
       // return slot, so that we can elide the copy when returning this
       // variable (C++0x [class.copy]p34).
       address = ReturnValue;
+      AllocaAddr = ReturnValue;
 
       if (const RecordType *RecordTy = Ty->getAs<RecordType>()) {
         const auto *RD = RecordTy->getDecl();
@@ -1503,7 +1505,8 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
           // applied.
           llvm::Value *Zero = Builder.getFalse();
           Address NRVOFlag =
-            CreateTempAlloca(Zero->getType(), CharUnits::One(), "nrvo");
+              CreateTempAlloca(Zero->getType(), CharUnits::One(), "nrvo",
+                               /*ArraySize=*/nullptr, &AllocaAddr);
           EnsureInsertPoint();
           Builder.CreateStore(Zero, NRVOFlag);
 
@@ -1605,10 +1608,11 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
     DI->setLocation(D.getLocation());
 
     // If NRVO, use a pointer to the return address.
-    if (UsePointerValue)
+    if (UsePointerValue) {
       DebugAddr = ReturnValuePointer;
-
-    (void)DI->EmitDeclareOfAutoVariable(&D, DebugAddr.getPointer(), Builder,
+      AllocaAddr = ReturnValuePointer;
+    }
+    (void)DI->EmitDeclareOfAutoVariable(&D, AllocaAddr.getPointer(), Builder,
                                         UsePointerValue);
   }
 
@@ -2450,6 +2454,7 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
   }
 
   Address DeclPtr = Address::invalid();
+  Address AllocaPtr = Address::invalid();
   bool DoStore = false;
   bool IsScalar = hasScalarEvaluationKind(Ty);
   // If we already have a pointer to the argument, reuse the input pointer.
@@ -2464,6 +2469,7 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
     // from the default address space.
     auto AllocaAS = CGM.getASTAllocaAddressSpace();
     auto *V = DeclPtr.getPointer();
+    AllocaPtr = DeclPtr;
     auto SrcLangAS = getLangOpts().OpenCL ? LangAS::opencl_private : AllocaAS;
     auto DestLangAS =
         getLangOpts().OpenCL ? LangAS::opencl_private : LangAS::Default;
@@ -2500,10 +2506,11 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
             : Address::invalid();
     if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
       DeclPtr = OpenMPLocalAddr;
+      AllocaPtr = DeclPtr;
     } else {
       // Otherwise, create a temporary to hold the value.
       DeclPtr = CreateMemTemp(Ty, getContext().getDeclAlign(&D),
-                              D.getName() + ".addr");
+                              D.getName() + ".addr", &AllocaPtr);
     }
     DoStore = true;
   }
@@ -2579,7 +2586,7 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
   if (CGDebugInfo *DI = getDebugInfo()) {
     if (CGM.getCodeGenOpts().hasReducedDebugInfo() && !CurFuncIsThunk) {
       llvm::DILocalVariable *DILocalVar = DI->EmitDeclareOfArgVariable(
-          &D, DeclPtr.getPointer(), ArgNo, Builder);
+          &D, AllocaPtr.getPointer(), ArgNo, Builder);
       if (const auto *Var = dyn_cast_or_null<ParmVarDecl>(&D))
         DI->getParamDbgMappings().insert({Var, DILocalVar});
     }

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

@@ -0,0 +1,60 @@
+// RUN: %clang_cc1 %s -o - -O0 -emit-llvm                                     \
+// RUN:            -triple spir64-unknown-unknown                             \
+// RUN:            -aux-triple x86_64-unknown-linux-gnu                       \
+// RUN:            -fsycl-is-device                                           \
+// RUN:            -finclude-default-header                                   \
+// RUN:            -debug-info-kind=limited -gno-column-info                  \
+// RUN:   | FileCheck %s
+//
+// In spir functions, validate the llvm.dbg.declare intrinsics created for
+// parameters and locals refer to the stack allocation in the alloca address
+// space.
+//
+
+#define KERNEL __attribute__((sycl_kernel))
+
+template <typename KernelName, typename KernelType>
+KERNEL void parallel_for(const KernelType &KernelFunc) {
+  KernelFunc();
+}
+
+void my_kernel(int my_param) {
+  int my_local = 0;
+  my_local = my_param;
+}
+
+int my_host() {
+  parallel_for<class K>([=]() { my_kernel(42); });
+  return 0;
+}
+
+// CHECK:      define {{.*}}spir_func void @_Z9my_kerneli(
+// CHECK-SAME    i32 %my_param
+// CHECK-SAME:   !dbg [[MY_KERNEL:![0-9]+]]
+// CHECK-SAME: {
+// CHECK:        %my_param.addr = alloca i32, align 4
+// CHECK:        %my_local = alloca i32, align 4
+// CHECK:        call void @llvm.dbg.declare(
+// CHECK-SAME:     metadata i32* %my_param.addr,
+// CHECK-SAME:     metadata [[MY_PARAM:![0-9]+]],
+// CHECK-SAME:     metadata !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
+// CHECK-SAME:     )
+// CHECK:        call void @llvm.dbg.declare(
+// CHECK-SAME:     metadata i32* %my_local,
+// CHECK-SAME:     metadata [[MY_LOCAL:![0-9]+]],
+// CHECK-SAME:     metadata !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)
+// CHECK-SAME:     )
+// CHECK:      }
+
+// CHECK:      [[MY_KERNEL]] = distinct !DISubprogram(
+// CHECK-SAME:   name: "my_kernel"
+// CHECK-SAME:   )
+// CHECK:      [[MY_PARAM]] = !DILocalVariable(
+// CHECK-SAME:   name: "my_param"
+// CHECK-SAME:   arg: 1
+// CHECK-SAME:   scope: [[MY_KERNEL]]
+// CHECK-SAME:   )
+// CHECK:      [[MY_LOCAL]] = !DILocalVariable(
+// CHECK-SAME:   name: "my_local"
+// CHECK-SAME:   scope: [[MY_KERNEL]]
+// CHECK-SAME:   )