Skip to content

Commit 6541ffb

Browse files
ronliebronlieb
authored
Merge branch 'amd-debug' into amd-staging (#636)
2 parents 4b7ddf4 + 5cf3ecb commit 6541ffb

31 files changed

+2875
-70103
lines changed

‎clang/lib/CodeGen/CGDebugInfo.cpp‎

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -5598,7 +5598,7 @@ llvm::DILocalVariable *CGDebugInfo::EmitDeclareForHeterogeneousDwarf(
55985598
} else if (const auto *RT = dyn_cast<RecordType>(VD->getType())) {
55995599
// If VD is an anonymous union then Storage represents value for
56005600
// all union fields.
5601-
const RecordDecl *RD = RT->getOriginalDecl()->getDefinitionOrSelf();
5601+
const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
56025602
if (RD->isUnion() && RD->isAnonymousStructOrUnion()) {
56035603
llvm::DIExprBuilder UnionExprBuilder{ExprBuilder};
56045604
llvm::DIExpression *UnionDIExpression = UnionExprBuilder.intoExpression();
@@ -6161,8 +6161,9 @@ llvm::DIGlobalVariableExpression *CGDebugInfo::CollectAnonRecordDecls(
61616161
// Ignore unnamed fields, but recurse into anonymous records.
61626162
if (FieldName.empty()) {
61636163
if (const auto *RT = dyn_cast<RecordType>(Field->getType()))
6164-
GVE = CollectAnonRecordDecls(RT->getDecl()->getDefinitionOrSelf(), Unit,
6165-
LineNo, LinkageName, MS, Var, DContext);
6164+
GVE = CollectAnonRecordDecls(
6165+
RT->getDecl()->getDefinitionOrSelf(), Unit, LineNo,
6166+
LinkageName, MS, Var, DContext);
61666167
continue;
61676168
}
61686169
// Use VarDecl's Tag, Scope and Line number.
@@ -6191,7 +6192,7 @@ CGDebugInfo::CollectAnonRecordDeclsForHeterogeneousDwarf(
61916192
if (FieldName.empty()) {
61926193
if (const auto *RT = dyn_cast<RecordType>(Field->getType()))
61936194
GVE = CollectAnonRecordDeclsForHeterogeneousDwarf(
6194-
RT->getOriginalDecl()->getDefinitionOrSelf(), Unit, LineNo,
6195+
RT->getDecl()->getDefinitionOrSelf(), Unit, LineNo,
61956196
LinkageName, MS, Var, DContext);
61966197
continue;
61976198
}
@@ -6547,7 +6548,7 @@ void CGDebugInfo::EmitGlobalVariableForHeterogeneousDwarf(
65476548
llvm::dwarf::MemorySpace MS = getDWARFMemorySpace(D);
65486549
if (T->isUnionType() && DeclName.empty()) {
65496550
const RecordDecl *RD =
6550-
T->castAs<RecordType>()->getOriginalDecl()->getDefinitionOrSelf();
6551+
T->castAs<RecordType>()->getDecl()->getDefinitionOrSelf();
65516552
assert(RD->isAnonymousStructOrUnion() &&
65526553
"unnamed non-anonymous struct or union?");
65536554
// FIXME(KZHURAVL): No tests for this path.

‎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
@@ -953,7 +953,7 @@ NewOpResult DwarfExpression::convertValueKind(const NewOpResult &Res,
953953
}
954954

955955
if (Res.VK == ValueKind::LocationDesc && ReqVK == ValueKind::Value) {
956-
readToValue(Res.Ty);
956+
readToValue(Res);
957957
return {Res.Ty, ValueKind::Value, Res.DivergentAddrSpace};
958958
}
959959

@@ -973,11 +973,12 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Arg Arg,
973973
// address spaces, e.g. LDS. Generate a 'DW_OP_constu' with a dummy
974974
// constant value (0) for now.
975975
unsigned AMDGPUGlobalAddrSpace = 1;
976+
unsigned AMDGPUConstantAddrSpace = 4;
976977
if ((AP.TM.getTargetTriple().getArch() == Triple::amdgcn) &&
977-
(GV->getAddressSpace() != AMDGPUGlobalAddrSpace)) {
978+
(GV->getAddressSpace() != AMDGPUGlobalAddrSpace &&
979+
GV->getAddressSpace() != AMDGPUConstantAddrSpace)) {
978980
emitConstu(0);
979-
emitOp(dwarf::DW_OP_stack_value);
980-
return NewOpResult{Arg.getResultType(), ValueKind::LocationDesc};
981+
return NewOpResult{Arg.getResultType(), ValueKind::Value};
981982
}
982983

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

990991
CU.getDwarfDebug().addArangeLabel(SymbolCU(&CU, AP.getSymbol(GV)));
991992
emitOpAddress(GV);
992-
emitOp(dwarf::DW_OP_stack_value);
993-
return NewOpResult{Arg.getResultType(), ValueKind::LocationDesc};
993+
return NewOpResult{Arg.getResultType(), ValueKind::Value};
994994
}
995995

996996
if (Entry.isLocation()) {
@@ -1152,7 +1152,7 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Convert Convert,
11521152
// If we're not dealing with the divergent address space case, Convert
11531153
// requires a value operand.
11541154
if (Child->VK == ValueKind::LocationDesc)
1155-
readToValue(Child->Ty);
1155+
readToValue(*Child);
11561156

11571157
uint64_t ToBits = DestTy->getPrimitiveSizeInBits().getFixedValue();
11581158
uint64_t FromBits = Child->Ty->getPrimitiveSizeInBits().getFixedValue();
@@ -1189,7 +1189,7 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::SExt SExt,
11891189

11901190
std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
11911191
ChildrenT Children) {
1192-
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc,
1192+
auto Child = traverse(Children[0].get(), ValueKind::Value,
11931193
/*PermitDivergentAddrSpace=*/true);
11941194
if (!Child)
11951195
return std::nullopt;
@@ -1204,11 +1204,6 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
12041204
unsigned PointerLLVMAddrSpace = Child->DivergentAddrSpace
12051205
? *Child->DivergentAddrSpace
12061206
: PointerResultType->getAddressSpace();
1207-
uint64_t PointerSizeInBits =
1208-
AP.getDataLayout().getPointerSizeInBits(PointerLLVMAddrSpace);
1209-
assert(PointerSizeInBits % 8 == 0 && "Expected multiple of 8");
1210-
1211-
uint64_t PointerSizeInBytes = PointerSizeInBits / 8;
12121207
auto PointerDWARFAddrSpace = AP.TM.mapToDWARFAddrSpace(PointerLLVMAddrSpace);
12131208
if (!PointerDWARFAddrSpace) {
12141209
LLVM_DEBUG(dbgs() << "Failed to lower DIOpDeref of pointer to addrspace("
@@ -1217,8 +1212,6 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Deref Deref,
12171212
return std::nullopt;
12181213
}
12191214

1220-
emitOp(dwarf::DW_OP_deref_size);
1221-
emitData1(PointerSizeInBytes);
12221215
emitConstu(*PointerDWARFAddrSpace);
12231216
emitUserOp(dwarf::DW_OP_LLVM_form_aspace_address);
12241217

@@ -1236,13 +1229,13 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Read Read,
12361229
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc);
12371230
if (!Child)
12381231
return std::nullopt;
1239-
readToValue(Children[0].get());
1232+
readToValue(*Child);
12401233
return NewOpResult{Child->Ty, ValueKind::Value};
12411234
}
12421235

12431236
std::optional<NewOpResult>
12441237
DwarfExpression::traverse(DIOp::Reinterpret Reinterpret, ChildrenT Children) {
1245-
auto Child = traverse(Children[0].get(), ValueKind::LocationDesc,
1238+
auto Child = traverse(Children[0].get(), /*ReqVK=*/std::nullopt,
12461239
/*PermitDivergentAddrSpace=*/true);
12471240
if (!Child)
12481241
return Child;
@@ -1325,26 +1318,21 @@ std::optional<NewOpResult> DwarfExpression::traverse(DIOp::Fragment Fragment,
13251318
return std::nullopt;
13261319
}
13271320

1328-
void DwarfExpression::readToValue(Type *Ty) {
1329-
uint64_t PrimitiveSizeInBits = Ty->getPrimitiveSizeInBits();
1330-
assert(PrimitiveSizeInBits != 0 && "Expected primitive type");
1331-
1332-
uint64_t ByteAlignedPrimitiveSizeInBits = alignTo<8>(PrimitiveSizeInBits);
1333-
uint64_t PrimitiveSizeInBytes = ByteAlignedPrimitiveSizeInBits / 8;
1334-
bool NeedsMask = ByteAlignedPrimitiveSizeInBits != PrimitiveSizeInBits;
1321+
void DwarfExpression::readToValue(const OpResult &R) {
1322+
const DataLayout &DL = AP.getDataLayout();
1323+
uint64_t SizeInBits = R.Ty->isPointerTy() && R.DivergentAddrSpace
1324+
? DL.getPointerSizeInBits(*R.DivergentAddrSpace)
1325+
: DL.getTypeSizeInBits(R.Ty).getFixedValue();
1326+
uint64_t ByteAlignedSizeInBits = alignTo<8>(SizeInBits);
1327+
uint64_t SizeInBytes = ByteAlignedSizeInBits / 8;
1328+
bool NeedsMask = ByteAlignedSizeInBits != SizeInBits;
13351329

13361330
emitOp(dwarf::DW_OP_deref_size);
1337-
emitData1(PrimitiveSizeInBytes);
1331+
emitData1(SizeInBytes);
13381332

13391333
if (NeedsMask) {
1340-
uint64_t Mask = (1ULL << PrimitiveSizeInBits) - 1ULL;
1334+
uint64_t Mask = (1ULL << SizeInBits) - 1ULL;
13411335
emitConstu(Mask);
13421336
emitOp(dwarf::DW_OP_and);
13431337
}
13441338
}
1345-
1346-
void DwarfExpression::readToValue(DwarfExpression::Node *OpNode) {
1347-
assert(OpNode->isLowered() && "Expected lowered node");
1348-
assert(OpNode->getResultType() && "Expected non-null result type");
1349-
readToValue(OpNode->getResultType());
1350-
}

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

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

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

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

0 commit comments

Comments
 (0)