mirror of
https://github.com/RPCS3/llvm-mirror.git
synced 2024-11-23 11:13:28 +01:00
80428fb35f
Lots of headers pass around MemoryBuffer objects, but very few open
them. Let those that do include FileSystem.h.
Saves ~250 includes of Chrono.h & FileSystem.h:
$ diff -u thedeps-before.txt thedeps-after.txt | grep '^[-+] ' | sort | uniq -c | sort -nr
254 - ../llvm/include/llvm/Support/FileSystem.h
253 - ../llvm/include/llvm/Support/Chrono.h
237 - ../llvm/include/llvm/Support/NativeFormatting.h
237 - ../llvm/include/llvm/Support/FormatProviders.h
192 - ../llvm/include/llvm/ADT/StringSwitch.h
190 - ../llvm/include/llvm/Support/FormatVariadicDetails.h
...
This requires duplicating the file_t typedef, which is unfortunate. I
sunk the choice of mapping mode down into the cpp file using variable
template specializations instead of class members in headers.
319 lines
12 KiB
C++
319 lines
12 KiB
C++
//===- AMDGPUMetadataVerifier.cpp - MsgPack Types ---------------*- C++ -*-===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
//
|
|
/// \file
|
|
/// Implements a verifier for AMDGPU HSA metadata.
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
|
|
#include "llvm/ADT/StringSwitch.h"
|
|
#include "llvm/Support/AMDGPUMetadata.h"
|
|
|
|
namespace llvm {
|
|
namespace AMDGPU {
|
|
namespace HSAMD {
|
|
namespace V3 {
|
|
|
|
bool MetadataVerifier::verifyScalar(
|
|
msgpack::DocNode &Node, msgpack::Type SKind,
|
|
function_ref<bool(msgpack::DocNode &)> verifyValue) {
|
|
if (!Node.isScalar())
|
|
return false;
|
|
if (Node.getKind() != SKind) {
|
|
if (Strict)
|
|
return false;
|
|
// If we are not strict, we interpret string values as "implicitly typed"
|
|
// and attempt to coerce them to the expected type here.
|
|
if (Node.getKind() != msgpack::Type::String)
|
|
return false;
|
|
StringRef StringValue = Node.getString();
|
|
Node.fromString(StringValue);
|
|
if (Node.getKind() != SKind)
|
|
return false;
|
|
}
|
|
if (verifyValue)
|
|
return verifyValue(Node);
|
|
return true;
|
|
}
|
|
|
|
bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
|
|
if (!verifyScalar(Node, msgpack::Type::UInt))
|
|
if (!verifyScalar(Node, msgpack::Type::Int))
|
|
return false;
|
|
return true;
|
|
}
|
|
|
|
bool MetadataVerifier::verifyArray(
|
|
msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
|
|
Optional<size_t> Size) {
|
|
if (!Node.isArray())
|
|
return false;
|
|
auto &Array = Node.getArray();
|
|
if (Size && Array.size() != *Size)
|
|
return false;
|
|
for (auto &Item : Array)
|
|
if (!verifyNode(Item))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
bool MetadataVerifier::verifyEntry(
|
|
msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
|
|
function_ref<bool(msgpack::DocNode &)> verifyNode) {
|
|
auto Entry = MapNode.find(Key);
|
|
if (Entry == MapNode.end())
|
|
return !Required;
|
|
return verifyNode(Entry->second);
|
|
}
|
|
|
|
bool MetadataVerifier::verifyScalarEntry(
|
|
msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
|
|
msgpack::Type SKind,
|
|
function_ref<bool(msgpack::DocNode &)> verifyValue) {
|
|
return verifyEntry(MapNode, Key, Required, [=](msgpack::DocNode &Node) {
|
|
return verifyScalar(Node, SKind, verifyValue);
|
|
});
|
|
}
|
|
|
|
bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
|
|
StringRef Key, bool Required) {
|
|
return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {
|
|
return verifyInteger(Node);
|
|
});
|
|
}
|
|
|
|
bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {
|
|
if (!Node.isMap())
|
|
return false;
|
|
auto &ArgsMap = Node.getMap();
|
|
|
|
if (!verifyScalarEntry(ArgsMap, ".name", false,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".type_name", false,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyIntegerEntry(ArgsMap, ".size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(ArgsMap, ".offset", true))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("by_value", true)
|
|
.Case("global_buffer", true)
|
|
.Case("dynamic_shared_pointer", true)
|
|
.Case("sampler", true)
|
|
.Case("image", true)
|
|
.Case("pipe", true)
|
|
.Case("queue", true)
|
|
.Case("hidden_global_offset_x", true)
|
|
.Case("hidden_global_offset_y", true)
|
|
.Case("hidden_global_offset_z", true)
|
|
.Case("hidden_none", true)
|
|
.Case("hidden_printf_buffer", true)
|
|
.Case("hidden_hostcall_buffer", true)
|
|
.Case("hidden_default_queue", true)
|
|
.Case("hidden_completion_action", true)
|
|
.Case("hidden_multigrid_sync_arg", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".value_type", true,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("struct", true)
|
|
.Case("i8", true)
|
|
.Case("u8", true)
|
|
.Case("i16", true)
|
|
.Case("u16", true)
|
|
.Case("f16", true)
|
|
.Case("i32", true)
|
|
.Case("u32", true)
|
|
.Case("f32", true)
|
|
.Case("i64", true)
|
|
.Case("u64", true)
|
|
.Case("f64", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".address_space", false,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("private", true)
|
|
.Case("global", true)
|
|
.Case("constant", true)
|
|
.Case("local", true)
|
|
.Case("generic", true)
|
|
.Case("region", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".access", false,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("read_only", true)
|
|
.Case("write_only", true)
|
|
.Case("read_write", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("read_only", true)
|
|
.Case("write_only", true)
|
|
.Case("read_write", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".is_const", false,
|
|
msgpack::Type::Boolean))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
|
|
msgpack::Type::Boolean))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
|
|
msgpack::Type::Boolean))
|
|
return false;
|
|
if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
|
|
msgpack::Type::Boolean))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
|
|
if (!Node.isMap())
|
|
return false;
|
|
auto &KernelMap = Node.getMap();
|
|
|
|
if (!verifyScalarEntry(KernelMap, ".name", true,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyScalarEntry(KernelMap, ".symbol", true,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyScalarEntry(KernelMap, ".language", false,
|
|
msgpack::Type::String,
|
|
[](msgpack::DocNode &SNode) {
|
|
return StringSwitch<bool>(SNode.getString())
|
|
.Case("OpenCL C", true)
|
|
.Case("OpenCL C++", true)
|
|
.Case("HCC", true)
|
|
.Case("HIP", true)
|
|
.Case("OpenMP", true)
|
|
.Case("Assembler", true)
|
|
.Default(false);
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(
|
|
KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
|
|
return verifyArray(
|
|
Node,
|
|
[this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
|
|
return verifyArray(Node, [this](msgpack::DocNode &Node) {
|
|
return verifyKernelArgs(Node);
|
|
});
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
|
|
[this](msgpack::DocNode &Node) {
|
|
return verifyArray(Node,
|
|
[this](msgpack::DocNode &Node) {
|
|
return verifyInteger(Node);
|
|
},
|
|
3);
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
|
|
[this](msgpack::DocNode &Node) {
|
|
return verifyArray(Node,
|
|
[this](msgpack::DocNode &Node) {
|
|
return verifyInteger(Node);
|
|
},
|
|
3);
|
|
}))
|
|
return false;
|
|
if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
|
|
msgpack::Type::String))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".max_flat_workgroup_size", true))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
|
|
return false;
|
|
if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
|
|
if (!HSAMetadataRoot.isMap())
|
|
return false;
|
|
auto &RootMap = HSAMetadataRoot.getMap();
|
|
|
|
if (!verifyEntry(
|
|
RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
|
|
return verifyArray(
|
|
Node,
|
|
[this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(
|
|
RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
|
|
return verifyArray(Node, [this](msgpack::DocNode &Node) {
|
|
return verifyScalar(Node, msgpack::Type::String);
|
|
});
|
|
}))
|
|
return false;
|
|
if (!verifyEntry(RootMap, "amdhsa.kernels", true,
|
|
[this](msgpack::DocNode &Node) {
|
|
return verifyArray(Node, [this](msgpack::DocNode &Node) {
|
|
return verifyKernel(Node);
|
|
});
|
|
}))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
} // end namespace V3
|
|
} // end namespace HSAMD
|
|
} // end namespace AMDGPU
|
|
} // end namespace llvm
|