| //===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===// |
| // |
| // The LLVM Compiler Infrastructure |
| // |
| // This file is distributed under the University of Illinois Open Source |
| // License. See LICENSE.TXT for details. |
| // |
| //===----------------------------------------------------------------------===// |
| // |
| /// \file |
| /// Implements a verifier for AMDGPU HSA metadata. |
| // |
| //===----------------------------------------------------------------------===// |
| |
| #include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h" |
| #include "llvm/Support/AMDGPUMetadata.h" |
| |
| namespace llvm { |
| namespace AMDGPU { |
| namespace HSAMD { |
| namespace V3 { |
| |
| bool MetadataVerifier::verifyScalar( |
| msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind, |
| function_ref<bool(msgpack::ScalarNode &)> verifyValue) { |
| auto ScalarPtr = dyn_cast<msgpack::ScalarNode>(&Node); |
| if (!ScalarPtr) |
| return false; |
| auto &Scalar = *ScalarPtr; |
| // Do not output extraneous tags for types we know from the spec. |
| Scalar.IgnoreTag = true; |
| if (Scalar.getScalarKind() != SKind) { |
| 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) |
| return false; |
| std::string StringValue = Scalar.getString(); |
| Scalar.setScalarKind(SKind); |
| if (Scalar.inputYAML(StringValue) != StringRef()) |
| return false; |
| } |
| if (verifyValue) |
| return verifyValue(Scalar); |
| return true; |
| } |
| |
| bool MetadataVerifier::verifyInteger(msgpack::Node &Node) { |
| if (!verifyScalar(Node, msgpack::ScalarNode::SK_UInt)) |
| if (!verifyScalar(Node, msgpack::ScalarNode::SK_Int)) |
| return false; |
| return true; |
| } |
| |
| bool MetadataVerifier::verifyArray( |
| msgpack::Node &Node, function_ref<bool(msgpack::Node &)> verifyNode, |
| Optional<size_t> Size) { |
| auto ArrayPtr = dyn_cast<msgpack::ArrayNode>(&Node); |
| if (!ArrayPtr) |
| return false; |
| auto &Array = *ArrayPtr; |
| if (Size && Array.size() != *Size) |
| return false; |
| for (auto &Item : Array) |
| if (!verifyNode(*Item.get())) |
| return false; |
| |
| return true; |
| } |
| |
| bool MetadataVerifier::verifyEntry( |
| msgpack::MapNode &MapNode, StringRef Key, bool Required, |
| function_ref<bool(msgpack::Node &)> verifyNode) { |
| auto Entry = MapNode.find(Key); |
| if (Entry == MapNode.end()) |
| return !Required; |
| return verifyNode(*Entry->second.get()); |
| } |
| |
| bool MetadataVerifier::verifyScalarEntry( |
| msgpack::MapNode &MapNode, StringRef Key, bool Required, |
| msgpack::ScalarNode::ScalarKind SKind, |
| function_ref<bool(msgpack::ScalarNode &)> verifyValue) { |
| return verifyEntry(MapNode, Key, Required, [=](msgpack::Node &Node) { |
| return verifyScalar(Node, SKind, verifyValue); |
| }); |
| } |
| |
| bool MetadataVerifier::verifyIntegerEntry(msgpack::MapNode &MapNode, |
| StringRef Key, bool Required) { |
| return verifyEntry(MapNode, Key, Required, [this](msgpack::Node &Node) { |
| return verifyInteger(Node); |
| }); |
| } |
| |
| bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) { |
| auto ArgsMapPtr = dyn_cast<msgpack::MapNode>(&Node); |
| if (!ArgsMapPtr) |
| return false; |
| auto &ArgsMap = *ArgsMapPtr; |
| |
| if (!verifyScalarEntry(ArgsMap, ".name", false, |
| msgpack::ScalarNode::SK_String)) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".type_name", false, |
| msgpack::ScalarNode::SK_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) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("by_value", true) |
| .Case("global_buffer", true) |
| .Case("dynamic_shared_pointer", true) |
| .Case("sampler", true) |
| .Case("image", true) |
| .Case("pipe", true) |
| .Case("queue", true) |
| .Case("hidden_global_offset_x", true) |
| .Case("hidden_global_offset_y", true) |
| .Case("hidden_global_offset_z", true) |
| .Case("hidden_none", true) |
| .Case("hidden_printf_buffer", true) |
| .Case("hidden_default_queue", true) |
| .Case("hidden_completion_action", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".value_type", true, |
| msgpack::ScalarNode::SK_String, |
| [](msgpack::ScalarNode &SNode) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("struct", true) |
| .Case("i8", true) |
| .Case("u8", true) |
| .Case("i16", true) |
| .Case("u16", true) |
| .Case("f16", true) |
| .Case("i32", true) |
| .Case("u32", true) |
| .Case("f32", true) |
| .Case("i64", true) |
| .Case("u64", true) |
| .Case("f64", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false)) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".address_space", false, |
| msgpack::ScalarNode::SK_String, |
| [](msgpack::ScalarNode &SNode) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("private", true) |
| .Case("global", true) |
| .Case("constant", true) |
| .Case("local", true) |
| .Case("generic", true) |
| .Case("region", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".access", false, |
| msgpack::ScalarNode::SK_String, |
| [](msgpack::ScalarNode &SNode) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("read_only", true) |
| .Case("write_only", true) |
| .Case("read_write", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".actual_access", false, |
| msgpack::ScalarNode::SK_String, |
| [](msgpack::ScalarNode &SNode) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("read_only", true) |
| .Case("write_only", true) |
| .Case("read_write", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".is_const", false, |
| msgpack::ScalarNode::SK_Boolean)) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".is_restrict", false, |
| msgpack::ScalarNode::SK_Boolean)) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".is_volatile", false, |
| msgpack::ScalarNode::SK_Boolean)) |
| return false; |
| if (!verifyScalarEntry(ArgsMap, ".is_pipe", false, |
| msgpack::ScalarNode::SK_Boolean)) |
| return false; |
| |
| return true; |
| } |
| |
| bool MetadataVerifier::verifyKernel(msgpack::Node &Node) { |
| auto KernelMapPtr = dyn_cast<msgpack::MapNode>(&Node); |
| if (!KernelMapPtr) |
| return false; |
| auto &KernelMap = *KernelMapPtr; |
| |
| if (!verifyScalarEntry(KernelMap, ".name", true, |
| msgpack::ScalarNode::SK_String)) |
| return false; |
| if (!verifyScalarEntry(KernelMap, ".symbol", true, |
| msgpack::ScalarNode::SK_String)) |
| return false; |
| if (!verifyScalarEntry(KernelMap, ".language", false, |
| msgpack::ScalarNode::SK_String, |
| [](msgpack::ScalarNode &SNode) { |
| return StringSwitch<bool>(SNode.getString()) |
| .Case("OpenCL C", true) |
| .Case("OpenCL C++", true) |
| .Case("HCC", true) |
| .Case("HIP", true) |
| .Case("OpenMP", true) |
| .Case("Assembler", true) |
| .Default(false); |
| })) |
| return false; |
| if (!verifyEntry( |
| KernelMap, ".language_version", false, [this](msgpack::Node &Node) { |
| return verifyArray( |
| Node, |
| [this](msgpack::Node &Node) { return verifyInteger(Node); }, 2); |
| })) |
| return false; |
| if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::Node &Node) { |
| return verifyArray(Node, [this](msgpack::Node &Node) { |
| return verifyKernelArgs(Node); |
| }); |
| })) |
| return false; |
| if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false, |
| [this](msgpack::Node &Node) { |
| return verifyArray(Node, |
| [this](msgpack::Node &Node) { |
| return verifyInteger(Node); |
| }, |
| 3); |
| })) |
| return false; |
| if (!verifyEntry(KernelMap, ".workgroup_size_hint", false, |
| [this](msgpack::Node &Node) { |
| return verifyArray(Node, |
| [this](msgpack::Node &Node) { |
| return verifyInteger(Node); |
| }, |
| 3); |
| })) |
| return false; |
| if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false, |
| msgpack::ScalarNode::SK_String)) |
| return false; |
| if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false, |
| msgpack::ScalarNode::SK_String)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false)) |
| return false; |
| if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false)) |
| return false; |
| |
| return true; |
| } |
| |
| bool MetadataVerifier::verify(msgpack::Node &HSAMetadataRoot) { |
| auto RootMapPtr = dyn_cast<msgpack::MapNode>(&HSAMetadataRoot); |
| if (!RootMapPtr) |
| return false; |
| auto &RootMap = *RootMapPtr; |
| |
| if (!verifyEntry( |
| RootMap, "amdhsa.version", true, [this](msgpack::Node &Node) { |
| return verifyArray( |
| Node, |
| [this](msgpack::Node &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); |
| }); |
| })) |
| return false; |
| if (!verifyEntry(RootMap, "amdhsa.kernels", true, |
| [this](msgpack::Node &Node) { |
| return verifyArray(Node, [this](msgpack::Node &Node) { |
| return verifyKernel(Node); |
| }); |
| })) |
| return false; |
| |
| return true; |
| } |
| |
| } // end namespace V3 |
| } // end namespace HSAMD |
| } // end namespace AMDGPU |
| } // end namespace llvm |
