LLVM 19.0.0git
AMDGPUHSAMetadataStreamer.cpp
Go to the documentation of this file.
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
16#include "AMDGPU.h"
17#include "GCNSubtarget.h"
20#include "SIProgramInfo.h"
21#include "llvm/IR/Module.h"
22#include "llvm/MC/MCContext.h"
23#include "llvm/MC/MCExpr.h"
24using namespace llvm;
25
26static std::pair<Type *, Align> getArgumentTypeAlign(const Argument &Arg,
27 const DataLayout &DL) {
28 Type *Ty = Arg.getType();
29 MaybeAlign ArgAlign;
30 if (Arg.hasByRefAttr()) {
31 Ty = Arg.getParamByRefType();
32 ArgAlign = Arg.getParamAlign();
33 }
34
35 if (!ArgAlign)
36 ArgAlign = DL.getABITypeAlign(Ty);
37
38 return std::pair(Ty, *ArgAlign);
39}
40
41namespace llvm {
42
44 "amdgpu-dump-hsa-metadata",
45 cl::desc("Dump AMDGPU HSA Metadata"));
47 "amdgpu-verify-hsa-metadata",
48 cl::desc("Verify AMDGPU HSA Metadata"));
49
50namespace AMDGPU {
51namespace HSAMD {
52
53//===----------------------------------------------------------------------===//
54// HSAMetadataStreamerV4
55//===----------------------------------------------------------------------===//
56
57void MetadataStreamerMsgPackV4::dump(StringRef HSAMetadataString) const {
58 errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
59}
60
61void MetadataStreamerMsgPackV4::verify(StringRef HSAMetadataString) const {
62 errs() << "AMDGPU HSA Metadata Parser Test: ";
63
64 msgpack::Document FromHSAMetadataString;
65
66 if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
67 errs() << "FAIL\n";
68 return;
69 }
70
71 std::string ToHSAMetadataString;
72 raw_string_ostream StrOS(ToHSAMetadataString);
73 FromHSAMetadataString.toYAML(StrOS);
74
75 errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
76 if (HSAMetadataString != ToHSAMetadataString) {
77 errs() << "Original input: " << HSAMetadataString << '\n'
78 << "Produced output: " << StrOS.str() << '\n';
79 }
80}
81
82std::optional<StringRef>
85 .Case("read_only", StringRef("read_only"))
86 .Case("write_only", StringRef("write_only"))
87 .Case("read_write", StringRef("read_write"))
88 .Default(std::nullopt);
89}
90
92 unsigned AddressSpace) const {
93 switch (AddressSpace) {
95 return StringRef("private");
97 return StringRef("global");
99 return StringRef("constant");
101 return StringRef("local");
103 return StringRef("generic");
105 return StringRef("region");
106 default:
107 return std::nullopt;
108 }
109}
110
113 StringRef BaseTypeName) const {
114 if (TypeQual.contains("pipe"))
115 return "pipe";
116
117 return StringSwitch<StringRef>(BaseTypeName)
118 .Case("image1d_t", "image")
119 .Case("image1d_array_t", "image")
120 .Case("image1d_buffer_t", "image")
121 .Case("image2d_t", "image")
122 .Case("image2d_array_t", "image")
123 .Case("image2d_array_depth_t", "image")
124 .Case("image2d_array_msaa_t", "image")
125 .Case("image2d_array_msaa_depth_t", "image")
126 .Case("image2d_depth_t", "image")
127 .Case("image2d_msaa_t", "image")
128 .Case("image2d_msaa_depth_t", "image")
129 .Case("image3d_t", "image")
130 .Case("sampler_t", "sampler")
131 .Case("queue_t", "queue")
132 .Default(isa<PointerType>(Ty)
134 ? "dynamic_shared_pointer"
135 : "global_buffer")
136 : "by_value");
137}
138
140 bool Signed) const {
141 switch (Ty->getTypeID()) {
142 case Type::IntegerTyID: {
143 if (!Signed)
144 return (Twine('u') + getTypeName(Ty, true)).str();
145
146 auto BitWidth = Ty->getIntegerBitWidth();
147 switch (BitWidth) {
148 case 8:
149 return "char";
150 case 16:
151 return "short";
152 case 32:
153 return "int";
154 case 64:
155 return "long";
156 default:
157 return (Twine('i') + Twine(BitWidth)).str();
158 }
159 }
160 case Type::HalfTyID:
161 return "half";
162 case Type::FloatTyID:
163 return "float";
164 case Type::DoubleTyID:
165 return "double";
167 auto VecTy = cast<FixedVectorType>(Ty);
168 auto ElTy = VecTy->getElementType();
169 auto NumElements = VecTy->getNumElements();
170 return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
171 }
172 default:
173 return "unknown";
174 }
175}
176
179 auto Dims = HSAMetadataDoc->getArrayNode();
180 if (Node->getNumOperands() != 3)
181 return Dims;
182
183 for (auto &Op : Node->operands())
184 Dims.push_back(Dims.getDocument()->getNode(
185 uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
186 return Dims;
187}
188
190 auto Version = HSAMetadataDoc->getArrayNode();
191 Version.push_back(Version.getDocument()->getNode(VersionMajorV4));
192 Version.push_back(Version.getDocument()->getNode(VersionMinorV4));
193 getRootMetadata("amdhsa.version") = Version;
194}
195
197 const IsaInfo::AMDGPUTargetID &TargetID) {
198 getRootMetadata("amdhsa.target") =
199 HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true);
200}
201
203 auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
204 if (!Node)
205 return;
206
207 auto Printf = HSAMetadataDoc->getArrayNode();
208 for (auto *Op : Node->operands())
209 if (Op->getNumOperands())
210 Printf.push_back(Printf.getDocument()->getNode(
211 cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
212 getRootMetadata("amdhsa.printf") = Printf;
213}
214
216 msgpack::MapDocNode Kern) {
217 // TODO: What about other languages?
218 auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
219 if (!Node || !Node->getNumOperands())
220 return;
221 auto Op0 = Node->getOperand(0);
222 if (Op0->getNumOperands() <= 1)
223 return;
224
225 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
227 LanguageVersion.push_back(Kern.getDocument()->getNode(
228 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
229 LanguageVersion.push_back(Kern.getDocument()->getNode(
230 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
231 Kern[".language_version"] = LanguageVersion;
232}
233
235 msgpack::MapDocNode Kern) {
236
237 if (auto Node = Func.getMetadata("reqd_work_group_size"))
238 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
239 if (auto Node = Func.getMetadata("work_group_size_hint"))
240 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
241 if (auto Node = Func.getMetadata("vec_type_hint")) {
242 Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
244 cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
245 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
246 /*Copy=*/true);
247 }
248 if (Func.hasFnAttribute("runtime-handle")) {
249 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
250 Func.getFnAttribute("runtime-handle").getValueAsString().str(),
251 /*Copy=*/true);
252 }
253 if (Func.hasFnAttribute("device-init"))
254 Kern[".kind"] = Kern.getDocument()->getNode("init");
255 else if (Func.hasFnAttribute("device-fini"))
256 Kern[".kind"] = Kern.getDocument()->getNode("fini");
257}
258
260 msgpack::MapDocNode Kern) {
261 auto &Func = MF.getFunction();
262 unsigned Offset = 0;
263 auto Args = HSAMetadataDoc->getArrayNode();
264 for (auto &Arg : Func.args())
265 emitKernelArg(Arg, Offset, Args);
266
267 emitHiddenKernelArgs(MF, Offset, Args);
268
269 Kern[".args"] = Args;
270}
271
273 unsigned &Offset,
275 auto Func = Arg.getParent();
276 auto ArgNo = Arg.getArgNo();
277 const MDNode *Node;
278
280 Node = Func->getMetadata("kernel_arg_name");
281 if (Node && ArgNo < Node->getNumOperands())
282 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
283 else if (Arg.hasName())
284 Name = Arg.getName();
285
286 StringRef TypeName;
287 Node = Func->getMetadata("kernel_arg_type");
288 if (Node && ArgNo < Node->getNumOperands())
289 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
290
291 StringRef BaseTypeName;
292 Node = Func->getMetadata("kernel_arg_base_type");
293 if (Node && ArgNo < Node->getNumOperands())
294 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
295
296 StringRef ActAccQual;
297 // Do we really need NoAlias check here?
298 if (Arg.getType()->isPointerTy() && Arg.hasNoAliasAttr()) {
299 if (Arg.onlyReadsMemory())
300 ActAccQual = "read_only";
301 else if (Arg.hasAttribute(Attribute::WriteOnly))
302 ActAccQual = "write_only";
303 }
304
305 StringRef AccQual;
306 Node = Func->getMetadata("kernel_arg_access_qual");
307 if (Node && ArgNo < Node->getNumOperands())
308 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
309
310 StringRef TypeQual;
311 Node = Func->getMetadata("kernel_arg_type_qual");
312 if (Node && ArgNo < Node->getNumOperands())
313 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
314
315 const DataLayout &DL = Func->getParent()->getDataLayout();
316
317 MaybeAlign PointeeAlign;
318 Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
319
320 // FIXME: Need to distinguish in memory alignment from pointer alignment.
321 if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
322 if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
323 PointeeAlign = Arg.getParamAlign().valueOrOne();
324 }
325
326 // There's no distinction between byval aggregates and raw aggregates.
327 Type *ArgTy;
328 Align ArgAlign;
329 std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
330
331 emitKernelArg(DL, ArgTy, ArgAlign,
332 getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args,
333 PointeeAlign, Name, TypeName, BaseTypeName, ActAccQual,
334 AccQual, TypeQual);
335}
336
338 const DataLayout &DL, Type *Ty, Align Alignment, StringRef ValueKind,
339 unsigned &Offset, msgpack::ArrayDocNode Args, MaybeAlign PointeeAlign,
340 StringRef Name, StringRef TypeName, StringRef BaseTypeName,
341 StringRef ActAccQual, StringRef AccQual, StringRef TypeQual) {
342 auto Arg = Args.getDocument()->getMapNode();
343
344 if (!Name.empty())
345 Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
346 if (!TypeName.empty())
347 Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
348 auto Size = DL.getTypeAllocSize(Ty);
349 Arg[".size"] = Arg.getDocument()->getNode(Size);
350 Offset = alignTo(Offset, Alignment);
351 Arg[".offset"] = Arg.getDocument()->getNode(Offset);
352 Offset += Size;
353 Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
354 if (PointeeAlign)
355 Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
356
357 if (auto PtrTy = dyn_cast<PointerType>(Ty))
358 if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
359 // Limiting address space to emit only for a certain ValueKind.
360 if (ValueKind == "global_buffer" || ValueKind == "dynamic_shared_pointer")
361 Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier,
362 /*Copy=*/true);
363
364 if (auto AQ = getAccessQualifier(AccQual))
365 Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
366
367 if (auto AAQ = getAccessQualifier(ActAccQual))
368 Arg[".actual_access"] = Arg.getDocument()->getNode(*AAQ, /*Copy=*/true);
369
370 SmallVector<StringRef, 1> SplitTypeQuals;
371 TypeQual.split(SplitTypeQuals, " ", -1, false);
372 for (StringRef Key : SplitTypeQuals) {
373 if (Key == "const")
374 Arg[".is_const"] = Arg.getDocument()->getNode(true);
375 else if (Key == "restrict")
376 Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
377 else if (Key == "volatile")
378 Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
379 else if (Key == "pipe")
380 Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
381 }
382
383 Args.push_back(Arg);
384}
385
387 const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
388 auto &Func = MF.getFunction();
389 const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
390
391 unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
392 if (!HiddenArgNumBytes)
393 return;
394
395 const Module *M = Func.getParent();
396 auto &DL = M->getDataLayout();
397 auto Int64Ty = Type::getInt64Ty(Func.getContext());
398
399 Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
400
401 if (HiddenArgNumBytes >= 8)
402 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset,
403 Args);
404 if (HiddenArgNumBytes >= 16)
405 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
406 Args);
407 if (HiddenArgNumBytes >= 24)
408 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
409 Args);
410
411 auto Int8PtrTy =
412 PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
413
414 if (HiddenArgNumBytes >= 32) {
415 // We forbid the use of features requiring hostcall when compiling OpenCL
416 // before code object V5, which makes the mutual exclusion between the
417 // "printf buffer" and "hostcall buffer" here sound.
418 if (M->getNamedMetadata("llvm.printf.fmts"))
419 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
420 Args);
421 else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr"))
422 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
423 Args);
424 else
425 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
426 }
427
428 // Emit "default queue" and "completion action" arguments if enqueue kernel is
429 // used, otherwise emit dummy "none" arguments.
430 if (HiddenArgNumBytes >= 40) {
431 if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
432 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
433 Args);
434 } else {
435 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
436 }
437 }
438
439 if (HiddenArgNumBytes >= 48) {
440 if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
441 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_completion_action", Offset,
442 Args);
443 } else {
444 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
445 }
446 }
447
448 // Emit the pointer argument for multi-grid object.
449 if (HiddenArgNumBytes >= 56) {
450 if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
451 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset,
452 Args);
453 } else {
454 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
455 }
456 }
457}
458
461 const SIProgramInfo &ProgramInfo,
462 unsigned CodeObjectVersion) const {
463 const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
465 const Function &F = MF.getFunction();
466
467 auto GetMCExprValue = [&MF](const MCExpr *Value) {
468 int64_t Val;
469 if (!Value->evaluateAsAbsolute(Val)) {
470 MCContext &Ctx = MF.getContext();
471 Ctx.reportError(SMLoc(), "could not resolve expression when required.");
472 Val = 0;
473 }
474 return static_cast<uint64_t>(Val);
475 };
476
477 auto Kern = HSAMetadataDoc->getMapNode();
478
479 Align MaxKernArgAlign;
480 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
481 STM.getKernArgSegmentSize(F, MaxKernArgAlign));
482 Kern[".group_segment_fixed_size"] =
483 Kern.getDocument()->getNode(ProgramInfo.LDSSize);
484 Kern[".private_segment_fixed_size"] =
485 Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.ScratchSize));
486 if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5) {
487 Kern[".uses_dynamic_stack"] = Kern.getDocument()->getNode(
488 static_cast<bool>(GetMCExprValue(ProgramInfo.DynamicCallStack)));
489 }
490
491 if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5 && STM.supportsWGP())
492 Kern[".workgroup_processor_mode"] =
493 Kern.getDocument()->getNode(ProgramInfo.WgpMode);
494
495 // FIXME: The metadata treats the minimum as 16?
496 Kern[".kernarg_segment_align"] =
497 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
498 Kern[".wavefront_size"] =
499 Kern.getDocument()->getNode(STM.getWavefrontSize());
500 Kern[".sgpr_count"] =
501 Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumSGPR));
502 Kern[".vgpr_count"] =
503 Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumVGPR));
504
505 // Only add AGPR count to metadata for supported devices
506 if (STM.hasMAIInsts()) {
507 Kern[".agpr_count"] =
508 Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumAccVGPR));
509 }
510
511 Kern[".max_flat_workgroup_size"] =
512 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
513 unsigned NumWGX = MFI.getMaxNumWorkGroupsX();
514 unsigned NumWGY = MFI.getMaxNumWorkGroupsY();
515 unsigned NumWGZ = MFI.getMaxNumWorkGroupsZ();
516 if (NumWGX != 0 && NumWGY != 0 && NumWGZ != 0) {
517 Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
518 Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
519 Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
520 }
521 Kern[".sgpr_spill_count"] =
522 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
523 Kern[".vgpr_spill_count"] =
524 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
525
526 return Kern;
527}
528
530 return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
531}
532
534 const IsaInfo::AMDGPUTargetID &TargetID) {
535 emitVersion();
536 emitTargetID(TargetID);
538 getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
539}
540
542 std::string HSAMetadataString;
543 raw_string_ostream StrOS(HSAMetadataString);
544 HSAMetadataDoc->toYAML(StrOS);
545
546 if (DumpHSAMetadata)
547 dump(StrOS.str());
549 verify(StrOS.str());
550}
551
553 const SIProgramInfo &ProgramInfo) {
554 auto &Func = MF.getFunction();
555 if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL &&
556 Func.getCallingConv() != CallingConv::SPIR_KERNEL)
557 return;
558
559 auto CodeObjectVersion =
560 AMDGPU::getAMDHSACodeObjectVersion(*Func.getParent());
561 auto Kern = getHSAKernelProps(MF, ProgramInfo, CodeObjectVersion);
562
563 auto Kernels =
564 getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
565
566 {
567 Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
568 Kern[".symbol"] = Kern.getDocument()->getNode(
569 (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
570 emitKernelLanguage(Func, Kern);
571 emitKernelAttrs(Func, Kern);
572 emitKernelArgs(MF, Kern);
573 }
574
575 Kernels.push_back(Kern);
576}
577
578//===----------------------------------------------------------------------===//
579// HSAMetadataStreamerV5
580//===----------------------------------------------------------------------===//
581
583 auto Version = HSAMetadataDoc->getArrayNode();
584 Version.push_back(Version.getDocument()->getNode(VersionMajorV5));
585 Version.push_back(Version.getDocument()->getNode(VersionMinorV5));
586 getRootMetadata("amdhsa.version") = Version;
587}
588
590 const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
591 auto &Func = MF.getFunction();
592 const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
593
594 // No implicit kernel argument is used.
595 if (ST.getImplicitArgNumBytes(Func) == 0)
596 return;
597
598 const Module *M = Func.getParent();
599 auto &DL = M->getDataLayout();
601
602 auto Int64Ty = Type::getInt64Ty(Func.getContext());
603 auto Int32Ty = Type::getInt32Ty(Func.getContext());
604 auto Int16Ty = Type::getInt16Ty(Func.getContext());
605
606 Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
607 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset, Args);
608 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_y", Offset, Args);
609 emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_z", Offset, Args);
610
611 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_x", Offset, Args);
612 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_y", Offset, Args);
613 emitKernelArg(DL, Int16Ty, Align(2), "hidden_group_size_z", Offset, Args);
614
615 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_x", Offset, Args);
616 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_y", Offset, Args);
617 emitKernelArg(DL, Int16Ty, Align(2), "hidden_remainder_z", Offset, Args);
618
619 // Reserved for hidden_tool_correlation_id.
620 Offset += 8;
621
622 Offset += 8; // Reserved.
623
624 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset, Args);
625 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset, Args);
626 emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset, Args);
627
628 emitKernelArg(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args);
629
630 Offset += 6; // Reserved.
631 auto Int8PtrTy =
632 PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
633
634 if (M->getNamedMetadata("llvm.printf.fmts")) {
635 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
636 Args);
637 } else {
638 Offset += 8; // Skipped.
639 }
640
641 if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) {
642 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
643 Args);
644 } else {
645 Offset += 8; // Skipped.
646 }
647
648 if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
649 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset,
650 Args);
651 } else {
652 Offset += 8; // Skipped.
653 }
654
655 if (!Func.hasFnAttribute("amdgpu-no-heap-ptr"))
656 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_heap_v1", Offset, Args);
657 else
658 Offset += 8; // Skipped.
659
660 if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
661 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
662 Args);
663 } else {
664 Offset += 8; // Skipped.
665 }
666
667 if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
668 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_completion_action", Offset,
669 Args);
670 } else {
671 Offset += 8; // Skipped.
672 }
673
674 // Emit argument for hidden dynamic lds size
675 if (MFI.isDynamicLDSUsed()) {
676 emitKernelArg(DL, Int32Ty, Align(4), "hidden_dynamic_lds_size", Offset,
677 Args);
678 } else {
679 Offset += 4; // skipped
680 }
681
682 Offset += 68; // Reserved.
683
684 // hidden_private_base and hidden_shared_base are only when the subtarget has
685 // ApertureRegs.
686 if (!ST.hasApertureRegs()) {
687 emitKernelArg(DL, Int32Ty, Align(4), "hidden_private_base", Offset, Args);
688 emitKernelArg(DL, Int32Ty, Align(4), "hidden_shared_base", Offset, Args);
689 } else {
690 Offset += 8; // Skipped.
691 }
692
693 if (MFI.getUserSGPRInfo().hasQueuePtr())
694 emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_queue_ptr", Offset, Args);
695}
696
698 msgpack::MapDocNode Kern) {
700
701 if (Func.getFnAttribute("uniform-work-group-size").getValueAsBool())
702 Kern[".uniform_work_group_size"] = Kern.getDocument()->getNode(1);
703}
704
705//===----------------------------------------------------------------------===//
706// HSAMetadataStreamerV6
707//===----------------------------------------------------------------------===//
708
710 auto Version = HSAMetadataDoc->getArrayNode();
711 Version.push_back(Version.getDocument()->getNode(VersionMajorV6));
712 Version.push_back(Version.getDocument()->getNode(VersionMinorV6));
713 getRootMetadata("amdhsa.version") = Version;
714}
715
716} // end namespace HSAMD
717} // end namespace AMDGPU
718} // end namespace llvm
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static std::pair< Type *, Align > getArgumentTypeAlign(const Argument &Arg, const DataLayout &DL)
AMDGPU HSA Metadata Streamer.
Given that RA is a live value
std::string Name
uint64_t Size
AMD GCN specific subclass of TargetSubtarget.
#define F(x, y, z)
Definition: MD5.cpp:55
Module.h This file contains the declarations for the Module class.
ppc ctr loops verify
Defines struct to track resource usage and hardware flags for kernels and entry functions.
unsigned getKernArgSegmentSize(const Function &F, Align &MaxAlign) const
unsigned getWavefrontSize() const
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict)
Emit HSA Metadata.
std::optional< StringRef > getAddressSpaceQualifier(unsigned AddressSpace) const
msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const
std::optional< StringRef > getAccessQualifier(StringRef AccQual) const
void emitKernelLanguage(const Function &Func, msgpack::MapDocNode Kern)
msgpack::MapDocNode getHSAKernelProps(const MachineFunction &MF, const SIProgramInfo &ProgramInfo, unsigned CodeObjectVersion) const
void emitKernelAttrs(const Function &Func, msgpack::MapDocNode Kern) override
std::unique_ptr< msgpack::Document > HSAMetadataDoc
bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override
void emitKernelArgs(const MachineFunction &MF, msgpack::MapDocNode Kern)
std::string getTypeName(Type *Ty, bool Signed) const
void dump(StringRef HSAMetadataString) const
void begin(const Module &Mod, const IsaInfo::AMDGPUTargetID &TargetID) override
void emitKernelArg(const Argument &Arg, unsigned &Offset, msgpack::ArrayDocNode Args)
void emitTargetID(const IsaInfo::AMDGPUTargetID &TargetID)
void verify(StringRef HSAMetadataString) const
StringRef getValueKind(Type *Ty, StringRef TypeQual, StringRef BaseTypeName) const
void emitHiddenKernelArgs(const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) override
void emitKernel(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) override
void emitHiddenKernelArgs(const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) override
void emitKernelAttrs(const Function &Func, msgpack::MapDocNode Kern) override
This class represents an incoming formal argument to a Function.
Definition: Argument.h:31
Type * getParamByRefType() const
If this is a byref argument, return its type.
Definition: Function.cpp:242
bool hasNoAliasAttr() const
Return true if this argument has the noalias attribute.
Definition: Function.cpp:280
bool hasByRefAttr() const
Return true if this argument has the byref attribute.
Definition: Function.cpp:146
bool onlyReadsMemory() const
Return true if this argument has the readonly or readnone attribute.
Definition: Function.cpp:316
bool hasAttribute(Attribute::AttrKind Kind) const
Check if an argument has a given attribute.
Definition: Function.cpp:346
const Function * getParent() const
Definition: Argument.h:43
unsigned getArgNo() const
Return the index of this formal argument in its containing function.
Definition: Argument.h:49
MaybeAlign getParamAlign() const
If this is a byval or inalloca argument, return its alignment.
Definition: Function.cpp:223
This class represents an Operation in the Expression.
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:110
bool hasMAIInsts() const
Definition: GCNSubtarget.h:809
bool supportsWGP() const
Definition: GCNSubtarget.h:346
Context object for machine code objects.
Definition: MCContext.h:82
void reportError(SMLoc L, const Twine &Msg)
Definition: MCContext.cpp:1093
Base class for the full range of assembler expressions which are needed for parsing.
Definition: MCExpr.h:35
Metadata node.
Definition: Metadata.h:1067
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
MCContext & getContext() const
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.
Definition: Module.h:65
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()
Represents a location in source code.
Definition: SMLoc.h:23
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1209
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:50
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition: StringRef.h:693
bool contains(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition: StringRef.h:417
A switch()-like statement whose cases are string literals.
Definition: StringSwitch.h:44
StringSwitch & Case(StringLiteral S, T Value)
Definition: StringSwitch.h:69
R Default(T Value)
Definition: StringSwitch.h:182
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
Definition: Twine.h:81
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
unsigned getIntegerBitWidth() const
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:255
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
@ HalfTyID
16-bit floating point type
Definition: Type.h:56
@ FloatTyID
32-bit floating point type
Definition: Type.h:58
@ IntegerTyID
Arbitrary bit width integers.
Definition: Type.h:71
@ FixedVectorTyID
Fixed width SIMD vector type.
Definition: Type.h:76
@ DoubleTyID
64-bit floating point type
Definition: Type.h:59
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.
Definition: Type.h:137
LLVM Value Representation.
Definition: Value.h:74
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:255
bool hasName() const
Definition: Value.h:261
StringRef getName() const
Return a constant reference to the value's name.
Definition: Value.cpp:309
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 DocNode that is a map.
A raw_ostream that writes to an std::string.
Definition: raw_ostream.h:661
std::string & str()
Returns the string's reference.
Definition: raw_ostream.h:679
unsigned LanguageVersion(SourceLanguage L)
Definition: Dwarf.cpp:412
@ 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.
ValueKind
Value kinds.
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.
Definition: CallingConv.h:200
@ SPIR_KERNEL
Used for SPIR kernel functions.
Definition: CallingConv.h:144
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
@ Offset
Definition: DWP.cpp:456
static cl::opt< bool > VerifyHSAMetadata("amdgpu-verify-hsa-metadata", cl::desc("Verify AMDGPU HSA Metadata"))
AddressSpace
Definition: NVPTXBaseInfo.h:21
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.
Definition: TypeName.h:27
@ 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.
Definition: Alignment.h:155
constexpr unsigned BitWidth
Definition: BitmaskEnum.h:191
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Definition: Alignment.h:117
Align valueOrOne() const
For convenience, returns a valid alignment or 1 if undefined.
Definition: Alignment.h:141
Track resource usage for kernels / entry functions.
Definition: SIProgramInfo.h:31
const MCExpr * NumSGPR
Definition: SIProgramInfo.h:70
const MCExpr * NumAccVGPR
Definition: SIProgramInfo.h:67
const MCExpr * DynamicCallStack
Definition: SIProgramInfo.h:87
const MCExpr * NumVGPR
Definition: SIProgramInfo.h:65
const MCExpr * ScratchSize
Definition: SIProgramInfo.h:44