34 ArgAlign =
DL.getABITypeAlign(Ty);
36 return std::pair(Ty, *ArgAlign);
42 "amdgpu-dump-hsa-metadata",
43 cl::desc(
"Dump AMDGPU HSA Metadata"));
45 "amdgpu-verify-hsa-metadata",
46 cl::desc(
"Verify AMDGPU HSA Metadata"));
56 errs() <<
"AMDGPU HSA Metadata:\n" << HSAMetadataString <<
'\n';
60 errs() <<
"AMDGPU HSA Metadata Parser Test: ";
64 if (!FromHSAMetadataString.
fromYAML(HSAMetadataString)) {
69 std::string ToHSAMetadataString;
71 FromHSAMetadataString.
toYAML(StrOS);
73 errs() << (HSAMetadataString == StrOS.
str() ?
"PASS" :
"FAIL") <<
'\n';
74 if (HSAMetadataString != ToHSAMetadataString) {
75 errs() <<
"Original input: " << HSAMetadataString <<
'\n'
76 <<
"Produced output: " << StrOS.
str() <<
'\n';
80std::optional<StringRef>
83 .Case(
"read_only",
StringRef(
"read_only"))
84 .Case(
"write_only",
StringRef(
"write_only"))
85 .Case(
"read_write",
StringRef(
"read_write"))
116 .
Case(
"image1d_t",
"image")
117 .
Case(
"image1d_array_t",
"image")
118 .
Case(
"image1d_buffer_t",
"image")
119 .
Case(
"image2d_t",
"image")
120 .
Case(
"image2d_array_t",
"image")
121 .
Case(
"image2d_array_depth_t",
"image")
122 .
Case(
"image2d_array_msaa_t",
"image")
123 .
Case(
"image2d_array_msaa_depth_t",
"image")
124 .
Case(
"image2d_depth_t",
"image")
125 .
Case(
"image2d_msaa_t",
"image")
126 .
Case(
"image2d_msaa_depth_t",
"image")
127 .
Case(
"image3d_t",
"image")
128 .
Case(
"sampler_t",
"sampler")
129 .
Case(
"queue_t",
"queue")
132 ?
"dynamic_shared_pointer"
165 auto VecTy = cast<FixedVectorType>(Ty);
166 auto ElTy = VecTy->getElementType();
167 auto NumElements = VecTy->getNumElements();
178 if (Node->getNumOperands() != 3)
181 for (
auto &
Op : Node->operands())
182 Dims.push_back(Dims.getDocument()->getNode(
183 uint64_t(mdconst::extract<ConstantInt>(
Op)->getZExtValue())));
189 Version.push_back(Version.getDocument()->getNode(
VersionMajorV4));
190 Version.push_back(Version.getDocument()->getNode(
VersionMinorV4));
201 auto Node =
Mod.getNamedMetadata(
"llvm.printf.fmts");
206 for (
auto *
Op : Node->operands())
208 Printf.push_back(Printf.getDocument()->getNode(
209 cast<MDString>(
Op->getOperand(0))->getString(),
true));
216 auto Node = Func.getParent()->getNamedMetadata(
"opencl.ocl.version");
217 if (!Node || !Node->getNumOperands())
219 auto Op0 = Node->getOperand(0);
220 if (Op0->getNumOperands() <= 1)
226 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
228 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
235 if (
auto Node = Func.getMetadata(
"reqd_work_group_size"))
237 if (
auto Node = Func.getMetadata(
"work_group_size_hint"))
239 if (
auto Node = Func.getMetadata(
"vec_type_hint")) {
242 cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
243 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
246 if (Func.hasFnAttribute(
"runtime-handle")) {
248 Func.getFnAttribute(
"runtime-handle").getValueAsString().str(),
251 if (Func.hasFnAttribute(
"device-init"))
253 else if (Func.hasFnAttribute(
"device-fini"))
262 for (
auto &Arg : Func.args())
267 Kern[
".args"] = Args;
278 Node = Func->getMetadata(
"kernel_arg_name");
279 if (Node && ArgNo < Node->getNumOperands())
280 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
285 Node = Func->getMetadata(
"kernel_arg_type");
286 if (Node && ArgNo < Node->getNumOperands())
287 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
290 Node = Func->getMetadata(
"kernel_arg_base_type");
291 if (Node && ArgNo < Node->getNumOperands())
292 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
298 ActAccQual =
"read_only";
300 ActAccQual =
"write_only";
304 Node = Func->getMetadata(
"kernel_arg_access_qual");
305 if (Node && ArgNo < Node->getNumOperands())
306 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
309 Node = Func->getMetadata(
"kernel_arg_type_qual");
310 if (Node && ArgNo < Node->getNumOperands())
311 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
313 const DataLayout &
DL = Func->getParent()->getDataLayout();
319 if (
auto PtrTy = dyn_cast<PointerType>(Ty)) {
331 PointeeAlign,
Name, TypeName, BaseTypeName, ActAccQual,
340 auto Arg = Args.getDocument()->getMapNode();
343 Arg[
".name"] = Arg.getDocument()->getNode(
Name,
true);
344 if (!TypeName.empty())
345 Arg[
".type_name"] = Arg.getDocument()->getNode(TypeName,
true);
346 auto Size =
DL.getTypeAllocSize(Ty);
347 Arg[
".size"] = Arg.getDocument()->getNode(
Size);
349 Arg[
".offset"] = Arg.getDocument()->getNode(
Offset);
351 Arg[
".value_kind"] = Arg.getDocument()->getNode(
ValueKind,
true);
353 Arg[
".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
355 if (
auto PtrTy = dyn_cast<PointerType>(Ty))
359 Arg[
".address_space"] = Arg.getDocument()->getNode(*Qualifier,
363 Arg[
".access"] = Arg.getDocument()->getNode(*AQ,
true);
366 Arg[
".actual_access"] = Arg.getDocument()->getNode(*AAQ,
true);
369 TypeQual.
split(SplitTypeQuals,
" ", -1,
false);
372 Arg[
".is_const"] = Arg.getDocument()->getNode(
true);
373 else if (Key ==
"restrict")
374 Arg[
".is_restrict"] = Arg.getDocument()->getNode(
true);
375 else if (Key ==
"volatile")
376 Arg[
".is_volatile"] = Arg.getDocument()->getNode(
true);
377 else if (Key ==
"pipe")
378 Arg[
".is_pipe"] = Arg.getDocument()->getNode(
true);
389 unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
390 if (!HiddenArgNumBytes)
393 const Module *M = Func.getParent();
394 auto &
DL = M->getDataLayout();
399 if (HiddenArgNumBytes >= 8)
402 if (HiddenArgNumBytes >= 16)
405 if (HiddenArgNumBytes >= 24)
412 if (HiddenArgNumBytes >= 32) {
416 if (M->getNamedMetadata(
"llvm.printf.fmts"))
419 else if (!Func.hasFnAttribute(
"amdgpu-no-hostcall-ptr"))
428 if (HiddenArgNumBytes >= 40) {
429 if (!Func.hasFnAttribute(
"amdgpu-no-default-queue")) {
437 if (HiddenArgNumBytes >= 48) {
438 if (!Func.hasFnAttribute(
"amdgpu-no-completion-action")) {
447 if (HiddenArgNumBytes >= 56) {
448 if (!Func.hasFnAttribute(
"amdgpu-no-multigrid-sync-arg")) {
460 unsigned CodeObjectVersion)
const {
467 Align MaxKernArgAlign;
468 Kern[
".kernarg_segment_size"] = Kern.getDocument()->getNode(
470 Kern[
".group_segment_fixed_size"] =
471 Kern.getDocument()->getNode(ProgramInfo.
LDSSize);
472 Kern[
".private_segment_fixed_size"] =
473 Kern.getDocument()->getNode(ProgramInfo.
ScratchSize);
475 Kern[
".uses_dynamic_stack"] =
479 Kern[
".workgroup_processor_mode"] =
480 Kern.getDocument()->getNode(ProgramInfo.
WgpMode);
483 Kern[
".kernarg_segment_align"] =
484 Kern.getDocument()->getNode(std::max(
Align(4), MaxKernArgAlign).
value());
485 Kern[
".wavefront_size"] =
487 Kern[
".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.
NumSGPR);
488 Kern[
".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.
NumVGPR);
492 Kern[
".agpr_count"] = Kern.getDocument()->getNode(ProgramInfo.
NumAccVGPR);
495 Kern[
".max_flat_workgroup_size"] =
496 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
497 unsigned NumWGX = MFI.getMaxNumWorkGroupsX();
498 unsigned NumWGY = MFI.getMaxNumWorkGroupsY();
499 unsigned NumWGZ = MFI.getMaxNumWorkGroupsZ();
500 if (NumWGX != 0 && NumWGY != 0 && NumWGZ != 0) {
501 Kern[
".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
502 Kern[
".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
503 Kern[
".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
505 Kern[
".sgpr_spill_count"] =
506 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
507 Kern[
".vgpr_spill_count"] =
508 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
526 std::string HSAMetadataString;
543 auto CodeObjectVersion =
553 (
Twine(Func.getName()) +
Twine(
".kd")).str(),
true);
559 Kernels.push_back(Kern);
568 Version.push_back(Version.getDocument()->getNode(
VersionMajorV5));
569 Version.push_back(Version.getDocument()->getNode(
VersionMinorV5));
579 if (ST.getImplicitArgNumBytes(Func) == 0)
582 const Module *M = Func.getParent();
583 auto &
DL = M->getDataLayout();
618 if (M->getNamedMetadata(
"llvm.printf.fmts")) {
625 if (!Func.hasFnAttribute(
"amdgpu-no-hostcall-ptr")) {
632 if (!Func.hasFnAttribute(
"amdgpu-no-multigrid-sync-arg")) {
639 if (!Func.hasFnAttribute(
"amdgpu-no-heap-ptr"))
644 if (!Func.hasFnAttribute(
"amdgpu-no-default-queue")) {
651 if (!Func.hasFnAttribute(
"amdgpu-no-completion-action")) {
670 if (!ST.hasApertureRegs()) {
685 if (Func.getFnAttribute(
"uniform-work-group-size").getValueAsBool())
695 Version.push_back(Version.getDocument()->getNode(
VersionMajorV6));
696 Version.push_back(Version.getDocument()->getNode(
VersionMinorV6));
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
Given that RA is a live value
AMD GCN specific subclass of TargetSubtarget.
Module.h This file contains the declarations for the Module class.
Defines struct to track resource usage and hardware flags for kernels and entry functions.
bool isDynamicLDSUsed() const
unsigned getKernArgSegmentSize(const Function &F, Align &MaxAlign) const
unsigned getWavefrontSize() const
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict)
Emit HSA Metadata.
std::string toString() const
This class represents an incoming formal argument to a Function.
Type * getParamByRefType() const
If this is a byref argument, return its type.
bool hasNoAliasAttr() const
Return true if this argument has the noalias attribute.
bool hasByRefAttr() const
Return true if this argument has the byref attribute.
bool onlyReadsMemory() const
Return true if this argument has the readonly or readnone attribute.
bool hasAttribute(Attribute::AttrKind Kind) const
Check if an argument has a given attribute.
const Function * getParent() const
unsigned getArgNo() const
Return the index of this formal argument in its containing function.
MaybeAlign getParamAlign() const
If this is a byval or inalloca argument, return its alignment.
This class represents an Operation in the Expression.
uint64_t getNumOperands() const
A parsed version of the target data layout string in and methods for querying it.
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
Function & getFunction()
Return the LLVM function that this machine code represents.
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
A Module instance is used to store all the information related to an LLVM module.
static PointerType * get(Type *ElementType, unsigned AddressSpace)
This constructs a pointer to an object of the specified type in a numbered address space.
This class keeps track of the SPI_SP_INPUT_ADDR config register, which tells the hardware which inter...
GCNUserSGPRUsageInfo & getUserSGPRInfo()
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
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.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
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.
unsigned getIntegerBitWidth() const
bool isPointerTy() const
True if this is an instance of PointerType.
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ IntegerTyID
Arbitrary bit width integers.
@ FixedVectorTyID
Fixed width SIMD vector type.
@ DoubleTyID
64-bit floating point type
static IntegerType * getInt16Ty(LLVMContext &C)
static IntegerType * getInt32Ty(LLVMContext &C)
static IntegerType * getInt64Ty(LLVMContext &C)
TypeID getTypeID() const
Return the type id for the type.
Type * getType() const
All values are typed, get the type of this value.
StringRef getName() const
Return a constant reference to the value's name.
A DocNode that is an array.
ArrayDocNode & getArray(bool Convert=false)
Get an ArrayDocNode for an array node.
Document * getDocument() const
Simple in-memory representation of a document of msgpack objects with ability to find and create arra...
DocNode getNode()
Create a nil node associated with this Document.
ArrayDocNode getArrayNode()
Create an empty Array node associated with this Document.
void toYAML(raw_ostream &OS)
Convert MsgPack Document to YAML text.
bool fromYAML(StringRef S)
Read YAML text into the MsgPack document. Returns false on failure.
A raw_ostream that writes to an std::string.
std::string & str()
Returns the string's reference.
unsigned LanguageVersion(SourceLanguage L)
@ REGION_ADDRESS
Address space for region memory. (GDS)
@ LOCAL_ADDRESS
Address space for local memory.
@ CONSTANT_ADDRESS
Address space for constant memory (VTX2).
@ FLAT_ADDRESS
Address space for flat memory.
@ GLOBAL_ADDRESS
Address space for global memory (RAT0, VTX0).
@ PRIVATE_ADDRESS
Address space for private memory.
constexpr uint32_t VersionMajorV5
HSA metadata major version for code object V5.
constexpr uint32_t VersionMinorV4
HSA metadata minor version for code object V4.
constexpr uint32_t VersionMinorV5
HSA metadata minor version for code object V5.
constexpr uint32_t VersionMinorV6
HSA metadata minor version for code object V6.
constexpr uint32_t VersionMajorV6
HSA metadata major version for code object V6.
constexpr uint32_t VersionMajorV4
HSA metadata major version for code object V4.
unsigned getAMDHSACodeObjectVersion(const Module &M)
@ AMDGPU_KERNEL
Used for AMDGPU code object kernels.
@ SPIR_KERNEL
Used for SPIR kernel functions.
This is an optimization pass for GlobalISel generic memory operations.
static cl::opt< bool > VerifyHSAMetadata("amdgpu-verify-hsa-metadata", cl::desc("Verify AMDGPU HSA Metadata"))
static cl::opt< bool > DumpHSAMetadata("amdgpu-dump-hsa-metadata", cl::desc("Dump AMDGPU HSA Metadata"))
raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
StringRef getTypeName()
We provide a function which tries to compute the (demangled) name of a type statically.
@ Mod
The access may modify the value stored in memory.
uint64_t alignTo(uint64_t Size, Align A)
Returns a multiple of A needed to store Size bytes.
constexpr unsigned BitWidth
This struct is a compact representation of a valid (non-zero power of two) alignment.
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Align valueOrOne() const
For convenience, returns a valid alignment or 1 if undefined.
Track resource usage for kernels / entry functions.