Skip to content

Commit f1f7225

Browse files
epilkepilk
authored
[HeterogeneousDwarf] Fix constant (AS4) global emission (#4541)
Fixes SWDEV-560695
1 parent 1f47806 commit f1f7225

File tree

5 files changed

+67
-58
lines changed

5 files changed

+67
-58
lines changed

‎clang/test/CodeGenHIP/debug-info-diop-in-diexpression_dwarf.hip‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -62,28 +62,28 @@ __device__ void func1(int Arg) {
6262
// CHECK: DW_AT_type ("int")
6363
// CHECK: DW_AT_external (true)
6464
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_global)
65-
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
65+
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
6666

6767
// CHECK: DW_TAG_variable
6868
// CHECK: DW_AT_name ("GlobalDeviceB")
6969
// CHECK: DW_AT_type ("int")
7070
// CHECK: DW_AT_external (true)
7171
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_global)
72-
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
72+
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
7373

7474
// CHECK: DW_TAG_variable
7575
// CHECK: DW_AT_name ("GlobalConstantA")
7676
// CHECK: DW_AT_type ("int")
7777
// CHECK: DW_AT_external (true)
7878
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_constant)
79-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
79+
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
8080

8181
// CHECK: DW_TAG_variable
8282
// CHECK: DW_AT_name ("GlobalConstantB")
8383
// CHECK: DW_AT_type ("int")
8484
// CHECK: DW_AT_external (true)
8585
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_constant)
86-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
86+
// CHECK: DW_AT_location (DW_OP_addr 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
8787

8888
// CHECK: DW_TAG_subprogram
8989
// CHECK: DW_AT_linkage_name ("_Z7kernel1i")
@@ -94,46 +94,46 @@ __device__ void func1(int Arg) {
9494
// CHECK: DW_AT_name ("KernelVarSharedA")
9595
// CHECK: DW_AT_type ("int")
9696
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_group)
97-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit0, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit3, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
97+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit0, DW_OP_plus, DW_OP_lit3, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
9898

9999
// CHECK: DW_TAG_variable
100100
// CHECK: DW_AT_name ("KernelVarSharedB")
101101
// CHECK: DW_AT_type ("int")
102102
// CHECK: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_group)
103-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit4, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit3, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
103+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit4, DW_OP_plus, DW_OP_lit3, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
104104

105105
// CHECK: DW_TAG_formal_parameter
106-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit0, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
106+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_lit0, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
107107
// CHECK: DW_AT_name ("Arg")
108108
// CHECK: DW_AT_type ("int")
109109

110110
// CHECK: DW_TAG_variable
111-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit4, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
111+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_lit4, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
112112
// CHECK: DW_AT_name ("KernelVarA")
113113
// CHECK: DW_AT_type ("int")
114114

115115
// CHECK: DW_TAG_variable
116-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit8, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
116+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_lit8, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
117117
// CHECK: DW_AT_name ("KernelVarB")
118118
// CHECK: DW_AT_type ("int")
119119

120120
// CHECK: DW_TAG_variable
121-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit16, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
121+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_lit16, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
122122
// CHECK: DW_AT_name ("KernelVarSharedAPointer")
123123
// CHECK: DW_AT_type ("int *")
124124

125125
// CHECK: DW_TAG_variable
126-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit24, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
126+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_lit24, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
127127
// CHECK: DW_AT_name ("KernelVarSharedBPointer")
128128
// CHECK: DW_AT_type ("int *")
129129

130130
// CHECK: DW_TAG_variable
131-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_constu 0x20, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
131+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_constu 0x20, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
132132
// CHECK: DW_AT_name ("KernelVarAPointer")
133133
// CHECK: DW_AT_type ("int *")
134134

135135
// CHECK: DW_TAG_variable
136-
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_constu 0x28, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
136+
// CHECK: DW_AT_location (DW_OP_lit0, DW_OP_lit6, DW_OP_shr, DW_OP_constu 0x28, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
137137
// CHECK: DW_AT_name ("KernelVarBPointer")
138138
// CHECK: DW_AT_type ("int *")
139139
// CHECK: NULL
@@ -144,27 +144,27 @@ __device__ void func1(int Arg) {
144144
// CHECK: DW_AT_external (true)
145145

146146
// CHECK: DW_TAG_formal_parameter
147-
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit0, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
147+
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit0, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
148148
// CHECK: DW_AT_name ("Arg")
149149
// CHECK: DW_AT_type ("int")
150150

151151
// CHECK: DW_TAG_variable
152-
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit4, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
152+
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit4, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
153153
// CHECK: DW_AT_name ("FuncVarA")
154154
// CHECK: DW_AT_type ("int")
155155

156156
// CHECK: DW_TAG_variable
157-
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit8, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
157+
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit8, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
158158
// CHECK: DW_AT_name ("FuncVarB")
159159
// CHECK: DW_AT_type ("int")
160160

161161
// CHECK: DW_TAG_variable
162-
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit16, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
162+
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit16, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
163163
// CHECK: DW_AT_name ("FuncVarAPointer")
164164
// CHECK: DW_AT_type ("int *")
165165

166166
// CHECK: DW_TAG_variable
167-
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit24, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x4, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
167+
// CHECK: DW_AT_location (DW_OP_regx 0x40, DW_OP_deref_size 0x4, DW_OP_lit6, DW_OP_shr, DW_OP_lit24, DW_OP_plus, DW_OP_lit5, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
168168
// CHECK: DW_AT_name ("FuncVarBPointer")
169169
// CHECK: DW_AT_type ("int *")
170170
// CHECK: NULL

‎llvm/lib/CodeGen/AsmPrinter/DwarfExpression.cpp‎

Lines changed: 20 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -929,7 +929,7 @@ NewOpResult DwarfExpression::convertValueKind(const NewOpResult &Res,
929929
}
930930

931931
if (Res.VK == ValueKind::LocationDesc && ReqVK == ValueKind::Value) {
932-
readToValue(Res.Ty);
932+
readToValue(Res);
933933
return {Res.Ty, ValueKind::Value, Res.DivergentAddrSpace};
934934
}
935935

@@ -949,11 +949,12 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Arg Arg,
949949
// address spaces, e.g. LDS. Generate a 'DW_OP_constu' with a dummy
950950
// constant value (0) for now.
951951
unsigned AMDGPUGlobalAddrSpace = 1;
952+
unsigned AMDGPUConstantAddrSpace = 4;
952953
if ((AP.TM.getTargetTriple().getArch() == Triple::amdgcn) &&
953-
(GV->getAddressSpace() != AMDGPUGlobalAddrSpace)) {
954+
(GV->getAddressSpace() != AMDGPUGlobalAddrSpace &&
955+
GV->getAddressSpace() != AMDGPUConstantAddrSpace)) {
954956
emitConstu(0);
955-
emitOp(dwarf::DW_OP_stack_value);
956-
return NewOpResult{Arg.getResultType(), ValueKind::LocationDesc};
957+
return NewOpResult{Arg.getResultType(), ValueKind::Value};
957958
}
958959

959960
// TODO: We only support PIC reloc-model and non-TLS globals so far, see
@@ -965,8 +966,7 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Arg Arg,
965966

966967
CU.getDwarfDebug().addArangeLabel(SymbolCU(&CU, AP.getSymbol(GV)));
967968
emitOpAddress(GV);
968-
emitOp(dwarf::DW_OP_stack_value);
969-
return NewOpResult{Arg.getResultType(), ValueKind::LocationDesc};
969+
return NewOpResult{Arg.getResultType(), ValueKind::Value};
970970
}
971971

972972
if (Entry.isLocation()) {
@@ -1128,7 +1128,7 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Convert Convert,
11281128
// If we're not dealing with the divergent address space case, Convert
11291129
// requires a value operand.
11301130
if (Child->VK == ValueKind::LocationDesc)
1131-
readToValue(Child->Ty);
1131+
readToValue(*Child);
11321132

11331133
uint64_t ToBits = DestTy->getPrimitiveSizeInBits().getFixedValue();
11341134
uint64_t FromBits = Child->Ty->getPrimitiveSizeInBits().getFixedValue();
@@ -1165,7 +1165,7 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::SExt SExt,
11651165

11661166
std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
11671167
ChildrenT Children) {
1168-
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc,
1168+
auto Child = traverse(Children[0].get(), ValueKind::Value,
11691169
/*PermitDivergentAddrSpace=*/true);
11701170
if (!Child)
11711171
return std::nullopt;
@@ -1180,11 +1180,6 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
11801180
unsigned PointerLLVMAddrSpace = Child->DivergentAddrSpace
11811181
? *Child->DivergentAddrSpace
11821182
: PointerResultType->getAddressSpace();
1183-
uint64_t PointerSizeInBits =
1184-
AP.getDataLayout().getPointerSizeInBits(PointerLLVMAddrSpace);
1185-
assert(PointerSizeInBits % 8 == 0 && "Expected multiple of 8");
1186-
1187-
uint64_t PointerSizeInBytes = PointerSizeInBits / 8;
11881183
auto PointerDWARFAddrSpace = AP.TM.mapToDWARFAddrSpace(PointerLLVMAddrSpace);
11891184
if (!PointerDWARFAddrSpace) {
11901185
LLVM_DEBUG(dbgs() << "Failed to lower DIOpDeref of pointer to addrspace("
@@ -1193,8 +1188,6 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
11931188
return std::nullopt;
11941189
}
11951190

1196-
emitOp(dwarf::DW_OP_deref_size);
1197-
emitData1(PointerSizeInBytes);
11981191
emitConstu(*PointerDWARFAddrSpace);
11991192
emitUserOp(dwarf::DW_OP_LLVM_form_aspace_address);
12001193

@@ -1212,13 +1205,13 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Read Read,
12121205
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc);
12131206
if (!Child)
12141207
return std::nullopt;
1215-
readToValue(Children[0].get());
1208+
readToValue(*Child);
12161209
return NewOpResult{Child->Ty, ValueKind::Value};
12171210
}
12181211

12191212
std::optional<NewOpResult>
12201213
DwarfExpression::traverse(DIOp::Reinterpret Reinterpret, ChildrenT Children) {
1221-
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc,
1214+
auto Child = traverse(Children[0].get(), /*ReqVK=*/std::nullopt,
12221215
/*PermitDivergentAddrSpace=*/true);
12231216
if (!Child)
12241217
return Child;
@@ -1301,26 +1294,21 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Fragment Fragment,
13011294
return std::nullopt;
13021295
}
13031296

1304-
void DwarfExpression::readToValue(Type *Ty) {
1305-
uint64_t PrimitiveSizeInBits = Ty->getPrimitiveSizeInBits();
1306-
assert(PrimitiveSizeInBits != 0 && "Expected primitive type");
1307-
1308-
uint64_t ByteAlignedPrimitiveSizeInBits = alignTo<8>(PrimitiveSizeInBits);
1309-
uint64_t PrimitiveSizeInBytes = ByteAlignedPrimitiveSizeInBits / 8;
1310-
bool NeedsMask = ByteAlignedPrimitiveSizeInBits != PrimitiveSizeInBits;
1297+
void DwarfExpression::readToValue(const OpResult &R) {
1298+
const DataLayout &DL = AP.getDataLayout();
1299+
uint64_t SizeInBits = R.Ty->isPointerTy() && R.DivergentAddrSpace
1300+
? DL.getPointerSizeInBits(*R.DivergentAddrSpace)
1301+
: DL.getTypeSizeInBits(R.Ty).getFixedValue();
1302+
uint64_t ByteAlignedSizeInBits = alignTo<8>(SizeInBits);
1303+
uint64_t SizeInBytes = ByteAlignedSizeInBits / 8;
1304+
bool NeedsMask = ByteAlignedSizeInBits != SizeInBits;
13111305

13121306
emitOp(dwarf::DW_OP_deref_size);
1313-
emitData1(PrimitiveSizeInBytes);
1307+
emitData1(SizeInBytes);
13141308

13151309
if (NeedsMask) {
1316-
uint64_t Mask = (1ULL << PrimitiveSizeInBits) - 1ULL;
1310+
uint64_t Mask = (1ULL << SizeInBits) - 1ULL;
13171311
emitConstu(Mask);
13181312
emitOp(dwarf::DW_OP_and);
13191313
}
13201314
}
1321-
1322-
void DwarfExpression::readToValue(DwarfExpression::Node *OpNode) {
1323-
assert(OpNode->isLowered() && "Expected lowered node");
1324-
assert(OpNode->getResultType() && "Expected non-null result type");
1325-
readToValue(OpNode->getResultType());
1326-
}

‎llvm/lib/CodeGen/AsmPrinter/DwarfExpression.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -401,8 +401,7 @@ class DwarfExpression {
401401
/// stack to RequiredVK. Nop if Res.VK is RequiredVK.
402402
OpResult convertValueKind(const OpResult &Res, ValueKind RequiredVK);
403403

404-
void readToValue(Type *Ty);
405-
void readToValue(Node *OpNode);
404+
void readToValue(const OpResult &R);
406405

407406
using ChildrenT = ArrayRef<std::unique_ptr<Node>>;
408407

‎llvm/test/DebugInfo/AMDGPU/heterogeneous-dwarf-diop-diexpression-address-spaces.ll‎

Lines changed: 23 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,22 @@
11
; RUN: llc -O0 -mcpu=gfx1030 -mtriple=amdgcn-amd-amdhsa -filetype=obj -o - < %s | llvm-dwarfdump --debug-info - | FileCheck %s
22

3+
@GlobMutable = protected addrspace(1) global i32 0, align 4, !dbg !39
4+
; CHECK-LABEL: DW_AT_name ("GlobMutable")
5+
; CHECK-NEXT: DW_AT_type
6+
; CHECK-NEXT: DW_AT_external
7+
; CHECK-NEXT: DW_AT_decl_file
8+
; CHECK-NEXT: DW_AT_decl_line
9+
; CHECK-NEXT: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_global)
10+
; CHECK-NEXT: DW_AT_location (DW_OP_addrx 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
11+
12+
@GlobConst = internal addrspace(4) constant i32 0, align 4, !dbg !41
13+
; CHECK-LABEL: DW_AT_name ("GlobConst")
14+
; CHECK-NEXT: DW_AT_type
15+
; CHECK-NEXT: DW_AT_decl_file
16+
; CHECK-NEXT: DW_AT_decl_line
17+
; CHECK-NEXT: DW_AT_LLVM_memory_space (DW_MSPACE_LLVM_constant)
18+
; CHECK-NEXT: DW_AT_location (DW_OP_addrx 0x1, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
19+
320
; CHECK-LABEL: DW_AT_name ("test_loc_single")
421
define void @test_loc_single(ptr addrspace(3) %ptr) #0 !dbg !9 {
522
; Verify that the right address class attribute is attached to the variable's
@@ -122,7 +139,7 @@ attributes #0 = { "frame-pointer"="all" }
122139
!llvm.module.flags = !{!2, !3, !4, !5, !6, !7}
123140
!llvm.ident = !{!8}
124141

125-
!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus_14, file: !1, producer: "clang version 19.0.0", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
142+
!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus_14, file: !1, producer: "clang version 19.0.0", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None, globals: !38)
126143
!1 = !DIFile(filename: "t.cpp", directory: "/")
127144
!2 = !{i32 1, !"amdhsa_code_object_version", i32 500}
128145
!3 = !{i32 7, !"Dwarf Version", i32 5}
@@ -160,3 +177,8 @@ attributes #0 = { "frame-pointer"="all" }
160177
!35 = !{!36}
161178
!36 = !DILocalVariable(name: "not_divergent", scope: !34, file: !1, line: 1, type: !14)
162179
!37 = !DILocation(line: 1, column: 1, scope: !34)
180+
!38 = !{!39, !41}
181+
!39 = !DIGlobalVariableExpression(var: !40, expr: !DIExpression(DIOpArg(0, ptr addrspace(1)), DIOpDeref(i32)))
182+
!40 = distinct !DIGlobalVariable(name: "GlobMutable", linkageName: "GlobMutable", scope: !0, file: !1, line: 1, type: !15, isLocal: false, isDefinition: true, memorySpace: DW_MSPACE_LLVM_global)
183+
!41 = !DIGlobalVariableExpression(var: !42, expr: !DIExpression(DIOpArg(0, ptr addrspace(4)), DIOpDeref(i32)))
184+
!42 = distinct !DIGlobalVariable(name: "GlobConst", linkageName: "GlobConst", scope: !0, file: !1, line: 1, type: !15, isLocal: true, isDefinition: true, memorySpace: DW_MSPACE_LLVM_constant)

‎llvm/test/DebugInfo/heterogeneous-diop-in-diexpression.ll‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -6,21 +6,21 @@
66

77
; DWARF: DW_TAG_variable
88
; DWARF: DW_AT_name [DW_FORM_strx1] (indexed ([[#%x,]]) string = "glob")
9-
; DWARF: DW_AT_location [DW_FORM_exprloc] (DW_OP_addrx 0x0, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
9+
; DWARF: DW_AT_location [DW_FORM_exprloc] (DW_OP_addrx 0x0, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address)
1010

1111
; DWARF: DW_TAG_variable
1212
; DWARF: DW_AT_name [DW_FORM_strx1] (indexed ([[#%x,]]) string = "glob_fragmented")
13-
; DWARF: DW_AT_location [DW_FORM_exprloc] (DW_OP_addrx 0x1, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2, DW_OP_addrx 0x2, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2)
13+
; DWARF: DW_AT_location [DW_FORM_exprloc] (DW_OP_addrx 0x1, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2, DW_OP_addrx 0x2, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2)
1414

1515
; DWARF: DW_TAG_variable
1616
; DWARF: DW_AT_location [DW_FORM_loclistx] (indexed (0x[[#%x,]]) loclist = 0x[[#%x,]]:
17-
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -4, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address
17+
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -4, DW_OP_plus, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address
1818
; DWARF: DW_AT_name [DW_FORM_strx1] (indexed ([[#%x,]]) string = "var")
1919

2020
; DWARF: DW_TAG_variable
2121
; DWARF: DW_AT_location [DW_FORM_loclistx] (indexed (0x[[#%x,]]) loclist = 0x[[#%x,]]:
22-
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -8, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address
23-
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -8, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2, DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -6, DW_OP_plus, DW_OP_stack_value, DW_OP_deref_size 0x8, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2)
22+
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -8, DW_OP_plus, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address
23+
; DWARF: [0x[[#%x,]], 0x[[#%x,]]) ".text": DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -8, DW_OP_plus, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2, DW_OP_reg6 RBP, DW_OP_deref_size 0x8, DW_OP_consts -6, DW_OP_plus, DW_OP_lit0, DW_OP_LLVM_user DW_OP_LLVM_form_aspace_address, DW_OP_piece 0x2)
2424
; DWARF: DW_AT_name [DW_FORM_strx1] (indexed ([[#%x,]]) string = "var_fragmented")
2525

2626
; ModuleID = '<stdin>'

0 commit comments

Comments
 (0)