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(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)
445 std::string HSAMetadataString;
446 if (
toString(HSAMetadata, HSAMetadataString))
450 dump(HSAMetadataString);
452 verify(HSAMetadataString);
461 auto CodeProps = getHSACodeProps(MF, ProgramInfo);
462 auto DebugProps = getHSADebugProps(MF, ProgramInfo);
467 Kernel.mName = std::string(Func.getName());
469 emitKernelLanguage(Func);
470 emitKernelAttrs(Func);
471 emitKernelArgs(Func);
481 errs() <<
"AMDGPU HSA Metadata:\n" << HSAMetadataString <<
'\n';
485 errs() <<
"AMDGPU HSA Metadata Parser Test: ";
489 if (!FromHSAMetadataString.
fromYAML(HSAMetadataString)) {
494 std::string ToHSAMetadataString;
496 FromHSAMetadataString.
toYAML(StrOS);
498 errs() << (HSAMetadataString == StrOS.
str() ?
"PASS" :
"FAIL") <<
'\n';
499 if (HSAMetadataString != ToHSAMetadataString) {
500 errs() <<
"Original input: " << HSAMetadataString <<
'\n'
501 <<
"Produced output: " << StrOS.
str() <<
'\n';
508 .Case(
"read_only",
StringRef(
"read_only"))
509 .Case(
"write_only",
StringRef(
"write_only"))
510 .Case(
"read_write",
StringRef(
"read_write"))
540 .
Case(
"image1d_t",
"image")
541 .
Case(
"image1d_array_t",
"image")
542 .
Case(
"image1d_buffer_t",
"image")
543 .
Case(
"image2d_t",
"image")
544 .
Case(
"image2d_array_t",
"image")
545 .
Case(
"image2d_array_depth_t",
"image")
546 .
Case(
"image2d_array_msaa_t",
"image")
547 .
Case(
"image2d_array_msaa_depth_t",
"image")
548 .
Case(
"image2d_depth_t",
"image")
549 .
Case(
"image2d_msaa_t",
"image")
550 .
Case(
"image2d_msaa_depth_t",
"image")
551 .
Case(
"image3d_t",
"image")
552 .
Case(
"sampler_t",
"sampler")
553 .
Case(
"queue_t",
"queue")
556 ?
"dynamic_shared_pointer"
588 auto VecTy = cast<FixedVectorType>(Ty);
589 auto ElTy = VecTy->getElementType();
590 auto NumElements = VecTy->getNumElements();
601 if (Node->getNumOperands() != 3)
604 for (
auto &
Op : Node->operands())
605 Dims.push_back(Dims.getDocument()->getNode(
606 uint64_t(mdconst::extract<ConstantInt>(
Op)->getZExtValue())));
618 auto Node =
Mod.getNamedMetadata(
"llvm.printf.fmts");
623 for (
auto Op : Node->operands())
624 if (
Op->getNumOperands())
626 cast<MDString>(
Op->getOperand(0))->getString(),
true));
633 auto Node = Func.getParent()->getNamedMetadata(
"opencl.ocl.version");
634 if (!Node || !Node->getNumOperands())
636 auto Op0 = Node->getOperand(0);
637 if (Op0->getNumOperands() <= 1)
643 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
645 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
652 if (
auto Node = Func.getMetadata(
"reqd_work_group_size"))
654 if (
auto Node = Func.getMetadata(
"work_group_size_hint"))
656 if (
auto Node = Func.getMetadata(
"vec_type_hint")) {
659 cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
660 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
663 if (Func.hasFnAttribute(
"runtime-handle")) {
665 Func.getFnAttribute(
"runtime-handle").getValueAsString().str(),
674 for (
auto &
Arg : Func.args())
679 Kern[
".args"] =
Args;
684 auto Func =
Arg.getParent();
685 auto ArgNo =
Arg.getArgNo();
689 Node = Func->getMetadata(
"kernel_arg_name");
690 if (Node && ArgNo < Node->getNumOperands())
691 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
692 else if (
Arg.hasName())
696 Node = Func->getMetadata(
"kernel_arg_type");
697 if (Node && ArgNo < Node->getNumOperands())
698 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
701 Node = Func->getMetadata(
"kernel_arg_base_type");
702 if (Node && ArgNo < Node->getNumOperands())
703 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
706 if (
Arg.getType()->isPointerTy() &&
Arg.onlyReadsMemory() &&
707 Arg.hasNoAliasAttr()) {
710 Node = Func->getMetadata(
"kernel_arg_access_qual");
711 if (Node && ArgNo < Node->getNumOperands())
712 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
716 Node = Func->getMetadata(
"kernel_arg_type_qual");
717 if (Node && ArgNo < Node->getNumOperands())
718 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
720 const DataLayout &
DL = Func->getParent()->getDataLayout();
723 Type *Ty =
Arg.hasByRefAttr() ?
Arg.getParamByRefType() :
Arg.getType();
726 if (
auto PtrTy = dyn_cast<PointerType>(Ty)) {
729 PtrTy->getElementType());
748 auto Arg =
Args.getDocument()->getMapNode();
751 Arg[
".name"] =
Arg.getDocument()->getNode(
Name,
true);
754 auto Size =
DL.getTypeAllocSize(Ty);
755 Arg[
".size"] =
Arg.getDocument()->getNode(
Size);
763 if (
auto PtrTy = dyn_cast<PointerType>(Ty))
765 Arg[
".address_space"] =
Arg.getDocument()->getNode(*Qualifier,
true);
768 Arg[
".access"] =
Arg.getDocument()->getNode(*AQ,
true);
773 TypeQual.
split(SplitTypeQuals,
" ", -1,
false);
776 Arg[
".is_const"] =
Arg.getDocument()->getNode(
true);
777 else if (
Key ==
"restrict")
778 Arg[
".is_restrict"] =
Arg.getDocument()->getNode(
true);
779 else if (
Key ==
"volatile")
780 Arg[
".is_volatile"] =
Arg.getDocument()->getNode(
true);
781 else if (
Key ==
"pipe")
782 Arg[
".is_pipe"] =
Arg.getDocument()->getNode(
true);
791 int HiddenArgNumBytes =
794 if (!HiddenArgNumBytes)
797 auto &
DL = Func.getParent()->getDataLayout();
800 if (HiddenArgNumBytes >= 8)
803 if (HiddenArgNumBytes >= 16)
806 if (HiddenArgNumBytes >= 24)
815 if (HiddenArgNumBytes >= 32) {
816 if (Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"))
819 else if (Func.getParent()->getFunction(
"__ockl_hostcall_internal")) {
822 assert(!Func.getParent()->getNamedMetadata(
"llvm.printf.fmts"));
831 if (HiddenArgNumBytes >= 48) {
832 if (Func.hasFnAttribute(
"calls-enqueue-kernel")) {
844 if (HiddenArgNumBytes >= 56)
858 Align MaxKernArgAlign;
859 Kern[
".kernarg_segment_size"] = Kern.getDocument()->getNode(
861 Kern[
".group_segment_fixed_size"] =
862 Kern.getDocument()->getNode(ProgramInfo.
LDSSize);
863 Kern[
".private_segment_fixed_size"] =
864 Kern.getDocument()->getNode(ProgramInfo.
ScratchSize);
865 Kern[
".kernarg_segment_align"] =
866 Kern.getDocument()->getNode(
std::max(
Align(4), MaxKernArgAlign).value());
867 Kern[
".wavefront_size"] =
869 Kern[
".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.
NumSGPR);
870 Kern[
".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.
NumVGPR);
871 Kern[
".max_flat_workgroup_size"] =
872 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
873 Kern[
".sgpr_spill_count"] =
874 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
875 Kern[
".vgpr_spill_count"] =
876 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
893 std::string HSAMetadataString;
917 (
Twine(Func.getName()) +
Twine(
".kd")).str(),
true);
930 void MetadataStreamerV4::emitVersion() {
945 emitTargetID(TargetID);