scott.linder created this revision.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
scott.linder requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added a subscriber: sstefan1.
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.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D88978
Files:
clang/lib/CodeGen/CGDebugInfo.cpp
clang/lib/CodeGen/CGDecl.cpp
clang/test/CodeGenHIP/debug-info-address-class.hip
Index: clang/test/CodeGenHIP/debug-info-address-class.hip
===================================================================
--- clang/test/CodeGenHIP/debug-info-address-class.hip
+++ clang/test/CodeGenHIP/debug-info-address-class.hip
@@ -16,16 +16,13 @@
__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;
}
Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ clang/lib/CodeGen/CGDecl.cpp
@@ -1576,7 +1576,7 @@
// 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());
@@ -2417,11 +2417,12 @@
}
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);
@@ -2466,11 +2467,11 @@
? 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;
}
@@ -2545,7 +2546,7 @@
// 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);
}
}
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -4186,7 +4186,8 @@
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
Index: clang/test/CodeGenHIP/debug-info-address-class.hip
===================================================================
--- clang/test/CodeGenHIP/debug-info-address-class.hip
+++ clang/test/CodeGenHIP/debug-info-address-class.hip
@@ -16,16 +16,13 @@
__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;
}
Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ clang/lib/CodeGen/CGDecl.cpp
@@ -1576,7 +1576,7 @@
// 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());
@@ -2417,11 +2417,12 @@
}
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);
@@ -2466,11 +2467,11 @@
? 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;
}
@@ -2545,7 +2546,7 @@
// 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);
}
}
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -4186,7 +4186,8 @@
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
