1 //===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===// 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 /// \file 10 /// AMDGPU HSA Metadata Streamer. 11 /// 12 // 13 //===----------------------------------------------------------------------===// 14 15 #include "AMDGPUHSAMetadataStreamer.h" 16 #include "AMDGPU.h" 17 #include "GCNSubtarget.h" 18 #include "MCTargetDesc/AMDGPUTargetStreamer.h" 19 #include "SIMachineFunctionInfo.h" 20 #include "SIProgramInfo.h" 21 #include "llvm/IR/Module.h" 22 using namespace llvm; 23 24 static std::pair<Type *, Align> getArgumentTypeAlign(const Argument &Arg, 25 const DataLayout &DL) { 26 Type *Ty = Arg.getType(); 27 MaybeAlign ArgAlign; 28 if (Arg.hasByRefAttr()) { 29 Ty = Arg.getParamByRefType(); 30 ArgAlign = Arg.getParamAlign(); 31 } 32 33 if (!ArgAlign) 34 ArgAlign = DL.getABITypeAlign(Ty); 35 36 return std::make_pair(Ty, *ArgAlign); 37 } 38 39 namespace llvm { 40 41 static cl::opt<bool> DumpHSAMetadata( 42 "amdgpu-dump-hsa-metadata", 43 cl::desc("Dump AMDGPU HSA Metadata")); 44 static cl::opt<bool> VerifyHSAMetadata( 45 "amdgpu-verify-hsa-metadata", 46 cl::desc("Verify AMDGPU HSA Metadata")); 47 48 namespace AMDGPU { 49 namespace HSAMD { 50 51 //===----------------------------------------------------------------------===// 52 // HSAMetadataStreamerV2 53 //===----------------------------------------------------------------------===// 54 void MetadataStreamerV2::dump(StringRef HSAMetadataString) const { 55 errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n'; 56 } 57 58 void MetadataStreamerV2::verify(StringRef HSAMetadataString) const { 59 errs() << "AMDGPU HSA Metadata Parser Test: "; 60 61 HSAMD::Metadata FromHSAMetadataString; 62 if (fromString(HSAMetadataString, FromHSAMetadataString)) { 63 errs() << "FAIL\n"; 64 return; 65 } 66 67 std::string ToHSAMetadataString; 68 if (toString(FromHSAMetadataString, ToHSAMetadataString)) { 69 errs() << "FAIL\n"; 70 return; 71 } 72 73 errs() << (HSAMetadataString == ToHSAMetadataString ? "PASS" : "FAIL") 74 << '\n'; 75 if (HSAMetadataString != ToHSAMetadataString) { 76 errs() << "Original input: " << HSAMetadataString << '\n' 77 << "Produced output: " << ToHSAMetadataString << '\n'; 78 } 79 } 80 81 AccessQualifier 82 MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const { 83 if (AccQual.empty()) 84 return AccessQualifier::Unknown; 85 86 return StringSwitch<AccessQualifier>(AccQual) 87 .Case("read_only", AccessQualifier::ReadOnly) 88 .Case("write_only", AccessQualifier::WriteOnly) 89 .Case("read_write", AccessQualifier::ReadWrite) 90 .Default(AccessQualifier::Default); 91 } 92 93 AddressSpaceQualifier 94 MetadataStreamerV2::getAddressSpaceQualifier( 95 unsigned AddressSpace) const { 96 switch (AddressSpace) { 97 case AMDGPUAS::PRIVATE_ADDRESS: 98 return AddressSpaceQualifier::Private; 99 case AMDGPUAS::GLOBAL_ADDRESS: 100 return AddressSpaceQualifier::Global; 101 case AMDGPUAS::CONSTANT_ADDRESS: 102 return AddressSpaceQualifier::Constant; 103 case AMDGPUAS::LOCAL_ADDRESS: 104 return AddressSpaceQualifier::Local; 105 case AMDGPUAS::FLAT_ADDRESS: 106 return AddressSpaceQualifier::Generic; 107 case AMDGPUAS::REGION_ADDRESS: 108 return AddressSpaceQualifier::Region; 109 default: 110 return AddressSpaceQualifier::Unknown; 111 } 112 } 113 114 ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual, 115 StringRef BaseTypeName) const { 116 if (TypeQual.contains("pipe")) 117 return ValueKind::Pipe; 118 119 return StringSwitch<ValueKind>(BaseTypeName) 120 .Case("image1d_t", ValueKind::Image) 121 .Case("image1d_array_t", ValueKind::Image) 122 .Case("image1d_buffer_t", ValueKind::Image) 123 .Case("image2d_t", ValueKind::Image) 124 .Case("image2d_array_t", ValueKind::Image) 125 .Case("image2d_array_depth_t", ValueKind::Image) 126 .Case("image2d_array_msaa_t", ValueKind::Image) 127 .Case("image2d_array_msaa_depth_t", ValueKind::Image) 128 .Case("image2d_depth_t", ValueKind::Image) 129 .Case("image2d_msaa_t", ValueKind::Image) 130 .Case("image2d_msaa_depth_t", ValueKind::Image) 131 .Case("image3d_t", ValueKind::Image) 132 .Case("sampler_t", ValueKind::Sampler) 133 .Case("queue_t", ValueKind::Queue) 134 .Default(isa<PointerType>(Ty) ? 135 (Ty->getPointerAddressSpace() == 136 AMDGPUAS::LOCAL_ADDRESS ? 137 ValueKind::DynamicSharedPointer : 138 ValueKind::GlobalBuffer) : 139 ValueKind::ByValue); 140 } 141 142 std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const { 143 switch (Ty->getTypeID()) { 144 case Type::IntegerTyID: { 145 if (!Signed) 146 return (Twine('u') + getTypeName(Ty, true)).str(); 147 148 auto BitWidth = Ty->getIntegerBitWidth(); 149 switch (BitWidth) { 150 case 8: 151 return "char"; 152 case 16: 153 return "short"; 154 case 32: 155 return "int"; 156 case 64: 157 return "long"; 158 default: 159 return (Twine('i') + Twine(BitWidth)).str(); 160 } 161 } 162 case Type::HalfTyID: 163 return "half"; 164 case Type::FloatTyID: 165 return "float"; 166 case Type::DoubleTyID: 167 return "double"; 168 case Type::FixedVectorTyID: { 169 auto VecTy = cast<FixedVectorType>(Ty); 170 auto ElTy = VecTy->getElementType(); 171 auto NumElements = VecTy->getNumElements(); 172 return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str(); 173 } 174 default: 175 return "unknown"; 176 } 177 } 178 179 std::vector<uint32_t> 180 MetadataStreamerV2::getWorkGroupDimensions(MDNode *Node) const { 181 std::vector<uint32_t> Dims; 182 if (Node->getNumOperands() != 3) 183 return Dims; 184 185 for (auto &Op : Node->operands()) 186 Dims.push_back(mdconst::extract<ConstantInt>(Op)->getZExtValue()); 187 return Dims; 188 } 189 190 Kernel::CodeProps::Metadata 191 MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF, 192 const SIProgramInfo &ProgramInfo) const { 193 const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>(); 194 const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>(); 195 HSAMD::Kernel::CodeProps::Metadata HSACodeProps; 196 const Function &F = MF.getFunction(); 197 198 assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL || 199 F.getCallingConv() == CallingConv::SPIR_KERNEL); 200 201 Align MaxKernArgAlign; 202 HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F, 203 MaxKernArgAlign); 204 HSACodeProps.mKernargSegmentAlign = 205 std::max(MaxKernArgAlign, Align(4)).value(); 206 207 HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize; 208 HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize; 209 HSACodeProps.mWavefrontSize = STM.getWavefrontSize(); 210 HSACodeProps.mNumSGPRs = ProgramInfo.NumSGPR; 211 HSACodeProps.mNumVGPRs = ProgramInfo.NumVGPR; 212 HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize(); 213 HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack; 214 HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled(); 215 HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs(); 216 HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs(); 217 218 return HSACodeProps; 219 } 220 221 Kernel::DebugProps::Metadata 222 MetadataStreamerV2::getHSADebugProps(const MachineFunction &MF, 223 const SIProgramInfo &ProgramInfo) const { 224 return HSAMD::Kernel::DebugProps::Metadata(); 225 } 226 227 void MetadataStreamerV2::emitVersion() { 228 auto &Version = HSAMetadata.mVersion; 229 230 Version.push_back(VersionMajorV2); 231 Version.push_back(VersionMinorV2); 232 } 233 234 void MetadataStreamerV2::emitPrintf(const Module &Mod) { 235 auto &Printf = HSAMetadata.mPrintf; 236 237 auto Node = Mod.getNamedMetadata("llvm.printf.fmts"); 238 if (!Node) 239 return; 240 241 for (auto Op : Node->operands()) 242 if (Op->getNumOperands()) 243 Printf.push_back( 244 std::string(cast<MDString>(Op->getOperand(0))->getString())); 245 } 246 247 void MetadataStreamerV2::emitKernelLanguage(const Function &Func) { 248 auto &Kernel = HSAMetadata.mKernels.back(); 249 250 // TODO: What about other languages? 251 auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version"); 252 if (!Node || !Node->getNumOperands()) 253 return; 254 auto Op0 = Node->getOperand(0); 255 if (Op0->getNumOperands() <= 1) 256 return; 257 258 Kernel.mLanguage = "OpenCL C"; 259 Kernel.mLanguageVersion.push_back( 260 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()); 261 Kernel.mLanguageVersion.push_back( 262 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()); 263 } 264 265 void MetadataStreamerV2::emitKernelAttrs(const Function &Func) { 266 auto &Attrs = HSAMetadata.mKernels.back().mAttrs; 267 268 if (auto Node = Func.getMetadata("reqd_work_group_size")) 269 Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node); 270 if (auto Node = Func.getMetadata("work_group_size_hint")) 271 Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node); 272 if (auto Node = Func.getMetadata("vec_type_hint")) { 273 Attrs.mVecTypeHint = getTypeName( 274 cast<ValueAsMetadata>(Node->getOperand(0))->getType(), 275 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()); 276 } 277 if (Func.hasFnAttribute("runtime-handle")) { 278 Attrs.mRuntimeHandle = 279 Func.getFnAttribute("runtime-handle").getValueAsString().str(); 280 } 281 } 282 283 void MetadataStreamerV2::emitKernelArgs(const Function &Func, 284 const GCNSubtarget &ST) { 285 for (auto &Arg : Func.args()) 286 emitKernelArg(Arg); 287 288 emitHiddenKernelArgs(Func, ST); 289 } 290 291 void MetadataStreamerV2::emitKernelArg(const Argument &Arg) { 292 auto Func = Arg.getParent(); 293 auto ArgNo = Arg.getArgNo(); 294 const MDNode *Node; 295 296 StringRef Name; 297 Node = Func->getMetadata("kernel_arg_name"); 298 if (Node && ArgNo < Node->getNumOperands()) 299 Name = cast<MDString>(Node->getOperand(ArgNo))->getString(); 300 else if (Arg.hasName()) 301 Name = Arg.getName(); 302 303 StringRef TypeName; 304 Node = Func->getMetadata("kernel_arg_type"); 305 if (Node && ArgNo < Node->getNumOperands()) 306 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString(); 307 308 StringRef BaseTypeName; 309 Node = Func->getMetadata("kernel_arg_base_type"); 310 if (Node && ArgNo < Node->getNumOperands()) 311 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString(); 312 313 StringRef AccQual; 314 if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() && 315 Arg.hasNoAliasAttr()) { 316 AccQual = "read_only"; 317 } else { 318 Node = Func->getMetadata("kernel_arg_access_qual"); 319 if (Node && ArgNo < Node->getNumOperands()) 320 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString(); 321 } 322 323 StringRef TypeQual; 324 Node = Func->getMetadata("kernel_arg_type_qual"); 325 if (Node && ArgNo < Node->getNumOperands()) 326 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString(); 327 328 const DataLayout &DL = Func->getParent()->getDataLayout(); 329 330 MaybeAlign PointeeAlign; 331 if (auto PtrTy = dyn_cast<PointerType>(Arg.getType())) { 332 if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) { 333 // FIXME: Should report this for all address spaces 334 PointeeAlign = Arg.getParamAlign().valueOrOne(); 335 } 336 } 337 338 Type *ArgTy; 339 Align ArgAlign; 340 std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL); 341 342 emitKernelArg(DL, ArgTy, ArgAlign, 343 getValueKind(ArgTy, TypeQual, BaseTypeName), PointeeAlign, Name, 344 TypeName, BaseTypeName, AccQual, TypeQual); 345 } 346 347 void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty, 348 Align Alignment, ValueKind ValueKind, 349 MaybeAlign PointeeAlign, StringRef Name, 350 StringRef TypeName, 351 StringRef BaseTypeName, 352 StringRef AccQual, StringRef TypeQual) { 353 HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata()); 354 auto &Arg = HSAMetadata.mKernels.back().mArgs.back(); 355 356 Arg.mName = std::string(Name); 357 Arg.mTypeName = std::string(TypeName); 358 Arg.mSize = DL.getTypeAllocSize(Ty); 359 Arg.mAlign = Alignment.value(); 360 Arg.mValueKind = ValueKind; 361 Arg.mPointeeAlign = PointeeAlign ? PointeeAlign->value() : 0; 362 363 if (auto PtrTy = dyn_cast<PointerType>(Ty)) 364 Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace()); 365 366 Arg.mAccQual = getAccessQualifier(AccQual); 367 368 // TODO: Emit Arg.mActualAccQual. 369 370 SmallVector<StringRef, 1> SplitTypeQuals; 371 TypeQual.split(SplitTypeQuals, " ", -1, false); 372 for (StringRef Key : SplitTypeQuals) { 373 auto P = StringSwitch<bool*>(Key) 374 .Case("const", &Arg.mIsConst) 375 .Case("restrict", &Arg.mIsRestrict) 376 .Case("volatile", &Arg.mIsVolatile) 377 .Case("pipe", &Arg.mIsPipe) 378 .Default(nullptr); 379 if (P) 380 *P = true; 381 } 382 } 383 384 void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func, 385 const GCNSubtarget &ST) { 386 unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func); 387 if (!HiddenArgNumBytes) 388 return; 389 390 auto &DL = Func.getParent()->getDataLayout(); 391 auto Int64Ty = Type::getInt64Ty(Func.getContext()); 392 393 if (HiddenArgNumBytes >= 8) 394 emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetX); 395 if (HiddenArgNumBytes >= 16) 396 emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetY); 397 if (HiddenArgNumBytes >= 24) 398 emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetZ); 399 400 auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(), 401 AMDGPUAS::GLOBAL_ADDRESS); 402 403 // Emit "printf buffer" argument if printf is used, otherwise emit dummy 404 // "none" argument. 405 if (HiddenArgNumBytes >= 32) { 406 if (Func.getParent()->getNamedMetadata("llvm.printf.fmts")) 407 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer); 408 else if (Func.getParent()->getFunction("__ockl_hostcall_internal")) { 409 // The printf runtime binding pass should have ensured that hostcall and 410 // printf are not used in the same module. 411 assert(!Func.getParent()->getNamedMetadata("llvm.printf.fmts")); 412 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer); 413 } else 414 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone); 415 } 416 417 // Emit "default queue" and "completion action" arguments if enqueue kernel is 418 // used, otherwise emit dummy "none" arguments. 419 if (HiddenArgNumBytes >= 48) { 420 if (Func.hasFnAttribute("calls-enqueue-kernel")) { 421 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenDefaultQueue); 422 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenCompletionAction); 423 } else { 424 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone); 425 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone); 426 } 427 } 428 429 // Emit the pointer argument for multi-grid object. 430 if (HiddenArgNumBytes >= 56) 431 emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenMultiGridSyncArg); 432 } 433 434 bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) { 435 return TargetStreamer.EmitHSAMetadata(getHSAMetadata()); 436 } 437 438 void MetadataStreamerV2::begin(const Module &Mod, 439 const IsaInfo::AMDGPUTargetID &TargetID) { 440 emitVersion(); 441 emitPrintf(Mod); 442 } 443 444 void MetadataStreamerV2::end() { 445 std::string HSAMetadataString; 446 if (toString(HSAMetadata, HSAMetadataString)) 447 return; 448 449 if (DumpHSAMetadata) 450 dump(HSAMetadataString); 451 if (VerifyHSAMetadata) 452 verify(HSAMetadataString); 453 } 454 455 void MetadataStreamerV2::emitKernel(const MachineFunction &MF, 456 const SIProgramInfo &ProgramInfo) { 457 auto &Func = MF.getFunction(); 458 if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL) 459 return; 460 461 auto CodeProps = getHSACodeProps(MF, ProgramInfo); 462 auto DebugProps = getHSADebugProps(MF, ProgramInfo); 463 464 HSAMetadata.mKernels.push_back(Kernel::Metadata()); 465 auto &Kernel = HSAMetadata.mKernels.back(); 466 467 const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>(); 468 Kernel.mName = std::string(Func.getName()); 469 Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str(); 470 emitKernelLanguage(Func); 471 emitKernelAttrs(Func); 472 emitKernelArgs(Func, ST); 473 HSAMetadata.mKernels.back().mCodeProps = CodeProps; 474 HSAMetadata.mKernels.back().mDebugProps = DebugProps; 475 } 476 477 //===----------------------------------------------------------------------===// 478 // HSAMetadataStreamerV3 479 //===----------------------------------------------------------------------===// 480 481 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const { 482 errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n'; 483 } 484 485 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const { 486 errs() << "AMDGPU HSA Metadata Parser Test: "; 487 488 msgpack::Document FromHSAMetadataString; 489 490 if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) { 491 errs() << "FAIL\n"; 492 return; 493 } 494 495 std::string ToHSAMetadataString; 496 raw_string_ostream StrOS(ToHSAMetadataString); 497 FromHSAMetadataString.toYAML(StrOS); 498 499 errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n'; 500 if (HSAMetadataString != ToHSAMetadataString) { 501 errs() << "Original input: " << HSAMetadataString << '\n' 502 << "Produced output: " << StrOS.str() << '\n'; 503 } 504 } 505 506 Optional<StringRef> 507 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const { 508 return StringSwitch<Optional<StringRef>>(AccQual) 509 .Case("read_only", StringRef("read_only")) 510 .Case("write_only", StringRef("write_only")) 511 .Case("read_write", StringRef("read_write")) 512 .Default(None); 513 } 514 515 Optional<StringRef> 516 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const { 517 switch (AddressSpace) { 518 case AMDGPUAS::PRIVATE_ADDRESS: 519 return StringRef("private"); 520 case AMDGPUAS::GLOBAL_ADDRESS: 521 return StringRef("global"); 522 case AMDGPUAS::CONSTANT_ADDRESS: 523 return StringRef("constant"); 524 case AMDGPUAS::LOCAL_ADDRESS: 525 return StringRef("local"); 526 case AMDGPUAS::FLAT_ADDRESS: 527 return StringRef("generic"); 528 case AMDGPUAS::REGION_ADDRESS: 529 return StringRef("region"); 530 default: 531 return None; 532 } 533 } 534 535 StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual, 536 StringRef BaseTypeName) const { 537 if (TypeQual.contains("pipe")) 538 return "pipe"; 539 540 return StringSwitch<StringRef>(BaseTypeName) 541 .Case("image1d_t", "image") 542 .Case("image1d_array_t", "image") 543 .Case("image1d_buffer_t", "image") 544 .Case("image2d_t", "image") 545 .Case("image2d_array_t", "image") 546 .Case("image2d_array_depth_t", "image") 547 .Case("image2d_array_msaa_t", "image") 548 .Case("image2d_array_msaa_depth_t", "image") 549 .Case("image2d_depth_t", "image") 550 .Case("image2d_msaa_t", "image") 551 .Case("image2d_msaa_depth_t", "image") 552 .Case("image3d_t", "image") 553 .Case("sampler_t", "sampler") 554 .Case("queue_t", "queue") 555 .Default(isa<PointerType>(Ty) 556 ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS 557 ? "dynamic_shared_pointer" 558 : "global_buffer") 559 : "by_value"); 560 } 561 562 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const { 563 switch (Ty->getTypeID()) { 564 case Type::IntegerTyID: { 565 if (!Signed) 566 return (Twine('u') + getTypeName(Ty, true)).str(); 567 568 auto BitWidth = Ty->getIntegerBitWidth(); 569 switch (BitWidth) { 570 case 8: 571 return "char"; 572 case 16: 573 return "short"; 574 case 32: 575 return "int"; 576 case 64: 577 return "long"; 578 default: 579 return (Twine('i') + Twine(BitWidth)).str(); 580 } 581 } 582 case Type::HalfTyID: 583 return "half"; 584 case Type::FloatTyID: 585 return "float"; 586 case Type::DoubleTyID: 587 return "double"; 588 case Type::FixedVectorTyID: { 589 auto VecTy = cast<FixedVectorType>(Ty); 590 auto ElTy = VecTy->getElementType(); 591 auto NumElements = VecTy->getNumElements(); 592 return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str(); 593 } 594 default: 595 return "unknown"; 596 } 597 } 598 599 msgpack::ArrayDocNode 600 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const { 601 auto Dims = HSAMetadataDoc->getArrayNode(); 602 if (Node->getNumOperands() != 3) 603 return Dims; 604 605 for (auto &Op : Node->operands()) 606 Dims.push_back(Dims.getDocument()->getNode( 607 uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue()))); 608 return Dims; 609 } 610 611 void MetadataStreamerV3::emitVersion() { 612 auto Version = HSAMetadataDoc->getArrayNode(); 613 Version.push_back(Version.getDocument()->getNode(VersionMajorV3)); 614 Version.push_back(Version.getDocument()->getNode(VersionMinorV3)); 615 getRootMetadata("amdhsa.version") = Version; 616 } 617 618 void MetadataStreamerV3::emitPrintf(const Module &Mod) { 619 auto Node = Mod.getNamedMetadata("llvm.printf.fmts"); 620 if (!Node) 621 return; 622 623 auto Printf = HSAMetadataDoc->getArrayNode(); 624 for (auto Op : Node->operands()) 625 if (Op->getNumOperands()) 626 Printf.push_back(Printf.getDocument()->getNode( 627 cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true)); 628 getRootMetadata("amdhsa.printf") = Printf; 629 } 630 631 void MetadataStreamerV3::emitKernelLanguage(const Function &Func, 632 msgpack::MapDocNode Kern) { 633 // TODO: What about other languages? 634 auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version"); 635 if (!Node || !Node->getNumOperands()) 636 return; 637 auto Op0 = Node->getOperand(0); 638 if (Op0->getNumOperands() <= 1) 639 return; 640 641 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C"); 642 auto LanguageVersion = Kern.getDocument()->getArrayNode(); 643 LanguageVersion.push_back(Kern.getDocument()->getNode( 644 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue())); 645 LanguageVersion.push_back(Kern.getDocument()->getNode( 646 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue())); 647 Kern[".language_version"] = LanguageVersion; 648 } 649 650 void MetadataStreamerV3::emitKernelAttrs(const Function &Func, 651 msgpack::MapDocNode Kern) { 652 653 if (auto Node = Func.getMetadata("reqd_work_group_size")) 654 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node); 655 if (auto Node = Func.getMetadata("work_group_size_hint")) 656 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node); 657 if (auto Node = Func.getMetadata("vec_type_hint")) { 658 Kern[".vec_type_hint"] = Kern.getDocument()->getNode( 659 getTypeName( 660 cast<ValueAsMetadata>(Node->getOperand(0))->getType(), 661 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()), 662 /*Copy=*/true); 663 } 664 if (Func.hasFnAttribute("runtime-handle")) { 665 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode( 666 Func.getFnAttribute("runtime-handle").getValueAsString().str(), 667 /*Copy=*/true); 668 } 669 if (Func.hasFnAttribute("device-init")) 670 Kern[".kind"] = Kern.getDocument()->getNode("init"); 671 else if (Func.hasFnAttribute("device-fini")) 672 Kern[".kind"] = Kern.getDocument()->getNode("fini"); 673 } 674 675 void MetadataStreamerV3::emitKernelArgs(const MachineFunction &MF, 676 msgpack::MapDocNode Kern) { 677 auto &Func = MF.getFunction(); 678 unsigned Offset = 0; 679 auto Args = HSAMetadataDoc->getArrayNode(); 680 for (auto &Arg : Func.args()) 681 emitKernelArg(Arg, Offset, Args); 682 683 emitHiddenKernelArgs(MF, Offset, Args); 684 685 Kern[".args"] = Args; 686 } 687 688 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset, 689 msgpack::ArrayDocNode Args) { 690 auto Func = Arg.getParent(); 691 auto ArgNo = Arg.getArgNo(); 692 const MDNode *Node; 693 694 StringRef Name; 695 Node = Func->getMetadata("kernel_arg_name"); 696 if (Node && ArgNo < Node->getNumOperands()) 697 Name = cast<MDString>(Node->getOperand(ArgNo))->getString(); 698 else if (Arg.hasName()) 699 Name = Arg.getName(); 700 701 StringRef TypeName; 702 Node = Func->getMetadata("kernel_arg_type"); 703 if (Node && ArgNo < Node->getNumOperands()) 704 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString(); 705 706 StringRef BaseTypeName; 707 Node = Func->getMetadata("kernel_arg_base_type"); 708 if (Node && ArgNo < Node->getNumOperands()) 709 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString(); 710 711 StringRef AccQual; 712 if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() && 713 Arg.hasNoAliasAttr()) { 714 AccQual = "read_only"; 715 } else { 716 Node = Func->getMetadata("kernel_arg_access_qual"); 717 if (Node && ArgNo < Node->getNumOperands()) 718 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString(); 719 } 720 721 StringRef TypeQual; 722 Node = Func->getMetadata("kernel_arg_type_qual"); 723 if (Node && ArgNo < Node->getNumOperands()) 724 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString(); 725 726 const DataLayout &DL = Func->getParent()->getDataLayout(); 727 728 MaybeAlign PointeeAlign; 729 Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType(); 730 731 // FIXME: Need to distinguish in memory alignment from pointer alignment. 732 if (auto PtrTy = dyn_cast<PointerType>(Ty)) { 733 if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) 734 PointeeAlign = Arg.getParamAlign().valueOrOne(); 735 } 736 737 // There's no distinction between byval aggregates and raw aggregates. 738 Type *ArgTy; 739 Align ArgAlign; 740 std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL); 741 742 emitKernelArg(DL, ArgTy, ArgAlign, 743 getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args, 744 PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual); 745 } 746 747 void MetadataStreamerV3::emitKernelArg( 748 const DataLayout &DL, Type *Ty, Align Alignment, StringRef ValueKind, 749 unsigned &Offset, msgpack::ArrayDocNode Args, MaybeAlign PointeeAlign, 750 StringRef Name, StringRef TypeName, StringRef BaseTypeName, 751 StringRef AccQual, StringRef TypeQual) { 752 auto Arg = Args.getDocument()->getMapNode(); 753 754 if (!Name.empty()) 755 Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true); 756 if (!TypeName.empty()) 757 Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true); 758 auto Size = DL.getTypeAllocSize(Ty); 759 Arg[".size"] = Arg.getDocument()->getNode(Size); 760 Offset = alignTo(Offset, Alignment); 761 Arg[".offset"] = Arg.getDocument()->getNode(Offset); 762 Offset += Size; 763 Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true); 764 if (PointeeAlign) 765 Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value()); 766 767 if (auto PtrTy = dyn_cast<PointerType>(Ty)) 768 if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace())) 769 Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true); 770 771 if (auto AQ = getAccessQualifier(AccQual)) 772 Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true); 773 774 // TODO: Emit Arg[".actual_access"]. 775 776 SmallVector<StringRef, 1> SplitTypeQuals; 777 TypeQual.split(SplitTypeQuals, " ", -1, false); 778 for (StringRef Key : SplitTypeQuals) { 779 if (Key == "const") 780 Arg[".is_const"] = Arg.getDocument()->getNode(true); 781 else if (Key == "restrict") 782 Arg[".is_restrict"] = Arg.getDocument()->getNode(true); 783 else if (Key == "volatile") 784 Arg[".is_volatile"] = Arg.getDocument()->getNode(true); 785 else if (Key == "pipe") 786 Arg[".is_pipe"] = Arg.getDocument()->getNode(true); 787 } 788 789 Args.push_back(Arg); 790 } 791 792 void MetadataStreamerV3::emitHiddenKernelArgs(const MachineFunction &MF, 793 unsigned &Offset, 794 msgpack::ArrayDocNode Args) { 795 auto &Func = MF.getFunction(); 796 const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>(); 797 798 unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func); 799 if (!HiddenArgNumBytes) 800 return; 801 802 const Module *M = Func.getParent(); 803 auto &DL = M->getDataLayout(); 804 auto Int64Ty = Type::getInt64Ty(Func.getContext()); 805 806 if (HiddenArgNumBytes >= 8) 807 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset, 808 Args); 809 if (HiddenArgNumBytes >= 16) 810 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset, 811 Args); 812 if (HiddenArgNumBytes >= 24) 813 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset, 814 Args); 815 816 auto Int8PtrTy = 817 Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS); 818 819 // Emit "printf buffer" argument if printf is used, emit "hostcall buffer" 820 // if "hostcall" module flag is set, otherwise emit dummy "none" argument. 821 if (HiddenArgNumBytes >= 32) { 822 if (M->getNamedMetadata("llvm.printf.fmts")) 823 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset, 824 Args); 825 else if (M->getModuleFlag("amdgpu_hostcall")) { 826 // The printf runtime binding pass should have ensured that hostcall and 827 // printf are not used in the same module. 828 assert(!M->getNamedMetadata("llvm.printf.fmts")); 829 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset, 830 Args); 831 } else 832 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args); 833 } 834 835 // Emit "default queue" and "completion action" arguments if enqueue kernel is 836 // used, otherwise emit dummy "none" arguments. 837 if (HiddenArgNumBytes >= 48) { 838 if (Func.hasFnAttribute("calls-enqueue-kernel")) { 839 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset, 840 Args); 841 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_completion_action", Offset, 842 Args); 843 } else { 844 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args); 845 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args); 846 } 847 } 848 849 // Emit the pointer argument for multi-grid object. 850 if (HiddenArgNumBytes >= 56) 851 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset, 852 Args); 853 } 854 855 msgpack::MapDocNode 856 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF, 857 const SIProgramInfo &ProgramInfo) const { 858 const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>(); 859 const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>(); 860 const Function &F = MF.getFunction(); 861 862 auto Kern = HSAMetadataDoc->getMapNode(); 863 864 Align MaxKernArgAlign; 865 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode( 866 STM.getKernArgSegmentSize(F, MaxKernArgAlign)); 867 Kern[".group_segment_fixed_size"] = 868 Kern.getDocument()->getNode(ProgramInfo.LDSSize); 869 Kern[".private_segment_fixed_size"] = 870 Kern.getDocument()->getNode(ProgramInfo.ScratchSize); 871 872 // FIXME: The metadata treats the minimum as 16? 873 Kern[".kernarg_segment_align"] = 874 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value()); 875 Kern[".wavefront_size"] = 876 Kern.getDocument()->getNode(STM.getWavefrontSize()); 877 Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR); 878 Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR); 879 Kern[".max_flat_workgroup_size"] = 880 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize()); 881 Kern[".sgpr_spill_count"] = 882 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs()); 883 Kern[".vgpr_spill_count"] = 884 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs()); 885 886 return Kern; 887 } 888 889 bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) { 890 return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true); 891 } 892 893 void MetadataStreamerV3::begin(const Module &Mod, 894 const IsaInfo::AMDGPUTargetID &TargetID) { 895 emitVersion(); 896 emitPrintf(Mod); 897 getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode(); 898 } 899 900 void MetadataStreamerV3::end() { 901 std::string HSAMetadataString; 902 raw_string_ostream StrOS(HSAMetadataString); 903 HSAMetadataDoc->toYAML(StrOS); 904 905 if (DumpHSAMetadata) 906 dump(StrOS.str()); 907 if (VerifyHSAMetadata) 908 verify(StrOS.str()); 909 } 910 911 void MetadataStreamerV3::emitKernel(const MachineFunction &MF, 912 const SIProgramInfo &ProgramInfo) { 913 auto &Func = MF.getFunction(); 914 auto Kern = getHSAKernelProps(MF, ProgramInfo); 915 916 assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL || 917 Func.getCallingConv() == CallingConv::SPIR_KERNEL); 918 919 auto Kernels = 920 getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true); 921 922 { 923 Kern[".name"] = Kern.getDocument()->getNode(Func.getName()); 924 Kern[".symbol"] = Kern.getDocument()->getNode( 925 (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true); 926 emitKernelLanguage(Func, Kern); 927 emitKernelAttrs(Func, Kern); 928 emitKernelArgs(MF, Kern); 929 } 930 931 Kernels.push_back(Kern); 932 } 933 934 //===----------------------------------------------------------------------===// 935 // HSAMetadataStreamerV4 936 //===----------------------------------------------------------------------===// 937 938 void MetadataStreamerV4::emitVersion() { 939 auto Version = HSAMetadataDoc->getArrayNode(); 940 Version.push_back(Version.getDocument()->getNode(VersionMajorV4)); 941 Version.push_back(Version.getDocument()->getNode(VersionMinorV4)); 942 getRootMetadata("amdhsa.version") = Version; 943 } 944 945 void MetadataStreamerV4::emitTargetID(const IsaInfo::AMDGPUTargetID &TargetID) { 946 getRootMetadata("amdhsa.target") = 947 HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true); 948 } 949 950 void MetadataStreamerV4::begin(const Module &Mod, 951 const IsaInfo::AMDGPUTargetID &TargetID) { 952 emitVersion(); 953 emitTargetID(TargetID); 954 emitPrintf(Mod); 955 getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode(); 956 } 957 958 //===----------------------------------------------------------------------===// 959 // HSAMetadataStreamerV5 960 //===----------------------------------------------------------------------===// 961 962 void MetadataStreamerV5::emitVersion() { 963 auto Version = HSAMetadataDoc->getArrayNode(); 964 Version.push_back(Version.getDocument()->getNode(VersionMajorV5)); 965 Version.push_back(Version.getDocument()->getNode(VersionMinorV5)); 966 getRootMetadata("amdhsa.version") = Version; 967 } 968 969 void MetadataStreamerV5::emitHiddenKernelArgs(const MachineFunction &MF, 970 unsigned &Offset, 971 msgpack::ArrayDocNode Args) { 972 auto &Func = MF.getFunction(); 973 const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>(); 974 const Module *M = Func.getParent(); 975 auto &DL = M->getDataLayout(); 976 977 auto Int64Ty = Type::getInt64Ty(Func.getContext()); 978 auto Int32Ty = Type::getInt32Ty(Func.getContext()); 979 auto Int16Ty = Type::getInt16Ty(Func.getContext()); 980 981 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset, Args); 982 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_y", Offset, Args); 983 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_z", Offset, Args); 984 985 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_x", Offset, Args); 986 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_y", Offset, Args); 987 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_z", Offset, Args); 988 989 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_x", Offset, Args); 990 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_y", Offset, Args); 991 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_z", Offset, Args); 992 993 // Reserved for hidden_tool_correlation_id. 994 Offset += 8; 995 996 Offset += 8; // Reserved. 997 998 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset, Args); 999 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset, Args); 1000 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset, Args); 1001 1002 emitKernelArg(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args); 1003 1004 Offset += 6; // Reserved. 1005 auto Int8PtrTy = 1006 Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS); 1007 1008 if (M->getNamedMetadata("llvm.printf.fmts")) { 1009 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset, 1010 Args); 1011 } else 1012 Offset += 8; // Skipped. 1013 1014 if (M->getModuleFlag("amdgpu_hostcall")) { 1015 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset, 1016 Args); 1017 } else 1018 Offset += 8; // Skipped. 1019 1020 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset, 1021 Args); 1022 1023 // Ignore temporarily until it is implemented. 1024 // emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_heap_v1", Offset, Args); 1025 Offset += 8; 1026 1027 if (Func.hasFnAttribute("calls-enqueue-kernel")) { 1028 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset, 1029 Args); 1030 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_completion_action", Offset, 1031 Args); 1032 } else 1033 Offset += 16; // Skipped. 1034 1035 Offset += 72; // Reserved. 1036 1037 // hidden_private_base and hidden_shared_base are only used by GFX8. 1038 if (ST.getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS) { 1039 emitKernelArg(DL, Int32Ty, Align(4), "hidden_private_base", Offset, Args); 1040 emitKernelArg(DL, Int32Ty, Align(4), "hidden_shared_base", Offset, Args); 1041 } else 1042 Offset += 8; // Skipped. 1043 1044 const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>(); 1045 if (MFI.hasQueuePtr()) 1046 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_queue_ptr", Offset, Args); 1047 } 1048 1049 } // end namespace HSAMD 1050 } // end namespace AMDGPU 1051 } // end namespace llvm 1052