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 "ttmlir/Target/Common/types_generated.h"
#include "ttmlir/Target/TTMetal/types_generated.h"
 
namespace tt {
namespace target {
namespace metal {
 
struct NocConfig;
struct NocConfigBuilder;
 
struct ComputeConfig;
struct ComputeConfigBuilder;
 
struct EthernetConfig;
struct EthernetConfigBuilder;
 
struct KernelSource;
struct KernelSourceBuilder;
 
struct KernelBinary;
struct KernelBinaryBuilder;
 
struct KernelArgCBPort;
struct KernelArgCBPortBuilder;
 
struct KernelArgBufferAddress;
struct KernelArgBufferAddressBuilder;
 
struct KernelArgSemaphore;
struct KernelArgSemaphoreBuilder;
 
struct KernelArg;
struct KernelArgBuilder;
 
struct KernelArgs;
struct KernelArgsBuilder;
 
struct KernelConfig;
struct KernelConfigBuilder;
 
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 UnpackToDestMode : uint8_t {
  Fp32 = 0,
  Default = 1,
  MIN = Fp32,
  MAX = Default
};
 
inline const UnpackToDestMode (&EnumValuesUnpackToDestMode())[2] {
  static const UnpackToDestMode values[] = {
    UnpackToDestMode::Fp32,
    UnpackToDestMode::Default
  };
  return values;
}
 
inline const char * const *EnumNamesUnpackToDestMode() {
  static const char * const names[3] = {
    "Fp32",
    "Default",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameUnpackToDestMode(UnpackToDestMode e) {
  if (::flatbuffers::IsOutRange(e, UnpackToDestMode::Fp32, UnpackToDestMode::Default)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesUnpackToDestMode()[index];
}
 
enum class KernelConfigType : uint8_t {
  NONE = 0,
  NocConfig = 1,
  ComputeConfig = 2,
  EthernetConfig = 3,
  MIN = NONE,
  MAX = EthernetConfig
};
 
inline const KernelConfigType (&EnumValuesKernelConfigType())[4] {
  static const KernelConfigType values[] = {
    KernelConfigType::NONE,
    KernelConfigType::NocConfig,
    KernelConfigType::ComputeConfig,
    KernelConfigType::EthernetConfig
  };
  return values;
}
 
inline const char * const *EnumNamesKernelConfigType() {
  static const char * const names[5] = {
    "NONE",
    "NocConfig",
    "ComputeConfig",
    "EthernetConfig",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameKernelConfigType(KernelConfigType e) {
  if (::flatbuffers::IsOutRange(e, KernelConfigType::NONE, KernelConfigType::EthernetConfig)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesKernelConfigType()[index];
}
 
template<typename T> struct KernelConfigTypeTraits {
  static const KernelConfigType enum_value = KernelConfigType::NONE;
};
 
template<> struct KernelConfigTypeTraits<tt::target::metal::NocConfig> {
  static const KernelConfigType enum_value = KernelConfigType::NocConfig;
};
 
template<> struct KernelConfigTypeTraits<tt::target::metal::ComputeConfig> {
  static const KernelConfigType enum_value = KernelConfigType::ComputeConfig;
};
 
template<> struct KernelConfigTypeTraits<tt::target::metal::EthernetConfig> {
  static const KernelConfigType enum_value = KernelConfigType::EthernetConfig;
};
 
bool VerifyKernelConfigType(::flatbuffers::Verifier &verifier, const void *obj, KernelConfigType type);
bool VerifyKernelConfigTypeVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelConfigType> *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 KernelArgType : uint8_t {
  NONE = 0,
  KernelArgCBPort = 1,
  KernelArgBufferAddress = 2,
  KernelArgSemaphore = 3,
  MIN = NONE,
  MAX = KernelArgSemaphore
};
 
inline const KernelArgType (&EnumValuesKernelArgType())[4] {
  static const KernelArgType values[] = {
    KernelArgType::NONE,
    KernelArgType::KernelArgCBPort,
    KernelArgType::KernelArgBufferAddress,
    KernelArgType::KernelArgSemaphore
  };
  return values;
}
 
inline const char * const *EnumNamesKernelArgType() {
  static const char * const names[5] = {
    "NONE",
    "KernelArgCBPort",
    "KernelArgBufferAddress",
    "KernelArgSemaphore",
    nullptr
  };
  return names;
}
 
inline const char *EnumNameKernelArgType(KernelArgType e) {
  if (::flatbuffers::IsOutRange(e, KernelArgType::NONE, KernelArgType::KernelArgSemaphore)) return "";
  const size_t index = static_cast<size_t>(e);
  return EnumNamesKernelArgType()[index];
}
 
template<typename T> struct KernelArgTypeTraits {
  static const KernelArgType enum_value = KernelArgType::NONE;
};
 
template<> struct KernelArgTypeTraits<tt::target::metal::KernelArgCBPort> {
  static const KernelArgType enum_value = KernelArgType::KernelArgCBPort;
};
 
template<> struct KernelArgTypeTraits<tt::target::metal::KernelArgBufferAddress> {
  static const KernelArgType enum_value = KernelArgType::KernelArgBufferAddress;
};
 
template<> struct KernelArgTypeTraits<tt::target::metal::KernelArgSemaphore> {
  static const KernelArgType enum_value = KernelArgType::KernelArgSemaphore;
};
 
bool VerifyKernelArgType(::flatbuffers::Verifier &verifier, const void *obj, KernelArgType type);
bool VerifyKernelArgTypeVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelArgType> *types);
 
struct NocConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef NocConfigBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.NocConfig";
  }
  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 ComputeConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef ComputeConfigBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.ComputeConfig";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_MATH_FIDELITY = 4,
    VT_FP32_DEST_ACC_EN = 6,
    VT_MATH_APPROX_MODE = 8,
    VT_UNPACK_TO_DEST_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 math_approx_mode() const {
    return GetField<uint8_t>(VT_MATH_APPROX_MODE, 0) != 0;
  }
  const ::flatbuffers::Vector<tt::target::metal::UnpackToDestMode> *unpack_to_dest_mode() const {
    return GetPointer<const ::flatbuffers::Vector<tt::target::metal::UnpackToDestMode> *>(VT_UNPACK_TO_DEST_MODE);
  }
  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_MATH_APPROX_MODE, 1) &&
           VerifyOffset(verifier, VT_UNPACK_TO_DEST_MODE) &&
           verifier.VerifyVector(unpack_to_dest_mode()) &&
           verifier.EndTable();
  }
};
 
struct ComputeConfigBuilder {
  typedef ComputeConfig Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_math_fidelity(tt::target::MathFidelity math_fidelity) {
    fbb_.AddElement<uint8_t>(ComputeConfig::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>(ComputeConfig::VT_FP32_DEST_ACC_EN, static_cast<uint8_t>(fp32_dest_acc_en), 0);
  }
  void add_math_approx_mode(bool math_approx_mode) {
    fbb_.AddElement<uint8_t>(ComputeConfig::VT_MATH_APPROX_MODE, static_cast<uint8_t>(math_approx_mode), 0);
  }
  void add_unpack_to_dest_mode(::flatbuffers::Offset<::flatbuffers::Vector<tt::target::metal::UnpackToDestMode>> unpack_to_dest_mode) {
    fbb_.AddOffset(ComputeConfig::VT_UNPACK_TO_DEST_MODE, unpack_to_dest_mode);
  }
  explicit ComputeConfigBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<ComputeConfig> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<ComputeConfig>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<ComputeConfig> CreateComputeConfig(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::MathFidelity math_fidelity = tt::target::MathFidelity::LoFi,
    bool fp32_dest_acc_en = false,
    bool math_approx_mode = false,
    ::flatbuffers::Offset<::flatbuffers::Vector<tt::target::metal::UnpackToDestMode>> unpack_to_dest_mode = 0) {
  ComputeConfigBuilder builder_(_fbb);
  builder_.add_unpack_to_dest_mode(unpack_to_dest_mode);
  builder_.add_math_approx_mode(math_approx_mode);
  builder_.add_fp32_dest_acc_en(fp32_dest_acc_en);
  builder_.add_math_fidelity(math_fidelity);
  return builder_.Finish();
}
 
struct ComputeConfig::Traits {
  using type = ComputeConfig;
  static auto constexpr Create = CreateComputeConfig;
};
 
inline ::flatbuffers::Offset<ComputeConfig> CreateComputeConfigDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::MathFidelity math_fidelity = tt::target::MathFidelity::LoFi,
    bool fp32_dest_acc_en = false,
    bool math_approx_mode = false,
    const std::vector<tt::target::metal::UnpackToDestMode> *unpack_to_dest_mode = nullptr) {
  auto unpack_to_dest_mode__ = unpack_to_dest_mode ? _fbb.CreateVector<tt::target::metal::UnpackToDestMode>(*unpack_to_dest_mode) : 0;
  return tt::target::metal::CreateComputeConfig(
      _fbb,
      math_fidelity,
      fp32_dest_acc_en,
      math_approx_mode,
      unpack_to_dest_mode__);
}
 
struct EthernetConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef EthernetConfigBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.EthernetConfig";
  }
  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;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelSource";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_SOURCE = 4
  };
  const ::flatbuffers::String *source() const {
    return GetPointer<const ::flatbuffers::String *>(VT_SOURCE);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyOffset(verifier, VT_SOURCE) &&
           verifier.VerifyString(source()) &&
           verifier.EndTable();
  }
};
 
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);
  }
  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) {
  KernelSourceBuilder builder_(_fbb);
  builder_.add_source(source);
  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) {
  auto source__ = source ? _fbb.CreateString(source) : 0;
  return tt::target::metal::CreateKernelSource(
      _fbb,
      source__);
}
 
struct KernelBinary FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelBinaryBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelBinary";
  }
  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 KernelArgCBPort FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelArgCBPortBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelArgCBPort";
  }
  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 KernelArgCBPortBuilder {
  typedef KernelArgCBPort Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_operand_idx(uint32_t operand_idx) {
    fbb_.AddElement<uint32_t>(KernelArgCBPort::VT_OPERAND_IDX, operand_idx, 0);
  }
  explicit KernelArgCBPortBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelArgCBPort> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelArgCBPort>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelArgCBPort> CreateKernelArgCBPort(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    uint32_t operand_idx = 0) {
  KernelArgCBPortBuilder builder_(_fbb);
  builder_.add_operand_idx(operand_idx);
  return builder_.Finish();
}
 
struct KernelArgCBPort::Traits {
  using type = KernelArgCBPort;
  static auto constexpr Create = CreateKernelArgCBPort;
};
 
struct KernelArgBufferAddress FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelArgBufferAddressBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelArgBufferAddress";
  }
  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 KernelArgBufferAddressBuilder {
  typedef KernelArgBufferAddress Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_operand_idx(uint32_t operand_idx) {
    fbb_.AddElement<uint32_t>(KernelArgBufferAddress::VT_OPERAND_IDX, operand_idx, 0);
  }
  explicit KernelArgBufferAddressBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelArgBufferAddress> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelArgBufferAddress>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelArgBufferAddress> CreateKernelArgBufferAddress(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    uint32_t operand_idx = 0) {
  KernelArgBufferAddressBuilder builder_(_fbb);
  builder_.add_operand_idx(operand_idx);
  return builder_.Finish();
}
 
struct KernelArgBufferAddress::Traits {
  using type = KernelArgBufferAddress;
  static auto constexpr Create = CreateKernelArgBufferAddress;
};
 
struct KernelArgSemaphore FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelArgSemaphoreBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelArgSemaphore";
  }
  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 KernelArgSemaphoreBuilder {
  typedef KernelArgSemaphore Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_initial_value(uint32_t initial_value) {
    fbb_.AddElement<uint32_t>(KernelArgSemaphore::VT_INITIAL_VALUE, initial_value, 0);
  }
  void add_core_type(tt::target::metal::CoreType core_type) {
    fbb_.AddElement<uint16_t>(KernelArgSemaphore::VT_CORE_TYPE, static_cast<uint16_t>(core_type), 0);
  }
  explicit KernelArgSemaphoreBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelArgSemaphore> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelArgSemaphore>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelArgSemaphore> CreateKernelArgSemaphore(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    uint32_t initial_value = 0,
    tt::target::metal::CoreType core_type = tt::target::metal::CoreType::WORKER) {
  KernelArgSemaphoreBuilder builder_(_fbb);
  builder_.add_initial_value(initial_value);
  builder_.add_core_type(core_type);
  return builder_.Finish();
}
 
struct KernelArgSemaphore::Traits {
  using type = KernelArgSemaphore;
  static auto constexpr Create = CreateKernelArgSemaphore;
};
 
struct KernelArg FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelArgBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelArg";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_ARG_TYPE = 4,
    VT_ARG = 6
  };
  tt::target::metal::KernelArgType arg_type() const {
    return static_cast<tt::target::metal::KernelArgType>(GetField<uint8_t>(VT_ARG_TYPE, 0));
  }
  const void *arg() const {
    return GetPointer<const void *>(VT_ARG);
  }
  template<typename T> const T *arg_as() const;
  const tt::target::metal::KernelArgCBPort *arg_as_KernelArgCBPort() const {
    return arg_type() == tt::target::metal::KernelArgType::KernelArgCBPort ? static_cast<const tt::target::metal::KernelArgCBPort *>(arg()) : nullptr;
  }
  const tt::target::metal::KernelArgBufferAddress *arg_as_KernelArgBufferAddress() const {
    return arg_type() == tt::target::metal::KernelArgType::KernelArgBufferAddress ? static_cast<const tt::target::metal::KernelArgBufferAddress *>(arg()) : nullptr;
  }
  const tt::target::metal::KernelArgSemaphore *arg_as_KernelArgSemaphore() const {
    return arg_type() == tt::target::metal::KernelArgType::KernelArgSemaphore ? static_cast<const tt::target::metal::KernelArgSemaphore *>(arg()) : nullptr;
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyField<uint8_t>(verifier, VT_ARG_TYPE, 1) &&
           VerifyOffset(verifier, VT_ARG) &&
           VerifyKernelArgType(verifier, arg(), arg_type()) &&
           verifier.EndTable();
  }
};
 
template<> inline const tt::target::metal::KernelArgCBPort *KernelArg::arg_as<tt::target::metal::KernelArgCBPort>() const {
  return arg_as_KernelArgCBPort();
}
 
template<> inline const tt::target::metal::KernelArgBufferAddress *KernelArg::arg_as<tt::target::metal::KernelArgBufferAddress>() const {
  return arg_as_KernelArgBufferAddress();
}
 
template<> inline const tt::target::metal::KernelArgSemaphore *KernelArg::arg_as<tt::target::metal::KernelArgSemaphore>() const {
  return arg_as_KernelArgSemaphore();
}
 
struct KernelArgBuilder {
  typedef KernelArg Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_arg_type(tt::target::metal::KernelArgType arg_type) {
    fbb_.AddElement<uint8_t>(KernelArg::VT_ARG_TYPE, static_cast<uint8_t>(arg_type), 0);
  }
  void add_arg(::flatbuffers::Offset<void> arg) {
    fbb_.AddOffset(KernelArg::VT_ARG, arg);
  }
  explicit KernelArgBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelArg> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelArg>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelArg> CreateKernelArg(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    tt::target::metal::KernelArgType arg_type = tt::target::metal::KernelArgType::NONE,
    ::flatbuffers::Offset<void> arg = 0) {
  KernelArgBuilder builder_(_fbb);
  builder_.add_arg(arg);
  builder_.add_arg_type(arg_type);
  return builder_.Finish();
}
 
struct KernelArg::Traits {
  using type = KernelArg;
  static auto constexpr Create = CreateKernelArg;
};
 
struct KernelArgs FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelArgsBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelArgs";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_RT_ARGS = 4,
    VT_CT_ARGS = 6
  };
  const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *rt_args() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *>(VT_RT_ARGS);
  }
  const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *ct_args() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *>(VT_CT_ARGS);
  }
  bool Verify(::flatbuffers::Verifier &verifier) const {
    return VerifyTableStart(verifier) &&
           VerifyOffset(verifier, VT_RT_ARGS) &&
           verifier.VerifyVector(rt_args()) &&
           verifier.VerifyVectorOfTables(rt_args()) &&
           VerifyOffset(verifier, VT_CT_ARGS) &&
           verifier.VerifyVector(ct_args()) &&
           verifier.VerifyVectorOfTables(ct_args()) &&
           verifier.EndTable();
  }
};
 
struct KernelArgsBuilder {
  typedef KernelArgs Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_rt_args(::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>>> rt_args) {
    fbb_.AddOffset(KernelArgs::VT_RT_ARGS, rt_args);
  }
  void add_ct_args(::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>>> ct_args) {
    fbb_.AddOffset(KernelArgs::VT_CT_ARGS, ct_args);
  }
  explicit KernelArgsBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelArgs> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelArgs>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelArgs> CreateKernelArgs(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    ::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>>> rt_args = 0,
    ::flatbuffers::Offset<::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelArg>>> ct_args = 0) {
  KernelArgsBuilder builder_(_fbb);
  builder_.add_ct_args(ct_args);
  builder_.add_rt_args(rt_args);
  return builder_.Finish();
}
 
struct KernelArgs::Traits {
  using type = KernelArgs;
  static auto constexpr Create = CreateKernelArgs;
};
 
inline ::flatbuffers::Offset<KernelArgs> CreateKernelArgsDirect(
    ::flatbuffers::FlatBufferBuilder &_fbb,
    const std::vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *rt_args = nullptr,
    const std::vector<::flatbuffers::Offset<tt::target::metal::KernelArg>> *ct_args = nullptr) {
  auto rt_args__ = rt_args ? _fbb.CreateVector<::flatbuffers::Offset<tt::target::metal::KernelArg>>(*rt_args) : 0;
  auto ct_args__ = ct_args ? _fbb.CreateVector<::flatbuffers::Offset<tt::target::metal::KernelArg>>(*ct_args) : 0;
  return tt::target::metal::CreateKernelArgs(
      _fbb,
      rt_args__,
      ct_args__);
}
 
struct KernelConfig FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef KernelConfigBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.KernelConfig";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_KERNEL_TYPE = 4,
    VT_KERNEL = 6,
    VT_CORE_RANGE_SET = 8,
    VT_ARGS = 10,
    VT_TYPE_TYPE = 12,
    VT_TYPE = 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 tt::target::metal::KernelArgs *args() const {
    return GetPointer<const tt::target::metal::KernelArgs *>(VT_ARGS);
  }
  tt::target::metal::KernelConfigType type_type() const {
    return static_cast<tt::target::metal::KernelConfigType>(GetField<uint8_t>(VT_TYPE_TYPE, 0));
  }
  const void *type() const {
    return GetPointer<const void *>(VT_TYPE);
  }
  template<typename T> const T *type_as() const;
  const tt::target::metal::NocConfig *type_as_NocConfig() const {
    return type_type() == tt::target::metal::KernelConfigType::NocConfig ? static_cast<const tt::target::metal::NocConfig *>(type()) : nullptr;
  }
  const tt::target::metal::ComputeConfig *type_as_ComputeConfig() const {
    return type_type() == tt::target::metal::KernelConfigType::ComputeConfig ? static_cast<const tt::target::metal::ComputeConfig *>(type()) : nullptr;
  }
  const tt::target::metal::EthernetConfig *type_as_EthernetConfig() const {
    return type_type() == tt::target::metal::KernelConfigType::EthernetConfig ? static_cast<const tt::target::metal::EthernetConfig *>(type()) : nullptr;
  }
  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_ARGS) &&
           verifier.VerifyTable(args()) &&
           VerifyField<uint8_t>(verifier, VT_TYPE_TYPE, 1) &&
           VerifyOffset(verifier, VT_TYPE) &&
           VerifyKernelConfigType(verifier, type(), type_type()) &&
           VerifyOffset(verifier, VT_DEBUG_INFO) &&
           verifier.VerifyString(debug_info()) &&
           verifier.EndTable();
  }
};
 
template<> inline const tt::target::metal::KernelSource *KernelConfig::kernel_as<tt::target::metal::KernelSource>() const {
  return kernel_as_KernelSource();
}
 
template<> inline const tt::target::metal::KernelBinary *KernelConfig::kernel_as<tt::target::metal::KernelBinary>() const {
  return kernel_as_KernelBinary();
}
 
template<> inline const tt::target::metal::NocConfig *KernelConfig::type_as<tt::target::metal::NocConfig>() const {
  return type_as_NocConfig();
}
 
template<> inline const tt::target::metal::ComputeConfig *KernelConfig::type_as<tt::target::metal::ComputeConfig>() const {
  return type_as_ComputeConfig();
}
 
template<> inline const tt::target::metal::EthernetConfig *KernelConfig::type_as<tt::target::metal::EthernetConfig>() const {
  return type_as_EthernetConfig();
}
 
struct KernelConfigBuilder {
  typedef KernelConfig Table;
  ::flatbuffers::FlatBufferBuilder &fbb_;
  ::flatbuffers::uoffset_t start_;
  void add_kernel_type(tt::target::metal::Kernel kernel_type) {
    fbb_.AddElement<uint8_t>(KernelConfig::VT_KERNEL_TYPE, static_cast<uint8_t>(kernel_type), 0);
  }
  void add_kernel(::flatbuffers::Offset<void> kernel) {
    fbb_.AddOffset(KernelConfig::VT_KERNEL, kernel);
  }
  void add_core_range_set(::flatbuffers::Offset<::flatbuffers::Vector<const tt::target::Dim2dRange *>> core_range_set) {
    fbb_.AddOffset(KernelConfig::VT_CORE_RANGE_SET, core_range_set);
  }
  void add_args(::flatbuffers::Offset<tt::target::metal::KernelArgs> args) {
    fbb_.AddOffset(KernelConfig::VT_ARGS, args);
  }
  void add_type_type(tt::target::metal::KernelConfigType type_type) {
    fbb_.AddElement<uint8_t>(KernelConfig::VT_TYPE_TYPE, static_cast<uint8_t>(type_type), 0);
  }
  void add_type(::flatbuffers::Offset<void> type) {
    fbb_.AddOffset(KernelConfig::VT_TYPE, type);
  }
  void add_debug_info(::flatbuffers::Offset<::flatbuffers::String> debug_info) {
    fbb_.AddOffset(KernelConfig::VT_DEBUG_INFO, debug_info);
  }
  explicit KernelConfigBuilder(::flatbuffers::FlatBufferBuilder &_fbb)
        : fbb_(_fbb) {
    start_ = fbb_.StartTable();
  }
  ::flatbuffers::Offset<KernelConfig> Finish() {
    const auto end = fbb_.EndTable(start_);
    auto o = ::flatbuffers::Offset<KernelConfig>(end);
    return o;
  }
};
 
inline ::flatbuffers::Offset<KernelConfig> CreateKernelConfig(
    ::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<tt::target::metal::KernelArgs> args = 0,
    tt::target::metal::KernelConfigType type_type = tt::target::metal::KernelConfigType::NONE,
    ::flatbuffers::Offset<void> type = 0,
    ::flatbuffers::Offset<::flatbuffers::String> debug_info = 0) {
  KernelConfigBuilder builder_(_fbb);
  builder_.add_debug_info(debug_info);
  builder_.add_type(type);
  builder_.add_args(args);
  builder_.add_core_range_set(core_range_set);
  builder_.add_kernel(kernel);
  builder_.add_type_type(type_type);
  builder_.add_kernel_type(kernel_type);
  return builder_.Finish();
}
 
struct KernelConfig::Traits {
  using type = KernelConfig;
  static auto constexpr Create = CreateKernelConfig;
};
 
inline ::flatbuffers::Offset<KernelConfig> CreateKernelConfigDirect(
    ::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,
    ::flatbuffers::Offset<tt::target::metal::KernelArgs> args = 0,
    tt::target::metal::KernelConfigType type_type = tt::target::metal::KernelConfigType::NONE,
    ::flatbuffers::Offset<void> type = 0,
    const char *debug_info = nullptr) {
  auto core_range_set__ = core_range_set ? _fbb.CreateVectorOfStructs<tt::target::Dim2dRange>(*core_range_set) : 0;
  auto debug_info__ = debug_info ? _fbb.CreateString(debug_info) : 0;
  return tt::target::metal::CreateKernelConfig(
      _fbb,
      kernel_type,
      kernel,
      core_range_set__,
      args,
      type_type,
      type,
      debug_info__);
}
 
struct ProgramDesc FLATBUFFERS_FINAL_CLASS : private ::flatbuffers::Table {
  typedef ProgramDescBuilder Builder;
  struct Traits;
  static FLATBUFFERS_CONSTEXPR_CPP11 const char *GetFullyQualifiedName() {
    return "tt.target.metal.ProgramDesc";
  }
  enum FlatBuffersVTableOffset FLATBUFFERS_VTABLE_UNDERLYING_TYPE {
    VT_KERNELS = 4
  };
  const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelConfig>> *kernels() const {
    return GetPointer<const ::flatbuffers::Vector<::flatbuffers::Offset<tt::target::metal::KernelConfig>> *>(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::KernelConfig>>> 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::KernelConfig>>> 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::KernelConfig>> *kernels = nullptr) {
  auto kernels__ = kernels ? _fbb.CreateVector<::flatbuffers::Offset<tt::target::metal::KernelConfig>>(*kernels) : 0;
  return tt::target::metal::CreateProgramDesc(
      _fbb,
      kernels__);
}
 
inline bool VerifyKernelConfigType(::flatbuffers::Verifier &verifier, const void *obj, KernelConfigType type) {
  switch (type) {
    case KernelConfigType::NONE: {
      return true;
    }
    case KernelConfigType::NocConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::NocConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelConfigType::ComputeConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::ComputeConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelConfigType::EthernetConfig: {
      auto ptr = reinterpret_cast<const tt::target::metal::EthernetConfig *>(obj);
      return verifier.VerifyTable(ptr);
    }
    default: return true;
  }
}
 
inline bool VerifyKernelConfigTypeVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelConfigType> *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 (!VerifyKernelConfigType(
        verifier,  values->Get(i), types->GetEnum<KernelConfigType>(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 VerifyKernelArgType(::flatbuffers::Verifier &verifier, const void *obj, KernelArgType type) {
  switch (type) {
    case KernelArgType::NONE: {
      return true;
    }
    case KernelArgType::KernelArgCBPort: {
      auto ptr = reinterpret_cast<const tt::target::metal::KernelArgCBPort *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelArgType::KernelArgBufferAddress: {
      auto ptr = reinterpret_cast<const tt::target::metal::KernelArgBufferAddress *>(obj);
      return verifier.VerifyTable(ptr);
    }
    case KernelArgType::KernelArgSemaphore: {
      auto ptr = reinterpret_cast<const tt::target::metal::KernelArgSemaphore *>(obj);
      return verifier.VerifyTable(ptr);
    }
    default: return true;
  }
}
 
inline bool VerifyKernelArgTypeVector(::flatbuffers::Verifier &verifier, const ::flatbuffers::Vector<::flatbuffers::Offset<void>> *values, const ::flatbuffers::Vector<KernelArgType> *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 (!VerifyKernelArgType(
        verifier,  values->Get(i), types->GetEnum<KernelArgType>(i))) {
      return false;
    }
  }
  return true;
}
 
}  // namespace metal
}  // namespace target
}  // namespace tt
 
#endif  // FLATBUFFERS_GENERATED_PROGRAM_TT_TARGET_METAL_H_