LLVM 22.0.0git
AMDGPUMetadataVerifier.cpp
Go to the documentation of this file.
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-exception
6//
7//===----------------------------------------------------------------------===//
8//
9/// \file
10/// 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 V3
325} // end namespace HSAMD
326} // end namespace AMDGPU
327} // end namespace llvm
AMDGPUMetadataVerifier.h
This is a verifier for AMDGPU HSA metadata, which can verify both well-typed metadata and untyped met...
MsgPackDocument.h
This file declares a class that exposes a simple in-memory representation of a document of MsgPack ob...
STLExtras.h
This file contains some templates that are useful if you are working with the STL at all.
StringSwitch.h
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
llvm::AMDGPU::HSAMD::V3::MetadataVerifier::verify
LLVM_ABI bool verify(msgpack::DocNode &HSAMetadataRoot)
Verify given HSA metadata.
Definition AMDGPUMetadataVerifier.cpp:294
llvm::msgpack::DocNode
A node in a MsgPack Document.
Definition MsgPackDocument.h:39
llvm::msgpack::DocNode::getMap
MapDocNode & getMap(bool Convert=false)
Get a MapDocNode for a map node.
Definition MsgPackDocument.h:146
llvm::COFF::Entry
@ Entry
Definition COFF.h:862
llvm::msgpack::Type
Type
MessagePack types as defined in the standard, with the exception of Integer being divided into a sign...
Definition MsgPackReader.h:54
llvm::rdf::Node
NodeAddr< NodeBase * > Node
Definition RDFGraph.h:381
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition AddressRanges.h:18
llvm::all_of
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
Definition STLExtras.h:1725
Generated on for LLVM by doxygen 1.14.0