@@ -53,7 +53,7 @@ struct AMDGPUImplicitArgsTyCOV4 {
53
53
uint8_t Unused[56 ];
54
54
};
55
55
56
- uint32_t getImplicitArgsSize (uint16_t Version) {
56
+ inline uint32_t getImplicitArgsSize (uint16_t Version) {
57
57
return Version < ELF::ELFABIVERSION_AMDGPU_HSA_V5
58
58
? sizeof (AMDGPUImplicitArgsTyCOV4)
59
59
: sizeof (AMDGPUImplicitArgsTy);
@@ -173,44 +173,44 @@ class KernelInfoReader {
173
173
if (!V.first .isString ())
174
174
return Error::success ();
175
175
176
- const auto isKey = [](const msgpack::DocNode &DK, StringRef SK) {
176
+ const auto IsKey = [](const msgpack::DocNode &DK, StringRef SK) {
177
177
return DK.getString () == SK;
178
178
};
179
179
180
- const auto getSequenceOfThreeInts = [](msgpack::DocNode &DN,
180
+ const auto GetSequenceOfThreeInts = [](msgpack::DocNode &DN,
181
181
uint32_t *Vals) {
182
182
assert (DN.isArray () && " MsgPack DocNode is an array node" );
183
183
auto DNA = DN.getArray ();
184
184
assert (DNA.size () == 3 && " ArrayNode has at most three elements" );
185
185
186
- int i = 0 ;
186
+ int I = 0 ;
187
187
for (auto DNABegin = DNA.begin (), DNAEnd = DNA.end (); DNABegin != DNAEnd;
188
188
++DNABegin) {
189
- Vals[i ++] = DNABegin->getUInt ();
189
+ Vals[I ++] = DNABegin->getUInt ();
190
190
}
191
191
};
192
192
193
- if (isKey (V.first , " .name" )) {
193
+ if (IsKey (V.first , " .name" )) {
194
194
KernelName = V.second .toString ();
195
- } else if (isKey (V.first , " .sgpr_count" )) {
195
+ } else if (IsKey (V.first , " .sgpr_count" )) {
196
196
KernelData.SGPRCount = V.second .getUInt ();
197
- } else if (isKey (V.first , " .sgpr_spill_count" )) {
197
+ } else if (IsKey (V.first , " .sgpr_spill_count" )) {
198
198
KernelData.SGPRSpillCount = V.second .getUInt ();
199
- } else if (isKey (V.first , " .vgpr_count" )) {
199
+ } else if (IsKey (V.first , " .vgpr_count" )) {
200
200
KernelData.VGPRCount = V.second .getUInt ();
201
- } else if (isKey (V.first , " .vgpr_spill_count" )) {
201
+ } else if (IsKey (V.first , " .vgpr_spill_count" )) {
202
202
KernelData.VGPRSpillCount = V.second .getUInt ();
203
- } else if (isKey (V.first , " .private_segment_fixed_size" )) {
203
+ } else if (IsKey (V.first , " .private_segment_fixed_size" )) {
204
204
KernelData.PrivateSegmentSize = V.second .getUInt ();
205
- } else if (isKey (V.first , " .group_segement_fixed_size" )) {
205
+ } else if (IsKey (V.first , " .group_segement_fixed_size" )) {
206
206
KernelData.GroupSegmentList = V.second .getUInt ();
207
- } else if (isKey (V.first , " .reqd_workgroup_size" )) {
208
- getSequenceOfThreeInts (V.second , KernelData.RequestedWorkgroupSize );
209
- } else if (isKey (V.first , " .workgroup_size_hint" )) {
210
- getSequenceOfThreeInts (V.second , KernelData.WorkgroupSizeHint );
211
- } else if (isKey (V.first , " .wavefront_size" )) {
207
+ } else if (IsKey (V.first , " .reqd_workgroup_size" )) {
208
+ GetSequenceOfThreeInts (V.second , KernelData.RequestedWorkgroupSize );
209
+ } else if (IsKey (V.first , " .workgroup_size_hint" )) {
210
+ GetSequenceOfThreeInts (V.second , KernelData.WorkgroupSizeHint );
211
+ } else if (IsKey (V.first , " .wavefront_size" )) {
212
212
KernelData.WavefronSize = V.second .getUInt ();
213
- } else if (isKey (V.first , " .max_flat_workgroup_size" )) {
213
+ } else if (IsKey (V.first , " .max_flat_workgroup_size" )) {
214
214
KernelData.MaxFlatWorkgroupSize = V.second .getUInt ();
215
215
}
216
216
@@ -273,9 +273,10 @@ class KernelInfoReader {
273
273
274
274
// / Reads the AMDGPU specific metadata from the ELF file and propagates the
275
275
// / KernelInfoMap
276
- Error readAMDGPUMetaDataFromImage (MemoryBufferRef MemBuffer,
277
- StringMap<KernelMetaDataTy> &KernelInfoMap,
278
- uint16_t &ELFABIVersion) {
276
+ inline Error
277
+ readAMDGPUMetaDataFromImage (MemoryBufferRef MemBuffer,
278
+ StringMap<KernelMetaDataTy> &KernelInfoMap,
279
+ uint16_t &ELFABIVersion) {
279
280
Error Err = Error::success (); // Used later as out-parameter
280
281
281
282
auto ELFOrError = object::ELF64LEFile::create (MemBuffer.getBuffer ());
0 commit comments