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
285
286 return true;
287}
288
289bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
290 if (!HSAMetadataRoot.isMap())
291 return false;
292 auto &RootMap = HSAMetadataRoot.getMap();
293
294 if (!verifyEntry(
295 RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
296 return verifyArray(
297 Node,
298 [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
299 }))
300 return false;
301 if (!verifyEntry(
302 RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
303 return verifyArray(Node, [this](msgpack::DocNode &Node) {
304 return verifyScalar(Node, msgpack::Type::String);
305 });
306 }))
307 return false;
308 if (!verifyEntry(RootMap, "amdhsa.kernels", true,
309 [this](msgpack::DocNode &Node) {
310 return verifyArray(Node, [this](msgpack::DocNode &Node) {
311 return verifyKernel(Node);
312 });
313 }))
314 return false;
315
316 return true;
317}
318
319} // end namespace V3
320} // end namespace HSAMD
321} // end namespace AMDGPU
322} // end namespace llvm
AMDGPUMetadataVerifier.h
This is a verifier for AMDGPU HSA metadata, which can verify both well-typed metadata and untyped met...
Size
uint64_t Size
Definition: ELFObjHandler.cpp:81
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...
Node::getKind
Kind getKind() const
Definition: ItaniumDemangle.h:258
llvm::AMDGPU::HSAMD::V3::MetadataVerifier::verify
LLVM_ABI bool verify(msgpack::DocNode &HSAMetadataRoot)
Verify given HSA metadata.
Definition: AMDGPUMetadataVerifier.cpp:289
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
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:1744
Generated on Wed Sep 3 2025 02:35:15 for LLVM by doxygen 1.9.6