// Generated by the protocol buffer compiler.  DO NOT EDIT!
// NO CHECKED-IN PROTOBUF GENCODE
// source: xla/service/gpu/gpu_norm_runner.proto
// Protobuf C++ Version: 6.31.1

#ifndef xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto_2epb_2eh
#define xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto_2epb_2eh

#include <limits>
#include <string>
#include <type_traits>
#include <utility>

#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/gpu/backend_configs.pb.h"
#include "xla/xla_data.pb.h"
// @@protoc_insertion_point(includes)

// Must be included last.
#include "google/protobuf/port_def.inc"

#define PROTOBUF_INTERNAL_EXPORT_xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto PROTOBUF_EXPORT

namespace google {
namespace protobuf {
namespace internal {
template <typename T>
::absl::string_view GetAnyMessageName();
}  // namespace internal
}  // namespace protobuf
}  // namespace google

// Internal implementation detail -- do not use these members.
struct PROTOBUF_EXPORT TableStruct_xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto {
  static const ::uint32_t offsets[];
};
extern "C" {
PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable descriptor_table_xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto;
}  // extern "C"
namespace xla {
namespace gpu {
class GpuNormDescriptorProto;
struct GpuNormDescriptorProtoDefaultTypeInternal;
PROTOBUF_EXPORT extern GpuNormDescriptorProtoDefaultTypeInternal _GpuNormDescriptorProto_default_instance_;
PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull GpuNormDescriptorProto_class_data_;
}  // namespace gpu
}  // namespace xla
namespace google {
namespace protobuf {
}  // namespace protobuf
}  // namespace google

namespace xla {
namespace gpu {

// ===================================================================


// -------------------------------------------------------------------

class PROTOBUF_EXPORT GpuNormDescriptorProto final : public ::google::protobuf::Message
/* @@protoc_insertion_point(class_definition:xla.gpu.GpuNormDescriptorProto) */ {
 public:
  inline GpuNormDescriptorProto() : GpuNormDescriptorProto(nullptr) {}
  ~GpuNormDescriptorProto() PROTOBUF_FINAL;

#if defined(PROTOBUF_CUSTOM_VTABLE)
  void operator delete(GpuNormDescriptorProto* PROTOBUF_NONNULL msg, std::destroying_delete_t) {
    SharedDtor(*msg);
    ::google::protobuf::internal::SizedDelete(msg, sizeof(GpuNormDescriptorProto));
  }
#endif

  template <typename = void>
  explicit PROTOBUF_CONSTEXPR GpuNormDescriptorProto(::google::protobuf::internal::ConstantInitialized);

  inline GpuNormDescriptorProto(const GpuNormDescriptorProto& from) : GpuNormDescriptorProto(nullptr, from) {}
  inline GpuNormDescriptorProto(GpuNormDescriptorProto&& from) noexcept
      : GpuNormDescriptorProto(nullptr, ::std::move(from)) {}
  inline GpuNormDescriptorProto& operator=(const GpuNormDescriptorProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline GpuNormDescriptorProto& operator=(GpuNormDescriptorProto&& 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 GpuNormDescriptorProto& default_instance() {
    return *reinterpret_cast<const GpuNormDescriptorProto*>(
        &_GpuNormDescriptorProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages = 0;
  friend void swap(GpuNormDescriptorProto& a, GpuNormDescriptorProto& b) { a.Swap(&b); }
  inline void Swap(GpuNormDescriptorProto* 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(GpuNormDescriptorProto* PROTOBUF_NONNULL other) {
    if (other == this) return;
    ABSL_DCHECK(GetArena() == other->GetArena());
    InternalSwap(other);
  }

  // implements Message ----------------------------------------------

  GpuNormDescriptorProto* PROTOBUF_NONNULL New(::google::protobuf::Arena* PROTOBUF_NULLABLE arena = nullptr) const {
    return ::google::protobuf::Message::DefaultConstruct<GpuNormDescriptorProto>(arena);
  }
  using ::google::protobuf::Message::CopyFrom;
  void CopyFrom(const GpuNormDescriptorProto& from);
  using ::google::protobuf::Message::MergeFrom;
  void MergeFrom(const GpuNormDescriptorProto& from) { GpuNormDescriptorProto::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(GpuNormDescriptorProto* PROTOBUF_NONNULL other);
 private:
  template <typename T>
  friend ::absl::string_view(::google::protobuf::internal::GetAnyMessageName)();
  static ::absl::string_view FullMessageName() { return "xla.gpu.GpuNormDescriptorProto"; }

 protected:
  explicit GpuNormDescriptorProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena);
  GpuNormDescriptorProto(::google::protobuf::Arena* PROTOBUF_NULLABLE arena, const GpuNormDescriptorProto& from);
  GpuNormDescriptorProto(
      ::google::protobuf::Arena* PROTOBUF_NULLABLE arena, GpuNormDescriptorProto&& from) noexcept
      : GpuNormDescriptorProto(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 {
    kBackendConfigFieldNumber = 1,
    kXShapeFieldNumber = 2,
    kScaleShapeFieldNumber = 3,
    kBiasShapeFieldNumber = 4,
    kYOrDxShapeFieldNumber = 5,
    kExpectationShapeFieldNumber = 6,
    kNormFactorShapeFieldNumber = 7,
    kDyShapeFieldNumber = 8,
    kDscaleShapeFieldNumber = 9,
    kDbiasShapeFieldNumber = 10,
    kScratchShapeFieldNumber = 11,
  };
  // .xla.gpu.CudnnNormBackendConfig backend_config = 1;
  bool has_backend_config() const;
  void clear_backend_config() ;
  const ::xla::gpu::CudnnNormBackendConfig& backend_config() const;
  [[nodiscard]] ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE release_backend_config();
  ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NONNULL mutable_backend_config();
  void set_allocated_backend_config(::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_backend_config(::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE value);
  ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE unsafe_arena_release_backend_config();

  private:
  const ::xla::gpu::CudnnNormBackendConfig& _internal_backend_config() const;
  ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NONNULL _internal_mutable_backend_config();

  public:
  // .xla.ShapeProto x_shape = 2;
  bool has_x_shape() const;
  void clear_x_shape() ;
  const ::xla::ShapeProto& x_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_x_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_x_shape();
  void set_allocated_x_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_x_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_x_shape();

  private:
  const ::xla::ShapeProto& _internal_x_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_x_shape();

  public:
  // .xla.ShapeProto scale_shape = 3;
  bool has_scale_shape() const;
  void clear_scale_shape() ;
  const ::xla::ShapeProto& scale_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_scale_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_scale_shape();
  void set_allocated_scale_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_scale_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_scale_shape();

  private:
  const ::xla::ShapeProto& _internal_scale_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_scale_shape();

  public:
  // optional .xla.ShapeProto bias_shape = 4;
  bool has_bias_shape() const;
  void clear_bias_shape() ;
  const ::xla::ShapeProto& bias_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_bias_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_bias_shape();
  void set_allocated_bias_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_bias_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_bias_shape();

  private:
  const ::xla::ShapeProto& _internal_bias_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_bias_shape();

  public:
  // .xla.ShapeProto y_or_dx_shape = 5;
  bool has_y_or_dx_shape() const;
  void clear_y_or_dx_shape() ;
  const ::xla::ShapeProto& y_or_dx_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_y_or_dx_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_y_or_dx_shape();
  void set_allocated_y_or_dx_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_y_or_dx_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_y_or_dx_shape();

  private:
  const ::xla::ShapeProto& _internal_y_or_dx_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_y_or_dx_shape();

  public:
  // optional .xla.ShapeProto expectation_shape = 6;
  bool has_expectation_shape() const;
  void clear_expectation_shape() ;
  const ::xla::ShapeProto& expectation_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_expectation_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_expectation_shape();
  void set_allocated_expectation_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_expectation_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_expectation_shape();

  private:
  const ::xla::ShapeProto& _internal_expectation_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_expectation_shape();

  public:
  // optional .xla.ShapeProto norm_factor_shape = 7;
  bool has_norm_factor_shape() const;
  void clear_norm_factor_shape() ;
  const ::xla::ShapeProto& norm_factor_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_norm_factor_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_norm_factor_shape();
  void set_allocated_norm_factor_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_norm_factor_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_norm_factor_shape();

  private:
  const ::xla::ShapeProto& _internal_norm_factor_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_norm_factor_shape();

  public:
  // optional .xla.ShapeProto dy_shape = 8;
  bool has_dy_shape() const;
  void clear_dy_shape() ;
  const ::xla::ShapeProto& dy_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_dy_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_dy_shape();
  void set_allocated_dy_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_dy_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_dy_shape();

  private:
  const ::xla::ShapeProto& _internal_dy_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_dy_shape();

  public:
  // optional .xla.ShapeProto dscale_shape = 9;
  bool has_dscale_shape() const;
  void clear_dscale_shape() ;
  const ::xla::ShapeProto& dscale_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_dscale_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_dscale_shape();
  void set_allocated_dscale_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_dscale_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_dscale_shape();

  private:
  const ::xla::ShapeProto& _internal_dscale_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_dscale_shape();

  public:
  // optional .xla.ShapeProto dbias_shape = 10;
  bool has_dbias_shape() const;
  void clear_dbias_shape() ;
  const ::xla::ShapeProto& dbias_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_dbias_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_dbias_shape();
  void set_allocated_dbias_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_dbias_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_dbias_shape();

  private:
  const ::xla::ShapeProto& _internal_dbias_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_dbias_shape();

  public:
  // optional .xla.ShapeProto scratch_shape = 11;
  bool has_scratch_shape() const;
  void clear_scratch_shape() ;
  const ::xla::ShapeProto& scratch_shape() const;
  [[nodiscard]] ::xla::ShapeProto* PROTOBUF_NULLABLE release_scratch_shape();
  ::xla::ShapeProto* PROTOBUF_NONNULL mutable_scratch_shape();
  void set_allocated_scratch_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  void unsafe_arena_set_allocated_scratch_shape(::xla::ShapeProto* PROTOBUF_NULLABLE value);
  ::xla::ShapeProto* PROTOBUF_NULLABLE unsafe_arena_release_scratch_shape();

  private:
  const ::xla::ShapeProto& _internal_scratch_shape() const;
  ::xla::ShapeProto* PROTOBUF_NONNULL _internal_mutable_scratch_shape();

  public:
  // @@protoc_insertion_point(class_scope:xla.gpu.GpuNormDescriptorProto)
 private:
  class _Internal;
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<4, 11,
                                   11, 0,
                                   2>
      _table_;

  friend class ::google::protobuf::MessageLite;
  friend class ::google::protobuf::Arena;
  template <typename T>
  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 GpuNormDescriptorProto& from_msg);
    ::google::protobuf::internal::HasBits<1> _has_bits_;
    ::google::protobuf::internal::CachedSize _cached_size_;
    ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE backend_config_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE x_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE scale_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE bias_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE y_or_dx_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE expectation_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE norm_factor_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE dy_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE dscale_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE dbias_shape_;
    ::xla::ShapeProto* PROTOBUF_NULLABLE scratch_shape_;
    PROTOBUF_TSAN_DECLARE_MEMBER
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto;
};

PROTOBUF_EXPORT extern const ::google::protobuf::internal::ClassDataFull GpuNormDescriptorProto_class_data_;

// ===================================================================




// ===================================================================


#ifdef __GNUC__
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wstrict-aliasing"
#endif  // __GNUC__
// -------------------------------------------------------------------

// GpuNormDescriptorProto

// .xla.gpu.CudnnNormBackendConfig backend_config = 1;
inline bool GpuNormDescriptorProto::has_backend_config() const {
  bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.backend_config_ != nullptr);
  return value;
}
inline const ::xla::gpu::CudnnNormBackendConfig& GpuNormDescriptorProto::_internal_backend_config() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::gpu::CudnnNormBackendConfig* p = _impl_.backend_config_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::gpu::CudnnNormBackendConfig&>(::xla::gpu::_CudnnNormBackendConfig_default_instance_);
}
inline const ::xla::gpu::CudnnNormBackendConfig& GpuNormDescriptorProto::backend_config() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.backend_config)
  return _internal_backend_config();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_backend_config(
    ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.backend_config_);
  }
  _impl_.backend_config_ = reinterpret_cast<::xla::gpu::CudnnNormBackendConfig*>(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.GpuNormDescriptorProto.backend_config)
}
inline ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_backend_config() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000001u;
  ::xla::gpu::CudnnNormBackendConfig* released = _impl_.backend_config_;
  _impl_.backend_config_ = 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::CudnnNormBackendConfig* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_backend_config() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.backend_config)

  _impl_._has_bits_[0] &= ~0x00000001u;
  ::xla::gpu::CudnnNormBackendConfig* temp = _impl_.backend_config_;
  _impl_.backend_config_ = nullptr;
  return temp;
}
inline ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_backend_config() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.backend_config_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::gpu::CudnnNormBackendConfig>(GetArena());
    _impl_.backend_config_ = reinterpret_cast<::xla::gpu::CudnnNormBackendConfig*>(p);
  }
  return _impl_.backend_config_;
}
inline ::xla::gpu::CudnnNormBackendConfig* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_backend_config()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000001u;
  ::xla::gpu::CudnnNormBackendConfig* _msg = _internal_mutable_backend_config();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.backend_config)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_backend_config(::xla::gpu::CudnnNormBackendConfig* 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_.backend_config_);
  }

  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_.backend_config_ = reinterpret_cast<::xla::gpu::CudnnNormBackendConfig*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.backend_config)
}

// .xla.ShapeProto x_shape = 2;
inline bool GpuNormDescriptorProto::has_x_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000002u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.x_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_x_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.x_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::x_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.x_shape)
  return _internal_x_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_x_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.x_shape_);
  }
  _impl_.x_shape_ = reinterpret_cast<::xla::ShapeProto*>(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.GpuNormDescriptorProto.x_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_x_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000002u;
  ::xla::ShapeProto* released = _impl_.x_shape_;
  _impl_.x_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_x_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.x_shape)

  _impl_._has_bits_[0] &= ~0x00000002u;
  ::xla::ShapeProto* temp = _impl_.x_shape_;
  _impl_.x_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_x_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.x_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.x_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.x_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_x_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000002u;
  ::xla::ShapeProto* _msg = _internal_mutable_x_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.x_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_x_shape(::xla::ShapeProto* 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_.x_shape_);
  }

  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_.x_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.x_shape)
}

// .xla.ShapeProto scale_shape = 3;
inline bool GpuNormDescriptorProto::has_scale_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000004u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.scale_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_scale_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.scale_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::scale_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.scale_shape)
  return _internal_scale_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_scale_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.scale_shape_);
  }
  _impl_.scale_shape_ = reinterpret_cast<::xla::ShapeProto*>(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.GpuNormDescriptorProto.scale_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_scale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000004u;
  ::xla::ShapeProto* released = _impl_.scale_shape_;
  _impl_.scale_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_scale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.scale_shape)

  _impl_._has_bits_[0] &= ~0x00000004u;
  ::xla::ShapeProto* temp = _impl_.scale_shape_;
  _impl_.scale_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_scale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.scale_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.scale_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.scale_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_scale_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000004u;
  ::xla::ShapeProto* _msg = _internal_mutable_scale_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.scale_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_scale_shape(::xla::ShapeProto* 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_.scale_shape_);
  }

  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] |= 0x00000004u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000004u;
  }

  _impl_.scale_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.scale_shape)
}

// optional .xla.ShapeProto bias_shape = 4;
inline bool GpuNormDescriptorProto::has_bias_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000008u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.bias_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_bias_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.bias_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::bias_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.bias_shape)
  return _internal_bias_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_bias_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.bias_shape_);
  }
  _impl_.bias_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000008u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000008u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.bias_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_bias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000008u;
  ::xla::ShapeProto* released = _impl_.bias_shape_;
  _impl_.bias_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_bias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.bias_shape)

  _impl_._has_bits_[0] &= ~0x00000008u;
  ::xla::ShapeProto* temp = _impl_.bias_shape_;
  _impl_.bias_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_bias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.bias_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.bias_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.bias_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_bias_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000008u;
  ::xla::ShapeProto* _msg = _internal_mutable_bias_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.bias_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_bias_shape(::xla::ShapeProto* 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_.bias_shape_);
  }

  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] |= 0x00000008u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000008u;
  }

  _impl_.bias_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.bias_shape)
}

// .xla.ShapeProto y_or_dx_shape = 5;
inline bool GpuNormDescriptorProto::has_y_or_dx_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000010u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.y_or_dx_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_y_or_dx_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.y_or_dx_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::y_or_dx_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.y_or_dx_shape)
  return _internal_y_or_dx_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_y_or_dx_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.y_or_dx_shape_);
  }
  _impl_.y_or_dx_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000010u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000010u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.y_or_dx_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_y_or_dx_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000010u;
  ::xla::ShapeProto* released = _impl_.y_or_dx_shape_;
  _impl_.y_or_dx_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_y_or_dx_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.y_or_dx_shape)

  _impl_._has_bits_[0] &= ~0x00000010u;
  ::xla::ShapeProto* temp = _impl_.y_or_dx_shape_;
  _impl_.y_or_dx_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_y_or_dx_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.y_or_dx_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.y_or_dx_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.y_or_dx_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_y_or_dx_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000010u;
  ::xla::ShapeProto* _msg = _internal_mutable_y_or_dx_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.y_or_dx_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_y_or_dx_shape(::xla::ShapeProto* 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_.y_or_dx_shape_);
  }

  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] |= 0x00000010u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000010u;
  }

  _impl_.y_or_dx_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.y_or_dx_shape)
}

// optional .xla.ShapeProto expectation_shape = 6;
inline bool GpuNormDescriptorProto::has_expectation_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000020u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.expectation_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_expectation_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.expectation_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::expectation_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.expectation_shape)
  return _internal_expectation_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_expectation_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.expectation_shape_);
  }
  _impl_.expectation_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000020u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000020u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.expectation_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_expectation_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000020u;
  ::xla::ShapeProto* released = _impl_.expectation_shape_;
  _impl_.expectation_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_expectation_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.expectation_shape)

  _impl_._has_bits_[0] &= ~0x00000020u;
  ::xla::ShapeProto* temp = _impl_.expectation_shape_;
  _impl_.expectation_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_expectation_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.expectation_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.expectation_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.expectation_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_expectation_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000020u;
  ::xla::ShapeProto* _msg = _internal_mutable_expectation_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.expectation_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_expectation_shape(::xla::ShapeProto* 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_.expectation_shape_);
  }

  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] |= 0x00000020u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000020u;
  }

  _impl_.expectation_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.expectation_shape)
}

// optional .xla.ShapeProto norm_factor_shape = 7;
inline bool GpuNormDescriptorProto::has_norm_factor_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000040u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.norm_factor_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_norm_factor_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.norm_factor_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::norm_factor_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.norm_factor_shape)
  return _internal_norm_factor_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_norm_factor_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.norm_factor_shape_);
  }
  _impl_.norm_factor_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000040u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000040u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.norm_factor_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_norm_factor_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000040u;
  ::xla::ShapeProto* released = _impl_.norm_factor_shape_;
  _impl_.norm_factor_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_norm_factor_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.norm_factor_shape)

  _impl_._has_bits_[0] &= ~0x00000040u;
  ::xla::ShapeProto* temp = _impl_.norm_factor_shape_;
  _impl_.norm_factor_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_norm_factor_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.norm_factor_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.norm_factor_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.norm_factor_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_norm_factor_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000040u;
  ::xla::ShapeProto* _msg = _internal_mutable_norm_factor_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.norm_factor_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_norm_factor_shape(::xla::ShapeProto* 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_.norm_factor_shape_);
  }

  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] |= 0x00000040u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000040u;
  }

  _impl_.norm_factor_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.norm_factor_shape)
}

// optional .xla.ShapeProto dy_shape = 8;
inline bool GpuNormDescriptorProto::has_dy_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000080u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.dy_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_dy_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.dy_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::dy_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.dy_shape)
  return _internal_dy_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_dy_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dy_shape_);
  }
  _impl_.dy_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000080u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000080u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.dy_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_dy_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000080u;
  ::xla::ShapeProto* released = _impl_.dy_shape_;
  _impl_.dy_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_dy_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.dy_shape)

  _impl_._has_bits_[0] &= ~0x00000080u;
  ::xla::ShapeProto* temp = _impl_.dy_shape_;
  _impl_.dy_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_dy_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.dy_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.dy_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.dy_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_dy_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000080u;
  ::xla::ShapeProto* _msg = _internal_mutable_dy_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.dy_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_dy_shape(::xla::ShapeProto* 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_.dy_shape_);
  }

  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] |= 0x00000080u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000080u;
  }

  _impl_.dy_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.dy_shape)
}

// optional .xla.ShapeProto dscale_shape = 9;
inline bool GpuNormDescriptorProto::has_dscale_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000100u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.dscale_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_dscale_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.dscale_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::dscale_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.dscale_shape)
  return _internal_dscale_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_dscale_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dscale_shape_);
  }
  _impl_.dscale_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000100u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000100u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.dscale_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_dscale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000100u;
  ::xla::ShapeProto* released = _impl_.dscale_shape_;
  _impl_.dscale_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_dscale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.dscale_shape)

  _impl_._has_bits_[0] &= ~0x00000100u;
  ::xla::ShapeProto* temp = _impl_.dscale_shape_;
  _impl_.dscale_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_dscale_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.dscale_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.dscale_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.dscale_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_dscale_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000100u;
  ::xla::ShapeProto* _msg = _internal_mutable_dscale_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.dscale_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_dscale_shape(::xla::ShapeProto* 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_.dscale_shape_);
  }

  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] |= 0x00000100u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000100u;
  }

  _impl_.dscale_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.dscale_shape)
}

// optional .xla.ShapeProto dbias_shape = 10;
inline bool GpuNormDescriptorProto::has_dbias_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000200u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.dbias_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_dbias_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.dbias_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::dbias_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.dbias_shape)
  return _internal_dbias_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_dbias_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.dbias_shape_);
  }
  _impl_.dbias_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000200u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000200u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.dbias_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_dbias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000200u;
  ::xla::ShapeProto* released = _impl_.dbias_shape_;
  _impl_.dbias_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_dbias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.dbias_shape)

  _impl_._has_bits_[0] &= ~0x00000200u;
  ::xla::ShapeProto* temp = _impl_.dbias_shape_;
  _impl_.dbias_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_dbias_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.dbias_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.dbias_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.dbias_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_dbias_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000200u;
  ::xla::ShapeProto* _msg = _internal_mutable_dbias_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.dbias_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_dbias_shape(::xla::ShapeProto* 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_.dbias_shape_);
  }

  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] |= 0x00000200u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000200u;
  }

  _impl_.dbias_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.dbias_shape)
}

// optional .xla.ShapeProto scratch_shape = 11;
inline bool GpuNormDescriptorProto::has_scratch_shape() const {
  bool value = (_impl_._has_bits_[0] & 0x00000400u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.scratch_shape_ != nullptr);
  return value;
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::_internal_scratch_shape() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::xla::ShapeProto* p = _impl_.scratch_shape_;
  return p != nullptr ? *p : reinterpret_cast<const ::xla::ShapeProto&>(::xla::_ShapeProto_default_instance_);
}
inline const ::xla::ShapeProto& GpuNormDescriptorProto::scratch_shape() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:xla.gpu.GpuNormDescriptorProto.scratch_shape)
  return _internal_scratch_shape();
}
inline void GpuNormDescriptorProto::unsafe_arena_set_allocated_scratch_shape(
    ::xla::ShapeProto* PROTOBUF_NULLABLE value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.scratch_shape_);
  }
  _impl_.scratch_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000400u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000400u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:xla.gpu.GpuNormDescriptorProto.scratch_shape)
}
inline ::xla::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::release_scratch_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000400u;
  ::xla::ShapeProto* released = _impl_.scratch_shape_;
  _impl_.scratch_shape_ = 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::ShapeProto* PROTOBUF_NULLABLE GpuNormDescriptorProto::unsafe_arena_release_scratch_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:xla.gpu.GpuNormDescriptorProto.scratch_shape)

  _impl_._has_bits_[0] &= ~0x00000400u;
  ::xla::ShapeProto* temp = _impl_.scratch_shape_;
  _impl_.scratch_shape_ = nullptr;
  return temp;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::_internal_mutable_scratch_shape() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.scratch_shape_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::xla::ShapeProto>(GetArena());
    _impl_.scratch_shape_ = reinterpret_cast<::xla::ShapeProto*>(p);
  }
  return _impl_.scratch_shape_;
}
inline ::xla::ShapeProto* PROTOBUF_NONNULL GpuNormDescriptorProto::mutable_scratch_shape()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000400u;
  ::xla::ShapeProto* _msg = _internal_mutable_scratch_shape();
  // @@protoc_insertion_point(field_mutable:xla.gpu.GpuNormDescriptorProto.scratch_shape)
  return _msg;
}
inline void GpuNormDescriptorProto::set_allocated_scratch_shape(::xla::ShapeProto* 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_.scratch_shape_);
  }

  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] |= 0x00000400u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000400u;
  }

  _impl_.scratch_shape_ = reinterpret_cast<::xla::ShapeProto*>(value);
  // @@protoc_insertion_point(field_set_allocated:xla.gpu.GpuNormDescriptorProto.scratch_shape)
}

#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_2fservice_2fgpu_2fgpu_5fnorm_5frunner_2eproto_2epb_2eh