// Generated by the protocol buffer compiler. DO NOT EDIT! // NO CHECKED-IN PROTOBUF GENCODE // source: xla/service/gpu/autotuning/gpu_autotuning.proto // Protobuf C++ Version: 5.28.3 #ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto_2epb_2eh #define GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto_2epb_2eh #include #include #include #include #include "google/protobuf/runtime_version.h" #if PROTOBUF_VERSION != 5028003 #error "Protobuf C++ gencode is built with an incompatible version of" #error "Protobuf C++ headers/runtime. See" #error "https://protobuf.dev/support/cross-version-runtime-guarantee/#cpp" #endif #include "google/protobuf/io/coded_stream.h" #include "google/protobuf/arena.h" #include "google/protobuf/arenastring.h" #include "google/protobuf/generated_message_tctable_decl.h" #include "google/protobuf/generated_message_util.h" #include "google/protobuf/metadata_lite.h" #include "google/protobuf/generated_message_reflection.h" #include "google/protobuf/message.h" #include "google/protobuf/repeated_field.h" // IWYU pragma: export #include "google/protobuf/extension_set.h" // IWYU pragma: export #include "google/protobuf/unknown_field_set.h" #include "xla/autotuning.pb.h" #include "xla/service/gpu/backend_configs.pb.h" #include "xla/service/hlo.pb.h" #include "xla/xla_data.pb.h" // @@protoc_insertion_point(includes) // Must be included last. #include "google/protobuf/port_def.inc" #define PROTOBUF_INTERNAL_EXPORT_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto PROTOBUF_EXPORT namespace google { namespace protobuf { namespace internal { class AnyMetadata; } // namespace internal } // namespace protobuf } // namespace google // Internal implementation detail -- do not use these members. struct PROTOBUF_EXPORT TableStruct_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto { static const ::uint32_t offsets[]; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable descriptor_table_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto; namespace xla { namespace gpu { class AlgorithmDenylist; struct AlgorithmDenylistDefaultTypeInternal; PROTOBUF_EXPORT extern AlgorithmDenylistDefaultTypeInternal _AlgorithmDenylist_default_instance_; class AlgorithmDenylistEntry; struct AlgorithmDenylistEntryDefaultTypeInternal; PROTOBUF_EXPORT extern AlgorithmDenylistEntryDefaultTypeInternal _AlgorithmDenylistEntry_default_instance_; class ConvInstructionLog; struct ConvInstructionLogDefaultTypeInternal; PROTOBUF_EXPORT extern ConvInstructionLogDefaultTypeInternal _ConvInstructionLog_default_instance_; class DenylistedAlgorithm; struct DenylistedAlgorithmDefaultTypeInternal; PROTOBUF_EXPORT extern DenylistedAlgorithmDefaultTypeInternal _DenylistedAlgorithm_default_instance_; } // namespace gpu } // namespace xla namespace google { namespace protobuf { } // namespace protobuf } // namespace google namespace xla { namespace gpu { // =================================================================== // ------------------------------------------------------------------- class PROTOBUF_EXPORT DenylistedAlgorithm final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DenylistedAlgorithm) */ { public: inline DenylistedAlgorithm() : DenylistedAlgorithm(nullptr) {} ~DenylistedAlgorithm() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR DenylistedAlgorithm( ::google::protobuf::internal::ConstantInitialized); inline DenylistedAlgorithm(const DenylistedAlgorithm& from) : DenylistedAlgorithm(nullptr, from) {} inline DenylistedAlgorithm(DenylistedAlgorithm&& from) noexcept : DenylistedAlgorithm(nullptr, std::move(from)) {} inline DenylistedAlgorithm& operator=(const DenylistedAlgorithm& from) { CopyFrom(from); return *this; } inline DenylistedAlgorithm& operator=(DenylistedAlgorithm&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const DenylistedAlgorithm& default_instance() { return *internal_default_instance(); } static inline const DenylistedAlgorithm* internal_default_instance() { return reinterpret_cast( &_DenylistedAlgorithm_default_instance_); } static constexpr int kIndexInFileMessages = 1; friend void swap(DenylistedAlgorithm& a, DenylistedAlgorithm& b) { a.Swap(&b); } inline void Swap(DenylistedAlgorithm* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DenylistedAlgorithm* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DenylistedAlgorithm* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DenylistedAlgorithm& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DenylistedAlgorithm& from) { DenylistedAlgorithm::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(DenylistedAlgorithm* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.DenylistedAlgorithm"; } protected: explicit DenylistedAlgorithm(::google::protobuf::Arena* arena); DenylistedAlgorithm(::google::protobuf::Arena* arena, const DenylistedAlgorithm& from); DenylistedAlgorithm(::google::protobuf::Arena* arena, DenylistedAlgorithm&& from) noexcept : DenylistedAlgorithm(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kIdFieldNumber = 1, kTensorOpsFieldNumber = 2, }; // int64 id = 1; void clear_id() ; ::int64_t id() const; void set_id(::int64_t value); private: ::int64_t _internal_id() const; void _internal_set_id(::int64_t value); public: // bool tensor_ops = 2; void clear_tensor_ops() ; bool tensor_ops() const; void set_tensor_ops(bool value); private: bool _internal_tensor_ops() const; void _internal_set_tensor_ops(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.DenylistedAlgorithm) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 1, 2, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const DenylistedAlgorithm& from_msg); ::int64_t id_; bool tensor_ops_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ConvInstructionLog final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ConvInstructionLog) */ { public: inline ConvInstructionLog() : ConvInstructionLog(nullptr) {} ~ConvInstructionLog() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ConvInstructionLog( ::google::protobuf::internal::ConstantInitialized); inline ConvInstructionLog(const ConvInstructionLog& from) : ConvInstructionLog(nullptr, from) {} inline ConvInstructionLog(ConvInstructionLog&& from) noexcept : ConvInstructionLog(nullptr, std::move(from)) {} inline ConvInstructionLog& operator=(const ConvInstructionLog& from) { CopyFrom(from); return *this; } inline ConvInstructionLog& operator=(ConvInstructionLog&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const ConvInstructionLog& default_instance() { return *internal_default_instance(); } static inline const ConvInstructionLog* internal_default_instance() { return reinterpret_cast( &_ConvInstructionLog_default_instance_); } static constexpr int kIndexInFileMessages = 0; friend void swap(ConvInstructionLog& a, ConvInstructionLog& b) { a.Swap(&b); } inline void Swap(ConvInstructionLog* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(ConvInstructionLog* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ConvInstructionLog* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ConvInstructionLog& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ConvInstructionLog& from) { ConvInstructionLog::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(ConvInstructionLog* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ConvInstructionLog"; } protected: explicit ConvInstructionLog(::google::protobuf::Arena* arena); ConvInstructionLog(::google::protobuf::Arena* arena, const ConvInstructionLog& from); ConvInstructionLog(::google::protobuf::Arena* arena, ConvInstructionLog&& from) noexcept : ConvInstructionLog(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kOperandShapesFieldNumber = 2, kResultAddressesFieldNumber = 3, kOperandAddressesFieldNumber = 4, kInstructionFieldNumber = 1, }; // repeated .xla.ShapeProto operand_shapes = 2; int operand_shapes_size() const; private: int _internal_operand_shapes_size() const; public: void clear_operand_shapes() ; ::xla::ShapeProto* mutable_operand_shapes(int index); ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>* mutable_operand_shapes(); private: const ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>& _internal_operand_shapes() const; ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>* _internal_mutable_operand_shapes(); public: const ::xla::ShapeProto& operand_shapes(int index) const; ::xla::ShapeProto* add_operand_shapes(); const ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>& operand_shapes() const; // repeated uint64 result_addresses = 3; int result_addresses_size() const; private: int _internal_result_addresses_size() const; public: void clear_result_addresses() ; ::uint64_t result_addresses(int index) const; void set_result_addresses(int index, ::uint64_t value); void add_result_addresses(::uint64_t value); const ::google::protobuf::RepeatedField<::uint64_t>& result_addresses() const; ::google::protobuf::RepeatedField<::uint64_t>* mutable_result_addresses(); private: const ::google::protobuf::RepeatedField<::uint64_t>& _internal_result_addresses() const; ::google::protobuf::RepeatedField<::uint64_t>* _internal_mutable_result_addresses(); public: // repeated uint64 operand_addresses = 4; int operand_addresses_size() const; private: int _internal_operand_addresses_size() const; public: void clear_operand_addresses() ; ::uint64_t operand_addresses(int index) const; void set_operand_addresses(int index, ::uint64_t value); void add_operand_addresses(::uint64_t value); const ::google::protobuf::RepeatedField<::uint64_t>& operand_addresses() const; ::google::protobuf::RepeatedField<::uint64_t>* mutable_operand_addresses(); private: const ::google::protobuf::RepeatedField<::uint64_t>& _internal_operand_addresses() const; ::google::protobuf::RepeatedField<::uint64_t>* _internal_mutable_operand_addresses(); public: // .xla.HloInstructionProto instruction = 1; bool has_instruction() const; void clear_instruction() ; const ::xla::HloInstructionProto& instruction() const; PROTOBUF_NODISCARD ::xla::HloInstructionProto* release_instruction(); ::xla::HloInstructionProto* mutable_instruction(); void set_allocated_instruction(::xla::HloInstructionProto* value); void unsafe_arena_set_allocated_instruction(::xla::HloInstructionProto* value); ::xla::HloInstructionProto* unsafe_arena_release_instruction(); private: const ::xla::HloInstructionProto& _internal_instruction() const; ::xla::HloInstructionProto* _internal_mutable_instruction(); public: // @@protoc_insertion_point(class_scope:xla.gpu.ConvInstructionLog) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 4, 2, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const ConvInstructionLog& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::RepeatedPtrField< ::xla::ShapeProto > operand_shapes_; ::google::protobuf::RepeatedField<::uint64_t> result_addresses_; mutable ::google::protobuf::internal::CachedSize _result_addresses_cached_byte_size_; ::google::protobuf::RepeatedField<::uint64_t> operand_addresses_; mutable ::google::protobuf::internal::CachedSize _operand_addresses_cached_byte_size_; ::xla::HloInstructionProto* instruction_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT AlgorithmDenylistEntry final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.AlgorithmDenylistEntry) */ { public: inline AlgorithmDenylistEntry() : AlgorithmDenylistEntry(nullptr) {} ~AlgorithmDenylistEntry() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR AlgorithmDenylistEntry( ::google::protobuf::internal::ConstantInitialized); inline AlgorithmDenylistEntry(const AlgorithmDenylistEntry& from) : AlgorithmDenylistEntry(nullptr, from) {} inline AlgorithmDenylistEntry(AlgorithmDenylistEntry&& from) noexcept : AlgorithmDenylistEntry(nullptr, std::move(from)) {} inline AlgorithmDenylistEntry& operator=(const AlgorithmDenylistEntry& from) { CopyFrom(from); return *this; } inline AlgorithmDenylistEntry& operator=(AlgorithmDenylistEntry&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const AlgorithmDenylistEntry& default_instance() { return *internal_default_instance(); } static inline const AlgorithmDenylistEntry* internal_default_instance() { return reinterpret_cast( &_AlgorithmDenylistEntry_default_instance_); } static constexpr int kIndexInFileMessages = 2; friend void swap(AlgorithmDenylistEntry& a, AlgorithmDenylistEntry& b) { a.Swap(&b); } inline void Swap(AlgorithmDenylistEntry* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(AlgorithmDenylistEntry* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- AlgorithmDenylistEntry* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const AlgorithmDenylistEntry& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const AlgorithmDenylistEntry& from) { AlgorithmDenylistEntry::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(AlgorithmDenylistEntry* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.AlgorithmDenylistEntry"; } protected: explicit AlgorithmDenylistEntry(::google::protobuf::Arena* arena); AlgorithmDenylistEntry(::google::protobuf::Arena* arena, const AlgorithmDenylistEntry& from); AlgorithmDenylistEntry(::google::protobuf::Arena* arena, AlgorithmDenylistEntry&& from) noexcept : AlgorithmDenylistEntry(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kAlgosFieldNumber = 4, kHloFieldNumber = 1, kBlasVersionFieldNumber = 5, kCcFieldNumber = 2, kCudnnVersionFieldNumber = 3, kBackendConfigFieldNumber = 6, }; // repeated .xla.gpu.DenylistedAlgorithm algos = 4; int algos_size() const; private: int _internal_algos_size() const; public: void clear_algos() ; ::xla::gpu::DenylistedAlgorithm* mutable_algos(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>* mutable_algos(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>& _internal_algos() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>* _internal_mutable_algos(); public: const ::xla::gpu::DenylistedAlgorithm& algos(int index) const; ::xla::gpu::DenylistedAlgorithm* add_algos(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>& algos() const; // string hlo = 1; void clear_hlo() ; const std::string& hlo() const; template void set_hlo(Arg_&& arg, Args_... args); std::string* mutable_hlo(); PROTOBUF_NODISCARD std::string* release_hlo(); void set_allocated_hlo(std::string* value); private: const std::string& _internal_hlo() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_hlo( const std::string& value); std::string* _internal_mutable_hlo(); public: // string blas_version = 5; void clear_blas_version() ; const std::string& blas_version() const; template void set_blas_version(Arg_&& arg, Args_... args); std::string* mutable_blas_version(); PROTOBUF_NODISCARD std::string* release_blas_version(); void set_allocated_blas_version(std::string* value); private: const std::string& _internal_blas_version() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_blas_version( const std::string& value); std::string* _internal_mutable_blas_version(); public: // .xla.ComputeCapability cc = 2; bool has_cc() const; void clear_cc() ; const ::xla::ComputeCapability& cc() const; PROTOBUF_NODISCARD ::xla::ComputeCapability* release_cc(); ::xla::ComputeCapability* mutable_cc(); void set_allocated_cc(::xla::ComputeCapability* value); void unsafe_arena_set_allocated_cc(::xla::ComputeCapability* value); ::xla::ComputeCapability* unsafe_arena_release_cc(); private: const ::xla::ComputeCapability& _internal_cc() const; ::xla::ComputeCapability* _internal_mutable_cc(); public: // .xla.CudnnVersion cudnn_version = 3; bool has_cudnn_version() const; void clear_cudnn_version() ; const ::xla::CudnnVersion& cudnn_version() const; PROTOBUF_NODISCARD ::xla::CudnnVersion* release_cudnn_version(); ::xla::CudnnVersion* mutable_cudnn_version(); void set_allocated_cudnn_version(::xla::CudnnVersion* value); void unsafe_arena_set_allocated_cudnn_version(::xla::CudnnVersion* value); ::xla::CudnnVersion* unsafe_arena_release_cudnn_version(); private: const ::xla::CudnnVersion& _internal_cudnn_version() const; ::xla::CudnnVersion* _internal_mutable_cudnn_version(); public: // .xla.gpu.GpuBackendConfig backend_config = 6; bool has_backend_config() const; void clear_backend_config() ; const ::xla::gpu::GpuBackendConfig& backend_config() const; PROTOBUF_NODISCARD ::xla::gpu::GpuBackendConfig* release_backend_config(); ::xla::gpu::GpuBackendConfig* mutable_backend_config(); void set_allocated_backend_config(::xla::gpu::GpuBackendConfig* value); void unsafe_arena_set_allocated_backend_config(::xla::gpu::GpuBackendConfig* value); ::xla::gpu::GpuBackendConfig* unsafe_arena_release_backend_config(); private: const ::xla::gpu::GpuBackendConfig& _internal_backend_config() const; ::xla::gpu::GpuBackendConfig* _internal_mutable_backend_config(); public: // @@protoc_insertion_point(class_scope:xla.gpu.AlgorithmDenylistEntry) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 6, 4, 54, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const AlgorithmDenylistEntry& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::RepeatedPtrField< ::xla::gpu::DenylistedAlgorithm > algos_; ::google::protobuf::internal::ArenaStringPtr hlo_; ::google::protobuf::internal::ArenaStringPtr blas_version_; ::xla::ComputeCapability* cc_; ::xla::CudnnVersion* cudnn_version_; ::xla::gpu::GpuBackendConfig* backend_config_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT AlgorithmDenylist final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.AlgorithmDenylist) */ { public: inline AlgorithmDenylist() : AlgorithmDenylist(nullptr) {} ~AlgorithmDenylist() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR AlgorithmDenylist( ::google::protobuf::internal::ConstantInitialized); inline AlgorithmDenylist(const AlgorithmDenylist& from) : AlgorithmDenylist(nullptr, from) {} inline AlgorithmDenylist(AlgorithmDenylist&& from) noexcept : AlgorithmDenylist(nullptr, std::move(from)) {} inline AlgorithmDenylist& operator=(const AlgorithmDenylist& from) { CopyFrom(from); return *this; } inline AlgorithmDenylist& operator=(AlgorithmDenylist&& from) noexcept { if (this == &from) return *this; if (GetArena() == from.GetArena() #ifdef PROTOBUF_FORCE_COPY_IN_MOVE && GetArena() != nullptr #endif // !PROTOBUF_FORCE_COPY_IN_MOVE ) { InternalSwap(&from); } else { CopyFrom(from); } return *this; } inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance); } inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* GetReflection() { return default_instance().GetMetadata().reflection; } static const AlgorithmDenylist& default_instance() { return *internal_default_instance(); } static inline const AlgorithmDenylist* internal_default_instance() { return reinterpret_cast( &_AlgorithmDenylist_default_instance_); } static constexpr int kIndexInFileMessages = 3; friend void swap(AlgorithmDenylist& a, AlgorithmDenylist& b) { a.Swap(&b); } inline void Swap(AlgorithmDenylist* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() != nullptr && GetArena() == other->GetArena()) { #else // PROTOBUF_FORCE_COPY_IN_SWAP if (GetArena() == other->GetArena()) { #endif // !PROTOBUF_FORCE_COPY_IN_SWAP InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(AlgorithmDenylist* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- AlgorithmDenylist* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const AlgorithmDenylist& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const AlgorithmDenylist& from) { AlgorithmDenylist::MergeImpl(*this, from); } private: static void MergeImpl( ::google::protobuf::MessageLite& to_msg, const ::google::protobuf::MessageLite& from_msg); public: bool IsInitialized() const { return true; } ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) private: static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg); static ::uint8_t* _InternalSerialize( const MessageLite& msg, ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* _InternalSerialize( ::uint8_t* target, ::google::protobuf::io::EpsCopyOutputStream* stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* arena); void SharedDtor(); void InternalSwap(AlgorithmDenylist* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.AlgorithmDenylist"; } protected: explicit AlgorithmDenylist(::google::protobuf::Arena* arena); AlgorithmDenylist(::google::protobuf::Arena* arena, const AlgorithmDenylist& from); AlgorithmDenylist(::google::protobuf::Arena* arena, AlgorithmDenylist&& from) noexcept : AlgorithmDenylist(arena) { *this = ::std::move(from); } const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL; static const ::google::protobuf::Message::ClassDataFull _class_data_; public: ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kEntriesFieldNumber = 1, }; // repeated .xla.gpu.AlgorithmDenylistEntry entries = 1; int entries_size() const; private: int _internal_entries_size() const; public: void clear_entries() ; ::xla::gpu::AlgorithmDenylistEntry* mutable_entries(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>* mutable_entries(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>& _internal_entries() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>* _internal_mutable_entries(); public: const ::xla::gpu::AlgorithmDenylistEntry& entries(int index) const; ::xla::gpu::AlgorithmDenylistEntry* add_entries(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>& entries() const; // @@protoc_insertion_point(class_scope:xla.gpu.AlgorithmDenylist) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 0, 1, 1, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_( ::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena); inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* arena, const Impl_& from, const AlgorithmDenylist& from_msg); ::google::protobuf::RepeatedPtrField< ::xla::gpu::AlgorithmDenylistEntry > entries_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto; }; // =================================================================== // =================================================================== #ifdef __GNUC__ #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif // __GNUC__ // ------------------------------------------------------------------- // ConvInstructionLog // .xla.HloInstructionProto instruction = 1; inline bool ConvInstructionLog::has_instruction() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.instruction_ != nullptr); return value; } inline const ::xla::HloInstructionProto& ConvInstructionLog::_internal_instruction() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::HloInstructionProto* p = _impl_.instruction_; return p != nullptr ? *p : reinterpret_cast(::xla::_HloInstructionProto_default_instance_); } inline const ::xla::HloInstructionProto& ConvInstructionLog::instruction() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ConvInstructionLog.instruction) return _internal_instruction(); } inline void ConvInstructionLog::unsafe_arena_set_allocated_instruction(::xla::HloInstructionProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.instruction_); } _impl_.instruction_ = reinterpret_cast<::xla::HloInstructionProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ConvInstructionLog.instruction) } inline ::xla::HloInstructionProto* ConvInstructionLog::release_instruction() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::HloInstructionProto* released = _impl_.instruction_; _impl_.instruction_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::HloInstructionProto* ConvInstructionLog::unsafe_arena_release_instruction() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ConvInstructionLog.instruction) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::HloInstructionProto* temp = _impl_.instruction_; _impl_.instruction_ = nullptr; return temp; } inline ::xla::HloInstructionProto* ConvInstructionLog::_internal_mutable_instruction() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.instruction_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::HloInstructionProto>(GetArena()); _impl_.instruction_ = reinterpret_cast<::xla::HloInstructionProto*>(p); } return _impl_.instruction_; } inline ::xla::HloInstructionProto* ConvInstructionLog::mutable_instruction() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::HloInstructionProto* _msg = _internal_mutable_instruction(); // @@protoc_insertion_point(field_mutable:xla.gpu.ConvInstructionLog.instruction) return _msg; } inline void ConvInstructionLog::set_allocated_instruction(::xla::HloInstructionProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.instruction_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.instruction_ = reinterpret_cast<::xla::HloInstructionProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.ConvInstructionLog.instruction) } // repeated .xla.ShapeProto operand_shapes = 2; inline int ConvInstructionLog::_internal_operand_shapes_size() const { return _internal_operand_shapes().size(); } inline int ConvInstructionLog::operand_shapes_size() const { return _internal_operand_shapes_size(); } inline ::xla::ShapeProto* ConvInstructionLog::mutable_operand_shapes(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.ConvInstructionLog.operand_shapes) return _internal_mutable_operand_shapes()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>* ConvInstructionLog::mutable_operand_shapes() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.ConvInstructionLog.operand_shapes) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_operand_shapes(); } inline const ::xla::ShapeProto& ConvInstructionLog::operand_shapes(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ConvInstructionLog.operand_shapes) return _internal_operand_shapes().Get(index); } inline ::xla::ShapeProto* ConvInstructionLog::add_operand_shapes() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::ShapeProto* _add = _internal_mutable_operand_shapes()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.ConvInstructionLog.operand_shapes) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>& ConvInstructionLog::operand_shapes() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.ConvInstructionLog.operand_shapes) return _internal_operand_shapes(); } inline const ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>& ConvInstructionLog::_internal_operand_shapes() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.operand_shapes_; } inline ::google::protobuf::RepeatedPtrField<::xla::ShapeProto>* ConvInstructionLog::_internal_mutable_operand_shapes() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.operand_shapes_; } // repeated uint64 result_addresses = 3; inline int ConvInstructionLog::_internal_result_addresses_size() const { return _internal_result_addresses().size(); } inline int ConvInstructionLog::result_addresses_size() const { return _internal_result_addresses_size(); } inline void ConvInstructionLog::clear_result_addresses() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.result_addresses_.Clear(); } inline ::uint64_t ConvInstructionLog::result_addresses(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.ConvInstructionLog.result_addresses) return _internal_result_addresses().Get(index); } inline void ConvInstructionLog::set_result_addresses(int index, ::uint64_t value) { _internal_mutable_result_addresses()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.ConvInstructionLog.result_addresses) } inline void ConvInstructionLog::add_result_addresses(::uint64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_result_addresses()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.ConvInstructionLog.result_addresses) } inline const ::google::protobuf::RepeatedField<::uint64_t>& ConvInstructionLog::result_addresses() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.ConvInstructionLog.result_addresses) return _internal_result_addresses(); } inline ::google::protobuf::RepeatedField<::uint64_t>* ConvInstructionLog::mutable_result_addresses() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.ConvInstructionLog.result_addresses) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_result_addresses(); } inline const ::google::protobuf::RepeatedField<::uint64_t>& ConvInstructionLog::_internal_result_addresses() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.result_addresses_; } inline ::google::protobuf::RepeatedField<::uint64_t>* ConvInstructionLog::_internal_mutable_result_addresses() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.result_addresses_; } // repeated uint64 operand_addresses = 4; inline int ConvInstructionLog::_internal_operand_addresses_size() const { return _internal_operand_addresses().size(); } inline int ConvInstructionLog::operand_addresses_size() const { return _internal_operand_addresses_size(); } inline void ConvInstructionLog::clear_operand_addresses() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.operand_addresses_.Clear(); } inline ::uint64_t ConvInstructionLog::operand_addresses(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.ConvInstructionLog.operand_addresses) return _internal_operand_addresses().Get(index); } inline void ConvInstructionLog::set_operand_addresses(int index, ::uint64_t value) { _internal_mutable_operand_addresses()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.ConvInstructionLog.operand_addresses) } inline void ConvInstructionLog::add_operand_addresses(::uint64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_operand_addresses()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.ConvInstructionLog.operand_addresses) } inline const ::google::protobuf::RepeatedField<::uint64_t>& ConvInstructionLog::operand_addresses() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.ConvInstructionLog.operand_addresses) return _internal_operand_addresses(); } inline ::google::protobuf::RepeatedField<::uint64_t>* ConvInstructionLog::mutable_operand_addresses() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.ConvInstructionLog.operand_addresses) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_operand_addresses(); } inline const ::google::protobuf::RepeatedField<::uint64_t>& ConvInstructionLog::_internal_operand_addresses() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.operand_addresses_; } inline ::google::protobuf::RepeatedField<::uint64_t>* ConvInstructionLog::_internal_mutable_operand_addresses() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.operand_addresses_; } // ------------------------------------------------------------------- // DenylistedAlgorithm // int64 id = 1; inline void DenylistedAlgorithm::clear_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.id_ = ::int64_t{0}; } inline ::int64_t DenylistedAlgorithm::id() const { // @@protoc_insertion_point(field_get:xla.gpu.DenylistedAlgorithm.id) return _internal_id(); } inline void DenylistedAlgorithm::set_id(::int64_t value) { _internal_set_id(value); // @@protoc_insertion_point(field_set:xla.gpu.DenylistedAlgorithm.id) } inline ::int64_t DenylistedAlgorithm::_internal_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.id_; } inline void DenylistedAlgorithm::_internal_set_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.id_ = value; } // bool tensor_ops = 2; inline void DenylistedAlgorithm::clear_tensor_ops() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.tensor_ops_ = false; } inline bool DenylistedAlgorithm::tensor_ops() const { // @@protoc_insertion_point(field_get:xla.gpu.DenylistedAlgorithm.tensor_ops) return _internal_tensor_ops(); } inline void DenylistedAlgorithm::set_tensor_ops(bool value) { _internal_set_tensor_ops(value); // @@protoc_insertion_point(field_set:xla.gpu.DenylistedAlgorithm.tensor_ops) } inline bool DenylistedAlgorithm::_internal_tensor_ops() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.tensor_ops_; } inline void DenylistedAlgorithm::_internal_set_tensor_ops(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.tensor_ops_ = value; } // ------------------------------------------------------------------- // AlgorithmDenylistEntry // string hlo = 1; inline void AlgorithmDenylistEntry::clear_hlo() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.hlo_.ClearToEmpty(); } inline const std::string& AlgorithmDenylistEntry::hlo() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.hlo) return _internal_hlo(); } template inline PROTOBUF_ALWAYS_INLINE void AlgorithmDenylistEntry::set_hlo(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.hlo_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.AlgorithmDenylistEntry.hlo) } inline std::string* AlgorithmDenylistEntry::mutable_hlo() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_hlo(); // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.hlo) return _s; } inline const std::string& AlgorithmDenylistEntry::_internal_hlo() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.hlo_.Get(); } inline void AlgorithmDenylistEntry::_internal_set_hlo(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.hlo_.Set(value, GetArena()); } inline std::string* AlgorithmDenylistEntry::_internal_mutable_hlo() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.hlo_.Mutable( GetArena()); } inline std::string* AlgorithmDenylistEntry::release_hlo() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.AlgorithmDenylistEntry.hlo) return _impl_.hlo_.Release(); } inline void AlgorithmDenylistEntry::set_allocated_hlo(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.hlo_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.hlo_.IsDefault()) { _impl_.hlo_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.AlgorithmDenylistEntry.hlo) } // .xla.gpu.GpuBackendConfig backend_config = 6; inline bool AlgorithmDenylistEntry::has_backend_config() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.backend_config_ != nullptr); return value; } inline const ::xla::gpu::GpuBackendConfig& AlgorithmDenylistEntry::_internal_backend_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::GpuBackendConfig* p = _impl_.backend_config_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_GpuBackendConfig_default_instance_); } inline const ::xla::gpu::GpuBackendConfig& AlgorithmDenylistEntry::backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.backend_config) return _internal_backend_config(); } inline void AlgorithmDenylistEntry::unsafe_arena_set_allocated_backend_config(::xla::gpu::GpuBackendConfig* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.backend_config_); } _impl_.backend_config_ = reinterpret_cast<::xla::gpu::GpuBackendConfig*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.AlgorithmDenylistEntry.backend_config) } inline ::xla::gpu::GpuBackendConfig* AlgorithmDenylistEntry::release_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::GpuBackendConfig* released = _impl_.backend_config_; _impl_.backend_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::gpu::GpuBackendConfig* AlgorithmDenylistEntry::unsafe_arena_release_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.AlgorithmDenylistEntry.backend_config) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::GpuBackendConfig* temp = _impl_.backend_config_; _impl_.backend_config_ = nullptr; return temp; } inline ::xla::gpu::GpuBackendConfig* AlgorithmDenylistEntry::_internal_mutable_backend_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.backend_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::GpuBackendConfig>(GetArena()); _impl_.backend_config_ = reinterpret_cast<::xla::gpu::GpuBackendConfig*>(p); } return _impl_.backend_config_; } inline ::xla::gpu::GpuBackendConfig* AlgorithmDenylistEntry::mutable_backend_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::gpu::GpuBackendConfig* _msg = _internal_mutable_backend_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.backend_config) return _msg; } inline void AlgorithmDenylistEntry::set_allocated_backend_config(::xla::gpu::GpuBackendConfig* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.backend_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } _impl_.backend_config_ = reinterpret_cast<::xla::gpu::GpuBackendConfig*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.AlgorithmDenylistEntry.backend_config) } // .xla.ComputeCapability cc = 2; inline bool AlgorithmDenylistEntry::has_cc() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.cc_ != nullptr); return value; } inline const ::xla::ComputeCapability& AlgorithmDenylistEntry::_internal_cc() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ComputeCapability* p = _impl_.cc_; return p != nullptr ? *p : reinterpret_cast(::xla::_ComputeCapability_default_instance_); } inline const ::xla::ComputeCapability& AlgorithmDenylistEntry::cc() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.cc) return _internal_cc(); } inline void AlgorithmDenylistEntry::unsafe_arena_set_allocated_cc(::xla::ComputeCapability* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cc_); } _impl_.cc_ = reinterpret_cast<::xla::ComputeCapability*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.AlgorithmDenylistEntry.cc) } inline ::xla::ComputeCapability* AlgorithmDenylistEntry::release_cc() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::ComputeCapability* released = _impl_.cc_; _impl_.cc_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::ComputeCapability* AlgorithmDenylistEntry::unsafe_arena_release_cc() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.AlgorithmDenylistEntry.cc) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::ComputeCapability* temp = _impl_.cc_; _impl_.cc_ = nullptr; return temp; } inline ::xla::ComputeCapability* AlgorithmDenylistEntry::_internal_mutable_cc() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.cc_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ComputeCapability>(GetArena()); _impl_.cc_ = reinterpret_cast<::xla::ComputeCapability*>(p); } return _impl_.cc_; } inline ::xla::ComputeCapability* AlgorithmDenylistEntry::mutable_cc() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::ComputeCapability* _msg = _internal_mutable_cc(); // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.cc) return _msg; } inline void AlgorithmDenylistEntry::set_allocated_cc(::xla::ComputeCapability* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cc_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.cc_ = reinterpret_cast<::xla::ComputeCapability*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.AlgorithmDenylistEntry.cc) } // .xla.CudnnVersion cudnn_version = 3; inline bool AlgorithmDenylistEntry::has_cudnn_version() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.cudnn_version_ != nullptr); return value; } inline const ::xla::CudnnVersion& AlgorithmDenylistEntry::_internal_cudnn_version() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::CudnnVersion* p = _impl_.cudnn_version_; return p != nullptr ? *p : reinterpret_cast(::xla::_CudnnVersion_default_instance_); } inline const ::xla::CudnnVersion& AlgorithmDenylistEntry::cudnn_version() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.cudnn_version) return _internal_cudnn_version(); } inline void AlgorithmDenylistEntry::unsafe_arena_set_allocated_cudnn_version(::xla::CudnnVersion* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cudnn_version_); } _impl_.cudnn_version_ = reinterpret_cast<::xla::CudnnVersion*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.AlgorithmDenylistEntry.cudnn_version) } inline ::xla::CudnnVersion* AlgorithmDenylistEntry::release_cudnn_version() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::CudnnVersion* released = _impl_.cudnn_version_; _impl_.cudnn_version_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::CudnnVersion* AlgorithmDenylistEntry::unsafe_arena_release_cudnn_version() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.AlgorithmDenylistEntry.cudnn_version) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::CudnnVersion* temp = _impl_.cudnn_version_; _impl_.cudnn_version_ = nullptr; return temp; } inline ::xla::CudnnVersion* AlgorithmDenylistEntry::_internal_mutable_cudnn_version() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.cudnn_version_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::CudnnVersion>(GetArena()); _impl_.cudnn_version_ = reinterpret_cast<::xla::CudnnVersion*>(p); } return _impl_.cudnn_version_; } inline ::xla::CudnnVersion* AlgorithmDenylistEntry::mutable_cudnn_version() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::CudnnVersion* _msg = _internal_mutable_cudnn_version(); // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.cudnn_version) return _msg; } inline void AlgorithmDenylistEntry::set_allocated_cudnn_version(::xla::CudnnVersion* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cudnn_version_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.cudnn_version_ = reinterpret_cast<::xla::CudnnVersion*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.AlgorithmDenylistEntry.cudnn_version) } // string blas_version = 5; inline void AlgorithmDenylistEntry::clear_blas_version() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.blas_version_.ClearToEmpty(); } inline const std::string& AlgorithmDenylistEntry::blas_version() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.blas_version) return _internal_blas_version(); } template inline PROTOBUF_ALWAYS_INLINE void AlgorithmDenylistEntry::set_blas_version(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.blas_version_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.AlgorithmDenylistEntry.blas_version) } inline std::string* AlgorithmDenylistEntry::mutable_blas_version() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_blas_version(); // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.blas_version) return _s; } inline const std::string& AlgorithmDenylistEntry::_internal_blas_version() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.blas_version_.Get(); } inline void AlgorithmDenylistEntry::_internal_set_blas_version(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.blas_version_.Set(value, GetArena()); } inline std::string* AlgorithmDenylistEntry::_internal_mutable_blas_version() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.blas_version_.Mutable( GetArena()); } inline std::string* AlgorithmDenylistEntry::release_blas_version() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.AlgorithmDenylistEntry.blas_version) return _impl_.blas_version_.Release(); } inline void AlgorithmDenylistEntry::set_allocated_blas_version(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.blas_version_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.blas_version_.IsDefault()) { _impl_.blas_version_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.AlgorithmDenylistEntry.blas_version) } // repeated .xla.gpu.DenylistedAlgorithm algos = 4; inline int AlgorithmDenylistEntry::_internal_algos_size() const { return _internal_algos().size(); } inline int AlgorithmDenylistEntry::algos_size() const { return _internal_algos_size(); } inline void AlgorithmDenylistEntry::clear_algos() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.algos_.Clear(); } inline ::xla::gpu::DenylistedAlgorithm* AlgorithmDenylistEntry::mutable_algos(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylistEntry.algos) return _internal_mutable_algos()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>* AlgorithmDenylistEntry::mutable_algos() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.AlgorithmDenylistEntry.algos) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_algos(); } inline const ::xla::gpu::DenylistedAlgorithm& AlgorithmDenylistEntry::algos(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylistEntry.algos) return _internal_algos().Get(index); } inline ::xla::gpu::DenylistedAlgorithm* AlgorithmDenylistEntry::add_algos() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::DenylistedAlgorithm* _add = _internal_mutable_algos()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.AlgorithmDenylistEntry.algos) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>& AlgorithmDenylistEntry::algos() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.AlgorithmDenylistEntry.algos) return _internal_algos(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>& AlgorithmDenylistEntry::_internal_algos() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.algos_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::DenylistedAlgorithm>* AlgorithmDenylistEntry::_internal_mutable_algos() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.algos_; } // ------------------------------------------------------------------- // AlgorithmDenylist // repeated .xla.gpu.AlgorithmDenylistEntry entries = 1; inline int AlgorithmDenylist::_internal_entries_size() const { return _internal_entries().size(); } inline int AlgorithmDenylist::entries_size() const { return _internal_entries_size(); } inline void AlgorithmDenylist::clear_entries() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.entries_.Clear(); } inline ::xla::gpu::AlgorithmDenylistEntry* AlgorithmDenylist::mutable_entries(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.AlgorithmDenylist.entries) return _internal_mutable_entries()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>* AlgorithmDenylist::mutable_entries() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.AlgorithmDenylist.entries) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_entries(); } inline const ::xla::gpu::AlgorithmDenylistEntry& AlgorithmDenylist::entries(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.AlgorithmDenylist.entries) return _internal_entries().Get(index); } inline ::xla::gpu::AlgorithmDenylistEntry* AlgorithmDenylist::add_entries() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::AlgorithmDenylistEntry* _add = _internal_mutable_entries()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.AlgorithmDenylist.entries) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>& AlgorithmDenylist::entries() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.AlgorithmDenylist.entries) return _internal_entries(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>& AlgorithmDenylist::_internal_entries() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.entries_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::AlgorithmDenylistEntry>* AlgorithmDenylist::_internal_mutable_entries() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.entries_; } #ifdef __GNUC__ #pragma GCC diagnostic pop #endif // __GNUC__ // @@protoc_insertion_point(namespace_scope) } // namespace gpu } // namespace xla // @@protoc_insertion_point(global_scope) #include "google/protobuf/port_undef.inc" #endif // GOOGLE_PROTOBUF_INCLUDED_xla_2fservice_2fgpu_2fautotuning_2fgpu_5fautotuning_2eproto_2epb_2eh