[clang] 6278682 - In spir functions, llvm.dbg.declare intrinsics created
Ammarguellat, Zahira via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 8 04:36:35 PST 2021
https://reviews.llvm.org/D112963
-----Original Message-----
From: Aaron Ballman <aaron at aaronballman.com>
Sent: Monday, November 8, 2021 6:51 AM
To: Ammarguellat, Zahira <zahira.ammarguellat at intel.com>; Zahira Ammarguellat <llvmlistbot at llvm.org>
Cc: cfe-commits <cfe-commits at lists.llvm.org>
Subject: Re: [clang] 6278682 - In spir functions, llvm.dbg.declare intrinsics created
Hello! Was this code reviewed anywhere? I can't seem to spot a review for it, so wondering if I missed something.
Thanks!
~Aaron
On Fri, Nov 5, 2021 at 6:08 PM Zahira Ammarguellat via cfe-commits <cfe-commits at lists.llvm.org> wrote:
>
>
> Author: Zahira Ammarguellat
> Date: 2021-11-05T15:08:09-07:00
> New Revision: 627868263cd4d57c230b61904483a3dad9e1a1da
>
> URL:
> https://github.com/llvm/llvm-project/commit/627868263cd4d57c230b619044
> 83a3dad9e1a1da
> DIFF:
> https://github.com/llvm/llvm-project/commit/627868263cd4d57c230b619044
> 83a3dad9e1a1da.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: )
>
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
More information about the cfe-commits
mailing list