[llvm] r356081 - [AMDGPU] Switched HSA metadata to use MsgPackDocument
Tim Renouf via llvm-commits
llvm-commits at lists.llvm.org
Wed Mar 13 11:55:51 PDT 2019
Author: tpr
Date: Wed Mar 13 11:55:50 2019
New Revision: 356081
URL: http://llvm.org/viewvc/llvm-project?rev=356081&view=rev
Log:
[AMDGPU] Switched HSA metadata to use MsgPackDocument
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
Modified:
llvm/trunk/include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h
llvm/trunk/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll
llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll
llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
llvm/trunk/test/MC/AMDGPU/hsa-v3.s
llvm/trunk/tools/llvm-readobj/ELFDumper.cpp
Modified: llvm/trunk/include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h (original)
+++ llvm/trunk/include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h Wed Mar 13 11:55:50 2019
@@ -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
Modified: llvm/trunk/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp (original)
+++ llvm/trunk/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp Wed Mar 13 11:55:50 2019
@@ -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)
- return false;
- auto &Scalar = *ScalarPtr;
- // Do not output extraneous tags for types we know from the spec.
- Scalar.IgnoreTag = true;
- if (Scalar.getScalarKind() != SKind) {
+ msgpack::DocNode &Node, msgpack::Type SKind,
+ function_ref<bool(msgpack::DocNode &)> verifyValue) {
+ if (!Node.isScalar())
+ return false;
+ 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(
}))
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(
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(
}))
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(
}))
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(
}))
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(msgp
}))
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(msgp
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);
});
}))
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp Wed Mar 13 11:55:50 2019
@@ -489,20 +489,16 @@ void MetadataStreamerV3::dump(StringRef
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::getTypeN
}
}
-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(cons
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::emitKernelLangu
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(
// 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(c
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(c
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::emitHiddenKerne
}
}
-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
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h Wed Mar 13 11:55:50 2019
@@ -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:
Modified: llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp Wed Mar 13 11:55:50 2019
@@ -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::EmitHSAMetada
}
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::EmitHSAMet
}
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::EmitISAVer
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::EmitHSAMet
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;
Modified: llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h (original)
+++ llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h Wed Mar 13 11:55:50 2019
@@ -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;
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll Wed Mar 13 11:55:50 2019
@@ -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
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll Wed Mar 13 11:55:50 2019
@@ -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 {
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll Wed Mar 13 11:55:50 2019
@@ -16,576 +16,581 @@
@__test_block_invoke_kernel_runtime_handle = external addrspace(1) externally_initialized constant i8 addrspace(1)*
-; CHECK: ---
-; CHECK: amdhsa.kernels:
-; CHECK: .symbol: test_char.kd
-; CHECK: .name: test_char
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK-NOT: .value_kind: hidden_default_queue
-; CHECK-NOT: .value_kind: hidden_completion_action
+; CHECK: ---
+; CHECK-NEXT: amdhsa.kernels:
+; CHECK-NEXT: - .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-NOT: .value_kind: hidden_default_queue
+; CHECK-NOT: .value_kind: hidden_completion_action
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_char
+; CHECK: .symbol: test_char.kd
define amdgpu_kernel void @test_char(i8 %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
!kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_ushort2.kd
-; CHECK: .name: test_ushort2
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: ushort2
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: u16
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; 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: 4
+; CHECK-NEXT: .type_name: ushort2
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: u16
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .offset: 40
+; 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_ushort2
+; CHECK: .symbol: test_ushort2.kd
define amdgpu_kernel void @test_ushort2(<2 x i16> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
!kernel_arg_base_type !10 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_int3.kd
-; CHECK: .name: test_int3
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int3
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 16
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 40
-; 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: 16
+; CHECK-NEXT: .type_name: int3
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 24
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 32
+; 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: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_int3
+; CHECK: .symbol: test_int3.kd
define amdgpu_kernel void @test_int3(<3 x i32> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
!kernel_arg_base_type !11 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_ulong4.kd
-; CHECK: .name: test_ulong4
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: ulong4
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 32
-; CHECK-NEXT: .value_type: u64
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 40
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 56
-; 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: 32
+; CHECK-NEXT: .type_name: ulong4
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: u64
+; CHECK-NEXT: - .offset: 32
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 48
+; 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: 56
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_ulong4
+; CHECK: .symbol: test_ulong4.kd
define amdgpu_kernel void @test_ulong4(<4 x i64> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
!kernel_arg_base_type !12 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_half8.kd
-; CHECK: .name: test_half8
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: half8
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 16
-; CHECK-NEXT: .value_type: f16
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 40
-; 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: 16
+; CHECK-NEXT: .type_name: half8
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: f16
+; CHECK-NEXT: - .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 24
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 32
+; 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: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_half8
+; CHECK: .symbol: test_half8.kd
define amdgpu_kernel void @test_half8(<8 x half> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
!kernel_arg_base_type !13 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_float16.kd
-; CHECK: .name: test_float16
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: float16
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 64
-; CHECK-NEXT: .value_type: f32
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 64
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 72
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 80
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 88
-; 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: 64
+; CHECK-NEXT: .type_name: float16
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: f32
+; CHECK-NEXT: - .offset: 64
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 72
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 80
+; 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: 88
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_float16
+; CHECK: .symbol: test_float16.kd
define amdgpu_kernel void @test_float16(<16 x float> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
!kernel_arg_base_type !14 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_double16.kd
-; CHECK: .name: test_double16
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: double16
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 128
-; CHECK-NEXT: .value_type: f64
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 128
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 136
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 144
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 152
-; 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: 128
+; CHECK-NEXT: .type_name: double16
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: f64
+; CHECK-NEXT: - .offset: 128
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 136
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 144
+; 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: 152
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_double16
+; CHECK: .symbol: test_double16.kd
define amdgpu_kernel void @test_double16(<16 x double> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
!kernel_arg_base_type !15 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_pointer.kd
-; CHECK: .name: test_pointer
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: global
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_pointer
+; CHECK: .symbol: test_pointer.kd
define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16
!kernel_arg_base_type !16 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_image.kd
-; CHECK: .name: test_image
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: image2d_t
-; CHECK-NEXT: .value_kind: image
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: global
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: image2d_t
+; CHECK-NEXT: .value_kind: image
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_image
+; CHECK: .symbol: test_image.kd
define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
!kernel_arg_base_type !17 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_sampler.kd
-; CHECK: .name: test_sampler
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: sampler_t
-; CHECK-NEXT: .value_kind: sampler
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; 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: 4
+; CHECK-NEXT: .type_name: sampler_t
+; CHECK-NEXT: .value_kind: sampler
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_sampler
+; CHECK: .symbol: test_sampler.kd
define amdgpu_kernel void @test_sampler(i32 %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18
!kernel_arg_base_type !18 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_queue.kd
-; CHECK: .name: test_queue
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: queue_t
-; CHECK-NEXT: .value_kind: queue
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: global
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: queue_t
+; CHECK-NEXT: .value_kind: queue
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_queue
+; CHECK: .symbol: test_queue.kd
define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
!kernel_arg_base_type !19 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_struct.kd
-; CHECK: .name: test_struct
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: struct A
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: private
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: private
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: struct A
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_struct
+; CHECK: .symbol: test_struct.kd
define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
!kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_i128.kd
-; CHECK: .name: test_i128
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: i128
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 16
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 40
-; 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: 16
+; CHECK-NEXT: .type_name: i128
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: struct
+; CHECK-NEXT: - .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 24
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 32
+; 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: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_i128
+; CHECK: .symbol: test_i128.kd
define amdgpu_kernel void @test_i128(i128 %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
!kernel_arg_base_type !21 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_multi_arg.kd
-; CHECK: .name: test_multi_arg
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .type_name: short2
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 4
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i16
-; CHECK-NEXT: .name: b
-; CHECK-NEXT: - .type_name: char3
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 8
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .name: c
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 40
-; 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: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .name: b
+; CHECK-NEXT: .offset: 4
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: short2
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i16
+; CHECK-NEXT: - .name: c
+; CHECK-NEXT: .offset: 8
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: char3
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 24
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 32
+; 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: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_multi_arg
+; CHECK: .symbol: test_multi_arg.kd
define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c)
!kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
!kernel_arg_base_type !24 !kernel_arg_type_qual !25 {
ret void
}
-; CHECK: .symbol: test_addr_space.kd
-; CHECK: .name: test_addr_space
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: g
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: c
-; CHECK-NEXT: .offset: 8
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: constant
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: l
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .pointee_align: 4
-; CHECK-NEXT: .address_space: local
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: g
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .address_space: constant
+; CHECK-NEXT: .name: c
+; CHECK-NEXT: .offset: 8
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: l
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .pointee_align: 4
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_addr_space
+; CHECK: .symbol: test_addr_space.kd
define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
i32 addrspace(4)* %c,
i32 addrspace(3)* %l)
@@ -594,55 +599,55 @@ define amdgpu_kernel void @test_addr_spa
ret void
}
-; CHECK: .symbol: test_type_qual.kd
-; CHECK: .name: test_type_qual
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .is_volatile: true
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: b
-; CHECK-NEXT: .is_const: true
-; CHECK-NEXT: .offset: 8
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .is_restrict: true
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
-; CHECK-NEXT: .value_kind: pipe
-; CHECK-NEXT: .name: c
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .is_pipe: true
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .is_volatile: true
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .is_const: true
+; CHECK-NEXT: .is_restrict: true
+; CHECK-NEXT: .name: b
+; CHECK-NEXT: .offset: 8
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .is_pipe: true
+; CHECK-NEXT: .name: c
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)*'
+; CHECK-NEXT: .value_kind: pipe
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_type_qual
+; CHECK: .symbol: test_type_qual.kd
define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a,
i32 addrspace(1)* %b,
%opencl.pipe_t addrspace(1)* %c)
@@ -651,54 +656,54 @@ define amdgpu_kernel void @test_type_qua
ret void
}
-; CHECK: .symbol: test_access_qual.kd
-; CHECK: .name: test_access_qual
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: image1d_t
-; CHECK-NEXT: .value_kind: image
-; CHECK-NEXT: .name: ro
-; CHECK-NEXT: .access: read_only
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: image2d_t
-; CHECK-NEXT: .value_kind: image
-; CHECK-NEXT: .name: wo
-; CHECK-NEXT: .access: write_only
-; CHECK-NEXT: .offset: 8
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: image3d_t
-; CHECK-NEXT: .value_kind: image
-; CHECK-NEXT: .name: rw
-; CHECK-NEXT: .access: read_write
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .access: read_only
+; CHECK-NEXT: .address_space: global
+; CHECK-NEXT: .name: ro
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: image1d_t
+; CHECK-NEXT: .value_kind: image
+; CHECK-NEXT: .value_type: struct
+; CHECK-NEXT: - .access: write_only
+; CHECK-NEXT: .address_space: global
+; CHECK-NEXT: .name: wo
+; CHECK-NEXT: .offset: 8
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: image2d_t
+; CHECK-NEXT: .value_kind: image
+; CHECK-NEXT: .value_type: struct
+; CHECK-NEXT: - .access: read_write
+; CHECK-NEXT: .address_space: global
+; CHECK-NEXT: .name: rw
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: image3d_t
+; CHECK-NEXT: .value_kind: image
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_access_qual
+; CHECK: .symbol: test_access_qual.kd
define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro,
%opencl.image2d_t addrspace(1)* %wo,
%opencl.image3d_t addrspace(1)* %rw)
@@ -707,300 +712,300 @@ define amdgpu_kernel void @test_access_q
ret void
}
-; CHECK: .symbol: test_vec_type_hint_half.kd
-; CHECK: .name: test_vec_type_hint_half
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: half
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_half
+; CHECK: .symbol: test_vec_type_hint_half.kd
+; CHECK: .vec_type_hint: half
define amdgpu_kernel void @test_vec_type_hint_half(i32 %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 !vec_type_hint !26 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_float.kd
-; CHECK: .name: test_vec_type_hint_float
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: float
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_float
+; CHECK: .symbol: test_vec_type_hint_float.kd
+; CHECK: .vec_type_hint: float
define amdgpu_kernel void @test_vec_type_hint_float(i32 %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 !vec_type_hint !27 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_double.kd
-; CHECK: .name: test_vec_type_hint_double
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: double
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_double
+; CHECK: .symbol: test_vec_type_hint_double.kd
+; CHECK: .vec_type_hint: double
define amdgpu_kernel void @test_vec_type_hint_double(i32 %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 !vec_type_hint !28 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_char.kd
-; CHECK: .name: test_vec_type_hint_char
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: char
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_char
+; CHECK: .symbol: test_vec_type_hint_char.kd
+; CHECK: .vec_type_hint: char
define amdgpu_kernel void @test_vec_type_hint_char(i32 %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 !vec_type_hint !29 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_short.kd
-; CHECK: .name: test_vec_type_hint_short
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: short
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_short
+; CHECK: .symbol: test_vec_type_hint_short.kd
+; CHECK: .vec_type_hint: short
define amdgpu_kernel void @test_vec_type_hint_short(i32 %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 !vec_type_hint !30 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_long.kd
-; CHECK: .name: test_vec_type_hint_long
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: long
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_long
+; CHECK: .symbol: test_vec_type_hint_long.kd
+; CHECK: .vec_type_hint: long
define amdgpu_kernel void @test_vec_type_hint_long(i32 %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 !vec_type_hint !31 {
ret void
}
-; CHECK: .symbol: test_vec_type_hint_unknown.kd
-; CHECK: .name: test_vec_type_hint_unknown
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: unknown
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_vec_type_hint_unknown
+; CHECK: .symbol: test_vec_type_hint_unknown.kd
+; CHECK: .vec_type_hint: unknown
define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %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 !vec_type_hint !32 {
ret void
}
-; CHECK: .reqd_workgroup_size:
-; CHECK-NEXT: - 1
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 4
-; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd
-; CHECK: .name: test_reqd_wgs_vec_type_hint
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: int
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_reqd_wgs_vec_type_hint
+; CHECK: .reqd_workgroup_size:
+; CHECK-NEXT: - 1
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 4
+; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd
+; CHECK: .vec_type_hint: int
define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %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 !vec_type_hint !5
@@ -1008,41 +1013,41 @@ define amdgpu_kernel void @test_reqd_wgs
ret void
}
-; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
-; CHECK: .workgroup_size_hint:
-; CHECK-NEXT: - 8
-; CHECK-NEXT: - 16
-; CHECK-NEXT: - 32
-; CHECK: .name: test_wgs_hint_vec_type_hint
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: int
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i32
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
-; CHECK: .vec_type_hint: uint4
+; CHECK: - .args:
+; CHECK-NEXT: - .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: int
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_wgs_hint_vec_type_hint
+; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
+; CHECK: .vec_type_hint: uint4
+; CHECK: .workgroup_size_hint:
+; CHECK-NEXT: - 8
+; CHECK-NEXT: - 16
+; CHECK-NEXT: - 32
define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %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 !vec_type_hint !7
@@ -1050,147 +1055,147 @@ define amdgpu_kernel void @test_wgs_hint
ret void
}
-; CHECK: .symbol: test_arg_ptr_to_ptr.kd
-; CHECK: .name: test_arg_ptr_to_ptr
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'int addrspace(5)* addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .address_space: global
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'int addrspace(5)* addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_arg_ptr_to_ptr
+; CHECK: .symbol: test_arg_ptr_to_ptr.kd
define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a)
!kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
!kernel_arg_base_type !80 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_arg_struct_contains_ptr.kd
-; CHECK: .name: test_arg_struct_contains_ptr
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: struct B
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: private
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: private
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: struct B
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_arg_struct_contains_ptr
+; CHECK: .symbol: test_arg_struct_contains_ptr.kd
define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B addrspace(5)* byval %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82
!kernel_arg_base_type !82 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_arg_vector_of_ptr.kd
-; CHECK: .name: test_arg_vector_of_ptr
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 16
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 40
-; 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: 16
+; CHECK-NEXT: .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: i32
+; CHECK-NEXT: - .offset: 16
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 24
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 32
+; 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: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_arg_vector_of_ptr
+; CHECK: .symbol: test_arg_vector_of_ptr.kd
define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a)
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
!kernel_arg_base_type !83 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: test_arg_unknown_builtin_type.kd
-; CHECK: .name: test_arg_unknown_builtin_type
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: clk_event_t
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .address_space: global
-; 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_printf_buffer
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: clk_event_t
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: struct
+; 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_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_arg_unknown_builtin_type
+; CHECK: .symbol: test_arg_unknown_builtin_type.kd
define amdgpu_kernel void @test_arg_unknown_builtin_type(
%opencl.clk_event_t addrspace(1)* %a)
!kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84
@@ -1198,85 +1203,85 @@ define amdgpu_kernel void @test_arg_unkn
ret void
}
-; CHECK: .symbol: test_pointee_align.kd
-; CHECK: .name: test_pointee_align
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: 'long addrspace(5)*'
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .name: a
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: .address_space: global
-; CHECK-NEXT: - .type_name: 'char addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: b
-; CHECK-NEXT: .offset: 8
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 1
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .type_name: 'char2 addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: c
-; CHECK-NEXT: .offset: 12
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 2
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .type_name: 'char3 addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: d
-; CHECK-NEXT: .offset: 16
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 4
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .type_name: 'char4 addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: e
-; CHECK-NEXT: .offset: 20
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 4
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .type_name: 'char8 addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: f
-; CHECK-NEXT: .offset: 24
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 8
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .type_name: 'char16 addrspace(5)*'
-; CHECK-NEXT: .value_kind: dynamic_shared_pointer
-; CHECK-NEXT: .name: g
-; CHECK-NEXT: .offset: 28
-; CHECK-NEXT: .size: 4
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .pointee_align: 16
-; CHECK-NEXT: .address_space: local
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 40
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 56
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .address_space: global
+; CHECK-NEXT: .name: a
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .type_name: 'long addrspace(5)*'
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: b
+; CHECK-NEXT: .offset: 8
+; CHECK-NEXT: .pointee_align: 1
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: c
+; CHECK-NEXT: .offset: 12
+; CHECK-NEXT: .pointee_align: 2
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: d
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .pointee_align: 4
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: e
+; CHECK-NEXT: .offset: 20
+; CHECK-NEXT: .pointee_align: 4
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: f
+; CHECK-NEXT: .offset: 24
+; CHECK-NEXT: .pointee_align: 8
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .address_space: local
+; CHECK-NEXT: .name: g
+; CHECK-NEXT: .offset: 28
+; CHECK-NEXT: .pointee_align: 16
+; CHECK-NEXT: .size: 4
+; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
+; CHECK-NEXT: .value_kind: dynamic_shared_pointer
+; CHECK-NEXT: .value_type: i8
+; CHECK-NEXT: - .offset: 32
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 48
+; 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: 56
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: test_pointee_align
+; CHECK: .symbol: test_pointee_align.kd
define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
i8 addrspace(3)* %b,
<2 x i8> addrspace(3)* %c,
@@ -1289,37 +1294,37 @@ define amdgpu_kernel void @test_pointee_
ret void
}
-; CHECK: .symbol: __test_block_invoke_kernel.kd
-; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
-; CHECK: .name: __test_block_invoke_kernel
-; CHECK: .language: OpenCL C
-; CHECK: .language_version:
-; CHECK-NEXT: - 2
-; CHECK-NEXT: - 0
-; CHECK: .args:
-; CHECK-NEXT: - .type_name: __block_literal
-; CHECK-NEXT: .value_kind: by_value
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 25
-; CHECK-NEXT: .value_type: struct
-; CHECK-NEXT: .name: arg
-; CHECK-NEXT: - .value_kind: hidden_global_offset_x
-; CHECK-NEXT: .offset: 32
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_y
-; CHECK-NEXT: .offset: 40
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_global_offset_z
-; CHECK-NEXT: .offset: 48
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i64
-; CHECK-NEXT: - .value_kind: hidden_printf_buffer
-; CHECK-NEXT: .offset: 56
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i8
-; CHECK-NEXT: .address_space: global
+; CHECK: - .args:
+; CHECK-NEXT: - .name: arg
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 25
+; CHECK-NEXT: .type_name: __block_literal
+; CHECK-NEXT: .value_kind: by_value
+; CHECK-NEXT: .value_type: struct
+; CHECK-NEXT: - .offset: 32
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_x
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 40
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_global_offset_y
+; CHECK-NEXT: .value_type: i64
+; CHECK-NEXT: - .offset: 48
+; 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: 56
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: hidden_printf_buffer
+; CHECK-NEXT: .value_type: i8
+; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
+; CHECK: .language: OpenCL C
+; CHECK-NEXT: .language_version:
+; CHECK-NEXT: - 2
+; CHECK-NEXT: - 0
+; CHECK: .name: __test_block_invoke_kernel
+; CHECK: .symbol: __test_block_invoke_kernel.kd
define amdgpu_kernel void @__test_block_invoke_kernel(
<{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
@@ -1327,70 +1332,70 @@ define amdgpu_kernel void @__test_block_
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_printf_buffer
-; 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_printf_buffer
+; 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) #1
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
!kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
ret void
}
-; CHECK: .symbol: unknown_addrspace_kernarg.kd
-; CHECK: .name: unknown_addrspace_kernarg
-; CHECK: .args:
-; CHECK-NEXT: .value_kind: global_buffer
-; CHECK-NEXT: .offset: 0
-; CHECK-NEXT: .size: 8
-; CHECK-NEXT: .value_type: i32
-; CHECK-NEXT: .name: ptr
+; CHECK: - .args:
+; CHECK-NEXT: - .name: ptr
+; CHECK-NEXT: .offset: 0
+; CHECK-NEXT: .size: 8
+; CHECK-NEXT: .value_kind: global_buffer
+; CHECK-NEXT: .value_type: i32
+; CHECK: .name: unknown_addrspace_kernarg
+; CHECK: .symbol: unknown_addrspace_kernarg.kd
define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
ret void
}
-; 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: amdhsa.version:
+; CHECK-NEXT: - 1
+; CHECK-NEXT: - 0
attributes #0 = { "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
attributes #1 = { "calls-enqueue-kernel" }
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll Wed Mar 13 11:55:50 2019
@@ -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,
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll Wed Mar 13 11:55:50 2019
@@ -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,
Modified: llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll Wed Mar 13 11:55:50 2019
@@ -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
Modified: llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s (original)
+++ llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s Wed Mar 13 11:55:50 2019
@@ -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 at 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 at 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
Modified: llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s (original)
+++ llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s Wed Mar 13 11:55:50 2019
@@ -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 at 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 at 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
Modified: llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s (original)
+++ llvm/trunk/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s Wed Mar 13 11:55:50 2019
@@ -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 at 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 at 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
Modified: llvm/trunk/test/MC/AMDGPU/hsa-v3.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/hsa-v3.s?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/hsa-v3.s (original)
+++ llvm/trunk/test/MC/AMDGPU/hsa-v3.s Wed Mar 13 11:55:50 2019
@@ -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 at 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 at 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 at 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 at kd'
+// ASM: .vgpr_count: 40
+// ASM: .wavefront_size: 128
+// ASM: amdhsa.version:
+// ASM-NEXT: - 3
+// ASM-NEXT: - 0
+// ASM: .end_amdgpu_metadata
Modified: llvm/trunk/tools/llvm-readobj/ELFDumper.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/tools/llvm-readobj/ELFDumper.cpp?rev=356081&r1=356080&r2=356081&view=diff
==============================================================================
--- llvm/trunk/tools/llvm-readobj/ELFDumper.cpp (original)
+++ llvm/trunk/tools/llvm-readobj/ELFDumper.cpp Wed Mar 13 11:55:50 2019
@@ -3917,29 +3917,24 @@ static AMDGPUNote getAMDGPUNote(uint32_t
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>
More information about the llvm-commits
mailing list