Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ clang/lib/CodeGen/CGDecl.cpp
@@ -1608,8 +1608,9 @@
     if (UsePointerValue)
       DebugAddr = ReturnValuePointer;
 
-    (void)DI->EmitDeclareOfAutoVariable(&D, DebugAddr.getPointer(), Builder,
-                                        UsePointerValue);
+    (void)DI->EmitDeclareOfAutoVariable(
+        &D, DebugAddr.getPointer()->stripPointerCasts(), Builder,
+        UsePointerValue);
   }
 
   if (D.hasAttr<AnnotateAttr>() && HaveInsertPoint())
@@ -2450,6 +2451,7 @@
   }
 
   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 +2466,7 @@
     // 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 +2503,11 @@
             : 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 +2583,7 @@
   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});
     }
Index: clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp
===================================================================
--- /dev/null
+++ 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:   )