LLVM  12.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 #include "llvm/ADT/StringSwitch.h"
17 
18 namespace llvm {
19 namespace AMDGPU {
20 namespace HSAMD {
21 namespace V3 {
22 
23 bool MetadataVerifier::verifyScalar(
24  msgpack::DocNode &Node, msgpack::Type SKind,
25  function_ref<bool(msgpack::DocNode &)> verifyValue) {
26  if (!Node.isScalar())
27  return false;
28  if (Node.getKind() != SKind) {
29  if (Strict)
30  return false;
31  // If we are not strict, we interpret string values as "implicitly typed"
32  // and attempt to coerce them to the expected type here.
33  if (Node.getKind() != msgpack::Type::String)
34  return false;
35  StringRef StringValue = Node.getString();
36  Node.fromString(StringValue);
37  if (Node.getKind() != SKind)
38  return false;
39  }
40  if (verifyValue)
41  return verifyValue(Node);
42  return true;
43 }
44 
45 bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
46  if (!verifyScalar(Node, msgpack::Type::UInt))
47  if (!verifyScalar(Node, msgpack::Type::Int))
48  return false;
49  return true;
50 }
51 
52 bool MetadataVerifier::verifyArray(
53  msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
54  Optional<size_t> Size) {
55  if (!Node.isArray())
56  return false;
57  auto &Array = Node.getArray();
58  if (Size && Array.size() != *Size)
59  return false;
60  for (auto &Item : Array)
61  if (!verifyNode(Item))
62  return false;
63 
64  return true;
65 }
66 
67 bool 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 
76 bool 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, [=](msgpack::DocNode &Node) {
81  return verifyScalar(Node, SKind, verifyValue);
82  });
83 }
84 
85 bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
86  StringRef Key, bool Required) {
87  return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {
88  return verifyInteger(Node);
89  });
90 }
91 
92 bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {
93  if (!Node.isMap())
94  return false;
95  auto &ArgsMap = Node.getMap();
96 
97  if (!verifyScalarEntry(ArgsMap, ".name", false,
99  return false;
100  if (!verifyScalarEntry(ArgsMap, ".type_name", false,
102  return false;
103  if (!verifyIntegerEntry(ArgsMap, ".size", true))
104  return false;
105  if (!verifyIntegerEntry(ArgsMap, ".offset", true))
106  return false;
107  if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
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_global_offset_x", true)
119  .Case("hidden_global_offset_y", true)
120  .Case("hidden_global_offset_z", true)
121  .Case("hidden_none", true)
122  .Case("hidden_printf_buffer", true)
123  .Case("hidden_hostcall_buffer", true)
124  .Case("hidden_default_queue", true)
125  .Case("hidden_completion_action", true)
126  .Case("hidden_multigrid_sync_arg", true)
127  .Default(false);
128  }))
129  return false;
130  if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
131  return false;
132  if (!verifyScalarEntry(ArgsMap, ".address_space", false,
134  [](msgpack::DocNode &SNode) {
135  return StringSwitch<bool>(SNode.getString())
136  .Case("private", true)
137  .Case("global", true)
138  .Case("constant", true)
139  .Case("local", true)
140  .Case("generic", true)
141  .Case("region", true)
142  .Default(false);
143  }))
144  return false;
145  if (!verifyScalarEntry(ArgsMap, ".access", false,
147  [](msgpack::DocNode &SNode) {
148  return StringSwitch<bool>(SNode.getString())
149  .Case("read_only", true)
150  .Case("write_only", true)
151  .Case("read_write", true)
152  .Default(false);
153  }))
154  return false;
155  if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
157  [](msgpack::DocNode &SNode) {
158  return StringSwitch<bool>(SNode.getString())
159  .Case("read_only", true)
160  .Case("write_only", true)
161  .Case("read_write", true)
162  .Default(false);
163  }))
164  return false;
165  if (!verifyScalarEntry(ArgsMap, ".is_const", false,
167  return false;
168  if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
170  return false;
171  if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
173  return false;
174  if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
176  return false;
177 
178  return true;
179 }
180 
181 bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
182  if (!Node.isMap())
183  return false;
184  auto &KernelMap = Node.getMap();
185 
186  if (!verifyScalarEntry(KernelMap, ".name", true,
188  return false;
189  if (!verifyScalarEntry(KernelMap, ".symbol", true,
191  return false;
192  if (!verifyScalarEntry(KernelMap, ".language", false,
194  [](msgpack::DocNode &SNode) {
195  return StringSwitch<bool>(SNode.getString())
196  .Case("OpenCL C", true)
197  .Case("OpenCL C++", true)
198  .Case("HCC", true)
199  .Case("HIP", true)
200  .Case("OpenMP", true)
201  .Case("Assembler", true)
202  .Default(false);
203  }))
204  return false;
205  if (!verifyEntry(
206  KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
207  return verifyArray(
208  Node,
209  [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
210  }))
211  return false;
212  if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
213  return verifyArray(Node, [this](msgpack::DocNode &Node) {
214  return verifyKernelArgs(Node);
215  });
216  }))
217  return false;
218  if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
219  [this](msgpack::DocNode &Node) {
220  return verifyArray(Node,
221  [this](msgpack::DocNode &Node) {
222  return verifyInteger(Node);
223  },
224  3);
225  }))
226  return false;
227  if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
228  [this](msgpack::DocNode &Node) {
229  return verifyArray(Node,
230  [this](msgpack::DocNode &Node) {
231  return verifyInteger(Node);
232  },
233  3);
234  }))
235  return false;
236  if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
238  return false;
239  if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
241  return false;
242  if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
243  return false;
244  if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
245  return false;
246  if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
247  return false;
248  if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
249  return false;
250  if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
251  return false;
252  if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
253  return false;
254  if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
255  return false;
256  if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
257  return false;
258  if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
259  return false;
260  if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
261  return false;
262 
263  return true;
264 }
265 
267  if (!HSAMetadataRoot.isMap())
268  return false;
269  auto &RootMap = HSAMetadataRoot.getMap();
270 
271  if (!verifyEntry(
272  RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
273  return verifyArray(
274  Node,
275  [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
276  }))
277  return false;
278  if (!verifyEntry(
279  RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
280  return verifyArray(Node, [this](msgpack::DocNode &Node) {
281  return verifyScalar(Node, msgpack::Type::String);
282  });
283  }))
284  return false;
285  if (!verifyEntry(RootMap, "amdhsa.kernels", true,
286  [this](msgpack::DocNode &Node) {
287  return verifyArray(Node, [this](msgpack::DocNode &Node) {
288  return verifyKernel(Node);
289  });
290  }))
291  return false;
292 
293  return true;
294 }
295 
296 } // end namespace V3
297 } // end namespace HSAMD
298 } // end namespace AMDGPU
299 } // end namespace llvm
Type
MessagePack types as defined in the standard, with the exception of Integer being divided into a sign...
Definition: MsgPackReader.h:49
This class represents lattice values for constants.
Definition: AllocatorList.h:23
MapDocNode & getMap(bool Convert=false)
Get a MapDocNode for a map node.
Key
PAL metadata keys.
bool verify(msgpack::DocNode &HSAMetadataRoot)
Verify given HSA metadata.
This is a verifier for AMDGPU HSA metadata, which can verify both well-typed metadata and untyped met...
uint32_t Size
Definition: Profile.cpp:46
AMDGPU metadata definitions and in-memory representations.
A node in a MsgPack Document.