// Generated by the protocol buffer compiler.  DO NOT EDIT!
// source: xla/backends/cpu/runtime/thunk.proto

#ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto
#define GOOGLE_PROTOBUF_INCLUDED_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto

#include <limits>
#include <string>

#include <google/protobuf/port_def.inc>
#if PROTOBUF_VERSION < 3021000
#error This file was generated by a newer version of protoc which is
#error incompatible with your Protocol Buffer headers. Please update
#error your headers.
#endif
#if 3021009 < PROTOBUF_MIN_PROTOC_VERSION
#error This file was generated by an older version of protoc which is
#error incompatible with your Protocol Buffer headers. Please
#error regenerate this file with a newer version of protoc.
#endif

#include <google/protobuf/port_undef.inc>
#include <google/protobuf/io/coded_stream.h>
#include <google/protobuf/arena.h>
#include <google/protobuf/arenastring.h>
#include <google/protobuf/generated_message_bases.h>
#include <google/protobuf/generated_message_util.h>
#include <google/protobuf/metadata_lite.h>
#include <google/protobuf/generated_message_reflection.h>
#include <google/protobuf/message.h>
#include <google/protobuf/repeated_field.h>  // IWYU pragma: export
#include <google/protobuf/extension_set.h>  // IWYU pragma: export
#include <google/protobuf/generated_enum_reflection.h>
#include <google/protobuf/unknown_field_set.h>
#include "xla/service/hlo.pb.h"
#include "xla/xla_data.pb.h"
// @@protoc_insertion_point(includes)
#include <google/protobuf/port_def.inc>
#define PROTOBUF_INTERNAL_EXPORT_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto
PROTOBUF_NAMESPACE_OPEN
namespace internal {
class AnyMetadata;
}  // namespace internal
PROTOBUF_NAMESPACE_CLOSE

// Internal implementation detail -- do not use these members.
struct TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto {
  static const uint32_t offsets[];
};
extern const ::PROTOBUF_NAMESPACE_ID::internal::DescriptorTable descriptor_table_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
namespace xla {
namespace cpu {
class AllGatherThunkProto;
struct AllGatherThunkProtoDefaultTypeInternal;
extern AllGatherThunkProtoDefaultTypeInternal _AllGatherThunkProto_default_instance_;
class AllReduceThunkProto;
struct AllReduceThunkProtoDefaultTypeInternal;
extern AllReduceThunkProtoDefaultTypeInternal _AllReduceThunkProto_default_instance_;
class AllToAllThunkProto;
struct AllToAllThunkProtoDefaultTypeInternal;
extern AllToAllThunkProtoDefaultTypeInternal _AllToAllThunkProto_default_instance_;
class BoolOptional;
struct BoolOptionalDefaultTypeInternal;
extern BoolOptionalDefaultTypeInternal _BoolOptional_default_instance_;
class BufferAllocationSliceProto;
struct BufferAllocationSliceProtoDefaultTypeInternal;
extern BufferAllocationSliceProtoDefaultTypeInternal _BufferAllocationSliceProto_default_instance_;
class CallThunkProto;
struct CallThunkProtoDefaultTypeInternal;
extern CallThunkProtoDefaultTypeInternal _CallThunkProto_default_instance_;
class CollectivePermuteThunkProto;
struct CollectivePermuteThunkProtoDefaultTypeInternal;
extern CollectivePermuteThunkProtoDefaultTypeInternal _CollectivePermuteThunkProto_default_instance_;
class CollectivePermuteThunkProto_SourceTargetPairProto;
struct CollectivePermuteThunkProto_SourceTargetPairProtoDefaultTypeInternal;
extern CollectivePermuteThunkProto_SourceTargetPairProtoDefaultTypeInternal _CollectivePermuteThunkProto_SourceTargetPairProto_default_instance_;
class CollectiveThunkProto;
struct CollectiveThunkProtoDefaultTypeInternal;
extern CollectiveThunkProtoDefaultTypeInternal _CollectiveThunkProto_default_instance_;
class ConditionalThunkProto;
struct ConditionalThunkProtoDefaultTypeInternal;
extern ConditionalThunkProtoDefaultTypeInternal _ConditionalThunkProto_default_instance_;
class ConvolutionThunkProto;
struct ConvolutionThunkProtoDefaultTypeInternal;
extern ConvolutionThunkProtoDefaultTypeInternal _ConvolutionThunkProto_default_instance_;
class ConvolutionThunkProto_Options;
struct ConvolutionThunkProto_OptionsDefaultTypeInternal;
extern ConvolutionThunkProto_OptionsDefaultTypeInternal _ConvolutionThunkProto_Options_default_instance_;
class CopyThunkProto;
struct CopyThunkProtoDefaultTypeInternal;
extern CopyThunkProtoDefaultTypeInternal _CopyThunkProto_default_instance_;
class CustomCallThunkProto;
struct CustomCallThunkProtoDefaultTypeInternal;
extern CustomCallThunkProtoDefaultTypeInternal _CustomCallThunkProto_default_instance_;
class CustomCallThunkProto_OpBuffers;
struct CustomCallThunkProto_OpBuffersDefaultTypeInternal;
extern CustomCallThunkProto_OpBuffersDefaultTypeInternal _CustomCallThunkProto_OpBuffers_default_instance_;
class DotThunkProto;
struct DotThunkProtoDefaultTypeInternal;
extern DotThunkProtoDefaultTypeInternal _DotThunkProto_default_instance_;
class FftThunkProto;
struct FftThunkProtoDefaultTypeInternal;
extern FftThunkProtoDefaultTypeInternal _FftThunkProto_default_instance_;
class InfeedThunkProto;
struct InfeedThunkProtoDefaultTypeInternal;
extern InfeedThunkProtoDefaultTypeInternal _InfeedThunkProto_default_instance_;
class InfeedThunkProto_InfeedResource;
struct InfeedThunkProto_InfeedResourceDefaultTypeInternal;
extern InfeedThunkProto_InfeedResourceDefaultTypeInternal _InfeedThunkProto_InfeedResource_default_instance_;
class InfoProto;
struct InfoProtoDefaultTypeInternal;
extern InfoProtoDefaultTypeInternal _InfoProto_default_instance_;
class Int64Optional;
struct Int64OptionalDefaultTypeInternal;
extern Int64OptionalDefaultTypeInternal _Int64Optional_default_instance_;
class KernelThunkProto;
struct KernelThunkProtoDefaultTypeInternal;
extern KernelThunkProtoDefaultTypeInternal _KernelThunkProto_default_instance_;
class KernelThunkProto_ThreadDim;
struct KernelThunkProto_ThreadDimDefaultTypeInternal;
extern KernelThunkProto_ThreadDimDefaultTypeInternal _KernelThunkProto_ThreadDim_default_instance_;
class OpBuffersProto;
struct OpBuffersProtoDefaultTypeInternal;
extern OpBuffersProtoDefaultTypeInternal _OpBuffersProto_default_instance_;
class OpParamsProto;
struct OpParamsProtoDefaultTypeInternal;
extern OpParamsProtoDefaultTypeInternal _OpParamsProto_default_instance_;
class OpResourcesProto;
struct OpResourcesProtoDefaultTypeInternal;
extern OpResourcesProtoDefaultTypeInternal _OpResourcesProto_default_instance_;
class OutfeedThunkProto;
struct OutfeedThunkProtoDefaultTypeInternal;
extern OutfeedThunkProtoDefaultTypeInternal _OutfeedThunkProto_default_instance_;
class OutfeedThunkProto_OutfeedResource;
struct OutfeedThunkProto_OutfeedResourceDefaultTypeInternal;
extern OutfeedThunkProto_OutfeedResourceDefaultTypeInternal _OutfeedThunkProto_OutfeedResource_default_instance_;
class PartitionIdThunkProto;
struct PartitionIdThunkProtoDefaultTypeInternal;
extern PartitionIdThunkProtoDefaultTypeInternal _PartitionIdThunkProto_default_instance_;
class ReduceScatterThunkProto;
struct ReduceScatterThunkProtoDefaultTypeInternal;
extern ReduceScatterThunkProtoDefaultTypeInternal _ReduceScatterThunkProto_default_instance_;
class ReplicaIdThunkProto;
struct ReplicaIdThunkProtoDefaultTypeInternal;
extern ReplicaIdThunkProtoDefaultTypeInternal _ReplicaIdThunkProto_default_instance_;
class ResourceOptional;
struct ResourceOptionalDefaultTypeInternal;
extern ResourceOptionalDefaultTypeInternal _ResourceOptional_default_instance_;
class ResourceProto;
struct ResourceProtoDefaultTypeInternal;
extern ResourceProtoDefaultTypeInternal _ResourceProto_default_instance_;
class RngGetAndUpdateStateThunkProto;
struct RngGetAndUpdateStateThunkProtoDefaultTypeInternal;
extern RngGetAndUpdateStateThunkProtoDefaultTypeInternal _RngGetAndUpdateStateThunkProto_default_instance_;
class ShapeBufferAllocationSliceProto;
struct ShapeBufferAllocationSliceProtoDefaultTypeInternal;
extern ShapeBufferAllocationSliceProtoDefaultTypeInternal _ShapeBufferAllocationSliceProto_default_instance_;
class SortDirectionOptional;
struct SortDirectionOptionalDefaultTypeInternal;
extern SortDirectionOptionalDefaultTypeInternal _SortDirectionOptional_default_instance_;
class SortThunkProto;
struct SortThunkProtoDefaultTypeInternal;
extern SortThunkProtoDefaultTypeInternal _SortThunkProto_default_instance_;
class ThunkProto;
struct ThunkProtoDefaultTypeInternal;
extern ThunkProtoDefaultTypeInternal _ThunkProto_default_instance_;
class ThunkSequenceProto;
struct ThunkSequenceProtoDefaultTypeInternal;
extern ThunkSequenceProtoDefaultTypeInternal _ThunkSequenceProto_default_instance_;
class TopKThunkProto;
struct TopKThunkProtoDefaultTypeInternal;
extern TopKThunkProtoDefaultTypeInternal _TopKThunkProto_default_instance_;
class WhileThunkProto;
struct WhileThunkProtoDefaultTypeInternal;
extern WhileThunkProtoDefaultTypeInternal _WhileThunkProto_default_instance_;
class XnnConvolutionThunkProto;
struct XnnConvolutionThunkProtoDefaultTypeInternal;
extern XnnConvolutionThunkProtoDefaultTypeInternal _XnnConvolutionThunkProto_default_instance_;
class XnnDotThunkProto;
struct XnnDotThunkProtoDefaultTypeInternal;
extern XnnDotThunkProtoDefaultTypeInternal _XnnDotThunkProto_default_instance_;
class XnnFusionThunkProto;
struct XnnFusionThunkProtoDefaultTypeInternal;
extern XnnFusionThunkProtoDefaultTypeInternal _XnnFusionThunkProto_default_instance_;
class XnnFusionThunkProto_Options;
struct XnnFusionThunkProto_OptionsDefaultTypeInternal;
extern XnnFusionThunkProto_OptionsDefaultTypeInternal _XnnFusionThunkProto_Options_default_instance_;
}  // namespace cpu
}  // namespace xla
PROTOBUF_NAMESPACE_OPEN
template<> ::xla::cpu::AllGatherThunkProto* Arena::CreateMaybeMessage<::xla::cpu::AllGatherThunkProto>(Arena*);
template<> ::xla::cpu::AllReduceThunkProto* Arena::CreateMaybeMessage<::xla::cpu::AllReduceThunkProto>(Arena*);
template<> ::xla::cpu::AllToAllThunkProto* Arena::CreateMaybeMessage<::xla::cpu::AllToAllThunkProto>(Arena*);
template<> ::xla::cpu::BoolOptional* Arena::CreateMaybeMessage<::xla::cpu::BoolOptional>(Arena*);
template<> ::xla::cpu::BufferAllocationSliceProto* Arena::CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(Arena*);
template<> ::xla::cpu::CallThunkProto* Arena::CreateMaybeMessage<::xla::cpu::CallThunkProto>(Arena*);
template<> ::xla::cpu::CollectivePermuteThunkProto* Arena::CreateMaybeMessage<::xla::cpu::CollectivePermuteThunkProto>(Arena*);
template<> ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* Arena::CreateMaybeMessage<::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto>(Arena*);
template<> ::xla::cpu::CollectiveThunkProto* Arena::CreateMaybeMessage<::xla::cpu::CollectiveThunkProto>(Arena*);
template<> ::xla::cpu::ConditionalThunkProto* Arena::CreateMaybeMessage<::xla::cpu::ConditionalThunkProto>(Arena*);
template<> ::xla::cpu::ConvolutionThunkProto* Arena::CreateMaybeMessage<::xla::cpu::ConvolutionThunkProto>(Arena*);
template<> ::xla::cpu::ConvolutionThunkProto_Options* Arena::CreateMaybeMessage<::xla::cpu::ConvolutionThunkProto_Options>(Arena*);
template<> ::xla::cpu::CopyThunkProto* Arena::CreateMaybeMessage<::xla::cpu::CopyThunkProto>(Arena*);
template<> ::xla::cpu::CustomCallThunkProto* Arena::CreateMaybeMessage<::xla::cpu::CustomCallThunkProto>(Arena*);
template<> ::xla::cpu::CustomCallThunkProto_OpBuffers* Arena::CreateMaybeMessage<::xla::cpu::CustomCallThunkProto_OpBuffers>(Arena*);
template<> ::xla::cpu::DotThunkProto* Arena::CreateMaybeMessage<::xla::cpu::DotThunkProto>(Arena*);
template<> ::xla::cpu::FftThunkProto* Arena::CreateMaybeMessage<::xla::cpu::FftThunkProto>(Arena*);
template<> ::xla::cpu::InfeedThunkProto* Arena::CreateMaybeMessage<::xla::cpu::InfeedThunkProto>(Arena*);
template<> ::xla::cpu::InfeedThunkProto_InfeedResource* Arena::CreateMaybeMessage<::xla::cpu::InfeedThunkProto_InfeedResource>(Arena*);
template<> ::xla::cpu::InfoProto* Arena::CreateMaybeMessage<::xla::cpu::InfoProto>(Arena*);
template<> ::xla::cpu::Int64Optional* Arena::CreateMaybeMessage<::xla::cpu::Int64Optional>(Arena*);
template<> ::xla::cpu::KernelThunkProto* Arena::CreateMaybeMessage<::xla::cpu::KernelThunkProto>(Arena*);
template<> ::xla::cpu::KernelThunkProto_ThreadDim* Arena::CreateMaybeMessage<::xla::cpu::KernelThunkProto_ThreadDim>(Arena*);
template<> ::xla::cpu::OpBuffersProto* Arena::CreateMaybeMessage<::xla::cpu::OpBuffersProto>(Arena*);
template<> ::xla::cpu::OpParamsProto* Arena::CreateMaybeMessage<::xla::cpu::OpParamsProto>(Arena*);
template<> ::xla::cpu::OpResourcesProto* Arena::CreateMaybeMessage<::xla::cpu::OpResourcesProto>(Arena*);
template<> ::xla::cpu::OutfeedThunkProto* Arena::CreateMaybeMessage<::xla::cpu::OutfeedThunkProto>(Arena*);
template<> ::xla::cpu::OutfeedThunkProto_OutfeedResource* Arena::CreateMaybeMessage<::xla::cpu::OutfeedThunkProto_OutfeedResource>(Arena*);
template<> ::xla::cpu::PartitionIdThunkProto* Arena::CreateMaybeMessage<::xla::cpu::PartitionIdThunkProto>(Arena*);
template<> ::xla::cpu::ReduceScatterThunkProto* Arena::CreateMaybeMessage<::xla::cpu::ReduceScatterThunkProto>(Arena*);
template<> ::xla::cpu::ReplicaIdThunkProto* Arena::CreateMaybeMessage<::xla::cpu::ReplicaIdThunkProto>(Arena*);
template<> ::xla::cpu::ResourceOptional* Arena::CreateMaybeMessage<::xla::cpu::ResourceOptional>(Arena*);
template<> ::xla::cpu::ResourceProto* Arena::CreateMaybeMessage<::xla::cpu::ResourceProto>(Arena*);
template<> ::xla::cpu::RngGetAndUpdateStateThunkProto* Arena::CreateMaybeMessage<::xla::cpu::RngGetAndUpdateStateThunkProto>(Arena*);
template<> ::xla::cpu::ShapeBufferAllocationSliceProto* Arena::CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(Arena*);
template<> ::xla::cpu::SortDirectionOptional* Arena::CreateMaybeMessage<::xla::cpu::SortDirectionOptional>(Arena*);
template<> ::xla::cpu::SortThunkProto* Arena::CreateMaybeMessage<::xla::cpu::SortThunkProto>(Arena*);
template<> ::xla::cpu::ThunkProto* Arena::CreateMaybeMessage<::xla::cpu::ThunkProto>(Arena*);
template<> ::xla::cpu::ThunkSequenceProto* Arena::CreateMaybeMessage<::xla::cpu::ThunkSequenceProto>(Arena*);
template<> ::xla::cpu::TopKThunkProto* Arena::CreateMaybeMessage<::xla::cpu::TopKThunkProto>(Arena*);
template<> ::xla::cpu::WhileThunkProto* Arena::CreateMaybeMessage<::xla::cpu::WhileThunkProto>(Arena*);
template<> ::xla::cpu::XnnConvolutionThunkProto* Arena::CreateMaybeMessage<::xla::cpu::XnnConvolutionThunkProto>(Arena*);
template<> ::xla::cpu::XnnDotThunkProto* Arena::CreateMaybeMessage<::xla::cpu::XnnDotThunkProto>(Arena*);
template<> ::xla::cpu::XnnFusionThunkProto* Arena::CreateMaybeMessage<::xla::cpu::XnnFusionThunkProto>(Arena*);
template<> ::xla::cpu::XnnFusionThunkProto_Options* Arena::CreateMaybeMessage<::xla::cpu::XnnFusionThunkProto_Options>(Arena*);
PROTOBUF_NAMESPACE_CLOSE
namespace xla {
namespace cpu {

enum ResourceProto_Kind : int {
  ResourceProto_Kind_UNKNOWN = 0,
  ResourceProto_Kind_TOKEN = 1,
  ResourceProto_Kind_COLLECTIVE_COMMUNICATOR = 2,
  ResourceProto_Kind_ResourceProto_Kind_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::min(),
  ResourceProto_Kind_ResourceProto_Kind_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::max()
};
bool ResourceProto_Kind_IsValid(int value);
constexpr ResourceProto_Kind ResourceProto_Kind_Kind_MIN = ResourceProto_Kind_UNKNOWN;
constexpr ResourceProto_Kind ResourceProto_Kind_Kind_MAX = ResourceProto_Kind_COLLECTIVE_COMMUNICATOR;
constexpr int ResourceProto_Kind_Kind_ARRAYSIZE = ResourceProto_Kind_Kind_MAX + 1;

const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* ResourceProto_Kind_descriptor();
template<typename T>
inline const std::string& ResourceProto_Kind_Name(T enum_t_value) {
  static_assert(::std::is_same<T, ResourceProto_Kind>::value ||
    ::std::is_integral<T>::value,
    "Incorrect type passed to function ResourceProto_Kind_Name.");
  return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum(
    ResourceProto_Kind_descriptor(), enum_t_value);
}
inline bool ResourceProto_Kind_Parse(
    ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, ResourceProto_Kind* value) {
  return ::PROTOBUF_NAMESPACE_ID::internal::ParseNamedEnum<ResourceProto_Kind>(
    ResourceProto_Kind_descriptor(), name, value);
}
enum SortDirectionProto : int {
  UNKNOWN = 0,
  ASCENDING = 1,
  DESCENDING = 2,
  SortDirectionProto_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::min(),
  SortDirectionProto_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::max()
};
bool SortDirectionProto_IsValid(int value);
constexpr SortDirectionProto SortDirectionProto_MIN = UNKNOWN;
constexpr SortDirectionProto SortDirectionProto_MAX = DESCENDING;
constexpr int SortDirectionProto_ARRAYSIZE = SortDirectionProto_MAX + 1;

const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* SortDirectionProto_descriptor();
template<typename T>
inline const std::string& SortDirectionProto_Name(T enum_t_value) {
  static_assert(::std::is_same<T, SortDirectionProto>::value ||
    ::std::is_integral<T>::value,
    "Incorrect type passed to function SortDirectionProto_Name.");
  return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum(
    SortDirectionProto_descriptor(), enum_t_value);
}
inline bool SortDirectionProto_Parse(
    ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, SortDirectionProto* value) {
  return ::PROTOBUF_NAMESPACE_ID::internal::ParseNamedEnum<SortDirectionProto>(
    SortDirectionProto_descriptor(), name, value);
}
// ===================================================================

class ResourceProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ResourceProto) */ {
 public:
  inline ResourceProto() : ResourceProto(nullptr) {}
  ~ResourceProto() override;
  explicit PROTOBUF_CONSTEXPR ResourceProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ResourceProto(const ResourceProto& from);
  ResourceProto(ResourceProto&& from) noexcept
    : ResourceProto() {
    *this = ::std::move(from);
  }

  inline ResourceProto& operator=(const ResourceProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ResourceProto& operator=(ResourceProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ResourceProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ResourceProto* internal_default_instance() {
    return reinterpret_cast<const ResourceProto*>(
               &_ResourceProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    0;

  friend void swap(ResourceProto& a, ResourceProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ResourceProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ResourceProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ResourceProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ResourceProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ResourceProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ResourceProto& from) {
    ResourceProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ResourceProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ResourceProto";
  }
  protected:
  explicit ResourceProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef ResourceProto_Kind Kind;
  static constexpr Kind UNKNOWN =
    ResourceProto_Kind_UNKNOWN;
  static constexpr Kind TOKEN =
    ResourceProto_Kind_TOKEN;
  static constexpr Kind COLLECTIVE_COMMUNICATOR =
    ResourceProto_Kind_COLLECTIVE_COMMUNICATOR;
  static inline bool Kind_IsValid(int value) {
    return ResourceProto_Kind_IsValid(value);
  }
  static constexpr Kind Kind_MIN =
    ResourceProto_Kind_Kind_MIN;
  static constexpr Kind Kind_MAX =
    ResourceProto_Kind_Kind_MAX;
  static constexpr int Kind_ARRAYSIZE =
    ResourceProto_Kind_Kind_ARRAYSIZE;
  static inline const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor*
  Kind_descriptor() {
    return ResourceProto_Kind_descriptor();
  }
  template<typename T>
  static inline const std::string& Kind_Name(T enum_t_value) {
    static_assert(::std::is_same<T, Kind>::value ||
      ::std::is_integral<T>::value,
      "Incorrect type passed to function Kind_Name.");
    return ResourceProto_Kind_Name(enum_t_value);
  }
  static inline bool Kind_Parse(::PROTOBUF_NAMESPACE_ID::ConstStringParam name,
      Kind* value) {
    return ResourceProto_Kind_Parse(name, value);
  }

  // accessors -------------------------------------------------------

  enum : int {
    kKindFieldNumber = 1,
  };
  // .xla.cpu.ResourceProto.Kind kind = 1;
  void clear_kind();
  ::xla::cpu::ResourceProto_Kind kind() const;
  void set_kind(::xla::cpu::ResourceProto_Kind value);
  private:
  ::xla::cpu::ResourceProto_Kind _internal_kind() const;
  void _internal_set_kind(::xla::cpu::ResourceProto_Kind value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.ResourceProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int kind_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class BoolOptional final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.BoolOptional) */ {
 public:
  inline BoolOptional() : BoolOptional(nullptr) {}
  ~BoolOptional() override;
  explicit PROTOBUF_CONSTEXPR BoolOptional(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  BoolOptional(const BoolOptional& from);
  BoolOptional(BoolOptional&& from) noexcept
    : BoolOptional() {
    *this = ::std::move(from);
  }

  inline BoolOptional& operator=(const BoolOptional& from) {
    CopyFrom(from);
    return *this;
  }
  inline BoolOptional& operator=(BoolOptional&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const BoolOptional& default_instance() {
    return *internal_default_instance();
  }
  static inline const BoolOptional* internal_default_instance() {
    return reinterpret_cast<const BoolOptional*>(
               &_BoolOptional_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    1;

  friend void swap(BoolOptional& a, BoolOptional& b) {
    a.Swap(&b);
  }
  inline void Swap(BoolOptional* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(BoolOptional* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  BoolOptional* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<BoolOptional>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const BoolOptional& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const BoolOptional& from) {
    BoolOptional::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(BoolOptional* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.BoolOptional";
  }
  protected:
  explicit BoolOptional(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kValueFieldNumber = 1,
    kContainsValueFieldNumber = 2,
  };
  // bool value = 1;
  void clear_value();
  bool value() const;
  void set_value(bool value);
  private:
  bool _internal_value() const;
  void _internal_set_value(bool value);
  public:

  // bool contains_value = 2;
  void clear_contains_value();
  bool contains_value() const;
  void set_contains_value(bool value);
  private:
  bool _internal_contains_value() const;
  void _internal_set_contains_value(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.BoolOptional)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    bool value_;
    bool contains_value_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ResourceOptional final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ResourceOptional) */ {
 public:
  inline ResourceOptional() : ResourceOptional(nullptr) {}
  ~ResourceOptional() override;
  explicit PROTOBUF_CONSTEXPR ResourceOptional(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ResourceOptional(const ResourceOptional& from);
  ResourceOptional(ResourceOptional&& from) noexcept
    : ResourceOptional() {
    *this = ::std::move(from);
  }

  inline ResourceOptional& operator=(const ResourceOptional& from) {
    CopyFrom(from);
    return *this;
  }
  inline ResourceOptional& operator=(ResourceOptional&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ResourceOptional& default_instance() {
    return *internal_default_instance();
  }
  static inline const ResourceOptional* internal_default_instance() {
    return reinterpret_cast<const ResourceOptional*>(
               &_ResourceOptional_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    2;

  friend void swap(ResourceOptional& a, ResourceOptional& b) {
    a.Swap(&b);
  }
  inline void Swap(ResourceOptional* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ResourceOptional* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ResourceOptional* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ResourceOptional>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ResourceOptional& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ResourceOptional& from) {
    ResourceOptional::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ResourceOptional* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ResourceOptional";
  }
  protected:
  explicit ResourceOptional(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kValueFieldNumber = 1,
    kContainsValueFieldNumber = 2,
  };
  // .xla.cpu.ResourceProto value = 1;
  bool has_value() const;
  private:
  bool _internal_has_value() const;
  public:
  void clear_value();
  const ::xla::cpu::ResourceProto& value() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceProto* release_value();
  ::xla::cpu::ResourceProto* mutable_value();
  void set_allocated_value(::xla::cpu::ResourceProto* value);
  private:
  const ::xla::cpu::ResourceProto& _internal_value() const;
  ::xla::cpu::ResourceProto* _internal_mutable_value();
  public:
  void unsafe_arena_set_allocated_value(
      ::xla::cpu::ResourceProto* value);
  ::xla::cpu::ResourceProto* unsafe_arena_release_value();

  // bool contains_value = 2;
  void clear_contains_value();
  bool contains_value() const;
  void set_contains_value(bool value);
  private:
  bool _internal_contains_value() const;
  void _internal_set_contains_value(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.ResourceOptional)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ResourceProto* value_;
    bool contains_value_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class Int64Optional final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.Int64Optional) */ {
 public:
  inline Int64Optional() : Int64Optional(nullptr) {}
  ~Int64Optional() override;
  explicit PROTOBUF_CONSTEXPR Int64Optional(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  Int64Optional(const Int64Optional& from);
  Int64Optional(Int64Optional&& from) noexcept
    : Int64Optional() {
    *this = ::std::move(from);
  }

  inline Int64Optional& operator=(const Int64Optional& from) {
    CopyFrom(from);
    return *this;
  }
  inline Int64Optional& operator=(Int64Optional&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const Int64Optional& default_instance() {
    return *internal_default_instance();
  }
  static inline const Int64Optional* internal_default_instance() {
    return reinterpret_cast<const Int64Optional*>(
               &_Int64Optional_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    3;

  friend void swap(Int64Optional& a, Int64Optional& b) {
    a.Swap(&b);
  }
  inline void Swap(Int64Optional* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(Int64Optional* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  Int64Optional* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<Int64Optional>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const Int64Optional& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const Int64Optional& from) {
    Int64Optional::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(Int64Optional* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.Int64Optional";
  }
  protected:
  explicit Int64Optional(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kValueFieldNumber = 1,
    kContainsValueFieldNumber = 2,
  };
  // int64 value = 1;
  void clear_value();
  int64_t value() const;
  void set_value(int64_t value);
  private:
  int64_t _internal_value() const;
  void _internal_set_value(int64_t value);
  public:

  // bool contains_value = 2;
  void clear_contains_value();
  bool contains_value() const;
  void set_contains_value(bool value);
  private:
  bool _internal_contains_value() const;
  void _internal_set_contains_value(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.Int64Optional)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int64_t value_;
    bool contains_value_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class SortDirectionOptional final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.SortDirectionOptional) */ {
 public:
  inline SortDirectionOptional() : SortDirectionOptional(nullptr) {}
  ~SortDirectionOptional() override;
  explicit PROTOBUF_CONSTEXPR SortDirectionOptional(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  SortDirectionOptional(const SortDirectionOptional& from);
  SortDirectionOptional(SortDirectionOptional&& from) noexcept
    : SortDirectionOptional() {
    *this = ::std::move(from);
  }

  inline SortDirectionOptional& operator=(const SortDirectionOptional& from) {
    CopyFrom(from);
    return *this;
  }
  inline SortDirectionOptional& operator=(SortDirectionOptional&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const SortDirectionOptional& default_instance() {
    return *internal_default_instance();
  }
  static inline const SortDirectionOptional* internal_default_instance() {
    return reinterpret_cast<const SortDirectionOptional*>(
               &_SortDirectionOptional_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    4;

  friend void swap(SortDirectionOptional& a, SortDirectionOptional& b) {
    a.Swap(&b);
  }
  inline void Swap(SortDirectionOptional* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(SortDirectionOptional* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  SortDirectionOptional* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<SortDirectionOptional>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const SortDirectionOptional& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const SortDirectionOptional& from) {
    SortDirectionOptional::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(SortDirectionOptional* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.SortDirectionOptional";
  }
  protected:
  explicit SortDirectionOptional(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kValueFieldNumber = 1,
    kContainsValueFieldNumber = 2,
  };
  // .xla.cpu.SortDirectionProto value = 1;
  void clear_value();
  ::xla::cpu::SortDirectionProto value() const;
  void set_value(::xla::cpu::SortDirectionProto value);
  private:
  ::xla::cpu::SortDirectionProto _internal_value() const;
  void _internal_set_value(::xla::cpu::SortDirectionProto value);
  public:

  // bool contains_value = 2;
  void clear_contains_value();
  bool contains_value() const;
  void set_contains_value(bool value);
  private:
  bool _internal_contains_value() const;
  void _internal_set_contains_value(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.SortDirectionOptional)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int value_;
    bool contains_value_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class BufferAllocationSliceProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.BufferAllocationSliceProto) */ {
 public:
  inline BufferAllocationSliceProto() : BufferAllocationSliceProto(nullptr) {}
  ~BufferAllocationSliceProto() override;
  explicit PROTOBUF_CONSTEXPR BufferAllocationSliceProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  BufferAllocationSliceProto(const BufferAllocationSliceProto& from);
  BufferAllocationSliceProto(BufferAllocationSliceProto&& from) noexcept
    : BufferAllocationSliceProto() {
    *this = ::std::move(from);
  }

  inline BufferAllocationSliceProto& operator=(const BufferAllocationSliceProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline BufferAllocationSliceProto& operator=(BufferAllocationSliceProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const BufferAllocationSliceProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const BufferAllocationSliceProto* internal_default_instance() {
    return reinterpret_cast<const BufferAllocationSliceProto*>(
               &_BufferAllocationSliceProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    5;

  friend void swap(BufferAllocationSliceProto& a, BufferAllocationSliceProto& b) {
    a.Swap(&b);
  }
  inline void Swap(BufferAllocationSliceProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(BufferAllocationSliceProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  BufferAllocationSliceProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<BufferAllocationSliceProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const BufferAllocationSliceProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const BufferAllocationSliceProto& from) {
    BufferAllocationSliceProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(BufferAllocationSliceProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.BufferAllocationSliceProto";
  }
  protected:
  explicit BufferAllocationSliceProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kOffsetFieldNumber = 1,
    kSizeFieldNumber = 2,
    kBufferAllocationIndexFieldNumber = 3,
  };
  // int64 offset = 1;
  void clear_offset();
  int64_t offset() const;
  void set_offset(int64_t value);
  private:
  int64_t _internal_offset() const;
  void _internal_set_offset(int64_t value);
  public:

  // int64 size = 2;
  void clear_size();
  int64_t size() const;
  void set_size(int64_t value);
  private:
  int64_t _internal_size() const;
  void _internal_set_size(int64_t value);
  public:

  // int64 buffer_allocation_index = 3;
  void clear_buffer_allocation_index();
  int64_t buffer_allocation_index() const;
  void set_buffer_allocation_index(int64_t value);
  private:
  int64_t _internal_buffer_allocation_index() const;
  void _internal_set_buffer_allocation_index(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.BufferAllocationSliceProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int64_t offset_;
    int64_t size_;
    int64_t buffer_allocation_index_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ShapeBufferAllocationSliceProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ShapeBufferAllocationSliceProto) */ {
 public:
  inline ShapeBufferAllocationSliceProto() : ShapeBufferAllocationSliceProto(nullptr) {}
  ~ShapeBufferAllocationSliceProto() override;
  explicit PROTOBUF_CONSTEXPR ShapeBufferAllocationSliceProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ShapeBufferAllocationSliceProto(const ShapeBufferAllocationSliceProto& from);
  ShapeBufferAllocationSliceProto(ShapeBufferAllocationSliceProto&& from) noexcept
    : ShapeBufferAllocationSliceProto() {
    *this = ::std::move(from);
  }

  inline ShapeBufferAllocationSliceProto& operator=(const ShapeBufferAllocationSliceProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ShapeBufferAllocationSliceProto& operator=(ShapeBufferAllocationSliceProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ShapeBufferAllocationSliceProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ShapeBufferAllocationSliceProto* internal_default_instance() {
    return reinterpret_cast<const ShapeBufferAllocationSliceProto*>(
               &_ShapeBufferAllocationSliceProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    6;

  friend void swap(ShapeBufferAllocationSliceProto& a, ShapeBufferAllocationSliceProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ShapeBufferAllocationSliceProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ShapeBufferAllocationSliceProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ShapeBufferAllocationSliceProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ShapeBufferAllocationSliceProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ShapeBufferAllocationSliceProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ShapeBufferAllocationSliceProto& from) {
    ShapeBufferAllocationSliceProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ShapeBufferAllocationSliceProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ShapeBufferAllocationSliceProto";
  }
  protected:
  explicit ShapeBufferAllocationSliceProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kShapeFieldNumber = 1,
    kSliceFieldNumber = 2,
  };
  // .xla.ShapeProto shape = 1;
  bool has_shape() const;
  private:
  bool _internal_has_shape() const;
  public:
  void clear_shape();
  const ::xla::ShapeProto& shape() const;
  PROTOBUF_NODISCARD ::xla::ShapeProto* release_shape();
  ::xla::ShapeProto* mutable_shape();
  void set_allocated_shape(::xla::ShapeProto* shape);
  private:
  const ::xla::ShapeProto& _internal_shape() const;
  ::xla::ShapeProto* _internal_mutable_shape();
  public:
  void unsafe_arena_set_allocated_shape(
      ::xla::ShapeProto* shape);
  ::xla::ShapeProto* unsafe_arena_release_shape();

  // .xla.cpu.BufferAllocationSliceProto slice = 2;
  bool has_slice() const;
  private:
  bool _internal_has_slice() const;
  public:
  void clear_slice();
  const ::xla::cpu::BufferAllocationSliceProto& slice() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_slice();
  ::xla::cpu::BufferAllocationSliceProto* mutable_slice();
  void set_allocated_slice(::xla::cpu::BufferAllocationSliceProto* slice);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_slice() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_slice();
  public:
  void unsafe_arena_set_allocated_slice(
      ::xla::cpu::BufferAllocationSliceProto* slice);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_slice();

  // @@protoc_insertion_point(class_scope:xla.cpu.ShapeBufferAllocationSliceProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::ShapeProto* shape_;
    ::xla::cpu::BufferAllocationSliceProto* slice_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class OpParamsProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.OpParamsProto) */ {
 public:
  inline OpParamsProto() : OpParamsProto(nullptr) {}
  ~OpParamsProto() override;
  explicit PROTOBUF_CONSTEXPR OpParamsProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  OpParamsProto(const OpParamsProto& from);
  OpParamsProto(OpParamsProto&& from) noexcept
    : OpParamsProto() {
    *this = ::std::move(from);
  }

  inline OpParamsProto& operator=(const OpParamsProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline OpParamsProto& operator=(OpParamsProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const OpParamsProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const OpParamsProto* internal_default_instance() {
    return reinterpret_cast<const OpParamsProto*>(
               &_OpParamsProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    7;

  friend void swap(OpParamsProto& a, OpParamsProto& b) {
    a.Swap(&b);
  }
  inline void Swap(OpParamsProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(OpParamsProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  OpParamsProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<OpParamsProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const OpParamsProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const OpParamsProto& from) {
    OpParamsProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(OpParamsProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.OpParamsProto";
  }
  protected:
  explicit OpParamsProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kReplicaGroupFieldNumber = 4,
    kUseGlobalDeviceIdsFieldNumber = 3,
    kOpIdFieldNumber = 1,
    kHasChannelIdFieldNumber = 2,
  };
  // repeated .xla.ReplicaGroup replica_group = 4;
  int replica_group_size() const;
  private:
  int _internal_replica_group_size() const;
  public:
  void clear_replica_group();
  ::xla::ReplicaGroup* mutable_replica_group(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::ReplicaGroup >*
      mutable_replica_group();
  private:
  const ::xla::ReplicaGroup& _internal_replica_group(int index) const;
  ::xla::ReplicaGroup* _internal_add_replica_group();
  public:
  const ::xla::ReplicaGroup& replica_group(int index) const;
  ::xla::ReplicaGroup* add_replica_group();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::ReplicaGroup >&
      replica_group() const;

  // .xla.cpu.BoolOptional use_global_device_ids = 3;
  bool has_use_global_device_ids() const;
  private:
  bool _internal_has_use_global_device_ids() const;
  public:
  void clear_use_global_device_ids();
  const ::xla::cpu::BoolOptional& use_global_device_ids() const;
  PROTOBUF_NODISCARD ::xla::cpu::BoolOptional* release_use_global_device_ids();
  ::xla::cpu::BoolOptional* mutable_use_global_device_ids();
  void set_allocated_use_global_device_ids(::xla::cpu::BoolOptional* use_global_device_ids);
  private:
  const ::xla::cpu::BoolOptional& _internal_use_global_device_ids() const;
  ::xla::cpu::BoolOptional* _internal_mutable_use_global_device_ids();
  public:
  void unsafe_arena_set_allocated_use_global_device_ids(
      ::xla::cpu::BoolOptional* use_global_device_ids);
  ::xla::cpu::BoolOptional* unsafe_arena_release_use_global_device_ids();

  // int64 op_id = 1;
  void clear_op_id();
  int64_t op_id() const;
  void set_op_id(int64_t value);
  private:
  int64_t _internal_op_id() const;
  void _internal_set_op_id(int64_t value);
  public:

  // bool has_channel_id = 2;
  void clear_has_channel_id();
  bool has_channel_id() const;
  void set_has_channel_id(bool value);
  private:
  bool _internal_has_channel_id() const;
  void _internal_set_has_channel_id(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.OpParamsProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::ReplicaGroup > replica_group_;
    ::xla::cpu::BoolOptional* use_global_device_ids_;
    int64_t op_id_;
    bool has_channel_id_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class OpBuffersProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.OpBuffersProto) */ {
 public:
  inline OpBuffersProto() : OpBuffersProto(nullptr) {}
  ~OpBuffersProto() override;
  explicit PROTOBUF_CONSTEXPR OpBuffersProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  OpBuffersProto(const OpBuffersProto& from);
  OpBuffersProto(OpBuffersProto&& from) noexcept
    : OpBuffersProto() {
    *this = ::std::move(from);
  }

  inline OpBuffersProto& operator=(const OpBuffersProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline OpBuffersProto& operator=(OpBuffersProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const OpBuffersProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const OpBuffersProto* internal_default_instance() {
    return reinterpret_cast<const OpBuffersProto*>(
               &_OpBuffersProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    8;

  friend void swap(OpBuffersProto& a, OpBuffersProto& b) {
    a.Swap(&b);
  }
  inline void Swap(OpBuffersProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(OpBuffersProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  OpBuffersProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<OpBuffersProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const OpBuffersProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const OpBuffersProto& from) {
    OpBuffersProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(OpBuffersProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.OpBuffersProto";
  }
  protected:
  explicit OpBuffersProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kSourceShapesBufferSlicesFieldNumber = 1,
    kDestinationShapesBufferSlicesFieldNumber = 2,
  };
  // repeated .xla.cpu.ShapeBufferAllocationSliceProto source_shapes_buffer_slices = 1;
  int source_shapes_buffer_slices_size() const;
  private:
  int _internal_source_shapes_buffer_slices_size() const;
  public:
  void clear_source_shapes_buffer_slices();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_source_shapes_buffer_slices(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_source_shapes_buffer_slices();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_source_shapes_buffer_slices(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_source_shapes_buffer_slices();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& source_shapes_buffer_slices(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_source_shapes_buffer_slices();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      source_shapes_buffer_slices() const;

  // repeated .xla.cpu.ShapeBufferAllocationSliceProto destination_shapes_buffer_slices = 2;
  int destination_shapes_buffer_slices_size() const;
  private:
  int _internal_destination_shapes_buffer_slices_size() const;
  public:
  void clear_destination_shapes_buffer_slices();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_destination_shapes_buffer_slices(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_destination_shapes_buffer_slices();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_destination_shapes_buffer_slices(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_destination_shapes_buffer_slices();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& destination_shapes_buffer_slices(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_destination_shapes_buffer_slices();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      destination_shapes_buffer_slices() const;

  // @@protoc_insertion_point(class_scope:xla.cpu.OpBuffersProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > source_shapes_buffer_slices_;
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > destination_shapes_buffer_slices_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class OpResourcesProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.OpResourcesProto) */ {
 public:
  inline OpResourcesProto() : OpResourcesProto(nullptr) {}
  ~OpResourcesProto() override;
  explicit PROTOBUF_CONSTEXPR OpResourcesProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  OpResourcesProto(const OpResourcesProto& from);
  OpResourcesProto(OpResourcesProto&& from) noexcept
    : OpResourcesProto() {
    *this = ::std::move(from);
  }

  inline OpResourcesProto& operator=(const OpResourcesProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline OpResourcesProto& operator=(OpResourcesProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const OpResourcesProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const OpResourcesProto* internal_default_instance() {
    return reinterpret_cast<const OpResourcesProto*>(
               &_OpResourcesProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    9;

  friend void swap(OpResourcesProto& a, OpResourcesProto& b) {
    a.Swap(&b);
  }
  inline void Swap(OpResourcesProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(OpResourcesProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  OpResourcesProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<OpResourcesProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const OpResourcesProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const OpResourcesProto& from) {
    OpResourcesProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(OpResourcesProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.OpResourcesProto";
  }
  protected:
  explicit OpResourcesProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kCommunicatorResourceFieldNumber = 1,
  };
  // .xla.cpu.ResourceOptional communicator_resource = 1;
  bool has_communicator_resource() const;
  private:
  bool _internal_has_communicator_resource() const;
  public:
  void clear_communicator_resource();
  const ::xla::cpu::ResourceOptional& communicator_resource() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceOptional* release_communicator_resource();
  ::xla::cpu::ResourceOptional* mutable_communicator_resource();
  void set_allocated_communicator_resource(::xla::cpu::ResourceOptional* communicator_resource);
  private:
  const ::xla::cpu::ResourceOptional& _internal_communicator_resource() const;
  ::xla::cpu::ResourceOptional* _internal_mutable_communicator_resource();
  public:
  void unsafe_arena_set_allocated_communicator_resource(
      ::xla::cpu::ResourceOptional* communicator_resource);
  ::xla::cpu::ResourceOptional* unsafe_arena_release_communicator_resource();

  // @@protoc_insertion_point(class_scope:xla.cpu.OpResourcesProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ResourceOptional* communicator_resource_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class AllGatherThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase /* @@protoc_insertion_point(class_definition:xla.cpu.AllGatherThunkProto) */ {
 public:
  inline AllGatherThunkProto() : AllGatherThunkProto(nullptr) {}
  explicit PROTOBUF_CONSTEXPR AllGatherThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  AllGatherThunkProto(const AllGatherThunkProto& from);
  AllGatherThunkProto(AllGatherThunkProto&& from) noexcept
    : AllGatherThunkProto() {
    *this = ::std::move(from);
  }

  inline AllGatherThunkProto& operator=(const AllGatherThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline AllGatherThunkProto& operator=(AllGatherThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const AllGatherThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const AllGatherThunkProto* internal_default_instance() {
    return reinterpret_cast<const AllGatherThunkProto*>(
               &_AllGatherThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    10;

  friend void swap(AllGatherThunkProto& a, AllGatherThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(AllGatherThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(AllGatherThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  AllGatherThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<AllGatherThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::CopyFrom;
  inline void CopyFrom(const AllGatherThunkProto& from) {
    ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::CopyImpl(*this, from);
  }
  using ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::MergeFrom;
  void MergeFrom(const AllGatherThunkProto& from) {
    ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::MergeImpl(*this, from);
  }
  public:

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.AllGatherThunkProto";
  }
  protected:
  explicit AllGatherThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  // @@protoc_insertion_point(class_scope:xla.cpu.AllGatherThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
  };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class AllReduceThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.AllReduceThunkProto) */ {
 public:
  inline AllReduceThunkProto() : AllReduceThunkProto(nullptr) {}
  ~AllReduceThunkProto() override;
  explicit PROTOBUF_CONSTEXPR AllReduceThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  AllReduceThunkProto(const AllReduceThunkProto& from);
  AllReduceThunkProto(AllReduceThunkProto&& from) noexcept
    : AllReduceThunkProto() {
    *this = ::std::move(from);
  }

  inline AllReduceThunkProto& operator=(const AllReduceThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline AllReduceThunkProto& operator=(AllReduceThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const AllReduceThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const AllReduceThunkProto* internal_default_instance() {
    return reinterpret_cast<const AllReduceThunkProto*>(
               &_AllReduceThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    11;

  friend void swap(AllReduceThunkProto& a, AllReduceThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(AllReduceThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(AllReduceThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  AllReduceThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<AllReduceThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const AllReduceThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const AllReduceThunkProto& from) {
    AllReduceThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(AllReduceThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.AllReduceThunkProto";
  }
  protected:
  explicit AllReduceThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kReductionKindFieldNumber = 1,
    kSingleReplicaFieldNumber = 2,
  };
  // string reduction_kind = 1;
  void clear_reduction_kind();
  const std::string& reduction_kind() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_reduction_kind(ArgT0&& arg0, ArgT... args);
  std::string* mutable_reduction_kind();
  PROTOBUF_NODISCARD std::string* release_reduction_kind();
  void set_allocated_reduction_kind(std::string* reduction_kind);
  private:
  const std::string& _internal_reduction_kind() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_reduction_kind(const std::string& value);
  std::string* _internal_mutable_reduction_kind();
  public:

  // bool single_replica = 2;
  void clear_single_replica();
  bool single_replica() const;
  void set_single_replica(bool value);
  private:
  bool _internal_single_replica() const;
  void _internal_set_single_replica(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.AllReduceThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr reduction_kind_;
    bool single_replica_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class AllToAllThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase /* @@protoc_insertion_point(class_definition:xla.cpu.AllToAllThunkProto) */ {
 public:
  inline AllToAllThunkProto() : AllToAllThunkProto(nullptr) {}
  explicit PROTOBUF_CONSTEXPR AllToAllThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  AllToAllThunkProto(const AllToAllThunkProto& from);
  AllToAllThunkProto(AllToAllThunkProto&& from) noexcept
    : AllToAllThunkProto() {
    *this = ::std::move(from);
  }

  inline AllToAllThunkProto& operator=(const AllToAllThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline AllToAllThunkProto& operator=(AllToAllThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const AllToAllThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const AllToAllThunkProto* internal_default_instance() {
    return reinterpret_cast<const AllToAllThunkProto*>(
               &_AllToAllThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    12;

  friend void swap(AllToAllThunkProto& a, AllToAllThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(AllToAllThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(AllToAllThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  AllToAllThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<AllToAllThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::CopyFrom;
  inline void CopyFrom(const AllToAllThunkProto& from) {
    ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::CopyImpl(*this, from);
  }
  using ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::MergeFrom;
  void MergeFrom(const AllToAllThunkProto& from) {
    ::PROTOBUF_NAMESPACE_ID::internal::ZeroFieldsBase::MergeImpl(*this, from);
  }
  public:

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.AllToAllThunkProto";
  }
  protected:
  explicit AllToAllThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  // @@protoc_insertion_point(class_scope:xla.cpu.AllToAllThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
  };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ReduceScatterThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ReduceScatterThunkProto) */ {
 public:
  inline ReduceScatterThunkProto() : ReduceScatterThunkProto(nullptr) {}
  ~ReduceScatterThunkProto() override;
  explicit PROTOBUF_CONSTEXPR ReduceScatterThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ReduceScatterThunkProto(const ReduceScatterThunkProto& from);
  ReduceScatterThunkProto(ReduceScatterThunkProto&& from) noexcept
    : ReduceScatterThunkProto() {
    *this = ::std::move(from);
  }

  inline ReduceScatterThunkProto& operator=(const ReduceScatterThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ReduceScatterThunkProto& operator=(ReduceScatterThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ReduceScatterThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ReduceScatterThunkProto* internal_default_instance() {
    return reinterpret_cast<const ReduceScatterThunkProto*>(
               &_ReduceScatterThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    13;

  friend void swap(ReduceScatterThunkProto& a, ReduceScatterThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ReduceScatterThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ReduceScatterThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ReduceScatterThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ReduceScatterThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ReduceScatterThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ReduceScatterThunkProto& from) {
    ReduceScatterThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ReduceScatterThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ReduceScatterThunkProto";
  }
  protected:
  explicit ReduceScatterThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kReductionKindFieldNumber = 1,
  };
  // string reduction_kind = 1;
  void clear_reduction_kind();
  const std::string& reduction_kind() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_reduction_kind(ArgT0&& arg0, ArgT... args);
  std::string* mutable_reduction_kind();
  PROTOBUF_NODISCARD std::string* release_reduction_kind();
  void set_allocated_reduction_kind(std::string* reduction_kind);
  private:
  const std::string& _internal_reduction_kind() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_reduction_kind(const std::string& value);
  std::string* _internal_mutable_reduction_kind();
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.ReduceScatterThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr reduction_kind_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CollectivePermuteThunkProto_SourceTargetPairProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto) */ {
 public:
  inline CollectivePermuteThunkProto_SourceTargetPairProto() : CollectivePermuteThunkProto_SourceTargetPairProto(nullptr) {}
  ~CollectivePermuteThunkProto_SourceTargetPairProto() override;
  explicit PROTOBUF_CONSTEXPR CollectivePermuteThunkProto_SourceTargetPairProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CollectivePermuteThunkProto_SourceTargetPairProto(const CollectivePermuteThunkProto_SourceTargetPairProto& from);
  CollectivePermuteThunkProto_SourceTargetPairProto(CollectivePermuteThunkProto_SourceTargetPairProto&& from) noexcept
    : CollectivePermuteThunkProto_SourceTargetPairProto() {
    *this = ::std::move(from);
  }

  inline CollectivePermuteThunkProto_SourceTargetPairProto& operator=(const CollectivePermuteThunkProto_SourceTargetPairProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CollectivePermuteThunkProto_SourceTargetPairProto& operator=(CollectivePermuteThunkProto_SourceTargetPairProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CollectivePermuteThunkProto_SourceTargetPairProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const CollectivePermuteThunkProto_SourceTargetPairProto* internal_default_instance() {
    return reinterpret_cast<const CollectivePermuteThunkProto_SourceTargetPairProto*>(
               &_CollectivePermuteThunkProto_SourceTargetPairProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    14;

  friend void swap(CollectivePermuteThunkProto_SourceTargetPairProto& a, CollectivePermuteThunkProto_SourceTargetPairProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CollectivePermuteThunkProto_SourceTargetPairProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CollectivePermuteThunkProto_SourceTargetPairProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CollectivePermuteThunkProto_SourceTargetPairProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CollectivePermuteThunkProto_SourceTargetPairProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CollectivePermuteThunkProto_SourceTargetPairProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CollectivePermuteThunkProto_SourceTargetPairProto& from) {
    CollectivePermuteThunkProto_SourceTargetPairProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CollectivePermuteThunkProto_SourceTargetPairProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto";
  }
  protected:
  explicit CollectivePermuteThunkProto_SourceTargetPairProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kSourceFieldNumber = 1,
    kTargetFieldNumber = 2,
  };
  // int64 source = 1;
  void clear_source();
  int64_t source() const;
  void set_source(int64_t value);
  private:
  int64_t _internal_source() const;
  void _internal_set_source(int64_t value);
  public:

  // int64 target = 2;
  void clear_target();
  int64_t target() const;
  void set_target(int64_t value);
  private:
  int64_t _internal_target() const;
  void _internal_set_target(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int64_t source_;
    int64_t target_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CollectivePermuteThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CollectivePermuteThunkProto) */ {
 public:
  inline CollectivePermuteThunkProto() : CollectivePermuteThunkProto(nullptr) {}
  ~CollectivePermuteThunkProto() override;
  explicit PROTOBUF_CONSTEXPR CollectivePermuteThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CollectivePermuteThunkProto(const CollectivePermuteThunkProto& from);
  CollectivePermuteThunkProto(CollectivePermuteThunkProto&& from) noexcept
    : CollectivePermuteThunkProto() {
    *this = ::std::move(from);
  }

  inline CollectivePermuteThunkProto& operator=(const CollectivePermuteThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CollectivePermuteThunkProto& operator=(CollectivePermuteThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CollectivePermuteThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const CollectivePermuteThunkProto* internal_default_instance() {
    return reinterpret_cast<const CollectivePermuteThunkProto*>(
               &_CollectivePermuteThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    15;

  friend void swap(CollectivePermuteThunkProto& a, CollectivePermuteThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CollectivePermuteThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CollectivePermuteThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CollectivePermuteThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CollectivePermuteThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CollectivePermuteThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CollectivePermuteThunkProto& from) {
    CollectivePermuteThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CollectivePermuteThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CollectivePermuteThunkProto";
  }
  protected:
  explicit CollectivePermuteThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef CollectivePermuteThunkProto_SourceTargetPairProto SourceTargetPairProto;

  // accessors -------------------------------------------------------

  enum : int {
    kSourceTargetPairsFieldNumber = 1,
  };
  // repeated .xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto source_target_pairs = 1;
  int source_target_pairs_size() const;
  private:
  int _internal_source_target_pairs_size() const;
  public:
  void clear_source_target_pairs();
  ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* mutable_source_target_pairs(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto >*
      mutable_source_target_pairs();
  private:
  const ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto& _internal_source_target_pairs(int index) const;
  ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* _internal_add_source_target_pairs();
  public:
  const ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto& source_target_pairs(int index) const;
  ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* add_source_target_pairs();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto >&
      source_target_pairs() const;

  // @@protoc_insertion_point(class_scope:xla.cpu.CollectivePermuteThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto > source_target_pairs_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CollectiveThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CollectiveThunkProto) */ {
 public:
  inline CollectiveThunkProto() : CollectiveThunkProto(nullptr) {}
  ~CollectiveThunkProto() override;
  explicit PROTOBUF_CONSTEXPR CollectiveThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CollectiveThunkProto(const CollectiveThunkProto& from);
  CollectiveThunkProto(CollectiveThunkProto&& from) noexcept
    : CollectiveThunkProto() {
    *this = ::std::move(from);
  }

  inline CollectiveThunkProto& operator=(const CollectiveThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CollectiveThunkProto& operator=(CollectiveThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CollectiveThunkProto& default_instance() {
    return *internal_default_instance();
  }
  enum ImplCase {
    kAllGatherThunk = 4,
    kAllReduceThunk = 5,
    kAllToAllThunk = 6,
    kReduceScatterThunk = 7,
    kCollectivePermuteThunk = 8,
    IMPL_NOT_SET = 0,
  };

  static inline const CollectiveThunkProto* internal_default_instance() {
    return reinterpret_cast<const CollectiveThunkProto*>(
               &_CollectiveThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    16;

  friend void swap(CollectiveThunkProto& a, CollectiveThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CollectiveThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CollectiveThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CollectiveThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CollectiveThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CollectiveThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CollectiveThunkProto& from) {
    CollectiveThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CollectiveThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CollectiveThunkProto";
  }
  protected:
  explicit CollectiveThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kOpParamsFieldNumber = 1,
    kOpBuffersFieldNumber = 2,
    kOpResourcesFieldNumber = 3,
    kAllGatherThunkFieldNumber = 4,
    kAllReduceThunkFieldNumber = 5,
    kAllToAllThunkFieldNumber = 6,
    kReduceScatterThunkFieldNumber = 7,
    kCollectivePermuteThunkFieldNumber = 8,
  };
  // .xla.cpu.OpParamsProto op_params = 1;
  bool has_op_params() const;
  private:
  bool _internal_has_op_params() const;
  public:
  void clear_op_params();
  const ::xla::cpu::OpParamsProto& op_params() const;
  PROTOBUF_NODISCARD ::xla::cpu::OpParamsProto* release_op_params();
  ::xla::cpu::OpParamsProto* mutable_op_params();
  void set_allocated_op_params(::xla::cpu::OpParamsProto* op_params);
  private:
  const ::xla::cpu::OpParamsProto& _internal_op_params() const;
  ::xla::cpu::OpParamsProto* _internal_mutable_op_params();
  public:
  void unsafe_arena_set_allocated_op_params(
      ::xla::cpu::OpParamsProto* op_params);
  ::xla::cpu::OpParamsProto* unsafe_arena_release_op_params();

  // .xla.cpu.OpBuffersProto op_buffers = 2;
  bool has_op_buffers() const;
  private:
  bool _internal_has_op_buffers() const;
  public:
  void clear_op_buffers();
  const ::xla::cpu::OpBuffersProto& op_buffers() const;
  PROTOBUF_NODISCARD ::xla::cpu::OpBuffersProto* release_op_buffers();
  ::xla::cpu::OpBuffersProto* mutable_op_buffers();
  void set_allocated_op_buffers(::xla::cpu::OpBuffersProto* op_buffers);
  private:
  const ::xla::cpu::OpBuffersProto& _internal_op_buffers() const;
  ::xla::cpu::OpBuffersProto* _internal_mutable_op_buffers();
  public:
  void unsafe_arena_set_allocated_op_buffers(
      ::xla::cpu::OpBuffersProto* op_buffers);
  ::xla::cpu::OpBuffersProto* unsafe_arena_release_op_buffers();

  // .xla.cpu.OpResourcesProto op_resources = 3;
  bool has_op_resources() const;
  private:
  bool _internal_has_op_resources() const;
  public:
  void clear_op_resources();
  const ::xla::cpu::OpResourcesProto& op_resources() const;
  PROTOBUF_NODISCARD ::xla::cpu::OpResourcesProto* release_op_resources();
  ::xla::cpu::OpResourcesProto* mutable_op_resources();
  void set_allocated_op_resources(::xla::cpu::OpResourcesProto* op_resources);
  private:
  const ::xla::cpu::OpResourcesProto& _internal_op_resources() const;
  ::xla::cpu::OpResourcesProto* _internal_mutable_op_resources();
  public:
  void unsafe_arena_set_allocated_op_resources(
      ::xla::cpu::OpResourcesProto* op_resources);
  ::xla::cpu::OpResourcesProto* unsafe_arena_release_op_resources();

  // .xla.cpu.AllGatherThunkProto all_gather_thunk = 4;
  bool has_all_gather_thunk() const;
  private:
  bool _internal_has_all_gather_thunk() const;
  public:
  void clear_all_gather_thunk();
  const ::xla::cpu::AllGatherThunkProto& all_gather_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::AllGatherThunkProto* release_all_gather_thunk();
  ::xla::cpu::AllGatherThunkProto* mutable_all_gather_thunk();
  void set_allocated_all_gather_thunk(::xla::cpu::AllGatherThunkProto* all_gather_thunk);
  private:
  const ::xla::cpu::AllGatherThunkProto& _internal_all_gather_thunk() const;
  ::xla::cpu::AllGatherThunkProto* _internal_mutable_all_gather_thunk();
  public:
  void unsafe_arena_set_allocated_all_gather_thunk(
      ::xla::cpu::AllGatherThunkProto* all_gather_thunk);
  ::xla::cpu::AllGatherThunkProto* unsafe_arena_release_all_gather_thunk();

  // .xla.cpu.AllReduceThunkProto all_reduce_thunk = 5;
  bool has_all_reduce_thunk() const;
  private:
  bool _internal_has_all_reduce_thunk() const;
  public:
  void clear_all_reduce_thunk();
  const ::xla::cpu::AllReduceThunkProto& all_reduce_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::AllReduceThunkProto* release_all_reduce_thunk();
  ::xla::cpu::AllReduceThunkProto* mutable_all_reduce_thunk();
  void set_allocated_all_reduce_thunk(::xla::cpu::AllReduceThunkProto* all_reduce_thunk);
  private:
  const ::xla::cpu::AllReduceThunkProto& _internal_all_reduce_thunk() const;
  ::xla::cpu::AllReduceThunkProto* _internal_mutable_all_reduce_thunk();
  public:
  void unsafe_arena_set_allocated_all_reduce_thunk(
      ::xla::cpu::AllReduceThunkProto* all_reduce_thunk);
  ::xla::cpu::AllReduceThunkProto* unsafe_arena_release_all_reduce_thunk();

  // .xla.cpu.AllToAllThunkProto all_to_all_thunk = 6;
  bool has_all_to_all_thunk() const;
  private:
  bool _internal_has_all_to_all_thunk() const;
  public:
  void clear_all_to_all_thunk();
  const ::xla::cpu::AllToAllThunkProto& all_to_all_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::AllToAllThunkProto* release_all_to_all_thunk();
  ::xla::cpu::AllToAllThunkProto* mutable_all_to_all_thunk();
  void set_allocated_all_to_all_thunk(::xla::cpu::AllToAllThunkProto* all_to_all_thunk);
  private:
  const ::xla::cpu::AllToAllThunkProto& _internal_all_to_all_thunk() const;
  ::xla::cpu::AllToAllThunkProto* _internal_mutable_all_to_all_thunk();
  public:
  void unsafe_arena_set_allocated_all_to_all_thunk(
      ::xla::cpu::AllToAllThunkProto* all_to_all_thunk);
  ::xla::cpu::AllToAllThunkProto* unsafe_arena_release_all_to_all_thunk();

  // .xla.cpu.ReduceScatterThunkProto reduce_scatter_thunk = 7;
  bool has_reduce_scatter_thunk() const;
  private:
  bool _internal_has_reduce_scatter_thunk() const;
  public:
  void clear_reduce_scatter_thunk();
  const ::xla::cpu::ReduceScatterThunkProto& reduce_scatter_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::ReduceScatterThunkProto* release_reduce_scatter_thunk();
  ::xla::cpu::ReduceScatterThunkProto* mutable_reduce_scatter_thunk();
  void set_allocated_reduce_scatter_thunk(::xla::cpu::ReduceScatterThunkProto* reduce_scatter_thunk);
  private:
  const ::xla::cpu::ReduceScatterThunkProto& _internal_reduce_scatter_thunk() const;
  ::xla::cpu::ReduceScatterThunkProto* _internal_mutable_reduce_scatter_thunk();
  public:
  void unsafe_arena_set_allocated_reduce_scatter_thunk(
      ::xla::cpu::ReduceScatterThunkProto* reduce_scatter_thunk);
  ::xla::cpu::ReduceScatterThunkProto* unsafe_arena_release_reduce_scatter_thunk();

  // .xla.cpu.CollectivePermuteThunkProto collective_permute_thunk = 8;
  bool has_collective_permute_thunk() const;
  private:
  bool _internal_has_collective_permute_thunk() const;
  public:
  void clear_collective_permute_thunk();
  const ::xla::cpu::CollectivePermuteThunkProto& collective_permute_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::CollectivePermuteThunkProto* release_collective_permute_thunk();
  ::xla::cpu::CollectivePermuteThunkProto* mutable_collective_permute_thunk();
  void set_allocated_collective_permute_thunk(::xla::cpu::CollectivePermuteThunkProto* collective_permute_thunk);
  private:
  const ::xla::cpu::CollectivePermuteThunkProto& _internal_collective_permute_thunk() const;
  ::xla::cpu::CollectivePermuteThunkProto* _internal_mutable_collective_permute_thunk();
  public:
  void unsafe_arena_set_allocated_collective_permute_thunk(
      ::xla::cpu::CollectivePermuteThunkProto* collective_permute_thunk);
  ::xla::cpu::CollectivePermuteThunkProto* unsafe_arena_release_collective_permute_thunk();

  void clear_impl();
  ImplCase impl_case() const;
  // @@protoc_insertion_point(class_scope:xla.cpu.CollectiveThunkProto)
 private:
  class _Internal;
  void set_has_all_gather_thunk();
  void set_has_all_reduce_thunk();
  void set_has_all_to_all_thunk();
  void set_has_reduce_scatter_thunk();
  void set_has_collective_permute_thunk();

  inline bool has_impl() const;
  inline void clear_has_impl();

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::OpParamsProto* op_params_;
    ::xla::cpu::OpBuffersProto* op_buffers_;
    ::xla::cpu::OpResourcesProto* op_resources_;
    union ImplUnion {
      constexpr ImplUnion() : _constinit_{} {}
        ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_;
      ::xla::cpu::AllGatherThunkProto* all_gather_thunk_;
      ::xla::cpu::AllReduceThunkProto* all_reduce_thunk_;
      ::xla::cpu::AllToAllThunkProto* all_to_all_thunk_;
      ::xla::cpu::ReduceScatterThunkProto* reduce_scatter_thunk_;
      ::xla::cpu::CollectivePermuteThunkProto* collective_permute_thunk_;
    } impl_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
    uint32_t _oneof_case_[1];

  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CallThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CallThunkProto) */ {
 public:
  inline CallThunkProto() : CallThunkProto(nullptr) {}
  ~CallThunkProto() override;
  explicit PROTOBUF_CONSTEXPR CallThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CallThunkProto(const CallThunkProto& from);
  CallThunkProto(CallThunkProto&& from) noexcept
    : CallThunkProto() {
    *this = ::std::move(from);
  }

  inline CallThunkProto& operator=(const CallThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CallThunkProto& operator=(CallThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CallThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const CallThunkProto* internal_default_instance() {
    return reinterpret_cast<const CallThunkProto*>(
               &_CallThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    17;

  friend void swap(CallThunkProto& a, CallThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CallThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CallThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CallThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CallThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CallThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CallThunkProto& from) {
    CallThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CallThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CallThunkProto";
  }
  protected:
  explicit CallThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kCalledSequenceFieldNumber = 1,
  };
  // .xla.cpu.ThunkSequenceProto called_sequence = 1;
  bool has_called_sequence() const;
  private:
  bool _internal_has_called_sequence() const;
  public:
  void clear_called_sequence();
  const ::xla::cpu::ThunkSequenceProto& called_sequence() const;
  PROTOBUF_NODISCARD ::xla::cpu::ThunkSequenceProto* release_called_sequence();
  ::xla::cpu::ThunkSequenceProto* mutable_called_sequence();
  void set_allocated_called_sequence(::xla::cpu::ThunkSequenceProto* called_sequence);
  private:
  const ::xla::cpu::ThunkSequenceProto& _internal_called_sequence() const;
  ::xla::cpu::ThunkSequenceProto* _internal_mutable_called_sequence();
  public:
  void unsafe_arena_set_allocated_called_sequence(
      ::xla::cpu::ThunkSequenceProto* called_sequence);
  ::xla::cpu::ThunkSequenceProto* unsafe_arena_release_called_sequence();

  // @@protoc_insertion_point(class_scope:xla.cpu.CallThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ThunkSequenceProto* called_sequence_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ConditionalThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ConditionalThunkProto) */ {
 public:
  inline ConditionalThunkProto() : ConditionalThunkProto(nullptr) {}
  ~ConditionalThunkProto() override;
  explicit PROTOBUF_CONSTEXPR ConditionalThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ConditionalThunkProto(const ConditionalThunkProto& from);
  ConditionalThunkProto(ConditionalThunkProto&& from) noexcept
    : ConditionalThunkProto() {
    *this = ::std::move(from);
  }

  inline ConditionalThunkProto& operator=(const ConditionalThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ConditionalThunkProto& operator=(ConditionalThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ConditionalThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ConditionalThunkProto* internal_default_instance() {
    return reinterpret_cast<const ConditionalThunkProto*>(
               &_ConditionalThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    18;

  friend void swap(ConditionalThunkProto& a, ConditionalThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ConditionalThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ConditionalThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ConditionalThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ConditionalThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ConditionalThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ConditionalThunkProto& from) {
    ConditionalThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ConditionalThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ConditionalThunkProto";
  }
  protected:
  explicit ConditionalThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kBranchSequencesFieldNumber = 1,
    kBranchIndexBufferFieldNumber = 2,
  };
  // repeated .xla.cpu.ThunkSequenceProto branch_sequences = 1;
  int branch_sequences_size() const;
  private:
  int _internal_branch_sequences_size() const;
  public:
  void clear_branch_sequences();
  ::xla::cpu::ThunkSequenceProto* mutable_branch_sequences(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkSequenceProto >*
      mutable_branch_sequences();
  private:
  const ::xla::cpu::ThunkSequenceProto& _internal_branch_sequences(int index) const;
  ::xla::cpu::ThunkSequenceProto* _internal_add_branch_sequences();
  public:
  const ::xla::cpu::ThunkSequenceProto& branch_sequences(int index) const;
  ::xla::cpu::ThunkSequenceProto* add_branch_sequences();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkSequenceProto >&
      branch_sequences() const;

  // .xla.cpu.BufferAllocationSliceProto branch_index_buffer = 2;
  bool has_branch_index_buffer() const;
  private:
  bool _internal_has_branch_index_buffer() const;
  public:
  void clear_branch_index_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& branch_index_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_branch_index_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_branch_index_buffer();
  void set_allocated_branch_index_buffer(::xla::cpu::BufferAllocationSliceProto* branch_index_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_branch_index_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_branch_index_buffer();
  public:
  void unsafe_arena_set_allocated_branch_index_buffer(
      ::xla::cpu::BufferAllocationSliceProto* branch_index_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_branch_index_buffer();

  // @@protoc_insertion_point(class_scope:xla.cpu.ConditionalThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkSequenceProto > branch_sequences_;
    ::xla::cpu::BufferAllocationSliceProto* branch_index_buffer_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ConvolutionThunkProto_Options final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ConvolutionThunkProto.Options) */ {
 public:
  inline ConvolutionThunkProto_Options() : ConvolutionThunkProto_Options(nullptr) {}
  ~ConvolutionThunkProto_Options() override;
  explicit PROTOBUF_CONSTEXPR ConvolutionThunkProto_Options(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ConvolutionThunkProto_Options(const ConvolutionThunkProto_Options& from);
  ConvolutionThunkProto_Options(ConvolutionThunkProto_Options&& from) noexcept
    : ConvolutionThunkProto_Options() {
    *this = ::std::move(from);
  }

  inline ConvolutionThunkProto_Options& operator=(const ConvolutionThunkProto_Options& from) {
    CopyFrom(from);
    return *this;
  }
  inline ConvolutionThunkProto_Options& operator=(ConvolutionThunkProto_Options&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ConvolutionThunkProto_Options& default_instance() {
    return *internal_default_instance();
  }
  static inline const ConvolutionThunkProto_Options* internal_default_instance() {
    return reinterpret_cast<const ConvolutionThunkProto_Options*>(
               &_ConvolutionThunkProto_Options_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    19;

  friend void swap(ConvolutionThunkProto_Options& a, ConvolutionThunkProto_Options& b) {
    a.Swap(&b);
  }
  inline void Swap(ConvolutionThunkProto_Options* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ConvolutionThunkProto_Options* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ConvolutionThunkProto_Options* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ConvolutionThunkProto_Options>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ConvolutionThunkProto_Options& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ConvolutionThunkProto_Options& from) {
    ConvolutionThunkProto_Options::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ConvolutionThunkProto_Options* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ConvolutionThunkProto.Options";
  }
  protected:
  explicit ConvolutionThunkProto_Options(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kMultiThreadedFieldNumber = 1,
  };
  // bool multi_threaded = 1;
  void clear_multi_threaded();
  bool multi_threaded() const;
  void set_multi_threaded(bool value);
  private:
  bool _internal_multi_threaded() const;
  void _internal_set_multi_threaded(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.ConvolutionThunkProto.Options)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    bool multi_threaded_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ConvolutionThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ConvolutionThunkProto) */ {
 public:
  inline ConvolutionThunkProto() : ConvolutionThunkProto(nullptr) {}
  ~ConvolutionThunkProto() override;
  explicit PROTOBUF_CONSTEXPR ConvolutionThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ConvolutionThunkProto(const ConvolutionThunkProto& from);
  ConvolutionThunkProto(ConvolutionThunkProto&& from) noexcept
    : ConvolutionThunkProto() {
    *this = ::std::move(from);
  }

  inline ConvolutionThunkProto& operator=(const ConvolutionThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ConvolutionThunkProto& operator=(ConvolutionThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ConvolutionThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ConvolutionThunkProto* internal_default_instance() {
    return reinterpret_cast<const ConvolutionThunkProto*>(
               &_ConvolutionThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    20;

  friend void swap(ConvolutionThunkProto& a, ConvolutionThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ConvolutionThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ConvolutionThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ConvolutionThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ConvolutionThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ConvolutionThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ConvolutionThunkProto& from) {
    ConvolutionThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ConvolutionThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ConvolutionThunkProto";
  }
  protected:
  explicit ConvolutionThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef ConvolutionThunkProto_Options Options;

  // accessors -------------------------------------------------------

  enum : int {
    kDimensionNumbersFieldNumber = 1,
    kWindowFieldNumber = 2,
    kInputBufferShapeFieldNumber = 4,
    kKernelBufferShapeFieldNumber = 5,
    kOutputBufferShapeFieldNumber = 6,
    kOptionsFieldNumber = 7,
    kFeatureGroupCountFieldNumber = 3,
  };
  // .xla.ConvolutionDimensionNumbers dimension_numbers = 1;
  bool has_dimension_numbers() const;
  private:
  bool _internal_has_dimension_numbers() const;
  public:
  void clear_dimension_numbers();
  const ::xla::ConvolutionDimensionNumbers& dimension_numbers() const;
  PROTOBUF_NODISCARD ::xla::ConvolutionDimensionNumbers* release_dimension_numbers();
  ::xla::ConvolutionDimensionNumbers* mutable_dimension_numbers();
  void set_allocated_dimension_numbers(::xla::ConvolutionDimensionNumbers* dimension_numbers);
  private:
  const ::xla::ConvolutionDimensionNumbers& _internal_dimension_numbers() const;
  ::xla::ConvolutionDimensionNumbers* _internal_mutable_dimension_numbers();
  public:
  void unsafe_arena_set_allocated_dimension_numbers(
      ::xla::ConvolutionDimensionNumbers* dimension_numbers);
  ::xla::ConvolutionDimensionNumbers* unsafe_arena_release_dimension_numbers();

  // .xla.Window window = 2;
  bool has_window() const;
  private:
  bool _internal_has_window() const;
  public:
  void clear_window();
  const ::xla::Window& window() const;
  PROTOBUF_NODISCARD ::xla::Window* release_window();
  ::xla::Window* mutable_window();
  void set_allocated_window(::xla::Window* window);
  private:
  const ::xla::Window& _internal_window() const;
  ::xla::Window* _internal_mutable_window();
  public:
  void unsafe_arena_set_allocated_window(
      ::xla::Window* window);
  ::xla::Window* unsafe_arena_release_window();

  // .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
  bool has_input_buffer_shape() const;
  private:
  bool _internal_has_input_buffer_shape() const;
  public:
  void clear_input_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& input_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_input_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_input_buffer_shape();
  void set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_input_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_input_buffer_shape();
  public:
  void unsafe_arena_set_allocated_input_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_input_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto kernel_buffer_shape = 5;
  bool has_kernel_buffer_shape() const;
  private:
  bool _internal_has_kernel_buffer_shape() const;
  public:
  void clear_kernel_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& kernel_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_kernel_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_kernel_buffer_shape();
  void set_allocated_kernel_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_kernel_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_kernel_buffer_shape();
  public:
  void unsafe_arena_set_allocated_kernel_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_kernel_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 6;
  bool has_output_buffer_shape() const;
  private:
  bool _internal_has_output_buffer_shape() const;
  public:
  void clear_output_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& output_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_output_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_output_buffer_shape();
  void set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_output_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_output_buffer_shape();
  public:
  void unsafe_arena_set_allocated_output_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_output_buffer_shape();

  // .xla.cpu.ConvolutionThunkProto.Options options = 7;
  bool has_options() const;
  private:
  bool _internal_has_options() const;
  public:
  void clear_options();
  const ::xla::cpu::ConvolutionThunkProto_Options& options() const;
  PROTOBUF_NODISCARD ::xla::cpu::ConvolutionThunkProto_Options* release_options();
  ::xla::cpu::ConvolutionThunkProto_Options* mutable_options();
  void set_allocated_options(::xla::cpu::ConvolutionThunkProto_Options* options);
  private:
  const ::xla::cpu::ConvolutionThunkProto_Options& _internal_options() const;
  ::xla::cpu::ConvolutionThunkProto_Options* _internal_mutable_options();
  public:
  void unsafe_arena_set_allocated_options(
      ::xla::cpu::ConvolutionThunkProto_Options* options);
  ::xla::cpu::ConvolutionThunkProto_Options* unsafe_arena_release_options();

  // int64 feature_group_count = 3;
  void clear_feature_group_count();
  int64_t feature_group_count() const;
  void set_feature_group_count(int64_t value);
  private:
  int64_t _internal_feature_group_count() const;
  void _internal_set_feature_group_count(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.ConvolutionThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::ConvolutionDimensionNumbers* dimension_numbers_;
    ::xla::Window* window_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape_;
    ::xla::cpu::ConvolutionThunkProto_Options* options_;
    int64_t feature_group_count_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class SortThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.SortThunkProto) */ {
 public:
  inline SortThunkProto() : SortThunkProto(nullptr) {}
  ~SortThunkProto() override;
  explicit PROTOBUF_CONSTEXPR SortThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  SortThunkProto(const SortThunkProto& from);
  SortThunkProto(SortThunkProto&& from) noexcept
    : SortThunkProto() {
    *this = ::std::move(from);
  }

  inline SortThunkProto& operator=(const SortThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline SortThunkProto& operator=(SortThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const SortThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const SortThunkProto* internal_default_instance() {
    return reinterpret_cast<const SortThunkProto*>(
               &_SortThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    21;

  friend void swap(SortThunkProto& a, SortThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(SortThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(SortThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  SortThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<SortThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const SortThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const SortThunkProto& from) {
    SortThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(SortThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.SortThunkProto";
  }
  protected:
  explicit SortThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kInputsShapesFieldNumber = 5,
    kComparatorNameFieldNumber = 4,
    kDirectionFieldNumber = 3,
    kDimensionFieldNumber = 1,
    kIsStableFieldNumber = 2,
  };
  // repeated .xla.cpu.ShapeBufferAllocationSliceProto inputs_shapes = 5;
  int inputs_shapes_size() const;
  private:
  int _internal_inputs_shapes_size() const;
  public:
  void clear_inputs_shapes();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_inputs_shapes(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_inputs_shapes();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_inputs_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_inputs_shapes();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& inputs_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_inputs_shapes();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      inputs_shapes() const;

  // string comparator_name = 4;
  void clear_comparator_name();
  const std::string& comparator_name() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_comparator_name(ArgT0&& arg0, ArgT... args);
  std::string* mutable_comparator_name();
  PROTOBUF_NODISCARD std::string* release_comparator_name();
  void set_allocated_comparator_name(std::string* comparator_name);
  private:
  const std::string& _internal_comparator_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_comparator_name(const std::string& value);
  std::string* _internal_mutable_comparator_name();
  public:

  // .xla.cpu.SortDirectionOptional direction = 3;
  bool has_direction() const;
  private:
  bool _internal_has_direction() const;
  public:
  void clear_direction();
  const ::xla::cpu::SortDirectionOptional& direction() const;
  PROTOBUF_NODISCARD ::xla::cpu::SortDirectionOptional* release_direction();
  ::xla::cpu::SortDirectionOptional* mutable_direction();
  void set_allocated_direction(::xla::cpu::SortDirectionOptional* direction);
  private:
  const ::xla::cpu::SortDirectionOptional& _internal_direction() const;
  ::xla::cpu::SortDirectionOptional* _internal_mutable_direction();
  public:
  void unsafe_arena_set_allocated_direction(
      ::xla::cpu::SortDirectionOptional* direction);
  ::xla::cpu::SortDirectionOptional* unsafe_arena_release_direction();

  // int64 dimension = 1;
  void clear_dimension();
  int64_t dimension() const;
  void set_dimension(int64_t value);
  private:
  int64_t _internal_dimension() const;
  void _internal_set_dimension(int64_t value);
  public:

  // bool is_stable = 2;
  void clear_is_stable();
  bool is_stable() const;
  void set_is_stable(bool value);
  private:
  bool _internal_is_stable() const;
  void _internal_set_is_stable(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.SortThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > inputs_shapes_;
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr comparator_name_;
    ::xla::cpu::SortDirectionOptional* direction_;
    int64_t dimension_;
    bool is_stable_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class XnnDotThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.XnnDotThunkProto) */ {
 public:
  inline XnnDotThunkProto() : XnnDotThunkProto(nullptr) {}
  ~XnnDotThunkProto() override;
  explicit PROTOBUF_CONSTEXPR XnnDotThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  XnnDotThunkProto(const XnnDotThunkProto& from);
  XnnDotThunkProto(XnnDotThunkProto&& from) noexcept
    : XnnDotThunkProto() {
    *this = ::std::move(from);
  }

  inline XnnDotThunkProto& operator=(const XnnDotThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline XnnDotThunkProto& operator=(XnnDotThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const XnnDotThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const XnnDotThunkProto* internal_default_instance() {
    return reinterpret_cast<const XnnDotThunkProto*>(
               &_XnnDotThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    22;

  friend void swap(XnnDotThunkProto& a, XnnDotThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(XnnDotThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(XnnDotThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  XnnDotThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<XnnDotThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const XnnDotThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const XnnDotThunkProto& from) {
    XnnDotThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(XnnDotThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.XnnDotThunkProto";
  }
  protected:
  explicit XnnDotThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kDotDimensionsFieldNumber = 1,
    kLhsBufferShapeFieldNumber = 2,
    kRhsBufferShapeFieldNumber = 3,
    kOutBufferShapeFieldNumber = 4,
  };
  // .xla.DotDimensionNumbers dot_dimensions = 1;
  bool has_dot_dimensions() const;
  private:
  bool _internal_has_dot_dimensions() const;
  public:
  void clear_dot_dimensions();
  const ::xla::DotDimensionNumbers& dot_dimensions() const;
  PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_dot_dimensions();
  ::xla::DotDimensionNumbers* mutable_dot_dimensions();
  void set_allocated_dot_dimensions(::xla::DotDimensionNumbers* dot_dimensions);
  private:
  const ::xla::DotDimensionNumbers& _internal_dot_dimensions() const;
  ::xla::DotDimensionNumbers* _internal_mutable_dot_dimensions();
  public:
  void unsafe_arena_set_allocated_dot_dimensions(
      ::xla::DotDimensionNumbers* dot_dimensions);
  ::xla::DotDimensionNumbers* unsafe_arena_release_dot_dimensions();

  // .xla.cpu.ShapeBufferAllocationSliceProto lhs_buffer_shape = 2;
  bool has_lhs_buffer_shape() const;
  private:
  bool _internal_has_lhs_buffer_shape() const;
  public:
  void clear_lhs_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& lhs_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_lhs_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_lhs_buffer_shape();
  void set_allocated_lhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_lhs_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_lhs_buffer_shape();
  public:
  void unsafe_arena_set_allocated_lhs_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_lhs_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto rhs_buffer_shape = 3;
  bool has_rhs_buffer_shape() const;
  private:
  bool _internal_has_rhs_buffer_shape() const;
  public:
  void clear_rhs_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& rhs_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_rhs_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_rhs_buffer_shape();
  void set_allocated_rhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_rhs_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_rhs_buffer_shape();
  public:
  void unsafe_arena_set_allocated_rhs_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_rhs_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto out_buffer_shape = 4;
  bool has_out_buffer_shape() const;
  private:
  bool _internal_has_out_buffer_shape() const;
  public:
  void clear_out_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& out_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_out_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_out_buffer_shape();
  void set_allocated_out_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_out_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_out_buffer_shape();
  public:
  void unsafe_arena_set_allocated_out_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_out_buffer_shape();

  // @@protoc_insertion_point(class_scope:xla.cpu.XnnDotThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::DotDimensionNumbers* dot_dimensions_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class XnnConvolutionThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.XnnConvolutionThunkProto) */ {
 public:
  inline XnnConvolutionThunkProto() : XnnConvolutionThunkProto(nullptr) {}
  ~XnnConvolutionThunkProto() override;
  explicit PROTOBUF_CONSTEXPR XnnConvolutionThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  XnnConvolutionThunkProto(const XnnConvolutionThunkProto& from);
  XnnConvolutionThunkProto(XnnConvolutionThunkProto&& from) noexcept
    : XnnConvolutionThunkProto() {
    *this = ::std::move(from);
  }

  inline XnnConvolutionThunkProto& operator=(const XnnConvolutionThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline XnnConvolutionThunkProto& operator=(XnnConvolutionThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const XnnConvolutionThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const XnnConvolutionThunkProto* internal_default_instance() {
    return reinterpret_cast<const XnnConvolutionThunkProto*>(
               &_XnnConvolutionThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    23;

  friend void swap(XnnConvolutionThunkProto& a, XnnConvolutionThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(XnnConvolutionThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(XnnConvolutionThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  XnnConvolutionThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<XnnConvolutionThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const XnnConvolutionThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const XnnConvolutionThunkProto& from) {
    XnnConvolutionThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(XnnConvolutionThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.XnnConvolutionThunkProto";
  }
  protected:
  explicit XnnConvolutionThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kDimensionNumbersFieldNumber = 1,
    kWindowFieldNumber = 2,
    kInputBufferShapeFieldNumber = 4,
    kKernelBufferShapeFieldNumber = 5,
    kOutputBufferShapeFieldNumber = 6,
    kFeatureGroupCountFieldNumber = 3,
  };
  // .xla.ConvolutionDimensionNumbers dimension_numbers = 1;
  bool has_dimension_numbers() const;
  private:
  bool _internal_has_dimension_numbers() const;
  public:
  void clear_dimension_numbers();
  const ::xla::ConvolutionDimensionNumbers& dimension_numbers() const;
  PROTOBUF_NODISCARD ::xla::ConvolutionDimensionNumbers* release_dimension_numbers();
  ::xla::ConvolutionDimensionNumbers* mutable_dimension_numbers();
  void set_allocated_dimension_numbers(::xla::ConvolutionDimensionNumbers* dimension_numbers);
  private:
  const ::xla::ConvolutionDimensionNumbers& _internal_dimension_numbers() const;
  ::xla::ConvolutionDimensionNumbers* _internal_mutable_dimension_numbers();
  public:
  void unsafe_arena_set_allocated_dimension_numbers(
      ::xla::ConvolutionDimensionNumbers* dimension_numbers);
  ::xla::ConvolutionDimensionNumbers* unsafe_arena_release_dimension_numbers();

  // .xla.Window window = 2;
  bool has_window() const;
  private:
  bool _internal_has_window() const;
  public:
  void clear_window();
  const ::xla::Window& window() const;
  PROTOBUF_NODISCARD ::xla::Window* release_window();
  ::xla::Window* mutable_window();
  void set_allocated_window(::xla::Window* window);
  private:
  const ::xla::Window& _internal_window() const;
  ::xla::Window* _internal_mutable_window();
  public:
  void unsafe_arena_set_allocated_window(
      ::xla::Window* window);
  ::xla::Window* unsafe_arena_release_window();

  // .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
  bool has_input_buffer_shape() const;
  private:
  bool _internal_has_input_buffer_shape() const;
  public:
  void clear_input_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& input_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_input_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_input_buffer_shape();
  void set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_input_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_input_buffer_shape();
  public:
  void unsafe_arena_set_allocated_input_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_input_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto kernel_buffer_shape = 5;
  bool has_kernel_buffer_shape() const;
  private:
  bool _internal_has_kernel_buffer_shape() const;
  public:
  void clear_kernel_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& kernel_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_kernel_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_kernel_buffer_shape();
  void set_allocated_kernel_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_kernel_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_kernel_buffer_shape();
  public:
  void unsafe_arena_set_allocated_kernel_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_kernel_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 6;
  bool has_output_buffer_shape() const;
  private:
  bool _internal_has_output_buffer_shape() const;
  public:
  void clear_output_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& output_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_output_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_output_buffer_shape();
  void set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_output_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_output_buffer_shape();
  public:
  void unsafe_arena_set_allocated_output_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_output_buffer_shape();

  // int64 feature_group_count = 3;
  void clear_feature_group_count();
  int64_t feature_group_count() const;
  void set_feature_group_count(int64_t value);
  private:
  int64_t _internal_feature_group_count() const;
  void _internal_set_feature_group_count(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.XnnConvolutionThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::ConvolutionDimensionNumbers* dimension_numbers_;
    ::xla::Window* window_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape_;
    int64_t feature_group_count_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class XnnFusionThunkProto_Options final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.XnnFusionThunkProto.Options) */ {
 public:
  inline XnnFusionThunkProto_Options() : XnnFusionThunkProto_Options(nullptr) {}
  ~XnnFusionThunkProto_Options() override;
  explicit PROTOBUF_CONSTEXPR XnnFusionThunkProto_Options(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  XnnFusionThunkProto_Options(const XnnFusionThunkProto_Options& from);
  XnnFusionThunkProto_Options(XnnFusionThunkProto_Options&& from) noexcept
    : XnnFusionThunkProto_Options() {
    *this = ::std::move(from);
  }

  inline XnnFusionThunkProto_Options& operator=(const XnnFusionThunkProto_Options& from) {
    CopyFrom(from);
    return *this;
  }
  inline XnnFusionThunkProto_Options& operator=(XnnFusionThunkProto_Options&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const XnnFusionThunkProto_Options& default_instance() {
    return *internal_default_instance();
  }
  static inline const XnnFusionThunkProto_Options* internal_default_instance() {
    return reinterpret_cast<const XnnFusionThunkProto_Options*>(
               &_XnnFusionThunkProto_Options_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    24;

  friend void swap(XnnFusionThunkProto_Options& a, XnnFusionThunkProto_Options& b) {
    a.Swap(&b);
  }
  inline void Swap(XnnFusionThunkProto_Options* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(XnnFusionThunkProto_Options* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  XnnFusionThunkProto_Options* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<XnnFusionThunkProto_Options>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const XnnFusionThunkProto_Options& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const XnnFusionThunkProto_Options& from) {
    XnnFusionThunkProto_Options::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(XnnFusionThunkProto_Options* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.XnnFusionThunkProto.Options";
  }
  protected:
  explicit XnnFusionThunkProto_Options(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kUseThreadpoolFieldNumber = 1,
  };
  // bool use_threadpool = 1;
  void clear_use_threadpool();
  bool use_threadpool() const;
  void set_use_threadpool(bool value);
  private:
  bool _internal_use_threadpool() const;
  void _internal_set_use_threadpool(bool value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.XnnFusionThunkProto.Options)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    bool use_threadpool_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class XnnFusionThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.XnnFusionThunkProto) */ {
 public:
  inline XnnFusionThunkProto() : XnnFusionThunkProto(nullptr) {}
  ~XnnFusionThunkProto() override;
  explicit PROTOBUF_CONSTEXPR XnnFusionThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  XnnFusionThunkProto(const XnnFusionThunkProto& from);
  XnnFusionThunkProto(XnnFusionThunkProto&& from) noexcept
    : XnnFusionThunkProto() {
    *this = ::std::move(from);
  }

  inline XnnFusionThunkProto& operator=(const XnnFusionThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline XnnFusionThunkProto& operator=(XnnFusionThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const XnnFusionThunkProto& default_instance() {
    return *internal_default_instance();
  }
  enum ImplCase {
    kXnnDotThunk = 2,
    kXnnConvolutionThunk = 3,
    IMPL_NOT_SET = 0,
  };

  static inline const XnnFusionThunkProto* internal_default_instance() {
    return reinterpret_cast<const XnnFusionThunkProto*>(
               &_XnnFusionThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    25;

  friend void swap(XnnFusionThunkProto& a, XnnFusionThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(XnnFusionThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(XnnFusionThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  XnnFusionThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<XnnFusionThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const XnnFusionThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const XnnFusionThunkProto& from) {
    XnnFusionThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(XnnFusionThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.XnnFusionThunkProto";
  }
  protected:
  explicit XnnFusionThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef XnnFusionThunkProto_Options Options;

  // accessors -------------------------------------------------------

  enum : int {
    kOptionsFieldNumber = 1,
    kXnnDotThunkFieldNumber = 2,
    kXnnConvolutionThunkFieldNumber = 3,
  };
  // .xla.cpu.XnnFusionThunkProto.Options options = 1;
  bool has_options() const;
  private:
  bool _internal_has_options() const;
  public:
  void clear_options();
  const ::xla::cpu::XnnFusionThunkProto_Options& options() const;
  PROTOBUF_NODISCARD ::xla::cpu::XnnFusionThunkProto_Options* release_options();
  ::xla::cpu::XnnFusionThunkProto_Options* mutable_options();
  void set_allocated_options(::xla::cpu::XnnFusionThunkProto_Options* options);
  private:
  const ::xla::cpu::XnnFusionThunkProto_Options& _internal_options() const;
  ::xla::cpu::XnnFusionThunkProto_Options* _internal_mutable_options();
  public:
  void unsafe_arena_set_allocated_options(
      ::xla::cpu::XnnFusionThunkProto_Options* options);
  ::xla::cpu::XnnFusionThunkProto_Options* unsafe_arena_release_options();

  // .xla.cpu.XnnDotThunkProto xnn_dot_thunk = 2;
  bool has_xnn_dot_thunk() const;
  private:
  bool _internal_has_xnn_dot_thunk() const;
  public:
  void clear_xnn_dot_thunk();
  const ::xla::cpu::XnnDotThunkProto& xnn_dot_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::XnnDotThunkProto* release_xnn_dot_thunk();
  ::xla::cpu::XnnDotThunkProto* mutable_xnn_dot_thunk();
  void set_allocated_xnn_dot_thunk(::xla::cpu::XnnDotThunkProto* xnn_dot_thunk);
  private:
  const ::xla::cpu::XnnDotThunkProto& _internal_xnn_dot_thunk() const;
  ::xla::cpu::XnnDotThunkProto* _internal_mutable_xnn_dot_thunk();
  public:
  void unsafe_arena_set_allocated_xnn_dot_thunk(
      ::xla::cpu::XnnDotThunkProto* xnn_dot_thunk);
  ::xla::cpu::XnnDotThunkProto* unsafe_arena_release_xnn_dot_thunk();

  // .xla.cpu.XnnConvolutionThunkProto xnn_convolution_thunk = 3;
  bool has_xnn_convolution_thunk() const;
  private:
  bool _internal_has_xnn_convolution_thunk() const;
  public:
  void clear_xnn_convolution_thunk();
  const ::xla::cpu::XnnConvolutionThunkProto& xnn_convolution_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::XnnConvolutionThunkProto* release_xnn_convolution_thunk();
  ::xla::cpu::XnnConvolutionThunkProto* mutable_xnn_convolution_thunk();
  void set_allocated_xnn_convolution_thunk(::xla::cpu::XnnConvolutionThunkProto* xnn_convolution_thunk);
  private:
  const ::xla::cpu::XnnConvolutionThunkProto& _internal_xnn_convolution_thunk() const;
  ::xla::cpu::XnnConvolutionThunkProto* _internal_mutable_xnn_convolution_thunk();
  public:
  void unsafe_arena_set_allocated_xnn_convolution_thunk(
      ::xla::cpu::XnnConvolutionThunkProto* xnn_convolution_thunk);
  ::xla::cpu::XnnConvolutionThunkProto* unsafe_arena_release_xnn_convolution_thunk();

  void clear_impl();
  ImplCase impl_case() const;
  // @@protoc_insertion_point(class_scope:xla.cpu.XnnFusionThunkProto)
 private:
  class _Internal;
  void set_has_xnn_dot_thunk();
  void set_has_xnn_convolution_thunk();

  inline bool has_impl() const;
  inline void clear_has_impl();

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::XnnFusionThunkProto_Options* options_;
    union ImplUnion {
      constexpr ImplUnion() : _constinit_{} {}
        ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_;
      ::xla::cpu::XnnDotThunkProto* xnn_dot_thunk_;
      ::xla::cpu::XnnConvolutionThunkProto* xnn_convolution_thunk_;
    } impl_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
    uint32_t _oneof_case_[1];

  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class DotThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.DotThunkProto) */ {
 public:
  inline DotThunkProto() : DotThunkProto(nullptr) {}
  ~DotThunkProto() override;
  explicit PROTOBUF_CONSTEXPR DotThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  DotThunkProto(const DotThunkProto& from);
  DotThunkProto(DotThunkProto&& from) noexcept
    : DotThunkProto() {
    *this = ::std::move(from);
  }

  inline DotThunkProto& operator=(const DotThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline DotThunkProto& operator=(DotThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const DotThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const DotThunkProto* internal_default_instance() {
    return reinterpret_cast<const DotThunkProto*>(
               &_DotThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    26;

  friend void swap(DotThunkProto& a, DotThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(DotThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(DotThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  DotThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<DotThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const DotThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const DotThunkProto& from) {
    DotThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(DotThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.DotThunkProto";
  }
  protected:
  explicit DotThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kDotDimensionsFieldNumber = 1,
    kLhsBufferShapeFieldNumber = 2,
    kRhsBufferShapeFieldNumber = 3,
    kOutBufferShapeFieldNumber = 4,
  };
  // .xla.DotDimensionNumbers dot_dimensions = 1;
  bool has_dot_dimensions() const;
  private:
  bool _internal_has_dot_dimensions() const;
  public:
  void clear_dot_dimensions();
  const ::xla::DotDimensionNumbers& dot_dimensions() const;
  PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_dot_dimensions();
  ::xla::DotDimensionNumbers* mutable_dot_dimensions();
  void set_allocated_dot_dimensions(::xla::DotDimensionNumbers* dot_dimensions);
  private:
  const ::xla::DotDimensionNumbers& _internal_dot_dimensions() const;
  ::xla::DotDimensionNumbers* _internal_mutable_dot_dimensions();
  public:
  void unsafe_arena_set_allocated_dot_dimensions(
      ::xla::DotDimensionNumbers* dot_dimensions);
  ::xla::DotDimensionNumbers* unsafe_arena_release_dot_dimensions();

  // .xla.cpu.ShapeBufferAllocationSliceProto lhs_buffer_shape = 2;
  bool has_lhs_buffer_shape() const;
  private:
  bool _internal_has_lhs_buffer_shape() const;
  public:
  void clear_lhs_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& lhs_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_lhs_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_lhs_buffer_shape();
  void set_allocated_lhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_lhs_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_lhs_buffer_shape();
  public:
  void unsafe_arena_set_allocated_lhs_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_lhs_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto rhs_buffer_shape = 3;
  bool has_rhs_buffer_shape() const;
  private:
  bool _internal_has_rhs_buffer_shape() const;
  public:
  void clear_rhs_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& rhs_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_rhs_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_rhs_buffer_shape();
  void set_allocated_rhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_rhs_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_rhs_buffer_shape();
  public:
  void unsafe_arena_set_allocated_rhs_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_rhs_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto out_buffer_shape = 4;
  bool has_out_buffer_shape() const;
  private:
  bool _internal_has_out_buffer_shape() const;
  public:
  void clear_out_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& out_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_out_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_out_buffer_shape();
  void set_allocated_out_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_out_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_out_buffer_shape();
  public:
  void unsafe_arena_set_allocated_out_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_out_buffer_shape();

  // @@protoc_insertion_point(class_scope:xla.cpu.DotThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::DotDimensionNumbers* dot_dimensions_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class RngGetAndUpdateStateThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.RngGetAndUpdateStateThunkProto) */ {
 public:
  inline RngGetAndUpdateStateThunkProto() : RngGetAndUpdateStateThunkProto(nullptr) {}
  ~RngGetAndUpdateStateThunkProto() override;
  explicit PROTOBUF_CONSTEXPR RngGetAndUpdateStateThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  RngGetAndUpdateStateThunkProto(const RngGetAndUpdateStateThunkProto& from);
  RngGetAndUpdateStateThunkProto(RngGetAndUpdateStateThunkProto&& from) noexcept
    : RngGetAndUpdateStateThunkProto() {
    *this = ::std::move(from);
  }

  inline RngGetAndUpdateStateThunkProto& operator=(const RngGetAndUpdateStateThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline RngGetAndUpdateStateThunkProto& operator=(RngGetAndUpdateStateThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const RngGetAndUpdateStateThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const RngGetAndUpdateStateThunkProto* internal_default_instance() {
    return reinterpret_cast<const RngGetAndUpdateStateThunkProto*>(
               &_RngGetAndUpdateStateThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    27;

  friend void swap(RngGetAndUpdateStateThunkProto& a, RngGetAndUpdateStateThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(RngGetAndUpdateStateThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(RngGetAndUpdateStateThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  RngGetAndUpdateStateThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<RngGetAndUpdateStateThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const RngGetAndUpdateStateThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const RngGetAndUpdateStateThunkProto& from) {
    RngGetAndUpdateStateThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(RngGetAndUpdateStateThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.RngGetAndUpdateStateThunkProto";
  }
  protected:
  explicit RngGetAndUpdateStateThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kStateBufferFieldNumber = 2,
    kDeltaFieldNumber = 1,
  };
  // .xla.cpu.BufferAllocationSliceProto state_buffer = 2;
  bool has_state_buffer() const;
  private:
  bool _internal_has_state_buffer() const;
  public:
  void clear_state_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& state_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_state_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_state_buffer();
  void set_allocated_state_buffer(::xla::cpu::BufferAllocationSliceProto* state_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_state_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_state_buffer();
  public:
  void unsafe_arena_set_allocated_state_buffer(
      ::xla::cpu::BufferAllocationSliceProto* state_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_state_buffer();

  // int64 delta = 1;
  void clear_delta();
  int64_t delta() const;
  void set_delta(int64_t value);
  private:
  int64_t _internal_delta() const;
  void _internal_set_delta(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.RngGetAndUpdateStateThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::BufferAllocationSliceProto* state_buffer_;
    int64_t delta_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class TopKThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.TopKThunkProto) */ {
 public:
  inline TopKThunkProto() : TopKThunkProto(nullptr) {}
  ~TopKThunkProto() override;
  explicit PROTOBUF_CONSTEXPR TopKThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  TopKThunkProto(const TopKThunkProto& from);
  TopKThunkProto(TopKThunkProto&& from) noexcept
    : TopKThunkProto() {
    *this = ::std::move(from);
  }

  inline TopKThunkProto& operator=(const TopKThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline TopKThunkProto& operator=(TopKThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const TopKThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const TopKThunkProto* internal_default_instance() {
    return reinterpret_cast<const TopKThunkProto*>(
               &_TopKThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    28;

  friend void swap(TopKThunkProto& a, TopKThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(TopKThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(TopKThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  TopKThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<TopKThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const TopKThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const TopKThunkProto& from) {
    TopKThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(TopKThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.TopKThunkProto";
  }
  protected:
  explicit TopKThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kValuesBufferFieldNumber = 4,
    kOutputBufferFieldNumber = 5,
    kIndicesBufferFieldNumber = 6,
    kBatchSizeFieldNumber = 1,
    kInputSizeFieldNumber = 2,
    kKFieldNumber = 3,
  };
  // .xla.cpu.BufferAllocationSliceProto values_buffer = 4;
  bool has_values_buffer() const;
  private:
  bool _internal_has_values_buffer() const;
  public:
  void clear_values_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& values_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_values_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_values_buffer();
  void set_allocated_values_buffer(::xla::cpu::BufferAllocationSliceProto* values_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_values_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_values_buffer();
  public:
  void unsafe_arena_set_allocated_values_buffer(
      ::xla::cpu::BufferAllocationSliceProto* values_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_values_buffer();

  // .xla.cpu.BufferAllocationSliceProto output_buffer = 5;
  bool has_output_buffer() const;
  private:
  bool _internal_has_output_buffer() const;
  public:
  void clear_output_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& output_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_output_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_output_buffer();
  void set_allocated_output_buffer(::xla::cpu::BufferAllocationSliceProto* output_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_output_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_output_buffer();
  public:
  void unsafe_arena_set_allocated_output_buffer(
      ::xla::cpu::BufferAllocationSliceProto* output_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_output_buffer();

  // .xla.cpu.BufferAllocationSliceProto indices_buffer = 6;
  bool has_indices_buffer() const;
  private:
  bool _internal_has_indices_buffer() const;
  public:
  void clear_indices_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& indices_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_indices_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_indices_buffer();
  void set_allocated_indices_buffer(::xla::cpu::BufferAllocationSliceProto* indices_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_indices_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_indices_buffer();
  public:
  void unsafe_arena_set_allocated_indices_buffer(
      ::xla::cpu::BufferAllocationSliceProto* indices_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_indices_buffer();

  // int64 batch_size = 1;
  void clear_batch_size();
  int64_t batch_size() const;
  void set_batch_size(int64_t value);
  private:
  int64_t _internal_batch_size() const;
  void _internal_set_batch_size(int64_t value);
  public:

  // int64 input_size = 2;
  void clear_input_size();
  int64_t input_size() const;
  void set_input_size(int64_t value);
  private:
  int64_t _internal_input_size() const;
  void _internal_set_input_size(int64_t value);
  public:

  // int64 k = 3;
  void clear_k();
  int64_t k() const;
  void set_k(int64_t value);
  private:
  int64_t _internal_k() const;
  void _internal_set_k(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.TopKThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::BufferAllocationSliceProto* values_buffer_;
    ::xla::cpu::BufferAllocationSliceProto* output_buffer_;
    ::xla::cpu::BufferAllocationSliceProto* indices_buffer_;
    int64_t batch_size_;
    int64_t input_size_;
    int64_t k_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class WhileThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.WhileThunkProto) */ {
 public:
  inline WhileThunkProto() : WhileThunkProto(nullptr) {}
  ~WhileThunkProto() override;
  explicit PROTOBUF_CONSTEXPR WhileThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  WhileThunkProto(const WhileThunkProto& from);
  WhileThunkProto(WhileThunkProto&& from) noexcept
    : WhileThunkProto() {
    *this = ::std::move(from);
  }

  inline WhileThunkProto& operator=(const WhileThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline WhileThunkProto& operator=(WhileThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const WhileThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const WhileThunkProto* internal_default_instance() {
    return reinterpret_cast<const WhileThunkProto*>(
               &_WhileThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    29;

  friend void swap(WhileThunkProto& a, WhileThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(WhileThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(WhileThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  WhileThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<WhileThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const WhileThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const WhileThunkProto& from) {
    WhileThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(WhileThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.WhileThunkProto";
  }
  protected:
  explicit WhileThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kCondSequenceFieldNumber = 1,
    kBodySequenceFieldNumber = 2,
    kTripCountFieldNumber = 3,
    kCondBufferFieldNumber = 4,
  };
  // .xla.cpu.ThunkSequenceProto cond_sequence = 1;
  bool has_cond_sequence() const;
  private:
  bool _internal_has_cond_sequence() const;
  public:
  void clear_cond_sequence();
  const ::xla::cpu::ThunkSequenceProto& cond_sequence() const;
  PROTOBUF_NODISCARD ::xla::cpu::ThunkSequenceProto* release_cond_sequence();
  ::xla::cpu::ThunkSequenceProto* mutable_cond_sequence();
  void set_allocated_cond_sequence(::xla::cpu::ThunkSequenceProto* cond_sequence);
  private:
  const ::xla::cpu::ThunkSequenceProto& _internal_cond_sequence() const;
  ::xla::cpu::ThunkSequenceProto* _internal_mutable_cond_sequence();
  public:
  void unsafe_arena_set_allocated_cond_sequence(
      ::xla::cpu::ThunkSequenceProto* cond_sequence);
  ::xla::cpu::ThunkSequenceProto* unsafe_arena_release_cond_sequence();

  // .xla.cpu.ThunkSequenceProto body_sequence = 2;
  bool has_body_sequence() const;
  private:
  bool _internal_has_body_sequence() const;
  public:
  void clear_body_sequence();
  const ::xla::cpu::ThunkSequenceProto& body_sequence() const;
  PROTOBUF_NODISCARD ::xla::cpu::ThunkSequenceProto* release_body_sequence();
  ::xla::cpu::ThunkSequenceProto* mutable_body_sequence();
  void set_allocated_body_sequence(::xla::cpu::ThunkSequenceProto* body_sequence);
  private:
  const ::xla::cpu::ThunkSequenceProto& _internal_body_sequence() const;
  ::xla::cpu::ThunkSequenceProto* _internal_mutable_body_sequence();
  public:
  void unsafe_arena_set_allocated_body_sequence(
      ::xla::cpu::ThunkSequenceProto* body_sequence);
  ::xla::cpu::ThunkSequenceProto* unsafe_arena_release_body_sequence();

  // .xla.cpu.Int64Optional trip_count = 3;
  bool has_trip_count() const;
  private:
  bool _internal_has_trip_count() const;
  public:
  void clear_trip_count();
  const ::xla::cpu::Int64Optional& trip_count() const;
  PROTOBUF_NODISCARD ::xla::cpu::Int64Optional* release_trip_count();
  ::xla::cpu::Int64Optional* mutable_trip_count();
  void set_allocated_trip_count(::xla::cpu::Int64Optional* trip_count);
  private:
  const ::xla::cpu::Int64Optional& _internal_trip_count() const;
  ::xla::cpu::Int64Optional* _internal_mutable_trip_count();
  public:
  void unsafe_arena_set_allocated_trip_count(
      ::xla::cpu::Int64Optional* trip_count);
  ::xla::cpu::Int64Optional* unsafe_arena_release_trip_count();

  // .xla.cpu.BufferAllocationSliceProto cond_buffer = 4;
  bool has_cond_buffer() const;
  private:
  bool _internal_has_cond_buffer() const;
  public:
  void clear_cond_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& cond_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_cond_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_cond_buffer();
  void set_allocated_cond_buffer(::xla::cpu::BufferAllocationSliceProto* cond_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_cond_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_cond_buffer();
  public:
  void unsafe_arena_set_allocated_cond_buffer(
      ::xla::cpu::BufferAllocationSliceProto* cond_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_cond_buffer();

  // @@protoc_insertion_point(class_scope:xla.cpu.WhileThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ThunkSequenceProto* cond_sequence_;
    ::xla::cpu::ThunkSequenceProto* body_sequence_;
    ::xla::cpu::Int64Optional* trip_count_;
    ::xla::cpu::BufferAllocationSliceProto* cond_buffer_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class KernelThunkProto_ThreadDim final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.KernelThunkProto.ThreadDim) */ {
 public:
  inline KernelThunkProto_ThreadDim() : KernelThunkProto_ThreadDim(nullptr) {}
  ~KernelThunkProto_ThreadDim() override;
  explicit PROTOBUF_CONSTEXPR KernelThunkProto_ThreadDim(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  KernelThunkProto_ThreadDim(const KernelThunkProto_ThreadDim& from);
  KernelThunkProto_ThreadDim(KernelThunkProto_ThreadDim&& from) noexcept
    : KernelThunkProto_ThreadDim() {
    *this = ::std::move(from);
  }

  inline KernelThunkProto_ThreadDim& operator=(const KernelThunkProto_ThreadDim& from) {
    CopyFrom(from);
    return *this;
  }
  inline KernelThunkProto_ThreadDim& operator=(KernelThunkProto_ThreadDim&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const KernelThunkProto_ThreadDim& default_instance() {
    return *internal_default_instance();
  }
  static inline const KernelThunkProto_ThreadDim* internal_default_instance() {
    return reinterpret_cast<const KernelThunkProto_ThreadDim*>(
               &_KernelThunkProto_ThreadDim_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    30;

  friend void swap(KernelThunkProto_ThreadDim& a, KernelThunkProto_ThreadDim& b) {
    a.Swap(&b);
  }
  inline void Swap(KernelThunkProto_ThreadDim* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(KernelThunkProto_ThreadDim* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  KernelThunkProto_ThreadDim* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<KernelThunkProto_ThreadDim>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const KernelThunkProto_ThreadDim& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const KernelThunkProto_ThreadDim& from) {
    KernelThunkProto_ThreadDim::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(KernelThunkProto_ThreadDim* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.KernelThunkProto.ThreadDim";
  }
  protected:
  explicit KernelThunkProto_ThreadDim(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kXFieldNumber = 1,
    kYFieldNumber = 2,
    kZFieldNumber = 3,
  };
  // int64 x = 1;
  void clear_x();
  int64_t x() const;
  void set_x(int64_t value);
  private:
  int64_t _internal_x() const;
  void _internal_set_x(int64_t value);
  public:

  // int64 y = 2;
  void clear_y();
  int64_t y() const;
  void set_y(int64_t value);
  private:
  int64_t _internal_y() const;
  void _internal_set_y(int64_t value);
  public:

  // int64 z = 3;
  void clear_z();
  int64_t z() const;
  void set_z(int64_t value);
  private:
  int64_t _internal_z() const;
  void _internal_set_z(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.KernelThunkProto.ThreadDim)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    int64_t x_;
    int64_t y_;
    int64_t z_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class KernelThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.KernelThunkProto) */ {
 public:
  inline KernelThunkProto() : KernelThunkProto(nullptr) {}
  ~KernelThunkProto() override;
  explicit PROTOBUF_CONSTEXPR KernelThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  KernelThunkProto(const KernelThunkProto& from);
  KernelThunkProto(KernelThunkProto&& from) noexcept
    : KernelThunkProto() {
    *this = ::std::move(from);
  }

  inline KernelThunkProto& operator=(const KernelThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline KernelThunkProto& operator=(KernelThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const KernelThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const KernelThunkProto* internal_default_instance() {
    return reinterpret_cast<const KernelThunkProto*>(
               &_KernelThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    31;

  friend void swap(KernelThunkProto& a, KernelThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(KernelThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(KernelThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  KernelThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<KernelThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const KernelThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const KernelThunkProto& from) {
    KernelThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(KernelThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.KernelThunkProto";
  }
  protected:
  explicit KernelThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef KernelThunkProto_ThreadDim ThreadDim;

  // accessors -------------------------------------------------------

  enum : int {
    kInvariantArgumentsFieldNumber = 3,
    kArgumentsBuffersFieldNumber = 5,
    kResultsBuffersFieldNumber = 6,
    kKernelNameFieldNumber = 1,
    kThreadDimFieldNumber = 2,
    kMinAlignmentFieldNumber = 4,
  };
  // repeated int64 invariant_arguments = 3;
  int invariant_arguments_size() const;
  private:
  int _internal_invariant_arguments_size() const;
  public:
  void clear_invariant_arguments();
  private:
  int64_t _internal_invariant_arguments(int index) const;
  const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
      _internal_invariant_arguments() const;
  void _internal_add_invariant_arguments(int64_t value);
  ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
      _internal_mutable_invariant_arguments();
  public:
  int64_t invariant_arguments(int index) const;
  void set_invariant_arguments(int index, int64_t value);
  void add_invariant_arguments(int64_t value);
  const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
      invariant_arguments() const;
  ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
      mutable_invariant_arguments();

  // repeated .xla.cpu.BufferAllocationSliceProto arguments_buffers = 5;
  int arguments_buffers_size() const;
  private:
  int _internal_arguments_buffers_size() const;
  public:
  void clear_arguments_buffers();
  ::xla::cpu::BufferAllocationSliceProto* mutable_arguments_buffers(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >*
      mutable_arguments_buffers();
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_arguments_buffers(int index) const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_add_arguments_buffers();
  public:
  const ::xla::cpu::BufferAllocationSliceProto& arguments_buffers(int index) const;
  ::xla::cpu::BufferAllocationSliceProto* add_arguments_buffers();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >&
      arguments_buffers() const;

  // repeated .xla.cpu.BufferAllocationSliceProto results_buffers = 6;
  int results_buffers_size() const;
  private:
  int _internal_results_buffers_size() const;
  public:
  void clear_results_buffers();
  ::xla::cpu::BufferAllocationSliceProto* mutable_results_buffers(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >*
      mutable_results_buffers();
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_results_buffers(int index) const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_add_results_buffers();
  public:
  const ::xla::cpu::BufferAllocationSliceProto& results_buffers(int index) const;
  ::xla::cpu::BufferAllocationSliceProto* add_results_buffers();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >&
      results_buffers() const;

  // string kernel_name = 1;
  void clear_kernel_name();
  const std::string& kernel_name() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_kernel_name(ArgT0&& arg0, ArgT... args);
  std::string* mutable_kernel_name();
  PROTOBUF_NODISCARD std::string* release_kernel_name();
  void set_allocated_kernel_name(std::string* kernel_name);
  private:
  const std::string& _internal_kernel_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_kernel_name(const std::string& value);
  std::string* _internal_mutable_kernel_name();
  public:

  // .xla.cpu.KernelThunkProto.ThreadDim thread_dim = 2;
  bool has_thread_dim() const;
  private:
  bool _internal_has_thread_dim() const;
  public:
  void clear_thread_dim();
  const ::xla::cpu::KernelThunkProto_ThreadDim& thread_dim() const;
  PROTOBUF_NODISCARD ::xla::cpu::KernelThunkProto_ThreadDim* release_thread_dim();
  ::xla::cpu::KernelThunkProto_ThreadDim* mutable_thread_dim();
  void set_allocated_thread_dim(::xla::cpu::KernelThunkProto_ThreadDim* thread_dim);
  private:
  const ::xla::cpu::KernelThunkProto_ThreadDim& _internal_thread_dim() const;
  ::xla::cpu::KernelThunkProto_ThreadDim* _internal_mutable_thread_dim();
  public:
  void unsafe_arena_set_allocated_thread_dim(
      ::xla::cpu::KernelThunkProto_ThreadDim* thread_dim);
  ::xla::cpu::KernelThunkProto_ThreadDim* unsafe_arena_release_thread_dim();

  // .xla.cpu.Int64Optional min_alignment = 4;
  bool has_min_alignment() const;
  private:
  bool _internal_has_min_alignment() const;
  public:
  void clear_min_alignment();
  const ::xla::cpu::Int64Optional& min_alignment() const;
  PROTOBUF_NODISCARD ::xla::cpu::Int64Optional* release_min_alignment();
  ::xla::cpu::Int64Optional* mutable_min_alignment();
  void set_allocated_min_alignment(::xla::cpu::Int64Optional* min_alignment);
  private:
  const ::xla::cpu::Int64Optional& _internal_min_alignment() const;
  ::xla::cpu::Int64Optional* _internal_mutable_min_alignment();
  public:
  void unsafe_arena_set_allocated_min_alignment(
      ::xla::cpu::Int64Optional* min_alignment);
  ::xla::cpu::Int64Optional* unsafe_arena_release_min_alignment();

  // @@protoc_insertion_point(class_scope:xla.cpu.KernelThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t > invariant_arguments_;
    mutable std::atomic<int> _invariant_arguments_cached_byte_size_;
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto > arguments_buffers_;
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto > results_buffers_;
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr kernel_name_;
    ::xla::cpu::KernelThunkProto_ThreadDim* thread_dim_;
    ::xla::cpu::Int64Optional* min_alignment_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CopyThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CopyThunkProto) */ {
 public:
  inline CopyThunkProto() : CopyThunkProto(nullptr) {}
  ~CopyThunkProto() override;
  explicit PROTOBUF_CONSTEXPR CopyThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CopyThunkProto(const CopyThunkProto& from);
  CopyThunkProto(CopyThunkProto&& from) noexcept
    : CopyThunkProto() {
    *this = ::std::move(from);
  }

  inline CopyThunkProto& operator=(const CopyThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CopyThunkProto& operator=(CopyThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CopyThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const CopyThunkProto* internal_default_instance() {
    return reinterpret_cast<const CopyThunkProto*>(
               &_CopyThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    32;

  friend void swap(CopyThunkProto& a, CopyThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CopyThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CopyThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CopyThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CopyThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CopyThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CopyThunkProto& from) {
    CopyThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CopyThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CopyThunkProto";
  }
  protected:
  explicit CopyThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kSrcBufferShapeFieldNumber = 1,
    kDstBufferShapeFieldNumber = 2,
  };
  // .xla.cpu.ShapeBufferAllocationSliceProto src_buffer_shape = 1;
  bool has_src_buffer_shape() const;
  private:
  bool _internal_has_src_buffer_shape() const;
  public:
  void clear_src_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& src_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_src_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_src_buffer_shape();
  void set_allocated_src_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* src_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_src_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_src_buffer_shape();
  public:
  void unsafe_arena_set_allocated_src_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* src_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_src_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto dst_buffer_shape = 2;
  bool has_dst_buffer_shape() const;
  private:
  bool _internal_has_dst_buffer_shape() const;
  public:
  void clear_dst_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& dst_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_dst_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_dst_buffer_shape();
  void set_allocated_dst_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* dst_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_dst_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_dst_buffer_shape();
  public:
  void unsafe_arena_set_allocated_dst_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* dst_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_dst_buffer_shape();

  // @@protoc_insertion_point(class_scope:xla.cpu.CopyThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ShapeBufferAllocationSliceProto* src_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* dst_buffer_shape_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class FftThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.FftThunkProto) */ {
 public:
  inline FftThunkProto() : FftThunkProto(nullptr) {}
  ~FftThunkProto() override;
  explicit PROTOBUF_CONSTEXPR FftThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  FftThunkProto(const FftThunkProto& from);
  FftThunkProto(FftThunkProto&& from) noexcept
    : FftThunkProto() {
    *this = ::std::move(from);
  }

  inline FftThunkProto& operator=(const FftThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline FftThunkProto& operator=(FftThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const FftThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const FftThunkProto* internal_default_instance() {
    return reinterpret_cast<const FftThunkProto*>(
               &_FftThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    33;

  friend void swap(FftThunkProto& a, FftThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(FftThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(FftThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  FftThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<FftThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const FftThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const FftThunkProto& from) {
    FftThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(FftThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.FftThunkProto";
  }
  protected:
  explicit FftThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kFftLengthFieldNumber = 3,
    kInputBufferShapeFieldNumber = 4,
    kOutputBufferShapeFieldNumber = 5,
    kIsMultiThreadEigenFieldNumber = 1,
    kFftTypeFieldNumber = 2,
  };
  // repeated int64 fft_length = 3;
  int fft_length_size() const;
  private:
  int _internal_fft_length_size() const;
  public:
  void clear_fft_length();
  private:
  int64_t _internal_fft_length(int index) const;
  const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
      _internal_fft_length() const;
  void _internal_add_fft_length(int64_t value);
  ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
      _internal_mutable_fft_length();
  public:
  int64_t fft_length(int index) const;
  void set_fft_length(int index, int64_t value);
  void add_fft_length(int64_t value);
  const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
      fft_length() const;
  ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
      mutable_fft_length();

  // .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
  bool has_input_buffer_shape() const;
  private:
  bool _internal_has_input_buffer_shape() const;
  public:
  void clear_input_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& input_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_input_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_input_buffer_shape();
  void set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_input_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_input_buffer_shape();
  public:
  void unsafe_arena_set_allocated_input_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_input_buffer_shape();

  // .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 5;
  bool has_output_buffer_shape() const;
  private:
  bool _internal_has_output_buffer_shape() const;
  public:
  void clear_output_buffer_shape();
  const ::xla::cpu::ShapeBufferAllocationSliceProto& output_buffer_shape() const;
  PROTOBUF_NODISCARD ::xla::cpu::ShapeBufferAllocationSliceProto* release_output_buffer_shape();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_output_buffer_shape();
  void set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_output_buffer_shape() const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_mutable_output_buffer_shape();
  public:
  void unsafe_arena_set_allocated_output_buffer_shape(
      ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape);
  ::xla::cpu::ShapeBufferAllocationSliceProto* unsafe_arena_release_output_buffer_shape();

  // bool is_multi_thread_eigen = 1;
  void clear_is_multi_thread_eigen();
  bool is_multi_thread_eigen() const;
  void set_is_multi_thread_eigen(bool value);
  private:
  bool _internal_is_multi_thread_eigen() const;
  void _internal_set_is_multi_thread_eigen(bool value);
  public:

  // int32 fft_type = 2;
  void clear_fft_type();
  int32_t fft_type() const;
  void set_fft_type(int32_t value);
  private:
  int32_t _internal_fft_type() const;
  void _internal_set_fft_type(int32_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.FftThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t > fft_length_;
    mutable std::atomic<int> _fft_length_cached_byte_size_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape_;
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape_;
    bool is_multi_thread_eigen_;
    int32_t fft_type_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class InfeedThunkProto_InfeedResource final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.InfeedThunkProto.InfeedResource) */ {
 public:
  inline InfeedThunkProto_InfeedResource() : InfeedThunkProto_InfeedResource(nullptr) {}
  ~InfeedThunkProto_InfeedResource() override;
  explicit PROTOBUF_CONSTEXPR InfeedThunkProto_InfeedResource(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  InfeedThunkProto_InfeedResource(const InfeedThunkProto_InfeedResource& from);
  InfeedThunkProto_InfeedResource(InfeedThunkProto_InfeedResource&& from) noexcept
    : InfeedThunkProto_InfeedResource() {
    *this = ::std::move(from);
  }

  inline InfeedThunkProto_InfeedResource& operator=(const InfeedThunkProto_InfeedResource& from) {
    CopyFrom(from);
    return *this;
  }
  inline InfeedThunkProto_InfeedResource& operator=(InfeedThunkProto_InfeedResource&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const InfeedThunkProto_InfeedResource& default_instance() {
    return *internal_default_instance();
  }
  static inline const InfeedThunkProto_InfeedResource* internal_default_instance() {
    return reinterpret_cast<const InfeedThunkProto_InfeedResource*>(
               &_InfeedThunkProto_InfeedResource_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    34;

  friend void swap(InfeedThunkProto_InfeedResource& a, InfeedThunkProto_InfeedResource& b) {
    a.Swap(&b);
  }
  inline void Swap(InfeedThunkProto_InfeedResource* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(InfeedThunkProto_InfeedResource* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  InfeedThunkProto_InfeedResource* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<InfeedThunkProto_InfeedResource>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const InfeedThunkProto_InfeedResource& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const InfeedThunkProto_InfeedResource& from) {
    InfeedThunkProto_InfeedResource::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(InfeedThunkProto_InfeedResource* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.InfeedThunkProto.InfeedResource";
  }
  protected:
  explicit InfeedThunkProto_InfeedResource(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kConsumeTokenFieldNumber = 1,
    kProduceTokenFieldNumber = 2,
  };
  // .xla.cpu.ResourceOptional consume_token = 1;
  bool has_consume_token() const;
  private:
  bool _internal_has_consume_token() const;
  public:
  void clear_consume_token();
  const ::xla::cpu::ResourceOptional& consume_token() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceOptional* release_consume_token();
  ::xla::cpu::ResourceOptional* mutable_consume_token();
  void set_allocated_consume_token(::xla::cpu::ResourceOptional* consume_token);
  private:
  const ::xla::cpu::ResourceOptional& _internal_consume_token() const;
  ::xla::cpu::ResourceOptional* _internal_mutable_consume_token();
  public:
  void unsafe_arena_set_allocated_consume_token(
      ::xla::cpu::ResourceOptional* consume_token);
  ::xla::cpu::ResourceOptional* unsafe_arena_release_consume_token();

  // .xla.cpu.ResourceOptional produce_token = 2;
  bool has_produce_token() const;
  private:
  bool _internal_has_produce_token() const;
  public:
  void clear_produce_token();
  const ::xla::cpu::ResourceOptional& produce_token() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceOptional* release_produce_token();
  ::xla::cpu::ResourceOptional* mutable_produce_token();
  void set_allocated_produce_token(::xla::cpu::ResourceOptional* produce_token);
  private:
  const ::xla::cpu::ResourceOptional& _internal_produce_token() const;
  ::xla::cpu::ResourceOptional* _internal_mutable_produce_token();
  public:
  void unsafe_arena_set_allocated_produce_token(
      ::xla::cpu::ResourceOptional* produce_token);
  ::xla::cpu::ResourceOptional* unsafe_arena_release_produce_token();

  // @@protoc_insertion_point(class_scope:xla.cpu.InfeedThunkProto.InfeedResource)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ResourceOptional* consume_token_;
    ::xla::cpu::ResourceOptional* produce_token_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class InfeedThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.InfeedThunkProto) */ {
 public:
  inline InfeedThunkProto() : InfeedThunkProto(nullptr) {}
  ~InfeedThunkProto() override;
  explicit PROTOBUF_CONSTEXPR InfeedThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  InfeedThunkProto(const InfeedThunkProto& from);
  InfeedThunkProto(InfeedThunkProto&& from) noexcept
    : InfeedThunkProto() {
    *this = ::std::move(from);
  }

  inline InfeedThunkProto& operator=(const InfeedThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline InfeedThunkProto& operator=(InfeedThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const InfeedThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const InfeedThunkProto* internal_default_instance() {
    return reinterpret_cast<const InfeedThunkProto*>(
               &_InfeedThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    35;

  friend void swap(InfeedThunkProto& a, InfeedThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(InfeedThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(InfeedThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  InfeedThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<InfeedThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const InfeedThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const InfeedThunkProto& from) {
    InfeedThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(InfeedThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.InfeedThunkProto";
  }
  protected:
  explicit InfeedThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef InfeedThunkProto_InfeedResource InfeedResource;

  // accessors -------------------------------------------------------

  enum : int {
    kInfeedBuffersShapesFieldNumber = 2,
    kInfeedResourcesFieldNumber = 1,
  };
  // repeated .xla.cpu.ShapeBufferAllocationSliceProto infeed_buffers_shapes = 2;
  int infeed_buffers_shapes_size() const;
  private:
  int _internal_infeed_buffers_shapes_size() const;
  public:
  void clear_infeed_buffers_shapes();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_infeed_buffers_shapes(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_infeed_buffers_shapes();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_infeed_buffers_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_infeed_buffers_shapes();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& infeed_buffers_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_infeed_buffers_shapes();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      infeed_buffers_shapes() const;

  // .xla.cpu.InfeedThunkProto.InfeedResource infeed_resources = 1;
  bool has_infeed_resources() const;
  private:
  bool _internal_has_infeed_resources() const;
  public:
  void clear_infeed_resources();
  const ::xla::cpu::InfeedThunkProto_InfeedResource& infeed_resources() const;
  PROTOBUF_NODISCARD ::xla::cpu::InfeedThunkProto_InfeedResource* release_infeed_resources();
  ::xla::cpu::InfeedThunkProto_InfeedResource* mutable_infeed_resources();
  void set_allocated_infeed_resources(::xla::cpu::InfeedThunkProto_InfeedResource* infeed_resources);
  private:
  const ::xla::cpu::InfeedThunkProto_InfeedResource& _internal_infeed_resources() const;
  ::xla::cpu::InfeedThunkProto_InfeedResource* _internal_mutable_infeed_resources();
  public:
  void unsafe_arena_set_allocated_infeed_resources(
      ::xla::cpu::InfeedThunkProto_InfeedResource* infeed_resources);
  ::xla::cpu::InfeedThunkProto_InfeedResource* unsafe_arena_release_infeed_resources();

  // @@protoc_insertion_point(class_scope:xla.cpu.InfeedThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > infeed_buffers_shapes_;
    ::xla::cpu::InfeedThunkProto_InfeedResource* infeed_resources_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class OutfeedThunkProto_OutfeedResource final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.OutfeedThunkProto.OutfeedResource) */ {
 public:
  inline OutfeedThunkProto_OutfeedResource() : OutfeedThunkProto_OutfeedResource(nullptr) {}
  ~OutfeedThunkProto_OutfeedResource() override;
  explicit PROTOBUF_CONSTEXPR OutfeedThunkProto_OutfeedResource(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  OutfeedThunkProto_OutfeedResource(const OutfeedThunkProto_OutfeedResource& from);
  OutfeedThunkProto_OutfeedResource(OutfeedThunkProto_OutfeedResource&& from) noexcept
    : OutfeedThunkProto_OutfeedResource() {
    *this = ::std::move(from);
  }

  inline OutfeedThunkProto_OutfeedResource& operator=(const OutfeedThunkProto_OutfeedResource& from) {
    CopyFrom(from);
    return *this;
  }
  inline OutfeedThunkProto_OutfeedResource& operator=(OutfeedThunkProto_OutfeedResource&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const OutfeedThunkProto_OutfeedResource& default_instance() {
    return *internal_default_instance();
  }
  static inline const OutfeedThunkProto_OutfeedResource* internal_default_instance() {
    return reinterpret_cast<const OutfeedThunkProto_OutfeedResource*>(
               &_OutfeedThunkProto_OutfeedResource_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    36;

  friend void swap(OutfeedThunkProto_OutfeedResource& a, OutfeedThunkProto_OutfeedResource& b) {
    a.Swap(&b);
  }
  inline void Swap(OutfeedThunkProto_OutfeedResource* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(OutfeedThunkProto_OutfeedResource* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  OutfeedThunkProto_OutfeedResource* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<OutfeedThunkProto_OutfeedResource>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const OutfeedThunkProto_OutfeedResource& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const OutfeedThunkProto_OutfeedResource& from) {
    OutfeedThunkProto_OutfeedResource::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(OutfeedThunkProto_OutfeedResource* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.OutfeedThunkProto.OutfeedResource";
  }
  protected:
  explicit OutfeedThunkProto_OutfeedResource(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kConsumeTokenFieldNumber = 1,
    kProduceTokenFieldNumber = 2,
  };
  // .xla.cpu.ResourceOptional consume_token = 1;
  bool has_consume_token() const;
  private:
  bool _internal_has_consume_token() const;
  public:
  void clear_consume_token();
  const ::xla::cpu::ResourceOptional& consume_token() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceOptional* release_consume_token();
  ::xla::cpu::ResourceOptional* mutable_consume_token();
  void set_allocated_consume_token(::xla::cpu::ResourceOptional* consume_token);
  private:
  const ::xla::cpu::ResourceOptional& _internal_consume_token() const;
  ::xla::cpu::ResourceOptional* _internal_mutable_consume_token();
  public:
  void unsafe_arena_set_allocated_consume_token(
      ::xla::cpu::ResourceOptional* consume_token);
  ::xla::cpu::ResourceOptional* unsafe_arena_release_consume_token();

  // .xla.cpu.ResourceOptional produce_token = 2;
  bool has_produce_token() const;
  private:
  bool _internal_has_produce_token() const;
  public:
  void clear_produce_token();
  const ::xla::cpu::ResourceOptional& produce_token() const;
  PROTOBUF_NODISCARD ::xla::cpu::ResourceOptional* release_produce_token();
  ::xla::cpu::ResourceOptional* mutable_produce_token();
  void set_allocated_produce_token(::xla::cpu::ResourceOptional* produce_token);
  private:
  const ::xla::cpu::ResourceOptional& _internal_produce_token() const;
  ::xla::cpu::ResourceOptional* _internal_mutable_produce_token();
  public:
  void unsafe_arena_set_allocated_produce_token(
      ::xla::cpu::ResourceOptional* produce_token);
  ::xla::cpu::ResourceOptional* unsafe_arena_release_produce_token();

  // @@protoc_insertion_point(class_scope:xla.cpu.OutfeedThunkProto.OutfeedResource)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::ResourceOptional* consume_token_;
    ::xla::cpu::ResourceOptional* produce_token_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class OutfeedThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.OutfeedThunkProto) */ {
 public:
  inline OutfeedThunkProto() : OutfeedThunkProto(nullptr) {}
  ~OutfeedThunkProto() override;
  explicit PROTOBUF_CONSTEXPR OutfeedThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  OutfeedThunkProto(const OutfeedThunkProto& from);
  OutfeedThunkProto(OutfeedThunkProto&& from) noexcept
    : OutfeedThunkProto() {
    *this = ::std::move(from);
  }

  inline OutfeedThunkProto& operator=(const OutfeedThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline OutfeedThunkProto& operator=(OutfeedThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const OutfeedThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const OutfeedThunkProto* internal_default_instance() {
    return reinterpret_cast<const OutfeedThunkProto*>(
               &_OutfeedThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    37;

  friend void swap(OutfeedThunkProto& a, OutfeedThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(OutfeedThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(OutfeedThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  OutfeedThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<OutfeedThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const OutfeedThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const OutfeedThunkProto& from) {
    OutfeedThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(OutfeedThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.OutfeedThunkProto";
  }
  protected:
  explicit OutfeedThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef OutfeedThunkProto_OutfeedResource OutfeedResource;

  // accessors -------------------------------------------------------

  enum : int {
    kOutfeedBuffersShapesFieldNumber = 2,
    kOutfeedResourcesFieldNumber = 1,
  };
  // repeated .xla.cpu.ShapeBufferAllocationSliceProto outfeed_buffers_shapes = 2;
  int outfeed_buffers_shapes_size() const;
  private:
  int _internal_outfeed_buffers_shapes_size() const;
  public:
  void clear_outfeed_buffers_shapes();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_outfeed_buffers_shapes(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_outfeed_buffers_shapes();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_outfeed_buffers_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_outfeed_buffers_shapes();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& outfeed_buffers_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_outfeed_buffers_shapes();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      outfeed_buffers_shapes() const;

  // .xla.cpu.OutfeedThunkProto.OutfeedResource outfeed_resources = 1;
  bool has_outfeed_resources() const;
  private:
  bool _internal_has_outfeed_resources() const;
  public:
  void clear_outfeed_resources();
  const ::xla::cpu::OutfeedThunkProto_OutfeedResource& outfeed_resources() const;
  PROTOBUF_NODISCARD ::xla::cpu::OutfeedThunkProto_OutfeedResource* release_outfeed_resources();
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* mutable_outfeed_resources();
  void set_allocated_outfeed_resources(::xla::cpu::OutfeedThunkProto_OutfeedResource* outfeed_resources);
  private:
  const ::xla::cpu::OutfeedThunkProto_OutfeedResource& _internal_outfeed_resources() const;
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* _internal_mutable_outfeed_resources();
  public:
  void unsafe_arena_set_allocated_outfeed_resources(
      ::xla::cpu::OutfeedThunkProto_OutfeedResource* outfeed_resources);
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* unsafe_arena_release_outfeed_resources();

  // @@protoc_insertion_point(class_scope:xla.cpu.OutfeedThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > outfeed_buffers_shapes_;
    ::xla::cpu::OutfeedThunkProto_OutfeedResource* outfeed_resources_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CustomCallThunkProto_OpBuffers final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CustomCallThunkProto.OpBuffers) */ {
 public:
  inline CustomCallThunkProto_OpBuffers() : CustomCallThunkProto_OpBuffers(nullptr) {}
  ~CustomCallThunkProto_OpBuffers() override;
  explicit PROTOBUF_CONSTEXPR CustomCallThunkProto_OpBuffers(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CustomCallThunkProto_OpBuffers(const CustomCallThunkProto_OpBuffers& from);
  CustomCallThunkProto_OpBuffers(CustomCallThunkProto_OpBuffers&& from) noexcept
    : CustomCallThunkProto_OpBuffers() {
    *this = ::std::move(from);
  }

  inline CustomCallThunkProto_OpBuffers& operator=(const CustomCallThunkProto_OpBuffers& from) {
    CopyFrom(from);
    return *this;
  }
  inline CustomCallThunkProto_OpBuffers& operator=(CustomCallThunkProto_OpBuffers&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CustomCallThunkProto_OpBuffers& default_instance() {
    return *internal_default_instance();
  }
  static inline const CustomCallThunkProto_OpBuffers* internal_default_instance() {
    return reinterpret_cast<const CustomCallThunkProto_OpBuffers*>(
               &_CustomCallThunkProto_OpBuffers_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    38;

  friend void swap(CustomCallThunkProto_OpBuffers& a, CustomCallThunkProto_OpBuffers& b) {
    a.Swap(&b);
  }
  inline void Swap(CustomCallThunkProto_OpBuffers* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CustomCallThunkProto_OpBuffers* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CustomCallThunkProto_OpBuffers* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CustomCallThunkProto_OpBuffers>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CustomCallThunkProto_OpBuffers& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CustomCallThunkProto_OpBuffers& from) {
    CustomCallThunkProto_OpBuffers::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CustomCallThunkProto_OpBuffers* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CustomCallThunkProto.OpBuffers";
  }
  protected:
  explicit CustomCallThunkProto_OpBuffers(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kArgumentsShapesFieldNumber = 1,
    kResultsShapesFieldNumber = 2,
  };
  // repeated .xla.cpu.ShapeBufferAllocationSliceProto arguments_shapes = 1;
  int arguments_shapes_size() const;
  private:
  int _internal_arguments_shapes_size() const;
  public:
  void clear_arguments_shapes();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_arguments_shapes(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_arguments_shapes();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_arguments_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_arguments_shapes();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& arguments_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_arguments_shapes();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      arguments_shapes() const;

  // repeated .xla.cpu.ShapeBufferAllocationSliceProto results_shapes = 2;
  int results_shapes_size() const;
  private:
  int _internal_results_shapes_size() const;
  public:
  void clear_results_shapes();
  ::xla::cpu::ShapeBufferAllocationSliceProto* mutable_results_shapes(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
      mutable_results_shapes();
  private:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& _internal_results_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* _internal_add_results_shapes();
  public:
  const ::xla::cpu::ShapeBufferAllocationSliceProto& results_shapes(int index) const;
  ::xla::cpu::ShapeBufferAllocationSliceProto* add_results_shapes();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
      results_shapes() const;

  // @@protoc_insertion_point(class_scope:xla.cpu.CustomCallThunkProto.OpBuffers)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > arguments_shapes_;
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto > results_shapes_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class CustomCallThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.CustomCallThunkProto) */ {
 public:
  inline CustomCallThunkProto() : CustomCallThunkProto(nullptr) {}
  ~CustomCallThunkProto() override;
  explicit PROTOBUF_CONSTEXPR CustomCallThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  CustomCallThunkProto(const CustomCallThunkProto& from);
  CustomCallThunkProto(CustomCallThunkProto&& from) noexcept
    : CustomCallThunkProto() {
    *this = ::std::move(from);
  }

  inline CustomCallThunkProto& operator=(const CustomCallThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline CustomCallThunkProto& operator=(CustomCallThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const CustomCallThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const CustomCallThunkProto* internal_default_instance() {
    return reinterpret_cast<const CustomCallThunkProto*>(
               &_CustomCallThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    39;

  friend void swap(CustomCallThunkProto& a, CustomCallThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(CustomCallThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(CustomCallThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  CustomCallThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<CustomCallThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const CustomCallThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const CustomCallThunkProto& from) {
    CustomCallThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(CustomCallThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.CustomCallThunkProto";
  }
  protected:
  explicit CustomCallThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  typedef CustomCallThunkProto_OpBuffers OpBuffers;

  // accessors -------------------------------------------------------

  enum : int {
    kTargetNameFieldNumber = 2,
    kBackendConfigFieldNumber = 3,
    kOpBuffersFieldNumber = 4,
    kApiVersionFieldNumber = 1,
  };
  // string target_name = 2;
  void clear_target_name();
  const std::string& target_name() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_target_name(ArgT0&& arg0, ArgT... args);
  std::string* mutable_target_name();
  PROTOBUF_NODISCARD std::string* release_target_name();
  void set_allocated_target_name(std::string* target_name);
  private:
  const std::string& _internal_target_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_target_name(const std::string& value);
  std::string* _internal_mutable_target_name();
  public:

  // string backend_config = 3;
  void clear_backend_config();
  const std::string& backend_config() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_backend_config(ArgT0&& arg0, ArgT... args);
  std::string* mutable_backend_config();
  PROTOBUF_NODISCARD std::string* release_backend_config();
  void set_allocated_backend_config(std::string* backend_config);
  private:
  const std::string& _internal_backend_config() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_backend_config(const std::string& value);
  std::string* _internal_mutable_backend_config();
  public:

  // .xla.cpu.CustomCallThunkProto.OpBuffers op_buffers = 4;
  bool has_op_buffers() const;
  private:
  bool _internal_has_op_buffers() const;
  public:
  void clear_op_buffers();
  const ::xla::cpu::CustomCallThunkProto_OpBuffers& op_buffers() const;
  PROTOBUF_NODISCARD ::xla::cpu::CustomCallThunkProto_OpBuffers* release_op_buffers();
  ::xla::cpu::CustomCallThunkProto_OpBuffers* mutable_op_buffers();
  void set_allocated_op_buffers(::xla::cpu::CustomCallThunkProto_OpBuffers* op_buffers);
  private:
  const ::xla::cpu::CustomCallThunkProto_OpBuffers& _internal_op_buffers() const;
  ::xla::cpu::CustomCallThunkProto_OpBuffers* _internal_mutable_op_buffers();
  public:
  void unsafe_arena_set_allocated_op_buffers(
      ::xla::cpu::CustomCallThunkProto_OpBuffers* op_buffers);
  ::xla::cpu::CustomCallThunkProto_OpBuffers* unsafe_arena_release_op_buffers();

  // .xla.CustomCallApiVersion api_version = 1;
  void clear_api_version();
  ::xla::CustomCallApiVersion api_version() const;
  void set_api_version(::xla::CustomCallApiVersion value);
  private:
  ::xla::CustomCallApiVersion _internal_api_version() const;
  void _internal_set_api_version(::xla::CustomCallApiVersion value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.CustomCallThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr target_name_;
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr backend_config_;
    ::xla::cpu::CustomCallThunkProto_OpBuffers* op_buffers_;
    int api_version_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class PartitionIdThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.PartitionIdThunkProto) */ {
 public:
  inline PartitionIdThunkProto() : PartitionIdThunkProto(nullptr) {}
  ~PartitionIdThunkProto() override;
  explicit PROTOBUF_CONSTEXPR PartitionIdThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  PartitionIdThunkProto(const PartitionIdThunkProto& from);
  PartitionIdThunkProto(PartitionIdThunkProto&& from) noexcept
    : PartitionIdThunkProto() {
    *this = ::std::move(from);
  }

  inline PartitionIdThunkProto& operator=(const PartitionIdThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline PartitionIdThunkProto& operator=(PartitionIdThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const PartitionIdThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const PartitionIdThunkProto* internal_default_instance() {
    return reinterpret_cast<const PartitionIdThunkProto*>(
               &_PartitionIdThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    40;

  friend void swap(PartitionIdThunkProto& a, PartitionIdThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(PartitionIdThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(PartitionIdThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  PartitionIdThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<PartitionIdThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const PartitionIdThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const PartitionIdThunkProto& from) {
    PartitionIdThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(PartitionIdThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.PartitionIdThunkProto";
  }
  protected:
  explicit PartitionIdThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kLogicalIdBufferFieldNumber = 1,
  };
  // .xla.cpu.BufferAllocationSliceProto logical_id_buffer = 1;
  bool has_logical_id_buffer() const;
  private:
  bool _internal_has_logical_id_buffer() const;
  public:
  void clear_logical_id_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& logical_id_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_logical_id_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_logical_id_buffer();
  void set_allocated_logical_id_buffer(::xla::cpu::BufferAllocationSliceProto* logical_id_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_logical_id_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_logical_id_buffer();
  public:
  void unsafe_arena_set_allocated_logical_id_buffer(
      ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_logical_id_buffer();

  // @@protoc_insertion_point(class_scope:xla.cpu.PartitionIdThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ReplicaIdThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ReplicaIdThunkProto) */ {
 public:
  inline ReplicaIdThunkProto() : ReplicaIdThunkProto(nullptr) {}
  ~ReplicaIdThunkProto() override;
  explicit PROTOBUF_CONSTEXPR ReplicaIdThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ReplicaIdThunkProto(const ReplicaIdThunkProto& from);
  ReplicaIdThunkProto(ReplicaIdThunkProto&& from) noexcept
    : ReplicaIdThunkProto() {
    *this = ::std::move(from);
  }

  inline ReplicaIdThunkProto& operator=(const ReplicaIdThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ReplicaIdThunkProto& operator=(ReplicaIdThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ReplicaIdThunkProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ReplicaIdThunkProto* internal_default_instance() {
    return reinterpret_cast<const ReplicaIdThunkProto*>(
               &_ReplicaIdThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    41;

  friend void swap(ReplicaIdThunkProto& a, ReplicaIdThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ReplicaIdThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ReplicaIdThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ReplicaIdThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ReplicaIdThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ReplicaIdThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ReplicaIdThunkProto& from) {
    ReplicaIdThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ReplicaIdThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ReplicaIdThunkProto";
  }
  protected:
  explicit ReplicaIdThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kLogicalIdBufferFieldNumber = 1,
  };
  // .xla.cpu.BufferAllocationSliceProto logical_id_buffer = 1;
  bool has_logical_id_buffer() const;
  private:
  bool _internal_has_logical_id_buffer() const;
  public:
  void clear_logical_id_buffer();
  const ::xla::cpu::BufferAllocationSliceProto& logical_id_buffer() const;
  PROTOBUF_NODISCARD ::xla::cpu::BufferAllocationSliceProto* release_logical_id_buffer();
  ::xla::cpu::BufferAllocationSliceProto* mutable_logical_id_buffer();
  void set_allocated_logical_id_buffer(::xla::cpu::BufferAllocationSliceProto* logical_id_buffer);
  private:
  const ::xla::cpu::BufferAllocationSliceProto& _internal_logical_id_buffer() const;
  ::xla::cpu::BufferAllocationSliceProto* _internal_mutable_logical_id_buffer();
  public:
  void unsafe_arena_set_allocated_logical_id_buffer(
      ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer);
  ::xla::cpu::BufferAllocationSliceProto* unsafe_arena_release_logical_id_buffer();

  // @@protoc_insertion_point(class_scope:xla.cpu.ReplicaIdThunkProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class InfoProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.InfoProto) */ {
 public:
  inline InfoProto() : InfoProto(nullptr) {}
  ~InfoProto() override;
  explicit PROTOBUF_CONSTEXPR InfoProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  InfoProto(const InfoProto& from);
  InfoProto(InfoProto&& from) noexcept
    : InfoProto() {
    *this = ::std::move(from);
  }

  inline InfoProto& operator=(const InfoProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline InfoProto& operator=(InfoProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const InfoProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const InfoProto* internal_default_instance() {
    return reinterpret_cast<const InfoProto*>(
               &_InfoProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    42;

  friend void swap(InfoProto& a, InfoProto& b) {
    a.Swap(&b);
  }
  inline void Swap(InfoProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(InfoProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  InfoProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<InfoProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const InfoProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const InfoProto& from) {
    InfoProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(InfoProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.InfoProto";
  }
  protected:
  explicit InfoProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kOpNameFieldNumber = 1,
    kModuleNameFieldNumber = 2,
    kModuleIdFieldNumber = 3,
  };
  // string op_name = 1;
  void clear_op_name();
  const std::string& op_name() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_op_name(ArgT0&& arg0, ArgT... args);
  std::string* mutable_op_name();
  PROTOBUF_NODISCARD std::string* release_op_name();
  void set_allocated_op_name(std::string* op_name);
  private:
  const std::string& _internal_op_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_op_name(const std::string& value);
  std::string* _internal_mutable_op_name();
  public:

  // string module_name = 2;
  void clear_module_name();
  const std::string& module_name() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_module_name(ArgT0&& arg0, ArgT... args);
  std::string* mutable_module_name();
  PROTOBUF_NODISCARD std::string* release_module_name();
  void set_allocated_module_name(std::string* module_name);
  private:
  const std::string& _internal_module_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_module_name(const std::string& value);
  std::string* _internal_mutable_module_name();
  public:

  // int64 module_id = 3;
  void clear_module_id();
  int64_t module_id() const;
  void set_module_id(int64_t value);
  private:
  int64_t _internal_module_id() const;
  void _internal_set_module_id(int64_t value);
  public:

  // @@protoc_insertion_point(class_scope:xla.cpu.InfoProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr op_name_;
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr module_name_;
    int64_t module_id_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ThunkProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ThunkProto) */ {
 public:
  inline ThunkProto() : ThunkProto(nullptr) {}
  ~ThunkProto() override;
  explicit PROTOBUF_CONSTEXPR ThunkProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ThunkProto(const ThunkProto& from);
  ThunkProto(ThunkProto&& from) noexcept
    : ThunkProto() {
    *this = ::std::move(from);
  }

  inline ThunkProto& operator=(const ThunkProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ThunkProto& operator=(ThunkProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ThunkProto& default_instance() {
    return *internal_default_instance();
  }
  enum ImplCase {
    kCallThunk = 3,
    kConditionalThunk = 4,
    kSortThunk = 5,
    kXnnFusionThunk = 6,
    kDotThunk = 7,
    kRngGetAndUpdateStateThunk = 8,
    kTopKThunk = 9,
    kWhileThunk = 10,
    kKernelThunk = 11,
    kCopyThunk = 12,
    kFftThunk = 13,
    kInfeedThunk = 14,
    kOutfeedThunk = 15,
    kCustomCallThunk = 16,
    kConvolutionThunk = 17,
    kCollectiveThunk = 18,
    kPartitionIdThunk = 19,
    kReplicaIdThunk = 20,
    IMPL_NOT_SET = 0,
  };

  static inline const ThunkProto* internal_default_instance() {
    return reinterpret_cast<const ThunkProto*>(
               &_ThunkProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    43;

  friend void swap(ThunkProto& a, ThunkProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ThunkProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ThunkProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ThunkProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ThunkProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ThunkProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ThunkProto& from) {
    ThunkProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ThunkProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ThunkProto";
  }
  protected:
  explicit ThunkProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kKindFieldNumber = 1,
    kInfoFieldNumber = 2,
    kCallThunkFieldNumber = 3,
    kConditionalThunkFieldNumber = 4,
    kSortThunkFieldNumber = 5,
    kXnnFusionThunkFieldNumber = 6,
    kDotThunkFieldNumber = 7,
    kRngGetAndUpdateStateThunkFieldNumber = 8,
    kTopKThunkFieldNumber = 9,
    kWhileThunkFieldNumber = 10,
    kKernelThunkFieldNumber = 11,
    kCopyThunkFieldNumber = 12,
    kFftThunkFieldNumber = 13,
    kInfeedThunkFieldNumber = 14,
    kOutfeedThunkFieldNumber = 15,
    kCustomCallThunkFieldNumber = 16,
    kConvolutionThunkFieldNumber = 17,
    kCollectiveThunkFieldNumber = 18,
    kPartitionIdThunkFieldNumber = 19,
    kReplicaIdThunkFieldNumber = 20,
  };
  // string kind = 1;
  void clear_kind();
  const std::string& kind() const;
  template <typename ArgT0 = const std::string&, typename... ArgT>
  void set_kind(ArgT0&& arg0, ArgT... args);
  std::string* mutable_kind();
  PROTOBUF_NODISCARD std::string* release_kind();
  void set_allocated_kind(std::string* kind);
  private:
  const std::string& _internal_kind() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_kind(const std::string& value);
  std::string* _internal_mutable_kind();
  public:

  // .xla.cpu.InfoProto info = 2;
  bool has_info() const;
  private:
  bool _internal_has_info() const;
  public:
  void clear_info();
  const ::xla::cpu::InfoProto& info() const;
  PROTOBUF_NODISCARD ::xla::cpu::InfoProto* release_info();
  ::xla::cpu::InfoProto* mutable_info();
  void set_allocated_info(::xla::cpu::InfoProto* info);
  private:
  const ::xla::cpu::InfoProto& _internal_info() const;
  ::xla::cpu::InfoProto* _internal_mutable_info();
  public:
  void unsafe_arena_set_allocated_info(
      ::xla::cpu::InfoProto* info);
  ::xla::cpu::InfoProto* unsafe_arena_release_info();

  // .xla.cpu.CallThunkProto call_thunk = 3;
  bool has_call_thunk() const;
  private:
  bool _internal_has_call_thunk() const;
  public:
  void clear_call_thunk();
  const ::xla::cpu::CallThunkProto& call_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::CallThunkProto* release_call_thunk();
  ::xla::cpu::CallThunkProto* mutable_call_thunk();
  void set_allocated_call_thunk(::xla::cpu::CallThunkProto* call_thunk);
  private:
  const ::xla::cpu::CallThunkProto& _internal_call_thunk() const;
  ::xla::cpu::CallThunkProto* _internal_mutable_call_thunk();
  public:
  void unsafe_arena_set_allocated_call_thunk(
      ::xla::cpu::CallThunkProto* call_thunk);
  ::xla::cpu::CallThunkProto* unsafe_arena_release_call_thunk();

  // .xla.cpu.ConditionalThunkProto conditional_thunk = 4;
  bool has_conditional_thunk() const;
  private:
  bool _internal_has_conditional_thunk() const;
  public:
  void clear_conditional_thunk();
  const ::xla::cpu::ConditionalThunkProto& conditional_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::ConditionalThunkProto* release_conditional_thunk();
  ::xla::cpu::ConditionalThunkProto* mutable_conditional_thunk();
  void set_allocated_conditional_thunk(::xla::cpu::ConditionalThunkProto* conditional_thunk);
  private:
  const ::xla::cpu::ConditionalThunkProto& _internal_conditional_thunk() const;
  ::xla::cpu::ConditionalThunkProto* _internal_mutable_conditional_thunk();
  public:
  void unsafe_arena_set_allocated_conditional_thunk(
      ::xla::cpu::ConditionalThunkProto* conditional_thunk);
  ::xla::cpu::ConditionalThunkProto* unsafe_arena_release_conditional_thunk();

  // .xla.cpu.SortThunkProto sort_thunk = 5;
  bool has_sort_thunk() const;
  private:
  bool _internal_has_sort_thunk() const;
  public:
  void clear_sort_thunk();
  const ::xla::cpu::SortThunkProto& sort_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::SortThunkProto* release_sort_thunk();
  ::xla::cpu::SortThunkProto* mutable_sort_thunk();
  void set_allocated_sort_thunk(::xla::cpu::SortThunkProto* sort_thunk);
  private:
  const ::xla::cpu::SortThunkProto& _internal_sort_thunk() const;
  ::xla::cpu::SortThunkProto* _internal_mutable_sort_thunk();
  public:
  void unsafe_arena_set_allocated_sort_thunk(
      ::xla::cpu::SortThunkProto* sort_thunk);
  ::xla::cpu::SortThunkProto* unsafe_arena_release_sort_thunk();

  // .xla.cpu.XnnFusionThunkProto xnn_fusion_thunk = 6;
  bool has_xnn_fusion_thunk() const;
  private:
  bool _internal_has_xnn_fusion_thunk() const;
  public:
  void clear_xnn_fusion_thunk();
  const ::xla::cpu::XnnFusionThunkProto& xnn_fusion_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::XnnFusionThunkProto* release_xnn_fusion_thunk();
  ::xla::cpu::XnnFusionThunkProto* mutable_xnn_fusion_thunk();
  void set_allocated_xnn_fusion_thunk(::xla::cpu::XnnFusionThunkProto* xnn_fusion_thunk);
  private:
  const ::xla::cpu::XnnFusionThunkProto& _internal_xnn_fusion_thunk() const;
  ::xla::cpu::XnnFusionThunkProto* _internal_mutable_xnn_fusion_thunk();
  public:
  void unsafe_arena_set_allocated_xnn_fusion_thunk(
      ::xla::cpu::XnnFusionThunkProto* xnn_fusion_thunk);
  ::xla::cpu::XnnFusionThunkProto* unsafe_arena_release_xnn_fusion_thunk();

  // .xla.cpu.DotThunkProto dot_thunk = 7;
  bool has_dot_thunk() const;
  private:
  bool _internal_has_dot_thunk() const;
  public:
  void clear_dot_thunk();
  const ::xla::cpu::DotThunkProto& dot_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::DotThunkProto* release_dot_thunk();
  ::xla::cpu::DotThunkProto* mutable_dot_thunk();
  void set_allocated_dot_thunk(::xla::cpu::DotThunkProto* dot_thunk);
  private:
  const ::xla::cpu::DotThunkProto& _internal_dot_thunk() const;
  ::xla::cpu::DotThunkProto* _internal_mutable_dot_thunk();
  public:
  void unsafe_arena_set_allocated_dot_thunk(
      ::xla::cpu::DotThunkProto* dot_thunk);
  ::xla::cpu::DotThunkProto* unsafe_arena_release_dot_thunk();

  // .xla.cpu.RngGetAndUpdateStateThunkProto rng_get_and_update_state_thunk = 8;
  bool has_rng_get_and_update_state_thunk() const;
  private:
  bool _internal_has_rng_get_and_update_state_thunk() const;
  public:
  void clear_rng_get_and_update_state_thunk();
  const ::xla::cpu::RngGetAndUpdateStateThunkProto& rng_get_and_update_state_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::RngGetAndUpdateStateThunkProto* release_rng_get_and_update_state_thunk();
  ::xla::cpu::RngGetAndUpdateStateThunkProto* mutable_rng_get_and_update_state_thunk();
  void set_allocated_rng_get_and_update_state_thunk(::xla::cpu::RngGetAndUpdateStateThunkProto* rng_get_and_update_state_thunk);
  private:
  const ::xla::cpu::RngGetAndUpdateStateThunkProto& _internal_rng_get_and_update_state_thunk() const;
  ::xla::cpu::RngGetAndUpdateStateThunkProto* _internal_mutable_rng_get_and_update_state_thunk();
  public:
  void unsafe_arena_set_allocated_rng_get_and_update_state_thunk(
      ::xla::cpu::RngGetAndUpdateStateThunkProto* rng_get_and_update_state_thunk);
  ::xla::cpu::RngGetAndUpdateStateThunkProto* unsafe_arena_release_rng_get_and_update_state_thunk();

  // .xla.cpu.TopKThunkProto top_k_thunk = 9;
  bool has_top_k_thunk() const;
  private:
  bool _internal_has_top_k_thunk() const;
  public:
  void clear_top_k_thunk();
  const ::xla::cpu::TopKThunkProto& top_k_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::TopKThunkProto* release_top_k_thunk();
  ::xla::cpu::TopKThunkProto* mutable_top_k_thunk();
  void set_allocated_top_k_thunk(::xla::cpu::TopKThunkProto* top_k_thunk);
  private:
  const ::xla::cpu::TopKThunkProto& _internal_top_k_thunk() const;
  ::xla::cpu::TopKThunkProto* _internal_mutable_top_k_thunk();
  public:
  void unsafe_arena_set_allocated_top_k_thunk(
      ::xla::cpu::TopKThunkProto* top_k_thunk);
  ::xla::cpu::TopKThunkProto* unsafe_arena_release_top_k_thunk();

  // .xla.cpu.WhileThunkProto while_thunk = 10;
  bool has_while_thunk() const;
  private:
  bool _internal_has_while_thunk() const;
  public:
  void clear_while_thunk();
  const ::xla::cpu::WhileThunkProto& while_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::WhileThunkProto* release_while_thunk();
  ::xla::cpu::WhileThunkProto* mutable_while_thunk();
  void set_allocated_while_thunk(::xla::cpu::WhileThunkProto* while_thunk);
  private:
  const ::xla::cpu::WhileThunkProto& _internal_while_thunk() const;
  ::xla::cpu::WhileThunkProto* _internal_mutable_while_thunk();
  public:
  void unsafe_arena_set_allocated_while_thunk(
      ::xla::cpu::WhileThunkProto* while_thunk);
  ::xla::cpu::WhileThunkProto* unsafe_arena_release_while_thunk();

  // .xla.cpu.KernelThunkProto kernel_thunk = 11;
  bool has_kernel_thunk() const;
  private:
  bool _internal_has_kernel_thunk() const;
  public:
  void clear_kernel_thunk();
  const ::xla::cpu::KernelThunkProto& kernel_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::KernelThunkProto* release_kernel_thunk();
  ::xla::cpu::KernelThunkProto* mutable_kernel_thunk();
  void set_allocated_kernel_thunk(::xla::cpu::KernelThunkProto* kernel_thunk);
  private:
  const ::xla::cpu::KernelThunkProto& _internal_kernel_thunk() const;
  ::xla::cpu::KernelThunkProto* _internal_mutable_kernel_thunk();
  public:
  void unsafe_arena_set_allocated_kernel_thunk(
      ::xla::cpu::KernelThunkProto* kernel_thunk);
  ::xla::cpu::KernelThunkProto* unsafe_arena_release_kernel_thunk();

  // .xla.cpu.CopyThunkProto copy_thunk = 12;
  bool has_copy_thunk() const;
  private:
  bool _internal_has_copy_thunk() const;
  public:
  void clear_copy_thunk();
  const ::xla::cpu::CopyThunkProto& copy_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::CopyThunkProto* release_copy_thunk();
  ::xla::cpu::CopyThunkProto* mutable_copy_thunk();
  void set_allocated_copy_thunk(::xla::cpu::CopyThunkProto* copy_thunk);
  private:
  const ::xla::cpu::CopyThunkProto& _internal_copy_thunk() const;
  ::xla::cpu::CopyThunkProto* _internal_mutable_copy_thunk();
  public:
  void unsafe_arena_set_allocated_copy_thunk(
      ::xla::cpu::CopyThunkProto* copy_thunk);
  ::xla::cpu::CopyThunkProto* unsafe_arena_release_copy_thunk();

  // .xla.cpu.FftThunkProto fft_thunk = 13;
  bool has_fft_thunk() const;
  private:
  bool _internal_has_fft_thunk() const;
  public:
  void clear_fft_thunk();
  const ::xla::cpu::FftThunkProto& fft_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::FftThunkProto* release_fft_thunk();
  ::xla::cpu::FftThunkProto* mutable_fft_thunk();
  void set_allocated_fft_thunk(::xla::cpu::FftThunkProto* fft_thunk);
  private:
  const ::xla::cpu::FftThunkProto& _internal_fft_thunk() const;
  ::xla::cpu::FftThunkProto* _internal_mutable_fft_thunk();
  public:
  void unsafe_arena_set_allocated_fft_thunk(
      ::xla::cpu::FftThunkProto* fft_thunk);
  ::xla::cpu::FftThunkProto* unsafe_arena_release_fft_thunk();

  // .xla.cpu.InfeedThunkProto infeed_thunk = 14;
  bool has_infeed_thunk() const;
  private:
  bool _internal_has_infeed_thunk() const;
  public:
  void clear_infeed_thunk();
  const ::xla::cpu::InfeedThunkProto& infeed_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::InfeedThunkProto* release_infeed_thunk();
  ::xla::cpu::InfeedThunkProto* mutable_infeed_thunk();
  void set_allocated_infeed_thunk(::xla::cpu::InfeedThunkProto* infeed_thunk);
  private:
  const ::xla::cpu::InfeedThunkProto& _internal_infeed_thunk() const;
  ::xla::cpu::InfeedThunkProto* _internal_mutable_infeed_thunk();
  public:
  void unsafe_arena_set_allocated_infeed_thunk(
      ::xla::cpu::InfeedThunkProto* infeed_thunk);
  ::xla::cpu::InfeedThunkProto* unsafe_arena_release_infeed_thunk();

  // .xla.cpu.OutfeedThunkProto outfeed_thunk = 15;
  bool has_outfeed_thunk() const;
  private:
  bool _internal_has_outfeed_thunk() const;
  public:
  void clear_outfeed_thunk();
  const ::xla::cpu::OutfeedThunkProto& outfeed_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::OutfeedThunkProto* release_outfeed_thunk();
  ::xla::cpu::OutfeedThunkProto* mutable_outfeed_thunk();
  void set_allocated_outfeed_thunk(::xla::cpu::OutfeedThunkProto* outfeed_thunk);
  private:
  const ::xla::cpu::OutfeedThunkProto& _internal_outfeed_thunk() const;
  ::xla::cpu::OutfeedThunkProto* _internal_mutable_outfeed_thunk();
  public:
  void unsafe_arena_set_allocated_outfeed_thunk(
      ::xla::cpu::OutfeedThunkProto* outfeed_thunk);
  ::xla::cpu::OutfeedThunkProto* unsafe_arena_release_outfeed_thunk();

  // .xla.cpu.CustomCallThunkProto custom_call_thunk = 16;
  bool has_custom_call_thunk() const;
  private:
  bool _internal_has_custom_call_thunk() const;
  public:
  void clear_custom_call_thunk();
  const ::xla::cpu::CustomCallThunkProto& custom_call_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::CustomCallThunkProto* release_custom_call_thunk();
  ::xla::cpu::CustomCallThunkProto* mutable_custom_call_thunk();
  void set_allocated_custom_call_thunk(::xla::cpu::CustomCallThunkProto* custom_call_thunk);
  private:
  const ::xla::cpu::CustomCallThunkProto& _internal_custom_call_thunk() const;
  ::xla::cpu::CustomCallThunkProto* _internal_mutable_custom_call_thunk();
  public:
  void unsafe_arena_set_allocated_custom_call_thunk(
      ::xla::cpu::CustomCallThunkProto* custom_call_thunk);
  ::xla::cpu::CustomCallThunkProto* unsafe_arena_release_custom_call_thunk();

  // .xla.cpu.ConvolutionThunkProto convolution_thunk = 17;
  bool has_convolution_thunk() const;
  private:
  bool _internal_has_convolution_thunk() const;
  public:
  void clear_convolution_thunk();
  const ::xla::cpu::ConvolutionThunkProto& convolution_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::ConvolutionThunkProto* release_convolution_thunk();
  ::xla::cpu::ConvolutionThunkProto* mutable_convolution_thunk();
  void set_allocated_convolution_thunk(::xla::cpu::ConvolutionThunkProto* convolution_thunk);
  private:
  const ::xla::cpu::ConvolutionThunkProto& _internal_convolution_thunk() const;
  ::xla::cpu::ConvolutionThunkProto* _internal_mutable_convolution_thunk();
  public:
  void unsafe_arena_set_allocated_convolution_thunk(
      ::xla::cpu::ConvolutionThunkProto* convolution_thunk);
  ::xla::cpu::ConvolutionThunkProto* unsafe_arena_release_convolution_thunk();

  // .xla.cpu.CollectiveThunkProto collective_thunk = 18;
  bool has_collective_thunk() const;
  private:
  bool _internal_has_collective_thunk() const;
  public:
  void clear_collective_thunk();
  const ::xla::cpu::CollectiveThunkProto& collective_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::CollectiveThunkProto* release_collective_thunk();
  ::xla::cpu::CollectiveThunkProto* mutable_collective_thunk();
  void set_allocated_collective_thunk(::xla::cpu::CollectiveThunkProto* collective_thunk);
  private:
  const ::xla::cpu::CollectiveThunkProto& _internal_collective_thunk() const;
  ::xla::cpu::CollectiveThunkProto* _internal_mutable_collective_thunk();
  public:
  void unsafe_arena_set_allocated_collective_thunk(
      ::xla::cpu::CollectiveThunkProto* collective_thunk);
  ::xla::cpu::CollectiveThunkProto* unsafe_arena_release_collective_thunk();

  // .xla.cpu.PartitionIdThunkProto partition_id_thunk = 19;
  bool has_partition_id_thunk() const;
  private:
  bool _internal_has_partition_id_thunk() const;
  public:
  void clear_partition_id_thunk();
  const ::xla::cpu::PartitionIdThunkProto& partition_id_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::PartitionIdThunkProto* release_partition_id_thunk();
  ::xla::cpu::PartitionIdThunkProto* mutable_partition_id_thunk();
  void set_allocated_partition_id_thunk(::xla::cpu::PartitionIdThunkProto* partition_id_thunk);
  private:
  const ::xla::cpu::PartitionIdThunkProto& _internal_partition_id_thunk() const;
  ::xla::cpu::PartitionIdThunkProto* _internal_mutable_partition_id_thunk();
  public:
  void unsafe_arena_set_allocated_partition_id_thunk(
      ::xla::cpu::PartitionIdThunkProto* partition_id_thunk);
  ::xla::cpu::PartitionIdThunkProto* unsafe_arena_release_partition_id_thunk();

  // .xla.cpu.ReplicaIdThunkProto replica_id_thunk = 20;
  bool has_replica_id_thunk() const;
  private:
  bool _internal_has_replica_id_thunk() const;
  public:
  void clear_replica_id_thunk();
  const ::xla::cpu::ReplicaIdThunkProto& replica_id_thunk() const;
  PROTOBUF_NODISCARD ::xla::cpu::ReplicaIdThunkProto* release_replica_id_thunk();
  ::xla::cpu::ReplicaIdThunkProto* mutable_replica_id_thunk();
  void set_allocated_replica_id_thunk(::xla::cpu::ReplicaIdThunkProto* replica_id_thunk);
  private:
  const ::xla::cpu::ReplicaIdThunkProto& _internal_replica_id_thunk() const;
  ::xla::cpu::ReplicaIdThunkProto* _internal_mutable_replica_id_thunk();
  public:
  void unsafe_arena_set_allocated_replica_id_thunk(
      ::xla::cpu::ReplicaIdThunkProto* replica_id_thunk);
  ::xla::cpu::ReplicaIdThunkProto* unsafe_arena_release_replica_id_thunk();

  void clear_impl();
  ImplCase impl_case() const;
  // @@protoc_insertion_point(class_scope:xla.cpu.ThunkProto)
 private:
  class _Internal;
  void set_has_call_thunk();
  void set_has_conditional_thunk();
  void set_has_sort_thunk();
  void set_has_xnn_fusion_thunk();
  void set_has_dot_thunk();
  void set_has_rng_get_and_update_state_thunk();
  void set_has_top_k_thunk();
  void set_has_while_thunk();
  void set_has_kernel_thunk();
  void set_has_copy_thunk();
  void set_has_fft_thunk();
  void set_has_infeed_thunk();
  void set_has_outfeed_thunk();
  void set_has_custom_call_thunk();
  void set_has_convolution_thunk();
  void set_has_collective_thunk();
  void set_has_partition_id_thunk();
  void set_has_replica_id_thunk();

  inline bool has_impl() const;
  inline void clear_has_impl();

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr kind_;
    ::xla::cpu::InfoProto* info_;
    union ImplUnion {
      constexpr ImplUnion() : _constinit_{} {}
        ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_;
      ::xla::cpu::CallThunkProto* call_thunk_;
      ::xla::cpu::ConditionalThunkProto* conditional_thunk_;
      ::xla::cpu::SortThunkProto* sort_thunk_;
      ::xla::cpu::XnnFusionThunkProto* xnn_fusion_thunk_;
      ::xla::cpu::DotThunkProto* dot_thunk_;
      ::xla::cpu::RngGetAndUpdateStateThunkProto* rng_get_and_update_state_thunk_;
      ::xla::cpu::TopKThunkProto* top_k_thunk_;
      ::xla::cpu::WhileThunkProto* while_thunk_;
      ::xla::cpu::KernelThunkProto* kernel_thunk_;
      ::xla::cpu::CopyThunkProto* copy_thunk_;
      ::xla::cpu::FftThunkProto* fft_thunk_;
      ::xla::cpu::InfeedThunkProto* infeed_thunk_;
      ::xla::cpu::OutfeedThunkProto* outfeed_thunk_;
      ::xla::cpu::CustomCallThunkProto* custom_call_thunk_;
      ::xla::cpu::ConvolutionThunkProto* convolution_thunk_;
      ::xla::cpu::CollectiveThunkProto* collective_thunk_;
      ::xla::cpu::PartitionIdThunkProto* partition_id_thunk_;
      ::xla::cpu::ReplicaIdThunkProto* replica_id_thunk_;
    } impl_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
    uint32_t _oneof_case_[1];

  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// -------------------------------------------------------------------

class ThunkSequenceProto final :
    public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.cpu.ThunkSequenceProto) */ {
 public:
  inline ThunkSequenceProto() : ThunkSequenceProto(nullptr) {}
  ~ThunkSequenceProto() override;
  explicit PROTOBUF_CONSTEXPR ThunkSequenceProto(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized);

  ThunkSequenceProto(const ThunkSequenceProto& from);
  ThunkSequenceProto(ThunkSequenceProto&& from) noexcept
    : ThunkSequenceProto() {
    *this = ::std::move(from);
  }

  inline ThunkSequenceProto& operator=(const ThunkSequenceProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ThunkSequenceProto& operator=(ThunkSequenceProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetOwningArena() == from.GetOwningArena()
  #ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetOwningArena() != nullptr
  #endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ThunkSequenceProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ThunkSequenceProto* internal_default_instance() {
    return reinterpret_cast<const ThunkSequenceProto*>(
               &_ThunkSequenceProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages =
    44;

  friend void swap(ThunkSequenceProto& a, ThunkSequenceProto& b) {
    a.Swap(&b);
  }
  inline void Swap(ThunkSequenceProto* other) {
    if (other == this) return;
  #ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() != nullptr &&
        GetOwningArena() == other->GetOwningArena()) {
   #else  // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetOwningArena() == other->GetOwningArena()) {
  #endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ThunkSequenceProto* other) {
    if (other == this) return;
    GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  ThunkSequenceProto* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final {
    return CreateMaybeMessage<ThunkSequenceProto>(arena);
  }
  using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom;
  void CopyFrom(const ThunkSequenceProto& from);
  using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom;
  void MergeFrom( const ThunkSequenceProto& from) {
    ThunkSequenceProto::MergeImpl(*this, from);
  }
  private:
  static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg);
  public:
  PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final;
  bool IsInitialized() const final;

  size_t ByteSizeLong() const final;
  const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final;
  uint8_t* _InternalSerialize(
      uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final;
  int GetCachedSize() const final { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned);
  void SharedDtor();
  void SetCachedSize(int size) const final;
  void InternalSwap(ThunkSequenceProto* other);

  private:
  friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata;
  static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() {
    return "xla.cpu.ThunkSequenceProto";
  }
  protected:
  explicit ThunkSequenceProto(::PROTOBUF_NAMESPACE_ID::Arena* arena,
                       bool is_message_owned = false);
  public:

  static const ClassData _class_data_;
  const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final;

  ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final;

  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------

  enum : int {
    kThunksFieldNumber = 1,
  };
  // repeated .xla.cpu.ThunkProto thunks = 1;
  int thunks_size() const;
  private:
  int _internal_thunks_size() const;
  public:
  void clear_thunks();
  ::xla::cpu::ThunkProto* mutable_thunks(int index);
  ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkProto >*
      mutable_thunks();
  private:
  const ::xla::cpu::ThunkProto& _internal_thunks(int index) const;
  ::xla::cpu::ThunkProto* _internal_add_thunks();
  public:
  const ::xla::cpu::ThunkProto& thunks(int index) const;
  ::xla::cpu::ThunkProto* add_thunks();
  const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkProto >&
      thunks() const;

  // @@protoc_insertion_point(class_scope:xla.cpu.ThunkSequenceProto)
 private:
  class _Internal;

  template <typename T> friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper;
  typedef void InternalArenaConstructable_;
  typedef void DestructorSkippable_;
  struct Impl_ {
    ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkProto > thunks_;
    mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto;
};
// ===================================================================


// ===================================================================

#ifdef __GNUC__
  #pragma GCC diagnostic push
  #pragma GCC diagnostic ignored "-Wstrict-aliasing"
#endif  // __GNUC__
// ResourceProto

// .xla.cpu.ResourceProto.Kind kind = 1;
inline void ResourceProto::clear_kind() {
  _impl_.kind_ = 0;
}
inline ::xla::cpu::ResourceProto_Kind ResourceProto::_internal_kind() const {
  return static_cast< ::xla::cpu::ResourceProto_Kind >(_impl_.kind_);
}
inline ::xla::cpu::ResourceProto_Kind ResourceProto::kind() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ResourceProto.kind)
  return _internal_kind();
}
inline void ResourceProto::_internal_set_kind(::xla::cpu::ResourceProto_Kind value) {
  
  _impl_.kind_ = value;
}
inline void ResourceProto::set_kind(::xla::cpu::ResourceProto_Kind value) {
  _internal_set_kind(value);
  // @@protoc_insertion_point(field_set:xla.cpu.ResourceProto.kind)
}

// -------------------------------------------------------------------

// BoolOptional

// bool value = 1;
inline void BoolOptional::clear_value() {
  _impl_.value_ = false;
}
inline bool BoolOptional::_internal_value() const {
  return _impl_.value_;
}
inline bool BoolOptional::value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.BoolOptional.value)
  return _internal_value();
}
inline void BoolOptional::_internal_set_value(bool value) {
  
  _impl_.value_ = value;
}
inline void BoolOptional::set_value(bool value) {
  _internal_set_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.BoolOptional.value)
}

// bool contains_value = 2;
inline void BoolOptional::clear_contains_value() {
  _impl_.contains_value_ = false;
}
inline bool BoolOptional::_internal_contains_value() const {
  return _impl_.contains_value_;
}
inline bool BoolOptional::contains_value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.BoolOptional.contains_value)
  return _internal_contains_value();
}
inline void BoolOptional::_internal_set_contains_value(bool value) {
  
  _impl_.contains_value_ = value;
}
inline void BoolOptional::set_contains_value(bool value) {
  _internal_set_contains_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.BoolOptional.contains_value)
}

// -------------------------------------------------------------------

// ResourceOptional

// .xla.cpu.ResourceProto value = 1;
inline bool ResourceOptional::_internal_has_value() const {
  return this != internal_default_instance() && _impl_.value_ != nullptr;
}
inline bool ResourceOptional::has_value() const {
  return _internal_has_value();
}
inline void ResourceOptional::clear_value() {
  if (GetArenaForAllocation() == nullptr && _impl_.value_ != nullptr) {
    delete _impl_.value_;
  }
  _impl_.value_ = nullptr;
}
inline const ::xla::cpu::ResourceProto& ResourceOptional::_internal_value() const {
  const ::xla::cpu::ResourceProto* p = _impl_.value_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceProto&>(
      ::xla::cpu::_ResourceProto_default_instance_);
}
inline const ::xla::cpu::ResourceProto& ResourceOptional::value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ResourceOptional.value)
  return _internal_value();
}
inline void ResourceOptional::unsafe_arena_set_allocated_value(
    ::xla::cpu::ResourceProto* value) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.value_);
  }
  _impl_.value_ = value;
  if (value) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ResourceOptional.value)
}
inline ::xla::cpu::ResourceProto* ResourceOptional::release_value() {
  
  ::xla::cpu::ResourceProto* temp = _impl_.value_;
  _impl_.value_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceProto* ResourceOptional::unsafe_arena_release_value() {
  // @@protoc_insertion_point(field_release:xla.cpu.ResourceOptional.value)
  
  ::xla::cpu::ResourceProto* temp = _impl_.value_;
  _impl_.value_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceProto* ResourceOptional::_internal_mutable_value() {
  
  if (_impl_.value_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceProto>(GetArenaForAllocation());
    _impl_.value_ = p;
  }
  return _impl_.value_;
}
inline ::xla::cpu::ResourceProto* ResourceOptional::mutable_value() {
  ::xla::cpu::ResourceProto* _msg = _internal_mutable_value();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ResourceOptional.value)
  return _msg;
}
inline void ResourceOptional::set_allocated_value(::xla::cpu::ResourceProto* value) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.value_;
  }
  if (value) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(value);
    if (message_arena != submessage_arena) {
      value = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, value, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.value_ = value;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ResourceOptional.value)
}

// bool contains_value = 2;
inline void ResourceOptional::clear_contains_value() {
  _impl_.contains_value_ = false;
}
inline bool ResourceOptional::_internal_contains_value() const {
  return _impl_.contains_value_;
}
inline bool ResourceOptional::contains_value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ResourceOptional.contains_value)
  return _internal_contains_value();
}
inline void ResourceOptional::_internal_set_contains_value(bool value) {
  
  _impl_.contains_value_ = value;
}
inline void ResourceOptional::set_contains_value(bool value) {
  _internal_set_contains_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.ResourceOptional.contains_value)
}

// -------------------------------------------------------------------

// Int64Optional

// int64 value = 1;
inline void Int64Optional::clear_value() {
  _impl_.value_ = int64_t{0};
}
inline int64_t Int64Optional::_internal_value() const {
  return _impl_.value_;
}
inline int64_t Int64Optional::value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.Int64Optional.value)
  return _internal_value();
}
inline void Int64Optional::_internal_set_value(int64_t value) {
  
  _impl_.value_ = value;
}
inline void Int64Optional::set_value(int64_t value) {
  _internal_set_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.Int64Optional.value)
}

// bool contains_value = 2;
inline void Int64Optional::clear_contains_value() {
  _impl_.contains_value_ = false;
}
inline bool Int64Optional::_internal_contains_value() const {
  return _impl_.contains_value_;
}
inline bool Int64Optional::contains_value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.Int64Optional.contains_value)
  return _internal_contains_value();
}
inline void Int64Optional::_internal_set_contains_value(bool value) {
  
  _impl_.contains_value_ = value;
}
inline void Int64Optional::set_contains_value(bool value) {
  _internal_set_contains_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.Int64Optional.contains_value)
}

// -------------------------------------------------------------------

// SortDirectionOptional

// .xla.cpu.SortDirectionProto value = 1;
inline void SortDirectionOptional::clear_value() {
  _impl_.value_ = 0;
}
inline ::xla::cpu::SortDirectionProto SortDirectionOptional::_internal_value() const {
  return static_cast< ::xla::cpu::SortDirectionProto >(_impl_.value_);
}
inline ::xla::cpu::SortDirectionProto SortDirectionOptional::value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortDirectionOptional.value)
  return _internal_value();
}
inline void SortDirectionOptional::_internal_set_value(::xla::cpu::SortDirectionProto value) {
  
  _impl_.value_ = value;
}
inline void SortDirectionOptional::set_value(::xla::cpu::SortDirectionProto value) {
  _internal_set_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.SortDirectionOptional.value)
}

// bool contains_value = 2;
inline void SortDirectionOptional::clear_contains_value() {
  _impl_.contains_value_ = false;
}
inline bool SortDirectionOptional::_internal_contains_value() const {
  return _impl_.contains_value_;
}
inline bool SortDirectionOptional::contains_value() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortDirectionOptional.contains_value)
  return _internal_contains_value();
}
inline void SortDirectionOptional::_internal_set_contains_value(bool value) {
  
  _impl_.contains_value_ = value;
}
inline void SortDirectionOptional::set_contains_value(bool value) {
  _internal_set_contains_value(value);
  // @@protoc_insertion_point(field_set:xla.cpu.SortDirectionOptional.contains_value)
}

// -------------------------------------------------------------------

// BufferAllocationSliceProto

// int64 offset = 1;
inline void BufferAllocationSliceProto::clear_offset() {
  _impl_.offset_ = int64_t{0};
}
inline int64_t BufferAllocationSliceProto::_internal_offset() const {
  return _impl_.offset_;
}
inline int64_t BufferAllocationSliceProto::offset() const {
  // @@protoc_insertion_point(field_get:xla.cpu.BufferAllocationSliceProto.offset)
  return _internal_offset();
}
inline void BufferAllocationSliceProto::_internal_set_offset(int64_t value) {
  
  _impl_.offset_ = value;
}
inline void BufferAllocationSliceProto::set_offset(int64_t value) {
  _internal_set_offset(value);
  // @@protoc_insertion_point(field_set:xla.cpu.BufferAllocationSliceProto.offset)
}

// int64 size = 2;
inline void BufferAllocationSliceProto::clear_size() {
  _impl_.size_ = int64_t{0};
}
inline int64_t BufferAllocationSliceProto::_internal_size() const {
  return _impl_.size_;
}
inline int64_t BufferAllocationSliceProto::size() const {
  // @@protoc_insertion_point(field_get:xla.cpu.BufferAllocationSliceProto.size)
  return _internal_size();
}
inline void BufferAllocationSliceProto::_internal_set_size(int64_t value) {
  
  _impl_.size_ = value;
}
inline void BufferAllocationSliceProto::set_size(int64_t value) {
  _internal_set_size(value);
  // @@protoc_insertion_point(field_set:xla.cpu.BufferAllocationSliceProto.size)
}

// int64 buffer_allocation_index = 3;
inline void BufferAllocationSliceProto::clear_buffer_allocation_index() {
  _impl_.buffer_allocation_index_ = int64_t{0};
}
inline int64_t BufferAllocationSliceProto::_internal_buffer_allocation_index() const {
  return _impl_.buffer_allocation_index_;
}
inline int64_t BufferAllocationSliceProto::buffer_allocation_index() const {
  // @@protoc_insertion_point(field_get:xla.cpu.BufferAllocationSliceProto.buffer_allocation_index)
  return _internal_buffer_allocation_index();
}
inline void BufferAllocationSliceProto::_internal_set_buffer_allocation_index(int64_t value) {
  
  _impl_.buffer_allocation_index_ = value;
}
inline void BufferAllocationSliceProto::set_buffer_allocation_index(int64_t value) {
  _internal_set_buffer_allocation_index(value);
  // @@protoc_insertion_point(field_set:xla.cpu.BufferAllocationSliceProto.buffer_allocation_index)
}

// -------------------------------------------------------------------

// ShapeBufferAllocationSliceProto

// .xla.ShapeProto shape = 1;
inline bool ShapeBufferAllocationSliceProto::_internal_has_shape() const {
  return this != internal_default_instance() && _impl_.shape_ != nullptr;
}
inline bool ShapeBufferAllocationSliceProto::has_shape() const {
  return _internal_has_shape();
}
inline const ::xla::ShapeProto& ShapeBufferAllocationSliceProto::_internal_shape() const {
  const ::xla::ShapeProto* p = _impl_.shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(
      ::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& ShapeBufferAllocationSliceProto::shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ShapeBufferAllocationSliceProto.shape)
  return _internal_shape();
}
inline void ShapeBufferAllocationSliceProto::unsafe_arena_set_allocated_shape(
    ::xla::ShapeProto* shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.shape_);
  }
  _impl_.shape_ = shape;
  if (shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ShapeBufferAllocationSliceProto.shape)
}
inline ::xla::ShapeProto* ShapeBufferAllocationSliceProto::release_shape() {
  
  ::xla::ShapeProto* temp = _impl_.shape_;
  _impl_.shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::ShapeProto* ShapeBufferAllocationSliceProto::unsafe_arena_release_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.ShapeBufferAllocationSliceProto.shape)
  
  ::xla::ShapeProto* temp = _impl_.shape_;
  _impl_.shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* ShapeBufferAllocationSliceProto::_internal_mutable_shape() {
  
  if (_impl_.shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::ShapeProto>(GetArenaForAllocation());
    _impl_.shape_ = p;
  }
  return _impl_.shape_;
}
inline ::xla::ShapeProto* ShapeBufferAllocationSliceProto::mutable_shape() {
  ::xla::ShapeProto* _msg = _internal_mutable_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ShapeBufferAllocationSliceProto.shape)
  return _msg;
}
inline void ShapeBufferAllocationSliceProto::set_allocated_shape(::xla::ShapeProto* shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.shape_);
  }
  if (shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(shape));
    if (message_arena != submessage_arena) {
      shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.shape_ = shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ShapeBufferAllocationSliceProto.shape)
}

// .xla.cpu.BufferAllocationSliceProto slice = 2;
inline bool ShapeBufferAllocationSliceProto::_internal_has_slice() const {
  return this != internal_default_instance() && _impl_.slice_ != nullptr;
}
inline bool ShapeBufferAllocationSliceProto::has_slice() const {
  return _internal_has_slice();
}
inline void ShapeBufferAllocationSliceProto::clear_slice() {
  if (GetArenaForAllocation() == nullptr && _impl_.slice_ != nullptr) {
    delete _impl_.slice_;
  }
  _impl_.slice_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& ShapeBufferAllocationSliceProto::_internal_slice() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.slice_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& ShapeBufferAllocationSliceProto::slice() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ShapeBufferAllocationSliceProto.slice)
  return _internal_slice();
}
inline void ShapeBufferAllocationSliceProto::unsafe_arena_set_allocated_slice(
    ::xla::cpu::BufferAllocationSliceProto* slice) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.slice_);
  }
  _impl_.slice_ = slice;
  if (slice) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ShapeBufferAllocationSliceProto.slice)
}
inline ::xla::cpu::BufferAllocationSliceProto* ShapeBufferAllocationSliceProto::release_slice() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.slice_;
  _impl_.slice_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ShapeBufferAllocationSliceProto::unsafe_arena_release_slice() {
  // @@protoc_insertion_point(field_release:xla.cpu.ShapeBufferAllocationSliceProto.slice)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.slice_;
  _impl_.slice_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ShapeBufferAllocationSliceProto::_internal_mutable_slice() {
  
  if (_impl_.slice_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.slice_ = p;
  }
  return _impl_.slice_;
}
inline ::xla::cpu::BufferAllocationSliceProto* ShapeBufferAllocationSliceProto::mutable_slice() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_slice();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ShapeBufferAllocationSliceProto.slice)
  return _msg;
}
inline void ShapeBufferAllocationSliceProto::set_allocated_slice(::xla::cpu::BufferAllocationSliceProto* slice) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.slice_;
  }
  if (slice) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(slice);
    if (message_arena != submessage_arena) {
      slice = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, slice, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.slice_ = slice;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ShapeBufferAllocationSliceProto.slice)
}

// -------------------------------------------------------------------

// OpParamsProto

// int64 op_id = 1;
inline void OpParamsProto::clear_op_id() {
  _impl_.op_id_ = int64_t{0};
}
inline int64_t OpParamsProto::_internal_op_id() const {
  return _impl_.op_id_;
}
inline int64_t OpParamsProto::op_id() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpParamsProto.op_id)
  return _internal_op_id();
}
inline void OpParamsProto::_internal_set_op_id(int64_t value) {
  
  _impl_.op_id_ = value;
}
inline void OpParamsProto::set_op_id(int64_t value) {
  _internal_set_op_id(value);
  // @@protoc_insertion_point(field_set:xla.cpu.OpParamsProto.op_id)
}

// bool has_channel_id = 2;
inline void OpParamsProto::clear_has_channel_id() {
  _impl_.has_channel_id_ = false;
}
inline bool OpParamsProto::_internal_has_channel_id() const {
  return _impl_.has_channel_id_;
}
inline bool OpParamsProto::has_channel_id() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpParamsProto.has_channel_id)
  return _internal_has_channel_id();
}
inline void OpParamsProto::_internal_set_has_channel_id(bool value) {
  
  _impl_.has_channel_id_ = value;
}
inline void OpParamsProto::set_has_channel_id(bool value) {
  _internal_set_has_channel_id(value);
  // @@protoc_insertion_point(field_set:xla.cpu.OpParamsProto.has_channel_id)
}

// .xla.cpu.BoolOptional use_global_device_ids = 3;
inline bool OpParamsProto::_internal_has_use_global_device_ids() const {
  return this != internal_default_instance() && _impl_.use_global_device_ids_ != nullptr;
}
inline bool OpParamsProto::has_use_global_device_ids() const {
  return _internal_has_use_global_device_ids();
}
inline void OpParamsProto::clear_use_global_device_ids() {
  if (GetArenaForAllocation() == nullptr && _impl_.use_global_device_ids_ != nullptr) {
    delete _impl_.use_global_device_ids_;
  }
  _impl_.use_global_device_ids_ = nullptr;
}
inline const ::xla::cpu::BoolOptional& OpParamsProto::_internal_use_global_device_ids() const {
  const ::xla::cpu::BoolOptional* p = _impl_.use_global_device_ids_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BoolOptional&>(
      ::xla::cpu::_BoolOptional_default_instance_);
}
inline const ::xla::cpu::BoolOptional& OpParamsProto::use_global_device_ids() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpParamsProto.use_global_device_ids)
  return _internal_use_global_device_ids();
}
inline void OpParamsProto::unsafe_arena_set_allocated_use_global_device_ids(
    ::xla::cpu::BoolOptional* use_global_device_ids) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.use_global_device_ids_);
  }
  _impl_.use_global_device_ids_ = use_global_device_ids;
  if (use_global_device_ids) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.OpParamsProto.use_global_device_ids)
}
inline ::xla::cpu::BoolOptional* OpParamsProto::release_use_global_device_ids() {
  
  ::xla::cpu::BoolOptional* temp = _impl_.use_global_device_ids_;
  _impl_.use_global_device_ids_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BoolOptional* OpParamsProto::unsafe_arena_release_use_global_device_ids() {
  // @@protoc_insertion_point(field_release:xla.cpu.OpParamsProto.use_global_device_ids)
  
  ::xla::cpu::BoolOptional* temp = _impl_.use_global_device_ids_;
  _impl_.use_global_device_ids_ = nullptr;
  return temp;
}
inline ::xla::cpu::BoolOptional* OpParamsProto::_internal_mutable_use_global_device_ids() {
  
  if (_impl_.use_global_device_ids_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BoolOptional>(GetArenaForAllocation());
    _impl_.use_global_device_ids_ = p;
  }
  return _impl_.use_global_device_ids_;
}
inline ::xla::cpu::BoolOptional* OpParamsProto::mutable_use_global_device_ids() {
  ::xla::cpu::BoolOptional* _msg = _internal_mutable_use_global_device_ids();
  // @@protoc_insertion_point(field_mutable:xla.cpu.OpParamsProto.use_global_device_ids)
  return _msg;
}
inline void OpParamsProto::set_allocated_use_global_device_ids(::xla::cpu::BoolOptional* use_global_device_ids) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.use_global_device_ids_;
  }
  if (use_global_device_ids) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(use_global_device_ids);
    if (message_arena != submessage_arena) {
      use_global_device_ids = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, use_global_device_ids, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.use_global_device_ids_ = use_global_device_ids;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.OpParamsProto.use_global_device_ids)
}

// repeated .xla.ReplicaGroup replica_group = 4;
inline int OpParamsProto::_internal_replica_group_size() const {
  return _impl_.replica_group_.size();
}
inline int OpParamsProto::replica_group_size() const {
  return _internal_replica_group_size();
}
inline ::xla::ReplicaGroup* OpParamsProto::mutable_replica_group(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.OpParamsProto.replica_group)
  return _impl_.replica_group_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::ReplicaGroup >*
OpParamsProto::mutable_replica_group() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.OpParamsProto.replica_group)
  return &_impl_.replica_group_;
}
inline const ::xla::ReplicaGroup& OpParamsProto::_internal_replica_group(int index) const {
  return _impl_.replica_group_.Get(index);
}
inline const ::xla::ReplicaGroup& OpParamsProto::replica_group(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpParamsProto.replica_group)
  return _internal_replica_group(index);
}
inline ::xla::ReplicaGroup* OpParamsProto::_internal_add_replica_group() {
  return _impl_.replica_group_.Add();
}
inline ::xla::ReplicaGroup* OpParamsProto::add_replica_group() {
  ::xla::ReplicaGroup* _add = _internal_add_replica_group();
  // @@protoc_insertion_point(field_add:xla.cpu.OpParamsProto.replica_group)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::ReplicaGroup >&
OpParamsProto::replica_group() const {
  // @@protoc_insertion_point(field_list:xla.cpu.OpParamsProto.replica_group)
  return _impl_.replica_group_;
}

// -------------------------------------------------------------------

// OpBuffersProto

// repeated .xla.cpu.ShapeBufferAllocationSliceProto source_shapes_buffer_slices = 1;
inline int OpBuffersProto::_internal_source_shapes_buffer_slices_size() const {
  return _impl_.source_shapes_buffer_slices_.size();
}
inline int OpBuffersProto::source_shapes_buffer_slices_size() const {
  return _internal_source_shapes_buffer_slices_size();
}
inline void OpBuffersProto::clear_source_shapes_buffer_slices() {
  _impl_.source_shapes_buffer_slices_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::mutable_source_shapes_buffer_slices(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.OpBuffersProto.source_shapes_buffer_slices)
  return _impl_.source_shapes_buffer_slices_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
OpBuffersProto::mutable_source_shapes_buffer_slices() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.OpBuffersProto.source_shapes_buffer_slices)
  return &_impl_.source_shapes_buffer_slices_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OpBuffersProto::_internal_source_shapes_buffer_slices(int index) const {
  return _impl_.source_shapes_buffer_slices_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OpBuffersProto::source_shapes_buffer_slices(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpBuffersProto.source_shapes_buffer_slices)
  return _internal_source_shapes_buffer_slices(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::_internal_add_source_shapes_buffer_slices() {
  return _impl_.source_shapes_buffer_slices_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::add_source_shapes_buffer_slices() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_source_shapes_buffer_slices();
  // @@protoc_insertion_point(field_add:xla.cpu.OpBuffersProto.source_shapes_buffer_slices)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
OpBuffersProto::source_shapes_buffer_slices() const {
  // @@protoc_insertion_point(field_list:xla.cpu.OpBuffersProto.source_shapes_buffer_slices)
  return _impl_.source_shapes_buffer_slices_;
}

// repeated .xla.cpu.ShapeBufferAllocationSliceProto destination_shapes_buffer_slices = 2;
inline int OpBuffersProto::_internal_destination_shapes_buffer_slices_size() const {
  return _impl_.destination_shapes_buffer_slices_.size();
}
inline int OpBuffersProto::destination_shapes_buffer_slices_size() const {
  return _internal_destination_shapes_buffer_slices_size();
}
inline void OpBuffersProto::clear_destination_shapes_buffer_slices() {
  _impl_.destination_shapes_buffer_slices_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::mutable_destination_shapes_buffer_slices(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.OpBuffersProto.destination_shapes_buffer_slices)
  return _impl_.destination_shapes_buffer_slices_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
OpBuffersProto::mutable_destination_shapes_buffer_slices() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.OpBuffersProto.destination_shapes_buffer_slices)
  return &_impl_.destination_shapes_buffer_slices_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OpBuffersProto::_internal_destination_shapes_buffer_slices(int index) const {
  return _impl_.destination_shapes_buffer_slices_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OpBuffersProto::destination_shapes_buffer_slices(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpBuffersProto.destination_shapes_buffer_slices)
  return _internal_destination_shapes_buffer_slices(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::_internal_add_destination_shapes_buffer_slices() {
  return _impl_.destination_shapes_buffer_slices_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OpBuffersProto::add_destination_shapes_buffer_slices() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_destination_shapes_buffer_slices();
  // @@protoc_insertion_point(field_add:xla.cpu.OpBuffersProto.destination_shapes_buffer_slices)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
OpBuffersProto::destination_shapes_buffer_slices() const {
  // @@protoc_insertion_point(field_list:xla.cpu.OpBuffersProto.destination_shapes_buffer_slices)
  return _impl_.destination_shapes_buffer_slices_;
}

// -------------------------------------------------------------------

// OpResourcesProto

// .xla.cpu.ResourceOptional communicator_resource = 1;
inline bool OpResourcesProto::_internal_has_communicator_resource() const {
  return this != internal_default_instance() && _impl_.communicator_resource_ != nullptr;
}
inline bool OpResourcesProto::has_communicator_resource() const {
  return _internal_has_communicator_resource();
}
inline void OpResourcesProto::clear_communicator_resource() {
  if (GetArenaForAllocation() == nullptr && _impl_.communicator_resource_ != nullptr) {
    delete _impl_.communicator_resource_;
  }
  _impl_.communicator_resource_ = nullptr;
}
inline const ::xla::cpu::ResourceOptional& OpResourcesProto::_internal_communicator_resource() const {
  const ::xla::cpu::ResourceOptional* p = _impl_.communicator_resource_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceOptional&>(
      ::xla::cpu::_ResourceOptional_default_instance_);
}
inline const ::xla::cpu::ResourceOptional& OpResourcesProto::communicator_resource() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OpResourcesProto.communicator_resource)
  return _internal_communicator_resource();
}
inline void OpResourcesProto::unsafe_arena_set_allocated_communicator_resource(
    ::xla::cpu::ResourceOptional* communicator_resource) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.communicator_resource_);
  }
  _impl_.communicator_resource_ = communicator_resource;
  if (communicator_resource) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.OpResourcesProto.communicator_resource)
}
inline ::xla::cpu::ResourceOptional* OpResourcesProto::release_communicator_resource() {
  
  ::xla::cpu::ResourceOptional* temp = _impl_.communicator_resource_;
  _impl_.communicator_resource_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceOptional* OpResourcesProto::unsafe_arena_release_communicator_resource() {
  // @@protoc_insertion_point(field_release:xla.cpu.OpResourcesProto.communicator_resource)
  
  ::xla::cpu::ResourceOptional* temp = _impl_.communicator_resource_;
  _impl_.communicator_resource_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceOptional* OpResourcesProto::_internal_mutable_communicator_resource() {
  
  if (_impl_.communicator_resource_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceOptional>(GetArenaForAllocation());
    _impl_.communicator_resource_ = p;
  }
  return _impl_.communicator_resource_;
}
inline ::xla::cpu::ResourceOptional* OpResourcesProto::mutable_communicator_resource() {
  ::xla::cpu::ResourceOptional* _msg = _internal_mutable_communicator_resource();
  // @@protoc_insertion_point(field_mutable:xla.cpu.OpResourcesProto.communicator_resource)
  return _msg;
}
inline void OpResourcesProto::set_allocated_communicator_resource(::xla::cpu::ResourceOptional* communicator_resource) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.communicator_resource_;
  }
  if (communicator_resource) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(communicator_resource);
    if (message_arena != submessage_arena) {
      communicator_resource = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, communicator_resource, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.communicator_resource_ = communicator_resource;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.OpResourcesProto.communicator_resource)
}

// -------------------------------------------------------------------

// AllGatherThunkProto

// -------------------------------------------------------------------

// AllReduceThunkProto

// string reduction_kind = 1;
inline void AllReduceThunkProto::clear_reduction_kind() {
  _impl_.reduction_kind_.ClearToEmpty();
}
inline const std::string& AllReduceThunkProto::reduction_kind() const {
  // @@protoc_insertion_point(field_get:xla.cpu.AllReduceThunkProto.reduction_kind)
  return _internal_reduction_kind();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void AllReduceThunkProto::set_reduction_kind(ArgT0&& arg0, ArgT... args) {
 
 _impl_.reduction_kind_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.AllReduceThunkProto.reduction_kind)
}
inline std::string* AllReduceThunkProto::mutable_reduction_kind() {
  std::string* _s = _internal_mutable_reduction_kind();
  // @@protoc_insertion_point(field_mutable:xla.cpu.AllReduceThunkProto.reduction_kind)
  return _s;
}
inline const std::string& AllReduceThunkProto::_internal_reduction_kind() const {
  return _impl_.reduction_kind_.Get();
}
inline void AllReduceThunkProto::_internal_set_reduction_kind(const std::string& value) {
  
  _impl_.reduction_kind_.Set(value, GetArenaForAllocation());
}
inline std::string* AllReduceThunkProto::_internal_mutable_reduction_kind() {
  
  return _impl_.reduction_kind_.Mutable(GetArenaForAllocation());
}
inline std::string* AllReduceThunkProto::release_reduction_kind() {
  // @@protoc_insertion_point(field_release:xla.cpu.AllReduceThunkProto.reduction_kind)
  return _impl_.reduction_kind_.Release();
}
inline void AllReduceThunkProto::set_allocated_reduction_kind(std::string* reduction_kind) {
  if (reduction_kind != nullptr) {
    
  } else {
    
  }
  _impl_.reduction_kind_.SetAllocated(reduction_kind, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.reduction_kind_.IsDefault()) {
    _impl_.reduction_kind_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.AllReduceThunkProto.reduction_kind)
}

// bool single_replica = 2;
inline void AllReduceThunkProto::clear_single_replica() {
  _impl_.single_replica_ = false;
}
inline bool AllReduceThunkProto::_internal_single_replica() const {
  return _impl_.single_replica_;
}
inline bool AllReduceThunkProto::single_replica() const {
  // @@protoc_insertion_point(field_get:xla.cpu.AllReduceThunkProto.single_replica)
  return _internal_single_replica();
}
inline void AllReduceThunkProto::_internal_set_single_replica(bool value) {
  
  _impl_.single_replica_ = value;
}
inline void AllReduceThunkProto::set_single_replica(bool value) {
  _internal_set_single_replica(value);
  // @@protoc_insertion_point(field_set:xla.cpu.AllReduceThunkProto.single_replica)
}

// -------------------------------------------------------------------

// AllToAllThunkProto

// -------------------------------------------------------------------

// ReduceScatterThunkProto

// string reduction_kind = 1;
inline void ReduceScatterThunkProto::clear_reduction_kind() {
  _impl_.reduction_kind_.ClearToEmpty();
}
inline const std::string& ReduceScatterThunkProto::reduction_kind() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ReduceScatterThunkProto.reduction_kind)
  return _internal_reduction_kind();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void ReduceScatterThunkProto::set_reduction_kind(ArgT0&& arg0, ArgT... args) {
 
 _impl_.reduction_kind_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.ReduceScatterThunkProto.reduction_kind)
}
inline std::string* ReduceScatterThunkProto::mutable_reduction_kind() {
  std::string* _s = _internal_mutable_reduction_kind();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ReduceScatterThunkProto.reduction_kind)
  return _s;
}
inline const std::string& ReduceScatterThunkProto::_internal_reduction_kind() const {
  return _impl_.reduction_kind_.Get();
}
inline void ReduceScatterThunkProto::_internal_set_reduction_kind(const std::string& value) {
  
  _impl_.reduction_kind_.Set(value, GetArenaForAllocation());
}
inline std::string* ReduceScatterThunkProto::_internal_mutable_reduction_kind() {
  
  return _impl_.reduction_kind_.Mutable(GetArenaForAllocation());
}
inline std::string* ReduceScatterThunkProto::release_reduction_kind() {
  // @@protoc_insertion_point(field_release:xla.cpu.ReduceScatterThunkProto.reduction_kind)
  return _impl_.reduction_kind_.Release();
}
inline void ReduceScatterThunkProto::set_allocated_reduction_kind(std::string* reduction_kind) {
  if (reduction_kind != nullptr) {
    
  } else {
    
  }
  _impl_.reduction_kind_.SetAllocated(reduction_kind, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.reduction_kind_.IsDefault()) {
    _impl_.reduction_kind_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ReduceScatterThunkProto.reduction_kind)
}

// -------------------------------------------------------------------

// CollectivePermuteThunkProto_SourceTargetPairProto

// int64 source = 1;
inline void CollectivePermuteThunkProto_SourceTargetPairProto::clear_source() {
  _impl_.source_ = int64_t{0};
}
inline int64_t CollectivePermuteThunkProto_SourceTargetPairProto::_internal_source() const {
  return _impl_.source_;
}
inline int64_t CollectivePermuteThunkProto_SourceTargetPairProto::source() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto.source)
  return _internal_source();
}
inline void CollectivePermuteThunkProto_SourceTargetPairProto::_internal_set_source(int64_t value) {
  
  _impl_.source_ = value;
}
inline void CollectivePermuteThunkProto_SourceTargetPairProto::set_source(int64_t value) {
  _internal_set_source(value);
  // @@protoc_insertion_point(field_set:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto.source)
}

// int64 target = 2;
inline void CollectivePermuteThunkProto_SourceTargetPairProto::clear_target() {
  _impl_.target_ = int64_t{0};
}
inline int64_t CollectivePermuteThunkProto_SourceTargetPairProto::_internal_target() const {
  return _impl_.target_;
}
inline int64_t CollectivePermuteThunkProto_SourceTargetPairProto::target() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto.target)
  return _internal_target();
}
inline void CollectivePermuteThunkProto_SourceTargetPairProto::_internal_set_target(int64_t value) {
  
  _impl_.target_ = value;
}
inline void CollectivePermuteThunkProto_SourceTargetPairProto::set_target(int64_t value) {
  _internal_set_target(value);
  // @@protoc_insertion_point(field_set:xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto.target)
}

// -------------------------------------------------------------------

// CollectivePermuteThunkProto

// repeated .xla.cpu.CollectivePermuteThunkProto.SourceTargetPairProto source_target_pairs = 1;
inline int CollectivePermuteThunkProto::_internal_source_target_pairs_size() const {
  return _impl_.source_target_pairs_.size();
}
inline int CollectivePermuteThunkProto::source_target_pairs_size() const {
  return _internal_source_target_pairs_size();
}
inline void CollectivePermuteThunkProto::clear_source_target_pairs() {
  _impl_.source_target_pairs_.Clear();
}
inline ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* CollectivePermuteThunkProto::mutable_source_target_pairs(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectivePermuteThunkProto.source_target_pairs)
  return _impl_.source_target_pairs_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto >*
CollectivePermuteThunkProto::mutable_source_target_pairs() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.CollectivePermuteThunkProto.source_target_pairs)
  return &_impl_.source_target_pairs_;
}
inline const ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto& CollectivePermuteThunkProto::_internal_source_target_pairs(int index) const {
  return _impl_.source_target_pairs_.Get(index);
}
inline const ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto& CollectivePermuteThunkProto::source_target_pairs(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectivePermuteThunkProto.source_target_pairs)
  return _internal_source_target_pairs(index);
}
inline ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* CollectivePermuteThunkProto::_internal_add_source_target_pairs() {
  return _impl_.source_target_pairs_.Add();
}
inline ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* CollectivePermuteThunkProto::add_source_target_pairs() {
  ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto* _add = _internal_add_source_target_pairs();
  // @@protoc_insertion_point(field_add:xla.cpu.CollectivePermuteThunkProto.source_target_pairs)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::CollectivePermuteThunkProto_SourceTargetPairProto >&
CollectivePermuteThunkProto::source_target_pairs() const {
  // @@protoc_insertion_point(field_list:xla.cpu.CollectivePermuteThunkProto.source_target_pairs)
  return _impl_.source_target_pairs_;
}

// -------------------------------------------------------------------

// CollectiveThunkProto

// .xla.cpu.OpParamsProto op_params = 1;
inline bool CollectiveThunkProto::_internal_has_op_params() const {
  return this != internal_default_instance() && _impl_.op_params_ != nullptr;
}
inline bool CollectiveThunkProto::has_op_params() const {
  return _internal_has_op_params();
}
inline void CollectiveThunkProto::clear_op_params() {
  if (GetArenaForAllocation() == nullptr && _impl_.op_params_ != nullptr) {
    delete _impl_.op_params_;
  }
  _impl_.op_params_ = nullptr;
}
inline const ::xla::cpu::OpParamsProto& CollectiveThunkProto::_internal_op_params() const {
  const ::xla::cpu::OpParamsProto* p = _impl_.op_params_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::OpParamsProto&>(
      ::xla::cpu::_OpParamsProto_default_instance_);
}
inline const ::xla::cpu::OpParamsProto& CollectiveThunkProto::op_params() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.op_params)
  return _internal_op_params();
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_op_params(
    ::xla::cpu::OpParamsProto* op_params) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.op_params_);
  }
  _impl_.op_params_ = op_params;
  if (op_params) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.op_params)
}
inline ::xla::cpu::OpParamsProto* CollectiveThunkProto::release_op_params() {
  
  ::xla::cpu::OpParamsProto* temp = _impl_.op_params_;
  _impl_.op_params_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::OpParamsProto* CollectiveThunkProto::unsafe_arena_release_op_params() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.op_params)
  
  ::xla::cpu::OpParamsProto* temp = _impl_.op_params_;
  _impl_.op_params_ = nullptr;
  return temp;
}
inline ::xla::cpu::OpParamsProto* CollectiveThunkProto::_internal_mutable_op_params() {
  
  if (_impl_.op_params_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::OpParamsProto>(GetArenaForAllocation());
    _impl_.op_params_ = p;
  }
  return _impl_.op_params_;
}
inline ::xla::cpu::OpParamsProto* CollectiveThunkProto::mutable_op_params() {
  ::xla::cpu::OpParamsProto* _msg = _internal_mutable_op_params();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.op_params)
  return _msg;
}
inline void CollectiveThunkProto::set_allocated_op_params(::xla::cpu::OpParamsProto* op_params) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.op_params_;
  }
  if (op_params) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(op_params);
    if (message_arena != submessage_arena) {
      op_params = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, op_params, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.op_params_ = op_params;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CollectiveThunkProto.op_params)
}

// .xla.cpu.OpBuffersProto op_buffers = 2;
inline bool CollectiveThunkProto::_internal_has_op_buffers() const {
  return this != internal_default_instance() && _impl_.op_buffers_ != nullptr;
}
inline bool CollectiveThunkProto::has_op_buffers() const {
  return _internal_has_op_buffers();
}
inline void CollectiveThunkProto::clear_op_buffers() {
  if (GetArenaForAllocation() == nullptr && _impl_.op_buffers_ != nullptr) {
    delete _impl_.op_buffers_;
  }
  _impl_.op_buffers_ = nullptr;
}
inline const ::xla::cpu::OpBuffersProto& CollectiveThunkProto::_internal_op_buffers() const {
  const ::xla::cpu::OpBuffersProto* p = _impl_.op_buffers_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::OpBuffersProto&>(
      ::xla::cpu::_OpBuffersProto_default_instance_);
}
inline const ::xla::cpu::OpBuffersProto& CollectiveThunkProto::op_buffers() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.op_buffers)
  return _internal_op_buffers();
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_op_buffers(
    ::xla::cpu::OpBuffersProto* op_buffers) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.op_buffers_);
  }
  _impl_.op_buffers_ = op_buffers;
  if (op_buffers) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.op_buffers)
}
inline ::xla::cpu::OpBuffersProto* CollectiveThunkProto::release_op_buffers() {
  
  ::xla::cpu::OpBuffersProto* temp = _impl_.op_buffers_;
  _impl_.op_buffers_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::OpBuffersProto* CollectiveThunkProto::unsafe_arena_release_op_buffers() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.op_buffers)
  
  ::xla::cpu::OpBuffersProto* temp = _impl_.op_buffers_;
  _impl_.op_buffers_ = nullptr;
  return temp;
}
inline ::xla::cpu::OpBuffersProto* CollectiveThunkProto::_internal_mutable_op_buffers() {
  
  if (_impl_.op_buffers_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::OpBuffersProto>(GetArenaForAllocation());
    _impl_.op_buffers_ = p;
  }
  return _impl_.op_buffers_;
}
inline ::xla::cpu::OpBuffersProto* CollectiveThunkProto::mutable_op_buffers() {
  ::xla::cpu::OpBuffersProto* _msg = _internal_mutable_op_buffers();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.op_buffers)
  return _msg;
}
inline void CollectiveThunkProto::set_allocated_op_buffers(::xla::cpu::OpBuffersProto* op_buffers) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.op_buffers_;
  }
  if (op_buffers) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(op_buffers);
    if (message_arena != submessage_arena) {
      op_buffers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, op_buffers, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.op_buffers_ = op_buffers;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CollectiveThunkProto.op_buffers)
}

// .xla.cpu.OpResourcesProto op_resources = 3;
inline bool CollectiveThunkProto::_internal_has_op_resources() const {
  return this != internal_default_instance() && _impl_.op_resources_ != nullptr;
}
inline bool CollectiveThunkProto::has_op_resources() const {
  return _internal_has_op_resources();
}
inline void CollectiveThunkProto::clear_op_resources() {
  if (GetArenaForAllocation() == nullptr && _impl_.op_resources_ != nullptr) {
    delete _impl_.op_resources_;
  }
  _impl_.op_resources_ = nullptr;
}
inline const ::xla::cpu::OpResourcesProto& CollectiveThunkProto::_internal_op_resources() const {
  const ::xla::cpu::OpResourcesProto* p = _impl_.op_resources_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::OpResourcesProto&>(
      ::xla::cpu::_OpResourcesProto_default_instance_);
}
inline const ::xla::cpu::OpResourcesProto& CollectiveThunkProto::op_resources() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.op_resources)
  return _internal_op_resources();
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_op_resources(
    ::xla::cpu::OpResourcesProto* op_resources) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.op_resources_);
  }
  _impl_.op_resources_ = op_resources;
  if (op_resources) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.op_resources)
}
inline ::xla::cpu::OpResourcesProto* CollectiveThunkProto::release_op_resources() {
  
  ::xla::cpu::OpResourcesProto* temp = _impl_.op_resources_;
  _impl_.op_resources_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::OpResourcesProto* CollectiveThunkProto::unsafe_arena_release_op_resources() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.op_resources)
  
  ::xla::cpu::OpResourcesProto* temp = _impl_.op_resources_;
  _impl_.op_resources_ = nullptr;
  return temp;
}
inline ::xla::cpu::OpResourcesProto* CollectiveThunkProto::_internal_mutable_op_resources() {
  
  if (_impl_.op_resources_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::OpResourcesProto>(GetArenaForAllocation());
    _impl_.op_resources_ = p;
  }
  return _impl_.op_resources_;
}
inline ::xla::cpu::OpResourcesProto* CollectiveThunkProto::mutable_op_resources() {
  ::xla::cpu::OpResourcesProto* _msg = _internal_mutable_op_resources();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.op_resources)
  return _msg;
}
inline void CollectiveThunkProto::set_allocated_op_resources(::xla::cpu::OpResourcesProto* op_resources) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.op_resources_;
  }
  if (op_resources) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(op_resources);
    if (message_arena != submessage_arena) {
      op_resources = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, op_resources, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.op_resources_ = op_resources;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CollectiveThunkProto.op_resources)
}

// .xla.cpu.AllGatherThunkProto all_gather_thunk = 4;
inline bool CollectiveThunkProto::_internal_has_all_gather_thunk() const {
  return impl_case() == kAllGatherThunk;
}
inline bool CollectiveThunkProto::has_all_gather_thunk() const {
  return _internal_has_all_gather_thunk();
}
inline void CollectiveThunkProto::set_has_all_gather_thunk() {
  _impl_._oneof_case_[0] = kAllGatherThunk;
}
inline void CollectiveThunkProto::clear_all_gather_thunk() {
  if (_internal_has_all_gather_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.all_gather_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::AllGatherThunkProto* CollectiveThunkProto::release_all_gather_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.all_gather_thunk)
  if (_internal_has_all_gather_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllGatherThunkProto* temp = _impl_.impl_.all_gather_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.all_gather_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::AllGatherThunkProto& CollectiveThunkProto::_internal_all_gather_thunk() const {
  return _internal_has_all_gather_thunk()
      ? *_impl_.impl_.all_gather_thunk_
      : reinterpret_cast< ::xla::cpu::AllGatherThunkProto&>(::xla::cpu::_AllGatherThunkProto_default_instance_);
}
inline const ::xla::cpu::AllGatherThunkProto& CollectiveThunkProto::all_gather_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.all_gather_thunk)
  return _internal_all_gather_thunk();
}
inline ::xla::cpu::AllGatherThunkProto* CollectiveThunkProto::unsafe_arena_release_all_gather_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.CollectiveThunkProto.all_gather_thunk)
  if (_internal_has_all_gather_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllGatherThunkProto* temp = _impl_.impl_.all_gather_thunk_;
    _impl_.impl_.all_gather_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_all_gather_thunk(::xla::cpu::AllGatherThunkProto* all_gather_thunk) {
  clear_impl();
  if (all_gather_thunk) {
    set_has_all_gather_thunk();
    _impl_.impl_.all_gather_thunk_ = all_gather_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.all_gather_thunk)
}
inline ::xla::cpu::AllGatherThunkProto* CollectiveThunkProto::_internal_mutable_all_gather_thunk() {
  if (!_internal_has_all_gather_thunk()) {
    clear_impl();
    set_has_all_gather_thunk();
    _impl_.impl_.all_gather_thunk_ = CreateMaybeMessage< ::xla::cpu::AllGatherThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.all_gather_thunk_;
}
inline ::xla::cpu::AllGatherThunkProto* CollectiveThunkProto::mutable_all_gather_thunk() {
  ::xla::cpu::AllGatherThunkProto* _msg = _internal_mutable_all_gather_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.all_gather_thunk)
  return _msg;
}

// .xla.cpu.AllReduceThunkProto all_reduce_thunk = 5;
inline bool CollectiveThunkProto::_internal_has_all_reduce_thunk() const {
  return impl_case() == kAllReduceThunk;
}
inline bool CollectiveThunkProto::has_all_reduce_thunk() const {
  return _internal_has_all_reduce_thunk();
}
inline void CollectiveThunkProto::set_has_all_reduce_thunk() {
  _impl_._oneof_case_[0] = kAllReduceThunk;
}
inline void CollectiveThunkProto::clear_all_reduce_thunk() {
  if (_internal_has_all_reduce_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.all_reduce_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::AllReduceThunkProto* CollectiveThunkProto::release_all_reduce_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.all_reduce_thunk)
  if (_internal_has_all_reduce_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllReduceThunkProto* temp = _impl_.impl_.all_reduce_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.all_reduce_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::AllReduceThunkProto& CollectiveThunkProto::_internal_all_reduce_thunk() const {
  return _internal_has_all_reduce_thunk()
      ? *_impl_.impl_.all_reduce_thunk_
      : reinterpret_cast< ::xla::cpu::AllReduceThunkProto&>(::xla::cpu::_AllReduceThunkProto_default_instance_);
}
inline const ::xla::cpu::AllReduceThunkProto& CollectiveThunkProto::all_reduce_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.all_reduce_thunk)
  return _internal_all_reduce_thunk();
}
inline ::xla::cpu::AllReduceThunkProto* CollectiveThunkProto::unsafe_arena_release_all_reduce_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.CollectiveThunkProto.all_reduce_thunk)
  if (_internal_has_all_reduce_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllReduceThunkProto* temp = _impl_.impl_.all_reduce_thunk_;
    _impl_.impl_.all_reduce_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_all_reduce_thunk(::xla::cpu::AllReduceThunkProto* all_reduce_thunk) {
  clear_impl();
  if (all_reduce_thunk) {
    set_has_all_reduce_thunk();
    _impl_.impl_.all_reduce_thunk_ = all_reduce_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.all_reduce_thunk)
}
inline ::xla::cpu::AllReduceThunkProto* CollectiveThunkProto::_internal_mutable_all_reduce_thunk() {
  if (!_internal_has_all_reduce_thunk()) {
    clear_impl();
    set_has_all_reduce_thunk();
    _impl_.impl_.all_reduce_thunk_ = CreateMaybeMessage< ::xla::cpu::AllReduceThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.all_reduce_thunk_;
}
inline ::xla::cpu::AllReduceThunkProto* CollectiveThunkProto::mutable_all_reduce_thunk() {
  ::xla::cpu::AllReduceThunkProto* _msg = _internal_mutable_all_reduce_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.all_reduce_thunk)
  return _msg;
}

// .xla.cpu.AllToAllThunkProto all_to_all_thunk = 6;
inline bool CollectiveThunkProto::_internal_has_all_to_all_thunk() const {
  return impl_case() == kAllToAllThunk;
}
inline bool CollectiveThunkProto::has_all_to_all_thunk() const {
  return _internal_has_all_to_all_thunk();
}
inline void CollectiveThunkProto::set_has_all_to_all_thunk() {
  _impl_._oneof_case_[0] = kAllToAllThunk;
}
inline void CollectiveThunkProto::clear_all_to_all_thunk() {
  if (_internal_has_all_to_all_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.all_to_all_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::AllToAllThunkProto* CollectiveThunkProto::release_all_to_all_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.all_to_all_thunk)
  if (_internal_has_all_to_all_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllToAllThunkProto* temp = _impl_.impl_.all_to_all_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.all_to_all_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::AllToAllThunkProto& CollectiveThunkProto::_internal_all_to_all_thunk() const {
  return _internal_has_all_to_all_thunk()
      ? *_impl_.impl_.all_to_all_thunk_
      : reinterpret_cast< ::xla::cpu::AllToAllThunkProto&>(::xla::cpu::_AllToAllThunkProto_default_instance_);
}
inline const ::xla::cpu::AllToAllThunkProto& CollectiveThunkProto::all_to_all_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.all_to_all_thunk)
  return _internal_all_to_all_thunk();
}
inline ::xla::cpu::AllToAllThunkProto* CollectiveThunkProto::unsafe_arena_release_all_to_all_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.CollectiveThunkProto.all_to_all_thunk)
  if (_internal_has_all_to_all_thunk()) {
    clear_has_impl();
    ::xla::cpu::AllToAllThunkProto* temp = _impl_.impl_.all_to_all_thunk_;
    _impl_.impl_.all_to_all_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_all_to_all_thunk(::xla::cpu::AllToAllThunkProto* all_to_all_thunk) {
  clear_impl();
  if (all_to_all_thunk) {
    set_has_all_to_all_thunk();
    _impl_.impl_.all_to_all_thunk_ = all_to_all_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.all_to_all_thunk)
}
inline ::xla::cpu::AllToAllThunkProto* CollectiveThunkProto::_internal_mutable_all_to_all_thunk() {
  if (!_internal_has_all_to_all_thunk()) {
    clear_impl();
    set_has_all_to_all_thunk();
    _impl_.impl_.all_to_all_thunk_ = CreateMaybeMessage< ::xla::cpu::AllToAllThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.all_to_all_thunk_;
}
inline ::xla::cpu::AllToAllThunkProto* CollectiveThunkProto::mutable_all_to_all_thunk() {
  ::xla::cpu::AllToAllThunkProto* _msg = _internal_mutable_all_to_all_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.all_to_all_thunk)
  return _msg;
}

// .xla.cpu.ReduceScatterThunkProto reduce_scatter_thunk = 7;
inline bool CollectiveThunkProto::_internal_has_reduce_scatter_thunk() const {
  return impl_case() == kReduceScatterThunk;
}
inline bool CollectiveThunkProto::has_reduce_scatter_thunk() const {
  return _internal_has_reduce_scatter_thunk();
}
inline void CollectiveThunkProto::set_has_reduce_scatter_thunk() {
  _impl_._oneof_case_[0] = kReduceScatterThunk;
}
inline void CollectiveThunkProto::clear_reduce_scatter_thunk() {
  if (_internal_has_reduce_scatter_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.reduce_scatter_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::ReduceScatterThunkProto* CollectiveThunkProto::release_reduce_scatter_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.reduce_scatter_thunk)
  if (_internal_has_reduce_scatter_thunk()) {
    clear_has_impl();
    ::xla::cpu::ReduceScatterThunkProto* temp = _impl_.impl_.reduce_scatter_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.reduce_scatter_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::ReduceScatterThunkProto& CollectiveThunkProto::_internal_reduce_scatter_thunk() const {
  return _internal_has_reduce_scatter_thunk()
      ? *_impl_.impl_.reduce_scatter_thunk_
      : reinterpret_cast< ::xla::cpu::ReduceScatterThunkProto&>(::xla::cpu::_ReduceScatterThunkProto_default_instance_);
}
inline const ::xla::cpu::ReduceScatterThunkProto& CollectiveThunkProto::reduce_scatter_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.reduce_scatter_thunk)
  return _internal_reduce_scatter_thunk();
}
inline ::xla::cpu::ReduceScatterThunkProto* CollectiveThunkProto::unsafe_arena_release_reduce_scatter_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.CollectiveThunkProto.reduce_scatter_thunk)
  if (_internal_has_reduce_scatter_thunk()) {
    clear_has_impl();
    ::xla::cpu::ReduceScatterThunkProto* temp = _impl_.impl_.reduce_scatter_thunk_;
    _impl_.impl_.reduce_scatter_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_reduce_scatter_thunk(::xla::cpu::ReduceScatterThunkProto* reduce_scatter_thunk) {
  clear_impl();
  if (reduce_scatter_thunk) {
    set_has_reduce_scatter_thunk();
    _impl_.impl_.reduce_scatter_thunk_ = reduce_scatter_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.reduce_scatter_thunk)
}
inline ::xla::cpu::ReduceScatterThunkProto* CollectiveThunkProto::_internal_mutable_reduce_scatter_thunk() {
  if (!_internal_has_reduce_scatter_thunk()) {
    clear_impl();
    set_has_reduce_scatter_thunk();
    _impl_.impl_.reduce_scatter_thunk_ = CreateMaybeMessage< ::xla::cpu::ReduceScatterThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.reduce_scatter_thunk_;
}
inline ::xla::cpu::ReduceScatterThunkProto* CollectiveThunkProto::mutable_reduce_scatter_thunk() {
  ::xla::cpu::ReduceScatterThunkProto* _msg = _internal_mutable_reduce_scatter_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.reduce_scatter_thunk)
  return _msg;
}

// .xla.cpu.CollectivePermuteThunkProto collective_permute_thunk = 8;
inline bool CollectiveThunkProto::_internal_has_collective_permute_thunk() const {
  return impl_case() == kCollectivePermuteThunk;
}
inline bool CollectiveThunkProto::has_collective_permute_thunk() const {
  return _internal_has_collective_permute_thunk();
}
inline void CollectiveThunkProto::set_has_collective_permute_thunk() {
  _impl_._oneof_case_[0] = kCollectivePermuteThunk;
}
inline void CollectiveThunkProto::clear_collective_permute_thunk() {
  if (_internal_has_collective_permute_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.collective_permute_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::CollectivePermuteThunkProto* CollectiveThunkProto::release_collective_permute_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.CollectiveThunkProto.collective_permute_thunk)
  if (_internal_has_collective_permute_thunk()) {
    clear_has_impl();
    ::xla::cpu::CollectivePermuteThunkProto* temp = _impl_.impl_.collective_permute_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.collective_permute_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::CollectivePermuteThunkProto& CollectiveThunkProto::_internal_collective_permute_thunk() const {
  return _internal_has_collective_permute_thunk()
      ? *_impl_.impl_.collective_permute_thunk_
      : reinterpret_cast< ::xla::cpu::CollectivePermuteThunkProto&>(::xla::cpu::_CollectivePermuteThunkProto_default_instance_);
}
inline const ::xla::cpu::CollectivePermuteThunkProto& CollectiveThunkProto::collective_permute_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CollectiveThunkProto.collective_permute_thunk)
  return _internal_collective_permute_thunk();
}
inline ::xla::cpu::CollectivePermuteThunkProto* CollectiveThunkProto::unsafe_arena_release_collective_permute_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.CollectiveThunkProto.collective_permute_thunk)
  if (_internal_has_collective_permute_thunk()) {
    clear_has_impl();
    ::xla::cpu::CollectivePermuteThunkProto* temp = _impl_.impl_.collective_permute_thunk_;
    _impl_.impl_.collective_permute_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void CollectiveThunkProto::unsafe_arena_set_allocated_collective_permute_thunk(::xla::cpu::CollectivePermuteThunkProto* collective_permute_thunk) {
  clear_impl();
  if (collective_permute_thunk) {
    set_has_collective_permute_thunk();
    _impl_.impl_.collective_permute_thunk_ = collective_permute_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CollectiveThunkProto.collective_permute_thunk)
}
inline ::xla::cpu::CollectivePermuteThunkProto* CollectiveThunkProto::_internal_mutable_collective_permute_thunk() {
  if (!_internal_has_collective_permute_thunk()) {
    clear_impl();
    set_has_collective_permute_thunk();
    _impl_.impl_.collective_permute_thunk_ = CreateMaybeMessage< ::xla::cpu::CollectivePermuteThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.collective_permute_thunk_;
}
inline ::xla::cpu::CollectivePermuteThunkProto* CollectiveThunkProto::mutable_collective_permute_thunk() {
  ::xla::cpu::CollectivePermuteThunkProto* _msg = _internal_mutable_collective_permute_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CollectiveThunkProto.collective_permute_thunk)
  return _msg;
}

inline bool CollectiveThunkProto::has_impl() const {
  return impl_case() != IMPL_NOT_SET;
}
inline void CollectiveThunkProto::clear_has_impl() {
  _impl_._oneof_case_[0] = IMPL_NOT_SET;
}
inline CollectiveThunkProto::ImplCase CollectiveThunkProto::impl_case() const {
  return CollectiveThunkProto::ImplCase(_impl_._oneof_case_[0]);
}
// -------------------------------------------------------------------

// CallThunkProto

// .xla.cpu.ThunkSequenceProto called_sequence = 1;
inline bool CallThunkProto::_internal_has_called_sequence() const {
  return this != internal_default_instance() && _impl_.called_sequence_ != nullptr;
}
inline bool CallThunkProto::has_called_sequence() const {
  return _internal_has_called_sequence();
}
inline void CallThunkProto::clear_called_sequence() {
  if (GetArenaForAllocation() == nullptr && _impl_.called_sequence_ != nullptr) {
    delete _impl_.called_sequence_;
  }
  _impl_.called_sequence_ = nullptr;
}
inline const ::xla::cpu::ThunkSequenceProto& CallThunkProto::_internal_called_sequence() const {
  const ::xla::cpu::ThunkSequenceProto* p = _impl_.called_sequence_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ThunkSequenceProto&>(
      ::xla::cpu::_ThunkSequenceProto_default_instance_);
}
inline const ::xla::cpu::ThunkSequenceProto& CallThunkProto::called_sequence() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CallThunkProto.called_sequence)
  return _internal_called_sequence();
}
inline void CallThunkProto::unsafe_arena_set_allocated_called_sequence(
    ::xla::cpu::ThunkSequenceProto* called_sequence) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.called_sequence_);
  }
  _impl_.called_sequence_ = called_sequence;
  if (called_sequence) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CallThunkProto.called_sequence)
}
inline ::xla::cpu::ThunkSequenceProto* CallThunkProto::release_called_sequence() {
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.called_sequence_;
  _impl_.called_sequence_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* CallThunkProto::unsafe_arena_release_called_sequence() {
  // @@protoc_insertion_point(field_release:xla.cpu.CallThunkProto.called_sequence)
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.called_sequence_;
  _impl_.called_sequence_ = nullptr;
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* CallThunkProto::_internal_mutable_called_sequence() {
  
  if (_impl_.called_sequence_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ThunkSequenceProto>(GetArenaForAllocation());
    _impl_.called_sequence_ = p;
  }
  return _impl_.called_sequence_;
}
inline ::xla::cpu::ThunkSequenceProto* CallThunkProto::mutable_called_sequence() {
  ::xla::cpu::ThunkSequenceProto* _msg = _internal_mutable_called_sequence();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CallThunkProto.called_sequence)
  return _msg;
}
inline void CallThunkProto::set_allocated_called_sequence(::xla::cpu::ThunkSequenceProto* called_sequence) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.called_sequence_;
  }
  if (called_sequence) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(called_sequence);
    if (message_arena != submessage_arena) {
      called_sequence = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, called_sequence, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.called_sequence_ = called_sequence;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CallThunkProto.called_sequence)
}

// -------------------------------------------------------------------

// ConditionalThunkProto

// repeated .xla.cpu.ThunkSequenceProto branch_sequences = 1;
inline int ConditionalThunkProto::_internal_branch_sequences_size() const {
  return _impl_.branch_sequences_.size();
}
inline int ConditionalThunkProto::branch_sequences_size() const {
  return _internal_branch_sequences_size();
}
inline void ConditionalThunkProto::clear_branch_sequences() {
  _impl_.branch_sequences_.Clear();
}
inline ::xla::cpu::ThunkSequenceProto* ConditionalThunkProto::mutable_branch_sequences(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConditionalThunkProto.branch_sequences)
  return _impl_.branch_sequences_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkSequenceProto >*
ConditionalThunkProto::mutable_branch_sequences() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.ConditionalThunkProto.branch_sequences)
  return &_impl_.branch_sequences_;
}
inline const ::xla::cpu::ThunkSequenceProto& ConditionalThunkProto::_internal_branch_sequences(int index) const {
  return _impl_.branch_sequences_.Get(index);
}
inline const ::xla::cpu::ThunkSequenceProto& ConditionalThunkProto::branch_sequences(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConditionalThunkProto.branch_sequences)
  return _internal_branch_sequences(index);
}
inline ::xla::cpu::ThunkSequenceProto* ConditionalThunkProto::_internal_add_branch_sequences() {
  return _impl_.branch_sequences_.Add();
}
inline ::xla::cpu::ThunkSequenceProto* ConditionalThunkProto::add_branch_sequences() {
  ::xla::cpu::ThunkSequenceProto* _add = _internal_add_branch_sequences();
  // @@protoc_insertion_point(field_add:xla.cpu.ConditionalThunkProto.branch_sequences)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkSequenceProto >&
ConditionalThunkProto::branch_sequences() const {
  // @@protoc_insertion_point(field_list:xla.cpu.ConditionalThunkProto.branch_sequences)
  return _impl_.branch_sequences_;
}

// .xla.cpu.BufferAllocationSliceProto branch_index_buffer = 2;
inline bool ConditionalThunkProto::_internal_has_branch_index_buffer() const {
  return this != internal_default_instance() && _impl_.branch_index_buffer_ != nullptr;
}
inline bool ConditionalThunkProto::has_branch_index_buffer() const {
  return _internal_has_branch_index_buffer();
}
inline void ConditionalThunkProto::clear_branch_index_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.branch_index_buffer_ != nullptr) {
    delete _impl_.branch_index_buffer_;
  }
  _impl_.branch_index_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& ConditionalThunkProto::_internal_branch_index_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.branch_index_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& ConditionalThunkProto::branch_index_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConditionalThunkProto.branch_index_buffer)
  return _internal_branch_index_buffer();
}
inline void ConditionalThunkProto::unsafe_arena_set_allocated_branch_index_buffer(
    ::xla::cpu::BufferAllocationSliceProto* branch_index_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.branch_index_buffer_);
  }
  _impl_.branch_index_buffer_ = branch_index_buffer;
  if (branch_index_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConditionalThunkProto.branch_index_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* ConditionalThunkProto::release_branch_index_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.branch_index_buffer_;
  _impl_.branch_index_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ConditionalThunkProto::unsafe_arena_release_branch_index_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConditionalThunkProto.branch_index_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.branch_index_buffer_;
  _impl_.branch_index_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ConditionalThunkProto::_internal_mutable_branch_index_buffer() {
  
  if (_impl_.branch_index_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.branch_index_buffer_ = p;
  }
  return _impl_.branch_index_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* ConditionalThunkProto::mutable_branch_index_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_branch_index_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConditionalThunkProto.branch_index_buffer)
  return _msg;
}
inline void ConditionalThunkProto::set_allocated_branch_index_buffer(::xla::cpu::BufferAllocationSliceProto* branch_index_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.branch_index_buffer_;
  }
  if (branch_index_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(branch_index_buffer);
    if (message_arena != submessage_arena) {
      branch_index_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, branch_index_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.branch_index_buffer_ = branch_index_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConditionalThunkProto.branch_index_buffer)
}

// -------------------------------------------------------------------

// ConvolutionThunkProto_Options

// bool multi_threaded = 1;
inline void ConvolutionThunkProto_Options::clear_multi_threaded() {
  _impl_.multi_threaded_ = false;
}
inline bool ConvolutionThunkProto_Options::_internal_multi_threaded() const {
  return _impl_.multi_threaded_;
}
inline bool ConvolutionThunkProto_Options::multi_threaded() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.Options.multi_threaded)
  return _internal_multi_threaded();
}
inline void ConvolutionThunkProto_Options::_internal_set_multi_threaded(bool value) {
  
  _impl_.multi_threaded_ = value;
}
inline void ConvolutionThunkProto_Options::set_multi_threaded(bool value) {
  _internal_set_multi_threaded(value);
  // @@protoc_insertion_point(field_set:xla.cpu.ConvolutionThunkProto.Options.multi_threaded)
}

// -------------------------------------------------------------------

// ConvolutionThunkProto

// .xla.ConvolutionDimensionNumbers dimension_numbers = 1;
inline bool ConvolutionThunkProto::_internal_has_dimension_numbers() const {
  return this != internal_default_instance() && _impl_.dimension_numbers_ != nullptr;
}
inline bool ConvolutionThunkProto::has_dimension_numbers() const {
  return _internal_has_dimension_numbers();
}
inline const ::xla::ConvolutionDimensionNumbers& ConvolutionThunkProto::_internal_dimension_numbers() const {
  const ::xla::ConvolutionDimensionNumbers* p = _impl_.dimension_numbers_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ConvolutionDimensionNumbers&>(
      ::xla::_ConvolutionDimensionNumbers_default_instance_);
}
inline const ::xla::ConvolutionDimensionNumbers& ConvolutionThunkProto::dimension_numbers() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.dimension_numbers)
  return _internal_dimension_numbers();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_dimension_numbers(
    ::xla::ConvolutionDimensionNumbers* dimension_numbers) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dimension_numbers_);
  }
  _impl_.dimension_numbers_ = dimension_numbers;
  if (dimension_numbers) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.dimension_numbers)
}
inline ::xla::ConvolutionDimensionNumbers* ConvolutionThunkProto::release_dimension_numbers() {
  
  ::xla::ConvolutionDimensionNumbers* temp = _impl_.dimension_numbers_;
  _impl_.dimension_numbers_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::ConvolutionDimensionNumbers* ConvolutionThunkProto::unsafe_arena_release_dimension_numbers() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.dimension_numbers)
  
  ::xla::ConvolutionDimensionNumbers* temp = _impl_.dimension_numbers_;
  _impl_.dimension_numbers_ = nullptr;
  return temp;
}
inline ::xla::ConvolutionDimensionNumbers* ConvolutionThunkProto::_internal_mutable_dimension_numbers() {
  
  if (_impl_.dimension_numbers_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::ConvolutionDimensionNumbers>(GetArenaForAllocation());
    _impl_.dimension_numbers_ = p;
  }
  return _impl_.dimension_numbers_;
}
inline ::xla::ConvolutionDimensionNumbers* ConvolutionThunkProto::mutable_dimension_numbers() {
  ::xla::ConvolutionDimensionNumbers* _msg = _internal_mutable_dimension_numbers();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.dimension_numbers)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_dimension_numbers(::xla::ConvolutionDimensionNumbers* dimension_numbers) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dimension_numbers_);
  }
  if (dimension_numbers) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(dimension_numbers));
    if (message_arena != submessage_arena) {
      dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, dimension_numbers, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.dimension_numbers_ = dimension_numbers;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.dimension_numbers)
}

// .xla.Window window = 2;
inline bool ConvolutionThunkProto::_internal_has_window() const {
  return this != internal_default_instance() && _impl_.window_ != nullptr;
}
inline bool ConvolutionThunkProto::has_window() const {
  return _internal_has_window();
}
inline const ::xla::Window& ConvolutionThunkProto::_internal_window() const {
  const ::xla::Window* p = _impl_.window_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::Window&>(
      ::xla::_Window_default_instance_);
}
inline const ::xla::Window& ConvolutionThunkProto::window() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.window)
  return _internal_window();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_window(
    ::xla::Window* window) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.window_);
  }
  _impl_.window_ = window;
  if (window) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.window)
}
inline ::xla::Window* ConvolutionThunkProto::release_window() {
  
  ::xla::Window* temp = _impl_.window_;
  _impl_.window_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::Window* ConvolutionThunkProto::unsafe_arena_release_window() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.window)
  
  ::xla::Window* temp = _impl_.window_;
  _impl_.window_ = nullptr;
  return temp;
}
inline ::xla::Window* ConvolutionThunkProto::_internal_mutable_window() {
  
  if (_impl_.window_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::Window>(GetArenaForAllocation());
    _impl_.window_ = p;
  }
  return _impl_.window_;
}
inline ::xla::Window* ConvolutionThunkProto::mutable_window() {
  ::xla::Window* _msg = _internal_mutable_window();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.window)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_window(::xla::Window* window) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.window_);
  }
  if (window) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(window));
    if (message_arena != submessage_arena) {
      window = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, window, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.window_ = window;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.window)
}

// int64 feature_group_count = 3;
inline void ConvolutionThunkProto::clear_feature_group_count() {
  _impl_.feature_group_count_ = int64_t{0};
}
inline int64_t ConvolutionThunkProto::_internal_feature_group_count() const {
  return _impl_.feature_group_count_;
}
inline int64_t ConvolutionThunkProto::feature_group_count() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.feature_group_count)
  return _internal_feature_group_count();
}
inline void ConvolutionThunkProto::_internal_set_feature_group_count(int64_t value) {
  
  _impl_.feature_group_count_ = value;
}
inline void ConvolutionThunkProto::set_feature_group_count(int64_t value) {
  _internal_set_feature_group_count(value);
  // @@protoc_insertion_point(field_set:xla.cpu.ConvolutionThunkProto.feature_group_count)
}

// .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
inline bool ConvolutionThunkProto::_internal_has_input_buffer_shape() const {
  return this != internal_default_instance() && _impl_.input_buffer_shape_ != nullptr;
}
inline bool ConvolutionThunkProto::has_input_buffer_shape() const {
  return _internal_has_input_buffer_shape();
}
inline void ConvolutionThunkProto::clear_input_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.input_buffer_shape_ != nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  _impl_.input_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::_internal_input_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.input_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::input_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.input_buffer_shape)
  return _internal_input_buffer_shape();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_input_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.input_buffer_shape_);
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  if (input_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.input_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::release_input_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::unsafe_arena_release_input_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.input_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::_internal_mutable_input_buffer_shape() {
  
  if (_impl_.input_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.input_buffer_shape_ = p;
  }
  return _impl_.input_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::mutable_input_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_input_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.input_buffer_shape)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  if (input_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(input_buffer_shape);
    if (message_arena != submessage_arena) {
      input_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, input_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.input_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto kernel_buffer_shape = 5;
inline bool ConvolutionThunkProto::_internal_has_kernel_buffer_shape() const {
  return this != internal_default_instance() && _impl_.kernel_buffer_shape_ != nullptr;
}
inline bool ConvolutionThunkProto::has_kernel_buffer_shape() const {
  return _internal_has_kernel_buffer_shape();
}
inline void ConvolutionThunkProto::clear_kernel_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.kernel_buffer_shape_ != nullptr) {
    delete _impl_.kernel_buffer_shape_;
  }
  _impl_.kernel_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::_internal_kernel_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.kernel_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::kernel_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.kernel_buffer_shape)
  return _internal_kernel_buffer_shape();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_kernel_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.kernel_buffer_shape_);
  }
  _impl_.kernel_buffer_shape_ = kernel_buffer_shape;
  if (kernel_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.kernel_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::release_kernel_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.kernel_buffer_shape_;
  _impl_.kernel_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::unsafe_arena_release_kernel_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.kernel_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.kernel_buffer_shape_;
  _impl_.kernel_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::_internal_mutable_kernel_buffer_shape() {
  
  if (_impl_.kernel_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.kernel_buffer_shape_ = p;
  }
  return _impl_.kernel_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::mutable_kernel_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_kernel_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.kernel_buffer_shape)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_kernel_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.kernel_buffer_shape_;
  }
  if (kernel_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(kernel_buffer_shape);
    if (message_arena != submessage_arena) {
      kernel_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, kernel_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.kernel_buffer_shape_ = kernel_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.kernel_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 6;
inline bool ConvolutionThunkProto::_internal_has_output_buffer_shape() const {
  return this != internal_default_instance() && _impl_.output_buffer_shape_ != nullptr;
}
inline bool ConvolutionThunkProto::has_output_buffer_shape() const {
  return _internal_has_output_buffer_shape();
}
inline void ConvolutionThunkProto::clear_output_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.output_buffer_shape_ != nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  _impl_.output_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::_internal_output_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.output_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& ConvolutionThunkProto::output_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.output_buffer_shape)
  return _internal_output_buffer_shape();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_output_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.output_buffer_shape_);
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  if (output_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.output_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::release_output_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::unsafe_arena_release_output_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.output_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::_internal_mutable_output_buffer_shape() {
  
  if (_impl_.output_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.output_buffer_shape_ = p;
  }
  return _impl_.output_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* ConvolutionThunkProto::mutable_output_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_output_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.output_buffer_shape)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  if (output_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(output_buffer_shape);
    if (message_arena != submessage_arena) {
      output_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, output_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.output_buffer_shape)
}

// .xla.cpu.ConvolutionThunkProto.Options options = 7;
inline bool ConvolutionThunkProto::_internal_has_options() const {
  return this != internal_default_instance() && _impl_.options_ != nullptr;
}
inline bool ConvolutionThunkProto::has_options() const {
  return _internal_has_options();
}
inline void ConvolutionThunkProto::clear_options() {
  if (GetArenaForAllocation() == nullptr && _impl_.options_ != nullptr) {
    delete _impl_.options_;
  }
  _impl_.options_ = nullptr;
}
inline const ::xla::cpu::ConvolutionThunkProto_Options& ConvolutionThunkProto::_internal_options() const {
  const ::xla::cpu::ConvolutionThunkProto_Options* p = _impl_.options_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ConvolutionThunkProto_Options&>(
      ::xla::cpu::_ConvolutionThunkProto_Options_default_instance_);
}
inline const ::xla::cpu::ConvolutionThunkProto_Options& ConvolutionThunkProto::options() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ConvolutionThunkProto.options)
  return _internal_options();
}
inline void ConvolutionThunkProto::unsafe_arena_set_allocated_options(
    ::xla::cpu::ConvolutionThunkProto_Options* options) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.options_);
  }
  _impl_.options_ = options;
  if (options) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ConvolutionThunkProto.options)
}
inline ::xla::cpu::ConvolutionThunkProto_Options* ConvolutionThunkProto::release_options() {
  
  ::xla::cpu::ConvolutionThunkProto_Options* temp = _impl_.options_;
  _impl_.options_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ConvolutionThunkProto_Options* ConvolutionThunkProto::unsafe_arena_release_options() {
  // @@protoc_insertion_point(field_release:xla.cpu.ConvolutionThunkProto.options)
  
  ::xla::cpu::ConvolutionThunkProto_Options* temp = _impl_.options_;
  _impl_.options_ = nullptr;
  return temp;
}
inline ::xla::cpu::ConvolutionThunkProto_Options* ConvolutionThunkProto::_internal_mutable_options() {
  
  if (_impl_.options_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ConvolutionThunkProto_Options>(GetArenaForAllocation());
    _impl_.options_ = p;
  }
  return _impl_.options_;
}
inline ::xla::cpu::ConvolutionThunkProto_Options* ConvolutionThunkProto::mutable_options() {
  ::xla::cpu::ConvolutionThunkProto_Options* _msg = _internal_mutable_options();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ConvolutionThunkProto.options)
  return _msg;
}
inline void ConvolutionThunkProto::set_allocated_options(::xla::cpu::ConvolutionThunkProto_Options* options) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.options_;
  }
  if (options) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(options);
    if (message_arena != submessage_arena) {
      options = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, options, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.options_ = options;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ConvolutionThunkProto.options)
}

// -------------------------------------------------------------------

// SortThunkProto

// int64 dimension = 1;
inline void SortThunkProto::clear_dimension() {
  _impl_.dimension_ = int64_t{0};
}
inline int64_t SortThunkProto::_internal_dimension() const {
  return _impl_.dimension_;
}
inline int64_t SortThunkProto::dimension() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortThunkProto.dimension)
  return _internal_dimension();
}
inline void SortThunkProto::_internal_set_dimension(int64_t value) {
  
  _impl_.dimension_ = value;
}
inline void SortThunkProto::set_dimension(int64_t value) {
  _internal_set_dimension(value);
  // @@protoc_insertion_point(field_set:xla.cpu.SortThunkProto.dimension)
}

// bool is_stable = 2;
inline void SortThunkProto::clear_is_stable() {
  _impl_.is_stable_ = false;
}
inline bool SortThunkProto::_internal_is_stable() const {
  return _impl_.is_stable_;
}
inline bool SortThunkProto::is_stable() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortThunkProto.is_stable)
  return _internal_is_stable();
}
inline void SortThunkProto::_internal_set_is_stable(bool value) {
  
  _impl_.is_stable_ = value;
}
inline void SortThunkProto::set_is_stable(bool value) {
  _internal_set_is_stable(value);
  // @@protoc_insertion_point(field_set:xla.cpu.SortThunkProto.is_stable)
}

// .xla.cpu.SortDirectionOptional direction = 3;
inline bool SortThunkProto::_internal_has_direction() const {
  return this != internal_default_instance() && _impl_.direction_ != nullptr;
}
inline bool SortThunkProto::has_direction() const {
  return _internal_has_direction();
}
inline void SortThunkProto::clear_direction() {
  if (GetArenaForAllocation() == nullptr && _impl_.direction_ != nullptr) {
    delete _impl_.direction_;
  }
  _impl_.direction_ = nullptr;
}
inline const ::xla::cpu::SortDirectionOptional& SortThunkProto::_internal_direction() const {
  const ::xla::cpu::SortDirectionOptional* p = _impl_.direction_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::SortDirectionOptional&>(
      ::xla::cpu::_SortDirectionOptional_default_instance_);
}
inline const ::xla::cpu::SortDirectionOptional& SortThunkProto::direction() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortThunkProto.direction)
  return _internal_direction();
}
inline void SortThunkProto::unsafe_arena_set_allocated_direction(
    ::xla::cpu::SortDirectionOptional* direction) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.direction_);
  }
  _impl_.direction_ = direction;
  if (direction) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.SortThunkProto.direction)
}
inline ::xla::cpu::SortDirectionOptional* SortThunkProto::release_direction() {
  
  ::xla::cpu::SortDirectionOptional* temp = _impl_.direction_;
  _impl_.direction_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::SortDirectionOptional* SortThunkProto::unsafe_arena_release_direction() {
  // @@protoc_insertion_point(field_release:xla.cpu.SortThunkProto.direction)
  
  ::xla::cpu::SortDirectionOptional* temp = _impl_.direction_;
  _impl_.direction_ = nullptr;
  return temp;
}
inline ::xla::cpu::SortDirectionOptional* SortThunkProto::_internal_mutable_direction() {
  
  if (_impl_.direction_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::SortDirectionOptional>(GetArenaForAllocation());
    _impl_.direction_ = p;
  }
  return _impl_.direction_;
}
inline ::xla::cpu::SortDirectionOptional* SortThunkProto::mutable_direction() {
  ::xla::cpu::SortDirectionOptional* _msg = _internal_mutable_direction();
  // @@protoc_insertion_point(field_mutable:xla.cpu.SortThunkProto.direction)
  return _msg;
}
inline void SortThunkProto::set_allocated_direction(::xla::cpu::SortDirectionOptional* direction) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.direction_;
  }
  if (direction) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(direction);
    if (message_arena != submessage_arena) {
      direction = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, direction, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.direction_ = direction;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.SortThunkProto.direction)
}

// string comparator_name = 4;
inline void SortThunkProto::clear_comparator_name() {
  _impl_.comparator_name_.ClearToEmpty();
}
inline const std::string& SortThunkProto::comparator_name() const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortThunkProto.comparator_name)
  return _internal_comparator_name();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void SortThunkProto::set_comparator_name(ArgT0&& arg0, ArgT... args) {
 
 _impl_.comparator_name_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.SortThunkProto.comparator_name)
}
inline std::string* SortThunkProto::mutable_comparator_name() {
  std::string* _s = _internal_mutable_comparator_name();
  // @@protoc_insertion_point(field_mutable:xla.cpu.SortThunkProto.comparator_name)
  return _s;
}
inline const std::string& SortThunkProto::_internal_comparator_name() const {
  return _impl_.comparator_name_.Get();
}
inline void SortThunkProto::_internal_set_comparator_name(const std::string& value) {
  
  _impl_.comparator_name_.Set(value, GetArenaForAllocation());
}
inline std::string* SortThunkProto::_internal_mutable_comparator_name() {
  
  return _impl_.comparator_name_.Mutable(GetArenaForAllocation());
}
inline std::string* SortThunkProto::release_comparator_name() {
  // @@protoc_insertion_point(field_release:xla.cpu.SortThunkProto.comparator_name)
  return _impl_.comparator_name_.Release();
}
inline void SortThunkProto::set_allocated_comparator_name(std::string* comparator_name) {
  if (comparator_name != nullptr) {
    
  } else {
    
  }
  _impl_.comparator_name_.SetAllocated(comparator_name, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.comparator_name_.IsDefault()) {
    _impl_.comparator_name_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.SortThunkProto.comparator_name)
}

// repeated .xla.cpu.ShapeBufferAllocationSliceProto inputs_shapes = 5;
inline int SortThunkProto::_internal_inputs_shapes_size() const {
  return _impl_.inputs_shapes_.size();
}
inline int SortThunkProto::inputs_shapes_size() const {
  return _internal_inputs_shapes_size();
}
inline void SortThunkProto::clear_inputs_shapes() {
  _impl_.inputs_shapes_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* SortThunkProto::mutable_inputs_shapes(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.SortThunkProto.inputs_shapes)
  return _impl_.inputs_shapes_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
SortThunkProto::mutable_inputs_shapes() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.SortThunkProto.inputs_shapes)
  return &_impl_.inputs_shapes_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& SortThunkProto::_internal_inputs_shapes(int index) const {
  return _impl_.inputs_shapes_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& SortThunkProto::inputs_shapes(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.SortThunkProto.inputs_shapes)
  return _internal_inputs_shapes(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* SortThunkProto::_internal_add_inputs_shapes() {
  return _impl_.inputs_shapes_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* SortThunkProto::add_inputs_shapes() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_inputs_shapes();
  // @@protoc_insertion_point(field_add:xla.cpu.SortThunkProto.inputs_shapes)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
SortThunkProto::inputs_shapes() const {
  // @@protoc_insertion_point(field_list:xla.cpu.SortThunkProto.inputs_shapes)
  return _impl_.inputs_shapes_;
}

// -------------------------------------------------------------------

// XnnDotThunkProto

// .xla.DotDimensionNumbers dot_dimensions = 1;
inline bool XnnDotThunkProto::_internal_has_dot_dimensions() const {
  return this != internal_default_instance() && _impl_.dot_dimensions_ != nullptr;
}
inline bool XnnDotThunkProto::has_dot_dimensions() const {
  return _internal_has_dot_dimensions();
}
inline const ::xla::DotDimensionNumbers& XnnDotThunkProto::_internal_dot_dimensions() const {
  const ::xla::DotDimensionNumbers* p = _impl_.dot_dimensions_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::DotDimensionNumbers&>(
      ::xla::_DotDimensionNumbers_default_instance_);
}
inline const ::xla::DotDimensionNumbers& XnnDotThunkProto::dot_dimensions() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnDotThunkProto.dot_dimensions)
  return _internal_dot_dimensions();
}
inline void XnnDotThunkProto::unsafe_arena_set_allocated_dot_dimensions(
    ::xla::DotDimensionNumbers* dot_dimensions) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimensions_);
  }
  _impl_.dot_dimensions_ = dot_dimensions;
  if (dot_dimensions) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnDotThunkProto.dot_dimensions)
}
inline ::xla::DotDimensionNumbers* XnnDotThunkProto::release_dot_dimensions() {
  
  ::xla::DotDimensionNumbers* temp = _impl_.dot_dimensions_;
  _impl_.dot_dimensions_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::DotDimensionNumbers* XnnDotThunkProto::unsafe_arena_release_dot_dimensions() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnDotThunkProto.dot_dimensions)
  
  ::xla::DotDimensionNumbers* temp = _impl_.dot_dimensions_;
  _impl_.dot_dimensions_ = nullptr;
  return temp;
}
inline ::xla::DotDimensionNumbers* XnnDotThunkProto::_internal_mutable_dot_dimensions() {
  
  if (_impl_.dot_dimensions_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation());
    _impl_.dot_dimensions_ = p;
  }
  return _impl_.dot_dimensions_;
}
inline ::xla::DotDimensionNumbers* XnnDotThunkProto::mutable_dot_dimensions() {
  ::xla::DotDimensionNumbers* _msg = _internal_mutable_dot_dimensions();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnDotThunkProto.dot_dimensions)
  return _msg;
}
inline void XnnDotThunkProto::set_allocated_dot_dimensions(::xla::DotDimensionNumbers* dot_dimensions) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimensions_);
  }
  if (dot_dimensions) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(dot_dimensions));
    if (message_arena != submessage_arena) {
      dot_dimensions = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, dot_dimensions, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.dot_dimensions_ = dot_dimensions;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnDotThunkProto.dot_dimensions)
}

// .xla.cpu.ShapeBufferAllocationSliceProto lhs_buffer_shape = 2;
inline bool XnnDotThunkProto::_internal_has_lhs_buffer_shape() const {
  return this != internal_default_instance() && _impl_.lhs_buffer_shape_ != nullptr;
}
inline bool XnnDotThunkProto::has_lhs_buffer_shape() const {
  return _internal_has_lhs_buffer_shape();
}
inline void XnnDotThunkProto::clear_lhs_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.lhs_buffer_shape_ != nullptr) {
    delete _impl_.lhs_buffer_shape_;
  }
  _impl_.lhs_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::_internal_lhs_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.lhs_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::lhs_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnDotThunkProto.lhs_buffer_shape)
  return _internal_lhs_buffer_shape();
}
inline void XnnDotThunkProto::unsafe_arena_set_allocated_lhs_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.lhs_buffer_shape_);
  }
  _impl_.lhs_buffer_shape_ = lhs_buffer_shape;
  if (lhs_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnDotThunkProto.lhs_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::release_lhs_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.lhs_buffer_shape_;
  _impl_.lhs_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::unsafe_arena_release_lhs_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnDotThunkProto.lhs_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.lhs_buffer_shape_;
  _impl_.lhs_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::_internal_mutable_lhs_buffer_shape() {
  
  if (_impl_.lhs_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.lhs_buffer_shape_ = p;
  }
  return _impl_.lhs_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::mutable_lhs_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_lhs_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnDotThunkProto.lhs_buffer_shape)
  return _msg;
}
inline void XnnDotThunkProto::set_allocated_lhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.lhs_buffer_shape_;
  }
  if (lhs_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(lhs_buffer_shape);
    if (message_arena != submessage_arena) {
      lhs_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, lhs_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.lhs_buffer_shape_ = lhs_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnDotThunkProto.lhs_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto rhs_buffer_shape = 3;
inline bool XnnDotThunkProto::_internal_has_rhs_buffer_shape() const {
  return this != internal_default_instance() && _impl_.rhs_buffer_shape_ != nullptr;
}
inline bool XnnDotThunkProto::has_rhs_buffer_shape() const {
  return _internal_has_rhs_buffer_shape();
}
inline void XnnDotThunkProto::clear_rhs_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.rhs_buffer_shape_ != nullptr) {
    delete _impl_.rhs_buffer_shape_;
  }
  _impl_.rhs_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::_internal_rhs_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.rhs_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::rhs_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnDotThunkProto.rhs_buffer_shape)
  return _internal_rhs_buffer_shape();
}
inline void XnnDotThunkProto::unsafe_arena_set_allocated_rhs_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.rhs_buffer_shape_);
  }
  _impl_.rhs_buffer_shape_ = rhs_buffer_shape;
  if (rhs_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnDotThunkProto.rhs_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::release_rhs_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.rhs_buffer_shape_;
  _impl_.rhs_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::unsafe_arena_release_rhs_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnDotThunkProto.rhs_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.rhs_buffer_shape_;
  _impl_.rhs_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::_internal_mutable_rhs_buffer_shape() {
  
  if (_impl_.rhs_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.rhs_buffer_shape_ = p;
  }
  return _impl_.rhs_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::mutable_rhs_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_rhs_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnDotThunkProto.rhs_buffer_shape)
  return _msg;
}
inline void XnnDotThunkProto::set_allocated_rhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.rhs_buffer_shape_;
  }
  if (rhs_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(rhs_buffer_shape);
    if (message_arena != submessage_arena) {
      rhs_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, rhs_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.rhs_buffer_shape_ = rhs_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnDotThunkProto.rhs_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto out_buffer_shape = 4;
inline bool XnnDotThunkProto::_internal_has_out_buffer_shape() const {
  return this != internal_default_instance() && _impl_.out_buffer_shape_ != nullptr;
}
inline bool XnnDotThunkProto::has_out_buffer_shape() const {
  return _internal_has_out_buffer_shape();
}
inline void XnnDotThunkProto::clear_out_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.out_buffer_shape_ != nullptr) {
    delete _impl_.out_buffer_shape_;
  }
  _impl_.out_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::_internal_out_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.out_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnDotThunkProto::out_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnDotThunkProto.out_buffer_shape)
  return _internal_out_buffer_shape();
}
inline void XnnDotThunkProto::unsafe_arena_set_allocated_out_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.out_buffer_shape_);
  }
  _impl_.out_buffer_shape_ = out_buffer_shape;
  if (out_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnDotThunkProto.out_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::release_out_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.out_buffer_shape_;
  _impl_.out_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::unsafe_arena_release_out_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnDotThunkProto.out_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.out_buffer_shape_;
  _impl_.out_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::_internal_mutable_out_buffer_shape() {
  
  if (_impl_.out_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.out_buffer_shape_ = p;
  }
  return _impl_.out_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnDotThunkProto::mutable_out_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_out_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnDotThunkProto.out_buffer_shape)
  return _msg;
}
inline void XnnDotThunkProto::set_allocated_out_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.out_buffer_shape_;
  }
  if (out_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(out_buffer_shape);
    if (message_arena != submessage_arena) {
      out_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, out_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.out_buffer_shape_ = out_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnDotThunkProto.out_buffer_shape)
}

// -------------------------------------------------------------------

// XnnConvolutionThunkProto

// .xla.ConvolutionDimensionNumbers dimension_numbers = 1;
inline bool XnnConvolutionThunkProto::_internal_has_dimension_numbers() const {
  return this != internal_default_instance() && _impl_.dimension_numbers_ != nullptr;
}
inline bool XnnConvolutionThunkProto::has_dimension_numbers() const {
  return _internal_has_dimension_numbers();
}
inline const ::xla::ConvolutionDimensionNumbers& XnnConvolutionThunkProto::_internal_dimension_numbers() const {
  const ::xla::ConvolutionDimensionNumbers* p = _impl_.dimension_numbers_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ConvolutionDimensionNumbers&>(
      ::xla::_ConvolutionDimensionNumbers_default_instance_);
}
inline const ::xla::ConvolutionDimensionNumbers& XnnConvolutionThunkProto::dimension_numbers() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.dimension_numbers)
  return _internal_dimension_numbers();
}
inline void XnnConvolutionThunkProto::unsafe_arena_set_allocated_dimension_numbers(
    ::xla::ConvolutionDimensionNumbers* dimension_numbers) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dimension_numbers_);
  }
  _impl_.dimension_numbers_ = dimension_numbers;
  if (dimension_numbers) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnConvolutionThunkProto.dimension_numbers)
}
inline ::xla::ConvolutionDimensionNumbers* XnnConvolutionThunkProto::release_dimension_numbers() {
  
  ::xla::ConvolutionDimensionNumbers* temp = _impl_.dimension_numbers_;
  _impl_.dimension_numbers_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::ConvolutionDimensionNumbers* XnnConvolutionThunkProto::unsafe_arena_release_dimension_numbers() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnConvolutionThunkProto.dimension_numbers)
  
  ::xla::ConvolutionDimensionNumbers* temp = _impl_.dimension_numbers_;
  _impl_.dimension_numbers_ = nullptr;
  return temp;
}
inline ::xla::ConvolutionDimensionNumbers* XnnConvolutionThunkProto::_internal_mutable_dimension_numbers() {
  
  if (_impl_.dimension_numbers_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::ConvolutionDimensionNumbers>(GetArenaForAllocation());
    _impl_.dimension_numbers_ = p;
  }
  return _impl_.dimension_numbers_;
}
inline ::xla::ConvolutionDimensionNumbers* XnnConvolutionThunkProto::mutable_dimension_numbers() {
  ::xla::ConvolutionDimensionNumbers* _msg = _internal_mutable_dimension_numbers();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnConvolutionThunkProto.dimension_numbers)
  return _msg;
}
inline void XnnConvolutionThunkProto::set_allocated_dimension_numbers(::xla::ConvolutionDimensionNumbers* dimension_numbers) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dimension_numbers_);
  }
  if (dimension_numbers) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(dimension_numbers));
    if (message_arena != submessage_arena) {
      dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, dimension_numbers, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.dimension_numbers_ = dimension_numbers;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnConvolutionThunkProto.dimension_numbers)
}

// .xla.Window window = 2;
inline bool XnnConvolutionThunkProto::_internal_has_window() const {
  return this != internal_default_instance() && _impl_.window_ != nullptr;
}
inline bool XnnConvolutionThunkProto::has_window() const {
  return _internal_has_window();
}
inline const ::xla::Window& XnnConvolutionThunkProto::_internal_window() const {
  const ::xla::Window* p = _impl_.window_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::Window&>(
      ::xla::_Window_default_instance_);
}
inline const ::xla::Window& XnnConvolutionThunkProto::window() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.window)
  return _internal_window();
}
inline void XnnConvolutionThunkProto::unsafe_arena_set_allocated_window(
    ::xla::Window* window) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.window_);
  }
  _impl_.window_ = window;
  if (window) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnConvolutionThunkProto.window)
}
inline ::xla::Window* XnnConvolutionThunkProto::release_window() {
  
  ::xla::Window* temp = _impl_.window_;
  _impl_.window_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::Window* XnnConvolutionThunkProto::unsafe_arena_release_window() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnConvolutionThunkProto.window)
  
  ::xla::Window* temp = _impl_.window_;
  _impl_.window_ = nullptr;
  return temp;
}
inline ::xla::Window* XnnConvolutionThunkProto::_internal_mutable_window() {
  
  if (_impl_.window_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::Window>(GetArenaForAllocation());
    _impl_.window_ = p;
  }
  return _impl_.window_;
}
inline ::xla::Window* XnnConvolutionThunkProto::mutable_window() {
  ::xla::Window* _msg = _internal_mutable_window();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnConvolutionThunkProto.window)
  return _msg;
}
inline void XnnConvolutionThunkProto::set_allocated_window(::xla::Window* window) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.window_);
  }
  if (window) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(window));
    if (message_arena != submessage_arena) {
      window = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, window, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.window_ = window;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnConvolutionThunkProto.window)
}

// int64 feature_group_count = 3;
inline void XnnConvolutionThunkProto::clear_feature_group_count() {
  _impl_.feature_group_count_ = int64_t{0};
}
inline int64_t XnnConvolutionThunkProto::_internal_feature_group_count() const {
  return _impl_.feature_group_count_;
}
inline int64_t XnnConvolutionThunkProto::feature_group_count() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.feature_group_count)
  return _internal_feature_group_count();
}
inline void XnnConvolutionThunkProto::_internal_set_feature_group_count(int64_t value) {
  
  _impl_.feature_group_count_ = value;
}
inline void XnnConvolutionThunkProto::set_feature_group_count(int64_t value) {
  _internal_set_feature_group_count(value);
  // @@protoc_insertion_point(field_set:xla.cpu.XnnConvolutionThunkProto.feature_group_count)
}

// .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
inline bool XnnConvolutionThunkProto::_internal_has_input_buffer_shape() const {
  return this != internal_default_instance() && _impl_.input_buffer_shape_ != nullptr;
}
inline bool XnnConvolutionThunkProto::has_input_buffer_shape() const {
  return _internal_has_input_buffer_shape();
}
inline void XnnConvolutionThunkProto::clear_input_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.input_buffer_shape_ != nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  _impl_.input_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::_internal_input_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.input_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::input_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.input_buffer_shape)
  return _internal_input_buffer_shape();
}
inline void XnnConvolutionThunkProto::unsafe_arena_set_allocated_input_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.input_buffer_shape_);
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  if (input_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnConvolutionThunkProto.input_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::release_input_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::unsafe_arena_release_input_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnConvolutionThunkProto.input_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::_internal_mutable_input_buffer_shape() {
  
  if (_impl_.input_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.input_buffer_shape_ = p;
  }
  return _impl_.input_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::mutable_input_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_input_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnConvolutionThunkProto.input_buffer_shape)
  return _msg;
}
inline void XnnConvolutionThunkProto::set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  if (input_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(input_buffer_shape);
    if (message_arena != submessage_arena) {
      input_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, input_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnConvolutionThunkProto.input_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto kernel_buffer_shape = 5;
inline bool XnnConvolutionThunkProto::_internal_has_kernel_buffer_shape() const {
  return this != internal_default_instance() && _impl_.kernel_buffer_shape_ != nullptr;
}
inline bool XnnConvolutionThunkProto::has_kernel_buffer_shape() const {
  return _internal_has_kernel_buffer_shape();
}
inline void XnnConvolutionThunkProto::clear_kernel_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.kernel_buffer_shape_ != nullptr) {
    delete _impl_.kernel_buffer_shape_;
  }
  _impl_.kernel_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::_internal_kernel_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.kernel_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::kernel_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.kernel_buffer_shape)
  return _internal_kernel_buffer_shape();
}
inline void XnnConvolutionThunkProto::unsafe_arena_set_allocated_kernel_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.kernel_buffer_shape_);
  }
  _impl_.kernel_buffer_shape_ = kernel_buffer_shape;
  if (kernel_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnConvolutionThunkProto.kernel_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::release_kernel_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.kernel_buffer_shape_;
  _impl_.kernel_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::unsafe_arena_release_kernel_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnConvolutionThunkProto.kernel_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.kernel_buffer_shape_;
  _impl_.kernel_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::_internal_mutable_kernel_buffer_shape() {
  
  if (_impl_.kernel_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.kernel_buffer_shape_ = p;
  }
  return _impl_.kernel_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::mutable_kernel_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_kernel_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnConvolutionThunkProto.kernel_buffer_shape)
  return _msg;
}
inline void XnnConvolutionThunkProto::set_allocated_kernel_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* kernel_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.kernel_buffer_shape_;
  }
  if (kernel_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(kernel_buffer_shape);
    if (message_arena != submessage_arena) {
      kernel_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, kernel_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.kernel_buffer_shape_ = kernel_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnConvolutionThunkProto.kernel_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 6;
inline bool XnnConvolutionThunkProto::_internal_has_output_buffer_shape() const {
  return this != internal_default_instance() && _impl_.output_buffer_shape_ != nullptr;
}
inline bool XnnConvolutionThunkProto::has_output_buffer_shape() const {
  return _internal_has_output_buffer_shape();
}
inline void XnnConvolutionThunkProto::clear_output_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.output_buffer_shape_ != nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  _impl_.output_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::_internal_output_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.output_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& XnnConvolutionThunkProto::output_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnConvolutionThunkProto.output_buffer_shape)
  return _internal_output_buffer_shape();
}
inline void XnnConvolutionThunkProto::unsafe_arena_set_allocated_output_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.output_buffer_shape_);
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  if (output_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnConvolutionThunkProto.output_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::release_output_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::unsafe_arena_release_output_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnConvolutionThunkProto.output_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::_internal_mutable_output_buffer_shape() {
  
  if (_impl_.output_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.output_buffer_shape_ = p;
  }
  return _impl_.output_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* XnnConvolutionThunkProto::mutable_output_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_output_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnConvolutionThunkProto.output_buffer_shape)
  return _msg;
}
inline void XnnConvolutionThunkProto::set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  if (output_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(output_buffer_shape);
    if (message_arena != submessage_arena) {
      output_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, output_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnConvolutionThunkProto.output_buffer_shape)
}

// -------------------------------------------------------------------

// XnnFusionThunkProto_Options

// bool use_threadpool = 1;
inline void XnnFusionThunkProto_Options::clear_use_threadpool() {
  _impl_.use_threadpool_ = false;
}
inline bool XnnFusionThunkProto_Options::_internal_use_threadpool() const {
  return _impl_.use_threadpool_;
}
inline bool XnnFusionThunkProto_Options::use_threadpool() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnFusionThunkProto.Options.use_threadpool)
  return _internal_use_threadpool();
}
inline void XnnFusionThunkProto_Options::_internal_set_use_threadpool(bool value) {
  
  _impl_.use_threadpool_ = value;
}
inline void XnnFusionThunkProto_Options::set_use_threadpool(bool value) {
  _internal_set_use_threadpool(value);
  // @@protoc_insertion_point(field_set:xla.cpu.XnnFusionThunkProto.Options.use_threadpool)
}

// -------------------------------------------------------------------

// XnnFusionThunkProto

// .xla.cpu.XnnFusionThunkProto.Options options = 1;
inline bool XnnFusionThunkProto::_internal_has_options() const {
  return this != internal_default_instance() && _impl_.options_ != nullptr;
}
inline bool XnnFusionThunkProto::has_options() const {
  return _internal_has_options();
}
inline void XnnFusionThunkProto::clear_options() {
  if (GetArenaForAllocation() == nullptr && _impl_.options_ != nullptr) {
    delete _impl_.options_;
  }
  _impl_.options_ = nullptr;
}
inline const ::xla::cpu::XnnFusionThunkProto_Options& XnnFusionThunkProto::_internal_options() const {
  const ::xla::cpu::XnnFusionThunkProto_Options* p = _impl_.options_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::XnnFusionThunkProto_Options&>(
      ::xla::cpu::_XnnFusionThunkProto_Options_default_instance_);
}
inline const ::xla::cpu::XnnFusionThunkProto_Options& XnnFusionThunkProto::options() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnFusionThunkProto.options)
  return _internal_options();
}
inline void XnnFusionThunkProto::unsafe_arena_set_allocated_options(
    ::xla::cpu::XnnFusionThunkProto_Options* options) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.options_);
  }
  _impl_.options_ = options;
  if (options) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnFusionThunkProto.options)
}
inline ::xla::cpu::XnnFusionThunkProto_Options* XnnFusionThunkProto::release_options() {
  
  ::xla::cpu::XnnFusionThunkProto_Options* temp = _impl_.options_;
  _impl_.options_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::XnnFusionThunkProto_Options* XnnFusionThunkProto::unsafe_arena_release_options() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnFusionThunkProto.options)
  
  ::xla::cpu::XnnFusionThunkProto_Options* temp = _impl_.options_;
  _impl_.options_ = nullptr;
  return temp;
}
inline ::xla::cpu::XnnFusionThunkProto_Options* XnnFusionThunkProto::_internal_mutable_options() {
  
  if (_impl_.options_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::XnnFusionThunkProto_Options>(GetArenaForAllocation());
    _impl_.options_ = p;
  }
  return _impl_.options_;
}
inline ::xla::cpu::XnnFusionThunkProto_Options* XnnFusionThunkProto::mutable_options() {
  ::xla::cpu::XnnFusionThunkProto_Options* _msg = _internal_mutable_options();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnFusionThunkProto.options)
  return _msg;
}
inline void XnnFusionThunkProto::set_allocated_options(::xla::cpu::XnnFusionThunkProto_Options* options) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.options_;
  }
  if (options) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(options);
    if (message_arena != submessage_arena) {
      options = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, options, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.options_ = options;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.XnnFusionThunkProto.options)
}

// .xla.cpu.XnnDotThunkProto xnn_dot_thunk = 2;
inline bool XnnFusionThunkProto::_internal_has_xnn_dot_thunk() const {
  return impl_case() == kXnnDotThunk;
}
inline bool XnnFusionThunkProto::has_xnn_dot_thunk() const {
  return _internal_has_xnn_dot_thunk();
}
inline void XnnFusionThunkProto::set_has_xnn_dot_thunk() {
  _impl_._oneof_case_[0] = kXnnDotThunk;
}
inline void XnnFusionThunkProto::clear_xnn_dot_thunk() {
  if (_internal_has_xnn_dot_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.xnn_dot_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::XnnDotThunkProto* XnnFusionThunkProto::release_xnn_dot_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnFusionThunkProto.xnn_dot_thunk)
  if (_internal_has_xnn_dot_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnDotThunkProto* temp = _impl_.impl_.xnn_dot_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.xnn_dot_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::XnnDotThunkProto& XnnFusionThunkProto::_internal_xnn_dot_thunk() const {
  return _internal_has_xnn_dot_thunk()
      ? *_impl_.impl_.xnn_dot_thunk_
      : reinterpret_cast< ::xla::cpu::XnnDotThunkProto&>(::xla::cpu::_XnnDotThunkProto_default_instance_);
}
inline const ::xla::cpu::XnnDotThunkProto& XnnFusionThunkProto::xnn_dot_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnFusionThunkProto.xnn_dot_thunk)
  return _internal_xnn_dot_thunk();
}
inline ::xla::cpu::XnnDotThunkProto* XnnFusionThunkProto::unsafe_arena_release_xnn_dot_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.XnnFusionThunkProto.xnn_dot_thunk)
  if (_internal_has_xnn_dot_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnDotThunkProto* temp = _impl_.impl_.xnn_dot_thunk_;
    _impl_.impl_.xnn_dot_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void XnnFusionThunkProto::unsafe_arena_set_allocated_xnn_dot_thunk(::xla::cpu::XnnDotThunkProto* xnn_dot_thunk) {
  clear_impl();
  if (xnn_dot_thunk) {
    set_has_xnn_dot_thunk();
    _impl_.impl_.xnn_dot_thunk_ = xnn_dot_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnFusionThunkProto.xnn_dot_thunk)
}
inline ::xla::cpu::XnnDotThunkProto* XnnFusionThunkProto::_internal_mutable_xnn_dot_thunk() {
  if (!_internal_has_xnn_dot_thunk()) {
    clear_impl();
    set_has_xnn_dot_thunk();
    _impl_.impl_.xnn_dot_thunk_ = CreateMaybeMessage< ::xla::cpu::XnnDotThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.xnn_dot_thunk_;
}
inline ::xla::cpu::XnnDotThunkProto* XnnFusionThunkProto::mutable_xnn_dot_thunk() {
  ::xla::cpu::XnnDotThunkProto* _msg = _internal_mutable_xnn_dot_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnFusionThunkProto.xnn_dot_thunk)
  return _msg;
}

// .xla.cpu.XnnConvolutionThunkProto xnn_convolution_thunk = 3;
inline bool XnnFusionThunkProto::_internal_has_xnn_convolution_thunk() const {
  return impl_case() == kXnnConvolutionThunk;
}
inline bool XnnFusionThunkProto::has_xnn_convolution_thunk() const {
  return _internal_has_xnn_convolution_thunk();
}
inline void XnnFusionThunkProto::set_has_xnn_convolution_thunk() {
  _impl_._oneof_case_[0] = kXnnConvolutionThunk;
}
inline void XnnFusionThunkProto::clear_xnn_convolution_thunk() {
  if (_internal_has_xnn_convolution_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.xnn_convolution_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::XnnConvolutionThunkProto* XnnFusionThunkProto::release_xnn_convolution_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.XnnFusionThunkProto.xnn_convolution_thunk)
  if (_internal_has_xnn_convolution_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnConvolutionThunkProto* temp = _impl_.impl_.xnn_convolution_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.xnn_convolution_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::XnnConvolutionThunkProto& XnnFusionThunkProto::_internal_xnn_convolution_thunk() const {
  return _internal_has_xnn_convolution_thunk()
      ? *_impl_.impl_.xnn_convolution_thunk_
      : reinterpret_cast< ::xla::cpu::XnnConvolutionThunkProto&>(::xla::cpu::_XnnConvolutionThunkProto_default_instance_);
}
inline const ::xla::cpu::XnnConvolutionThunkProto& XnnFusionThunkProto::xnn_convolution_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.XnnFusionThunkProto.xnn_convolution_thunk)
  return _internal_xnn_convolution_thunk();
}
inline ::xla::cpu::XnnConvolutionThunkProto* XnnFusionThunkProto::unsafe_arena_release_xnn_convolution_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.XnnFusionThunkProto.xnn_convolution_thunk)
  if (_internal_has_xnn_convolution_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnConvolutionThunkProto* temp = _impl_.impl_.xnn_convolution_thunk_;
    _impl_.impl_.xnn_convolution_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void XnnFusionThunkProto::unsafe_arena_set_allocated_xnn_convolution_thunk(::xla::cpu::XnnConvolutionThunkProto* xnn_convolution_thunk) {
  clear_impl();
  if (xnn_convolution_thunk) {
    set_has_xnn_convolution_thunk();
    _impl_.impl_.xnn_convolution_thunk_ = xnn_convolution_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.XnnFusionThunkProto.xnn_convolution_thunk)
}
inline ::xla::cpu::XnnConvolutionThunkProto* XnnFusionThunkProto::_internal_mutable_xnn_convolution_thunk() {
  if (!_internal_has_xnn_convolution_thunk()) {
    clear_impl();
    set_has_xnn_convolution_thunk();
    _impl_.impl_.xnn_convolution_thunk_ = CreateMaybeMessage< ::xla::cpu::XnnConvolutionThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.xnn_convolution_thunk_;
}
inline ::xla::cpu::XnnConvolutionThunkProto* XnnFusionThunkProto::mutable_xnn_convolution_thunk() {
  ::xla::cpu::XnnConvolutionThunkProto* _msg = _internal_mutable_xnn_convolution_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.XnnFusionThunkProto.xnn_convolution_thunk)
  return _msg;
}

inline bool XnnFusionThunkProto::has_impl() const {
  return impl_case() != IMPL_NOT_SET;
}
inline void XnnFusionThunkProto::clear_has_impl() {
  _impl_._oneof_case_[0] = IMPL_NOT_SET;
}
inline XnnFusionThunkProto::ImplCase XnnFusionThunkProto::impl_case() const {
  return XnnFusionThunkProto::ImplCase(_impl_._oneof_case_[0]);
}
// -------------------------------------------------------------------

// DotThunkProto

// .xla.DotDimensionNumbers dot_dimensions = 1;
inline bool DotThunkProto::_internal_has_dot_dimensions() const {
  return this != internal_default_instance() && _impl_.dot_dimensions_ != nullptr;
}
inline bool DotThunkProto::has_dot_dimensions() const {
  return _internal_has_dot_dimensions();
}
inline const ::xla::DotDimensionNumbers& DotThunkProto::_internal_dot_dimensions() const {
  const ::xla::DotDimensionNumbers* p = _impl_.dot_dimensions_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::DotDimensionNumbers&>(
      ::xla::_DotDimensionNumbers_default_instance_);
}
inline const ::xla::DotDimensionNumbers& DotThunkProto::dot_dimensions() const {
  // @@protoc_insertion_point(field_get:xla.cpu.DotThunkProto.dot_dimensions)
  return _internal_dot_dimensions();
}
inline void DotThunkProto::unsafe_arena_set_allocated_dot_dimensions(
    ::xla::DotDimensionNumbers* dot_dimensions) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimensions_);
  }
  _impl_.dot_dimensions_ = dot_dimensions;
  if (dot_dimensions) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.DotThunkProto.dot_dimensions)
}
inline ::xla::DotDimensionNumbers* DotThunkProto::release_dot_dimensions() {
  
  ::xla::DotDimensionNumbers* temp = _impl_.dot_dimensions_;
  _impl_.dot_dimensions_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::DotDimensionNumbers* DotThunkProto::unsafe_arena_release_dot_dimensions() {
  // @@protoc_insertion_point(field_release:xla.cpu.DotThunkProto.dot_dimensions)
  
  ::xla::DotDimensionNumbers* temp = _impl_.dot_dimensions_;
  _impl_.dot_dimensions_ = nullptr;
  return temp;
}
inline ::xla::DotDimensionNumbers* DotThunkProto::_internal_mutable_dot_dimensions() {
  
  if (_impl_.dot_dimensions_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation());
    _impl_.dot_dimensions_ = p;
  }
  return _impl_.dot_dimensions_;
}
inline ::xla::DotDimensionNumbers* DotThunkProto::mutable_dot_dimensions() {
  ::xla::DotDimensionNumbers* _msg = _internal_mutable_dot_dimensions();
  // @@protoc_insertion_point(field_mutable:xla.cpu.DotThunkProto.dot_dimensions)
  return _msg;
}
inline void DotThunkProto::set_allocated_dot_dimensions(::xla::DotDimensionNumbers* dot_dimensions) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimensions_);
  }
  if (dot_dimensions) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(
                reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(dot_dimensions));
    if (message_arena != submessage_arena) {
      dot_dimensions = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, dot_dimensions, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.dot_dimensions_ = dot_dimensions;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.DotThunkProto.dot_dimensions)
}

// .xla.cpu.ShapeBufferAllocationSliceProto lhs_buffer_shape = 2;
inline bool DotThunkProto::_internal_has_lhs_buffer_shape() const {
  return this != internal_default_instance() && _impl_.lhs_buffer_shape_ != nullptr;
}
inline bool DotThunkProto::has_lhs_buffer_shape() const {
  return _internal_has_lhs_buffer_shape();
}
inline void DotThunkProto::clear_lhs_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.lhs_buffer_shape_ != nullptr) {
    delete _impl_.lhs_buffer_shape_;
  }
  _impl_.lhs_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::_internal_lhs_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.lhs_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::lhs_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.DotThunkProto.lhs_buffer_shape)
  return _internal_lhs_buffer_shape();
}
inline void DotThunkProto::unsafe_arena_set_allocated_lhs_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.lhs_buffer_shape_);
  }
  _impl_.lhs_buffer_shape_ = lhs_buffer_shape;
  if (lhs_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.DotThunkProto.lhs_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::release_lhs_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.lhs_buffer_shape_;
  _impl_.lhs_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::unsafe_arena_release_lhs_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.DotThunkProto.lhs_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.lhs_buffer_shape_;
  _impl_.lhs_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::_internal_mutable_lhs_buffer_shape() {
  
  if (_impl_.lhs_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.lhs_buffer_shape_ = p;
  }
  return _impl_.lhs_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::mutable_lhs_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_lhs_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.DotThunkProto.lhs_buffer_shape)
  return _msg;
}
inline void DotThunkProto::set_allocated_lhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* lhs_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.lhs_buffer_shape_;
  }
  if (lhs_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(lhs_buffer_shape);
    if (message_arena != submessage_arena) {
      lhs_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, lhs_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.lhs_buffer_shape_ = lhs_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.DotThunkProto.lhs_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto rhs_buffer_shape = 3;
inline bool DotThunkProto::_internal_has_rhs_buffer_shape() const {
  return this != internal_default_instance() && _impl_.rhs_buffer_shape_ != nullptr;
}
inline bool DotThunkProto::has_rhs_buffer_shape() const {
  return _internal_has_rhs_buffer_shape();
}
inline void DotThunkProto::clear_rhs_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.rhs_buffer_shape_ != nullptr) {
    delete _impl_.rhs_buffer_shape_;
  }
  _impl_.rhs_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::_internal_rhs_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.rhs_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::rhs_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.DotThunkProto.rhs_buffer_shape)
  return _internal_rhs_buffer_shape();
}
inline void DotThunkProto::unsafe_arena_set_allocated_rhs_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.rhs_buffer_shape_);
  }
  _impl_.rhs_buffer_shape_ = rhs_buffer_shape;
  if (rhs_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.DotThunkProto.rhs_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::release_rhs_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.rhs_buffer_shape_;
  _impl_.rhs_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::unsafe_arena_release_rhs_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.DotThunkProto.rhs_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.rhs_buffer_shape_;
  _impl_.rhs_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::_internal_mutable_rhs_buffer_shape() {
  
  if (_impl_.rhs_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.rhs_buffer_shape_ = p;
  }
  return _impl_.rhs_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::mutable_rhs_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_rhs_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.DotThunkProto.rhs_buffer_shape)
  return _msg;
}
inline void DotThunkProto::set_allocated_rhs_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* rhs_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.rhs_buffer_shape_;
  }
  if (rhs_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(rhs_buffer_shape);
    if (message_arena != submessage_arena) {
      rhs_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, rhs_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.rhs_buffer_shape_ = rhs_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.DotThunkProto.rhs_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto out_buffer_shape = 4;
inline bool DotThunkProto::_internal_has_out_buffer_shape() const {
  return this != internal_default_instance() && _impl_.out_buffer_shape_ != nullptr;
}
inline bool DotThunkProto::has_out_buffer_shape() const {
  return _internal_has_out_buffer_shape();
}
inline void DotThunkProto::clear_out_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.out_buffer_shape_ != nullptr) {
    delete _impl_.out_buffer_shape_;
  }
  _impl_.out_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::_internal_out_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.out_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& DotThunkProto::out_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.DotThunkProto.out_buffer_shape)
  return _internal_out_buffer_shape();
}
inline void DotThunkProto::unsafe_arena_set_allocated_out_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.out_buffer_shape_);
  }
  _impl_.out_buffer_shape_ = out_buffer_shape;
  if (out_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.DotThunkProto.out_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::release_out_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.out_buffer_shape_;
  _impl_.out_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::unsafe_arena_release_out_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.DotThunkProto.out_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.out_buffer_shape_;
  _impl_.out_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::_internal_mutable_out_buffer_shape() {
  
  if (_impl_.out_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.out_buffer_shape_ = p;
  }
  return _impl_.out_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* DotThunkProto::mutable_out_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_out_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.DotThunkProto.out_buffer_shape)
  return _msg;
}
inline void DotThunkProto::set_allocated_out_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* out_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.out_buffer_shape_;
  }
  if (out_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(out_buffer_shape);
    if (message_arena != submessage_arena) {
      out_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, out_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.out_buffer_shape_ = out_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.DotThunkProto.out_buffer_shape)
}

// -------------------------------------------------------------------

// RngGetAndUpdateStateThunkProto

// int64 delta = 1;
inline void RngGetAndUpdateStateThunkProto::clear_delta() {
  _impl_.delta_ = int64_t{0};
}
inline int64_t RngGetAndUpdateStateThunkProto::_internal_delta() const {
  return _impl_.delta_;
}
inline int64_t RngGetAndUpdateStateThunkProto::delta() const {
  // @@protoc_insertion_point(field_get:xla.cpu.RngGetAndUpdateStateThunkProto.delta)
  return _internal_delta();
}
inline void RngGetAndUpdateStateThunkProto::_internal_set_delta(int64_t value) {
  
  _impl_.delta_ = value;
}
inline void RngGetAndUpdateStateThunkProto::set_delta(int64_t value) {
  _internal_set_delta(value);
  // @@protoc_insertion_point(field_set:xla.cpu.RngGetAndUpdateStateThunkProto.delta)
}

// .xla.cpu.BufferAllocationSliceProto state_buffer = 2;
inline bool RngGetAndUpdateStateThunkProto::_internal_has_state_buffer() const {
  return this != internal_default_instance() && _impl_.state_buffer_ != nullptr;
}
inline bool RngGetAndUpdateStateThunkProto::has_state_buffer() const {
  return _internal_has_state_buffer();
}
inline void RngGetAndUpdateStateThunkProto::clear_state_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.state_buffer_ != nullptr) {
    delete _impl_.state_buffer_;
  }
  _impl_.state_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& RngGetAndUpdateStateThunkProto::_internal_state_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.state_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& RngGetAndUpdateStateThunkProto::state_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.RngGetAndUpdateStateThunkProto.state_buffer)
  return _internal_state_buffer();
}
inline void RngGetAndUpdateStateThunkProto::unsafe_arena_set_allocated_state_buffer(
    ::xla::cpu::BufferAllocationSliceProto* state_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.state_buffer_);
  }
  _impl_.state_buffer_ = state_buffer;
  if (state_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.RngGetAndUpdateStateThunkProto.state_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* RngGetAndUpdateStateThunkProto::release_state_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.state_buffer_;
  _impl_.state_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* RngGetAndUpdateStateThunkProto::unsafe_arena_release_state_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.RngGetAndUpdateStateThunkProto.state_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.state_buffer_;
  _impl_.state_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* RngGetAndUpdateStateThunkProto::_internal_mutable_state_buffer() {
  
  if (_impl_.state_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.state_buffer_ = p;
  }
  return _impl_.state_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* RngGetAndUpdateStateThunkProto::mutable_state_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_state_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.RngGetAndUpdateStateThunkProto.state_buffer)
  return _msg;
}
inline void RngGetAndUpdateStateThunkProto::set_allocated_state_buffer(::xla::cpu::BufferAllocationSliceProto* state_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.state_buffer_;
  }
  if (state_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(state_buffer);
    if (message_arena != submessage_arena) {
      state_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, state_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.state_buffer_ = state_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.RngGetAndUpdateStateThunkProto.state_buffer)
}

// -------------------------------------------------------------------

// TopKThunkProto

// int64 batch_size = 1;
inline void TopKThunkProto::clear_batch_size() {
  _impl_.batch_size_ = int64_t{0};
}
inline int64_t TopKThunkProto::_internal_batch_size() const {
  return _impl_.batch_size_;
}
inline int64_t TopKThunkProto::batch_size() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.batch_size)
  return _internal_batch_size();
}
inline void TopKThunkProto::_internal_set_batch_size(int64_t value) {
  
  _impl_.batch_size_ = value;
}
inline void TopKThunkProto::set_batch_size(int64_t value) {
  _internal_set_batch_size(value);
  // @@protoc_insertion_point(field_set:xla.cpu.TopKThunkProto.batch_size)
}

// int64 input_size = 2;
inline void TopKThunkProto::clear_input_size() {
  _impl_.input_size_ = int64_t{0};
}
inline int64_t TopKThunkProto::_internal_input_size() const {
  return _impl_.input_size_;
}
inline int64_t TopKThunkProto::input_size() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.input_size)
  return _internal_input_size();
}
inline void TopKThunkProto::_internal_set_input_size(int64_t value) {
  
  _impl_.input_size_ = value;
}
inline void TopKThunkProto::set_input_size(int64_t value) {
  _internal_set_input_size(value);
  // @@protoc_insertion_point(field_set:xla.cpu.TopKThunkProto.input_size)
}

// int64 k = 3;
inline void TopKThunkProto::clear_k() {
  _impl_.k_ = int64_t{0};
}
inline int64_t TopKThunkProto::_internal_k() const {
  return _impl_.k_;
}
inline int64_t TopKThunkProto::k() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.k)
  return _internal_k();
}
inline void TopKThunkProto::_internal_set_k(int64_t value) {
  
  _impl_.k_ = value;
}
inline void TopKThunkProto::set_k(int64_t value) {
  _internal_set_k(value);
  // @@protoc_insertion_point(field_set:xla.cpu.TopKThunkProto.k)
}

// .xla.cpu.BufferAllocationSliceProto values_buffer = 4;
inline bool TopKThunkProto::_internal_has_values_buffer() const {
  return this != internal_default_instance() && _impl_.values_buffer_ != nullptr;
}
inline bool TopKThunkProto::has_values_buffer() const {
  return _internal_has_values_buffer();
}
inline void TopKThunkProto::clear_values_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.values_buffer_ != nullptr) {
    delete _impl_.values_buffer_;
  }
  _impl_.values_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::_internal_values_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.values_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::values_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.values_buffer)
  return _internal_values_buffer();
}
inline void TopKThunkProto::unsafe_arena_set_allocated_values_buffer(
    ::xla::cpu::BufferAllocationSliceProto* values_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.values_buffer_);
  }
  _impl_.values_buffer_ = values_buffer;
  if (values_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.TopKThunkProto.values_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::release_values_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.values_buffer_;
  _impl_.values_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::unsafe_arena_release_values_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.TopKThunkProto.values_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.values_buffer_;
  _impl_.values_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::_internal_mutable_values_buffer() {
  
  if (_impl_.values_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.values_buffer_ = p;
  }
  return _impl_.values_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::mutable_values_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_values_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.TopKThunkProto.values_buffer)
  return _msg;
}
inline void TopKThunkProto::set_allocated_values_buffer(::xla::cpu::BufferAllocationSliceProto* values_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.values_buffer_;
  }
  if (values_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(values_buffer);
    if (message_arena != submessage_arena) {
      values_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, values_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.values_buffer_ = values_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.TopKThunkProto.values_buffer)
}

// .xla.cpu.BufferAllocationSliceProto output_buffer = 5;
inline bool TopKThunkProto::_internal_has_output_buffer() const {
  return this != internal_default_instance() && _impl_.output_buffer_ != nullptr;
}
inline bool TopKThunkProto::has_output_buffer() const {
  return _internal_has_output_buffer();
}
inline void TopKThunkProto::clear_output_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.output_buffer_ != nullptr) {
    delete _impl_.output_buffer_;
  }
  _impl_.output_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::_internal_output_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.output_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::output_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.output_buffer)
  return _internal_output_buffer();
}
inline void TopKThunkProto::unsafe_arena_set_allocated_output_buffer(
    ::xla::cpu::BufferAllocationSliceProto* output_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.output_buffer_);
  }
  _impl_.output_buffer_ = output_buffer;
  if (output_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.TopKThunkProto.output_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::release_output_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.output_buffer_;
  _impl_.output_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::unsafe_arena_release_output_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.TopKThunkProto.output_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.output_buffer_;
  _impl_.output_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::_internal_mutable_output_buffer() {
  
  if (_impl_.output_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.output_buffer_ = p;
  }
  return _impl_.output_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::mutable_output_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_output_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.TopKThunkProto.output_buffer)
  return _msg;
}
inline void TopKThunkProto::set_allocated_output_buffer(::xla::cpu::BufferAllocationSliceProto* output_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.output_buffer_;
  }
  if (output_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(output_buffer);
    if (message_arena != submessage_arena) {
      output_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, output_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.output_buffer_ = output_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.TopKThunkProto.output_buffer)
}

// .xla.cpu.BufferAllocationSliceProto indices_buffer = 6;
inline bool TopKThunkProto::_internal_has_indices_buffer() const {
  return this != internal_default_instance() && _impl_.indices_buffer_ != nullptr;
}
inline bool TopKThunkProto::has_indices_buffer() const {
  return _internal_has_indices_buffer();
}
inline void TopKThunkProto::clear_indices_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.indices_buffer_ != nullptr) {
    delete _impl_.indices_buffer_;
  }
  _impl_.indices_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::_internal_indices_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.indices_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& TopKThunkProto::indices_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.TopKThunkProto.indices_buffer)
  return _internal_indices_buffer();
}
inline void TopKThunkProto::unsafe_arena_set_allocated_indices_buffer(
    ::xla::cpu::BufferAllocationSliceProto* indices_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.indices_buffer_);
  }
  _impl_.indices_buffer_ = indices_buffer;
  if (indices_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.TopKThunkProto.indices_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::release_indices_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.indices_buffer_;
  _impl_.indices_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::unsafe_arena_release_indices_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.TopKThunkProto.indices_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.indices_buffer_;
  _impl_.indices_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::_internal_mutable_indices_buffer() {
  
  if (_impl_.indices_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.indices_buffer_ = p;
  }
  return _impl_.indices_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* TopKThunkProto::mutable_indices_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_indices_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.TopKThunkProto.indices_buffer)
  return _msg;
}
inline void TopKThunkProto::set_allocated_indices_buffer(::xla::cpu::BufferAllocationSliceProto* indices_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.indices_buffer_;
  }
  if (indices_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(indices_buffer);
    if (message_arena != submessage_arena) {
      indices_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, indices_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.indices_buffer_ = indices_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.TopKThunkProto.indices_buffer)
}

// -------------------------------------------------------------------

// WhileThunkProto

// .xla.cpu.ThunkSequenceProto cond_sequence = 1;
inline bool WhileThunkProto::_internal_has_cond_sequence() const {
  return this != internal_default_instance() && _impl_.cond_sequence_ != nullptr;
}
inline bool WhileThunkProto::has_cond_sequence() const {
  return _internal_has_cond_sequence();
}
inline void WhileThunkProto::clear_cond_sequence() {
  if (GetArenaForAllocation() == nullptr && _impl_.cond_sequence_ != nullptr) {
    delete _impl_.cond_sequence_;
  }
  _impl_.cond_sequence_ = nullptr;
}
inline const ::xla::cpu::ThunkSequenceProto& WhileThunkProto::_internal_cond_sequence() const {
  const ::xla::cpu::ThunkSequenceProto* p = _impl_.cond_sequence_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ThunkSequenceProto&>(
      ::xla::cpu::_ThunkSequenceProto_default_instance_);
}
inline const ::xla::cpu::ThunkSequenceProto& WhileThunkProto::cond_sequence() const {
  // @@protoc_insertion_point(field_get:xla.cpu.WhileThunkProto.cond_sequence)
  return _internal_cond_sequence();
}
inline void WhileThunkProto::unsafe_arena_set_allocated_cond_sequence(
    ::xla::cpu::ThunkSequenceProto* cond_sequence) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.cond_sequence_);
  }
  _impl_.cond_sequence_ = cond_sequence;
  if (cond_sequence) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.WhileThunkProto.cond_sequence)
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::release_cond_sequence() {
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.cond_sequence_;
  _impl_.cond_sequence_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::unsafe_arena_release_cond_sequence() {
  // @@protoc_insertion_point(field_release:xla.cpu.WhileThunkProto.cond_sequence)
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.cond_sequence_;
  _impl_.cond_sequence_ = nullptr;
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::_internal_mutable_cond_sequence() {
  
  if (_impl_.cond_sequence_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ThunkSequenceProto>(GetArenaForAllocation());
    _impl_.cond_sequence_ = p;
  }
  return _impl_.cond_sequence_;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::mutable_cond_sequence() {
  ::xla::cpu::ThunkSequenceProto* _msg = _internal_mutable_cond_sequence();
  // @@protoc_insertion_point(field_mutable:xla.cpu.WhileThunkProto.cond_sequence)
  return _msg;
}
inline void WhileThunkProto::set_allocated_cond_sequence(::xla::cpu::ThunkSequenceProto* cond_sequence) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.cond_sequence_;
  }
  if (cond_sequence) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(cond_sequence);
    if (message_arena != submessage_arena) {
      cond_sequence = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, cond_sequence, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.cond_sequence_ = cond_sequence;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.WhileThunkProto.cond_sequence)
}

// .xla.cpu.ThunkSequenceProto body_sequence = 2;
inline bool WhileThunkProto::_internal_has_body_sequence() const {
  return this != internal_default_instance() && _impl_.body_sequence_ != nullptr;
}
inline bool WhileThunkProto::has_body_sequence() const {
  return _internal_has_body_sequence();
}
inline void WhileThunkProto::clear_body_sequence() {
  if (GetArenaForAllocation() == nullptr && _impl_.body_sequence_ != nullptr) {
    delete _impl_.body_sequence_;
  }
  _impl_.body_sequence_ = nullptr;
}
inline const ::xla::cpu::ThunkSequenceProto& WhileThunkProto::_internal_body_sequence() const {
  const ::xla::cpu::ThunkSequenceProto* p = _impl_.body_sequence_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ThunkSequenceProto&>(
      ::xla::cpu::_ThunkSequenceProto_default_instance_);
}
inline const ::xla::cpu::ThunkSequenceProto& WhileThunkProto::body_sequence() const {
  // @@protoc_insertion_point(field_get:xla.cpu.WhileThunkProto.body_sequence)
  return _internal_body_sequence();
}
inline void WhileThunkProto::unsafe_arena_set_allocated_body_sequence(
    ::xla::cpu::ThunkSequenceProto* body_sequence) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.body_sequence_);
  }
  _impl_.body_sequence_ = body_sequence;
  if (body_sequence) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.WhileThunkProto.body_sequence)
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::release_body_sequence() {
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.body_sequence_;
  _impl_.body_sequence_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::unsafe_arena_release_body_sequence() {
  // @@protoc_insertion_point(field_release:xla.cpu.WhileThunkProto.body_sequence)
  
  ::xla::cpu::ThunkSequenceProto* temp = _impl_.body_sequence_;
  _impl_.body_sequence_ = nullptr;
  return temp;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::_internal_mutable_body_sequence() {
  
  if (_impl_.body_sequence_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ThunkSequenceProto>(GetArenaForAllocation());
    _impl_.body_sequence_ = p;
  }
  return _impl_.body_sequence_;
}
inline ::xla::cpu::ThunkSequenceProto* WhileThunkProto::mutable_body_sequence() {
  ::xla::cpu::ThunkSequenceProto* _msg = _internal_mutable_body_sequence();
  // @@protoc_insertion_point(field_mutable:xla.cpu.WhileThunkProto.body_sequence)
  return _msg;
}
inline void WhileThunkProto::set_allocated_body_sequence(::xla::cpu::ThunkSequenceProto* body_sequence) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.body_sequence_;
  }
  if (body_sequence) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(body_sequence);
    if (message_arena != submessage_arena) {
      body_sequence = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, body_sequence, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.body_sequence_ = body_sequence;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.WhileThunkProto.body_sequence)
}

// .xla.cpu.Int64Optional trip_count = 3;
inline bool WhileThunkProto::_internal_has_trip_count() const {
  return this != internal_default_instance() && _impl_.trip_count_ != nullptr;
}
inline bool WhileThunkProto::has_trip_count() const {
  return _internal_has_trip_count();
}
inline void WhileThunkProto::clear_trip_count() {
  if (GetArenaForAllocation() == nullptr && _impl_.trip_count_ != nullptr) {
    delete _impl_.trip_count_;
  }
  _impl_.trip_count_ = nullptr;
}
inline const ::xla::cpu::Int64Optional& WhileThunkProto::_internal_trip_count() const {
  const ::xla::cpu::Int64Optional* p = _impl_.trip_count_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::Int64Optional&>(
      ::xla::cpu::_Int64Optional_default_instance_);
}
inline const ::xla::cpu::Int64Optional& WhileThunkProto::trip_count() const {
  // @@protoc_insertion_point(field_get:xla.cpu.WhileThunkProto.trip_count)
  return _internal_trip_count();
}
inline void WhileThunkProto::unsafe_arena_set_allocated_trip_count(
    ::xla::cpu::Int64Optional* trip_count) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.trip_count_);
  }
  _impl_.trip_count_ = trip_count;
  if (trip_count) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.WhileThunkProto.trip_count)
}
inline ::xla::cpu::Int64Optional* WhileThunkProto::release_trip_count() {
  
  ::xla::cpu::Int64Optional* temp = _impl_.trip_count_;
  _impl_.trip_count_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::Int64Optional* WhileThunkProto::unsafe_arena_release_trip_count() {
  // @@protoc_insertion_point(field_release:xla.cpu.WhileThunkProto.trip_count)
  
  ::xla::cpu::Int64Optional* temp = _impl_.trip_count_;
  _impl_.trip_count_ = nullptr;
  return temp;
}
inline ::xla::cpu::Int64Optional* WhileThunkProto::_internal_mutable_trip_count() {
  
  if (_impl_.trip_count_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::Int64Optional>(GetArenaForAllocation());
    _impl_.trip_count_ = p;
  }
  return _impl_.trip_count_;
}
inline ::xla::cpu::Int64Optional* WhileThunkProto::mutable_trip_count() {
  ::xla::cpu::Int64Optional* _msg = _internal_mutable_trip_count();
  // @@protoc_insertion_point(field_mutable:xla.cpu.WhileThunkProto.trip_count)
  return _msg;
}
inline void WhileThunkProto::set_allocated_trip_count(::xla::cpu::Int64Optional* trip_count) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.trip_count_;
  }
  if (trip_count) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(trip_count);
    if (message_arena != submessage_arena) {
      trip_count = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, trip_count, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.trip_count_ = trip_count;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.WhileThunkProto.trip_count)
}

// .xla.cpu.BufferAllocationSliceProto cond_buffer = 4;
inline bool WhileThunkProto::_internal_has_cond_buffer() const {
  return this != internal_default_instance() && _impl_.cond_buffer_ != nullptr;
}
inline bool WhileThunkProto::has_cond_buffer() const {
  return _internal_has_cond_buffer();
}
inline void WhileThunkProto::clear_cond_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.cond_buffer_ != nullptr) {
    delete _impl_.cond_buffer_;
  }
  _impl_.cond_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& WhileThunkProto::_internal_cond_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.cond_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& WhileThunkProto::cond_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.WhileThunkProto.cond_buffer)
  return _internal_cond_buffer();
}
inline void WhileThunkProto::unsafe_arena_set_allocated_cond_buffer(
    ::xla::cpu::BufferAllocationSliceProto* cond_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.cond_buffer_);
  }
  _impl_.cond_buffer_ = cond_buffer;
  if (cond_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.WhileThunkProto.cond_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* WhileThunkProto::release_cond_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.cond_buffer_;
  _impl_.cond_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* WhileThunkProto::unsafe_arena_release_cond_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.WhileThunkProto.cond_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.cond_buffer_;
  _impl_.cond_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* WhileThunkProto::_internal_mutable_cond_buffer() {
  
  if (_impl_.cond_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.cond_buffer_ = p;
  }
  return _impl_.cond_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* WhileThunkProto::mutable_cond_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_cond_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.WhileThunkProto.cond_buffer)
  return _msg;
}
inline void WhileThunkProto::set_allocated_cond_buffer(::xla::cpu::BufferAllocationSliceProto* cond_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.cond_buffer_;
  }
  if (cond_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(cond_buffer);
    if (message_arena != submessage_arena) {
      cond_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, cond_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.cond_buffer_ = cond_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.WhileThunkProto.cond_buffer)
}

// -------------------------------------------------------------------

// KernelThunkProto_ThreadDim

// int64 x = 1;
inline void KernelThunkProto_ThreadDim::clear_x() {
  _impl_.x_ = int64_t{0};
}
inline int64_t KernelThunkProto_ThreadDim::_internal_x() const {
  return _impl_.x_;
}
inline int64_t KernelThunkProto_ThreadDim::x() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.ThreadDim.x)
  return _internal_x();
}
inline void KernelThunkProto_ThreadDim::_internal_set_x(int64_t value) {
  
  _impl_.x_ = value;
}
inline void KernelThunkProto_ThreadDim::set_x(int64_t value) {
  _internal_set_x(value);
  // @@protoc_insertion_point(field_set:xla.cpu.KernelThunkProto.ThreadDim.x)
}

// int64 y = 2;
inline void KernelThunkProto_ThreadDim::clear_y() {
  _impl_.y_ = int64_t{0};
}
inline int64_t KernelThunkProto_ThreadDim::_internal_y() const {
  return _impl_.y_;
}
inline int64_t KernelThunkProto_ThreadDim::y() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.ThreadDim.y)
  return _internal_y();
}
inline void KernelThunkProto_ThreadDim::_internal_set_y(int64_t value) {
  
  _impl_.y_ = value;
}
inline void KernelThunkProto_ThreadDim::set_y(int64_t value) {
  _internal_set_y(value);
  // @@protoc_insertion_point(field_set:xla.cpu.KernelThunkProto.ThreadDim.y)
}

// int64 z = 3;
inline void KernelThunkProto_ThreadDim::clear_z() {
  _impl_.z_ = int64_t{0};
}
inline int64_t KernelThunkProto_ThreadDim::_internal_z() const {
  return _impl_.z_;
}
inline int64_t KernelThunkProto_ThreadDim::z() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.ThreadDim.z)
  return _internal_z();
}
inline void KernelThunkProto_ThreadDim::_internal_set_z(int64_t value) {
  
  _impl_.z_ = value;
}
inline void KernelThunkProto_ThreadDim::set_z(int64_t value) {
  _internal_set_z(value);
  // @@protoc_insertion_point(field_set:xla.cpu.KernelThunkProto.ThreadDim.z)
}

// -------------------------------------------------------------------

// KernelThunkProto

// string kernel_name = 1;
inline void KernelThunkProto::clear_kernel_name() {
  _impl_.kernel_name_.ClearToEmpty();
}
inline const std::string& KernelThunkProto::kernel_name() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.kernel_name)
  return _internal_kernel_name();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void KernelThunkProto::set_kernel_name(ArgT0&& arg0, ArgT... args) {
 
 _impl_.kernel_name_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.KernelThunkProto.kernel_name)
}
inline std::string* KernelThunkProto::mutable_kernel_name() {
  std::string* _s = _internal_mutable_kernel_name();
  // @@protoc_insertion_point(field_mutable:xla.cpu.KernelThunkProto.kernel_name)
  return _s;
}
inline const std::string& KernelThunkProto::_internal_kernel_name() const {
  return _impl_.kernel_name_.Get();
}
inline void KernelThunkProto::_internal_set_kernel_name(const std::string& value) {
  
  _impl_.kernel_name_.Set(value, GetArenaForAllocation());
}
inline std::string* KernelThunkProto::_internal_mutable_kernel_name() {
  
  return _impl_.kernel_name_.Mutable(GetArenaForAllocation());
}
inline std::string* KernelThunkProto::release_kernel_name() {
  // @@protoc_insertion_point(field_release:xla.cpu.KernelThunkProto.kernel_name)
  return _impl_.kernel_name_.Release();
}
inline void KernelThunkProto::set_allocated_kernel_name(std::string* kernel_name) {
  if (kernel_name != nullptr) {
    
  } else {
    
  }
  _impl_.kernel_name_.SetAllocated(kernel_name, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.kernel_name_.IsDefault()) {
    _impl_.kernel_name_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.KernelThunkProto.kernel_name)
}

// .xla.cpu.KernelThunkProto.ThreadDim thread_dim = 2;
inline bool KernelThunkProto::_internal_has_thread_dim() const {
  return this != internal_default_instance() && _impl_.thread_dim_ != nullptr;
}
inline bool KernelThunkProto::has_thread_dim() const {
  return _internal_has_thread_dim();
}
inline void KernelThunkProto::clear_thread_dim() {
  if (GetArenaForAllocation() == nullptr && _impl_.thread_dim_ != nullptr) {
    delete _impl_.thread_dim_;
  }
  _impl_.thread_dim_ = nullptr;
}
inline const ::xla::cpu::KernelThunkProto_ThreadDim& KernelThunkProto::_internal_thread_dim() const {
  const ::xla::cpu::KernelThunkProto_ThreadDim* p = _impl_.thread_dim_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::KernelThunkProto_ThreadDim&>(
      ::xla::cpu::_KernelThunkProto_ThreadDim_default_instance_);
}
inline const ::xla::cpu::KernelThunkProto_ThreadDim& KernelThunkProto::thread_dim() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.thread_dim)
  return _internal_thread_dim();
}
inline void KernelThunkProto::unsafe_arena_set_allocated_thread_dim(
    ::xla::cpu::KernelThunkProto_ThreadDim* thread_dim) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.thread_dim_);
  }
  _impl_.thread_dim_ = thread_dim;
  if (thread_dim) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.KernelThunkProto.thread_dim)
}
inline ::xla::cpu::KernelThunkProto_ThreadDim* KernelThunkProto::release_thread_dim() {
  
  ::xla::cpu::KernelThunkProto_ThreadDim* temp = _impl_.thread_dim_;
  _impl_.thread_dim_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::KernelThunkProto_ThreadDim* KernelThunkProto::unsafe_arena_release_thread_dim() {
  // @@protoc_insertion_point(field_release:xla.cpu.KernelThunkProto.thread_dim)
  
  ::xla::cpu::KernelThunkProto_ThreadDim* temp = _impl_.thread_dim_;
  _impl_.thread_dim_ = nullptr;
  return temp;
}
inline ::xla::cpu::KernelThunkProto_ThreadDim* KernelThunkProto::_internal_mutable_thread_dim() {
  
  if (_impl_.thread_dim_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::KernelThunkProto_ThreadDim>(GetArenaForAllocation());
    _impl_.thread_dim_ = p;
  }
  return _impl_.thread_dim_;
}
inline ::xla::cpu::KernelThunkProto_ThreadDim* KernelThunkProto::mutable_thread_dim() {
  ::xla::cpu::KernelThunkProto_ThreadDim* _msg = _internal_mutable_thread_dim();
  // @@protoc_insertion_point(field_mutable:xla.cpu.KernelThunkProto.thread_dim)
  return _msg;
}
inline void KernelThunkProto::set_allocated_thread_dim(::xla::cpu::KernelThunkProto_ThreadDim* thread_dim) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.thread_dim_;
  }
  if (thread_dim) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(thread_dim);
    if (message_arena != submessage_arena) {
      thread_dim = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, thread_dim, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.thread_dim_ = thread_dim;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.KernelThunkProto.thread_dim)
}

// repeated int64 invariant_arguments = 3;
inline int KernelThunkProto::_internal_invariant_arguments_size() const {
  return _impl_.invariant_arguments_.size();
}
inline int KernelThunkProto::invariant_arguments_size() const {
  return _internal_invariant_arguments_size();
}
inline void KernelThunkProto::clear_invariant_arguments() {
  _impl_.invariant_arguments_.Clear();
}
inline int64_t KernelThunkProto::_internal_invariant_arguments(int index) const {
  return _impl_.invariant_arguments_.Get(index);
}
inline int64_t KernelThunkProto::invariant_arguments(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.invariant_arguments)
  return _internal_invariant_arguments(index);
}
inline void KernelThunkProto::set_invariant_arguments(int index, int64_t value) {
  _impl_.invariant_arguments_.Set(index, value);
  // @@protoc_insertion_point(field_set:xla.cpu.KernelThunkProto.invariant_arguments)
}
inline void KernelThunkProto::_internal_add_invariant_arguments(int64_t value) {
  _impl_.invariant_arguments_.Add(value);
}
inline void KernelThunkProto::add_invariant_arguments(int64_t value) {
  _internal_add_invariant_arguments(value);
  // @@protoc_insertion_point(field_add:xla.cpu.KernelThunkProto.invariant_arguments)
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
KernelThunkProto::_internal_invariant_arguments() const {
  return _impl_.invariant_arguments_;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
KernelThunkProto::invariant_arguments() const {
  // @@protoc_insertion_point(field_list:xla.cpu.KernelThunkProto.invariant_arguments)
  return _internal_invariant_arguments();
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
KernelThunkProto::_internal_mutable_invariant_arguments() {
  return &_impl_.invariant_arguments_;
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
KernelThunkProto::mutable_invariant_arguments() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.KernelThunkProto.invariant_arguments)
  return _internal_mutable_invariant_arguments();
}

// .xla.cpu.Int64Optional min_alignment = 4;
inline bool KernelThunkProto::_internal_has_min_alignment() const {
  return this != internal_default_instance() && _impl_.min_alignment_ != nullptr;
}
inline bool KernelThunkProto::has_min_alignment() const {
  return _internal_has_min_alignment();
}
inline void KernelThunkProto::clear_min_alignment() {
  if (GetArenaForAllocation() == nullptr && _impl_.min_alignment_ != nullptr) {
    delete _impl_.min_alignment_;
  }
  _impl_.min_alignment_ = nullptr;
}
inline const ::xla::cpu::Int64Optional& KernelThunkProto::_internal_min_alignment() const {
  const ::xla::cpu::Int64Optional* p = _impl_.min_alignment_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::Int64Optional&>(
      ::xla::cpu::_Int64Optional_default_instance_);
}
inline const ::xla::cpu::Int64Optional& KernelThunkProto::min_alignment() const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.min_alignment)
  return _internal_min_alignment();
}
inline void KernelThunkProto::unsafe_arena_set_allocated_min_alignment(
    ::xla::cpu::Int64Optional* min_alignment) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.min_alignment_);
  }
  _impl_.min_alignment_ = min_alignment;
  if (min_alignment) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.KernelThunkProto.min_alignment)
}
inline ::xla::cpu::Int64Optional* KernelThunkProto::release_min_alignment() {
  
  ::xla::cpu::Int64Optional* temp = _impl_.min_alignment_;
  _impl_.min_alignment_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::Int64Optional* KernelThunkProto::unsafe_arena_release_min_alignment() {
  // @@protoc_insertion_point(field_release:xla.cpu.KernelThunkProto.min_alignment)
  
  ::xla::cpu::Int64Optional* temp = _impl_.min_alignment_;
  _impl_.min_alignment_ = nullptr;
  return temp;
}
inline ::xla::cpu::Int64Optional* KernelThunkProto::_internal_mutable_min_alignment() {
  
  if (_impl_.min_alignment_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::Int64Optional>(GetArenaForAllocation());
    _impl_.min_alignment_ = p;
  }
  return _impl_.min_alignment_;
}
inline ::xla::cpu::Int64Optional* KernelThunkProto::mutable_min_alignment() {
  ::xla::cpu::Int64Optional* _msg = _internal_mutable_min_alignment();
  // @@protoc_insertion_point(field_mutable:xla.cpu.KernelThunkProto.min_alignment)
  return _msg;
}
inline void KernelThunkProto::set_allocated_min_alignment(::xla::cpu::Int64Optional* min_alignment) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.min_alignment_;
  }
  if (min_alignment) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(min_alignment);
    if (message_arena != submessage_arena) {
      min_alignment = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, min_alignment, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.min_alignment_ = min_alignment;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.KernelThunkProto.min_alignment)
}

// repeated .xla.cpu.BufferAllocationSliceProto arguments_buffers = 5;
inline int KernelThunkProto::_internal_arguments_buffers_size() const {
  return _impl_.arguments_buffers_.size();
}
inline int KernelThunkProto::arguments_buffers_size() const {
  return _internal_arguments_buffers_size();
}
inline void KernelThunkProto::clear_arguments_buffers() {
  _impl_.arguments_buffers_.Clear();
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::mutable_arguments_buffers(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.KernelThunkProto.arguments_buffers)
  return _impl_.arguments_buffers_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >*
KernelThunkProto::mutable_arguments_buffers() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.KernelThunkProto.arguments_buffers)
  return &_impl_.arguments_buffers_;
}
inline const ::xla::cpu::BufferAllocationSliceProto& KernelThunkProto::_internal_arguments_buffers(int index) const {
  return _impl_.arguments_buffers_.Get(index);
}
inline const ::xla::cpu::BufferAllocationSliceProto& KernelThunkProto::arguments_buffers(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.arguments_buffers)
  return _internal_arguments_buffers(index);
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::_internal_add_arguments_buffers() {
  return _impl_.arguments_buffers_.Add();
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::add_arguments_buffers() {
  ::xla::cpu::BufferAllocationSliceProto* _add = _internal_add_arguments_buffers();
  // @@protoc_insertion_point(field_add:xla.cpu.KernelThunkProto.arguments_buffers)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >&
KernelThunkProto::arguments_buffers() const {
  // @@protoc_insertion_point(field_list:xla.cpu.KernelThunkProto.arguments_buffers)
  return _impl_.arguments_buffers_;
}

// repeated .xla.cpu.BufferAllocationSliceProto results_buffers = 6;
inline int KernelThunkProto::_internal_results_buffers_size() const {
  return _impl_.results_buffers_.size();
}
inline int KernelThunkProto::results_buffers_size() const {
  return _internal_results_buffers_size();
}
inline void KernelThunkProto::clear_results_buffers() {
  _impl_.results_buffers_.Clear();
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::mutable_results_buffers(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.KernelThunkProto.results_buffers)
  return _impl_.results_buffers_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >*
KernelThunkProto::mutable_results_buffers() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.KernelThunkProto.results_buffers)
  return &_impl_.results_buffers_;
}
inline const ::xla::cpu::BufferAllocationSliceProto& KernelThunkProto::_internal_results_buffers(int index) const {
  return _impl_.results_buffers_.Get(index);
}
inline const ::xla::cpu::BufferAllocationSliceProto& KernelThunkProto::results_buffers(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.KernelThunkProto.results_buffers)
  return _internal_results_buffers(index);
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::_internal_add_results_buffers() {
  return _impl_.results_buffers_.Add();
}
inline ::xla::cpu::BufferAllocationSliceProto* KernelThunkProto::add_results_buffers() {
  ::xla::cpu::BufferAllocationSliceProto* _add = _internal_add_results_buffers();
  // @@protoc_insertion_point(field_add:xla.cpu.KernelThunkProto.results_buffers)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::BufferAllocationSliceProto >&
KernelThunkProto::results_buffers() const {
  // @@protoc_insertion_point(field_list:xla.cpu.KernelThunkProto.results_buffers)
  return _impl_.results_buffers_;
}

// -------------------------------------------------------------------

// CopyThunkProto

// .xla.cpu.ShapeBufferAllocationSliceProto src_buffer_shape = 1;
inline bool CopyThunkProto::_internal_has_src_buffer_shape() const {
  return this != internal_default_instance() && _impl_.src_buffer_shape_ != nullptr;
}
inline bool CopyThunkProto::has_src_buffer_shape() const {
  return _internal_has_src_buffer_shape();
}
inline void CopyThunkProto::clear_src_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.src_buffer_shape_ != nullptr) {
    delete _impl_.src_buffer_shape_;
  }
  _impl_.src_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CopyThunkProto::_internal_src_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.src_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CopyThunkProto::src_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CopyThunkProto.src_buffer_shape)
  return _internal_src_buffer_shape();
}
inline void CopyThunkProto::unsafe_arena_set_allocated_src_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* src_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.src_buffer_shape_);
  }
  _impl_.src_buffer_shape_ = src_buffer_shape;
  if (src_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CopyThunkProto.src_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::release_src_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.src_buffer_shape_;
  _impl_.src_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::unsafe_arena_release_src_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.CopyThunkProto.src_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.src_buffer_shape_;
  _impl_.src_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::_internal_mutable_src_buffer_shape() {
  
  if (_impl_.src_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.src_buffer_shape_ = p;
  }
  return _impl_.src_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::mutable_src_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_src_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CopyThunkProto.src_buffer_shape)
  return _msg;
}
inline void CopyThunkProto::set_allocated_src_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* src_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.src_buffer_shape_;
  }
  if (src_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(src_buffer_shape);
    if (message_arena != submessage_arena) {
      src_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, src_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.src_buffer_shape_ = src_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CopyThunkProto.src_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto dst_buffer_shape = 2;
inline bool CopyThunkProto::_internal_has_dst_buffer_shape() const {
  return this != internal_default_instance() && _impl_.dst_buffer_shape_ != nullptr;
}
inline bool CopyThunkProto::has_dst_buffer_shape() const {
  return _internal_has_dst_buffer_shape();
}
inline void CopyThunkProto::clear_dst_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.dst_buffer_shape_ != nullptr) {
    delete _impl_.dst_buffer_shape_;
  }
  _impl_.dst_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CopyThunkProto::_internal_dst_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.dst_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CopyThunkProto::dst_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CopyThunkProto.dst_buffer_shape)
  return _internal_dst_buffer_shape();
}
inline void CopyThunkProto::unsafe_arena_set_allocated_dst_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* dst_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dst_buffer_shape_);
  }
  _impl_.dst_buffer_shape_ = dst_buffer_shape;
  if (dst_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CopyThunkProto.dst_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::release_dst_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.dst_buffer_shape_;
  _impl_.dst_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::unsafe_arena_release_dst_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.CopyThunkProto.dst_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.dst_buffer_shape_;
  _impl_.dst_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::_internal_mutable_dst_buffer_shape() {
  
  if (_impl_.dst_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.dst_buffer_shape_ = p;
  }
  return _impl_.dst_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CopyThunkProto::mutable_dst_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_dst_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CopyThunkProto.dst_buffer_shape)
  return _msg;
}
inline void CopyThunkProto::set_allocated_dst_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* dst_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.dst_buffer_shape_;
  }
  if (dst_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(dst_buffer_shape);
    if (message_arena != submessage_arena) {
      dst_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, dst_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.dst_buffer_shape_ = dst_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CopyThunkProto.dst_buffer_shape)
}

// -------------------------------------------------------------------

// FftThunkProto

// bool is_multi_thread_eigen = 1;
inline void FftThunkProto::clear_is_multi_thread_eigen() {
  _impl_.is_multi_thread_eigen_ = false;
}
inline bool FftThunkProto::_internal_is_multi_thread_eigen() const {
  return _impl_.is_multi_thread_eigen_;
}
inline bool FftThunkProto::is_multi_thread_eigen() const {
  // @@protoc_insertion_point(field_get:xla.cpu.FftThunkProto.is_multi_thread_eigen)
  return _internal_is_multi_thread_eigen();
}
inline void FftThunkProto::_internal_set_is_multi_thread_eigen(bool value) {
  
  _impl_.is_multi_thread_eigen_ = value;
}
inline void FftThunkProto::set_is_multi_thread_eigen(bool value) {
  _internal_set_is_multi_thread_eigen(value);
  // @@protoc_insertion_point(field_set:xla.cpu.FftThunkProto.is_multi_thread_eigen)
}

// int32 fft_type = 2;
inline void FftThunkProto::clear_fft_type() {
  _impl_.fft_type_ = 0;
}
inline int32_t FftThunkProto::_internal_fft_type() const {
  return _impl_.fft_type_;
}
inline int32_t FftThunkProto::fft_type() const {
  // @@protoc_insertion_point(field_get:xla.cpu.FftThunkProto.fft_type)
  return _internal_fft_type();
}
inline void FftThunkProto::_internal_set_fft_type(int32_t value) {
  
  _impl_.fft_type_ = value;
}
inline void FftThunkProto::set_fft_type(int32_t value) {
  _internal_set_fft_type(value);
  // @@protoc_insertion_point(field_set:xla.cpu.FftThunkProto.fft_type)
}

// repeated int64 fft_length = 3;
inline int FftThunkProto::_internal_fft_length_size() const {
  return _impl_.fft_length_.size();
}
inline int FftThunkProto::fft_length_size() const {
  return _internal_fft_length_size();
}
inline void FftThunkProto::clear_fft_length() {
  _impl_.fft_length_.Clear();
}
inline int64_t FftThunkProto::_internal_fft_length(int index) const {
  return _impl_.fft_length_.Get(index);
}
inline int64_t FftThunkProto::fft_length(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.FftThunkProto.fft_length)
  return _internal_fft_length(index);
}
inline void FftThunkProto::set_fft_length(int index, int64_t value) {
  _impl_.fft_length_.Set(index, value);
  // @@protoc_insertion_point(field_set:xla.cpu.FftThunkProto.fft_length)
}
inline void FftThunkProto::_internal_add_fft_length(int64_t value) {
  _impl_.fft_length_.Add(value);
}
inline void FftThunkProto::add_fft_length(int64_t value) {
  _internal_add_fft_length(value);
  // @@protoc_insertion_point(field_add:xla.cpu.FftThunkProto.fft_length)
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
FftThunkProto::_internal_fft_length() const {
  return _impl_.fft_length_;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >&
FftThunkProto::fft_length() const {
  // @@protoc_insertion_point(field_list:xla.cpu.FftThunkProto.fft_length)
  return _internal_fft_length();
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
FftThunkProto::_internal_mutable_fft_length() {
  return &_impl_.fft_length_;
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >*
FftThunkProto::mutable_fft_length() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.FftThunkProto.fft_length)
  return _internal_mutable_fft_length();
}

// .xla.cpu.ShapeBufferAllocationSliceProto input_buffer_shape = 4;
inline bool FftThunkProto::_internal_has_input_buffer_shape() const {
  return this != internal_default_instance() && _impl_.input_buffer_shape_ != nullptr;
}
inline bool FftThunkProto::has_input_buffer_shape() const {
  return _internal_has_input_buffer_shape();
}
inline void FftThunkProto::clear_input_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.input_buffer_shape_ != nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  _impl_.input_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& FftThunkProto::_internal_input_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.input_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& FftThunkProto::input_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.FftThunkProto.input_buffer_shape)
  return _internal_input_buffer_shape();
}
inline void FftThunkProto::unsafe_arena_set_allocated_input_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.input_buffer_shape_);
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  if (input_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.FftThunkProto.input_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::release_input_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::unsafe_arena_release_input_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.FftThunkProto.input_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.input_buffer_shape_;
  _impl_.input_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::_internal_mutable_input_buffer_shape() {
  
  if (_impl_.input_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.input_buffer_shape_ = p;
  }
  return _impl_.input_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::mutable_input_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_input_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.FftThunkProto.input_buffer_shape)
  return _msg;
}
inline void FftThunkProto::set_allocated_input_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* input_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.input_buffer_shape_;
  }
  if (input_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(input_buffer_shape);
    if (message_arena != submessage_arena) {
      input_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, input_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.input_buffer_shape_ = input_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.FftThunkProto.input_buffer_shape)
}

// .xla.cpu.ShapeBufferAllocationSliceProto output_buffer_shape = 5;
inline bool FftThunkProto::_internal_has_output_buffer_shape() const {
  return this != internal_default_instance() && _impl_.output_buffer_shape_ != nullptr;
}
inline bool FftThunkProto::has_output_buffer_shape() const {
  return _internal_has_output_buffer_shape();
}
inline void FftThunkProto::clear_output_buffer_shape() {
  if (GetArenaForAllocation() == nullptr && _impl_.output_buffer_shape_ != nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  _impl_.output_buffer_shape_ = nullptr;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& FftThunkProto::_internal_output_buffer_shape() const {
  const ::xla::cpu::ShapeBufferAllocationSliceProto* p = _impl_.output_buffer_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ShapeBufferAllocationSliceProto&>(
      ::xla::cpu::_ShapeBufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& FftThunkProto::output_buffer_shape() const {
  // @@protoc_insertion_point(field_get:xla.cpu.FftThunkProto.output_buffer_shape)
  return _internal_output_buffer_shape();
}
inline void FftThunkProto::unsafe_arena_set_allocated_output_buffer_shape(
    ::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.output_buffer_shape_);
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  if (output_buffer_shape) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.FftThunkProto.output_buffer_shape)
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::release_output_buffer_shape() {
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::unsafe_arena_release_output_buffer_shape() {
  // @@protoc_insertion_point(field_release:xla.cpu.FftThunkProto.output_buffer_shape)
  
  ::xla::cpu::ShapeBufferAllocationSliceProto* temp = _impl_.output_buffer_shape_;
  _impl_.output_buffer_shape_ = nullptr;
  return temp;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::_internal_mutable_output_buffer_shape() {
  
  if (_impl_.output_buffer_shape_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ShapeBufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.output_buffer_shape_ = p;
  }
  return _impl_.output_buffer_shape_;
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* FftThunkProto::mutable_output_buffer_shape() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _msg = _internal_mutable_output_buffer_shape();
  // @@protoc_insertion_point(field_mutable:xla.cpu.FftThunkProto.output_buffer_shape)
  return _msg;
}
inline void FftThunkProto::set_allocated_output_buffer_shape(::xla::cpu::ShapeBufferAllocationSliceProto* output_buffer_shape) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.output_buffer_shape_;
  }
  if (output_buffer_shape) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(output_buffer_shape);
    if (message_arena != submessage_arena) {
      output_buffer_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, output_buffer_shape, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.output_buffer_shape_ = output_buffer_shape;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.FftThunkProto.output_buffer_shape)
}

// -------------------------------------------------------------------

// InfeedThunkProto_InfeedResource

// .xla.cpu.ResourceOptional consume_token = 1;
inline bool InfeedThunkProto_InfeedResource::_internal_has_consume_token() const {
  return this != internal_default_instance() && _impl_.consume_token_ != nullptr;
}
inline bool InfeedThunkProto_InfeedResource::has_consume_token() const {
  return _internal_has_consume_token();
}
inline void InfeedThunkProto_InfeedResource::clear_consume_token() {
  if (GetArenaForAllocation() == nullptr && _impl_.consume_token_ != nullptr) {
    delete _impl_.consume_token_;
  }
  _impl_.consume_token_ = nullptr;
}
inline const ::xla::cpu::ResourceOptional& InfeedThunkProto_InfeedResource::_internal_consume_token() const {
  const ::xla::cpu::ResourceOptional* p = _impl_.consume_token_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceOptional&>(
      ::xla::cpu::_ResourceOptional_default_instance_);
}
inline const ::xla::cpu::ResourceOptional& InfeedThunkProto_InfeedResource::consume_token() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfeedThunkProto.InfeedResource.consume_token)
  return _internal_consume_token();
}
inline void InfeedThunkProto_InfeedResource::unsafe_arena_set_allocated_consume_token(
    ::xla::cpu::ResourceOptional* consume_token) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.consume_token_);
  }
  _impl_.consume_token_ = consume_token;
  if (consume_token) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.InfeedThunkProto.InfeedResource.consume_token)
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::release_consume_token() {
  
  ::xla::cpu::ResourceOptional* temp = _impl_.consume_token_;
  _impl_.consume_token_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::unsafe_arena_release_consume_token() {
  // @@protoc_insertion_point(field_release:xla.cpu.InfeedThunkProto.InfeedResource.consume_token)
  
  ::xla::cpu::ResourceOptional* temp = _impl_.consume_token_;
  _impl_.consume_token_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::_internal_mutable_consume_token() {
  
  if (_impl_.consume_token_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceOptional>(GetArenaForAllocation());
    _impl_.consume_token_ = p;
  }
  return _impl_.consume_token_;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::mutable_consume_token() {
  ::xla::cpu::ResourceOptional* _msg = _internal_mutable_consume_token();
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfeedThunkProto.InfeedResource.consume_token)
  return _msg;
}
inline void InfeedThunkProto_InfeedResource::set_allocated_consume_token(::xla::cpu::ResourceOptional* consume_token) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.consume_token_;
  }
  if (consume_token) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(consume_token);
    if (message_arena != submessage_arena) {
      consume_token = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, consume_token, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.consume_token_ = consume_token;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.InfeedThunkProto.InfeedResource.consume_token)
}

// .xla.cpu.ResourceOptional produce_token = 2;
inline bool InfeedThunkProto_InfeedResource::_internal_has_produce_token() const {
  return this != internal_default_instance() && _impl_.produce_token_ != nullptr;
}
inline bool InfeedThunkProto_InfeedResource::has_produce_token() const {
  return _internal_has_produce_token();
}
inline void InfeedThunkProto_InfeedResource::clear_produce_token() {
  if (GetArenaForAllocation() == nullptr && _impl_.produce_token_ != nullptr) {
    delete _impl_.produce_token_;
  }
  _impl_.produce_token_ = nullptr;
}
inline const ::xla::cpu::ResourceOptional& InfeedThunkProto_InfeedResource::_internal_produce_token() const {
  const ::xla::cpu::ResourceOptional* p = _impl_.produce_token_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceOptional&>(
      ::xla::cpu::_ResourceOptional_default_instance_);
}
inline const ::xla::cpu::ResourceOptional& InfeedThunkProto_InfeedResource::produce_token() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfeedThunkProto.InfeedResource.produce_token)
  return _internal_produce_token();
}
inline void InfeedThunkProto_InfeedResource::unsafe_arena_set_allocated_produce_token(
    ::xla::cpu::ResourceOptional* produce_token) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.produce_token_);
  }
  _impl_.produce_token_ = produce_token;
  if (produce_token) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.InfeedThunkProto.InfeedResource.produce_token)
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::release_produce_token() {
  
  ::xla::cpu::ResourceOptional* temp = _impl_.produce_token_;
  _impl_.produce_token_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::unsafe_arena_release_produce_token() {
  // @@protoc_insertion_point(field_release:xla.cpu.InfeedThunkProto.InfeedResource.produce_token)
  
  ::xla::cpu::ResourceOptional* temp = _impl_.produce_token_;
  _impl_.produce_token_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::_internal_mutable_produce_token() {
  
  if (_impl_.produce_token_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceOptional>(GetArenaForAllocation());
    _impl_.produce_token_ = p;
  }
  return _impl_.produce_token_;
}
inline ::xla::cpu::ResourceOptional* InfeedThunkProto_InfeedResource::mutable_produce_token() {
  ::xla::cpu::ResourceOptional* _msg = _internal_mutable_produce_token();
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfeedThunkProto.InfeedResource.produce_token)
  return _msg;
}
inline void InfeedThunkProto_InfeedResource::set_allocated_produce_token(::xla::cpu::ResourceOptional* produce_token) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.produce_token_;
  }
  if (produce_token) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(produce_token);
    if (message_arena != submessage_arena) {
      produce_token = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, produce_token, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.produce_token_ = produce_token;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.InfeedThunkProto.InfeedResource.produce_token)
}

// -------------------------------------------------------------------

// InfeedThunkProto

// .xla.cpu.InfeedThunkProto.InfeedResource infeed_resources = 1;
inline bool InfeedThunkProto::_internal_has_infeed_resources() const {
  return this != internal_default_instance() && _impl_.infeed_resources_ != nullptr;
}
inline bool InfeedThunkProto::has_infeed_resources() const {
  return _internal_has_infeed_resources();
}
inline void InfeedThunkProto::clear_infeed_resources() {
  if (GetArenaForAllocation() == nullptr && _impl_.infeed_resources_ != nullptr) {
    delete _impl_.infeed_resources_;
  }
  _impl_.infeed_resources_ = nullptr;
}
inline const ::xla::cpu::InfeedThunkProto_InfeedResource& InfeedThunkProto::_internal_infeed_resources() const {
  const ::xla::cpu::InfeedThunkProto_InfeedResource* p = _impl_.infeed_resources_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::InfeedThunkProto_InfeedResource&>(
      ::xla::cpu::_InfeedThunkProto_InfeedResource_default_instance_);
}
inline const ::xla::cpu::InfeedThunkProto_InfeedResource& InfeedThunkProto::infeed_resources() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfeedThunkProto.infeed_resources)
  return _internal_infeed_resources();
}
inline void InfeedThunkProto::unsafe_arena_set_allocated_infeed_resources(
    ::xla::cpu::InfeedThunkProto_InfeedResource* infeed_resources) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.infeed_resources_);
  }
  _impl_.infeed_resources_ = infeed_resources;
  if (infeed_resources) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.InfeedThunkProto.infeed_resources)
}
inline ::xla::cpu::InfeedThunkProto_InfeedResource* InfeedThunkProto::release_infeed_resources() {
  
  ::xla::cpu::InfeedThunkProto_InfeedResource* temp = _impl_.infeed_resources_;
  _impl_.infeed_resources_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::InfeedThunkProto_InfeedResource* InfeedThunkProto::unsafe_arena_release_infeed_resources() {
  // @@protoc_insertion_point(field_release:xla.cpu.InfeedThunkProto.infeed_resources)
  
  ::xla::cpu::InfeedThunkProto_InfeedResource* temp = _impl_.infeed_resources_;
  _impl_.infeed_resources_ = nullptr;
  return temp;
}
inline ::xla::cpu::InfeedThunkProto_InfeedResource* InfeedThunkProto::_internal_mutable_infeed_resources() {
  
  if (_impl_.infeed_resources_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::InfeedThunkProto_InfeedResource>(GetArenaForAllocation());
    _impl_.infeed_resources_ = p;
  }
  return _impl_.infeed_resources_;
}
inline ::xla::cpu::InfeedThunkProto_InfeedResource* InfeedThunkProto::mutable_infeed_resources() {
  ::xla::cpu::InfeedThunkProto_InfeedResource* _msg = _internal_mutable_infeed_resources();
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfeedThunkProto.infeed_resources)
  return _msg;
}
inline void InfeedThunkProto::set_allocated_infeed_resources(::xla::cpu::InfeedThunkProto_InfeedResource* infeed_resources) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.infeed_resources_;
  }
  if (infeed_resources) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(infeed_resources);
    if (message_arena != submessage_arena) {
      infeed_resources = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, infeed_resources, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.infeed_resources_ = infeed_resources;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.InfeedThunkProto.infeed_resources)
}

// repeated .xla.cpu.ShapeBufferAllocationSliceProto infeed_buffers_shapes = 2;
inline int InfeedThunkProto::_internal_infeed_buffers_shapes_size() const {
  return _impl_.infeed_buffers_shapes_.size();
}
inline int InfeedThunkProto::infeed_buffers_shapes_size() const {
  return _internal_infeed_buffers_shapes_size();
}
inline void InfeedThunkProto::clear_infeed_buffers_shapes() {
  _impl_.infeed_buffers_shapes_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* InfeedThunkProto::mutable_infeed_buffers_shapes(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfeedThunkProto.infeed_buffers_shapes)
  return _impl_.infeed_buffers_shapes_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
InfeedThunkProto::mutable_infeed_buffers_shapes() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.InfeedThunkProto.infeed_buffers_shapes)
  return &_impl_.infeed_buffers_shapes_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& InfeedThunkProto::_internal_infeed_buffers_shapes(int index) const {
  return _impl_.infeed_buffers_shapes_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& InfeedThunkProto::infeed_buffers_shapes(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfeedThunkProto.infeed_buffers_shapes)
  return _internal_infeed_buffers_shapes(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* InfeedThunkProto::_internal_add_infeed_buffers_shapes() {
  return _impl_.infeed_buffers_shapes_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* InfeedThunkProto::add_infeed_buffers_shapes() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_infeed_buffers_shapes();
  // @@protoc_insertion_point(field_add:xla.cpu.InfeedThunkProto.infeed_buffers_shapes)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
InfeedThunkProto::infeed_buffers_shapes() const {
  // @@protoc_insertion_point(field_list:xla.cpu.InfeedThunkProto.infeed_buffers_shapes)
  return _impl_.infeed_buffers_shapes_;
}

// -------------------------------------------------------------------

// OutfeedThunkProto_OutfeedResource

// .xla.cpu.ResourceOptional consume_token = 1;
inline bool OutfeedThunkProto_OutfeedResource::_internal_has_consume_token() const {
  return this != internal_default_instance() && _impl_.consume_token_ != nullptr;
}
inline bool OutfeedThunkProto_OutfeedResource::has_consume_token() const {
  return _internal_has_consume_token();
}
inline void OutfeedThunkProto_OutfeedResource::clear_consume_token() {
  if (GetArenaForAllocation() == nullptr && _impl_.consume_token_ != nullptr) {
    delete _impl_.consume_token_;
  }
  _impl_.consume_token_ = nullptr;
}
inline const ::xla::cpu::ResourceOptional& OutfeedThunkProto_OutfeedResource::_internal_consume_token() const {
  const ::xla::cpu::ResourceOptional* p = _impl_.consume_token_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceOptional&>(
      ::xla::cpu::_ResourceOptional_default_instance_);
}
inline const ::xla::cpu::ResourceOptional& OutfeedThunkProto_OutfeedResource::consume_token() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OutfeedThunkProto.OutfeedResource.consume_token)
  return _internal_consume_token();
}
inline void OutfeedThunkProto_OutfeedResource::unsafe_arena_set_allocated_consume_token(
    ::xla::cpu::ResourceOptional* consume_token) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.consume_token_);
  }
  _impl_.consume_token_ = consume_token;
  if (consume_token) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.OutfeedThunkProto.OutfeedResource.consume_token)
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::release_consume_token() {
  
  ::xla::cpu::ResourceOptional* temp = _impl_.consume_token_;
  _impl_.consume_token_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::unsafe_arena_release_consume_token() {
  // @@protoc_insertion_point(field_release:xla.cpu.OutfeedThunkProto.OutfeedResource.consume_token)
  
  ::xla::cpu::ResourceOptional* temp = _impl_.consume_token_;
  _impl_.consume_token_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::_internal_mutable_consume_token() {
  
  if (_impl_.consume_token_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceOptional>(GetArenaForAllocation());
    _impl_.consume_token_ = p;
  }
  return _impl_.consume_token_;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::mutable_consume_token() {
  ::xla::cpu::ResourceOptional* _msg = _internal_mutable_consume_token();
  // @@protoc_insertion_point(field_mutable:xla.cpu.OutfeedThunkProto.OutfeedResource.consume_token)
  return _msg;
}
inline void OutfeedThunkProto_OutfeedResource::set_allocated_consume_token(::xla::cpu::ResourceOptional* consume_token) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.consume_token_;
  }
  if (consume_token) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(consume_token);
    if (message_arena != submessage_arena) {
      consume_token = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, consume_token, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.consume_token_ = consume_token;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.OutfeedThunkProto.OutfeedResource.consume_token)
}

// .xla.cpu.ResourceOptional produce_token = 2;
inline bool OutfeedThunkProto_OutfeedResource::_internal_has_produce_token() const {
  return this != internal_default_instance() && _impl_.produce_token_ != nullptr;
}
inline bool OutfeedThunkProto_OutfeedResource::has_produce_token() const {
  return _internal_has_produce_token();
}
inline void OutfeedThunkProto_OutfeedResource::clear_produce_token() {
  if (GetArenaForAllocation() == nullptr && _impl_.produce_token_ != nullptr) {
    delete _impl_.produce_token_;
  }
  _impl_.produce_token_ = nullptr;
}
inline const ::xla::cpu::ResourceOptional& OutfeedThunkProto_OutfeedResource::_internal_produce_token() const {
  const ::xla::cpu::ResourceOptional* p = _impl_.produce_token_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::ResourceOptional&>(
      ::xla::cpu::_ResourceOptional_default_instance_);
}
inline const ::xla::cpu::ResourceOptional& OutfeedThunkProto_OutfeedResource::produce_token() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OutfeedThunkProto.OutfeedResource.produce_token)
  return _internal_produce_token();
}
inline void OutfeedThunkProto_OutfeedResource::unsafe_arena_set_allocated_produce_token(
    ::xla::cpu::ResourceOptional* produce_token) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.produce_token_);
  }
  _impl_.produce_token_ = produce_token;
  if (produce_token) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.OutfeedThunkProto.OutfeedResource.produce_token)
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::release_produce_token() {
  
  ::xla::cpu::ResourceOptional* temp = _impl_.produce_token_;
  _impl_.produce_token_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::unsafe_arena_release_produce_token() {
  // @@protoc_insertion_point(field_release:xla.cpu.OutfeedThunkProto.OutfeedResource.produce_token)
  
  ::xla::cpu::ResourceOptional* temp = _impl_.produce_token_;
  _impl_.produce_token_ = nullptr;
  return temp;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::_internal_mutable_produce_token() {
  
  if (_impl_.produce_token_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::ResourceOptional>(GetArenaForAllocation());
    _impl_.produce_token_ = p;
  }
  return _impl_.produce_token_;
}
inline ::xla::cpu::ResourceOptional* OutfeedThunkProto_OutfeedResource::mutable_produce_token() {
  ::xla::cpu::ResourceOptional* _msg = _internal_mutable_produce_token();
  // @@protoc_insertion_point(field_mutable:xla.cpu.OutfeedThunkProto.OutfeedResource.produce_token)
  return _msg;
}
inline void OutfeedThunkProto_OutfeedResource::set_allocated_produce_token(::xla::cpu::ResourceOptional* produce_token) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.produce_token_;
  }
  if (produce_token) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(produce_token);
    if (message_arena != submessage_arena) {
      produce_token = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, produce_token, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.produce_token_ = produce_token;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.OutfeedThunkProto.OutfeedResource.produce_token)
}

// -------------------------------------------------------------------

// OutfeedThunkProto

// .xla.cpu.OutfeedThunkProto.OutfeedResource outfeed_resources = 1;
inline bool OutfeedThunkProto::_internal_has_outfeed_resources() const {
  return this != internal_default_instance() && _impl_.outfeed_resources_ != nullptr;
}
inline bool OutfeedThunkProto::has_outfeed_resources() const {
  return _internal_has_outfeed_resources();
}
inline void OutfeedThunkProto::clear_outfeed_resources() {
  if (GetArenaForAllocation() == nullptr && _impl_.outfeed_resources_ != nullptr) {
    delete _impl_.outfeed_resources_;
  }
  _impl_.outfeed_resources_ = nullptr;
}
inline const ::xla::cpu::OutfeedThunkProto_OutfeedResource& OutfeedThunkProto::_internal_outfeed_resources() const {
  const ::xla::cpu::OutfeedThunkProto_OutfeedResource* p = _impl_.outfeed_resources_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::OutfeedThunkProto_OutfeedResource&>(
      ::xla::cpu::_OutfeedThunkProto_OutfeedResource_default_instance_);
}
inline const ::xla::cpu::OutfeedThunkProto_OutfeedResource& OutfeedThunkProto::outfeed_resources() const {
  // @@protoc_insertion_point(field_get:xla.cpu.OutfeedThunkProto.outfeed_resources)
  return _internal_outfeed_resources();
}
inline void OutfeedThunkProto::unsafe_arena_set_allocated_outfeed_resources(
    ::xla::cpu::OutfeedThunkProto_OutfeedResource* outfeed_resources) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.outfeed_resources_);
  }
  _impl_.outfeed_resources_ = outfeed_resources;
  if (outfeed_resources) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.OutfeedThunkProto.outfeed_resources)
}
inline ::xla::cpu::OutfeedThunkProto_OutfeedResource* OutfeedThunkProto::release_outfeed_resources() {
  
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* temp = _impl_.outfeed_resources_;
  _impl_.outfeed_resources_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::OutfeedThunkProto_OutfeedResource* OutfeedThunkProto::unsafe_arena_release_outfeed_resources() {
  // @@protoc_insertion_point(field_release:xla.cpu.OutfeedThunkProto.outfeed_resources)
  
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* temp = _impl_.outfeed_resources_;
  _impl_.outfeed_resources_ = nullptr;
  return temp;
}
inline ::xla::cpu::OutfeedThunkProto_OutfeedResource* OutfeedThunkProto::_internal_mutable_outfeed_resources() {
  
  if (_impl_.outfeed_resources_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::OutfeedThunkProto_OutfeedResource>(GetArenaForAllocation());
    _impl_.outfeed_resources_ = p;
  }
  return _impl_.outfeed_resources_;
}
inline ::xla::cpu::OutfeedThunkProto_OutfeedResource* OutfeedThunkProto::mutable_outfeed_resources() {
  ::xla::cpu::OutfeedThunkProto_OutfeedResource* _msg = _internal_mutable_outfeed_resources();
  // @@protoc_insertion_point(field_mutable:xla.cpu.OutfeedThunkProto.outfeed_resources)
  return _msg;
}
inline void OutfeedThunkProto::set_allocated_outfeed_resources(::xla::cpu::OutfeedThunkProto_OutfeedResource* outfeed_resources) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.outfeed_resources_;
  }
  if (outfeed_resources) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(outfeed_resources);
    if (message_arena != submessage_arena) {
      outfeed_resources = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, outfeed_resources, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.outfeed_resources_ = outfeed_resources;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.OutfeedThunkProto.outfeed_resources)
}

// repeated .xla.cpu.ShapeBufferAllocationSliceProto outfeed_buffers_shapes = 2;
inline int OutfeedThunkProto::_internal_outfeed_buffers_shapes_size() const {
  return _impl_.outfeed_buffers_shapes_.size();
}
inline int OutfeedThunkProto::outfeed_buffers_shapes_size() const {
  return _internal_outfeed_buffers_shapes_size();
}
inline void OutfeedThunkProto::clear_outfeed_buffers_shapes() {
  _impl_.outfeed_buffers_shapes_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OutfeedThunkProto::mutable_outfeed_buffers_shapes(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.OutfeedThunkProto.outfeed_buffers_shapes)
  return _impl_.outfeed_buffers_shapes_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
OutfeedThunkProto::mutable_outfeed_buffers_shapes() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.OutfeedThunkProto.outfeed_buffers_shapes)
  return &_impl_.outfeed_buffers_shapes_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OutfeedThunkProto::_internal_outfeed_buffers_shapes(int index) const {
  return _impl_.outfeed_buffers_shapes_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& OutfeedThunkProto::outfeed_buffers_shapes(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.OutfeedThunkProto.outfeed_buffers_shapes)
  return _internal_outfeed_buffers_shapes(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OutfeedThunkProto::_internal_add_outfeed_buffers_shapes() {
  return _impl_.outfeed_buffers_shapes_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* OutfeedThunkProto::add_outfeed_buffers_shapes() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_outfeed_buffers_shapes();
  // @@protoc_insertion_point(field_add:xla.cpu.OutfeedThunkProto.outfeed_buffers_shapes)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
OutfeedThunkProto::outfeed_buffers_shapes() const {
  // @@protoc_insertion_point(field_list:xla.cpu.OutfeedThunkProto.outfeed_buffers_shapes)
  return _impl_.outfeed_buffers_shapes_;
}

// -------------------------------------------------------------------

// CustomCallThunkProto_OpBuffers

// repeated .xla.cpu.ShapeBufferAllocationSliceProto arguments_shapes = 1;
inline int CustomCallThunkProto_OpBuffers::_internal_arguments_shapes_size() const {
  return _impl_.arguments_shapes_.size();
}
inline int CustomCallThunkProto_OpBuffers::arguments_shapes_size() const {
  return _internal_arguments_shapes_size();
}
inline void CustomCallThunkProto_OpBuffers::clear_arguments_shapes() {
  _impl_.arguments_shapes_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::mutable_arguments_shapes(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.CustomCallThunkProto.OpBuffers.arguments_shapes)
  return _impl_.arguments_shapes_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
CustomCallThunkProto_OpBuffers::mutable_arguments_shapes() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.CustomCallThunkProto.OpBuffers.arguments_shapes)
  return &_impl_.arguments_shapes_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CustomCallThunkProto_OpBuffers::_internal_arguments_shapes(int index) const {
  return _impl_.arguments_shapes_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CustomCallThunkProto_OpBuffers::arguments_shapes(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.OpBuffers.arguments_shapes)
  return _internal_arguments_shapes(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::_internal_add_arguments_shapes() {
  return _impl_.arguments_shapes_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::add_arguments_shapes() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_arguments_shapes();
  // @@protoc_insertion_point(field_add:xla.cpu.CustomCallThunkProto.OpBuffers.arguments_shapes)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
CustomCallThunkProto_OpBuffers::arguments_shapes() const {
  // @@protoc_insertion_point(field_list:xla.cpu.CustomCallThunkProto.OpBuffers.arguments_shapes)
  return _impl_.arguments_shapes_;
}

// repeated .xla.cpu.ShapeBufferAllocationSliceProto results_shapes = 2;
inline int CustomCallThunkProto_OpBuffers::_internal_results_shapes_size() const {
  return _impl_.results_shapes_.size();
}
inline int CustomCallThunkProto_OpBuffers::results_shapes_size() const {
  return _internal_results_shapes_size();
}
inline void CustomCallThunkProto_OpBuffers::clear_results_shapes() {
  _impl_.results_shapes_.Clear();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::mutable_results_shapes(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.CustomCallThunkProto.OpBuffers.results_shapes)
  return _impl_.results_shapes_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >*
CustomCallThunkProto_OpBuffers::mutable_results_shapes() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.CustomCallThunkProto.OpBuffers.results_shapes)
  return &_impl_.results_shapes_;
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CustomCallThunkProto_OpBuffers::_internal_results_shapes(int index) const {
  return _impl_.results_shapes_.Get(index);
}
inline const ::xla::cpu::ShapeBufferAllocationSliceProto& CustomCallThunkProto_OpBuffers::results_shapes(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.OpBuffers.results_shapes)
  return _internal_results_shapes(index);
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::_internal_add_results_shapes() {
  return _impl_.results_shapes_.Add();
}
inline ::xla::cpu::ShapeBufferAllocationSliceProto* CustomCallThunkProto_OpBuffers::add_results_shapes() {
  ::xla::cpu::ShapeBufferAllocationSliceProto* _add = _internal_add_results_shapes();
  // @@protoc_insertion_point(field_add:xla.cpu.CustomCallThunkProto.OpBuffers.results_shapes)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ShapeBufferAllocationSliceProto >&
CustomCallThunkProto_OpBuffers::results_shapes() const {
  // @@protoc_insertion_point(field_list:xla.cpu.CustomCallThunkProto.OpBuffers.results_shapes)
  return _impl_.results_shapes_;
}

// -------------------------------------------------------------------

// CustomCallThunkProto

// .xla.CustomCallApiVersion api_version = 1;
inline void CustomCallThunkProto::clear_api_version() {
  _impl_.api_version_ = 0;
}
inline ::xla::CustomCallApiVersion CustomCallThunkProto::_internal_api_version() const {
  return static_cast< ::xla::CustomCallApiVersion >(_impl_.api_version_);
}
inline ::xla::CustomCallApiVersion CustomCallThunkProto::api_version() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.api_version)
  return _internal_api_version();
}
inline void CustomCallThunkProto::_internal_set_api_version(::xla::CustomCallApiVersion value) {
  
  _impl_.api_version_ = value;
}
inline void CustomCallThunkProto::set_api_version(::xla::CustomCallApiVersion value) {
  _internal_set_api_version(value);
  // @@protoc_insertion_point(field_set:xla.cpu.CustomCallThunkProto.api_version)
}

// string target_name = 2;
inline void CustomCallThunkProto::clear_target_name() {
  _impl_.target_name_.ClearToEmpty();
}
inline const std::string& CustomCallThunkProto::target_name() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.target_name)
  return _internal_target_name();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void CustomCallThunkProto::set_target_name(ArgT0&& arg0, ArgT... args) {
 
 _impl_.target_name_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.CustomCallThunkProto.target_name)
}
inline std::string* CustomCallThunkProto::mutable_target_name() {
  std::string* _s = _internal_mutable_target_name();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CustomCallThunkProto.target_name)
  return _s;
}
inline const std::string& CustomCallThunkProto::_internal_target_name() const {
  return _impl_.target_name_.Get();
}
inline void CustomCallThunkProto::_internal_set_target_name(const std::string& value) {
  
  _impl_.target_name_.Set(value, GetArenaForAllocation());
}
inline std::string* CustomCallThunkProto::_internal_mutable_target_name() {
  
  return _impl_.target_name_.Mutable(GetArenaForAllocation());
}
inline std::string* CustomCallThunkProto::release_target_name() {
  // @@protoc_insertion_point(field_release:xla.cpu.CustomCallThunkProto.target_name)
  return _impl_.target_name_.Release();
}
inline void CustomCallThunkProto::set_allocated_target_name(std::string* target_name) {
  if (target_name != nullptr) {
    
  } else {
    
  }
  _impl_.target_name_.SetAllocated(target_name, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.target_name_.IsDefault()) {
    _impl_.target_name_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CustomCallThunkProto.target_name)
}

// string backend_config = 3;
inline void CustomCallThunkProto::clear_backend_config() {
  _impl_.backend_config_.ClearToEmpty();
}
inline const std::string& CustomCallThunkProto::backend_config() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.backend_config)
  return _internal_backend_config();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void CustomCallThunkProto::set_backend_config(ArgT0&& arg0, ArgT... args) {
 
 _impl_.backend_config_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.CustomCallThunkProto.backend_config)
}
inline std::string* CustomCallThunkProto::mutable_backend_config() {
  std::string* _s = _internal_mutable_backend_config();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CustomCallThunkProto.backend_config)
  return _s;
}
inline const std::string& CustomCallThunkProto::_internal_backend_config() const {
  return _impl_.backend_config_.Get();
}
inline void CustomCallThunkProto::_internal_set_backend_config(const std::string& value) {
  
  _impl_.backend_config_.Set(value, GetArenaForAllocation());
}
inline std::string* CustomCallThunkProto::_internal_mutable_backend_config() {
  
  return _impl_.backend_config_.Mutable(GetArenaForAllocation());
}
inline std::string* CustomCallThunkProto::release_backend_config() {
  // @@protoc_insertion_point(field_release:xla.cpu.CustomCallThunkProto.backend_config)
  return _impl_.backend_config_.Release();
}
inline void CustomCallThunkProto::set_allocated_backend_config(std::string* backend_config) {
  if (backend_config != nullptr) {
    
  } else {
    
  }
  _impl_.backend_config_.SetAllocated(backend_config, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.backend_config_.IsDefault()) {
    _impl_.backend_config_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CustomCallThunkProto.backend_config)
}

// .xla.cpu.CustomCallThunkProto.OpBuffers op_buffers = 4;
inline bool CustomCallThunkProto::_internal_has_op_buffers() const {
  return this != internal_default_instance() && _impl_.op_buffers_ != nullptr;
}
inline bool CustomCallThunkProto::has_op_buffers() const {
  return _internal_has_op_buffers();
}
inline void CustomCallThunkProto::clear_op_buffers() {
  if (GetArenaForAllocation() == nullptr && _impl_.op_buffers_ != nullptr) {
    delete _impl_.op_buffers_;
  }
  _impl_.op_buffers_ = nullptr;
}
inline const ::xla::cpu::CustomCallThunkProto_OpBuffers& CustomCallThunkProto::_internal_op_buffers() const {
  const ::xla::cpu::CustomCallThunkProto_OpBuffers* p = _impl_.op_buffers_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::CustomCallThunkProto_OpBuffers&>(
      ::xla::cpu::_CustomCallThunkProto_OpBuffers_default_instance_);
}
inline const ::xla::cpu::CustomCallThunkProto_OpBuffers& CustomCallThunkProto::op_buffers() const {
  // @@protoc_insertion_point(field_get:xla.cpu.CustomCallThunkProto.op_buffers)
  return _internal_op_buffers();
}
inline void CustomCallThunkProto::unsafe_arena_set_allocated_op_buffers(
    ::xla::cpu::CustomCallThunkProto_OpBuffers* op_buffers) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.op_buffers_);
  }
  _impl_.op_buffers_ = op_buffers;
  if (op_buffers) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.CustomCallThunkProto.op_buffers)
}
inline ::xla::cpu::CustomCallThunkProto_OpBuffers* CustomCallThunkProto::release_op_buffers() {
  
  ::xla::cpu::CustomCallThunkProto_OpBuffers* temp = _impl_.op_buffers_;
  _impl_.op_buffers_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::CustomCallThunkProto_OpBuffers* CustomCallThunkProto::unsafe_arena_release_op_buffers() {
  // @@protoc_insertion_point(field_release:xla.cpu.CustomCallThunkProto.op_buffers)
  
  ::xla::cpu::CustomCallThunkProto_OpBuffers* temp = _impl_.op_buffers_;
  _impl_.op_buffers_ = nullptr;
  return temp;
}
inline ::xla::cpu::CustomCallThunkProto_OpBuffers* CustomCallThunkProto::_internal_mutable_op_buffers() {
  
  if (_impl_.op_buffers_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::CustomCallThunkProto_OpBuffers>(GetArenaForAllocation());
    _impl_.op_buffers_ = p;
  }
  return _impl_.op_buffers_;
}
inline ::xla::cpu::CustomCallThunkProto_OpBuffers* CustomCallThunkProto::mutable_op_buffers() {
  ::xla::cpu::CustomCallThunkProto_OpBuffers* _msg = _internal_mutable_op_buffers();
  // @@protoc_insertion_point(field_mutable:xla.cpu.CustomCallThunkProto.op_buffers)
  return _msg;
}
inline void CustomCallThunkProto::set_allocated_op_buffers(::xla::cpu::CustomCallThunkProto_OpBuffers* op_buffers) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.op_buffers_;
  }
  if (op_buffers) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(op_buffers);
    if (message_arena != submessage_arena) {
      op_buffers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, op_buffers, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.op_buffers_ = op_buffers;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.CustomCallThunkProto.op_buffers)
}

// -------------------------------------------------------------------

// PartitionIdThunkProto

// .xla.cpu.BufferAllocationSliceProto logical_id_buffer = 1;
inline bool PartitionIdThunkProto::_internal_has_logical_id_buffer() const {
  return this != internal_default_instance() && _impl_.logical_id_buffer_ != nullptr;
}
inline bool PartitionIdThunkProto::has_logical_id_buffer() const {
  return _internal_has_logical_id_buffer();
}
inline void PartitionIdThunkProto::clear_logical_id_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.logical_id_buffer_ != nullptr) {
    delete _impl_.logical_id_buffer_;
  }
  _impl_.logical_id_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& PartitionIdThunkProto::_internal_logical_id_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.logical_id_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& PartitionIdThunkProto::logical_id_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.PartitionIdThunkProto.logical_id_buffer)
  return _internal_logical_id_buffer();
}
inline void PartitionIdThunkProto::unsafe_arena_set_allocated_logical_id_buffer(
    ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.logical_id_buffer_);
  }
  _impl_.logical_id_buffer_ = logical_id_buffer;
  if (logical_id_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.PartitionIdThunkProto.logical_id_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* PartitionIdThunkProto::release_logical_id_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.logical_id_buffer_;
  _impl_.logical_id_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* PartitionIdThunkProto::unsafe_arena_release_logical_id_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.PartitionIdThunkProto.logical_id_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.logical_id_buffer_;
  _impl_.logical_id_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* PartitionIdThunkProto::_internal_mutable_logical_id_buffer() {
  
  if (_impl_.logical_id_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.logical_id_buffer_ = p;
  }
  return _impl_.logical_id_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* PartitionIdThunkProto::mutable_logical_id_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_logical_id_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.PartitionIdThunkProto.logical_id_buffer)
  return _msg;
}
inline void PartitionIdThunkProto::set_allocated_logical_id_buffer(::xla::cpu::BufferAllocationSliceProto* logical_id_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.logical_id_buffer_;
  }
  if (logical_id_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(logical_id_buffer);
    if (message_arena != submessage_arena) {
      logical_id_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, logical_id_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.logical_id_buffer_ = logical_id_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.PartitionIdThunkProto.logical_id_buffer)
}

// -------------------------------------------------------------------

// ReplicaIdThunkProto

// .xla.cpu.BufferAllocationSliceProto logical_id_buffer = 1;
inline bool ReplicaIdThunkProto::_internal_has_logical_id_buffer() const {
  return this != internal_default_instance() && _impl_.logical_id_buffer_ != nullptr;
}
inline bool ReplicaIdThunkProto::has_logical_id_buffer() const {
  return _internal_has_logical_id_buffer();
}
inline void ReplicaIdThunkProto::clear_logical_id_buffer() {
  if (GetArenaForAllocation() == nullptr && _impl_.logical_id_buffer_ != nullptr) {
    delete _impl_.logical_id_buffer_;
  }
  _impl_.logical_id_buffer_ = nullptr;
}
inline const ::xla::cpu::BufferAllocationSliceProto& ReplicaIdThunkProto::_internal_logical_id_buffer() const {
  const ::xla::cpu::BufferAllocationSliceProto* p = _impl_.logical_id_buffer_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::BufferAllocationSliceProto&>(
      ::xla::cpu::_BufferAllocationSliceProto_default_instance_);
}
inline const ::xla::cpu::BufferAllocationSliceProto& ReplicaIdThunkProto::logical_id_buffer() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ReplicaIdThunkProto.logical_id_buffer)
  return _internal_logical_id_buffer();
}
inline void ReplicaIdThunkProto::unsafe_arena_set_allocated_logical_id_buffer(
    ::xla::cpu::BufferAllocationSliceProto* logical_id_buffer) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.logical_id_buffer_);
  }
  _impl_.logical_id_buffer_ = logical_id_buffer;
  if (logical_id_buffer) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ReplicaIdThunkProto.logical_id_buffer)
}
inline ::xla::cpu::BufferAllocationSliceProto* ReplicaIdThunkProto::release_logical_id_buffer() {
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.logical_id_buffer_;
  _impl_.logical_id_buffer_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ReplicaIdThunkProto::unsafe_arena_release_logical_id_buffer() {
  // @@protoc_insertion_point(field_release:xla.cpu.ReplicaIdThunkProto.logical_id_buffer)
  
  ::xla::cpu::BufferAllocationSliceProto* temp = _impl_.logical_id_buffer_;
  _impl_.logical_id_buffer_ = nullptr;
  return temp;
}
inline ::xla::cpu::BufferAllocationSliceProto* ReplicaIdThunkProto::_internal_mutable_logical_id_buffer() {
  
  if (_impl_.logical_id_buffer_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::BufferAllocationSliceProto>(GetArenaForAllocation());
    _impl_.logical_id_buffer_ = p;
  }
  return _impl_.logical_id_buffer_;
}
inline ::xla::cpu::BufferAllocationSliceProto* ReplicaIdThunkProto::mutable_logical_id_buffer() {
  ::xla::cpu::BufferAllocationSliceProto* _msg = _internal_mutable_logical_id_buffer();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ReplicaIdThunkProto.logical_id_buffer)
  return _msg;
}
inline void ReplicaIdThunkProto::set_allocated_logical_id_buffer(::xla::cpu::BufferAllocationSliceProto* logical_id_buffer) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.logical_id_buffer_;
  }
  if (logical_id_buffer) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(logical_id_buffer);
    if (message_arena != submessage_arena) {
      logical_id_buffer = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, logical_id_buffer, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.logical_id_buffer_ = logical_id_buffer;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ReplicaIdThunkProto.logical_id_buffer)
}

// -------------------------------------------------------------------

// InfoProto

// string op_name = 1;
inline void InfoProto::clear_op_name() {
  _impl_.op_name_.ClearToEmpty();
}
inline const std::string& InfoProto::op_name() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfoProto.op_name)
  return _internal_op_name();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void InfoProto::set_op_name(ArgT0&& arg0, ArgT... args) {
 
 _impl_.op_name_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.InfoProto.op_name)
}
inline std::string* InfoProto::mutable_op_name() {
  std::string* _s = _internal_mutable_op_name();
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfoProto.op_name)
  return _s;
}
inline const std::string& InfoProto::_internal_op_name() const {
  return _impl_.op_name_.Get();
}
inline void InfoProto::_internal_set_op_name(const std::string& value) {
  
  _impl_.op_name_.Set(value, GetArenaForAllocation());
}
inline std::string* InfoProto::_internal_mutable_op_name() {
  
  return _impl_.op_name_.Mutable(GetArenaForAllocation());
}
inline std::string* InfoProto::release_op_name() {
  // @@protoc_insertion_point(field_release:xla.cpu.InfoProto.op_name)
  return _impl_.op_name_.Release();
}
inline void InfoProto::set_allocated_op_name(std::string* op_name) {
  if (op_name != nullptr) {
    
  } else {
    
  }
  _impl_.op_name_.SetAllocated(op_name, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.op_name_.IsDefault()) {
    _impl_.op_name_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.InfoProto.op_name)
}

// string module_name = 2;
inline void InfoProto::clear_module_name() {
  _impl_.module_name_.ClearToEmpty();
}
inline const std::string& InfoProto::module_name() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfoProto.module_name)
  return _internal_module_name();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void InfoProto::set_module_name(ArgT0&& arg0, ArgT... args) {
 
 _impl_.module_name_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.InfoProto.module_name)
}
inline std::string* InfoProto::mutable_module_name() {
  std::string* _s = _internal_mutable_module_name();
  // @@protoc_insertion_point(field_mutable:xla.cpu.InfoProto.module_name)
  return _s;
}
inline const std::string& InfoProto::_internal_module_name() const {
  return _impl_.module_name_.Get();
}
inline void InfoProto::_internal_set_module_name(const std::string& value) {
  
  _impl_.module_name_.Set(value, GetArenaForAllocation());
}
inline std::string* InfoProto::_internal_mutable_module_name() {
  
  return _impl_.module_name_.Mutable(GetArenaForAllocation());
}
inline std::string* InfoProto::release_module_name() {
  // @@protoc_insertion_point(field_release:xla.cpu.InfoProto.module_name)
  return _impl_.module_name_.Release();
}
inline void InfoProto::set_allocated_module_name(std::string* module_name) {
  if (module_name != nullptr) {
    
  } else {
    
  }
  _impl_.module_name_.SetAllocated(module_name, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.module_name_.IsDefault()) {
    _impl_.module_name_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.InfoProto.module_name)
}

// int64 module_id = 3;
inline void InfoProto::clear_module_id() {
  _impl_.module_id_ = int64_t{0};
}
inline int64_t InfoProto::_internal_module_id() const {
  return _impl_.module_id_;
}
inline int64_t InfoProto::module_id() const {
  // @@protoc_insertion_point(field_get:xla.cpu.InfoProto.module_id)
  return _internal_module_id();
}
inline void InfoProto::_internal_set_module_id(int64_t value) {
  
  _impl_.module_id_ = value;
}
inline void InfoProto::set_module_id(int64_t value) {
  _internal_set_module_id(value);
  // @@protoc_insertion_point(field_set:xla.cpu.InfoProto.module_id)
}

// -------------------------------------------------------------------

// ThunkProto

// string kind = 1;
inline void ThunkProto::clear_kind() {
  _impl_.kind_.ClearToEmpty();
}
inline const std::string& ThunkProto::kind() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.kind)
  return _internal_kind();
}
template <typename ArgT0, typename... ArgT>
inline PROTOBUF_ALWAYS_INLINE
void ThunkProto::set_kind(ArgT0&& arg0, ArgT... args) {
 
 _impl_.kind_.Set(static_cast<ArgT0 &&>(arg0), args..., GetArenaForAllocation());
  // @@protoc_insertion_point(field_set:xla.cpu.ThunkProto.kind)
}
inline std::string* ThunkProto::mutable_kind() {
  std::string* _s = _internal_mutable_kind();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.kind)
  return _s;
}
inline const std::string& ThunkProto::_internal_kind() const {
  return _impl_.kind_.Get();
}
inline void ThunkProto::_internal_set_kind(const std::string& value) {
  
  _impl_.kind_.Set(value, GetArenaForAllocation());
}
inline std::string* ThunkProto::_internal_mutable_kind() {
  
  return _impl_.kind_.Mutable(GetArenaForAllocation());
}
inline std::string* ThunkProto::release_kind() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.kind)
  return _impl_.kind_.Release();
}
inline void ThunkProto::set_allocated_kind(std::string* kind) {
  if (kind != nullptr) {
    
  } else {
    
  }
  _impl_.kind_.SetAllocated(kind, GetArenaForAllocation());
#ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
  if (_impl_.kind_.IsDefault()) {
    _impl_.kind_.Set("", GetArenaForAllocation());
  }
#endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ThunkProto.kind)
}

// .xla.cpu.InfoProto info = 2;
inline bool ThunkProto::_internal_has_info() const {
  return this != internal_default_instance() && _impl_.info_ != nullptr;
}
inline bool ThunkProto::has_info() const {
  return _internal_has_info();
}
inline void ThunkProto::clear_info() {
  if (GetArenaForAllocation() == nullptr && _impl_.info_ != nullptr) {
    delete _impl_.info_;
  }
  _impl_.info_ = nullptr;
}
inline const ::xla::cpu::InfoProto& ThunkProto::_internal_info() const {
  const ::xla::cpu::InfoProto* p = _impl_.info_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::cpu::InfoProto&>(
      ::xla::cpu::_InfoProto_default_instance_);
}
inline const ::xla::cpu::InfoProto& ThunkProto::info() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.info)
  return _internal_info();
}
inline void ThunkProto::unsafe_arena_set_allocated_info(
    ::xla::cpu::InfoProto* info) {
  if (GetArenaForAllocation() == nullptr) {
    delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.info_);
  }
  _impl_.info_ = info;
  if (info) {
    
  } else {
    
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.info)
}
inline ::xla::cpu::InfoProto* ThunkProto::release_info() {
  
  ::xla::cpu::InfoProto* temp = _impl_.info_;
  _impl_.info_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old =  reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp);
  temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  if (GetArenaForAllocation() == nullptr) { delete old; }
#else  // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArenaForAllocation() != nullptr) {
    temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return temp;
}
inline ::xla::cpu::InfoProto* ThunkProto::unsafe_arena_release_info() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.info)
  
  ::xla::cpu::InfoProto* temp = _impl_.info_;
  _impl_.info_ = nullptr;
  return temp;
}
inline ::xla::cpu::InfoProto* ThunkProto::_internal_mutable_info() {
  
  if (_impl_.info_ == nullptr) {
    auto* p = CreateMaybeMessage<::xla::cpu::InfoProto>(GetArenaForAllocation());
    _impl_.info_ = p;
  }
  return _impl_.info_;
}
inline ::xla::cpu::InfoProto* ThunkProto::mutable_info() {
  ::xla::cpu::InfoProto* _msg = _internal_mutable_info();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.info)
  return _msg;
}
inline void ThunkProto::set_allocated_info(::xla::cpu::InfoProto* info) {
  ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation();
  if (message_arena == nullptr) {
    delete _impl_.info_;
  }
  if (info) {
    ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena =
        ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(info);
    if (message_arena != submessage_arena) {
      info = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage(
          message_arena, info, submessage_arena);
    }
    
  } else {
    
  }
  _impl_.info_ = info;
  // @@protoc_insertion_point(field_set_allocated:xla.cpu.ThunkProto.info)
}

// .xla.cpu.CallThunkProto call_thunk = 3;
inline bool ThunkProto::_internal_has_call_thunk() const {
  return impl_case() == kCallThunk;
}
inline bool ThunkProto::has_call_thunk() const {
  return _internal_has_call_thunk();
}
inline void ThunkProto::set_has_call_thunk() {
  _impl_._oneof_case_[0] = kCallThunk;
}
inline void ThunkProto::clear_call_thunk() {
  if (_internal_has_call_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.call_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::CallThunkProto* ThunkProto::release_call_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.call_thunk)
  if (_internal_has_call_thunk()) {
    clear_has_impl();
    ::xla::cpu::CallThunkProto* temp = _impl_.impl_.call_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.call_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::CallThunkProto& ThunkProto::_internal_call_thunk() const {
  return _internal_has_call_thunk()
      ? *_impl_.impl_.call_thunk_
      : reinterpret_cast< ::xla::cpu::CallThunkProto&>(::xla::cpu::_CallThunkProto_default_instance_);
}
inline const ::xla::cpu::CallThunkProto& ThunkProto::call_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.call_thunk)
  return _internal_call_thunk();
}
inline ::xla::cpu::CallThunkProto* ThunkProto::unsafe_arena_release_call_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.call_thunk)
  if (_internal_has_call_thunk()) {
    clear_has_impl();
    ::xla::cpu::CallThunkProto* temp = _impl_.impl_.call_thunk_;
    _impl_.impl_.call_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_call_thunk(::xla::cpu::CallThunkProto* call_thunk) {
  clear_impl();
  if (call_thunk) {
    set_has_call_thunk();
    _impl_.impl_.call_thunk_ = call_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.call_thunk)
}
inline ::xla::cpu::CallThunkProto* ThunkProto::_internal_mutable_call_thunk() {
  if (!_internal_has_call_thunk()) {
    clear_impl();
    set_has_call_thunk();
    _impl_.impl_.call_thunk_ = CreateMaybeMessage< ::xla::cpu::CallThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.call_thunk_;
}
inline ::xla::cpu::CallThunkProto* ThunkProto::mutable_call_thunk() {
  ::xla::cpu::CallThunkProto* _msg = _internal_mutable_call_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.call_thunk)
  return _msg;
}

// .xla.cpu.ConditionalThunkProto conditional_thunk = 4;
inline bool ThunkProto::_internal_has_conditional_thunk() const {
  return impl_case() == kConditionalThunk;
}
inline bool ThunkProto::has_conditional_thunk() const {
  return _internal_has_conditional_thunk();
}
inline void ThunkProto::set_has_conditional_thunk() {
  _impl_._oneof_case_[0] = kConditionalThunk;
}
inline void ThunkProto::clear_conditional_thunk() {
  if (_internal_has_conditional_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.conditional_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::ConditionalThunkProto* ThunkProto::release_conditional_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.conditional_thunk)
  if (_internal_has_conditional_thunk()) {
    clear_has_impl();
    ::xla::cpu::ConditionalThunkProto* temp = _impl_.impl_.conditional_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.conditional_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::ConditionalThunkProto& ThunkProto::_internal_conditional_thunk() const {
  return _internal_has_conditional_thunk()
      ? *_impl_.impl_.conditional_thunk_
      : reinterpret_cast< ::xla::cpu::ConditionalThunkProto&>(::xla::cpu::_ConditionalThunkProto_default_instance_);
}
inline const ::xla::cpu::ConditionalThunkProto& ThunkProto::conditional_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.conditional_thunk)
  return _internal_conditional_thunk();
}
inline ::xla::cpu::ConditionalThunkProto* ThunkProto::unsafe_arena_release_conditional_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.conditional_thunk)
  if (_internal_has_conditional_thunk()) {
    clear_has_impl();
    ::xla::cpu::ConditionalThunkProto* temp = _impl_.impl_.conditional_thunk_;
    _impl_.impl_.conditional_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_conditional_thunk(::xla::cpu::ConditionalThunkProto* conditional_thunk) {
  clear_impl();
  if (conditional_thunk) {
    set_has_conditional_thunk();
    _impl_.impl_.conditional_thunk_ = conditional_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.conditional_thunk)
}
inline ::xla::cpu::ConditionalThunkProto* ThunkProto::_internal_mutable_conditional_thunk() {
  if (!_internal_has_conditional_thunk()) {
    clear_impl();
    set_has_conditional_thunk();
    _impl_.impl_.conditional_thunk_ = CreateMaybeMessage< ::xla::cpu::ConditionalThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.conditional_thunk_;
}
inline ::xla::cpu::ConditionalThunkProto* ThunkProto::mutable_conditional_thunk() {
  ::xla::cpu::ConditionalThunkProto* _msg = _internal_mutable_conditional_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.conditional_thunk)
  return _msg;
}

// .xla.cpu.SortThunkProto sort_thunk = 5;
inline bool ThunkProto::_internal_has_sort_thunk() const {
  return impl_case() == kSortThunk;
}
inline bool ThunkProto::has_sort_thunk() const {
  return _internal_has_sort_thunk();
}
inline void ThunkProto::set_has_sort_thunk() {
  _impl_._oneof_case_[0] = kSortThunk;
}
inline void ThunkProto::clear_sort_thunk() {
  if (_internal_has_sort_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.sort_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::SortThunkProto* ThunkProto::release_sort_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.sort_thunk)
  if (_internal_has_sort_thunk()) {
    clear_has_impl();
    ::xla::cpu::SortThunkProto* temp = _impl_.impl_.sort_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.sort_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::SortThunkProto& ThunkProto::_internal_sort_thunk() const {
  return _internal_has_sort_thunk()
      ? *_impl_.impl_.sort_thunk_
      : reinterpret_cast< ::xla::cpu::SortThunkProto&>(::xla::cpu::_SortThunkProto_default_instance_);
}
inline const ::xla::cpu::SortThunkProto& ThunkProto::sort_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.sort_thunk)
  return _internal_sort_thunk();
}
inline ::xla::cpu::SortThunkProto* ThunkProto::unsafe_arena_release_sort_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.sort_thunk)
  if (_internal_has_sort_thunk()) {
    clear_has_impl();
    ::xla::cpu::SortThunkProto* temp = _impl_.impl_.sort_thunk_;
    _impl_.impl_.sort_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_sort_thunk(::xla::cpu::SortThunkProto* sort_thunk) {
  clear_impl();
  if (sort_thunk) {
    set_has_sort_thunk();
    _impl_.impl_.sort_thunk_ = sort_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.sort_thunk)
}
inline ::xla::cpu::SortThunkProto* ThunkProto::_internal_mutable_sort_thunk() {
  if (!_internal_has_sort_thunk()) {
    clear_impl();
    set_has_sort_thunk();
    _impl_.impl_.sort_thunk_ = CreateMaybeMessage< ::xla::cpu::SortThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.sort_thunk_;
}
inline ::xla::cpu::SortThunkProto* ThunkProto::mutable_sort_thunk() {
  ::xla::cpu::SortThunkProto* _msg = _internal_mutable_sort_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.sort_thunk)
  return _msg;
}

// .xla.cpu.XnnFusionThunkProto xnn_fusion_thunk = 6;
inline bool ThunkProto::_internal_has_xnn_fusion_thunk() const {
  return impl_case() == kXnnFusionThunk;
}
inline bool ThunkProto::has_xnn_fusion_thunk() const {
  return _internal_has_xnn_fusion_thunk();
}
inline void ThunkProto::set_has_xnn_fusion_thunk() {
  _impl_._oneof_case_[0] = kXnnFusionThunk;
}
inline void ThunkProto::clear_xnn_fusion_thunk() {
  if (_internal_has_xnn_fusion_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.xnn_fusion_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::XnnFusionThunkProto* ThunkProto::release_xnn_fusion_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.xnn_fusion_thunk)
  if (_internal_has_xnn_fusion_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnFusionThunkProto* temp = _impl_.impl_.xnn_fusion_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.xnn_fusion_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::XnnFusionThunkProto& ThunkProto::_internal_xnn_fusion_thunk() const {
  return _internal_has_xnn_fusion_thunk()
      ? *_impl_.impl_.xnn_fusion_thunk_
      : reinterpret_cast< ::xla::cpu::XnnFusionThunkProto&>(::xla::cpu::_XnnFusionThunkProto_default_instance_);
}
inline const ::xla::cpu::XnnFusionThunkProto& ThunkProto::xnn_fusion_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.xnn_fusion_thunk)
  return _internal_xnn_fusion_thunk();
}
inline ::xla::cpu::XnnFusionThunkProto* ThunkProto::unsafe_arena_release_xnn_fusion_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.xnn_fusion_thunk)
  if (_internal_has_xnn_fusion_thunk()) {
    clear_has_impl();
    ::xla::cpu::XnnFusionThunkProto* temp = _impl_.impl_.xnn_fusion_thunk_;
    _impl_.impl_.xnn_fusion_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_xnn_fusion_thunk(::xla::cpu::XnnFusionThunkProto* xnn_fusion_thunk) {
  clear_impl();
  if (xnn_fusion_thunk) {
    set_has_xnn_fusion_thunk();
    _impl_.impl_.xnn_fusion_thunk_ = xnn_fusion_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.xnn_fusion_thunk)
}
inline ::xla::cpu::XnnFusionThunkProto* ThunkProto::_internal_mutable_xnn_fusion_thunk() {
  if (!_internal_has_xnn_fusion_thunk()) {
    clear_impl();
    set_has_xnn_fusion_thunk();
    _impl_.impl_.xnn_fusion_thunk_ = CreateMaybeMessage< ::xla::cpu::XnnFusionThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.xnn_fusion_thunk_;
}
inline ::xla::cpu::XnnFusionThunkProto* ThunkProto::mutable_xnn_fusion_thunk() {
  ::xla::cpu::XnnFusionThunkProto* _msg = _internal_mutable_xnn_fusion_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.xnn_fusion_thunk)
  return _msg;
}

// .xla.cpu.DotThunkProto dot_thunk = 7;
inline bool ThunkProto::_internal_has_dot_thunk() const {
  return impl_case() == kDotThunk;
}
inline bool ThunkProto::has_dot_thunk() const {
  return _internal_has_dot_thunk();
}
inline void ThunkProto::set_has_dot_thunk() {
  _impl_._oneof_case_[0] = kDotThunk;
}
inline void ThunkProto::clear_dot_thunk() {
  if (_internal_has_dot_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.dot_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::DotThunkProto* ThunkProto::release_dot_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.dot_thunk)
  if (_internal_has_dot_thunk()) {
    clear_has_impl();
    ::xla::cpu::DotThunkProto* temp = _impl_.impl_.dot_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.dot_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::DotThunkProto& ThunkProto::_internal_dot_thunk() const {
  return _internal_has_dot_thunk()
      ? *_impl_.impl_.dot_thunk_
      : reinterpret_cast< ::xla::cpu::DotThunkProto&>(::xla::cpu::_DotThunkProto_default_instance_);
}
inline const ::xla::cpu::DotThunkProto& ThunkProto::dot_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.dot_thunk)
  return _internal_dot_thunk();
}
inline ::xla::cpu::DotThunkProto* ThunkProto::unsafe_arena_release_dot_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.dot_thunk)
  if (_internal_has_dot_thunk()) {
    clear_has_impl();
    ::xla::cpu::DotThunkProto* temp = _impl_.impl_.dot_thunk_;
    _impl_.impl_.dot_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_dot_thunk(::xla::cpu::DotThunkProto* dot_thunk) {
  clear_impl();
  if (dot_thunk) {
    set_has_dot_thunk();
    _impl_.impl_.dot_thunk_ = dot_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.dot_thunk)
}
inline ::xla::cpu::DotThunkProto* ThunkProto::_internal_mutable_dot_thunk() {
  if (!_internal_has_dot_thunk()) {
    clear_impl();
    set_has_dot_thunk();
    _impl_.impl_.dot_thunk_ = CreateMaybeMessage< ::xla::cpu::DotThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.dot_thunk_;
}
inline ::xla::cpu::DotThunkProto* ThunkProto::mutable_dot_thunk() {
  ::xla::cpu::DotThunkProto* _msg = _internal_mutable_dot_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.dot_thunk)
  return _msg;
}

// .xla.cpu.RngGetAndUpdateStateThunkProto rng_get_and_update_state_thunk = 8;
inline bool ThunkProto::_internal_has_rng_get_and_update_state_thunk() const {
  return impl_case() == kRngGetAndUpdateStateThunk;
}
inline bool ThunkProto::has_rng_get_and_update_state_thunk() const {
  return _internal_has_rng_get_and_update_state_thunk();
}
inline void ThunkProto::set_has_rng_get_and_update_state_thunk() {
  _impl_._oneof_case_[0] = kRngGetAndUpdateStateThunk;
}
inline void ThunkProto::clear_rng_get_and_update_state_thunk() {
  if (_internal_has_rng_get_and_update_state_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.rng_get_and_update_state_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::RngGetAndUpdateStateThunkProto* ThunkProto::release_rng_get_and_update_state_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.rng_get_and_update_state_thunk)
  if (_internal_has_rng_get_and_update_state_thunk()) {
    clear_has_impl();
    ::xla::cpu::RngGetAndUpdateStateThunkProto* temp = _impl_.impl_.rng_get_and_update_state_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.rng_get_and_update_state_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::RngGetAndUpdateStateThunkProto& ThunkProto::_internal_rng_get_and_update_state_thunk() const {
  return _internal_has_rng_get_and_update_state_thunk()
      ? *_impl_.impl_.rng_get_and_update_state_thunk_
      : reinterpret_cast< ::xla::cpu::RngGetAndUpdateStateThunkProto&>(::xla::cpu::_RngGetAndUpdateStateThunkProto_default_instance_);
}
inline const ::xla::cpu::RngGetAndUpdateStateThunkProto& ThunkProto::rng_get_and_update_state_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.rng_get_and_update_state_thunk)
  return _internal_rng_get_and_update_state_thunk();
}
inline ::xla::cpu::RngGetAndUpdateStateThunkProto* ThunkProto::unsafe_arena_release_rng_get_and_update_state_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.rng_get_and_update_state_thunk)
  if (_internal_has_rng_get_and_update_state_thunk()) {
    clear_has_impl();
    ::xla::cpu::RngGetAndUpdateStateThunkProto* temp = _impl_.impl_.rng_get_and_update_state_thunk_;
    _impl_.impl_.rng_get_and_update_state_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_rng_get_and_update_state_thunk(::xla::cpu::RngGetAndUpdateStateThunkProto* rng_get_and_update_state_thunk) {
  clear_impl();
  if (rng_get_and_update_state_thunk) {
    set_has_rng_get_and_update_state_thunk();
    _impl_.impl_.rng_get_and_update_state_thunk_ = rng_get_and_update_state_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.rng_get_and_update_state_thunk)
}
inline ::xla::cpu::RngGetAndUpdateStateThunkProto* ThunkProto::_internal_mutable_rng_get_and_update_state_thunk() {
  if (!_internal_has_rng_get_and_update_state_thunk()) {
    clear_impl();
    set_has_rng_get_and_update_state_thunk();
    _impl_.impl_.rng_get_and_update_state_thunk_ = CreateMaybeMessage< ::xla::cpu::RngGetAndUpdateStateThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.rng_get_and_update_state_thunk_;
}
inline ::xla::cpu::RngGetAndUpdateStateThunkProto* ThunkProto::mutable_rng_get_and_update_state_thunk() {
  ::xla::cpu::RngGetAndUpdateStateThunkProto* _msg = _internal_mutable_rng_get_and_update_state_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.rng_get_and_update_state_thunk)
  return _msg;
}

// .xla.cpu.TopKThunkProto top_k_thunk = 9;
inline bool ThunkProto::_internal_has_top_k_thunk() const {
  return impl_case() == kTopKThunk;
}
inline bool ThunkProto::has_top_k_thunk() const {
  return _internal_has_top_k_thunk();
}
inline void ThunkProto::set_has_top_k_thunk() {
  _impl_._oneof_case_[0] = kTopKThunk;
}
inline void ThunkProto::clear_top_k_thunk() {
  if (_internal_has_top_k_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.top_k_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::TopKThunkProto* ThunkProto::release_top_k_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.top_k_thunk)
  if (_internal_has_top_k_thunk()) {
    clear_has_impl();
    ::xla::cpu::TopKThunkProto* temp = _impl_.impl_.top_k_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.top_k_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::TopKThunkProto& ThunkProto::_internal_top_k_thunk() const {
  return _internal_has_top_k_thunk()
      ? *_impl_.impl_.top_k_thunk_
      : reinterpret_cast< ::xla::cpu::TopKThunkProto&>(::xla::cpu::_TopKThunkProto_default_instance_);
}
inline const ::xla::cpu::TopKThunkProto& ThunkProto::top_k_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.top_k_thunk)
  return _internal_top_k_thunk();
}
inline ::xla::cpu::TopKThunkProto* ThunkProto::unsafe_arena_release_top_k_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.top_k_thunk)
  if (_internal_has_top_k_thunk()) {
    clear_has_impl();
    ::xla::cpu::TopKThunkProto* temp = _impl_.impl_.top_k_thunk_;
    _impl_.impl_.top_k_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_top_k_thunk(::xla::cpu::TopKThunkProto* top_k_thunk) {
  clear_impl();
  if (top_k_thunk) {
    set_has_top_k_thunk();
    _impl_.impl_.top_k_thunk_ = top_k_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.top_k_thunk)
}
inline ::xla::cpu::TopKThunkProto* ThunkProto::_internal_mutable_top_k_thunk() {
  if (!_internal_has_top_k_thunk()) {
    clear_impl();
    set_has_top_k_thunk();
    _impl_.impl_.top_k_thunk_ = CreateMaybeMessage< ::xla::cpu::TopKThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.top_k_thunk_;
}
inline ::xla::cpu::TopKThunkProto* ThunkProto::mutable_top_k_thunk() {
  ::xla::cpu::TopKThunkProto* _msg = _internal_mutable_top_k_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.top_k_thunk)
  return _msg;
}

// .xla.cpu.WhileThunkProto while_thunk = 10;
inline bool ThunkProto::_internal_has_while_thunk() const {
  return impl_case() == kWhileThunk;
}
inline bool ThunkProto::has_while_thunk() const {
  return _internal_has_while_thunk();
}
inline void ThunkProto::set_has_while_thunk() {
  _impl_._oneof_case_[0] = kWhileThunk;
}
inline void ThunkProto::clear_while_thunk() {
  if (_internal_has_while_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.while_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::WhileThunkProto* ThunkProto::release_while_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.while_thunk)
  if (_internal_has_while_thunk()) {
    clear_has_impl();
    ::xla::cpu::WhileThunkProto* temp = _impl_.impl_.while_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.while_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::WhileThunkProto& ThunkProto::_internal_while_thunk() const {
  return _internal_has_while_thunk()
      ? *_impl_.impl_.while_thunk_
      : reinterpret_cast< ::xla::cpu::WhileThunkProto&>(::xla::cpu::_WhileThunkProto_default_instance_);
}
inline const ::xla::cpu::WhileThunkProto& ThunkProto::while_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.while_thunk)
  return _internal_while_thunk();
}
inline ::xla::cpu::WhileThunkProto* ThunkProto::unsafe_arena_release_while_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.while_thunk)
  if (_internal_has_while_thunk()) {
    clear_has_impl();
    ::xla::cpu::WhileThunkProto* temp = _impl_.impl_.while_thunk_;
    _impl_.impl_.while_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_while_thunk(::xla::cpu::WhileThunkProto* while_thunk) {
  clear_impl();
  if (while_thunk) {
    set_has_while_thunk();
    _impl_.impl_.while_thunk_ = while_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.while_thunk)
}
inline ::xla::cpu::WhileThunkProto* ThunkProto::_internal_mutable_while_thunk() {
  if (!_internal_has_while_thunk()) {
    clear_impl();
    set_has_while_thunk();
    _impl_.impl_.while_thunk_ = CreateMaybeMessage< ::xla::cpu::WhileThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.while_thunk_;
}
inline ::xla::cpu::WhileThunkProto* ThunkProto::mutable_while_thunk() {
  ::xla::cpu::WhileThunkProto* _msg = _internal_mutable_while_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.while_thunk)
  return _msg;
}

// .xla.cpu.KernelThunkProto kernel_thunk = 11;
inline bool ThunkProto::_internal_has_kernel_thunk() const {
  return impl_case() == kKernelThunk;
}
inline bool ThunkProto::has_kernel_thunk() const {
  return _internal_has_kernel_thunk();
}
inline void ThunkProto::set_has_kernel_thunk() {
  _impl_._oneof_case_[0] = kKernelThunk;
}
inline void ThunkProto::clear_kernel_thunk() {
  if (_internal_has_kernel_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.kernel_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::KernelThunkProto* ThunkProto::release_kernel_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.kernel_thunk)
  if (_internal_has_kernel_thunk()) {
    clear_has_impl();
    ::xla::cpu::KernelThunkProto* temp = _impl_.impl_.kernel_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.kernel_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::KernelThunkProto& ThunkProto::_internal_kernel_thunk() const {
  return _internal_has_kernel_thunk()
      ? *_impl_.impl_.kernel_thunk_
      : reinterpret_cast< ::xla::cpu::KernelThunkProto&>(::xla::cpu::_KernelThunkProto_default_instance_);
}
inline const ::xla::cpu::KernelThunkProto& ThunkProto::kernel_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.kernel_thunk)
  return _internal_kernel_thunk();
}
inline ::xla::cpu::KernelThunkProto* ThunkProto::unsafe_arena_release_kernel_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.kernel_thunk)
  if (_internal_has_kernel_thunk()) {
    clear_has_impl();
    ::xla::cpu::KernelThunkProto* temp = _impl_.impl_.kernel_thunk_;
    _impl_.impl_.kernel_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_kernel_thunk(::xla::cpu::KernelThunkProto* kernel_thunk) {
  clear_impl();
  if (kernel_thunk) {
    set_has_kernel_thunk();
    _impl_.impl_.kernel_thunk_ = kernel_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.kernel_thunk)
}
inline ::xla::cpu::KernelThunkProto* ThunkProto::_internal_mutable_kernel_thunk() {
  if (!_internal_has_kernel_thunk()) {
    clear_impl();
    set_has_kernel_thunk();
    _impl_.impl_.kernel_thunk_ = CreateMaybeMessage< ::xla::cpu::KernelThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.kernel_thunk_;
}
inline ::xla::cpu::KernelThunkProto* ThunkProto::mutable_kernel_thunk() {
  ::xla::cpu::KernelThunkProto* _msg = _internal_mutable_kernel_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.kernel_thunk)
  return _msg;
}

// .xla.cpu.CopyThunkProto copy_thunk = 12;
inline bool ThunkProto::_internal_has_copy_thunk() const {
  return impl_case() == kCopyThunk;
}
inline bool ThunkProto::has_copy_thunk() const {
  return _internal_has_copy_thunk();
}
inline void ThunkProto::set_has_copy_thunk() {
  _impl_._oneof_case_[0] = kCopyThunk;
}
inline void ThunkProto::clear_copy_thunk() {
  if (_internal_has_copy_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.copy_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::CopyThunkProto* ThunkProto::release_copy_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.copy_thunk)
  if (_internal_has_copy_thunk()) {
    clear_has_impl();
    ::xla::cpu::CopyThunkProto* temp = _impl_.impl_.copy_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.copy_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::CopyThunkProto& ThunkProto::_internal_copy_thunk() const {
  return _internal_has_copy_thunk()
      ? *_impl_.impl_.copy_thunk_
      : reinterpret_cast< ::xla::cpu::CopyThunkProto&>(::xla::cpu::_CopyThunkProto_default_instance_);
}
inline const ::xla::cpu::CopyThunkProto& ThunkProto::copy_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.copy_thunk)
  return _internal_copy_thunk();
}
inline ::xla::cpu::CopyThunkProto* ThunkProto::unsafe_arena_release_copy_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.copy_thunk)
  if (_internal_has_copy_thunk()) {
    clear_has_impl();
    ::xla::cpu::CopyThunkProto* temp = _impl_.impl_.copy_thunk_;
    _impl_.impl_.copy_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_copy_thunk(::xla::cpu::CopyThunkProto* copy_thunk) {
  clear_impl();
  if (copy_thunk) {
    set_has_copy_thunk();
    _impl_.impl_.copy_thunk_ = copy_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.copy_thunk)
}
inline ::xla::cpu::CopyThunkProto* ThunkProto::_internal_mutable_copy_thunk() {
  if (!_internal_has_copy_thunk()) {
    clear_impl();
    set_has_copy_thunk();
    _impl_.impl_.copy_thunk_ = CreateMaybeMessage< ::xla::cpu::CopyThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.copy_thunk_;
}
inline ::xla::cpu::CopyThunkProto* ThunkProto::mutable_copy_thunk() {
  ::xla::cpu::CopyThunkProto* _msg = _internal_mutable_copy_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.copy_thunk)
  return _msg;
}

// .xla.cpu.FftThunkProto fft_thunk = 13;
inline bool ThunkProto::_internal_has_fft_thunk() const {
  return impl_case() == kFftThunk;
}
inline bool ThunkProto::has_fft_thunk() const {
  return _internal_has_fft_thunk();
}
inline void ThunkProto::set_has_fft_thunk() {
  _impl_._oneof_case_[0] = kFftThunk;
}
inline void ThunkProto::clear_fft_thunk() {
  if (_internal_has_fft_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.fft_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::FftThunkProto* ThunkProto::release_fft_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.fft_thunk)
  if (_internal_has_fft_thunk()) {
    clear_has_impl();
    ::xla::cpu::FftThunkProto* temp = _impl_.impl_.fft_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.fft_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::FftThunkProto& ThunkProto::_internal_fft_thunk() const {
  return _internal_has_fft_thunk()
      ? *_impl_.impl_.fft_thunk_
      : reinterpret_cast< ::xla::cpu::FftThunkProto&>(::xla::cpu::_FftThunkProto_default_instance_);
}
inline const ::xla::cpu::FftThunkProto& ThunkProto::fft_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.fft_thunk)
  return _internal_fft_thunk();
}
inline ::xla::cpu::FftThunkProto* ThunkProto::unsafe_arena_release_fft_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.fft_thunk)
  if (_internal_has_fft_thunk()) {
    clear_has_impl();
    ::xla::cpu::FftThunkProto* temp = _impl_.impl_.fft_thunk_;
    _impl_.impl_.fft_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_fft_thunk(::xla::cpu::FftThunkProto* fft_thunk) {
  clear_impl();
  if (fft_thunk) {
    set_has_fft_thunk();
    _impl_.impl_.fft_thunk_ = fft_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.fft_thunk)
}
inline ::xla::cpu::FftThunkProto* ThunkProto::_internal_mutable_fft_thunk() {
  if (!_internal_has_fft_thunk()) {
    clear_impl();
    set_has_fft_thunk();
    _impl_.impl_.fft_thunk_ = CreateMaybeMessage< ::xla::cpu::FftThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.fft_thunk_;
}
inline ::xla::cpu::FftThunkProto* ThunkProto::mutable_fft_thunk() {
  ::xla::cpu::FftThunkProto* _msg = _internal_mutable_fft_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.fft_thunk)
  return _msg;
}

// .xla.cpu.InfeedThunkProto infeed_thunk = 14;
inline bool ThunkProto::_internal_has_infeed_thunk() const {
  return impl_case() == kInfeedThunk;
}
inline bool ThunkProto::has_infeed_thunk() const {
  return _internal_has_infeed_thunk();
}
inline void ThunkProto::set_has_infeed_thunk() {
  _impl_._oneof_case_[0] = kInfeedThunk;
}
inline void ThunkProto::clear_infeed_thunk() {
  if (_internal_has_infeed_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.infeed_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::InfeedThunkProto* ThunkProto::release_infeed_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.infeed_thunk)
  if (_internal_has_infeed_thunk()) {
    clear_has_impl();
    ::xla::cpu::InfeedThunkProto* temp = _impl_.impl_.infeed_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.infeed_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::InfeedThunkProto& ThunkProto::_internal_infeed_thunk() const {
  return _internal_has_infeed_thunk()
      ? *_impl_.impl_.infeed_thunk_
      : reinterpret_cast< ::xla::cpu::InfeedThunkProto&>(::xla::cpu::_InfeedThunkProto_default_instance_);
}
inline const ::xla::cpu::InfeedThunkProto& ThunkProto::infeed_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.infeed_thunk)
  return _internal_infeed_thunk();
}
inline ::xla::cpu::InfeedThunkProto* ThunkProto::unsafe_arena_release_infeed_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.infeed_thunk)
  if (_internal_has_infeed_thunk()) {
    clear_has_impl();
    ::xla::cpu::InfeedThunkProto* temp = _impl_.impl_.infeed_thunk_;
    _impl_.impl_.infeed_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_infeed_thunk(::xla::cpu::InfeedThunkProto* infeed_thunk) {
  clear_impl();
  if (infeed_thunk) {
    set_has_infeed_thunk();
    _impl_.impl_.infeed_thunk_ = infeed_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.infeed_thunk)
}
inline ::xla::cpu::InfeedThunkProto* ThunkProto::_internal_mutable_infeed_thunk() {
  if (!_internal_has_infeed_thunk()) {
    clear_impl();
    set_has_infeed_thunk();
    _impl_.impl_.infeed_thunk_ = CreateMaybeMessage< ::xla::cpu::InfeedThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.infeed_thunk_;
}
inline ::xla::cpu::InfeedThunkProto* ThunkProto::mutable_infeed_thunk() {
  ::xla::cpu::InfeedThunkProto* _msg = _internal_mutable_infeed_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.infeed_thunk)
  return _msg;
}

// .xla.cpu.OutfeedThunkProto outfeed_thunk = 15;
inline bool ThunkProto::_internal_has_outfeed_thunk() const {
  return impl_case() == kOutfeedThunk;
}
inline bool ThunkProto::has_outfeed_thunk() const {
  return _internal_has_outfeed_thunk();
}
inline void ThunkProto::set_has_outfeed_thunk() {
  _impl_._oneof_case_[0] = kOutfeedThunk;
}
inline void ThunkProto::clear_outfeed_thunk() {
  if (_internal_has_outfeed_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.outfeed_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::OutfeedThunkProto* ThunkProto::release_outfeed_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.outfeed_thunk)
  if (_internal_has_outfeed_thunk()) {
    clear_has_impl();
    ::xla::cpu::OutfeedThunkProto* temp = _impl_.impl_.outfeed_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.outfeed_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::OutfeedThunkProto& ThunkProto::_internal_outfeed_thunk() const {
  return _internal_has_outfeed_thunk()
      ? *_impl_.impl_.outfeed_thunk_
      : reinterpret_cast< ::xla::cpu::OutfeedThunkProto&>(::xla::cpu::_OutfeedThunkProto_default_instance_);
}
inline const ::xla::cpu::OutfeedThunkProto& ThunkProto::outfeed_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.outfeed_thunk)
  return _internal_outfeed_thunk();
}
inline ::xla::cpu::OutfeedThunkProto* ThunkProto::unsafe_arena_release_outfeed_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.outfeed_thunk)
  if (_internal_has_outfeed_thunk()) {
    clear_has_impl();
    ::xla::cpu::OutfeedThunkProto* temp = _impl_.impl_.outfeed_thunk_;
    _impl_.impl_.outfeed_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_outfeed_thunk(::xla::cpu::OutfeedThunkProto* outfeed_thunk) {
  clear_impl();
  if (outfeed_thunk) {
    set_has_outfeed_thunk();
    _impl_.impl_.outfeed_thunk_ = outfeed_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.outfeed_thunk)
}
inline ::xla::cpu::OutfeedThunkProto* ThunkProto::_internal_mutable_outfeed_thunk() {
  if (!_internal_has_outfeed_thunk()) {
    clear_impl();
    set_has_outfeed_thunk();
    _impl_.impl_.outfeed_thunk_ = CreateMaybeMessage< ::xla::cpu::OutfeedThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.outfeed_thunk_;
}
inline ::xla::cpu::OutfeedThunkProto* ThunkProto::mutable_outfeed_thunk() {
  ::xla::cpu::OutfeedThunkProto* _msg = _internal_mutable_outfeed_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.outfeed_thunk)
  return _msg;
}

// .xla.cpu.CustomCallThunkProto custom_call_thunk = 16;
inline bool ThunkProto::_internal_has_custom_call_thunk() const {
  return impl_case() == kCustomCallThunk;
}
inline bool ThunkProto::has_custom_call_thunk() const {
  return _internal_has_custom_call_thunk();
}
inline void ThunkProto::set_has_custom_call_thunk() {
  _impl_._oneof_case_[0] = kCustomCallThunk;
}
inline void ThunkProto::clear_custom_call_thunk() {
  if (_internal_has_custom_call_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.custom_call_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::CustomCallThunkProto* ThunkProto::release_custom_call_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.custom_call_thunk)
  if (_internal_has_custom_call_thunk()) {
    clear_has_impl();
    ::xla::cpu::CustomCallThunkProto* temp = _impl_.impl_.custom_call_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.custom_call_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::CustomCallThunkProto& ThunkProto::_internal_custom_call_thunk() const {
  return _internal_has_custom_call_thunk()
      ? *_impl_.impl_.custom_call_thunk_
      : reinterpret_cast< ::xla::cpu::CustomCallThunkProto&>(::xla::cpu::_CustomCallThunkProto_default_instance_);
}
inline const ::xla::cpu::CustomCallThunkProto& ThunkProto::custom_call_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.custom_call_thunk)
  return _internal_custom_call_thunk();
}
inline ::xla::cpu::CustomCallThunkProto* ThunkProto::unsafe_arena_release_custom_call_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.custom_call_thunk)
  if (_internal_has_custom_call_thunk()) {
    clear_has_impl();
    ::xla::cpu::CustomCallThunkProto* temp = _impl_.impl_.custom_call_thunk_;
    _impl_.impl_.custom_call_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_custom_call_thunk(::xla::cpu::CustomCallThunkProto* custom_call_thunk) {
  clear_impl();
  if (custom_call_thunk) {
    set_has_custom_call_thunk();
    _impl_.impl_.custom_call_thunk_ = custom_call_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.custom_call_thunk)
}
inline ::xla::cpu::CustomCallThunkProto* ThunkProto::_internal_mutable_custom_call_thunk() {
  if (!_internal_has_custom_call_thunk()) {
    clear_impl();
    set_has_custom_call_thunk();
    _impl_.impl_.custom_call_thunk_ = CreateMaybeMessage< ::xla::cpu::CustomCallThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.custom_call_thunk_;
}
inline ::xla::cpu::CustomCallThunkProto* ThunkProto::mutable_custom_call_thunk() {
  ::xla::cpu::CustomCallThunkProto* _msg = _internal_mutable_custom_call_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.custom_call_thunk)
  return _msg;
}

// .xla.cpu.ConvolutionThunkProto convolution_thunk = 17;
inline bool ThunkProto::_internal_has_convolution_thunk() const {
  return impl_case() == kConvolutionThunk;
}
inline bool ThunkProto::has_convolution_thunk() const {
  return _internal_has_convolution_thunk();
}
inline void ThunkProto::set_has_convolution_thunk() {
  _impl_._oneof_case_[0] = kConvolutionThunk;
}
inline void ThunkProto::clear_convolution_thunk() {
  if (_internal_has_convolution_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.convolution_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::ConvolutionThunkProto* ThunkProto::release_convolution_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.convolution_thunk)
  if (_internal_has_convolution_thunk()) {
    clear_has_impl();
    ::xla::cpu::ConvolutionThunkProto* temp = _impl_.impl_.convolution_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.convolution_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::ConvolutionThunkProto& ThunkProto::_internal_convolution_thunk() const {
  return _internal_has_convolution_thunk()
      ? *_impl_.impl_.convolution_thunk_
      : reinterpret_cast< ::xla::cpu::ConvolutionThunkProto&>(::xla::cpu::_ConvolutionThunkProto_default_instance_);
}
inline const ::xla::cpu::ConvolutionThunkProto& ThunkProto::convolution_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.convolution_thunk)
  return _internal_convolution_thunk();
}
inline ::xla::cpu::ConvolutionThunkProto* ThunkProto::unsafe_arena_release_convolution_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.convolution_thunk)
  if (_internal_has_convolution_thunk()) {
    clear_has_impl();
    ::xla::cpu::ConvolutionThunkProto* temp = _impl_.impl_.convolution_thunk_;
    _impl_.impl_.convolution_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_convolution_thunk(::xla::cpu::ConvolutionThunkProto* convolution_thunk) {
  clear_impl();
  if (convolution_thunk) {
    set_has_convolution_thunk();
    _impl_.impl_.convolution_thunk_ = convolution_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.convolution_thunk)
}
inline ::xla::cpu::ConvolutionThunkProto* ThunkProto::_internal_mutable_convolution_thunk() {
  if (!_internal_has_convolution_thunk()) {
    clear_impl();
    set_has_convolution_thunk();
    _impl_.impl_.convolution_thunk_ = CreateMaybeMessage< ::xla::cpu::ConvolutionThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.convolution_thunk_;
}
inline ::xla::cpu::ConvolutionThunkProto* ThunkProto::mutable_convolution_thunk() {
  ::xla::cpu::ConvolutionThunkProto* _msg = _internal_mutable_convolution_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.convolution_thunk)
  return _msg;
}

// .xla.cpu.CollectiveThunkProto collective_thunk = 18;
inline bool ThunkProto::_internal_has_collective_thunk() const {
  return impl_case() == kCollectiveThunk;
}
inline bool ThunkProto::has_collective_thunk() const {
  return _internal_has_collective_thunk();
}
inline void ThunkProto::set_has_collective_thunk() {
  _impl_._oneof_case_[0] = kCollectiveThunk;
}
inline void ThunkProto::clear_collective_thunk() {
  if (_internal_has_collective_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.collective_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::CollectiveThunkProto* ThunkProto::release_collective_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.collective_thunk)
  if (_internal_has_collective_thunk()) {
    clear_has_impl();
    ::xla::cpu::CollectiveThunkProto* temp = _impl_.impl_.collective_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.collective_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::CollectiveThunkProto& ThunkProto::_internal_collective_thunk() const {
  return _internal_has_collective_thunk()
      ? *_impl_.impl_.collective_thunk_
      : reinterpret_cast< ::xla::cpu::CollectiveThunkProto&>(::xla::cpu::_CollectiveThunkProto_default_instance_);
}
inline const ::xla::cpu::CollectiveThunkProto& ThunkProto::collective_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.collective_thunk)
  return _internal_collective_thunk();
}
inline ::xla::cpu::CollectiveThunkProto* ThunkProto::unsafe_arena_release_collective_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.collective_thunk)
  if (_internal_has_collective_thunk()) {
    clear_has_impl();
    ::xla::cpu::CollectiveThunkProto* temp = _impl_.impl_.collective_thunk_;
    _impl_.impl_.collective_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_collective_thunk(::xla::cpu::CollectiveThunkProto* collective_thunk) {
  clear_impl();
  if (collective_thunk) {
    set_has_collective_thunk();
    _impl_.impl_.collective_thunk_ = collective_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.collective_thunk)
}
inline ::xla::cpu::CollectiveThunkProto* ThunkProto::_internal_mutable_collective_thunk() {
  if (!_internal_has_collective_thunk()) {
    clear_impl();
    set_has_collective_thunk();
    _impl_.impl_.collective_thunk_ = CreateMaybeMessage< ::xla::cpu::CollectiveThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.collective_thunk_;
}
inline ::xla::cpu::CollectiveThunkProto* ThunkProto::mutable_collective_thunk() {
  ::xla::cpu::CollectiveThunkProto* _msg = _internal_mutable_collective_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.collective_thunk)
  return _msg;
}

// .xla.cpu.PartitionIdThunkProto partition_id_thunk = 19;
inline bool ThunkProto::_internal_has_partition_id_thunk() const {
  return impl_case() == kPartitionIdThunk;
}
inline bool ThunkProto::has_partition_id_thunk() const {
  return _internal_has_partition_id_thunk();
}
inline void ThunkProto::set_has_partition_id_thunk() {
  _impl_._oneof_case_[0] = kPartitionIdThunk;
}
inline void ThunkProto::clear_partition_id_thunk() {
  if (_internal_has_partition_id_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.partition_id_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::PartitionIdThunkProto* ThunkProto::release_partition_id_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.partition_id_thunk)
  if (_internal_has_partition_id_thunk()) {
    clear_has_impl();
    ::xla::cpu::PartitionIdThunkProto* temp = _impl_.impl_.partition_id_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.partition_id_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::PartitionIdThunkProto& ThunkProto::_internal_partition_id_thunk() const {
  return _internal_has_partition_id_thunk()
      ? *_impl_.impl_.partition_id_thunk_
      : reinterpret_cast< ::xla::cpu::PartitionIdThunkProto&>(::xla::cpu::_PartitionIdThunkProto_default_instance_);
}
inline const ::xla::cpu::PartitionIdThunkProto& ThunkProto::partition_id_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.partition_id_thunk)
  return _internal_partition_id_thunk();
}
inline ::xla::cpu::PartitionIdThunkProto* ThunkProto::unsafe_arena_release_partition_id_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.partition_id_thunk)
  if (_internal_has_partition_id_thunk()) {
    clear_has_impl();
    ::xla::cpu::PartitionIdThunkProto* temp = _impl_.impl_.partition_id_thunk_;
    _impl_.impl_.partition_id_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_partition_id_thunk(::xla::cpu::PartitionIdThunkProto* partition_id_thunk) {
  clear_impl();
  if (partition_id_thunk) {
    set_has_partition_id_thunk();
    _impl_.impl_.partition_id_thunk_ = partition_id_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.partition_id_thunk)
}
inline ::xla::cpu::PartitionIdThunkProto* ThunkProto::_internal_mutable_partition_id_thunk() {
  if (!_internal_has_partition_id_thunk()) {
    clear_impl();
    set_has_partition_id_thunk();
    _impl_.impl_.partition_id_thunk_ = CreateMaybeMessage< ::xla::cpu::PartitionIdThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.partition_id_thunk_;
}
inline ::xla::cpu::PartitionIdThunkProto* ThunkProto::mutable_partition_id_thunk() {
  ::xla::cpu::PartitionIdThunkProto* _msg = _internal_mutable_partition_id_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.partition_id_thunk)
  return _msg;
}

// .xla.cpu.ReplicaIdThunkProto replica_id_thunk = 20;
inline bool ThunkProto::_internal_has_replica_id_thunk() const {
  return impl_case() == kReplicaIdThunk;
}
inline bool ThunkProto::has_replica_id_thunk() const {
  return _internal_has_replica_id_thunk();
}
inline void ThunkProto::set_has_replica_id_thunk() {
  _impl_._oneof_case_[0] = kReplicaIdThunk;
}
inline void ThunkProto::clear_replica_id_thunk() {
  if (_internal_has_replica_id_thunk()) {
    if (GetArenaForAllocation() == nullptr) {
      delete _impl_.impl_.replica_id_thunk_;
    }
    clear_has_impl();
  }
}
inline ::xla::cpu::ReplicaIdThunkProto* ThunkProto::release_replica_id_thunk() {
  // @@protoc_insertion_point(field_release:xla.cpu.ThunkProto.replica_id_thunk)
  if (_internal_has_replica_id_thunk()) {
    clear_has_impl();
    ::xla::cpu::ReplicaIdThunkProto* temp = _impl_.impl_.replica_id_thunk_;
    if (GetArenaForAllocation() != nullptr) {
      temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp);
    }
    _impl_.impl_.replica_id_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::xla::cpu::ReplicaIdThunkProto& ThunkProto::_internal_replica_id_thunk() const {
  return _internal_has_replica_id_thunk()
      ? *_impl_.impl_.replica_id_thunk_
      : reinterpret_cast< ::xla::cpu::ReplicaIdThunkProto&>(::xla::cpu::_ReplicaIdThunkProto_default_instance_);
}
inline const ::xla::cpu::ReplicaIdThunkProto& ThunkProto::replica_id_thunk() const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkProto.replica_id_thunk)
  return _internal_replica_id_thunk();
}
inline ::xla::cpu::ReplicaIdThunkProto* ThunkProto::unsafe_arena_release_replica_id_thunk() {
  // @@protoc_insertion_point(field_unsafe_arena_release:xla.cpu.ThunkProto.replica_id_thunk)
  if (_internal_has_replica_id_thunk()) {
    clear_has_impl();
    ::xla::cpu::ReplicaIdThunkProto* temp = _impl_.impl_.replica_id_thunk_;
    _impl_.impl_.replica_id_thunk_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void ThunkProto::unsafe_arena_set_allocated_replica_id_thunk(::xla::cpu::ReplicaIdThunkProto* replica_id_thunk) {
  clear_impl();
  if (replica_id_thunk) {
    set_has_replica_id_thunk();
    _impl_.impl_.replica_id_thunk_ = replica_id_thunk;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.cpu.ThunkProto.replica_id_thunk)
}
inline ::xla::cpu::ReplicaIdThunkProto* ThunkProto::_internal_mutable_replica_id_thunk() {
  if (!_internal_has_replica_id_thunk()) {
    clear_impl();
    set_has_replica_id_thunk();
    _impl_.impl_.replica_id_thunk_ = CreateMaybeMessage< ::xla::cpu::ReplicaIdThunkProto >(GetArenaForAllocation());
  }
  return _impl_.impl_.replica_id_thunk_;
}
inline ::xla::cpu::ReplicaIdThunkProto* ThunkProto::mutable_replica_id_thunk() {
  ::xla::cpu::ReplicaIdThunkProto* _msg = _internal_mutable_replica_id_thunk();
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkProto.replica_id_thunk)
  return _msg;
}

inline bool ThunkProto::has_impl() const {
  return impl_case() != IMPL_NOT_SET;
}
inline void ThunkProto::clear_has_impl() {
  _impl_._oneof_case_[0] = IMPL_NOT_SET;
}
inline ThunkProto::ImplCase ThunkProto::impl_case() const {
  return ThunkProto::ImplCase(_impl_._oneof_case_[0]);
}
// -------------------------------------------------------------------

// ThunkSequenceProto

// repeated .xla.cpu.ThunkProto thunks = 1;
inline int ThunkSequenceProto::_internal_thunks_size() const {
  return _impl_.thunks_.size();
}
inline int ThunkSequenceProto::thunks_size() const {
  return _internal_thunks_size();
}
inline void ThunkSequenceProto::clear_thunks() {
  _impl_.thunks_.Clear();
}
inline ::xla::cpu::ThunkProto* ThunkSequenceProto::mutable_thunks(int index) {
  // @@protoc_insertion_point(field_mutable:xla.cpu.ThunkSequenceProto.thunks)
  return _impl_.thunks_.Mutable(index);
}
inline ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkProto >*
ThunkSequenceProto::mutable_thunks() {
  // @@protoc_insertion_point(field_mutable_list:xla.cpu.ThunkSequenceProto.thunks)
  return &_impl_.thunks_;
}
inline const ::xla::cpu::ThunkProto& ThunkSequenceProto::_internal_thunks(int index) const {
  return _impl_.thunks_.Get(index);
}
inline const ::xla::cpu::ThunkProto& ThunkSequenceProto::thunks(int index) const {
  // @@protoc_insertion_point(field_get:xla.cpu.ThunkSequenceProto.thunks)
  return _internal_thunks(index);
}
inline ::xla::cpu::ThunkProto* ThunkSequenceProto::_internal_add_thunks() {
  return _impl_.thunks_.Add();
}
inline ::xla::cpu::ThunkProto* ThunkSequenceProto::add_thunks() {
  ::xla::cpu::ThunkProto* _add = _internal_add_thunks();
  // @@protoc_insertion_point(field_add:xla.cpu.ThunkSequenceProto.thunks)
  return _add;
}
inline const ::PROTOBUF_NAMESPACE_ID::RepeatedPtrField< ::xla::cpu::ThunkProto >&
ThunkSequenceProto::thunks() const {
  // @@protoc_insertion_point(field_list:xla.cpu.ThunkSequenceProto.thunks)
  return _impl_.thunks_;
}

#ifdef __GNUC__
  #pragma GCC diagnostic pop
#endif  // __GNUC__
// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------

// -------------------------------------------------------------------


// @@protoc_insertion_point(namespace_scope)

}  // namespace cpu
}  // namespace xla

PROTOBUF_NAMESPACE_OPEN

template <> struct is_proto_enum< ::xla::cpu::ResourceProto_Kind> : ::std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor< ::xla::cpu::ResourceProto_Kind>() {
  return ::xla::cpu::ResourceProto_Kind_descriptor();
}
template <> struct is_proto_enum< ::xla::cpu::SortDirectionProto> : ::std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor< ::xla::cpu::SortDirectionProto>() {
  return ::xla::cpu::SortDirectionProto_descriptor();
}

PROTOBUF_NAMESPACE_CLOSE

// @@protoc_insertion_point(global_scope)

#include <google/protobuf/port_undef.inc>
#endif  // GOOGLE_PROTOBUF_INCLUDED_GOOGLE_PROTOBUF_INCLUDED_xla_2fbackends_2fcpu_2fruntime_2fthunk_2eproto