328 lines · cpp
1//===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9/// \file10/// Implements a verifier for AMDGPU HSA metadata.11//12//===----------------------------------------------------------------------===//13 14#include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"15 16#include "llvm/ADT/STLExtras.h"17#include "llvm/ADT/StringSwitch.h"18#include "llvm/BinaryFormat/MsgPackDocument.h"19 20namespace llvm {21namespace AMDGPU {22namespace HSAMD {23namespace V3 {24 25bool MetadataVerifier::verifyScalar(26 msgpack::DocNode &Node, msgpack::Type SKind,27 function_ref<bool(msgpack::DocNode &)> verifyValue) {28 if (!Node.isScalar())29 return false;30 if (Node.getKind() != SKind) {31 if (Strict)32 return false;33 // If we are not strict, we interpret string values as "implicitly typed"34 // and attempt to coerce them to the expected type here.35 if (Node.getKind() != msgpack::Type::String)36 return false;37 StringRef StringValue = Node.getString();38 Node.fromString(StringValue);39 if (Node.getKind() != SKind)40 return false;41 }42 if (verifyValue)43 return verifyValue(Node);44 return true;45}46 47bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {48 if (!verifyScalar(Node, msgpack::Type::UInt))49 if (!verifyScalar(Node, msgpack::Type::Int))50 return false;51 return true;52}53 54bool MetadataVerifier::verifyArray(55 msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,56 std::optional<size_t> Size) {57 if (!Node.isArray())58 return false;59 auto &Array = Node.getArray();60 if (Size && Array.size() != *Size)61 return false;62 return llvm::all_of(Array, verifyNode);63}64 65bool MetadataVerifier::verifyEntry(66 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,67 function_ref<bool(msgpack::DocNode &)> verifyNode) {68 auto Entry = MapNode.find(Key);69 if (Entry == MapNode.end())70 return !Required;71 return verifyNode(Entry->second);72}73 74bool MetadataVerifier::verifyScalarEntry(75 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,76 msgpack::Type SKind,77 function_ref<bool(msgpack::DocNode &)> verifyValue) {78 return verifyEntry(MapNode, Key, Required,79 [this, SKind, verifyValue](msgpack::DocNode &Node) {80 return verifyScalar(Node, SKind, verifyValue);81 });82}83 84bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,85 StringRef Key, bool Required) {86 return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {87 return verifyInteger(Node);88 });89}90 91bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {92 if (!Node.isMap())93 return false;94 auto &ArgsMap = Node.getMap();95 96 if (!verifyScalarEntry(ArgsMap, ".name", false,97 msgpack::Type::String))98 return false;99 if (!verifyScalarEntry(ArgsMap, ".type_name", false,100 msgpack::Type::String))101 return false;102 if (!verifyIntegerEntry(ArgsMap, ".size", true))103 return false;104 if (!verifyIntegerEntry(ArgsMap, ".offset", true))105 return false;106 if (!verifyScalarEntry(ArgsMap, ".value_kind", true, msgpack::Type::String,107 [](msgpack::DocNode &SNode) {108 return StringSwitch<bool>(SNode.getString())109 .Case("by_value", true)110 .Case("global_buffer", true)111 .Case("dynamic_shared_pointer", true)112 .Case("sampler", true)113 .Case("image", true)114 .Case("pipe", true)115 .Case("queue", true)116 .Case("hidden_block_count_x", true)117 .Case("hidden_block_count_y", true)118 .Case("hidden_block_count_z", true)119 .Case("hidden_group_size_x", true)120 .Case("hidden_group_size_y", true)121 .Case("hidden_group_size_z", true)122 .Case("hidden_remainder_x", true)123 .Case("hidden_remainder_y", true)124 .Case("hidden_remainder_z", true)125 .Case("hidden_global_offset_x", true)126 .Case("hidden_global_offset_y", true)127 .Case("hidden_global_offset_z", true)128 .Case("hidden_grid_dims", true)129 .Case("hidden_none", true)130 .Case("hidden_printf_buffer", true)131 .Case("hidden_hostcall_buffer", true)132 .Case("hidden_heap_v1", true)133 .Case("hidden_default_queue", true)134 .Case("hidden_completion_action", true)135 .Case("hidden_multigrid_sync_arg", true)136 .Case("hidden_dynamic_lds_size", true)137 .Case("hidden_private_base", true)138 .Case("hidden_shared_base", true)139 .Case("hidden_queue_ptr", true)140 .Default(false);141 }))142 return false;143 if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))144 return false;145 if (!verifyScalarEntry(ArgsMap, ".address_space", false,146 msgpack::Type::String,147 [](msgpack::DocNode &SNode) {148 return StringSwitch<bool>(SNode.getString())149 .Case("private", true)150 .Case("global", true)151 .Case("constant", true)152 .Case("local", true)153 .Case("generic", true)154 .Case("region", true)155 .Default(false);156 }))157 return false;158 if (!verifyScalarEntry(ArgsMap, ".access", false,159 msgpack::Type::String,160 [](msgpack::DocNode &SNode) {161 return StringSwitch<bool>(SNode.getString())162 .Case("read_only", true)163 .Case("write_only", true)164 .Case("read_write", true)165 .Default(false);166 }))167 return false;168 if (!verifyScalarEntry(ArgsMap, ".actual_access", false,169 msgpack::Type::String,170 [](msgpack::DocNode &SNode) {171 return StringSwitch<bool>(SNode.getString())172 .Case("read_only", true)173 .Case("write_only", true)174 .Case("read_write", true)175 .Default(false);176 }))177 return false;178 if (!verifyScalarEntry(ArgsMap, ".is_const", false,179 msgpack::Type::Boolean))180 return false;181 if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,182 msgpack::Type::Boolean))183 return false;184 if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,185 msgpack::Type::Boolean))186 return false;187 if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,188 msgpack::Type::Boolean))189 return false;190 191 return true;192}193 194bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {195 if (!Node.isMap())196 return false;197 auto &KernelMap = Node.getMap();198 199 if (!verifyScalarEntry(KernelMap, ".name", true,200 msgpack::Type::String))201 return false;202 if (!verifyScalarEntry(KernelMap, ".symbol", true,203 msgpack::Type::String))204 return false;205 if (!verifyScalarEntry(KernelMap, ".language", false,206 msgpack::Type::String,207 [](msgpack::DocNode &SNode) {208 return StringSwitch<bool>(SNode.getString())209 .Case("OpenCL C", true)210 .Case("OpenCL C++", true)211 .Case("HCC", true)212 .Case("HIP", true)213 .Case("OpenMP", true)214 .Case("Assembler", true)215 .Default(false);216 }))217 return false;218 if (!verifyEntry(219 KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {220 return verifyArray(221 Node,222 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);223 }))224 return false;225 if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {226 return verifyArray(Node, [this](msgpack::DocNode &Node) {227 return verifyKernelArgs(Node);228 });229 }))230 return false;231 if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,232 [this](msgpack::DocNode &Node) {233 return verifyArray(Node,234 [this](msgpack::DocNode &Node) {235 return verifyInteger(Node);236 },237 3);238 }))239 return false;240 if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,241 [this](msgpack::DocNode &Node) {242 return verifyArray(Node,243 [this](msgpack::DocNode &Node) {244 return verifyInteger(Node);245 },246 3);247 }))248 return false;249 if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,250 msgpack::Type::String))251 return false;252 if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,253 msgpack::Type::String))254 return false;255 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))256 return false;257 if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))258 return false;259 if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))260 return false;261 if (!verifyScalarEntry(KernelMap, ".uses_dynamic_stack", false,262 msgpack::Type::Boolean))263 return false;264 if (!verifyIntegerEntry(KernelMap, ".workgroup_processor_mode", false))265 return false;266 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))267 return false;268 if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))269 return false;270 if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))271 return false;272 if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))273 return false;274 if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))275 return false;276 if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))277 return false;278 if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))279 return false;280 if (!verifyIntegerEntry(KernelMap, ".uniform_work_group_size", false))281 return false;282 if (!verifyEntry(283 KernelMap, ".cluster_dims", false, [this](msgpack::DocNode &Node) {284 return verifyArray(285 Node,286 [this](msgpack::DocNode &Node) { return verifyInteger(Node); },287 3);288 }))289 return false;290 291 return true;292}293 294bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {295 if (!HSAMetadataRoot.isMap())296 return false;297 auto &RootMap = HSAMetadataRoot.getMap();298 299 if (!verifyEntry(300 RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {301 return verifyArray(302 Node,303 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);304 }))305 return false;306 if (!verifyEntry(307 RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {308 return verifyArray(Node, [this](msgpack::DocNode &Node) {309 return verifyScalar(Node, msgpack::Type::String);310 });311 }))312 return false;313 if (!verifyEntry(RootMap, "amdhsa.kernels", true,314 [this](msgpack::DocNode &Node) {315 return verifyArray(Node, [this](msgpack::DocNode &Node) {316 return verifyKernel(Node);317 });318 }))319 return false;320 321 return true;322}323 324} // end namespace V3325} // end namespace HSAMD326} // end namespace AMDGPU327} // end namespace llvm328