[clang] 6278682 - In spir functions, llvm.dbg.declare intrinsics created
Zahira Ammarguellat via cfe-commits
cfe-commits at lists.llvm.org
Fri Nov 5 15:08:21 PDT 2021
Author: Zahira Ammarguellat
Date: 2021-11-05T15:08:09-07:00
New Revision: 627868263cd4d57c230b61904483a3dad9e1a1da
URL: https://github.com/llvm/llvm-project/commit/627868263cd4d57c230b61904483a3dad9e1a1da
DIFF: https://github.com/llvm/llvm-project/commit/627868263cd4d57c230b61904483a3dad9e1a1da.diff
LOG: In spir functions, llvm.dbg.declare intrinsics created
for parameters and locals need to refer to the stack
allocation in the alloca address space.
Added:
clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp
Modified:
clang/lib/CodeGen/CGDecl.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index dfb74a3fc6547..941671c614824 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/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});
}
diff --git a/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp b/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp
new file mode 100644
index 0000000000000..e6efa92716fbc
--- /dev/null
+++ b/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: )
More information about the cfe-commits
mailing list