RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2025-01-31 12:41:49 +01:00
Code Issues Projects Releases Wiki Activity

[AMDGPU] Switched HSA metadata to use MsgPackDocument

Browse Source 
Summary:
MsgPackDocument is the lighter-weight replacement for MsgPackTypes. This
commit switches AMDGPU HSA metadata processing to use MsgPackDocument
instead of MsgPackTypes.

Differential Revision: https://reviews.llvm.org/D57024

Change-Id: I0751668013abe8c87db01db1170831a76079b3a6
llvm-svn: 356081
This commit is contained in:
Tim Renouf 2019-03-13 18:55:50 +00:00
parent 2600aff087
commit 0740b7d5b8
17 changed files with 1729 additions and 1713 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

30
include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h
View File

@ -16,7 +16,7 @@
#ifndef LLVM_BINARYFORMAT_AMDGPUMETADATAVERIFIER_H
#define LLVM_BINARYFORMAT_AMDGPUMETADATAVERIFIER_H
#include "llvm/BinaryFormat/MsgPackTypes.h"
#include "llvm/BinaryFormat/MsgPackDocument.h"
namespace llvm {
namespace AMDGPU {
@ -33,22 +33,22 @@ namespace V3 {
class MetadataVerifier {
bool Strict;
bool verifyScalar(msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind,
function_ref<bool(msgpack::ScalarNode &)> verifyValue = {});
bool verifyInteger(msgpack::Node &Node);
bool verifyArray(msgpack::Node &Node,
function_ref<bool(msgpack::Node &)> verifyNode,
bool verifyScalar(msgpack::DocNode &Node, msgpack::Type SKind,
function_ref<bool(msgpack::DocNode &)> verifyValue = {});
bool verifyInteger(msgpack::DocNode &Node);
bool verifyArray(msgpack::DocNode &Node,
function_ref<bool(msgpack::DocNode &)> verifyNode,
Optional<size_t> Size = None);
bool verifyEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required,
function_ref<bool(msgpack::Node &)> verifyNode);
bool verifyEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
function_ref<bool(msgpack::DocNode &)> verifyNode);
bool
verifyScalarEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required,
msgpack::ScalarNode::ScalarKind SKind,
function_ref<bool(msgpack::ScalarNode &)> verifyValue = {});
bool verifyIntegerEntry(msgpack::MapNode &MapNode, StringRef Key,
verifyScalarEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
msgpack::Type SKind,
function_ref<bool(msgpack::DocNode &)> verifyValue = {});
bool verifyIntegerEntry(msgpack::MapDocNode &MapNode, StringRef Key,
bool Required);
bool verifyKernelArgs(msgpack::Node &Node);
bool verifyKernel(msgpack::Node &Node);
bool verifyKernelArgs(msgpack::DocNode &Node);
bool verifyKernel(msgpack::DocNode &Node);
public:
/// Construct a MetadataVerifier, specifying whether it will operate in \p
@ -58,7 +58,7 @@ public:
/// Verify given HSA metadata.
///
/// \returns True when successful, false when metadata is invalid.
bool verify(msgpack::Node &HSAMetadataRoot);
bool verify(msgpack::DocNode &HSAMetadataRoot);
};
} // end namespace V3

150
lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
View File

@ -20,98 +20,92 @@ namespace HSAMD {
namespace V3 {
bool MetadataVerifier::verifyScalar(
msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind,
function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
auto ScalarPtr = dyn_cast<msgpack::ScalarNode>(&Node);
if (!ScalarPtr)
msgpack::DocNode &Node, msgpack::Type SKind,
function_ref<bool(msgpack::DocNode &)> verifyValue) {
if (!Node.isScalar())
return false;
auto &Scalar = *ScalarPtr;
// Do not output extraneous tags for types we know from the spec.
Scalar.IgnoreTag = true;
if (Scalar.getScalarKind() != SKind) {
if (Node.getKind() != SKind) {
if (Strict)
return false;
// If we are not strict, we interpret string values as "implicitly typed"
// and attempt to coerce them to the expected type here.
if (Scalar.getScalarKind() != msgpack::ScalarNode::SK_String)
if (Node.getKind() != msgpack::Type::String)
return false;
std::string StringValue = Scalar.getString();
Scalar.setScalarKind(SKind);
if (Scalar.inputYAML(StringValue) != StringRef())
StringRef StringValue = Node.getString();
Node.fromString(StringValue);
if (Node.getKind() != SKind)
return false;
}
if (verifyValue)
return verifyValue(Scalar);
return verifyValue(Node);
return true;
}
bool MetadataVerifier::verifyInteger(msgpack::Node &Node) {
if (!verifyScalar(Node, msgpack::ScalarNode::SK_UInt))
if (!verifyScalar(Node, msgpack::ScalarNode::SK_Int))
bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
if (!verifyScalar(Node, msgpack::Type::UInt))
if (!verifyScalar(Node, msgpack::Type::Int))
return false;
return true;
}
bool MetadataVerifier::verifyArray(
msgpack::Node &Node, function_ref<bool(msgpack::Node &)> verifyNode,
msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
Optional<size_t> Size) {
auto ArrayPtr = dyn_cast<msgpack::ArrayNode>(&Node);
if (!ArrayPtr)
if (!Node.isArray())
return false;
auto &Array = *ArrayPtr;
auto &Array = Node.getArray();
if (Size && Array.size() != *Size)
return false;
for (auto &Item : Array)
if (!verifyNode(*Item.get()))
if (!verifyNode(Item))
return false;
return true;
}
bool MetadataVerifier::verifyEntry(
msgpack::MapNode &MapNode, StringRef Key, bool Required,
function_ref<bool(msgpack::Node &)> verifyNode) {
msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
function_ref<bool(msgpack::DocNode &)> verifyNode) {
auto Entry = MapNode.find(Key);
if (Entry == MapNode.end())
return !Required;
return verifyNode(*Entry->second.get());
return verifyNode(Entry->second);
}
bool MetadataVerifier::verifyScalarEntry(
msgpack::MapNode &MapNode, StringRef Key, bool Required,
msgpack::ScalarNode::ScalarKind SKind,
function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
return verifyEntry(MapNode, Key, Required, [=](msgpack::Node &Node) {
msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
msgpack::Type SKind,
function_ref<bool(msgpack::DocNode &)> verifyValue) {
return verifyEntry(MapNode, Key, Required, [=](msgpack::DocNode &Node) {
return verifyScalar(Node, SKind, verifyValue);
});
}
bool MetadataVerifier::verifyIntegerEntry(msgpack::MapNode &MapNode,
bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
StringRef Key, bool Required) {
return verifyEntry(MapNode, Key, Required, [this](msgpack::Node &Node) {
return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {
return verifyInteger(Node);
});
}
bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
auto ArgsMapPtr = dyn_cast<msgpack::MapNode>(&Node);
if (!ArgsMapPtr)
bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {
if (!Node.isMap())
return false;
auto &ArgsMap = *ArgsMapPtr;
auto &ArgsMap = Node.getMap();
if (!verifyScalarEntry(ArgsMap, ".name", false,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyScalarEntry(ArgsMap, ".type_name", false,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyIntegerEntry(ArgsMap, ".size", true))
return false;
if (!verifyIntegerEntry(ArgsMap, ".offset", true))
return false;
if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("by_value", true)
.Case("global_buffer", true)
@ -131,8 +125,8 @@ bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
}))
return false;
if (!verifyScalarEntry(ArgsMap, ".value_type", true,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("struct", true)
.Case("i8", true)
@ -152,8 +146,8 @@ bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
return false;
if (!verifyScalarEntry(ArgsMap, ".address_space", false,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("private", true)
.Case("global", true)
@ -165,8 +159,8 @@ bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
}))
return false;
if (!verifyScalarEntry(ArgsMap, ".access", false,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("read_only", true)
.Case("write_only", true)
@ -175,8 +169,8 @@ bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
}))
return false;
if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("read_only", true)
.Case("write_only", true)
@ -185,36 +179,35 @@ bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
}))
return false;
if (!verifyScalarEntry(ArgsMap, ".is_const", false,
msgpack::ScalarNode::SK_Boolean))
msgpack::Type::Boolean))
return false;
if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
msgpack::ScalarNode::SK_Boolean))
msgpack::Type::Boolean))
return false;
if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
msgpack::ScalarNode::SK_Boolean))
msgpack::Type::Boolean))
return false;
if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
msgpack::ScalarNode::SK_Boolean))
msgpack::Type::Boolean))
return false;
return true;
}
bool MetadataVerifier::verifyKernel(msgpack::Node &Node) {
auto KernelMapPtr = dyn_cast<msgpack::MapNode>(&Node);
if (!KernelMapPtr)
bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
if (!Node.isMap())
return false;
auto &KernelMap = *KernelMapPtr;
auto &KernelMap = Node.getMap();
if (!verifyScalarEntry(KernelMap, ".name", true,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyScalarEntry(KernelMap, ".symbol", true,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyScalarEntry(KernelMap, ".language", false,
msgpack::ScalarNode::SK_String,
[](msgpack::ScalarNode &SNode) {
msgpack::Type::String,
[](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString())
.Case("OpenCL C", true)
.Case("OpenCL C++", true)
@ -226,41 +219,41 @@ bool MetadataVerifier::verifyKernel(msgpack::Node &Node) {
}))
return false;
if (!verifyEntry(
KernelMap, ".language_version", false, [this](msgpack::Node &Node) {
KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
return verifyArray(
Node,
[this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
[this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
}))
return false;
if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::Node &Node) {
return verifyArray(Node, [this](msgpack::Node &Node) {
if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
return verifyArray(Node, [this](msgpack::DocNode &Node) {
return verifyKernelArgs(Node);
});
}))
return false;
if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
[this](msgpack::Node &Node) {
[this](msgpack::DocNode &Node) {
return verifyArray(Node,
[this](msgpack::Node &Node) {
[this](msgpack::DocNode &Node) {
return verifyInteger(Node);
},
3);
}))
return false;
if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
[this](msgpack::Node &Node) {
[this](msgpack::DocNode &Node) {
return verifyArray(Node,
[this](msgpack::Node &Node) {
[this](msgpack::DocNode &Node) {
return verifyInteger(Node);
},
3);
}))
return false;
if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
msgpack::ScalarNode::SK_String))
msgpack::Type::String))
return false;
if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
return false;
@ -286,29 +279,28 @@ bool MetadataVerifier::verifyKernel(msgpack::Node &Node) {
return true;
}
bool MetadataVerifier::verify(msgpack::Node &HSAMetadataRoot) {
auto RootMapPtr = dyn_cast<msgpack::MapNode>(&HSAMetadataRoot);
if (!RootMapPtr)
bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
if (!HSAMetadataRoot.isMap())
return false;
auto &RootMap = *RootMapPtr;
auto &RootMap = HSAMetadataRoot.getMap();
if (!verifyEntry(
RootMap, "amdhsa.version", true, [this](msgpack::Node &Node) {
RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
return verifyArray(
Node,
[this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
[this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
}))
return false;
if (!verifyEntry(
RootMap, "amdhsa.printf", false, [this](msgpack::Node &Node) {
return verifyArray(Node, [this](msgpack::Node &Node) {
return verifyScalar(Node, msgpack::ScalarNode::SK_String);
RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
return verifyArray(Node, [this](msgpack::DocNode &Node) {
return verifyScalar(Node, msgpack::Type::String);
});
}))
return false;
if (!verifyEntry(RootMap, "amdhsa.kernels", true,
[this](msgpack::Node &Node) {
return verifyArray(Node, [this](msgpack::Node &Node) {
[this](msgpack::DocNode &Node) {
return verifyArray(Node, [this](msgpack::DocNode &Node) {
return verifyKernel(Node);
});
}))

175
lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
View File

@ -489,20 +489,16 @@ void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
errs() << "AMDGPU HSA Metadata Parser Test: ";
std::shared_ptr<msgpack::Node> FromHSAMetadataString =
std::make_shared<msgpack::MapNode>();
msgpack::Document FromHSAMetadataString;
yaml::Input YIn(HSAMetadataString);
YIn >> FromHSAMetadataString;
if (YIn.error()) {
if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
errs() << "FAIL\n";
return;
}
std::string ToHSAMetadataString;
raw_string_ostream StrOS(ToHSAMetadataString);
yaml::Output YOut(StrOS);
YOut << FromHSAMetadataString;
FromHSAMetadataString.toYAML(StrOS);
errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
if (HSAMetadataString != ToHSAMetadataString) {
@ -636,23 +632,23 @@ std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
}
}
std::shared_ptr<msgpack::ArrayNode>
msgpack::ArrayDocNode
MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
auto Dims = std::make_shared<msgpack::ArrayNode>();
auto Dims = HSAMetadataDoc->getArrayNode();
if (Node->getNumOperands() != 3)
return Dims;
for (auto &Op : Node->operands())
Dims->push_back(std::make_shared<msgpack::ScalarNode>(
mdconst::extract<ConstantInt>(Op)->getZExtValue()));
Dims.push_back(Dims.getDocument()->getNode(
uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
return Dims;
}
void MetadataStreamerV3::emitVersion() {
auto Version = std::make_shared<msgpack::ArrayNode>();
Version->push_back(std::make_shared<msgpack::ScalarNode>(V3::VersionMajor));
Version->push_back(std::make_shared<msgpack::ScalarNode>(V3::VersionMinor));
getRootMetadata("amdhsa.version") = std::move(Version);
auto Version = HSAMetadataDoc->getArrayNode();
Version.push_back(Version.getDocument()->getNode(VersionMajor));
Version.push_back(Version.getDocument()->getNode(VersionMinor));
getRootMetadata("amdhsa.version") = Version;
}
void MetadataStreamerV3::emitPrintf(const Module &Mod) {
@ -660,16 +656,16 @@ void MetadataStreamerV3::emitPrintf(const Module &Mod) {
if (!Node)
return;
auto Printf = std::make_shared<msgpack::ArrayNode>();
auto Printf = HSAMetadataDoc->getArrayNode();
for (auto Op : Node->operands())
if (Op->getNumOperands())
Printf->push_back(std::make_shared<msgpack::ScalarNode>(
cast<MDString>(Op->getOperand(0))->getString()));
getRootMetadata("amdhsa.printf") = std::move(Printf);
Printf.push_back(Printf.getDocument()->getNode(
cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
getRootMetadata("amdhsa.printf") = Printf;
}
void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
msgpack::MapNode &Kern) {
msgpack::MapDocNode Kern) {
// TODO: What about other languages?
auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
if (!Node || !Node->getNumOperands())
@ -678,50 +674,53 @@ void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
if (Op0->getNumOperands() <= 1)
return;
Kern[".language"] = std::make_shared<msgpack::ScalarNode>("OpenCL C");
auto LanguageVersion = std::make_shared<msgpack::ArrayNode>();
LanguageVersion->push_back(std::make_shared<msgpack::ScalarNode>(
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(std::make_shared<msgpack::ScalarNode>(
LanguageVersion.push_back(Kern.getDocument()->getNode(
mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
Kern[".language_version"] = std::move(LanguageVersion);
Kern[".language_version"] = LanguageVersion;
}
void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
msgpack::MapNode &Kern) {
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"] = std::make_shared<msgpack::ScalarNode>(getTypeName(
cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()));
Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
getTypeName(
cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
/*Copy=*/true);
}
if (Func.hasFnAttribute("runtime-handle")) {
Kern[".device_enqueue_symbol"] = std::make_shared<msgpack::ScalarNode>(
Func.getFnAttribute("runtime-handle").getValueAsString().str());
Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
Func.getFnAttribute("runtime-handle").getValueAsString().str(),
/*Copy=*/true);
}
}
void MetadataStreamerV3::emitKernelArgs(const Function &Func,
msgpack::MapNode &Kern) {
msgpack::MapDocNode Kern) {
unsigned Offset = 0;
auto Args = std::make_shared<msgpack::ArrayNode>();
auto Args = HSAMetadataDoc->getArrayNode();
for (auto &Arg : Func.args())
emitKernelArg(Arg, Offset, *Args);
emitKernelArg(Arg, Offset, Args);
emitHiddenKernelArgs(Func, Offset, *Args);
emitHiddenKernelArgs(Func, Offset, Args);
// TODO: What about other languages?
if (Func.getParent()->getNamedMetadata("opencl.ocl.version")) {
auto &DL = Func.getParent()->getDataLayout();
auto Int64Ty = Type::getInt64Ty(Func.getContext());
emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, *Args);
emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, *Args);
emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, *Args);
emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, Args);
emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, Args);
emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, Args);
auto Int8PtrTy =
Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
@ -729,26 +728,26 @@ void MetadataStreamerV3::emitKernelArgs(const Function &Func,
// Emit "printf buffer" argument if printf is used, otherwise emit dummy
// "none" argument.
if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, Args);
else
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
// Emit "default queue" and "completion action" arguments if enqueue kernel
// is used, otherwise emit dummy "none" arguments.
if (Func.hasFnAttribute("calls-enqueue-kernel")) {
emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
} else {
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
}
}
Kern[".args"] = std::move(Args);
Kern[".args"] = Args;
}
void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
msgpack::ArrayNode &Args) {
msgpack::ArrayDocNode Args) {
auto Func = Arg.getParent();
auto ArgNo = Arg.getArgNo();
const MDNode *Node;
@ -805,36 +804,35 @@ void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
StringRef ValueKind, unsigned &Offset,
msgpack::ArrayNode &Args,
msgpack::ArrayDocNode Args,
unsigned PointeeAlign, StringRef Name,
StringRef TypeName,
StringRef BaseTypeName,
StringRef AccQual, StringRef TypeQual) {
auto ArgPtr = std::make_shared<msgpack::MapNode>();
auto &Arg = *ArgPtr;
auto Arg = Args.getDocument()->getMapNode();
if (!Name.empty())
Arg[".name"] = std::make_shared<msgpack::ScalarNode>(Name);
Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
if (!TypeName.empty())
Arg[".type_name"] = std::make_shared<msgpack::ScalarNode>(TypeName);
Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
auto Size = DL.getTypeAllocSize(Ty);
auto Align = DL.getABITypeAlignment(Ty);
Arg[".size"] = std::make_shared<msgpack::ScalarNode>(Size);
Arg[".size"] = Arg.getDocument()->getNode(Size);
Offset = alignTo(Offset, Align);
Arg[".offset"] = std::make_shared<msgpack::ScalarNode>(Offset);
Arg[".offset"] = Arg.getDocument()->getNode(Offset);
Offset += Size;
Arg[".value_kind"] = std::make_shared<msgpack::ScalarNode>(ValueKind);
Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
Arg[".value_type"] =
std::make_shared<msgpack::ScalarNode>(getValueType(Ty, BaseTypeName));
Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
if (PointeeAlign)
Arg[".pointee_align"] = std::make_shared<msgpack::ScalarNode>(PointeeAlign);
Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign);
if (auto PtrTy = dyn_cast<PointerType>(Ty))
if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
Arg[".address_space"] = std::make_shared<msgpack::ScalarNode>(*Qualifier);
Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
if (auto AQ = getAccessQualifier(AccQual))
Arg[".access"] = std::make_shared<msgpack::ScalarNode>(*AQ);
Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
// TODO: Emit Arg[".actual_access"].
@ -842,21 +840,21 @@ void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
TypeQual.split(SplitTypeQuals, " ", -1, false);
for (StringRef Key : SplitTypeQuals) {
if (Key == "const")
Arg[".is_const"] = std::make_shared<msgpack::ScalarNode>(true);
Arg[".is_const"] = Arg.getDocument()->getNode(true);
else if (Key == "restrict")
Arg[".is_restrict"] = std::make_shared<msgpack::ScalarNode>(true);
Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
else if (Key == "volatile")
Arg[".is_volatile"] = std::make_shared<msgpack::ScalarNode>(true);
Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
else if (Key == "pipe")
Arg[".is_pipe"] = std::make_shared<msgpack::ScalarNode>(true);
Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
}
Args.push_back(std::move(ArgPtr));
Args.push_back(Arg);
}
void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
unsigned &Offset,
msgpack::ArrayNode &Args) {
msgpack::ArrayDocNode Args) {
int HiddenArgNumBytes =
getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
@ -898,54 +896,52 @@ void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
}
}
std::shared_ptr<msgpack::MapNode>
msgpack::MapDocNode
MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo) const {
const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
const Function &F = MF.getFunction();
auto HSAKernelProps = std::make_shared<msgpack::MapNode>();
auto &Kern = *HSAKernelProps;
auto Kern = HSAMetadataDoc->getMapNode();
unsigned MaxKernArgAlign;
Kern[".kernarg_segment_size"] = std::make_shared<msgpack::ScalarNode>(
Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
STM.getKernArgSegmentSize(F, MaxKernArgAlign));
Kern[".group_segment_fixed_size"] =
std::make_shared<msgpack::ScalarNode>(ProgramInfo.LDSSize);
Kern.getDocument()->getNode(ProgramInfo.LDSSize);
Kern[".private_segment_fixed_size"] =
std::make_shared<msgpack::ScalarNode>(ProgramInfo.ScratchSize);
Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
Kern[".kernarg_segment_align"] =
std::make_shared<msgpack::ScalarNode>(std::max(uint32_t(4), MaxKernArgAlign));
Kern.getDocument()->getNode(std::max(uint32_t(4), MaxKernArgAlign));
Kern[".wavefront_size"] =
std::make_shared<msgpack::ScalarNode>(STM.getWavefrontSize());
Kern[".sgpr_count"] = std::make_shared<msgpack::ScalarNode>(ProgramInfo.NumSGPR);
Kern[".vgpr_count"] = std::make_shared<msgpack::ScalarNode>(ProgramInfo.NumVGPR);
Kern.getDocument()->getNode(STM.getWavefrontSize());
Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
Kern[".max_flat_workgroup_size"] =
std::make_shared<msgpack::ScalarNode>(MFI.getMaxFlatWorkGroupSize());
Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
Kern[".sgpr_spill_count"] =
std::make_shared<msgpack::ScalarNode>(MFI.getNumSpilledSGPRs());
Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
Kern[".vgpr_spill_count"] =
std::make_shared<msgpack::ScalarNode>(MFI.getNumSpilledVGPRs());
Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
return HSAKernelProps;
return Kern;
}
bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
return TargetStreamer.EmitHSAMetadata(getHSAMetadataRoot(), true);
return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
}
void MetadataStreamerV3::begin(const Module &Mod) {
emitVersion();
emitPrintf(Mod);
getRootMetadata("amdhsa.kernels").reset(new msgpack::ArrayNode());
getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
}
void MetadataStreamerV3::end() {
std::string HSAMetadataString;
raw_string_ostream StrOS(HSAMetadataString);
yaml::Output YOut(StrOS);
YOut << HSAMetadataRoot;
HSAMetadataDoc->toYAML(StrOS);
if (DumpHSAMetadata)
dump(StrOS.str());
@ -956,25 +952,24 @@ void MetadataStreamerV3::end() {
void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo) {
auto &Func = MF.getFunction();
auto KernelProps = getHSAKernelProps(MF, ProgramInfo);
auto Kern = getHSAKernelProps(MF, ProgramInfo);
assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
Func.getCallingConv() == CallingConv::SPIR_KERNEL);
auto &KernelsNode = getRootMetadata("amdhsa.kernels");
auto Kernels = cast<msgpack::ArrayNode>(KernelsNode.get());
auto Kernels =
getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
{
auto &Kern = *KernelProps;
Kern[".name"] = std::make_shared<msgpack::ScalarNode>(Func.getName());
Kern[".symbol"] = std::make_shared<msgpack::ScalarNode>(
(Twine(Func.getName()) + Twine(".kd")).str());
Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
Kern[".symbol"] = Kern.getDocument()->getNode(
(Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
emitKernelLanguage(Func, Kern);
emitKernelAttrs(Func, Kern);
emitKernelArgs(Func, Kern);
}
Kernels->push_back(std::move(KernelProps));
Kernels.push_back(Kern);
}
} // end namespace HSAMD

34
lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
View File

@ -18,7 +18,7 @@
#include "AMDGPU.h"
#include "AMDKernelCodeT.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/BinaryFormat/MsgPackTypes.h"
#include "llvm/BinaryFormat/MsgPackDocument.h"
#include "llvm/Support/AMDGPUMetadata.h"
namespace llvm {
@ -51,8 +51,8 @@ public:
class MetadataStreamerV3 final : public MetadataStreamer {
private:
std::shared_ptr<msgpack::Node> HSAMetadataRoot =
std::make_shared<msgpack::MapNode>();
std::unique_ptr<msgpack::Document> HSAMetadataDoc =
llvm::make_unique<msgpack::Document>();
void dump(StringRef HSAMetadataString) const;
@ -69,41 +69,39 @@ private:
std::string getTypeName(Type *Ty, bool Signed) const;
std::shared_ptr<msgpack::ArrayNode>
getWorkGroupDimensions(MDNode *Node) const;
msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const;
std::shared_ptr<msgpack::MapNode>
getHSAKernelProps(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo) const;
msgpack::MapDocNode getHSAKernelProps(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo) const;
void emitVersion();
void emitPrintf(const Module &Mod);
void emitKernelLanguage(const Function &Func, msgpack::MapNode &Kern);
void emitKernelLanguage(const Function &Func, msgpack::MapDocNode Kern);
void emitKernelAttrs(const Function &Func, msgpack::MapNode &Kern);
void emitKernelAttrs(const Function &Func, msgpack::MapDocNode Kern);
void emitKernelArgs(const Function &Func, msgpack::MapNode &Kern);
void emitKernelArgs(const Function &Func, msgpack::MapDocNode Kern);
void emitKernelArg(const Argument &Arg, unsigned &Offset,
msgpack::ArrayNode &Args);
msgpack::ArrayDocNode Args);
void emitKernelArg(const DataLayout &DL, Type *Ty, StringRef ValueKind,
unsigned &Offset, msgpack::ArrayNode &Args,
unsigned &Offset, msgpack::ArrayDocNode Args,
unsigned PointeeAlign = 0, StringRef Name = "",
StringRef TypeName = "", StringRef BaseTypeName = "",
StringRef AccQual = "", StringRef TypeQual = "");
void emitHiddenKernelArgs(const Function &Func, unsigned &Offset,
msgpack::ArrayNode &Args);
msgpack::ArrayDocNode Args);
std::shared_ptr<msgpack::Node> &getRootMetadata(StringRef Key) {
return (*cast<msgpack::MapNode>(HSAMetadataRoot.get()))[Key];
msgpack::DocNode &getRootMetadata(StringRef Key) {
return HSAMetadataDoc->getRoot().getMap(/*Convert=*/true)[Key];
}
std::shared_ptr<msgpack::Node> &getHSAMetadataRoot() {
return HSAMetadataRoot;
msgpack::DocNode &getHSAMetadataRoot() {
return HSAMetadataDoc->getRoot();
}
public:

28
lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
View File

@ -18,7 +18,6 @@
#include "llvm/ADT/Twine.h"
#include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
#include "llvm/BinaryFormat/ELF.h"
#include "llvm/BinaryFormat/MsgPackTypes.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/Metadata.h"
@ -51,12 +50,10 @@ bool AMDGPUTargetStreamer::EmitHSAMetadataV2(StringRef HSAMetadataString) {
}
bool AMDGPUTargetStreamer::EmitHSAMetadataV3(StringRef HSAMetadataString) {
std::shared_ptr<msgpack::Node> HSAMetadataRoot;
yaml::Input YIn(HSAMetadataString);
YIn >> HSAMetadataRoot;
if (YIn.error())
msgpack::Document HSAMetadataDoc;
if (!HSAMetadataDoc.fromYAML(HSAMetadataString))
return false;
return EmitHSAMetadata(HSAMetadataRoot, false);
return EmitHSAMetadata(HSAMetadataDoc, false);
}
StringRef AMDGPUTargetStreamer::getArchNameFromElfMach(unsigned ElfMach) {
@ -213,15 +210,14 @@ bool AMDGPUTargetAsmStreamer::EmitHSAMetadata(
}
bool AMDGPUTargetAsmStreamer::EmitHSAMetadata(
std::shared_ptr<msgpack::Node> &HSAMetadataRoot, bool Strict) {
msgpack::Document &HSAMetadataDoc, bool Strict) {
V3::MetadataVerifier Verifier(Strict);
if (!Verifier.verify(*HSAMetadataRoot))
if (!Verifier.verify(HSAMetadataDoc.getRoot()))
return false;
std::string HSAMetadataString;
raw_string_ostream StrOS(HSAMetadataString);
yaml::Output YOut(StrOS);
YOut << HSAMetadataRoot;
HSAMetadataDoc.toYAML(StrOS);
OS << '\t' << V3::AssemblerDirectiveBegin << '\n';
OS << StrOS.str() << '\n';
@ -481,16 +477,14 @@ bool AMDGPUTargetELFStreamer::EmitISAVersion(StringRef IsaVersionString) {
return true;
}
bool AMDGPUTargetELFStreamer::EmitHSAMetadata(
std::shared_ptr<msgpack::Node> &HSAMetadataRoot, bool Strict) {
bool AMDGPUTargetELFStreamer::EmitHSAMetadata(msgpack::Document &HSAMetadataDoc,
bool Strict) {
V3::MetadataVerifier Verifier(Strict);
if (!Verifier.verify(*HSAMetadataRoot))
if (!Verifier.verify(HSAMetadataDoc.getRoot()))
return false;
std::string HSAMetadataString;
raw_string_ostream StrOS(HSAMetadataString);
msgpack::Writer MPWriter(StrOS);
HSAMetadataRoot->write(MPWriter);
HSAMetadataDoc.writeToBlob(HSAMetadataString);
// Create two labels to mark the beginning and end of the desc field
// and a MCExpr to calculate the size of the desc field.
@ -504,7 +498,7 @@ bool AMDGPUTargetELFStreamer::EmitHSAMetadata(
EmitNote(ElfNote::NoteNameV3, DescSZ, ELF::NT_AMDGPU_METADATA,
[&](MCELFStreamer &OS) {
OS.EmitLabel(DescBegin);
OS.EmitBytes(StrOS.str());
OS.EmitBytes(HSAMetadataString);
OS.EmitLabel(DescEnd);
});
return true;

11
lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
View File

@ -10,7 +10,7 @@
#define LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUTARGETSTREAMER_H
#include "AMDKernelCodeT.h"
#include "llvm/BinaryFormat/MsgPackTypes.h"
#include "llvm/BinaryFormat/MsgPackDocument.h"
#include "llvm/MC/MCStreamer.h"
#include "llvm/MC/MCSubtargetInfo.h"
#include "llvm/Support/AMDGPUMetadata.h"
@ -64,8 +64,7 @@ public:
/// the \p HSAMetadata structure is updated with the correct types.
///
/// \returns True on success, false on failure.
virtual bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
bool Strict) = 0;
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) = 0;
/// \returns True on success, false on failure.
virtual bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) = 0;
@ -105,8 +104,7 @@ public:
bool EmitISAVersion(StringRef IsaVersionString) override;
/// \returns True on success, false on failure.
bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
bool Strict) override;
bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) override;
/// \returns True on success, false on failure.
bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) override;
@ -149,8 +147,7 @@ public:
bool EmitISAVersion(StringRef IsaVersionString) override;
/// \returns True on success, false on failure.
bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
bool Strict) override;
bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) override;
/// \returns True on success, false on failure.
bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) override;

40
test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll
View File

@ -1,25 +1,25 @@
; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck %s
; CHECK: .symbol: test_ro_arg.kd
; CHECK: .name: test_ro_arg
; CHECK: .args:
; CHECK-NEXT: - .type_name: 'float*'
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .name: in
; CHECK-NEXT: .access: read_only
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .is_const: true
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .is_restrict: true
; CHECK-NEXT: .value_type: f32
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .type_name: 'float*'
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .name: out
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: f32
; CHECK-NEXT: .address_space: global
; CHECK: - .args:
; CHECK-NEXT: - .access: read_only
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: .is_const: true
; CHECK-NEXT: .is_restrict: true
; CHECK-NEXT: .name: in
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: 'float*'
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f32
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: out
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: 'float*'
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f32
; CHECK: .name: test_ro_arg
; CHECK: .symbol: test_ro_arg.kd
define amdgpu_kernel void @test_ro_arg(float addrspace(1)* noalias readonly %in, float addrspace(1)* %out)
!kernel_arg_addr_space !0 !kernel_arg_access_qual !1 !kernel_arg_type !2

134
test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll
View File

@ -1,81 +1,81 @@
; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: .symbol: test_non_enqueue_kernel_caller.kd
; CHECK: .name: test_non_enqueue_kernel_caller
; CHECK: .language: OpenCL C
; CHECK: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK: .args:
; CHECK-NEXT: - .type_name: char
; CHECK-NEXT: .value_kind: by_value
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 1
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .name: a
; CHECK-NEXT: - .value_kind: hidden_global_offset_x
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_y
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_z
; CHECK-NEXT: .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: - .args:
; CHECK-NEXT: - .name: a
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 1
; CHECK-NEXT: .type_name: char
; CHECK-NEXT: .value_kind: by_value
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: .value_type: i64
; CHECK-NOT: .value_kind: hidden_default_queue
; CHECK-NOT: .value_kind: hidden_completion_action
; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK: .name: test_non_enqueue_kernel_caller
; CHECK: .symbol: test_non_enqueue_kernel_caller.kd
define amdgpu_kernel void @test_non_enqueue_kernel_caller(i8 %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
ret void
}
; CHECK: .symbol: test_enqueue_kernel_caller.kd
; CHECK: .name: test_enqueue_kernel_caller
; CHECK: .language: OpenCL C
; CHECK: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK: .args:
; CHECK-NEXT: - .type_name: char
; CHECK-NEXT: .value_kind: by_value
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 1
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .name: a
; CHECK-NEXT: - .value_kind: hidden_global_offset_x
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_y
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_z
; CHECK-NEXT: .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_none
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: hidden_default_queue
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: hidden_completion_action
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK: - .args:
; CHECK-NEXT: - .name: a
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 1
; CHECK-NEXT: .type_name: char
; CHECK-NEXT: .value_kind: by_value
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_default_queue
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_completion_action
; CHECK-NEXT: .value_type: i8
; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK: .name: test_enqueue_kernel_caller
; CHECK: .symbol: test_enqueue_kernel_caller.kd
define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 {

2313
test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll
View File

File diff suppressed because it is too large Load Diff

104
test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll
View File

@ -2,56 +2,60 @@
; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX803 --check-prefix=NOTES %s
; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: .symbol: test.kd
; CHECK: .name: test
; CHECK: .args:
; CHECK-NEXT: - .value_kind: global_buffer
; CHECK-NEXT: .name: r
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: global_buffer
; CHECK-NEXT: .name: a
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: global_buffer
; CHECK-NEXT: .name: b
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: hidden_global_offset_x
; CHECK-NEXT: .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_y
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_global_offset_z
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .value_kind: hidden_none
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: hidden_none
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: - .value_kind: hidden_none
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .address_space: global
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: - .args:
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: r
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: a
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: b
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK: .name: test
; CHECK: .symbol: test.kd
define amdgpu_kernel void @test(
half addrspace(1)* %r,
half addrspace(1)* %a,

140
test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll
View File

@ -15,59 +15,93 @@
%opencl.image2d_msaa_depth_t = type opaque
%opencl.image3d_t = type opaque
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: .symbol: test.kd
; CHECK: .name: test
; CHECK: .args:
; CHECK: - .type_name: image1d_t
; CHECK: .value_kind: image
; CHECK: .name: a
; CHECK: .size: 8
; CHECK: - .type_name: image1d_array_t
; CHECK: .value_kind: image
; CHECK: .name: b
; CHECK: .size: 8
; CHECK: - .type_name: image1d_buffer_t
; CHECK: .value_kind: image
; CHECK: .name: c
; CHECK: .size: 8
; CHECK: - .type_name: image2d_t
; CHECK: .value_kind: image
; CHECK: .name: d
; CHECK: .size: 8
; CHECK: - .type_name: image2d_array_t
; CHECK: .value_kind: image
; CHECK: .name: e
; CHECK: .size: 8
; CHECK: - .type_name: image2d_array_depth_t
; CHECK: .value_kind: image
; CHECK: .name: f
; CHECK: .size: 8
; CHECK: - .type_name: image2d_array_msaa_t
; CHECK: .value_kind: image
; CHECK: .name: g
; CHECK: .size: 8
; CHECK: - .type_name: image2d_array_msaa_depth_t
; CHECK: .value_kind: image
; CHECK: .name: h
; CHECK: .size: 8
; CHECK: - .type_name: image2d_depth_t
; CHECK: .value_kind: image
; CHECK: .name: i
; CHECK: .size: 8
; CHECK: - .type_name: image2d_msaa_t
; CHECK: .value_kind: image
; CHECK: .name: j
; CHECK: .size: 8
; CHECK: - .type_name: image2d_msaa_depth_t
; CHECK: .value_kind: image
; CHECK: .name: k
; CHECK: .size: 8
; CHECK: - .type_name: image3d_t
; CHECK: .value_kind: image
; CHECK: .name: l
; CHECK: .size: 8
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: - .args:
; CHECK: - .address_space: global
; CHECK: .name: a
; CHECK: .offset: 0
; CHECK: .size: 8
; CHECK: .type_name: image1d_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: b
; CHECK: .offset: 8
; CHECK: .size: 8
; CHECK: .type_name: image1d_array_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: c
; CHECK: .offset: 16
; CHECK: .size: 8
; CHECK: .type_name: image1d_buffer_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: d
; CHECK: .offset: 24
; CHECK: .size: 8
; CHECK: .type_name: image2d_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: e
; CHECK: .offset: 32
; CHECK: .size: 8
; CHECK: .type_name: image2d_array_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: f
; CHECK: .offset: 40
; CHECK: .size: 8
; CHECK: .type_name: image2d_array_depth_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: g
; CHECK: .offset: 48
; CHECK: .size: 8
; CHECK: .type_name: image2d_array_msaa_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: h
; CHECK: .offset: 56
; CHECK: .size: 8
; CHECK: .type_name: image2d_array_msaa_depth_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: i
; CHECK: .offset: 64
; CHECK: .size: 8
; CHECK: .type_name: image2d_depth_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: j
; CHECK: .offset: 72
; CHECK: .size: 8
; CHECK: .type_name: image2d_msaa_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: k
; CHECK: .offset: 80
; CHECK: .size: 8
; CHECK: .type_name: image2d_msaa_depth_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
; CHECK: - .address_space: global
; CHECK: .name: l
; CHECK: .offset: 88
; CHECK: .size: 8
; CHECK: .type_name: image3d_t
; CHECK: .value_kind: image
; CHECK: .value_type: struct
define amdgpu_kernel void @test(%opencl.image1d_t addrspace(1)* %a,
%opencl.image1d_array_t addrspace(1)* %b,
%opencl.image1d_buffer_t addrspace(1)* %c,

25
test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll
View File

@ -7,16 +7,17 @@
; CHECK: ---
; CHECK: amdhsa.kernels:
; CHECK: - .max_flat_workgroup_size: 256
; CHECK: .kernarg_segment_size: 24
; CHECK: .private_segment_fixed_size: 0
; CHECK: .wavefront_size: 64
; CHECK: .symbol: test.kd
; CHECK: .name: test
; CHECK: .sgpr_count: 8
; CHECK: .kernarg_segment_align: 8
; CHECK: .vgpr_count: 6
; CHECK: .group_segment_fixed_size: 0
; CHECK: - .args:
; CHECK: .group_segment_fixed_size: 0
; CHECK: .kernarg_segment_align: 8
; CHECK: .kernarg_segment_size: 24
; CHECK: .max_flat_workgroup_size: 256
; CHECK: .name: test
; CHECK: .private_segment_fixed_size: 0
; CHECK: .sgpr_count: 8
; CHECK: .symbol: test.kd
; CHECK: .vgpr_count: 6
; CHECK: .wavefront_size: 64
define amdgpu_kernel void @test(
half addrspace(1)* %r,
half addrspace(1)* %a,
@ -29,11 +30,11 @@ entry:
ret void
}
; CHECK: .symbol: num_spilled_sgprs.kd
; CHECK: .name: num_spilled_sgprs
; GFX700: .sgpr_spill_count: 40
; GFX803: .sgpr_spill_count: 24
; GFX900: .sgpr_spill_count: 24
; CHECK: .symbol: num_spilled_sgprs.kd
define amdgpu_kernel void @num_spilled_sgprs(
i32 addrspace(1)* %out0, i32 addrspace(1)* %out1, [8 x i32],
i32 addrspace(1)* %out2, i32 addrspace(1)* %out3, [8 x i32],
@ -67,8 +68,8 @@ entry:
ret void
}
; CHECK: .symbol: num_spilled_vgprs.kd
; CHECK: .name: num_spilled_vgprs
; CHECK: .symbol: num_spilled_vgprs.kd
; CHECK: .vgpr_spill_count: 14
define amdgpu_kernel void @num_spilled_vgprs() #1 {
%val0 = load volatile float, float addrspace(1)* @var

92
test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
View File

@ -2,52 +2,52 @@
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
// CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels:
// CHECK-NEXT: - .max_flat_workgroup_size: 256
// CHECK-NEXT: .wavefront_size: 128
// CHECK-NEXT: .symbol: 'test_kernel@kd'
// CHECK-NEXT: .kernarg_segment_size: 8
// CHECK-NEXT: .private_segment_fixed_size: 32
// CHECK-NEXT: .name: test_kernel
// CHECK-NEXT: .language: OpenCL C
// CHECK-NEXT: .sgpr_count: 14
// CHECK-NEXT: .kernarg_segment_align: 64
// CHECK-NEXT: .vgpr_count: 40
// CHECK-NEXT: .group_segment_fixed_size: 16
// CHECK-NEXT: .language_version:
// CHECK-NEXT: - 2
// CHECK-NEXT: - 0
// CHECK-NEXT: .args:
// CHECK-NEXT: - .type_name: char
// CHECK-NEXT: .value_kind: by_value
// CHECK-NEXT: .offset: 1
// CHECK-NEXT: .size: 1
// CHECK-NEXT: .value_type: i8
// CHECK-NEXT: - .value_kind: hidden_global_offset_x
// CHECK-NEXT: .offset: 8
// CHECK-NEXT: .size: 8
// CHECK-NEXT: .value_type: i64
// CHECK-NEXT: - .value_kind: hidden_global_offset_y
// CHECK-NEXT: .offset: 8
// CHECK-NEXT: .size: 8
// CHECK-NEXT: .value_type: i64
// CHECK-NEXT: - .value_kind: hidden_global_offset_z
// CHECK-NEXT: .offset: 8
// CHECK-NEXT: .size: 8
// CHECK-NEXT: .value_type: i64
// CHECK-NEXT: - .value_kind: hidden_printf_buffer
// CHECK-NEXT: .offset: 8
// CHECK-NEXT: .size: 8
// CHECK-NEXT: .value_type: i8
// CHECK-NEXT: .address_space: global
// CHECK: amdhsa.version:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 0
// CHECK: amdhsa.printf:
// CHECK-NEXT: - '1:1:4:%d\n'
// CHECK-NEXt: - '2:1:8:%g\n'
// CHECK: .end_amdgpu_metadata
; CHECK: .amdgpu_metadata
; CHECK: amdhsa.kernels:
; CHECK-NEXT: - .args:
; CHECK-NEXT: - .offset: 1
; CHECK-NEXT: .size: 1
; CHECK-NEXT: .type_name: char
; CHECK-NEXT: .value_kind: by_value
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: .group_segment_fixed_size: 16
; CHECK-NEXT: .kernarg_segment_align: 64
; CHECK-NEXT: .kernarg_segment_size: 8
; CHECK-NEXT: .language: OpenCL C
; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2
; CHECK-NEXT: - 0
; CHECK-NEXT: .max_flat_workgroup_size: 256
; CHECK-NEXT: .name: test_kernel
; CHECK-NEXT: .private_segment_fixed_size: 32
; CHECK-NEXT: .sgpr_count: 14
; CHECK-NEXT: .symbol: 'test_kernel@kd'
; CHECK-NEXT: .vgpr_count: 40
; CHECK-NEXT: .wavefront_size: 128
; CHECK-NEXT: amdhsa.printf:
; CHECK-NEXT: - '1:1:4:%d\n'
; CHECK-NEXT: - '2:1:8:%g\n'
; CHECK-NEXT: amdhsa.version:
; CHECK-NEXT: - 1
; CHECK-NEXT: - 0
; CHECK: .end_amdgpu_metadata
.amdgpu_metadata
amdhsa.version:
- 1

63
test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
View File

@ -2,37 +2,38 @@
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
// CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels:
// CHECK: - .max_flat_workgroup_size: 256
// CHECK: .wavefront_size: 128
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .reqd_workgroup_size:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 2
// CHECK-NEXT: - 4
// CHECK: .kernarg_segment_size: 8
// CHECK: .private_segment_fixed_size: 32
// CHECK: .workgroup_size_hint:
// CHECK-NEXT: - 8
// CHECK-NEXT: - 16
// CHECK-NEXT: - 32
// CHECK: .name: test_kernel
// CHECK: .language: OpenCL C
// CHECK: .sgpr_count: 14
// CHECK: .kernarg_segment_align: 64
// CHECK: .vgpr_count: 40
// CHECK: .language_version:
// CHECK-NEXT: - 2
// CHECK-NEXT: - 0
// CHECK: .vec_type_hint: int
// CHECK: amdhsa.version:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 0
// CHECK: amdhsa.printf:
// CHECK: - '1:1:4:%d\n'
// CHECK: - '2:1:8:%g\n'
// CHECK: .end_amdgpu_metadata
// CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels:
// CHECK: - .group_segment_fixed_size: 16
// CHECK: .kernarg_segment_align: 64
// CHECK: .kernarg_segment_size: 8
// CHECK: .language: OpenCL C
// CHECK: .language_version:
// CHECK-NEXT: - 2
// CHECK-NEXT: - 0
// CHECK: .max_flat_workgroup_size: 256
// CHECK: .name: test_kernel
// CHECK: .private_segment_fixed_size: 32
// CHECK: .reqd_workgroup_size:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 2
// CHECK-NEXT: - 4
// CHECK: .sgpr_count: 14
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .vec_type_hint: int
// CHECK: .vgpr_count: 40
// CHECK: .wavefront_size: 128
// CHECK: .workgroup_size_hint:
// CHECK-NEXT: - 8
// CHECK-NEXT: - 16
// CHECK-NEXT: - 32
// CHECK: amdhsa.printf:
// CHECK: - '1:1:4:%d\n'
// CHECK: - '2:1:8:%g\n'
// CHECK: amdhsa.version:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 0
// CHECK: .end_amdgpu_metadata
.amdgpu_metadata
amdhsa.version:
- 1

34
test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
View File

@ -2,23 +2,23 @@
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
// CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels:
// CHECK: - .sgpr_count: 40
// CHECK: .max_flat_workgroup_size: 256
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .kernarg_segment_size: 24
// CHECK: .group_segment_fixed_size: 24
// CHECK: .private_segment_fixed_size: 16
// CHECK: .vgpr_count: 14
// CHECK: .vgpr_spill_count: 1
// CHECK: .kernarg_segment_align: 16
// CHECK: .sgpr_spill_count: 1
// CHECK: .wavefront_size: 64
// CHECK: .name: test_kernel
// CHECK: amdhsa.version:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 0
// CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels:
// CHECK: - .group_segment_fixed_size: 24
// CHECK: .kernarg_segment_align: 16
// CHECK: .kernarg_segment_size: 24
// CHECK: .max_flat_workgroup_size: 256
// CHECK: .name: test_kernel
// CHECK: .private_segment_fixed_size: 16
// CHECK: .sgpr_count: 40
// CHECK: .sgpr_spill_count: 1
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .vgpr_count: 14
// CHECK: .vgpr_spill_count: 1
// CHECK: .wavefront_size: 64
// CHECK: amdhsa.version:
// CHECK-NEXT: - 1
// CHECK-NEXT: - 0
.amdgpu_metadata
amdhsa.version:
- 1

52
test/MC/AMDGPU/hsa-v3.s
View File

@ -249,29 +249,29 @@ v_mov_b32_e32 v16, s3
.max_flat_workgroup_size: 256
.end_amdgpu_metadata
// ASM: .amdgpu_metadata
// ASM: amdhsa.kernels:
// ASM: - .sgpr_count: 14
// ASM: .max_flat_workgroup_size: 256
// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'
// ASM: .kernarg_segment_size: 8
// ASM: .group_segment_fixed_size: 16
// ASM: .private_segment_fixed_size: 32
// ASM: .vgpr_count: 40
// ASM: .kernarg_segment_align: 64
// ASM: .wavefront_size: 128
// ASM: .name: amd_kernel_code_t_test_all
// ASM: - .sgpr_count: 14
// ASM: .max_flat_workgroup_size: 256
// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'
// ASM: .kernarg_segment_size: 8
// ASM: .group_segment_fixed_size: 16
// ASM: .private_segment_fixed_size: 32
// ASM: .vgpr_count: 40
// ASM: .kernarg_segment_align: 64
// ASM: .wavefront_size: 128
// ASM: .name: amd_kernel_code_t_minimal
// ASM: amdhsa.version:
// ASM-NEXT: - 3
// ASM-NEXT: - 0
// ASM: .end_amdgpu_metadata
// ASM: .amdgpu_metadata
// ASM: amdhsa.kernels:
// ASM: - .group_segment_fixed_size: 16
// ASM: .kernarg_segment_align: 64
// ASM: .kernarg_segment_size: 8
// ASM: .max_flat_workgroup_size: 256
// ASM: .name: amd_kernel_code_t_test_all
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: - .group_segment_fixed_size: 16
// ASM: .kernarg_segment_align: 64
// ASM: .kernarg_segment_size: 8
// ASM: .max_flat_workgroup_size: 256
// ASM: .name: amd_kernel_code_t_minimal
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:
// ASM-NEXT: - 3
// ASM-NEXT: - 0
// ASM: .end_amdgpu_metadata

17
tools/llvm-readobj/ELFDumper.cpp
View File

@ -3917,29 +3917,24 @@ static AMDGPUNote getAMDGPUNote(uint32_t NoteType, ArrayRef<uint8_t> Desc) {
switch (NoteType) {
default:
return {"", ""};
case ELF::NT_AMDGPU_METADATA:
case ELF::NT_AMDGPU_METADATA: {
auto MsgPackString =
StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());
msgpack::Reader MsgPackReader(MsgPackString);
auto OptMsgPackNodeOrErr = msgpack::Node::read(MsgPackReader);
if (errorToBool(OptMsgPackNodeOrErr.takeError()))
msgpack::Document MsgPackDoc;
if (!MsgPackDoc.readFromBlob(MsgPackString, /*Multi=*/false))
return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
auto &OptMsgPackNode = *OptMsgPackNodeOrErr;
if (!OptMsgPackNode)
return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
auto &MsgPackNode = *OptMsgPackNode;
AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
if (!Verifier.verify(*MsgPackNode))
if (!Verifier.verify(MsgPackDoc.getRoot()))
return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
std::string HSAMetadataString;
raw_string_ostream StrOS(HSAMetadataString);
yaml::Output YOut(StrOS);
YOut << MsgPackNode;
MsgPackDoc.toYAML(StrOS);
return {"AMDGPU Metadata", StrOS.str()};
}
}
}
template <class ELFT>