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.ammarguel...@intel.com> wrote: > > https://reviews.llvm.org/D112963 > > -----Original Message----- > From: Aaron Ballman <aa...@aaronballman.com> > Sent: Monday, November 8, 2021 6:51 AM > To: Ammarguellat, Zahira <zahira.ammarguel...@intel.com>; Zahira Ammarguellat > <llvmlist...@llvm.org> > Cc: cfe-commits <cfe-commits@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@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@lists.llvm.org > > https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
