28 if (
Arg.hasByRefAttr()) {
29 Ty =
Arg.getParamByRefType();
30 ArgAlign =
Arg.getParamAlign();
34 ArgAlign =
DL.getABITypeAlign(Ty);
36 return std::make_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"));
54 void MetadataStreamerV2::dump(
StringRef HSAMetadataString)
const {
55 errs() <<
"AMDGPU HSA Metadata:\n" << HSAMetadataString <<
'\n';
58 void MetadataStreamerV2::verify(
StringRef HSAMetadataString)
const {
59 errs() <<
"AMDGPU HSA Metadata Parser Test: ";
61 HSAMD::Metadata FromHSAMetadataString;
62 if (
fromString(std::string(HSAMetadataString), FromHSAMetadataString)) {
67 std::string ToHSAMetadataString;
68 if (
toString(FromHSAMetadataString, ToHSAMetadataString)) {
73 errs() << (HSAMetadataString == ToHSAMetadataString ?
"PASS" :
"FAIL")
75 if (HSAMetadataString != ToHSAMetadataString) {
76 errs() <<
"Original input: " << HSAMetadataString <<
'\n' 77 <<
"Produced output: " << ToHSAMetadataString <<
'\n';
94 MetadataStreamerV2::getAddressSpaceQualifier(
134 .
Default(isa<PointerType>(Ty) ?
142 std::string MetadataStreamerV2::getTypeName(
Type *Ty,
bool Signed)
const {
146 return (
Twine(
'u') + getTypeName(Ty,
true)).str();
169 auto VecTy = cast<FixedVectorType>(Ty);
170 auto ElTy = VecTy->getElementType();
171 auto NumElements = VecTy->getNumElements();
179 std::vector<uint32_t>
180 MetadataStreamerV2::getWorkGroupDimensions(
MDNode *Node)
const {
181 std::vector<uint32_t> Dims;
182 if (Node->getNumOperands() != 3)
185 for (
auto &
Op : Node->operands())
186 Dims.push_back(mdconst::extract<ConstantInt>(
Op)->getZExtValue());
190 Kernel::CodeProps::Metadata
195 HSAMD::Kernel::CodeProps::Metadata HSACodeProps;
201 Align MaxKernArgAlign;
204 HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.
LDSSize;
205 HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.
ScratchSize;
206 HSACodeProps.mKernargSegmentAlign =
209 HSACodeProps.mNumSGPRs = ProgramInfo.
NumSGPR;
210 HSACodeProps.mNumVGPRs = ProgramInfo.
NumVGPR;
211 HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize();
214 HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
215 HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
220 Kernel::DebugProps::Metadata
223 return HSAMD::Kernel::DebugProps::Metadata();
226 void MetadataStreamerV2::emitVersion() {
233 void MetadataStreamerV2::emitPrintf(
const Module &Mod) {
236 auto Node =
Mod.getNamedMetadata(
"llvm.printf.fmts");
240 for (
auto Op : Node->operands())
241 if (
Op->getNumOperands())
243 std::string(cast<MDString>(
Op->getOperand(0))->getString()));
246 void MetadataStreamerV2::emitKernelLanguage(
const Function &Func) {
250 auto Node =
Func.getParent()->getNamedMetadata(
"opencl.ocl.version");
251 if (!Node || !Node->getNumOperands())
253 auto Op0 = Node->getOperand(0);
254 if (Op0->getNumOperands() <= 1)
257 Kernel.mLanguage =
"OpenCL C";
258 Kernel.mLanguageVersion.push_back(
259 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
260 Kernel.mLanguageVersion.push_back(
261 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
264 void MetadataStreamerV2::emitKernelAttrs(
const Function &Func) {
267 if (
auto Node =
Func.getMetadata(
"reqd_work_group_size"))
268 Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
269 if (
auto Node =
Func.getMetadata(
"work_group_size_hint"))
270 Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
271 if (
auto Node =
Func.getMetadata(
"vec_type_hint")) {
272 Attrs.mVecTypeHint = getTypeName(
273 cast<ValueAsMetadata>(Node->getOperand(0))->
getType(),
274 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
276 if (
Func.hasFnAttribute(
"runtime-handle")) {
277 Attrs.mRuntimeHandle =
278 Func.getFnAttribute(
"runtime-handle").getValueAsString().str();
282 void MetadataStreamerV2::emitKernelArgs(
const Function &Func) {
286 emitHiddenKernelArgs(Func);
289 void MetadataStreamerV2::emitKernelArg(
const Argument &
Arg) {
291 auto ArgNo =
Arg.getArgNo();
295 Node =
Func->getMetadata(
"kernel_arg_name");
296 if (Node && ArgNo < Node->getNumOperands())
297 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
298 else if (
Arg.hasName())
302 Node =
Func->getMetadata(
"kernel_arg_type");
303 if (Node && ArgNo < Node->getNumOperands())
304 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
307 Node =
Func->getMetadata(
"kernel_arg_base_type");
308 if (Node && ArgNo < Node->getNumOperands())
309 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
312 if (
Arg.getType()->isPointerTy() &&
Arg.onlyReadsMemory() &&
313 Arg.hasNoAliasAttr()) {
316 Node =
Func->getMetadata(
"kernel_arg_access_qual");
317 if (Node && ArgNo < Node->getNumOperands())
318 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
322 Node =
Func->getMetadata(
"kernel_arg_type_qual");
323 if (Node && ArgNo < Node->getNumOperands())
324 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
329 if (
auto PtrTy = dyn_cast<PointerType>(
Arg.getType())) {
333 PtrTy->getElementType());
341 emitKernelArg(
DL, ArgTy, ArgAlign,
352 HSAMetadata.
mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
353 auto &
Arg = HSAMetadata.
mKernels.back().mArgs.back();
357 Arg.mSize =
DL.getTypeAllocSize(Ty);
362 if (
auto PtrTy = dyn_cast<PointerType>(Ty))
363 Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
370 TypeQual.
split(SplitTypeQuals,
" ", -1,
false);
374 .
Case(
"restrict", &
Arg.mIsRestrict)
375 .
Case(
"volatile", &
Arg.mIsVolatile)
383 void MetadataStreamerV2::emitHiddenKernelArgs(
const Function &Func) {
384 int HiddenArgNumBytes =
387 if (!HiddenArgNumBytes)
390 auto &
DL =
Func.getParent()->getDataLayout();
393 if (HiddenArgNumBytes >= 8)
395 if (HiddenArgNumBytes >= 16)
397 if (HiddenArgNumBytes >= 24)
405 if (HiddenArgNumBytes >= 32) {
406 if (
Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"))
408 else if (
Func.getParent()->getFunction(
"__ockl_hostcall_internal")) {
411 assert(!
Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"));
419 if (HiddenArgNumBytes >= 48) {
420 if (
Func.hasFnAttribute(
"calls-enqueue-kernel")) {
430 if (HiddenArgNumBytes >= 56)
444 std::string HSAMetadataString;
445 if (
toString(HSAMetadata, HSAMetadataString))
449 dump(HSAMetadataString);
451 verify(HSAMetadataString);
460 auto CodeProps = getHSACodeProps(MF, ProgramInfo);
461 auto DebugProps = getHSADebugProps(MF, ProgramInfo);
466 Kernel.mName = std::string(Func.getName());
468 emitKernelLanguage(Func);
469 emitKernelAttrs(Func);
470 emitKernelArgs(Func);
479 void MetadataStreamerV3::dump(
StringRef HSAMetadataString)
const {
480 errs() <<
"AMDGPU HSA Metadata:\n" << HSAMetadataString <<
'\n';
483 void MetadataStreamerV3::verify(
StringRef HSAMetadataString)
const {
484 errs() <<
"AMDGPU HSA Metadata Parser Test: ";
488 if (!FromHSAMetadataString.
fromYAML(HSAMetadataString)) {
493 std::string ToHSAMetadataString;
495 FromHSAMetadataString.
toYAML(StrOS);
497 errs() << (HSAMetadataString == StrOS.
str() ?
"PASS" :
"FAIL") <<
'\n';
498 if (HSAMetadataString != ToHSAMetadataString) {
499 errs() <<
"Original input: " << HSAMetadataString <<
'\n' 500 <<
"Produced output: " << StrOS.
str() <<
'\n';
507 .Case(
"read_only",
StringRef(
"read_only"))
508 .Case(
"write_only",
StringRef(
"write_only"))
509 .Case(
"read_write",
StringRef(
"read_write"))
514 MetadataStreamerV3::getAddressSpaceQualifier(
unsigned AddressSpace)
const {
539 .
Case(
"image1d_t",
"image")
540 .
Case(
"image1d_array_t",
"image")
541 .
Case(
"image1d_buffer_t",
"image")
542 .
Case(
"image2d_t",
"image")
543 .
Case(
"image2d_array_t",
"image")
544 .
Case(
"image2d_array_depth_t",
"image")
545 .
Case(
"image2d_array_msaa_t",
"image")
546 .
Case(
"image2d_array_msaa_depth_t",
"image")
547 .
Case(
"image2d_depth_t",
"image")
548 .
Case(
"image2d_msaa_t",
"image")
549 .
Case(
"image2d_msaa_depth_t",
"image")
550 .
Case(
"image3d_t",
"image")
551 .
Case(
"sampler_t",
"sampler")
552 .
Case(
"queue_t",
"queue")
555 ?
"dynamic_shared_pointer" 560 std::string MetadataStreamerV3::getTypeName(
Type *Ty,
bool Signed)
const {
564 return (
Twine(
'u') + getTypeName(Ty,
true)).str();
587 auto VecTy = cast<FixedVectorType>(Ty);
588 auto ElTy = VecTy->getElementType();
589 auto NumElements = VecTy->getNumElements();
598 MetadataStreamerV3::getWorkGroupDimensions(
MDNode *Node)
const {
599 auto Dims = HSAMetadataDoc->getArrayNode();
600 if (Node->getNumOperands() != 3)
603 for (
auto &
Op : Node->operands())
604 Dims.push_back(Dims.getDocument()->getNode(
605 uint64_t(mdconst::extract<ConstantInt>(
Op)->getZExtValue())));
609 void MetadataStreamerV3::emitVersion() {
610 auto Version = HSAMetadataDoc->getArrayNode();
613 getRootMetadata(
"amdhsa.version") =
Version;
616 void MetadataStreamerV3::emitPrintf(
const Module &Mod) {
617 auto Node =
Mod.getNamedMetadata(
"llvm.printf.fmts");
621 auto Printf = HSAMetadataDoc->getArrayNode();
622 for (
auto Op : Node->operands())
623 if (
Op->getNumOperands())
625 cast<MDString>(
Op->getOperand(0))->getString(),
true));
626 getRootMetadata(
"amdhsa.printf") =
Printf;
629 void MetadataStreamerV3::emitKernelLanguage(
const Function &Func,
632 auto Node =
Func.getParent()->getNamedMetadata(
"opencl.ocl.version");
633 if (!Node || !Node->getNumOperands())
635 auto Op0 = Node->getOperand(0);
636 if (Op0->getNumOperands() <= 1)
642 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
644 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
648 void MetadataStreamerV3::emitKernelAttrs(
const Function &Func,
651 if (
auto Node =
Func.getMetadata(
"reqd_work_group_size"))
652 Kern[
".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
653 if (
auto Node =
Func.getMetadata(
"work_group_size_hint"))
654 Kern[
".workgroup_size_hint"] = getWorkGroupDimensions(Node);
655 if (
auto Node =
Func.getMetadata(
"vec_type_hint")) {
658 cast<ValueAsMetadata>(Node->getOperand(0))->
getType(),
659 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
662 if (
Func.hasFnAttribute(
"runtime-handle")) {
664 Func.getFnAttribute(
"runtime-handle").getValueAsString().str(),
669 void MetadataStreamerV3::emitKernelArgs(
const Function &Func,
672 auto Args = HSAMetadataDoc->getArrayNode();
678 Kern[
".args"] =
Args;
684 auto ArgNo =
Arg.getArgNo();
688 Node =
Func->getMetadata(
"kernel_arg_name");
689 if (Node && ArgNo < Node->getNumOperands())
690 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
691 else if (
Arg.hasName())
695 Node =
Func->getMetadata(
"kernel_arg_type");
696 if (Node && ArgNo < Node->getNumOperands())
697 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
700 Node =
Func->getMetadata(
"kernel_arg_base_type");
701 if (Node && ArgNo < Node->getNumOperands())
702 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
705 if (
Arg.getType()->isPointerTy() &&
Arg.onlyReadsMemory() &&
706 Arg.hasNoAliasAttr()) {
709 Node =
Func->getMetadata(
"kernel_arg_access_qual");
710 if (Node && ArgNo < Node->getNumOperands())
711 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
715 Node =
Func->getMetadata(
"kernel_arg_type_qual");
716 if (Node && ArgNo < Node->getNumOperands())
717 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
722 Type *Ty =
Arg.hasByRefAttr() ?
Arg.getParamByRefType() :
Arg.getType();
725 if (
auto PtrTy = dyn_cast<PointerType>(Ty)) {
728 PtrTy->getElementType());
737 emitKernelArg(
DL, ArgTy, ArgAlign,
738 getValueKind(ArgTy, TypeQual, BaseTypeName),
Offset,
Args,
742 void MetadataStreamerV3::emitKernelArg(
747 auto Arg =
Args.getDocument()->getMapNode();
750 Arg[
".name"] =
Arg.getDocument()->getNode(
Name,
true);
753 auto Size =
DL.getTypeAllocSize(Ty);
754 Arg[
".size"] =
Arg.getDocument()->getNode(
Size);
762 if (
auto PtrTy = dyn_cast<PointerType>(Ty))
763 if (
auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
764 Arg[
".address_space"] =
Arg.getDocument()->getNode(*Qualifier,
true);
766 if (
auto AQ = getAccessQualifier(
AccQual))
767 Arg[
".access"] =
Arg.getDocument()->getNode(*AQ,
true);
772 TypeQual.
split(SplitTypeQuals,
" ", -1,
false);
775 Arg[
".is_const"] =
Arg.getDocument()->getNode(
true);
776 else if (
Key ==
"restrict")
777 Arg[
".is_restrict"] =
Arg.getDocument()->getNode(
true);
778 else if (
Key ==
"volatile")
779 Arg[
".is_volatile"] =
Arg.getDocument()->getNode(
true);
780 else if (
Key ==
"pipe")
781 Arg[
".is_pipe"] =
Arg.getDocument()->getNode(
true);
787 void MetadataStreamerV3::emitHiddenKernelArgs(
const Function &Func,
790 int HiddenArgNumBytes =
793 if (!HiddenArgNumBytes)
796 auto &
DL =
Func.getParent()->getDataLayout();
799 if (HiddenArgNumBytes >= 8)
800 emitKernelArg(
DL, Int64Ty,
Align(8),
"hidden_global_offset_x",
Offset,
802 if (HiddenArgNumBytes >= 16)
803 emitKernelArg(
DL, Int64Ty,
Align(8),
"hidden_global_offset_y",
Offset,
805 if (HiddenArgNumBytes >= 24)
806 emitKernelArg(
DL, Int64Ty,
Align(8),
"hidden_global_offset_z",
Offset,
814 if (HiddenArgNumBytes >= 32) {
815 if (
Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"))
816 emitKernelArg(
DL, Int8PtrTy,
Align(8),
"hidden_printf_buffer",
Offset,
818 else if (
Func.getParent()->getFunction(
"__ockl_hostcall_internal")) {
821 assert(!
Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"));
822 emitKernelArg(
DL, Int8PtrTy,
Align(8),
"hidden_hostcall_buffer",
Offset,
830 if (HiddenArgNumBytes >= 48) {
831 if (
Func.hasFnAttribute(
"calls-enqueue-kernel")) {
832 emitKernelArg(
DL, Int8PtrTy,
Align(8),
"hidden_default_queue",
Offset,
834 emitKernelArg(
DL, Int8PtrTy,
Align(8),
"hidden_completion_action",
Offset,
843 if (HiddenArgNumBytes >= 56)
844 emitKernelArg(
DL, Int8PtrTy,
Align(8),
"hidden_multigrid_sync_arg",
Offset,
855 auto Kern = HSAMetadataDoc->getMapNode();
857 Align MaxKernArgAlign;
860 Kern[
".group_segment_fixed_size"] =
862 Kern[
".private_segment_fixed_size"] =
864 Kern[
".kernarg_segment_align"] =
866 Kern[
".wavefront_size"] =
870 Kern[
".max_flat_workgroup_size"] =
872 Kern[
".sgpr_spill_count"] =
874 Kern[
".vgpr_spill_count"] =
887 getRootMetadata(
"amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
891 std::string HSAMetadataString;
893 HSAMetadataDoc->toYAML(StrOS);
904 auto Kern = getHSAKernelProps(MF, ProgramInfo);
910 getRootMetadata(
"amdhsa.kernels").
getArray(
true);
915 (
Twine(Func.getName()) +
Twine(
".kd")).str(),
true);
916 emitKernelLanguage(Func, Kern);
917 emitKernelAttrs(Func, Kern);
918 emitKernelArgs(Func, Kern);
A parsed version of the target data layout string in and methods for querying it.
This class represents an incoming formal argument to a Function.
LLVM_NODISCARD std::string str() const
str - Get the contents as an std::string.
AMDGPU specific subclass of TargetSubtarget.
This class represents lattice values for constants.
A Module instance is used to store all the information related to an LLVM module.
32-bit floating point type
SPIR_KERNEL - Calling convention for SPIR kernel functions.
Document * getDocument() const
constexpr char PointeeAlign[]
Key for Kernel::Arg::Metadata::mPointeeAlign.
void toYAML(raw_ostream &OS)
Convert MsgPack Document to YAML text.
Address space for constant memory (VTX2).
Address space for local memory.
Function * Kernel
Summary of a kernel (=entry point for target offloading).
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
unsigned getKernArgSegmentSize(const Function &F, Align &MaxAlign) const
16-bit floating point type
static IntegerType * getInt64Ty(LLVMContext &C)
StringSwitch & Case(StringLiteral S, T Value)
Track resource usage for kernels / entry functions.
Function & getFunction()
Return the LLVM function that this machine code represents.
A DocNode that is an array.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
TypeID getTypeID() const
Return the type id for the type.
constexpr char Printf[]
Key for HSA::Metadata::mPrintf.
Defines struct to track resource usage and hardware flags for kernels and entry functions.
LLVM_NODISCARD R Default(T Value)
Fixed width SIMD vector type.
static constexpr size_t npos
bool fromYAML(StringRef S)
Read YAML text into the MsgPack document. Returns false on failure.
std::error_code fromString(std::string String, Metadata &HSAMetadata)
Converts String to HSAMetadata.
static cl::opt< bool > VerifyHSAMetadata("amdgpu-verify-hsa-metadata", cl::desc("Verify AMDGPU HSA Metadata"))
DocNode getNode()
Create a nil node associated with this Document.
uint64_t value() const
This is a hole in the type system and should not be abused.
constexpr char Attrs[]
Key for Kernel::Metadata::mAttrs.
Address space for private memory.
ArrayDocNode getArrayNode()
Create an empty Array node associated with this Document.
Arbitrary bit width integers.
A switch()-like statement whose cases are string literals.
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
The instances of the Type class are immutable: once they are created, they are never changed.
static cl::opt< bool > DumpHSAMetadata("amdgpu-dump-hsa-metadata", cl::desc("Dump AMDGPU HSA Metadata"))
Address space for global memory (RAT0, VTX0).
LLVM_NODISCARD size_t find(char C, size_t From=0) const
Search for the first character C in the string.
ArrayDocNode & getArray(bool Convert=false)
Get an ArrayDocNode for an array node.
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
constexpr char TypeName[]
Key for Kernel::Arg::Metadata::mTypeName.
amdgpu Simplify well known AMD library false FunctionCallee Value * Arg
raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
constexpr unsigned BitWidth
Address space for region memory. (GDS)
static wasm::ValType getType(const TargetRegisterClass *RC)
static PointerType * getInt8PtrTy(LLVMContext &C, unsigned AS=0)
std::string & str()
Flushes the stream contents to the target string and returns the string's reference.
This struct is a compact representation of a valid (non-zero power of two) alignment.
AccessQualifier
Access qualifiers.
unsigned getWavefrontSize() const
Align max(MaybeAlign Lhs, Align Rhs)
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Module.h This file contains the declarations for the Module class.
LLVM_NODISCARD std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
constexpr char CodeProps[]
Key for Kernel::Metadata::mCodeProps.
The access may modify the value stored in memory.
constexpr char LanguageVersion[]
Key for Kernel::Metadata::mLanguageVersion.
Address space for flat memory.
constexpr uint32_t VersionMinor
HSA metadata minor version.
Simple in-memory representation of a document of msgpack objects with ability to find and create arra...
constexpr char Kernels[]
Key for HSA::Metadata::mKernels.
std::error_code toString(Metadata HSAMetadata, std::string &String)
Converts HSAMetadata to String.
This class keeps track of the SPI_SP_INPUT_ADDR config register, which tells the hardware which inter...
uint64_t alignTo(uint64_t Size, Align A)
Returns a multiple of A needed to store Size bytes.
unsigned getIntegerBitWidth() const
constexpr char DebugProps[]
Key for Kernel::Metadata::mDebugProps.
AddressSpaceQualifier
Address space qualifiers.
int getIntegerAttribute(const Function &F, StringRef Name, int Default)
bool isXNACKEnabled() const
Calling convention for AMDGPU code object kernels.
constexpr char AccQual[]
Key for Kernel::Arg::Metadata::mAccQual.
64-bit floating point type
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict)=0
Emit HSA Metadata.
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
A raw_ostream that writes to an std::string.
StringRef - Represent a constant reference to a string, i.e.
constexpr char Args[]
Key for Kernel::Metadata::mArgs.
constexpr uint32_t VersionMajor
HSA metadata major version.
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL