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/ADT/StringSwitch.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::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,
  98                          msgpack::Type::String))
  99     return false;
 100   if (!verifyScalarEntry(ArgsMap, ".type_name", false,
 101                          msgpack::Type::String))
 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,
 108                          msgpack::Type::String,
 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,
 133                          msgpack::Type::String,
 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,
 146                          msgpack::Type::String,
 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,
 156                          msgpack::Type::String,
 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,
 166                          msgpack::Type::Boolean))
 167     return false;
 168   if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
 169                          msgpack::Type::Boolean))
 170     return false;
 171   if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
 172                          msgpack::Type::Boolean))
 173     return false;
 174   if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
 175                          msgpack::Type::Boolean))
 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,
 187                          msgpack::Type::String))
 188     return false;
 189   if (!verifyScalarEntry(KernelMap, ".symbol", true,
 190                          msgpack::Type::String))
 191     return false;
 192   if (!verifyScalarEntry(KernelMap, ".language", false,
 193                          msgpack::Type::String,
 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,
 237                          msgpack::Type::String))
 238     return false;
 239   if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
 240                          msgpack::Type::String))
 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
 266 bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
 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