Index: include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h =================================================================== --- include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h +++ include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h @@ -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 @@ class MetadataVerifier { bool Strict; - bool verifyScalar(msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind, - function_ref verifyValue = {}); - bool verifyInteger(msgpack::Node &Node); - bool verifyArray(msgpack::Node &Node, - function_ref verifyNode, + bool verifyScalar(msgpack::DocNode &Node, msgpack::Type SKind, + function_ref verifyValue = {}); + bool verifyInteger(msgpack::DocNode &Node); + bool verifyArray(msgpack::DocNode &Node, + function_ref verifyNode, Optional Size = None); - bool verifyEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required, - function_ref verifyNode); + bool verifyEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required, + function_ref verifyNode); bool - verifyScalarEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required, - msgpack::ScalarNode::ScalarKind SKind, - function_ref verifyValue = {}); - bool verifyIntegerEntry(msgpack::MapNode &MapNode, StringRef Key, + verifyScalarEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required, + msgpack::Type SKind, + function_ref 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 @@ /// 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 Index: lib/BinaryFormat/AMDGPUMetadataVerifier.cpp =================================================================== --- lib/BinaryFormat/AMDGPUMetadataVerifier.cpp +++ lib/BinaryFormat/AMDGPUMetadataVerifier.cpp @@ -20,98 +20,92 @@ namespace V3 { bool MetadataVerifier::verifyScalar( - msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind, - function_ref verifyValue) { - auto ScalarPtr = dyn_cast(&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 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 verifyNode, + msgpack::DocNode &Node, function_ref verifyNode, Optional Size) { - auto ArrayPtr = dyn_cast(&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 verifyNode) { + msgpack::MapDocNode &MapNode, StringRef Key, bool Required, + function_ref 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 verifyValue) { - return verifyEntry(MapNode, Key, Required, [=](msgpack::Node &Node) { + msgpack::MapDocNode &MapNode, StringRef Key, bool Required, + msgpack::Type SKind, + function_ref 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(&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(SNode.getString()) .Case("by_value", true) .Case("global_buffer", true) @@ -131,8 +125,8 @@ })) return false; if (!verifyScalarEntry(ArgsMap, ".value_type", true, - msgpack::ScalarNode::SK_String, - [](msgpack::ScalarNode &SNode) { + msgpack::Type::String, + [](msgpack::DocNode &SNode) { return StringSwitch(SNode.getString()) .Case("struct", true) .Case("i8", true) @@ -152,8 +146,8 @@ 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(SNode.getString()) .Case("private", true) .Case("global", true) @@ -165,8 +159,8 @@ })) return false; if (!verifyScalarEntry(ArgsMap, ".access", false, - msgpack::ScalarNode::SK_String, - [](msgpack::ScalarNode &SNode) { + msgpack::Type::String, + [](msgpack::DocNode &SNode) { return StringSwitch(SNode.getString()) .Case("read_only", true) .Case("write_only", true) @@ -175,8 +169,8 @@ })) return false; if (!verifyScalarEntry(ArgsMap, ".actual_access", false, - msgpack::ScalarNode::SK_String, - [](msgpack::ScalarNode &SNode) { + msgpack::Type::String, + [](msgpack::DocNode &SNode) { return StringSwitch(SNode.getString()) .Case("read_only", true) .Case("write_only", true) @@ -185,36 +179,35 @@ })) 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(&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(SNode.getString()) .Case("OpenCL C", true) .Case("OpenCL C++", true) @@ -226,41 +219,41 @@ })) 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 @@ return true; } -bool MetadataVerifier::verify(msgpack::Node &HSAMetadataRoot) { - auto RootMapPtr = dyn_cast(&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); }); })) Index: lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h =================================================================== --- lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h +++ lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h @@ -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 @@ class MetadataStreamerV3 final : public MetadataStreamer { private: - std::shared_ptr HSAMetadataRoot = - std::make_shared(); + std::unique_ptr HSAMetadataDoc = + std::unique_ptr(new msgpack::Document); void dump(StringRef HSAMetadataString) const; @@ -69,41 +69,39 @@ std::string getTypeName(Type *Ty, bool Signed) const; - std::shared_ptr - getWorkGroupDimensions(MDNode *Node) const; + msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const; - std::shared_ptr - 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 &getRootMetadata(StringRef Key) { - return (*cast(HSAMetadataRoot.get()))[Key]; + msgpack::DocNode &getRootMetadata(StringRef Key) { + return HSAMetadataDoc->getRoot().getMap(/*Convert=*/true)[Key]; } - std::shared_ptr &getHSAMetadataRoot() { - return HSAMetadataRoot; + msgpack::DocNode &getHSAMetadataRoot() { + return HSAMetadataDoc->getRoot(); } public: Index: lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp =================================================================== --- lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp +++ lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp @@ -505,20 +505,16 @@ void MetadataStreamerV3::verify(StringRef HSAMetadataString) const { errs() << "AMDGPU HSA Metadata Parser Test: "; - std::shared_ptr FromHSAMetadataString = - std::make_shared(); + 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) { @@ -652,23 +648,23 @@ } } -std::shared_ptr +msgpack::ArrayDocNode MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const { - auto Dims = std::make_shared(); + auto Dims = HSAMetadataDoc->getArrayNode(); if (Node->getNumOperands() != 3) return Dims; for (auto &Op : Node->operands()) - Dims->push_back(std::make_shared( - mdconst::extract(Op)->getZExtValue())); + Dims.push_back(Dims.getDocument()->getNode( + uint64_t(mdconst::extract(Op)->getZExtValue()))); return Dims; } void MetadataStreamerV3::emitVersion() { - auto Version = std::make_shared(); - Version->push_back(std::make_shared(V3::VersionMajor)); - Version->push_back(std::make_shared(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) { @@ -676,16 +672,16 @@ if (!Node) return; - auto Printf = std::make_shared(); + auto Printf = HSAMetadataDoc->getArrayNode(); for (auto Op : Node->operands()) if (Op->getNumOperands()) - Printf->push_back(std::make_shared( - cast(Op->getOperand(0))->getString())); - getRootMetadata("amdhsa.printf") = std::move(Printf); + Printf.push_back(Printf.getDocument()->getNode( + cast(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()) @@ -694,50 +690,53 @@ if (Op0->getNumOperands() <= 1) return; - Kern[".language"] = std::make_shared("OpenCL C"); - auto LanguageVersion = std::make_shared(); - LanguageVersion->push_back(std::make_shared( + Kern[".language"] = Kern.getDocument()->getNode("OpenCL C"); + auto LanguageVersion = Kern.getDocument()->getArrayNode(); + LanguageVersion.push_back(Kern.getDocument()->getNode( mdconst::extract(Op0->getOperand(0))->getZExtValue())); - LanguageVersion->push_back(std::make_shared( + LanguageVersion.push_back(Kern.getDocument()->getNode( mdconst::extract(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(getTypeName( - cast(Node->getOperand(0))->getType(), - mdconst::extract(Node->getOperand(1))->getZExtValue())); + Kern[".vec_type_hint"] = Kern.getDocument()->getNode( + getTypeName( + cast(Node->getOperand(0))->getType(), + mdconst::extract(Node->getOperand(1))->getZExtValue()), + /*Copy=*/true); } if (Func.hasFnAttribute("runtime-handle")) { - Kern[".device_enqueue_symbol"] = std::make_shared( - 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(); + 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); @@ -745,26 +744,26 @@ // 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; @@ -821,36 +820,35 @@ 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(); - auto &Arg = *ArgPtr; + auto Arg = Args.getDocument()->getMapNode(); if (!Name.empty()) - Arg[".name"] = std::make_shared(Name); + Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true); if (!TypeName.empty()) - Arg[".type_name"] = std::make_shared(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(Size); + Arg[".size"] = Arg.getDocument()->getNode(Size); Offset = alignTo(Offset, Align); - Arg[".offset"] = std::make_shared(Offset); + Arg[".offset"] = Arg.getDocument()->getNode(Offset); Offset += Size; - Arg[".value_kind"] = std::make_shared(ValueKind); + Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true); Arg[".value_type"] = - std::make_shared(getValueType(Ty, BaseTypeName)); + Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true); if (PointeeAlign) - Arg[".pointee_align"] = std::make_shared(PointeeAlign); + Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign); if (auto PtrTy = dyn_cast(Ty)) if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace())) - Arg[".address_space"] = std::make_shared(*Qualifier); + Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true); if (auto AQ = getAccessQualifier(AccQual)) - Arg[".access"] = std::make_shared(*AQ); + Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true); // TODO: Emit Arg[".actual_access"]. @@ -858,21 +856,21 @@ TypeQual.split(SplitTypeQuals, " ", -1, false); for (StringRef Key : SplitTypeQuals) { if (Key == "const") - Arg[".is_const"] = std::make_shared(true); + Arg[".is_const"] = Arg.getDocument()->getNode(true); else if (Key == "restrict") - Arg[".is_restrict"] = std::make_shared(true); + Arg[".is_restrict"] = Arg.getDocument()->getNode(true); else if (Key == "volatile") - Arg[".is_volatile"] = std::make_shared(true); + Arg[".is_volatile"] = Arg.getDocument()->getNode(true); else if (Key == "pipe") - Arg[".is_pipe"] = std::make_shared(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); @@ -914,54 +912,52 @@ } } -std::shared_ptr +msgpack::MapDocNode MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) const { const GCNSubtarget &STM = MF.getSubtarget(); const SIMachineFunctionInfo &MFI = *MF.getInfo(); const Function &F = MF.getFunction(); - auto HSAKernelProps = std::make_shared(); - auto &Kern = *HSAKernelProps; + auto Kern = HSAMetadataDoc->getMapNode(); unsigned MaxKernArgAlign; - Kern[".kernarg_segment_size"] = std::make_shared( + Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode( STM.getKernArgSegmentSize(F, MaxKernArgAlign)); Kern[".group_segment_fixed_size"] = - std::make_shared(ProgramInfo.LDSSize); + Kern.getDocument()->getNode(ProgramInfo.LDSSize); Kern[".private_segment_fixed_size"] = - std::make_shared(ProgramInfo.ScratchSize); + Kern.getDocument()->getNode(ProgramInfo.ScratchSize); Kern[".kernarg_segment_align"] = - std::make_shared(std::max(uint32_t(4), MaxKernArgAlign)); + Kern.getDocument()->getNode(std::max(uint32_t(4), MaxKernArgAlign)); Kern[".wavefront_size"] = - std::make_shared(STM.getWavefrontSize()); - Kern[".sgpr_count"] = std::make_shared(ProgramInfo.NumSGPR); - Kern[".vgpr_count"] = std::make_shared(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(MFI.getMaxFlatWorkGroupSize()); + Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize()); Kern[".sgpr_spill_count"] = - std::make_shared(MFI.getNumSpilledSGPRs()); + Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs()); Kern[".vgpr_spill_count"] = - std::make_shared(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()); @@ -972,25 +968,24 @@ 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(KernelsNode.get()); + auto Kernels = + getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true); { - auto &Kern = *KernelProps; - Kern[".name"] = std::make_shared(Func.getName()); - Kern[".symbol"] = std::make_shared( - (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 Index: lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h =================================================================== --- lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h +++ lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h @@ -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 @@ /// the \p HSAMetadata structure is updated with the correct types. /// /// \returns True on success, false on failure. - virtual bool EmitHSAMetadata(std::shared_ptr &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 @@ bool EmitISAVersion(StringRef IsaVersionString) override; /// \returns True on success, false on failure. - bool EmitHSAMetadata(std::shared_ptr &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 @@ bool EmitISAVersion(StringRef IsaVersionString) override; /// \returns True on success, false on failure. - bool EmitHSAMetadata(std::shared_ptr &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; Index: lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp =================================================================== --- lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp +++ lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp @@ -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::EmitHSAMetadataV3(StringRef HSAMetadataString) { - std::shared_ptr HSAMetadataRoot; - yaml::Input YIn(HSAMetadataString); - YIn >> HSAMetadataRoot; - if (YIn.error()) + msgpack::Document HSAMetadataDoc; + if (!HSAMetadataDoc.fromYAML(HSAMetadataString)) return false; - return EmitHSAMetadata(HSAMetadataRoot, false); + return EmitHSAMetadata(HSAMetadataDoc, false); } StringRef AMDGPUTargetStreamer::getArchNameFromElfMach(unsigned ElfMach) { @@ -213,15 +210,14 @@ } bool AMDGPUTargetAsmStreamer::EmitHSAMetadata( - std::shared_ptr &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 @@ return true; } -bool AMDGPUTargetELFStreamer::EmitHSAMetadata( - std::shared_ptr &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 @@ 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; Index: test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll @@ -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 Index: test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll @@ -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 { Index: test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll @@ -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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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" } Index: test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll @@ -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, Index: test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll @@ -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, Index: test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll =================================================================== --- test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll +++ test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll @@ -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 @@ 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 @@ 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 Index: test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s =================================================================== --- test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s +++ test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s @@ -2,52 +2,52 @@ // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s -// CHECK: .amdgpu_metadata -// CHECK: amdhsa.kernels: -// CHECK-NEXT: - .max_flat_workgroup_size: 256 -// CHECK-NEXT: .wavefront_size: 128 -// CHECK-NEXT: .symbol: 'test_kernel@kd' -// CHECK-NEXT: .kernarg_segment_size: 8 -// CHECK-NEXT: .private_segment_fixed_size: 32 -// CHECK-NEXT: .name: test_kernel -// CHECK-NEXT: .language: OpenCL C -// CHECK-NEXT: .sgpr_count: 14 -// CHECK-NEXT: .kernarg_segment_align: 64 -// CHECK-NEXT: .vgpr_count: 40 -// CHECK-NEXT: .group_segment_fixed_size: 16 -// CHECK-NEXT: .language_version: -// CHECK-NEXT: - 2 -// CHECK-NEXT: - 0 -// CHECK-NEXT: .args: -// CHECK-NEXT: - .type_name: char -// CHECK-NEXT: .value_kind: by_value -// CHECK-NEXT: .offset: 1 -// CHECK-NEXT: .size: 1 -// CHECK-NEXT: .value_type: i8 -// CHECK-NEXT: - .value_kind: hidden_global_offset_x -// CHECK-NEXT: .offset: 8 -// CHECK-NEXT: .size: 8 -// CHECK-NEXT: .value_type: i64 -// CHECK-NEXT: - .value_kind: hidden_global_offset_y -// CHECK-NEXT: .offset: 8 -// CHECK-NEXT: .size: 8 -// CHECK-NEXT: .value_type: i64 -// CHECK-NEXT: - .value_kind: hidden_global_offset_z -// CHECK-NEXT: .offset: 8 -// CHECK-NEXT: .size: 8 -// CHECK-NEXT: .value_type: i64 -// CHECK-NEXT: - .value_kind: hidden_printf_buffer -// CHECK-NEXT: .offset: 8 -// CHECK-NEXT: .size: 8 -// CHECK-NEXT: .value_type: i8 -// CHECK-NEXT: .address_space: global -// CHECK: amdhsa.version: -// CHECK-NEXT: - 1 -// CHECK-NEXT: - 0 -// CHECK: amdhsa.printf: -// CHECK-NEXT: - '1:1:4:%d\n' -// CHECK-NEXt: - '2:1:8:%g\n' -// CHECK: .end_amdgpu_metadata +; CHECK: .amdgpu_metadata +; CHECK: amdhsa.kernels: +; CHECK-NEXT: - .args: +; CHECK-NEXT: - .offset: 1 +; CHECK-NEXT: .size: 1 +; CHECK-NEXT: .type_name: char +; CHECK-NEXT: .value_kind: by_value +; CHECK-NEXT: .value_type: i8 +; CHECK-NEXT: - .offset: 8 +; CHECK-NEXT: .size: 8 +; CHECK-NEXT: .value_kind: hidden_global_offset_x +; CHECK-NEXT: .value_type: i64 +; CHECK-NEXT: - .offset: 8 +; CHECK-NEXT: .size: 8 +; CHECK-NEXT: .value_kind: hidden_global_offset_y +; CHECK-NEXT: .value_type: i64 +; CHECK-NEXT: - .offset: 8 +; CHECK-NEXT: .size: 8 +; CHECK-NEXT: .value_kind: hidden_global_offset_z +; CHECK-NEXT: .value_type: i64 +; CHECK-NEXT: - .address_space: global +; CHECK-NEXT: .offset: 8 +; CHECK-NEXT: .size: 8 +; CHECK-NEXT: .value_kind: hidden_printf_buffer +; CHECK-NEXT: .value_type: i8 +; CHECK-NEXT: .group_segment_fixed_size: 16 +; CHECK-NEXT: .kernarg_segment_align: 64 +; CHECK-NEXT: .kernarg_segment_size: 8 +; CHECK-NEXT: .language: OpenCL C +; CHECK-NEXT: .language_version: +; CHECK-NEXT: - 2 +; CHECK-NEXT: - 0 +; CHECK-NEXT: .max_flat_workgroup_size: 256 +; CHECK-NEXT: .name: test_kernel +; CHECK-NEXT: .private_segment_fixed_size: 32 +; CHECK-NEXT: .sgpr_count: 14 +; CHECK-NEXT: .symbol: 'test_kernel@kd' +; CHECK-NEXT: .vgpr_count: 40 +; CHECK-NEXT: .wavefront_size: 128 +; CHECK-NEXT: amdhsa.printf: +; CHECK-NEXT: - '1:1:4:%d\n' +; CHECK-NEXT: - '2:1:8:%g\n' +; CHECK-NEXT: amdhsa.version: +; CHECK-NEXT: - 1 +; CHECK-NEXT: - 0 +; CHECK: .end_amdgpu_metadata .amdgpu_metadata amdhsa.version: - 1 Index: test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s =================================================================== --- test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s +++ test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s @@ -2,37 +2,38 @@ // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s -// CHECK: .amdgpu_metadata -// CHECK: amdhsa.kernels: -// CHECK: - .max_flat_workgroup_size: 256 -// CHECK: .wavefront_size: 128 -// CHECK: .symbol: 'test_kernel@kd' -// CHECK: .reqd_workgroup_size: -// CHECK-NEXT: - 1 -// CHECK-NEXT: - 2 -// CHECK-NEXT: - 4 -// CHECK: .kernarg_segment_size: 8 -// CHECK: .private_segment_fixed_size: 32 -// CHECK: .workgroup_size_hint: -// CHECK-NEXT: - 8 -// CHECK-NEXT: - 16 -// CHECK-NEXT: - 32 -// CHECK: .name: test_kernel -// CHECK: .language: OpenCL C -// CHECK: .sgpr_count: 14 -// CHECK: .kernarg_segment_align: 64 -// CHECK: .vgpr_count: 40 -// CHECK: .language_version: -// CHECK-NEXT: - 2 -// CHECK-NEXT: - 0 -// CHECK: .vec_type_hint: int -// CHECK: amdhsa.version: -// CHECK-NEXT: - 1 -// CHECK-NEXT: - 0 -// CHECK: amdhsa.printf: -// CHECK: - '1:1:4:%d\n' -// CHECK: - '2:1:8:%g\n' -// CHECK: .end_amdgpu_metadata +// CHECK: .amdgpu_metadata +// CHECK: amdhsa.kernels: +// CHECK: - .group_segment_fixed_size: 16 +// CHECK: .kernarg_segment_align: 64 +// CHECK: .kernarg_segment_size: 8 +// CHECK: .language: OpenCL C +// CHECK: .language_version: +// CHECK-NEXT: - 2 +// CHECK-NEXT: - 0 +// CHECK: .max_flat_workgroup_size: 256 +// CHECK: .name: test_kernel +// CHECK: .private_segment_fixed_size: 32 +// CHECK: .reqd_workgroup_size: +// CHECK-NEXT: - 1 +// CHECK-NEXT: - 2 +// CHECK-NEXT: - 4 +// CHECK: .sgpr_count: 14 +// CHECK: .symbol: 'test_kernel@kd' +// CHECK: .vec_type_hint: int +// CHECK: .vgpr_count: 40 +// CHECK: .wavefront_size: 128 +// CHECK: .workgroup_size_hint: +// CHECK-NEXT: - 8 +// CHECK-NEXT: - 16 +// CHECK-NEXT: - 32 +// CHECK: amdhsa.printf: +// CHECK: - '1:1:4:%d\n' +// CHECK: - '2:1:8:%g\n' +// CHECK: amdhsa.version: +// CHECK-NEXT: - 1 +// CHECK-NEXT: - 0 +// CHECK: .end_amdgpu_metadata .amdgpu_metadata amdhsa.version: - 1 Index: test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s =================================================================== --- test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s +++ test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s @@ -2,23 +2,23 @@ // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s -// CHECK: .amdgpu_metadata -// CHECK: amdhsa.kernels: -// CHECK: - .sgpr_count: 40 -// CHECK: .max_flat_workgroup_size: 256 -// CHECK: .symbol: 'test_kernel@kd' -// CHECK: .kernarg_segment_size: 24 -// CHECK: .group_segment_fixed_size: 24 -// CHECK: .private_segment_fixed_size: 16 -// CHECK: .vgpr_count: 14 -// CHECK: .vgpr_spill_count: 1 -// CHECK: .kernarg_segment_align: 16 -// CHECK: .sgpr_spill_count: 1 -// CHECK: .wavefront_size: 64 -// CHECK: .name: test_kernel -// CHECK: amdhsa.version: -// CHECK-NEXT: - 1 -// CHECK-NEXT: - 0 +// CHECK: .amdgpu_metadata +// CHECK: amdhsa.kernels: +// CHECK: - .group_segment_fixed_size: 24 +// CHECK: .kernarg_segment_align: 16 +// CHECK: .kernarg_segment_size: 24 +// CHECK: .max_flat_workgroup_size: 256 +// CHECK: .name: test_kernel +// CHECK: .private_segment_fixed_size: 16 +// CHECK: .sgpr_count: 40 +// CHECK: .sgpr_spill_count: 1 +// CHECK: .symbol: 'test_kernel@kd' +// CHECK: .vgpr_count: 14 +// CHECK: .vgpr_spill_count: 1 +// CHECK: .wavefront_size: 64 +// CHECK: amdhsa.version: +// CHECK-NEXT: - 1 +// CHECK-NEXT: - 0 .amdgpu_metadata amdhsa.version: - 1 Index: test/MC/AMDGPU/hsa-v3.s =================================================================== --- test/MC/AMDGPU/hsa-v3.s +++ test/MC/AMDGPU/hsa-v3.s @@ -243,29 +243,29 @@ .max_flat_workgroup_size: 256 .end_amdgpu_metadata -// ASM: .amdgpu_metadata -// ASM: amdhsa.kernels: -// ASM: - .sgpr_count: 14 -// ASM: .max_flat_workgroup_size: 256 -// ASM: .symbol: 'amd_kernel_code_t_test_all@kd' -// ASM: .kernarg_segment_size: 8 -// ASM: .group_segment_fixed_size: 16 -// ASM: .private_segment_fixed_size: 32 -// ASM: .vgpr_count: 40 -// ASM: .kernarg_segment_align: 64 -// ASM: .wavefront_size: 128 -// ASM: .name: amd_kernel_code_t_test_all -// ASM: - .sgpr_count: 14 -// ASM: .max_flat_workgroup_size: 256 -// ASM: .symbol: 'amd_kernel_code_t_minimal@kd' -// ASM: .kernarg_segment_size: 8 -// ASM: .group_segment_fixed_size: 16 -// ASM: .private_segment_fixed_size: 32 -// ASM: .vgpr_count: 40 -// ASM: .kernarg_segment_align: 64 -// ASM: .wavefront_size: 128 -// ASM: .name: amd_kernel_code_t_minimal -// ASM: amdhsa.version: -// ASM-NEXT: - 3 -// ASM-NEXT: - 0 -// ASM: .end_amdgpu_metadata +// ASM: .amdgpu_metadata +// ASM: amdhsa.kernels: +// ASM: - .group_segment_fixed_size: 16 +// ASM: .kernarg_segment_align: 64 +// ASM: .kernarg_segment_size: 8 +// ASM: .max_flat_workgroup_size: 256 +// ASM: .name: amd_kernel_code_t_test_all +// ASM: .private_segment_fixed_size: 32 +// ASM: .sgpr_count: 14 +// ASM: .symbol: 'amd_kernel_code_t_test_all@kd' +// ASM: .vgpr_count: 40 +// ASM: .wavefront_size: 128 +// ASM: - .group_segment_fixed_size: 16 +// ASM: .kernarg_segment_align: 64 +// ASM: .kernarg_segment_size: 8 +// ASM: .max_flat_workgroup_size: 256 +// ASM: .name: amd_kernel_code_t_minimal +// ASM: .private_segment_fixed_size: 32 +// ASM: .sgpr_count: 14 +// ASM: .symbol: 'amd_kernel_code_t_minimal@kd' +// ASM: .vgpr_count: 40 +// ASM: .wavefront_size: 128 +// ASM: amdhsa.version: +// ASM-NEXT: - 3 +// ASM-NEXT: - 0 +// ASM: .end_amdgpu_metadata Index: tools/llvm-readobj/ELFDumper.cpp =================================================================== --- tools/llvm-readobj/ELFDumper.cpp +++ tools/llvm-readobj/ELFDumper.cpp @@ -3880,29 +3880,24 @@ switch (NoteType) { default: return {"", ""}; - case ELF::NT_AMDGPU_METADATA: + case ELF::NT_AMDGPU_METADATA: { auto MsgPackString = StringRef(reinterpret_cast(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