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
