// Generated by the protocol buffer compiler. DO NOT EDIT! // NO CHECKED-IN PROTOBUF GENCODE // source: xla/backends/gpu/runtime/thunk.proto // Protobuf C++ Version: 5.28.3 #ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto_2epb_2eh #define GOOGLE_PROTOBUF_INCLUDED_xla_2fbackends_2fgpu_2fruntime_2fthunk_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/service/buffer_assignment.pb.h" #include "xla/service/gpu/launch_dimensions.pb.h" #include "xla/stream_executor/gpu/gpu_blas_lt.pb.h" #include "xla/stream_executor/launch_dim.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_2fbackends_2fgpu_2fruntime_2fthunk_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_2fbackends_2fgpu_2fruntime_2fthunk_2eproto { static const ::uint32_t offsets[]; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable descriptor_table_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; namespace xla { namespace gpu { class ConditionalThunkProto; struct ConditionalThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern ConditionalThunkProtoDefaultTypeInternal _ConditionalThunkProto_default_instance_; class CopyThunkProto; struct CopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern CopyThunkProtoDefaultTypeInternal _CopyThunkProto_default_instance_; class DeviceToDeviceCopyThunkProto; struct DeviceToDeviceCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern DeviceToDeviceCopyThunkProtoDefaultTypeInternal _DeviceToDeviceCopyThunkProto_default_instance_; class DeviceToHostCopyThunkProto; struct DeviceToHostCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern DeviceToHostCopyThunkProtoDefaultTypeInternal _DeviceToHostCopyThunkProto_default_instance_; class GemmThunkProto; struct GemmThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern GemmThunkProtoDefaultTypeInternal _GemmThunkProto_default_instance_; class HostToDeviceCopyThunkProto; struct HostToDeviceCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern HostToDeviceCopyThunkProtoDefaultTypeInternal _HostToDeviceCopyThunkProto_default_instance_; class KernelThunkProto; struct KernelThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern KernelThunkProtoDefaultTypeInternal _KernelThunkProto_default_instance_; class PartitionIdThunkProto; struct PartitionIdThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern PartitionIdThunkProtoDefaultTypeInternal _PartitionIdThunkProto_default_instance_; class ReplicaIdThunkProto; struct ReplicaIdThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern ReplicaIdThunkProtoDefaultTypeInternal _ReplicaIdThunkProto_default_instance_; class SequentialThunkProto; struct SequentialThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern SequentialThunkProtoDefaultTypeInternal _SequentialThunkProto_default_instance_; class ThunkInfoProto; struct ThunkInfoProtoDefaultTypeInternal; PROTOBUF_EXPORT extern ThunkInfoProtoDefaultTypeInternal _ThunkInfoProto_default_instance_; class ThunkProto; struct ThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern ThunkProtoDefaultTypeInternal _ThunkProto_default_instance_; class TriangularSolveThunkProto; struct TriangularSolveThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern TriangularSolveThunkProtoDefaultTypeInternal _TriangularSolveThunkProto_default_instance_; class WaitForStreamsThunkProto; struct WaitForStreamsThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern WaitForStreamsThunkProtoDefaultTypeInternal _WaitForStreamsThunkProto_default_instance_; class WhileThunkProto; struct WhileThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern WhileThunkProtoDefaultTypeInternal _WhileThunkProto_default_instance_; } // namespace gpu } // namespace xla namespace google { namespace protobuf { } // namespace protobuf } // namespace google namespace xla { namespace gpu { // =================================================================== // ------------------------------------------------------------------- class PROTOBUF_EXPORT WaitForStreamsThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.WaitForStreamsThunkProto) */ { public: inline WaitForStreamsThunkProto() : WaitForStreamsThunkProto(nullptr) {} ~WaitForStreamsThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR WaitForStreamsThunkProto( ::google::protobuf::internal::ConstantInitialized); inline WaitForStreamsThunkProto(const WaitForStreamsThunkProto& from) : WaitForStreamsThunkProto(nullptr, from) {} inline WaitForStreamsThunkProto(WaitForStreamsThunkProto&& from) noexcept : WaitForStreamsThunkProto(nullptr, std::move(from)) {} inline WaitForStreamsThunkProto& operator=(const WaitForStreamsThunkProto& from) { CopyFrom(from); return *this; } inline WaitForStreamsThunkProto& operator=(WaitForStreamsThunkProto&& 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 WaitForStreamsThunkProto& default_instance() { return *internal_default_instance(); } static inline const WaitForStreamsThunkProto* internal_default_instance() { return reinterpret_cast( &_WaitForStreamsThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 9; friend void swap(WaitForStreamsThunkProto& a, WaitForStreamsThunkProto& b) { a.Swap(&b); } inline void Swap(WaitForStreamsThunkProto* 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(WaitForStreamsThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- WaitForStreamsThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const WaitForStreamsThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const WaitForStreamsThunkProto& from) { WaitForStreamsThunkProto::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(WaitForStreamsThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.WaitForStreamsThunkProto"; } protected: explicit WaitForStreamsThunkProto(::google::protobuf::Arena* arena); WaitForStreamsThunkProto(::google::protobuf::Arena* arena, const WaitForStreamsThunkProto& from); WaitForStreamsThunkProto(::google::protobuf::Arena* arena, WaitForStreamsThunkProto&& from) noexcept : WaitForStreamsThunkProto(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 { kStreamIdFieldNumber = 1, kWaitForStreamIdFieldNumber = 2, }; // int64 stream_id = 1; void clear_stream_id() ; ::int64_t stream_id() const; void set_stream_id(::int64_t value); private: ::int64_t _internal_stream_id() const; void _internal_set_stream_id(::int64_t value); public: // int64 wait_for_stream_id = 2; void clear_wait_for_stream_id() ; ::int64_t wait_for_stream_id() const; void set_wait_for_stream_id(::int64_t value); private: ::int64_t _internal_wait_for_stream_id() const; void _internal_set_wait_for_stream_id(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.WaitForStreamsThunkProto) 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 WaitForStreamsThunkProto& from_msg); ::int64_t stream_id_; ::int64_t wait_for_stream_id_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ThunkInfoProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ThunkInfoProto) */ { public: inline ThunkInfoProto() : ThunkInfoProto(nullptr) {} ~ThunkInfoProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ThunkInfoProto( ::google::protobuf::internal::ConstantInitialized); inline ThunkInfoProto(const ThunkInfoProto& from) : ThunkInfoProto(nullptr, from) {} inline ThunkInfoProto(ThunkInfoProto&& from) noexcept : ThunkInfoProto(nullptr, std::move(from)) {} inline ThunkInfoProto& operator=(const ThunkInfoProto& from) { CopyFrom(from); return *this; } inline ThunkInfoProto& operator=(ThunkInfoProto&& 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 ThunkInfoProto& default_instance() { return *internal_default_instance(); } static inline const ThunkInfoProto* internal_default_instance() { return reinterpret_cast( &_ThunkInfoProto_default_instance_); } static constexpr int kIndexInFileMessages = 0; friend void swap(ThunkInfoProto& a, ThunkInfoProto& b) { a.Swap(&b); } inline void Swap(ThunkInfoProto* 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(ThunkInfoProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ThunkInfoProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ThunkInfoProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ThunkInfoProto& from) { ThunkInfoProto::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(ThunkInfoProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ThunkInfoProto"; } protected: explicit ThunkInfoProto(::google::protobuf::Arena* arena); ThunkInfoProto(::google::protobuf::Arena* arena, const ThunkInfoProto& from); ThunkInfoProto(::google::protobuf::Arena* arena, ThunkInfoProto&& from) noexcept : ThunkInfoProto(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 { kProfileAnnotationFieldNumber = 1, kExecutionStreamIdFieldNumber = 2, }; // string profile_annotation = 1; void clear_profile_annotation() ; const std::string& profile_annotation() const; template void set_profile_annotation(Arg_&& arg, Args_... args); std::string* mutable_profile_annotation(); PROTOBUF_NODISCARD std::string* release_profile_annotation(); void set_allocated_profile_annotation(std::string* value); private: const std::string& _internal_profile_annotation() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_profile_annotation( const std::string& value); std::string* _internal_mutable_profile_annotation(); public: // int64 execution_stream_id = 2; void clear_execution_stream_id() ; ::int64_t execution_stream_id() const; void set_execution_stream_id(::int64_t value); private: ::int64_t _internal_execution_stream_id() const; void _internal_set_execution_stream_id(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.ThunkInfoProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 1, 2, 0, 49, 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 ThunkInfoProto& from_msg); ::google::protobuf::internal::ArenaStringPtr profile_annotation_; ::int64_t execution_stream_id_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT TriangularSolveThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.TriangularSolveThunkProto) */ { public: inline TriangularSolveThunkProto() : TriangularSolveThunkProto(nullptr) {} ~TriangularSolveThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR TriangularSolveThunkProto( ::google::protobuf::internal::ConstantInitialized); inline TriangularSolveThunkProto(const TriangularSolveThunkProto& from) : TriangularSolveThunkProto(nullptr, from) {} inline TriangularSolveThunkProto(TriangularSolveThunkProto&& from) noexcept : TriangularSolveThunkProto(nullptr, std::move(from)) {} inline TriangularSolveThunkProto& operator=(const TriangularSolveThunkProto& from) { CopyFrom(from); return *this; } inline TriangularSolveThunkProto& operator=(TriangularSolveThunkProto&& 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 TriangularSolveThunkProto& default_instance() { return *internal_default_instance(); } static inline const TriangularSolveThunkProto* internal_default_instance() { return reinterpret_cast( &_TriangularSolveThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 10; friend void swap(TriangularSolveThunkProto& a, TriangularSolveThunkProto& b) { a.Swap(&b); } inline void Swap(TriangularSolveThunkProto* 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(TriangularSolveThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- TriangularSolveThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const TriangularSolveThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const TriangularSolveThunkProto& from) { TriangularSolveThunkProto::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(TriangularSolveThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.TriangularSolveThunkProto"; } protected: explicit TriangularSolveThunkProto(::google::protobuf::Arena* arena); TriangularSolveThunkProto(::google::protobuf::Arena* arena, const TriangularSolveThunkProto& from); TriangularSolveThunkProto(::google::protobuf::Arena* arena, TriangularSolveThunkProto&& from) noexcept : TriangularSolveThunkProto(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 { kOptionsFieldNumber = 1, kABufferFieldNumber = 2, kBBufferFieldNumber = 3, kTempBufferFieldNumber = 4, kBatchSizeFieldNumber = 6, kMFieldNumber = 7, kNFieldNumber = 8, kABatchStrideFieldNumber = 9, kBBatchStrideFieldNumber = 10, kTypeFieldNumber = 5, }; // .xla.TriangularSolveOptions options = 1; bool has_options() const; void clear_options() ; const ::xla::TriangularSolveOptions& options() const; PROTOBUF_NODISCARD ::xla::TriangularSolveOptions* release_options(); ::xla::TriangularSolveOptions* mutable_options(); void set_allocated_options(::xla::TriangularSolveOptions* value); void unsafe_arena_set_allocated_options(::xla::TriangularSolveOptions* value); ::xla::TriangularSolveOptions* unsafe_arena_release_options(); private: const ::xla::TriangularSolveOptions& _internal_options() const; ::xla::TriangularSolveOptions* _internal_mutable_options(); public: // .xla.buffer_assignment.BufferAllocationSliceProto a_buffer = 2; bool has_a_buffer() const; void clear_a_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& a_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_a_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_a_buffer(); void set_allocated_a_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_a_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_a_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_a_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_a_buffer(); public: // .xla.buffer_assignment.BufferAllocationSliceProto b_buffer = 3; bool has_b_buffer() const; void clear_b_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& b_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_b_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_b_buffer(); void set_allocated_b_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_b_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_b_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_b_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_b_buffer(); public: // .xla.buffer_assignment.BufferAllocationSliceProto temp_buffer = 4; bool has_temp_buffer() const; void clear_temp_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& temp_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_temp_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_temp_buffer(); void set_allocated_temp_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_temp_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_temp_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_temp_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_temp_buffer(); public: // int64 batch_size = 6; void clear_batch_size() ; ::int64_t batch_size() const; void set_batch_size(::int64_t value); private: ::int64_t _internal_batch_size() const; void _internal_set_batch_size(::int64_t value); public: // int64 m = 7; void clear_m() ; ::int64_t m() const; void set_m(::int64_t value); private: ::int64_t _internal_m() const; void _internal_set_m(::int64_t value); public: // int64 n = 8; void clear_n() ; ::int64_t n() const; void set_n(::int64_t value); private: ::int64_t _internal_n() const; void _internal_set_n(::int64_t value); public: // int64 a_batch_stride = 9; void clear_a_batch_stride() ; ::int64_t a_batch_stride() const; void set_a_batch_stride(::int64_t value); private: ::int64_t _internal_a_batch_stride() const; void _internal_set_a_batch_stride(::int64_t value); public: // int64 b_batch_stride = 10; void clear_b_batch_stride() ; ::int64_t b_batch_stride() const; void set_b_batch_stride(::int64_t value); private: ::int64_t _internal_b_batch_stride() const; void _internal_set_b_batch_stride(::int64_t value); public: // .xla.PrimitiveType type = 5; void clear_type() ; ::xla::PrimitiveType type() const; void set_type(::xla::PrimitiveType value); private: ::xla::PrimitiveType _internal_type() const; void _internal_set_type(::xla::PrimitiveType value); public: // @@protoc_insertion_point(class_scope:xla.gpu.TriangularSolveThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 4, 10, 4, 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 TriangularSolveThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::TriangularSolveOptions* options_; ::xla::buffer_assignment::BufferAllocationSliceProto* a_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* b_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* temp_buffer_; ::int64_t batch_size_; ::int64_t m_; ::int64_t n_; ::int64_t a_batch_stride_; ::int64_t b_batch_stride_; int type_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ReplicaIdThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ReplicaIdThunkProto) */ { public: inline ReplicaIdThunkProto() : ReplicaIdThunkProto(nullptr) {} ~ReplicaIdThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ReplicaIdThunkProto( ::google::protobuf::internal::ConstantInitialized); inline ReplicaIdThunkProto(const ReplicaIdThunkProto& from) : ReplicaIdThunkProto(nullptr, from) {} inline ReplicaIdThunkProto(ReplicaIdThunkProto&& from) noexcept : ReplicaIdThunkProto(nullptr, std::move(from)) {} inline ReplicaIdThunkProto& operator=(const ReplicaIdThunkProto& from) { CopyFrom(from); return *this; } inline ReplicaIdThunkProto& operator=(ReplicaIdThunkProto&& from) noexcept { if (this == &from) return *this; if (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 ReplicaIdThunkProto& default_instance() { return *internal_default_instance(); } static inline const ReplicaIdThunkProto* internal_default_instance() { return reinterpret_cast( &_ReplicaIdThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 11; friend void swap(ReplicaIdThunkProto& a, ReplicaIdThunkProto& b) { a.Swap(&b); } inline void Swap(ReplicaIdThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(ReplicaIdThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ReplicaIdThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ReplicaIdThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ReplicaIdThunkProto& from) { ReplicaIdThunkProto::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(ReplicaIdThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ReplicaIdThunkProto"; } protected: explicit ReplicaIdThunkProto(::google::protobuf::Arena* arena); ReplicaIdThunkProto(::google::protobuf::Arena* arena, const ReplicaIdThunkProto& from); ReplicaIdThunkProto(::google::protobuf::Arena* arena, ReplicaIdThunkProto&& from) noexcept : ReplicaIdThunkProto(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 { kDestBufferFieldNumber = 1, }; // .xla.buffer_assignment.BufferAllocationSliceProto dest_buffer = 1; bool has_dest_buffer() const; void clear_dest_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& dest_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_dest_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_dest_buffer(); void set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_dest_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_dest_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_dest_buffer(); public: // @@protoc_insertion_point(class_scope:xla.gpu.ReplicaIdThunkProto) 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 ReplicaIdThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::buffer_assignment::BufferAllocationSliceProto* dest_buffer_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT PartitionIdThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.PartitionIdThunkProto) */ { public: inline PartitionIdThunkProto() : PartitionIdThunkProto(nullptr) {} ~PartitionIdThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR PartitionIdThunkProto( ::google::protobuf::internal::ConstantInitialized); inline PartitionIdThunkProto(const PartitionIdThunkProto& from) : PartitionIdThunkProto(nullptr, from) {} inline PartitionIdThunkProto(PartitionIdThunkProto&& from) noexcept : PartitionIdThunkProto(nullptr, std::move(from)) {} inline PartitionIdThunkProto& operator=(const PartitionIdThunkProto& from) { CopyFrom(from); return *this; } inline PartitionIdThunkProto& operator=(PartitionIdThunkProto&& from) noexcept { if (this == &from) return *this; if (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 PartitionIdThunkProto& default_instance() { return *internal_default_instance(); } static inline const PartitionIdThunkProto* internal_default_instance() { return reinterpret_cast( &_PartitionIdThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 12; friend void swap(PartitionIdThunkProto& a, PartitionIdThunkProto& b) { a.Swap(&b); } inline void Swap(PartitionIdThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(PartitionIdThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- PartitionIdThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const PartitionIdThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const PartitionIdThunkProto& from) { PartitionIdThunkProto::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(PartitionIdThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.PartitionIdThunkProto"; } protected: explicit PartitionIdThunkProto(::google::protobuf::Arena* arena); PartitionIdThunkProto(::google::protobuf::Arena* arena, const PartitionIdThunkProto& from); PartitionIdThunkProto(::google::protobuf::Arena* arena, PartitionIdThunkProto&& from) noexcept : PartitionIdThunkProto(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 { kDestBufferFieldNumber = 1, }; // .xla.buffer_assignment.BufferAllocationSliceProto dest_buffer = 1; bool has_dest_buffer() const; void clear_dest_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& dest_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_dest_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_dest_buffer(); void set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_dest_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_dest_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_dest_buffer(); public: // @@protoc_insertion_point(class_scope:xla.gpu.PartitionIdThunkProto) 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 PartitionIdThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::buffer_assignment::BufferAllocationSliceProto* dest_buffer_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT CopyThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.CopyThunkProto) */ { public: inline CopyThunkProto() : CopyThunkProto(nullptr) {} ~CopyThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR CopyThunkProto( ::google::protobuf::internal::ConstantInitialized); inline CopyThunkProto(const CopyThunkProto& from) : CopyThunkProto(nullptr, from) {} inline CopyThunkProto(CopyThunkProto&& from) noexcept : CopyThunkProto(nullptr, std::move(from)) {} inline CopyThunkProto& operator=(const CopyThunkProto& from) { CopyFrom(from); return *this; } inline CopyThunkProto& operator=(CopyThunkProto&& from) noexcept { if (this == &from) return *this; if (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 CopyThunkProto& default_instance() { return *internal_default_instance(); } static inline const CopyThunkProto* internal_default_instance() { return reinterpret_cast( &_CopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 1; friend void swap(CopyThunkProto& a, CopyThunkProto& b) { a.Swap(&b); } inline void Swap(CopyThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(CopyThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CopyThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const CopyThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const CopyThunkProto& from) { CopyThunkProto::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(CopyThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.CopyThunkProto"; } protected: explicit CopyThunkProto(::google::protobuf::Arena* arena); CopyThunkProto(::google::protobuf::Arena* arena, const CopyThunkProto& from); CopyThunkProto(::google::protobuf::Arena* arena, CopyThunkProto&& from) noexcept : CopyThunkProto(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 { kSourceBufferFieldNumber = 1, kDestinationBufferFieldNumber = 2, kMemSizeFieldNumber = 3, }; // .xla.buffer_assignment.BufferAllocationSliceProto source_buffer = 1; bool has_source_buffer() const; void clear_source_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& source_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_source_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_source_buffer(); void set_allocated_source_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_source_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_source_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_source_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_source_buffer(); public: // .xla.buffer_assignment.BufferAllocationSliceProto destination_buffer = 2; bool has_destination_buffer() const; void clear_destination_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& destination_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_destination_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_destination_buffer(); void set_allocated_destination_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_destination_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_destination_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_destination_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_destination_buffer(); public: // int64 mem_size = 3; void clear_mem_size() ; ::int64_t mem_size() const; void set_mem_size(::int64_t value); private: ::int64_t _internal_mem_size() const; void _internal_set_mem_size(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.CopyThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 3, 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 CopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::buffer_assignment::BufferAllocationSliceProto* source_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* destination_buffer_; ::int64_t mem_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT HostToDeviceCopyThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.HostToDeviceCopyThunkProto) */ { public: inline HostToDeviceCopyThunkProto() : HostToDeviceCopyThunkProto(nullptr) {} ~HostToDeviceCopyThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR HostToDeviceCopyThunkProto( ::google::protobuf::internal::ConstantInitialized); inline HostToDeviceCopyThunkProto(const HostToDeviceCopyThunkProto& from) : HostToDeviceCopyThunkProto(nullptr, from) {} inline HostToDeviceCopyThunkProto(HostToDeviceCopyThunkProto&& from) noexcept : HostToDeviceCopyThunkProto(nullptr, std::move(from)) {} inline HostToDeviceCopyThunkProto& operator=(const HostToDeviceCopyThunkProto& from) { CopyFrom(from); return *this; } inline HostToDeviceCopyThunkProto& operator=(HostToDeviceCopyThunkProto&& 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 HostToDeviceCopyThunkProto& default_instance() { return *internal_default_instance(); } static inline const HostToDeviceCopyThunkProto* internal_default_instance() { return reinterpret_cast( &_HostToDeviceCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 3; friend void swap(HostToDeviceCopyThunkProto& a, HostToDeviceCopyThunkProto& b) { a.Swap(&b); } inline void Swap(HostToDeviceCopyThunkProto* 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(HostToDeviceCopyThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- HostToDeviceCopyThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const HostToDeviceCopyThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const HostToDeviceCopyThunkProto& from) { HostToDeviceCopyThunkProto::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(HostToDeviceCopyThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.HostToDeviceCopyThunkProto"; } protected: explicit HostToDeviceCopyThunkProto(::google::protobuf::Arena* arena); HostToDeviceCopyThunkProto(::google::protobuf::Arena* arena, const HostToDeviceCopyThunkProto& from); HostToDeviceCopyThunkProto(::google::protobuf::Arena* arena, HostToDeviceCopyThunkProto&& from) noexcept : HostToDeviceCopyThunkProto(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 { kCopyThunkFieldNumber = 1, }; // .xla.gpu.CopyThunkProto copy_thunk = 1; bool has_copy_thunk() const; void clear_copy_thunk() ; const ::xla::gpu::CopyThunkProto& copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::CopyThunkProto* release_copy_thunk(); ::xla::gpu::CopyThunkProto* mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); ::xla::gpu::CopyThunkProto* unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* _internal_mutable_copy_thunk(); public: // @@protoc_insertion_point(class_scope:xla.gpu.HostToDeviceCopyThunkProto) 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 HostToDeviceCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT GemmThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.GemmThunkProto) */ { public: inline GemmThunkProto() : GemmThunkProto(nullptr) {} ~GemmThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR GemmThunkProto( ::google::protobuf::internal::ConstantInitialized); inline GemmThunkProto(const GemmThunkProto& from) : GemmThunkProto(nullptr, from) {} inline GemmThunkProto(GemmThunkProto&& from) noexcept : GemmThunkProto(nullptr, std::move(from)) {} inline GemmThunkProto& operator=(const GemmThunkProto& from) { CopyFrom(from); return *this; } inline GemmThunkProto& operator=(GemmThunkProto&& 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 GemmThunkProto& default_instance() { return *internal_default_instance(); } static inline const GemmThunkProto* internal_default_instance() { return reinterpret_cast( &_GemmThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 8; friend void swap(GemmThunkProto& a, GemmThunkProto& b) { a.Swap(&b); } inline void Swap(GemmThunkProto* 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(GemmThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- GemmThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const GemmThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const GemmThunkProto& from) { GemmThunkProto::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(GemmThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.GemmThunkProto"; } protected: explicit GemmThunkProto(::google::protobuf::Arena* arena); GemmThunkProto(::google::protobuf::Arena* arena, const GemmThunkProto& from); GemmThunkProto(::google::protobuf::Arena* arena, GemmThunkProto&& from) noexcept : GemmThunkProto(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 { kGemmConfigFieldNumber = 1, kLhsBufferFieldNumber = 2, kRhsBufferFieldNumber = 3, kOutputBufferFieldNumber = 4, kWorkspaceFieldNumber = 5, kDeterministicFieldNumber = 6, }; // .xla.GemmConfigProto gemm_config = 1; bool has_gemm_config() const; void clear_gemm_config() ; const ::xla::GemmConfigProto& gemm_config() const; PROTOBUF_NODISCARD ::xla::GemmConfigProto* release_gemm_config(); ::xla::GemmConfigProto* mutable_gemm_config(); void set_allocated_gemm_config(::xla::GemmConfigProto* value); void unsafe_arena_set_allocated_gemm_config(::xla::GemmConfigProto* value); ::xla::GemmConfigProto* unsafe_arena_release_gemm_config(); private: const ::xla::GemmConfigProto& _internal_gemm_config() const; ::xla::GemmConfigProto* _internal_mutable_gemm_config(); public: // .xla.buffer_assignment.BufferAllocationSliceProto lhs_buffer = 2; bool has_lhs_buffer() const; void clear_lhs_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& lhs_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_lhs_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_lhs_buffer(); void set_allocated_lhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_lhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_lhs_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_lhs_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_lhs_buffer(); public: // .xla.buffer_assignment.BufferAllocationSliceProto rhs_buffer = 3; bool has_rhs_buffer() const; void clear_rhs_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& rhs_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_rhs_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_rhs_buffer(); void set_allocated_rhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_rhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_rhs_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_rhs_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_rhs_buffer(); public: // .xla.buffer_assignment.BufferAllocationSliceProto output_buffer = 4; bool has_output_buffer() const; void clear_output_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& output_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_output_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_output_buffer(); void set_allocated_output_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_output_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_output_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_output_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_output_buffer(); public: // optional .xla.buffer_assignment.BufferAllocationSliceProto workspace = 5; bool has_workspace() const; void clear_workspace() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& workspace() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_workspace(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_workspace(); void set_allocated_workspace(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_workspace(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_workspace(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_workspace() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_workspace(); public: // bool deterministic = 6; void clear_deterministic() ; bool deterministic() const; void set_deterministic(bool value); private: bool _internal_deterministic() const; void _internal_set_deterministic(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.GemmThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 6, 5, 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 GemmThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::GemmConfigProto* gemm_config_; ::xla::buffer_assignment::BufferAllocationSliceProto* lhs_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* rhs_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* output_buffer_; ::xla::buffer_assignment::BufferAllocationSliceProto* workspace_; bool deterministic_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT DeviceToHostCopyThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DeviceToHostCopyThunkProto) */ { public: inline DeviceToHostCopyThunkProto() : DeviceToHostCopyThunkProto(nullptr) {} ~DeviceToHostCopyThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR DeviceToHostCopyThunkProto( ::google::protobuf::internal::ConstantInitialized); inline DeviceToHostCopyThunkProto(const DeviceToHostCopyThunkProto& from) : DeviceToHostCopyThunkProto(nullptr, from) {} inline DeviceToHostCopyThunkProto(DeviceToHostCopyThunkProto&& from) noexcept : DeviceToHostCopyThunkProto(nullptr, std::move(from)) {} inline DeviceToHostCopyThunkProto& operator=(const DeviceToHostCopyThunkProto& from) { CopyFrom(from); return *this; } inline DeviceToHostCopyThunkProto& operator=(DeviceToHostCopyThunkProto&& 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 DeviceToHostCopyThunkProto& default_instance() { return *internal_default_instance(); } static inline const DeviceToHostCopyThunkProto* internal_default_instance() { return reinterpret_cast( &_DeviceToHostCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 2; friend void swap(DeviceToHostCopyThunkProto& a, DeviceToHostCopyThunkProto& b) { a.Swap(&b); } inline void Swap(DeviceToHostCopyThunkProto* 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(DeviceToHostCopyThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DeviceToHostCopyThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DeviceToHostCopyThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DeviceToHostCopyThunkProto& from) { DeviceToHostCopyThunkProto::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(DeviceToHostCopyThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.DeviceToHostCopyThunkProto"; } protected: explicit DeviceToHostCopyThunkProto(::google::protobuf::Arena* arena); DeviceToHostCopyThunkProto(::google::protobuf::Arena* arena, const DeviceToHostCopyThunkProto& from); DeviceToHostCopyThunkProto(::google::protobuf::Arena* arena, DeviceToHostCopyThunkProto&& from) noexcept : DeviceToHostCopyThunkProto(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 { kCopyThunkFieldNumber = 1, }; // .xla.gpu.CopyThunkProto copy_thunk = 1; bool has_copy_thunk() const; void clear_copy_thunk() ; const ::xla::gpu::CopyThunkProto& copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::CopyThunkProto* release_copy_thunk(); ::xla::gpu::CopyThunkProto* mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); ::xla::gpu::CopyThunkProto* unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* _internal_mutable_copy_thunk(); public: // @@protoc_insertion_point(class_scope:xla.gpu.DeviceToHostCopyThunkProto) 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 DeviceToHostCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT DeviceToDeviceCopyThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DeviceToDeviceCopyThunkProto) */ { public: inline DeviceToDeviceCopyThunkProto() : DeviceToDeviceCopyThunkProto(nullptr) {} ~DeviceToDeviceCopyThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR DeviceToDeviceCopyThunkProto( ::google::protobuf::internal::ConstantInitialized); inline DeviceToDeviceCopyThunkProto(const DeviceToDeviceCopyThunkProto& from) : DeviceToDeviceCopyThunkProto(nullptr, from) {} inline DeviceToDeviceCopyThunkProto(DeviceToDeviceCopyThunkProto&& from) noexcept : DeviceToDeviceCopyThunkProto(nullptr, std::move(from)) {} inline DeviceToDeviceCopyThunkProto& operator=(const DeviceToDeviceCopyThunkProto& from) { CopyFrom(from); return *this; } inline DeviceToDeviceCopyThunkProto& operator=(DeviceToDeviceCopyThunkProto&& 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 DeviceToDeviceCopyThunkProto& default_instance() { return *internal_default_instance(); } static inline const DeviceToDeviceCopyThunkProto* internal_default_instance() { return reinterpret_cast( &_DeviceToDeviceCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 4; friend void swap(DeviceToDeviceCopyThunkProto& a, DeviceToDeviceCopyThunkProto& b) { a.Swap(&b); } inline void Swap(DeviceToDeviceCopyThunkProto* 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(DeviceToDeviceCopyThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DeviceToDeviceCopyThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DeviceToDeviceCopyThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DeviceToDeviceCopyThunkProto& from) { DeviceToDeviceCopyThunkProto::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(DeviceToDeviceCopyThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.DeviceToDeviceCopyThunkProto"; } protected: explicit DeviceToDeviceCopyThunkProto(::google::protobuf::Arena* arena); DeviceToDeviceCopyThunkProto(::google::protobuf::Arena* arena, const DeviceToDeviceCopyThunkProto& from); DeviceToDeviceCopyThunkProto(::google::protobuf::Arena* arena, DeviceToDeviceCopyThunkProto&& from) noexcept : DeviceToDeviceCopyThunkProto(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 { kCopyThunkFieldNumber = 1, }; // .xla.gpu.CopyThunkProto copy_thunk = 1; bool has_copy_thunk() const; void clear_copy_thunk() ; const ::xla::gpu::CopyThunkProto& copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::CopyThunkProto* release_copy_thunk(); ::xla::gpu::CopyThunkProto* mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); ::xla::gpu::CopyThunkProto* unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* _internal_mutable_copy_thunk(); public: // @@protoc_insertion_point(class_scope:xla.gpu.DeviceToDeviceCopyThunkProto) 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 DeviceToDeviceCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT KernelThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.KernelThunkProto) */ { public: inline KernelThunkProto() : KernelThunkProto(nullptr) {} ~KernelThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR KernelThunkProto( ::google::protobuf::internal::ConstantInitialized); inline KernelThunkProto(const KernelThunkProto& from) : KernelThunkProto(nullptr, from) {} inline KernelThunkProto(KernelThunkProto&& from) noexcept : KernelThunkProto(nullptr, std::move(from)) {} inline KernelThunkProto& operator=(const KernelThunkProto& from) { CopyFrom(from); return *this; } inline KernelThunkProto& operator=(KernelThunkProto&& from) noexcept { if (this == &from) return *this; if (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 KernelThunkProto& default_instance() { return *internal_default_instance(); } static inline const KernelThunkProto* internal_default_instance() { return reinterpret_cast( &_KernelThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 7; friend void swap(KernelThunkProto& a, KernelThunkProto& b) { a.Swap(&b); } inline void Swap(KernelThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(KernelThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- KernelThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const KernelThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const KernelThunkProto& from) { KernelThunkProto::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(KernelThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.KernelThunkProto"; } protected: explicit KernelThunkProto(::google::protobuf::Arena* arena); KernelThunkProto(::google::protobuf::Arena* arena, const KernelThunkProto& from); KernelThunkProto(::google::protobuf::Arena* arena, KernelThunkProto&& from) noexcept : KernelThunkProto(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 { kArgsFieldNumber = 1, kWrittenFieldNumber = 2, kKernelNameFieldNumber = 3, kLaunchDimensionsFieldNumber = 4, kClusterDimFieldNumber = 5, kShmemBytesFieldNumber = 6, }; // repeated .xla.buffer_assignment.BufferAllocationSliceProto args = 1; int args_size() const; private: int _internal_args_size() const; public: void clear_args() ; ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_args(int index); ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>* mutable_args(); private: const ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>& _internal_args() const; ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>* _internal_mutable_args(); public: const ::xla::buffer_assignment::BufferAllocationSliceProto& args(int index) const; ::xla::buffer_assignment::BufferAllocationSliceProto* add_args(); const ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>& args() const; // repeated bool written = 2; int written_size() const; private: int _internal_written_size() const; public: void clear_written() ; bool written(int index) const; void set_written(int index, bool value); void add_written(bool value); const ::google::protobuf::RepeatedField& written() const; ::google::protobuf::RepeatedField* mutable_written(); private: const ::google::protobuf::RepeatedField& _internal_written() const; ::google::protobuf::RepeatedField* _internal_mutable_written(); public: // string kernel_name = 3; void clear_kernel_name() ; const std::string& kernel_name() const; template void set_kernel_name(Arg_&& arg, Args_... args); std::string* mutable_kernel_name(); PROTOBUF_NODISCARD std::string* release_kernel_name(); void set_allocated_kernel_name(std::string* value); private: const std::string& _internal_kernel_name() const; inline PROTOBUF_ALWAYS_INLINE void _internal_set_kernel_name( const std::string& value); std::string* _internal_mutable_kernel_name(); public: // .xla.gpu.LaunchDimensionsProto launch_dimensions = 4; bool has_launch_dimensions() const; void clear_launch_dimensions() ; const ::xla::gpu::LaunchDimensionsProto& launch_dimensions() const; PROTOBUF_NODISCARD ::xla::gpu::LaunchDimensionsProto* release_launch_dimensions(); ::xla::gpu::LaunchDimensionsProto* mutable_launch_dimensions(); void set_allocated_launch_dimensions(::xla::gpu::LaunchDimensionsProto* value); void unsafe_arena_set_allocated_launch_dimensions(::xla::gpu::LaunchDimensionsProto* value); ::xla::gpu::LaunchDimensionsProto* unsafe_arena_release_launch_dimensions(); private: const ::xla::gpu::LaunchDimensionsProto& _internal_launch_dimensions() const; ::xla::gpu::LaunchDimensionsProto* _internal_mutable_launch_dimensions(); public: // optional .stream_executor.ClusterDimProto cluster_dim = 5; bool has_cluster_dim() const; void clear_cluster_dim() ; const ::stream_executor::ClusterDimProto& cluster_dim() const; PROTOBUF_NODISCARD ::stream_executor::ClusterDimProto* release_cluster_dim(); ::stream_executor::ClusterDimProto* mutable_cluster_dim(); void set_allocated_cluster_dim(::stream_executor::ClusterDimProto* value); void unsafe_arena_set_allocated_cluster_dim(::stream_executor::ClusterDimProto* value); ::stream_executor::ClusterDimProto* unsafe_arena_release_cluster_dim(); private: const ::stream_executor::ClusterDimProto& _internal_cluster_dim() const; ::stream_executor::ClusterDimProto* _internal_mutable_cluster_dim(); public: // int64 shmem_bytes = 6; void clear_shmem_bytes() ; ::int64_t shmem_bytes() const; void set_shmem_bytes(::int64_t value); private: ::int64_t _internal_shmem_bytes() const; void _internal_set_shmem_bytes(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.KernelThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 3, 6, 3, 44, 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 KernelThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::RepeatedPtrField< ::xla::buffer_assignment::BufferAllocationSliceProto > args_; ::google::protobuf::RepeatedField written_; ::google::protobuf::internal::ArenaStringPtr kernel_name_; ::xla::gpu::LaunchDimensionsProto* launch_dimensions_; ::stream_executor::ClusterDimProto* cluster_dim_; ::int64_t shmem_bytes_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ConditionalThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ConditionalThunkProto) */ { public: inline ConditionalThunkProto() : ConditionalThunkProto(nullptr) {} ~ConditionalThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ConditionalThunkProto( ::google::protobuf::internal::ConstantInitialized); inline ConditionalThunkProto(const ConditionalThunkProto& from) : ConditionalThunkProto(nullptr, from) {} inline ConditionalThunkProto(ConditionalThunkProto&& from) noexcept : ConditionalThunkProto(nullptr, std::move(from)) {} inline ConditionalThunkProto& operator=(const ConditionalThunkProto& from) { CopyFrom(from); return *this; } inline ConditionalThunkProto& operator=(ConditionalThunkProto&& from) noexcept { if (this == &from) return *this; if (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 ConditionalThunkProto& default_instance() { return *internal_default_instance(); } static inline const ConditionalThunkProto* internal_default_instance() { return reinterpret_cast( &_ConditionalThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 5; friend void swap(ConditionalThunkProto& a, ConditionalThunkProto& b) { a.Swap(&b); } inline void Swap(ConditionalThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(ConditionalThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ConditionalThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ConditionalThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ConditionalThunkProto& from) { ConditionalThunkProto::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(ConditionalThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ConditionalThunkProto"; } protected: explicit ConditionalThunkProto(::google::protobuf::Arena* arena); ConditionalThunkProto(::google::protobuf::Arena* arena, const ConditionalThunkProto& from); ConditionalThunkProto(::google::protobuf::Arena* arena, ConditionalThunkProto&& from) noexcept : ConditionalThunkProto(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 { kBranchThunksFieldNumber = 2, kBranchIndexBufferFieldNumber = 1, kBranchIndexIsBoolFieldNumber = 3, }; // repeated .xla.gpu.SequentialThunkProto branch_thunks = 2; int branch_thunks_size() const; private: int _internal_branch_thunks_size() const; public: void clear_branch_thunks() ; ::xla::gpu::SequentialThunkProto* mutable_branch_thunks(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>* mutable_branch_thunks(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>& _internal_branch_thunks() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>* _internal_mutable_branch_thunks(); public: const ::xla::gpu::SequentialThunkProto& branch_thunks(int index) const; ::xla::gpu::SequentialThunkProto* add_branch_thunks(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>& branch_thunks() const; // .xla.buffer_assignment.BufferAllocationSliceProto branch_index_buffer = 1; bool has_branch_index_buffer() const; void clear_branch_index_buffer() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& branch_index_buffer() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_branch_index_buffer(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_branch_index_buffer(); void set_allocated_branch_index_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_branch_index_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_branch_index_buffer(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_branch_index_buffer() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_branch_index_buffer(); public: // bool branch_index_is_bool = 3; void clear_branch_index_is_bool() ; bool branch_index_is_bool() const; void set_branch_index_is_bool(bool value); private: bool _internal_branch_index_is_bool() const; void _internal_set_branch_index_is_bool(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.ConditionalThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 3, 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 ConditionalThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::RepeatedPtrField< ::xla::gpu::SequentialThunkProto > branch_thunks_; ::xla::buffer_assignment::BufferAllocationSliceProto* branch_index_buffer_; bool branch_index_is_bool_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT SequentialThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.SequentialThunkProto) */ { public: inline SequentialThunkProto() : SequentialThunkProto(nullptr) {} ~SequentialThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR SequentialThunkProto( ::google::protobuf::internal::ConstantInitialized); inline SequentialThunkProto(const SequentialThunkProto& from) : SequentialThunkProto(nullptr, from) {} inline SequentialThunkProto(SequentialThunkProto&& from) noexcept : SequentialThunkProto(nullptr, std::move(from)) {} inline SequentialThunkProto& operator=(const SequentialThunkProto& from) { CopyFrom(from); return *this; } inline SequentialThunkProto& operator=(SequentialThunkProto&& 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 SequentialThunkProto& default_instance() { return *internal_default_instance(); } static inline const SequentialThunkProto* internal_default_instance() { return reinterpret_cast( &_SequentialThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 14; friend void swap(SequentialThunkProto& a, SequentialThunkProto& b) { a.Swap(&b); } inline void Swap(SequentialThunkProto* 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(SequentialThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- SequentialThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const SequentialThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const SequentialThunkProto& from) { SequentialThunkProto::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(SequentialThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.SequentialThunkProto"; } protected: explicit SequentialThunkProto(::google::protobuf::Arena* arena); SequentialThunkProto(::google::protobuf::Arena* arena, const SequentialThunkProto& from); SequentialThunkProto(::google::protobuf::Arena* arena, SequentialThunkProto&& from) noexcept : SequentialThunkProto(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 { kThunksFieldNumber = 1, }; // repeated .xla.gpu.ThunkProto thunks = 1; int thunks_size() const; private: int _internal_thunks_size() const; public: void clear_thunks() ; ::xla::gpu::ThunkProto* mutable_thunks(int index); ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>* mutable_thunks(); private: const ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>& _internal_thunks() const; ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>* _internal_mutable_thunks(); public: const ::xla::gpu::ThunkProto& thunks(int index) const; ::xla::gpu::ThunkProto* add_thunks(); const ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>& thunks() const; // @@protoc_insertion_point(class_scope:xla.gpu.SequentialThunkProto) 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 SequentialThunkProto& from_msg); ::google::protobuf::RepeatedPtrField< ::xla::gpu::ThunkProto > thunks_; mutable ::google::protobuf::internal::CachedSize _cached_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT ThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.ThunkProto) */ { public: inline ThunkProto() : ThunkProto(nullptr) {} ~ThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR ThunkProto( ::google::protobuf::internal::ConstantInitialized); inline ThunkProto(const ThunkProto& from) : ThunkProto(nullptr, from) {} inline ThunkProto(ThunkProto&& from) noexcept : ThunkProto(nullptr, std::move(from)) {} inline ThunkProto& operator=(const ThunkProto& from) { CopyFrom(from); return *this; } inline ThunkProto& operator=(ThunkProto&& from) noexcept { if (this == &from) return *this; if (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 ThunkProto& default_instance() { return *internal_default_instance(); } enum ImplCase { kSequentialThunk = 2, kCopyThunk = 3, kDeviceToHostCopyThunk = 4, kHostToDeviceCopyThunk = 5, kConditionalThunk = 6, kWhileThunk = 7, kKernelThunk = 8, kGemmThunk = 9, kWaitForStreamsThunk = 10, kTriangularSolveThunk = 11, kReplicaIdThunk = 12, kDeviceToDeviceCopyThunk = 13, kPartitionIdThunk = 14, IMPL_NOT_SET = 0, }; static inline const ThunkProto* internal_default_instance() { return reinterpret_cast( &_ThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 13; friend void swap(ThunkProto& a, ThunkProto& b) { a.Swap(&b); } inline void Swap(ThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(ThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- ThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const ThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const ThunkProto& from) { ThunkProto::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(ThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.ThunkProto"; } protected: explicit ThunkProto(::google::protobuf::Arena* arena); ThunkProto(::google::protobuf::Arena* arena, const ThunkProto& from); ThunkProto(::google::protobuf::Arena* arena, ThunkProto&& from) noexcept : ThunkProto(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 { kThunkInfoFieldNumber = 1, kSequentialThunkFieldNumber = 2, kCopyThunkFieldNumber = 3, kDeviceToHostCopyThunkFieldNumber = 4, kHostToDeviceCopyThunkFieldNumber = 5, kConditionalThunkFieldNumber = 6, kWhileThunkFieldNumber = 7, kKernelThunkFieldNumber = 8, kGemmThunkFieldNumber = 9, kWaitForStreamsThunkFieldNumber = 10, kTriangularSolveThunkFieldNumber = 11, kReplicaIdThunkFieldNumber = 12, kDeviceToDeviceCopyThunkFieldNumber = 13, kPartitionIdThunkFieldNumber = 14, }; // .xla.gpu.ThunkInfoProto thunk_info = 1; bool has_thunk_info() const; void clear_thunk_info() ; const ::xla::gpu::ThunkInfoProto& thunk_info() const; PROTOBUF_NODISCARD ::xla::gpu::ThunkInfoProto* release_thunk_info(); ::xla::gpu::ThunkInfoProto* mutable_thunk_info(); void set_allocated_thunk_info(::xla::gpu::ThunkInfoProto* value); void unsafe_arena_set_allocated_thunk_info(::xla::gpu::ThunkInfoProto* value); ::xla::gpu::ThunkInfoProto* unsafe_arena_release_thunk_info(); private: const ::xla::gpu::ThunkInfoProto& _internal_thunk_info() const; ::xla::gpu::ThunkInfoProto* _internal_mutable_thunk_info(); public: // .xla.gpu.SequentialThunkProto sequential_thunk = 2; bool has_sequential_thunk() const; private: bool _internal_has_sequential_thunk() const; public: void clear_sequential_thunk() ; const ::xla::gpu::SequentialThunkProto& sequential_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::SequentialThunkProto* release_sequential_thunk(); ::xla::gpu::SequentialThunkProto* mutable_sequential_thunk(); void set_allocated_sequential_thunk(::xla::gpu::SequentialThunkProto* value); void unsafe_arena_set_allocated_sequential_thunk(::xla::gpu::SequentialThunkProto* value); ::xla::gpu::SequentialThunkProto* unsafe_arena_release_sequential_thunk(); private: const ::xla::gpu::SequentialThunkProto& _internal_sequential_thunk() const; ::xla::gpu::SequentialThunkProto* _internal_mutable_sequential_thunk(); public: // .xla.gpu.CopyThunkProto copy_thunk = 3; bool has_copy_thunk() const; private: bool _internal_has_copy_thunk() const; public: void clear_copy_thunk() ; const ::xla::gpu::CopyThunkProto& copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::CopyThunkProto* release_copy_thunk(); ::xla::gpu::CopyThunkProto* mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value); ::xla::gpu::CopyThunkProto* unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* _internal_mutable_copy_thunk(); public: // .xla.gpu.DeviceToHostCopyThunkProto device_to_host_copy_thunk = 4; bool has_device_to_host_copy_thunk() const; private: bool _internal_has_device_to_host_copy_thunk() const; public: void clear_device_to_host_copy_thunk() ; const ::xla::gpu::DeviceToHostCopyThunkProto& device_to_host_copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::DeviceToHostCopyThunkProto* release_device_to_host_copy_thunk(); ::xla::gpu::DeviceToHostCopyThunkProto* mutable_device_to_host_copy_thunk(); void set_allocated_device_to_host_copy_thunk(::xla::gpu::DeviceToHostCopyThunkProto* value); void unsafe_arena_set_allocated_device_to_host_copy_thunk(::xla::gpu::DeviceToHostCopyThunkProto* value); ::xla::gpu::DeviceToHostCopyThunkProto* unsafe_arena_release_device_to_host_copy_thunk(); private: const ::xla::gpu::DeviceToHostCopyThunkProto& _internal_device_to_host_copy_thunk() const; ::xla::gpu::DeviceToHostCopyThunkProto* _internal_mutable_device_to_host_copy_thunk(); public: // .xla.gpu.HostToDeviceCopyThunkProto host_to_device_copy_thunk = 5; bool has_host_to_device_copy_thunk() const; private: bool _internal_has_host_to_device_copy_thunk() const; public: void clear_host_to_device_copy_thunk() ; const ::xla::gpu::HostToDeviceCopyThunkProto& host_to_device_copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::HostToDeviceCopyThunkProto* release_host_to_device_copy_thunk(); ::xla::gpu::HostToDeviceCopyThunkProto* mutable_host_to_device_copy_thunk(); void set_allocated_host_to_device_copy_thunk(::xla::gpu::HostToDeviceCopyThunkProto* value); void unsafe_arena_set_allocated_host_to_device_copy_thunk(::xla::gpu::HostToDeviceCopyThunkProto* value); ::xla::gpu::HostToDeviceCopyThunkProto* unsafe_arena_release_host_to_device_copy_thunk(); private: const ::xla::gpu::HostToDeviceCopyThunkProto& _internal_host_to_device_copy_thunk() const; ::xla::gpu::HostToDeviceCopyThunkProto* _internal_mutable_host_to_device_copy_thunk(); public: // .xla.gpu.ConditionalThunkProto conditional_thunk = 6; bool has_conditional_thunk() const; private: bool _internal_has_conditional_thunk() const; public: void clear_conditional_thunk() ; const ::xla::gpu::ConditionalThunkProto& conditional_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::ConditionalThunkProto* release_conditional_thunk(); ::xla::gpu::ConditionalThunkProto* mutable_conditional_thunk(); void set_allocated_conditional_thunk(::xla::gpu::ConditionalThunkProto* value); void unsafe_arena_set_allocated_conditional_thunk(::xla::gpu::ConditionalThunkProto* value); ::xla::gpu::ConditionalThunkProto* unsafe_arena_release_conditional_thunk(); private: const ::xla::gpu::ConditionalThunkProto& _internal_conditional_thunk() const; ::xla::gpu::ConditionalThunkProto* _internal_mutable_conditional_thunk(); public: // .xla.gpu.WhileThunkProto while_thunk = 7; bool has_while_thunk() const; private: bool _internal_has_while_thunk() const; public: void clear_while_thunk() ; const ::xla::gpu::WhileThunkProto& while_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::WhileThunkProto* release_while_thunk(); ::xla::gpu::WhileThunkProto* mutable_while_thunk(); void set_allocated_while_thunk(::xla::gpu::WhileThunkProto* value); void unsafe_arena_set_allocated_while_thunk(::xla::gpu::WhileThunkProto* value); ::xla::gpu::WhileThunkProto* unsafe_arena_release_while_thunk(); private: const ::xla::gpu::WhileThunkProto& _internal_while_thunk() const; ::xla::gpu::WhileThunkProto* _internal_mutable_while_thunk(); public: // .xla.gpu.KernelThunkProto kernel_thunk = 8; bool has_kernel_thunk() const; private: bool _internal_has_kernel_thunk() const; public: void clear_kernel_thunk() ; const ::xla::gpu::KernelThunkProto& kernel_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::KernelThunkProto* release_kernel_thunk(); ::xla::gpu::KernelThunkProto* mutable_kernel_thunk(); void set_allocated_kernel_thunk(::xla::gpu::KernelThunkProto* value); void unsafe_arena_set_allocated_kernel_thunk(::xla::gpu::KernelThunkProto* value); ::xla::gpu::KernelThunkProto* unsafe_arena_release_kernel_thunk(); private: const ::xla::gpu::KernelThunkProto& _internal_kernel_thunk() const; ::xla::gpu::KernelThunkProto* _internal_mutable_kernel_thunk(); public: // .xla.gpu.GemmThunkProto gemm_thunk = 9; bool has_gemm_thunk() const; private: bool _internal_has_gemm_thunk() const; public: void clear_gemm_thunk() ; const ::xla::gpu::GemmThunkProto& gemm_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::GemmThunkProto* release_gemm_thunk(); ::xla::gpu::GemmThunkProto* mutable_gemm_thunk(); void set_allocated_gemm_thunk(::xla::gpu::GemmThunkProto* value); void unsafe_arena_set_allocated_gemm_thunk(::xla::gpu::GemmThunkProto* value); ::xla::gpu::GemmThunkProto* unsafe_arena_release_gemm_thunk(); private: const ::xla::gpu::GemmThunkProto& _internal_gemm_thunk() const; ::xla::gpu::GemmThunkProto* _internal_mutable_gemm_thunk(); public: // .xla.gpu.WaitForStreamsThunkProto wait_for_streams_thunk = 10; bool has_wait_for_streams_thunk() const; private: bool _internal_has_wait_for_streams_thunk() const; public: void clear_wait_for_streams_thunk() ; const ::xla::gpu::WaitForStreamsThunkProto& wait_for_streams_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::WaitForStreamsThunkProto* release_wait_for_streams_thunk(); ::xla::gpu::WaitForStreamsThunkProto* mutable_wait_for_streams_thunk(); void set_allocated_wait_for_streams_thunk(::xla::gpu::WaitForStreamsThunkProto* value); void unsafe_arena_set_allocated_wait_for_streams_thunk(::xla::gpu::WaitForStreamsThunkProto* value); ::xla::gpu::WaitForStreamsThunkProto* unsafe_arena_release_wait_for_streams_thunk(); private: const ::xla::gpu::WaitForStreamsThunkProto& _internal_wait_for_streams_thunk() const; ::xla::gpu::WaitForStreamsThunkProto* _internal_mutable_wait_for_streams_thunk(); public: // .xla.gpu.TriangularSolveThunkProto triangular_solve_thunk = 11; bool has_triangular_solve_thunk() const; private: bool _internal_has_triangular_solve_thunk() const; public: void clear_triangular_solve_thunk() ; const ::xla::gpu::TriangularSolveThunkProto& triangular_solve_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::TriangularSolveThunkProto* release_triangular_solve_thunk(); ::xla::gpu::TriangularSolveThunkProto* mutable_triangular_solve_thunk(); void set_allocated_triangular_solve_thunk(::xla::gpu::TriangularSolveThunkProto* value); void unsafe_arena_set_allocated_triangular_solve_thunk(::xla::gpu::TriangularSolveThunkProto* value); ::xla::gpu::TriangularSolveThunkProto* unsafe_arena_release_triangular_solve_thunk(); private: const ::xla::gpu::TriangularSolveThunkProto& _internal_triangular_solve_thunk() const; ::xla::gpu::TriangularSolveThunkProto* _internal_mutable_triangular_solve_thunk(); public: // .xla.gpu.ReplicaIdThunkProto replica_id_thunk = 12; bool has_replica_id_thunk() const; private: bool _internal_has_replica_id_thunk() const; public: void clear_replica_id_thunk() ; const ::xla::gpu::ReplicaIdThunkProto& replica_id_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::ReplicaIdThunkProto* release_replica_id_thunk(); ::xla::gpu::ReplicaIdThunkProto* mutable_replica_id_thunk(); void set_allocated_replica_id_thunk(::xla::gpu::ReplicaIdThunkProto* value); void unsafe_arena_set_allocated_replica_id_thunk(::xla::gpu::ReplicaIdThunkProto* value); ::xla::gpu::ReplicaIdThunkProto* unsafe_arena_release_replica_id_thunk(); private: const ::xla::gpu::ReplicaIdThunkProto& _internal_replica_id_thunk() const; ::xla::gpu::ReplicaIdThunkProto* _internal_mutable_replica_id_thunk(); public: // .xla.gpu.DeviceToDeviceCopyThunkProto device_to_device_copy_thunk = 13; bool has_device_to_device_copy_thunk() const; private: bool _internal_has_device_to_device_copy_thunk() const; public: void clear_device_to_device_copy_thunk() ; const ::xla::gpu::DeviceToDeviceCopyThunkProto& device_to_device_copy_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::DeviceToDeviceCopyThunkProto* release_device_to_device_copy_thunk(); ::xla::gpu::DeviceToDeviceCopyThunkProto* mutable_device_to_device_copy_thunk(); void set_allocated_device_to_device_copy_thunk(::xla::gpu::DeviceToDeviceCopyThunkProto* value); void unsafe_arena_set_allocated_device_to_device_copy_thunk(::xla::gpu::DeviceToDeviceCopyThunkProto* value); ::xla::gpu::DeviceToDeviceCopyThunkProto* unsafe_arena_release_device_to_device_copy_thunk(); private: const ::xla::gpu::DeviceToDeviceCopyThunkProto& _internal_device_to_device_copy_thunk() const; ::xla::gpu::DeviceToDeviceCopyThunkProto* _internal_mutable_device_to_device_copy_thunk(); public: // .xla.gpu.PartitionIdThunkProto partition_id_thunk = 14; bool has_partition_id_thunk() const; private: bool _internal_has_partition_id_thunk() const; public: void clear_partition_id_thunk() ; const ::xla::gpu::PartitionIdThunkProto& partition_id_thunk() const; PROTOBUF_NODISCARD ::xla::gpu::PartitionIdThunkProto* release_partition_id_thunk(); ::xla::gpu::PartitionIdThunkProto* mutable_partition_id_thunk(); void set_allocated_partition_id_thunk(::xla::gpu::PartitionIdThunkProto* value); void unsafe_arena_set_allocated_partition_id_thunk(::xla::gpu::PartitionIdThunkProto* value); ::xla::gpu::PartitionIdThunkProto* unsafe_arena_release_partition_id_thunk(); private: const ::xla::gpu::PartitionIdThunkProto& _internal_partition_id_thunk() const; ::xla::gpu::PartitionIdThunkProto* _internal_mutable_partition_id_thunk(); public: void clear_impl(); ImplCase impl_case() const; // @@protoc_insertion_point(class_scope:xla.gpu.ThunkProto) private: class _Internal; void set_has_sequential_thunk(); void set_has_copy_thunk(); void set_has_device_to_host_copy_thunk(); void set_has_host_to_device_copy_thunk(); void set_has_conditional_thunk(); void set_has_while_thunk(); void set_has_kernel_thunk(); void set_has_gemm_thunk(); void set_has_wait_for_streams_thunk(); void set_has_triangular_solve_thunk(); void set_has_replica_id_thunk(); void set_has_device_to_device_copy_thunk(); void set_has_partition_id_thunk(); inline bool has_impl() const; inline void clear_has_impl(); friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 0, 14, 14, 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 ThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::ThunkInfoProto* thunk_info_; union ImplUnion { constexpr ImplUnion() : _constinit_{} {} ::google::protobuf::internal::ConstantInitialized _constinit_; ::xla::gpu::SequentialThunkProto* sequential_thunk_; ::xla::gpu::CopyThunkProto* copy_thunk_; ::xla::gpu::DeviceToHostCopyThunkProto* device_to_host_copy_thunk_; ::xla::gpu::HostToDeviceCopyThunkProto* host_to_device_copy_thunk_; ::xla::gpu::ConditionalThunkProto* conditional_thunk_; ::xla::gpu::WhileThunkProto* while_thunk_; ::xla::gpu::KernelThunkProto* kernel_thunk_; ::xla::gpu::GemmThunkProto* gemm_thunk_; ::xla::gpu::WaitForStreamsThunkProto* wait_for_streams_thunk_; ::xla::gpu::TriangularSolveThunkProto* triangular_solve_thunk_; ::xla::gpu::ReplicaIdThunkProto* replica_id_thunk_; ::xla::gpu::DeviceToDeviceCopyThunkProto* device_to_device_copy_thunk_; ::xla::gpu::PartitionIdThunkProto* partition_id_thunk_; } impl_; ::uint32_t _oneof_case_[1]; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // ------------------------------------------------------------------- class PROTOBUF_EXPORT WhileThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.WhileThunkProto) */ { public: inline WhileThunkProto() : WhileThunkProto(nullptr) {} ~WhileThunkProto() PROTOBUF_FINAL; template explicit PROTOBUF_CONSTEXPR WhileThunkProto( ::google::protobuf::internal::ConstantInitialized); inline WhileThunkProto(const WhileThunkProto& from) : WhileThunkProto(nullptr, from) {} inline WhileThunkProto(WhileThunkProto&& from) noexcept : WhileThunkProto(nullptr, std::move(from)) {} inline WhileThunkProto& operator=(const WhileThunkProto& from) { CopyFrom(from); return *this; } inline WhileThunkProto& operator=(WhileThunkProto&& from) noexcept { if (this == &from) return *this; if (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 WhileThunkProto& default_instance() { return *internal_default_instance(); } static inline const WhileThunkProto* internal_default_instance() { return reinterpret_cast( &_WhileThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 6; friend void swap(WhileThunkProto& a, WhileThunkProto& b) { a.Swap(&b); } inline void Swap(WhileThunkProto* other) { if (other == this) return; #ifdef PROTOBUF_FORCE_COPY_IN_SWAP if (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(WhileThunkProto* other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- WhileThunkProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const WhileThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const WhileThunkProto& from) { WhileThunkProto::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(WhileThunkProto* other); private: friend class ::google::protobuf::internal::AnyMetadata; static ::absl::string_view FullMessageName() { return "xla.gpu.WhileThunkProto"; } protected: explicit WhileThunkProto(::google::protobuf::Arena* arena); WhileThunkProto(::google::protobuf::Arena* arena, const WhileThunkProto& from); WhileThunkProto(::google::protobuf::Arena* arena, WhileThunkProto&& from) noexcept : WhileThunkProto(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 { kConditionResultBufferIndexFieldNumber = 1, kConditionThunkSequenceFieldNumber = 2, kBodyThunkSequenceFieldNumber = 3, kTripCountFieldNumber = 4, }; // .xla.buffer_assignment.BufferAllocationSliceProto condition_result_buffer_index = 1; bool has_condition_result_buffer_index() const; void clear_condition_result_buffer_index() ; const ::xla::buffer_assignment::BufferAllocationSliceProto& condition_result_buffer_index() const; PROTOBUF_NODISCARD ::xla::buffer_assignment::BufferAllocationSliceProto* release_condition_result_buffer_index(); ::xla::buffer_assignment::BufferAllocationSliceProto* mutable_condition_result_buffer_index(); void set_allocated_condition_result_buffer_index(::xla::buffer_assignment::BufferAllocationSliceProto* value); void unsafe_arena_set_allocated_condition_result_buffer_index(::xla::buffer_assignment::BufferAllocationSliceProto* value); ::xla::buffer_assignment::BufferAllocationSliceProto* unsafe_arena_release_condition_result_buffer_index(); private: const ::xla::buffer_assignment::BufferAllocationSliceProto& _internal_condition_result_buffer_index() const; ::xla::buffer_assignment::BufferAllocationSliceProto* _internal_mutable_condition_result_buffer_index(); public: // .xla.gpu.SequentialThunkProto condition_thunk_sequence = 2; bool has_condition_thunk_sequence() const; void clear_condition_thunk_sequence() ; const ::xla::gpu::SequentialThunkProto& condition_thunk_sequence() const; PROTOBUF_NODISCARD ::xla::gpu::SequentialThunkProto* release_condition_thunk_sequence(); ::xla::gpu::SequentialThunkProto* mutable_condition_thunk_sequence(); void set_allocated_condition_thunk_sequence(::xla::gpu::SequentialThunkProto* value); void unsafe_arena_set_allocated_condition_thunk_sequence(::xla::gpu::SequentialThunkProto* value); ::xla::gpu::SequentialThunkProto* unsafe_arena_release_condition_thunk_sequence(); private: const ::xla::gpu::SequentialThunkProto& _internal_condition_thunk_sequence() const; ::xla::gpu::SequentialThunkProto* _internal_mutable_condition_thunk_sequence(); public: // .xla.gpu.SequentialThunkProto body_thunk_sequence = 3; bool has_body_thunk_sequence() const; void clear_body_thunk_sequence() ; const ::xla::gpu::SequentialThunkProto& body_thunk_sequence() const; PROTOBUF_NODISCARD ::xla::gpu::SequentialThunkProto* release_body_thunk_sequence(); ::xla::gpu::SequentialThunkProto* mutable_body_thunk_sequence(); void set_allocated_body_thunk_sequence(::xla::gpu::SequentialThunkProto* value); void unsafe_arena_set_allocated_body_thunk_sequence(::xla::gpu::SequentialThunkProto* value); ::xla::gpu::SequentialThunkProto* unsafe_arena_release_body_thunk_sequence(); private: const ::xla::gpu::SequentialThunkProto& _internal_body_thunk_sequence() const; ::xla::gpu::SequentialThunkProto* _internal_mutable_body_thunk_sequence(); public: // optional int64 trip_count = 4; bool has_trip_count() const; void clear_trip_count() ; ::int64_t trip_count() const; void set_trip_count(::int64_t value); private: ::int64_t _internal_trip_count() const; void _internal_set_trip_count(::int64_t value); public: // @@protoc_insertion_point(class_scope:xla.gpu.WhileThunkProto) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable< 2, 4, 3, 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 WhileThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; mutable ::google::protobuf::internal::CachedSize _cached_size_; ::xla::buffer_assignment::BufferAllocationSliceProto* condition_result_buffer_index_; ::xla::gpu::SequentialThunkProto* condition_thunk_sequence_; ::xla::gpu::SequentialThunkProto* body_thunk_sequence_; ::int64_t trip_count_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fthunk_2eproto; }; // =================================================================== // =================================================================== #ifdef __GNUC__ #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif // __GNUC__ // ------------------------------------------------------------------- // ThunkInfoProto // string profile_annotation = 1; inline void ThunkInfoProto::clear_profile_annotation() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.profile_annotation_.ClearToEmpty(); } inline const std::string& ThunkInfoProto::profile_annotation() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkInfoProto.profile_annotation) return _internal_profile_annotation(); } template inline PROTOBUF_ALWAYS_INLINE void ThunkInfoProto::set_profile_annotation(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.profile_annotation_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.ThunkInfoProto.profile_annotation) } inline std::string* ThunkInfoProto::mutable_profile_annotation() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_profile_annotation(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkInfoProto.profile_annotation) return _s; } inline const std::string& ThunkInfoProto::_internal_profile_annotation() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.profile_annotation_.Get(); } inline void ThunkInfoProto::_internal_set_profile_annotation(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.profile_annotation_.Set(value, GetArena()); } inline std::string* ThunkInfoProto::_internal_mutable_profile_annotation() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.profile_annotation_.Mutable( GetArena()); } inline std::string* ThunkInfoProto::release_profile_annotation() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ThunkInfoProto.profile_annotation) return _impl_.profile_annotation_.Release(); } inline void ThunkInfoProto::set_allocated_profile_annotation(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.profile_annotation_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.profile_annotation_.IsDefault()) { _impl_.profile_annotation_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.ThunkInfoProto.profile_annotation) } // int64 execution_stream_id = 2; inline void ThunkInfoProto::clear_execution_stream_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.execution_stream_id_ = ::int64_t{0}; } inline ::int64_t ThunkInfoProto::execution_stream_id() const { // @@protoc_insertion_point(field_get:xla.gpu.ThunkInfoProto.execution_stream_id) return _internal_execution_stream_id(); } inline void ThunkInfoProto::set_execution_stream_id(::int64_t value) { _internal_set_execution_stream_id(value); // @@protoc_insertion_point(field_set:xla.gpu.ThunkInfoProto.execution_stream_id) } inline ::int64_t ThunkInfoProto::_internal_execution_stream_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.execution_stream_id_; } inline void ThunkInfoProto::_internal_set_execution_stream_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.execution_stream_id_ = value; } // ------------------------------------------------------------------- // CopyThunkProto // .xla.buffer_assignment.BufferAllocationSliceProto source_buffer = 1; inline bool CopyThunkProto::has_source_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.source_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& CopyThunkProto::_internal_source_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.source_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& CopyThunkProto::source_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CopyThunkProto.source_buffer) return _internal_source_buffer(); } inline void CopyThunkProto::unsafe_arena_set_allocated_source_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_buffer_); } _impl_.source_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.CopyThunkProto.source_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::release_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.source_buffer_; _impl_.source_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::unsafe_arena_release_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CopyThunkProto.source_buffer) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.source_buffer_; _impl_.source_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::_internal_mutable_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.source_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.source_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.source_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::mutable_source_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_source_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.CopyThunkProto.source_buffer) return _msg; } inline void CopyThunkProto::set_allocated_source_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_buffer_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.source_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CopyThunkProto.source_buffer) } // .xla.buffer_assignment.BufferAllocationSliceProto destination_buffer = 2; inline bool CopyThunkProto::has_destination_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.destination_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& CopyThunkProto::_internal_destination_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.destination_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& CopyThunkProto::destination_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.CopyThunkProto.destination_buffer) return _internal_destination_buffer(); } inline void CopyThunkProto::unsafe_arena_set_allocated_destination_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.destination_buffer_); } _impl_.destination_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.CopyThunkProto.destination_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::release_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.destination_buffer_; _impl_.destination_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::unsafe_arena_release_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.CopyThunkProto.destination_buffer) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.destination_buffer_; _impl_.destination_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::_internal_mutable_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.destination_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.destination_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.destination_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* CopyThunkProto::mutable_destination_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_destination_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.CopyThunkProto.destination_buffer) return _msg; } inline void CopyThunkProto::set_allocated_destination_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.destination_buffer_); } 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_.destination_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CopyThunkProto.destination_buffer) } // int64 mem_size = 3; inline void CopyThunkProto::clear_mem_size() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mem_size_ = ::int64_t{0}; } inline ::int64_t CopyThunkProto::mem_size() const { // @@protoc_insertion_point(field_get:xla.gpu.CopyThunkProto.mem_size) return _internal_mem_size(); } inline void CopyThunkProto::set_mem_size(::int64_t value) { _internal_set_mem_size(value); // @@protoc_insertion_point(field_set:xla.gpu.CopyThunkProto.mem_size) } inline ::int64_t CopyThunkProto::_internal_mem_size() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.mem_size_; } inline void CopyThunkProto::_internal_set_mem_size(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mem_size_ = value; } // ------------------------------------------------------------------- // DeviceToHostCopyThunkProto // .xla.gpu.CopyThunkProto copy_thunk = 1; inline bool DeviceToHostCopyThunkProto::has_copy_thunk() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.copy_thunk_ != nullptr); return value; } inline void DeviceToHostCopyThunkProto::clear_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ != nullptr) _impl_.copy_thunk_->Clear(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const ::xla::gpu::CopyThunkProto& DeviceToHostCopyThunkProto::_internal_copy_thunk() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::CopyThunkProto* p = _impl_.copy_thunk_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_CopyThunkProto_default_instance_); } inline const ::xla::gpu::CopyThunkProto& DeviceToHostCopyThunkProto::copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.DeviceToHostCopyThunkProto.copy_thunk) return _internal_copy_thunk(); } inline void DeviceToHostCopyThunkProto::unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.copy_thunk_); } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(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.DeviceToHostCopyThunkProto.copy_thunk) } inline ::xla::gpu::CopyThunkProto* DeviceToHostCopyThunkProto::release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* released = _impl_.copy_thunk_; _impl_.copy_thunk_ = 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::CopyThunkProto* DeviceToHostCopyThunkProto::unsafe_arena_release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.DeviceToHostCopyThunkProto.copy_thunk) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* temp = _impl_.copy_thunk_; _impl_.copy_thunk_ = nullptr; return temp; } inline ::xla::gpu::CopyThunkProto* DeviceToHostCopyThunkProto::_internal_mutable_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CopyThunkProto>(GetArena()); _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(p); } return _impl_.copy_thunk_; } inline ::xla::gpu::CopyThunkProto* DeviceToHostCopyThunkProto::mutable_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::gpu::CopyThunkProto* _msg = _internal_mutable_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.DeviceToHostCopyThunkProto.copy_thunk) return _msg; } inline void DeviceToHostCopyThunkProto::set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.copy_thunk_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.DeviceToHostCopyThunkProto.copy_thunk) } // ------------------------------------------------------------------- // HostToDeviceCopyThunkProto // .xla.gpu.CopyThunkProto copy_thunk = 1; inline bool HostToDeviceCopyThunkProto::has_copy_thunk() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.copy_thunk_ != nullptr); return value; } inline void HostToDeviceCopyThunkProto::clear_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ != nullptr) _impl_.copy_thunk_->Clear(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const ::xla::gpu::CopyThunkProto& HostToDeviceCopyThunkProto::_internal_copy_thunk() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::CopyThunkProto* p = _impl_.copy_thunk_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_CopyThunkProto_default_instance_); } inline const ::xla::gpu::CopyThunkProto& HostToDeviceCopyThunkProto::copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.HostToDeviceCopyThunkProto.copy_thunk) return _internal_copy_thunk(); } inline void HostToDeviceCopyThunkProto::unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.copy_thunk_); } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(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.HostToDeviceCopyThunkProto.copy_thunk) } inline ::xla::gpu::CopyThunkProto* HostToDeviceCopyThunkProto::release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* released = _impl_.copy_thunk_; _impl_.copy_thunk_ = 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::CopyThunkProto* HostToDeviceCopyThunkProto::unsafe_arena_release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.HostToDeviceCopyThunkProto.copy_thunk) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* temp = _impl_.copy_thunk_; _impl_.copy_thunk_ = nullptr; return temp; } inline ::xla::gpu::CopyThunkProto* HostToDeviceCopyThunkProto::_internal_mutable_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CopyThunkProto>(GetArena()); _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(p); } return _impl_.copy_thunk_; } inline ::xla::gpu::CopyThunkProto* HostToDeviceCopyThunkProto::mutable_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::gpu::CopyThunkProto* _msg = _internal_mutable_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.HostToDeviceCopyThunkProto.copy_thunk) return _msg; } inline void HostToDeviceCopyThunkProto::set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.copy_thunk_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.HostToDeviceCopyThunkProto.copy_thunk) } // ------------------------------------------------------------------- // DeviceToDeviceCopyThunkProto // .xla.gpu.CopyThunkProto copy_thunk = 1; inline bool DeviceToDeviceCopyThunkProto::has_copy_thunk() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.copy_thunk_ != nullptr); return value; } inline void DeviceToDeviceCopyThunkProto::clear_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ != nullptr) _impl_.copy_thunk_->Clear(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const ::xla::gpu::CopyThunkProto& DeviceToDeviceCopyThunkProto::_internal_copy_thunk() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::CopyThunkProto* p = _impl_.copy_thunk_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_CopyThunkProto_default_instance_); } inline const ::xla::gpu::CopyThunkProto& DeviceToDeviceCopyThunkProto::copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.DeviceToDeviceCopyThunkProto.copy_thunk) return _internal_copy_thunk(); } inline void DeviceToDeviceCopyThunkProto::unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.copy_thunk_); } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(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.DeviceToDeviceCopyThunkProto.copy_thunk) } inline ::xla::gpu::CopyThunkProto* DeviceToDeviceCopyThunkProto::release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* released = _impl_.copy_thunk_; _impl_.copy_thunk_ = 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::CopyThunkProto* DeviceToDeviceCopyThunkProto::unsafe_arena_release_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.DeviceToDeviceCopyThunkProto.copy_thunk) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::CopyThunkProto* temp = _impl_.copy_thunk_; _impl_.copy_thunk_ = nullptr; return temp; } inline ::xla::gpu::CopyThunkProto* DeviceToDeviceCopyThunkProto::_internal_mutable_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.copy_thunk_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CopyThunkProto>(GetArena()); _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(p); } return _impl_.copy_thunk_; } inline ::xla::gpu::CopyThunkProto* DeviceToDeviceCopyThunkProto::mutable_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::gpu::CopyThunkProto* _msg = _internal_mutable_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.DeviceToDeviceCopyThunkProto.copy_thunk) return _msg; } inline void DeviceToDeviceCopyThunkProto::set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.copy_thunk_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.copy_thunk_ = reinterpret_cast<::xla::gpu::CopyThunkProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.DeviceToDeviceCopyThunkProto.copy_thunk) } // ------------------------------------------------------------------- // ConditionalThunkProto // .xla.buffer_assignment.BufferAllocationSliceProto branch_index_buffer = 1; inline bool ConditionalThunkProto::has_branch_index_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.branch_index_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& ConditionalThunkProto::_internal_branch_index_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.branch_index_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& ConditionalThunkProto::branch_index_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ConditionalThunkProto.branch_index_buffer) return _internal_branch_index_buffer(); } inline void ConditionalThunkProto::unsafe_arena_set_allocated_branch_index_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.branch_index_buffer_); } _impl_.branch_index_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.ConditionalThunkProto.branch_index_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ConditionalThunkProto::release_branch_index_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.branch_index_buffer_; _impl_.branch_index_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* ConditionalThunkProto::unsafe_arena_release_branch_index_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ConditionalThunkProto.branch_index_buffer) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.branch_index_buffer_; _impl_.branch_index_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ConditionalThunkProto::_internal_mutable_branch_index_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.branch_index_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.branch_index_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.branch_index_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ConditionalThunkProto::mutable_branch_index_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_branch_index_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.ConditionalThunkProto.branch_index_buffer) return _msg; } inline void ConditionalThunkProto::set_allocated_branch_index_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.branch_index_buffer_); } 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_.branch_index_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.ConditionalThunkProto.branch_index_buffer) } // repeated .xla.gpu.SequentialThunkProto branch_thunks = 2; inline int ConditionalThunkProto::_internal_branch_thunks_size() const { return _internal_branch_thunks().size(); } inline int ConditionalThunkProto::branch_thunks_size() const { return _internal_branch_thunks_size(); } inline void ConditionalThunkProto::clear_branch_thunks() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.branch_thunks_.Clear(); } inline ::xla::gpu::SequentialThunkProto* ConditionalThunkProto::mutable_branch_thunks(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.ConditionalThunkProto.branch_thunks) return _internal_mutable_branch_thunks()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>* ConditionalThunkProto::mutable_branch_thunks() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.ConditionalThunkProto.branch_thunks) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_branch_thunks(); } inline const ::xla::gpu::SequentialThunkProto& ConditionalThunkProto::branch_thunks(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ConditionalThunkProto.branch_thunks) return _internal_branch_thunks().Get(index); } inline ::xla::gpu::SequentialThunkProto* ConditionalThunkProto::add_branch_thunks() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::SequentialThunkProto* _add = _internal_mutable_branch_thunks()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.ConditionalThunkProto.branch_thunks) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>& ConditionalThunkProto::branch_thunks() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.ConditionalThunkProto.branch_thunks) return _internal_branch_thunks(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>& ConditionalThunkProto::_internal_branch_thunks() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.branch_thunks_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::SequentialThunkProto>* ConditionalThunkProto::_internal_mutable_branch_thunks() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.branch_thunks_; } // bool branch_index_is_bool = 3; inline void ConditionalThunkProto::clear_branch_index_is_bool() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.branch_index_is_bool_ = false; } inline bool ConditionalThunkProto::branch_index_is_bool() const { // @@protoc_insertion_point(field_get:xla.gpu.ConditionalThunkProto.branch_index_is_bool) return _internal_branch_index_is_bool(); } inline void ConditionalThunkProto::set_branch_index_is_bool(bool value) { _internal_set_branch_index_is_bool(value); // @@protoc_insertion_point(field_set:xla.gpu.ConditionalThunkProto.branch_index_is_bool) } inline bool ConditionalThunkProto::_internal_branch_index_is_bool() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.branch_index_is_bool_; } inline void ConditionalThunkProto::_internal_set_branch_index_is_bool(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.branch_index_is_bool_ = value; } // ------------------------------------------------------------------- // WhileThunkProto // .xla.buffer_assignment.BufferAllocationSliceProto condition_result_buffer_index = 1; inline bool WhileThunkProto::has_condition_result_buffer_index() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.condition_result_buffer_index_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& WhileThunkProto::_internal_condition_result_buffer_index() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.condition_result_buffer_index_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& WhileThunkProto::condition_result_buffer_index() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.WhileThunkProto.condition_result_buffer_index) return _internal_condition_result_buffer_index(); } inline void WhileThunkProto::unsafe_arena_set_allocated_condition_result_buffer_index(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.condition_result_buffer_index_); } _impl_.condition_result_buffer_index_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.WhileThunkProto.condition_result_buffer_index) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* WhileThunkProto::release_condition_result_buffer_index() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.condition_result_buffer_index_; _impl_.condition_result_buffer_index_ = 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::buffer_assignment::BufferAllocationSliceProto* WhileThunkProto::unsafe_arena_release_condition_result_buffer_index() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.WhileThunkProto.condition_result_buffer_index) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.condition_result_buffer_index_; _impl_.condition_result_buffer_index_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* WhileThunkProto::_internal_mutable_condition_result_buffer_index() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.condition_result_buffer_index_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.condition_result_buffer_index_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.condition_result_buffer_index_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* WhileThunkProto::mutable_condition_result_buffer_index() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_condition_result_buffer_index(); // @@protoc_insertion_point(field_mutable:xla.gpu.WhileThunkProto.condition_result_buffer_index) return _msg; } inline void WhileThunkProto::set_allocated_condition_result_buffer_index(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.condition_result_buffer_index_); } 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_.condition_result_buffer_index_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.WhileThunkProto.condition_result_buffer_index) } // .xla.gpu.SequentialThunkProto condition_thunk_sequence = 2; inline bool WhileThunkProto::has_condition_thunk_sequence() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.condition_thunk_sequence_ != nullptr); return value; } inline void WhileThunkProto::clear_condition_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.condition_thunk_sequence_ != nullptr) _impl_.condition_thunk_sequence_->Clear(); _impl_._has_bits_[0] &= ~0x00000002u; } inline const ::xla::gpu::SequentialThunkProto& WhileThunkProto::_internal_condition_thunk_sequence() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::SequentialThunkProto* p = _impl_.condition_thunk_sequence_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_SequentialThunkProto_default_instance_); } inline const ::xla::gpu::SequentialThunkProto& WhileThunkProto::condition_thunk_sequence() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.WhileThunkProto.condition_thunk_sequence) return _internal_condition_thunk_sequence(); } inline void WhileThunkProto::unsafe_arena_set_allocated_condition_thunk_sequence(::xla::gpu::SequentialThunkProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.condition_thunk_sequence_); } _impl_.condition_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(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.WhileThunkProto.condition_thunk_sequence) } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::release_condition_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::gpu::SequentialThunkProto* released = _impl_.condition_thunk_sequence_; _impl_.condition_thunk_sequence_ = 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::SequentialThunkProto* WhileThunkProto::unsafe_arena_release_condition_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.WhileThunkProto.condition_thunk_sequence) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::gpu::SequentialThunkProto* temp = _impl_.condition_thunk_sequence_; _impl_.condition_thunk_sequence_ = nullptr; return temp; } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::_internal_mutable_condition_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.condition_thunk_sequence_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::SequentialThunkProto>(GetArena()); _impl_.condition_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(p); } return _impl_.condition_thunk_sequence_; } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::mutable_condition_thunk_sequence() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::gpu::SequentialThunkProto* _msg = _internal_mutable_condition_thunk_sequence(); // @@protoc_insertion_point(field_mutable:xla.gpu.WhileThunkProto.condition_thunk_sequence) return _msg; } inline void WhileThunkProto::set_allocated_condition_thunk_sequence(::xla::gpu::SequentialThunkProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.condition_thunk_sequence_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000002u; } else { _impl_._has_bits_[0] &= ~0x00000002u; } _impl_.condition_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.WhileThunkProto.condition_thunk_sequence) } // .xla.gpu.SequentialThunkProto body_thunk_sequence = 3; inline bool WhileThunkProto::has_body_thunk_sequence() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.body_thunk_sequence_ != nullptr); return value; } inline void WhileThunkProto::clear_body_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.body_thunk_sequence_ != nullptr) _impl_.body_thunk_sequence_->Clear(); _impl_._has_bits_[0] &= ~0x00000004u; } inline const ::xla::gpu::SequentialThunkProto& WhileThunkProto::_internal_body_thunk_sequence() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::SequentialThunkProto* p = _impl_.body_thunk_sequence_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_SequentialThunkProto_default_instance_); } inline const ::xla::gpu::SequentialThunkProto& WhileThunkProto::body_thunk_sequence() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.WhileThunkProto.body_thunk_sequence) return _internal_body_thunk_sequence(); } inline void WhileThunkProto::unsafe_arena_set_allocated_body_thunk_sequence(::xla::gpu::SequentialThunkProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.body_thunk_sequence_); } _impl_.body_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(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.WhileThunkProto.body_thunk_sequence) } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::release_body_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::SequentialThunkProto* released = _impl_.body_thunk_sequence_; _impl_.body_thunk_sequence_ = 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::SequentialThunkProto* WhileThunkProto::unsafe_arena_release_body_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.WhileThunkProto.body_thunk_sequence) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::SequentialThunkProto* temp = _impl_.body_thunk_sequence_; _impl_.body_thunk_sequence_ = nullptr; return temp; } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::_internal_mutable_body_thunk_sequence() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.body_thunk_sequence_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::SequentialThunkProto>(GetArena()); _impl_.body_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(p); } return _impl_.body_thunk_sequence_; } inline ::xla::gpu::SequentialThunkProto* WhileThunkProto::mutable_body_thunk_sequence() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::gpu::SequentialThunkProto* _msg = _internal_mutable_body_thunk_sequence(); // @@protoc_insertion_point(field_mutable:xla.gpu.WhileThunkProto.body_thunk_sequence) return _msg; } inline void WhileThunkProto::set_allocated_body_thunk_sequence(::xla::gpu::SequentialThunkProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.body_thunk_sequence_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000004u; } else { _impl_._has_bits_[0] &= ~0x00000004u; } _impl_.body_thunk_sequence_ = reinterpret_cast<::xla::gpu::SequentialThunkProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.WhileThunkProto.body_thunk_sequence) } // optional int64 trip_count = 4; inline bool WhileThunkProto::has_trip_count() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; return value; } inline void WhileThunkProto::clear_trip_count() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.trip_count_ = ::int64_t{0}; _impl_._has_bits_[0] &= ~0x00000008u; } inline ::int64_t WhileThunkProto::trip_count() const { // @@protoc_insertion_point(field_get:xla.gpu.WhileThunkProto.trip_count) return _internal_trip_count(); } inline void WhileThunkProto::set_trip_count(::int64_t value) { _internal_set_trip_count(value); _impl_._has_bits_[0] |= 0x00000008u; // @@protoc_insertion_point(field_set:xla.gpu.WhileThunkProto.trip_count) } inline ::int64_t WhileThunkProto::_internal_trip_count() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.trip_count_; } inline void WhileThunkProto::_internal_set_trip_count(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.trip_count_ = value; } // ------------------------------------------------------------------- // KernelThunkProto // repeated .xla.buffer_assignment.BufferAllocationSliceProto args = 1; inline int KernelThunkProto::_internal_args_size() const { return _internal_args().size(); } inline int KernelThunkProto::args_size() const { return _internal_args_size(); } inline ::xla::buffer_assignment::BufferAllocationSliceProto* KernelThunkProto::mutable_args(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.KernelThunkProto.args) return _internal_mutable_args()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>* KernelThunkProto::mutable_args() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.KernelThunkProto.args) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_args(); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& KernelThunkProto::args(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.args) return _internal_args().Get(index); } inline ::xla::buffer_assignment::BufferAllocationSliceProto* KernelThunkProto::add_args() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::buffer_assignment::BufferAllocationSliceProto* _add = _internal_mutable_args()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.KernelThunkProto.args) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>& KernelThunkProto::args() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.KernelThunkProto.args) return _internal_args(); } inline const ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>& KernelThunkProto::_internal_args() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.args_; } inline ::google::protobuf::RepeatedPtrField<::xla::buffer_assignment::BufferAllocationSliceProto>* KernelThunkProto::_internal_mutable_args() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.args_; } // repeated bool written = 2; inline int KernelThunkProto::_internal_written_size() const { return _internal_written().size(); } inline int KernelThunkProto::written_size() const { return _internal_written_size(); } inline void KernelThunkProto::clear_written() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.written_.Clear(); } inline bool KernelThunkProto::written(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.written) return _internal_written().Get(index); } inline void KernelThunkProto::set_written(int index, bool value) { _internal_mutable_written()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.KernelThunkProto.written) } inline void KernelThunkProto::add_written(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_written()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.KernelThunkProto.written) } inline const ::google::protobuf::RepeatedField& KernelThunkProto::written() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.KernelThunkProto.written) return _internal_written(); } inline ::google::protobuf::RepeatedField* KernelThunkProto::mutable_written() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.KernelThunkProto.written) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_written(); } inline const ::google::protobuf::RepeatedField& KernelThunkProto::_internal_written() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.written_; } inline ::google::protobuf::RepeatedField* KernelThunkProto::_internal_mutable_written() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.written_; } // string kernel_name = 3; inline void KernelThunkProto::clear_kernel_name() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_name_.ClearToEmpty(); } inline const std::string& KernelThunkProto::kernel_name() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.kernel_name) return _internal_kernel_name(); } template inline PROTOBUF_ALWAYS_INLINE void KernelThunkProto::set_kernel_name(Arg_&& arg, Args_... args) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_name_.Set(static_cast(arg), args..., GetArena()); // @@protoc_insertion_point(field_set:xla.gpu.KernelThunkProto.kernel_name) } inline std::string* KernelThunkProto::mutable_kernel_name() ABSL_ATTRIBUTE_LIFETIME_BOUND { std::string* _s = _internal_mutable_kernel_name(); // @@protoc_insertion_point(field_mutable:xla.gpu.KernelThunkProto.kernel_name) return _s; } inline const std::string& KernelThunkProto::_internal_kernel_name() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.kernel_name_.Get(); } inline void KernelThunkProto::_internal_set_kernel_name(const std::string& value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_name_.Set(value, GetArena()); } inline std::string* KernelThunkProto::_internal_mutable_kernel_name() { ::google::protobuf::internal::TSanWrite(&_impl_); return _impl_.kernel_name_.Mutable( GetArena()); } inline std::string* KernelThunkProto::release_kernel_name() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.KernelThunkProto.kernel_name) return _impl_.kernel_name_.Release(); } inline void KernelThunkProto::set_allocated_kernel_name(std::string* value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.kernel_name_.SetAllocated(value, GetArena()); #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING if (_impl_.kernel_name_.IsDefault()) { _impl_.kernel_name_.Set("", GetArena()); } #endif // PROTOBUF_FORCE_COPY_DEFAULT_STRING // @@protoc_insertion_point(field_set_allocated:xla.gpu.KernelThunkProto.kernel_name) } // .xla.gpu.LaunchDimensionsProto launch_dimensions = 4; inline bool KernelThunkProto::has_launch_dimensions() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.launch_dimensions_ != nullptr); return value; } inline const ::xla::gpu::LaunchDimensionsProto& KernelThunkProto::_internal_launch_dimensions() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::LaunchDimensionsProto* p = _impl_.launch_dimensions_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_LaunchDimensionsProto_default_instance_); } inline const ::xla::gpu::LaunchDimensionsProto& KernelThunkProto::launch_dimensions() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.launch_dimensions) return _internal_launch_dimensions(); } inline void KernelThunkProto::unsafe_arena_set_allocated_launch_dimensions(::xla::gpu::LaunchDimensionsProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.launch_dimensions_); } _impl_.launch_dimensions_ = reinterpret_cast<::xla::gpu::LaunchDimensionsProto*>(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.KernelThunkProto.launch_dimensions) } inline ::xla::gpu::LaunchDimensionsProto* KernelThunkProto::release_launch_dimensions() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::LaunchDimensionsProto* released = _impl_.launch_dimensions_; _impl_.launch_dimensions_ = 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::LaunchDimensionsProto* KernelThunkProto::unsafe_arena_release_launch_dimensions() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.KernelThunkProto.launch_dimensions) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::LaunchDimensionsProto* temp = _impl_.launch_dimensions_; _impl_.launch_dimensions_ = nullptr; return temp; } inline ::xla::gpu::LaunchDimensionsProto* KernelThunkProto::_internal_mutable_launch_dimensions() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.launch_dimensions_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::LaunchDimensionsProto>(GetArena()); _impl_.launch_dimensions_ = reinterpret_cast<::xla::gpu::LaunchDimensionsProto*>(p); } return _impl_.launch_dimensions_; } inline ::xla::gpu::LaunchDimensionsProto* KernelThunkProto::mutable_launch_dimensions() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::gpu::LaunchDimensionsProto* _msg = _internal_mutable_launch_dimensions(); // @@protoc_insertion_point(field_mutable:xla.gpu.KernelThunkProto.launch_dimensions) return _msg; } inline void KernelThunkProto::set_allocated_launch_dimensions(::xla::gpu::LaunchDimensionsProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.launch_dimensions_); } 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_.launch_dimensions_ = reinterpret_cast<::xla::gpu::LaunchDimensionsProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.KernelThunkProto.launch_dimensions) } // optional .stream_executor.ClusterDimProto cluster_dim = 5; inline bool KernelThunkProto::has_cluster_dim() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.cluster_dim_ != nullptr); return value; } inline const ::stream_executor::ClusterDimProto& KernelThunkProto::_internal_cluster_dim() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::stream_executor::ClusterDimProto* p = _impl_.cluster_dim_; return p != nullptr ? *p : reinterpret_cast(::stream_executor::_ClusterDimProto_default_instance_); } inline const ::stream_executor::ClusterDimProto& KernelThunkProto::cluster_dim() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.cluster_dim) return _internal_cluster_dim(); } inline void KernelThunkProto::unsafe_arena_set_allocated_cluster_dim(::stream_executor::ClusterDimProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cluster_dim_); } _impl_.cluster_dim_ = reinterpret_cast<::stream_executor::ClusterDimProto*>(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.KernelThunkProto.cluster_dim) } inline ::stream_executor::ClusterDimProto* KernelThunkProto::release_cluster_dim() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::stream_executor::ClusterDimProto* released = _impl_.cluster_dim_; _impl_.cluster_dim_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::stream_executor::ClusterDimProto* KernelThunkProto::unsafe_arena_release_cluster_dim() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.KernelThunkProto.cluster_dim) _impl_._has_bits_[0] &= ~0x00000002u; ::stream_executor::ClusterDimProto* temp = _impl_.cluster_dim_; _impl_.cluster_dim_ = nullptr; return temp; } inline ::stream_executor::ClusterDimProto* KernelThunkProto::_internal_mutable_cluster_dim() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.cluster_dim_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::stream_executor::ClusterDimProto>(GetArena()); _impl_.cluster_dim_ = reinterpret_cast<::stream_executor::ClusterDimProto*>(p); } return _impl_.cluster_dim_; } inline ::stream_executor::ClusterDimProto* KernelThunkProto::mutable_cluster_dim() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::stream_executor::ClusterDimProto* _msg = _internal_mutable_cluster_dim(); // @@protoc_insertion_point(field_mutable:xla.gpu.KernelThunkProto.cluster_dim) return _msg; } inline void KernelThunkProto::set_allocated_cluster_dim(::stream_executor::ClusterDimProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.cluster_dim_); } 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_.cluster_dim_ = reinterpret_cast<::stream_executor::ClusterDimProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.KernelThunkProto.cluster_dim) } // int64 shmem_bytes = 6; inline void KernelThunkProto::clear_shmem_bytes() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.shmem_bytes_ = ::int64_t{0}; } inline ::int64_t KernelThunkProto::shmem_bytes() const { // @@protoc_insertion_point(field_get:xla.gpu.KernelThunkProto.shmem_bytes) return _internal_shmem_bytes(); } inline void KernelThunkProto::set_shmem_bytes(::int64_t value) { _internal_set_shmem_bytes(value); // @@protoc_insertion_point(field_set:xla.gpu.KernelThunkProto.shmem_bytes) } inline ::int64_t KernelThunkProto::_internal_shmem_bytes() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.shmem_bytes_; } inline void KernelThunkProto::_internal_set_shmem_bytes(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.shmem_bytes_ = value; } // ------------------------------------------------------------------- // GemmThunkProto // .xla.GemmConfigProto gemm_config = 1; inline bool GemmThunkProto::has_gemm_config() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.gemm_config_ != nullptr); return value; } inline const ::xla::GemmConfigProto& GemmThunkProto::_internal_gemm_config() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::GemmConfigProto* p = _impl_.gemm_config_; return p != nullptr ? *p : reinterpret_cast(::xla::_GemmConfigProto_default_instance_); } inline const ::xla::GemmConfigProto& GemmThunkProto::gemm_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.gemm_config) return _internal_gemm_config(); } inline void GemmThunkProto::unsafe_arena_set_allocated_gemm_config(::xla::GemmConfigProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.gemm_config_); } _impl_.gemm_config_ = reinterpret_cast<::xla::GemmConfigProto*>(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.GemmThunkProto.gemm_config) } inline ::xla::GemmConfigProto* GemmThunkProto::release_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::GemmConfigProto* released = _impl_.gemm_config_; _impl_.gemm_config_ = nullptr; #ifdef PROTOBUF_FORCE_COPY_IN_RELEASE auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } #else // PROTOBUF_FORCE_COPY_IN_RELEASE if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } #endif // !PROTOBUF_FORCE_COPY_IN_RELEASE return released; } inline ::xla::GemmConfigProto* GemmThunkProto::unsafe_arena_release_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmThunkProto.gemm_config) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::GemmConfigProto* temp = _impl_.gemm_config_; _impl_.gemm_config_ = nullptr; return temp; } inline ::xla::GemmConfigProto* GemmThunkProto::_internal_mutable_gemm_config() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.gemm_config_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::GemmConfigProto>(GetArena()); _impl_.gemm_config_ = reinterpret_cast<::xla::GemmConfigProto*>(p); } return _impl_.gemm_config_; } inline ::xla::GemmConfigProto* GemmThunkProto::mutable_gemm_config() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::GemmConfigProto* _msg = _internal_mutable_gemm_config(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmThunkProto.gemm_config) return _msg; } inline void GemmThunkProto::set_allocated_gemm_config(::xla::GemmConfigProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.gemm_config_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.gemm_config_ = reinterpret_cast<::xla::GemmConfigProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmThunkProto.gemm_config) } // .xla.buffer_assignment.BufferAllocationSliceProto lhs_buffer = 2; inline bool GemmThunkProto::has_lhs_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.lhs_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::_internal_lhs_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.lhs_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::lhs_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.lhs_buffer) return _internal_lhs_buffer(); } inline void GemmThunkProto::unsafe_arena_set_allocated_lhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.lhs_buffer_); } _impl_.lhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.GemmThunkProto.lhs_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::release_lhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.lhs_buffer_; _impl_.lhs_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::unsafe_arena_release_lhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmThunkProto.lhs_buffer) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.lhs_buffer_; _impl_.lhs_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::_internal_mutable_lhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.lhs_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.lhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.lhs_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::mutable_lhs_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_lhs_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmThunkProto.lhs_buffer) return _msg; } inline void GemmThunkProto::set_allocated_lhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.lhs_buffer_); } 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_.lhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmThunkProto.lhs_buffer) } // .xla.buffer_assignment.BufferAllocationSliceProto rhs_buffer = 3; inline bool GemmThunkProto::has_rhs_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.rhs_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::_internal_rhs_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.rhs_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::rhs_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.rhs_buffer) return _internal_rhs_buffer(); } inline void GemmThunkProto::unsafe_arena_set_allocated_rhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.rhs_buffer_); } _impl_.rhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.GemmThunkProto.rhs_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::release_rhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.rhs_buffer_; _impl_.rhs_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::unsafe_arena_release_rhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmThunkProto.rhs_buffer) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.rhs_buffer_; _impl_.rhs_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::_internal_mutable_rhs_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.rhs_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.rhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.rhs_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::mutable_rhs_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_rhs_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmThunkProto.rhs_buffer) return _msg; } inline void GemmThunkProto::set_allocated_rhs_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.rhs_buffer_); } 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_.rhs_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmThunkProto.rhs_buffer) } // .xla.buffer_assignment.BufferAllocationSliceProto output_buffer = 4; inline bool GemmThunkProto::has_output_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; PROTOBUF_ASSUME(!value || _impl_.output_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::_internal_output_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.output_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::output_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.output_buffer) return _internal_output_buffer(); } inline void GemmThunkProto::unsafe_arena_set_allocated_output_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.output_buffer_); } _impl_.output_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmThunkProto.output_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::release_output_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.output_buffer_; _impl_.output_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::unsafe_arena_release_output_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmThunkProto.output_buffer) _impl_._has_bits_[0] &= ~0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.output_buffer_; _impl_.output_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::_internal_mutable_output_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.output_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.output_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.output_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::mutable_output_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_output_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmThunkProto.output_buffer) return _msg; } inline void GemmThunkProto::set_allocated_output_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.output_buffer_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } _impl_.output_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmThunkProto.output_buffer) } // optional .xla.buffer_assignment.BufferAllocationSliceProto workspace = 5; inline bool GemmThunkProto::has_workspace() const { bool value = (_impl_._has_bits_[0] & 0x00000010u) != 0; PROTOBUF_ASSUME(!value || _impl_.workspace_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::_internal_workspace() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.workspace_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& GemmThunkProto::workspace() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.workspace) return _internal_workspace(); } inline void GemmThunkProto::unsafe_arena_set_allocated_workspace(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.workspace_); } _impl_.workspace_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GemmThunkProto.workspace) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::release_workspace() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000010u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.workspace_; _impl_.workspace_ = 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::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::unsafe_arena_release_workspace() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.GemmThunkProto.workspace) _impl_._has_bits_[0] &= ~0x00000010u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.workspace_; _impl_.workspace_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::_internal_mutable_workspace() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.workspace_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.workspace_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.workspace_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* GemmThunkProto::mutable_workspace() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000010u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_workspace(); // @@protoc_insertion_point(field_mutable:xla.gpu.GemmThunkProto.workspace) return _msg; } inline void GemmThunkProto::set_allocated_workspace(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.workspace_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000010u; } else { _impl_._has_bits_[0] &= ~0x00000010u; } _impl_.workspace_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.GemmThunkProto.workspace) } // bool deterministic = 6; inline void GemmThunkProto::clear_deterministic() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.deterministic_ = false; } inline bool GemmThunkProto::deterministic() const { // @@protoc_insertion_point(field_get:xla.gpu.GemmThunkProto.deterministic) return _internal_deterministic(); } inline void GemmThunkProto::set_deterministic(bool value) { _internal_set_deterministic(value); // @@protoc_insertion_point(field_set:xla.gpu.GemmThunkProto.deterministic) } inline bool GemmThunkProto::_internal_deterministic() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.deterministic_; } inline void GemmThunkProto::_internal_set_deterministic(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.deterministic_ = value; } // ------------------------------------------------------------------- // WaitForStreamsThunkProto // int64 stream_id = 1; inline void WaitForStreamsThunkProto::clear_stream_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.stream_id_ = ::int64_t{0}; } inline ::int64_t WaitForStreamsThunkProto::stream_id() const { // @@protoc_insertion_point(field_get:xla.gpu.WaitForStreamsThunkProto.stream_id) return _internal_stream_id(); } inline void WaitForStreamsThunkProto::set_stream_id(::int64_t value) { _internal_set_stream_id(value); // @@protoc_insertion_point(field_set:xla.gpu.WaitForStreamsThunkProto.stream_id) } inline ::int64_t WaitForStreamsThunkProto::_internal_stream_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.stream_id_; } inline void WaitForStreamsThunkProto::_internal_set_stream_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.stream_id_ = value; } // int64 wait_for_stream_id = 2; inline void WaitForStreamsThunkProto::clear_wait_for_stream_id() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.wait_for_stream_id_ = ::int64_t{0}; } inline ::int64_t WaitForStreamsThunkProto::wait_for_stream_id() const { // @@protoc_insertion_point(field_get:xla.gpu.WaitForStreamsThunkProto.wait_for_stream_id) return _internal_wait_for_stream_id(); } inline void WaitForStreamsThunkProto::set_wait_for_stream_id(::int64_t value) { _internal_set_wait_for_stream_id(value); // @@protoc_insertion_point(field_set:xla.gpu.WaitForStreamsThunkProto.wait_for_stream_id) } inline ::int64_t WaitForStreamsThunkProto::_internal_wait_for_stream_id() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.wait_for_stream_id_; } inline void WaitForStreamsThunkProto::_internal_set_wait_for_stream_id(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.wait_for_stream_id_ = value; } // ------------------------------------------------------------------- // TriangularSolveThunkProto // .xla.TriangularSolveOptions options = 1; inline bool TriangularSolveThunkProto::has_options() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.options_ != nullptr); return value; } inline const ::xla::TriangularSolveOptions& TriangularSolveThunkProto::_internal_options() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::TriangularSolveOptions* p = _impl_.options_; return p != nullptr ? *p : reinterpret_cast(::xla::_TriangularSolveOptions_default_instance_); } inline const ::xla::TriangularSolveOptions& TriangularSolveThunkProto::options() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.options) return _internal_options(); } inline void TriangularSolveThunkProto::unsafe_arena_set_allocated_options(::xla::TriangularSolveOptions* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.options_); } _impl_.options_ = reinterpret_cast<::xla::TriangularSolveOptions*>(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.TriangularSolveThunkProto.options) } inline ::xla::TriangularSolveOptions* TriangularSolveThunkProto::release_options() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::TriangularSolveOptions* released = _impl_.options_; _impl_.options_ = 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::TriangularSolveOptions* TriangularSolveThunkProto::unsafe_arena_release_options() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.TriangularSolveThunkProto.options) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::TriangularSolveOptions* temp = _impl_.options_; _impl_.options_ = nullptr; return temp; } inline ::xla::TriangularSolveOptions* TriangularSolveThunkProto::_internal_mutable_options() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.options_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::TriangularSolveOptions>(GetArena()); _impl_.options_ = reinterpret_cast<::xla::TriangularSolveOptions*>(p); } return _impl_.options_; } inline ::xla::TriangularSolveOptions* TriangularSolveThunkProto::mutable_options() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::TriangularSolveOptions* _msg = _internal_mutable_options(); // @@protoc_insertion_point(field_mutable:xla.gpu.TriangularSolveThunkProto.options) return _msg; } inline void TriangularSolveThunkProto::set_allocated_options(::xla::TriangularSolveOptions* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.options_); } 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_.options_ = reinterpret_cast<::xla::TriangularSolveOptions*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.TriangularSolveThunkProto.options) } // .xla.buffer_assignment.BufferAllocationSliceProto a_buffer = 2; inline bool TriangularSolveThunkProto::has_a_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.a_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::_internal_a_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.a_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::a_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.a_buffer) return _internal_a_buffer(); } inline void TriangularSolveThunkProto::unsafe_arena_set_allocated_a_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.a_buffer_); } _impl_.a_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.TriangularSolveThunkProto.a_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::release_a_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.a_buffer_; _impl_.a_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::unsafe_arena_release_a_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.TriangularSolveThunkProto.a_buffer) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.a_buffer_; _impl_.a_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::_internal_mutable_a_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.a_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.a_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.a_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::mutable_a_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_a_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.TriangularSolveThunkProto.a_buffer) return _msg; } inline void TriangularSolveThunkProto::set_allocated_a_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.a_buffer_); } 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_.a_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.TriangularSolveThunkProto.a_buffer) } // .xla.buffer_assignment.BufferAllocationSliceProto b_buffer = 3; inline bool TriangularSolveThunkProto::has_b_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.b_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::_internal_b_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.b_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::b_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.b_buffer) return _internal_b_buffer(); } inline void TriangularSolveThunkProto::unsafe_arena_set_allocated_b_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.b_buffer_); } _impl_.b_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.TriangularSolveThunkProto.b_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::release_b_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.b_buffer_; _impl_.b_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::unsafe_arena_release_b_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.TriangularSolveThunkProto.b_buffer) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.b_buffer_; _impl_.b_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::_internal_mutable_b_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.b_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.b_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.b_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::mutable_b_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_b_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.TriangularSolveThunkProto.b_buffer) return _msg; } inline void TriangularSolveThunkProto::set_allocated_b_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.b_buffer_); } 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_.b_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.TriangularSolveThunkProto.b_buffer) } // .xla.buffer_assignment.BufferAllocationSliceProto temp_buffer = 4; inline bool TriangularSolveThunkProto::has_temp_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0; PROTOBUF_ASSUME(!value || _impl_.temp_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::_internal_temp_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.temp_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& TriangularSolveThunkProto::temp_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.temp_buffer) return _internal_temp_buffer(); } inline void TriangularSolveThunkProto::unsafe_arena_set_allocated_temp_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.temp_buffer_); } _impl_.temp_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); if (value != nullptr) { _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.TriangularSolveThunkProto.temp_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::release_temp_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.temp_buffer_; _impl_.temp_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::unsafe_arena_release_temp_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.TriangularSolveThunkProto.temp_buffer) _impl_._has_bits_[0] &= ~0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.temp_buffer_; _impl_.temp_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::_internal_mutable_temp_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.temp_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.temp_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.temp_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* TriangularSolveThunkProto::mutable_temp_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000008u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_temp_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.TriangularSolveThunkProto.temp_buffer) return _msg; } inline void TriangularSolveThunkProto::set_allocated_temp_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.temp_buffer_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000008u; } else { _impl_._has_bits_[0] &= ~0x00000008u; } _impl_.temp_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.TriangularSolveThunkProto.temp_buffer) } // .xla.PrimitiveType type = 5; inline void TriangularSolveThunkProto::clear_type() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.type_ = 0; } inline ::xla::PrimitiveType TriangularSolveThunkProto::type() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.type) return _internal_type(); } inline void TriangularSolveThunkProto::set_type(::xla::PrimitiveType value) { _internal_set_type(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.type) } inline ::xla::PrimitiveType TriangularSolveThunkProto::_internal_type() const { ::google::protobuf::internal::TSanRead(&_impl_); return static_cast<::xla::PrimitiveType>(_impl_.type_); } inline void TriangularSolveThunkProto::_internal_set_type(::xla::PrimitiveType value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.type_ = value; } // int64 batch_size = 6; inline void TriangularSolveThunkProto::clear_batch_size() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.batch_size_ = ::int64_t{0}; } inline ::int64_t TriangularSolveThunkProto::batch_size() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.batch_size) return _internal_batch_size(); } inline void TriangularSolveThunkProto::set_batch_size(::int64_t value) { _internal_set_batch_size(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.batch_size) } inline ::int64_t TriangularSolveThunkProto::_internal_batch_size() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.batch_size_; } inline void TriangularSolveThunkProto::_internal_set_batch_size(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.batch_size_ = value; } // int64 m = 7; inline void TriangularSolveThunkProto::clear_m() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.m_ = ::int64_t{0}; } inline ::int64_t TriangularSolveThunkProto::m() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.m) return _internal_m(); } inline void TriangularSolveThunkProto::set_m(::int64_t value) { _internal_set_m(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.m) } inline ::int64_t TriangularSolveThunkProto::_internal_m() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.m_; } inline void TriangularSolveThunkProto::_internal_set_m(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.m_ = value; } // int64 n = 8; inline void TriangularSolveThunkProto::clear_n() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.n_ = ::int64_t{0}; } inline ::int64_t TriangularSolveThunkProto::n() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.n) return _internal_n(); } inline void TriangularSolveThunkProto::set_n(::int64_t value) { _internal_set_n(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.n) } inline ::int64_t TriangularSolveThunkProto::_internal_n() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.n_; } inline void TriangularSolveThunkProto::_internal_set_n(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.n_ = value; } // int64 a_batch_stride = 9; inline void TriangularSolveThunkProto::clear_a_batch_stride() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.a_batch_stride_ = ::int64_t{0}; } inline ::int64_t TriangularSolveThunkProto::a_batch_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.a_batch_stride) return _internal_a_batch_stride(); } inline void TriangularSolveThunkProto::set_a_batch_stride(::int64_t value) { _internal_set_a_batch_stride(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.a_batch_stride) } inline ::int64_t TriangularSolveThunkProto::_internal_a_batch_stride() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.a_batch_stride_; } inline void TriangularSolveThunkProto::_internal_set_a_batch_stride(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.a_batch_stride_ = value; } // int64 b_batch_stride = 10; inline void TriangularSolveThunkProto::clear_b_batch_stride() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.b_batch_stride_ = ::int64_t{0}; } inline ::int64_t TriangularSolveThunkProto::b_batch_stride() const { // @@protoc_insertion_point(field_get:xla.gpu.TriangularSolveThunkProto.b_batch_stride) return _internal_b_batch_stride(); } inline void TriangularSolveThunkProto::set_b_batch_stride(::int64_t value) { _internal_set_b_batch_stride(value); // @@protoc_insertion_point(field_set:xla.gpu.TriangularSolveThunkProto.b_batch_stride) } inline ::int64_t TriangularSolveThunkProto::_internal_b_batch_stride() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.b_batch_stride_; } inline void TriangularSolveThunkProto::_internal_set_b_batch_stride(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.b_batch_stride_ = value; } // ------------------------------------------------------------------- // ReplicaIdThunkProto // .xla.buffer_assignment.BufferAllocationSliceProto dest_buffer = 1; inline bool ReplicaIdThunkProto::has_dest_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.dest_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& ReplicaIdThunkProto::_internal_dest_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.dest_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& ReplicaIdThunkProto::dest_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ReplicaIdThunkProto.dest_buffer) return _internal_dest_buffer(); } inline void ReplicaIdThunkProto::unsafe_arena_set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dest_buffer_); } _impl_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.ReplicaIdThunkProto.dest_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ReplicaIdThunkProto::release_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.dest_buffer_; _impl_.dest_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* ReplicaIdThunkProto::unsafe_arena_release_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ReplicaIdThunkProto.dest_buffer) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.dest_buffer_; _impl_.dest_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ReplicaIdThunkProto::_internal_mutable_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.dest_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.dest_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* ReplicaIdThunkProto::mutable_dest_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_dest_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.ReplicaIdThunkProto.dest_buffer) return _msg; } inline void ReplicaIdThunkProto::set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dest_buffer_); } 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_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.ReplicaIdThunkProto.dest_buffer) } // ------------------------------------------------------------------- // PartitionIdThunkProto // .xla.buffer_assignment.BufferAllocationSliceProto dest_buffer = 1; inline bool PartitionIdThunkProto::has_dest_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.dest_buffer_ != nullptr); return value; } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& PartitionIdThunkProto::_internal_dest_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::buffer_assignment::BufferAllocationSliceProto* p = _impl_.dest_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::buffer_assignment::_BufferAllocationSliceProto_default_instance_); } inline const ::xla::buffer_assignment::BufferAllocationSliceProto& PartitionIdThunkProto::dest_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.PartitionIdThunkProto.dest_buffer) return _internal_dest_buffer(); } inline void PartitionIdThunkProto::unsafe_arena_set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dest_buffer_); } _impl_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(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.PartitionIdThunkProto.dest_buffer) } inline ::xla::buffer_assignment::BufferAllocationSliceProto* PartitionIdThunkProto::release_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* released = _impl_.dest_buffer_; _impl_.dest_buffer_ = 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::buffer_assignment::BufferAllocationSliceProto* PartitionIdThunkProto::unsafe_arena_release_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.PartitionIdThunkProto.dest_buffer) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* temp = _impl_.dest_buffer_; _impl_.dest_buffer_ = nullptr; return temp; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* PartitionIdThunkProto::_internal_mutable_dest_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.dest_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::buffer_assignment::BufferAllocationSliceProto>(GetArena()); _impl_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(p); } return _impl_.dest_buffer_; } inline ::xla::buffer_assignment::BufferAllocationSliceProto* PartitionIdThunkProto::mutable_dest_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::buffer_assignment::BufferAllocationSliceProto* _msg = _internal_mutable_dest_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.PartitionIdThunkProto.dest_buffer) return _msg; } inline void PartitionIdThunkProto::set_allocated_dest_buffer(::xla::buffer_assignment::BufferAllocationSliceProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dest_buffer_); } 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_.dest_buffer_ = reinterpret_cast<::xla::buffer_assignment::BufferAllocationSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.PartitionIdThunkProto.dest_buffer) } // ------------------------------------------------------------------- // ThunkProto // .xla.gpu.ThunkInfoProto thunk_info = 1; inline bool ThunkProto::has_thunk_info() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.thunk_info_ != nullptr); return value; } inline void ThunkProto::clear_thunk_info() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.thunk_info_ != nullptr) _impl_.thunk_info_->Clear(); _impl_._has_bits_[0] &= ~0x00000001u; } inline const ::xla::gpu::ThunkInfoProto& ThunkProto::_internal_thunk_info() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::ThunkInfoProto* p = _impl_.thunk_info_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_ThunkInfoProto_default_instance_); } inline const ::xla::gpu::ThunkInfoProto& ThunkProto::thunk_info() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.thunk_info) return _internal_thunk_info(); } inline void ThunkProto::unsafe_arena_set_allocated_thunk_info(::xla::gpu::ThunkInfoProto* value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.thunk_info_); } _impl_.thunk_info_ = reinterpret_cast<::xla::gpu::ThunkInfoProto*>(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.ThunkProto.thunk_info) } inline ::xla::gpu::ThunkInfoProto* ThunkProto::release_thunk_info() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::ThunkInfoProto* released = _impl_.thunk_info_; _impl_.thunk_info_ = 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::ThunkInfoProto* ThunkProto::unsafe_arena_release_thunk_info() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.thunk_info) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::gpu::ThunkInfoProto* temp = _impl_.thunk_info_; _impl_.thunk_info_ = nullptr; return temp; } inline ::xla::gpu::ThunkInfoProto* ThunkProto::_internal_mutable_thunk_info() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.thunk_info_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::ThunkInfoProto>(GetArena()); _impl_.thunk_info_ = reinterpret_cast<::xla::gpu::ThunkInfoProto*>(p); } return _impl_.thunk_info_; } inline ::xla::gpu::ThunkInfoProto* ThunkProto::mutable_thunk_info() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::gpu::ThunkInfoProto* _msg = _internal_mutable_thunk_info(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.thunk_info) return _msg; } inline void ThunkProto::set_allocated_thunk_info(::xla::gpu::ThunkInfoProto* value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete (_impl_.thunk_info_); } if (value != nullptr) { ::google::protobuf::Arena* submessage_arena = (value)->GetArena(); if (message_arena != submessage_arena) { value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena); } _impl_._has_bits_[0] |= 0x00000001u; } else { _impl_._has_bits_[0] &= ~0x00000001u; } _impl_.thunk_info_ = reinterpret_cast<::xla::gpu::ThunkInfoProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.ThunkProto.thunk_info) } // .xla.gpu.SequentialThunkProto sequential_thunk = 2; inline bool ThunkProto::has_sequential_thunk() const { return impl_case() == kSequentialThunk; } inline bool ThunkProto::_internal_has_sequential_thunk() const { return impl_case() == kSequentialThunk; } inline void ThunkProto::set_has_sequential_thunk() { _impl_._oneof_case_[0] = kSequentialThunk; } inline void ThunkProto::clear_sequential_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kSequentialThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.sequential_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.sequential_thunk_); } clear_has_impl(); } } inline ::xla::gpu::SequentialThunkProto* ThunkProto::release_sequential_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.sequential_thunk) if (impl_case() == kSequentialThunk) { clear_has_impl(); auto* temp = _impl_.impl_.sequential_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.sequential_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::SequentialThunkProto& ThunkProto::_internal_sequential_thunk() const { return impl_case() == kSequentialThunk ? *_impl_.impl_.sequential_thunk_ : reinterpret_cast<::xla::gpu::SequentialThunkProto&>(::xla::gpu::_SequentialThunkProto_default_instance_); } inline const ::xla::gpu::SequentialThunkProto& ThunkProto::sequential_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.sequential_thunk) return _internal_sequential_thunk(); } inline ::xla::gpu::SequentialThunkProto* ThunkProto::unsafe_arena_release_sequential_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.sequential_thunk) if (impl_case() == kSequentialThunk) { clear_has_impl(); auto* temp = _impl_.impl_.sequential_thunk_; _impl_.impl_.sequential_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_sequential_thunk(::xla::gpu::SequentialThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_sequential_thunk(); _impl_.impl_.sequential_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.sequential_thunk) } inline ::xla::gpu::SequentialThunkProto* ThunkProto::_internal_mutable_sequential_thunk() { if (impl_case() != kSequentialThunk) { clear_impl(); set_has_sequential_thunk(); _impl_.impl_.sequential_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::SequentialThunkProto>(GetArena()); } return _impl_.impl_.sequential_thunk_; } inline ::xla::gpu::SequentialThunkProto* ThunkProto::mutable_sequential_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::SequentialThunkProto* _msg = _internal_mutable_sequential_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.sequential_thunk) return _msg; } // .xla.gpu.CopyThunkProto copy_thunk = 3; inline bool ThunkProto::has_copy_thunk() const { return impl_case() == kCopyThunk; } inline bool ThunkProto::_internal_has_copy_thunk() const { return impl_case() == kCopyThunk; } inline void ThunkProto::set_has_copy_thunk() { _impl_._oneof_case_[0] = kCopyThunk; } inline void ThunkProto::clear_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kCopyThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.copy_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.copy_thunk_); } clear_has_impl(); } } inline ::xla::gpu::CopyThunkProto* ThunkProto::release_copy_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.copy_thunk) if (impl_case() == kCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.copy_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::CopyThunkProto& ThunkProto::_internal_copy_thunk() const { return impl_case() == kCopyThunk ? *_impl_.impl_.copy_thunk_ : reinterpret_cast<::xla::gpu::CopyThunkProto&>(::xla::gpu::_CopyThunkProto_default_instance_); } inline const ::xla::gpu::CopyThunkProto& ThunkProto::copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.copy_thunk) return _internal_copy_thunk(); } inline ::xla::gpu::CopyThunkProto* ThunkProto::unsafe_arena_release_copy_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.copy_thunk) if (impl_case() == kCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.copy_thunk_; _impl_.impl_.copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_copy_thunk(); _impl_.impl_.copy_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.copy_thunk) } inline ::xla::gpu::CopyThunkProto* ThunkProto::_internal_mutable_copy_thunk() { if (impl_case() != kCopyThunk) { clear_impl(); set_has_copy_thunk(); _impl_.impl_.copy_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CopyThunkProto>(GetArena()); } return _impl_.impl_.copy_thunk_; } inline ::xla::gpu::CopyThunkProto* ThunkProto::mutable_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::CopyThunkProto* _msg = _internal_mutable_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.copy_thunk) return _msg; } // .xla.gpu.DeviceToHostCopyThunkProto device_to_host_copy_thunk = 4; inline bool ThunkProto::has_device_to_host_copy_thunk() const { return impl_case() == kDeviceToHostCopyThunk; } inline bool ThunkProto::_internal_has_device_to_host_copy_thunk() const { return impl_case() == kDeviceToHostCopyThunk; } inline void ThunkProto::set_has_device_to_host_copy_thunk() { _impl_._oneof_case_[0] = kDeviceToHostCopyThunk; } inline void ThunkProto::clear_device_to_host_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kDeviceToHostCopyThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.device_to_host_copy_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.device_to_host_copy_thunk_); } clear_has_impl(); } } inline ::xla::gpu::DeviceToHostCopyThunkProto* ThunkProto::release_device_to_host_copy_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.device_to_host_copy_thunk) if (impl_case() == kDeviceToHostCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.device_to_host_copy_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.device_to_host_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::DeviceToHostCopyThunkProto& ThunkProto::_internal_device_to_host_copy_thunk() const { return impl_case() == kDeviceToHostCopyThunk ? *_impl_.impl_.device_to_host_copy_thunk_ : reinterpret_cast<::xla::gpu::DeviceToHostCopyThunkProto&>(::xla::gpu::_DeviceToHostCopyThunkProto_default_instance_); } inline const ::xla::gpu::DeviceToHostCopyThunkProto& ThunkProto::device_to_host_copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.device_to_host_copy_thunk) return _internal_device_to_host_copy_thunk(); } inline ::xla::gpu::DeviceToHostCopyThunkProto* ThunkProto::unsafe_arena_release_device_to_host_copy_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.device_to_host_copy_thunk) if (impl_case() == kDeviceToHostCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.device_to_host_copy_thunk_; _impl_.impl_.device_to_host_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_device_to_host_copy_thunk(::xla::gpu::DeviceToHostCopyThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_device_to_host_copy_thunk(); _impl_.impl_.device_to_host_copy_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.device_to_host_copy_thunk) } inline ::xla::gpu::DeviceToHostCopyThunkProto* ThunkProto::_internal_mutable_device_to_host_copy_thunk() { if (impl_case() != kDeviceToHostCopyThunk) { clear_impl(); set_has_device_to_host_copy_thunk(); _impl_.impl_.device_to_host_copy_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::DeviceToHostCopyThunkProto>(GetArena()); } return _impl_.impl_.device_to_host_copy_thunk_; } inline ::xla::gpu::DeviceToHostCopyThunkProto* ThunkProto::mutable_device_to_host_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::DeviceToHostCopyThunkProto* _msg = _internal_mutable_device_to_host_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.device_to_host_copy_thunk) return _msg; } // .xla.gpu.HostToDeviceCopyThunkProto host_to_device_copy_thunk = 5; inline bool ThunkProto::has_host_to_device_copy_thunk() const { return impl_case() == kHostToDeviceCopyThunk; } inline bool ThunkProto::_internal_has_host_to_device_copy_thunk() const { return impl_case() == kHostToDeviceCopyThunk; } inline void ThunkProto::set_has_host_to_device_copy_thunk() { _impl_._oneof_case_[0] = kHostToDeviceCopyThunk; } inline void ThunkProto::clear_host_to_device_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kHostToDeviceCopyThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.host_to_device_copy_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.host_to_device_copy_thunk_); } clear_has_impl(); } } inline ::xla::gpu::HostToDeviceCopyThunkProto* ThunkProto::release_host_to_device_copy_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.host_to_device_copy_thunk) if (impl_case() == kHostToDeviceCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.host_to_device_copy_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.host_to_device_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::HostToDeviceCopyThunkProto& ThunkProto::_internal_host_to_device_copy_thunk() const { return impl_case() == kHostToDeviceCopyThunk ? *_impl_.impl_.host_to_device_copy_thunk_ : reinterpret_cast<::xla::gpu::HostToDeviceCopyThunkProto&>(::xla::gpu::_HostToDeviceCopyThunkProto_default_instance_); } inline const ::xla::gpu::HostToDeviceCopyThunkProto& ThunkProto::host_to_device_copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.host_to_device_copy_thunk) return _internal_host_to_device_copy_thunk(); } inline ::xla::gpu::HostToDeviceCopyThunkProto* ThunkProto::unsafe_arena_release_host_to_device_copy_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.host_to_device_copy_thunk) if (impl_case() == kHostToDeviceCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.host_to_device_copy_thunk_; _impl_.impl_.host_to_device_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_host_to_device_copy_thunk(::xla::gpu::HostToDeviceCopyThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_host_to_device_copy_thunk(); _impl_.impl_.host_to_device_copy_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.host_to_device_copy_thunk) } inline ::xla::gpu::HostToDeviceCopyThunkProto* ThunkProto::_internal_mutable_host_to_device_copy_thunk() { if (impl_case() != kHostToDeviceCopyThunk) { clear_impl(); set_has_host_to_device_copy_thunk(); _impl_.impl_.host_to_device_copy_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::HostToDeviceCopyThunkProto>(GetArena()); } return _impl_.impl_.host_to_device_copy_thunk_; } inline ::xla::gpu::HostToDeviceCopyThunkProto* ThunkProto::mutable_host_to_device_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::HostToDeviceCopyThunkProto* _msg = _internal_mutable_host_to_device_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.host_to_device_copy_thunk) return _msg; } // .xla.gpu.ConditionalThunkProto conditional_thunk = 6; inline bool ThunkProto::has_conditional_thunk() const { return impl_case() == kConditionalThunk; } inline bool ThunkProto::_internal_has_conditional_thunk() const { return impl_case() == kConditionalThunk; } inline void ThunkProto::set_has_conditional_thunk() { _impl_._oneof_case_[0] = kConditionalThunk; } inline void ThunkProto::clear_conditional_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kConditionalThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.conditional_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.conditional_thunk_); } clear_has_impl(); } } inline ::xla::gpu::ConditionalThunkProto* ThunkProto::release_conditional_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.conditional_thunk) if (impl_case() == kConditionalThunk) { clear_has_impl(); auto* temp = _impl_.impl_.conditional_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.conditional_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::ConditionalThunkProto& ThunkProto::_internal_conditional_thunk() const { return impl_case() == kConditionalThunk ? *_impl_.impl_.conditional_thunk_ : reinterpret_cast<::xla::gpu::ConditionalThunkProto&>(::xla::gpu::_ConditionalThunkProto_default_instance_); } inline const ::xla::gpu::ConditionalThunkProto& ThunkProto::conditional_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.conditional_thunk) return _internal_conditional_thunk(); } inline ::xla::gpu::ConditionalThunkProto* ThunkProto::unsafe_arena_release_conditional_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.conditional_thunk) if (impl_case() == kConditionalThunk) { clear_has_impl(); auto* temp = _impl_.impl_.conditional_thunk_; _impl_.impl_.conditional_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_conditional_thunk(::xla::gpu::ConditionalThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_conditional_thunk(); _impl_.impl_.conditional_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.conditional_thunk) } inline ::xla::gpu::ConditionalThunkProto* ThunkProto::_internal_mutable_conditional_thunk() { if (impl_case() != kConditionalThunk) { clear_impl(); set_has_conditional_thunk(); _impl_.impl_.conditional_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::ConditionalThunkProto>(GetArena()); } return _impl_.impl_.conditional_thunk_; } inline ::xla::gpu::ConditionalThunkProto* ThunkProto::mutable_conditional_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::ConditionalThunkProto* _msg = _internal_mutable_conditional_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.conditional_thunk) return _msg; } // .xla.gpu.WhileThunkProto while_thunk = 7; inline bool ThunkProto::has_while_thunk() const { return impl_case() == kWhileThunk; } inline bool ThunkProto::_internal_has_while_thunk() const { return impl_case() == kWhileThunk; } inline void ThunkProto::set_has_while_thunk() { _impl_._oneof_case_[0] = kWhileThunk; } inline void ThunkProto::clear_while_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kWhileThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.while_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.while_thunk_); } clear_has_impl(); } } inline ::xla::gpu::WhileThunkProto* ThunkProto::release_while_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.while_thunk) if (impl_case() == kWhileThunk) { clear_has_impl(); auto* temp = _impl_.impl_.while_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.while_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::WhileThunkProto& ThunkProto::_internal_while_thunk() const { return impl_case() == kWhileThunk ? *_impl_.impl_.while_thunk_ : reinterpret_cast<::xla::gpu::WhileThunkProto&>(::xla::gpu::_WhileThunkProto_default_instance_); } inline const ::xla::gpu::WhileThunkProto& ThunkProto::while_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.while_thunk) return _internal_while_thunk(); } inline ::xla::gpu::WhileThunkProto* ThunkProto::unsafe_arena_release_while_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.while_thunk) if (impl_case() == kWhileThunk) { clear_has_impl(); auto* temp = _impl_.impl_.while_thunk_; _impl_.impl_.while_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_while_thunk(::xla::gpu::WhileThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_while_thunk(); _impl_.impl_.while_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.while_thunk) } inline ::xla::gpu::WhileThunkProto* ThunkProto::_internal_mutable_while_thunk() { if (impl_case() != kWhileThunk) { clear_impl(); set_has_while_thunk(); _impl_.impl_.while_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::WhileThunkProto>(GetArena()); } return _impl_.impl_.while_thunk_; } inline ::xla::gpu::WhileThunkProto* ThunkProto::mutable_while_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::WhileThunkProto* _msg = _internal_mutable_while_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.while_thunk) return _msg; } // .xla.gpu.KernelThunkProto kernel_thunk = 8; inline bool ThunkProto::has_kernel_thunk() const { return impl_case() == kKernelThunk; } inline bool ThunkProto::_internal_has_kernel_thunk() const { return impl_case() == kKernelThunk; } inline void ThunkProto::set_has_kernel_thunk() { _impl_._oneof_case_[0] = kKernelThunk; } inline void ThunkProto::clear_kernel_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kKernelThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.kernel_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.kernel_thunk_); } clear_has_impl(); } } inline ::xla::gpu::KernelThunkProto* ThunkProto::release_kernel_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.kernel_thunk) if (impl_case() == kKernelThunk) { clear_has_impl(); auto* temp = _impl_.impl_.kernel_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.kernel_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::KernelThunkProto& ThunkProto::_internal_kernel_thunk() const { return impl_case() == kKernelThunk ? *_impl_.impl_.kernel_thunk_ : reinterpret_cast<::xla::gpu::KernelThunkProto&>(::xla::gpu::_KernelThunkProto_default_instance_); } inline const ::xla::gpu::KernelThunkProto& ThunkProto::kernel_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.kernel_thunk) return _internal_kernel_thunk(); } inline ::xla::gpu::KernelThunkProto* ThunkProto::unsafe_arena_release_kernel_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.kernel_thunk) if (impl_case() == kKernelThunk) { clear_has_impl(); auto* temp = _impl_.impl_.kernel_thunk_; _impl_.impl_.kernel_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_kernel_thunk(::xla::gpu::KernelThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_kernel_thunk(); _impl_.impl_.kernel_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.kernel_thunk) } inline ::xla::gpu::KernelThunkProto* ThunkProto::_internal_mutable_kernel_thunk() { if (impl_case() != kKernelThunk) { clear_impl(); set_has_kernel_thunk(); _impl_.impl_.kernel_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::KernelThunkProto>(GetArena()); } return _impl_.impl_.kernel_thunk_; } inline ::xla::gpu::KernelThunkProto* ThunkProto::mutable_kernel_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::KernelThunkProto* _msg = _internal_mutable_kernel_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.kernel_thunk) return _msg; } // .xla.gpu.GemmThunkProto gemm_thunk = 9; inline bool ThunkProto::has_gemm_thunk() const { return impl_case() == kGemmThunk; } inline bool ThunkProto::_internal_has_gemm_thunk() const { return impl_case() == kGemmThunk; } inline void ThunkProto::set_has_gemm_thunk() { _impl_._oneof_case_[0] = kGemmThunk; } inline void ThunkProto::clear_gemm_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kGemmThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.gemm_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.gemm_thunk_); } clear_has_impl(); } } inline ::xla::gpu::GemmThunkProto* ThunkProto::release_gemm_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.gemm_thunk) if (impl_case() == kGemmThunk) { clear_has_impl(); auto* temp = _impl_.impl_.gemm_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.gemm_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::GemmThunkProto& ThunkProto::_internal_gemm_thunk() const { return impl_case() == kGemmThunk ? *_impl_.impl_.gemm_thunk_ : reinterpret_cast<::xla::gpu::GemmThunkProto&>(::xla::gpu::_GemmThunkProto_default_instance_); } inline const ::xla::gpu::GemmThunkProto& ThunkProto::gemm_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.gemm_thunk) return _internal_gemm_thunk(); } inline ::xla::gpu::GemmThunkProto* ThunkProto::unsafe_arena_release_gemm_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.gemm_thunk) if (impl_case() == kGemmThunk) { clear_has_impl(); auto* temp = _impl_.impl_.gemm_thunk_; _impl_.impl_.gemm_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_gemm_thunk(::xla::gpu::GemmThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_gemm_thunk(); _impl_.impl_.gemm_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.gemm_thunk) } inline ::xla::gpu::GemmThunkProto* ThunkProto::_internal_mutable_gemm_thunk() { if (impl_case() != kGemmThunk) { clear_impl(); set_has_gemm_thunk(); _impl_.impl_.gemm_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::GemmThunkProto>(GetArena()); } return _impl_.impl_.gemm_thunk_; } inline ::xla::gpu::GemmThunkProto* ThunkProto::mutable_gemm_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::GemmThunkProto* _msg = _internal_mutable_gemm_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.gemm_thunk) return _msg; } // .xla.gpu.WaitForStreamsThunkProto wait_for_streams_thunk = 10; inline bool ThunkProto::has_wait_for_streams_thunk() const { return impl_case() == kWaitForStreamsThunk; } inline bool ThunkProto::_internal_has_wait_for_streams_thunk() const { return impl_case() == kWaitForStreamsThunk; } inline void ThunkProto::set_has_wait_for_streams_thunk() { _impl_._oneof_case_[0] = kWaitForStreamsThunk; } inline void ThunkProto::clear_wait_for_streams_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kWaitForStreamsThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.wait_for_streams_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.wait_for_streams_thunk_); } clear_has_impl(); } } inline ::xla::gpu::WaitForStreamsThunkProto* ThunkProto::release_wait_for_streams_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.wait_for_streams_thunk) if (impl_case() == kWaitForStreamsThunk) { clear_has_impl(); auto* temp = _impl_.impl_.wait_for_streams_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.wait_for_streams_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::WaitForStreamsThunkProto& ThunkProto::_internal_wait_for_streams_thunk() const { return impl_case() == kWaitForStreamsThunk ? *_impl_.impl_.wait_for_streams_thunk_ : reinterpret_cast<::xla::gpu::WaitForStreamsThunkProto&>(::xla::gpu::_WaitForStreamsThunkProto_default_instance_); } inline const ::xla::gpu::WaitForStreamsThunkProto& ThunkProto::wait_for_streams_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.wait_for_streams_thunk) return _internal_wait_for_streams_thunk(); } inline ::xla::gpu::WaitForStreamsThunkProto* ThunkProto::unsafe_arena_release_wait_for_streams_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.wait_for_streams_thunk) if (impl_case() == kWaitForStreamsThunk) { clear_has_impl(); auto* temp = _impl_.impl_.wait_for_streams_thunk_; _impl_.impl_.wait_for_streams_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_wait_for_streams_thunk(::xla::gpu::WaitForStreamsThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_wait_for_streams_thunk(); _impl_.impl_.wait_for_streams_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.wait_for_streams_thunk) } inline ::xla::gpu::WaitForStreamsThunkProto* ThunkProto::_internal_mutable_wait_for_streams_thunk() { if (impl_case() != kWaitForStreamsThunk) { clear_impl(); set_has_wait_for_streams_thunk(); _impl_.impl_.wait_for_streams_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::WaitForStreamsThunkProto>(GetArena()); } return _impl_.impl_.wait_for_streams_thunk_; } inline ::xla::gpu::WaitForStreamsThunkProto* ThunkProto::mutable_wait_for_streams_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::WaitForStreamsThunkProto* _msg = _internal_mutable_wait_for_streams_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.wait_for_streams_thunk) return _msg; } // .xla.gpu.TriangularSolveThunkProto triangular_solve_thunk = 11; inline bool ThunkProto::has_triangular_solve_thunk() const { return impl_case() == kTriangularSolveThunk; } inline bool ThunkProto::_internal_has_triangular_solve_thunk() const { return impl_case() == kTriangularSolveThunk; } inline void ThunkProto::set_has_triangular_solve_thunk() { _impl_._oneof_case_[0] = kTriangularSolveThunk; } inline void ThunkProto::clear_triangular_solve_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kTriangularSolveThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.triangular_solve_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.triangular_solve_thunk_); } clear_has_impl(); } } inline ::xla::gpu::TriangularSolveThunkProto* ThunkProto::release_triangular_solve_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.triangular_solve_thunk) if (impl_case() == kTriangularSolveThunk) { clear_has_impl(); auto* temp = _impl_.impl_.triangular_solve_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.triangular_solve_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::TriangularSolveThunkProto& ThunkProto::_internal_triangular_solve_thunk() const { return impl_case() == kTriangularSolveThunk ? *_impl_.impl_.triangular_solve_thunk_ : reinterpret_cast<::xla::gpu::TriangularSolveThunkProto&>(::xla::gpu::_TriangularSolveThunkProto_default_instance_); } inline const ::xla::gpu::TriangularSolveThunkProto& ThunkProto::triangular_solve_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.triangular_solve_thunk) return _internal_triangular_solve_thunk(); } inline ::xla::gpu::TriangularSolveThunkProto* ThunkProto::unsafe_arena_release_triangular_solve_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.triangular_solve_thunk) if (impl_case() == kTriangularSolveThunk) { clear_has_impl(); auto* temp = _impl_.impl_.triangular_solve_thunk_; _impl_.impl_.triangular_solve_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_triangular_solve_thunk(::xla::gpu::TriangularSolveThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_triangular_solve_thunk(); _impl_.impl_.triangular_solve_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.triangular_solve_thunk) } inline ::xla::gpu::TriangularSolveThunkProto* ThunkProto::_internal_mutable_triangular_solve_thunk() { if (impl_case() != kTriangularSolveThunk) { clear_impl(); set_has_triangular_solve_thunk(); _impl_.impl_.triangular_solve_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::TriangularSolveThunkProto>(GetArena()); } return _impl_.impl_.triangular_solve_thunk_; } inline ::xla::gpu::TriangularSolveThunkProto* ThunkProto::mutable_triangular_solve_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::TriangularSolveThunkProto* _msg = _internal_mutable_triangular_solve_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.triangular_solve_thunk) return _msg; } // .xla.gpu.ReplicaIdThunkProto replica_id_thunk = 12; inline bool ThunkProto::has_replica_id_thunk() const { return impl_case() == kReplicaIdThunk; } inline bool ThunkProto::_internal_has_replica_id_thunk() const { return impl_case() == kReplicaIdThunk; } inline void ThunkProto::set_has_replica_id_thunk() { _impl_._oneof_case_[0] = kReplicaIdThunk; } inline void ThunkProto::clear_replica_id_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kReplicaIdThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.replica_id_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.replica_id_thunk_); } clear_has_impl(); } } inline ::xla::gpu::ReplicaIdThunkProto* ThunkProto::release_replica_id_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.replica_id_thunk) if (impl_case() == kReplicaIdThunk) { clear_has_impl(); auto* temp = _impl_.impl_.replica_id_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.replica_id_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::ReplicaIdThunkProto& ThunkProto::_internal_replica_id_thunk() const { return impl_case() == kReplicaIdThunk ? *_impl_.impl_.replica_id_thunk_ : reinterpret_cast<::xla::gpu::ReplicaIdThunkProto&>(::xla::gpu::_ReplicaIdThunkProto_default_instance_); } inline const ::xla::gpu::ReplicaIdThunkProto& ThunkProto::replica_id_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.replica_id_thunk) return _internal_replica_id_thunk(); } inline ::xla::gpu::ReplicaIdThunkProto* ThunkProto::unsafe_arena_release_replica_id_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.replica_id_thunk) if (impl_case() == kReplicaIdThunk) { clear_has_impl(); auto* temp = _impl_.impl_.replica_id_thunk_; _impl_.impl_.replica_id_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_replica_id_thunk(::xla::gpu::ReplicaIdThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_replica_id_thunk(); _impl_.impl_.replica_id_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.replica_id_thunk) } inline ::xla::gpu::ReplicaIdThunkProto* ThunkProto::_internal_mutable_replica_id_thunk() { if (impl_case() != kReplicaIdThunk) { clear_impl(); set_has_replica_id_thunk(); _impl_.impl_.replica_id_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::ReplicaIdThunkProto>(GetArena()); } return _impl_.impl_.replica_id_thunk_; } inline ::xla::gpu::ReplicaIdThunkProto* ThunkProto::mutable_replica_id_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::ReplicaIdThunkProto* _msg = _internal_mutable_replica_id_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.replica_id_thunk) return _msg; } // .xla.gpu.DeviceToDeviceCopyThunkProto device_to_device_copy_thunk = 13; inline bool ThunkProto::has_device_to_device_copy_thunk() const { return impl_case() == kDeviceToDeviceCopyThunk; } inline bool ThunkProto::_internal_has_device_to_device_copy_thunk() const { return impl_case() == kDeviceToDeviceCopyThunk; } inline void ThunkProto::set_has_device_to_device_copy_thunk() { _impl_._oneof_case_[0] = kDeviceToDeviceCopyThunk; } inline void ThunkProto::clear_device_to_device_copy_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kDeviceToDeviceCopyThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.device_to_device_copy_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.device_to_device_copy_thunk_); } clear_has_impl(); } } inline ::xla::gpu::DeviceToDeviceCopyThunkProto* ThunkProto::release_device_to_device_copy_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.device_to_device_copy_thunk) if (impl_case() == kDeviceToDeviceCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.device_to_device_copy_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.device_to_device_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::DeviceToDeviceCopyThunkProto& ThunkProto::_internal_device_to_device_copy_thunk() const { return impl_case() == kDeviceToDeviceCopyThunk ? *_impl_.impl_.device_to_device_copy_thunk_ : reinterpret_cast<::xla::gpu::DeviceToDeviceCopyThunkProto&>(::xla::gpu::_DeviceToDeviceCopyThunkProto_default_instance_); } inline const ::xla::gpu::DeviceToDeviceCopyThunkProto& ThunkProto::device_to_device_copy_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.device_to_device_copy_thunk) return _internal_device_to_device_copy_thunk(); } inline ::xla::gpu::DeviceToDeviceCopyThunkProto* ThunkProto::unsafe_arena_release_device_to_device_copy_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.device_to_device_copy_thunk) if (impl_case() == kDeviceToDeviceCopyThunk) { clear_has_impl(); auto* temp = _impl_.impl_.device_to_device_copy_thunk_; _impl_.impl_.device_to_device_copy_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_device_to_device_copy_thunk(::xla::gpu::DeviceToDeviceCopyThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_device_to_device_copy_thunk(); _impl_.impl_.device_to_device_copy_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.device_to_device_copy_thunk) } inline ::xla::gpu::DeviceToDeviceCopyThunkProto* ThunkProto::_internal_mutable_device_to_device_copy_thunk() { if (impl_case() != kDeviceToDeviceCopyThunk) { clear_impl(); set_has_device_to_device_copy_thunk(); _impl_.impl_.device_to_device_copy_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::DeviceToDeviceCopyThunkProto>(GetArena()); } return _impl_.impl_.device_to_device_copy_thunk_; } inline ::xla::gpu::DeviceToDeviceCopyThunkProto* ThunkProto::mutable_device_to_device_copy_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::DeviceToDeviceCopyThunkProto* _msg = _internal_mutable_device_to_device_copy_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.device_to_device_copy_thunk) return _msg; } // .xla.gpu.PartitionIdThunkProto partition_id_thunk = 14; inline bool ThunkProto::has_partition_id_thunk() const { return impl_case() == kPartitionIdThunk; } inline bool ThunkProto::_internal_has_partition_id_thunk() const { return impl_case() == kPartitionIdThunk; } inline void ThunkProto::set_has_partition_id_thunk() { _impl_._oneof_case_[0] = kPartitionIdThunk; } inline void ThunkProto::clear_partition_id_thunk() { ::google::protobuf::internal::TSanWrite(&_impl_); if (impl_case() == kPartitionIdThunk) { if (GetArena() == nullptr) { delete _impl_.impl_.partition_id_thunk_; } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) { ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.impl_.partition_id_thunk_); } clear_has_impl(); } } inline ::xla::gpu::PartitionIdThunkProto* ThunkProto::release_partition_id_thunk() { // @@protoc_insertion_point(field_release:xla.gpu.ThunkProto.partition_id_thunk) if (impl_case() == kPartitionIdThunk) { clear_has_impl(); auto* temp = _impl_.impl_.partition_id_thunk_; if (GetArena() != nullptr) { temp = ::google::protobuf::internal::DuplicateIfNonNull(temp); } _impl_.impl_.partition_id_thunk_ = nullptr; return temp; } else { return nullptr; } } inline const ::xla::gpu::PartitionIdThunkProto& ThunkProto::_internal_partition_id_thunk() const { return impl_case() == kPartitionIdThunk ? *_impl_.impl_.partition_id_thunk_ : reinterpret_cast<::xla::gpu::PartitionIdThunkProto&>(::xla::gpu::_PartitionIdThunkProto_default_instance_); } inline const ::xla::gpu::PartitionIdThunkProto& ThunkProto::partition_id_thunk() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.ThunkProto.partition_id_thunk) return _internal_partition_id_thunk(); } inline ::xla::gpu::PartitionIdThunkProto* ThunkProto::unsafe_arena_release_partition_id_thunk() { // @@protoc_insertion_point(field_unsafe_arena_release:xla.gpu.ThunkProto.partition_id_thunk) if (impl_case() == kPartitionIdThunk) { clear_has_impl(); auto* temp = _impl_.impl_.partition_id_thunk_; _impl_.impl_.partition_id_thunk_ = nullptr; return temp; } else { return nullptr; } } inline void ThunkProto::unsafe_arena_set_allocated_partition_id_thunk(::xla::gpu::PartitionIdThunkProto* value) { // We rely on the oneof clear method to free the earlier contents // of this oneof. We can directly use the pointer we're given to // set the new value. clear_impl(); if (value) { set_has_partition_id_thunk(); _impl_.impl_.partition_id_thunk_ = value; } // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.ThunkProto.partition_id_thunk) } inline ::xla::gpu::PartitionIdThunkProto* ThunkProto::_internal_mutable_partition_id_thunk() { if (impl_case() != kPartitionIdThunk) { clear_impl(); set_has_partition_id_thunk(); _impl_.impl_.partition_id_thunk_ = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::PartitionIdThunkProto>(GetArena()); } return _impl_.impl_.partition_id_thunk_; } inline ::xla::gpu::PartitionIdThunkProto* ThunkProto::mutable_partition_id_thunk() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::xla::gpu::PartitionIdThunkProto* _msg = _internal_mutable_partition_id_thunk(); // @@protoc_insertion_point(field_mutable:xla.gpu.ThunkProto.partition_id_thunk) return _msg; } inline bool ThunkProto::has_impl() const { return impl_case() != IMPL_NOT_SET; } inline void ThunkProto::clear_has_impl() { _impl_._oneof_case_[0] = IMPL_NOT_SET; } inline ThunkProto::ImplCase ThunkProto::impl_case() const { return ThunkProto::ImplCase(_impl_._oneof_case_[0]); } // ------------------------------------------------------------------- // SequentialThunkProto // repeated .xla.gpu.ThunkProto thunks = 1; inline int SequentialThunkProto::_internal_thunks_size() const { return _internal_thunks().size(); } inline int SequentialThunkProto::thunks_size() const { return _internal_thunks_size(); } inline void SequentialThunkProto::clear_thunks() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.thunks_.Clear(); } inline ::xla::gpu::ThunkProto* SequentialThunkProto::mutable_thunks(int index) ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable:xla.gpu.SequentialThunkProto.thunks) return _internal_mutable_thunks()->Mutable(index); } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>* SequentialThunkProto::mutable_thunks() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.SequentialThunkProto.thunks) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_thunks(); } inline const ::xla::gpu::ThunkProto& SequentialThunkProto::thunks(int index) const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.SequentialThunkProto.thunks) return _internal_thunks().Get(index); } inline ::xla::gpu::ThunkProto* SequentialThunkProto::add_thunks() ABSL_ATTRIBUTE_LIFETIME_BOUND { ::google::protobuf::internal::TSanWrite(&_impl_); ::xla::gpu::ThunkProto* _add = _internal_mutable_thunks()->Add(); // @@protoc_insertion_point(field_add:xla.gpu.SequentialThunkProto.thunks) return _add; } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>& SequentialThunkProto::thunks() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.SequentialThunkProto.thunks) return _internal_thunks(); } inline const ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>& SequentialThunkProto::_internal_thunks() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.thunks_; } inline ::google::protobuf::RepeatedPtrField<::xla::gpu::ThunkProto>* SequentialThunkProto::_internal_mutable_thunks() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.thunks_; } #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_2fbackends_2fgpu_2fruntime_2fthunk_2eproto_2epb_2eh