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.
// 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();
// 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);
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);
}
}
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.
// 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;
: 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;
}
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});
}
--- /dev/null
+// 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: )