In SPIR functions validate the llvm.dbg.declare intrinsics for locals and parameters refer to stack allocation in the alloca address space.

[erichkeane]

clang/lib/CodeGen/CGDecl.cpp

1617

I don't believe this is the correct answer here. I think we need to store the AllocaAddress out of the CreateTempAlloca call we do in the NRVO case above. In the other cases, the AllocaAddr and the DebugAddr should be the same, so there is no need to strip casts.

You probably want to do it more like how the below function is done, the AllocaAddr variable is already there, it just isn't set on every branch.

[erichkeane]

clang/lib/CodeGen/CGDecl.cpp

1617

I don't believe this is correct either, AllocaAddr is set in a couple of cases. There are two options I see here as potential:
1- make sure AllocaAddr is always updated to the correct thing in each of the above branches.
2- Only use DebugAddr when allocaAddr is not set.

But doing it off of the logic from above (in this case, the incorrect logic, but logic from above all the same) is too fragile.

[zahiraam]

clang/lib/CodeGen/CGDecl.cpp

1617

I opted for option 2, mostly because I am not sure what should the value of AllocaAddr be in case NRVO is true. I see that AllocaAddr is calculated from the CreateTempAlloc function for the other case.

[erichkeane]

clang/lib/CodeGen/CGDecl.cpp

1508–1509

I suspect in the N RVO case we want the AllocAddr from the optional 5th parameter here.

[zahiraam]

clang/lib/CodeGen/CGDecl.cpp

1508–1509

Yes I saw that, but even that call is under a condition that in some OpenMP test case is not satisfied and I wind up with an AllocAddr invalid!

[erichkeane]

clang/lib/CodeGen/CGDecl.cpp

1436

I think there are a couple of places this needs to be set in order to be 'correct', I'll comment below:

1449–1450

Probably should be the same as address here.

1495–1496

For the 'else' cases to all the below 'ifs', it is the same as address here too.

1508–1509

Should be the last param here as said above.

1612

This is the last case, if it pops into this UsePointerValue, we should do the same as DebugAddr here.

[erichkeane]

clang/lib/CodeGen/CGDecl.cpp

1517

This looks wrong, it looks like it overwrites the value put in on 1507, doesn't it? Should you put this on 1496 so that it gets overwritten by the CreateTempAlloc instead of the reverse?

1612

Missed this one, which might be why you were tempted to have the NRVO case wrong above.

[erichkeane]

I'm OK with this, but think that @rjmccall or @bader should have a chance to comment. So please don't commit for 48 hrs.

[zahiraam]

rG627868263cd4

clang/lib/CodeGen/CGDecl.cpp

CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
   // If the type is variably-modified, emit all the VLA sizes for it.
   if (Ty->isVariablyModifiedType())
     EmitVariablyModifiedType(Ty);
 
   auto *DI = getDebugInfo();
   bool EmitDebugInfo = DI && CGM.getCodeGenOpts().hasReducedDebugInfo();
 
   Address address = Address::invalid();
   Address AllocaAddr = Address::invalid();
   Address OpenMPLocalAddr = Address::invalid();
   if (CGM.getLangOpts().OpenMPIRBuilder)
     OpenMPLocalAddr = OMPBuilderCBHelpers::getAddressOfLocalVariable(*this, &D);
   else
     OpenMPLocalAddr =
         getLangOpts().OpenMP
             ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D)
             : Address::invalid();
 
   bool NRVO = getLangOpts().ElideConstructors && D.isNRVOVariable();
 
   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.
     //
     // TODO: We should constant-evaluate the initializer of any variable,
     // as long as it is initialized by a constant expression. Currently,
     // isConstantInitializer produces wrong answers for structs with
     // reference or bitfield members, and a few other cases, and checking
if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
     // A normal fixed sized variable becomes an alloca in the entry block,
     // unless:
     // - it's an NRVO variable.
     // - we are compiling OpenMP and it's an OpenMP local variable.
     if (NRVO) {
       // The named return value optimization: allocate this variable in the
       // 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();
         const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD);
         if ((CXXRD && !CXXRD->hasTrivialDestructor()) ||
             RD->isNonTrivialToPrimitiveDestroy()) {
           // Create a flag that is used to indicate when the NRVO was applied
           // to this variable. Set it to zero to indicate that NRVO was not
           // 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);
 
           // Record the NRVO flag for this variable.
           NRVOFlags[&D] = NRVOFlag.getPointer();
           emission.NRVOFlag = NRVOFlag.getPointer();
         }
       }
     } else {
       CharUnits allocaAlignment;
       llvm::Type *allocaTy;
       if (isEscapingByRef) {
         auto &byrefInfo = getBlockByrefInfo(&D);
         allocaTy = byrefInfo.Type;
         allocaAlignment = byrefInfo.ByrefAlignment;
       } else {
CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
 
   // Emit debug info for local var declaration.
   if (EmitDebugInfo && HaveInsertPoint()) {
     Address DebugAddr = address;
     bool UsePointerValue = NRVO && ReturnValuePointer.isValid();
     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);
   }
 
   if (D.hasAttr<AnnotateAttr>() && HaveInsertPoint())
     EmitVarAnnotations(&D, address.getPointer());
 
   // Make sure we call @llvm.lifetime.end.
   if (emission.useLifetimeMarkers())
     EHStack.pushCleanup<CallLifetimeEnd>(NormalEHLifetimeMarker,
                                          emission.getOriginalAllocatedAddress(),
if (BlockInfo) {
                            ? Builder.CreateLoad(Arg.getIndirectAddress())
                            : Arg.getDirectValue();
       setBlockContextParameter(IPD, ArgNo, V);
       return;
     }
   }
 
   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.
   if (Arg.isIndirect()) {
     DeclPtr = Arg.getIndirectAddress();
     // If we have a prettier pointer type at this point, bitcast to that.
     unsigned AS = DeclPtr.getType()->getAddressSpace();
     llvm::Type *IRTy = ConvertTypeForMem(Ty)->getPointerTo(AS);
     if (DeclPtr.getType() != IRTy)
       DeclPtr = Builder.CreateBitCast(DeclPtr, IRTy, D.getName());
     // Indirect argument is in alloca address space, which may be different
     // 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;
     if (SrcLangAS != DestLangAS) {
       assert(getContext().getTargetAddressSpace(SrcLangAS) ==
              CGM.getDataLayout().getAllocaAddrSpace());
       auto DestAS = getContext().getTargetAddressSpace(DestLangAS);
       auto *T = V->getType()->getPointerElementType()->getPointerTo(DestAS);
void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
   } else {
     // Check if the parameter address is controlled by OpenMP runtime.
     Address OpenMPLocalAddr =
         getLangOpts().OpenMP
             ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D)
             : 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;
   }
 
   llvm::Value *ArgVal = (DoStore ? Arg.getDirectValue() : nullptr);
 
   LValue lv = MakeAddrLValue(DeclPtr, Ty);
   if (IsScalar) {
if (DoStore)
     EmitStoreOfScalar(ArgVal, lv, /* isInitialization */ true);
 
   setAddrOfLocalVar(&D, DeclPtr);
 
   // Emit debug info for param declarations in non-thunk functions.
   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});
     }
   }
 
   if (D.hasAttr<AnnotateAttr>())
     EmitVarAnnotations(&D, DeclPtr.getPointer());
 

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

+// 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:   )