Lines Matching +full:node +full:- +full:version
1 //===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===//
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7 //===----------------------------------------------------------------------===//
13 //===----------------------------------------------------------------------===//
44 "amdgpu-dump-hsa-metadata",
47 "amdgpu-verify-hsa-metadata",
52 //===----------------------------------------------------------------------===//
54 //===----------------------------------------------------------------------===//
132 ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
140 switch (Ty->getTypeID()) {
145 auto BitWidth = Ty->getIntegerBitWidth();
167 auto *ElTy = VecTy->getElementType();
168 auto NumElements = VecTy->getNumElements();
177 MetadataStreamerMsgPackV4::getWorkGroupDimensions(MDNode *Node) const {
178 auto Dims = HSAMetadataDoc->getArrayNode();
179 if (Node->getNumOperands() != 3)
182 for (auto &Op : Node->operands())
183 Dims.push_back(Dims.getDocument()->getNode(
184 uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
189 auto Version = HSAMetadataDoc->getArrayNode();
190 Version.push_back(Version.getDocument()->getNode(VersionMajorV4));
191 Version.push_back(Version.getDocument()->getNode(VersionMinorV4));
192 getRootMetadata("amdhsa.version") = Version;
198 HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true);
202 auto *Node = Mod.getNamedMetadata("llvm.printf.fmts");
203 if (!Node)
206 auto Printf = HSAMetadataDoc->getArrayNode();
207 for (auto *Op : Node->operands())
208 if (Op->getNumOperands())
209 Printf.push_back(Printf.getDocument()->getNode(
210 cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
217 auto *Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
218 if (!Node || !Node->getNumOperands())
220 auto *Op0 = Node->getOperand(0);
221 if (Op0->getNumOperands() <= 1)
224 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
225 auto LanguageVersion = Kern.getDocument()->getArrayNode();
226 LanguageVersion.push_back(Kern.getDocument()->getNode(
227 mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
228 LanguageVersion.push_back(Kern.getDocument()->getNode(
229 mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
236 if (auto *Node = Func.getMetadata("reqd_work_group_size"))
237 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
238 if (auto *Node = Func.getMetadata("work_group_size_hint"))
239 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
240 if (auto *Node = Func.getMetadata("vec_type_hint")) {
241 Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
243 cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
244 mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
247 if (Func.hasFnAttribute("runtime-handle")) {
248 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
249 Func.getFnAttribute("runtime-handle").getValueAsString().str(),
252 if (Func.hasFnAttribute("device-init"))
253 Kern[".kind"] = Kern.getDocument()->getNode("init");
254 else if (Func.hasFnAttribute("device-fini"))
255 Kern[".kind"] = Kern.getDocument()->getNode("fini");
262 auto Args = HSAMetadataDoc->getArrayNode();
264 if (Arg.hasAttribute("amdgpu-hidden-argument"))
280 const MDNode *Node;
283 Node = Func->getMetadata("kernel_arg_name");
284 if (Node && ArgNo < Node->getNumOperands())
285 Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
290 Node = Func->getMetadata("kernel_arg_type");
291 if (Node && ArgNo < Node->getNumOperands())
292 TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
295 Node = Func->getMetadata("kernel_arg_base_type");
296 if (Node && ArgNo < Node->getNumOperands())
297 BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
301 if (Arg.getType()->isPointerTy() && Arg.hasNoAliasAttr()) {
309 Node = Func->getMetadata("kernel_arg_access_qual");
310 if (Node && ArgNo < Node->getNumOperands())
311 AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
314 Node = Func->getMetadata("kernel_arg_type_qual");
315 if (Node && ArgNo < Node->getNumOperands())
316 TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
318 const DataLayout &DL = Func->getDataLayout();
325 if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
345 auto Arg = Args.getDocument()->getMapNode();
348 Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
350 Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
352 Arg[".size"] = Arg.getDocument()->getNode(Size);
354 Arg[".offset"] = Arg.getDocument()->getNode(Offset);
356 Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
358 Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
361 if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
364 Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier,
368 Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
371 Arg[".actual_access"] = Arg.getDocument()->getNode(*AAQ, /*Copy=*/true);
374 TypeQual.split(SplitTypeQuals, " ", -1, false);
377 Arg[".is_const"] = Arg.getDocument()->getNode(true);
379 Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
381 Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
383 Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
399 auto &DL = M->getDataLayout();
421 if (M->getNamedMetadata("llvm.printf.fmts"))
424 else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr"))
434 if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
443 if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
451 // Emit the pointer argument for multi-grid object.
453 if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
470 auto Kern = HSAMetadataDoc->getMapNode();
473 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
476 Kern.getDocument()->getNode(ProgramInfo.LDSSize);
477 DelayedExprs->assignDocNode(Kern[".private_segment_fixed_size"],
480 DelayedExprs->assignDocNode(Kern[".uses_dynamic_stack"],
487 Kern.getDocument()->getNode(ProgramInfo.WgpMode);
491 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
493 Kern.getDocument()->getNode(STM.getWavefrontSize());
494 DelayedExprs->assignDocNode(Kern[".sgpr_count"], msgpack::Type::UInt,
496 DelayedExprs->assignDocNode(Kern[".vgpr_count"], msgpack::Type::UInt,
501 DelayedExprs->assignDocNode(Kern[".agpr_count"], msgpack::Type::UInt,
506 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
514 Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
517 Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
520 Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
523 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
525 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
531 DelayedExprs->resolveDelayedExpressions();
540 getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
541 DelayedExprs->clear();
545 DelayedExprs->resolveDelayedExpressions();
548 HSAMetadataDoc->toYAML(StrOS);
571 Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
572 Kern[".symbol"] = Kern.getDocument()->getNode(
582 //===----------------------------------------------------------------------===//
584 //===----------------------------------------------------------------------===//
587 auto Version = HSAMetadataDoc->getArrayNode();
588 Version.push_back(Version.getDocument()->getNode(VersionMajorV5));
589 Version.push_back(Version.getDocument()->getNode(VersionMinorV5));
590 getRootMetadata("amdhsa.version") = Version;
603 auto &DL = M->getDataLayout();
638 if (M->getNamedMetadata("llvm.printf.fmts")) {
645 if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) {
652 if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
659 if (!Func.hasFnAttribute("amdgpu-no-heap-ptr"))
664 if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
671 if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
705 if (Func.getFnAttribute("uniform-work-group-size").getValueAsBool())
706 Kern[".uniform_work_group_size"] = Kern.getDocument()->getNode(1);
709 //===----------------------------------------------------------------------===//
711 //===----------------------------------------------------------------------===//
714 auto Version = HSAMetadataDoc->getArrayNode();
715 Version.push_back(Version.getDocument()->getNode(VersionMajorV6));
716 Version.push_back(Version.getDocument()->getNode(VersionMinorV6));
717 getRootMetadata("amdhsa.version") = Version;