Coverage Report

Created: 2019-07-24 05:18

/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