program_generated.h

Go to the documentation of this file.

// automatically generated by the FlatBuffers compiler, do not modify
 
 
#ifndef FLATBUFFERS_GENERATED_PROGRAM_TT_TARGET_METAL_H_
#define FLATBUFFERS_GENERATED_PROGRAM_TT_TARGET_METAL_H_
 
#include "flatbuffers/flatbuffers.h"
 
// Ensure the included flatbuffers.h is the same version as when this file was
// generated, otherwise it may not be compatible.
static_assert(FLATBUFFERS_VERSION_MAJOR == 24 &&
              FLATBUFFERS_VERSION_MINOR == 3 &&
              FLATBUFFERS_VERSION_REVISION == 25,
             "Non-compatible flatbuffers version included");
 
#include "types_generated.h"
 
namespace tt {
namespace target {
namespace metal {
 
struct NocConfig;
struct NocConfigBuilder;
 
struct TensixConfig;
struct TensixConfigBuilder;
 
struct EthernetConfig;
struct EthernetConfigBuilder;
 
struct KernelSource;
struct KernelSourceBuilder;
 
struct KernelBinary;
struct KernelBinaryBuilder;
 
struct RuntimeArgTensorAddress;
struct RuntimeArgTensorAddressBuilder;
 
struct RuntimeArgSemaphoreAddress;
struct RuntimeArgSemaphoreAddressBuilder;
 
struct KernelDesc;
struct KernelDescBuilder;
 
struct ProgramDesc;
struct ProgramDescBuilder;
 
enum class NocIndex : uint16_t {
  Noc0 = 0,
  Noc1 = 1,
  MIN = Noc0,
  MAX = Noc1
};
 
inline const NocIndex (&EnumValuesNocIndex())[2] {
  static const NocIndex values[] = {
    NocIndex::Noc0,
    NocIndex::Noc1
  };
  return values;
}
 
inline const char * const *EnumNamesNocIndex() {
  static const char * const names[3] = {
    "Noc0",
    "Noc1",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameNocIndex(NocIndex e) {
  if (::flatbuffers::IsOutRange(e, NocIndex::Noc0, NocIndex::Noc1)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesNocIndex()[index];
}
 
enum class EthType : uint16_t {
  Sender = 0,
  Receiver = 1,
  MIN = Sender,
  MAX = Receiver
};
 
inline const EthType (&EnumValuesEthType())[2] {
  static const EthType values[] = {
    EthType::Sender,
    EthType::Receiver
  };
  return values;
}
 
inline const char * const *EnumNamesEthType() {
  static const char * const names[3] = {
    "Sender",
    "Receiver",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameEthType(EthType e) {
  if (::flatbuffers::IsOutRange(e, EthType::Sender, EthType::Receiver)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesEthType()[index];
}
 
enum class KernelConfig : uint8_t {
  NONE = 0,
  NocConfig = 1,
  TensixConfig = 2,
  EthernetConfig = 3,
  MIN = NONE,
  MAX = EthernetConfig
};
 
inline const KernelConfig (&EnumValuesKernelConfig())[4] {
  static const KernelConfig values[] = {
    KernelConfig::NONE,
    KernelConfig::NocConfig,
    KernelConfig::TensixConfig,
    KernelConfig::EthernetConfig
  };
  return values;
}
 
inline const char * const *EnumNamesKernelConfig() {
  static const char * const names[5] = {
    "NONE",
    "NocConfig",
    "TensixConfig",
    "EthernetConfig",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameKernelConfig(KernelConfig e) {
  if (::flatbuffers::IsOutRange(e, KernelConfig::NONE, KernelConfig::EthernetConfig)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesKernelConfig()[index];
}
 
template<typename T> struct KernelConfigTraits {
  static const KernelConfig enum_value = KernelConfig::NONE;
};
 
template<> struct KernelConfigTraits<tt::target::metal::NocConfig> {
  static const KernelConfig enum_value = KernelConfig::NocConfig;
};
 
template<> struct KernelConfigTraits<tt::target::metal::TensixConfig> {
  static const KernelConfig enum_value = KernelConfig::TensixConfig;
};
 
template<> struct KernelConfigTraits<tt::target::metal::EthernetConfig> {
  static const KernelConfig enum_value = KernelConfig::EthernetConfig;
};
 
bool VerifyKernelConfig(::flatbuffers::Verifier &verifier, const void *obj, KernelConfig type);
bool VerifyKernelConfigVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelConfig> *types);
 
enum class BinaryType : uint16_t {
  BRISC = 0,
  NCRISC = 1,
  TRISC0 = 2,
  TRISC1 = 3,
  TRISC2 = 4,
  ERISC = 5,
  MIN = BRISC,
  MAX = ERISC
};
 
inline const BinaryType (&EnumValuesBinaryType())[6] {
  static const BinaryType values[] = {
    BinaryType::BRISC,
    BinaryType::NCRISC,
    BinaryType::TRISC0,
    BinaryType::TRISC1,
    BinaryType::TRISC2,
    BinaryType::ERISC
  };
  return values;
}
 
inline const char * const *EnumNamesBinaryType() {
  static const char * const names[7] = {
    "BRISC",
    "NCRISC",
    "TRISC0",
    "TRISC1",
    "TRISC2",
    "ERISC",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameBinaryType(BinaryType e) {
  if (::flatbuffers::IsOutRange(e, BinaryType::BRISC, BinaryType::ERISC)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesBinaryType()[index];
}
 
enum class CoreType : uint16_t {
  WORKER = 0,
  ETH = 1,
  MIN = WORKER,
  MAX = ETH
};
 
inline const CoreType (&EnumValuesCoreType())[2] {
  static const CoreType values[] = {
    CoreType::WORKER,
    CoreType::ETH
  };
  return values;
}
 
inline const char * const *EnumNamesCoreType() {
  static const char * const names[3] = {
    "WORKER",
    "ETH",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameCoreType(CoreType e) {
  if (::flatbuffers::IsOutRange(e, CoreType::WORKER, CoreType::ETH)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesCoreType()[index];
}
 
enum class Kernel : uint8_t {
  NONE = 0,
  KernelSource = 1,
  KernelBinary = 2,
  MIN = NONE,
  MAX = KernelBinary
};
 
inline const Kernel (&EnumValuesKernel())[3] {
  static const Kernel values[] = {
    Kernel::NONE,
    Kernel::KernelSource,
    Kernel::KernelBinary
  };
  return values;
}
 
inline const char * const *EnumNamesKernel() {
  static const char * const names[4] = {
    "NONE",
    "KernelSource",
    "KernelBinary",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameKernel(Kernel e) {
  if (::flatbuffers::IsOutRange(e, Kernel::NONE, Kernel::KernelBinary)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesKernel()[index];
}
 
template<typename T> struct KernelTraits {
  static const Kernel enum_value = Kernel::NONE;
};
 
template<> struct KernelTraits<tt::target::metal::KernelSource> {
  static const Kernel enum_value = Kernel::KernelSource;
};
 
template<> struct KernelTraits<tt::target::metal::KernelBinary> {
  static const Kernel enum_value = Kernel::KernelBinary;
};
 
bool VerifyKernel(::flatbuffers::Verifier &verifier, const void *obj, Kernel type);
bool VerifyKernelVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<Kernel> *types);
 
enum class RuntimeArg : uint8_t {
  NONE = 0,
  RuntimeArgTensorAddress = 1,
  RuntimeArgSemaphoreAddress = 2,
  MIN = NONE,
  MAX = RuntimeArgSemaphoreAddress
};
 
inline const RuntimeArg (&EnumValuesRuntimeArg())[3] {
  static const RuntimeArg values[] = {
    RuntimeArg::NONE,
    RuntimeArg::RuntimeArgTensorAddress,
    RuntimeArg::RuntimeArgSemaphoreAddress
  };
  return values;
}
 
inline const char * const *EnumNamesRuntimeArg() {
  static const char * const names[4] = {
    "NONE",
    "RuntimeArgTensorAddress",
    "RuntimeArgSemaphoreAddress",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameRuntimeArg(RuntimeArg e) {
  if (::flatbuffers::IsOutRange(e, RuntimeArg::NONE, RuntimeArg::RuntimeArgSemaphoreAddress)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesRuntimeArg()[index];
}
 
template<typename T> struct RuntimeArgTraits {
  static const RuntimeArg enum_value = RuntimeArg::NONE;
};
 
template<> struct RuntimeArgTraits<tt::target::metal::RuntimeArgTensorAddress> {
  static const RuntimeArg enum_value = RuntimeArg::RuntimeArgTensorAddress;
};
 
template<> struct RuntimeArgTraits<tt::target::metal::RuntimeArgSemaphoreAddress> {
  static const RuntimeArg enum_value = RuntimeArg::RuntimeArgSemaphoreAddress;
};
 
bool VerifyRuntimeArg(::flatbuffers::Verifier &verifier, const void *obj, RuntimeArg type);
bool VerifyRuntimeArgVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<RuntimeArg> *types);
 
struct NocConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef NocConfigBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_NOC_INDEX = 4
  };
  tt::target::metal::NocIndex noc_index() const {
    return static_cast<tt::target::metal::NocIndex>(GetField<uint16_t>(VT_NOC_INDEX, 0));
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint16_t>(verifier, VT_NOC_INDEX, 2) &&
           verifier.EndTable();
  }
};
 
struct NocConfigBuilder {
  typedef NocConfig Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_noc_index(tt::target::metal::NocIndex noc_index) {
    fbb_.AddElement<uint16_t>(NocConfig::VT_NOC_INDEX, static_cast<uint16_t>(noc_index), 0);
  }
  explicit NocConfigBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<NocConfig> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<NocConfig>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<NocConfig> CreateNocConfig(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::NocIndex noc_index = tt::target::metal::NocIndex::Noc0) {
  NocConfigBuilder builder_(_fbb);
  builder_.add_noc_index(noc_index);
  return builder_.Finish();
}
 
struct NocConfig::Traits {
  using type = NocConfig;
  static auto constexpr Create = CreateNocConfig;
};
 
struct TensixConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef TensixConfigBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_MATH_FIDELITY = 4,
    VT_FP32_DEST_ACC_EN = 6,
    VT_PRESERVE_FP32_PRECISION = 8,
    VT_MATH_APPROX_MODE = 10
  };
  tt::target::MathFidelity math_fidelity() const {
    return static_cast<tt::target::MathFidelity>(GetField<uint8_t>(VT_MATH_FIDELITY, 0));
  }
  bool fp32_dest_acc_en() const {
    return GetField<uint8_t>(VT_FP32_DEST_ACC_EN, 0) != 0;
  }
  bool preserve_fp32_precision() const {
    return GetField<uint8_t>(VT_PRESERVE_FP32_PRECISION, 0) != 0;
  }
  bool math_approx_mode() const {
    return GetField<uint8_t>(VT_MATH_APPROX_MODE, 0) != 0;
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint8_t>(verifier, VT_MATH_FIDELITY, 1) &&
           VerifyField<uint8_t>(verifier, VT_FP32_DEST_ACC_EN, 1) &&
           VerifyField<uint8_t>(verifier, VT_PRESERVE_FP32_PRECISION, 1) &&
           VerifyField<uint8_t>(verifier, VT_MATH_APPROX_MODE, 1) &&
           verifier.EndTable();
  }
};
 
struct TensixConfigBuilder {
  typedef TensixConfig Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_math_fidelity(tt::target::MathFidelity math_fidelity) {
    fbb_.AddElement<uint8_t>(TensixConfig::VT_MATH_FIDELITY, static_cast<uint8_t>(math_fidelity), 0);
  }
  void add_fp32_dest_acc_en(bool fp32_dest_acc_en) {
    fbb_.AddElement<uint8_t>(TensixConfig::VT_FP32_DEST_ACC_EN, static_cast<uint8_t>(fp32_dest_acc_en), 0);
  }
  void add_preserve_fp32_precision(bool preserve_fp32_precision) {
    fbb_.AddElement<uint8_t>(TensixConfig::VT_PRESERVE_FP32_PRECISION, static_cast<uint8_t>(preserve_fp32_precision), 0);
  }
  void add_math_approx_mode(bool math_approx_mode) {
    fbb_.AddElement<uint8_t>(TensixConfig::VT_MATH_APPROX_MODE, static_cast<uint8_t>(math_approx_mode), 0);
  }
  explicit TensixConfigBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<TensixConfig> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<TensixConfig>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<TensixConfig> CreateTensixConfig(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::MathFidelity math_fidelity = tt::target::MathFidelity::LoFi,
    bool fp32_dest_acc_en = false,
    bool preserve_fp32_precision = false,
    bool math_approx_mode = false) {
  TensixConfigBuilder builder_(_fbb);
  builder_.add_math_approx_mode(math_approx_mode);
  builder_.add_preserve_fp32_precision(preserve_fp32_precision);
  builder_.add_fp32_dest_acc_en(fp32_dest_acc_en);
  builder_.add_math_fidelity(math_fidelity);
  return builder_.Finish();
}
 
struct TensixConfig::Traits {
  using type = TensixConfig;
  static auto constexpr Create = CreateTensixConfig;
};
 
struct EthernetConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef EthernetConfigBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_ETH_TYPE = 4,
    VT_NOC_INDEX = 6
  };
  tt::target::metal::EthType eth_type() const {
    return static_cast<tt::target::metal::EthType>(GetField<uint16_t>(VT_ETH_TYPE, 0));
  }
  tt::target::metal::NocIndex noc_index() const {
    return static_cast<tt::target::metal::NocIndex>(GetField<uint16_t>(VT_NOC_INDEX, 0));
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint16_t>(verifier, VT_ETH_TYPE, 2) &&
           VerifyField<uint16_t>(verifier, VT_NOC_INDEX, 2) &&
           verifier.EndTable();
  }
};
 
struct EthernetConfigBuilder {
  typedef EthernetConfig Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_eth_type(tt::target::metal::EthType eth_type) {
    fbb_.AddElement<uint16_t>(EthernetConfig::VT_ETH_TYPE, static_cast<uint16_t>(eth_type), 0);
  }
  void add_noc_index(tt::target::metal::NocIndex noc_index) {
    fbb_.AddElement<uint16_t>(EthernetConfig::VT_NOC_INDEX, static_cast<uint16_t>(noc_index), 0);
  }
  explicit EthernetConfigBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<EthernetConfig> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<EthernetConfig>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<EthernetConfig> CreateEthernetConfig(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::EthType eth_type = tt::target::metal::EthType::Sender,
    tt::target::metal::NocIndex noc_index = tt::target::metal::NocIndex::Noc0) {
  EthernetConfigBuilder builder_(_fbb);
  builder_.add_noc_index(noc_index);
  builder_.add_eth_type(eth_type);
  return builder_.Finish();
}
 
struct EthernetConfig::Traits {
  using type = EthernetConfig;
  static auto constexpr Create = CreateEthernetConfig;
};
 
struct KernelSource FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelSourceBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_SOURCE = 4,
    VT_CONFIG_TYPE = 6,
    VT_CONFIG = 8
  };
  const ::flatbuffers::String *source() const {
    return GetPointer<const ::flatbuffers::String *>(VT_SOURCE);
  }
  tt::target::metal::KernelConfig config_type() const {
    return static_cast<tt::target::metal::KernelConfig>(GetField<uint8_t>(VT_CONFIG_TYPE, 0));
  }
  const void *config() const {
    return GetPointer<const void *>(VT_CONFIG);
  }
  template<typename T> const T *config_as() const;
  const tt::target::metal::NocConfig *config_as_NocConfig() const {
    return config_type() == tt::target::metal::KernelConfig::NocConfig ? static_cast<const tt::target::metal::NocConfig *>(config()) : nullptr;
  }
  const tt::target::metal::TensixConfig *config_as_TensixConfig() const {
    return config_type() == tt::target::metal::KernelConfig::TensixConfig ? static_cast<const tt::target::metal::TensixConfig *>(config()) : nullptr;
  }
  const tt::target::metal::EthernetConfig *config_as_EthernetConfig() const {
    return config_type() == tt::target::metal::KernelConfig::EthernetConfig ? static_cast<const tt::target::metal::EthernetConfig *>(config()) : nullptr;
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyOffset(verifier, VT_SOURCE) &&
           verifier.VerifyString(source()) &&
           VerifyField<uint8_t>(verifier, VT_CONFIG_TYPE, 1) &&
           VerifyOffset(verifier, VT_CONFIG) &&
           VerifyKernelConfig(verifier, config(), config_type()) &&
           verifier.EndTable();
  }
};
 
template<> inline const tt::target::metal::NocConfig *KernelSource::config_as<tt::target::metal::NocConfig>() const {
  return config_as_NocConfig();
}
 
template<> inline const tt::target::metal::TensixConfig *KernelSource::config_as<tt::target::metal::TensixConfig>() const {
  return config_as_TensixConfig();
}
 
template<> inline const tt::target::metal::EthernetConfig *KernelSource::config_as<tt::target::metal::EthernetConfig>() const {
  return config_as_EthernetConfig();
}
 
struct KernelSourceBuilder {
  typedef KernelSource Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_source(::flatbuffers::Offset<::flatbuffers::String> source) {
    fbb_.AddOffset(KernelSource::VT_SOURCE, source);
  }
  void add_config_type(tt::target::metal::KernelConfig config_type) {
    fbb_.AddElement<uint8_t>(KernelSource::VT_CONFIG_TYPE, static_cast<uint8_t>(config_type), 0);
  }
  void add_config(::flatbuffers::Offset<void> config) {
    fbb_.AddOffset(KernelSource::VT_CONFIG, config);
  }
  explicit KernelSourceBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelSource> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelSource>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelSource> CreateKernelSource(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    ::flatbuffers::Offset<::flatbuffers::String> source = 0,
    tt::target::metal::KernelConfig config_type = tt::target::metal::KernelConfig::NONE,
    ::flatbuffers::Offset<void> config = 0) {
  KernelSourceBuilder builder_(_fbb);
  builder_.add_config(config);
  builder_.add_source(source);
  builder_.add_config_type(config_type);
  return builder_.Finish();
}
 
struct KernelSource::Traits {
  using type = KernelSource;
  static auto constexpr Create = CreateKernelSource;
};
 
inline ::flatbuffers::Offset<KernelSource> CreateKernelSourceDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    const char *source = nullptr,
    tt::target::metal::KernelConfig config_type = tt::target::metal::KernelConfig::NONE,
    ::flatbuffers::Offset<void> config = 0) {
  auto source__ = source ? _fbb.CreateString(source) : 0;
  return tt::target::metal::CreateKernelSource(
      _fbb,
      source__,
      config_type,
      config);
}
 
struct KernelBinary FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelBinaryBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_CORE_TYPE = 4,
    VT_DATA = 6,
    VT_DEBUG_SOURCE = 8
  };
  tt::target::metal::BinaryType core_type() const {
    return static_cast<tt::target::metal::BinaryType>(GetField<uint16_t>(VT_CORE_TYPE, 0));
  }
  const ::flatbuffers::Vector<uint8_t> *data() const {
    return GetPointer<const ::flatbuffers::Vector<uint8_t> *>(VT_DATA);
  }
  const ::flatbuffers::String *debug_source() const {
    return GetPointer<const ::flatbuffers::String *>(VT_DEBUG_SOURCE);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint16_t>(verifier, VT_CORE_TYPE, 2) &&
           VerifyOffset(verifier, VT_DATA) &&
           verifier.VerifyVector(data()) &&
           VerifyOffset(verifier, VT_DEBUG_SOURCE) &&
           verifier.VerifyString(debug_source()) &&
           verifier.EndTable();
  }
};
 
struct KernelBinaryBuilder {
  typedef KernelBinary Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_core_type(tt::target::metal::BinaryType core_type) {
    fbb_.AddElement<uint16_t>(KernelBinary::VT_CORE_TYPE, static_cast<uint16_t>(core_type), 0);
  }
  void add_data(::flatbuffers::Offset<::flatbuffers::Vector<uint8_t>> data) {
    fbb_.AddOffset(KernelBinary::VT_DATA, data);
  }
  void add_debug_source(::flatbuffers::Offset<::flatbuffers::String> debug_source) {
    fbb_.AddOffset(KernelBinary::VT_DEBUG_SOURCE, debug_source);
  }
  explicit KernelBinaryBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelBinary> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelBinary>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelBinary> CreateKernelBinary(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::BinaryType core_type = tt::target::metal::BinaryType::BRISC,
    ::flatbuffers::Offset<::flatbuffers::Vector<uint8_t>> data = 0,
    ::flatbuffers::Offset<::flatbuffers::String> debug_source = 0) {
  KernelBinaryBuilder builder_(_fbb);
  builder_.add_debug_source(debug_source);
  builder_.add_data(data);
  builder_.add_core_type(core_type);
  return builder_.Finish();
}
 
struct KernelBinary::Traits {
  using type = KernelBinary;
  static auto constexpr Create = CreateKernelBinary;
};
 
inline ::flatbuffers::Offset<KernelBinary> CreateKernelBinaryDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::BinaryType core_type = tt::target::metal::BinaryType::BRISC,
    const std::vector<uint8_t> *data = nullptr,
    const char *debug_source = nullptr) {
  auto data__ = data ? _fbb.CreateVector<uint8_t>(*data) : 0;
  auto debug_source__ = debug_source ? _fbb.CreateString(debug_source) : 0;
  return tt::target::metal::CreateKernelBinary(
      _fbb,
      core_type,
      data__,
      debug_source__);
}
 
struct RuntimeArgTensorAddress FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef RuntimeArgTensorAddressBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_OPERAND_IDX = 4
  };
  uint32_t operand_idx() const {
    return GetField<uint32_t>(VT_OPERAND_IDX, 0);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint32_t>(verifier, VT_OPERAND_IDX, 4) &&
           verifier.EndTable();
  }
};
 
struct RuntimeArgTensorAddressBuilder {
  typedef RuntimeArgTensorAddress Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_operand_idx(uint32_t operand_idx) {
    fbb_.AddElement<uint32_t>(RuntimeArgTensorAddress::VT_OPERAND_IDX, operand_idx, 0);
  }
  explicit RuntimeArgTensorAddressBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<RuntimeArgTensorAddress> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<RuntimeArgTensorAddress>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<RuntimeArgTensorAddress> CreateRuntimeArgTensorAddress(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    uint32_t operand_idx = 0) {
  RuntimeArgTensorAddressBuilder builder_(_fbb);
  builder_.add_operand_idx(operand_idx);
  return builder_.Finish();
}
 
struct RuntimeArgTensorAddress::Traits {
  using type = RuntimeArgTensorAddress;
  static auto constexpr Create = CreateRuntimeArgTensorAddress;
};
 
struct RuntimeArgSemaphoreAddress FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef RuntimeArgSemaphoreAddressBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_INITIAL_VALUE = 4,
    VT_CORE_TYPE = 6
  };
  uint32_t initial_value() const {
    return GetField<uint32_t>(VT_INITIAL_VALUE, 0);
  }
  tt::target::metal::CoreType core_type() const {
    return static_cast<tt::target::metal::CoreType>(GetField<uint16_t>(VT_CORE_TYPE, 0));
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint32_t>(verifier, VT_INITIAL_VALUE, 4) &&
           VerifyField<uint16_t>(verifier, VT_CORE_TYPE, 2) &&
           verifier.EndTable();
  }
};
 
struct RuntimeArgSemaphoreAddressBuilder {
  typedef RuntimeArgSemaphoreAddress Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_initial_value(uint32_t initial_value) {
    fbb_.AddElement<uint32_t>(RuntimeArgSemaphoreAddress::VT_INITIAL_VALUE, initial_value, 0);
  }
  void add_core_type(tt::target::metal::CoreType core_type) {
    fbb_.AddElement<uint16_t>(RuntimeArgSemaphoreAddress::VT_CORE_TYPE, static_cast<uint16_t>(core_type), 0);
  }
  explicit RuntimeArgSemaphoreAddressBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<RuntimeArgSemaphoreAddress> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<RuntimeArgSemaphoreAddress>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<RuntimeArgSemaphoreAddress> CreateRuntimeArgSemaphoreAddress(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    uint32_t initial_value = 0,
    tt::target::metal::CoreType core_type = tt::target::metal::CoreType::WORKER) {
  RuntimeArgSemaphoreAddressBuilder builder_(_fbb);
  builder_.add_initial_value(initial_value);
  builder_.add_core_type(core_type);
  return builder_.Finish();
}
 
struct RuntimeArgSemaphoreAddress::Traits {
  using type = RuntimeArgSemaphoreAddress;
  static auto constexpr Create = CreateRuntimeArgSemaphoreAddress;
};
 
struct KernelDesc FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelDescBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_KERNEL_TYPE = 4,
    VT_KERNEL = 6,
    VT_CORE_RANGE_SET = 8,
    VT_CBS = 10,
    VT_RUNTIME_ARGS_TYPE = 12,
    VT_RUNTIME_ARGS = 14,
    VT_DEBUG_INFO = 16
  };
  tt::target::metal::Kernel kernel_type() const {
    return static_cast<tt::target::metal::Kernel>(GetField<uint8_t>(VT_KERNEL_TYPE, 0));
  }
  const void *kernel() const {
    return GetPointer<const void *>(VT_KERNEL);
  }
  template<typename T> const T *kernel_as() const;
  const tt::target::metal::KernelSource *kernel_as_KernelSource() const {
    return kernel_type() == tt::target::metal::Kernel::KernelSource ? static_cast<const tt::target::metal::KernelSource *>(kernel()) : nullptr;
  }
  const tt::target::metal::KernelBinary *kernel_as_KernelBinary() const {
    return kernel_type() == tt::target::metal::Kernel::KernelBinary ? static_cast<const tt::target::metal::KernelBinary *>(kernel()) : nullptr;
  }
  const ::flatbuffers::Vector<const tt::target::Dim2dRange *> *core_range_set() const {
    return GetPointer<const ::flatbuffers::Vector<const tt::target::Dim2dRange *> *>(VT_CORE_RANGE_SET);
  }
  const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::CBRef>> *cbs() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::CBRef>> *>(VT_CBS);
  }
  const ::flatbuffers::Vector<tt::target::metal::RuntimeArg> *runtime_args_type() const {
    return GetPointer<const ::flatbuffers::Vector<tt::target::metal::RuntimeArg> *>(VT_RUNTIME_ARGS_TYPE);
  }
  const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *runtime_args() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *>(VT_RUNTIME_ARGS);
  }
  const ::flatbuffers::String *debug_info() const {
    return GetPointer<const ::flatbuffers::String *>(VT_DEBUG_INFO);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint8_t>(verifier, VT_KERNEL_TYPE, 1) &&
           VerifyOffset(verifier, VT_KERNEL) &&
           VerifyKernel(verifier, kernel(), kernel_type()) &&
           VerifyOffset(verifier, VT_CORE_RANGE_SET) &&
           verifier.VerifyVector(core_range_set()) &&
           VerifyOffset(verifier, VT_CBS) &&
           verifier.VerifyVector(cbs()) &&
           verifier.VerifyVectorOfTables(cbs()) &&
           VerifyOffset(verifier, VT_RUNTIME_ARGS_TYPE) &&
           verifier.VerifyVector(runtime_args_type()) &&
           VerifyOffset(verifier, VT_RUNTIME_ARGS) &&
           verifier.VerifyVector(runtime_args()) &&
           VerifyRuntimeArgVector(verifier, runtime_args(), runtime_args_type()) &&
           VerifyOffset(verifier, VT_DEBUG_INFO) &&
           verifier.VerifyString(debug_info()) &&
           verifier.EndTable();
  }
};
 
template<> inline const tt::target::metal::KernelSource *KernelDesc::kernel_as<tt::target::metal::KernelSource>() const {
  return kernel_as_KernelSource();
}
 
template<> inline const tt::target::metal::KernelBinary *KernelDesc::kernel_as<tt::target::metal::KernelBinary>() const {
  return kernel_as_KernelBinary();
}
 
struct KernelDescBuilder {
  typedef KernelDesc Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_kernel_type(tt::target::metal::Kernel kernel_type) {
    fbb_.AddElement<uint8_t>(KernelDesc::VT_KERNEL_TYPE, static_cast<uint8_t>(kernel_type), 0);
  }
  void add_kernel(::flatbuffers::Offset<void> kernel) {
    fbb_.AddOffset(KernelDesc::VT_KERNEL, kernel);
  }
  void add_core_range_set(::flatbuffers::Offset<::flatbuffers::Vector<const tt::target::Dim2dRange *>> core_range_set) {
    fbb_.AddOffset(KernelDesc::VT_CORE_RANGE_SET, core_range_set);
  }
  void add_cbs(::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::CBRef>>> cbs) {
    fbb_.AddOffset(KernelDesc::VT_CBS, cbs);
  }
  void add_runtime_args_type(::flatbuffers::Offset<::flatbuffers::Vector<tt::target::metal::RuntimeArg>> runtime_args_type) {
    fbb_.AddOffset(KernelDesc::VT_RUNTIME_ARGS_TYPE, runtime_args_type);
  }
  void add_runtime_args(::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<void>>> runtime_args) {
    fbb_.AddOffset(KernelDesc::VT_RUNTIME_ARGS, runtime_args);
  }
  void add_debug_info(::flatbuffers::Offset<::flatbuffers::String> debug_info) {
    fbb_.AddOffset(KernelDesc::VT_DEBUG_INFO, debug_info);
  }
  explicit KernelDescBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelDesc> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelDesc>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelDesc> CreateKernelDesc(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::Kernel kernel_type = tt::target::metal::Kernel::NONE,
    ::flatbuffers::Offset<void> kernel = 0,
    ::flatbuffers::Offset<::flatbuffers::Vector<const tt::target::Dim2dRange *>> core_range_set = 0,
    ::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::CBRef>>> cbs = 0,
    ::flatbuffers::Offset<::flatbuffers::Vector<tt::target::metal::RuntimeArg>> runtime_args_type = 0,
    ::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<void>>> runtime_args = 0,
    ::flatbuffers::Offset<::flatbuffers::String> debug_info = 0) {
  KernelDescBuilder builder_(_fbb);
  builder_.add_debug_info(debug_info);
  builder_.add_runtime_args(runtime_args);
  builder_.add_runtime_args_type(runtime_args_type);
  builder_.add_cbs(cbs);
  builder_.add_core_range_set(core_range_set);
  builder_.add_kernel(kernel);
  builder_.add_kernel_type(kernel_type);
  return builder_.Finish();
}
 
struct KernelDesc::Traits {
  using type = KernelDesc;
  static auto constexpr Create = CreateKernelDesc;
};
 
inline ::flatbuffers::Offset<KernelDesc> CreateKernelDescDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::Kernel kernel_type = tt::target::metal::Kernel::NONE,
    ::flatbuffers::Offset<void> kernel = 0,
    const std::vector<tt::target::Dim2dRange> *core_range_set = nullptr,
    const std::vector<::flatbuffers::Offset<tt::target::CBRef>> *cbs = nullptr,
    const std::vector<tt::target::metal::RuntimeArg> *runtime_args_type = nullptr,
    const std::vector<::flatbuffers::Offset<void>> *runtime_args = nullptr,
    const char *debug_info = nullptr) {
  auto core_range_set__ = core_range_set ? _fbb.CreateVectorOfStructs<tt::target::Dim2dRange>(*core_range_set) : 0;
  auto cbs__ = cbs ? _fbb.CreateVector<::flatbuffers::Offset<tt::target::CBRef>>(*cbs) : 0;
  auto runtime_args_type__ = runtime_args_type ? _fbb.CreateVector<tt::target::metal::RuntimeArg>(*runtime_args_type) : 0;
  auto runtime_args__ = runtime_args ? _fbb.CreateVector<::flatbuffers::Offset<void>>(*runtime_args) : 0;
  auto debug_info__ = debug_info ? _fbb.CreateString(debug_info) : 0;
  return tt::target::metal::CreateKernelDesc(
      _fbb,
      kernel_type,
      kernel,
      core_range_set__,
      cbs__,
      runtime_args_type__,
      runtime_args__,
      debug_info__);
}
 
struct ProgramDesc FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef ProgramDescBuilder Builder;
  struct Traits;
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_KERNELS = 4
  };
  const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelDesc>> *kernels() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelDesc>> *>(VT_KERNELS);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyOffset(verifier, VT_KERNELS) &&
           verifier.VerifyVector(kernels()) &&
           verifier.VerifyVectorOfTables(kernels()) &&
           verifier.EndTable();
  }
};
 
struct ProgramDescBuilder {
  typedef ProgramDesc Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_kernels(::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelDesc>>> kernels) {
    fbb_.AddOffset(ProgramDesc::VT_KERNELS, kernels);
  }
  explicit ProgramDescBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<ProgramDesc> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<ProgramDesc>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<ProgramDesc> CreateProgramDesc(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    ::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelDesc>>> kernels = 0) {
  ProgramDescBuilder builder_(_fbb);
  builder_.add_kernels(kernels);
  return builder_.Finish();
}
 
struct ProgramDesc::Traits {
  using type = ProgramDesc;
  static auto constexpr Create = CreateProgramDesc;
};
 
inline ::flatbuffers::Offset<ProgramDesc> CreateProgramDescDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    const std::vector<::flatbuffers::Offset<tt::target::metal::KernelDesc>> *kernels = nullptr) {
  auto kernels__ = kernels ? _fbb.CreateVector<::flatbuffers::Offset<tt::target::metal::KernelDesc>>(*kernels) : 0;
  return tt::target::metal::CreateProgramDesc(
      _fbb,
      kernels__);
}
 
inline bool VerifyKernelConfig(::flatbuffers::Verifier &verifier, const void *obj, KernelConfig type) {
  switch (type) {
    case KernelConfig::NONE: {
      return true;
    }
    case KernelConfig::NocConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::NocConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelConfig::TensixConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::TensixConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelConfig::EthernetConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::EthernetConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    default: return true;
  }
}
 
inline bool VerifyKernelConfigVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelConfig> *types) {
  if (!values || !types) return !values && !types;
  if (values->size() != types->size()) return false;
  for (::flatbuffers::uoffset_t i = 0; i < values->size(); ++i) {
    if (!VerifyKernelConfig(
        verifier,  values->Get(i), types->GetEnum<KernelConfig>(i))) {
      return false;
    }
  }
  return true;
}
 
inline bool VerifyKernel(::flatbuffers::Verifier &verifier, const void *obj, Kernel type) {
  switch (type) {
    case Kernel::NONE: {
      return true;
    }
    case Kernel::KernelSource: {
      auto ptr = reinterpret_cast<const tt::target::metal::KernelSource *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case Kernel::KernelBinary: {
      auto ptr = reinterpret_cast<const tt::target::metal::KernelBinary *>(obj);
      return verifier.VerifyTable(ptr);
    }
    default: return true;
  }
}
 
inline bool VerifyKernelVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<Kernel> *types) {
  if (!values || !types) return !values && !types;
  if (values->size() != types->size()) return false;
  for (::flatbuffers::uoffset_t i = 0; i < values->size(); ++i) {
    if (!VerifyKernel(
        verifier,  values->Get(i), types->GetEnum<Kernel>(i))) {
      return false;
    }
  }
  return true;
}
 
inline bool VerifyRuntimeArg(::flatbuffers::Verifier &verifier, const void *obj, RuntimeArg type) {
  switch (type) {
    case RuntimeArg::NONE: {
      return true;
    }
    case RuntimeArg::RuntimeArgTensorAddress: {
      auto ptr = reinterpret_cast<const tt::target::metal::RuntimeArgTensorAddress *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case RuntimeArg::RuntimeArgSemaphoreAddress: {
      auto ptr = reinterpret_cast<const tt::target::metal::RuntimeArgSemaphoreAddress *>(obj);
      return verifier.VerifyTable(ptr);
    }
    default: return true;
  }
}
 
inline bool VerifyRuntimeArgVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<RuntimeArg> *types) {
  if (!values || !types) return !values && !types;
  if (values->size() != types->size()) return false;
  for (::flatbuffers::uoffset_t i = 0; i < values->size(); ++i) {
    if (!VerifyRuntimeArg(
        verifier,  values->Get(i), types->GetEnum<RuntimeArg>(i))) {
      return false;
    }
  }
  return true;
}
 
}  // namespace metal
}  // namespace target
}  // namespace tt
 
#endif  // FLATBUFFERS_GENERATED_PROGRAM_TT_TARGET_METAL_H_