mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-16 11:36:46 +00:00
62e5168889
Tracks the registers that explicit and hidden arguments are preloaded to with new code object metadata. IR arguments may be split across multiple parts by isel, and SGPR tuple alignment means that an argument may be spread across multiple registers. To support this, some of the utilities for hidden kernel arguments are moved to `AMDGPUArgumentUsageInfo.h`. Additional bookkeeping is also needed for tracking purposes.
1003 lines
36 KiB
C++
1003 lines
36 KiB
C++
//===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
//
|
|
/// \file
|
|
/// AMDGPU HSA Metadata Streamer.
|
|
///
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#include "AMDGPUHSAMetadataStreamer.h"
|
|
#include "AMDGPU.h"
|
|
#include "GCNSubtarget.h"
|
|
#include "MCTargetDesc/AMDGPUInstPrinter.h"
|
|
#include "MCTargetDesc/AMDGPUTargetStreamer.h"
|
|
#include "SIMachineFunctionInfo.h"
|
|
#include "SIProgramInfo.h"
|
|
#include "llvm/IR/Module.h"
|
|
#include "llvm/MC/MCContext.h"
|
|
#include "llvm/MC/MCExpr.h"
|
|
#include "llvm/Target/TargetLoweringObjectFile.h"
|
|
|
|
using namespace llvm;
|
|
|
|
static std::pair<Type *, Align> getArgumentTypeAlign(const Argument &Arg,
|
|
const DataLayout &DL) {
|
|
Type *Ty = Arg.getType();
|
|
MaybeAlign ArgAlign;
|
|
if (Arg.hasByRefAttr()) {
|
|
Ty = Arg.getParamByRefType();
|
|
ArgAlign = Arg.getParamAlign();
|
|
}
|
|
|
|
if (!ArgAlign)
|
|
ArgAlign = DL.getABITypeAlign(Ty);
|
|
|
|
return std::pair(Ty, *ArgAlign);
|
|
}
|
|
|
|
/// Find the mangled symbol name for the runtime handle for \p EnqueuedBlock
|
|
static std::string getEnqueuedBlockSymbolName(const AMDGPUTargetMachine &TM,
|
|
const Function &EnqueuedBlock) {
|
|
const MDNode *Associated =
|
|
EnqueuedBlock.getMetadata(LLVMContext::MD_associated);
|
|
if (!Associated)
|
|
return "";
|
|
|
|
auto *VM = cast<ValueAsMetadata>(Associated->getOperand(0));
|
|
auto *RuntimeHandle =
|
|
dyn_cast<GlobalVariable>(VM->getValue()->stripPointerCasts());
|
|
if (!RuntimeHandle ||
|
|
RuntimeHandle->getSection() != ".amdgpu.kernel.runtime.handle")
|
|
return "";
|
|
|
|
SmallString<128> Name;
|
|
TM.getNameWithPrefix(Name, RuntimeHandle,
|
|
TM.getObjFileLowering()->getMangler());
|
|
return Name.str().str();
|
|
}
|
|
|
|
namespace llvm {
|
|
|
|
static cl::opt<bool> DumpHSAMetadata(
|
|
"amdgpu-dump-hsa-metadata",
|
|
cl::desc("Dump AMDGPU HSA Metadata"));
|
|
static cl::opt<bool> VerifyHSAMetadata(
|
|
"amdgpu-verify-hsa-metadata",
|
|
cl::desc("Verify AMDGPU HSA Metadata"));
|
|
|
|
namespace AMDGPU::HSAMD {
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
// HSAMetadataStreamerV4
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
void MetadataStreamerMsgPackV4::dump(StringRef HSAMetadataString) const {
|
|
errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::verify(StringRef HSAMetadataString) const {
|
|
errs() << "AMDGPU HSA Metadata Parser Test: ";
|
|
|
|
msgpack::Document FromHSAMetadataString;
|
|
|
|
if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
|
|
errs() << "FAIL\n";
|
|
return;
|
|
}
|
|
|
|
std::string ToHSAMetadataString;
|
|
raw_string_ostream StrOS(ToHSAMetadataString);
|
|
FromHSAMetadataString.toYAML(StrOS);
|
|
|
|
errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
|
|
if (HSAMetadataString != ToHSAMetadataString) {
|
|
errs() << "Original input: " << HSAMetadataString << '\n'
|
|
<< "Produced output: " << StrOS.str() << '\n';
|
|
}
|
|
}
|
|
|
|
std::optional<StringRef>
|
|
MetadataStreamerMsgPackV4::getAccessQualifier(StringRef AccQual) const {
|
|
return StringSwitch<std::optional<StringRef>>(AccQual)
|
|
.Case("read_only", StringRef("read_only"))
|
|
.Case("write_only", StringRef("write_only"))
|
|
.Case("read_write", StringRef("read_write"))
|
|
.Default(std::nullopt);
|
|
}
|
|
|
|
std::optional<StringRef> MetadataStreamerMsgPackV4::getAddressSpaceQualifier(
|
|
unsigned AddressSpace) const {
|
|
switch (AddressSpace) {
|
|
case AMDGPUAS::PRIVATE_ADDRESS:
|
|
return StringRef("private");
|
|
case AMDGPUAS::GLOBAL_ADDRESS:
|
|
return StringRef("global");
|
|
case AMDGPUAS::CONSTANT_ADDRESS:
|
|
return StringRef("constant");
|
|
case AMDGPUAS::LOCAL_ADDRESS:
|
|
return StringRef("local");
|
|
case AMDGPUAS::FLAT_ADDRESS:
|
|
return StringRef("generic");
|
|
case AMDGPUAS::REGION_ADDRESS:
|
|
return StringRef("region");
|
|
default:
|
|
return std::nullopt;
|
|
}
|
|
}
|
|
|
|
StringRef
|
|
MetadataStreamerMsgPackV4::getValueKind(Type *Ty, StringRef TypeQual,
|
|
StringRef BaseTypeName) const {
|
|
if (TypeQual.contains("pipe"))
|
|
return "pipe";
|
|
|
|
return StringSwitch<StringRef>(BaseTypeName)
|
|
.Case("image1d_t", "image")
|
|
.Case("image1d_array_t", "image")
|
|
.Case("image1d_buffer_t", "image")
|
|
.Case("image2d_t", "image")
|
|
.Case("image2d_array_t", "image")
|
|
.Case("image2d_array_depth_t", "image")
|
|
.Case("image2d_array_msaa_t", "image")
|
|
.Case("image2d_array_msaa_depth_t", "image")
|
|
.Case("image2d_depth_t", "image")
|
|
.Case("image2d_msaa_t", "image")
|
|
.Case("image2d_msaa_depth_t", "image")
|
|
.Case("image3d_t", "image")
|
|
.Case("sampler_t", "sampler")
|
|
.Case("queue_t", "queue")
|
|
.Default(isa<PointerType>(Ty)
|
|
? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
|
|
? "dynamic_shared_pointer"
|
|
: "global_buffer")
|
|
: "by_value");
|
|
}
|
|
|
|
std::string MetadataStreamerMsgPackV4::getTypeName(Type *Ty,
|
|
bool Signed) const {
|
|
switch (Ty->getTypeID()) {
|
|
case Type::IntegerTyID: {
|
|
if (!Signed)
|
|
return (Twine('u') + getTypeName(Ty, true)).str();
|
|
|
|
auto BitWidth = Ty->getIntegerBitWidth();
|
|
switch (BitWidth) {
|
|
case 8:
|
|
return "char";
|
|
case 16:
|
|
return "short";
|
|
case 32:
|
|
return "int";
|
|
case 64:
|
|
return "long";
|
|
default:
|
|
return (Twine('i') + Twine(BitWidth)).str();
|
|
}
|
|
}
|
|
case Type::HalfTyID:
|
|
return "half";
|
|
case Type::FloatTyID:
|
|
return "float";
|
|
case Type::DoubleTyID:
|
|
return "double";
|
|
case Type::FixedVectorTyID: {
|
|
auto *VecTy = cast<FixedVectorType>(Ty);
|
|
auto *ElTy = VecTy->getElementType();
|
|
auto NumElements = VecTy->getNumElements();
|
|
return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
|
|
}
|
|
default:
|
|
return "unknown";
|
|
}
|
|
}
|
|
|
|
msgpack::ArrayDocNode
|
|
MetadataStreamerMsgPackV4::getWorkGroupDimensions(MDNode *Node) const {
|
|
auto Dims = HSAMetadataDoc->getArrayNode();
|
|
if (Node->getNumOperands() != 3)
|
|
return Dims;
|
|
|
|
for (auto &Op : Node->operands())
|
|
Dims.push_back(Dims.getDocument()->getNode(
|
|
uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
|
|
return Dims;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitVersion() {
|
|
auto Version = HSAMetadataDoc->getArrayNode();
|
|
Version.push_back(Version.getDocument()->getNode(VersionMajorV4));
|
|
Version.push_back(Version.getDocument()->getNode(VersionMinorV4));
|
|
getRootMetadata("amdhsa.version") = Version;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitTargetID(
|
|
const IsaInfo::AMDGPUTargetID &TargetID) {
|
|
getRootMetadata("amdhsa.target") =
|
|
HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitPrintf(const Module &Mod) {
|
|
auto *Node = Mod.getNamedMetadata("llvm.printf.fmts");
|
|
if (!Node)
|
|
return;
|
|
|
|
auto Printf = HSAMetadataDoc->getArrayNode();
|
|
for (auto *Op : Node->operands())
|
|
if (Op->getNumOperands())
|
|
Printf.push_back(Printf.getDocument()->getNode(
|
|
cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
|
|
getRootMetadata("amdhsa.printf") = Printf;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernelLanguage(const Function &Func,
|
|
msgpack::MapDocNode Kern) {
|
|
// TODO: What about other languages?
|
|
auto *Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
|
|
if (!Node || !Node->getNumOperands())
|
|
return;
|
|
auto *Op0 = Node->getOperand(0);
|
|
if (Op0->getNumOperands() <= 1)
|
|
return;
|
|
|
|
Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
|
|
auto LanguageVersion = Kern.getDocument()->getArrayNode();
|
|
LanguageVersion.push_back(Kern.getDocument()->getNode(
|
|
mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
|
|
LanguageVersion.push_back(Kern.getDocument()->getNode(
|
|
mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
|
|
Kern[".language_version"] = LanguageVersion;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernelAttrs(const AMDGPUTargetMachine &TM,
|
|
const Function &Func,
|
|
msgpack::MapDocNode Kern) {
|
|
|
|
if (auto *Node = Func.getMetadata("reqd_work_group_size"))
|
|
Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
|
|
if (auto *Node = Func.getMetadata("work_group_size_hint"))
|
|
Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
|
|
if (auto *Node = Func.getMetadata("vec_type_hint")) {
|
|
Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
|
|
getTypeName(
|
|
cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
|
|
mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
|
|
/*Copy=*/true);
|
|
}
|
|
|
|
std::string HandleName = getEnqueuedBlockSymbolName(TM, Func);
|
|
if (!HandleName.empty()) {
|
|
Kern[".device_enqueue_symbol"] =
|
|
Kern.getDocument()->getNode(std::move(HandleName), /*Copy=*/true);
|
|
}
|
|
|
|
if (Func.hasFnAttribute("device-init"))
|
|
Kern[".kind"] = Kern.getDocument()->getNode("init");
|
|
else if (Func.hasFnAttribute("device-fini"))
|
|
Kern[".kind"] = Kern.getDocument()->getNode("fini");
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernelArgs(const MachineFunction &MF,
|
|
msgpack::MapDocNode Kern) {
|
|
auto &Func = MF.getFunction();
|
|
unsigned Offset = 0;
|
|
auto Args = HSAMetadataDoc->getArrayNode();
|
|
for (auto &Arg : Func.args()) {
|
|
if (Arg.hasAttribute("amdgpu-hidden-argument"))
|
|
continue;
|
|
|
|
emitKernelArg(Arg, Offset, Args, MF);
|
|
}
|
|
|
|
emitHiddenKernelArgs(MF, Offset, Args);
|
|
|
|
Kern[".args"] = Args;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernelArg(const Argument &Arg,
|
|
unsigned &Offset,
|
|
msgpack::ArrayDocNode Args,
|
|
const MachineFunction &MF) {
|
|
const auto *Func = Arg.getParent();
|
|
auto ArgNo = Arg.getArgNo();
|
|
const MDNode *Node;
|
|
|
|
StringRef Name;
|
|
Node = Func->getMetadata("kernel_arg_name");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
else if (Arg.hasName())
|
|
Name = Arg.getName();
|
|
|
|
StringRef TypeName;
|
|
Node = Func->getMetadata("kernel_arg_type");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef BaseTypeName;
|
|
Node = Func->getMetadata("kernel_arg_base_type");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef ActAccQual;
|
|
// Do we really need NoAlias check here?
|
|
if (Arg.getType()->isPointerTy() && Arg.hasNoAliasAttr()) {
|
|
if (Arg.onlyReadsMemory())
|
|
ActAccQual = "read_only";
|
|
else if (Arg.hasAttribute(Attribute::WriteOnly))
|
|
ActAccQual = "write_only";
|
|
}
|
|
|
|
StringRef AccQual;
|
|
Node = Func->getMetadata("kernel_arg_access_qual");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef TypeQual;
|
|
Node = Func->getMetadata("kernel_arg_type_qual");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
const DataLayout &DL = Func->getDataLayout();
|
|
|
|
MaybeAlign PointeeAlign;
|
|
Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
|
|
|
|
// FIXME: Need to distinguish in memory alignment from pointer alignment.
|
|
if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
|
|
if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
|
|
PointeeAlign = Arg.getParamAlign().valueOrOne();
|
|
}
|
|
|
|
// There's no distinction between byval aggregates and raw aggregates.
|
|
Type *ArgTy;
|
|
Align ArgAlign;
|
|
std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
|
|
|
|
emitKernelArgImpl(DL, ArgTy, ArgAlign,
|
|
getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args,
|
|
"" /* PreloadRegisters */, PointeeAlign, Name, TypeName,
|
|
BaseTypeName, ActAccQual, AccQual, TypeQual);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernelArgImpl(
|
|
const DataLayout &DL, Type *Ty, Align Alignment, StringRef ValueKind,
|
|
unsigned &Offset, msgpack::ArrayDocNode Args, StringRef PreloadRegisters,
|
|
MaybeAlign PointeeAlign, StringRef Name, StringRef TypeName,
|
|
StringRef BaseTypeName, StringRef ActAccQual, StringRef AccQual,
|
|
StringRef TypeQual) {
|
|
auto Arg = Args.getDocument()->getMapNode();
|
|
|
|
if (!Name.empty())
|
|
Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
|
|
if (!TypeName.empty())
|
|
Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
|
|
auto Size = DL.getTypeAllocSize(Ty);
|
|
Arg[".size"] = Arg.getDocument()->getNode(Size);
|
|
Offset = alignTo(Offset, Alignment);
|
|
Arg[".offset"] = Arg.getDocument()->getNode(Offset);
|
|
Offset += Size;
|
|
Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
|
|
if (PointeeAlign)
|
|
Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
|
|
|
|
if (auto *PtrTy = dyn_cast<PointerType>(Ty))
|
|
if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
|
|
// Limiting address space to emit only for a certain ValueKind.
|
|
if (ValueKind == "global_buffer" || ValueKind == "dynamic_shared_pointer")
|
|
Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier,
|
|
/*Copy=*/true);
|
|
|
|
if (auto AQ = getAccessQualifier(AccQual))
|
|
Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
|
|
|
|
if (auto AAQ = getAccessQualifier(ActAccQual))
|
|
Arg[".actual_access"] = Arg.getDocument()->getNode(*AAQ, /*Copy=*/true);
|
|
|
|
SmallVector<StringRef, 1> SplitTypeQuals;
|
|
TypeQual.split(SplitTypeQuals, " ", -1, false);
|
|
for (StringRef Key : SplitTypeQuals) {
|
|
if (Key == "const")
|
|
Arg[".is_const"] = Arg.getDocument()->getNode(true);
|
|
else if (Key == "restrict")
|
|
Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
|
|
else if (Key == "volatile")
|
|
Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
|
|
else if (Key == "pipe")
|
|
Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
|
|
}
|
|
|
|
if (!PreloadRegisters.empty()) {
|
|
Arg[".preload_registers"] =
|
|
Arg.getDocument()->getNode(PreloadRegisters, /*Copy=*/true);
|
|
}
|
|
|
|
Args.push_back(Arg);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitHiddenKernelArgs(
|
|
const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
|
|
auto &Func = MF.getFunction();
|
|
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
|
|
|
|
unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
|
|
if (!HiddenArgNumBytes)
|
|
return;
|
|
|
|
const Module *M = Func.getParent();
|
|
auto &DL = M->getDataLayout();
|
|
auto *Int64Ty = Type::getInt64Ty(Func.getContext());
|
|
|
|
Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
|
|
|
|
if (HiddenArgNumBytes >= 8)
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset,
|
|
Args);
|
|
if (HiddenArgNumBytes >= 16)
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
|
|
Args);
|
|
if (HiddenArgNumBytes >= 24)
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
|
|
Args);
|
|
|
|
auto *Int8PtrTy =
|
|
PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
|
|
|
|
if (HiddenArgNumBytes >= 32) {
|
|
// We forbid the use of features requiring hostcall when compiling OpenCL
|
|
// before code object V5, which makes the mutual exclusion between the
|
|
// "printf buffer" and "hostcall buffer" here sound.
|
|
if (M->getNamedMetadata("llvm.printf.fmts"))
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
|
|
Args);
|
|
else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr"))
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer",
|
|
Offset, Args);
|
|
else
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
|
|
}
|
|
|
|
// Emit "default queue" and "completion action" arguments if enqueue kernel is
|
|
// used, otherwise emit dummy "none" arguments.
|
|
if (HiddenArgNumBytes >= 40) {
|
|
if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
|
|
Args);
|
|
} else {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
|
|
}
|
|
}
|
|
|
|
if (HiddenArgNumBytes >= 48) {
|
|
if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_completion_action",
|
|
Offset, Args);
|
|
} else {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
|
|
}
|
|
}
|
|
|
|
// Emit the pointer argument for multi-grid object.
|
|
if (HiddenArgNumBytes >= 56) {
|
|
if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg",
|
|
Offset, Args);
|
|
} else {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
|
|
}
|
|
}
|
|
}
|
|
|
|
msgpack::MapDocNode
|
|
MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
|
|
const SIProgramInfo &ProgramInfo,
|
|
unsigned CodeObjectVersion) const {
|
|
const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
|
|
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
|
|
const Function &F = MF.getFunction();
|
|
|
|
auto Kern = HSAMetadataDoc->getMapNode();
|
|
|
|
Align MaxKernArgAlign;
|
|
Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
|
|
STM.getKernArgSegmentSize(F, MaxKernArgAlign));
|
|
Kern[".group_segment_fixed_size"] =
|
|
Kern.getDocument()->getNode(ProgramInfo.LDSSize);
|
|
DelayedExprs->assignDocNode(Kern[".private_segment_fixed_size"],
|
|
msgpack::Type::UInt, ProgramInfo.ScratchSize);
|
|
if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5) {
|
|
DelayedExprs->assignDocNode(Kern[".uses_dynamic_stack"],
|
|
msgpack::Type::Boolean,
|
|
ProgramInfo.DynamicCallStack);
|
|
}
|
|
|
|
if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5 && STM.supportsWGP())
|
|
Kern[".workgroup_processor_mode"] =
|
|
Kern.getDocument()->getNode(ProgramInfo.WgpMode);
|
|
|
|
// FIXME: The metadata treats the minimum as 16?
|
|
Kern[".kernarg_segment_align"] =
|
|
Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
|
|
Kern[".wavefront_size"] =
|
|
Kern.getDocument()->getNode(STM.getWavefrontSize());
|
|
DelayedExprs->assignDocNode(Kern[".sgpr_count"], msgpack::Type::UInt,
|
|
ProgramInfo.NumSGPR);
|
|
DelayedExprs->assignDocNode(Kern[".vgpr_count"], msgpack::Type::UInt,
|
|
ProgramInfo.NumVGPR);
|
|
|
|
// Only add AGPR count to metadata for supported devices
|
|
if (STM.hasMAIInsts()) {
|
|
DelayedExprs->assignDocNode(Kern[".agpr_count"], msgpack::Type::UInt,
|
|
ProgramInfo.NumAccVGPR);
|
|
}
|
|
|
|
Kern[".max_flat_workgroup_size"] =
|
|
Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
|
|
|
|
uint32_t NumWGY = MFI.getMaxNumWorkGroupsY();
|
|
uint32_t NumWGZ = MFI.getMaxNumWorkGroupsZ();
|
|
uint32_t NumWGX = MFI.getMaxNumWorkGroupsX();
|
|
|
|
// TODO: Should consider 0 invalid and reject in IR verifier.
|
|
if (NumWGX != std::numeric_limits<uint32_t>::max() && NumWGX != 0)
|
|
Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
|
|
|
|
if (NumWGY != std::numeric_limits<uint32_t>::max() && NumWGY != 0)
|
|
Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
|
|
|
|
if (NumWGZ != std::numeric_limits<uint32_t>::max() && NumWGZ != 0)
|
|
Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
|
|
|
|
Kern[".sgpr_spill_count"] =
|
|
Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
|
|
Kern[".vgpr_spill_count"] =
|
|
Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
|
|
|
|
return Kern;
|
|
}
|
|
|
|
bool MetadataStreamerMsgPackV4::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
|
|
DelayedExprs->resolveDelayedExpressions();
|
|
return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::begin(const Module &Mod,
|
|
const IsaInfo::AMDGPUTargetID &TargetID) {
|
|
emitVersion();
|
|
emitTargetID(TargetID);
|
|
emitPrintf(Mod);
|
|
getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
|
|
DelayedExprs->clear();
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::end() {
|
|
DelayedExprs->resolveDelayedExpressions();
|
|
std::string HSAMetadataString;
|
|
raw_string_ostream StrOS(HSAMetadataString);
|
|
HSAMetadataDoc->toYAML(StrOS);
|
|
|
|
if (DumpHSAMetadata)
|
|
dump(StrOS.str());
|
|
if (VerifyHSAMetadata)
|
|
verify(StrOS.str());
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV4::emitKernel(const MachineFunction &MF,
|
|
const SIProgramInfo &ProgramInfo) {
|
|
auto &Func = MF.getFunction();
|
|
if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL &&
|
|
Func.getCallingConv() != CallingConv::SPIR_KERNEL)
|
|
return;
|
|
|
|
auto CodeObjectVersion =
|
|
AMDGPU::getAMDHSACodeObjectVersion(*Func.getParent());
|
|
auto Kern = getHSAKernelProps(MF, ProgramInfo, CodeObjectVersion);
|
|
|
|
auto Kernels =
|
|
getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
|
|
|
|
auto &TM = static_cast<const AMDGPUTargetMachine &>(MF.getTarget());
|
|
{
|
|
Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
|
|
Kern[".symbol"] = Kern.getDocument()->getNode(
|
|
(Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
|
|
emitKernelLanguage(Func, Kern);
|
|
emitKernelAttrs(TM, Func, Kern);
|
|
emitKernelArgs(MF, Kern);
|
|
}
|
|
|
|
Kernels.push_back(Kern);
|
|
}
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
// HSAMetadataStreamerV5
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
void MetadataStreamerMsgPackV5::emitVersion() {
|
|
auto Version = HSAMetadataDoc->getArrayNode();
|
|
Version.push_back(Version.getDocument()->getNode(VersionMajorV5));
|
|
Version.push_back(Version.getDocument()->getNode(VersionMinorV5));
|
|
getRootMetadata("amdhsa.version") = Version;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
|
|
const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
|
|
auto &Func = MF.getFunction();
|
|
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
|
|
|
|
// No implicit kernel argument is used.
|
|
if (ST.getImplicitArgNumBytes(Func) == 0)
|
|
return;
|
|
|
|
const Module *M = Func.getParent();
|
|
auto &DL = M->getDataLayout();
|
|
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
|
|
|
|
auto *Int64Ty = Type::getInt64Ty(Func.getContext());
|
|
auto *Int32Ty = Type::getInt32Ty(Func.getContext());
|
|
auto *Int16Ty = Type::getInt16Ty(Func.getContext());
|
|
|
|
Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_block_count_y", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_block_count_z", Offset,
|
|
Args);
|
|
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_group_size_x", Offset, Args);
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_group_size_y", Offset, Args);
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_group_size_z", Offset, Args);
|
|
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_remainder_x", Offset, Args);
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_remainder_y", Offset, Args);
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_remainder_z", Offset, Args);
|
|
|
|
// Reserved for hidden_tool_correlation_id.
|
|
Offset += 8;
|
|
|
|
Offset += 8; // Reserved.
|
|
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
|
|
Args);
|
|
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args);
|
|
|
|
Offset += 6; // Reserved.
|
|
auto *Int8PtrTy =
|
|
PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
|
|
|
|
if (M->getNamedMetadata("llvm.printf.fmts")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg",
|
|
Offset, Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-heap-ptr"))
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_heap_v1", Offset, Args);
|
|
else
|
|
Offset += 8; // Skipped.
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_completion_action",
|
|
Offset, Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
// Emit argument for hidden dynamic lds size
|
|
if (MFI.isDynamicLDSUsed()) {
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_dynamic_lds_size", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 4; // skipped
|
|
}
|
|
|
|
Offset += 68; // Reserved.
|
|
|
|
// hidden_private_base and hidden_shared_base are only when the subtarget has
|
|
// ApertureRegs.
|
|
if (!ST.hasApertureRegs()) {
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_private_base", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_shared_base", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (MFI.getUserSGPRInfo().hasQueuePtr())
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_queue_ptr", Offset,
|
|
Args);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV5::emitKernelAttrs(const AMDGPUTargetMachine &TM,
|
|
const Function &Func,
|
|
msgpack::MapDocNode Kern) {
|
|
MetadataStreamerMsgPackV4::emitKernelAttrs(TM, Func, Kern);
|
|
|
|
if (Func.getFnAttribute("uniform-work-group-size").getValueAsBool())
|
|
Kern[".uniform_work_group_size"] = Kern.getDocument()->getNode(1);
|
|
}
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
// HSAMetadataStreamerV6
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
void MetadataStreamerMsgPackV6::emitVersion() {
|
|
auto Version = HSAMetadataDoc->getArrayNode();
|
|
Version.push_back(Version.getDocument()->getNode(VersionMajorV6));
|
|
Version.push_back(Version.getDocument()->getNode(VersionMinorV6));
|
|
getRootMetadata("amdhsa.version") = Version;
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV6::emitHiddenKernelArgWithPreload(
|
|
const DataLayout &DL, Type *ArgTy, Align Alignment,
|
|
KernArgPreload::HiddenArg HiddenArg, StringRef ArgName, unsigned &Offset,
|
|
msgpack::ArrayDocNode Args, const AMDGPUFunctionArgInfo &ArgInfo) {
|
|
|
|
SmallString<16> PreloadStr;
|
|
auto PreloadDesc = ArgInfo.getHiddenArgPreloadDescriptor(HiddenArg);
|
|
if (PreloadDesc) {
|
|
const auto &Regs = (*PreloadDesc)->Regs;
|
|
for (unsigned I = 0; I < Regs.size(); ++I) {
|
|
if (I > 0)
|
|
PreloadStr += " ";
|
|
PreloadStr += AMDGPUInstPrinter::getRegisterName(Regs[I]);
|
|
}
|
|
}
|
|
emitKernelArgImpl(DL, ArgTy, Alignment, ArgName, Offset, Args, PreloadStr);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV6::emitHiddenKernelArgs(
|
|
const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
|
|
auto &Func = MF.getFunction();
|
|
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
|
|
|
|
// No implicit kernel argument is used.
|
|
if (ST.getImplicitArgNumBytes(Func) == 0)
|
|
return;
|
|
|
|
const Module *M = Func.getParent();
|
|
auto &DL = M->getDataLayout();
|
|
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
|
|
|
|
auto *Int64Ty = Type::getInt64Ty(Func.getContext());
|
|
auto *Int32Ty = Type::getInt32Ty(Func.getContext());
|
|
auto *Int16Ty = Type::getInt16Ty(Func.getContext());
|
|
|
|
Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
|
|
|
|
const AMDGPUFunctionArgInfo &ArgInfo = MFI.getArgInfo();
|
|
emitHiddenKernelArgWithPreload(DL, Int32Ty, Align(4),
|
|
KernArgPreload::HIDDEN_BLOCK_COUNT_X,
|
|
"hidden_block_count_x", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int32Ty, Align(4),
|
|
KernArgPreload::HIDDEN_BLOCK_COUNT_Y,
|
|
"hidden_block_count_y", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int32Ty, Align(4),
|
|
KernArgPreload::HIDDEN_BLOCK_COUNT_Z,
|
|
"hidden_block_count_z", Offset, Args, ArgInfo);
|
|
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_GROUP_SIZE_X,
|
|
"hidden_group_size_x", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_GROUP_SIZE_Y,
|
|
"hidden_group_size_y", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_GROUP_SIZE_Z,
|
|
"hidden_group_size_z", Offset, Args, ArgInfo);
|
|
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_REMAINDER_X,
|
|
"hidden_remainder_x", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_REMAINDER_Y,
|
|
"hidden_remainder_y", Offset, Args, ArgInfo);
|
|
emitHiddenKernelArgWithPreload(DL, Int16Ty, Align(2),
|
|
KernArgPreload::HIDDEN_REMAINDER_Z,
|
|
"hidden_remainder_z", Offset, Args, ArgInfo);
|
|
|
|
// Reserved for hidden_tool_correlation_id.
|
|
Offset += 8;
|
|
|
|
Offset += 8; // Reserved.
|
|
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
|
|
Args);
|
|
|
|
emitKernelArgImpl(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args);
|
|
|
|
Offset += 6; // Reserved.
|
|
auto *Int8PtrTy =
|
|
PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
|
|
|
|
if (M->getNamedMetadata("llvm.printf.fmts")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg",
|
|
Offset, Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-heap-ptr"))
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_heap_v1", Offset, Args);
|
|
else
|
|
Offset += 8; // Skipped.
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-default-queue")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (!Func.hasFnAttribute("amdgpu-no-completion-action")) {
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_completion_action",
|
|
Offset, Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
// Emit argument for hidden dynamic lds size
|
|
if (MFI.isDynamicLDSUsed()) {
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_dynamic_lds_size", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 4; // skipped
|
|
}
|
|
|
|
Offset += 68; // Reserved.
|
|
|
|
// hidden_private_base and hidden_shared_base are only when the subtarget has
|
|
// ApertureRegs.
|
|
if (!ST.hasApertureRegs()) {
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_private_base", Offset,
|
|
Args);
|
|
emitKernelArgImpl(DL, Int32Ty, Align(4), "hidden_shared_base", Offset,
|
|
Args);
|
|
} else {
|
|
Offset += 8; // Skipped.
|
|
}
|
|
|
|
if (MFI.getUserSGPRInfo().hasQueuePtr())
|
|
emitKernelArgImpl(DL, Int8PtrTy, Align(8), "hidden_queue_ptr", Offset,
|
|
Args);
|
|
}
|
|
|
|
void MetadataStreamerMsgPackV6::emitKernelArg(const Argument &Arg,
|
|
unsigned &Offset,
|
|
msgpack::ArrayDocNode Args,
|
|
const MachineFunction &MF) {
|
|
const auto *Func = Arg.getParent();
|
|
auto ArgNo = Arg.getArgNo();
|
|
const MDNode *Node;
|
|
|
|
StringRef Name;
|
|
Node = Func->getMetadata("kernel_arg_name");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
else if (Arg.hasName())
|
|
Name = Arg.getName();
|
|
|
|
StringRef TypeName;
|
|
Node = Func->getMetadata("kernel_arg_type");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef BaseTypeName;
|
|
Node = Func->getMetadata("kernel_arg_base_type");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef ActAccQual;
|
|
// Do we really need NoAlias check here?
|
|
if (Arg.getType()->isPointerTy() && Arg.hasNoAliasAttr()) {
|
|
if (Arg.onlyReadsMemory())
|
|
ActAccQual = "read_only";
|
|
else if (Arg.hasAttribute(Attribute::WriteOnly))
|
|
ActAccQual = "write_only";
|
|
}
|
|
|
|
StringRef AccQual;
|
|
Node = Func->getMetadata("kernel_arg_access_qual");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
StringRef TypeQual;
|
|
Node = Func->getMetadata("kernel_arg_type_qual");
|
|
if (Node && ArgNo < Node->getNumOperands())
|
|
TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
|
|
|
|
const DataLayout &DL = Func->getDataLayout();
|
|
|
|
MaybeAlign PointeeAlign;
|
|
Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
|
|
|
|
// FIXME: Need to distinguish in memory alignment from pointer alignment.
|
|
if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
|
|
if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
|
|
PointeeAlign = Arg.getParamAlign().valueOrOne();
|
|
}
|
|
|
|
const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
|
SmallString<8> PreloadRegisters;
|
|
if (MFI->getNumKernargPreloadedSGPRs()) {
|
|
assert(MF.getSubtarget<GCNSubtarget>().hasKernargPreload());
|
|
const auto &PreloadDescs =
|
|
MFI->getArgInfo().getPreloadDescriptorsForArgIdx(ArgNo);
|
|
for (auto &Desc : PreloadDescs) {
|
|
if (!PreloadRegisters.empty())
|
|
PreloadRegisters += " ";
|
|
|
|
for (unsigned I = 0; I < Desc->Regs.size(); ++I) {
|
|
if (I > 0)
|
|
PreloadRegisters += " ";
|
|
PreloadRegisters += AMDGPUInstPrinter::getRegisterName(Desc->Regs[I]);
|
|
}
|
|
}
|
|
}
|
|
|
|
// There's no distinction between byval aggregates and raw aggregates.
|
|
Type *ArgTy;
|
|
Align ArgAlign;
|
|
std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
|
|
|
|
emitKernelArgImpl(DL, ArgTy, ArgAlign,
|
|
getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args,
|
|
PreloadRegisters, PointeeAlign, Name, TypeName,
|
|
BaseTypeName, ActAccQual, AccQual, TypeQual);
|
|
}
|
|
|
|
} // end namespace AMDGPU::HSAMD
|
|
} // end namespace llvm
|