29 bool MetadataVerifier::verifyScalar(
31 function_ref<
bool(msgpack::DocNode &)> verifyValue) {
41 StringRef StringValue =
Node.getString();
42 Node.fromString(StringValue);
47 return verifyValue(
Node);
51 bool MetadataVerifier::verifyInteger(msgpack::DocNode &
Node) {
58 bool MetadataVerifier::verifyArray(
59 msgpack::DocNode &
Node, function_ref<
bool(msgpack::DocNode &)> verifyNode,
60 Optional<size_t> Size) {
64 if (Size &&
Array.size() != *Size)
69 bool MetadataVerifier::verifyEntry(
70 msgpack::MapDocNode &MapNode, StringRef
Key,
bool Required,
71 function_ref<
bool(msgpack::DocNode &)> verifyNode) {
72 auto Entry = MapNode.find(
Key);
73 if (Entry == MapNode.end())
75 return verifyNode(Entry->second);
78 bool MetadataVerifier::verifyScalarEntry(
79 msgpack::MapDocNode &MapNode, StringRef
Key,
bool Required,
81 function_ref<
bool(msgpack::DocNode &)> verifyValue) {
82 return verifyEntry(MapNode,
Key,
Required, [=](msgpack::DocNode &
Node) {
83 return verifyScalar(
Node, SKind, verifyValue);
87 bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
89 return verifyEntry(MapNode,
Key,
Required, [
this](msgpack::DocNode &
Node) {
90 return verifyInteger(
Node);
94 bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &
Node) {
97 auto &ArgsMap =
Node.getMap();
99 if (!verifyScalarEntry(ArgsMap,
".name",
false,
102 if (!verifyScalarEntry(ArgsMap,
".type_name",
false,
105 if (!verifyIntegerEntry(ArgsMap,
".size",
true))
107 if (!verifyIntegerEntry(ArgsMap,
".offset",
true))
110 [](msgpack::DocNode &SNode) {
111 return StringSwitch<bool>(SNode.getString())
112 .Case(
"by_value",
true)
113 .Case(
"global_buffer",
true)
114 .Case(
"dynamic_shared_pointer",
true)
115 .Case(
"sampler",
true)
119 .Case(
"hidden_block_count_x",
true)
120 .Case(
"hidden_block_count_y",
true)
121 .Case(
"hidden_block_count_z",
true)
122 .Case(
"hidden_group_size_x",
true)
123 .Case(
"hidden_group_size_y",
true)
124 .Case(
"hidden_group_size_z",
true)
125 .Case(
"hidden_remainder_x",
true)
126 .Case(
"hidden_remainder_y",
true)
127 .Case(
"hidden_remainder_z",
true)
128 .Case(
"hidden_global_offset_x",
true)
129 .Case(
"hidden_global_offset_y",
true)
130 .Case(
"hidden_global_offset_z",
true)
131 .Case(
"hidden_grid_dims",
true)
132 .Case(
"hidden_none",
true)
133 .Case(
"hidden_printf_buffer",
true)
134 .Case(
"hidden_hostcall_buffer",
true)
135 .Case(
"hidden_heap_v1",
true)
136 .Case(
"hidden_default_queue",
true)
137 .Case(
"hidden_completion_action",
true)
138 .Case(
"hidden_multigrid_sync_arg",
true)
139 .Case(
"hidden_private_base",
true)
140 .Case(
"hidden_shared_base",
true)
141 .Case(
"hidden_queue_ptr",
true)
145 if (!verifyIntegerEntry(ArgsMap,
".pointee_align",
false))
147 if (!verifyScalarEntry(ArgsMap,
".address_space",
false,
149 [](msgpack::DocNode &SNode) {
150 return StringSwitch<bool>(SNode.getString())
151 .Case(
"private",
true)
152 .Case(
"global",
true)
153 .Case(
"constant",
true)
155 .Case(
"generic",
true)
156 .Case(
"region",
true)
160 if (!verifyScalarEntry(ArgsMap,
".access",
false,
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)
170 if (!verifyScalarEntry(ArgsMap,
".actual_access",
false,
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)
180 if (!verifyScalarEntry(ArgsMap,
".is_const",
false,
183 if (!verifyScalarEntry(ArgsMap,
".is_restrict",
false,
186 if (!verifyScalarEntry(ArgsMap,
".is_volatile",
false,
189 if (!verifyScalarEntry(ArgsMap,
".is_pipe",
false,
196 bool MetadataVerifier::verifyKernel(msgpack::DocNode &
Node) {
199 auto &KernelMap =
Node.getMap();
201 if (!verifyScalarEntry(KernelMap,
".name",
true,
204 if (!verifyScalarEntry(KernelMap,
".symbol",
true,
207 if (!verifyScalarEntry(KernelMap,
".language",
false,
209 [](msgpack::DocNode &SNode) {
210 return StringSwitch<bool>(SNode.getString())
211 .Case(
"OpenCL C",
true)
212 .Case(
"OpenCL C++",
true)
215 .Case(
"OpenMP",
true)
216 .Case(
"Assembler",
true)
221 KernelMap,
".language_version",
false, [
this](msgpack::DocNode &
Node) {
224 [
this](msgpack::DocNode &
Node) {
return verifyInteger(
Node); }, 2);
227 if (!verifyEntry(KernelMap,
".args",
false, [
this](msgpack::DocNode &
Node) {
228 return verifyArray(
Node, [
this](msgpack::DocNode &
Node) {
229 return verifyKernelArgs(
Node);
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);
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);
251 if (!verifyScalarEntry(KernelMap,
".vec_type_hint",
false,
254 if (!verifyScalarEntry(KernelMap,
".device_enqueue_symbol",
false,
257 if (!verifyIntegerEntry(KernelMap,
".kernarg_segment_size",
true))
259 if (!verifyIntegerEntry(KernelMap,
".group_segment_fixed_size",
true))
261 if (!verifyIntegerEntry(KernelMap,
".private_segment_fixed_size",
true))
263 if (!verifyIntegerEntry(KernelMap,
".kernarg_segment_align",
true))
265 if (!verifyIntegerEntry(KernelMap,
".wavefront_size",
true))
267 if (!verifyIntegerEntry(KernelMap,
".sgpr_count",
true))
269 if (!verifyIntegerEntry(KernelMap,
".vgpr_count",
true))
271 if (!verifyIntegerEntry(KernelMap,
".max_flat_workgroup_size",
true))
273 if (!verifyIntegerEntry(KernelMap,
".sgpr_spill_count",
false))
275 if (!verifyIntegerEntry(KernelMap,
".vgpr_spill_count",
false))
282 if (!HSAMetadataRoot.
isMap())
284 auto &RootMap = HSAMetadataRoot.
getMap();
300 if (!verifyEntry(RootMap,
"amdhsa.kernels",
true,
303 return verifyKernel(
Node);