LLVM 17.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
15
16#include "llvm/ADT/STLExtras.h"
19
20#include <map>
21#include <utility>
22
23namespace llvm {
24namespace AMDGPU {
25namespace HSAMD {
26namespace V3 {
27
28bool MetadataVerifier::verifyScalar(
29 msgpack::DocNode &Node, msgpack::Type SKind,
30 function_ref<bool(msgpack::DocNode &)> verifyValue) {
31 if (!Node.isScalar())
32 return false;
33 if (Node.getKind() != SKind) {
34 if (Strict)
35 return false;
36 // If we are not strict, we interpret string values as "implicitly typed"
37 // and attempt to coerce them to the expected type here.
39 return false;
40 StringRef StringValue = Node.getString();
41 Node.fromString(StringValue);
42 if (Node.getKind() != SKind)
43 return false;
44 }
45 if (verifyValue)
46 return verifyValue(Node);
47 return true;
48}
49
50bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
51 if (!verifyScalar(Node, msgpack::Type::UInt))
52 if (!verifyScalar(Node, msgpack::Type::Int))
53 return false;
54 return true;
55}
56
57bool MetadataVerifier::verifyArray(
58 msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
59 std::optional<size_t> Size) {
60 if (!Node.isArray())
61 return false;
62 auto &Array = Node.getArray();
63 if (Size && Array.size() != *Size)
64 return false;
65 return llvm::all_of(Array, verifyNode);
66}
67
68bool MetadataVerifier::verifyEntry(
69 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
70 function_ref<bool(msgpack::DocNode &)> verifyNode) {
71 auto Entry = MapNode.find(Key);
72 if (Entry == MapNode.end())
73 return !Required;
74 return verifyNode(Entry->second);
75}
76
77bool MetadataVerifier::verifyScalarEntry(
78 msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
79 msgpack::Type SKind,
80 function_ref<bool(msgpack::DocNode &)> verifyValue) {
81 return verifyEntry(MapNode, Key, Required, [=](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,
100 return false;
101 if (!verifyScalarEntry(ArgsMap, ".type_name", false,
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_private_base", true)
139 .Case("hidden_shared_base", true)
140 .Case("hidden_queue_ptr", true)
141 .Default(false);
142 }))
143 return false;
144 if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
145 return false;
146 if (!verifyScalarEntry(ArgsMap, ".address_space", false,
148 [](msgpack::DocNode &SNode) {
149 return StringSwitch<bool>(SNode.getString())
150 .Case("private", true)
151 .Case("global", true)
152 .Case("constant", true)
153 .Case("local", true)
154 .Case("generic", true)
155 .Case("region", true)
156 .Default(false);
157 }))
158 return false;
159 if (!verifyScalarEntry(ArgsMap, ".access", false,
161 [](msgpack::DocNode &SNode) {
162 return StringSwitch<bool>(SNode.getString())
163 .Case("read_only", true)
164 .Case("write_only", true)
165 .Case("read_write", true)
166 .Default(false);
167 }))
168 return false;
169 if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
171 [](msgpack::DocNode &SNode) {
172 return StringSwitch<bool>(SNode.getString())
173 .Case("read_only", true)
174 .Case("write_only", true)
175 .Case("read_write", true)
176 .Default(false);
177 }))
178 return false;
179 if (!verifyScalarEntry(ArgsMap, ".is_const", false,
181 return false;
182 if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
184 return false;
185 if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
187 return false;
188 if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
190 return false;
191
192 return true;
193}
194
195bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
196 if (!Node.isMap())
197 return false;
198 auto &KernelMap = Node.getMap();
199
200 if (!verifyScalarEntry(KernelMap, ".name", true,
202 return false;
203 if (!verifyScalarEntry(KernelMap, ".symbol", true,
205 return false;
206 if (!verifyScalarEntry(KernelMap, ".language", false,
208 [](msgpack::DocNode &SNode) {
209 return StringSwitch<bool>(SNode.getString())
210 .Case("OpenCL C", true)
211 .Case("OpenCL C++", true)
212 .Case("HCC", true)
213 .Case("HIP", true)
214 .Case("OpenMP", true)
215 .Case("Assembler", true)
216 .Default(false);
217 }))
218 return false;
219 if (!verifyEntry(
220 KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
221 return verifyArray(
222 Node,
223 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
224 }))
225 return false;
226 if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
227 return verifyArray(Node, [this](msgpack::DocNode &Node) {
228 return verifyKernelArgs(Node);
229 });
230 }))
231 return false;
232 if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
233 [this](msgpack::DocNode &Node) {
234 return verifyArray(Node,
235 [this](msgpack::DocNode &Node) {
236 return verifyInteger(Node);
237 },
238 3);
239 }))
240 return false;
241 if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
242 [this](msgpack::DocNode &Node) {
243 return verifyArray(Node,
244 [this](msgpack::DocNode &Node) {
245 return verifyInteger(Node);
246 },
247 3);
248 }))
249 return false;
250 if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
252 return false;
253 if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
255 return false;
256 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
257 return false;
258 if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
259 return false;
260 if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
261 return false;
262 if (!verifyScalarEntry(KernelMap, ".uses_dynamic_stack", false,
264 return false;
265 if (!verifyIntegerEntry(KernelMap, ".workgroup_processor_mode", false))
266 return false;
267 if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
268 return false;
269 if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
270 return false;
271 if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
272 return false;
273 if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
274 return false;
275 if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
276 return false;
277 if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
278 return false;
279 if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
280 return false;
281 if (!verifyIntegerEntry(KernelMap, ".uniform_work_group_size", false))
282 return false;
283
284
285 return true;
286}
287
289 if (!HSAMetadataRoot.isMap())
290 return false;
291 auto &RootMap = HSAMetadataRoot.getMap();
292
293 if (!verifyEntry(
294 RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
295 return verifyArray(
296 Node,
297 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
298 }))
299 return false;
300 if (!verifyEntry(
301 RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
302 return verifyArray(Node, [this](msgpack::DocNode &Node) {
303 return verifyScalar(Node, msgpack::Type::String);
304 });
305 }))
306 return false;
307 if (!verifyEntry(RootMap, "amdhsa.kernels", true,
308 [this](msgpack::DocNode &Node) {
309 return verifyArray(Node, [this](msgpack::DocNode &Node) {
310 return verifyKernel(Node);
311 });
312 }))
313 return false;
314
315 return true;
316}
317
318} // end namespace V3
319} // end namespace HSAMD
320} // end namespace AMDGPU
321} // end namespace llvm
This is a verifier for AMDGPU HSA metadata, which can verify both well-typed metadata and untyped met...
uint64_t Size
This file declares a class that exposes a simple in-memory representation of a document of MsgPack ob...
This file contains some templates that are useful if you are working with the STL at all.
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
Kind getKind() const
bool verify(msgpack::DocNode &HSAMetadataRoot)
Verify given HSA metadata.
A node in a MsgPack Document.
MapDocNode & getMap(bool Convert=false)
Get a MapDocNode for a map node.
Type
MessagePack types as defined in the standard, with the exception of Integer being divided into a sign...
Definition: MsgPackReader.h:48
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
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:1735