Utility.cpp Source File (original) (raw)

23using namespace llvm;

39std::pair<Constant *, GlobalVariable *>

53 Triple.isNVPTX() ? "$offloading$entry_name" : ".offloading.entry_name";

56 auto *Str =

62 Str->setAlignment(Align(1));

65 NamedMDNode *MD = M.getOrInsertNamedMetadata("llvm.offloading.symbols");

72 ConstantInt::get(Int16Ty, 1),

73 ConstantInt::get(Int16Ty, Kind),

74 ConstantInt::get(Int32Ty, Flags),

77 ConstantInt::get(Int64Ty, Size),

78 ConstantInt::get(Int64Ty, Data),

82 return {EntryInitializer, Str};

83}

93 M, Kind, Addr, Name, Size, Flags, Data, AuxAddr);

96 Triple.isNVPTX() ? "$offloading$entry$" : ".offloading.entry.";

101 M.getDataLayout().getDefaultGlobalsAddressSpace());

102

103

105 Entry->setSection((SectionName + "$OE").str());

106 else

109 return Entry;

110}

111

112std::pair<GlobalVariable *, GlobalVariable *>

115

116 auto *ZeroInitilaizer =

122

123 auto *EntriesB =

124 new GlobalVariable(M, EntryType, true, Linkage, EntryInit,

127 auto *EntriesE =

128 new GlobalVariable(M, EntryType, true, Linkage, EntryInit,

131

133

134

135

136

139 ZeroInitilaizer, "__dummy." + SectionName);

143 } else {

144

145

146

147

148 EntriesB->setSection((SectionName + "$OA").str());

149 EntriesE->setSection((SectionName + "$OZ").str());

150 }

151

152 return std::make_pair(EntriesB, EntriesE);

153}

154

160

161

162 if (EnvArch != ImageArch)

163 return false;

164

165

168

169 if (!EnvTargetID.contains("xnack-"))

170 return false;

171 break;

173

174 if (!EnvTargetID.contains("xnack+"))

175 return false;

176 break;

179 default:

180 break;

181 }

182

183

186

187 if (!EnvTargetID.contains("sramecc-"))

188 return false;

189 break;

191

192 if (!EnvTargetID.contains("sramecc+"))

193 return false;

194 break;

197 break;

198 }

199

200 return true;

201}

202

203namespace {

204

205class KernelInfoReader {

206public:

208 : KernelInfoMap(KIM) {}

209

210

211

213 if (Note.getName() != "AMDGPU")

214 return Error::success();

215

217 "Parse AMDGPU MetaData");

220 StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());

222 if (!MsgPackDoc.readFromBlob(MsgPackString, false))

224

226 if (!Verifier.verify(MsgPackDoc.getRoot()))

228

229 auto RootMap = MsgPackDoc.getRoot().getMap(true);

230

231 if (auto Err = iterateAMDKernels(RootMap))

232 return Err;

233

235 }

236

237private:

238

239

241 extractKernelData(msgpack::MapDocNode::MapTy::value_type V,

242 std::string &KernelName,

244 if (!V.first.isString())

246

249 };

250

251 const auto GetSequenceOfThreeInts = [](msgpack::DocNode &DN,

252 uint32_t *Vals) {

253 assert(DN.isArray() && "MsgPack DocNode is an array node");

255 assert(DNA.size() == 3 && "ArrayNode has at most three elements");

256

257 int I = 0;

258 for (auto DNABegin = DNA.begin(), DNAEnd = DNA.end(); DNABegin != DNAEnd;

259 ++DNABegin) {

260 Vals[I++] = DNABegin->getUInt();

261 }

262 };

263

264 if (IsKey(V.first, ".name")) {

265 KernelName = V.second.toString();

266 } else if (IsKey(V.first, ".sgpr_count")) {

267 KernelData.SGPRCount = V.second.getUInt();

268 } else if (IsKey(V.first, ".sgpr_spill_count")) {

270 } else if (IsKey(V.first, ".vgpr_count")) {

271 KernelData.VGPRCount = V.second.getUInt();

272 } else if (IsKey(V.first, ".vgpr_spill_count")) {

274 } else if (IsKey(V.first, ".agpr_count")) {

275 KernelData.AGPRCount = V.second.getUInt();

276 } else if (IsKey(V.first, ".private_segment_fixed_size")) {

278 } else if (IsKey(V.first, ".group_segment_fixed_size")) {

280 } else if (IsKey(V.first, ".reqd_workgroup_size")) {

282 } else if (IsKey(V.first, ".workgroup_size_hint")) {

284 } else if (IsKey(V.first, ".wavefront_size")) {

286 } else if (IsKey(V.first, ".max_flat_workgroup_size")) {

288 }

289

291 }

292

293

294 Expectedmsgpack::ArrayDocNode getAMDKernelsArray(msgpack::MapDocNode &MDN) {

295 auto Res = MDN.find("amdhsa.kernels");

296 if (Res == MDN.end())

298 "Could not find amdhsa.kernels key");

299

300 auto Pair = *Res;

301 assert(Pair.second.isArray() &&

302 "AMDGPU kernel entries are arrays of entries");

303

304 return Pair.second.getArray();

305 }

306

307

308

309

310

311 Error generateKernelInfo(msgpack::ArrayDocNode::ArrayTy::iterator It) {

312 offloading::amdgpu::AMDGPUKernelMetaData KernelData;

313 std::string KernelName;

314 auto Entry = (*It).getMap();

316 if (auto Err = extractKernelData(*MI, KernelName, KernelData))

317 return Err;

318

319 KernelInfoMap.insert({KernelName, KernelData});

321 }

322

323

324 Error iterateAMDKernels(msgpack::MapDocNode &MDN) {

325 auto KernelsOrErr = getAMDKernelsArray(MDN);

326 if (auto Err = KernelsOrErr.takeError())

327 return Err;

328

329 auto KernelsArr = *KernelsOrErr;

330 for (auto It = KernelsArr.begin(), E = KernelsArr.end(); It != E; ++It) {

331 if (!It->isMap())

332 continue;

333

334

335

336 if (auto Err = generateKernelInfo(It))

337 return Err;

338 }

340 }

341

342

343 StringMapoffloading::amdgpu::AMDGPUKernelMetaData &KernelInfoMap;

344};

345}

346

352

354 if (auto Err = ELFOrError.takeError())

355 return Err;

356

359 if (!Sections)

361 KernelInfoReader Reader(KernelInfoMap);

362

363

364 auto Header = ELFObj.getHeader();

366 for (const auto &S : *Sections) {

368 continue;

369

370 for (const auto N : ELFObj.notes(S, Err)) {

371 if (Err)

372 return Err;

373

374 if ((Err = Reader.processNote(N, S.sh_addralign)))

375 return Err;

376 }

377 }

379}

381 std::unique_ptr &Img) {

382 constexpr char INTEL_ONEOMP_OFFLOAD_VERSION[] = "1.0";

383 constexpr int NT_INTEL_ONEOMP_OFFLOAD_VERSION = 1;

384 constexpr int NT_INTEL_ONEOMP_OFFLOAD_IMAGE_COUNT = 2;

385 constexpr int NT_INTEL_ONEOMP_OFFLOAD_IMAGE_AUX = 3;

386

387

388 std::vectorELFYAML::NoteEntry Notes;

389 std::string Version = toHex(INTEL_ONEOMP_OFFLOAD_VERSION);

392 NT_INTEL_ONEOMP_OFFLOAD_VERSION});

393

394

395

396

397 std::string AuxInfo;

398

399

402

403 unsigned ImageFmt = 1;

404

406 CompileOpts + Twine('\0') + LinkOpts)

407 .str());

410 NT_INTEL_ONEOMP_OFFLOAD_IMAGE_AUX});

411

412 std::string ImgCount = toHex(Twine(1).str());

415 NT_INTEL_ONEOMP_OFFLOAD_IMAGE_COUNT});

416

417 std::string YamlFile;

419

420

421

422

428

429

432 Section.AddressAlign = 0;

433 Section.Name = ".note.inteloneompoffload";

434 Section.Notes.emplace(std::move(Notes));

435

437 Object.Header = Header;

438 Object.Chunks.push_back(

439 std::make_uniqueELFYAML::NoteSection(std::move(Section)));

440

441

445 std::string Name = "__openmp_offload_spirv_0";

446 ImageSection.Name = Name;

449 Object.Chunks.push_back(

450 std::make_uniqueELFYAML::RawContentSection(std::move(ImageSection)));

453 Object, YamlFileStream,

455 if (Err)

456 return Err;

457

460}

assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")

static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")

This file contains the declarations for the subclasses of Constant, which represent the different fla...

This file declares classes for handling the YAML representation of ELF.

This file declares a class that exposes a simple in-memory representation of a document of MsgPack ob...

verify safepoint Safepoint IR Verifier

static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)

This static method is the primary way to construct an ArrayType.

static LLVM_ABI ConstantAggregateZero * get(Type *Ty)

static LLVM_ABI Constant * getString(LLVMContext &Context, StringRef Initializer, bool AddNull=true)

This method constructs a CDS and initializes it with a text string.

static LLVM_ABI Constant * getPointerBitCastOrAddrSpaceCast(Constant *C, Type *Ty)

Create a BitCast or AddrSpaceCast for a pointer type depending on the address space.

static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)

This is an important base class in LLVM.

static LLVM_ABI Constant * getNullValue(Type *Ty)

Constructor to create a '0' constant of arbitrary type.

Lightweight error class with error context and mandatory checking.

static ErrorSuccess success()

Create a success value.

Tagged union holding either a T or a Error.

Error takeError()

Take ownership of the stored error.

@ HiddenVisibility

The GV is hidden.

@ InternalLinkage

Rename collisions when linking (static functions).

@ WeakODRLinkage

Same, but only replaced by something equivalent.

@ ExternalLinkage

Externally visible function.

@ WeakAnyLinkage

Keep one copy of named function when linking (weak)

This is an important class for using LLVM in a threaded context.

static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)

StringRef getBuffer() const

static std::unique_ptr< MemoryBuffer > getMemBufferCopy(StringRef InputData, const Twine &BufferName="")

Open the specified memory range as a MemoryBuffer, copying the contents and taking ownership of it.

A Module instance is used to store all the information related to an LLVM module.

LLVM_ABI void addOperand(MDNode *M)

static PointerType * getUnqual(Type *ElementType)

This constructs a pointer to an object of the specified type in the default address space (address sp...

StringMap - This is an unconventional map that is specialized for handling keys that are "strings",...

StringRef - Represent a constant reference to a string, i.e.

std::pair< StringRef, StringRef > split(char Separator) const

Split into two substrings around the first occurrence of a separator character.

bool contains(StringRef Other) const

Return true if the given string is a substring of *this, and false otherwise.

Class to represent struct types.

static LLVM_ABI StructType * getTypeByName(LLVMContext &C, StringRef Name)

Return the type with the specified name, or null if there is none by that name.

static LLVM_ABI StructType * create(LLVMContext &Context, StringRef Name)

This creates an identified struct.

Triple - Helper class for working with autoconf configuration names.

bool isOSBinFormatCOFF() const

Tests whether the OS uses the COFF binary format.

bool isNVPTX() const

Tests whether the target is NVPTX (32- or 64-bit).

bool isOSBinFormatELF() const

Tests whether the OS uses the ELF binary format.

Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...

The instances of the Type class are immutable: once they are created, they are never changed.

static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)

static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)

static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)

Type * getType() const

All values are typed, get the type of this value.

A node in a MsgPack Document.

MapDocNode & getMap(bool Convert=false)

Get a MapDocNode for a map node.

ArrayDocNode & getArray(bool Convert=false)

Get an ArrayDocNode for an array node.

StringRef getString() const

Simple in-memory representation of a document of msgpack objects with ability to find and create arra...

DocNode & getRoot()

Get ref to the document's root element.

LLVM_ABI bool readFromBlob(StringRef Blob, bool Multi, function_ref< int(DocNode *DestNode, DocNode SrcNode, DocNode MapKey)> Merger=[](DocNode *DestNode, DocNode SrcNode, DocNode MapKey) { return -1;})

Read a document from a binary msgpack blob, merging into anything already in the Document.

MapTy::iterator find(DocNode Key)

const Elf_Ehdr & getHeader() const

static Expected< ELFFile > create(StringRef Object)

iterator_range< Elf_Note_Iterator > notes(const Elf_Phdr &Phdr, Error &Err) const

Get an iterator range over notes of a program header.

Expected< Elf_Shdr_Range > sections() const

static uint64_t getAlignment()

A raw_ostream that writes to an std::string.

Specialized YAMLIO scalar type for representing a binary blob.

@ C

The default llvm calling convention, compatible with C.

@ EF_AMDGPU_FEATURE_XNACK_ANY_V4

@ EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4

@ EF_AMDGPU_FEATURE_SRAMECC_OFF_V4

@ EF_AMDGPU_FEATURE_XNACK_UNSUPPORTED_V4

@ EF_AMDGPU_FEATURE_XNACK_OFF_V4

@ EF_AMDGPU_FEATURE_XNACK_V4

@ EF_AMDGPU_FEATURE_SRAMECC_V4

@ EF_AMDGPU_FEATURE_XNACK_ON_V4

@ EF_AMDGPU_FEATURE_SRAMECC_ANY_V4

@ EF_AMDGPU_FEATURE_SRAMECC_ON_V4

OffloadKind

The producer of the associated offloading image.

ELFFile< ELF64LE > ELF64LEFile

LLVM_ABI Error getAMDGPUMetaDataFromImage(MemoryBufferRef MemBuffer, StringMap< AMDGPUKernelMetaData > &KernelInfoMap, uint16_t &ELFABIVersion)

Reads AMDGPU specific metadata from the ELF file and propagates the KernelInfoMap.

Definition Utility.cpp:347

LLVM_ABI bool isImageCompatibleWithEnv(StringRef ImageArch, uint32_t ImageFlags, StringRef EnvTargetID)

Check if an image is compatible with current system's environment.

Definition Utility.cpp:155

LLVM_ABI Error containerizeOpenMPSPIRVImage(std::unique_ptr< MemoryBuffer > &Binary)

Containerizes an offloading binary into the ELF binary format expected by the Intel runtime offload p...

Definition Utility.cpp:380

LLVM_ABI GlobalVariable * emitOffloadingEntry(Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name, uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr=nullptr, StringRef SectionName="llvm_offload_entries")

Create an offloading section struct used to register this global at runtime.

Definition Utility.cpp:86

LLVM_ABI std::pair< Constant *, GlobalVariable * > getOffloadingEntryInitializer(Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name, uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr)

Create a constant struct initializer used to register this global at runtime.

Definition Utility.cpp:40

LLVM_ABI StructType * getEntryTy(Module &M)

Returns the type of the offloading entry we use to store kernels and globals that will be registered ...

Definition Utility.cpp:26

LLVM_ABI std::pair< GlobalVariable *, GlobalVariable * > getOffloadEntryArray(Module &M, StringRef SectionName="llvm_offload_entries")

Creates a pair of globals used to iterate the array of offloading entries by accessing the section va...

Definition Utility.cpp:113

LLVM_ABI bool yaml2elf(ELFYAML::Object &Doc, raw_ostream &Out, ErrorHandler EH, uint64_t MaxSize)

This is an optimization pass for GlobalISel generic memory operations.

ArrayRef< CharT > arrayRefFromStringRef(StringRef Input)

Construct a string ref from an array ref of unsigned chars.

FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty

LLVM_ABI std::error_code inconvertibleErrorCode()

The value returned by this function can be returned from convertToErrorCode for Error values where no...

Error createStringError(std::error_code EC, char const *Fmt, const Ts &... Vals)

Create formatted StringError object.

FunctionAddr VTableAddr uintptr_t uintptr_t Version

FunctionAddr VTableAddr uintptr_t uintptr_t Data

LLVM_ABI void appendToCompilerUsed(Module &M, ArrayRef< GlobalValue * > Values)

Adds global values to the llvm.compiler.used list.

void toHex(ArrayRef< uint8_t > Input, bool LowerCase, SmallVectorImpl< char > &Output)

Convert buffer Input to its hexadecimal representation. The returned string is double the size of Inp...

This struct is a compact representation of a valid (non-zero power of two) alignment.

llvm::yaml::Hex64 AddressAlign

std::optional< yaml::BinaryRef > Content

Elf_Note_Impl< ELFType< E, Is64 > > Note

This is the record of an object that just be registered with the offloading runtime.

Common declarations for yaml2obj.