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