LLVM  8.0.1
AMDGPUMetadataVerifier.cpp
Go to the documentation of this file.
1 //===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===//
2 //
3 // The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 //
10 /// \file
11 /// Implements a verifier for AMDGPU HSA metadata.
12 //
13 //===----------------------------------------------------------------------===//
14 
15 #include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
16 #include "llvm/Support/AMDGPUMetadata.h"
17 
18 namespace llvm {
19 namespace AMDGPU {
20 namespace HSAMD {
21 namespace V3 {
22 
23 bool MetadataVerifier::verifyScalar(
24  msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind,
25  function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
26  auto ScalarPtr = dyn_cast<msgpack::ScalarNode>(&Node);
27  if (!ScalarPtr)
28  return false;
29  auto &Scalar = *ScalarPtr;
30  // Do not output extraneous tags for types we know from the spec.
31  Scalar.IgnoreTag = true;
32  if (Scalar.getScalarKind() != 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 (Scalar.getScalarKind() != msgpack::ScalarNode::SK_String)
38  return false;
39  std::string StringValue = Scalar.getString();
40  Scalar.setScalarKind(SKind);
41  if (Scalar.inputYAML(StringValue) != StringRef())
42  return false;
43  }
44  if (verifyValue)
45  return verifyValue(Scalar);
46  return true;
47 }
48 
49 bool MetadataVerifier::verifyInteger(msgpack::Node &Node) {
50  if (!verifyScalar(Node, msgpack::ScalarNode::SK_UInt))
51  if (!verifyScalar(Node, msgpack::ScalarNode::SK_Int))
52  return false;
53  return true;
54 }
55 
56 bool MetadataVerifier::verifyArray(
57  msgpack::Node &Node, function_ref<bool(msgpack::Node &)> verifyNode,
58  Optional<size_t> Size) {
59  auto ArrayPtr = dyn_cast<msgpack::ArrayNode>(&Node);
60  if (!ArrayPtr)
61  return false;
62  auto &Array = *ArrayPtr;
63  if (Size && Array.size() != *Size)
64  return false;
65  for (auto &Item : Array)
66  if (!verifyNode(*Item.get()))
67  return false;
68 
69  return true;
70 }
71 
72 bool MetadataVerifier::verifyEntry(
73  msgpack::MapNode &MapNode, StringRef Key, bool Required,
74  function_ref<bool(msgpack::Node &)> verifyNode) {
75  auto Entry = MapNode.find(Key);
76  if (Entry == MapNode.end())
77  return !Required;
78  return verifyNode(*Entry->second.get());
79 }
80 
81 bool MetadataVerifier::verifyScalarEntry(
82  msgpack::MapNode &MapNode, StringRef Key, bool Required,
83  msgpack::ScalarNode::ScalarKind SKind,
84  function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
85  return verifyEntry(MapNode, Key, Required, [=](msgpack::Node &Node) {
86  return verifyScalar(Node, SKind, verifyValue);
87  });
88 }
89 
90 bool MetadataVerifier::verifyIntegerEntry(msgpack::MapNode &MapNode,
91  StringRef Key, bool Required) {
92  return verifyEntry(MapNode, Key, Required, [this](msgpack::Node &Node) {
93  return verifyInteger(Node);
94  });
95 }
96 
97 bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
98  auto ArgsMapPtr = dyn_cast<msgpack::MapNode>(&Node);
99  if (!ArgsMapPtr)
100  return false;
101  auto &ArgsMap = *ArgsMapPtr;
102 
103  if (!verifyScalarEntry(ArgsMap, ".name", false,
104  msgpack::ScalarNode::SK_String))
105  return false;
106  if (!verifyScalarEntry(ArgsMap, ".type_name", false,
107  msgpack::ScalarNode::SK_String))
108  return false;
109  if (!verifyIntegerEntry(ArgsMap, ".size", true))
110  return false;
111  if (!verifyIntegerEntry(ArgsMap, ".offset", true))
112  return false;
113  if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
114  msgpack::ScalarNode::SK_String,
115  [](msgpack::ScalarNode &SNode) {
116  return StringSwitch<bool>(SNode.getString())
117  .Case("by_value", true)
118  .Case("global_buffer", true)
119  .Case("dynamic_shared_pointer", true)
120  .Case("sampler", true)
121  .Case("image", true)
122  .Case("pipe", true)
123  .Case("queue", true)
124  .Case("hidden_global_offset_x", true)
125  .Case("hidden_global_offset_y", true)
126  .Case("hidden_global_offset_z", true)
127  .Case("hidden_none", true)
128  .Case("hidden_printf_buffer", true)
129  .Case("hidden_default_queue", true)
130  .Case("hidden_completion_action", true)
131  .Default(false);
132  }))
133  return false;
134  if (!verifyScalarEntry(ArgsMap, ".value_type", true,
135  msgpack::ScalarNode::SK_String,
136  [](msgpack::ScalarNode &SNode) {
137  return StringSwitch<bool>(SNode.getString())
138  .Case("struct", true)
139  .Case("i8", true)
140  .Case("u8", true)
141  .Case("i16", true)
142  .Case("u16", true)
143  .Case("f16", true)
144  .Case("i32", true)
145  .Case("u32", true)
146  .Case("f32", true)
147  .Case("i64", true)
148  .Case("u64", true)
149  .Case("f64", true)
150  .Default(false);
151  }))
152  return false;
153  if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
154  return false;
155  if (!verifyScalarEntry(ArgsMap, ".address_space", false,
156  msgpack::ScalarNode::SK_String,
157  [](msgpack::ScalarNode &SNode) {
158  return StringSwitch<bool>(SNode.getString())
159  .Case("private", true)
160  .Case("global", true)
161  .Case("constant", true)
162  .Case("local", true)
163  .Case("generic", true)
164  .Case("region", true)
165  .Default(false);
166  }))
167  return false;
168  if (!verifyScalarEntry(ArgsMap, ".access", false,
169  msgpack::ScalarNode::SK_String,
170  [](msgpack::ScalarNode &SNode) {
171  return StringSwitch<bool>(SNode.getString())
172  .Case("read_only", true)
173  .Case("write_only", true)
174  .Case("read_write", true)
175  .Default(false);
176  }))
177  return false;
178  if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
179  msgpack::ScalarNode::SK_String,
180  [](msgpack::ScalarNode &SNode) {
181  return StringSwitch<bool>(SNode.getString())
182  .Case("read_only", true)
183  .Case("write_only", true)
184  .Case("read_write", true)
185  .Default(false);
186  }))
187  return false;
188  if (!verifyScalarEntry(ArgsMap, ".is_const", false,
189  msgpack::ScalarNode::SK_Boolean))
190  return false;
191  if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
192  msgpack::ScalarNode::SK_Boolean))
193  return false;
194  if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
195  msgpack::ScalarNode::SK_Boolean))
196  return false;
197  if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
198  msgpack::ScalarNode::SK_Boolean))
199  return false;
200 
201  return true;
202 }
203 
204 bool MetadataVerifier::verifyKernel(msgpack::Node &Node) {
205  auto KernelMapPtr = dyn_cast<msgpack::MapNode>(&Node);
206  if (!KernelMapPtr)
207  return false;
208  auto &KernelMap = *KernelMapPtr;
209 
210  if (!verifyScalarEntry(KernelMap, ".name", true,
211  msgpack::ScalarNode::SK_String))
212  return false;
213  if (!verifyScalarEntry(KernelMap, ".symbol", true,
214  msgpack::ScalarNode::SK_String))
215  return false;
216  if (!verifyScalarEntry(KernelMap, ".language", false,
217  msgpack::ScalarNode::SK_String,
218  [](msgpack::ScalarNode &SNode) {
219  return StringSwitch<bool>(SNode.getString())
220  .Case("OpenCL C", true)
221  .Case("OpenCL C++", true)
222  .Case("HCC", true)
223  .Case("HIP", true)
224  .Case("OpenMP", true)
225  .Case("Assembler", true)
226  .Default(false);
227  }))
228  return false;
229  if (!verifyEntry(
230  KernelMap, ".language_version", false, [this](msgpack::Node &Node) {
231  return verifyArray(
232  Node,
233  [this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
234  }))
235  return false;
236  if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::Node &Node) {
237  return verifyArray(Node, [this](msgpack::Node &Node) {
238  return verifyKernelArgs(Node);
239  });
240  }))
241  return false;
242  if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
243  [this](msgpack::Node &Node) {
244  return verifyArray(Node,
245  [this](msgpack::Node &Node) {
246  return verifyInteger(Node);
247  },
248  3);
249  }))
250  return false;
251  if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
252  [this](msgpack::Node &Node) {
253  return verifyArray(Node,
254  [this](msgpack::Node &Node) {
255  return verifyInteger(Node);
256  },
257  3);
258  }))
259  return false;
260  if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
261  msgpack::ScalarNode::SK_String))
262  return false;
263  if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
264  msgpack::ScalarNode::SK_String))
265  return false;
266  if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
267  return false;
268  if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
269  return false;
270  if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
271  return false;
272  if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
273  return false;
274  if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
275  return false;
276  if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
277  return false;
278  if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
279  return false;
280  if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
281  return false;
282  if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
283  return false;
284  if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
285  return false;
286 
287  return true;
288 }
289 
290 bool MetadataVerifier::verify(msgpack::Node &HSAMetadataRoot) {
291  auto RootMapPtr = dyn_cast<msgpack::MapNode>(&HSAMetadataRoot);
292  if (!RootMapPtr)
293  return false;
294  auto &RootMap = *RootMapPtr;
295 
296  if (!verifyEntry(
297  RootMap, "amdhsa.version", true, [this](msgpack::Node &Node) {
298  return verifyArray(
299  Node,
300  [this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
301  }))
302  return false;
303  if (!verifyEntry(
304  RootMap, "amdhsa.printf", false, [this](msgpack::Node &Node) {
305  return verifyArray(Node, [this](msgpack::Node &Node) {
306  return verifyScalar(Node, msgpack::ScalarNode::SK_String);
307  });
308  }))
309  return false;
310  if (!verifyEntry(RootMap, "amdhsa.kernels", true,
311  [this](msgpack::Node &Node) {
312  return verifyArray(Node, [this](msgpack::Node &Node) {
313  return verifyKernel(Node);
314  });
315  }))
316  return false;
317 
318  return true;
319 }
320 
321 } // end namespace V3
322 } // end namespace HSAMD
323 } // end namespace AMDGPU
324 } // end namespace llvm
llvm
This class represents lattice values for constants.
Definition: AllocatorList.h:24
llvm::AMDGPU::PALMD::Key
Key
PAL metadata keys.
Definition: AMDGPUMetadata.h:460
llvm::AMDGPU::HSAMD::V3::MetadataVerifier::verify
bool verify(msgpack::Node &HSAMetadataRoot)
Verify given HSA metadata.
Definition: AMDGPUMetadataVerifier.cpp:290
llvm::msgpack::Node
Abstract base-class which can be any MessagePack type.
Definition: MsgPackTypes.h:39
AMDGPUMetadataVerifier.h
This is a verifier for AMDGPU HSA metadata, which can verify both well-typed metadata and untyped met...
llvm::dyn_cast
LLVM_NODISCARD std::enable_if<!is_simple_type< Y >::value, typename cast_retty< X, const Y >::ret_type >::type dyn_cast(const Y &Val)
Definition: Casting.h:323
Size
uint32_t Size
Definition: Profile.cpp:47
AMDGPUMetadata.h
AMDGPU metadata definitions and in-memory representations.
Generated on Sun Dec 20 2020 13:53:52 for LLVM by   doxygen 1.8.13