// Generated by the protocol buffer compiler. DO NOT EDIT! // NO CHECKED-IN PROTOBUF GENCODE // source: xla/backends/gpu/runtime/copy_thunk.proto // Protobuf C++ Version: 6.31.1 #ifndef xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto_2epb_2eh #define xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto_2epb_2eh #include #include #include #include #include "google/protobuf/runtime_version.h" #if PROTOBUF_VERSION != 6031001 #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/message_lite.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/shaped_slice.pb.h" // @@protoc_insertion_point(includes) // Must be included last. #include "google/protobuf/port_def.inc" #define PROTOBUF_INTERNAL_EXPORT_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto PROTOBUF_EXPORT namespace google { namespace protobuf { namespace internal { template ::absl::string_view GetAnyMessageName(); } // namespace internal } // namespace protobuf } // namespace google // Internal implementation detail -- do not use these members. struct PROTOBUF_EXPORT TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto { static const ::uint32_t offsets[]; }; extern "C" { PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable descriptor_table_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; } // extern "C" namespace xla { namespace gpu { class CopyThunkProto; struct CopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern CopyThunkProtoDefaultTypeInternal _CopyThunkProto_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull CopyThunkProto_class_data_; class DeviceToDeviceCopyThunkProto; struct DeviceToDeviceCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern DeviceToDeviceCopyThunkProtoDefaultTypeInternal _DeviceToDeviceCopyThunkProto_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DeviceToDeviceCopyThunkProto_class_data_; class DeviceToHostCopyThunkProto; struct DeviceToHostCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern DeviceToHostCopyThunkProtoDefaultTypeInternal _DeviceToHostCopyThunkProto_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DeviceToHostCopyThunkProto_class_data_; class DynamicMemcpyThunkProto; struct DynamicMemcpyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern DynamicMemcpyThunkProtoDefaultTypeInternal _DynamicMemcpyThunkProto_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DynamicMemcpyThunkProto_class_data_; class DynamicMemcpyThunkProto_Offsets; struct DynamicMemcpyThunkProto_OffsetsDefaultTypeInternal; PROTOBUF_EXPORT extern DynamicMemcpyThunkProto_OffsetsDefaultTypeInternal _DynamicMemcpyThunkProto_Offsets_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DynamicMemcpyThunkProto_Offsets_class_data_; class HostToDeviceCopyThunkProto; struct HostToDeviceCopyThunkProtoDefaultTypeInternal; PROTOBUF_EXPORT extern HostToDeviceCopyThunkProtoDefaultTypeInternal _HostToDeviceCopyThunkProto_default_instance_; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull HostToDeviceCopyThunkProto_class_data_; } // namespace gpu } // namespace xla namespace google { namespace protobuf { } // namespace protobuf } // namespace google namespace xla { namespace gpu { // =================================================================== // ------------------------------------------------------------------- class PROTOBUF_EXPORT DynamicMemcpyThunkProto_Offsets final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DynamicMemcpyThunkProto.Offsets) */ { public: inline DynamicMemcpyThunkProto_Offsets() : DynamicMemcpyThunkProto_Offsets(nullptr) {} ~DynamicMemcpyThunkProto_Offsets() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(DynamicMemcpyThunkProto_Offsets)); } #endif template explicit PROTOBUF_CONSTEXPR DynamicMemcpyThunkProto_Offsets(::google::protobuf::internal::ConstantInitialized); inline DynamicMemcpyThunkProto_Offsets(const DynamicMemcpyThunkProto_Offsets& from) : DynamicMemcpyThunkProto_Offsets(nullptr, from) {} inline DynamicMemcpyThunkProto_Offsets(DynamicMemcpyThunkProto_Offsets&& from) noexcept : DynamicMemcpyThunkProto_Offsets(nullptr, ::std::move(from)) {} inline DynamicMemcpyThunkProto_Offsets& operator=(const DynamicMemcpyThunkProto_Offsets& from) { CopyFrom(from); return *this; } inline DynamicMemcpyThunkProto_Offsets& operator=(DynamicMemcpyThunkProto_Offsets&& from) noexcept { if (this == &from) return *this; if (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const DynamicMemcpyThunkProto_Offsets& default_instance() { return *reinterpret_cast( &_DynamicMemcpyThunkProto_Offsets_default_instance_); } static constexpr int kIndexInFileMessages = 4; friend void swap(DynamicMemcpyThunkProto_Offsets& a, DynamicMemcpyThunkProto_Offsets& b) { a.Swap(&b); } inline void Swap(DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DynamicMemcpyThunkProto_Offsets& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DynamicMemcpyThunkProto_Offsets& from) { DynamicMemcpyThunkProto_Offsets::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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.DynamicMemcpyThunkProto.Offsets"; } protected: explicit DynamicMemcpyThunkProto_Offsets(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); DynamicMemcpyThunkProto_Offsets(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const DynamicMemcpyThunkProto_Offsets& from); DynamicMemcpyThunkProto_Offsets( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, DynamicMemcpyThunkProto_Offsets&& from) noexcept : DynamicMemcpyThunkProto_Offsets(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSrcOffsetsFieldNumber = 2, kDstOffsetsFieldNumber = 3, kDependsOnLoopFieldNumber = 1, }; // repeated int64 src_offsets = 2; int src_offsets_size() const; private: int _internal_src_offsets_size() const; public: void clear_src_offsets() ; ::int64_t src_offsets(int index) const; void set_src_offsets(int index, ::int64_t value); void add_src_offsets(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& src_offsets() const; ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL mutable_src_offsets(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_src_offsets() const; ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL _internal_mutable_src_offsets(); public: // repeated int64 dst_offsets = 3; int dst_offsets_size() const; private: int _internal_dst_offsets_size() const; public: void clear_dst_offsets() ; ::int64_t dst_offsets(int index) const; void set_dst_offsets(int index, ::int64_t value); void add_dst_offsets(::int64_t value); const ::google::protobuf::RepeatedField<::int64_t>& dst_offsets() const; ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL mutable_dst_offsets(); private: const ::google::protobuf::RepeatedField<::int64_t>& _internal_dst_offsets() const; ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL _internal_mutable_dst_offsets(); public: // bool depends_on_loop = 1; void clear_depends_on_loop() ; bool depends_on_loop() const; void set_depends_on_loop(bool value); private: bool _internal_depends_on_loop() const; void _internal_set_depends_on_loop(bool value); public: // @@protoc_insertion_point(class_scope:xla.gpu.DynamicMemcpyThunkProto.Offsets) private: class _Internal; friend class ::google::protobuf::internal::TcParser; static const ::google::protobuf::internal::TcParseTable<2, 3, 0, 0, 2> _table_; friend class ::google::protobuf::MessageLite; friend class ::google::protobuf::Arena; template friend class ::google::protobuf::Arena::InternalHelper; using InternalArenaConstructable_ = void; using DestructorSkippable_ = void; struct Impl_ { inline explicit constexpr Impl_(::google::protobuf::internal::ConstantInitialized) noexcept; inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const DynamicMemcpyThunkProto_Offsets& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::google::protobuf::RepeatedField<::int64_t> src_offsets_; ::google::protobuf::internal::CachedSize _src_offsets_cached_byte_size_; ::google::protobuf::RepeatedField<::int64_t> dst_offsets_; ::google::protobuf::internal::CachedSize _dst_offsets_cached_byte_size_; bool depends_on_loop_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DynamicMemcpyThunkProto_Offsets_class_data_; // ------------------------------------------------------------------- class PROTOBUF_EXPORT DynamicMemcpyThunkProto final : public ::google::protobuf::Message /* @@protoc_insertion_point(class_definition:xla.gpu.DynamicMemcpyThunkProto) */ { public: inline DynamicMemcpyThunkProto() : DynamicMemcpyThunkProto(nullptr) {} ~DynamicMemcpyThunkProto() PROTOBUF_FINAL; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(DynamicMemcpyThunkProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(DynamicMemcpyThunkProto)); } #endif template explicit PROTOBUF_CONSTEXPR DynamicMemcpyThunkProto(::google::protobuf::internal::ConstantInitialized); inline DynamicMemcpyThunkProto(const DynamicMemcpyThunkProto& from) : DynamicMemcpyThunkProto(nullptr, from) {} inline DynamicMemcpyThunkProto(DynamicMemcpyThunkProto&& from) noexcept : DynamicMemcpyThunkProto(nullptr, ::std::move(from)) {} inline DynamicMemcpyThunkProto& operator=(const DynamicMemcpyThunkProto& from) { CopyFrom(from); return *this; } inline DynamicMemcpyThunkProto& operator=(DynamicMemcpyThunkProto&& from) noexcept { if (this == &from) return *this; if (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const DynamicMemcpyThunkProto& default_instance() { return *reinterpret_cast( &_DynamicMemcpyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 5; friend void swap(DynamicMemcpyThunkProto& a, DynamicMemcpyThunkProto& b) { a.Swap(&b); } inline void Swap(DynamicMemcpyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DynamicMemcpyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DynamicMemcpyThunkProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { return ::google::protobuf::Message::DefaultConstruct(arena); } using ::google::protobuf::Message::CopyFrom; void CopyFrom(const DynamicMemcpyThunkProto& from); using ::google::protobuf::Message::MergeFrom; void MergeFrom(const DynamicMemcpyThunkProto& from) { DynamicMemcpyThunkProto::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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(DynamicMemcpyThunkProto* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.DynamicMemcpyThunkProto"; } protected: explicit DynamicMemcpyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); DynamicMemcpyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const DynamicMemcpyThunkProto& from); DynamicMemcpyThunkProto( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, DynamicMemcpyThunkProto&& from) noexcept : DynamicMemcpyThunkProto(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- using Offsets = DynamicMemcpyThunkProto_Offsets; // accessors ------------------------------------------------------- enum : int { kSourceBufferFieldNumber = 1, kDestinationBufferFieldNumber = 2, kOffsetsFieldNumber = 4, kMemSizeFieldNumber = 3, }; // .xla.ShapedSliceProto source_buffer = 1; bool has_source_buffer() const; void clear_source_buffer() ; const ::xla::ShapedSliceProto& source_buffer() const; [[nodiscard]] ::xla::ShapedSliceProto* PROTOBUF_NULLABLE release_source_buffer(); ::xla::ShapedSliceProto* PROTOBUF_NONNULL mutable_source_buffer(); void set_allocated_source_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_source_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); ::xla::ShapedSliceProto* PROTOBUF_NULLABLE unsafe_arena_release_source_buffer(); private: const ::xla::ShapedSliceProto& _internal_source_buffer() const; ::xla::ShapedSliceProto* PROTOBUF_NONNULL _internal_mutable_source_buffer(); public: // .xla.ShapedSliceProto destination_buffer = 2; bool has_destination_buffer() const; void clear_destination_buffer() ; const ::xla::ShapedSliceProto& destination_buffer() const; [[nodiscard]] ::xla::ShapedSliceProto* PROTOBUF_NULLABLE release_destination_buffer(); ::xla::ShapedSliceProto* PROTOBUF_NONNULL mutable_destination_buffer(); void set_allocated_destination_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_destination_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); ::xla::ShapedSliceProto* PROTOBUF_NULLABLE unsafe_arena_release_destination_buffer(); private: const ::xla::ShapedSliceProto& _internal_destination_buffer() const; ::xla::ShapedSliceProto* PROTOBUF_NONNULL _internal_mutable_destination_buffer(); public: // .xla.gpu.DynamicMemcpyThunkProto.Offsets offsets = 4; bool has_offsets() const; void clear_offsets() ; const ::xla::gpu::DynamicMemcpyThunkProto_Offsets& offsets() const; [[nodiscard]] ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE release_offsets(); ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL mutable_offsets(); void set_allocated_offsets(::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_offsets(::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE value); ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE unsafe_arena_release_offsets(); private: const ::xla::gpu::DynamicMemcpyThunkProto_Offsets& _internal_offsets() const; ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL _internal_mutable_offsets(); 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.DynamicMemcpyThunkProto) 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* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const DynamicMemcpyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::xla::ShapedSliceProto* PROTOBUF_NULLABLE source_buffer_; ::xla::ShapedSliceProto* PROTOBUF_NULLABLE destination_buffer_; ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE offsets_; ::int64_t mem_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DynamicMemcpyThunkProto_class_data_; // ------------------------------------------------------------------- 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; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(CopyThunkProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(CopyThunkProto)); } #endif 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 (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const CopyThunkProto& default_instance() { return *reinterpret_cast( &_CopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 0; friend void swap(CopyThunkProto& a, CopyThunkProto& b) { a.Swap(&b); } inline void Swap(CopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(CopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- CopyThunkProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { 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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(CopyThunkProto* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.CopyThunkProto"; } protected: explicit CopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); CopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const CopyThunkProto& from); CopyThunkProto( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, CopyThunkProto&& from) noexcept : CopyThunkProto(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::google::protobuf::Metadata GetMetadata() const; // nested types ---------------------------------------------------- // accessors ------------------------------------------------------- enum : int { kSourceBufferFieldNumber = 1, kDestinationBufferFieldNumber = 2, kMemSizeFieldNumber = 3, }; // .xla.ShapedSliceProto source_buffer = 1; bool has_source_buffer() const; void clear_source_buffer() ; const ::xla::ShapedSliceProto& source_buffer() const; [[nodiscard]] ::xla::ShapedSliceProto* PROTOBUF_NULLABLE release_source_buffer(); ::xla::ShapedSliceProto* PROTOBUF_NONNULL mutable_source_buffer(); void set_allocated_source_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_source_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); ::xla::ShapedSliceProto* PROTOBUF_NULLABLE unsafe_arena_release_source_buffer(); private: const ::xla::ShapedSliceProto& _internal_source_buffer() const; ::xla::ShapedSliceProto* PROTOBUF_NONNULL _internal_mutable_source_buffer(); public: // .xla.ShapedSliceProto destination_buffer = 2; bool has_destination_buffer() const; void clear_destination_buffer() ; const ::xla::ShapedSliceProto& destination_buffer() const; [[nodiscard]] ::xla::ShapedSliceProto* PROTOBUF_NULLABLE release_destination_buffer(); ::xla::ShapedSliceProto* PROTOBUF_NONNULL mutable_destination_buffer(); void set_allocated_destination_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_destination_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE value); ::xla::ShapedSliceProto* PROTOBUF_NULLABLE unsafe_arena_release_destination_buffer(); private: const ::xla::ShapedSliceProto& _internal_destination_buffer() const; ::xla::ShapedSliceProto* PROTOBUF_NONNULL _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* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const CopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::xla::ShapedSliceProto* PROTOBUF_NULLABLE source_buffer_; ::xla::ShapedSliceProto* PROTOBUF_NULLABLE destination_buffer_; ::int64_t mem_size_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull CopyThunkProto_class_data_; // ------------------------------------------------------------------- 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; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(HostToDeviceCopyThunkProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(HostToDeviceCopyThunkProto)); } #endif 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 (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const HostToDeviceCopyThunkProto& default_instance() { return *reinterpret_cast( &_HostToDeviceCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 2; friend void swap(HostToDeviceCopyThunkProto& a, HostToDeviceCopyThunkProto& b) { a.Swap(&b); } inline void Swap(HostToDeviceCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(HostToDeviceCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- HostToDeviceCopyThunkProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { 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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(HostToDeviceCopyThunkProto* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.HostToDeviceCopyThunkProto"; } protected: explicit HostToDeviceCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); HostToDeviceCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const HostToDeviceCopyThunkProto& from); HostToDeviceCopyThunkProto( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, HostToDeviceCopyThunkProto&& from) noexcept : HostToDeviceCopyThunkProto(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::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; [[nodiscard]] ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE release_copy_thunk(); ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL _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* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const HostToDeviceCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull HostToDeviceCopyThunkProto_class_data_; // ------------------------------------------------------------------- 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; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(DeviceToHostCopyThunkProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(DeviceToHostCopyThunkProto)); } #endif 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 (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const DeviceToHostCopyThunkProto& default_instance() { return *reinterpret_cast( &_DeviceToHostCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 1; friend void swap(DeviceToHostCopyThunkProto& a, DeviceToHostCopyThunkProto& b) { a.Swap(&b); } inline void Swap(DeviceToHostCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DeviceToHostCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DeviceToHostCopyThunkProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { 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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(DeviceToHostCopyThunkProto* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.DeviceToHostCopyThunkProto"; } protected: explicit DeviceToHostCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); DeviceToHostCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const DeviceToHostCopyThunkProto& from); DeviceToHostCopyThunkProto( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, DeviceToHostCopyThunkProto&& from) noexcept : DeviceToHostCopyThunkProto(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::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; [[nodiscard]] ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE release_copy_thunk(); ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL _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* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const DeviceToHostCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DeviceToHostCopyThunkProto_class_data_; // ------------------------------------------------------------------- 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; #if defined(PROTOBUF_CUSTOM_VTABLE) void operator delete(DeviceToDeviceCopyThunkProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) { SharedDtor(*msg); ::google::protobuf::internal::SizedDelete(msg, sizeof(DeviceToDeviceCopyThunkProto)); } #endif 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 (::google::protobuf::internal::CanMoveWithInternalSwap(GetArena(), from.GetArena())) { 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* PROTOBUF_NONNULL mutable_unknown_fields() ABSL_ATTRIBUTE_LIFETIME_BOUND { return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL descriptor() { return GetDescriptor(); } static const ::google::protobuf::Descriptor* PROTOBUF_NONNULL GetDescriptor() { return default_instance().GetMetadata().descriptor; } static const ::google::protobuf::Reflection* PROTOBUF_NONNULL GetReflection() { return default_instance().GetMetadata().reflection; } static const DeviceToDeviceCopyThunkProto& default_instance() { return *reinterpret_cast( &_DeviceToDeviceCopyThunkProto_default_instance_); } static constexpr int kIndexInFileMessages = 3; friend void swap(DeviceToDeviceCopyThunkProto& a, DeviceToDeviceCopyThunkProto& b) { a.Swap(&b); } inline void Swap(DeviceToDeviceCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; if (::google::protobuf::internal::CanUseInternalSwap(GetArena(), other->GetArena())) { InternalSwap(other); } else { ::google::protobuf::internal::GenericSwap(this, other); } } void UnsafeArenaSwap(DeviceToDeviceCopyThunkProto* PROTOBUF_NONNULL other) { if (other == this) return; ABSL_DCHECK(GetArena() == other->GetArena()); InternalSwap(other); } // implements Message ---------------------------------------------- DeviceToDeviceCopyThunkProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const { 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* PROTOBUF_NONNULL _InternalSerialize( const ::google::protobuf::MessageLite& msg, ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream); public: ::size_t ByteSizeLong() const { return ByteSizeLong(*this); } ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const { return _InternalSerialize(*this, target, stream); } #else // PROTOBUF_CUSTOM_VTABLE ::size_t ByteSizeLong() const final; ::uint8_t* PROTOBUF_NONNULL _InternalSerialize( ::uint8_t* PROTOBUF_NONNULL target, ::google::protobuf::io::EpsCopyOutputStream* PROTOBUF_NONNULL stream) const final; #endif // PROTOBUF_CUSTOM_VTABLE int GetCachedSize() const { return _impl_._cached_size_.Get(); } private: void SharedCtor(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static void SharedDtor(MessageLite& self); void InternalSwap(DeviceToDeviceCopyThunkProto* PROTOBUF_NONNULL other); private: template friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)(); static ::absl::string_view FullMessageName() { return "xla.gpu.DeviceToDeviceCopyThunkProto"; } protected: explicit DeviceToDeviceCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena); DeviceToDeviceCopyThunkProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const DeviceToDeviceCopyThunkProto& from); DeviceToDeviceCopyThunkProto( ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, DeviceToDeviceCopyThunkProto&& from) noexcept : DeviceToDeviceCopyThunkProto(arena) { *this = ::std::move(from); } const ::google::protobuf::internal::ClassData* PROTOBUF_NONNULL GetClassData() const PROTOBUF_FINAL; static void* PROTOBUF_NONNULL PlacementNew_( const void* PROTOBUF_NONNULL, void* PROTOBUF_NONNULL mem, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena); static constexpr auto InternalNewImpl_(); public: static constexpr auto InternalGenerateClassData_(); ::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; [[nodiscard]] ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE release_copy_thunk(); ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL mutable_copy_thunk(); void set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); void unsafe_arena_set_allocated_copy_thunk(::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE value); ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE unsafe_arena_release_copy_thunk(); private: const ::xla::gpu::CopyThunkProto& _internal_copy_thunk() const; ::xla::gpu::CopyThunkProto* PROTOBUF_NONNULL _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* PROTOBUF_NULLABLE arena); inline explicit Impl_( ::google::protobuf::internal::InternalVisibility visibility, ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const Impl_& from, const DeviceToDeviceCopyThunkProto& from_msg); ::google::protobuf::internal::HasBits<1> _has_bits_; ::google::protobuf::internal::CachedSize _cached_size_; ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE copy_thunk_; PROTOBUF_TSAN_DECLARE_MEMBER }; union { Impl_ _impl_; }; friend struct ::TableStruct_xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto; }; PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull DeviceToDeviceCopyThunkProto_class_data_; // =================================================================== // =================================================================== #ifdef __GNUC__ #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif // __GNUC__ // ------------------------------------------------------------------- // CopyThunkProto // .xla.ShapedSliceProto 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::ShapedSliceProto& CopyThunkProto::_internal_source_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ShapedSliceProto* p = _impl_.source_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::_ShapedSliceProto_default_instance_); } inline const ::xla::ShapedSliceProto& 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::ShapedSliceProto* PROTOBUF_NULLABLE value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_buffer_); } _impl_.source_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(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::ShapedSliceProto* PROTOBUF_NULLABLE CopyThunkProto::release_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::ShapedSliceProto* released = _impl_.source_buffer_; _impl_.source_buffer_ = nullptr; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE 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::ShapedSliceProto* temp = _impl_.source_buffer_; _impl_.source_buffer_ = nullptr; return temp; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL CopyThunkProto::_internal_mutable_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.source_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapedSliceProto>(GetArena()); _impl_.source_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(p); } return _impl_.source_buffer_; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL CopyThunkProto::mutable_source_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::ShapedSliceProto* _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::ShapedSliceProto* PROTOBUF_NULLABLE 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::Message*>(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::ShapedSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.CopyThunkProto.source_buffer) } // .xla.ShapedSliceProto 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::ShapedSliceProto& CopyThunkProto::_internal_destination_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ShapedSliceProto* p = _impl_.destination_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::_ShapedSliceProto_default_instance_); } inline const ::xla::ShapedSliceProto& 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::ShapedSliceProto* PROTOBUF_NULLABLE value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.destination_buffer_); } _impl_.destination_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(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::ShapedSliceProto* PROTOBUF_NULLABLE CopyThunkProto::release_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::ShapedSliceProto* released = _impl_.destination_buffer_; _impl_.destination_buffer_ = nullptr; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE 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::ShapedSliceProto* temp = _impl_.destination_buffer_; _impl_.destination_buffer_ = nullptr; return temp; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL CopyThunkProto::_internal_mutable_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.destination_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapedSliceProto>(GetArena()); _impl_.destination_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(p); } return _impl_.destination_buffer_; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL CopyThunkProto::mutable_destination_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::ShapedSliceProto* _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::ShapedSliceProto* PROTOBUF_NULLABLE 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::Message*>(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::ShapedSliceProto*>(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}; _impl_._has_bits_[0] &= ~0x00000004u; } 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); _impl_._has_bits_[0] |= 0x00000004u; // @@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* PROTOBUF_NULLABLE 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* PROTOBUF_NULLABLE 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; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE 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* PROTOBUF_NONNULL 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* PROTOBUF_NONNULL 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* PROTOBUF_NULLABLE value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_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* PROTOBUF_NULLABLE 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* PROTOBUF_NULLABLE 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; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE 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* PROTOBUF_NONNULL 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* PROTOBUF_NONNULL 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* PROTOBUF_NULLABLE value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_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* PROTOBUF_NULLABLE 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* PROTOBUF_NULLABLE 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; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::gpu::CopyThunkProto* PROTOBUF_NULLABLE 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* PROTOBUF_NONNULL 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* PROTOBUF_NONNULL 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* PROTOBUF_NULLABLE value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_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) } // ------------------------------------------------------------------- // DynamicMemcpyThunkProto_Offsets // bool depends_on_loop = 1; inline void DynamicMemcpyThunkProto_Offsets::clear_depends_on_loop() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.depends_on_loop_ = false; _impl_._has_bits_[0] &= ~0x00000001u; } inline bool DynamicMemcpyThunkProto_Offsets::depends_on_loop() const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.Offsets.depends_on_loop) return _internal_depends_on_loop(); } inline void DynamicMemcpyThunkProto_Offsets::set_depends_on_loop(bool value) { _internal_set_depends_on_loop(value); _impl_._has_bits_[0] |= 0x00000001u; // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyThunkProto.Offsets.depends_on_loop) } inline bool DynamicMemcpyThunkProto_Offsets::_internal_depends_on_loop() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.depends_on_loop_; } inline void DynamicMemcpyThunkProto_Offsets::_internal_set_depends_on_loop(bool value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.depends_on_loop_ = value; } // repeated int64 src_offsets = 2; inline int DynamicMemcpyThunkProto_Offsets::_internal_src_offsets_size() const { return _internal_src_offsets().size(); } inline int DynamicMemcpyThunkProto_Offsets::src_offsets_size() const { return _internal_src_offsets_size(); } inline void DynamicMemcpyThunkProto_Offsets::clear_src_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.src_offsets_.Clear(); } inline ::int64_t DynamicMemcpyThunkProto_Offsets::src_offsets(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.Offsets.src_offsets) return _internal_src_offsets().Get(index); } inline void DynamicMemcpyThunkProto_Offsets::set_src_offsets(int index, ::int64_t value) { _internal_mutable_src_offsets()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyThunkProto.Offsets.src_offsets) } inline void DynamicMemcpyThunkProto_Offsets::add_src_offsets(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_src_offsets()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.DynamicMemcpyThunkProto.Offsets.src_offsets) } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyThunkProto_Offsets::src_offsets() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.DynamicMemcpyThunkProto.Offsets.src_offsets) return _internal_src_offsets(); } inline ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL DynamicMemcpyThunkProto_Offsets::mutable_src_offsets() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.DynamicMemcpyThunkProto.Offsets.src_offsets) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_src_offsets(); } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyThunkProto_Offsets::_internal_src_offsets() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.src_offsets_; } inline ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL DynamicMemcpyThunkProto_Offsets::_internal_mutable_src_offsets() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.src_offsets_; } // repeated int64 dst_offsets = 3; inline int DynamicMemcpyThunkProto_Offsets::_internal_dst_offsets_size() const { return _internal_dst_offsets().size(); } inline int DynamicMemcpyThunkProto_Offsets::dst_offsets_size() const { return _internal_dst_offsets_size(); } inline void DynamicMemcpyThunkProto_Offsets::clear_dst_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.dst_offsets_.Clear(); } inline ::int64_t DynamicMemcpyThunkProto_Offsets::dst_offsets(int index) const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.Offsets.dst_offsets) return _internal_dst_offsets().Get(index); } inline void DynamicMemcpyThunkProto_Offsets::set_dst_offsets(int index, ::int64_t value) { _internal_mutable_dst_offsets()->Set(index, value); // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyThunkProto.Offsets.dst_offsets) } inline void DynamicMemcpyThunkProto_Offsets::add_dst_offsets(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _internal_mutable_dst_offsets()->Add(value); // @@protoc_insertion_point(field_add:xla.gpu.DynamicMemcpyThunkProto.Offsets.dst_offsets) } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyThunkProto_Offsets::dst_offsets() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_list:xla.gpu.DynamicMemcpyThunkProto.Offsets.dst_offsets) return _internal_dst_offsets(); } inline ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL DynamicMemcpyThunkProto_Offsets::mutable_dst_offsets() ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_mutable_list:xla.gpu.DynamicMemcpyThunkProto.Offsets.dst_offsets) ::google::protobuf::internal::TSanWrite(&_impl_); return _internal_mutable_dst_offsets(); } inline const ::google::protobuf::RepeatedField<::int64_t>& DynamicMemcpyThunkProto_Offsets::_internal_dst_offsets() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.dst_offsets_; } inline ::google::protobuf::RepeatedField<::int64_t>* PROTOBUF_NONNULL DynamicMemcpyThunkProto_Offsets::_internal_mutable_dst_offsets() { ::google::protobuf::internal::TSanRead(&_impl_); return &_impl_.dst_offsets_; } // ------------------------------------------------------------------- // DynamicMemcpyThunkProto // .xla.ShapedSliceProto source_buffer = 1; inline bool DynamicMemcpyThunkProto::has_source_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0; PROTOBUF_ASSUME(!value || _impl_.source_buffer_ != nullptr); return value; } inline const ::xla::ShapedSliceProto& DynamicMemcpyThunkProto::_internal_source_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ShapedSliceProto* p = _impl_.source_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::_ShapedSliceProto_default_instance_); } inline const ::xla::ShapedSliceProto& DynamicMemcpyThunkProto::source_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.source_buffer) return _internal_source_buffer(); } inline void DynamicMemcpyThunkProto::unsafe_arena_set_allocated_source_buffer( ::xla::ShapedSliceProto* PROTOBUF_NULLABLE value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.source_buffer_); } _impl_.source_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(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.DynamicMemcpyThunkProto.source_buffer) } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::release_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000001u; ::xla::ShapedSliceProto* released = _impl_.source_buffer_; _impl_.source_buffer_ = nullptr; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::unsafe_arena_release_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.DynamicMemcpyThunkProto.source_buffer) _impl_._has_bits_[0] &= ~0x00000001u; ::xla::ShapedSliceProto* temp = _impl_.source_buffer_; _impl_.source_buffer_ = nullptr; return temp; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL DynamicMemcpyThunkProto::_internal_mutable_source_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.source_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapedSliceProto>(GetArena()); _impl_.source_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(p); } return _impl_.source_buffer_; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL DynamicMemcpyThunkProto::mutable_source_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000001u; ::xla::ShapedSliceProto* _msg = _internal_mutable_source_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.DynamicMemcpyThunkProto.source_buffer) return _msg; } inline void DynamicMemcpyThunkProto::set_allocated_source_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE 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::Message*>(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::ShapedSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.DynamicMemcpyThunkProto.source_buffer) } // .xla.ShapedSliceProto destination_buffer = 2; inline bool DynamicMemcpyThunkProto::has_destination_buffer() const { bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0; PROTOBUF_ASSUME(!value || _impl_.destination_buffer_ != nullptr); return value; } inline const ::xla::ShapedSliceProto& DynamicMemcpyThunkProto::_internal_destination_buffer() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::ShapedSliceProto* p = _impl_.destination_buffer_; return p != nullptr ? *p : reinterpret_cast(::xla::_ShapedSliceProto_default_instance_); } inline const ::xla::ShapedSliceProto& DynamicMemcpyThunkProto::destination_buffer() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.destination_buffer) return _internal_destination_buffer(); } inline void DynamicMemcpyThunkProto::unsafe_arena_set_allocated_destination_buffer( ::xla::ShapedSliceProto* PROTOBUF_NULLABLE value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.destination_buffer_); } _impl_.destination_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(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.DynamicMemcpyThunkProto.destination_buffer) } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::release_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000002u; ::xla::ShapedSliceProto* released = _impl_.destination_buffer_; _impl_.destination_buffer_ = nullptr; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::ShapedSliceProto* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::unsafe_arena_release_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.DynamicMemcpyThunkProto.destination_buffer) _impl_._has_bits_[0] &= ~0x00000002u; ::xla::ShapedSliceProto* temp = _impl_.destination_buffer_; _impl_.destination_buffer_ = nullptr; return temp; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL DynamicMemcpyThunkProto::_internal_mutable_destination_buffer() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.destination_buffer_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapedSliceProto>(GetArena()); _impl_.destination_buffer_ = reinterpret_cast<::xla::ShapedSliceProto*>(p); } return _impl_.destination_buffer_; } inline ::xla::ShapedSliceProto* PROTOBUF_NONNULL DynamicMemcpyThunkProto::mutable_destination_buffer() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000002u; ::xla::ShapedSliceProto* _msg = _internal_mutable_destination_buffer(); // @@protoc_insertion_point(field_mutable:xla.gpu.DynamicMemcpyThunkProto.destination_buffer) return _msg; } inline void DynamicMemcpyThunkProto::set_allocated_destination_buffer(::xla::ShapedSliceProto* PROTOBUF_NULLABLE 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::Message*>(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::ShapedSliceProto*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.DynamicMemcpyThunkProto.destination_buffer) } // int64 mem_size = 3; inline void DynamicMemcpyThunkProto::clear_mem_size() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mem_size_ = ::int64_t{0}; _impl_._has_bits_[0] &= ~0x00000008u; } inline ::int64_t DynamicMemcpyThunkProto::mem_size() const { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.mem_size) return _internal_mem_size(); } inline void DynamicMemcpyThunkProto::set_mem_size(::int64_t value) { _internal_set_mem_size(value); _impl_._has_bits_[0] |= 0x00000008u; // @@protoc_insertion_point(field_set:xla.gpu.DynamicMemcpyThunkProto.mem_size) } inline ::int64_t DynamicMemcpyThunkProto::_internal_mem_size() const { ::google::protobuf::internal::TSanRead(&_impl_); return _impl_.mem_size_; } inline void DynamicMemcpyThunkProto::_internal_set_mem_size(::int64_t value) { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_.mem_size_ = value; } // .xla.gpu.DynamicMemcpyThunkProto.Offsets offsets = 4; inline bool DynamicMemcpyThunkProto::has_offsets() const { bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0; PROTOBUF_ASSUME(!value || _impl_.offsets_ != nullptr); return value; } inline void DynamicMemcpyThunkProto::clear_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.offsets_ != nullptr) _impl_.offsets_->Clear(); _impl_._has_bits_[0] &= ~0x00000004u; } inline const ::xla::gpu::DynamicMemcpyThunkProto_Offsets& DynamicMemcpyThunkProto::_internal_offsets() const { ::google::protobuf::internal::TSanRead(&_impl_); const ::xla::gpu::DynamicMemcpyThunkProto_Offsets* p = _impl_.offsets_; return p != nullptr ? *p : reinterpret_cast(::xla::gpu::_DynamicMemcpyThunkProto_Offsets_default_instance_); } inline const ::xla::gpu::DynamicMemcpyThunkProto_Offsets& DynamicMemcpyThunkProto::offsets() const ABSL_ATTRIBUTE_LIFETIME_BOUND { // @@protoc_insertion_point(field_get:xla.gpu.DynamicMemcpyThunkProto.offsets) return _internal_offsets(); } inline void DynamicMemcpyThunkProto::unsafe_arena_set_allocated_offsets( ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE value) { ::google::protobuf::internal::TSanWrite(&_impl_); if (GetArena() == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.offsets_); } _impl_.offsets_ = reinterpret_cast<::xla::gpu::DynamicMemcpyThunkProto_Offsets*>(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.DynamicMemcpyThunkProto.offsets) } inline ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::release_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::DynamicMemcpyThunkProto_Offsets* released = _impl_.offsets_; _impl_.offsets_ = nullptr; if (::google::protobuf::internal::DebugHardenForceCopyInRelease()) { auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released); released = ::google::protobuf::internal::DuplicateIfNonNull(released); if (GetArena() == nullptr) { delete old; } } else { if (GetArena() != nullptr) { released = ::google::protobuf::internal::DuplicateIfNonNull(released); } } return released; } inline ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE DynamicMemcpyThunkProto::unsafe_arena_release_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); // @@protoc_insertion_point(field_release:xla.gpu.DynamicMemcpyThunkProto.offsets) _impl_._has_bits_[0] &= ~0x00000004u; ::xla::gpu::DynamicMemcpyThunkProto_Offsets* temp = _impl_.offsets_; _impl_.offsets_ = nullptr; return temp; } inline ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL DynamicMemcpyThunkProto::_internal_mutable_offsets() { ::google::protobuf::internal::TSanWrite(&_impl_); if (_impl_.offsets_ == nullptr) { auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::DynamicMemcpyThunkProto_Offsets>(GetArena()); _impl_.offsets_ = reinterpret_cast<::xla::gpu::DynamicMemcpyThunkProto_Offsets*>(p); } return _impl_.offsets_; } inline ::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NONNULL DynamicMemcpyThunkProto::mutable_offsets() ABSL_ATTRIBUTE_LIFETIME_BOUND { _impl_._has_bits_[0] |= 0x00000004u; ::xla::gpu::DynamicMemcpyThunkProto_Offsets* _msg = _internal_mutable_offsets(); // @@protoc_insertion_point(field_mutable:xla.gpu.DynamicMemcpyThunkProto.offsets) return _msg; } inline void DynamicMemcpyThunkProto::set_allocated_offsets(::xla::gpu::DynamicMemcpyThunkProto_Offsets* PROTOBUF_NULLABLE value) { ::google::protobuf::Arena* message_arena = GetArena(); ::google::protobuf::internal::TSanWrite(&_impl_); if (message_arena == nullptr) { delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.offsets_); } 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_.offsets_ = reinterpret_cast<::xla::gpu::DynamicMemcpyThunkProto_Offsets*>(value); // @@protoc_insertion_point(field_set_allocated:xla.gpu.DynamicMemcpyThunkProto.offsets) } #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 // xla_2fbackends_2fgpu_2fruntime_2fcopy_5fthunk_2eproto_2epb_2eh