// Generated by the protocol buffer compiler. DO NOT EDIT! // source: xla/service/gpu/backend_configs.proto #ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto #define GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto #include #include #include #if PROTOBUF_VERSION < 3021000 #error This file was generated by a newer version of protoc which is #error incompatible with your Protocol Buffer headers. Please update #error your headers. #endif #if 3021009 < PROTOBUF_MIN_PROTOC_VERSION #error This file was generated by an older version of protoc which is #error incompatible with your Protocol Buffer headers. Please #error regenerate this file with a newer version of protoc. #endif #include #include #include #include #include #include #include #include #include // IWYU pragma: export #include // IWYU pragma: export #include #include #include "xla/autotuning.pb.h" #include "xla/tsl/protobuf/dnn.pb.h" #include "xla/xla_data.pb.h" // @@protoc_insertion_point(includes) #include #define PROTOBUF_INTERNAL_EXPORT_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto PROTOBUF_NAMESPACE_OPEN namespace internal { class AnyMetadata; } // namespace internal PROTOBUF_NAMESPACE_CLOSE // Internal implementation detail -- do not use these members. struct TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto { static const uint32_t offsets[]; }; extern const ::PROTOBUF_NAMESPACE_ID::internal::DescriptorTable descriptor_table_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; namespace xla { namespace gpu { class BitcastBackendConfig; struct BitcastBackendConfigDefaultTypeInternal; extern BitcastBackendConfigDefaultTypeInternal _BitcastBackendConfig_default_instance_; class BlockLevelFusionConfig; struct BlockLevelFusionConfigDefaultTypeInternal; extern BlockLevelFusionConfigDefaultTypeInternal _BlockLevelFusionConfig_default_instance_; class CollectiveBackendConfig; struct CollectiveBackendConfigDefaultTypeInternal; extern CollectiveBackendConfigDefaultTypeInternal _CollectiveBackendConfig_default_instance_; class CuDnnFusionConfig; struct CuDnnFusionConfigDefaultTypeInternal; extern CuDnnFusionConfigDefaultTypeInternal _CuDnnFusionConfig_default_instance_; class CudnnConvBackendConfig; struct CudnnConvBackendConfigDefaultTypeInternal; extern CudnnConvBackendConfigDefaultTypeInternal _CudnnConvBackendConfig_default_instance_; class CudnnNormBackendConfig; struct CudnnNormBackendConfigDefaultTypeInternal; extern CudnnNormBackendConfigDefaultTypeInternal _CudnnNormBackendConfig_default_instance_; class CudnnfMHABackendConfig; struct CudnnfMHABackendConfigDefaultTypeInternal; extern CudnnfMHABackendConfigDefaultTypeInternal _CudnnfMHABackendConfig_default_instance_; class CustomCallBackendConfig; struct CustomCallBackendConfigDefaultTypeInternal; extern CustomCallBackendConfigDefaultTypeInternal _CustomCallBackendConfig_default_instance_; class CustomFusionConfig; struct CustomFusionConfigDefaultTypeInternal; extern CustomFusionConfigDefaultTypeInternal _CustomFusionConfig_default_instance_; class FusionBackendConfig; struct FusionBackendConfigDefaultTypeInternal; extern FusionBackendConfigDefaultTypeInternal _FusionBackendConfig_default_instance_; class GemmBackendConfig; struct GemmBackendConfigDefaultTypeInternal; extern GemmBackendConfigDefaultTypeInternal _GemmBackendConfig_default_instance_; class GpuBackendConfig; struct GpuBackendConfigDefaultTypeInternal; extern GpuBackendConfigDefaultTypeInternal _GpuBackendConfig_default_instance_; class ReificationCost; struct ReificationCostDefaultTypeInternal; extern ReificationCostDefaultTypeInternal _ReificationCost_default_instance_; } // namespace gpu } // namespace xla PROTOBUF_NAMESPACE_OPEN template<> ::xla::gpu::BitcastBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::BitcastBackendConfig>(Arena*); template<> ::xla::gpu::BlockLevelFusionConfig* Arena::CreateMaybeMessage<::xla::gpu::BlockLevelFusionConfig>(Arena*); template<> ::xla::gpu::CollectiveBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::CollectiveBackendConfig>(Arena*); template<> ::xla::gpu::CuDnnFusionConfig* Arena::CreateMaybeMessage<::xla::gpu::CuDnnFusionConfig>(Arena*); template<> ::xla::gpu::CudnnConvBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::CudnnConvBackendConfig>(Arena*); template<> ::xla::gpu::CudnnNormBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::CudnnNormBackendConfig>(Arena*); template<> ::xla::gpu::CudnnfMHABackendConfig* Arena::CreateMaybeMessage<::xla::gpu::CudnnfMHABackendConfig>(Arena*); template<> ::xla::gpu::CustomCallBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::CustomCallBackendConfig>(Arena*); template<> ::xla::gpu::CustomFusionConfig* Arena::CreateMaybeMessage<::xla::gpu::CustomFusionConfig>(Arena*); template<> ::xla::gpu::FusionBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::FusionBackendConfig>(Arena*); template<> ::xla::gpu::GemmBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::GemmBackendConfig>(Arena*); template<> ::xla::gpu::GpuBackendConfig* Arena::CreateMaybeMessage<::xla::gpu::GpuBackendConfig>(Arena*); template<> ::xla::gpu::ReificationCost* Arena::CreateMaybeMessage<::xla::gpu::ReificationCost>(Arena*); PROTOBUF_NAMESPACE_CLOSE 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_GemmBackendConfig_Epilogue_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits::min(), GemmBackendConfig_Epilogue_GemmBackendConfig_Epilogue_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits::max() }; bool GemmBackendConfig_Epilogue_IsValid(int value); constexpr GemmBackendConfig_Epilogue GemmBackendConfig_Epilogue_Epilogue_MIN = GemmBackendConfig_Epilogue_DEFAULT; constexpr GemmBackendConfig_Epilogue GemmBackendConfig_Epilogue_Epilogue_MAX = GemmBackendConfig_Epilogue_BIAS_GELU_AUX; constexpr int GemmBackendConfig_Epilogue_Epilogue_ARRAYSIZE = GemmBackendConfig_Epilogue_Epilogue_MAX + 1; const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* GemmBackendConfig_Epilogue_descriptor(); template inline const std::string& GemmBackendConfig_Epilogue_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function GemmBackendConfig_Epilogue_Name."); return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum( GemmBackendConfig_Epilogue_descriptor(), enum_t_value); } inline bool GemmBackendConfig_Epilogue_Parse( ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, GemmBackendConfig_Epilogue* value) { return ::PROTOBUF_NAMESPACE_ID::internal::ParseNamedEnum( GemmBackendConfig_Epilogue_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::min(), CudnnNormBackendConfig_Kind_CudnnNormBackendConfig_Kind_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits::max() }; bool CudnnNormBackendConfig_Kind_IsValid(int value); constexpr CudnnNormBackendConfig_Kind CudnnNormBackendConfig_Kind_Kind_MIN = CudnnNormBackendConfig_Kind_LAYER_FWD_INFER; constexpr CudnnNormBackendConfig_Kind CudnnNormBackendConfig_Kind_Kind_MAX = CudnnNormBackendConfig_Kind_LAYER_BWD; constexpr int CudnnNormBackendConfig_Kind_Kind_ARRAYSIZE = CudnnNormBackendConfig_Kind_Kind_MAX + 1; const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* CudnnNormBackendConfig_Kind_descriptor(); template inline const std::string& CudnnNormBackendConfig_Kind_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function CudnnNormBackendConfig_Kind_Name."); return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum( CudnnNormBackendConfig_Kind_descriptor(), enum_t_value); } inline bool CudnnNormBackendConfig_Kind_Parse( ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, CudnnNormBackendConfig_Kind* value) { return ::PROTOBUF_NAMESPACE_ID::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::min(), CudnnfMHABackendConfig_MaskType_CudnnfMHABackendConfig_MaskType_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits::max() }; bool CudnnfMHABackendConfig_MaskType_IsValid(int value); constexpr CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig_MaskType_MaskType_MIN = CudnnfMHABackendConfig_MaskType_NO_MASK; constexpr CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig_MaskType_MaskType_MAX = CudnnfMHABackendConfig_MaskType_ALIBI; constexpr int CudnnfMHABackendConfig_MaskType_MaskType_ARRAYSIZE = CudnnfMHABackendConfig_MaskType_MaskType_MAX + 1; const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* CudnnfMHABackendConfig_MaskType_descriptor(); template inline const std::string& CudnnfMHABackendConfig_MaskType_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function CudnnfMHABackendConfig_MaskType_Name."); return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum( CudnnfMHABackendConfig_MaskType_descriptor(), enum_t_value); } inline bool CudnnfMHABackendConfig_MaskType_Parse( ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, CudnnfMHABackendConfig_MaskType* value) { return ::PROTOBUF_NAMESPACE_ID::internal::ParseNamedEnum( CudnnfMHABackendConfig_MaskType_descriptor(), name, value); } // =================================================================== class CudnnConvBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnConvBackendConfig) */ { public: inline CudnnConvBackendConfig() : CudnnConvBackendConfig(nullptr) {} ~CudnnConvBackendConfig() override; explicit PROTOBUF_CONSTEXPR CudnnConvBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CudnnConvBackendConfig(const CudnnConvBackendConfig& from); CudnnConvBackendConfig(CudnnConvBackendConfig&& from) noexcept : CudnnConvBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnConvBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnConvBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CudnnConvBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CudnnConvBackendConfig& from) { CudnnConvBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CudnnConvBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CudnnConvBackendConfig"; } protected: explicit CudnnConvBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSerializedGraphFieldNumber = 9, kAlgorithmFieldNumber = 6, kConvResultScaleFieldNumber = 4, kSideInputScaleFieldNumber = 5, kLeakyreluAlphaFieldNumber = 8, kActivationModeFieldNumber = 3, kReorderedInt8NchwVectFieldNumber = 7, }; // optional string serialized_graph = 9; bool has_serialized_graph() const; private: bool _internal_has_serialized_graph() const; public: void clear_serialized_graph(); const std::string& serialized_graph() const; template void set_serialized_graph(ArgT0&& arg0, ArgT... args); std::string* mutable_serialized_graph(); PROTOBUF_NODISCARD std::string* release_serialized_graph(); void set_allocated_serialized_graph(std::string* serialized_graph); 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; private: bool _internal_has_algorithm() const; public: 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* algorithm); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: void unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); // 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; private: bool _internal_has_reordered_int8_nchw_vect() const; public: 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(); template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::internal::HasBits<1> _has_bits_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; ::PROTOBUF_NAMESPACE_ID::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_{} {} ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_; bool reordered_int8_nchw_vect_; } filter_and_bias_reordering_oneof_; uint32_t _oneof_case_[1]; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class GemmBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.GemmBackendConfig) */ { public: inline GemmBackendConfig() : GemmBackendConfig(nullptr) {} ~GemmBackendConfig() override; explicit PROTOBUF_CONSTEXPR GemmBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); GemmBackendConfig(const GemmBackendConfig& from); GemmBackendConfig(GemmBackendConfig&& from) noexcept : GemmBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(GemmBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- GemmBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const GemmBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const GemmBackendConfig& from) { GemmBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(GemmBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.GemmBackendConfig"; } protected: explicit GemmBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- typedef GemmBackendConfig_Epilogue 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 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 ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Epilogue_descriptor() { return GemmBackendConfig_Epilogue_descriptor(); } template static inline const std::string& Epilogue_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function Epilogue_Name."); return GemmBackendConfig_Epilogue_Name(enum_t_value); } static inline bool Epilogue_Parse(::PROTOBUF_NAMESPACE_ID::ConstStringParam 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; private: bool _internal_has_dot_dimension_numbers() const; public: 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* dot_dimension_numbers); private: const ::xla::DotDimensionNumbers& _internal_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_dot_dimension_numbers(); public: void unsafe_arena_set_allocated_dot_dimension_numbers( ::xla::DotDimensionNumbers* dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_dot_dimension_numbers(); // .xla.PrecisionConfig precision_config = 12; bool has_precision_config() const; private: bool _internal_has_precision_config() const; public: 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* precision_config); private: const ::xla::PrecisionConfig& _internal_precision_config() const; ::xla::PrecisionConfig* _internal_mutable_precision_config(); public: void unsafe_arena_set_allocated_precision_config( ::xla::PrecisionConfig* precision_config); ::xla::PrecisionConfig* unsafe_arena_release_precision_config(); // 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; private: bool _internal_has_lhs_stride() const; public: 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; private: bool _internal_has_rhs_stride() const; public: 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; private: bool _internal_has_grad_x() const; public: 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; private: bool _internal_has_grad_y() const; public: 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; private: bool _internal_has_selected_algorithm() const; public: 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(); template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::internal::HasBits<1> _has_bits_; mutable ::PROTOBUF_NAMESPACE_ID::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_{} {} ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_; int64_t selected_algorithm_; } algorithm_; uint32_t _oneof_case_[1]; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class BitcastBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.BitcastBackendConfig) */ { public: inline BitcastBackendConfig() : BitcastBackendConfig(nullptr) {} ~BitcastBackendConfig() override; explicit PROTOBUF_CONSTEXPR BitcastBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); BitcastBackendConfig(const BitcastBackendConfig& from); BitcastBackendConfig(BitcastBackendConfig&& from) noexcept : BitcastBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(BitcastBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- BitcastBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const BitcastBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const BitcastBackendConfig& from) { BitcastBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(BitcastBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.BitcastBackendConfig"; } protected: explicit BitcastBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSourceLayoutFieldNumber = 1, kResultLayoutFieldNumber = 2, }; // .xla.LayoutProto source_layout = 1; bool has_source_layout() const; private: bool _internal_has_source_layout() const; public: 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* source_layout); private: const ::xla::LayoutProto& _internal_source_layout() const; ::xla::LayoutProto* _internal_mutable_source_layout(); public: void unsafe_arena_set_allocated_source_layout( ::xla::LayoutProto* source_layout); ::xla::LayoutProto* unsafe_arena_release_source_layout(); // .xla.LayoutProto result_layout = 2; bool has_result_layout() const; private: bool _internal_has_result_layout() const; public: 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* result_layout); private: const ::xla::LayoutProto& _internal_result_layout() const; ::xla::LayoutProto* _internal_mutable_result_layout(); public: void unsafe_arena_set_allocated_result_layout( ::xla::LayoutProto* result_layout); ::xla::LayoutProto* unsafe_arena_release_result_layout(); // @@protoc_insertion_point(class_scope:xla.gpu.BitcastBackendConfig) private: class _Internal; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::xla::LayoutProto* source_layout_; ::xla::LayoutProto* result_layout_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CollectiveBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CollectiveBackendConfig) */ { public: inline CollectiveBackendConfig() : CollectiveBackendConfig(nullptr) {} ~CollectiveBackendConfig() override; explicit PROTOBUF_CONSTEXPR CollectiveBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CollectiveBackendConfig(const CollectiveBackendConfig& from); CollectiveBackendConfig(CollectiveBackendConfig&& from) noexcept : CollectiveBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CollectiveBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CollectiveBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CollectiveBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CollectiveBackendConfig& from) { CollectiveBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CollectiveBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CollectiveBackendConfig"; } protected: explicit CollectiveBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kReificationCostFieldNumber = 4, kIsSyncFieldNumber = 1, kNoParallelCustomCallFieldNumber = 2, kIsPipelinedFieldNumber = 3, }; // .xla.gpu.ReificationCost reification_cost = 4; bool has_reification_cost() const; private: bool _internal_has_reification_cost() const; public: void clear_reification_cost(); const ::xla::gpu::ReificationCost& reification_cost() const; PROTOBUF_NODISCARD ::xla::gpu::ReificationCost* release_reification_cost(); ::xla::gpu::ReificationCost* mutable_reification_cost(); void set_allocated_reification_cost(::xla::gpu::ReificationCost* reification_cost); private: const ::xla::gpu::ReificationCost& _internal_reification_cost() const; ::xla::gpu::ReificationCost* _internal_mutable_reification_cost(); public: void unsafe_arena_set_allocated_reification_cost( ::xla::gpu::ReificationCost* reification_cost); ::xla::gpu::ReificationCost* unsafe_arena_release_reification_cost(); // 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 no_parallel_custom_call = 2; void clear_no_parallel_custom_call(); bool no_parallel_custom_call() const; void set_no_parallel_custom_call(bool value); private: bool _internal_no_parallel_custom_call() const; void _internal_set_no_parallel_custom_call(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: // @@protoc_insertion_point(class_scope:xla.gpu.CollectiveBackendConfig) private: class _Internal; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::xla::gpu::ReificationCost* reification_cost_; bool is_sync_; bool no_parallel_custom_call_; bool is_pipelined_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class ReificationCost final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ReificationCost) */ { public: inline ReificationCost() : ReificationCost(nullptr) {} ~ReificationCost() override; explicit PROTOBUF_CONSTEXPR ReificationCost(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); ReificationCost(const ReificationCost& from); ReificationCost(ReificationCost&& from) noexcept : ReificationCost() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(ReificationCost* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ReificationCost* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const ReificationCost& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const ReificationCost& from) { ReificationCost::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(ReificationCost* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.ReificationCost"; } protected: explicit ReificationCost(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kEndToEndCyclesFieldNumber = 1, kExecTimeUsFieldNumber = 2, kComputeTimeUsFieldNumber = 3, kMemoryAccessTimeUsFieldNumber = 4, }; // 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; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { double end_to_end_cycles_; double exec_time_us_; double compute_time_us_; double memory_access_time_us_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CustomFusionConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CustomFusionConfig) */ { public: inline CustomFusionConfig() : CustomFusionConfig(nullptr) {} ~CustomFusionConfig() override; explicit PROTOBUF_CONSTEXPR CustomFusionConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CustomFusionConfig(const CustomFusionConfig& from); CustomFusionConfig(CustomFusionConfig&& from) noexcept : CustomFusionConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CustomFusionConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CustomFusionConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CustomFusionConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CustomFusionConfig& from) { CustomFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CustomFusionConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CustomFusionConfig"; } protected: explicit CustomFusionConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kNameFieldNumber = 1, kKernelIndexFieldNumber = 2, }; // string name = 1; void clear_name(); const std::string& name() const; template void set_name(ArgT0&& arg0, ArgT... args); std::string* mutable_name(); PROTOBUF_NODISCARD std::string* release_name(); void set_allocated_name(std::string* name); 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; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr name_; int32_t kernel_index_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CuDnnFusionConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CuDnnFusionConfig) */ { public: inline CuDnnFusionConfig() : CuDnnFusionConfig(nullptr) {} ~CuDnnFusionConfig() override; explicit PROTOBUF_CONSTEXPR CuDnnFusionConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CuDnnFusionConfig(const CuDnnFusionConfig& from); CuDnnFusionConfig(CuDnnFusionConfig&& from) noexcept : CuDnnFusionConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CuDnnFusionConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CuDnnFusionConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CuDnnFusionConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CuDnnFusionConfig& from) { CuDnnFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CuDnnFusionConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CuDnnFusionConfig"; } protected: explicit CuDnnFusionConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { 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; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { int64_t plan_id_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class BlockLevelFusionConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.BlockLevelFusionConfig) */ { public: inline BlockLevelFusionConfig() : BlockLevelFusionConfig(nullptr) {} ~BlockLevelFusionConfig() override; explicit PROTOBUF_CONSTEXPR BlockLevelFusionConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); BlockLevelFusionConfig(const BlockLevelFusionConfig& from); BlockLevelFusionConfig(BlockLevelFusionConfig&& from) noexcept : BlockLevelFusionConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const BlockLevelFusionConfig& default_instance() { return *internal_default_instance(); } static inline const BlockLevelFusionConfig* internal_default_instance() { return reinterpret_cast( &_BlockLevelFusionConfig_default_instance_); } static constexpr int kIndexInFileMessages = 7; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(BlockLevelFusionConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- BlockLevelFusionConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const BlockLevelFusionConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const BlockLevelFusionConfig& from) { BlockLevelFusionConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(BlockLevelFusionConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.BlockLevelFusionConfig"; } protected: explicit BlockLevelFusionConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kOutputTileSizesFieldNumber = 1, kNumWarpsFieldNumber = 2, }; // repeated int64 output_tile_sizes = 1; int output_tile_sizes_size() const; private: int _internal_output_tile_sizes_size() const; public: void clear_output_tile_sizes(); private: int64_t _internal_output_tile_sizes(int index) const; const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& _internal_output_tile_sizes() const; void _internal_add_output_tile_sizes(int64_t value); ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* _internal_mutable_output_tile_sizes(); public: int64_t output_tile_sizes(int index) const; void set_output_tile_sizes(int index, int64_t value); void add_output_tile_sizes(int64_t value); const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& output_tile_sizes() const; ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* mutable_output_tile_sizes(); // 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: // @@protoc_insertion_point(class_scope:xla.gpu.BlockLevelFusionConfig) private: class _Internal; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t > output_tile_sizes_; mutable std::atomic _output_tile_sizes_cached_byte_size_; int64_t num_warps_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class FusionBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.FusionBackendConfig) */ { public: inline FusionBackendConfig() : FusionBackendConfig(nullptr) {} ~FusionBackendConfig() override; explicit PROTOBUF_CONSTEXPR FusionBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); FusionBackendConfig(const FusionBackendConfig& from); FusionBackendConfig(FusionBackendConfig&& from) noexcept : FusionBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const FusionBackendConfig& default_instance() { return *internal_default_instance(); } static inline const FusionBackendConfig* internal_default_instance() { return reinterpret_cast( &_FusionBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 8; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(FusionBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- FusionBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const FusionBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const FusionBackendConfig& from) { FusionBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(FusionBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.FusionBackendConfig"; } protected: explicit FusionBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kKindFieldNumber = 1, kTritonGemmConfigFieldNumber = 2, kReificationCostFieldNumber = 3, kCustomFusionConfigFieldNumber = 4, kCudnnFusionConfigFieldNumber = 5, kBlockLevelFusionConfigFieldNumber = 6, }; // string kind = 1; void clear_kind(); const std::string& kind() const; template void set_kind(ArgT0&& arg0, ArgT... args); std::string* mutable_kind(); PROTOBUF_NODISCARD std::string* release_kind(); void set_allocated_kind(std::string* kind); private: const std::string& _internal_kind() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_kind(const std::string& value); std::string* _internal_mutable_kind(); public: // .xla.AutotuneResult.TritonGemmKey triton_gemm_config = 2; bool has_triton_gemm_config() const; private: bool _internal_has_triton_gemm_config() const; public: 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* triton_gemm_config); private: const ::xla::AutotuneResult_TritonGemmKey& _internal_triton_gemm_config() const; ::xla::AutotuneResult_TritonGemmKey* _internal_mutable_triton_gemm_config(); public: void unsafe_arena_set_allocated_triton_gemm_config( ::xla::AutotuneResult_TritonGemmKey* triton_gemm_config); ::xla::AutotuneResult_TritonGemmKey* unsafe_arena_release_triton_gemm_config(); // .xla.gpu.ReificationCost reification_cost = 3; bool has_reification_cost() const; private: bool _internal_has_reification_cost() const; public: void clear_reification_cost(); const ::xla::gpu::ReificationCost& reification_cost() const; PROTOBUF_NODISCARD ::xla::gpu::ReificationCost* release_reification_cost(); ::xla::gpu::ReificationCost* mutable_reification_cost(); void set_allocated_reification_cost(::xla::gpu::ReificationCost* reification_cost); private: const ::xla::gpu::ReificationCost& _internal_reification_cost() const; ::xla::gpu::ReificationCost* _internal_mutable_reification_cost(); public: void unsafe_arena_set_allocated_reification_cost( ::xla::gpu::ReificationCost* reification_cost); ::xla::gpu::ReificationCost* unsafe_arena_release_reification_cost(); // .xla.gpu.CustomFusionConfig custom_fusion_config = 4; bool has_custom_fusion_config() const; private: bool _internal_has_custom_fusion_config() const; public: 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* custom_fusion_config); private: const ::xla::gpu::CustomFusionConfig& _internal_custom_fusion_config() const; ::xla::gpu::CustomFusionConfig* _internal_mutable_custom_fusion_config(); public: void unsafe_arena_set_allocated_custom_fusion_config( ::xla::gpu::CustomFusionConfig* custom_fusion_config); ::xla::gpu::CustomFusionConfig* unsafe_arena_release_custom_fusion_config(); // .xla.gpu.CuDnnFusionConfig cudnn_fusion_config = 5; bool has_cudnn_fusion_config() const; private: bool _internal_has_cudnn_fusion_config() const; public: 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* cudnn_fusion_config); private: const ::xla::gpu::CuDnnFusionConfig& _internal_cudnn_fusion_config() const; ::xla::gpu::CuDnnFusionConfig* _internal_mutable_cudnn_fusion_config(); public: void unsafe_arena_set_allocated_cudnn_fusion_config( ::xla::gpu::CuDnnFusionConfig* cudnn_fusion_config); ::xla::gpu::CuDnnFusionConfig* unsafe_arena_release_cudnn_fusion_config(); // .xla.gpu.BlockLevelFusionConfig block_level_fusion_config = 6; bool has_block_level_fusion_config() const; private: bool _internal_has_block_level_fusion_config() const; public: 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* 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: void unsafe_arena_set_allocated_block_level_fusion_config( ::xla::gpu::BlockLevelFusionConfig* block_level_fusion_config); ::xla::gpu::BlockLevelFusionConfig* unsafe_arena_release_block_level_fusion_config(); // @@protoc_insertion_point(class_scope:xla.gpu.FusionBackendConfig) private: class _Internal; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr kind_; ::xla::AutotuneResult_TritonGemmKey* triton_gemm_config_; ::xla::gpu::ReificationCost* reification_cost_; ::xla::gpu::CustomFusionConfig* custom_fusion_config_; ::xla::gpu::CuDnnFusionConfig* cudnn_fusion_config_; ::xla::gpu::BlockLevelFusionConfig* block_level_fusion_config_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CudnnNormBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnNormBackendConfig) */ { public: inline CudnnNormBackendConfig() : CudnnNormBackendConfig(nullptr) {} ~CudnnNormBackendConfig() override; explicit PROTOBUF_CONSTEXPR CudnnNormBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CudnnNormBackendConfig(const CudnnNormBackendConfig& from); CudnnNormBackendConfig(CudnnNormBackendConfig&& from) noexcept : CudnnNormBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CudnnNormBackendConfig& default_instance() { return *internal_default_instance(); } static inline const CudnnNormBackendConfig* internal_default_instance() { return reinterpret_cast( &_CudnnNormBackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 9; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnNormBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnNormBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CudnnNormBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CudnnNormBackendConfig& from) { CudnnNormBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CudnnNormBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CudnnNormBackendConfig"; } protected: explicit CudnnNormBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- typedef CudnnNormBackendConfig_Kind 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 ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Kind_descriptor() { return CudnnNormBackendConfig_Kind_descriptor(); } template static inline const std::string& Kind_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function Kind_Name."); return CudnnNormBackendConfig_Kind_Name(enum_t_value); } static inline bool Kind_Parse(::PROTOBUF_NAMESPACE_ID::ConstStringParam 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; private: bool _internal_has_algorithm() const; public: 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* algorithm); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: void unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); // 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; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::stream_executor::dnn::AlgorithmProto* algorithm_; double epsilon_; int kind_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CudnnfMHABackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CudnnfMHABackendConfig) */ { public: inline CudnnfMHABackendConfig() : CudnnfMHABackendConfig(nullptr) {} ~CudnnfMHABackendConfig() override; explicit PROTOBUF_CONSTEXPR CudnnfMHABackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CudnnfMHABackendConfig(const CudnnfMHABackendConfig& from); CudnnfMHABackendConfig(CudnnfMHABackendConfig&& from) noexcept : CudnnfMHABackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const CudnnfMHABackendConfig& default_instance() { return *internal_default_instance(); } static inline const CudnnfMHABackendConfig* internal_default_instance() { return reinterpret_cast( &_CudnnfMHABackendConfig_default_instance_); } static constexpr int kIndexInFileMessages = 10; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CudnnfMHABackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CudnnfMHABackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CudnnfMHABackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CudnnfMHABackendConfig& from) { CudnnfMHABackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CudnnfMHABackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CudnnfMHABackendConfig"; } protected: explicit CudnnfMHABackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- typedef CudnnfMHABackendConfig_MaskType 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 ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* MaskType_descriptor() { return CudnnfMHABackendConfig_MaskType_descriptor(); } template static inline const std::string& MaskType_Name(T enum_t_value) { static_assert(::std::is_same::value || ::std::is_integral::value, "Incorrect type passed to function MaskType_Name."); return CudnnfMHABackendConfig_MaskType_Name(enum_t_value); } static inline bool MaskType_Parse(::PROTOBUF_NAMESPACE_ID::ConstStringParam 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, kIsFlashAttentionFieldNumber = 20, kIsCausalMaskFieldNumber = 21, kForceDeterministicFieldNumber = 23, kMaskTypeFieldNumber = 22, kSlidingWindowLengthFieldNumber = 24, kMaxSegPerBatchFieldNumber = 25, }; // .stream_executor.dnn.AlgorithmProto algorithm = 8; bool has_algorithm() const; private: bool _internal_has_algorithm() const; public: 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* algorithm); private: const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const; ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm(); public: void unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm); ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm(); // .xla.DotDimensionNumbers bmm1_dot_dimension_numbers = 11; bool has_bmm1_dot_dimension_numbers() const; private: bool _internal_has_bmm1_dot_dimension_numbers() const; public: 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* bmm1_dot_dimension_numbers); private: const ::xla::DotDimensionNumbers& _internal_bmm1_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm1_dot_dimension_numbers(); public: void unsafe_arena_set_allocated_bmm1_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm1_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_dot_dimension_numbers(); // .xla.DotDimensionNumbers bmm2_dot_dimension_numbers = 12; bool has_bmm2_dot_dimension_numbers() const; private: bool _internal_has_bmm2_dot_dimension_numbers() const; public: 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* bmm2_dot_dimension_numbers); private: const ::xla::DotDimensionNumbers& _internal_bmm2_dot_dimension_numbers() const; ::xla::DotDimensionNumbers* _internal_mutable_bmm2_dot_dimension_numbers(); public: void unsafe_arena_set_allocated_bmm2_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm2_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_dot_dimension_numbers(); // .xla.ShapeProto intermediate_tensor_shape = 14; bool has_intermediate_tensor_shape() const; private: bool _internal_has_intermediate_tensor_shape() const; public: 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* intermediate_tensor_shape); private: const ::xla::ShapeProto& _internal_intermediate_tensor_shape() const; ::xla::ShapeProto* _internal_mutable_intermediate_tensor_shape(); public: void unsafe_arena_set_allocated_intermediate_tensor_shape( ::xla::ShapeProto* intermediate_tensor_shape); ::xla::ShapeProto* unsafe_arena_release_intermediate_tensor_shape(); // .xla.DotDimensionNumbers bmm1_grad_gemm1_dot_dimension_numbers = 16; bool has_bmm1_grad_gemm1_dot_dimension_numbers() const; private: bool _internal_has_bmm1_grad_gemm1_dot_dimension_numbers() const; public: 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* 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: void unsafe_arena_set_allocated_bmm1_grad_gemm1_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm1_grad_gemm1_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_grad_gemm1_dot_dimension_numbers(); // .xla.DotDimensionNumbers bmm1_grad_gemm2_dot_dimension_numbers = 17; bool has_bmm1_grad_gemm2_dot_dimension_numbers() const; private: bool _internal_has_bmm1_grad_gemm2_dot_dimension_numbers() const; public: 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* 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: void unsafe_arena_set_allocated_bmm1_grad_gemm2_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm1_grad_gemm2_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm1_grad_gemm2_dot_dimension_numbers(); // .xla.DotDimensionNumbers bmm2_grad_gemm1_dot_dimension_numbers = 18; bool has_bmm2_grad_gemm1_dot_dimension_numbers() const; private: bool _internal_has_bmm2_grad_gemm1_dot_dimension_numbers() const; public: 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* 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: void unsafe_arena_set_allocated_bmm2_grad_gemm1_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm2_grad_gemm1_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_grad_gemm1_dot_dimension_numbers(); // .xla.DotDimensionNumbers bmm2_grad_gemm2_dot_dimension_numbers = 19; bool has_bmm2_grad_gemm2_dot_dimension_numbers() const; private: bool _internal_has_bmm2_grad_gemm2_dot_dimension_numbers() const; public: 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* 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: void unsafe_arena_set_allocated_bmm2_grad_gemm2_dot_dimension_numbers( ::xla::DotDimensionNumbers* bmm2_grad_gemm2_dot_dimension_numbers); ::xla::DotDimensionNumbers* unsafe_arena_release_bmm2_grad_gemm2_dot_dimension_numbers(); // 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: // 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: // .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: // 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; template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::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_; bool is_flash_attention_; bool is_causal_mask_; bool force_deterministic_; int mask_type_; int32_t sliding_window_length_; int32_t max_seg_per_batch_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class CustomCallBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CustomCallBackendConfig) */ { public: inline CustomCallBackendConfig() : CustomCallBackendConfig(nullptr) {} ~CustomCallBackendConfig() override; explicit PROTOBUF_CONSTEXPR CustomCallBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); CustomCallBackendConfig(const CustomCallBackendConfig& from); CustomCallBackendConfig(CustomCallBackendConfig&& from) noexcept : CustomCallBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 = 11; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CustomCallBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CustomCallBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const CustomCallBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const CustomCallBackendConfig& from) { CustomCallBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(CustomCallBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.CustomCallBackendConfig"; } protected: explicit CustomCallBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kOpaqueFieldNumber = 1, kAttributesFieldNumber = 2, }; // string opaque = 1; bool has_opaque() const; private: bool _internal_has_opaque() const; public: void clear_opaque(); const std::string& opaque() const; template void set_opaque(ArgT0&& arg0, ArgT... args); std::string* mutable_opaque(); PROTOBUF_NODISCARD std::string* release_opaque(); void set_allocated_opaque(std::string* opaque); 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; private: bool _internal_has_attributes() const; public: void clear_attributes(); const std::string& attributes() const; template void set_attributes(ArgT0&& arg0, ArgT... args); std::string* mutable_attributes(); PROTOBUF_NODISCARD std::string* release_attributes(); void set_allocated_attributes(std::string* attributes); 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(); template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { union RawBackendConfigOneofUnion { constexpr RawBackendConfigOneofUnion() : _constinit_{} {} ::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized _constinit_; ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr opaque_; ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr attributes_; } raw_backend_config_oneof_; mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; uint32_t _oneof_case_[1]; }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto; }; // ------------------------------------------------------------------- class GpuBackendConfig final : public ::PROTOBUF_NAMESPACE_ID::Message /* @@protoc_insertion_point(class_definition:xla.gpu.GpuBackendConfig) */ { public: inline GpuBackendConfig() : GpuBackendConfig(nullptr) {} ~GpuBackendConfig() override; explicit PROTOBUF_CONSTEXPR GpuBackendConfig(::PROTOBUF_NAMESPACE_ID::internal::ConstantInitialized); GpuBackendConfig(const GpuBackendConfig& from); GpuBackendConfig(GpuBackendConfig&& from) noexcept : GpuBackendConfig() { *this = ::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 (GetOwningArena() == from.GetOwningArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetOwningArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* descriptor() { return GetDescriptor(); } static const ::PROTOBUF_NAMESPACE_ID::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::PROTOBUF_NAMESPACE_ID::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const 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 = 12; 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 (GetOwningArena() != nullptr && GetOwningArena() == other->GetOwningArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetOwningArena() == other->GetOwningArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::PROTOBUF_NAMESPACE_ID::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(GpuBackendConfig* other) { if (other == this) return; GOOGLE_DCHECK(GetOwningArena() == other->GetOwningArena()); InternalSwap(other); } // implements Message ---------------------------------------------- GpuBackendConfig* New(::PROTOBUF_NAMESPACE_ID::Arena* arena = nullptr) const final { return CreateMaybeMessage(arena); } using ::PROTOBUF_NAMESPACE_ID::Message::CopyFrom; void CopyFrom(const GpuBackendConfig& from); using ::PROTOBUF_NAMESPACE_ID::Message::MergeFrom; void MergeFrom( const GpuBackendConfig& from) { GpuBackendConfig::MergeImpl(*this, from); } private: static void MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBUF_NAMESPACE_ID::Message& from_msg); public: PROTOBUF_ATTRIBUTE_REINITIALIZES void Clear() final; bool IsInitialized() const final; size_t ByteSizeLong() const final; const char* _InternalParse(const char* ptr, ::PROTOBUF_NAMESPACE_ID::internal::ParseContext* ctx) final; uint8_t* _InternalSerialize( uint8_t* target, ::PROTOBUF_NAMESPACE_ID::io::EpsCopyOutputStream* stream) const final; int GetCachedSize() const final { return _impl_._cached_size_.Get(); } private: void SharedCtor(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned); void SharedDtor(); void SetCachedSize(int size) const final; void InternalSwap(GpuBackendConfig* other); private: friend class ::PROTOBUF_NAMESPACE_ID::internal::AnyMetadata; static ::PROTOBUF_NAMESPACE_ID::StringPiece FullMessageName() { return "xla.gpu.GpuBackendConfig"; } protected: explicit GpuBackendConfig(::PROTOBUF_NAMESPACE_ID::Arena* arena, bool is_message_owned = false); public: static const ClassData _class_data_; const ::PROTOBUF_NAMESPACE_ID::Message::ClassData*GetClassData() const final; ::PROTOBUF_NAMESPACE_ID::Metadata GetMetadata() const final; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kWaitOnOperationQueuesFieldNumber = 2, 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(); private: int64_t _internal_wait_on_operation_queues(int index) const; const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& _internal_wait_on_operation_queues() const; void _internal_add_wait_on_operation_queues(int64_t value); ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* _internal_mutable_wait_on_operation_queues(); public: 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 ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& wait_on_operation_queues() const; ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* mutable_wait_on_operation_queues(); // 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* 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: void unsafe_arena_set_allocated_cudnn_conv_backend_config( ::xla::gpu::CudnnConvBackendConfig* cudnn_conv_backend_config); ::xla::gpu::CudnnConvBackendConfig* unsafe_arena_release_cudnn_conv_backend_config(); // .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* gemm_backend_config); private: const ::xla::gpu::GemmBackendConfig& _internal_gemm_backend_config() const; ::xla::gpu::GemmBackendConfig* _internal_mutable_gemm_backend_config(); public: void unsafe_arena_set_allocated_gemm_backend_config( ::xla::gpu::GemmBackendConfig* gemm_backend_config); ::xla::gpu::GemmBackendConfig* unsafe_arena_release_gemm_backend_config(); // .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* bitcast_backend_config); private: const ::xla::gpu::BitcastBackendConfig& _internal_bitcast_backend_config() const; ::xla::gpu::BitcastBackendConfig* _internal_mutable_bitcast_backend_config(); public: void unsafe_arena_set_allocated_bitcast_backend_config( ::xla::gpu::BitcastBackendConfig* bitcast_backend_config); ::xla::gpu::BitcastBackendConfig* unsafe_arena_release_bitcast_backend_config(); // .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* collective_backend_config); private: const ::xla::gpu::CollectiveBackendConfig& _internal_collective_backend_config() const; ::xla::gpu::CollectiveBackendConfig* _internal_mutable_collective_backend_config(); public: void unsafe_arena_set_allocated_collective_backend_config( ::xla::gpu::CollectiveBackendConfig* collective_backend_config); ::xla::gpu::CollectiveBackendConfig* unsafe_arena_release_collective_backend_config(); // .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* fusion_backend_config); private: const ::xla::gpu::FusionBackendConfig& _internal_fusion_backend_config() const; ::xla::gpu::FusionBackendConfig* _internal_mutable_fusion_backend_config(); public: void unsafe_arena_set_allocated_fusion_backend_config( ::xla::gpu::FusionBackendConfig* fusion_backend_config); ::xla::gpu::FusionBackendConfig* unsafe_arena_release_fusion_backend_config(); // .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* 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: void unsafe_arena_set_allocated_cudnn_norm_backend_config( ::xla::gpu::CudnnNormBackendConfig* cudnn_norm_backend_config); ::xla::gpu::CudnnNormBackendConfig* unsafe_arena_release_cudnn_norm_backend_config(); // .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* 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: void unsafe_arena_set_allocated_cudnn_fmha_backend_config( ::xla::gpu::CudnnfMHABackendConfig* cudnn_fmha_backend_config); ::xla::gpu::CudnnfMHABackendConfig* unsafe_arena_release_cudnn_fmha_backend_config(); // .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* 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 unsafe_arena_set_allocated_custom_call_backend_config( ::xla::gpu::CustomCallBackendConfig* custom_call_backend_config); ::xla::gpu::CustomCallBackendConfig* unsafe_arena_release_custom_call_backend_config(); 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(); template friend class ::PROTOBUF_NAMESPACE_ID::Arena::InternalHelper; typedef void InternalArenaConstructable_; typedef void DestructorSkippable_; struct Impl_ { ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t > wait_on_operation_queues_; mutable std::atomic _wait_on_operation_queues_cached_byte_size_; int64_t operation_queue_id_; bool force_earliest_schedule_; union BackendConfigUnion { constexpr BackendConfigUnion() : _constinit_{} {} ::PROTOBUF_NAMESPACE_ID::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 ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_; uint32_t _oneof_case_[1]; }; 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::_internal_has_algorithm() const { return this != internal_default_instance() && _impl_.algorithm_ != nullptr; } inline bool CudnnConvBackendConfig::has_algorithm() const { return _internal_has_algorithm(); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnConvBackendConfig::_internal_algorithm() const { 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 { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnConvBackendConfig::unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = algorithm; if (algorithm) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnConvBackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::release_algorithm() { ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::unsafe_arena_release_algorithm() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnConvBackendConfig.algorithm) ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::_internal_mutable_algorithm() { if (_impl_.algorithm_ == nullptr) { auto* p = CreateMaybeMessage<::stream_executor::dnn::AlgorithmProto>(GetArenaForAllocation()); _impl_.algorithm_ = p; } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnConvBackendConfig::mutable_algorithm() { ::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* algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } if (algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(algorithm)); if (message_arena != submessage_arena) { algorithm = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, algorithm, submessage_arena); } } else { } _impl_.algorithm_ = algorithm; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnConvBackendConfig.algorithm) } // double conv_result_scale = 4; inline void CudnnConvBackendConfig::clear_conv_result_scale() { _impl_.conv_result_scale_ = 0; } inline double CudnnConvBackendConfig::_internal_conv_result_scale() const { return _impl_.conv_result_scale_; } 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::_internal_set_conv_result_scale(double value) { _impl_.conv_result_scale_ = value; } 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) } // .stream_executor.dnn.ActivationMode activation_mode = 3; inline void CudnnConvBackendConfig::clear_activation_mode() { _impl_.activation_mode_ = 0; } inline ::stream_executor::dnn::ActivationMode CudnnConvBackendConfig::_internal_activation_mode() const { return static_cast< ::stream_executor::dnn::ActivationMode >(_impl_.activation_mode_); } 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::_internal_set_activation_mode(::stream_executor::dnn::ActivationMode value) { _impl_.activation_mode_ = value; } 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) } // double side_input_scale = 5; inline void CudnnConvBackendConfig::clear_side_input_scale() { _impl_.side_input_scale_ = 0; } inline double CudnnConvBackendConfig::_internal_side_input_scale() const { return _impl_.side_input_scale_; } 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::_internal_set_side_input_scale(double value) { _impl_.side_input_scale_ = value; } 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) } // double leakyrelu_alpha = 8; inline void CudnnConvBackendConfig::clear_leakyrelu_alpha() { _impl_.leakyrelu_alpha_ = 0; } inline double CudnnConvBackendConfig::_internal_leakyrelu_alpha() const { return _impl_.leakyrelu_alpha_; } inline double CudnnConvBackendConfig::leakyrelu_alpha() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.leakyrelu_alpha) return _internal_leakyrelu_alpha(); } inline void CudnnConvBackendConfig::_internal_set_leakyrelu_alpha(double value) { _impl_.leakyrelu_alpha_ = value; } inline void CudnnConvBackendConfig::set_leakyrelu_alpha(double value) { _internal_set_leakyrelu_alpha(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.leakyrelu_alpha) } // bool reordered_int8_nchw_vect = 7; inline bool CudnnConvBackendConfig::_internal_has_reordered_int8_nchw_vect() const { return filter_and_bias_reordering_oneof_case() == kReorderedInt8NchwVect; } inline bool CudnnConvBackendConfig::has_reordered_int8_nchw_vect() const { return _internal_has_reordered_int8_nchw_vect(); } inline void CudnnConvBackendConfig::set_has_reordered_int8_nchw_vect() { _impl_._oneof_case_[0] = kReorderedInt8NchwVect; } inline void CudnnConvBackendConfig::clear_reordered_int8_nchw_vect() { if (_internal_has_reordered_int8_nchw_vect()) { _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_ = false; clear_has_filter_and_bias_reordering_oneof(); } } inline bool CudnnConvBackendConfig::_internal_reordered_int8_nchw_vect() const { if (_internal_has_reordered_int8_nchw_vect()) { return _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_; } return false; } inline void CudnnConvBackendConfig::_internal_set_reordered_int8_nchw_vect(bool value) { if (!_internal_has_reordered_int8_nchw_vect()) { clear_filter_and_bias_reordering_oneof(); set_has_reordered_int8_nchw_vect(); } _impl_.filter_and_bias_reordering_oneof_.reordered_int8_nchw_vect_ = value; } 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) { _internal_set_reordered_int8_nchw_vect(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.reordered_int8_nchw_vect) } // optional string serialized_graph = 9; inline bool CudnnConvBackendConfig::_internal_has_serialized_graph() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; return value; } inline bool CudnnConvBackendConfig::has_serialized_graph() const { return _internal_has_serialized_graph(); } inline void CudnnConvBackendConfig::clear_serialized_graph() { _impl_.serialized_graph_.ClearToEmpty(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const std::string& CudnnConvBackendConfig::serialized_graph() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnConvBackendConfig.serialized_graph) return _internal_serialized_graph(); } template inline PROTOBUF_ALWAYS_INLINE void CudnnConvBackendConfig::set_serialized_graph(ArgT0&& arg0, ArgT... args) { _impl_._has_bits_[0] |= 0x00000001u; _impl_.serialized_graph_.Set(static_cast(arg0), args..., GetArenaForAllocation()); // @@protoc_insertion_point(field_set:xla.gpu.CudnnConvBackendConfig.serialized_graph) } inline std::string* CudnnConvBackendConfig::mutable_serialized_graph() { 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 { return _impl_.serialized_graph_.Get(); } inline void CudnnConvBackendConfig::_internal_set_serialized_graph(const std::string& value) { _impl_._has_bits_[0] |= 0x00000001u; _impl_.serialized_graph_.Set(value, GetArenaForAllocation()); } inline std::string* CudnnConvBackendConfig::_internal_mutable_serialized_graph() { _impl_._has_bits_[0] |= 0x00000001u; return _impl_.serialized_graph_.Mutable(GetArenaForAllocation()); } inline std::string* CudnnConvBackendConfig::release_serialized_graph() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnConvBackendConfig.serialized_graph) if (!_internal_has_serialized_graph()) { return nullptr; } _impl_._has_bits_[0] &= ~0x00000001u; auto* p = _impl_.serialized_graph_.Release(); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.serialized_graph_.IsDefault()) { _impl_.serialized_graph_.Set("", GetArenaForAllocation()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING return p; } inline void CudnnConvBackendConfig::set_allocated_serialized_graph(std::string* serialized_graph) { if (serialized_graph != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.serialized_graph_.SetAllocated(serialized_graph, GetArenaForAllocation()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.serialized_graph_.IsDefault()) { _impl_.serialized_graph_.Set("", GetArenaForAllocation()); } #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::_internal_has_selected_algorithm() const { return algorithm_case() == kSelectedAlgorithm; } inline bool GemmBackendConfig::has_selected_algorithm() const { return _internal_has_selected_algorithm(); } inline void GemmBackendConfig::set_has_selected_algorithm() { _impl_._oneof_case_[0] = kSelectedAlgorithm; } inline void GemmBackendConfig::clear_selected_algorithm() { if (_internal_has_selected_algorithm()) { _impl_.algorithm_.selected_algorithm_ = int64_t{0}; clear_has_algorithm(); } } inline int64_t GemmBackendConfig::_internal_selected_algorithm() const { if (_internal_has_selected_algorithm()) { return _impl_.algorithm_.selected_algorithm_; } return int64_t{0}; } inline void GemmBackendConfig::_internal_set_selected_algorithm(int64_t value) { if (!_internal_has_selected_algorithm()) { clear_algorithm(); set_has_selected_algorithm(); } _impl_.algorithm_.selected_algorithm_ = value; } 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) { _internal_set_selected_algorithm(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.selected_algorithm) } // double alpha_real = 2; inline void GemmBackendConfig::clear_alpha_real() { _impl_.alpha_real_ = 0; } inline double GemmBackendConfig::_internal_alpha_real() const { return _impl_.alpha_real_; } inline double GemmBackendConfig::alpha_real() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.alpha_real) return _internal_alpha_real(); } inline void GemmBackendConfig::_internal_set_alpha_real(double value) { _impl_.alpha_real_ = value; } inline void GemmBackendConfig::set_alpha_real(double value) { _internal_set_alpha_real(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.alpha_real) } // double alpha_imag = 9; inline void GemmBackendConfig::clear_alpha_imag() { _impl_.alpha_imag_ = 0; } inline double GemmBackendConfig::_internal_alpha_imag() const { return _impl_.alpha_imag_; } inline double GemmBackendConfig::alpha_imag() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.alpha_imag) return _internal_alpha_imag(); } inline void GemmBackendConfig::_internal_set_alpha_imag(double value) { _impl_.alpha_imag_ = value; } inline void GemmBackendConfig::set_alpha_imag(double value) { _internal_set_alpha_imag(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.alpha_imag) } // double beta = 3; inline void GemmBackendConfig::clear_beta() { _impl_.beta_ = 0; } inline double GemmBackendConfig::_internal_beta() const { return _impl_.beta_; } inline double GemmBackendConfig::beta() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.beta) return _internal_beta(); } inline void GemmBackendConfig::_internal_set_beta(double value) { _impl_.beta_ = value; } inline void GemmBackendConfig::set_beta(double value) { _internal_set_beta(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.beta) } // .xla.DotDimensionNumbers dot_dimension_numbers = 7; inline bool GemmBackendConfig::_internal_has_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.dot_dimension_numbers_ != nullptr; } inline bool GemmBackendConfig::has_dot_dimension_numbers() const { return _internal_has_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& GemmBackendConfig::_internal_dot_dimension_numbers() const { 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 { // @@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* dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimension_numbers_); } _impl_.dot_dimension_numbers_ = dot_dimension_numbers; if (dot_dimension_numbers) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmBackendConfig.dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* GemmBackendConfig::release_dot_dimension_numbers() { ::xla::DotDimensionNumbers* temp = _impl_.dot_dimension_numbers_; _impl_.dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::unsafe_arena_release_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.GemmBackendConfig.dot_dimension_numbers) ::xla::DotDimensionNumbers* temp = _impl_.dot_dimension_numbers_; _impl_.dot_dimension_numbers_ = nullptr; return temp; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::_internal_mutable_dot_dimension_numbers() { if (_impl_.dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.dot_dimension_numbers_ = p; } return _impl_.dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* GemmBackendConfig::mutable_dot_dimension_numbers() { ::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* dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.dot_dimension_numbers_); } if (dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(dot_dimension_numbers)); if (message_arena != submessage_arena) { dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, dot_dimension_numbers, submessage_arena); } } else { } _impl_.dot_dimension_numbers_ = dot_dimension_numbers; // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmBackendConfig.dot_dimension_numbers) } // .xla.PrecisionConfig precision_config = 12; inline bool GemmBackendConfig::_internal_has_precision_config() const { return this != internal_default_instance() && _impl_.precision_config_ != nullptr; } inline bool GemmBackendConfig::has_precision_config() const { return _internal_has_precision_config(); } inline const ::xla::PrecisionConfig& GemmBackendConfig::_internal_precision_config() const { 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 { // @@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* precision_config) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.precision_config_); } _impl_.precision_config_ = precision_config; if (precision_config) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmBackendConfig.precision_config) } inline ::xla::PrecisionConfig* GemmBackendConfig::release_precision_config() { ::xla::PrecisionConfig* temp = _impl_.precision_config_; _impl_.precision_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::PrecisionConfig* GemmBackendConfig::unsafe_arena_release_precision_config() { // @@protoc_insertion_point(field_release:xla.gpu.GemmBackendConfig.precision_config) ::xla::PrecisionConfig* temp = _impl_.precision_config_; _impl_.precision_config_ = nullptr; return temp; } inline ::xla::PrecisionConfig* GemmBackendConfig::_internal_mutable_precision_config() { if (_impl_.precision_config_ == nullptr) { auto* p = CreateMaybeMessage<::xla::PrecisionConfig>(GetArenaForAllocation()); _impl_.precision_config_ = p; } return _impl_.precision_config_; } inline ::xla::PrecisionConfig* GemmBackendConfig::mutable_precision_config() { ::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* precision_config) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.precision_config_); } if (precision_config) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(precision_config)); if (message_arena != submessage_arena) { precision_config = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, precision_config, submessage_arena); } } else { } _impl_.precision_config_ = precision_config; // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmBackendConfig.precision_config) } // .xla.gpu.GemmBackendConfig.Epilogue epilogue = 13; inline void GemmBackendConfig::clear_epilogue() { _impl_.epilogue_ = 0; } inline ::xla::gpu::GemmBackendConfig_Epilogue GemmBackendConfig::_internal_epilogue() const { return static_cast< ::xla::gpu::GemmBackendConfig_Epilogue >(_impl_.epilogue_); } inline ::xla::gpu::GemmBackendConfig_Epilogue GemmBackendConfig::epilogue() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.epilogue) return _internal_epilogue(); } inline void GemmBackendConfig::_internal_set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value) { _impl_.epilogue_ = value; } inline void GemmBackendConfig::set_epilogue(::xla::gpu::GemmBackendConfig_Epilogue value) { _internal_set_epilogue(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.epilogue) } // optional int64 lhs_stride = 14; inline bool GemmBackendConfig::_internal_has_lhs_stride() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; return value; } inline bool GemmBackendConfig::has_lhs_stride() const { return _internal_has_lhs_stride(); } inline void GemmBackendConfig::clear_lhs_stride() { _impl_.lhs_stride_ = int64_t{0}; _impl_._has_bits_[0] &= ~0x00000001u; } inline int64_t GemmBackendConfig::_internal_lhs_stride() const { return _impl_.lhs_stride_; } inline int64_t GemmBackendConfig::lhs_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.lhs_stride) return _internal_lhs_stride(); } inline void GemmBackendConfig::_internal_set_lhs_stride(int64_t value) { _impl_._has_bits_[0] |= 0x00000001u; _impl_.lhs_stride_ = value; } inline void GemmBackendConfig::set_lhs_stride(int64_t value) { _internal_set_lhs_stride(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.lhs_stride) } // optional int64 rhs_stride = 15; inline bool GemmBackendConfig::_internal_has_rhs_stride() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; return value; } inline bool GemmBackendConfig::has_rhs_stride() const { return _internal_has_rhs_stride(); } inline void GemmBackendConfig::clear_rhs_stride() { _impl_.rhs_stride_ = int64_t{0}; _impl_._has_bits_[0] &= ~0x00000002u; } inline int64_t GemmBackendConfig::_internal_rhs_stride() const { return _impl_.rhs_stride_; } inline int64_t GemmBackendConfig::rhs_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.rhs_stride) return _internal_rhs_stride(); } inline void GemmBackendConfig::_internal_set_rhs_stride(int64_t value) { _impl_._has_bits_[0] |= 0x00000002u; _impl_.rhs_stride_ = value; } inline void GemmBackendConfig::set_rhs_stride(int64_t value) { _internal_set_rhs_stride(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.rhs_stride) } // optional bool grad_x = 16; inline bool GemmBackendConfig::_internal_has_grad_x() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; return value; } inline bool GemmBackendConfig::has_grad_x() const { return _internal_has_grad_x(); } inline void GemmBackendConfig::clear_grad_x() { _impl_.grad_x_ = false; _impl_._has_bits_[0] &= ~0x00000004u; } inline bool GemmBackendConfig::_internal_grad_x() const { return _impl_.grad_x_; } inline bool GemmBackendConfig::grad_x() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.grad_x) return _internal_grad_x(); } inline void GemmBackendConfig::_internal_set_grad_x(bool value) { _impl_._has_bits_[0] |= 0x00000004u; _impl_.grad_x_ = value; } inline void GemmBackendConfig::set_grad_x(bool value) { _internal_set_grad_x(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.grad_x) } // optional bool grad_y = 17; inline bool GemmBackendConfig::_internal_has_grad_y() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; return value; } inline bool GemmBackendConfig::has_grad_y() const { return _internal_has_grad_y(); } inline void GemmBackendConfig::clear_grad_y() { _impl_.grad_y_ = false; _impl_._has_bits_[0] &= ~0x00000008u; } inline bool GemmBackendConfig::_internal_grad_y() const { return _impl_.grad_y_; } inline bool GemmBackendConfig::grad_y() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.grad_y) return _internal_grad_y(); } inline void GemmBackendConfig::_internal_set_grad_y(bool value) { _impl_._has_bits_[0] |= 0x00000008u; _impl_.grad_y_ = value; } inline void GemmBackendConfig::set_grad_y(bool value) { _internal_set_grad_y(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmBackendConfig.grad_y) } // bool damax_output = 18; inline void GemmBackendConfig::clear_damax_output() { _impl_.damax_output_ = false; } inline bool GemmBackendConfig::_internal_damax_output() const { return _impl_.damax_output_; } inline bool GemmBackendConfig::damax_output() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmBackendConfig.damax_output) return _internal_damax_output(); } inline void GemmBackendConfig::_internal_set_damax_output(bool value) { _impl_.damax_output_ = value; } 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::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::_internal_has_source_layout() const { return this != internal_default_instance() && _impl_.source_layout_ != nullptr; } inline bool BitcastBackendConfig::has_source_layout() const { return _internal_has_source_layout(); } inline const ::xla::LayoutProto& BitcastBackendConfig::_internal_source_layout() const { 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 { // @@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* source_layout) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.source_layout_); } _impl_.source_layout_ = source_layout; if (source_layout) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.BitcastBackendConfig.source_layout) } inline ::xla::LayoutProto* BitcastBackendConfig::release_source_layout() { ::xla::LayoutProto* temp = _impl_.source_layout_; _impl_.source_layout_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::unsafe_arena_release_source_layout() { // @@protoc_insertion_point(field_release:xla.gpu.BitcastBackendConfig.source_layout) ::xla::LayoutProto* temp = _impl_.source_layout_; _impl_.source_layout_ = nullptr; return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::_internal_mutable_source_layout() { if (_impl_.source_layout_ == nullptr) { auto* p = CreateMaybeMessage<::xla::LayoutProto>(GetArenaForAllocation()); _impl_.source_layout_ = p; } return _impl_.source_layout_; } inline ::xla::LayoutProto* BitcastBackendConfig::mutable_source_layout() { ::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* source_layout) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.source_layout_); } if (source_layout) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(source_layout)); if (message_arena != submessage_arena) { source_layout = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, source_layout, submessage_arena); } } else { } _impl_.source_layout_ = source_layout; // @@protoc_insertion_point(field_set_allocated:xla.gpu.BitcastBackendConfig.source_layout) } // .xla.LayoutProto result_layout = 2; inline bool BitcastBackendConfig::_internal_has_result_layout() const { return this != internal_default_instance() && _impl_.result_layout_ != nullptr; } inline bool BitcastBackendConfig::has_result_layout() const { return _internal_has_result_layout(); } inline const ::xla::LayoutProto& BitcastBackendConfig::_internal_result_layout() const { 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 { // @@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* result_layout) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.result_layout_); } _impl_.result_layout_ = result_layout; if (result_layout) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.BitcastBackendConfig.result_layout) } inline ::xla::LayoutProto* BitcastBackendConfig::release_result_layout() { ::xla::LayoutProto* temp = _impl_.result_layout_; _impl_.result_layout_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::unsafe_arena_release_result_layout() { // @@protoc_insertion_point(field_release:xla.gpu.BitcastBackendConfig.result_layout) ::xla::LayoutProto* temp = _impl_.result_layout_; _impl_.result_layout_ = nullptr; return temp; } inline ::xla::LayoutProto* BitcastBackendConfig::_internal_mutable_result_layout() { if (_impl_.result_layout_ == nullptr) { auto* p = CreateMaybeMessage<::xla::LayoutProto>(GetArenaForAllocation()); _impl_.result_layout_ = p; } return _impl_.result_layout_; } inline ::xla::LayoutProto* BitcastBackendConfig::mutable_result_layout() { ::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* result_layout) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.result_layout_); } if (result_layout) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(result_layout)); if (message_arena != submessage_arena) { result_layout = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, result_layout, submessage_arena); } } else { } _impl_.result_layout_ = result_layout; // @@protoc_insertion_point(field_set_allocated:xla.gpu.BitcastBackendConfig.result_layout) } // ------------------------------------------------------------------- // CollectiveBackendConfig // bool is_sync = 1; inline void CollectiveBackendConfig::clear_is_sync() { _impl_.is_sync_ = false; } inline bool CollectiveBackendConfig::_internal_is_sync() const { return _impl_.is_sync_; } inline bool CollectiveBackendConfig::is_sync() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.is_sync) return _internal_is_sync(); } inline void CollectiveBackendConfig::_internal_set_is_sync(bool value) { _impl_.is_sync_ = value; } inline void CollectiveBackendConfig::set_is_sync(bool value) { _internal_set_is_sync(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.is_sync) } // bool no_parallel_custom_call = 2; inline void CollectiveBackendConfig::clear_no_parallel_custom_call() { _impl_.no_parallel_custom_call_ = false; } inline bool CollectiveBackendConfig::_internal_no_parallel_custom_call() const { return _impl_.no_parallel_custom_call_; } inline bool CollectiveBackendConfig::no_parallel_custom_call() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.no_parallel_custom_call) return _internal_no_parallel_custom_call(); } inline void CollectiveBackendConfig::_internal_set_no_parallel_custom_call(bool value) { _impl_.no_parallel_custom_call_ = value; } inline void CollectiveBackendConfig::set_no_parallel_custom_call(bool value) { _internal_set_no_parallel_custom_call(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.no_parallel_custom_call) } // bool is_pipelined = 3; inline void CollectiveBackendConfig::clear_is_pipelined() { _impl_.is_pipelined_ = false; } inline bool CollectiveBackendConfig::_internal_is_pipelined() const { return _impl_.is_pipelined_; } inline bool CollectiveBackendConfig::is_pipelined() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.is_pipelined) return _internal_is_pipelined(); } inline void CollectiveBackendConfig::_internal_set_is_pipelined(bool value) { _impl_.is_pipelined_ = value; } inline void CollectiveBackendConfig::set_is_pipelined(bool value) { _internal_set_is_pipelined(value); // @@protoc_insertion_point(field_set:xla.gpu.CollectiveBackendConfig.is_pipelined) } // .xla.gpu.ReificationCost reification_cost = 4; inline bool CollectiveBackendConfig::_internal_has_reification_cost() const { return this != internal_default_instance() && _impl_.reification_cost_ != nullptr; } inline bool CollectiveBackendConfig::has_reification_cost() const { return _internal_has_reification_cost(); } inline void CollectiveBackendConfig::clear_reification_cost() { if (GetArenaForAllocation() == nullptr && _impl_.reification_cost_ != nullptr) { delete _impl_.reification_cost_; } _impl_.reification_cost_ = nullptr; } inline const ::xla::gpu::ReificationCost& CollectiveBackendConfig::_internal_reification_cost() const { const ::xla::gpu::ReificationCost* p = _impl_.reification_cost_; return p != nullptr ? *p : reinterpret_cast( ::xla::gpu::_ReificationCost_default_instance_); } inline const ::xla::gpu::ReificationCost& CollectiveBackendConfig::reification_cost() const { // @@protoc_insertion_point(field_get:xla.gpu.CollectiveBackendConfig.reification_cost) return _internal_reification_cost(); } inline void CollectiveBackendConfig::unsafe_arena_set_allocated_reification_cost( ::xla::gpu::ReificationCost* reification_cost) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.reification_cost_); } _impl_.reification_cost_ = reification_cost; if (reification_cost) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CollectiveBackendConfig.reification_cost) } inline ::xla::gpu::ReificationCost* CollectiveBackendConfig::release_reification_cost() { ::xla::gpu::ReificationCost* temp = _impl_.reification_cost_; _impl_.reification_cost_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::gpu::ReificationCost* CollectiveBackendConfig::unsafe_arena_release_reification_cost() { // @@protoc_insertion_point(field_release:xla.gpu.CollectiveBackendConfig.reification_cost) ::xla::gpu::ReificationCost* temp = _impl_.reification_cost_; _impl_.reification_cost_ = nullptr; return temp; } inline ::xla::gpu::ReificationCost* CollectiveBackendConfig::_internal_mutable_reification_cost() { if (_impl_.reification_cost_ == nullptr) { auto* p = CreateMaybeMessage<::xla::gpu::ReificationCost>(GetArenaForAllocation()); _impl_.reification_cost_ = p; } return _impl_.reification_cost_; } inline ::xla::gpu::ReificationCost* CollectiveBackendConfig::mutable_reification_cost() { ::xla::gpu::ReificationCost* _msg = _internal_mutable_reification_cost(); // @@protoc_insertion_point(field_mutable:xla.gpu.CollectiveBackendConfig.reification_cost) return _msg; } inline void CollectiveBackendConfig::set_allocated_reification_cost(::xla::gpu::ReificationCost* reification_cost) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete _impl_.reification_cost_; } if (reification_cost) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(reification_cost); if (message_arena != submessage_arena) { reification_cost = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, reification_cost, submessage_arena); } } else { } _impl_.reification_cost_ = reification_cost; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CollectiveBackendConfig.reification_cost) } // ------------------------------------------------------------------- // ReificationCost // double end_to_end_cycles = 1; inline void ReificationCost::clear_end_to_end_cycles() { _impl_.end_to_end_cycles_ = 0; } inline double ReificationCost::_internal_end_to_end_cycles() const { return _impl_.end_to_end_cycles_; } 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::_internal_set_end_to_end_cycles(double value) { _impl_.end_to_end_cycles_ = value; } 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) } // double exec_time_us = 2; inline void ReificationCost::clear_exec_time_us() { _impl_.exec_time_us_ = 0; } inline double ReificationCost::_internal_exec_time_us() const { return _impl_.exec_time_us_; } 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::_internal_set_exec_time_us(double value) { _impl_.exec_time_us_ = value; } 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) } // double compute_time_us = 3; inline void ReificationCost::clear_compute_time_us() { _impl_.compute_time_us_ = 0; } inline double ReificationCost::_internal_compute_time_us() const { return _impl_.compute_time_us_; } 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::_internal_set_compute_time_us(double value) { _impl_.compute_time_us_ = value; } 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) } // double memory_access_time_us = 4; inline void ReificationCost::clear_memory_access_time_us() { _impl_.memory_access_time_us_ = 0; } inline double ReificationCost::_internal_memory_access_time_us() const { return _impl_.memory_access_time_us_; } 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::_internal_set_memory_access_time_us(double value) { _impl_.memory_access_time_us_ = value; } 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) } // ------------------------------------------------------------------- // CustomFusionConfig // string name = 1; inline void CustomFusionConfig::clear_name() { _impl_.name_.ClearToEmpty(); } inline const std::string& CustomFusionConfig::name() const { // @@protoc_insertion_point(field_get:xla.gpu.CustomFusionConfig.name) return _internal_name(); } template inline PROTOBUF_ALWAYS_INLINE void CustomFusionConfig::set_name(ArgT0&& arg0, ArgT... args) { _impl_.name_.Set(static_cast(arg0), args..., GetArenaForAllocation()); // @@protoc_insertion_point(field_set:xla.gpu.CustomFusionConfig.name) } inline std::string* CustomFusionConfig::mutable_name() { 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 { return _impl_.name_.Get(); } inline void CustomFusionConfig::_internal_set_name(const std::string& value) { _impl_.name_.Set(value, GetArenaForAllocation()); } inline std::string* CustomFusionConfig::_internal_mutable_name() { return _impl_.name_.Mutable(GetArenaForAllocation()); } inline std::string* CustomFusionConfig::release_name() { // @@protoc_insertion_point(field_release:xla.gpu.CustomFusionConfig.name) return _impl_.name_.Release(); } inline void CustomFusionConfig::set_allocated_name(std::string* name) { if (name != nullptr) { } else { } _impl_.name_.SetAllocated(name, GetArenaForAllocation()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.name_.IsDefault()) { _impl_.name_.Set("", GetArenaForAllocation()); } #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() { _impl_.kernel_index_ = 0; } inline int32_t CustomFusionConfig::_internal_kernel_index() const { return _impl_.kernel_index_; } inline int32_t CustomFusionConfig::kernel_index() const { // @@protoc_insertion_point(field_get:xla.gpu.CustomFusionConfig.kernel_index) return _internal_kernel_index(); } inline void CustomFusionConfig::_internal_set_kernel_index(int32_t value) { _impl_.kernel_index_ = value; } inline void CustomFusionConfig::set_kernel_index(int32_t value) { _internal_set_kernel_index(value); // @@protoc_insertion_point(field_set:xla.gpu.CustomFusionConfig.kernel_index) } // ------------------------------------------------------------------- // CuDnnFusionConfig // int64 plan_id = 1; inline void CuDnnFusionConfig::clear_plan_id() { _impl_.plan_id_ = int64_t{0}; } inline int64_t CuDnnFusionConfig::_internal_plan_id() const { return _impl_.plan_id_; } inline int64_t CuDnnFusionConfig::plan_id() const { // @@protoc_insertion_point(field_get:xla.gpu.CuDnnFusionConfig.plan_id) return _internal_plan_id(); } inline void CuDnnFusionConfig::_internal_set_plan_id(int64_t value) { _impl_.plan_id_ = value; } inline void CuDnnFusionConfig::set_plan_id(int64_t value) { _internal_set_plan_id(value); // @@protoc_insertion_point(field_set:xla.gpu.CuDnnFusionConfig.plan_id) } // ------------------------------------------------------------------- // BlockLevelFusionConfig // repeated int64 output_tile_sizes = 1; inline int BlockLevelFusionConfig::_internal_output_tile_sizes_size() const { return _impl_.output_tile_sizes_.size(); } inline int BlockLevelFusionConfig::output_tile_sizes_size() const { return _internal_output_tile_sizes_size(); } inline void BlockLevelFusionConfig::clear_output_tile_sizes() { _impl_.output_tile_sizes_.Clear(); } inline int64_t BlockLevelFusionConfig::_internal_output_tile_sizes(int index) const { return _impl_.output_tile_sizes_.Get(index); } inline int64_t BlockLevelFusionConfig::output_tile_sizes(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.output_tile_sizes) return _internal_output_tile_sizes(index); } inline void BlockLevelFusionConfig::set_output_tile_sizes(int index, int64_t value) { _impl_.output_tile_sizes_.Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.output_tile_sizes) } inline void BlockLevelFusionConfig::_internal_add_output_tile_sizes(int64_t value) { _impl_.output_tile_sizes_.Add(value); } inline void BlockLevelFusionConfig::add_output_tile_sizes(int64_t value) { _internal_add_output_tile_sizes(value); // @@protoc_insertion_point(field_add:xla.gpu.BlockLevelFusionConfig.output_tile_sizes) } inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& BlockLevelFusionConfig::_internal_output_tile_sizes() const { return _impl_.output_tile_sizes_; } inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& BlockLevelFusionConfig::output_tile_sizes() const { // @@protoc_insertion_point(field_list:xla.gpu.BlockLevelFusionConfig.output_tile_sizes) return _internal_output_tile_sizes(); } inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* BlockLevelFusionConfig::_internal_mutable_output_tile_sizes() { return &_impl_.output_tile_sizes_; } inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* BlockLevelFusionConfig::mutable_output_tile_sizes() { // @@protoc_insertion_point(field_mutable_list:xla.gpu.BlockLevelFusionConfig.output_tile_sizes) return _internal_mutable_output_tile_sizes(); } // int64 num_warps = 2; inline void BlockLevelFusionConfig::clear_num_warps() { _impl_.num_warps_ = int64_t{0}; } inline int64_t BlockLevelFusionConfig::_internal_num_warps() const { return _impl_.num_warps_; } inline int64_t BlockLevelFusionConfig::num_warps() const { // @@protoc_insertion_point(field_get:xla.gpu.BlockLevelFusionConfig.num_warps) return _internal_num_warps(); } inline void BlockLevelFusionConfig::_internal_set_num_warps(int64_t value) { _impl_.num_warps_ = value; } inline void BlockLevelFusionConfig::set_num_warps(int64_t value) { _internal_set_num_warps(value); // @@protoc_insertion_point(field_set:xla.gpu.BlockLevelFusionConfig.num_warps) } // ------------------------------------------------------------------- // FusionBackendConfig // string kind = 1; inline void FusionBackendConfig::clear_kind() { _impl_.kind_.ClearToEmpty(); } inline const std::string& FusionBackendConfig::kind() const { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.kind) return _internal_kind(); } template inline PROTOBUF_ALWAYS_INLINE void FusionBackendConfig::set_kind(ArgT0&& arg0, ArgT... args) { _impl_.kind_.Set(static_cast(arg0), args..., GetArenaForAllocation()); // @@protoc_insertion_point(field_set:xla.gpu.FusionBackendConfig.kind) } inline std::string* FusionBackendConfig::mutable_kind() { 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 { return _impl_.kind_.Get(); } inline void FusionBackendConfig::_internal_set_kind(const std::string& value) { _impl_.kind_.Set(value, GetArenaForAllocation()); } inline std::string* FusionBackendConfig::_internal_mutable_kind() { return _impl_.kind_.Mutable(GetArenaForAllocation()); } inline std::string* FusionBackendConfig::release_kind() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.kind) return _impl_.kind_.Release(); } inline void FusionBackendConfig::set_allocated_kind(std::string* kind) { if (kind != nullptr) { } else { } _impl_.kind_.SetAllocated(kind, GetArenaForAllocation()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.kind_.IsDefault()) { _impl_.kind_.Set("", GetArenaForAllocation()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.kind) } // .xla.AutotuneResult.TritonGemmKey triton_gemm_config = 2; inline bool FusionBackendConfig::_internal_has_triton_gemm_config() const { return this != internal_default_instance() && _impl_.triton_gemm_config_ != nullptr; } inline bool FusionBackendConfig::has_triton_gemm_config() const { return _internal_has_triton_gemm_config(); } inline const ::xla::AutotuneResult_TritonGemmKey& FusionBackendConfig::_internal_triton_gemm_config() const { 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 { // @@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* triton_gemm_config) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.triton_gemm_config_); } _impl_.triton_gemm_config_ = triton_gemm_config; if (triton_gemm_config) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.triton_gemm_config) } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::release_triton_gemm_config() { ::xla::AutotuneResult_TritonGemmKey* temp = _impl_.triton_gemm_config_; _impl_.triton_gemm_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::unsafe_arena_release_triton_gemm_config() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.triton_gemm_config) ::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() { if (_impl_.triton_gemm_config_ == nullptr) { auto* p = CreateMaybeMessage<::xla::AutotuneResult_TritonGemmKey>(GetArenaForAllocation()); _impl_.triton_gemm_config_ = p; } return _impl_.triton_gemm_config_; } inline ::xla::AutotuneResult_TritonGemmKey* FusionBackendConfig::mutable_triton_gemm_config() { ::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* triton_gemm_config) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.triton_gemm_config_); } if (triton_gemm_config) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(triton_gemm_config)); if (message_arena != submessage_arena) { triton_gemm_config = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, triton_gemm_config, submessage_arena); } } else { } _impl_.triton_gemm_config_ = triton_gemm_config; // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.triton_gemm_config) } // .xla.gpu.BlockLevelFusionConfig block_level_fusion_config = 6; inline bool FusionBackendConfig::_internal_has_block_level_fusion_config() const { return this != internal_default_instance() && _impl_.block_level_fusion_config_ != nullptr; } inline bool FusionBackendConfig::has_block_level_fusion_config() const { return _internal_has_block_level_fusion_config(); } inline void FusionBackendConfig::clear_block_level_fusion_config() { if (GetArenaForAllocation() == nullptr && _impl_.block_level_fusion_config_ != nullptr) { delete _impl_.block_level_fusion_config_; } _impl_.block_level_fusion_config_ = nullptr; } inline const ::xla::gpu::BlockLevelFusionConfig& FusionBackendConfig::_internal_block_level_fusion_config() const { 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 { // @@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* block_level_fusion_config) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.block_level_fusion_config_); } _impl_.block_level_fusion_config_ = block_level_fusion_config; if (block_level_fusion_config) { } else { } // @@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() { ::xla::gpu::BlockLevelFusionConfig* temp = _impl_.block_level_fusion_config_; _impl_.block_level_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::unsafe_arena_release_block_level_fusion_config() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.block_level_fusion_config) ::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() { if (_impl_.block_level_fusion_config_ == nullptr) { auto* p = CreateMaybeMessage<::xla::gpu::BlockLevelFusionConfig>(GetArenaForAllocation()); _impl_.block_level_fusion_config_ = p; } return _impl_.block_level_fusion_config_; } inline ::xla::gpu::BlockLevelFusionConfig* FusionBackendConfig::mutable_block_level_fusion_config() { ::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* block_level_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete _impl_.block_level_fusion_config_; } if (block_level_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(block_level_fusion_config); if (message_arena != submessage_arena) { block_level_fusion_config = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, block_level_fusion_config, submessage_arena); } } else { } _impl_.block_level_fusion_config_ = block_level_fusion_config; // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.block_level_fusion_config) } // .xla.gpu.CustomFusionConfig custom_fusion_config = 4; inline bool FusionBackendConfig::_internal_has_custom_fusion_config() const { return this != internal_default_instance() && _impl_.custom_fusion_config_ != nullptr; } inline bool FusionBackendConfig::has_custom_fusion_config() const { return _internal_has_custom_fusion_config(); } inline void FusionBackendConfig::clear_custom_fusion_config() { if (GetArenaForAllocation() == nullptr && _impl_.custom_fusion_config_ != nullptr) { delete _impl_.custom_fusion_config_; } _impl_.custom_fusion_config_ = nullptr; } inline const ::xla::gpu::CustomFusionConfig& FusionBackendConfig::_internal_custom_fusion_config() const { 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 { // @@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* custom_fusion_config) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.custom_fusion_config_); } _impl_.custom_fusion_config_ = custom_fusion_config; if (custom_fusion_config) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.custom_fusion_config) } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::release_custom_fusion_config() { ::xla::gpu::CustomFusionConfig* temp = _impl_.custom_fusion_config_; _impl_.custom_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::unsafe_arena_release_custom_fusion_config() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.custom_fusion_config) ::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() { if (_impl_.custom_fusion_config_ == nullptr) { auto* p = CreateMaybeMessage<::xla::gpu::CustomFusionConfig>(GetArenaForAllocation()); _impl_.custom_fusion_config_ = p; } return _impl_.custom_fusion_config_; } inline ::xla::gpu::CustomFusionConfig* FusionBackendConfig::mutable_custom_fusion_config() { ::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* custom_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete _impl_.custom_fusion_config_; } if (custom_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(custom_fusion_config); if (message_arena != submessage_arena) { custom_fusion_config = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, custom_fusion_config, submessage_arena); } } else { } _impl_.custom_fusion_config_ = custom_fusion_config; // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.custom_fusion_config) } // .xla.gpu.ReificationCost reification_cost = 3; inline bool FusionBackendConfig::_internal_has_reification_cost() const { return this != internal_default_instance() && _impl_.reification_cost_ != nullptr; } inline bool FusionBackendConfig::has_reification_cost() const { return _internal_has_reification_cost(); } inline void FusionBackendConfig::clear_reification_cost() { if (GetArenaForAllocation() == nullptr && _impl_.reification_cost_ != nullptr) { delete _impl_.reification_cost_; } _impl_.reification_cost_ = nullptr; } inline const ::xla::gpu::ReificationCost& FusionBackendConfig::_internal_reification_cost() const { const ::xla::gpu::ReificationCost* p = _impl_.reification_cost_; return p != nullptr ? *p : reinterpret_cast( ::xla::gpu::_ReificationCost_default_instance_); } inline const ::xla::gpu::ReificationCost& FusionBackendConfig::reification_cost() const { // @@protoc_insertion_point(field_get:xla.gpu.FusionBackendConfig.reification_cost) return _internal_reification_cost(); } inline void FusionBackendConfig::unsafe_arena_set_allocated_reification_cost( ::xla::gpu::ReificationCost* reification_cost) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.reification_cost_); } _impl_.reification_cost_ = reification_cost; if (reification_cost) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.reification_cost) } inline ::xla::gpu::ReificationCost* FusionBackendConfig::release_reification_cost() { ::xla::gpu::ReificationCost* temp = _impl_.reification_cost_; _impl_.reification_cost_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::gpu::ReificationCost* FusionBackendConfig::unsafe_arena_release_reification_cost() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.reification_cost) ::xla::gpu::ReificationCost* temp = _impl_.reification_cost_; _impl_.reification_cost_ = nullptr; return temp; } inline ::xla::gpu::ReificationCost* FusionBackendConfig::_internal_mutable_reification_cost() { if (_impl_.reification_cost_ == nullptr) { auto* p = CreateMaybeMessage<::xla::gpu::ReificationCost>(GetArenaForAllocation()); _impl_.reification_cost_ = p; } return _impl_.reification_cost_; } inline ::xla::gpu::ReificationCost* FusionBackendConfig::mutable_reification_cost() { ::xla::gpu::ReificationCost* _msg = _internal_mutable_reification_cost(); // @@protoc_insertion_point(field_mutable:xla.gpu.FusionBackendConfig.reification_cost) return _msg; } inline void FusionBackendConfig::set_allocated_reification_cost(::xla::gpu::ReificationCost* reification_cost) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete _impl_.reification_cost_; } if (reification_cost) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(reification_cost); if (message_arena != submessage_arena) { reification_cost = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, reification_cost, submessage_arena); } } else { } _impl_.reification_cost_ = reification_cost; // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.reification_cost) } // .xla.gpu.CuDnnFusionConfig cudnn_fusion_config = 5; inline bool FusionBackendConfig::_internal_has_cudnn_fusion_config() const { return this != internal_default_instance() && _impl_.cudnn_fusion_config_ != nullptr; } inline bool FusionBackendConfig::has_cudnn_fusion_config() const { return _internal_has_cudnn_fusion_config(); } inline void FusionBackendConfig::clear_cudnn_fusion_config() { if (GetArenaForAllocation() == nullptr && _impl_.cudnn_fusion_config_ != nullptr) { delete _impl_.cudnn_fusion_config_; } _impl_.cudnn_fusion_config_ = nullptr; } inline const ::xla::gpu::CuDnnFusionConfig& FusionBackendConfig::_internal_cudnn_fusion_config() const { 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 { // @@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* cudnn_fusion_config) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.cudnn_fusion_config_); } _impl_.cudnn_fusion_config_ = cudnn_fusion_config; if (cudnn_fusion_config) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.FusionBackendConfig.cudnn_fusion_config) } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::release_cudnn_fusion_config() { ::xla::gpu::CuDnnFusionConfig* temp = _impl_.cudnn_fusion_config_; _impl_.cudnn_fusion_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::unsafe_arena_release_cudnn_fusion_config() { // @@protoc_insertion_point(field_release:xla.gpu.FusionBackendConfig.cudnn_fusion_config) ::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() { if (_impl_.cudnn_fusion_config_ == nullptr) { auto* p = CreateMaybeMessage<::xla::gpu::CuDnnFusionConfig>(GetArenaForAllocation()); _impl_.cudnn_fusion_config_ = p; } return _impl_.cudnn_fusion_config_; } inline ::xla::gpu::CuDnnFusionConfig* FusionBackendConfig::mutable_cudnn_fusion_config() { ::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* cudnn_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete _impl_.cudnn_fusion_config_; } if (cudnn_fusion_config) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena(cudnn_fusion_config); if (message_arena != submessage_arena) { cudnn_fusion_config = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, cudnn_fusion_config, submessage_arena); } } else { } _impl_.cudnn_fusion_config_ = cudnn_fusion_config; // @@protoc_insertion_point(field_set_allocated:xla.gpu.FusionBackendConfig.cudnn_fusion_config) } // ------------------------------------------------------------------- // CudnnNormBackendConfig // double epsilon = 1; inline void CudnnNormBackendConfig::clear_epsilon() { _impl_.epsilon_ = 0; } inline double CudnnNormBackendConfig::_internal_epsilon() const { return _impl_.epsilon_; } inline double CudnnNormBackendConfig::epsilon() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.epsilon) return _internal_epsilon(); } inline void CudnnNormBackendConfig::_internal_set_epsilon(double value) { _impl_.epsilon_ = value; } inline void CudnnNormBackendConfig::set_epsilon(double value) { _internal_set_epsilon(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnNormBackendConfig.epsilon) } // .stream_executor.dnn.AlgorithmProto algorithm = 2; inline bool CudnnNormBackendConfig::_internal_has_algorithm() const { return this != internal_default_instance() && _impl_.algorithm_ != nullptr; } inline bool CudnnNormBackendConfig::has_algorithm() const { return _internal_has_algorithm(); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnNormBackendConfig::_internal_algorithm() const { 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 { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnNormBackendConfig::unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = algorithm; if (algorithm) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnNormBackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::release_algorithm() { ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::unsafe_arena_release_algorithm() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnNormBackendConfig.algorithm) ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::_internal_mutable_algorithm() { if (_impl_.algorithm_ == nullptr) { auto* p = CreateMaybeMessage<::stream_executor::dnn::AlgorithmProto>(GetArenaForAllocation()); _impl_.algorithm_ = p; } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnNormBackendConfig::mutable_algorithm() { ::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* algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } if (algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(algorithm)); if (message_arena != submessage_arena) { algorithm = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, algorithm, submessage_arena); } } else { } _impl_.algorithm_ = algorithm; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnNormBackendConfig.algorithm) } // .xla.gpu.CudnnNormBackendConfig.Kind kind = 3; inline void CudnnNormBackendConfig::clear_kind() { _impl_.kind_ = 0; } inline ::xla::gpu::CudnnNormBackendConfig_Kind CudnnNormBackendConfig::_internal_kind() const { return static_cast< ::xla::gpu::CudnnNormBackendConfig_Kind >(_impl_.kind_); } inline ::xla::gpu::CudnnNormBackendConfig_Kind CudnnNormBackendConfig::kind() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnNormBackendConfig.kind) return _internal_kind(); } inline void CudnnNormBackendConfig::_internal_set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value) { _impl_.kind_ = value; } inline void CudnnNormBackendConfig::set_kind(::xla::gpu::CudnnNormBackendConfig_Kind value) { _internal_set_kind(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnNormBackendConfig.kind) } // ------------------------------------------------------------------- // CudnnfMHABackendConfig // .stream_executor.dnn.AlgorithmProto algorithm = 8; inline bool CudnnfMHABackendConfig::_internal_has_algorithm() const { return this != internal_default_instance() && _impl_.algorithm_ != nullptr; } inline bool CudnnfMHABackendConfig::has_algorithm() const { return _internal_has_algorithm(); } inline const ::stream_executor::dnn::AlgorithmProto& CudnnfMHABackendConfig::_internal_algorithm() const { 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 { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.algorithm) return _internal_algorithm(); } inline void CudnnfMHABackendConfig::unsafe_arena_set_allocated_algorithm( ::stream_executor::dnn::AlgorithmProto* algorithm) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } _impl_.algorithm_ = algorithm; if (algorithm) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.algorithm) } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::release_algorithm() { ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::unsafe_arena_release_algorithm() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.algorithm) ::stream_executor::dnn::AlgorithmProto* temp = _impl_.algorithm_; _impl_.algorithm_ = nullptr; return temp; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::_internal_mutable_algorithm() { if (_impl_.algorithm_ == nullptr) { auto* p = CreateMaybeMessage<::stream_executor::dnn::AlgorithmProto>(GetArenaForAllocation()); _impl_.algorithm_ = p; } return _impl_.algorithm_; } inline ::stream_executor::dnn::AlgorithmProto* CudnnfMHABackendConfig::mutable_algorithm() { ::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* algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.algorithm_); } if (algorithm) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(algorithm)); if (message_arena != submessage_arena) { algorithm = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, algorithm, submessage_arena); } } else { } _impl_.algorithm_ = algorithm; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.algorithm) } // double fmha_scale = 10; inline void CudnnfMHABackendConfig::clear_fmha_scale() { _impl_.fmha_scale_ = 0; } inline double CudnnfMHABackendConfig::_internal_fmha_scale() const { return _impl_.fmha_scale_; } inline double CudnnfMHABackendConfig::fmha_scale() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.fmha_scale) return _internal_fmha_scale(); } inline void CudnnfMHABackendConfig::_internal_set_fmha_scale(double value) { _impl_.fmha_scale_ = value; } inline void CudnnfMHABackendConfig::set_fmha_scale(double value) { _internal_set_fmha_scale(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.fmha_scale) } // double dropout_rate = 13; inline void CudnnfMHABackendConfig::clear_dropout_rate() { _impl_.dropout_rate_ = 0; } inline double CudnnfMHABackendConfig::_internal_dropout_rate() const { return _impl_.dropout_rate_; } inline double CudnnfMHABackendConfig::dropout_rate() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.dropout_rate) return _internal_dropout_rate(); } inline void CudnnfMHABackendConfig::_internal_set_dropout_rate(double value) { _impl_.dropout_rate_ = value; } inline void CudnnfMHABackendConfig::set_dropout_rate(double value) { _internal_set_dropout_rate(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.dropout_rate) } // .xla.DotDimensionNumbers bmm1_dot_dimension_numbers = 11; inline bool CudnnfMHABackendConfig::_internal_has_bmm1_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm1_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm1_dot_dimension_numbers() const { return _internal_has_bmm1_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_dot_dimension_numbers() const { 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 { // @@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* bmm1_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_dot_dimension_numbers_); } _impl_.bmm1_dot_dimension_numbers_ = bmm1_dot_dimension_numbers; if (bmm1_dot_dimension_numbers) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm1_dot_dimension_numbers() { ::xla::DotDimensionNumbers* temp = _impl_.bmm1_dot_dimension_numbers_; _impl_.bmm1_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) ::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() { if (_impl_.bmm1_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm1_dot_dimension_numbers_ = p; } return _impl_.bmm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_dot_dimension_numbers() { ::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* bmm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_dot_dimension_numbers_); } if (bmm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm1_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm1_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm1_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm1_dot_dimension_numbers_ = bmm1_dot_dimension_numbers; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm1_dot_dimension_numbers) } // .xla.DotDimensionNumbers bmm2_dot_dimension_numbers = 12; inline bool CudnnfMHABackendConfig::_internal_has_bmm2_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm2_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm2_dot_dimension_numbers() const { return _internal_has_bmm2_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_dot_dimension_numbers() const { 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 { // @@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* bmm2_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_dot_dimension_numbers_); } _impl_.bmm2_dot_dimension_numbers_ = bmm2_dot_dimension_numbers; if (bmm2_dot_dimension_numbers) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::release_bmm2_dot_dimension_numbers() { ::xla::DotDimensionNumbers* temp = _impl_.bmm2_dot_dimension_numbers_; _impl_.bmm2_dot_dimension_numbers_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) ::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() { if (_impl_.bmm2_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm2_dot_dimension_numbers_ = p; } return _impl_.bmm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_dot_dimension_numbers() { ::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* bmm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_dot_dimension_numbers_); } if (bmm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm2_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm2_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm2_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm2_dot_dimension_numbers_ = bmm2_dot_dimension_numbers; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_dot_dimension_numbers) } // .xla.ShapeProto intermediate_tensor_shape = 14; inline bool CudnnfMHABackendConfig::_internal_has_intermediate_tensor_shape() const { return this != internal_default_instance() && _impl_.intermediate_tensor_shape_ != nullptr; } inline bool CudnnfMHABackendConfig::has_intermediate_tensor_shape() const { return _internal_has_intermediate_tensor_shape(); } inline const ::xla::ShapeProto& CudnnfMHABackendConfig::_internal_intermediate_tensor_shape() const { 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 { // @@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* intermediate_tensor_shape) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.intermediate_tensor_shape_); } _impl_.intermediate_tensor_shape_ = intermediate_tensor_shape; if (intermediate_tensor_shape) { } else { } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) } inline ::xla::ShapeProto* CudnnfMHABackendConfig::release_intermediate_tensor_shape() { ::xla::ShapeProto* temp = _impl_.intermediate_tensor_shape_; _impl_.intermediate_tensor_shape_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::unsafe_arena_release_intermediate_tensor_shape() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) ::xla::ShapeProto* temp = _impl_.intermediate_tensor_shape_; _impl_.intermediate_tensor_shape_ = nullptr; return temp; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::_internal_mutable_intermediate_tensor_shape() { if (_impl_.intermediate_tensor_shape_ == nullptr) { auto* p = CreateMaybeMessage<::xla::ShapeProto>(GetArenaForAllocation()); _impl_.intermediate_tensor_shape_ = p; } return _impl_.intermediate_tensor_shape_; } inline ::xla::ShapeProto* CudnnfMHABackendConfig::mutable_intermediate_tensor_shape() { ::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* intermediate_tensor_shape) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.intermediate_tensor_shape_); } if (intermediate_tensor_shape) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(intermediate_tensor_shape)); if (message_arena != submessage_arena) { intermediate_tensor_shape = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, intermediate_tensor_shape, submessage_arena); } } else { } _impl_.intermediate_tensor_shape_ = intermediate_tensor_shape; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.intermediate_tensor_shape) } // .xla.DotDimensionNumbers bmm1_grad_gemm1_dot_dimension_numbers = 16; inline bool CudnnfMHABackendConfig::_internal_has_bmm1_grad_gemm1_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm1_grad_gemm1_dot_dimension_numbers() const { return _internal_has_bmm1_grad_gemm1_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_grad_gemm1_dot_dimension_numbers() const { 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 { // @@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* bmm1_grad_gemm1_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_grad_gemm1_dot_dimension_numbers_); } _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = bmm1_grad_gemm1_dot_dimension_numbers; if (bmm1_grad_gemm1_dot_dimension_numbers) { } else { } // @@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() { ::xla::DotDimensionNumbers* temp = _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<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_grad_gemm1_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm1_dot_dimension_numbers) ::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() { if (_impl_.bmm1_grad_gemm1_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = p; } return _impl_.bmm1_grad_gemm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_grad_gemm1_dot_dimension_numbers() { ::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* bmm1_grad_gemm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_grad_gemm1_dot_dimension_numbers_); } if (bmm1_grad_gemm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm1_grad_gemm1_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm1_grad_gemm1_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm1_grad_gemm1_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm1_grad_gemm1_dot_dimension_numbers_ = bmm1_grad_gemm1_dot_dimension_numbers; // @@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::_internal_has_bmm1_grad_gemm2_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm1_grad_gemm2_dot_dimension_numbers() const { return _internal_has_bmm1_grad_gemm2_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm1_grad_gemm2_dot_dimension_numbers() const { 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 { // @@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* bmm1_grad_gemm2_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_grad_gemm2_dot_dimension_numbers_); } _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = bmm1_grad_gemm2_dot_dimension_numbers; if (bmm1_grad_gemm2_dot_dimension_numbers) { } else { } // @@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() { ::xla::DotDimensionNumbers* temp = _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<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm1_grad_gemm2_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm1_grad_gemm2_dot_dimension_numbers) ::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() { if (_impl_.bmm1_grad_gemm2_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = p; } return _impl_.bmm1_grad_gemm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm1_grad_gemm2_dot_dimension_numbers() { ::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* bmm1_grad_gemm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm1_grad_gemm2_dot_dimension_numbers_); } if (bmm1_grad_gemm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm1_grad_gemm2_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm1_grad_gemm2_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm1_grad_gemm2_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm1_grad_gemm2_dot_dimension_numbers_ = bmm1_grad_gemm2_dot_dimension_numbers; // @@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::_internal_has_bmm2_grad_gemm1_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm2_grad_gemm1_dot_dimension_numbers() const { return _internal_has_bmm2_grad_gemm1_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_grad_gemm1_dot_dimension_numbers() const { 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 { // @@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* bmm2_grad_gemm1_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_grad_gemm1_dot_dimension_numbers_); } _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = bmm2_grad_gemm1_dot_dimension_numbers; if (bmm2_grad_gemm1_dot_dimension_numbers) { } else { } // @@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() { ::xla::DotDimensionNumbers* temp = _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<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_grad_gemm1_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm1_dot_dimension_numbers) ::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() { if (_impl_.bmm2_grad_gemm1_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = p; } return _impl_.bmm2_grad_gemm1_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_grad_gemm1_dot_dimension_numbers() { ::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* bmm2_grad_gemm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_grad_gemm1_dot_dimension_numbers_); } if (bmm2_grad_gemm1_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm2_grad_gemm1_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm2_grad_gemm1_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm2_grad_gemm1_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm2_grad_gemm1_dot_dimension_numbers_ = bmm2_grad_gemm1_dot_dimension_numbers; // @@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::_internal_has_bmm2_grad_gemm2_dot_dimension_numbers() const { return this != internal_default_instance() && _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ != nullptr; } inline bool CudnnfMHABackendConfig::has_bmm2_grad_gemm2_dot_dimension_numbers() const { return _internal_has_bmm2_grad_gemm2_dot_dimension_numbers(); } inline const ::xla::DotDimensionNumbers& CudnnfMHABackendConfig::_internal_bmm2_grad_gemm2_dot_dimension_numbers() const { 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 { // @@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* bmm2_grad_gemm2_dot_dimension_numbers) { if (GetArenaForAllocation() == nullptr) { delete reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_grad_gemm2_dot_dimension_numbers_); } _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = bmm2_grad_gemm2_dot_dimension_numbers; if (bmm2_grad_gemm2_dot_dimension_numbers) { } else { } // @@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() { ::xla::DotDimensionNumbers* temp = _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<::PROTOBUF_NAMESPACE_ID::MessageLite*>(temp); temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); if (GetArenaForAllocation() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::internal::DuplicateIfNonNull(temp); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return temp; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::unsafe_arena_release_bmm2_grad_gemm2_dot_dimension_numbers() { // @@protoc_insertion_point(field_release:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) ::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() { if (_impl_.bmm2_grad_gemm2_dot_dimension_numbers_ == nullptr) { auto* p = CreateMaybeMessage<::xla::DotDimensionNumbers>(GetArenaForAllocation()); _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = p; } return _impl_.bmm2_grad_gemm2_dot_dimension_numbers_; } inline ::xla::DotDimensionNumbers* CudnnfMHABackendConfig::mutable_bmm2_grad_gemm2_dot_dimension_numbers() { ::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* bmm2_grad_gemm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* message_arena = GetArenaForAllocation(); if (message_arena == nullptr) { delete reinterpret_cast< ::PROTOBUF_NAMESPACE_ID::MessageLite*>(_impl_.bmm2_grad_gemm2_dot_dimension_numbers_); } if (bmm2_grad_gemm2_dot_dimension_numbers) { ::PROTOBUF_NAMESPACE_ID::Arena* submessage_arena = ::PROTOBUF_NAMESPACE_ID::Arena::InternalGetOwningArena( reinterpret_cast<::PROTOBUF_NAMESPACE_ID::MessageLite*>(bmm2_grad_gemm2_dot_dimension_numbers)); if (message_arena != submessage_arena) { bmm2_grad_gemm2_dot_dimension_numbers = ::PROTOBUF_NAMESPACE_ID::internal::GetOwnedMessage( message_arena, bmm2_grad_gemm2_dot_dimension_numbers, submessage_arena); } } else { } _impl_.bmm2_grad_gemm2_dot_dimension_numbers_ = bmm2_grad_gemm2_dot_dimension_numbers; // @@protoc_insertion_point(field_set_allocated:xla.gpu.CudnnfMHABackendConfig.bmm2_grad_gemm2_dot_dimension_numbers) } // int64 seed = 15; inline void CudnnfMHABackendConfig::clear_seed() { _impl_.seed_ = int64_t{0}; } inline int64_t CudnnfMHABackendConfig::_internal_seed() const { return _impl_.seed_; } inline int64_t CudnnfMHABackendConfig::seed() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.seed) return _internal_seed(); } inline void CudnnfMHABackendConfig::_internal_set_seed(int64_t value) { _impl_.seed_ = value; } inline void CudnnfMHABackendConfig::set_seed(int64_t value) { _internal_set_seed(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.seed) } // bool is_flash_attention = 20; inline void CudnnfMHABackendConfig::clear_is_flash_attention() { _impl_.is_flash_attention_ = false; } inline bool CudnnfMHABackendConfig::_internal_is_flash_attention() const { return _impl_.is_flash_attention_; } 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::_internal_set_is_flash_attention(bool value) { _impl_.is_flash_attention_ = value; } 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) } // bool is_causal_mask = 21; inline void CudnnfMHABackendConfig::clear_is_causal_mask() { _impl_.is_causal_mask_ = false; } inline bool CudnnfMHABackendConfig::_internal_is_causal_mask() const { return _impl_.is_causal_mask_; } 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::_internal_set_is_causal_mask(bool value) { _impl_.is_causal_mask_ = value; } 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) } // .xla.gpu.CudnnfMHABackendConfig.MaskType mask_type = 22; inline void CudnnfMHABackendConfig::clear_mask_type() { _impl_.mask_type_ = 0; } inline ::xla::gpu::CudnnfMHABackendConfig_MaskType CudnnfMHABackendConfig::_internal_mask_type() const { return static_cast< ::xla::gpu::CudnnfMHABackendConfig_MaskType >(_impl_.mask_type_); } 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::_internal_set_mask_type(::xla::gpu::CudnnfMHABackendConfig_MaskType value) { _impl_.mask_type_ = value; } 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) } // bool force_deterministic = 23; inline void CudnnfMHABackendConfig::clear_force_deterministic() { _impl_.force_deterministic_ = false; } inline bool CudnnfMHABackendConfig::_internal_force_deterministic() const { return _impl_.force_deterministic_; } inline bool CudnnfMHABackendConfig::force_deterministic() const { // @@protoc_insertion_point(field_get:xla.gpu.CudnnfMHABackendConfig.force_deterministic) return _internal_force_deterministic(); } inline void CudnnfMHABackendConfig::_internal_set_force_deterministic(bool value) { _impl_.force_deterministic_ = value; } inline void CudnnfMHABackendConfig::set_force_deterministic(bool value) { _internal_set_force_deterministic(value); // @@protoc_insertion_point(field_set:xla.gpu.CudnnfMHABackendConfig.force_deterministic) } // int32 sliding_window_length = 24; inline void CudnnfMHABackendConfig::clear_sliding_window_length() { _impl_.sliding_window_length_ = 0; } inline int32_t CudnnfMHABackendConfig::_internal_sliding_window_length() const { return _impl_.sliding_window_length_; } 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::_internal_set_sliding_window_length(int32_t value) { _impl_.sliding_window_length_ = value; } 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) } // int32 max_seg_per_batch = 25; inline void CudnnfMHABackendConfig::clear_max_seg_per_batch() { _impl_.max_seg_per_batch_ = 0; } inline int32_t CudnnfMHABackendConfig::_internal_max_seg_per_batch() const { return _impl_.max_seg_per_batch_; } 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::_internal_set_max_seg_per_batch(int32_t value) { _impl_.max_seg_per_batch_ = value; } 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) } // ------------------------------------------------------------------- // CustomCallBackendConfig // string opaque = 1; inline bool CustomCallBackendConfig::_internal_has_opaque() const { return raw_backend_config_oneof_case() == kOpaque; } inline bool CustomCallBackendConfig::has_opaque() const { return _internal_has_opaque(); } inline void CustomCallBackendConfig::set_has_opaque() { _impl_._oneof_case_[0] = kOpaque; } inline void CustomCallBackendConfig::clear_opaque() { if (_internal_has_opaque()) { _impl_.raw_backend_config_oneof_.opaque_.Destroy(); clear_has_raw_backend_config_oneof(); } } inline const std::string& CustomCallBackendConfig::opaque() const { // @@protoc_insertion_point(field_get:xla.gpu.CustomCallBackendConfig.opaque) return _internal_opaque(); } template inline void CustomCallBackendConfig::set_opaque(ArgT0&& arg0, ArgT... args) { if (!_internal_has_opaque()) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } _impl_.raw_backend_config_oneof_.opaque_.Set( static_cast(arg0), args..., GetArenaForAllocation()); // @@protoc_insertion_point(field_set:xla.gpu.CustomCallBackendConfig.opaque) } inline std::string* CustomCallBackendConfig::mutable_opaque() { 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 { if (_internal_has_opaque()) { return _impl_.raw_backend_config_oneof_.opaque_.Get(); } return ::PROTOBUF_NAMESPACE_ID::internal::GetEmptyStringAlreadyInited(); } inline void CustomCallBackendConfig::_internal_set_opaque(const std::string& value) { if (!_internal_has_opaque()) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } _impl_.raw_backend_config_oneof_.opaque_.Set(value, GetArenaForAllocation()); } inline std::string* CustomCallBackendConfig::_internal_mutable_opaque() { if (!_internal_has_opaque()) { clear_raw_backend_config_oneof(); set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitDefault(); } return _impl_.raw_backend_config_oneof_.opaque_.Mutable( GetArenaForAllocation()); } inline std::string* CustomCallBackendConfig::release_opaque() { // @@protoc_insertion_point(field_release:xla.gpu.CustomCallBackendConfig.opaque) if (_internal_has_opaque()) { clear_has_raw_backend_config_oneof(); return _impl_.raw_backend_config_oneof_.opaque_.Release(); } else { return nullptr; } } inline void CustomCallBackendConfig::set_allocated_opaque(std::string* opaque) { if (has_raw_backend_config_oneof()) { clear_raw_backend_config_oneof(); } if (opaque != nullptr) { set_has_opaque(); _impl_.raw_backend_config_oneof_.opaque_.InitAllocated(opaque, GetArenaForAllocation()); } // @@protoc_insertion_point(field_set_allocated:xla.gpu.CustomCallBackendConfig.opaque) } // string attributes = 2; inline bool CustomCallBackendConfig::_internal_has_attributes() const { return raw_backend_config_oneof_case() == kAttributes; } inline bool CustomCallBackendConfig::has_attributes() const { return _internal_has_attributes(); } inline void CustomCallBackendConfig::set_has_attributes() { _impl_._oneof_case_[0] = kAttributes; } inline void CustomCallBackendConfig::clear_attributes() { if (_internal_has_attributes()) { _impl_.raw_backend_config_oneof_.attributes_.Destroy(); clear_has_raw_backend_config_oneof(); } } inline const std::string& CustomCallBackendConfig::attributes() const { // @@protoc_insertion_point(field_get:xla.gpu.CustomCallBackendConfig.attributes) return _internal_attributes(); } template inline void CustomCallBackendConfig::set_attributes(ArgT0&& arg0, ArgT... args) { if (!_internal_has_attributes()) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } _impl_.raw_backend_config_oneof_.attributes_.Set( static_cast(arg0), args..., GetArenaForAllocation()); // @@protoc_insertion_point(field_set:xla.gpu.CustomCallBackendConfig.attributes) } inline std::string* CustomCallBackendConfig::mutable_attributes() { 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 { if (_internal_has_attributes()) { return _impl_.raw_backend_config_oneof_.attributes_.Get(); } return ::PROTOBUF_NAMESPACE_ID::internal::GetEmptyStringAlreadyInited(); } inline void CustomCallBackendConfig::_internal_set_attributes(const std::string& value) { if (!_internal_has_attributes()) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } _impl_.raw_backend_config_oneof_.attributes_.Set(value, GetArenaForAllocation()); } inline std::string* CustomCallBackendConfig::_internal_mutable_attributes() { if (!_internal_has_attributes()) { clear_raw_backend_config_oneof(); set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitDefault(); } return _impl_.raw_backend_config_oneof_.attributes_.Mutable( GetArenaForAllocation()); } inline std::string* CustomCallBackendConfig::release_attributes() { // @@protoc_insertion_point(field_release:xla.gpu.CustomCallBackendConfig.attributes) if (_internal_has_attributes()) { clear_has_raw_backend_config_oneof(); return _impl_.raw_backend_config_oneof_.attributes_.Release(); } else { return nullptr; } } inline void CustomCallBackendConfig::set_allocated_attributes(std::string* attributes) { if (has_raw_backend_config_oneof()) { clear_raw_backend_config_oneof(); } if (attributes != nullptr) { set_has_attributes(); _impl_.raw_backend_config_oneof_.attributes_.InitAllocated(attributes, GetArenaForAllocation()); } // @@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() { _impl_.operation_queue_id_ = int64_t{0}; } inline int64_t GpuBackendConfig::_internal_operation_queue_id() const { return _impl_.operation_queue_id_; } 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::_internal_set_operation_queue_id(int64_t value) { _impl_.operation_queue_id_ = value; } 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) } // repeated int64 wait_on_operation_queues = 2; inline int GpuBackendConfig::_internal_wait_on_operation_queues_size() const { return _impl_.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() { _impl_.wait_on_operation_queues_.Clear(); } inline int64_t GpuBackendConfig::_internal_wait_on_operation_queues(int index) const { return _impl_.wait_on_operation_queues_.Get(index); } 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(index); } inline void GpuBackendConfig::set_wait_on_operation_queues(int index, int64_t value) { _impl_.wait_on_operation_queues_.Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.GpuBackendConfig.wait_on_operation_queues) } inline void GpuBackendConfig::_internal_add_wait_on_operation_queues(int64_t value) { _impl_.wait_on_operation_queues_.Add(value); } inline void GpuBackendConfig::add_wait_on_operation_queues(int64_t value) { _internal_add_wait_on_operation_queues(value); // @@protoc_insertion_point(field_add:xla.gpu.GpuBackendConfig.wait_on_operation_queues) } inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& GpuBackendConfig::_internal_wait_on_operation_queues() const { return _impl_.wait_on_operation_queues_; } inline const ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >& GpuBackendConfig::wait_on_operation_queues() const { // @@protoc_insertion_point(field_list:xla.gpu.GpuBackendConfig.wait_on_operation_queues) return _internal_wait_on_operation_queues(); } inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* GpuBackendConfig::_internal_mutable_wait_on_operation_queues() { return &_impl_.wait_on_operation_queues_; } inline ::PROTOBUF_NAMESPACE_ID::RepeatedField< int64_t >* GpuBackendConfig::mutable_wait_on_operation_queues() { // @@protoc_insertion_point(field_mutable_list:xla.gpu.GpuBackendConfig.wait_on_operation_queues) return _internal_mutable_wait_on_operation_queues(); } // .xla.gpu.CudnnConvBackendConfig cudnn_conv_backend_config = 3; inline bool GpuBackendConfig::_internal_has_cudnn_conv_backend_config() const { return backend_config_case() == kCudnnConvBackendConfig; } inline bool GpuBackendConfig::has_cudnn_conv_backend_config() const { return _internal_has_cudnn_conv_backend_config(); } inline void GpuBackendConfig::set_has_cudnn_conv_backend_config() { _impl_._oneof_case_[0] = kCudnnConvBackendConfig; } inline void GpuBackendConfig::clear_cudnn_conv_backend_config() { if (_internal_has_cudnn_conv_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_cudnn_conv_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnConvBackendConfig* temp = _impl_.backend_config_.cudnn_conv_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_cudnn_conv_backend_config() ? *_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 { // @@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 (_internal_has_cudnn_conv_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnConvBackendConfig* 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* cudnn_conv_backend_config) { clear_backend_config(); if (cudnn_conv_backend_config) { set_has_cudnn_conv_backend_config(); _impl_.backend_config_.cudnn_conv_backend_config_ = cudnn_conv_backend_config; } // @@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 (!_internal_has_cudnn_conv_backend_config()) { clear_backend_config(); set_has_cudnn_conv_backend_config(); _impl_.backend_config_.cudnn_conv_backend_config_ = CreateMaybeMessage< ::xla::gpu::CudnnConvBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.cudnn_conv_backend_config_; } inline ::xla::gpu::CudnnConvBackendConfig* GpuBackendConfig::mutable_cudnn_conv_backend_config() { ::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::_internal_has_gemm_backend_config() const { return backend_config_case() == kGemmBackendConfig; } inline bool GpuBackendConfig::has_gemm_backend_config() const { return _internal_has_gemm_backend_config(); } inline void GpuBackendConfig::set_has_gemm_backend_config() { _impl_._oneof_case_[0] = kGemmBackendConfig; } inline void GpuBackendConfig::clear_gemm_backend_config() { if (_internal_has_gemm_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_gemm_backend_config()) { clear_has_backend_config(); ::xla::gpu::GemmBackendConfig* temp = _impl_.backend_config_.gemm_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_gemm_backend_config() ? *_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 { // @@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 (_internal_has_gemm_backend_config()) { clear_has_backend_config(); ::xla::gpu::GemmBackendConfig* 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* gemm_backend_config) { clear_backend_config(); if (gemm_backend_config) { set_has_gemm_backend_config(); _impl_.backend_config_.gemm_backend_config_ = gemm_backend_config; } // @@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 (!_internal_has_gemm_backend_config()) { clear_backend_config(); set_has_gemm_backend_config(); _impl_.backend_config_.gemm_backend_config_ = CreateMaybeMessage< ::xla::gpu::GemmBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.gemm_backend_config_; } inline ::xla::gpu::GemmBackendConfig* GpuBackendConfig::mutable_gemm_backend_config() { ::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::_internal_has_bitcast_backend_config() const { return backend_config_case() == kBitcastBackendConfig; } inline bool GpuBackendConfig::has_bitcast_backend_config() const { return _internal_has_bitcast_backend_config(); } inline void GpuBackendConfig::set_has_bitcast_backend_config() { _impl_._oneof_case_[0] = kBitcastBackendConfig; } inline void GpuBackendConfig::clear_bitcast_backend_config() { if (_internal_has_bitcast_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_bitcast_backend_config()) { clear_has_backend_config(); ::xla::gpu::BitcastBackendConfig* temp = _impl_.backend_config_.bitcast_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_bitcast_backend_config() ? *_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 { // @@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 (_internal_has_bitcast_backend_config()) { clear_has_backend_config(); ::xla::gpu::BitcastBackendConfig* 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* bitcast_backend_config) { clear_backend_config(); if (bitcast_backend_config) { set_has_bitcast_backend_config(); _impl_.backend_config_.bitcast_backend_config_ = bitcast_backend_config; } // @@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 (!_internal_has_bitcast_backend_config()) { clear_backend_config(); set_has_bitcast_backend_config(); _impl_.backend_config_.bitcast_backend_config_ = CreateMaybeMessage< ::xla::gpu::BitcastBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.bitcast_backend_config_; } inline ::xla::gpu::BitcastBackendConfig* GpuBackendConfig::mutable_bitcast_backend_config() { ::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::_internal_has_collective_backend_config() const { return backend_config_case() == kCollectiveBackendConfig; } inline bool GpuBackendConfig::has_collective_backend_config() const { return _internal_has_collective_backend_config(); } inline void GpuBackendConfig::set_has_collective_backend_config() { _impl_._oneof_case_[0] = kCollectiveBackendConfig; } inline void GpuBackendConfig::clear_collective_backend_config() { if (_internal_has_collective_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_collective_backend_config()) { clear_has_backend_config(); ::xla::gpu::CollectiveBackendConfig* temp = _impl_.backend_config_.collective_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_collective_backend_config() ? *_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 { // @@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 (_internal_has_collective_backend_config()) { clear_has_backend_config(); ::xla::gpu::CollectiveBackendConfig* 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* collective_backend_config) { clear_backend_config(); if (collective_backend_config) { set_has_collective_backend_config(); _impl_.backend_config_.collective_backend_config_ = collective_backend_config; } // @@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 (!_internal_has_collective_backend_config()) { clear_backend_config(); set_has_collective_backend_config(); _impl_.backend_config_.collective_backend_config_ = CreateMaybeMessage< ::xla::gpu::CollectiveBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.collective_backend_config_; } inline ::xla::gpu::CollectiveBackendConfig* GpuBackendConfig::mutable_collective_backend_config() { ::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::_internal_has_fusion_backend_config() const { return backend_config_case() == kFusionBackendConfig; } inline bool GpuBackendConfig::has_fusion_backend_config() const { return _internal_has_fusion_backend_config(); } inline void GpuBackendConfig::set_has_fusion_backend_config() { _impl_._oneof_case_[0] = kFusionBackendConfig; } inline void GpuBackendConfig::clear_fusion_backend_config() { if (_internal_has_fusion_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_fusion_backend_config()) { clear_has_backend_config(); ::xla::gpu::FusionBackendConfig* temp = _impl_.backend_config_.fusion_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_fusion_backend_config() ? *_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 { // @@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 (_internal_has_fusion_backend_config()) { clear_has_backend_config(); ::xla::gpu::FusionBackendConfig* 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* fusion_backend_config) { clear_backend_config(); if (fusion_backend_config) { set_has_fusion_backend_config(); _impl_.backend_config_.fusion_backend_config_ = fusion_backend_config; } // @@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 (!_internal_has_fusion_backend_config()) { clear_backend_config(); set_has_fusion_backend_config(); _impl_.backend_config_.fusion_backend_config_ = CreateMaybeMessage< ::xla::gpu::FusionBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.fusion_backend_config_; } inline ::xla::gpu::FusionBackendConfig* GpuBackendConfig::mutable_fusion_backend_config() { ::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::_internal_has_cudnn_norm_backend_config() const { return backend_config_case() == kCudnnNormBackendConfig; } inline bool GpuBackendConfig::has_cudnn_norm_backend_config() const { return _internal_has_cudnn_norm_backend_config(); } inline void GpuBackendConfig::set_has_cudnn_norm_backend_config() { _impl_._oneof_case_[0] = kCudnnNormBackendConfig; } inline void GpuBackendConfig::clear_cudnn_norm_backend_config() { if (_internal_has_cudnn_norm_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_cudnn_norm_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnNormBackendConfig* temp = _impl_.backend_config_.cudnn_norm_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_cudnn_norm_backend_config() ? *_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 { // @@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 (_internal_has_cudnn_norm_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnNormBackendConfig* 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* cudnn_norm_backend_config) { clear_backend_config(); if (cudnn_norm_backend_config) { set_has_cudnn_norm_backend_config(); _impl_.backend_config_.cudnn_norm_backend_config_ = cudnn_norm_backend_config; } // @@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 (!_internal_has_cudnn_norm_backend_config()) { clear_backend_config(); set_has_cudnn_norm_backend_config(); _impl_.backend_config_.cudnn_norm_backend_config_ = CreateMaybeMessage< ::xla::gpu::CudnnNormBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.cudnn_norm_backend_config_; } inline ::xla::gpu::CudnnNormBackendConfig* GpuBackendConfig::mutable_cudnn_norm_backend_config() { ::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::_internal_has_cudnn_fmha_backend_config() const { return backend_config_case() == kCudnnFmhaBackendConfig; } inline bool GpuBackendConfig::has_cudnn_fmha_backend_config() const { return _internal_has_cudnn_fmha_backend_config(); } inline void GpuBackendConfig::set_has_cudnn_fmha_backend_config() { _impl_._oneof_case_[0] = kCudnnFmhaBackendConfig; } inline void GpuBackendConfig::clear_cudnn_fmha_backend_config() { if (_internal_has_cudnn_fmha_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_cudnn_fmha_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnfMHABackendConfig* temp = _impl_.backend_config_.cudnn_fmha_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_cudnn_fmha_backend_config() ? *_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 { // @@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 (_internal_has_cudnn_fmha_backend_config()) { clear_has_backend_config(); ::xla::gpu::CudnnfMHABackendConfig* 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* cudnn_fmha_backend_config) { clear_backend_config(); if (cudnn_fmha_backend_config) { set_has_cudnn_fmha_backend_config(); _impl_.backend_config_.cudnn_fmha_backend_config_ = cudnn_fmha_backend_config; } // @@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 (!_internal_has_cudnn_fmha_backend_config()) { clear_backend_config(); set_has_cudnn_fmha_backend_config(); _impl_.backend_config_.cudnn_fmha_backend_config_ = CreateMaybeMessage< ::xla::gpu::CudnnfMHABackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.cudnn_fmha_backend_config_; } inline ::xla::gpu::CudnnfMHABackendConfig* GpuBackendConfig::mutable_cudnn_fmha_backend_config() { ::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::_internal_has_custom_call_backend_config() const { return backend_config_case() == kCustomCallBackendConfig; } inline bool GpuBackendConfig::has_custom_call_backend_config() const { return _internal_has_custom_call_backend_config(); } inline void GpuBackendConfig::set_has_custom_call_backend_config() { _impl_._oneof_case_[0] = kCustomCallBackendConfig; } inline void GpuBackendConfig::clear_custom_call_backend_config() { if (_internal_has_custom_call_backend_config()) { if (GetArenaForAllocation() == nullptr) { delete _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 (_internal_has_custom_call_backend_config()) { clear_has_backend_config(); ::xla::gpu::CustomCallBackendConfig* temp = _impl_.backend_config_.custom_call_backend_config_; if (GetArenaForAllocation() != nullptr) { temp = ::PROTOBUF_NAMESPACE_ID::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 _internal_has_custom_call_backend_config() ? *_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 { // @@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 (_internal_has_custom_call_backend_config()) { clear_has_backend_config(); ::xla::gpu::CustomCallBackendConfig* 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* custom_call_backend_config) { clear_backend_config(); if (custom_call_backend_config) { set_has_custom_call_backend_config(); _impl_.backend_config_.custom_call_backend_config_ = custom_call_backend_config; } // @@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 (!_internal_has_custom_call_backend_config()) { clear_backend_config(); set_has_custom_call_backend_config(); _impl_.backend_config_.custom_call_backend_config_ = CreateMaybeMessage< ::xla::gpu::CustomCallBackendConfig >(GetArenaForAllocation()); } return _impl_.backend_config_.custom_call_backend_config_; } inline ::xla::gpu::CustomCallBackendConfig* GpuBackendConfig::mutable_custom_call_backend_config() { ::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() { _impl_.force_earliest_schedule_ = false; } inline bool GpuBackendConfig::_internal_force_earliest_schedule() const { return _impl_.force_earliest_schedule_; } 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::_internal_set_force_earliest_schedule(bool value) { _impl_.force_earliest_schedule_ = value; } 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::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 PROTOBUF_NAMESPACE_OPEN 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::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(); } PROTOBUF_NAMESPACE_CLOSE // @@protoc_insertion_point(global_scope) #include #endif // GOOGLE_PROTOBUF_INCLUDED_GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fbackend_5fconfigs_2eproto