1//===- Utility.cpp ------ Collection of generic offloading utilities ------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8
9#include "llvm/Frontend/Offloading/Utility.h"
10#include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
11#include "llvm/BinaryFormat/ELF.h"
12#include "llvm/BinaryFormat/MsgPackDocument.h"
13#include "llvm/IR/Constants.h"
14#include "llvm/IR/GlobalValue.h"
15#include "llvm/IR/GlobalVariable.h"
16#include "llvm/IR/Value.h"
17#include "llvm/Object/ELFObjectFile.h"
18#include "llvm/Object/OffloadBinary.h"
19#include "llvm/ObjectYAML/yaml2obj.h"
20#include "llvm/Support/MemoryBufferRef.h"
21#include "llvm/Transforms/Utils/ModuleUtils.h"
22
23using namespace llvm;
24using namespace llvm::offloading;
25using namespace llvm::offloading::sycl;
26
27StructType *offloading::getEntryTy(Module &M) {
28 LLVMContext &C = M.getContext();
29 StructType *EntryTy =
30 StructType::getTypeByName(C, Name: "struct.__tgt_offload_entry");
31 if (!EntryTy)
32 EntryTy = StructType::create(
33 Name: "struct.__tgt_offload_entry", elt1: Type::getInt64Ty(C), elts: Type::getInt16Ty(C),
34 elts: Type::getInt16Ty(C), elts: Type::getInt32Ty(C), elts: PointerType::getUnqual(C),
35 elts: PointerType::getUnqual(C), elts: Type::getInt64Ty(C), elts: Type::getInt64Ty(C),
36 elts: PointerType::getUnqual(C));
37 return EntryTy;
38}
39
40std::pair<Constant *, GlobalVariable *>
41offloading::getOffloadingEntryInitializer(Module &M, object::OffloadKind Kind,
42 Constant *Addr, StringRef Name,
43 uint64_t Size, uint32_t Flags,
44 uint64_t Data, Constant *AuxAddr) {
45 const llvm::Triple &Triple = M.getTargetTriple();
46 Type *PtrTy = PointerType::getUnqual(C&: M.getContext());
47 Type *Int64Ty = Type::getInt64Ty(C&: M.getContext());
48 Type *Int32Ty = Type::getInt32Ty(C&: M.getContext());
49 Type *Int16Ty = Type::getInt16Ty(C&: M.getContext());
50
51 Constant *AddrName = ConstantDataArray::getString(Context&: M.getContext(), Initializer: Name);
52
53 StringRef Prefix =
54 Triple.isNVPTX() ? "$offloading$entry_name" : ".offloading.entry_name";
55
56 // Create the constant string used to look up the symbol in the device.
57 auto *Str =
58 new GlobalVariable(M, AddrName->getType(), /*isConstant=*/true,
59 GlobalValue::InternalLinkage, AddrName, Prefix);
60 StringRef SectionName = ".llvm.rodata.offloading";
61 Str->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
62 Str->setSection(SectionName);
63 Str->setAlignment(Align(1));
64
65 // Make a metadata node for these constants so it can be queried from IR.
66 NamedMDNode *MD = M.getOrInsertNamedMetadata(Name: "llvm.offloading.symbols");
67 Metadata *MDVals[] = {ConstantAsMetadata::get(C: Str)};
68 MD->addOperand(M: llvm::MDNode::get(Context&: M.getContext(), MDs: MDVals));
69
70 // Construct the offloading entry.
71 Constant *EntryData[] = {
72 ConstantExpr::getNullValue(Ty: Int64Ty),
73 ConstantInt::get(Ty: Int16Ty, V: 1),
74 ConstantInt::get(Ty: Int16Ty, V: Kind),
75 ConstantInt::get(Ty: Int32Ty, V: Flags),
76 ConstantExpr::getPointerBitCastOrAddrSpaceCast(C: Addr, Ty: PtrTy),
77 ConstantExpr::getPointerBitCastOrAddrSpaceCast(C: Str, Ty: PtrTy),
78 ConstantInt::get(Ty: Int64Ty, V: Size),
79 ConstantInt::get(Ty: Int64Ty, V: Data),
80 AuxAddr ? ConstantExpr::getPointerBitCastOrAddrSpaceCast(C: AuxAddr, Ty: PtrTy)
81 : ConstantExpr::getNullValue(Ty: PtrTy)};
82 Constant *EntryInitializer = ConstantStruct::get(T: getEntryTy(M), V: EntryData);
83 return {EntryInitializer, Str};
84}
85
86StringRef offloading::getOffloadEntrySection(Module &M) {
87 return M.getTargetTriple().isOSBinFormatMachO() ? "__LLVM,offload_entries"
88 : "llvm_offload_entries";
89}
90
91/// Returns the start/end symbol names for iterating offloading entries in a
92/// given section. Mach-O uses \1section$start$/\1section$end$ convention;
93/// ELF/COFF use __start_/__stop_ prefixes.
94static std::pair<std::string, std::string>
95getOffloadEntryBoundarySymbols(const Triple &T, StringRef SectionName) {
96 if (T.isOSBinFormatMachO()) {
97 std::string SymSection = SectionName.str();
98 std::replace(first: SymSection.begin(), last: SymSection.end(), old_value: ',', new_value: '$');
99 return {"\1section$start$" + SymSection, "\1section$end$" + SymSection};
100 }
101 return {("__start_" + SectionName).str(), ("__stop_" + SectionName).str()};
102}
103
104GlobalVariable *offloading::emitOffloadingEntry(
105 Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name,
106 uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr) {
107 const llvm::Triple &Triple = M.getTargetTriple();
108 StringRef SectionName = getOffloadEntrySection(M);
109
110 auto [EntryInitializer, NameGV] = getOffloadingEntryInitializer(
111 M, Kind, Addr, Name, Size, Flags, Data, AuxAddr);
112
113 StringRef Prefix =
114 Triple.isNVPTX() ? "$offloading$entry$" : ".offloading.entry.";
115 auto *Entry = new GlobalVariable(
116 M, getEntryTy(M),
117 /*isConstant=*/true, GlobalValue::WeakAnyLinkage, EntryInitializer,
118 Prefix + Name, nullptr, GlobalValue::NotThreadLocal,
119 M.getDataLayout().getDefaultGlobalsAddressSpace());
120
121 // The entry has to be created in the section the linker expects it to be.
122 if (Triple.isOSBinFormatCOFF())
123 Entry->setSection((SectionName + "$OE").str());
124 else
125 Entry->setSection(SectionName);
126 Entry->setAlignment(Align(object::OffloadBinary::getAlignment()));
127 return Entry;
128}
129
130std::pair<Constant *, Constant *> offloading::getOffloadEntryArray(Module &M) {
131 const llvm::Triple &Triple = M.getTargetTriple();
132 StringRef SectionName = getOffloadEntrySection(M);
133
134 constexpr unsigned COFFSentinelEntryCount = 1;
135 unsigned EntryCount =
136 Triple.isOSBinFormatCOFF() ? COFFSentinelEntryCount : 0u;
137 auto *ZeroInitializer =
138 ConstantAggregateZero::get(Ty: ArrayType::get(ElementType: getEntryTy(M), NumElements: EntryCount));
139 auto *EntryInit = Triple.isOSBinFormatCOFF() ? ZeroInitializer : nullptr;
140 auto *EntryType = ZeroInitializer->getType();
141 auto Linkage = Triple.isOSBinFormatCOFF() ? GlobalValue::WeakODRLinkage
142 : GlobalValue::ExternalLinkage;
143
144 auto [StartName, StopName] =
145 getOffloadEntryBoundarySymbols(T: Triple, SectionName);
146
147 auto *EntriesB = new GlobalVariable(M, EntryType, /*isConstant=*/true,
148 Linkage, EntryInit, StartName);
149 EntriesB->setVisibility(GlobalValue::HiddenVisibility);
150 auto *EntriesE = new GlobalVariable(M, EntryType, /*isConstant=*/true,
151 Linkage, EntryInit, StopName);
152 EntriesE->setVisibility(GlobalValue::HiddenVisibility);
153
154 if (Triple.isOSBinFormatELF()) {
155 // We assume that external begin/end symbols that we have created above will
156 // be defined by the linker. This is done whenever a section name with a
157 // valid C-identifier is present. We define a dummy variable here to force
158 // the linker to always provide these symbols.
159 auto *DummyEntry = new GlobalVariable(
160 M, ZeroInitializer->getType(), true, GlobalVariable::InternalLinkage,
161 ZeroInitializer, "__dummy." + SectionName);
162 DummyEntry->setSection(SectionName);
163 DummyEntry->setAlignment(Align(object::OffloadBinary::getAlignment()));
164 appendToUsed(M, Values: DummyEntry);
165 } else if (Triple.isOSBinFormatMachO()) {
166 // Mach-O needs a dummy variable in the section (like ELF) to ensure the
167 // linker provides the section boundary symbols. Mark it used so the
168 // section survives dead-stripping.
169 auto *DummyEntry = new GlobalVariable(
170 M, ZeroInitializer->getType(), true, GlobalVariable::InternalLinkage,
171 ZeroInitializer, "__dummy." + SectionName);
172 DummyEntry->setSection(SectionName);
173 DummyEntry->setAlignment(Align(object::OffloadBinary::getAlignment()));
174 appendToUsed(M, Values: DummyEntry);
175 } else {
176 // The COFF linker will merge sections containing a '$' together into a
177 // single section. The order of entries in this section will be sorted
178 // alphabetically by the characters following the '$' in the name. Set the
179 // sections here to ensure that the beginning and end symbols are sorted.
180 EntriesB->setSection((SectionName + "$OA").str());
181 EntriesE->setSection((SectionName + "$OZ").str());
182 EntriesB->setAlignment(Align(object::OffloadBinary::getAlignment()));
183 EntriesE->setAlignment(Align(object::OffloadBinary::getAlignment()));
184
185 // COFF lays out offload entries by sorted subsections: $OA is a synthetic
186 // begin sentinel, $OE contains real entries, and $OZ is a synthetic end
187 // sentinel. Keep the boundary sections non-empty so lld-link does not
188 // discard them under /opt:ref, but skip the begin sentinel for runtime
189 // users.
190 Type *Int32Ty = Type::getInt32Ty(C&: M.getContext());
191 Constant *Indices[] = {ConstantInt::get(Ty: Int32Ty, V: 0),
192 ConstantInt::get(Ty: Int32Ty, V: COFFSentinelEntryCount)};
193 Constant *BeginAfterSentinel = ConstantExpr::getInBoundsGetElementPtr(
194 Ty: EntriesB->getValueType(), C: EntriesB, IdxList: Indices);
195 return std::make_pair(x&: BeginAfterSentinel, y&: EntriesE);
196 }
197
198 return std::make_pair(x&: EntriesB, y&: EntriesE);
199}
200
201bool llvm::offloading::amdgpu::isImageCompatibleWithEnv(StringRef ImageArch,
202 uint32_t ImageFlags,
203 StringRef EnvTargetID) {
204 using namespace llvm::ELF;
205 StringRef EnvArch = EnvTargetID.split(Separator: ":").first;
206
207 // Trivial check if the base processors match.
208 if (EnvArch != ImageArch)
209 return false;
210
211 // Check if the image is requesting xnack on or off.
212 switch (ImageFlags & EF_AMDGPU_FEATURE_XNACK_V4) {
213 case EF_AMDGPU_FEATURE_XNACK_OFF_V4:
214 // The image is 'xnack-' so the environment must be 'xnack-'.
215 if (!EnvTargetID.contains(Other: "xnack-"))
216 return false;
217 break;
218 case EF_AMDGPU_FEATURE_XNACK_ON_V4:
219 // The image is 'xnack+' so the environment must be 'xnack+'.
220 if (!EnvTargetID.contains(Other: "xnack+"))
221 return false;
222 break;
223 case EF_AMDGPU_FEATURE_XNACK_UNSUPPORTED_V4:
224 case EF_AMDGPU_FEATURE_XNACK_ANY_V4:
225 default:
226 break;
227 }
228
229 // Check if the image is requesting sramecc on or off.
230 switch (ImageFlags & EF_AMDGPU_FEATURE_SRAMECC_V4) {
231 case EF_AMDGPU_FEATURE_SRAMECC_OFF_V4:
232 // The image is 'sramecc-' so the environment must be 'sramecc-'.
233 if (!EnvTargetID.contains(Other: "sramecc-"))
234 return false;
235 break;
236 case EF_AMDGPU_FEATURE_SRAMECC_ON_V4:
237 // The image is 'sramecc+' so the environment must be 'sramecc+'.
238 if (!EnvTargetID.contains(Other: "sramecc+"))
239 return false;
240 break;
241 case EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4:
242 case EF_AMDGPU_FEATURE_SRAMECC_ANY_V4:
243 break;
244 }
245
246 return true;
247}
248
249namespace {
250/// Reads the AMDGPU specific per-kernel-metadata from an image.
251class KernelInfoReader {
252public:
253 KernelInfoReader(StringMap<offloading::amdgpu::AMDGPUKernelMetaData> &KIM)
254 : KernelInfoMap(KIM) {}
255
256 /// Process ELF note to read AMDGPU metadata from respective information
257 /// fields.
258 Error processNote(const llvm::object::ELF64LE::Note &Note, size_t Align) {
259 if (Note.getName() != "AMDGPU")
260 return Error::success(); // We are not interested in other things
261
262 assert(Note.getType() == ELF::NT_AMDGPU_METADATA &&
263 "Parse AMDGPU MetaData");
264 auto Desc = Note.getDesc(Align);
265 StringRef MsgPackString =
266 StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());
267 msgpack::Document MsgPackDoc;
268 if (!MsgPackDoc.readFromBlob(Blob: MsgPackString, /*Multi=*/false))
269 return Error::success();
270
271 AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
272 if (!Verifier.verify(HSAMetadataRoot&: MsgPackDoc.getRoot()))
273 return Error::success();
274
275 auto RootMap = MsgPackDoc.getRoot().getMap(Convert: true);
276
277 if (auto Err = iterateAMDKernels(MDN&: RootMap))
278 return Err;
279
280 return Error::success();
281 }
282
283private:
284 /// Extracts the relevant information via simple string look-up in the msgpack
285 /// document elements.
286 Error
287 extractKernelData(msgpack::MapDocNode::MapTy::value_type V,
288 std::string &KernelName,
289 offloading::amdgpu::AMDGPUKernelMetaData &KernelData) {
290 if (!V.first.isString())
291 return Error::success();
292
293 const auto IsKey = [](const msgpack::DocNode &DK, StringRef SK) {
294 return DK.getString() == SK;
295 };
296
297 const auto GetSequenceOfThreeInts = [](msgpack::DocNode &DN,
298 uint32_t *Vals) {
299 assert(DN.isArray() && "MsgPack DocNode is an array node");
300 auto DNA = DN.getArray();
301 assert(DNA.size() == 3 && "ArrayNode has at most three elements");
302
303 int I = 0;
304 for (auto DNABegin = DNA.begin(), DNAEnd = DNA.end(); DNABegin != DNAEnd;
305 ++DNABegin) {
306 Vals[I++] = DNABegin->getUInt();
307 }
308 };
309
310 if (IsKey(V.first, ".name")) {
311 KernelName = V.second.toString();
312 } else if (IsKey(V.first, ".sgpr_count")) {
313 KernelData.SGPRCount = V.second.getUInt();
314 } else if (IsKey(V.first, ".sgpr_spill_count")) {
315 KernelData.SGPRSpillCount = V.second.getUInt();
316 } else if (IsKey(V.first, ".vgpr_count")) {
317 KernelData.VGPRCount = V.second.getUInt();
318 } else if (IsKey(V.first, ".vgpr_spill_count")) {
319 KernelData.VGPRSpillCount = V.second.getUInt();
320 } else if (IsKey(V.first, ".agpr_count")) {
321 KernelData.AGPRCount = V.second.getUInt();
322 } else if (IsKey(V.first, ".private_segment_fixed_size")) {
323 KernelData.PrivateSegmentSize = V.second.getUInt();
324 } else if (IsKey(V.first, ".group_segment_fixed_size")) {
325 KernelData.GroupSegmentList = V.second.getUInt();
326 } else if (IsKey(V.first, ".reqd_workgroup_size")) {
327 GetSequenceOfThreeInts(V.second, KernelData.RequestedWorkgroupSize);
328 } else if (IsKey(V.first, ".workgroup_size_hint")) {
329 GetSequenceOfThreeInts(V.second, KernelData.WorkgroupSizeHint);
330 } else if (IsKey(V.first, ".wavefront_size")) {
331 KernelData.WavefrontSize = V.second.getUInt();
332 } else if (IsKey(V.first, ".max_flat_workgroup_size")) {
333 KernelData.MaxFlatWorkgroupSize = V.second.getUInt();
334 } else if (IsKey(V.first, ".args")) {
335 auto ArgsArray = V.second.getArray();
336 for (auto ArgIt = ArgsArray.begin(), ArgEnd = ArgsArray.end();
337 ArgIt != ArgEnd; ++ArgIt) {
338 auto ArgMap = ArgIt->getMap();
339
340 auto OffsetIt = ArgMap.find(Key: ".offset");
341 if (OffsetIt == ArgMap.end())
342 return createStringError(
343 EC: inconvertibleErrorCode(),
344 S: "Missing required .offset key in kernel argument metadata map");
345
346 auto SizeIt = ArgMap.find(Key: ".size");
347 if (SizeIt == ArgMap.end())
348 return createStringError(
349 EC: inconvertibleErrorCode(),
350 S: "Missing required .size key in kernel argument metadata map");
351
352 KernelData.ArgMDs.emplace_back(Args&: OffsetIt->second.getUInt(),
353 Args&: SizeIt->second.getUInt());
354 }
355 }
356
357 return Error::success();
358 }
359
360 /// Get the "amdhsa.kernels" element from the msgpack Document
361 Expected<msgpack::ArrayDocNode> getAMDKernelsArray(msgpack::MapDocNode &MDN) {
362 auto Res = MDN.find(Key: "amdhsa.kernels");
363 if (Res == MDN.end())
364 return createStringError(EC: inconvertibleErrorCode(),
365 S: "Could not find amdhsa.kernels key");
366
367 auto Pair = *Res;
368 assert(Pair.second.isArray() &&
369 "AMDGPU kernel entries are arrays of entries");
370
371 return Pair.second.getArray();
372 }
373
374 /// Iterate all entries for one "amdhsa.kernels" entry. Each entry is a
375 /// MapDocNode that either maps a string to a single value (most of them) or
376 /// to another array of things. Currently, we only handle the case that maps
377 /// to scalar value.
378 Error generateKernelInfo(msgpack::ArrayDocNode::ArrayTy::iterator It) {
379 offloading::amdgpu::AMDGPUKernelMetaData KernelData;
380 std::string KernelName;
381 auto Entry = (*It).getMap();
382 for (auto MI = Entry.begin(), E = Entry.end(); MI != E; ++MI)
383 if (auto Err = extractKernelData(V: *MI, KernelName, KernelData))
384 return Err;
385
386 KernelInfoMap.insert(KV: {KernelName, KernelData});
387 return Error::success();
388 }
389
390 /// Go over the list of AMD kernels in the "amdhsa.kernels" entry
391 Error iterateAMDKernels(msgpack::MapDocNode &MDN) {
392 auto KernelsOrErr = getAMDKernelsArray(MDN);
393 if (auto Err = KernelsOrErr.takeError())
394 return Err;
395
396 auto KernelsArr = *KernelsOrErr;
397 for (auto It = KernelsArr.begin(), E = KernelsArr.end(); It != E; ++It) {
398 if (!It->isMap())
399 continue; // we expect <key,value> pairs
400
401 // Obtain the value for the different entries. Each array entry is a
402 // MapDocNode
403 if (auto Err = generateKernelInfo(It))
404 return Err;
405 }
406 return Error::success();
407 }
408
409 // Kernel names are the keys
410 StringMap<offloading::amdgpu::AMDGPUKernelMetaData> &KernelInfoMap;
411};
412} // namespace
413
414Error llvm::offloading::amdgpu::getAMDGPUMetaDataFromImage(
415 MemoryBufferRef MemBuffer,
416 StringMap<offloading::amdgpu::AMDGPUKernelMetaData> &KernelInfoMap,
417 uint16_t &ELFABIVersion) {
418 Error Err = Error::success(); // Used later as out-parameter
419
420 auto ELFOrError = object::ELF64LEFile::create(Object: MemBuffer.getBuffer());
421 if (auto Err = ELFOrError.takeError())
422 return Err;
423
424 const object::ELF64LEFile ELFObj = ELFOrError.get();
425 Expected<ArrayRef<object::ELF64LE::Shdr>> Sections = ELFObj.sections();
426 if (!Sections)
427 return Sections.takeError();
428 KernelInfoReader Reader(KernelInfoMap);
429
430 // Read the code object version from ELF image header
431 auto Header = ELFObj.getHeader();
432 ELFABIVersion = (uint8_t)(Header.e_ident[ELF::EI_ABIVERSION]);
433 for (const auto &S : *Sections) {
434 if (S.sh_type != ELF::SHT_NOTE)
435 continue;
436
437 for (const auto N : ELFObj.notes(Shdr: S, Err)) {
438 if (Err)
439 return Err;
440 // Fills the KernelInfoTabel entries in the reader
441 if ((Err = Reader.processNote(Note: N, Align: S.sh_addralign)))
442 return Err;
443 }
444 }
445 return Error::success();
446}
447
448Error offloading::containerizeImage(std::unique_ptr<MemoryBuffer> &Img,
449 llvm::Triple Triple,
450 object::ImageKind ImageKind,
451 object::OffloadKind OffloadKind,
452 int32_t ImageFlags,
453 MapVector<StringRef, StringRef> &MetaData) {
454 using namespace object;
455
456 // Create inner OffloadBinary containing the raw image.
457 OffloadBinary::OffloadingImage InnerImage;
458 InnerImage.TheImageKind = ImageKind;
459 InnerImage.TheOffloadKind = OffloadKind;
460 InnerImage.Flags = ImageFlags;
461
462 InnerImage.StringData["triple"] = Triple.getTriple();
463 for (const auto &[Key, Value] : MetaData)
464 InnerImage.StringData[Key] = Value;
465
466 InnerImage.Image = std::move(Img);
467
468 SmallString<0> InnerBinaryData = OffloadBinary::write(OffloadingData: InnerImage);
469
470 Img = MemoryBuffer::getMemBufferCopy(InputData: InnerBinaryData);
471 return Error::success();
472}
473
474Error offloading::intel::containerizeOpenMPSPIRVImage(
475 std::unique_ptr<MemoryBuffer> &Binary, llvm::Triple Triple,
476 StringRef CompileOpts, StringRef LinkOpts) {
477 constexpr char INTEL_ONEOMP_OFFLOAD_VERSION[] = "1.0";
478
479 assert(Triple.isSPIRV() && Triple.getVendor() == llvm::Triple::Intel &&
480 "Expected SPIR-V triple with Intel vendor");
481
482 MapVector<StringRef, StringRef> MetaData;
483 MetaData["version"] = INTEL_ONEOMP_OFFLOAD_VERSION;
484 if (!CompileOpts.empty())
485 MetaData["compile-opts"] = CompileOpts;
486 if (!LinkOpts.empty())
487 MetaData["link-opts"] = LinkOpts;
488
489 return containerizeImage(Img&: Binary, Triple, ImageKind: object::ImageKind::IMG_SPIRV,
490 OffloadKind: object::OffloadKind::OFK_OpenMP, /*ImageFlags=*/0,
491 MetaData);
492}
493
494void sycl::writeSymbolTable(ArrayRef<StringRef> Names, SmallString<0> &Out) {
495 uint32_t Count = Names.size();
496
497 // Compute the byte offset where string data begins: right after the header
498 // and the entry array.
499 uint32_t StringDataOffset =
500 sizeof(SymbolTableHeader) + Count * sizeof(SymbolTableEntry);
501
502 // Compute total size and reserve to prevent reallocation while writing
503 // entries via pointer (append() could otherwise invalidate the pointer).
504 uint32_t TotalSize = StringDataOffset;
505 for (StringRef N : Names)
506 TotalSize += N.size() + 1;
507 Out.reserve(N: TotalSize);
508 Out.resize(N: StringDataOffset);
509
510 // Write the header.
511 auto *Header = reinterpret_cast<SymbolTableHeader *>(Out.data());
512 Header->Count = Count;
513
514 // Write each entry and append the corresponding null-terminated name.
515 auto *Entries = reinterpret_cast<SymbolTableEntry *>(Header + 1);
516 uint32_t CurrentOffset = StringDataOffset;
517 for (uint32_t I = 0; I < Count; ++I) {
518 Entries[I].OffsetToSymbol = CurrentOffset;
519 Entries[I].SymbolSize = Names[I].size();
520 Out.append(RHS: Names[I]);
521 Out.push_back(Elt: '\0');
522 CurrentOffset += Names[I].size() + 1;
523 }
524}
525