[clang] 6278682 - In spir functions, llvm.dbg.declare intrinsics created
Aaron Ballman via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 8 05:38:15 PST 2021
Ah, I see what went wrong. That review was never signed up to the
cfe-commits mailing list, so the only people who ever saw that review
were the people explicitly on the reviewers and subscribers list. Next
time, please be sure to add cfe-commits to the subscribers list
explicitly.
~Aaron
On Mon, Nov 8, 2021 at 7:36 AM Ammarguellat, Zahira
<zahira.ammarguellat at intel.com> wrote:
>
> 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