1 //===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===//
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-exception
7 //===----------------------------------------------------------------------===//
10 /// Implements a verifier for AMDGPU HSA metadata.
12 //===----------------------------------------------------------------------===//
14 #include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
15 #include "llvm/Support/AMDGPUMetadata.h"
22 bool MetadataVerifier::verifyScalar(
23 msgpack::DocNode &Node, msgpack::Type SKind,
24 function_ref<bool(msgpack::DocNode &)> verifyValue) {
27 if (Node.getKind() != SKind) {
30 // If we are not strict, we interpret string values as "implicitly typed"
31 // and attempt to coerce them to the expected type here.
32 if (Node.getKind() != msgpack::Type::String)
34 StringRef StringValue = Node.getString();
35 Node.fromString(StringValue);
36 if (Node.getKind() != SKind)
40 return verifyValue(Node);
44 bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
45 if (!verifyScalar(Node, msgpack::Type::UInt))
46 if (!verifyScalar(Node, msgpack::Type::Int))
51 bool MetadataVerifier::verifyArray(
52 msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
53 Optional<size_t> Size) {
56 auto &Array = Node.getArray();
57 if (Size && Array.size() != *Size)
59 for (auto &Item : Array)
60 if (!verifyNode(Item))
66 bool MetadataVerifier::verifyEntry(
67 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
68 function_ref<bool(msgpack::DocNode &)> verifyNode) {
69 auto Entry = MapNode.find(Key);
70 if (Entry == MapNode.end())
72 return verifyNode(Entry->second);
75 bool MetadataVerifier::verifyScalarEntry(
76 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
78 function_ref<bool(msgpack::DocNode &)> verifyValue) {
79 return verifyEntry(MapNode, Key, Required, [=](msgpack::DocNode &Node) {
80 return verifyScalar(Node, SKind, verifyValue);
84 bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
85 StringRef Key, bool Required) {
86 return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {
87 return verifyInteger(Node);
91 bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {
94 auto &ArgsMap = Node.getMap();
96 if (!verifyScalarEntry(ArgsMap, ".name", false,
97 msgpack::Type::String))
99 if (!verifyScalarEntry(ArgsMap, ".type_name", false,
100 msgpack::Type::String))
102 if (!verifyIntegerEntry(ArgsMap, ".size", true))
104 if (!verifyIntegerEntry(ArgsMap, ".offset", true))
106 if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
107 msgpack::Type::String,
108 [](msgpack::DocNode &SNode) {
109 return StringSwitch<bool>(SNode.getString())
110 .Case("by_value", true)
111 .Case("global_buffer", true)
112 .Case("dynamic_shared_pointer", true)
113 .Case("sampler", true)
117 .Case("hidden_global_offset_x", true)
118 .Case("hidden_global_offset_y", true)
119 .Case("hidden_global_offset_z", true)
120 .Case("hidden_none", true)
121 .Case("hidden_printf_buffer", true)
122 .Case("hidden_hostcall_buffer", true)
123 .Case("hidden_default_queue", true)
124 .Case("hidden_completion_action", true)
125 .Case("hidden_multigrid_sync_arg", true)
129 if (!verifyScalarEntry(ArgsMap, ".value_type", true,
130 msgpack::Type::String,
131 [](msgpack::DocNode &SNode) {
132 return StringSwitch<bool>(SNode.getString())
133 .Case("struct", true)
148 if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
150 if (!verifyScalarEntry(ArgsMap, ".address_space", false,
151 msgpack::Type::String,
152 [](msgpack::DocNode &SNode) {
153 return StringSwitch<bool>(SNode.getString())
154 .Case("private", true)
155 .Case("global", true)
156 .Case("constant", true)
158 .Case("generic", true)
159 .Case("region", true)
163 if (!verifyScalarEntry(ArgsMap, ".access", false,
164 msgpack::Type::String,
165 [](msgpack::DocNode &SNode) {
166 return StringSwitch<bool>(SNode.getString())
167 .Case("read_only", true)
168 .Case("write_only", true)
169 .Case("read_write", true)
173 if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
174 msgpack::Type::String,
175 [](msgpack::DocNode &SNode) {
176 return StringSwitch<bool>(SNode.getString())
177 .Case("read_only", true)
178 .Case("write_only", true)
179 .Case("read_write", true)
183 if (!verifyScalarEntry(ArgsMap, ".is_const", false,
184 msgpack::Type::Boolean))
186 if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
187 msgpack::Type::Boolean))
189 if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
190 msgpack::Type::Boolean))
192 if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
193 msgpack::Type::Boolean))
199 bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
202 auto &KernelMap = Node.getMap();
204 if (!verifyScalarEntry(KernelMap, ".name", true,
205 msgpack::Type::String))
207 if (!verifyScalarEntry(KernelMap, ".symbol", true,
208 msgpack::Type::String))
210 if (!verifyScalarEntry(KernelMap, ".language", false,
211 msgpack::Type::String,
212 [](msgpack::DocNode &SNode) {
213 return StringSwitch<bool>(SNode.getString())
214 .Case("OpenCL C", true)
215 .Case("OpenCL C++", true)
218 .Case("OpenMP", true)
219 .Case("Assembler", true)
224 KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
227 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
230 if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
231 return verifyArray(Node, [this](msgpack::DocNode &Node) {
232 return verifyKernelArgs(Node);
236 if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
237 [this](msgpack::DocNode &Node) {
238 return verifyArray(Node,
239 [this](msgpack::DocNode &Node) {
240 return verifyInteger(Node);
245 if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
246 [this](msgpack::DocNode &Node) {
247 return verifyArray(Node,
248 [this](msgpack::DocNode &Node) {
249 return verifyInteger(Node);
254 if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
255 msgpack::Type::String))
257 if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
258 msgpack::Type::String))
260 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
262 if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
264 if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
266 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
268 if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
270 if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
272 if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
274 if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
276 if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
278 if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
284 bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
285 if (!HSAMetadataRoot.isMap())
287 auto &RootMap = HSAMetadataRoot.getMap();
290 RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
293 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
297 RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
298 return verifyArray(Node, [this](msgpack::DocNode &Node) {
299 return verifyScalar(Node, msgpack::Type::String);
303 if (!verifyEntry(RootMap, "amdhsa.kernels", true,
304 [this](msgpack::DocNode &Node) {
305 return verifyArray(Node, [this](msgpack::DocNode &Node) {
306 return verifyKernel(Node);
314 } // end namespace V3
315 } // end namespace HSAMD
316 } // end namespace AMDGPU
317 } // end namespace llvm