| //===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===// | |
| // | |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | |
| // See https://llvm.org/LICENSE.txt for license information. | |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | |
| // | |
| //===----------------------------------------------------------------------===// | |
| // | |
| /// \file | |
| /// Implements a verifier for AMDGPU HSA metadata. | |
| // | |
| //===----------------------------------------------------------------------===// | |
| #include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h" | |
| #include "llvm/ADT/STLExtras.h" | |
| #include "llvm/ADT/StringSwitch.h" | |
| #include "llvm/BinaryFormat/MsgPackDocument.h" | |
| namespace llvm { | |
| namespace AMDGPU { | |
| namespace HSAMD { | |
| namespace V3 { | |
| bool MetadataVerifier::verifyScalar( | |
| msgpack::DocNode &Node, msgpack::Type SKind, | |
| function_ref<bool(msgpack::DocNode &)> verifyValue) { | |
| if (!Node.isScalar()) | |
| return false; | |
| if (Node.getKind() != SKind) { | |
| if (Strict) | |
| return false; | |
| // If we are not strict, we interpret string values as "implicitly typed" | |
| // and attempt to coerce them to the expected type here. | |
| if (Node.getKind() != msgpack::Type::String) | |
| return false; | |
| StringRef StringValue = Node.getString(); | |
| Node.fromString(StringValue); | |
| if (Node.getKind() != SKind) | |
| return false; | |
| } | |
| if (verifyValue) | |
| return verifyValue(Node); | |
| return true; | |
| } | |
| 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::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode, | |
| std::optional<size_t> Size) { | |
| if (!Node.isArray()) | |
| return false; | |
| auto &Array = Node.getArray(); | |
| if (Size && Array.size() != *Size) | |
| return false; | |
| return llvm::all_of(Array, verifyNode); | |
| } | |
| bool MetadataVerifier::verifyEntry( | |
| msgpack::MapDocNode &MapNode, StringRef Key, bool Required, | |
| function_ref<bool(msgpack::DocNode &)> verifyNode) { | |
| auto Entry = MapNode.find(Key); | |
| if (Entry == MapNode.end()) | |
| return !Required; | |
| return verifyNode(Entry->second); | |
| } | |
| bool MetadataVerifier::verifyScalarEntry( | |
| msgpack::MapDocNode &MapNode, StringRef Key, bool Required, | |
| msgpack::Type SKind, | |
| function_ref<bool(msgpack::DocNode &)> verifyValue) { | |
| return verifyEntry(MapNode, Key, Required, | |
| [this, SKind, verifyValue](msgpack::DocNode &Node) { | |
| return verifyScalar(Node, SKind, verifyValue); | |
| }); | |
| } | |
| bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode, | |
| StringRef Key, bool Required) { | |
| return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) { | |
| return verifyInteger(Node); | |
| }); | |
| } | |
| bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) { | |
| if (!Node.isMap()) | |
| return false; | |
| auto &ArgsMap = Node.getMap(); | |
| if (!verifyScalarEntry(ArgsMap, ".name", false, | |
| msgpack::Type::String)) | |
| return false; | |
| if (!verifyScalarEntry(ArgsMap, ".type_name", false, | |
| 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::Type::String, | |
| [](msgpack::DocNode &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_block_count_x", true) | |
| .Case("hidden_block_count_y", true) | |
| .Case("hidden_block_count_z", true) | |
| .Case("hidden_group_size_x", true) | |
| .Case("hidden_group_size_y", true) | |
| .Case("hidden_group_size_z", true) | |
| .Case("hidden_remainder_x", true) | |
| .Case("hidden_remainder_y", true) | |
| .Case("hidden_remainder_z", true) | |
| .Case("hidden_global_offset_x", true) | |
| .Case("hidden_global_offset_y", true) | |
| .Case("hidden_global_offset_z", true) | |
| .Case("hidden_grid_dims", true) | |
| .Case("hidden_none", true) | |
| .Case("hidden_printf_buffer", true) | |
| .Case("hidden_hostcall_buffer", true) | |
| .Case("hidden_heap_v1", true) | |
| .Case("hidden_default_queue", true) | |
| .Case("hidden_completion_action", true) | |
| .Case("hidden_multigrid_sync_arg", true) | |
| .Case("hidden_dynamic_lds_size", true) | |
| .Case("hidden_private_base", true) | |
| .Case("hidden_shared_base", true) | |
| .Case("hidden_queue_ptr", true) | |
| .Default(false); | |
| })) | |
| return false; | |
| if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false)) | |
| return false; | |
| if (!verifyScalarEntry(ArgsMap, ".address_space", false, | |
| msgpack::Type::String, | |
| [](msgpack::DocNode &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::Type::String, | |
| [](msgpack::DocNode &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::Type::String, | |
| [](msgpack::DocNode &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::Type::Boolean)) | |
| return false; | |
| if (!verifyScalarEntry(ArgsMap, ".is_restrict", false, | |
| msgpack::Type::Boolean)) | |
| return false; | |
| if (!verifyScalarEntry(ArgsMap, ".is_volatile", false, | |
| msgpack::Type::Boolean)) | |
| return false; | |
| if (!verifyScalarEntry(ArgsMap, ".is_pipe", false, | |
| msgpack::Type::Boolean)) | |
| return false; | |
| return true; | |
| } | |
| bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) { | |
| if (!Node.isMap()) | |
| return false; | |
| auto &KernelMap = Node.getMap(); | |
| if (!verifyScalarEntry(KernelMap, ".name", true, | |
| msgpack::Type::String)) | |
| return false; | |
| if (!verifyScalarEntry(KernelMap, ".symbol", true, | |
| msgpack::Type::String)) | |
| return false; | |
| if (!verifyScalarEntry(KernelMap, ".language", false, | |
| msgpack::Type::String, | |
| [](msgpack::DocNode &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::DocNode &Node) { | |
| return verifyArray( | |
| Node, | |
| [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2); | |
| })) | |
| return false; | |
| 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::DocNode &Node) { | |
| return verifyArray(Node, | |
| [this](msgpack::DocNode &Node) { | |
| return verifyInteger(Node); | |
| }, | |
| 3); | |
| })) | |
| return false; | |
| if (!verifyEntry(KernelMap, ".workgroup_size_hint", false, | |
| [this](msgpack::DocNode &Node) { | |
| return verifyArray(Node, | |
| [this](msgpack::DocNode &Node) { | |
| return verifyInteger(Node); | |
| }, | |
| 3); | |
| })) | |
| return false; | |
| if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false, | |
| msgpack::Type::String)) | |
| return false; | |
| if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false, | |
| msgpack::Type::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 (!verifyScalarEntry(KernelMap, ".uses_dynamic_stack", false, | |
| msgpack::Type::Boolean)) | |
| return false; | |
| if (!verifyIntegerEntry(KernelMap, ".workgroup_processor_mode", false)) | |
| 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; | |
| if (!verifyIntegerEntry(KernelMap, ".uniform_work_group_size", false)) | |
| return false; | |
| if (!verifyEntry( | |
| KernelMap, ".cluster_dims", false, [this](msgpack::DocNode &Node) { | |
| return verifyArray( | |
| Node, | |
| [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, | |
| 3); | |
| })) | |
| return false; | |
| return true; | |
| } | |
| bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) { | |
| if (!HSAMetadataRoot.isMap()) | |
| return false; | |
| auto &RootMap = HSAMetadataRoot.getMap(); | |
| if (!verifyEntry( | |
| RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) { | |
| return verifyArray( | |
| Node, | |
| [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2); | |
| })) | |
| return false; | |
| if (!verifyEntry( | |
| 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::DocNode &Node) { | |
| return verifyArray(Node, [this](msgpack::DocNode &Node) { | |
| return verifyKernel(Node); | |
| }); | |
| })) | |
| return false; | |
| return true; | |
| } | |
| } // end namespace V3 | |
| } // end namespace HSAMD | |
| } // end namespace AMDGPU | |
| } // end namespace llvm |