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