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_default_queue", true)
123 .Case("hidden_completion_action", true)
124 .Case("hidden_multigrid_sync_arg", true)
128 if (!verifyScalarEntry(ArgsMap, ".value_type", true,
129 msgpack::Type::String,
130 [](msgpack::DocNode &SNode) {
131 return StringSwitch<bool>(SNode.getString())
132 .Case("struct", true)
147 if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
149 if (!verifyScalarEntry(ArgsMap, ".address_space", false,
150 msgpack::Type::String,
151 [](msgpack::DocNode &SNode) {
152 return StringSwitch<bool>(SNode.getString())
153 .Case("private", true)
154 .Case("global", true)
155 .Case("constant", true)
157 .Case("generic", true)
158 .Case("region", true)
162 if (!verifyScalarEntry(ArgsMap, ".access", false,
163 msgpack::Type::String,
164 [](msgpack::DocNode &SNode) {
165 return StringSwitch<bool>(SNode.getString())
166 .Case("read_only", true)
167 .Case("write_only", true)
168 .Case("read_write", true)
172 if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
173 msgpack::Type::String,
174 [](msgpack::DocNode &SNode) {
175 return StringSwitch<bool>(SNode.getString())
176 .Case("read_only", true)
177 .Case("write_only", true)
178 .Case("read_write", true)
182 if (!verifyScalarEntry(ArgsMap, ".is_const", false,
183 msgpack::Type::Boolean))
185 if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
186 msgpack::Type::Boolean))
188 if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
189 msgpack::Type::Boolean))
191 if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
192 msgpack::Type::Boolean))
198 bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
201 auto &KernelMap = Node.getMap();
203 if (!verifyScalarEntry(KernelMap, ".name", true,
204 msgpack::Type::String))
206 if (!verifyScalarEntry(KernelMap, ".symbol", true,
207 msgpack::Type::String))
209 if (!verifyScalarEntry(KernelMap, ".language", false,
210 msgpack::Type::String,
211 [](msgpack::DocNode &SNode) {
212 return StringSwitch<bool>(SNode.getString())
213 .Case("OpenCL C", true)
214 .Case("OpenCL C++", true)
217 .Case("OpenMP", true)
218 .Case("Assembler", true)
223 KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
226 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
229 if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
230 return verifyArray(Node, [this](msgpack::DocNode &Node) {
231 return verifyKernelArgs(Node);
235 if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
236 [this](msgpack::DocNode &Node) {
237 return verifyArray(Node,
238 [this](msgpack::DocNode &Node) {
239 return verifyInteger(Node);
244 if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
245 [this](msgpack::DocNode &Node) {
246 return verifyArray(Node,
247 [this](msgpack::DocNode &Node) {
248 return verifyInteger(Node);
253 if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
254 msgpack::Type::String))
256 if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
257 msgpack::Type::String))
259 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
261 if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
263 if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
265 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
267 if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
269 if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
271 if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
273 if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
275 if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
277 if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
283 bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
284 if (!HSAMetadataRoot.isMap())
286 auto &RootMap = HSAMetadataRoot.getMap();
289 RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
292 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
296 RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
297 return verifyArray(Node, [this](msgpack::DocNode &Node) {
298 return verifyScalar(Node, msgpack::Type::String);
302 if (!verifyEntry(RootMap, "amdhsa.kernels", true,
303 [this](msgpack::DocNode &Node) {
304 return verifyArray(Node, [this](msgpack::DocNode &Node) {
305 return verifyKernel(Node);
313 } // end namespace V3
314 } // end namespace HSAMD
315 } // end namespace AMDGPU
316 } // end namespace llvm