// Generated by the protocol buffer compiler. DO NOT EDIT! // NO CHECKED-IN PROTOBUF GENCODE // source: xla/service/gpu/backend_configs.proto // Protobuf C++ Version: 5.28.3 #ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto_2epb_2eh #define GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto_2epb_2eh #include #include #include #include #include "google/protobuf/runtime_version.h" #if PROTOBUF_VERSION != 5028003 #error "Protobuf C++ gencode is built with an incompatible version of" #error "Protobuf C++ headers/runtime. See" #error "https://protobuf.dev/support/cross-version-runtime-guarantee/#cpp" #endif #include "google/protobuf/io/coded_stream.h" #include "google/protobuf/arena.h" #include "google/protobuf/arenastring.h" #include "google/protobuf/generated_message_tctable_decl.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/autotuning.pb.h" #include "xla/tsl/protobuf/dnn.pb.h" #include "xla/xla_data.pb.h" // @@protoc_insertion_point(includes) // Must be included last. #include "google/protobuf/port_def.inc" #define PROTOBUF_INTERNAL_EXPORT_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto PROTOBUF_EXPORT namespace google { namespace protobuf { namespace internal { class AnyMetadata; } // namespace internal } // namespace protobuf } // namespace google // Internal implementation detail -- do not use these members. struct PROTOBUF_EXPORT TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto { static const ::uint32_t offsets[]; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable descriptor_table_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; namespace xla { namespace gpu { class BitcastBackendConfig; struct BitcastBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern BitcastBackendConfigDefaultTypeInternal _BitcastBackendConfig_default_instance_; class BlockLevelFusionConfig; struct BlockLevelFusionConfigDefaultTypeInternal; PROTOBUF_EXPORT extern BlockLevelFusionConfigDefaultTypeInternal _BlockLevelFusionConfig_default_instance_; class CollectiveBackendConfig; struct CollectiveBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CollectiveBackendConfigDefaultTypeInternal _CollectiveBackendConfig_default_instance_; class CuDnnFusionConfig; struct CuDnnFusionConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CuDnnFusionConfigDefaultTypeInternal _CuDnnFusionConfig_default_instance_; class CudnnConvBackendConfig; struct CudnnConvBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CudnnConvBackendConfigDefaultTypeInternal _CudnnConvBackendConfig_default_instance_; class CudnnNormBackendConfig; struct CudnnNormBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CudnnNormBackendConfigDefaultTypeInternal _CudnnNormBackendConfig_default_instance_; class CudnnfMHABackendConfig; struct CudnnfMHABackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CudnnfMHABackendConfigDefaultTypeInternal _CudnnfMHABackendConfig_default_instance_; class CustomCallBackendConfig; struct CustomCallBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CustomCallBackendConfigDefaultTypeInternal _CustomCallBackendConfig_default_instance_; class CustomFusionConfig; struct CustomFusionConfigDefaultTypeInternal; PROTOBUF_EXPORT extern CustomFusionConfigDefaultTypeInternal _CustomFusionConfig_default_instance_; class DynamicMemcpyConfig; struct DynamicMemcpyConfigDefaultTypeInternal; PROTOBUF_EXPORT extern DynamicMemcpyConfigDefaultTypeInternal _DynamicMemcpyConfig_default_instance_; class FusionBackendConfig; struct FusionBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern FusionBackendConfigDefaultTypeInternal _FusionBackendConfig_default_instance_; class GemmBackendConfig; struct GemmBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern GemmBackendConfigDefaultTypeInternal _GemmBackendConfig_default_instance_; class GpuBackendConfig; struct GpuBackendConfigDefaultTypeInternal; PROTOBUF_EXPORT extern GpuBackendConfigDefaultTypeInternal _GpuBackendConfig_default_instance_; class ReificationCost; struct ReificationCostDefaultTypeInternal; PROTOBUF_EXPORT extern ReificationCostDefaultTypeInternal _ReificationCost_default_instance_; class Tile; struct TileDefaultTypeInternal; PROTOBUF_EXPORT extern TileDefaultTypeInternal _Tile_default_instance_; } // namespace gpu } // namespace xla namespace google { namespace protobuf { } // namespace protobuf } // namespace google namespace xla { namespace gpu { enum GemmBackendConfig_Epilogue : int { GemmBackendConfig_Epilogue_DEFAULT = 0, GemmBackendConfig_Epilogue_BIAS = 1, GemmBackendConfig_Epilogue_RELU = 2, GemmBackendConfig_Epilogue_BIAS_RELU = 3, GemmBackendConfig_Epilogue_GELU = 4, GemmBackendConfig_Epilogue_GELU_AUX = 5, GemmBackendConfig_Epilogue_BIAS_GELU = 6, GemmBackendConfig_Epilogue_BIAS_GELU_AUX = 7, GemmBackendConfig_Epilogue_SILU = 8, GemmBackendConfig_Epilogue_BIAS_SILU = 9, GemmBackendConfig_Epilogue_GemmBackendConfig_Epilogue_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::min(), GemmBackendConfig_Epilogue_GemmBackendConfig_Epilogue_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::max(), }; PROTOBUF_EXPORT bool GemmBackendConfig_Epilogue_IsValid(int value); PROTOBUF_EXPORT extern const uint32_t GemmBackendConfig_Epilogue_internal_data_[]; constexpr GemmBackendConfig_Epilogue GemmBackendConfig_Epilogue_Epilogue_MIN = static_cast(0); constexpr GemmBackendConfig_Epilogue GemmBackendConfig_Epilogue_Epilogue_MAX = static_cast(9); constexpr int GemmBackendConfig_Epilogue_Epilogue_ARRAYSIZE = 9 + 1; PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor* GemmBackendConfig_Epilogue_descriptor(); template const std::string& GemmBackendConfig_Epilogue_Name(T value) { static_assert(std::is_same::value || std::is_integral::value, "Incorrect type passed to Epilogue_Name()."); return GemmBackendConfig_Epilogue_Name(static_cast(value)); } template <> inline const std::string& GemmBackendConfig_Epilogue_Name(GemmBackendConfig_Epilogue value) { return ::google::protobuf::internal::NameOfDenseEnum( static_cast(value)); } inline bool GemmBackendConfig_Epilogue_Parse(absl::string_view name, GemmBackendConfig_Epilogue* value) { return ::google::protobuf::internal::ParseNamedEnum( GemmBackendConfig_Epilogue_descriptor(), name, value); } enum CollectiveBackendConfig_CollectiveBackend : int { CollectiveBackendConfig_CollectiveBackend_DEFAULT = 0, CollectiveBackendConfig_CollectiveBackend_NVSHMEM = 1, CollectiveBackendConfig_CollectiveBackend_CollectiveBackendConfig_CollectiveBackend_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::min(), CollectiveBackendConfig_CollectiveBackend_CollectiveBackendConfig_CollectiveBackend_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::max(), }; PROTOBUF_EXPORT bool CollectiveBackendConfig_CollectiveBackend_IsValid(int value); PROTOBUF_EXPORT extern const uint32_t CollectiveBackendConfig_CollectiveBackend_internal_data_[]; constexpr CollectiveBackendConfig_CollectiveBackend CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_MIN = static_cast(0); constexpr CollectiveBackendConfig_CollectiveBackend CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_MAX = static_cast(1); constexpr int CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_ARRAYSIZE = 1 + 1; PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor* CollectiveBackendConfig_CollectiveBackend_descriptor(); template const std::string& CollectiveBackendConfig_CollectiveBackend_Name(T value) { static_assert(std::is_same::value || std::is_integral::value, "Incorrect type passed to CollectiveBackend_Name()."); return CollectiveBackendConfig_CollectiveBackend_Name(static_cast(value)); } template <> inline const std::string& CollectiveBackendConfig_CollectiveBackend_Name(CollectiveBackendConfig_CollectiveBackend value) { return ::google::protobuf::internal::NameOfDenseEnum( static_cast(value)); } inline bool CollectiveBackendConfig_CollectiveBackend_Parse(absl::string_view name, CollectiveBackendConfig_CollectiveBackend* value) { return ::google::protobuf::internal::ParseNamedEnum( CollectiveBackendConfig_CollectiveBackend_descriptor(), name, value); } enum CudnnNormBackendConfig_Kind : int { CudnnNormBackendConfig_Kind_LAYER_FWD_INFER = 0, CudnnNormBackendConfig_Kind_LAYER_FWD_TRAIN = 1, CudnnNormBackendConfig_Kind_LAYER_BWD = 2, CudnnNormBackendConfig_Kind_CudnnNormBackendConfig_Kind_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::min(), CudnnNormBackendConfig_Kind_CudnnNormBackendConfig_Kind_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::max(), }; PROTOBUF_EXPORT bool CudnnNormBackendConfig_Kind_IsValid(int value); PROTOBUF_EXPORT extern const uint32_t CudnnNormBackendConfig_Kind_internal_data_[]; constexpr CudnnNormBackendConfig_Kind CudnnNormBackendConfig_Kind_Kind_MIN = static_cast(0); constexpr CudnnNormBackendConfig_Kind CudnnNormBackendConfig_Kind_Kind_MAX = static_cast(2); constexpr int CudnnNormBackendConfig_Kind_Kind_ARRAYSIZE = 2 + 1; PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor* CudnnNormBackendConfig_Kind_descriptor(); template const std::string& CudnnNormBackendConfig_Kind_Name(T value) { static_assert(std::is_same::value || std::is_integral::value, "Incorrect type passed to Kind_Name()."); return CudnnNormBackendConfig_Kind_Name(static_cast(value)); } template <> inline const std::string& CudnnNormBackendConfig_Kind_Name(CudnnNormBackendConfig_Kind value) { return ::google::protobuf::internal::NameOfDenseEnum( static_cast(value)); } inline bool CudnnNormBackendConfig_Kind_Parse(absl::string_view name, CudnnNormBackendConfig_Kind* value) { return ::google::protobuf::internal::ParseNamedEnum( CudnnNormBackendConfig_Kind_descriptor(), name, value); } enum CudnnfMHABackendConfig_MaskType : int { CudnnfMHABackendConfig_MaskType_NO_MASK = 0, CudnnfMHABackendConfig_MaskType_PADDING = 1, CudnnfMHABackendConfig_MaskType_CAUSAL = 2, CudnnfMHABackendConfig_MaskType_PADDING_CAUSAL = 3, CudnnfMHABackendConfig_MaskType_ALIBI = 4, CudnnfMHABackendConfig_MaskType_CudnnfMHABackendConfig_MaskType_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::min(), CudnnfMHABackendConfig_MaskType_CudnnfMHABackendConfig_MaskType_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<::int32_t>::max(), }; PROTOBUF_EXPORT bool CudnnfMHABackendConfig_MaskType_IsValid(int value); PROTOBUF_EXPORT extern const uint32_t CudnnfMHABackendConfig_MaskType_internal_data_[]; constexpr CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig_MaskType_MaskType_MIN = static_cast(0); constexpr CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig_MaskType_MaskType_MAX = static_cast(4); constexpr int CudnnfMHABackendConfig_MaskType_MaskType_ARRAYSIZE = 4 + 1; PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor* CudnnfMHABackendConfig_MaskType_descriptor(); template const std::string& CudnnfMHABackendConfig_MaskType_Name(T value) { static_assert(std::is_same::value || std::is_integral::value, "Incorrect type passed to MaskType_Name()."); return CudnnfMHABackendConfig_MaskType_Name(static_cast(value)); } template <> inline const std::string& CudnnfMHABackendConfig_MaskType_Name(CudnnfMHABackendConfig_MaskType value) { return ::google::protobuf::internal::NameOfDenseEnum( static_cast(value)); } inline bool CudnnfMHABackendConfig_MaskType_Parse(absl::string_view name, CudnnfMHABackendConfig_MaskType* value) { return ::google::protobuf::internal::ParseNamedEnum( CudnnfMHABackendConfig_MaskType_descriptor(), name, value); } // =================================================================== // ------------------------------------------------------------------- class PROTOBUF_EXPORT Tile final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.Tile) */ { public: inline Tile() : Tile(nullptr) {} ~Tile() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR Tile( ::google::protobuf::internal::ConstantInitialized); inline Tile(const Tile& from) : Tile(nullptr, from) {} inline Tile(Tile&& from) noexcept : Tile(nullptr, std::move(from)) {} inline Tile& operator=(const Tile& from) { CopyFrom(from); return *this; } inline Tile& operator=(Tile&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const Tile& default_instance() { return *internal_default_instance(); } static inline const Tile* internal_default_instance() { return reinterpret_cast( &_Tile_default_instance_); } static constexpr int kIndexInFileMessages = 7; friend void swap(Tile& a, Tile& b) { a.Swap(&b); } inline void Swap(Tile* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(Tile* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- Tile* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const Tile& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const Tile& from) { Tile::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(Tile* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.Tile"; } protected: explicit Tile(::google::protobuf::Arena* arena); Tile(::google::protobuf::Arena* arena, const Tile& from); Tile(::google::protobuf::Arena* arena, Tile&& from) noexcept : Tile(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSizesFieldNumber = 1, }; // repeated int64 sizes = 1; int sizes_size() const; private: int _internal_sizes_size() const; public: void clear_sizes() ; ::int64_t sizes(int index) const; void set_sizes(int index, ::int64_t value); void add_sizes(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& sizes() const; ::google::protobuf::RepeatedField<::int64_t>* mutable_sizes(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_sizes() const; ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_sizes(); public: // @@protoc_insertion_point(class_scope:xla.gpu.Tile) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 0, 1, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const Tile& from_msg); ::google::protobuf::RepeatedField<::int64_t> sizes_; mutable ::google::protobuf::internal::CachedSize _sizes_cached_byte_size_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ReificationCost final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ReificationCost) */ { public: inline ReificationCost() : ReificationCost(nullptr) {} ~ReificationCost() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ReificationCost( ::google::protobuf::internal::ConstantInitialized); inline ReificationCost(const ReificationCost& from) : ReificationCost(nullptr, from) {} inline ReificationCost(ReificationCost&& from) noexcept : ReificationCost(nullptr, std::move(from)) {} inline ReificationCost& operator=(const ReificationCost& from) { CopyFrom(from); return *this; } inline ReificationCost& operator=(ReificationCost&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const ReificationCost& default_instance() { return *internal_default_instance(); } static inline const ReificationCost* internal_default_instance() { return reinterpret_cast( &_ReificationCost_default_instance_); } static constexpr int kIndexInFileMessages = 4; friend void swap(ReificationCost& a, ReificationCost& b) { a.Swap(&b); } inline void Swap(ReificationCost* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(ReificationCost* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ReificationCost* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ReificationCost& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ReificationCost& from) { ReificationCost::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(ReificationCost* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ReificationCost"; } protected: explicit ReificationCost(::google::protobuf::Arena* arena); ReificationCost(::google::protobuf::Arena* arena, const ReificationCost& from); ReificationCost(::google::protobuf::Arena* arena, ReificationCost&& from) noexcept : ReificationCost(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kNameFieldNumber = 5, kEndToEndCyclesFieldNumber = 1, kExecTimeUsFieldNumber = 2, kComputeTimeUsFieldNumber = 3, kMemoryAccessTimeUsFieldNumber = 4, }; // string name = 5; void clear_name() ; const std::string& name() const; template void set_name(Arg_&& arg, Args_... args); std::string* mutable_name(); PROTOBUF_NODISCARD std::string* release_name(); void set_allocated_name(std::string* value); private: const std::string& _internal_name() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_name( const std::string& value); std::string* _internal_mutable_name(); public: // double end_to_end_cycles = 1; void clear_end_to_end_cycles() ; double end_to_end_cycles() const; void set_end_to_end_cycles(double value); private: double _internal_end_to_end_cycles() const; void _internal_set_end_to_end_cycles(double value); public: // double exec_time_us = 2; void clear_exec_time_us() ; double exec_time_us() const; void set_exec_time_us(double value); private: double _internal_exec_time_us() const; void _internal_set_exec_time_us(double value); public: // double compute_time_us = 3; void clear_compute_time_us() ; double compute_time_us() const; void set_compute_time_us(double value); private: double _internal_compute_time_us() const; void _internal_set_compute_time_us(double value); public: // double memory_access_time_us = 4; void clear_memory_access_time_us() ; double memory_access_time_us() const; void set_memory_access_time_us(double value); private: double _internal_memory_access_time_us() const; void _internal_set_memory_access_time_us(double value); public: // @@protoc_insertion_point(class_scope:xla.gpu.ReificationCost) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 5, 0, 36, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const ReificationCost& from_msg); ::google::protobuf::internal::ArenaStringPtr name_; double end_to_end_cycles_; double exec_time_us_; double compute_time_us_; double memory_access_time_us_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT DynamicMemcpyConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DynamicMemcpyConfig) */ { public: inline DynamicMemcpyConfig() : DynamicMemcpyConfig(nullptr) {} ~DynamicMemcpyConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR DynamicMemcpyConfig( ::google::protobuf::internal::ConstantInitialized); inline DynamicMemcpyConfig(const DynamicMemcpyConfig& from) : DynamicMemcpyConfig(nullptr, from) {} inline DynamicMemcpyConfig(DynamicMemcpyConfig&& from) noexcept : DynamicMemcpyConfig(nullptr, std::move(from)) {} inline DynamicMemcpyConfig& operator=(const DynamicMemcpyConfig& from) { CopyFrom(from); return *this; } inline DynamicMemcpyConfig& operator=(DynamicMemcpyConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const DynamicMemcpyConfig& default_instance() { return *internal_default_instance(); } static inline const DynamicMemcpyConfig* internal_default_instance() { return reinterpret_cast( &_DynamicMemcpyConfig_default_instance_); } static constexpr int kIndexInFileMessages = 9; friend void swap(DynamicMemcpyConfig& a, DynamicMemcpyConfig& b) { a.Swap(&b); } inline void Swap(DynamicMemcpyConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DynamicMemcpyConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DynamicMemcpyConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DynamicMemcpyConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DynamicMemcpyConfig& from) { DynamicMemcpyConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(DynamicMemcpyConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.DynamicMemcpyConfig"; } protected: explicit DynamicMemcpyConfig(::google::protobuf::Arena* arena); DynamicMemcpyConfig(::google::protobuf::Arena* arena, const DynamicMemcpyConfig& from); DynamicMemcpyConfig(::google::protobuf::Arena* arena, DynamicMemcpyConfig&& from) noexcept : DynamicMemcpyConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSrcOffsetBytesFieldNumber = 2, kDstOffsetBytesFieldNumber = 3, kDependsOnLoopFieldNumber = 1, }; // repeated int64 src_offset_bytes = 2; int src_offset_bytes_size() const; private: int _internal_src_offset_bytes_size() const; public: void clear_src_offset_bytes() ; ::int64_t src_offset_bytes(int index) const; void set_src_offset_bytes(int index, ::int64_t value); void add_src_offset_bytes(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& src_offset_bytes() const; ::google::protobuf::RepeatedField<::int64_t>* mutable_src_offset_bytes(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_src_offset_bytes() const; ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_src_offset_bytes(); public: // repeated int64 dst_offset_bytes = 3; int dst_offset_bytes_size() const; private: int _internal_dst_offset_bytes_size() const; public: void clear_dst_offset_bytes() ; ::int64_t dst_offset_bytes(int index) const; void set_dst_offset_bytes(int index, ::int64_t value); void add_dst_offset_bytes(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& dst_offset_bytes() const; ::google::protobuf::RepeatedField<::int64_t>* mutable_dst_offset_bytes(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_dst_offset_bytes() const; ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_dst_offset_bytes(); public: // bool depends_on_loop = 1; void clear_depends_on_loop() ; bool depends_on_loop() const; void set_depends_on_loop(bool value); private: bool _internal_depends_on_loop() const; void _internal_set_depends_on_loop(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.DynamicMemcpyConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 3, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const DynamicMemcpyConfig& from_msg); ::google::protobuf::RepeatedField<::int64_t> src_offset_bytes_; mutable ::google::protobuf::internal::CachedSize _src_offset_bytes_cached_byte_size_; ::google::protobuf::RepeatedField<::int64_t> dst_offset_bytes_; mutable ::google::protobuf::internal::CachedSize _dst_offset_bytes_cached_byte_size_; bool depends_on_loop_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CustomFusionConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CustomFusionConfig) */ { public: inline CustomFusionConfig() : CustomFusionConfig(nullptr) {} ~CustomFusionConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CustomFusionConfig( ::google::protobuf::internal::ConstantInitialized); inline CustomFusionConfig(const CustomFusionConfig& from) : CustomFusionConfig(nullptr, from) {} inline CustomFusionConfig(CustomFusionConfig&& from) noexcept : CustomFusionConfig(nullptr, std::move(from)) {} inline CustomFusionConfig& operator=(const CustomFusionConfig& from) { CopyFrom(from); return *this; } inline CustomFusionConfig& operator=(CustomFusionConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CustomFusionConfig& default_instance() { return *internal_default_instance(); } static inline const CustomFusionConfig* internal_default_instance() { return reinterpret_cast( &_CustomFusionConfig_default_instance_); } static constexpr int kIndexInFileMessages = 5; friend void swap(CustomFusionConfig& a, CustomFusionConfig& b) { a.Swap(&b); } inline void Swap(CustomFusionConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CustomFusionConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CustomFusionConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CustomFusionConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CustomFusionConfig& from) { CustomFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CustomFusionConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CustomFusionConfig"; } protected: explicit CustomFusionConfig(::google::protobuf::Arena* arena); CustomFusionConfig(::google::protobuf::Arena* arena, const CustomFusionConfig& from); CustomFusionConfig(::google::protobuf::Arena* arena, CustomFusionConfig&& from) noexcept : CustomFusionConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kNameFieldNumber = 1, kKernelIndexFieldNumber = 2, }; // string name = 1; void clear_name() ; const std::string& name() const; template void set_name(Arg_&& arg, Args_... args); std::string* mutable_name(); PROTOBUF_NODISCARD std::string* release_name(); void set_allocated_name(std::string* value); private: const std::string& _internal_name() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_name( const std::string& value); std::string* _internal_mutable_name(); public: // int32 kernel_index = 2; void clear_kernel_index() ; ::int32_t kernel_index() const; void set_kernel_index(::int32_t value); private: ::int32_t _internal_kernel_index() const; void _internal_set_kernel_index(::int32_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CustomFusionConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 1, 2, 0, 39, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CustomFusionConfig& from_msg); ::google::protobuf::internal::ArenaStringPtr name_; ::int32_t kernel_index_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CustomCallBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CustomCallBackendConfig) */ { public: inline CustomCallBackendConfig() : CustomCallBackendConfig(nullptr) {} ~CustomCallBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CustomCallBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline CustomCallBackendConfig(const CustomCallBackendConfig& from) : CustomCallBackendConfig(nullptr, from) {} inline CustomCallBackendConfig(CustomCallBackendConfig&& from) noexcept : CustomCallBackendConfig(nullptr, std::move(from)) {} inline CustomCallBackendConfig& operator=(const CustomCallBackendConfig& from) { CopyFrom(from); return *this; } inline CustomCallBackendConfig& operator=(CustomCallBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CustomCallBackendConfig& default_instance() { return *internal_default_instance(); } enum RawBackendConfigOneofCase { kOpaque = 1, kAttributes = 2, RAW_BACKEND_CONFIG_ONEOF_NOT_SET = 0, }; static inline const CustomCallBackendConfig* internal_default_instance() { return reinterpret_cast( &_CustomCallBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 13; friend void swap(CustomCallBackendConfig& a, CustomCallBackendConfig& b) { a.Swap(&b); } inline void Swap(CustomCallBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CustomCallBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CustomCallBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CustomCallBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CustomCallBackendConfig& from) { CustomCallBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CustomCallBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CustomCallBackendConfig"; } protected: explicit CustomCallBackendConfig(::google::protobuf::Arena* arena); CustomCallBackendConfig(::google::protobuf::Arena* arena, const CustomCallBackendConfig& from); CustomCallBackendConfig(::google::protobuf::Arena* arena, CustomCallBackendConfig&& from) noexcept : CustomCallBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kOpaqueFieldNumber = 1, kAttributesFieldNumber = 2, }; // string opaque = 1; bool has_opaque() const; void clear_opaque() ; const std::string& opaque() const; template void set_opaque(Arg_&& arg, Args_... args); std::string* mutable_opaque(); PROTOBUF_NODISCARD std::string* release_opaque(); void set_allocated_opaque(std::string* value); private: const std::string& _internal_opaque() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_opaque( const std::string& value); std::string* _internal_mutable_opaque(); public: // string attributes = 2; bool has_attributes() const; void clear_attributes() ; const std::string& attributes() const; template void set_attributes(Arg_&& arg, Args_... args); std::string* mutable_attributes(); PROTOBUF_NODISCARD std::string* release_attributes(); void set_allocated_attributes(std::string* value); private: const std::string& _internal_attributes() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_attributes( const std::string& value); std::string* _internal_mutable_attributes(); public: void clear_raw_backend_config_oneof(); RawBackendConfigOneofCase raw_backend_config_oneof_case() const; // @@protoc_insertion_point(class_scope:xla.gpu.CustomCallBackendConfig) private: class _Internal; void set_has_opaque(); void set_has_attributes(); inline bool has_raw_backend_config_oneof() const; inline void clear_has_raw_backend_config_oneof(); friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 0, 2, 0, 56, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CustomCallBackendConfig& from_msg); union RawBackendConfigOneofUnion { constexpr RawBackendConfigOneofUnion() : _constinit_{} {} ::google::protobuf::internal::ConstantInitialized _constinit_; ::google::protobuf::internal::ArenaStringPtr opaque_; ::google::protobuf::internal::ArenaStringPtr attributes_; } raw_backend_config_oneof_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::uint32_t _oneof_case_[1]; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CuDnnFusionConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CuDnnFusionConfig) */ { public: inline CuDnnFusionConfig() : CuDnnFusionConfig(nullptr) {} ~CuDnnFusionConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CuDnnFusionConfig( ::google::protobuf::internal::ConstantInitialized); inline CuDnnFusionConfig(const CuDnnFusionConfig& from) : CuDnnFusionConfig(nullptr, from) {} inline CuDnnFusionConfig(CuDnnFusionConfig&& from) noexcept : CuDnnFusionConfig(nullptr, std::move(from)) {} inline CuDnnFusionConfig& operator=(const CuDnnFusionConfig& from) { CopyFrom(from); return *this; } inline CuDnnFusionConfig& operator=(CuDnnFusionConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CuDnnFusionConfig& default_instance() { return *internal_default_instance(); } static inline const CuDnnFusionConfig* internal_default_instance() { return reinterpret_cast( &_CuDnnFusionConfig_default_instance_); } static constexpr int kIndexInFileMessages = 6; friend void swap(CuDnnFusionConfig& a, CuDnnFusionConfig& b) { a.Swap(&b); } inline void Swap(CuDnnFusionConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CuDnnFusionConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CuDnnFusionConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CuDnnFusionConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CuDnnFusionConfig& from) { CuDnnFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CuDnnFusionConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CuDnnFusionConfig"; } protected: explicit CuDnnFusionConfig(::google::protobuf::Arena* arena); CuDnnFusionConfig(::google::protobuf::Arena* arena, const CuDnnFusionConfig& from); CuDnnFusionConfig(::google::protobuf::Arena* arena, CuDnnFusionConfig&& from) noexcept : CuDnnFusionConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kPlanIdFieldNumber = 1, }; // int64 plan_id = 1; void clear_plan_id() ; ::int64_t plan_id() const; void set_plan_id(::int64_t value); private: ::int64_t _internal_plan_id() const; void _internal_set_plan_id(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CuDnnFusionConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 0, 1, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CuDnnFusionConfig& from_msg); ::int64_t plan_id_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CollectiveBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CollectiveBackendConfig) */ { public: inline CollectiveBackendConfig() : CollectiveBackendConfig(nullptr) {} ~CollectiveBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CollectiveBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline CollectiveBackendConfig(const CollectiveBackendConfig& from) : CollectiveBackendConfig(nullptr, from) {} inline CollectiveBackendConfig(CollectiveBackendConfig&& from) noexcept : CollectiveBackendConfig(nullptr, std::move(from)) {} inline CollectiveBackendConfig& operator=(const CollectiveBackendConfig& from) { CopyFrom(from); return *this; } inline CollectiveBackendConfig& operator=(CollectiveBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CollectiveBackendConfig& default_instance() { return *internal_default_instance(); } static inline const CollectiveBackendConfig* internal_default_instance() { return reinterpret_cast( &_CollectiveBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 3; friend void swap(CollectiveBackendConfig& a, CollectiveBackendConfig& b) { a.Swap(&b); } inline void Swap(CollectiveBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CollectiveBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CollectiveBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CollectiveBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CollectiveBackendConfig& from) { CollectiveBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CollectiveBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CollectiveBackendConfig"; } protected: explicit CollectiveBackendConfig(::google::protobuf::Arena* arena); CollectiveBackendConfig(::google::protobuf::Arena* arena, const CollectiveBackendConfig& from); CollectiveBackendConfig(::google::protobuf::Arena* arena, CollectiveBackendConfig&& from) noexcept : CollectiveBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- using CollectiveBackend = CollectiveBackendConfig_CollectiveBackend; static constexpr CollectiveBackend DEFAULT = CollectiveBackendConfig_CollectiveBackend_DEFAULT; static constexpr CollectiveBackend NVSHMEM = CollectiveBackendConfig_CollectiveBackend_NVSHMEM; static inline bool CollectiveBackend_IsValid(int value) { return CollectiveBackendConfig_CollectiveBackend_IsValid(value); } static constexpr CollectiveBackend CollectiveBackend_MIN = CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_MIN; static constexpr CollectiveBackend CollectiveBackend_MAX = CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_MAX; static constexpr int CollectiveBackend_ARRAYSIZE = CollectiveBackendConfig_CollectiveBackend_CollectiveBackend_ARRAYSIZE; static inline const ::google::protobuf::EnumDescriptor* CollectiveBackend_descriptor() { return CollectiveBackendConfig_CollectiveBackend_descriptor(); } template static inline const std::string& CollectiveBackend_Name(T value) { return CollectiveBackendConfig_CollectiveBackend_Name(value); } static inline bool CollectiveBackend_Parse(absl::string_view name, CollectiveBackend* value) { return CollectiveBackendConfig_CollectiveBackend_Parse(name, value); } // accessors ------------------------------------------------------- enum : int { kIsSyncFieldNumber = 1, kIsPipelinedFieldNumber = 3, kBackendFieldNumber = 5, }; // bool is_sync = 1; void clear_is_sync() ; bool is_sync() const; void set_is_sync(bool value); private: bool _internal_is_sync() const; void _internal_set_is_sync(bool value); public: // bool is_pipelined = 3; void clear_is_pipelined() ; bool is_pipelined() const; void set_is_pipelined(bool value); private: bool _internal_is_pipelined() const; void _internal_set_is_pipelined(bool value); public: // .xla.gpu.CollectiveBackendConfig.CollectiveBackend backend = 5; void clear_backend() ; ::xla::gpu::CollectiveBackendConfig_CollectiveBackend backend() const; void set_backend(::xla::gpu::CollectiveBackendConfig_CollectiveBackend value); private: ::xla::gpu::CollectiveBackendConfig_CollectiveBackend _internal_backend() const; void _internal_set_backend(::xla::gpu::CollectiveBackendConfig_CollectiveBackend value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CollectiveBackendConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 3, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CollectiveBackendConfig& from_msg); bool is_sync_; bool is_pipelined_; int backend_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT GemmBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.GemmBackendConfig) */ { public: inline GemmBackendConfig() : GemmBackendConfig(nullptr) {} ~GemmBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR GemmBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline GemmBackendConfig(const GemmBackendConfig& from) : GemmBackendConfig(nullptr, from) {} inline GemmBackendConfig(GemmBackendConfig&& from) noexcept : GemmBackendConfig(nullptr, std::move(from)) {} inline GemmBackendConfig& operator=(const GemmBackendConfig& from) { CopyFrom(from); return *this; } inline GemmBackendConfig& operator=(GemmBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const GemmBackendConfig& default_instance() { return *internal_default_instance(); } enum AlgorithmCase { kSelectedAlgorithm = 1, ALGORITHM_NOT_SET = 0, }; static inline const GemmBackendConfig* internal_default_instance() { return reinterpret_cast( &_GemmBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 1; friend void swap(GemmBackendConfig& a, GemmBackendConfig& b) { a.Swap(&b); } inline void Swap(GemmBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(GemmBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- GemmBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const GemmBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const GemmBackendConfig& from) { GemmBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(GemmBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.GemmBackendConfig"; } protected: explicit GemmBackendConfig(::google::protobuf::Arena* arena); GemmBackendConfig(::google::protobuf::Arena* arena, const GemmBackendConfig& from); GemmBackendConfig(::google::protobuf::Arena* arena, GemmBackendConfig&& from) noexcept : GemmBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- using Epilogue = GemmBackendConfig_Epilogue; static constexpr Epilogue DEFAULT = GemmBackendConfig_Epilogue_DEFAULT; static constexpr Epilogue BIAS = GemmBackendConfig_Epilogue_BIAS; static constexpr Epilogue RELU = GemmBackendConfig_Epilogue_RELU; static constexpr Epilogue BIAS_RELU = GemmBackendConfig_Epilogue_BIAS_RELU; static constexpr Epilogue GELU = GemmBackendConfig_Epilogue_GELU; static constexpr Epilogue GELU_AUX = GemmBackendConfig_Epilogue_GELU_AUX; static constexpr Epilogue BIAS_GELU = GemmBackendConfig_Epilogue_BIAS_GELU; static constexpr Epilogue BIAS_GELU_AUX = GemmBackendConfig_Epilogue_BIAS_GELU_AUX; static constexpr Epilogue SILU = GemmBackendConfig_Epilogue_SILU; static constexpr Epilogue BIAS_SILU = GemmBackendConfig_Epilogue_BIAS_SILU; static inline bool Epilogue_IsValid(int value) { return GemmBackendConfig_Epilogue_IsValid(value); } static constexpr Epilogue Epilogue_MIN = GemmBackendConfig_Epilogue_Epilogue_MIN; static constexpr Epilogue Epilogue_MAX = GemmBackendConfig_Epilogue_Epilogue_MAX; static constexpr int Epilogue_ARRAYSIZE = GemmBackendConfig_Epilogue_Epilogue_ARRAYSIZE; static inline const ::google::protobuf::EnumDescriptor* Epilogue_descriptor() { return GemmBackendConfig_Epilogue_descriptor(); } template static inline const std::string& Epilogue_Name(T value) { return GemmBackendConfig_Epilogue_Name(value); } static inline bool Epilogue_Parse(absl::string_view name, Epilogue* value) { return GemmBackendConfig_Epilogue_Parse(name, value); } // accessors ------------------------------------------------------- enum : int { kDotDimensionNumbersFieldNumber = 7, kPrecisionConfigFieldNumber = 12, kAlphaRealFieldNumber = 2, kBetaFieldNumber = 3, kAlphaImagFieldNumber = 9, kLhsStrideFieldNumber = 14, kRhsStrideFieldNumber = 15, kEpilogueFieldNumber = 13, kGradXFieldNumber = 16, kGradYFieldNumber = 17, kDamaxOutputFieldNumber = 18, kSelectedAlgorithmFieldNumber = 1, }; // .xla.DotDimensionNumbers dot_dimension_numbers = 7; bool has_dot_dimension_numbers() const; void clear_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_dot_dimension_numbers(); void set_allocated_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_dot_dimension_numbers(); public: // .xla.PrecisionConfig precision_config = 12; bool has_precision_config() const; void clear_precision_config() ; const ::xla::PrecisionConfig& precision_config() const; PROTOBUF_NODISCARD ::xla::PrecisionConfig* release_precision_config(); ::xla::PrecisionConfig* mutable_precision_config(); void set_allocated_precision_config(::xla::PrecisionConfig* value); void unsafe_arena_set_allocated_precision_config(::xla::PrecisionConfig* value); ::xla::PrecisionConfig* unsafe_arena_release_precision_config(); private: const ::xla::PrecisionConfig& _internal_precision_config() const; ::xla::PrecisionConfig* _internal_mutable_precision_config(); public: // double alpha_real = 2; void clear_alpha_real() ; double alpha_real() const; void set_alpha_real(double value); private: double _internal_alpha_real() const; void _internal_set_alpha_real(double value); public: // double beta = 3; void clear_beta() ; double beta() const; void set_beta(double value); private: double _internal_beta() const; void _internal_set_beta(double value); public: // double alpha_imag = 9; void clear_alpha_imag() ; double alpha_imag() const; void set_alpha_imag(double value); private: double _internal_alpha_imag() const; void _internal_set_alpha_imag(double value); public: // optional int64 lhs_stride = 14; bool has_lhs_stride() const; void clear_lhs_stride() ; ::int64_t lhs_stride() const; void set_lhs_stride(::int64_t value); private: ::int64_t _internal_lhs_stride() const; void _internal_set_lhs_stride(::int64_t value); public: // optional int64 rhs_stride = 15; bool has_rhs_stride() const; void clear_rhs_stride() ; ::int64_t rhs_stride() const; void set_rhs_stride(::int64_t value); private: ::int64_t _internal_rhs_stride() const; void _internal_set_rhs_stride(::int64_t value); public: // .xla.gpu.GemmBackendConfig.Epilogue epilogue = 13; void clear_epilogue() ; ::xla::gpu::GemmBackendConfig_Epilogue epilogue() const; void set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value); private: ::xla::gpu::GemmBackendConfig_Epilogue _internal_epilogue() const; void _internal_set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value); public: // optional bool grad_x = 16; bool has_grad_x() const; void clear_grad_x() ; bool grad_x() const; void set_grad_x(bool value); private: bool _internal_grad_x() const; void _internal_set_grad_x(bool value); public: // optional bool grad_y = 17; bool has_grad_y() const; void clear_grad_y() ; bool grad_y() const; void set_grad_y(bool value); private: bool _internal_grad_y() const; void _internal_set_grad_y(bool value); public: // bool damax_output = 18; void clear_damax_output() ; bool damax_output() const; void set_damax_output(bool value); private: bool _internal_damax_output() const; void _internal_set_damax_output(bool value); public: // int64 selected_algorithm = 1; bool has_selected_algorithm() const; void clear_selected_algorithm() ; ::int64_t selected_algorithm() const; void set_selected_algorithm(::int64_t value); private: ::int64_t _internal_selected_algorithm() const; void _internal_set_selected_algorithm(::int64_t value); public: void clear_algorithm(); AlgorithmCase algorithm_case() const; // @@protoc_insertion_point(class_scope:xla.gpu.GemmBackendConfig) private: class _Internal; void set_has_selected_algorithm(); inline bool has_algorithm() const; inline void clear_has_algorithm(); friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 4, 12, 2, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const GemmBackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::DotDimensionNumbers* dot_dimension_numbers_; ::xla::PrecisionConfig* precision_config_; double alpha_real_; double beta_; double alpha_imag_; ::int64_t lhs_stride_; ::int64_t rhs_stride_; int epilogue_; bool grad_x_; bool grad_y_; bool damax_output_; union AlgorithmUnion { constexpr AlgorithmUnion() : _constinit_{} {} ::google::protobuf::internal::ConstantInitialized _constinit_; ::int64_t selected_algorithm_; } algorithm_; ::uint32_t _oneof_case_[1]; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT BlockLevelFusionConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.BlockLevelFusionConfig) */ { public: inline BlockLevelFusionConfig() : BlockLevelFusionConfig(nullptr) {} ~BlockLevelFusionConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR BlockLevelFusionConfig( ::google::protobuf::internal::ConstantInitialized); inline BlockLevelFusionConfig(const BlockLevelFusionConfig& from) : BlockLevelFusionConfig(nullptr, from) {} inline BlockLevelFusionConfig(BlockLevelFusionConfig&& from) noexcept : BlockLevelFusionConfig(nullptr, std::move(from)) {} inline BlockLevelFusionConfig& operator=(const BlockLevelFusionConfig& from) { CopyFrom(from); return *this; } inline BlockLevelFusionConfig& operator=(BlockLevelFusionConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const BlockLevelFusionConfig& default_instance() { return *internal_default_instance(); } static inline const BlockLevelFusionConfig* internal_default_instance() { return reinterpret_cast( &_BlockLevelFusionConfig_default_instance_); } static constexpr int kIndexInFileMessages = 8; friend void swap(BlockLevelFusionConfig& a, BlockLevelFusionConfig& b) { a.Swap(&b); } inline void Swap(BlockLevelFusionConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(BlockLevelFusionConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- BlockLevelFusionConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const BlockLevelFusionConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const BlockLevelFusionConfig& from) { BlockLevelFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(BlockLevelFusionConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.BlockLevelFusionConfig"; } protected: explicit BlockLevelFusionConfig(::google::protobuf::Arena* arena); BlockLevelFusionConfig(::google::protobuf::Arena* arena, const BlockLevelFusionConfig& from); BlockLevelFusionConfig(::google::protobuf::Arena* arena, BlockLevelFusionConfig&& from) noexcept : BlockLevelFusionConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kOutputTilesFieldNumber = 3, kNumWarpsFieldNumber = 2, kNumCtasFieldNumber = 4, kNumStagesFieldNumber = 5, kIsTmaAllowedFieldNumber = 6, }; // repeated .xla.gpu.Tile output_tiles = 3; int output_tiles_size() const; private: int _internal_output_tiles_size() const; public: void clear_output_tiles() ; ::xla::gpu::Tile* mutable_output_tiles(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>* mutable_output_tiles(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>& _internal_output_tiles() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>* _internal_mutable_output_tiles(); public: const ::xla::gpu::Tile& output_tiles(int index) const; ::xla::gpu::Tile* add_output_tiles(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>& output_tiles() const; // int64 num_warps = 2; void clear_num_warps() ; ::int64_t num_warps() const; void set_num_warps(::int64_t value); private: ::int64_t _internal_num_warps() const; void _internal_set_num_warps(::int64_t value); public: // int32 num_ctas = 4; void clear_num_ctas() ; ::int32_t num_ctas() const; void set_num_ctas(::int32_t value); private: ::int32_t _internal_num_ctas() const; void _internal_set_num_ctas(::int32_t value); public: // int32 num_stages = 5; void clear_num_stages() ; ::int32_t num_stages() const; void set_num_stages(::int32_t value); private: ::int32_t _internal_num_stages() const; void _internal_set_num_stages(::int32_t value); public: // bool is_tma_allowed = 6; void clear_is_tma_allowed() ; bool is_tma_allowed() const; void set_is_tma_allowed(bool value); private: bool _internal_is_tma_allowed() const; void _internal_set_is_tma_allowed(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.BlockLevelFusionConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 5, 1, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const BlockLevelFusionConfig& from_msg); ::google::protobuf::RepeatedPtrField< ::xla::gpu::Tile > output_tiles_; ::int64_t num_warps_; ::int32_t num_ctas_; ::int32_t num_stages_; bool is_tma_allowed_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT FusionBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.FusionBackendConfig) */ { public: inline FusionBackendConfig() : FusionBackendConfig(nullptr) {} ~FusionBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR FusionBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline FusionBackendConfig(const FusionBackendConfig& from) : FusionBackendConfig(nullptr, from) {} inline FusionBackendConfig(FusionBackendConfig&& from) noexcept : FusionBackendConfig(nullptr, std::move(from)) {} inline FusionBackendConfig& operator=(const FusionBackendConfig& from) { CopyFrom(from); return *this; } inline FusionBackendConfig& operator=(FusionBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const FusionBackendConfig& default_instance() { return *internal_default_instance(); } static inline const FusionBackendConfig* internal_default_instance() { return reinterpret_cast( &_FusionBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 10; friend void swap(FusionBackendConfig& a, FusionBackendConfig& b) { a.Swap(&b); } inline void Swap(FusionBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(FusionBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- FusionBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const FusionBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const FusionBackendConfig& from) { FusionBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(FusionBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.FusionBackendConfig"; } protected: explicit FusionBackendConfig(::google::protobuf::Arena* arena); FusionBackendConfig(::google::protobuf::Arena* arena, const FusionBackendConfig& from); FusionBackendConfig(::google::protobuf::Arena* arena, FusionBackendConfig&& from) noexcept : FusionBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kKindFieldNumber = 1, kTritonGemmConfigFieldNumber = 2, kCustomFusionConfigFieldNumber = 4, kCudnnFusionConfigFieldNumber = 5, kBlockLevelFusionConfigFieldNumber = 6, kDynamicMemcpyConfigFieldNumber = 7, }; // string kind = 1; void clear_kind() ; const std::string& kind() const; template void set_kind(Arg_&& arg, Args_... args); std::string* mutable_kind(); PROTOBUF_NODISCARD std::string* release_kind(); void set_allocated_kind(std::string* value); 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.AutotuneResult.TritonGemmKey triton_gemm_config = 2; bool has_triton_gemm_config() const; void clear_triton_gemm_config() ; const ::xla::AutotuneResult_TritonGemmKey& triton_gemm_config() const; PROTOBUF_NODISCARD ::xla::AutotuneResult_TritonGemmKey* release_triton_gemm_config(); ::xla::AutotuneResult_TritonGemmKey* mutable_triton_gemm_config(); void set_allocated_triton_gemm_config(::xla::AutotuneResult_TritonGemmKey* value); void unsafe_arena_set_allocated_triton_gemm_config(::xla::AutotuneResult_TritonGemmKey* value); ::xla::AutotuneResult_TritonGemmKey* unsafe_arena_release_triton_gemm_config(); private: const ::xla::AutotuneResult_TritonGemmKey& _internal_triton_gemm_config() const; ::xla::AutotuneResult_TritonGemmKey* _internal_mutable_triton_gemm_config(); public: // .xla.gpu.CustomFusionConfig custom_fusion_config = 4; bool has_custom_fusion_config() const; void clear_custom_fusion_config() ; const ::xla::gpu::CustomFusionConfig& custom_fusion_config() const; PROTOBUF_NODISCARD ::xla::gpu::CustomFusionConfig* release_custom_fusion_config(); ::xla::gpu::CustomFusionConfig* mutable_custom_fusion_config(); void set_allocated_custom_fusion_config(::xla::gpu::CustomFusionConfig* value); void unsafe_arena_set_allocated_custom_fusion_config(::xla::gpu::CustomFusionConfig* value); ::xla::gpu::CustomFusionConfig* unsafe_arena_release_custom_fusion_config(); private: const ::xla::gpu::CustomFusionConfig& _internal_custom_fusion_config() const; ::xla::gpu::CustomFusionConfig* _internal_mutable_custom_fusion_config(); public: // .xla.gpu.CuDnnFusionConfig cudnn_fusion_config = 5; bool has_cudnn_fusion_config() const; void clear_cudnn_fusion_config() ; const ::xla::gpu::CuDnnFusionConfig& cudnn_fusion_config() const; PROTOBUF_NODISCARD ::xla::gpu::CuDnnFusionConfig* release_cudnn_fusion_config(); ::xla::gpu::CuDnnFusionConfig* mutable_cudnn_fusion_config(); void set_allocated_cudnn_fusion_config(::xla::gpu::CuDnnFusionConfig* value); void unsafe_arena_set_allocated_cudnn_fusion_config(::xla::gpu::CuDnnFusionConfig* value); ::xla::gpu::CuDnnFusionConfig* unsafe_arena_release_cudnn_fusion_config(); private: const ::xla::gpu::CuDnnFusionConfig& _internal_cudnn_fusion_config() const; ::xla::gpu::CuDnnFusionConfig* _internal_mutable_cudnn_fusion_config(); public: // .xla.gpu.BlockLevelFusionConfig block_level_fusion_config = 6; bool has_block_level_fusion_config() const; void clear_block_level_fusion_config() ; const ::xla::gpu::BlockLevelFusionConfig& block_level_fusion_config() const; PROTOBUF_NODISCARD ::xla::gpu::BlockLevelFusionConfig* release_block_level_fusion_config(); ::xla::gpu::BlockLevelFusionConfig* mutable_block_level_fusion_config(); void set_allocated_block_level_fusion_config(::xla::gpu::BlockLevelFusionConfig* value); void unsafe_arena_set_allocated_block_level_fusion_config(::xla::gpu::BlockLevelFusionConfig* value); ::xla::gpu::BlockLevelFusionConfig* unsafe_arena_release_block_level_fusion_config(); private: const ::xla::gpu::BlockLevelFusionConfig& _internal_block_level_fusion_config() const; ::xla::gpu::BlockLevelFusionConfig* _internal_mutable_block_level_fusion_config(); public: // .xla.gpu.DynamicMemcpyConfig dynamic_memcpy_config = 7; bool has_dynamic_memcpy_config() const; void clear_dynamic_memcpy_config() ; const ::xla::gpu::DynamicMemcpyConfig& dynamic_memcpy_config() const; PROTOBUF_NODISCARD ::xla::gpu::DynamicMemcpyConfig* release_dynamic_memcpy_config(); ::xla::gpu::DynamicMemcpyConfig* mutable_dynamic_memcpy_config(); void set_allocated_dynamic_memcpy_config(::xla::gpu::DynamicMemcpyConfig* value); void unsafe_arena_set_allocated_dynamic_memcpy_config(::xla::gpu::DynamicMemcpyConfig* value); ::xla::gpu::DynamicMemcpyConfig* unsafe_arena_release_dynamic_memcpy_config(); private: const ::xla::gpu::DynamicMemcpyConfig& _internal_dynamic_memcpy_config() const; ::xla::gpu::DynamicMemcpyConfig* _internal_mutable_dynamic_memcpy_config(); public: // @@protoc_insertion_point(class_scope:xla.gpu.FusionBackendConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 6, 5, 40, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const FusionBackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::internal::ArenaStringPtr kind_; ::xla::AutotuneResult_TritonGemmKey* triton_gemm_config_; ::xla::gpu::CustomFusionConfig* custom_fusion_config_; ::xla::gpu::CuDnnFusionConfig* cudnn_fusion_config_; ::xla::gpu::BlockLevelFusionConfig* block_level_fusion_config_; ::xla::gpu::DynamicMemcpyConfig* dynamic_memcpy_config_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CudnnfMHABackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnfMHABackendConfig) */ { public: inline CudnnfMHABackendConfig() : CudnnfMHABackendConfig(nullptr) {} ~CudnnfMHABackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CudnnfMHABackendConfig( ::google::protobuf::internal::ConstantInitialized); inline CudnnfMHABackendConfig(const CudnnfMHABackendConfig& from) : CudnnfMHABackendConfig(nullptr, from) {} inline CudnnfMHABackendConfig(CudnnfMHABackendConfig&& from) noexcept : CudnnfMHABackendConfig(nullptr, std::move(from)) {} inline CudnnfMHABackendConfig& operator=(const CudnnfMHABackendConfig& from) { CopyFrom(from); return *this; } inline CudnnfMHABackendConfig& operator=(CudnnfMHABackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CudnnfMHABackendConfig& default_instance() { return *internal_default_instance(); } static inline const CudnnfMHABackendConfig* internal_default_instance() { return reinterpret_cast( &_CudnnfMHABackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 12; friend void swap(CudnnfMHABackendConfig& a, CudnnfMHABackendConfig& b) { a.Swap(&b); } inline void Swap(CudnnfMHABackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnfMHABackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnfMHABackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CudnnfMHABackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CudnnfMHABackendConfig& from) { CudnnfMHABackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CudnnfMHABackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CudnnfMHABackendConfig"; } protected: explicit CudnnfMHABackendConfig(::google::protobuf::Arena* arena); CudnnfMHABackendConfig(::google::protobuf::Arena* arena, const CudnnfMHABackendConfig& from); CudnnfMHABackendConfig(::google::protobuf::Arena* arena, CudnnfMHABackendConfig&& from) noexcept : CudnnfMHABackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- using MaskType = CudnnfMHABackendConfig_MaskType; static constexpr MaskType NO_MASK = CudnnfMHABackendConfig_MaskType_NO_MASK; static constexpr MaskType PADDING = CudnnfMHABackendConfig_MaskType_PADDING; static constexpr MaskType CAUSAL = CudnnfMHABackendConfig_MaskType_CAUSAL; static constexpr MaskType PADDING_CAUSAL = CudnnfMHABackendConfig_MaskType_PADDING_CAUSAL; static constexpr MaskType ALIBI = CudnnfMHABackendConfig_MaskType_ALIBI; static inline bool MaskType_IsValid(int value) { return CudnnfMHABackendConfig_MaskType_IsValid(value); } static constexpr MaskType MaskType_MIN = CudnnfMHABackendConfig_MaskType_MaskType_MIN; static constexpr MaskType MaskType_MAX = CudnnfMHABackendConfig_MaskType_MaskType_MAX; static constexpr int MaskType_ARRAYSIZE = CudnnfMHABackendConfig_MaskType_MaskType_ARRAYSIZE; static inline const ::google::protobuf::EnumDescriptor* MaskType_descriptor() { return CudnnfMHABackendConfig_MaskType_descriptor(); } template static inline const std::string& MaskType_Name(T value) { return CudnnfMHABackendConfig_MaskType_Name(value); } static inline bool MaskType_Parse(absl::string_view name, MaskType* value) { return CudnnfMHABackendConfig_MaskType_Parse(name, value); } // accessors ------------------------------------------------------- enum : int { kAlgorithmFieldNumber = 8, kBmm1DotDimensionNumbersFieldNumber = 11, kBmm2DotDimensionNumbersFieldNumber = 12, kIntermediateTensorShapeFieldNumber = 14, kBmm1GradGemm1DotDimensionNumbersFieldNumber = 16, kBmm1GradGemm2DotDimensionNumbersFieldNumber = 17, kBmm2GradGemm1DotDimensionNumbersFieldNumber = 18, kBmm2GradGemm2DotDimensionNumbersFieldNumber = 19, kFmhaScaleFieldNumber = 10, kDropoutRateFieldNumber = 13, kSeedFieldNumber = 15, kMaskTypeFieldNumber = 22, kIsFlashAttentionFieldNumber = 20, kIsCausalMaskFieldNumber = 21, kForceDeterministicFieldNumber = 23, kIsPagedAttentionFieldNumber = 26, kSlidingWindowLengthFieldNumber = 24, kMaxSegPerBatchFieldNumber = 25, }; // .stream_executor.dnn.AlgorithmProto algorithm = 8; bool has_algorithm() const; void clear_algorithm() ; const ::stream_executor::dnn::AlgorithmProto& algorithm() const; PROTOBUF_NODISCARD ::stream_executor::dnn::AlgorithmProto* release_algorithm(); ::stream_executor::dnn::AlgorithmProto* mutable_algorithm(); void set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); void unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: // .xla.DotDimensionNumbers bmm1_dot_dimension_numbers = 11; bool has_bmm1_dot_dimension_numbers() const; void clear_bmm1_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm1_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm1_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm1_dot_dimension_numbers(); void set_allocated_bmm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm1_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm1_dot_dimension_numbers(); public: // .xla.DotDimensionNumbers bmm2_dot_dimension_numbers = 12; bool has_bmm2_dot_dimension_numbers() const; void clear_bmm2_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm2_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm2_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm2_dot_dimension_numbers(); void set_allocated_bmm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm2_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm2_dot_dimension_numbers(); public: // .xla.ShapeProto intermediate_tensor_shape = 14; bool has_intermediate_tensor_shape() const; void clear_intermediate_tensor_shape() ; const ::xla::ShapeProto& intermediate_tensor_shape() const; PROTOBUF_NODISCARD ::xla::ShapeProto* release_intermediate_tensor_shape(); ::xla::ShapeProto* mutable_intermediate_tensor_shape(); void set_allocated_intermediate_tensor_shape(::xla::ShapeProto* value); void unsafe_arena_set_allocated_intermediate_tensor_shape(::xla::ShapeProto* value); ::xla::ShapeProto* unsafe_arena_release_intermediate_tensor_shape(); private: const ::xla::ShapeProto& _internal_intermediate_tensor_shape() const; ::xla::ShapeProto* _internal_mutable_intermediate_tensor_shape(); public: // .xla.DotDimensionNumbers bmm1_grad_gemm1_dot_dimension_numbers = 16; bool has_bmm1_grad_gemm1_dot_dimension_numbers() const; void clear_bmm1_grad_gemm1_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm1_grad_gemm1_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm1_grad_gemm1_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm1_grad_gemm1_dot_dimension_numbers(); void set_allocated_bmm1_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm1_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_grad_gemm1_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm1_grad_gemm1_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm1_grad_gemm1_dot_dimension_numbers(); public: // .xla.DotDimensionNumbers bmm1_grad_gemm2_dot_dimension_numbers = 17; bool has_bmm1_grad_gemm2_dot_dimension_numbers() const; void clear_bmm1_grad_gemm2_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm1_grad_gemm2_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm1_grad_gemm2_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm1_grad_gemm2_dot_dimension_numbers(); void set_allocated_bmm1_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm1_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_grad_gemm2_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm1_grad_gemm2_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm1_grad_gemm2_dot_dimension_numbers(); public: // .xla.DotDimensionNumbers bmm2_grad_gemm1_dot_dimension_numbers = 18; bool has_bmm2_grad_gemm1_dot_dimension_numbers() const; void clear_bmm2_grad_gemm1_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm2_grad_gemm1_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm2_grad_gemm1_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm2_grad_gemm1_dot_dimension_numbers(); void set_allocated_bmm2_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm2_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_grad_gemm1_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm2_grad_gemm1_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm2_grad_gemm1_dot_dimension_numbers(); public: // .xla.DotDimensionNumbers bmm2_grad_gemm2_dot_dimension_numbers = 19; bool has_bmm2_grad_gemm2_dot_dimension_numbers() const; void clear_bmm2_grad_gemm2_dot_dimension_numbers() ; const ::xla::DotDimensionNumbers& bmm2_grad_gemm2_dot_dimension_numbers() const; PROTOBUF_NODISCARD ::xla::DotDimensionNumbers* release_bmm2_grad_gemm2_dot_dimension_numbers(); ::xla::DotDimensionNumbers* mutable_bmm2_grad_gemm2_dot_dimension_numbers(); void set_allocated_bmm2_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); void unsafe_arena_set_allocated_bmm2_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_grad_gemm2_dot_dimension_numbers(); private: const ::xla::DotDimensionNumbers& _internal_bmm2_grad_gemm2_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm2_grad_gemm2_dot_dimension_numbers(); public: // double fmha_scale = 10; void clear_fmha_scale() ; double fmha_scale() const; void set_fmha_scale(double value); private: double _internal_fmha_scale() const; void _internal_set_fmha_scale(double value); public: // double dropout_rate = 13; void clear_dropout_rate() ; double dropout_rate() const; void set_dropout_rate(double value); private: double _internal_dropout_rate() const; void _internal_set_dropout_rate(double value); public: // int64 seed = 15; void clear_seed() ; ::int64_t seed() const; void set_seed(::int64_t value); private: ::int64_t _internal_seed() const; void _internal_set_seed(::int64_t value); public: // .xla.gpu.CudnnfMHABackendConfig.MaskType mask_type = 22; void clear_mask_type() ; ::xla::gpu::CudnnfMHABackendConfig_MaskType mask_type() const; void set_mask_type(::xla::gpu::CudnnfMHABackendConfig_MaskType value); private: ::xla::gpu::CudnnfMHABackendConfig_MaskType _internal_mask_type() const; void _internal_set_mask_type(::xla::gpu::CudnnfMHABackendConfig_MaskType value); public: // bool is_flash_attention = 20; void clear_is_flash_attention() ; bool is_flash_attention() const; void set_is_flash_attention(bool value); private: bool _internal_is_flash_attention() const; void _internal_set_is_flash_attention(bool value); public: // bool is_causal_mask = 21; void clear_is_causal_mask() ; bool is_causal_mask() const; void set_is_causal_mask(bool value); private: bool _internal_is_causal_mask() const; void _internal_set_is_causal_mask(bool value); public: // bool force_deterministic = 23; void clear_force_deterministic() ; bool force_deterministic() const; void set_force_deterministic(bool value); private: bool _internal_force_deterministic() const; void _internal_set_force_deterministic(bool value); public: // bool is_paged_attention = 26; void clear_is_paged_attention() ; bool is_paged_attention() const; void set_is_paged_attention(bool value); private: bool _internal_is_paged_attention() const; void _internal_set_is_paged_attention(bool value); public: // int32 sliding_window_length = 24; void clear_sliding_window_length() ; ::int32_t sliding_window_length() const; void set_sliding_window_length(::int32_t value); private: ::int32_t _internal_sliding_window_length() const; void _internal_set_sliding_window_length(::int32_t value); public: // int32 max_seg_per_batch = 25; void clear_max_seg_per_batch() ; ::int32_t max_seg_per_batch() const; void set_max_seg_per_batch(::int32_t value); private: ::int32_t _internal_max_seg_per_batch() const; void _internal_set_max_seg_per_batch(::int32_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CudnnfMHABackendConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 5, 18, 8, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CudnnfMHABackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::stream_executor::dnn::AlgorithmProto* algorithm_; ::xla::DotDimensionNumbers* bmm1_dot_dimension_numbers_; ::xla::DotDimensionNumbers* bmm2_dot_dimension_numbers_; ::xla::ShapeProto* intermediate_tensor_shape_; ::xla::DotDimensionNumbers* bmm1_grad_gemm1_dot_dimension_numbers_; ::xla::DotDimensionNumbers* bmm1_grad_gemm2_dot_dimension_numbers_; ::xla::DotDimensionNumbers* bmm2_grad_gemm1_dot_dimension_numbers_; ::xla::DotDimensionNumbers* bmm2_grad_gemm2_dot_dimension_numbers_; double fmha_scale_; double dropout_rate_; ::int64_t seed_; int mask_type_; bool is_flash_attention_; bool is_causal_mask_; bool force_deterministic_; bool is_paged_attention_; ::int32_t sliding_window_length_; ::int32_t max_seg_per_batch_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CudnnNormBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnNormBackendConfig) */ { public: inline CudnnNormBackendConfig() : CudnnNormBackendConfig(nullptr) {} ~CudnnNormBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CudnnNormBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline CudnnNormBackendConfig(const CudnnNormBackendConfig& from) : CudnnNormBackendConfig(nullptr, from) {} inline CudnnNormBackendConfig(CudnnNormBackendConfig&& from) noexcept : CudnnNormBackendConfig(nullptr, std::move(from)) {} inline CudnnNormBackendConfig& operator=(const CudnnNormBackendConfig& from) { CopyFrom(from); return *this; } inline CudnnNormBackendConfig& operator=(CudnnNormBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CudnnNormBackendConfig& default_instance() { return *internal_default_instance(); } static inline const CudnnNormBackendConfig* internal_default_instance() { return reinterpret_cast( &_CudnnNormBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 11; friend void swap(CudnnNormBackendConfig& a, CudnnNormBackendConfig& b) { a.Swap(&b); } inline void Swap(CudnnNormBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnNormBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnNormBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CudnnNormBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CudnnNormBackendConfig& from) { CudnnNormBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CudnnNormBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CudnnNormBackendConfig"; } protected: explicit CudnnNormBackendConfig(::google::protobuf::Arena* arena); CudnnNormBackendConfig(::google::protobuf::Arena* arena, const CudnnNormBackendConfig& from); CudnnNormBackendConfig(::google::protobuf::Arena* arena, CudnnNormBackendConfig&& from) noexcept : CudnnNormBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- using Kind = CudnnNormBackendConfig_Kind; static constexpr Kind LAYER_FWD_INFER = CudnnNormBackendConfig_Kind_LAYER_FWD_INFER; static constexpr Kind LAYER_FWD_TRAIN = CudnnNormBackendConfig_Kind_LAYER_FWD_TRAIN; static constexpr Kind LAYER_BWD = CudnnNormBackendConfig_Kind_LAYER_BWD; static inline bool Kind_IsValid(int value) { return CudnnNormBackendConfig_Kind_IsValid(value); } static constexpr Kind Kind_MIN = CudnnNormBackendConfig_Kind_Kind_MIN; static constexpr Kind Kind_MAX = CudnnNormBackendConfig_Kind_Kind_MAX; static constexpr int Kind_ARRAYSIZE = CudnnNormBackendConfig_Kind_Kind_ARRAYSIZE; static inline const ::google::protobuf::EnumDescriptor* Kind_descriptor() { return CudnnNormBackendConfig_Kind_descriptor(); } template static inline const std::string& Kind_Name(T value) { return CudnnNormBackendConfig_Kind_Name(value); } static inline bool Kind_Parse(absl::string_view name, Kind* value) { return CudnnNormBackendConfig_Kind_Parse(name, value); } // accessors ------------------------------------------------------- enum : int { kAlgorithmFieldNumber = 2, kEpsilonFieldNumber = 1, kKindFieldNumber = 3, }; // .stream_executor.dnn.AlgorithmProto algorithm = 2; bool has_algorithm() const; void clear_algorithm() ; const ::stream_executor::dnn::AlgorithmProto& algorithm() const; PROTOBUF_NODISCARD ::stream_executor::dnn::AlgorithmProto* release_algorithm(); ::stream_executor::dnn::AlgorithmProto* mutable_algorithm(); void set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); void unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: // double epsilon = 1; void clear_epsilon() ; double epsilon() const; void set_epsilon(double value); private: double _internal_epsilon() const; void _internal_set_epsilon(double value); public: // .xla.gpu.CudnnNormBackendConfig.Kind kind = 3; void clear_kind() ; ::xla::gpu::CudnnNormBackendConfig_Kind kind() const; void set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value); private: ::xla::gpu::CudnnNormBackendConfig_Kind _internal_kind() const; void _internal_set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CudnnNormBackendConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 3, 1, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CudnnNormBackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::stream_executor::dnn::AlgorithmProto* algorithm_; double epsilon_; int kind_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CudnnConvBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnConvBackendConfig) */ { public: inline CudnnConvBackendConfig() : CudnnConvBackendConfig(nullptr) {} ~CudnnConvBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CudnnConvBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline CudnnConvBackendConfig(const CudnnConvBackendConfig& from) : CudnnConvBackendConfig(nullptr, from) {} inline CudnnConvBackendConfig(CudnnConvBackendConfig&& from) noexcept : CudnnConvBackendConfig(nullptr, std::move(from)) {} inline CudnnConvBackendConfig& operator=(const CudnnConvBackendConfig& from) { CopyFrom(from); return *this; } inline CudnnConvBackendConfig& operator=(CudnnConvBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CudnnConvBackendConfig& default_instance() { return *internal_default_instance(); } enum FilterAndBiasReorderingOneofCase { kReorderedInt8NchwVect = 7, FILTER_AND_BIAS_REORDERING_ONEOF_NOT_SET = 0, }; static inline const CudnnConvBackendConfig* internal_default_instance() { return reinterpret_cast( &_CudnnConvBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 0; friend void swap(CudnnConvBackendConfig& a, CudnnConvBackendConfig& b) { a.Swap(&b); } inline void Swap(CudnnConvBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnConvBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnConvBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CudnnConvBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CudnnConvBackendConfig& from) { CudnnConvBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(CudnnConvBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CudnnConvBackendConfig"; } protected: explicit CudnnConvBackendConfig(::google::protobuf::Arena* arena); CudnnConvBackendConfig(::google::protobuf::Arena* arena, const CudnnConvBackendConfig& from); CudnnConvBackendConfig(::google::protobuf::Arena* arena, CudnnConvBackendConfig&& from) noexcept : CudnnConvBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSerializedGraphFieldNumber = 9, kAlgorithmFieldNumber = 6, kConvResultScaleFieldNumber = 4, kSideInputScaleFieldNumber = 5, kLeakyreluAlphaFieldNumber = 8, kActivationModeFieldNumber = 3, kReorderedInt8NchwVectFieldNumber = 7, }; // optional string serialized_graph = 9; bool has_serialized_graph() const; void clear_serialized_graph() ; const std::string& serialized_graph() const; template void set_serialized_graph(Arg_&& arg, Args_... args); std::string* mutable_serialized_graph(); PROTOBUF_NODISCARD std::string* release_serialized_graph(); void set_allocated_serialized_graph(std::string* value); private: const std::string& _internal_serialized_graph() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_serialized_graph( const std::string& value); std::string* _internal_mutable_serialized_graph(); public: // .stream_executor.dnn.AlgorithmProto algorithm = 6; bool has_algorithm() const; void clear_algorithm() ; const ::stream_executor::dnn::AlgorithmProto& algorithm() const; PROTOBUF_NODISCARD ::stream_executor::dnn::AlgorithmProto* release_algorithm(); ::stream_executor::dnn::AlgorithmProto* mutable_algorithm(); void set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); void unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: // double conv_result_scale = 4; void clear_conv_result_scale() ; double conv_result_scale() const; void set_conv_result_scale(double value); private: double _internal_conv_result_scale() const; void _internal_set_conv_result_scale(double value); public: // double side_input_scale = 5; void clear_side_input_scale() ; double side_input_scale() const; void set_side_input_scale(double value); private: double _internal_side_input_scale() const; void _internal_set_side_input_scale(double value); public: // double leakyrelu_alpha = 8; void clear_leakyrelu_alpha() ; double leakyrelu_alpha() const; void set_leakyrelu_alpha(double value); private: double _internal_leakyrelu_alpha() const; void _internal_set_leakyrelu_alpha(double value); public: // .stream_executor.dnn.ActivationMode activation_mode = 3; void clear_activation_mode() ; ::stream_executor::dnn::ActivationMode activation_mode() const; void set_activation_mode(::stream_executor::dnn::ActivationMode value); private: ::stream_executor::dnn::ActivationMode _internal_activation_mode() const; void _internal_set_activation_mode(::stream_executor::dnn::ActivationMode value); public: // bool reordered_int8_nchw_vect = 7; bool has_reordered_int8_nchw_vect() const; void clear_reordered_int8_nchw_vect() ; bool reordered_int8_nchw_vect() const; void set_reordered_int8_nchw_vect(bool value); private: bool _internal_reordered_int8_nchw_vect() const; void _internal_set_reordered_int8_nchw_vect(bool value); public: void clear_filter_and_bias_reordering_oneof(); FilterAndBiasReorderingOneofCase filter_and_bias_reordering_oneof_case() const; // @@protoc_insertion_point(class_scope:xla.gpu.CudnnConvBackendConfig) private: class _Internal; void set_has_reordered_int8_nchw_vect(); inline bool has_filter_and_bias_reordering_oneof() const; inline void clear_has_filter_and_bias_reordering_oneof(); friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 7, 1, 55, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const CudnnConvBackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::internal::ArenaStringPtr serialized_graph_; ::stream_executor::dnn::AlgorithmProto* algorithm_; double conv_result_scale_; double side_input_scale_; double leakyrelu_alpha_; int activation_mode_; union FilterAndBiasReorderingOneofUnion { constexpr FilterAndBiasReorderingOneofUnion() : _constinit_{} {} ::google::protobuf::internal::ConstantInitialized _constinit_; bool reordered_int8_nchw_vect_; } filter_and_bias_reordering_oneof_; ::uint32_t _oneof_case_[1]; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT BitcastBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.BitcastBackendConfig) */ { public: inline BitcastBackendConfig() : BitcastBackendConfig(nullptr) {} ~BitcastBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR BitcastBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline BitcastBackendConfig(const BitcastBackendConfig& from) : BitcastBackendConfig(nullptr, from) {} inline BitcastBackendConfig(BitcastBackendConfig&& from) noexcept : BitcastBackendConfig(nullptr, std::move(from)) {} inline BitcastBackendConfig& operator=(const BitcastBackendConfig& from) { CopyFrom(from); return *this; } inline BitcastBackendConfig& operator=(BitcastBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const BitcastBackendConfig& default_instance() { return *internal_default_instance(); } static inline const BitcastBackendConfig* internal_default_instance() { return reinterpret_cast( &_BitcastBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 2; friend void swap(BitcastBackendConfig& a, BitcastBackendConfig& b) { a.Swap(&b); } inline void Swap(BitcastBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(BitcastBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- BitcastBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const BitcastBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const BitcastBackendConfig& from) { BitcastBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(BitcastBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.BitcastBackendConfig"; } protected: explicit BitcastBackendConfig(::google::protobuf::Arena* arena); BitcastBackendConfig(::google::protobuf::Arena* arena, const BitcastBackendConfig& from); BitcastBackendConfig(::google::protobuf::Arena* arena, BitcastBackendConfig&& from) noexcept : BitcastBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSourceLayoutFieldNumber = 1, kResultLayoutFieldNumber = 2, }; // .xla.LayoutProto source_layout = 1; bool has_source_layout() const; void clear_source_layout() ; const ::xla::LayoutProto& source_layout() const; PROTOBUF_NODISCARD ::xla::LayoutProto* release_source_layout(); ::xla::LayoutProto* mutable_source_layout(); void set_allocated_source_layout(::xla::LayoutProto* value); void unsafe_arena_set_allocated_source_layout(::xla::LayoutProto* value); ::xla::LayoutProto* unsafe_arena_release_source_layout(); private: const ::xla::LayoutProto& _internal_source_layout() const; ::xla::LayoutProto* _internal_mutable_source_layout(); public: // .xla.LayoutProto result_layout = 2; bool has_result_layout() const; void clear_result_layout() ; const ::xla::LayoutProto& result_layout() const; PROTOBUF_NODISCARD ::xla::LayoutProto* release_result_layout(); ::xla::LayoutProto* mutable_result_layout(); void set_allocated_result_layout(::xla::LayoutProto* value); void unsafe_arena_set_allocated_result_layout(::xla::LayoutProto* value); ::xla::LayoutProto* unsafe_arena_release_result_layout(); private: const ::xla::LayoutProto& _internal_result_layout() const; ::xla::LayoutProto* _internal_mutable_result_layout(); public: // @@protoc_insertion_point(class_scope:xla.gpu.BitcastBackendConfig) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 1, 2, 2, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const BitcastBackendConfig& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::LayoutProto* source_layout_; ::xla::LayoutProto* result_layout_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT GpuBackendConfig final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.GpuBackendConfig) */ { public: inline GpuBackendConfig() : GpuBackendConfig(nullptr) {} ~GpuBackendConfig() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR GpuBackendConfig( ::google::protobuf::internal::ConstantInitialized); inline GpuBackendConfig(const GpuBackendConfig& from) : GpuBackendConfig(nullptr, from) {} inline GpuBackendConfig(GpuBackendConfig&& from) noexcept : GpuBackendConfig(nullptr, std::move(from)) {} inline GpuBackendConfig& operator=(const GpuBackendConfig& from) { CopyFrom(from); return *this; } inline GpuBackendConfig& operator=(GpuBackendConfig&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const GpuBackendConfig& default_instance() { return *internal_default_instance(); } enum BackendConfigCase { kCudnnConvBackendConfig = 3, kGemmBackendConfig = 4, kBitcastBackendConfig = 5, kCollectiveBackendConfig = 6, kFusionBackendConfig = 7, kCudnnNormBackendConfig = 8, kCudnnFmhaBackendConfig = 9, kCustomCallBackendConfig = 11, BACKEND_CONFIG_NOT_SET = 0, }; static inline const GpuBackendConfig* internal_default_instance() { return reinterpret_cast( &_GpuBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 14; friend void swap(GpuBackendConfig& a, GpuBackendConfig& b) { a.Swap(&b); } inline void Swap(GpuBackendConfig* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(GpuBackendConfig* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- GpuBackendConfig* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const GpuBackendConfig& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const GpuBackendConfig& from) { GpuBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(GpuBackendConfig* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.GpuBackendConfig"; } protected: explicit GpuBackendConfig(::google::protobuf::Arena* arena); GpuBackendConfig(::google::protobuf::Arena* arena, const GpuBackendConfig& from); GpuBackendConfig(::google::protobuf::Arena* arena, GpuBackendConfig&& from) noexcept : GpuBackendConfig(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kWaitOnOperationQueuesFieldNumber = 2, kReificationCostFieldNumber = 12, kOperationQueueIdFieldNumber = 1, kForceEarliestScheduleFieldNumber = 10, kCudnnConvBackendConfigFieldNumber = 3, kGemmBackendConfigFieldNumber = 4, kBitcastBackendConfigFieldNumber = 5, kCollectiveBackendConfigFieldNumber = 6, kFusionBackendConfigFieldNumber = 7, kCudnnNormBackendConfigFieldNumber = 8, kCudnnFmhaBackendConfigFieldNumber = 9, kCustomCallBackendConfigFieldNumber = 11, }; // repeated int64 wait_on_operation_queues = 2; int wait_on_operation_queues_size() const; private: int _internal_wait_on_operation_queues_size() const; public: void clear_wait_on_operation_queues() ; ::int64_t wait_on_operation_queues(int index) const; void set_wait_on_operation_queues(int index, ::int64_t value); void add_wait_on_operation_queues(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& wait_on_operation_queues() const; ::google::protobuf::RepeatedField<::int64_t>* mutable_wait_on_operation_queues(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_wait_on_operation_queues() const; ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_wait_on_operation_queues(); public: // repeated .xla.gpu.ReificationCost reification_cost = 12; int reification_cost_size() const; private: int _internal_reification_cost_size() const; public: void clear_reification_cost() ; ::xla::gpu::ReificationCost* mutable_reification_cost(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>* mutable_reification_cost(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>& _internal_reification_cost() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>* _internal_mutable_reification_cost(); public: const ::xla::gpu::ReificationCost& reification_cost(int index) const; ::xla::gpu::ReificationCost* add_reification_cost(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>& reification_cost() const; // int64 operation_queue_id = 1; void clear_operation_queue_id() ; ::int64_t operation_queue_id() const; void set_operation_queue_id(::int64_t value); private: ::int64_t _internal_operation_queue_id() const; void _internal_set_operation_queue_id(::int64_t value); public: // bool force_earliest_schedule = 10; void clear_force_earliest_schedule() ; bool force_earliest_schedule() const; void set_force_earliest_schedule(bool value); private: bool _internal_force_earliest_schedule() const; void _internal_set_force_earliest_schedule(bool value); public: // .xla.gpu.CudnnConvBackendConfig cudnn_conv_backend_config = 3; bool has_cudnn_conv_backend_config() const; private: bool _internal_has_cudnn_conv_backend_config() const; public: void clear_cudnn_conv_backend_config() ; const ::xla::gpu::CudnnConvBackendConfig& cudnn_conv_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::CudnnConvBackendConfig* release_cudnn_conv_backend_config(); ::xla::gpu::CudnnConvBackendConfig* mutable_cudnn_conv_backend_config(); void set_allocated_cudnn_conv_backend_config(::xla::gpu::CudnnConvBackendConfig* value); void unsafe_arena_set_allocated_cudnn_conv_backend_config(::xla::gpu::CudnnConvBackendConfig* value); ::xla::gpu::CudnnConvBackendConfig* unsafe_arena_release_cudnn_conv_backend_config(); private: const ::xla::gpu::CudnnConvBackendConfig& _internal_cudnn_conv_backend_config() const; ::xla::gpu::CudnnConvBackendConfig* _internal_mutable_cudnn_conv_backend_config(); public: // .xla.gpu.GemmBackendConfig gemm_backend_config = 4; bool has_gemm_backend_config() const; private: bool _internal_has_gemm_backend_config() const; public: void clear_gemm_backend_config() ; const ::xla::gpu::GemmBackendConfig& gemm_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::GemmBackendConfig* release_gemm_backend_config(); ::xla::gpu::GemmBackendConfig* mutable_gemm_backend_config(); void set_allocated_gemm_backend_config(::xla::gpu::GemmBackendConfig* value); void unsafe_arena_set_allocated_gemm_backend_config(::xla::gpu::GemmBackendConfig* value); ::xla::gpu::GemmBackendConfig* unsafe_arena_release_gemm_backend_config(); private: const ::xla::gpu::GemmBackendConfig& _internal_gemm_backend_config() const; ::xla::gpu::GemmBackendConfig* _internal_mutable_gemm_backend_config(); public: // .xla.gpu.BitcastBackendConfig bitcast_backend_config = 5; bool has_bitcast_backend_config() const; private: bool _internal_has_bitcast_backend_config() const; public: void clear_bitcast_backend_config() ; const ::xla::gpu::BitcastBackendConfig& bitcast_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::BitcastBackendConfig* release_bitcast_backend_config(); ::xla::gpu::BitcastBackendConfig* mutable_bitcast_backend_config(); void set_allocated_bitcast_backend_config(::xla::gpu::BitcastBackendConfig* value); void unsafe_arena_set_allocated_bitcast_backend_config(::xla::gpu::BitcastBackendConfig* value); ::xla::gpu::BitcastBackendConfig* unsafe_arena_release_bitcast_backend_config(); private: const ::xla::gpu::BitcastBackendConfig& _internal_bitcast_backend_config() const; ::xla::gpu::BitcastBackendConfig* _internal_mutable_bitcast_backend_config(); public: // .xla.gpu.CollectiveBackendConfig collective_backend_config = 6; bool has_collective_backend_config() const; private: bool _internal_has_collective_backend_config() const; public: void clear_collective_backend_config() ; const ::xla::gpu::CollectiveBackendConfig& collective_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::CollectiveBackendConfig* release_collective_backend_config(); ::xla::gpu::CollectiveBackendConfig* mutable_collective_backend_config(); void set_allocated_collective_backend_config(::xla::gpu::CollectiveBackendConfig* value); void unsafe_arena_set_allocated_collective_backend_config(::xla::gpu::CollectiveBackendConfig* value); ::xla::gpu::CollectiveBackendConfig* unsafe_arena_release_collective_backend_config(); private: const ::xla::gpu::CollectiveBackendConfig& _internal_collective_backend_config() const; ::xla::gpu::CollectiveBackendConfig* _internal_mutable_collective_backend_config(); public: // .xla.gpu.FusionBackendConfig fusion_backend_config = 7; bool has_fusion_backend_config() const; private: bool _internal_has_fusion_backend_config() const; public: void clear_fusion_backend_config() ; const ::xla::gpu::FusionBackendConfig& fusion_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::FusionBackendConfig* release_fusion_backend_config(); ::xla::gpu::FusionBackendConfig* mutable_fusion_backend_config(); void set_allocated_fusion_backend_config(::xla::gpu::FusionBackendConfig* value); void unsafe_arena_set_allocated_fusion_backend_config(::xla::gpu::FusionBackendConfig* value); ::xla::gpu::FusionBackendConfig* unsafe_arena_release_fusion_backend_config(); private: const ::xla::gpu::FusionBackendConfig& _internal_fusion_backend_config() const; ::xla::gpu::FusionBackendConfig* _internal_mutable_fusion_backend_config(); public: // .xla.gpu.CudnnNormBackendConfig cudnn_norm_backend_config = 8; bool has_cudnn_norm_backend_config() const; private: bool _internal_has_cudnn_norm_backend_config() const; public: void clear_cudnn_norm_backend_config() ; const ::xla::gpu::CudnnNormBackendConfig& cudnn_norm_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::CudnnNormBackendConfig* release_cudnn_norm_backend_config(); ::xla::gpu::CudnnNormBackendConfig* mutable_cudnn_norm_backend_config(); void set_allocated_cudnn_norm_backend_config(::xla::gpu::CudnnNormBackendConfig* value); void unsafe_arena_set_allocated_cudnn_norm_backend_config(::xla::gpu::CudnnNormBackendConfig* value); ::xla::gpu::CudnnNormBackendConfig* unsafe_arena_release_cudnn_norm_backend_config(); private: const ::xla::gpu::CudnnNormBackendConfig& _internal_cudnn_norm_backend_config() const; ::xla::gpu::CudnnNormBackendConfig* _internal_mutable_cudnn_norm_backend_config(); public: // .xla.gpu.CudnnfMHABackendConfig cudnn_fmha_backend_config = 9; bool has_cudnn_fmha_backend_config() const; private: bool _internal_has_cudnn_fmha_backend_config() const; public: void clear_cudnn_fmha_backend_config() ; const ::xla::gpu::CudnnfMHABackendConfig& cudnn_fmha_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::CudnnfMHABackendConfig* release_cudnn_fmha_backend_config(); ::xla::gpu::CudnnfMHABackendConfig* mutable_cudnn_fmha_backend_config(); void set_allocated_cudnn_fmha_backend_config(::xla::gpu::CudnnfMHABackendConfig* value); void unsafe_arena_set_allocated_cudnn_fmha_backend_config(::xla::gpu::CudnnfMHABackendConfig* value); ::xla::gpu::CudnnfMHABackendConfig* unsafe_arena_release_cudnn_fmha_backend_config(); private: const ::xla::gpu::CudnnfMHABackendConfig& _internal_cudnn_fmha_backend_config() const; ::xla::gpu::CudnnfMHABackendConfig* _internal_mutable_cudnn_fmha_backend_config(); public: // .xla.gpu.CustomCallBackendConfig custom_call_backend_config = 11; bool has_custom_call_backend_config() const; private: bool _internal_has_custom_call_backend_config() const; public: void clear_custom_call_backend_config() ; const ::xla::gpu::CustomCallBackendConfig& custom_call_backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::CustomCallBackendConfig* release_custom_call_backend_config(); ::xla::gpu::CustomCallBackendConfig* mutable_custom_call_backend_config(); void set_allocated_custom_call_backend_config(::xla::gpu::CustomCallBackendConfig* value); void unsafe_arena_set_allocated_custom_call_backend_config(::xla::gpu::CustomCallBackendConfig* value); ::xla::gpu::CustomCallBackendConfig* unsafe_arena_release_custom_call_backend_config(); private: const ::xla::gpu::CustomCallBackendConfig& _internal_custom_call_backend_config() const; ::xla::gpu::CustomCallBackendConfig* _internal_mutable_custom_call_backend_config(); public: void clear_backend_config(); BackendConfigCase backend_config_case() const; // @@protoc_insertion_point(class_scope:xla.gpu.GpuBackendConfig) private: class _Internal; void set_has_cudnn_conv_backend_config(); void set_has_gemm_backend_config(); void set_has_bitcast_backend_config(); void set_has_collective_backend_config(); void set_has_fusion_backend_config(); void set_has_cudnn_norm_backend_config(); void set_has_cudnn_fmha_backend_config(); void set_has_custom_call_backend_config(); inline bool has_backend_config() const; inline void clear_has_backend_config(); friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 4, 12, 9, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const GpuBackendConfig& from_msg); ::google::protobuf::RepeatedField<::int64_t> wait_on_operation_queues_; mutable ::google::protobuf::internal::CachedSize _wait_on_operation_queues_cached_byte_size_; ::google::protobuf::RepeatedPtrField< ::xla::gpu::ReificationCost > reification_cost_; ::int64_t operation_queue_id_; bool force_earliest_schedule_; union BackendConfigUnion { constexpr BackendConfigUnion() : _constinit_{} {} ::google::protobuf::internal::ConstantInitialized _constinit_; ::xla::gpu::CudnnConvBackendConfig* cudnn_conv_backend_config_; ::xla::gpu::GemmBackendConfig* gemm_backend_config_; ::xla::gpu::BitcastBackendConfig* bitcast_backend_config_; ::xla::gpu::CollectiveBackendConfig* collective_backend_config_; ::xla::gpu::FusionBackendConfig* fusion_backend_config_; ::xla::gpu::CudnnNormBackendConfig* cudnn_norm_backend_config_; ::xla::gpu::CudnnfMHABackendConfig* cudnn_fmha_backend_config_; ::xla::gpu::CustomCallBackendConfig* custom_call_backend_config_; } backend_config_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::uint32_t _oneof_case_[1]; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // =================================================================== // =================================================================== #ifdef __GNUC__ #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif // __GNUC__ // ------------------------------------------------------------------- // CudnnConvBackendConfig // .stream_executor.dnn.AlgorithmProto algorithm = 6; inline bool CudnnConvBackendConfig::has_algorithm() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.algorithm_ != nullptr); return value; } inline const ::stream_executor::dnn::AlgorithmProto& CudnnConvBackendConfig::_internal_algorithm() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::stream_executor::dnn::AlgorithmProto* p = _impl_.algorithm_; return p != nullptr ? *p : reinterpret_cast(::stream_executor::dnn::_AlgorithmProto_default_instance_); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnConvBackendConfig::algorithm() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnConvBackendConfig::unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnConvBackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::stream_executor::dnn::AlgorithmProto* released = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::unsafe_arena_release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnConvBackendConfig.algorithm) _impl_._has_bits_[0] &= ~0x00000002u; ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::_internal_mutable_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.algorithm_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::stream_executor::dnn::AlgorithmProto>(GetArena()); _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(p); } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::mutable_algorithm() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::stream_executor::dnn::AlgorithmProto* _msg = _internal_mutable_algorithm(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnConvBackendConfig.algorithm) return _msg; } inline void CudnnConvBackendConfig::set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnConvBackendConfig.algorithm) } // double conv_result_scale = 4; inline void CudnnConvBackendConfig::clear_conv_result_scale() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.conv_result_scale_ = 0; } inline double CudnnConvBackendConfig::conv_result_scale() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.conv_result_scale) return _internal_conv_result_scale(); } inline void CudnnConvBackendConfig::set_conv_result_scale(double value) { _internal_set_conv_result_scale(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.conv_result_scale) } inline double CudnnConvBackendConfig::_internal_conv_result_scale() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.conv_result_scale_; } inline void CudnnConvBackendConfig::_internal_set_conv_result_scale(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.conv_result_scale_ = value; } // .stream_executor.dnn.ActivationMode activation_mode = 3; inline void CudnnConvBackendConfig::clear_activation_mode() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.activation_mode_ = 0; } inline ::stream_executor::dnn::ActivationMode CudnnConvBackendConfig::activation_mode() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.activation_mode) return _internal_activation_mode(); } inline void CudnnConvBackendConfig::set_activation_mode(::stream_executor::dnn::ActivationMode value) { _internal_set_activation_mode(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.activation_mode) } inline ::stream_executor::dnn::ActivationMode CudnnConvBackendConfig::_internal_activation_mode() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::stream_executor::dnn::ActivationMode>(_impl_.activation_mode_); } inline void CudnnConvBackendConfig::_internal_set_activation_mode(::stream_executor::dnn::ActivationMode value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.activation_mode_ = value; } // double side_input_scale = 5; inline void CudnnConvBackendConfig::clear_side_input_scale() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.side_input_scale_ = 0; } inline double CudnnConvBackendConfig::side_input_scale() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.side_input_scale) return _internal_side_input_scale(); } inline void CudnnConvBackendConfig::set_side_input_scale(double value) { _internal_set_side_input_scale(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.side_input_scale) } inline double CudnnConvBackendConfig::_internal_side_input_scale() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.side_input_scale_; } inline void CudnnConvBackendConfig::_internal_set_side_input_scale(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.side_input_scale_ = value; } // double leakyrelu_alpha = 8; inline void CudnnConvBackendConfig::clear_leakyrelu_alpha() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.leakyrelu_alpha_ = 0; } inline double CudnnConvBackendConfig::leakyrelu_alpha() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.leakyrelu_alpha) return _internal_leakyrelu_alpha(); } inline void CudnnConvBackendConfig::set_leakyrelu_alpha(double value) { _internal_set_leakyrelu_alpha(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.leakyrelu_alpha) } inline double CudnnConvBackendConfig::_internal_leakyrelu_alpha() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.leakyrelu_alpha_; } inline void CudnnConvBackendConfig::_internal_set_leakyrelu_alpha(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.leakyrelu_alpha_ = value; } // bool reordered_int8_nchw_vect = 7; inline bool CudnnConvBackendConfig::has_reordered_int8_nchw_vect() const { return filter_and_bias_reordering_oneof_case() == kReorderedInt8NchwVect; } inline void CudnnConvBackendConfig::set_has_reordered_int8_nchw_vect() { _impl_._oneof_case_[0] = kReorderedInt8NchwVect; } inline void CudnnConvBackendConfig::clear_reordered_int8_nchw_vect() { ::google::protobuf::internal::TSanWrite(&_impl_); if (filter_and_bias_reordering_oneof_case() == kReorderedInt8NchwVect) { _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_ = false; clear_has_filter_and_bias_reordering_oneof(); } } inline bool CudnnConvBackendConfig::reordered_int8_nchw_vect() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.reordered_int8_nchw_vect) return _internal_reordered_int8_nchw_vect(); } inline void CudnnConvBackendConfig::set_reordered_int8_nchw_vect(bool value) { if (filter_and_bias_reordering_oneof_case() != kReorderedInt8NchwVect) { clear_filter_and_bias_reordering_oneof(); set_has_reordered_int8_nchw_vect(); } _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_ = value; // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.reordered_int8_nchw_vect) } inline bool CudnnConvBackendConfig::_internal_reordered_int8_nchw_vect() const { if (filter_and_bias_reordering_oneof_case() == kReorderedInt8NchwVect) { return _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_; } return false; } // optional string serialized_graph = 9; inline bool CudnnConvBackendConfig::has_serialized_graph() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; return value; } inline void CudnnConvBackendConfig::clear_serialized_graph() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.serialized_graph_.ClearToEmpty(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const std::string& CudnnConvBackendConfig::serialized_graph() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.serialized_graph) return _internal_serialized_graph(); } template inline PROTOBUF_ALWAYS_INLINE void CudnnConvBackendConfig::set_serialized_graph(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] |= 0x00000001u; _impl_.serialized_graph_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.serialized_graph) } inline std::string* CudnnConvBackendConfig::mutable_serialized_graph() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_serialized_graph(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnConvBackendConfig.serialized_graph) return _s; } inline const std::string& CudnnConvBackendConfig::_internal_serialized_graph() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.serialized_graph_.Get(); } inline void CudnnConvBackendConfig::_internal_set_serialized_graph(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] |= 0x00000001u; _impl_.serialized_graph_.Set(value, GetArena()); } inline std::string* CudnnConvBackendConfig::_internal_mutable_serialized_graph() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] |= 0x00000001u; return _impl_.serialized_graph_.Mutable( GetArena()); } inline std::string* CudnnConvBackendConfig::release_serialized_graph() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnConvBackendConfig.serialized_graph) if ((_impl_._has_bits_[0] & 0x00000001u) == 0) { return nullptr; } _impl_._has_bits_[0] &= ~0x00000001u; auto* released = _impl_.serialized_graph_.Release(); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING _impl_.serialized_graph_.Set("", GetArena()); #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING return released; } inline void CudnnConvBackendConfig::set_allocated_serialized_graph(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.serialized_graph_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.serialized_graph_.IsDefault()) { _impl_.serialized_graph_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnConvBackendConfig.serialized_graph) } inline bool CudnnConvBackendConfig::has_filter_and_bias_reordering_oneof() const { return filter_and_bias_reordering_oneof_case() != FILTER_AND_BIAS_REORDERING_ONEOF_NOT_SET; } inline void CudnnConvBackendConfig::clear_has_filter_and_bias_reordering_oneof() { _impl_._oneof_case_[0] = FILTER_AND_BIAS_REORDERING_ONEOF_NOT_SET; } inline CudnnConvBackendConfig::FilterAndBiasReorderingOneofCase CudnnConvBackendConfig::filter_and_bias_reordering_oneof_case() const { return CudnnConvBackendConfig::FilterAndBiasReorderingOneofCase(_impl_._oneof_case_[0]); } // ------------------------------------------------------------------- // GemmBackendConfig // int64 selected_algorithm = 1; inline bool GemmBackendConfig::has_selected_algorithm() const { return algorithm_case() == kSelectedAlgorithm; } inline void GemmBackendConfig::set_has_selected_algorithm() { _impl_._oneof_case_[0] = kSelectedAlgorithm; } inline void GemmBackendConfig::clear_selected_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); if (algorithm_case() == kSelectedAlgorithm) { _impl_.algorithm_.selected_algorithm_ = ::int64_t{0}; clear_has_algorithm(); } } inline ::int64_t GemmBackendConfig::selected_algorithm() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.selected_algorithm) return _internal_selected_algorithm(); } inline void GemmBackendConfig::set_selected_algorithm(::int64_t value) { if (algorithm_case() != kSelectedAlgorithm) { clear_algorithm(); set_has_selected_algorithm(); } _impl_.algorithm_.selected_algorithm_ = value; // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.selected_algorithm) } inline ::int64_t GemmBackendConfig::_internal_selected_algorithm() const { if (algorithm_case() == kSelectedAlgorithm) { return _impl_.algorithm_.selected_algorithm_; } return ::int64_t{0}; } // double alpha_real = 2; inline void GemmBackendConfig::clear_alpha_real() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.alpha_real_ = 0; } inline double GemmBackendConfig::alpha_real() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.alpha_real) return _internal_alpha_real(); } inline void GemmBackendConfig::set_alpha_real(double value) { _internal_set_alpha_real(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.alpha_real) } inline double GemmBackendConfig::_internal_alpha_real() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.alpha_real_; } inline void GemmBackendConfig::_internal_set_alpha_real(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.alpha_real_ = value; } // double alpha_imag = 9; inline void GemmBackendConfig::clear_alpha_imag() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.alpha_imag_ = 0; } inline double GemmBackendConfig::alpha_imag() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.alpha_imag) return _internal_alpha_imag(); } inline void GemmBackendConfig::set_alpha_imag(double value) { _internal_set_alpha_imag(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.alpha_imag) } inline double GemmBackendConfig::_internal_alpha_imag() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.alpha_imag_; } inline void GemmBackendConfig::_internal_set_alpha_imag(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.alpha_imag_ = value; } // double beta = 3; inline void GemmBackendConfig::clear_beta() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.beta_ = 0; } inline double GemmBackendConfig::beta() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.beta) return _internal_beta(); } inline void GemmBackendConfig::set_beta(double value) { _internal_set_beta(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.beta) } inline double GemmBackendConfig::_internal_beta() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.beta_; } inline void GemmBackendConfig::_internal_set_beta(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.beta_ = value; } // .xla.DotDimensionNumbers dot_dimension_numbers = 7; inline bool GemmBackendConfig::has_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& GemmBackendConfig::_internal_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& GemmBackendConfig::dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.dot_dimension_numbers) return _internal_dot_dimension_numbers(); } inline void GemmBackendConfig::unsafe_arena_set_allocated_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dot_dimension_numbers_); } _impl_.dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmBackendConfig.dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* GemmBackendConfig::release_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::DotDimensionNumbers* released = _impl_.dot_dimension_numbers_; _impl_.dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::unsafe_arena_release_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmBackendConfig.dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::DotDimensionNumbers* temp = _impl_.dot_dimension_numbers_; _impl_.dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::_internal_mutable_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::mutable_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmBackendConfig.dot_dimension_numbers) return _msg; } inline void GemmBackendConfig::set_allocated_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmBackendConfig.dot_dimension_numbers) } // .xla.PrecisionConfig precision_config = 12; inline bool GemmBackendConfig::has_precision_config() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.precision_config_ != nullptr); return value; } inline const ::xla::PrecisionConfig& GemmBackendConfig::_internal_precision_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::PrecisionConfig* p = _impl_.precision_config_; return p != nullptr ? *p : reinterpret_cast(::xla::_PrecisionConfig_default_instance_); } inline const ::xla::PrecisionConfig& GemmBackendConfig::precision_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.precision_config) return _internal_precision_config(); } inline void GemmBackendConfig::unsafe_arena_set_allocated_precision_config(::xla::PrecisionConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.precision_config_); } _impl_.precision_config_ = reinterpret_cast<::xla::PrecisionConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmBackendConfig.precision_config) } inline ::xla::PrecisionConfig* GemmBackendConfig::release_precision_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::PrecisionConfig* released = _impl_.precision_config_; _impl_.precision_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::PrecisionConfig* GemmBackendConfig::unsafe_arena_release_precision_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmBackendConfig.precision_config) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::PrecisionConfig* temp = _impl_.precision_config_; _impl_.precision_config_ = nullptr; return temp; } inline ::xla::PrecisionConfig* GemmBackendConfig::_internal_mutable_precision_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.precision_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::PrecisionConfig>(GetArena()); _impl_.precision_config_ = reinterpret_cast<::xla::PrecisionConfig*>(p); } return _impl_.precision_config_; } inline ::xla::PrecisionConfig* GemmBackendConfig::mutable_precision_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::PrecisionConfig* _msg = _internal_mutable_precision_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmBackendConfig.precision_config) return _msg; } inline void GemmBackendConfig::set_allocated_precision_config(::xla::PrecisionConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.precision_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.precision_config_ = reinterpret_cast<::xla::PrecisionConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmBackendConfig.precision_config) } // .xla.gpu.GemmBackendConfig.Epilogue epilogue = 13; inline void GemmBackendConfig::clear_epilogue() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.epilogue_ = 0; } inline ::xla::gpu::GemmBackendConfig_Epilogue GemmBackendConfig::epilogue() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.epilogue) return _internal_epilogue(); } inline void GemmBackendConfig::set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value) { _internal_set_epilogue(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.epilogue) } inline ::xla::gpu::GemmBackendConfig_Epilogue GemmBackendConfig::_internal_epilogue() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::xla::gpu::GemmBackendConfig_Epilogue>(_impl_.epilogue_); } inline void GemmBackendConfig::_internal_set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.epilogue_ = value; } // optional int64 lhs_stride = 14; inline bool GemmBackendConfig::has_lhs_stride() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; return value; } inline void GemmBackendConfig::clear_lhs_stride() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.lhs_stride_ = ::int64_t{0}; _impl_._has_bits_[0] &= ~0x00000004u; } inline ::int64_t GemmBackendConfig::lhs_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.lhs_stride) return _internal_lhs_stride(); } inline void GemmBackendConfig::set_lhs_stride(::int64_t value) { _internal_set_lhs_stride(value); _impl_._has_bits_[0] |= 0x00000004u; // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.lhs_stride) } inline ::int64_t GemmBackendConfig::_internal_lhs_stride() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.lhs_stride_; } inline void GemmBackendConfig::_internal_set_lhs_stride(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.lhs_stride_ = value; } // optional int64 rhs_stride = 15; inline bool GemmBackendConfig::has_rhs_stride() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; return value; } inline void GemmBackendConfig::clear_rhs_stride() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.rhs_stride_ = ::int64_t{0}; _impl_._has_bits_[0] &= ~0x00000008u; } inline ::int64_t GemmBackendConfig::rhs_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.rhs_stride) return _internal_rhs_stride(); } inline void GemmBackendConfig::set_rhs_stride(::int64_t value) { _internal_set_rhs_stride(value); _impl_._has_bits_[0] |= 0x00000008u; // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.rhs_stride) } inline ::int64_t GemmBackendConfig::_internal_rhs_stride() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.rhs_stride_; } inline void GemmBackendConfig::_internal_set_rhs_stride(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.rhs_stride_ = value; } // optional bool grad_x = 16; inline bool GemmBackendConfig::has_grad_x() const { bool value = (_impl_._has_bits_[0] & 0x00000010u) != 0; return value; } inline void GemmBackendConfig::clear_grad_x() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.grad_x_ = false; _impl_._has_bits_[0] &= ~0x00000010u; } inline bool GemmBackendConfig::grad_x() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.grad_x) return _internal_grad_x(); } inline void GemmBackendConfig::set_grad_x(bool value) { _internal_set_grad_x(value); _impl_._has_bits_[0] |= 0x00000010u; // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.grad_x) } inline bool GemmBackendConfig::_internal_grad_x() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.grad_x_; } inline void GemmBackendConfig::_internal_set_grad_x(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.grad_x_ = value; } // optional bool grad_y = 17; inline bool GemmBackendConfig::has_grad_y() const { bool value = (_impl_._has_bits_[0] & 0x00000020u) != 0; return value; } inline void GemmBackendConfig::clear_grad_y() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.grad_y_ = false; _impl_._has_bits_[0] &= ~0x00000020u; } inline bool GemmBackendConfig::grad_y() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.grad_y) return _internal_grad_y(); } inline void GemmBackendConfig::set_grad_y(bool value) { _internal_set_grad_y(value); _impl_._has_bits_[0] |= 0x00000020u; // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.grad_y) } inline bool GemmBackendConfig::_internal_grad_y() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.grad_y_; } inline void GemmBackendConfig::_internal_set_grad_y(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.grad_y_ = value; } // bool damax_output = 18; inline void GemmBackendConfig::clear_damax_output() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.damax_output_ = false; } inline bool GemmBackendConfig::damax_output() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.damax_output) return _internal_damax_output(); } inline void GemmBackendConfig::set_damax_output(bool value) { _internal_set_damax_output(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.damax_output) } inline bool GemmBackendConfig::_internal_damax_output() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.damax_output_; } inline void GemmBackendConfig::_internal_set_damax_output(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.damax_output_ = value; } inline bool GemmBackendConfig::has_algorithm() const { return algorithm_case() != ALGORITHM_NOT_SET; } inline void GemmBackendConfig::clear_has_algorithm() { _impl_._oneof_case_[0] = ALGORITHM_NOT_SET; } inline GemmBackendConfig::AlgorithmCase GemmBackendConfig::algorithm_case() const { return GemmBackendConfig::AlgorithmCase(_impl_._oneof_case_[0]); } // ------------------------------------------------------------------- // BitcastBackendConfig // .xla.LayoutProto source_layout = 1; inline bool BitcastBackendConfig::has_source_layout() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.source_layout_ != nullptr); return value; } inline const ::xla::LayoutProto& BitcastBackendConfig::_internal_source_layout() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::LayoutProto* p = _impl_.source_layout_; return p != nullptr ? *p : reinterpret_cast(::xla::_LayoutProto_default_instance_); } inline const ::xla::LayoutProto& BitcastBackendConfig::source_layout() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.BitcastBackendConfig.source_layout) return _internal_source_layout(); } inline void BitcastBackendConfig::unsafe_arena_set_allocated_source_layout(::xla::LayoutProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_layout_); } _impl_.source_layout_ = reinterpret_cast<::xla::LayoutProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.BitcastBackendConfig.source_layout) } inline ::xla::LayoutProto* BitcastBackendConfig::release_source_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::LayoutProto* released = _impl_.source_layout_; _impl_.source_layout_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::LayoutProto* BitcastBackendConfig::unsafe_arena_release_source_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.BitcastBackendConfig.source_layout) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::LayoutProto* temp = _impl_.source_layout_; _impl_.source_layout_ = nullptr; return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::_internal_mutable_source_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.source_layout_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::LayoutProto>(GetArena()); _impl_.source_layout_ = reinterpret_cast<::xla::LayoutProto*>(p); } return _impl_.source_layout_; } inline ::xla::LayoutProto* BitcastBackendConfig::mutable_source_layout() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::LayoutProto* _msg = _internal_mutable_source_layout(); // @@protoc_insertion_point(field_mutable:xla.gpu.BitcastBackendConfig.source_layout) return _msg; } inline void BitcastBackendConfig::set_allocated_source_layout(::xla::LayoutProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_layout_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.source_layout_ = reinterpret_cast<::xla::LayoutProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.BitcastBackendConfig.source_layout) } // .xla.LayoutProto result_layout = 2; inline bool BitcastBackendConfig::has_result_layout() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.result_layout_ != nullptr); return value; } inline const ::xla::LayoutProto& BitcastBackendConfig::_internal_result_layout() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::LayoutProto* p = _impl_.result_layout_; return p != nullptr ? *p : reinterpret_cast(::xla::_LayoutProto_default_instance_); } inline const ::xla::LayoutProto& BitcastBackendConfig::result_layout() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.BitcastBackendConfig.result_layout) return _internal_result_layout(); } inline void BitcastBackendConfig::unsafe_arena_set_allocated_result_layout(::xla::LayoutProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.result_layout_); } _impl_.result_layout_ = reinterpret_cast<::xla::LayoutProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.BitcastBackendConfig.result_layout) } inline ::xla::LayoutProto* BitcastBackendConfig::release_result_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::LayoutProto* released = _impl_.result_layout_; _impl_.result_layout_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::LayoutProto* BitcastBackendConfig::unsafe_arena_release_result_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.BitcastBackendConfig.result_layout) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::LayoutProto* temp = _impl_.result_layout_; _impl_.result_layout_ = nullptr; return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::_internal_mutable_result_layout() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.result_layout_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::LayoutProto>(GetArena()); _impl_.result_layout_ = reinterpret_cast<::xla::LayoutProto*>(p); } return _impl_.result_layout_; } inline ::xla::LayoutProto* BitcastBackendConfig::mutable_result_layout() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::LayoutProto* _msg = _internal_mutable_result_layout(); // @@protoc_insertion_point(field_mutable:xla.gpu.BitcastBackendConfig.result_layout) return _msg; } inline void BitcastBackendConfig::set_allocated_result_layout(::xla::LayoutProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.result_layout_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.result_layout_ = reinterpret_cast<::xla::LayoutProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.BitcastBackendConfig.result_layout) } // ------------------------------------------------------------------- // CollectiveBackendConfig // bool is_sync = 1; inline void CollectiveBackendConfig::clear_is_sync() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_sync_ = false; } inline bool CollectiveBackendConfig::is_sync() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.is_sync) return _internal_is_sync(); } inline void CollectiveBackendConfig::set_is_sync(bool value) { _internal_set_is_sync(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.is_sync) } inline bool CollectiveBackendConfig::_internal_is_sync() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_sync_; } inline void CollectiveBackendConfig::_internal_set_is_sync(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_sync_ = value; } // bool is_pipelined = 3; inline void CollectiveBackendConfig::clear_is_pipelined() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_pipelined_ = false; } inline bool CollectiveBackendConfig::is_pipelined() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.is_pipelined) return _internal_is_pipelined(); } inline void CollectiveBackendConfig::set_is_pipelined(bool value) { _internal_set_is_pipelined(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.is_pipelined) } inline bool CollectiveBackendConfig::_internal_is_pipelined() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_pipelined_; } inline void CollectiveBackendConfig::_internal_set_is_pipelined(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_pipelined_ = value; } // .xla.gpu.CollectiveBackendConfig.CollectiveBackend backend = 5; inline void CollectiveBackendConfig::clear_backend() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.backend_ = 0; } inline ::xla::gpu::CollectiveBackendConfig_CollectiveBackend CollectiveBackendConfig::backend() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.backend) return _internal_backend(); } inline void CollectiveBackendConfig::set_backend(::xla::gpu::CollectiveBackendConfig_CollectiveBackend value) { _internal_set_backend(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.backend) } inline ::xla::gpu::CollectiveBackendConfig_CollectiveBackend CollectiveBackendConfig::_internal_backend() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::xla::gpu::CollectiveBackendConfig_CollectiveBackend>(_impl_.backend_); } inline void CollectiveBackendConfig::_internal_set_backend(::xla::gpu::CollectiveBackendConfig_CollectiveBackend value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.backend_ = value; } // ------------------------------------------------------------------- // ReificationCost // double end_to_end_cycles = 1; inline void ReificationCost::clear_end_to_end_cycles() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.end_to_end_cycles_ = 0; } inline double ReificationCost::end_to_end_cycles() const { // @@protoc_insertion_point(field_get:xla.gpu.ReificationCost.end_to_end_cycles) return _internal_end_to_end_cycles(); } inline void ReificationCost::set_end_to_end_cycles(double value) { _internal_set_end_to_end_cycles(value); // @@protoc_insertion_point(field_set:xla.gpu.ReificationCost.end_to_end_cycles) } inline double ReificationCost::_internal_end_to_end_cycles() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.end_to_end_cycles_; } inline void ReificationCost::_internal_set_end_to_end_cycles(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.end_to_end_cycles_ = value; } // double exec_time_us = 2; inline void ReificationCost::clear_exec_time_us() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.exec_time_us_ = 0; } inline double ReificationCost::exec_time_us() const { // @@protoc_insertion_point(field_get:xla.gpu.ReificationCost.exec_time_us) return _internal_exec_time_us(); } inline void ReificationCost::set_exec_time_us(double value) { _internal_set_exec_time_us(value); // @@protoc_insertion_point(field_set:xla.gpu.ReificationCost.exec_time_us) } inline double ReificationCost::_internal_exec_time_us() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.exec_time_us_; } inline void ReificationCost::_internal_set_exec_time_us(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.exec_time_us_ = value; } // double compute_time_us = 3; inline void ReificationCost::clear_compute_time_us() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.compute_time_us_ = 0; } inline double ReificationCost::compute_time_us() const { // @@protoc_insertion_point(field_get:xla.gpu.ReificationCost.compute_time_us) return _internal_compute_time_us(); } inline void ReificationCost::set_compute_time_us(double value) { _internal_set_compute_time_us(value); // @@protoc_insertion_point(field_set:xla.gpu.ReificationCost.compute_time_us) } inline double ReificationCost::_internal_compute_time_us() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.compute_time_us_; } inline void ReificationCost::_internal_set_compute_time_us(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.compute_time_us_ = value; } // double memory_access_time_us = 4; inline void ReificationCost::clear_memory_access_time_us() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.memory_access_time_us_ = 0; } inline double ReificationCost::memory_access_time_us() const { // @@protoc_insertion_point(field_get:xla.gpu.ReificationCost.memory_access_time_us) return _internal_memory_access_time_us(); } inline void ReificationCost::set_memory_access_time_us(double value) { _internal_set_memory_access_time_us(value); // @@protoc_insertion_point(field_set:xla.gpu.ReificationCost.memory_access_time_us) } inline double ReificationCost::_internal_memory_access_time_us() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.memory_access_time_us_; } inline void ReificationCost::_internal_set_memory_access_time_us(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.memory_access_time_us_ = value; } // string name = 5; inline void ReificationCost::clear_name() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.ClearToEmpty(); } inline const std::string& ReificationCost::name() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ReificationCost.name) return _internal_name(); } template inline PROTOBUF_ALWAYS_INLINE void ReificationCost::set_name(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.ReificationCost.name) } inline std::string* ReificationCost::mutable_name() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_name(); // @@protoc_insertion_point(field_mutable:xla.gpu.ReificationCost.name) return _s; } inline const std::string& ReificationCost::_internal_name() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.name_.Get(); } inline void ReificationCost::_internal_set_name(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.Set(value, GetArena()); } inline std::string* ReificationCost::_internal_mutable_name() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.name_.Mutable( GetArena()); } inline std::string* ReificationCost::release_name() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ReificationCost.name) return _impl_.name_.Release(); } inline void ReificationCost::set_allocated_name(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.name_.IsDefault()) { _impl_.name_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.ReificationCost.name) } // ------------------------------------------------------------------- // CustomFusionConfig // string name = 1; inline void CustomFusionConfig::clear_name() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.ClearToEmpty(); } inline const std::string& CustomFusionConfig::name() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CustomFusionConfig.name) return _internal_name(); } template inline PROTOBUF_ALWAYS_INLINE void CustomFusionConfig::set_name(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.CustomFusionConfig.name) } inline std::string* CustomFusionConfig::mutable_name() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_name(); // @@protoc_insertion_point(field_mutable:xla.gpu.CustomFusionConfig.name) return _s; } inline const std::string& CustomFusionConfig::_internal_name() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.name_.Get(); } inline void CustomFusionConfig::_internal_set_name(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.Set(value, GetArena()); } inline std::string* CustomFusionConfig::_internal_mutable_name() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.name_.Mutable( GetArena()); } inline std::string* CustomFusionConfig::release_name() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CustomFusionConfig.name) return _impl_.name_.Release(); } inline void CustomFusionConfig::set_allocated_name(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.name_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.name_.IsDefault()) { _impl_.name_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.CustomFusionConfig.name) } // int32 kernel_index = 2; inline void CustomFusionConfig::clear_kernel_index() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_index_ = 0; } inline ::int32_t CustomFusionConfig::kernel_index() const { // @@protoc_insertion_point(field_get:xla.gpu.CustomFusionConfig.kernel_index) return _internal_kernel_index(); } inline void CustomFusionConfig::set_kernel_index(::int32_t value) { _internal_set_kernel_index(value); // @@protoc_insertion_point(field_set:xla.gpu.CustomFusionConfig.kernel_index) } inline ::int32_t CustomFusionConfig::_internal_kernel_index() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.kernel_index_; } inline void CustomFusionConfig::_internal_set_kernel_index(::int32_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_index_ = value; } // ------------------------------------------------------------------- // CuDnnFusionConfig // int64 plan_id = 1; inline void CuDnnFusionConfig::clear_plan_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.plan_id_ = ::int64_t{0}; } inline ::int64_t CuDnnFusionConfig::plan_id() const { // @@protoc_insertion_point(field_get:xla.gpu.CuDnnFusionConfig.plan_id) return _internal_plan_id(); } inline void CuDnnFusionConfig::set_plan_id(::int64_t value) { _internal_set_plan_id(value); // @@protoc_insertion_point(field_set:xla.gpu.CuDnnFusionConfig.plan_id) } inline ::int64_t CuDnnFusionConfig::_internal_plan_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.plan_id_; } inline void CuDnnFusionConfig::_internal_set_plan_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.plan_id_ = value; } // ------------------------------------------------------------------- // Tile // repeated int64 sizes = 1; inline int Tile::_internal_sizes_size() const { return _internal_sizes().size(); } inline int Tile::sizes_size() const { return _internal_sizes_size(); } inline void Tile::clear_sizes() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.sizes_.Clear(); } inline ::int64_t Tile::sizes(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.Tile.sizes) return _internal_sizes().Get(index); } inline void Tile::set_sizes(int index, ::int64_t value) { _internal_mutable_sizes()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.Tile.sizes) } inline void Tile::add_sizes(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_sizes()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.Tile.sizes) } inline const ::google::protobuf::RepeatedField<::int64_t>& Tile::sizes() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.Tile.sizes) return _internal_sizes(); } inline ::google::protobuf::RepeatedField<::int64_t>* Tile::mutable_sizes() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.Tile.sizes) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_sizes(); } inline const ::google::protobuf::RepeatedField<::int64_t>& Tile::_internal_sizes() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.sizes_; } inline ::google::protobuf::RepeatedField<::int64_t>* Tile::_internal_mutable_sizes() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.sizes_; } // ------------------------------------------------------------------- // BlockLevelFusionConfig // repeated .xla.gpu.Tile output_tiles = 3; inline int BlockLevelFusionConfig::_internal_output_tiles_size() const { return _internal_output_tiles().size(); } inline int BlockLevelFusionConfig::output_tiles_size() const { return _internal_output_tiles_size(); } inline void BlockLevelFusionConfig::clear_output_tiles() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.output_tiles_.Clear(); } inline ::xla::gpu::Tile* BlockLevelFusionConfig::mutable_output_tiles(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.BlockLevelFusionConfig.output_tiles) return _internal_mutable_output_tiles()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>* BlockLevelFusionConfig::mutable_output_tiles() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.BlockLevelFusionConfig.output_tiles) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_output_tiles(); } inline const ::xla::gpu::Tile& BlockLevelFusionConfig::output_tiles(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.output_tiles) return _internal_output_tiles().Get(index); } inline ::xla::gpu::Tile* BlockLevelFusionConfig::add_output_tiles() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::Tile* _add = _internal_mutable_output_tiles()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.BlockLevelFusionConfig.output_tiles) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>& BlockLevelFusionConfig::output_tiles() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.BlockLevelFusionConfig.output_tiles) return _internal_output_tiles(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>& BlockLevelFusionConfig::_internal_output_tiles() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.output_tiles_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::Tile>* BlockLevelFusionConfig::_internal_mutable_output_tiles() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.output_tiles_; } // int64 num_warps = 2; inline void BlockLevelFusionConfig::clear_num_warps() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_warps_ = ::int64_t{0}; } inline ::int64_t BlockLevelFusionConfig::num_warps() const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.num_warps) return _internal_num_warps(); } inline void BlockLevelFusionConfig::set_num_warps(::int64_t value) { _internal_set_num_warps(value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.num_warps) } inline ::int64_t BlockLevelFusionConfig::_internal_num_warps() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.num_warps_; } inline void BlockLevelFusionConfig::_internal_set_num_warps(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_warps_ = value; } // int32 num_ctas = 4; inline void BlockLevelFusionConfig::clear_num_ctas() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_ctas_ = 0; } inline ::int32_t BlockLevelFusionConfig::num_ctas() const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.num_ctas) return _internal_num_ctas(); } inline void BlockLevelFusionConfig::set_num_ctas(::int32_t value) { _internal_set_num_ctas(value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.num_ctas) } inline ::int32_t BlockLevelFusionConfig::_internal_num_ctas() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.num_ctas_; } inline void BlockLevelFusionConfig::_internal_set_num_ctas(::int32_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_ctas_ = value; } // int32 num_stages = 5; inline void BlockLevelFusionConfig::clear_num_stages() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_stages_ = 0; } inline ::int32_t BlockLevelFusionConfig::num_stages() const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.num_stages) return _internal_num_stages(); } inline void BlockLevelFusionConfig::set_num_stages(::int32_t value) { _internal_set_num_stages(value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.num_stages) } inline ::int32_t BlockLevelFusionConfig::_internal_num_stages() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.num_stages_; } inline void BlockLevelFusionConfig::_internal_set_num_stages(::int32_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.num_stages_ = value; } // bool is_tma_allowed = 6; inline void BlockLevelFusionConfig::clear_is_tma_allowed() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_tma_allowed_ = false; } inline bool BlockLevelFusionConfig::is_tma_allowed() const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.is_tma_allowed) return _internal_is_tma_allowed(); } inline void BlockLevelFusionConfig::set_is_tma_allowed(bool value) { _internal_set_is_tma_allowed(value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.is_tma_allowed) } inline bool BlockLevelFusionConfig::_internal_is_tma_allowed() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_tma_allowed_; } inline void BlockLevelFusionConfig::_internal_set_is_tma_allowed(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_tma_allowed_ = value; } // ------------------------------------------------------------------- // DynamicMemcpyConfig // bool depends_on_loop = 1; inline void DynamicMemcpyConfig::clear_depends_on_loop() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.depends_on_loop_ = false; } inline bool DynamicMemcpyConfig::depends_on_loop() const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyConfig.depends_on_loop) return _internal_depends_on_loop(); } inline void DynamicMemcpyConfig::set_depends_on_loop(bool value) { _internal_set_depends_on_loop(value); // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyConfig.depends_on_loop) } inline bool DynamicMemcpyConfig::_internal_depends_on_loop() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.depends_on_loop_; } inline void DynamicMemcpyConfig::_internal_set_depends_on_loop(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.depends_on_loop_ = value; } // repeated int64 src_offset_bytes = 2; inline int DynamicMemcpyConfig::_internal_src_offset_bytes_size() const { return _internal_src_offset_bytes().size(); } inline int DynamicMemcpyConfig::src_offset_bytes_size() const { return _internal_src_offset_bytes_size(); } inline void DynamicMemcpyConfig::clear_src_offset_bytes() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.src_offset_bytes_.Clear(); } inline ::int64_t DynamicMemcpyConfig::src_offset_bytes(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyConfig.src_offset_bytes) return _internal_src_offset_bytes().Get(index); } inline void DynamicMemcpyConfig::set_src_offset_bytes(int index, ::int64_t value) { _internal_mutable_src_offset_bytes()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyConfig.src_offset_bytes) } inline void DynamicMemcpyConfig::add_src_offset_bytes(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_src_offset_bytes()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.DynamicMemcpyConfig.src_offset_bytes) } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyConfig::src_offset_bytes() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.DynamicMemcpyConfig.src_offset_bytes) return _internal_src_offset_bytes(); } inline ::google::protobuf::RepeatedField<::int64_t>* DynamicMemcpyConfig::mutable_src_offset_bytes() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.DynamicMemcpyConfig.src_offset_bytes) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_src_offset_bytes(); } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyConfig::_internal_src_offset_bytes() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.src_offset_bytes_; } inline ::google::protobuf::RepeatedField<::int64_t>* DynamicMemcpyConfig::_internal_mutable_src_offset_bytes() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.src_offset_bytes_; } // repeated int64 dst_offset_bytes = 3; inline int DynamicMemcpyConfig::_internal_dst_offset_bytes_size() const { return _internal_dst_offset_bytes().size(); } inline int DynamicMemcpyConfig::dst_offset_bytes_size() const { return _internal_dst_offset_bytes_size(); } inline void DynamicMemcpyConfig::clear_dst_offset_bytes() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.dst_offset_bytes_.Clear(); } inline ::int64_t DynamicMemcpyConfig::dst_offset_bytes(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyConfig.dst_offset_bytes) return _internal_dst_offset_bytes().Get(index); } inline void DynamicMemcpyConfig::set_dst_offset_bytes(int index, ::int64_t value) { _internal_mutable_dst_offset_bytes()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyConfig.dst_offset_bytes) } inline void DynamicMemcpyConfig::add_dst_offset_bytes(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_dst_offset_bytes()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.DynamicMemcpyConfig.dst_offset_bytes) } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyConfig::dst_offset_bytes() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.DynamicMemcpyConfig.dst_offset_bytes) return _internal_dst_offset_bytes(); } inline ::google::protobuf::RepeatedField<::int64_t>* DynamicMemcpyConfig::mutable_dst_offset_bytes() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.DynamicMemcpyConfig.dst_offset_bytes) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_dst_offset_bytes(); } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyConfig::_internal_dst_offset_bytes() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.dst_offset_bytes_; } inline ::google::protobuf::RepeatedField<::int64_t>* DynamicMemcpyConfig::_internal_mutable_dst_offset_bytes() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.dst_offset_bytes_; } // ------------------------------------------------------------------- // FusionBackendConfig // string kind = 1; inline void FusionBackendConfig::clear_kind() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_.ClearToEmpty(); } inline const std::string& FusionBackendConfig::kind() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.kind) return _internal_kind(); } template inline PROTOBUF_ALWAYS_INLINE void FusionBackendConfig::set_kind(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.FusionBackendConfig.kind) } inline std::string* FusionBackendConfig::mutable_kind() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_kind(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.kind) return _s; } inline const std::string& FusionBackendConfig::_internal_kind() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.kind_.Get(); } inline void FusionBackendConfig::_internal_set_kind(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_.Set(value, GetArena()); } inline std::string* FusionBackendConfig::_internal_mutable_kind() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.kind_.Mutable( GetArena()); } inline std::string* FusionBackendConfig::release_kind() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.kind) return _impl_.kind_.Release(); } inline void FusionBackendConfig::set_allocated_kind(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.kind_.IsDefault()) { _impl_.kind_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.kind) } // .xla.AutotuneResult.TritonGemmKey triton_gemm_config = 2; inline bool FusionBackendConfig::has_triton_gemm_config() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.triton_gemm_config_ != nullptr); return value; } inline const ::xla::AutotuneResult_TritonGemmKey& FusionBackendConfig::_internal_triton_gemm_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::AutotuneResult_TritonGemmKey* p = _impl_.triton_gemm_config_; return p != nullptr ? *p : reinterpret_cast(::xla::_AutotuneResult_TritonGemmKey_default_instance_); } inline const ::xla::AutotuneResult_TritonGemmKey& FusionBackendConfig::triton_gemm_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.triton_gemm_config) return _internal_triton_gemm_config(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_triton_gemm_config(::xla::AutotuneResult_TritonGemmKey* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.triton_gemm_config_); } _impl_.triton_gemm_config_ = reinterpret_cast<::xla::AutotuneResult_TritonGemmKey*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.triton_gemm_config) } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::release_triton_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::AutotuneResult_TritonGemmKey* released = _impl_.triton_gemm_config_; _impl_.triton_gemm_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::unsafe_arena_release_triton_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.triton_gemm_config) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::AutotuneResult_TritonGemmKey* temp = _impl_.triton_gemm_config_; _impl_.triton_gemm_config_ = nullptr; return temp; } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::_internal_mutable_triton_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.triton_gemm_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::AutotuneResult_TritonGemmKey>(GetArena()); _impl_.triton_gemm_config_ = reinterpret_cast<::xla::AutotuneResult_TritonGemmKey*>(p); } return _impl_.triton_gemm_config_; } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::mutable_triton_gemm_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::AutotuneResult_TritonGemmKey* _msg = _internal_mutable_triton_gemm_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.triton_gemm_config) return _msg; } inline void FusionBackendConfig::set_allocated_triton_gemm_config(::xla::AutotuneResult_TritonGemmKey* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.triton_gemm_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.triton_gemm_config_ = reinterpret_cast<::xla::AutotuneResult_TritonGemmKey*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.triton_gemm_config) } // .xla.gpu.BlockLevelFusionConfig block_level_fusion_config = 6; inline bool FusionBackendConfig::has_block_level_fusion_config() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; PROTOBUF_ASSUME(!value || _impl_.block_level_fusion_config_ != nullptr); return value; } inline void FusionBackendConfig::clear_block_level_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.block_level_fusion_config_ != nullptr) _impl_.block_level_fusion_config_->Clear(); _impl_._has_bits_[0] &= ~0x00000008u; } inline const ::xla::gpu::BlockLevelFusionConfig& FusionBackendConfig::_internal_block_level_fusion_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::BlockLevelFusionConfig* p = _impl_.block_level_fusion_config_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_BlockLevelFusionConfig_default_instance_); } inline const ::xla::gpu::BlockLevelFusionConfig& FusionBackendConfig::block_level_fusion_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.block_level_fusion_config) return _internal_block_level_fusion_config(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_block_level_fusion_config(::xla::gpu::BlockLevelFusionConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.block_level_fusion_config_); } _impl_.block_level_fusion_config_ = reinterpret_cast<::xla::gpu::BlockLevelFusionConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.block_level_fusion_config) } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::release_block_level_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000008u; ::xla::gpu::BlockLevelFusionConfig* released = _impl_.block_level_fusion_config_; _impl_.block_level_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::unsafe_arena_release_block_level_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.block_level_fusion_config) _impl_._has_bits_[0] &= ~0x00000008u; ::xla::gpu::BlockLevelFusionConfig* temp = _impl_.block_level_fusion_config_; _impl_.block_level_fusion_config_ = nullptr; return temp; } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::_internal_mutable_block_level_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.block_level_fusion_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::BlockLevelFusionConfig>(GetArena()); _impl_.block_level_fusion_config_ = reinterpret_cast<::xla::gpu::BlockLevelFusionConfig*>(p); } return _impl_.block_level_fusion_config_; } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::mutable_block_level_fusion_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000008u; ::xla::gpu::BlockLevelFusionConfig* _msg = _internal_mutable_block_level_fusion_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.block_level_fusion_config) return _msg; } inline void FusionBackendConfig::set_allocated_block_level_fusion_config(::xla::gpu::BlockLevelFusionConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.block_level_fusion_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } _impl_.block_level_fusion_config_ = reinterpret_cast<::xla::gpu::BlockLevelFusionConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.block_level_fusion_config) } // .xla.gpu.CustomFusionConfig custom_fusion_config = 4; inline bool FusionBackendConfig::has_custom_fusion_config() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.custom_fusion_config_ != nullptr); return value; } inline void FusionBackendConfig::clear_custom_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.custom_fusion_config_ != nullptr) _impl_.custom_fusion_config_->Clear(); _impl_._has_bits_[0] &= ~0x00000002u; } inline const ::xla::gpu::CustomFusionConfig& FusionBackendConfig::_internal_custom_fusion_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::CustomFusionConfig* p = _impl_.custom_fusion_config_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_CustomFusionConfig_default_instance_); } inline const ::xla::gpu::CustomFusionConfig& FusionBackendConfig::custom_fusion_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.custom_fusion_config) return _internal_custom_fusion_config(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_custom_fusion_config(::xla::gpu::CustomFusionConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.custom_fusion_config_); } _impl_.custom_fusion_config_ = reinterpret_cast<::xla::gpu::CustomFusionConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.custom_fusion_config) } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::release_custom_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::gpu::CustomFusionConfig* released = _impl_.custom_fusion_config_; _impl_.custom_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::unsafe_arena_release_custom_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.custom_fusion_config) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::gpu::CustomFusionConfig* temp = _impl_.custom_fusion_config_; _impl_.custom_fusion_config_ = nullptr; return temp; } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::_internal_mutable_custom_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.custom_fusion_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CustomFusionConfig>(GetArena()); _impl_.custom_fusion_config_ = reinterpret_cast<::xla::gpu::CustomFusionConfig*>(p); } return _impl_.custom_fusion_config_; } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::mutable_custom_fusion_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::gpu::CustomFusionConfig* _msg = _internal_mutable_custom_fusion_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.custom_fusion_config) return _msg; } inline void FusionBackendConfig::set_allocated_custom_fusion_config(::xla::gpu::CustomFusionConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.custom_fusion_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.custom_fusion_config_ = reinterpret_cast<::xla::gpu::CustomFusionConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.custom_fusion_config) } // .xla.gpu.CuDnnFusionConfig cudnn_fusion_config = 5; inline bool FusionBackendConfig::has_cudnn_fusion_config() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.cudnn_fusion_config_ != nullptr); return value; } inline void FusionBackendConfig::clear_cudnn_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.cudnn_fusion_config_ != nullptr) _impl_.cudnn_fusion_config_->Clear(); _impl_._has_bits_[0] &= ~0x00000004u; } inline const ::xla::gpu::CuDnnFusionConfig& FusionBackendConfig::_internal_cudnn_fusion_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::CuDnnFusionConfig* p = _impl_.cudnn_fusion_config_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_CuDnnFusionConfig_default_instance_); } inline const ::xla::gpu::CuDnnFusionConfig& FusionBackendConfig::cudnn_fusion_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.cudnn_fusion_config) return _internal_cudnn_fusion_config(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_cudnn_fusion_config(::xla::gpu::CuDnnFusionConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cudnn_fusion_config_); } _impl_.cudnn_fusion_config_ = reinterpret_cast<::xla::gpu::CuDnnFusionConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.cudnn_fusion_config) } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::release_cudnn_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::CuDnnFusionConfig* released = _impl_.cudnn_fusion_config_; _impl_.cudnn_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::unsafe_arena_release_cudnn_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.cudnn_fusion_config) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::CuDnnFusionConfig* temp = _impl_.cudnn_fusion_config_; _impl_.cudnn_fusion_config_ = nullptr; return temp; } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::_internal_mutable_cudnn_fusion_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.cudnn_fusion_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CuDnnFusionConfig>(GetArena()); _impl_.cudnn_fusion_config_ = reinterpret_cast<::xla::gpu::CuDnnFusionConfig*>(p); } return _impl_.cudnn_fusion_config_; } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::mutable_cudnn_fusion_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::gpu::CuDnnFusionConfig* _msg = _internal_mutable_cudnn_fusion_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.cudnn_fusion_config) return _msg; } inline void FusionBackendConfig::set_allocated_cudnn_fusion_config(::xla::gpu::CuDnnFusionConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.cudnn_fusion_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } _impl_.cudnn_fusion_config_ = reinterpret_cast<::xla::gpu::CuDnnFusionConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.cudnn_fusion_config) } // .xla.gpu.DynamicMemcpyConfig dynamic_memcpy_config = 7; inline bool FusionBackendConfig::has_dynamic_memcpy_config() const { bool value = (_impl_._has_bits_[0] & 0x00000010u) != 0; PROTOBUF_ASSUME(!value || _impl_.dynamic_memcpy_config_ != nullptr); return value; } inline void FusionBackendConfig::clear_dynamic_memcpy_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.dynamic_memcpy_config_ != nullptr) _impl_.dynamic_memcpy_config_->Clear(); _impl_._has_bits_[0] &= ~0x00000010u; } inline const ::xla::gpu::DynamicMemcpyConfig& FusionBackendConfig::_internal_dynamic_memcpy_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::DynamicMemcpyConfig* p = _impl_.dynamic_memcpy_config_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_DynamicMemcpyConfig_default_instance_); } inline const ::xla::gpu::DynamicMemcpyConfig& FusionBackendConfig::dynamic_memcpy_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.dynamic_memcpy_config) return _internal_dynamic_memcpy_config(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_dynamic_memcpy_config(::xla::gpu::DynamicMemcpyConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dynamic_memcpy_config_); } _impl_.dynamic_memcpy_config_ = reinterpret_cast<::xla::gpu::DynamicMemcpyConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.dynamic_memcpy_config) } inline ::xla::gpu::DynamicMemcpyConfig* FusionBackendConfig::release_dynamic_memcpy_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000010u; ::xla::gpu::DynamicMemcpyConfig* released = _impl_.dynamic_memcpy_config_; _impl_.dynamic_memcpy_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::gpu::DynamicMemcpyConfig* FusionBackendConfig::unsafe_arena_release_dynamic_memcpy_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.dynamic_memcpy_config) _impl_._has_bits_[0] &= ~0x00000010u; ::xla::gpu::DynamicMemcpyConfig* temp = _impl_.dynamic_memcpy_config_; _impl_.dynamic_memcpy_config_ = nullptr; return temp; } inline ::xla::gpu::DynamicMemcpyConfig* FusionBackendConfig::_internal_mutable_dynamic_memcpy_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.dynamic_memcpy_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::DynamicMemcpyConfig>(GetArena()); _impl_.dynamic_memcpy_config_ = reinterpret_cast<::xla::gpu::DynamicMemcpyConfig*>(p); } return _impl_.dynamic_memcpy_config_; } inline ::xla::gpu::DynamicMemcpyConfig* FusionBackendConfig::mutable_dynamic_memcpy_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000010u; ::xla::gpu::DynamicMemcpyConfig* _msg = _internal_mutable_dynamic_memcpy_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.dynamic_memcpy_config) return _msg; } inline void FusionBackendConfig::set_allocated_dynamic_memcpy_config(::xla::gpu::DynamicMemcpyConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.dynamic_memcpy_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } _impl_.dynamic_memcpy_config_ = reinterpret_cast<::xla::gpu::DynamicMemcpyConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.dynamic_memcpy_config) } // ------------------------------------------------------------------- // CudnnNormBackendConfig // double epsilon = 1; inline void CudnnNormBackendConfig::clear_epsilon() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.epsilon_ = 0; } inline double CudnnNormBackendConfig::epsilon() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.epsilon) return _internal_epsilon(); } inline void CudnnNormBackendConfig::set_epsilon(double value) { _internal_set_epsilon(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnNormBackendConfig.epsilon) } inline double CudnnNormBackendConfig::_internal_epsilon() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.epsilon_; } inline void CudnnNormBackendConfig::_internal_set_epsilon(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.epsilon_ = value; } // .stream_executor.dnn.AlgorithmProto algorithm = 2; inline bool CudnnNormBackendConfig::has_algorithm() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.algorithm_ != nullptr); return value; } inline const ::stream_executor::dnn::AlgorithmProto& CudnnNormBackendConfig::_internal_algorithm() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::stream_executor::dnn::AlgorithmProto* p = _impl_.algorithm_; return p != nullptr ? *p : reinterpret_cast(::stream_executor::dnn::_AlgorithmProto_default_instance_); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnNormBackendConfig::algorithm() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnNormBackendConfig::unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnNormBackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::stream_executor::dnn::AlgorithmProto* released = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::unsafe_arena_release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnNormBackendConfig.algorithm) _impl_._has_bits_[0] &= ~0x00000001u; ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::_internal_mutable_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.algorithm_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::stream_executor::dnn::AlgorithmProto>(GetArena()); _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(p); } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::mutable_algorithm() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::stream_executor::dnn::AlgorithmProto* _msg = _internal_mutable_algorithm(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnNormBackendConfig.algorithm) return _msg; } inline void CudnnNormBackendConfig::set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnNormBackendConfig.algorithm) } // .xla.gpu.CudnnNormBackendConfig.Kind kind = 3; inline void CudnnNormBackendConfig::clear_kind() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_ = 0; } inline ::xla::gpu::CudnnNormBackendConfig_Kind CudnnNormBackendConfig::kind() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.kind) return _internal_kind(); } inline void CudnnNormBackendConfig::set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value) { _internal_set_kind(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnNormBackendConfig.kind) } inline ::xla::gpu::CudnnNormBackendConfig_Kind CudnnNormBackendConfig::_internal_kind() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::xla::gpu::CudnnNormBackendConfig_Kind>(_impl_.kind_); } inline void CudnnNormBackendConfig::_internal_set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kind_ = value; } // ------------------------------------------------------------------- // CudnnfMHABackendConfig // .stream_executor.dnn.AlgorithmProto algorithm = 8; inline bool CudnnfMHABackendConfig::has_algorithm() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.algorithm_ != nullptr); return value; } inline const ::stream_executor::dnn::AlgorithmProto& CudnnfMHABackendConfig::_internal_algorithm() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::stream_executor::dnn::AlgorithmProto* p = _impl_.algorithm_; return p != nullptr ? *p : reinterpret_cast(::stream_executor::dnn::_AlgorithmProto_default_instance_); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnfMHABackendConfig::algorithm() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::stream_executor::dnn::AlgorithmProto* released = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::unsafe_arena_release_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.algorithm) _impl_._has_bits_[0] &= ~0x00000001u; ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::_internal_mutable_algorithm() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.algorithm_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::stream_executor::dnn::AlgorithmProto>(GetArena()); _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(p); } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::mutable_algorithm() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::stream_executor::dnn::AlgorithmProto* _msg = _internal_mutable_algorithm(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.algorithm) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.algorithm_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.algorithm_ = reinterpret_cast<::stream_executor::dnn::AlgorithmProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.algorithm) } // double fmha_scale = 10; inline void CudnnfMHABackendConfig::clear_fmha_scale() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.fmha_scale_ = 0; } inline double CudnnfMHABackendConfig::fmha_scale() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.fmha_scale) return _internal_fmha_scale(); } inline void CudnnfMHABackendConfig::set_fmha_scale(double value) { _internal_set_fmha_scale(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.fmha_scale) } inline double CudnnfMHABackendConfig::_internal_fmha_scale() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.fmha_scale_; } inline void CudnnfMHABackendConfig::_internal_set_fmha_scale(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.fmha_scale_ = value; } // double dropout_rate = 13; inline void CudnnfMHABackendConfig::clear_dropout_rate() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.dropout_rate_ = 0; } inline double CudnnfMHABackendConfig::dropout_rate() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.dropout_rate) return _internal_dropout_rate(); } inline void CudnnfMHABackendConfig::set_dropout_rate(double value) { _internal_set_dropout_rate(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.dropout_rate) } inline double CudnnfMHABackendConfig::_internal_dropout_rate() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.dropout_rate_; } inline void CudnnfMHABackendConfig::_internal_set_dropout_rate(double value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.dropout_rate_ = value; } // .xla.DotDimensionNumbers bmm1_dot_dimension_numbers = 11; inline bool CudnnfMHABackendConfig::has_bmm1_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm1_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm1_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm1_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) return _internal_bmm1_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_dot_dimension_numbers_); } _impl_.bmm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::DotDimensionNumbers* released = _impl_.bmm1_dot_dimension_numbers_; _impl_.bmm1_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::DotDimensionNumbers* temp = _impl_.bmm1_dot_dimension_numbers_; _impl_.bmm1_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm1_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm1_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.bmm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) } // .xla.DotDimensionNumbers bmm2_dot_dimension_numbers = 12; inline bool CudnnfMHABackendConfig::has_bmm2_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm2_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm2_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm2_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) return _internal_bmm2_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_dot_dimension_numbers_); } _impl_.bmm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::DotDimensionNumbers* released = _impl_.bmm2_dot_dimension_numbers_; _impl_.bmm2_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::DotDimensionNumbers* temp = _impl_.bmm2_dot_dimension_numbers_; _impl_.bmm2_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm2_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm2_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } _impl_.bmm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) } // .xla.ShapeProto intermediate_tensor_shape = 14; inline bool CudnnfMHABackendConfig::has_intermediate_tensor_shape() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; PROTOBUF_ASSUME(!value || _impl_.intermediate_tensor_shape_ != nullptr); return value; } inline const ::xla::ShapeProto& CudnnfMHABackendConfig::_internal_intermediate_tensor_shape() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ShapeProto* p = _impl_.intermediate_tensor_shape_; return p != nullptr ? *p : reinterpret_cast(::xla::_ShapeProto_default_instance_); } inline const ::xla::ShapeProto& CudnnfMHABackendConfig::intermediate_tensor_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) return _internal_intermediate_tensor_shape(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_intermediate_tensor_shape(::xla::ShapeProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.intermediate_tensor_shape_); } _impl_.intermediate_tensor_shape_ = reinterpret_cast<::xla::ShapeProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) } inline ::xla::ShapeProto* CudnnfMHABackendConfig::release_intermediate_tensor_shape() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000008u; ::xla::ShapeProto* released = _impl_.intermediate_tensor_shape_; _impl_.intermediate_tensor_shape_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::unsafe_arena_release_intermediate_tensor_shape() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) _impl_._has_bits_[0] &= ~0x00000008u; ::xla::ShapeProto* temp = _impl_.intermediate_tensor_shape_; _impl_.intermediate_tensor_shape_ = nullptr; return temp; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::_internal_mutable_intermediate_tensor_shape() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.intermediate_tensor_shape_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena()); _impl_.intermediate_tensor_shape_ = reinterpret_cast<::xla::ShapeProto*>(p); } return _impl_.intermediate_tensor_shape_; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::mutable_intermediate_tensor_shape() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000008u; ::xla::ShapeProto* _msg = _internal_mutable_intermediate_tensor_shape(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_intermediate_tensor_shape(::xla::ShapeProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.intermediate_tensor_shape_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } _impl_.intermediate_tensor_shape_ = reinterpret_cast<::xla::ShapeProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) } // .xla.DotDimensionNumbers bmm1_grad_gemm1_dot_dimension_numbers = 16; inline bool CudnnfMHABackendConfig::has_bmm1_grad_gemm1_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000010u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_grad_gemm1_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm1_grad_gemm1_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm1_grad_gemm1_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) return _internal_bmm1_grad_gemm1_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm1_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_grad_gemm1_dot_dimension_numbers_); } _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm1_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000010u; ::xla::DotDimensionNumbers* released = _impl_.bmm1_grad_gemm1_dot_dimension_numbers_; _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000010u; ::xla::DotDimensionNumbers* temp = _impl_.bmm1_grad_gemm1_dot_dimension_numbers_; _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm1_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm1_grad_gemm1_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm1_grad_gemm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_grad_gemm1_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000010u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm1_grad_gemm1_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm1_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_grad_gemm1_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) } // .xla.DotDimensionNumbers bmm1_grad_gemm2_dot_dimension_numbers = 17; inline bool CudnnfMHABackendConfig::has_bmm1_grad_gemm2_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000020u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_grad_gemm2_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm1_grad_gemm2_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm1_grad_gemm2_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) return _internal_bmm1_grad_gemm2_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm1_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_grad_gemm2_dot_dimension_numbers_); } _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000020u; } else { _impl_._has_bits_[0] &= ~0x00000020u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm1_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000020u; ::xla::DotDimensionNumbers* released = _impl_.bmm1_grad_gemm2_dot_dimension_numbers_; _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000020u; ::xla::DotDimensionNumbers* temp = _impl_.bmm1_grad_gemm2_dot_dimension_numbers_; _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm1_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm1_grad_gemm2_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm1_grad_gemm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_grad_gemm2_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000020u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm1_grad_gemm2_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm1_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm1_grad_gemm2_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000020u; } else { _impl_._has_bits_[0] &= ~0x00000020u; } _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) } // .xla.DotDimensionNumbers bmm2_grad_gemm1_dot_dimension_numbers = 18; inline bool CudnnfMHABackendConfig::has_bmm2_grad_gemm1_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000040u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_grad_gemm1_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm2_grad_gemm1_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm2_grad_gemm1_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) return _internal_bmm2_grad_gemm1_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm2_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_grad_gemm1_dot_dimension_numbers_); } _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000040u; } else { _impl_._has_bits_[0] &= ~0x00000040u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm2_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000040u; ::xla::DotDimensionNumbers* released = _impl_.bmm2_grad_gemm1_dot_dimension_numbers_; _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000040u; ::xla::DotDimensionNumbers* temp = _impl_.bmm2_grad_gemm1_dot_dimension_numbers_; _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm2_grad_gemm1_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm2_grad_gemm1_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm2_grad_gemm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_grad_gemm1_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000040u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm2_grad_gemm1_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm2_grad_gemm1_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_grad_gemm1_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000040u; } else { _impl_._has_bits_[0] &= ~0x00000040u; } _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) } // .xla.DotDimensionNumbers bmm2_grad_gemm2_dot_dimension_numbers = 19; inline bool CudnnfMHABackendConfig::has_bmm2_grad_gemm2_dot_dimension_numbers() const { bool value = (_impl_._has_bits_[0] & 0x00000080u) != 0; PROTOBUF_ASSUME(!value || _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ != nullptr); return value; } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_grad_gemm2_dot_dimension_numbers() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::DotDimensionNumbers* p = _impl_.bmm2_grad_gemm2_dot_dimension_numbers_; return p != nullptr ? *p : reinterpret_cast(::xla::_DotDimensionNumbers_default_instance_); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::bmm2_grad_gemm2_dot_dimension_numbers() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) return _internal_bmm2_grad_gemm2_dot_dimension_numbers(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_bmm2_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_grad_gemm2_dot_dimension_numbers_); } _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000080u; } else { _impl_._has_bits_[0] &= ~0x00000080u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm2_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000080u; ::xla::DotDimensionNumbers* released = _impl_.bmm2_grad_gemm2_dot_dimension_numbers_; _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) _impl_._has_bits_[0] &= ~0x00000080u; ::xla::DotDimensionNumbers* temp = _impl_.bmm2_grad_gemm2_dot_dimension_numbers_; _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::_internal_mutable_bmm2_grad_gemm2_dot_dimension_numbers() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.bmm2_grad_gemm2_dot_dimension_numbers_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::DotDimensionNumbers>(GetArena()); _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(p); } return _impl_.bmm2_grad_gemm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_grad_gemm2_dot_dimension_numbers() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000080u; ::xla::DotDimensionNumbers* _msg = _internal_mutable_bmm2_grad_gemm2_dot_dimension_numbers(); // @@protoc_insertion_point(field_mutable:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) return _msg; } inline void CudnnfMHABackendConfig::set_allocated_bmm2_grad_gemm2_dot_dimension_numbers(::xla::DotDimensionNumbers* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bmm2_grad_gemm2_dot_dimension_numbers_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000080u; } else { _impl_._has_bits_[0] &= ~0x00000080u; } _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = reinterpret_cast<::xla::DotDimensionNumbers*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) } // int64 seed = 15; inline void CudnnfMHABackendConfig::clear_seed() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.seed_ = ::int64_t{0}; } inline ::int64_t CudnnfMHABackendConfig::seed() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.seed) return _internal_seed(); } inline void CudnnfMHABackendConfig::set_seed(::int64_t value) { _internal_set_seed(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.seed) } inline ::int64_t CudnnfMHABackendConfig::_internal_seed() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.seed_; } inline void CudnnfMHABackendConfig::_internal_set_seed(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.seed_ = value; } // bool is_flash_attention = 20; inline void CudnnfMHABackendConfig::clear_is_flash_attention() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_flash_attention_ = false; } inline bool CudnnfMHABackendConfig::is_flash_attention() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.is_flash_attention) return _internal_is_flash_attention(); } inline void CudnnfMHABackendConfig::set_is_flash_attention(bool value) { _internal_set_is_flash_attention(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.is_flash_attention) } inline bool CudnnfMHABackendConfig::_internal_is_flash_attention() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_flash_attention_; } inline void CudnnfMHABackendConfig::_internal_set_is_flash_attention(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_flash_attention_ = value; } // bool is_causal_mask = 21; inline void CudnnfMHABackendConfig::clear_is_causal_mask() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_causal_mask_ = false; } inline bool CudnnfMHABackendConfig::is_causal_mask() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.is_causal_mask) return _internal_is_causal_mask(); } inline void CudnnfMHABackendConfig::set_is_causal_mask(bool value) { _internal_set_is_causal_mask(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.is_causal_mask) } inline bool CudnnfMHABackendConfig::_internal_is_causal_mask() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_causal_mask_; } inline void CudnnfMHABackendConfig::_internal_set_is_causal_mask(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_causal_mask_ = value; } // .xla.gpu.CudnnfMHABackendConfig.MaskType mask_type = 22; inline void CudnnfMHABackendConfig::clear_mask_type() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mask_type_ = 0; } inline ::xla::gpu::CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig::mask_type() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.mask_type) return _internal_mask_type(); } inline void CudnnfMHABackendConfig::set_mask_type(::xla::gpu::CudnnfMHABackendConfig_MaskType value) { _internal_set_mask_type(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.mask_type) } inline ::xla::gpu::CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig::_internal_mask_type() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::xla::gpu::CudnnfMHABackendConfig_MaskType>(_impl_.mask_type_); } inline void CudnnfMHABackendConfig::_internal_set_mask_type(::xla::gpu::CudnnfMHABackendConfig_MaskType value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mask_type_ = value; } // bool force_deterministic = 23; inline void CudnnfMHABackendConfig::clear_force_deterministic() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.force_deterministic_ = false; } inline bool CudnnfMHABackendConfig::force_deterministic() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.force_deterministic) return _internal_force_deterministic(); } inline void CudnnfMHABackendConfig::set_force_deterministic(bool value) { _internal_set_force_deterministic(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.force_deterministic) } inline bool CudnnfMHABackendConfig::_internal_force_deterministic() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.force_deterministic_; } inline void CudnnfMHABackendConfig::_internal_set_force_deterministic(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.force_deterministic_ = value; } // int32 sliding_window_length = 24; inline void CudnnfMHABackendConfig::clear_sliding_window_length() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.sliding_window_length_ = 0; } inline ::int32_t CudnnfMHABackendConfig::sliding_window_length() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.sliding_window_length) return _internal_sliding_window_length(); } inline void CudnnfMHABackendConfig::set_sliding_window_length(::int32_t value) { _internal_set_sliding_window_length(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.sliding_window_length) } inline ::int32_t CudnnfMHABackendConfig::_internal_sliding_window_length() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.sliding_window_length_; } inline void CudnnfMHABackendConfig::_internal_set_sliding_window_length(::int32_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.sliding_window_length_ = value; } // int32 max_seg_per_batch = 25; inline void CudnnfMHABackendConfig::clear_max_seg_per_batch() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.max_seg_per_batch_ = 0; } inline ::int32_t CudnnfMHABackendConfig::max_seg_per_batch() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.max_seg_per_batch) return _internal_max_seg_per_batch(); } inline void CudnnfMHABackendConfig::set_max_seg_per_batch(::int32_t value) { _internal_set_max_seg_per_batch(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.max_seg_per_batch) } inline ::int32_t CudnnfMHABackendConfig::_internal_max_seg_per_batch() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.max_seg_per_batch_; } inline void CudnnfMHABackendConfig::_internal_set_max_seg_per_batch(::int32_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.max_seg_per_batch_ = value; } // bool is_paged_attention = 26; inline void CudnnfMHABackendConfig::clear_is_paged_attention() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_paged_attention_ = false; } inline bool CudnnfMHABackendConfig::is_paged_attention() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.is_paged_attention) return _internal_is_paged_attention(); } inline void CudnnfMHABackendConfig::set_is_paged_attention(bool value) { _internal_set_is_paged_attention(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.is_paged_attention) } inline bool CudnnfMHABackendConfig::_internal_is_paged_attention() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.is_paged_attention_; } inline void CudnnfMHABackendConfig::_internal_set_is_paged_attention(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.is_paged_attention_ = value; } // ------------------------------------------------------------------- // CustomCallBackendConfig // string opaque = 1; inline bool CustomCallBackendConfig::has_opaque() const { return raw_backend_config_oneof_case() == kOpaque; } inline void CustomCallBackendConfig::set_has_opaque() { _impl_._oneof_case_[0] = kOpaque; } inline void CustomCallBackendConfig::clear_opaque() { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() == kOpaque) { _impl_.raw_backend_config_oneof_.opaque_.Destroy(); clear_has_raw_backend_config_oneof(); } } inline const std::string& CustomCallBackendConfig::opaque() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CustomCallBackendConfig.opaque) return _internal_opaque(); } template inline PROTOBUF_ALWAYS_INLINE void CustomCallBackendConfig::set_opaque(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kOpaque) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } _impl_.raw_backend_config_oneof_.opaque_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.CustomCallBackendConfig.opaque) } inline std::string* CustomCallBackendConfig::mutable_opaque() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_opaque(); // @@protoc_insertion_point(field_mutable:xla.gpu.CustomCallBackendConfig.opaque) return _s; } inline const std::string& CustomCallBackendConfig::_internal_opaque() const { ::google::protobuf::internal::TSanRead(&_impl_); if (raw_backend_config_oneof_case() != kOpaque) { return ::google::protobuf::internal::GetEmptyStringAlreadyInited(); } return _impl_.raw_backend_config_oneof_.opaque_.Get(); } inline void CustomCallBackendConfig::_internal_set_opaque(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kOpaque) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } _impl_.raw_backend_config_oneof_.opaque_.Set(value, GetArena()); } inline std::string* CustomCallBackendConfig::_internal_mutable_opaque() { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kOpaque) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } return _impl_.raw_backend_config_oneof_.opaque_.Mutable( GetArena()); } inline std::string* CustomCallBackendConfig::release_opaque() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CustomCallBackendConfig.opaque) if (raw_backend_config_oneof_case() != kOpaque) { return nullptr; } clear_has_raw_backend_config_oneof(); return _impl_.raw_backend_config_oneof_.opaque_.Release(); } inline void CustomCallBackendConfig::set_allocated_opaque(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (has_raw_backend_config_oneof()) { clear_raw_backend_config_oneof(); } if (value != nullptr) { set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitAllocated(value, GetArena()); } // @@protoc_insertion_point(field_set_allocated:xla.gpu.CustomCallBackendConfig.opaque) } // string attributes = 2; inline bool CustomCallBackendConfig::has_attributes() const { return raw_backend_config_oneof_case() == kAttributes; } inline void CustomCallBackendConfig::set_has_attributes() { _impl_._oneof_case_[0] = kAttributes; } inline void CustomCallBackendConfig::clear_attributes() { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() == kAttributes) { _impl_.raw_backend_config_oneof_.attributes_.Destroy(); clear_has_raw_backend_config_oneof(); } } inline const std::string& CustomCallBackendConfig::attributes() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CustomCallBackendConfig.attributes) return _internal_attributes(); } template inline PROTOBUF_ALWAYS_INLINE void CustomCallBackendConfig::set_attributes(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kAttributes) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } _impl_.raw_backend_config_oneof_.attributes_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.CustomCallBackendConfig.attributes) } inline std::string* CustomCallBackendConfig::mutable_attributes() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_attributes(); // @@protoc_insertion_point(field_mutable:xla.gpu.CustomCallBackendConfig.attributes) return _s; } inline const std::string& CustomCallBackendConfig::_internal_attributes() const { ::google::protobuf::internal::TSanRead(&_impl_); if (raw_backend_config_oneof_case() != kAttributes) { return ::google::protobuf::internal::GetEmptyStringAlreadyInited(); } return _impl_.raw_backend_config_oneof_.attributes_.Get(); } inline void CustomCallBackendConfig::_internal_set_attributes(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kAttributes) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } _impl_.raw_backend_config_oneof_.attributes_.Set(value, GetArena()); } inline std::string* CustomCallBackendConfig::_internal_mutable_attributes() { ::google::protobuf::internal::TSanWrite(&_impl_); if (raw_backend_config_oneof_case() != kAttributes) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } return _impl_.raw_backend_config_oneof_.attributes_.Mutable( GetArena()); } inline std::string* CustomCallBackendConfig::release_attributes() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CustomCallBackendConfig.attributes) if (raw_backend_config_oneof_case() != kAttributes) { return nullptr; } clear_has_raw_backend_config_oneof(); return _impl_.raw_backend_config_oneof_.attributes_.Release(); } inline void CustomCallBackendConfig::set_allocated_attributes(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (has_raw_backend_config_oneof()) { clear_raw_backend_config_oneof(); } if (value != nullptr) { set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitAllocated(value, GetArena()); } // @@protoc_insertion_point(field_set_allocated:xla.gpu.CustomCallBackendConfig.attributes) } inline bool CustomCallBackendConfig::has_raw_backend_config_oneof() const { return raw_backend_config_oneof_case() != RAW_BACKEND_CONFIG_ONEOF_NOT_SET; } inline void CustomCallBackendConfig::clear_has_raw_backend_config_oneof() { _impl_._oneof_case_[0] = RAW_BACKEND_CONFIG_ONEOF_NOT_SET; } inline CustomCallBackendConfig::RawBackendConfigOneofCase CustomCallBackendConfig::raw_backend_config_oneof_case() const { return CustomCallBackendConfig::RawBackendConfigOneofCase(_impl_._oneof_case_[0]); } // ------------------------------------------------------------------- // GpuBackendConfig // int64 operation_queue_id = 1; inline void GpuBackendConfig::clear_operation_queue_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.operation_queue_id_ = ::int64_t{0}; } inline ::int64_t GpuBackendConfig::operation_queue_id() const { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.operation_queue_id) return _internal_operation_queue_id(); } inline void GpuBackendConfig::set_operation_queue_id(::int64_t value) { _internal_set_operation_queue_id(value); // @@protoc_insertion_point(field_set:xla.gpu.GpuBackendConfig.operation_queue_id) } inline ::int64_t GpuBackendConfig::_internal_operation_queue_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.operation_queue_id_; } inline void GpuBackendConfig::_internal_set_operation_queue_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.operation_queue_id_ = value; } // repeated int64 wait_on_operation_queues = 2; inline int GpuBackendConfig::_internal_wait_on_operation_queues_size() const { return _internal_wait_on_operation_queues().size(); } inline int GpuBackendConfig::wait_on_operation_queues_size() const { return _internal_wait_on_operation_queues_size(); } inline void GpuBackendConfig::clear_wait_on_operation_queues() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.wait_on_operation_queues_.Clear(); } inline ::int64_t GpuBackendConfig::wait_on_operation_queues(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.wait_on_operation_queues) return _internal_wait_on_operation_queues().Get(index); } inline void GpuBackendConfig::set_wait_on_operation_queues(int index, ::int64_t value) { _internal_mutable_wait_on_operation_queues()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.GpuBackendConfig.wait_on_operation_queues) } inline void GpuBackendConfig::add_wait_on_operation_queues(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_wait_on_operation_queues()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.GpuBackendConfig.wait_on_operation_queues) } inline const ::google::protobuf::RepeatedField<::int64_t>& GpuBackendConfig::wait_on_operation_queues() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.GpuBackendConfig.wait_on_operation_queues) return _internal_wait_on_operation_queues(); } inline ::google::protobuf::RepeatedField<::int64_t>* GpuBackendConfig::mutable_wait_on_operation_queues() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.GpuBackendConfig.wait_on_operation_queues) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_wait_on_operation_queues(); } inline const ::google::protobuf::RepeatedField<::int64_t>& GpuBackendConfig::_internal_wait_on_operation_queues() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.wait_on_operation_queues_; } inline ::google::protobuf::RepeatedField<::int64_t>* GpuBackendConfig::_internal_mutable_wait_on_operation_queues() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.wait_on_operation_queues_; } // .xla.gpu.CudnnConvBackendConfig cudnn_conv_backend_config = 3; inline bool GpuBackendConfig::has_cudnn_conv_backend_config() const { return backend_config_case() == kCudnnConvBackendConfig; } inline bool GpuBackendConfig::_internal_has_cudnn_conv_backend_config() const { return backend_config_case() == kCudnnConvBackendConfig; } inline void GpuBackendConfig::set_has_cudnn_conv_backend_config() { _impl_._oneof_case_[0] = kCudnnConvBackendConfig; } inline void GpuBackendConfig::clear_cudnn_conv_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kCudnnConvBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.cudnn_conv_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.cudnn_conv_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::CudnnConvBackendConfig* GpuBackendConfig::release_cudnn_conv_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.cudnn_conv_backend_config) if (backend_config_case() == kCudnnConvBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_conv_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.cudnn_conv_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CudnnConvBackendConfig& GpuBackendConfig::_internal_cudnn_conv_backend_config() const { return backend_config_case() == kCudnnConvBackendConfig ? *_impl_.backend_config_.cudnn_conv_backend_config_ : reinterpret_cast<::xla::gpu::CudnnConvBackendConfig&>(::xla::gpu::_CudnnConvBackendConfig_default_instance_); } inline const ::xla::gpu::CudnnConvBackendConfig& GpuBackendConfig::cudnn_conv_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.cudnn_conv_backend_config) return _internal_cudnn_conv_backend_config(); } inline ::xla::gpu::CudnnConvBackendConfig* GpuBackendConfig::unsafe_arena_release_cudnn_conv_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.cudnn_conv_backend_config) if (backend_config_case() == kCudnnConvBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_conv_backend_config_; _impl_.backend_config_.cudnn_conv_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_cudnn_conv_backend_config(::xla::gpu::CudnnConvBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_cudnn_conv_backend_config(); _impl_.backend_config_.cudnn_conv_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.cudnn_conv_backend_config) } inline ::xla::gpu::CudnnConvBackendConfig* GpuBackendConfig::_internal_mutable_cudnn_conv_backend_config() { if (backend_config_case() != kCudnnConvBackendConfig) { clear_backend_config(); set_has_cudnn_conv_backend_config(); _impl_.backend_config_.cudnn_conv_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CudnnConvBackendConfig>(GetArena()); } return _impl_.backend_config_.cudnn_conv_backend_config_; } inline ::xla::gpu::CudnnConvBackendConfig* GpuBackendConfig::mutable_cudnn_conv_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CudnnConvBackendConfig* _msg = _internal_mutable_cudnn_conv_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.cudnn_conv_backend_config) return _msg; } // .xla.gpu.GemmBackendConfig gemm_backend_config = 4; inline bool GpuBackendConfig::has_gemm_backend_config() const { return backend_config_case() == kGemmBackendConfig; } inline bool GpuBackendConfig::_internal_has_gemm_backend_config() const { return backend_config_case() == kGemmBackendConfig; } inline void GpuBackendConfig::set_has_gemm_backend_config() { _impl_._oneof_case_[0] = kGemmBackendConfig; } inline void GpuBackendConfig::clear_gemm_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kGemmBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.gemm_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.gemm_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::GemmBackendConfig* GpuBackendConfig::release_gemm_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.gemm_backend_config) if (backend_config_case() == kGemmBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.gemm_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.gemm_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::GemmBackendConfig& GpuBackendConfig::_internal_gemm_backend_config() const { return backend_config_case() == kGemmBackendConfig ? *_impl_.backend_config_.gemm_backend_config_ : reinterpret_cast<::xla::gpu::GemmBackendConfig&>(::xla::gpu::_GemmBackendConfig_default_instance_); } inline const ::xla::gpu::GemmBackendConfig& GpuBackendConfig::gemm_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.gemm_backend_config) return _internal_gemm_backend_config(); } inline ::xla::gpu::GemmBackendConfig* GpuBackendConfig::unsafe_arena_release_gemm_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.gemm_backend_config) if (backend_config_case() == kGemmBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.gemm_backend_config_; _impl_.backend_config_.gemm_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_gemm_backend_config(::xla::gpu::GemmBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_gemm_backend_config(); _impl_.backend_config_.gemm_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.gemm_backend_config) } inline ::xla::gpu::GemmBackendConfig* GpuBackendConfig::_internal_mutable_gemm_backend_config() { if (backend_config_case() != kGemmBackendConfig) { clear_backend_config(); set_has_gemm_backend_config(); _impl_.backend_config_.gemm_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::GemmBackendConfig>(GetArena()); } return _impl_.backend_config_.gemm_backend_config_; } inline ::xla::gpu::GemmBackendConfig* GpuBackendConfig::mutable_gemm_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::GemmBackendConfig* _msg = _internal_mutable_gemm_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.gemm_backend_config) return _msg; } // .xla.gpu.BitcastBackendConfig bitcast_backend_config = 5; inline bool GpuBackendConfig::has_bitcast_backend_config() const { return backend_config_case() == kBitcastBackendConfig; } inline bool GpuBackendConfig::_internal_has_bitcast_backend_config() const { return backend_config_case() == kBitcastBackendConfig; } inline void GpuBackendConfig::set_has_bitcast_backend_config() { _impl_._oneof_case_[0] = kBitcastBackendConfig; } inline void GpuBackendConfig::clear_bitcast_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kBitcastBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.bitcast_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.bitcast_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::BitcastBackendConfig* GpuBackendConfig::release_bitcast_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.bitcast_backend_config) if (backend_config_case() == kBitcastBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.bitcast_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.bitcast_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::BitcastBackendConfig& GpuBackendConfig::_internal_bitcast_backend_config() const { return backend_config_case() == kBitcastBackendConfig ? *_impl_.backend_config_.bitcast_backend_config_ : reinterpret_cast<::xla::gpu::BitcastBackendConfig&>(::xla::gpu::_BitcastBackendConfig_default_instance_); } inline const ::xla::gpu::BitcastBackendConfig& GpuBackendConfig::bitcast_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.bitcast_backend_config) return _internal_bitcast_backend_config(); } inline ::xla::gpu::BitcastBackendConfig* GpuBackendConfig::unsafe_arena_release_bitcast_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.bitcast_backend_config) if (backend_config_case() == kBitcastBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.bitcast_backend_config_; _impl_.backend_config_.bitcast_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_bitcast_backend_config(::xla::gpu::BitcastBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_bitcast_backend_config(); _impl_.backend_config_.bitcast_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.bitcast_backend_config) } inline ::xla::gpu::BitcastBackendConfig* GpuBackendConfig::_internal_mutable_bitcast_backend_config() { if (backend_config_case() != kBitcastBackendConfig) { clear_backend_config(); set_has_bitcast_backend_config(); _impl_.backend_config_.bitcast_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::BitcastBackendConfig>(GetArena()); } return _impl_.backend_config_.bitcast_backend_config_; } inline ::xla::gpu::BitcastBackendConfig* GpuBackendConfig::mutable_bitcast_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::BitcastBackendConfig* _msg = _internal_mutable_bitcast_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.bitcast_backend_config) return _msg; } // .xla.gpu.CollectiveBackendConfig collective_backend_config = 6; inline bool GpuBackendConfig::has_collective_backend_config() const { return backend_config_case() == kCollectiveBackendConfig; } inline bool GpuBackendConfig::_internal_has_collective_backend_config() const { return backend_config_case() == kCollectiveBackendConfig; } inline void GpuBackendConfig::set_has_collective_backend_config() { _impl_._oneof_case_[0] = kCollectiveBackendConfig; } inline void GpuBackendConfig::clear_collective_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kCollectiveBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.collective_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.collective_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::CollectiveBackendConfig* GpuBackendConfig::release_collective_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.collective_backend_config) if (backend_config_case() == kCollectiveBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.collective_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.collective_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CollectiveBackendConfig& GpuBackendConfig::_internal_collective_backend_config() const { return backend_config_case() == kCollectiveBackendConfig ? *_impl_.backend_config_.collective_backend_config_ : reinterpret_cast<::xla::gpu::CollectiveBackendConfig&>(::xla::gpu::_CollectiveBackendConfig_default_instance_); } inline const ::xla::gpu::CollectiveBackendConfig& GpuBackendConfig::collective_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.collective_backend_config) return _internal_collective_backend_config(); } inline ::xla::gpu::CollectiveBackendConfig* GpuBackendConfig::unsafe_arena_release_collective_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.collective_backend_config) if (backend_config_case() == kCollectiveBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.collective_backend_config_; _impl_.backend_config_.collective_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_collective_backend_config(::xla::gpu::CollectiveBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_collective_backend_config(); _impl_.backend_config_.collective_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.collective_backend_config) } inline ::xla::gpu::CollectiveBackendConfig* GpuBackendConfig::_internal_mutable_collective_backend_config() { if (backend_config_case() != kCollectiveBackendConfig) { clear_backend_config(); set_has_collective_backend_config(); _impl_.backend_config_.collective_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CollectiveBackendConfig>(GetArena()); } return _impl_.backend_config_.collective_backend_config_; } inline ::xla::gpu::CollectiveBackendConfig* GpuBackendConfig::mutable_collective_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CollectiveBackendConfig* _msg = _internal_mutable_collective_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.collective_backend_config) return _msg; } // .xla.gpu.FusionBackendConfig fusion_backend_config = 7; inline bool GpuBackendConfig::has_fusion_backend_config() const { return backend_config_case() == kFusionBackendConfig; } inline bool GpuBackendConfig::_internal_has_fusion_backend_config() const { return backend_config_case() == kFusionBackendConfig; } inline void GpuBackendConfig::set_has_fusion_backend_config() { _impl_._oneof_case_[0] = kFusionBackendConfig; } inline void GpuBackendConfig::clear_fusion_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kFusionBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.fusion_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.fusion_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::FusionBackendConfig* GpuBackendConfig::release_fusion_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.fusion_backend_config) if (backend_config_case() == kFusionBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.fusion_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.fusion_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::FusionBackendConfig& GpuBackendConfig::_internal_fusion_backend_config() const { return backend_config_case() == kFusionBackendConfig ? *_impl_.backend_config_.fusion_backend_config_ : reinterpret_cast<::xla::gpu::FusionBackendConfig&>(::xla::gpu::_FusionBackendConfig_default_instance_); } inline const ::xla::gpu::FusionBackendConfig& GpuBackendConfig::fusion_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.fusion_backend_config) return _internal_fusion_backend_config(); } inline ::xla::gpu::FusionBackendConfig* GpuBackendConfig::unsafe_arena_release_fusion_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.fusion_backend_config) if (backend_config_case() == kFusionBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.fusion_backend_config_; _impl_.backend_config_.fusion_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_fusion_backend_config(::xla::gpu::FusionBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_fusion_backend_config(); _impl_.backend_config_.fusion_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.fusion_backend_config) } inline ::xla::gpu::FusionBackendConfig* GpuBackendConfig::_internal_mutable_fusion_backend_config() { if (backend_config_case() != kFusionBackendConfig) { clear_backend_config(); set_has_fusion_backend_config(); _impl_.backend_config_.fusion_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::FusionBackendConfig>(GetArena()); } return _impl_.backend_config_.fusion_backend_config_; } inline ::xla::gpu::FusionBackendConfig* GpuBackendConfig::mutable_fusion_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::FusionBackendConfig* _msg = _internal_mutable_fusion_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.fusion_backend_config) return _msg; } // .xla.gpu.CudnnNormBackendConfig cudnn_norm_backend_config = 8; inline bool GpuBackendConfig::has_cudnn_norm_backend_config() const { return backend_config_case() == kCudnnNormBackendConfig; } inline bool GpuBackendConfig::_internal_has_cudnn_norm_backend_config() const { return backend_config_case() == kCudnnNormBackendConfig; } inline void GpuBackendConfig::set_has_cudnn_norm_backend_config() { _impl_._oneof_case_[0] = kCudnnNormBackendConfig; } inline void GpuBackendConfig::clear_cudnn_norm_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kCudnnNormBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.cudnn_norm_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.cudnn_norm_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::CudnnNormBackendConfig* GpuBackendConfig::release_cudnn_norm_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.cudnn_norm_backend_config) if (backend_config_case() == kCudnnNormBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_norm_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.cudnn_norm_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CudnnNormBackendConfig& GpuBackendConfig::_internal_cudnn_norm_backend_config() const { return backend_config_case() == kCudnnNormBackendConfig ? *_impl_.backend_config_.cudnn_norm_backend_config_ : reinterpret_cast<::xla::gpu::CudnnNormBackendConfig&>(::xla::gpu::_CudnnNormBackendConfig_default_instance_); } inline const ::xla::gpu::CudnnNormBackendConfig& GpuBackendConfig::cudnn_norm_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.cudnn_norm_backend_config) return _internal_cudnn_norm_backend_config(); } inline ::xla::gpu::CudnnNormBackendConfig* GpuBackendConfig::unsafe_arena_release_cudnn_norm_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.cudnn_norm_backend_config) if (backend_config_case() == kCudnnNormBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_norm_backend_config_; _impl_.backend_config_.cudnn_norm_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_cudnn_norm_backend_config(::xla::gpu::CudnnNormBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_cudnn_norm_backend_config(); _impl_.backend_config_.cudnn_norm_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.cudnn_norm_backend_config) } inline ::xla::gpu::CudnnNormBackendConfig* GpuBackendConfig::_internal_mutable_cudnn_norm_backend_config() { if (backend_config_case() != kCudnnNormBackendConfig) { clear_backend_config(); set_has_cudnn_norm_backend_config(); _impl_.backend_config_.cudnn_norm_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CudnnNormBackendConfig>(GetArena()); } return _impl_.backend_config_.cudnn_norm_backend_config_; } inline ::xla::gpu::CudnnNormBackendConfig* GpuBackendConfig::mutable_cudnn_norm_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CudnnNormBackendConfig* _msg = _internal_mutable_cudnn_norm_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.cudnn_norm_backend_config) return _msg; } // .xla.gpu.CudnnfMHABackendConfig cudnn_fmha_backend_config = 9; inline bool GpuBackendConfig::has_cudnn_fmha_backend_config() const { return backend_config_case() == kCudnnFmhaBackendConfig; } inline bool GpuBackendConfig::_internal_has_cudnn_fmha_backend_config() const { return backend_config_case() == kCudnnFmhaBackendConfig; } inline void GpuBackendConfig::set_has_cudnn_fmha_backend_config() { _impl_._oneof_case_[0] = kCudnnFmhaBackendConfig; } inline void GpuBackendConfig::clear_cudnn_fmha_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kCudnnFmhaBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.cudnn_fmha_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.cudnn_fmha_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::CudnnfMHABackendConfig* GpuBackendConfig::release_cudnn_fmha_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.cudnn_fmha_backend_config) if (backend_config_case() == kCudnnFmhaBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_fmha_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.cudnn_fmha_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CudnnfMHABackendConfig& GpuBackendConfig::_internal_cudnn_fmha_backend_config() const { return backend_config_case() == kCudnnFmhaBackendConfig ? *_impl_.backend_config_.cudnn_fmha_backend_config_ : reinterpret_cast<::xla::gpu::CudnnfMHABackendConfig&>(::xla::gpu::_CudnnfMHABackendConfig_default_instance_); } inline const ::xla::gpu::CudnnfMHABackendConfig& GpuBackendConfig::cudnn_fmha_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.cudnn_fmha_backend_config) return _internal_cudnn_fmha_backend_config(); } inline ::xla::gpu::CudnnfMHABackendConfig* GpuBackendConfig::unsafe_arena_release_cudnn_fmha_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.cudnn_fmha_backend_config) if (backend_config_case() == kCudnnFmhaBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.cudnn_fmha_backend_config_; _impl_.backend_config_.cudnn_fmha_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_cudnn_fmha_backend_config(::xla::gpu::CudnnfMHABackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_cudnn_fmha_backend_config(); _impl_.backend_config_.cudnn_fmha_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.cudnn_fmha_backend_config) } inline ::xla::gpu::CudnnfMHABackendConfig* GpuBackendConfig::_internal_mutable_cudnn_fmha_backend_config() { if (backend_config_case() != kCudnnFmhaBackendConfig) { clear_backend_config(); set_has_cudnn_fmha_backend_config(); _impl_.backend_config_.cudnn_fmha_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CudnnfMHABackendConfig>(GetArena()); } return _impl_.backend_config_.cudnn_fmha_backend_config_; } inline ::xla::gpu::CudnnfMHABackendConfig* GpuBackendConfig::mutable_cudnn_fmha_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CudnnfMHABackendConfig* _msg = _internal_mutable_cudnn_fmha_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.cudnn_fmha_backend_config) return _msg; } // .xla.gpu.CustomCallBackendConfig custom_call_backend_config = 11; inline bool GpuBackendConfig::has_custom_call_backend_config() const { return backend_config_case() == kCustomCallBackendConfig; } inline bool GpuBackendConfig::_internal_has_custom_call_backend_config() const { return backend_config_case() == kCustomCallBackendConfig; } inline void GpuBackendConfig::set_has_custom_call_backend_config() { _impl_._oneof_case_[0] = kCustomCallBackendConfig; } inline void GpuBackendConfig::clear_custom_call_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (backend_config_case() == kCustomCallBackendConfig) { if (GetArena() == nullptr) { delete _impl_.backend_config_.custom_call_backend_config_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.backend_config_.custom_call_backend_config_); } clear_has_backend_config(); } } inline ::xla::gpu::CustomCallBackendConfig* GpuBackendConfig::release_custom_call_backend_config() { // @@protoc_insertion_point(field_release:xla.gpu.GpuBackendConfig.custom_call_backend_config) if (backend_config_case() == kCustomCallBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.custom_call_backend_config_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.backend_config_.custom_call_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CustomCallBackendConfig& GpuBackendConfig::_internal_custom_call_backend_config() const { return backend_config_case() == kCustomCallBackendConfig ? *_impl_.backend_config_.custom_call_backend_config_ : reinterpret_cast<::xla::gpu::CustomCallBackendConfig&>(::xla::gpu::_CustomCallBackendConfig_default_instance_); } inline const ::xla::gpu::CustomCallBackendConfig& GpuBackendConfig::custom_call_backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.custom_call_backend_config) return _internal_custom_call_backend_config(); } inline ::xla::gpu::CustomCallBackendConfig* GpuBackendConfig::unsafe_arena_release_custom_call_backend_config() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.GpuBackendConfig.custom_call_backend_config) if (backend_config_case() == kCustomCallBackendConfig) { clear_has_backend_config(); auto* temp = _impl_.backend_config_.custom_call_backend_config_; _impl_.backend_config_.custom_call_backend_config_ = nullptr; return temp; } else { return nullptr; } } inline void GpuBackendConfig::unsafe_arena_set_allocated_custom_call_backend_config(::xla::gpu::CustomCallBackendConfig* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_backend_config(); if (value) { set_has_custom_call_backend_config(); _impl_.backend_config_.custom_call_backend_config_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuBackendConfig.custom_call_backend_config) } inline ::xla::gpu::CustomCallBackendConfig* GpuBackendConfig::_internal_mutable_custom_call_backend_config() { if (backend_config_case() != kCustomCallBackendConfig) { clear_backend_config(); set_has_custom_call_backend_config(); _impl_.backend_config_.custom_call_backend_config_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CustomCallBackendConfig>(GetArena()); } return _impl_.backend_config_.custom_call_backend_config_; } inline ::xla::gpu::CustomCallBackendConfig* GpuBackendConfig::mutable_custom_call_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CustomCallBackendConfig* _msg = _internal_mutable_custom_call_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.custom_call_backend_config) return _msg; } // bool force_earliest_schedule = 10; inline void GpuBackendConfig::clear_force_earliest_schedule() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.force_earliest_schedule_ = false; } inline bool GpuBackendConfig::force_earliest_schedule() const { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.force_earliest_schedule) return _internal_force_earliest_schedule(); } inline void GpuBackendConfig::set_force_earliest_schedule(bool value) { _internal_set_force_earliest_schedule(value); // @@protoc_insertion_point(field_set:xla.gpu.GpuBackendConfig.force_earliest_schedule) } inline bool GpuBackendConfig::_internal_force_earliest_schedule() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.force_earliest_schedule_; } inline void GpuBackendConfig::_internal_set_force_earliest_schedule(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.force_earliest_schedule_ = value; } // repeated .xla.gpu.ReificationCost reification_cost = 12; inline int GpuBackendConfig::_internal_reification_cost_size() const { return _internal_reification_cost().size(); } inline int GpuBackendConfig::reification_cost_size() const { return _internal_reification_cost_size(); } inline void GpuBackendConfig::clear_reification_cost() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.reification_cost_.Clear(); } inline ::xla::gpu::ReificationCost* GpuBackendConfig::mutable_reification_cost(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.GpuBackendConfig.reification_cost) return _internal_mutable_reification_cost()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>* GpuBackendConfig::mutable_reification_cost() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.GpuBackendConfig.reification_cost) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_reification_cost(); } inline const ::xla::gpu::ReificationCost& GpuBackendConfig::reification_cost(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GpuBackendConfig.reification_cost) return _internal_reification_cost().Get(index); } inline ::xla::gpu::ReificationCost* GpuBackendConfig::add_reification_cost() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::ReificationCost* _add = _internal_mutable_reification_cost()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.GpuBackendConfig.reification_cost) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>& GpuBackendConfig::reification_cost() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.GpuBackendConfig.reification_cost) return _internal_reification_cost(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>& GpuBackendConfig::_internal_reification_cost() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.reification_cost_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::ReificationCost>* GpuBackendConfig::_internal_mutable_reification_cost() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.reification_cost_; } inline bool GpuBackendConfig::has_backend_config() const { return backend_config_case() != BACKEND_CONFIG_NOT_SET; } inline void GpuBackendConfig::clear_has_backend_config() { _impl_._oneof_case_[0] = BACKEND_CONFIG_NOT_SET; } inline GpuBackendConfig::BackendConfigCase GpuBackendConfig::backend_config_case() const { return GpuBackendConfig::BackendConfigCase(_impl_._oneof_case_[0]); } #ifdef __GNUC__ #pragma GCC diagnostic pop #endif // __GNUC__ // @@protoc_insertion_point(namespace_scope) } // namespace gpu } // namespace xla namespace google { namespace protobuf { template <> struct is_proto_enum<::xla::gpu::GemmBackendConfig_Epilogue> : std::true_type {}; template <> inline const EnumDescriptor* GetEnumDescriptor<::xla::gpu::GemmBackendConfig_Epilogue>() { return ::xla::gpu::GemmBackendConfig_Epilogue_descriptor(); } template <> struct is_proto_enum<::xla::gpu::CollectiveBackendConfig_CollectiveBackend> : std::true_type {}; template <> inline const EnumDescriptor* GetEnumDescriptor<::xla::gpu::CollectiveBackendConfig_CollectiveBackend>() { return ::xla::gpu::CollectiveBackendConfig_CollectiveBackend_descriptor(); } template <> struct is_proto_enum<::xla::gpu::CudnnNormBackendConfig_Kind> : std::true_type {}; template <> inline const EnumDescriptor* GetEnumDescriptor<::xla::gpu::CudnnNormBackendConfig_Kind>() { return ::xla::gpu::CudnnNormBackendConfig_Kind_descriptor(); } template <> struct is_proto_enum<::xla::gpu::CudnnfMHABackendConfig_MaskType> : std::true_type {}; template <> inline const EnumDescriptor* GetEnumDescriptor<::xla::gpu::CudnnfMHABackendConfig_MaskType>() { return ::xla::gpu::CudnnfMHABackendConfig_MaskType_descriptor(); } } // namespace protobuf } // namespace google // @@protoc_insertion_point(global_scope) #include "google/protobuf/port_undef.inc" #endif // GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto_2epb_2eh