[WIP] Attach debug intrinsics to allocas, and use correct address space

[scott.linder]

Summary

A dbg.declare for a local/parameter describes the hardware location of
the source variable's value. This matches up with the semantics of the
alloca for the variable, whereas any addrspacecast inserted in order to
implement some source-level notion of address spaces does not.

When creating the dbg.declare intrinsic, attach it directly to the
alloca, not to any addrspacecast.

Update the DIExpression with the address space of the alloca, rather
than use the address space associated with the source level type.

[scott.linder]

I need to add more tests, but I wanted to float the idea of the change and get feedback first.

[aprantl]

clang/lib/CodeGen/CGDecl.cpp

1579

This is unintuitive — can you add a comment explaining why it may not be valid and why address should only be used then?

[scott.linder]

clang/lib/CodeGen/CGDecl.cpp

1579

This is kind of a cop-out on my part, the only path where this occurs is for OpenMP, and I think I just need to understand better what is happening. This also occurs for NRVO, but that is explicitly called out just below this. I'll try to understand this more completely and see if I can represent the possibilities more direclty.

Somewhat related, it is a bit unsettling reading through this, as the invariant seems to be that address.isValid() by the time the call to setAddrOfLocalVar is called, which makes sense but isn't explicit anywhere in the multiple nested ifs. I'll also add an assert of that.

[scott.linder]

@ABataev Sorry if I'm pulling you in without enough context/work on my end, but I wanted to ask how the Clang codegen for OpenMP locals works at a high level?

Is the idea that instead of an alloc the frontend can insert calls into the runtime in some cases, like __kmpc_alloc (e.g. for firstprivate as in https://reviews.llvm.org/D5140)?

If that is the case, I assume there is no equivalent to SROA/Mem2Reg here? I am trying to understand conceptually where the debug info for the source level local should be tied to in the IR, and at least for locals which use alloc it has turned out to be much simpler to tie the variable directly to the alloc itself rather than bitcasts and things which obscure the relationship.

[ABataev]

In D88978#2325982, @scott.linder wrote:

@ABataev Sorry if I'm pulling you in without enough context/work on my end, but I wanted to ask how the Clang codegen for OpenMP locals works at a high level?

Is the idea that instead of an alloc the frontend can insert calls into the runtime in some cases, like __kmpc_alloc (e.g. for firstprivate as in https://reviews.llvm.org/D5140)?

Yes, right.

If that is the case, I assume there is no equivalent to SROA/Mem2Reg here?

I assume, no.

I am trying to understand conceptually where the debug info for the source level local should be tied to in the IR, and at least for locals which use alloc it has turned out to be much simpler to tie the variable directly to the alloc itself rather than bitcasts and things which obscure the relationship.

[scott.linder]

In D88978#2325991, @ABataev wrote:

In D88978#2325982, @scott.linder wrote:

@ABataev Sorry if I'm pulling you in without enough context/work on my end, but I wanted to ask how the Clang codegen for OpenMP locals works at a high level?

Is the idea that instead of an alloc the frontend can insert calls into the runtime in some cases, like __kmpc_alloc (e.g. for firstprivate as in https://reviews.llvm.org/D5140)?

Yes, right.

If that is the case, I assume there is no equivalent to SROA/Mem2Reg here?

I assume, no.

I am trying to understand conceptually where the debug info for the source level local should be tied to in the IR, and at least for locals which use alloc it has turned out to be much simpler to tie the variable directly to the alloc itself rather than bitcasts and things which obscure the relationship.

Ok, thank you! I think the simplest thing is for me to update the patch to tie the debug info to the call to the runtime allocator, then.

[jdoerfert]

I have no idea what from the commit message what this has to do with the OpenMP code gen but I saw this randomly while looking for something else so here we go:

In D88978#2326036, @scott.linder wrote:

In D88978#2325991, @ABataev wrote:

In D88978#2325982, @scott.linder wrote:

@ABataev Sorry if I'm pulling you in without enough context/work on my end, but I wanted to ask how the Clang codegen for OpenMP locals works at a high level?

Is the idea that instead of an alloc the frontend can insert calls into the runtime in some cases, like __kmpc_alloc (e.g. for firstprivate as in https://reviews.llvm.org/D5140)?

Yes, right.

The frontend does *not* insert __kmpc_alloc calls for firstprivate, or almost anything else for that matter. Grep clang/lib and you can find 2 uses, both in very specialized cases not related to "regular" user allocations". alloca is used as with basically everything else: https://clang.godbolt.org/z/z8fEqG

If that is the case, I assume there is no equivalent to SROA/Mem2Reg here?

I assume, no.

I am trying to understand conceptually where the debug info for the source level local should be tied to in the IR, and at least for locals which use alloc it has turned out to be much simpler to tie the variable directly to the alloc itself rather than bitcasts and things which obscure the relationship.

Ok, thank you! I think the simplest thing is for me to update the patch to tie the debug info to the call to the runtime allocator, then.

SROA/mem2reg is happening as you expect it to. FWIW, we also have heap2stack and argument-promotion + constant prop for parallel regions implemented in the Attributor. That means we would/will apply SROA/mem2reg even if you have a runtime alloca and if the value is nominally "shared" but could be made firtprivate.

[ABataev]

In D88978#2343484, @jdoerfert wrote:

I have no idea what from the commit message what this has to do with the OpenMP code gen but I saw this randomly while looking for something else so here we go:

In D88978#2326036, @scott.linder wrote:

In D88978#2325991, @ABataev wrote:

In D88978#2325982, @scott.linder wrote:

@ABataev Sorry if I'm pulling you in without enough context/work on my end, but I wanted to ask how the Clang codegen for OpenMP locals works at a high level?

Is the idea that instead of an alloc the frontend can insert calls into the runtime in some cases, like __kmpc_alloc (e.g. for firstprivate as in https://reviews.llvm.org/D5140)?

Yes, right.

The frontend does *not* insert __kmpc_alloc calls for firstprivate, or almost anything else for that matter. Grep clang/lib and you can find 2 uses, both in very specialized cases not related to "regular" user allocations". alloca is used as with basically everything else: https://clang.godbolt.org/z/z8fEqG

They are inserted for pragma allocate and privates with allocate clauses.

If that is the case, I assume there is no equivalent to SROA/Mem2Reg here?

I assume, no.

I am trying to understand conceptually where the debug info for the source level local should be tied to in the IR, and at least for locals which use alloc it has turned out to be much simpler to tie the variable directly to the alloc itself rather than bitcasts and things which obscure the relationship.

Ok, thank you! I think the simplest thing is for me to update the patch to tie the debug info to the call to the runtime allocator, then.

SROA/mem2reg is happening as you expect it to. FWIW, we also have heap2stack and argument-promotion + constant prop for parallel regions implemented in the Attributor. That means we would/will apply SROA/mem2reg even if you have a runtime alloca and if the value is nominally "shared" but could be made firtprivate.

[arsenm]

Is this still needed?

[scott.linder]

In D88978#2660274, @arsenm wrote:

Is this still needed?

Yes, I just got a little bogged down in the OMP code and haven't gotten back to it to finish it up. I anticipate needing to do this to soon, though.

[arsenm]

Is this still needed?

[arsenm]

Please rebase if still relevant

[scott.linder]

I'll open a new review as part of upstreaming all of the debug info work

clang/lib/CodeGen/CGDebugInfo.cpp

llvm::DILocalVariable *CGDebugInfo::EmitDeclare(const VarDecl *VD,
   }
   SmallVector<int64_t, 13> Expr;
   llvm::DINode::DIFlags Flags = llvm::DINode::FlagZero;
   if (VD->isImplicit())
     Flags |= llvm::DINode::FlagArtificial;
 
   auto Align = getDeclAlignIfRequired(VD, CGM.getContext());
 
-  unsigned AddressSpace = CGM.getContext().getTargetAddressSpace(VD->getType());
+  unsigned AddressSpace =
+      llvm::cast<llvm::PointerType>(Storage->getType())->getAddressSpace();
   AppendAddressSpaceXDeref(AddressSpace, Expr);
 
   // If this is implicit parameter of CXXThis or ObjCSelf kind, then give it an
   // object pointer flag.
   if (const auto *IPD = dyn_cast<ImplicitParamDecl>(VD)) {
     if (IPD->getParameterKind() == ImplicitParamDecl::CXXThis ||
         IPD->getParameterKind() == ImplicitParamDecl::ObjCSelf)
       Flags |= llvm::DINode::FlagObjectPointer;

clang/lib/CodeGen/CGDecl.cpp

CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
   }
 
   setAddrOfLocalVar(&D, address);
   emission.Addr = address;
   emission.AllocaAddr = AllocaAddr;
 
   // Emit debug info for local var declaration.
   if (EmitDebugInfo && HaveInsertPoint()) {
-    Address DebugAddr = address;
+    Address DebugAddr = AllocaAddr.isValid() ? AllocaAddr : address;
     bool UsePointerValue = NRVO && ReturnValuePointer.isValid();
     DI->setLocation(D.getLocation());
 
     // If NRVO, use a pointer to the return address.
     if (UsePointerValue)
       DebugAddr = ReturnValuePointer;
 
     (void)DI->EmitDeclareOfAutoVariable(&D, DebugAddr.getPointer(), Builder,
if (BlockInfo) {
                            ? Builder.CreateLoad(Arg.getIndirectAddress())
                            : Arg.getDirectValue();
       setBlockContextParameter(IPD, ArgNo, V);
       return;
     }
   }
 
   Address DeclPtr = Address::invalid();
+  Address DebugAddr = Address::invalid();
   bool DoStore = false;
   bool IsScalar = hasScalarEvaluationKind(Ty);
   // If we already have a pointer to the argument, reuse the input pointer.
   if (Arg.isIndirect()) {
-    DeclPtr = Arg.getIndirectAddress();
+    DeclPtr = DebugAddr = Arg.getIndirectAddress();
     // If we have a prettier pointer type at this point, bitcast to that.
     unsigned AS = DeclPtr.getType()->getAddressSpace();
     llvm::Type *IRTy = ConvertTypeForMem(Ty)->getPointerTo(AS);
     if (DeclPtr.getType() != IRTy)
       DeclPtr = Builder.CreateBitCast(DeclPtr, IRTy, D.getName());
     // Indirect argument is in alloca address space, which may be different
     // from the default address space.
     auto AllocaAS = CGM.getASTAllocaAddressSpace();
if (Arg.isIndirect()) {
     }
   } else {
     // Check if the parameter address is controlled by OpenMP runtime.
     Address OpenMPLocalAddr =
         getLangOpts().OpenMP
             ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D)
             : Address::invalid();
     if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
-      DeclPtr = OpenMPLocalAddr;
+      DeclPtr = DebugAddr = OpenMPLocalAddr;
     } else {
       // Otherwise, create a temporary to hold the value.
       DeclPtr = CreateMemTemp(Ty, getContext().getDeclAlign(&D),
-                              D.getName() + ".addr");
+                              D.getName() + ".addr", &DebugAddr);
     }
     DoStore = true;
   }
 
   llvm::Value *ArgVal = (DoStore ? Arg.getDirectValue() : nullptr);
 
   LValue lv = MakeAddrLValue(DeclPtr, Ty);
   if (IsScalar) {
void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
   if (DoStore)
     EmitStoreOfScalar(ArgVal, lv, /* isInitialization */ true);
 
   setAddrOfLocalVar(&D, DeclPtr);
 
   // Emit debug info for param declarations in non-thunk functions.
   if (CGDebugInfo *DI = getDebugInfo()) {
     if (CGM.getCodeGenOpts().hasReducedDebugInfo() && !CurFuncIsThunk) {
-      DI->EmitDeclareOfArgVariable(&D, DeclPtr.getPointer(), ArgNo, Builder);
+      DI->EmitDeclareOfArgVariable(&D, DebugAddr.getPointer(), ArgNo, Builder);
     }
   }
 
   if (D.hasAttr<AnnotateAttr>())
     EmitVarAnnotations(&D, DeclPtr.getPointer());
 
   // We can only check return value nullability if all arguments to the
   // function satisfy their nullability preconditions. This makes it necessary

clang/test/CodeGenHIP/debug-info-address-class.hip

 // CHECK-DAG: ![[FILEVAR1:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR1]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
 __device__ __shared__ int FileVar1;
 // CHECK-DAG: ![[FILEVAR2:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR2]], expr: !DIExpression())
 __device__ __constant__ int FileVar2;
 
 __device__ void kernel1(
-    // FIXME This should be in the private address space.
     // CHECK-DAG: ![[ARG:[0-9]+]] = !DILocalVariable(name: "Arg", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
-    // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[ARG]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+    // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(5)* {{.*}}, metadata ![[ARG]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
     int Arg) {
     // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true)
     // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR0]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
   __shared__ int FuncVar0;
-
-  // FIXME This should be in the private address space.
   // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
-  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(5)* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   int FuncVar1;
 }