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