llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp

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