// Generated by the protocol buffer compiler.  DO NOT EDIT!
// NO CHECKED-IN PROTOBUF GENCODE
// source: xla/tsl/protobuf/dnn.proto
// Protobuf C++ Version: 5.28.3

#ifndef GOOGLE_PROTOBUF_INCLUDED_xla_2ftsl_2fprotobuf_2fdnn_2eproto_2epb_2eh
#define GOOGLE_PROTOBUF_INCLUDED_xla_2ftsl_2fprotobuf_2fdnn_2eproto_2epb_2eh

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

#include "google/protobuf/runtime_version.h"
#if PROTOBUF_VERSION != 5028003
#error "Protobuf C++ gencode is built with an incompatible version of"
#error "Protobuf C++ headers/runtime. See"
#error "https://protobuf.dev/support/cross-version-runtime-guarantee/#cpp"
#endif
#include "google/protobuf/io/coded_stream.h"
#include "google/protobuf/arena.h"
#include "google/protobuf/arenastring.h"
#include "google/protobuf/generated_message_tctable_decl.h"
#include "google/protobuf/generated_message_util.h"
#include "google/protobuf/metadata_lite.h"
#include "google/protobuf/generated_message_reflection.h"
#include "google/protobuf/message.h"
#include "google/protobuf/repeated_field.h"  // IWYU pragma: export
#include "google/protobuf/extension_set.h"  // IWYU pragma: export
#include "google/protobuf/map.h"  // IWYU pragma: export
#include "google/protobuf/map_entry.h"
#include "google/protobuf/map_field_inl.h"
#include "google/protobuf/generated_enum_reflection.h"
#include "google/protobuf/unknown_field_set.h"
#include "google/protobuf/wrappers.pb.h"
// @@protoc_insertion_point(includes)

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

#define PROTOBUF_INTERNAL_EXPORT_xla_2ftsl_2fprotobuf_2fdnn_2eproto PROTOBUF_EXPORT

namespace google {
namespace protobuf {
namespace internal {
class AnyMetadata;
}  // namespace internal
}  // namespace protobuf
}  // namespace google

// Internal implementation detail -- do not use these members.
struct PROTOBUF_EXPORT TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto {
  static const ::uint32_t offsets[];
};
PROTOBUF_EXPORT extern const ::google::protobuf::internal::DescriptorTable
    descriptor_table_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
namespace stream_executor {
namespace dnn {
class AlgorithmConfigProto;
struct AlgorithmConfigProtoDefaultTypeInternal;
PROTOBUF_EXPORT extern AlgorithmConfigProtoDefaultTypeInternal _AlgorithmConfigProto_default_instance_;
class AlgorithmProto;
struct AlgorithmProtoDefaultTypeInternal;
PROTOBUF_EXPORT extern AlgorithmProtoDefaultTypeInternal _AlgorithmProto_default_instance_;
class AlgorithmProto_TuningKnobsEntry_DoNotUse;
struct AlgorithmProto_TuningKnobsEntry_DoNotUseDefaultTypeInternal;
PROTOBUF_EXPORT extern AlgorithmProto_TuningKnobsEntry_DoNotUseDefaultTypeInternal _AlgorithmProto_TuningKnobsEntry_DoNotUse_default_instance_;
class ConvolutionDescriptorProto;
struct ConvolutionDescriptorProtoDefaultTypeInternal;
PROTOBUF_EXPORT extern ConvolutionDescriptorProtoDefaultTypeInternal _ConvolutionDescriptorProto_default_instance_;
class TensorDescriptorProto;
struct TensorDescriptorProtoDefaultTypeInternal;
PROTOBUF_EXPORT extern TensorDescriptorProtoDefaultTypeInternal _TensorDescriptorProto_default_instance_;
}  // namespace dnn
}  // namespace stream_executor
namespace google {
namespace protobuf {
}  // namespace protobuf
}  // namespace google

namespace stream_executor {
namespace dnn {
enum AlgorithmProto_MathType : int {
  AlgorithmProto_MathType_DEFAULT_MATH = 0,
  AlgorithmProto_MathType_TENSOR_OP_MATH = 1,
  AlgorithmProto_MathType_AlgorithmProto_MathType_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  AlgorithmProto_MathType_AlgorithmProto_MathType_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool AlgorithmProto_MathType_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t AlgorithmProto_MathType_internal_data_[];
constexpr AlgorithmProto_MathType AlgorithmProto_MathType_MathType_MIN = static_cast<AlgorithmProto_MathType>(0);
constexpr AlgorithmProto_MathType AlgorithmProto_MathType_MathType_MAX = static_cast<AlgorithmProto_MathType>(1);
constexpr int AlgorithmProto_MathType_MathType_ARRAYSIZE = 1 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
AlgorithmProto_MathType_descriptor();
template <typename T>
const std::string& AlgorithmProto_MathType_Name(T value) {
  static_assert(std::is_same<T, AlgorithmProto_MathType>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to MathType_Name().");
  return AlgorithmProto_MathType_Name(static_cast<AlgorithmProto_MathType>(value));
}
template <>
inline const std::string& AlgorithmProto_MathType_Name(AlgorithmProto_MathType value) {
  return ::google::protobuf::internal::NameOfDenseEnum<AlgorithmProto_MathType_descriptor,
                                                 0, 1>(
      static_cast<int>(value));
}
inline bool AlgorithmProto_MathType_Parse(absl::string_view name, AlgorithmProto_MathType* value) {
  return ::google::protobuf::internal::ParseNamedEnum<AlgorithmProto_MathType>(
      AlgorithmProto_MathType_descriptor(), name, value);
}
enum DataType : int {
  kFloat = 0,
  kDouble = 1,
  kHalf = 2,
  kInt8 = 3,
  kInt32 = 4,
  kComplexFloat = 5,
  kComplexDouble = 6,
  kBF16 = 7,
  kF8E5M2 = 8,
  kF8E4M3FN = 9,
  kF8E5M2FNUZ = 10,
  kF8E4M3FNUZ = 11,
  kInt64 = 12,
  kF8E4M3 = 13,
  kF8E3M4 = 14,
  kF4E2M1FN = 15,
  kF8E8M0FNU = 16,
  DataType_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  DataType_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool DataType_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t DataType_internal_data_[];
constexpr DataType DataType_MIN = static_cast<DataType>(0);
constexpr DataType DataType_MAX = static_cast<DataType>(16);
constexpr int DataType_ARRAYSIZE = 16 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
DataType_descriptor();
template <typename T>
const std::string& DataType_Name(T value) {
  static_assert(std::is_same<T, DataType>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to DataType_Name().");
  return DataType_Name(static_cast<DataType>(value));
}
template <>
inline const std::string& DataType_Name(DataType value) {
  return ::google::protobuf::internal::NameOfDenseEnum<DataType_descriptor,
                                                 0, 16>(
      static_cast<int>(value));
}
inline bool DataType_Parse(absl::string_view name, DataType* value) {
  return ::google::protobuf::internal::ParseNamedEnum<DataType>(
      DataType_descriptor(), name, value);
}
enum DataLayout : int {
  kYXDepthBatch = 0,
  kYXBatchDepth = 1,
  kBatchYXDepth = 2,
  kBatchDepthYX = 3,
  kBatchDepthYX4 = 4,
  kBatchDepthYX32 = 5,
  DataLayout_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  DataLayout_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool DataLayout_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t DataLayout_internal_data_[];
constexpr DataLayout DataLayout_MIN = static_cast<DataLayout>(0);
constexpr DataLayout DataLayout_MAX = static_cast<DataLayout>(5);
constexpr int DataLayout_ARRAYSIZE = 5 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
DataLayout_descriptor();
template <typename T>
const std::string& DataLayout_Name(T value) {
  static_assert(std::is_same<T, DataLayout>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to DataLayout_Name().");
  return DataLayout_Name(static_cast<DataLayout>(value));
}
template <>
inline const std::string& DataLayout_Name(DataLayout value) {
  return ::google::protobuf::internal::NameOfDenseEnum<DataLayout_descriptor,
                                                 0, 5>(
      static_cast<int>(value));
}
inline bool DataLayout_Parse(absl::string_view name, DataLayout* value) {
  return ::google::protobuf::internal::ParseNamedEnum<DataLayout>(
      DataLayout_descriptor(), name, value);
}
enum FilterLayout : int {
  kOutputInputYX = 0,
  kOutputYXInput = 1,
  kOutputInputYX4 = 2,
  kOutputInputYX32 = 5,
  kOutputInputYX32_CudnnReordered = 6,
  kInputYXOutput = 3,
  kYXInputOutput = 4,
  FilterLayout_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  FilterLayout_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool FilterLayout_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t FilterLayout_internal_data_[];
constexpr FilterLayout FilterLayout_MIN = static_cast<FilterLayout>(0);
constexpr FilterLayout FilterLayout_MAX = static_cast<FilterLayout>(6);
constexpr int FilterLayout_ARRAYSIZE = 6 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
FilterLayout_descriptor();
template <typename T>
const std::string& FilterLayout_Name(T value) {
  static_assert(std::is_same<T, FilterLayout>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to FilterLayout_Name().");
  return FilterLayout_Name(static_cast<FilterLayout>(value));
}
template <>
inline const std::string& FilterLayout_Name(FilterLayout value) {
  return ::google::protobuf::internal::NameOfDenseEnum<FilterLayout_descriptor,
                                                 0, 6>(
      static_cast<int>(value));
}
inline bool FilterLayout_Parse(absl::string_view name, FilterLayout* value) {
  return ::google::protobuf::internal::ParseNamedEnum<FilterLayout>(
      FilterLayout_descriptor(), name, value);
}
enum ActivationMode : int {
  kNone = 0,
  kSigmoid = 1,
  kRelu = 2,
  kRelu6 = 3,
  kReluX = 4,
  kTanh = 5,
  kBandPass = 6,
  kElu = 7,
  kLeakyRelu = 8,
  kGeluExact = 9,
  ActivationMode_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  ActivationMode_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool ActivationMode_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t ActivationMode_internal_data_[];
constexpr ActivationMode ActivationMode_MIN = static_cast<ActivationMode>(0);
constexpr ActivationMode ActivationMode_MAX = static_cast<ActivationMode>(9);
constexpr int ActivationMode_ARRAYSIZE = 9 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
ActivationMode_descriptor();
template <typename T>
const std::string& ActivationMode_Name(T value) {
  static_assert(std::is_same<T, ActivationMode>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to ActivationMode_Name().");
  return ActivationMode_Name(static_cast<ActivationMode>(value));
}
template <>
inline const std::string& ActivationMode_Name(ActivationMode value) {
  return ::google::protobuf::internal::NameOfDenseEnum<ActivationMode_descriptor,
                                                 0, 9>(
      static_cast<int>(value));
}
inline bool ActivationMode_Parse(absl::string_view name, ActivationMode* value) {
  return ::google::protobuf::internal::ParseNamedEnum<ActivationMode>(
      ActivationMode_descriptor(), name, value);
}
enum ConvolutionMode : int {
  CROSS_CORRELATION = 0,
  CONVOLUTION = 1,
  ConvolutionMode_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  ConvolutionMode_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool ConvolutionMode_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t ConvolutionMode_internal_data_[];
constexpr ConvolutionMode ConvolutionMode_MIN = static_cast<ConvolutionMode>(0);
constexpr ConvolutionMode ConvolutionMode_MAX = static_cast<ConvolutionMode>(1);
constexpr int ConvolutionMode_ARRAYSIZE = 1 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
ConvolutionMode_descriptor();
template <typename T>
const std::string& ConvolutionMode_Name(T value) {
  static_assert(std::is_same<T, ConvolutionMode>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to ConvolutionMode_Name().");
  return ConvolutionMode_Name(static_cast<ConvolutionMode>(value));
}
template <>
inline const std::string& ConvolutionMode_Name(ConvolutionMode value) {
  return ::google::protobuf::internal::NameOfDenseEnum<ConvolutionMode_descriptor,
                                                 0, 1>(
      static_cast<int>(value));
}
inline bool ConvolutionMode_Parse(absl::string_view name, ConvolutionMode* value) {
  return ::google::protobuf::internal::ParseNamedEnum<ConvolutionMode>(
      ConvolutionMode_descriptor(), name, value);
}
enum ConvolutionKind : int {
  INVALID = 0,
  FORWARD = 1,
  BACKWARD_FILTER = 2,
  BACKWARD_DATA = 3,
  FORWARD_BIAS_ACTIVATION = 4,
  FORWARD_GRAPH = 5,
  ConvolutionKind_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  ConvolutionKind_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool ConvolutionKind_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t ConvolutionKind_internal_data_[];
constexpr ConvolutionKind ConvolutionKind_MIN = static_cast<ConvolutionKind>(0);
constexpr ConvolutionKind ConvolutionKind_MAX = static_cast<ConvolutionKind>(5);
constexpr int ConvolutionKind_ARRAYSIZE = 5 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
ConvolutionKind_descriptor();
template <typename T>
const std::string& ConvolutionKind_Name(T value) {
  static_assert(std::is_same<T, ConvolutionKind>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to ConvolutionKind_Name().");
  return ConvolutionKind_Name(static_cast<ConvolutionKind>(value));
}
template <>
inline const std::string& ConvolutionKind_Name(ConvolutionKind value) {
  return ::google::protobuf::internal::NameOfDenseEnum<ConvolutionKind_descriptor,
                                                 0, 5>(
      static_cast<int>(value));
}
inline bool ConvolutionKind_Parse(absl::string_view name, ConvolutionKind* value) {
  return ::google::protobuf::internal::ParseNamedEnum<ConvolutionKind>(
      ConvolutionKind_descriptor(), name, value);
}
enum NormKind : int {
  LAYER_FWD_INFER = 0,
  LAYER_FWD_TRAIN = 1,
  LAYER_BWD = 2,
  NormKind_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  NormKind_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool NormKind_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t NormKind_internal_data_[];
constexpr NormKind NormKind_MIN = static_cast<NormKind>(0);
constexpr NormKind NormKind_MAX = static_cast<NormKind>(2);
constexpr int NormKind_ARRAYSIZE = 2 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
NormKind_descriptor();
template <typename T>
const std::string& NormKind_Name(T value) {
  static_assert(std::is_same<T, NormKind>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to NormKind_Name().");
  return NormKind_Name(static_cast<NormKind>(value));
}
template <>
inline const std::string& NormKind_Name(NormKind value) {
  return ::google::protobuf::internal::NameOfDenseEnum<NormKind_descriptor,
                                                 0, 2>(
      static_cast<int>(value));
}
inline bool NormKind_Parse(absl::string_view name, NormKind* value) {
  return ::google::protobuf::internal::ParseNamedEnum<NormKind>(
      NormKind_descriptor(), name, value);
}
enum FusedMHAKind : int {
  BMM1_OUTPUT_UNKNOWN = 0,
  BMM1_OUTPUT_INPUT_TYPE = 1,
  BMM1_OUTPUT_FLOAT = 2,
  FusedMHAKind_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  FusedMHAKind_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool FusedMHAKind_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t FusedMHAKind_internal_data_[];
constexpr FusedMHAKind FusedMHAKind_MIN = static_cast<FusedMHAKind>(0);
constexpr FusedMHAKind FusedMHAKind_MAX = static_cast<FusedMHAKind>(2);
constexpr int FusedMHAKind_ARRAYSIZE = 2 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
FusedMHAKind_descriptor();
template <typename T>
const std::string& FusedMHAKind_Name(T value) {
  static_assert(std::is_same<T, FusedMHAKind>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to FusedMHAKind_Name().");
  return FusedMHAKind_Name(static_cast<FusedMHAKind>(value));
}
template <>
inline const std::string& FusedMHAKind_Name(FusedMHAKind value) {
  return ::google::protobuf::internal::NameOfDenseEnum<FusedMHAKind_descriptor,
                                                 0, 2>(
      static_cast<int>(value));
}
inline bool FusedMHAKind_Parse(absl::string_view name, FusedMHAKind* value) {
  return ::google::protobuf::internal::ParseNamedEnum<FusedMHAKind>(
      FusedMHAKind_descriptor(), name, value);
}
enum FMHAMaskKind : int {
  NO_MASK = 0,
  PADDING = 1,
  CAUSAL = 2,
  PADDING_CAUSAL = 3,
  ALIBI = 4,
  FMHAMaskKind_INT_MIN_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::min(),
  FMHAMaskKind_INT_MAX_SENTINEL_DO_NOT_USE_ =
      std::numeric_limits<::int32_t>::max(),
};

PROTOBUF_EXPORT bool FMHAMaskKind_IsValid(int value);
PROTOBUF_EXPORT extern const uint32_t FMHAMaskKind_internal_data_[];
constexpr FMHAMaskKind FMHAMaskKind_MIN = static_cast<FMHAMaskKind>(0);
constexpr FMHAMaskKind FMHAMaskKind_MAX = static_cast<FMHAMaskKind>(4);
constexpr int FMHAMaskKind_ARRAYSIZE = 4 + 1;
PROTOBUF_EXPORT const ::google::protobuf::EnumDescriptor*
FMHAMaskKind_descriptor();
template <typename T>
const std::string& FMHAMaskKind_Name(T value) {
  static_assert(std::is_same<T, FMHAMaskKind>::value ||
                    std::is_integral<T>::value,
                "Incorrect type passed to FMHAMaskKind_Name().");
  return FMHAMaskKind_Name(static_cast<FMHAMaskKind>(value));
}
template <>
inline const std::string& FMHAMaskKind_Name(FMHAMaskKind value) {
  return ::google::protobuf::internal::NameOfDenseEnum<FMHAMaskKind_descriptor,
                                                 0, 4>(
      static_cast<int>(value));
}
inline bool FMHAMaskKind_Parse(absl::string_view name, FMHAMaskKind* value) {
  return ::google::protobuf::internal::ParseNamedEnum<FMHAMaskKind>(
      FMHAMaskKind_descriptor(), name, value);
}

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


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

class PROTOBUF_EXPORT TensorDescriptorProto final : public ::google::protobuf::Message
/* @@protoc_insertion_point(class_definition:stream_executor.dnn.TensorDescriptorProto) */ {
 public:
  inline TensorDescriptorProto() : TensorDescriptorProto(nullptr) {}
  ~TensorDescriptorProto() PROTOBUF_FINAL;
  template <typename = void>
  explicit PROTOBUF_CONSTEXPR TensorDescriptorProto(
      ::google::protobuf::internal::ConstantInitialized);

  inline TensorDescriptorProto(const TensorDescriptorProto& from) : TensorDescriptorProto(nullptr, from) {}
  inline TensorDescriptorProto(TensorDescriptorProto&& from) noexcept
      : TensorDescriptorProto(nullptr, std::move(from)) {}
  inline TensorDescriptorProto& operator=(const TensorDescriptorProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline TensorDescriptorProto& operator=(TensorDescriptorProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetArena() == from.GetArena()
#ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetArena() != nullptr
#endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance);
  }
  inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields()
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>();
  }

  static const ::google::protobuf::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::google::protobuf::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::google::protobuf::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const TensorDescriptorProto& default_instance() {
    return *internal_default_instance();
  }
  enum LayoutOneofCase {
    kDataLayout = 3,
    kFilterLayout = 4,
    LAYOUT_ONEOF_NOT_SET = 0,
  };
  static inline const TensorDescriptorProto* internal_default_instance() {
    return reinterpret_cast<const TensorDescriptorProto*>(
        &_TensorDescriptorProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages = 0;
  friend void swap(TensorDescriptorProto& a, TensorDescriptorProto& b) { a.Swap(&b); }
  inline void Swap(TensorDescriptorProto* other) {
    if (other == this) return;
#ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() != nullptr && GetArena() == other->GetArena()) {
#else   // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() == other->GetArena()) {
#endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::google::protobuf::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(TensorDescriptorProto* other) {
    if (other == this) return;
    ABSL_DCHECK(GetArena() == other->GetArena());
    InternalSwap(other);
  }

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

  TensorDescriptorProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL {
    return ::google::protobuf::Message::DefaultConstruct<TensorDescriptorProto>(arena);
  }
  using ::google::protobuf::Message::CopyFrom;
  void CopyFrom(const TensorDescriptorProto& from);
  using ::google::protobuf::Message::MergeFrom;
  void MergeFrom(const TensorDescriptorProto& from) { TensorDescriptorProto::MergeImpl(*this, from); }

  private:
  static void MergeImpl(
      ::google::protobuf::MessageLite& to_msg,
      const ::google::protobuf::MessageLite& from_msg);

  public:
  bool IsInitialized() const {
    return true;
  }
  ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL;
  #if defined(PROTOBUF_CUSTOM_VTABLE)
  private:
  static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg);
  static ::uint8_t* _InternalSerialize(
      const MessageLite& msg, ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream);

  public:
  ::size_t ByteSizeLong() const { return ByteSizeLong(*this); }
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const {
    return _InternalSerialize(*this, target, stream);
  }
  #else   // PROTOBUF_CUSTOM_VTABLE
  ::size_t ByteSizeLong() const final;
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const final;
  #endif  // PROTOBUF_CUSTOM_VTABLE
  int GetCachedSize() const { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::google::protobuf::Arena* arena);
  void SharedDtor();
  void InternalSwap(TensorDescriptorProto* other);
 private:
  friend class ::google::protobuf::internal::AnyMetadata;
  static ::absl::string_view FullMessageName() { return "stream_executor.dnn.TensorDescriptorProto"; }

 protected:
  explicit TensorDescriptorProto(::google::protobuf::Arena* arena);
  TensorDescriptorProto(::google::protobuf::Arena* arena, const TensorDescriptorProto& from);
  TensorDescriptorProto(::google::protobuf::Arena* arena, TensorDescriptorProto&& from) noexcept
      : TensorDescriptorProto(arena) {
    *this = ::std::move(from);
  }
  const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL;
  static const ::google::protobuf::Message::ClassDataFull _class_data_;

 public:
  ::google::protobuf::Metadata GetMetadata() const;
  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------
  enum : int {
    kDimensionsFieldNumber = 1,
    kDataTypeFieldNumber = 2,
    kDataLayoutFieldNumber = 3,
    kFilterLayoutFieldNumber = 4,
  };
  // repeated int64 dimensions = 1;
  int dimensions_size() const;
  private:
  int _internal_dimensions_size() const;

  public:
  void clear_dimensions() ;
  ::int64_t dimensions(int index) const;
  void set_dimensions(int index, ::int64_t value);
  void add_dimensions(::int64_t value);
  const ::google::protobuf::RepeatedField<::int64_t>& dimensions() const;
  ::google::protobuf::RepeatedField<::int64_t>* mutable_dimensions();

  private:
  const ::google::protobuf::RepeatedField<::int64_t>& _internal_dimensions() const;
  ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_dimensions();

  public:
  // .stream_executor.dnn.DataType data_type = 2;
  void clear_data_type() ;
  ::stream_executor::dnn::DataType data_type() const;
  void set_data_type(::stream_executor::dnn::DataType value);

  private:
  ::stream_executor::dnn::DataType _internal_data_type() const;
  void _internal_set_data_type(::stream_executor::dnn::DataType value);

  public:
  // .stream_executor.dnn.DataLayout data_layout = 3;
  bool has_data_layout() const;
  void clear_data_layout() ;
  ::stream_executor::dnn::DataLayout data_layout() const;
  void set_data_layout(::stream_executor::dnn::DataLayout value);

  private:
  ::stream_executor::dnn::DataLayout _internal_data_layout() const;
  void _internal_set_data_layout(::stream_executor::dnn::DataLayout value);

  public:
  // .stream_executor.dnn.FilterLayout filter_layout = 4;
  bool has_filter_layout() const;
  void clear_filter_layout() ;
  ::stream_executor::dnn::FilterLayout filter_layout() const;
  void set_filter_layout(::stream_executor::dnn::FilterLayout value);

  private:
  ::stream_executor::dnn::FilterLayout _internal_filter_layout() const;
  void _internal_set_filter_layout(::stream_executor::dnn::FilterLayout value);

  public:
  void clear_layout_oneof();
  LayoutOneofCase layout_oneof_case() const;
  // @@protoc_insertion_point(class_scope:stream_executor.dnn.TensorDescriptorProto)
 private:
  class _Internal;
  void set_has_data_layout();
  void set_has_filter_layout();
  inline bool has_layout_oneof() const;
  inline void clear_has_layout_oneof();
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<
      1, 4, 0,
      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* arena);
    inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility,
                          ::google::protobuf::Arena* arena, const Impl_& from,
                          const TensorDescriptorProto& from_msg);
    ::google::protobuf::RepeatedField<::int64_t> dimensions_;
    mutable ::google::protobuf::internal::CachedSize _dimensions_cached_byte_size_;
    int data_type_;
    union LayoutOneofUnion {
      constexpr LayoutOneofUnion() : _constinit_{} {}
      ::google::protobuf::internal::ConstantInitialized _constinit_;
      int data_layout_;
      int filter_layout_;
    } layout_oneof_;
    mutable ::google::protobuf::internal::CachedSize _cached_size_;
    ::uint32_t _oneof_case_[1];
    PROTOBUF_TSAN_DECLARE_MEMBER
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
};
// -------------------------------------------------------------------

class PROTOBUF_EXPORT ConvolutionDescriptorProto final : public ::google::protobuf::Message
/* @@protoc_insertion_point(class_definition:stream_executor.dnn.ConvolutionDescriptorProto) */ {
 public:
  inline ConvolutionDescriptorProto() : ConvolutionDescriptorProto(nullptr) {}
  ~ConvolutionDescriptorProto() PROTOBUF_FINAL;
  template <typename = void>
  explicit PROTOBUF_CONSTEXPR ConvolutionDescriptorProto(
      ::google::protobuf::internal::ConstantInitialized);

  inline ConvolutionDescriptorProto(const ConvolutionDescriptorProto& from) : ConvolutionDescriptorProto(nullptr, from) {}
  inline ConvolutionDescriptorProto(ConvolutionDescriptorProto&& from) noexcept
      : ConvolutionDescriptorProto(nullptr, std::move(from)) {}
  inline ConvolutionDescriptorProto& operator=(const ConvolutionDescriptorProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline ConvolutionDescriptorProto& operator=(ConvolutionDescriptorProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetArena() == from.GetArena()
#ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetArena() != nullptr
#endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance);
  }
  inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields()
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>();
  }

  static const ::google::protobuf::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::google::protobuf::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::google::protobuf::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const ConvolutionDescriptorProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const ConvolutionDescriptorProto* internal_default_instance() {
    return reinterpret_cast<const ConvolutionDescriptorProto*>(
        &_ConvolutionDescriptorProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages = 4;
  friend void swap(ConvolutionDescriptorProto& a, ConvolutionDescriptorProto& b) { a.Swap(&b); }
  inline void Swap(ConvolutionDescriptorProto* other) {
    if (other == this) return;
#ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() != nullptr && GetArena() == other->GetArena()) {
#else   // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() == other->GetArena()) {
#endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::google::protobuf::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(ConvolutionDescriptorProto* other) {
    if (other == this) return;
    ABSL_DCHECK(GetArena() == other->GetArena());
    InternalSwap(other);
  }

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

  ConvolutionDescriptorProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL {
    return ::google::protobuf::Message::DefaultConstruct<ConvolutionDescriptorProto>(arena);
  }
  using ::google::protobuf::Message::CopyFrom;
  void CopyFrom(const ConvolutionDescriptorProto& from);
  using ::google::protobuf::Message::MergeFrom;
  void MergeFrom(const ConvolutionDescriptorProto& from) { ConvolutionDescriptorProto::MergeImpl(*this, from); }

  private:
  static void MergeImpl(
      ::google::protobuf::MessageLite& to_msg,
      const ::google::protobuf::MessageLite& from_msg);

  public:
  bool IsInitialized() const {
    return true;
  }
  ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL;
  #if defined(PROTOBUF_CUSTOM_VTABLE)
  private:
  static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg);
  static ::uint8_t* _InternalSerialize(
      const MessageLite& msg, ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream);

  public:
  ::size_t ByteSizeLong() const { return ByteSizeLong(*this); }
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const {
    return _InternalSerialize(*this, target, stream);
  }
  #else   // PROTOBUF_CUSTOM_VTABLE
  ::size_t ByteSizeLong() const final;
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const final;
  #endif  // PROTOBUF_CUSTOM_VTABLE
  int GetCachedSize() const { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::google::protobuf::Arena* arena);
  void SharedDtor();
  void InternalSwap(ConvolutionDescriptorProto* other);
 private:
  friend class ::google::protobuf::internal::AnyMetadata;
  static ::absl::string_view FullMessageName() { return "stream_executor.dnn.ConvolutionDescriptorProto"; }

 protected:
  explicit ConvolutionDescriptorProto(::google::protobuf::Arena* arena);
  ConvolutionDescriptorProto(::google::protobuf::Arena* arena, const ConvolutionDescriptorProto& from);
  ConvolutionDescriptorProto(::google::protobuf::Arena* arena, ConvolutionDescriptorProto&& from) noexcept
      : ConvolutionDescriptorProto(arena) {
    *this = ::std::move(from);
  }
  const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL;
  static const ::google::protobuf::Message::ClassDataFull _class_data_;

 public:
  ::google::protobuf::Metadata GetMetadata() const;
  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------
  enum : int {
    kPaddingsFieldNumber = 1,
    kStridesFieldNumber = 2,
    kDilationsFieldNumber = 3,
    kNameFieldNumber = 7,
    kComputeModeFieldNumber = 4,
    kGroupCountFieldNumber = 5,
    kConvolutionModeFieldNumber = 6,
  };
  // repeated int64 paddings = 1;
  int paddings_size() const;
  private:
  int _internal_paddings_size() const;

  public:
  void clear_paddings() ;
  ::int64_t paddings(int index) const;
  void set_paddings(int index, ::int64_t value);
  void add_paddings(::int64_t value);
  const ::google::protobuf::RepeatedField<::int64_t>& paddings() const;
  ::google::protobuf::RepeatedField<::int64_t>* mutable_paddings();

  private:
  const ::google::protobuf::RepeatedField<::int64_t>& _internal_paddings() const;
  ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_paddings();

  public:
  // repeated int64 strides = 2;
  int strides_size() const;
  private:
  int _internal_strides_size() const;

  public:
  void clear_strides() ;
  ::int64_t strides(int index) const;
  void set_strides(int index, ::int64_t value);
  void add_strides(::int64_t value);
  const ::google::protobuf::RepeatedField<::int64_t>& strides() const;
  ::google::protobuf::RepeatedField<::int64_t>* mutable_strides();

  private:
  const ::google::protobuf::RepeatedField<::int64_t>& _internal_strides() const;
  ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_strides();

  public:
  // repeated int64 dilations = 3;
  int dilations_size() const;
  private:
  int _internal_dilations_size() const;

  public:
  void clear_dilations() ;
  ::int64_t dilations(int index) const;
  void set_dilations(int index, ::int64_t value);
  void add_dilations(::int64_t value);
  const ::google::protobuf::RepeatedField<::int64_t>& dilations() const;
  ::google::protobuf::RepeatedField<::int64_t>* mutable_dilations();

  private:
  const ::google::protobuf::RepeatedField<::int64_t>& _internal_dilations() const;
  ::google::protobuf::RepeatedField<::int64_t>* _internal_mutable_dilations();

  public:
  // string name = 7;
  void clear_name() ;
  const std::string& name() const;
  template <typename Arg_ = const std::string&, typename... Args_>
  void set_name(Arg_&& arg, Args_... args);
  std::string* mutable_name();
  PROTOBUF_NODISCARD std::string* release_name();
  void set_allocated_name(std::string* value);

  private:
  const std::string& _internal_name() const;
  inline PROTOBUF_ALWAYS_INLINE void _internal_set_name(
      const std::string& value);
  std::string* _internal_mutable_name();

  public:
  // .stream_executor.dnn.DataType compute_mode = 4;
  void clear_compute_mode() ;
  ::stream_executor::dnn::DataType compute_mode() const;
  void set_compute_mode(::stream_executor::dnn::DataType value);

  private:
  ::stream_executor::dnn::DataType _internal_compute_mode() const;
  void _internal_set_compute_mode(::stream_executor::dnn::DataType value);

  public:
  // int32 group_count = 5;
  void clear_group_count() ;
  ::int32_t group_count() const;
  void set_group_count(::int32_t value);

  private:
  ::int32_t _internal_group_count() const;
  void _internal_set_group_count(::int32_t value);

  public:
  // .stream_executor.dnn.ConvolutionMode convolution_mode = 6;
  void clear_convolution_mode() ;
  ::stream_executor::dnn::ConvolutionMode convolution_mode() const;
  void set_convolution_mode(::stream_executor::dnn::ConvolutionMode value);

  private:
  ::stream_executor::dnn::ConvolutionMode _internal_convolution_mode() const;
  void _internal_set_convolution_mode(::stream_executor::dnn::ConvolutionMode value);

  public:
  // @@protoc_insertion_point(class_scope:stream_executor.dnn.ConvolutionDescriptorProto)
 private:
  class _Internal;
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<
      3, 7, 0,
      59, 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* arena);
    inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility,
                          ::google::protobuf::Arena* arena, const Impl_& from,
                          const ConvolutionDescriptorProto& from_msg);
    ::google::protobuf::RepeatedField<::int64_t> paddings_;
    mutable ::google::protobuf::internal::CachedSize _paddings_cached_byte_size_;
    ::google::protobuf::RepeatedField<::int64_t> strides_;
    mutable ::google::protobuf::internal::CachedSize _strides_cached_byte_size_;
    ::google::protobuf::RepeatedField<::int64_t> dilations_;
    mutable ::google::protobuf::internal::CachedSize _dilations_cached_byte_size_;
    ::google::protobuf::internal::ArenaStringPtr name_;
    int compute_mode_;
    ::int32_t group_count_;
    int convolution_mode_;
    mutable ::google::protobuf::internal::CachedSize _cached_size_;
    PROTOBUF_TSAN_DECLARE_MEMBER
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
};
// -------------------------------------------------------------------

class AlgorithmProto_TuningKnobsEntry_DoNotUse final
    : public ::google::protobuf::internal::MapEntry<
          AlgorithmProto_TuningKnobsEntry_DoNotUse, ::int64_t, ::int64_t,
          ::google::protobuf::internal::WireFormatLite::TYPE_INT64,
          ::google::protobuf::internal::WireFormatLite::TYPE_INT64> {
 public:
  using SuperType = ::google::protobuf::internal::MapEntry<
      AlgorithmProto_TuningKnobsEntry_DoNotUse, ::int64_t, ::int64_t,
      ::google::protobuf::internal::WireFormatLite::TYPE_INT64,
      ::google::protobuf::internal::WireFormatLite::TYPE_INT64>;
  AlgorithmProto_TuningKnobsEntry_DoNotUse();
  template <typename = void>
  explicit PROTOBUF_CONSTEXPR AlgorithmProto_TuningKnobsEntry_DoNotUse(
      ::google::protobuf::internal::ConstantInitialized);
  explicit AlgorithmProto_TuningKnobsEntry_DoNotUse(::google::protobuf::Arena* arena);
  static const AlgorithmProto_TuningKnobsEntry_DoNotUse* internal_default_instance() {
    return reinterpret_cast<const AlgorithmProto_TuningKnobsEntry_DoNotUse*>(
        &_AlgorithmProto_TuningKnobsEntry_DoNotUse_default_instance_);
  }


 private:
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<
      1, 2, 0,
      0, 2>
      _table_;

  const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL;
  static const ::google::protobuf::Message::ClassDataFull _class_data_;
  friend struct ::TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
};
// -------------------------------------------------------------------

class PROTOBUF_EXPORT AlgorithmProto final : public ::google::protobuf::Message
/* @@protoc_insertion_point(class_definition:stream_executor.dnn.AlgorithmProto) */ {
 public:
  inline AlgorithmProto() : AlgorithmProto(nullptr) {}
  ~AlgorithmProto() PROTOBUF_FINAL;
  template <typename = void>
  explicit PROTOBUF_CONSTEXPR AlgorithmProto(
      ::google::protobuf::internal::ConstantInitialized);

  inline AlgorithmProto(const AlgorithmProto& from) : AlgorithmProto(nullptr, from) {}
  inline AlgorithmProto(AlgorithmProto&& from) noexcept
      : AlgorithmProto(nullptr, std::move(from)) {}
  inline AlgorithmProto& operator=(const AlgorithmProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline AlgorithmProto& operator=(AlgorithmProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetArena() == from.GetArena()
#ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetArena() != nullptr
#endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance);
  }
  inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields()
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>();
  }

  static const ::google::protobuf::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::google::protobuf::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::google::protobuf::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const AlgorithmProto& default_instance() {
    return *internal_default_instance();
  }
  static inline const AlgorithmProto* internal_default_instance() {
    return reinterpret_cast<const AlgorithmProto*>(
        &_AlgorithmProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages = 2;
  friend void swap(AlgorithmProto& a, AlgorithmProto& b) { a.Swap(&b); }
  inline void Swap(AlgorithmProto* other) {
    if (other == this) return;
#ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() != nullptr && GetArena() == other->GetArena()) {
#else   // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() == other->GetArena()) {
#endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::google::protobuf::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(AlgorithmProto* other) {
    if (other == this) return;
    ABSL_DCHECK(GetArena() == other->GetArena());
    InternalSwap(other);
  }

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

  AlgorithmProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL {
    return ::google::protobuf::Message::DefaultConstruct<AlgorithmProto>(arena);
  }
  using ::google::protobuf::Message::CopyFrom;
  void CopyFrom(const AlgorithmProto& from);
  using ::google::protobuf::Message::MergeFrom;
  void MergeFrom(const AlgorithmProto& from) { AlgorithmProto::MergeImpl(*this, from); }

  private:
  static void MergeImpl(
      ::google::protobuf::MessageLite& to_msg,
      const ::google::protobuf::MessageLite& from_msg);

  public:
  bool IsInitialized() const {
    return true;
  }
  ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL;
  #if defined(PROTOBUF_CUSTOM_VTABLE)
  private:
  static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg);
  static ::uint8_t* _InternalSerialize(
      const MessageLite& msg, ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream);

  public:
  ::size_t ByteSizeLong() const { return ByteSizeLong(*this); }
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const {
    return _InternalSerialize(*this, target, stream);
  }
  #else   // PROTOBUF_CUSTOM_VTABLE
  ::size_t ByteSizeLong() const final;
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const final;
  #endif  // PROTOBUF_CUSTOM_VTABLE
  int GetCachedSize() const { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::google::protobuf::Arena* arena);
  void SharedDtor();
  void InternalSwap(AlgorithmProto* other);
 private:
  friend class ::google::protobuf::internal::AnyMetadata;
  static ::absl::string_view FullMessageName() { return "stream_executor.dnn.AlgorithmProto"; }

 protected:
  explicit AlgorithmProto(::google::protobuf::Arena* arena);
  AlgorithmProto(::google::protobuf::Arena* arena, const AlgorithmProto& from);
  AlgorithmProto(::google::protobuf::Arena* arena, AlgorithmProto&& from) noexcept
      : AlgorithmProto(arena) {
    *this = ::std::move(from);
  }
  const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL;
  static const ::google::protobuf::Message::ClassDataFull _class_data_;

 public:
  ::google::protobuf::Metadata GetMetadata() const;
  // nested types ----------------------------------------------------
  using MathType = AlgorithmProto_MathType;
  static constexpr MathType DEFAULT_MATH = AlgorithmProto_MathType_DEFAULT_MATH;
  static constexpr MathType TENSOR_OP_MATH = AlgorithmProto_MathType_TENSOR_OP_MATH;
  static inline bool MathType_IsValid(int value) {
    return AlgorithmProto_MathType_IsValid(value);
  }
  static constexpr MathType MathType_MIN = AlgorithmProto_MathType_MathType_MIN;
  static constexpr MathType MathType_MAX = AlgorithmProto_MathType_MathType_MAX;
  static constexpr int MathType_ARRAYSIZE = AlgorithmProto_MathType_MathType_ARRAYSIZE;
  static inline const ::google::protobuf::EnumDescriptor* MathType_descriptor() {
    return AlgorithmProto_MathType_descriptor();
  }
  template <typename T>
  static inline const std::string& MathType_Name(T value) {
    return AlgorithmProto_MathType_Name(value);
  }
  static inline bool MathType_Parse(absl::string_view name, MathType* value) {
    return AlgorithmProto_MathType_Parse(name, value);
  }

  // accessors -------------------------------------------------------
  enum : int {
    kTuningKnobsFieldNumber = 4,
    kWorkspaceSizeFieldNumber = 6,
    kAlgoIdFieldNumber = 1,
    kMathTypeFieldNumber = 2,
    kIsCudnnFrontendFieldNumber = 5,
  };
  // map<int64, int64> tuning_knobs = 4;
  int tuning_knobs_size() const;
  private:
  int _internal_tuning_knobs_size() const;

  public:
  void clear_tuning_knobs() ;
  const ::google::protobuf::Map<::int64_t, ::int64_t>& tuning_knobs() const;
  ::google::protobuf::Map<::int64_t, ::int64_t>* mutable_tuning_knobs();

  private:
  const ::google::protobuf::Map<::int64_t, ::int64_t>& _internal_tuning_knobs() const;
  ::google::protobuf::Map<::int64_t, ::int64_t>* _internal_mutable_tuning_knobs();

  public:
  // .google.protobuf.UInt64Value workspace_size = 6;
  bool has_workspace_size() const;
  void clear_workspace_size() ;
  const ::google::protobuf::UInt64Value& workspace_size() const;
  PROTOBUF_NODISCARD ::google::protobuf::UInt64Value* release_workspace_size();
  ::google::protobuf::UInt64Value* mutable_workspace_size();
  void set_allocated_workspace_size(::google::protobuf::UInt64Value* value);
  void unsafe_arena_set_allocated_workspace_size(::google::protobuf::UInt64Value* value);
  ::google::protobuf::UInt64Value* unsafe_arena_release_workspace_size();

  private:
  const ::google::protobuf::UInt64Value& _internal_workspace_size() const;
  ::google::protobuf::UInt64Value* _internal_mutable_workspace_size();

  public:
  // int64 algo_id = 1;
  void clear_algo_id() ;
  ::int64_t algo_id() const;
  void set_algo_id(::int64_t value);

  private:
  ::int64_t _internal_algo_id() const;
  void _internal_set_algo_id(::int64_t value);

  public:
  // .stream_executor.dnn.AlgorithmProto.MathType math_type = 2;
  void clear_math_type() ;
  ::stream_executor::dnn::AlgorithmProto_MathType math_type() const;
  void set_math_type(::stream_executor::dnn::AlgorithmProto_MathType value);

  private:
  ::stream_executor::dnn::AlgorithmProto_MathType _internal_math_type() const;
  void _internal_set_math_type(::stream_executor::dnn::AlgorithmProto_MathType value);

  public:
  // bool is_cudnn_frontend = 5 [deprecated = true];
  [[deprecated]]  void clear_is_cudnn_frontend() ;
  [[deprecated]] bool is_cudnn_frontend() const;
  [[deprecated]] void set_is_cudnn_frontend(bool value);

  private:
  bool _internal_is_cudnn_frontend() const;
  void _internal_set_is_cudnn_frontend(bool value);

  public:
  // @@protoc_insertion_point(class_scope:stream_executor.dnn.AlgorithmProto)
 private:
  class _Internal;
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<
      3, 5, 2,
      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* arena);
    inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility,
                          ::google::protobuf::Arena* arena, const Impl_& from,
                          const AlgorithmProto& from_msg);
    ::google::protobuf::internal::HasBits<1> _has_bits_;
    mutable ::google::protobuf::internal::CachedSize _cached_size_;
    ::google::protobuf::internal::MapField<AlgorithmProto_TuningKnobsEntry_DoNotUse, ::int64_t, ::int64_t,
                      ::google::protobuf::internal::WireFormatLite::TYPE_INT64,
                      ::google::protobuf::internal::WireFormatLite::TYPE_INT64>
        tuning_knobs_;
    ::google::protobuf::UInt64Value* workspace_size_;
    ::int64_t algo_id_;
    int math_type_;
    bool is_cudnn_frontend_;
    PROTOBUF_TSAN_DECLARE_MEMBER
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
};
// -------------------------------------------------------------------

class PROTOBUF_EXPORT AlgorithmConfigProto final : public ::google::protobuf::Message
/* @@protoc_insertion_point(class_definition:stream_executor.dnn.AlgorithmConfigProto) */ {
 public:
  inline AlgorithmConfigProto() : AlgorithmConfigProto(nullptr) {}
  ~AlgorithmConfigProto() PROTOBUF_FINAL;
  template <typename = void>
  explicit PROTOBUF_CONSTEXPR AlgorithmConfigProto(
      ::google::protobuf::internal::ConstantInitialized);

  inline AlgorithmConfigProto(const AlgorithmConfigProto& from) : AlgorithmConfigProto(nullptr, from) {}
  inline AlgorithmConfigProto(AlgorithmConfigProto&& from) noexcept
      : AlgorithmConfigProto(nullptr, std::move(from)) {}
  inline AlgorithmConfigProto& operator=(const AlgorithmConfigProto& from) {
    CopyFrom(from);
    return *this;
  }
  inline AlgorithmConfigProto& operator=(AlgorithmConfigProto&& from) noexcept {
    if (this == &from) return *this;
    if (GetArena() == from.GetArena()
#ifdef PROTOBUF_FORCE_COPY_IN_MOVE
        && GetArena() != nullptr
#endif  // !PROTOBUF_FORCE_COPY_IN_MOVE
    ) {
      InternalSwap(&from);
    } else {
      CopyFrom(from);
    }
    return *this;
  }

  inline const ::google::protobuf::UnknownFieldSet& unknown_fields() const
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.unknown_fields<::google::protobuf::UnknownFieldSet>(::google::protobuf::UnknownFieldSet::default_instance);
  }
  inline ::google::protobuf::UnknownFieldSet* mutable_unknown_fields()
      ABSL_ATTRIBUTE_LIFETIME_BOUND {
    return _internal_metadata_.mutable_unknown_fields<::google::protobuf::UnknownFieldSet>();
  }

  static const ::google::protobuf::Descriptor* descriptor() {
    return GetDescriptor();
  }
  static const ::google::protobuf::Descriptor* GetDescriptor() {
    return default_instance().GetMetadata().descriptor;
  }
  static const ::google::protobuf::Reflection* GetReflection() {
    return default_instance().GetMetadata().reflection;
  }
  static const AlgorithmConfigProto& default_instance() {
    return *internal_default_instance();
  }
  enum OptionalAlgorithmCase {
    kAlgorithm = 1,
    OPTIONAL_ALGORITHM_NOT_SET = 0,
  };
  enum OptionalAlgorithmNoScratchCase {
    kAlgorithmNoScratch = 2,
    OPTIONAL_ALGORITHM_NO_SCRATCH_NOT_SET = 0,
  };
  enum OptionalScratchSizeCase {
    kScratchSize = 3,
    OPTIONAL_SCRATCH_SIZE_NOT_SET = 0,
  };
  static inline const AlgorithmConfigProto* internal_default_instance() {
    return reinterpret_cast<const AlgorithmConfigProto*>(
        &_AlgorithmConfigProto_default_instance_);
  }
  static constexpr int kIndexInFileMessages = 3;
  friend void swap(AlgorithmConfigProto& a, AlgorithmConfigProto& b) { a.Swap(&b); }
  inline void Swap(AlgorithmConfigProto* other) {
    if (other == this) return;
#ifdef PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() != nullptr && GetArena() == other->GetArena()) {
#else   // PROTOBUF_FORCE_COPY_IN_SWAP
    if (GetArena() == other->GetArena()) {
#endif  // !PROTOBUF_FORCE_COPY_IN_SWAP
      InternalSwap(other);
    } else {
      ::google::protobuf::internal::GenericSwap(this, other);
    }
  }
  void UnsafeArenaSwap(AlgorithmConfigProto* other) {
    if (other == this) return;
    ABSL_DCHECK(GetArena() == other->GetArena());
    InternalSwap(other);
  }

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

  AlgorithmConfigProto* New(::google::protobuf::Arena* arena = nullptr) const PROTOBUF_FINAL {
    return ::google::protobuf::Message::DefaultConstruct<AlgorithmConfigProto>(arena);
  }
  using ::google::protobuf::Message::CopyFrom;
  void CopyFrom(const AlgorithmConfigProto& from);
  using ::google::protobuf::Message::MergeFrom;
  void MergeFrom(const AlgorithmConfigProto& from) { AlgorithmConfigProto::MergeImpl(*this, from); }

  private:
  static void MergeImpl(
      ::google::protobuf::MessageLite& to_msg,
      const ::google::protobuf::MessageLite& from_msg);

  public:
  bool IsInitialized() const {
    return true;
  }
  ABSL_ATTRIBUTE_REINITIALIZES void Clear() PROTOBUF_FINAL;
  #if defined(PROTOBUF_CUSTOM_VTABLE)
  private:
  static ::size_t ByteSizeLong(const ::google::protobuf::MessageLite& msg);
  static ::uint8_t* _InternalSerialize(
      const MessageLite& msg, ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream);

  public:
  ::size_t ByteSizeLong() const { return ByteSizeLong(*this); }
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const {
    return _InternalSerialize(*this, target, stream);
  }
  #else   // PROTOBUF_CUSTOM_VTABLE
  ::size_t ByteSizeLong() const final;
  ::uint8_t* _InternalSerialize(
      ::uint8_t* target,
      ::google::protobuf::io::EpsCopyOutputStream* stream) const final;
  #endif  // PROTOBUF_CUSTOM_VTABLE
  int GetCachedSize() const { return _impl_._cached_size_.Get(); }

  private:
  void SharedCtor(::google::protobuf::Arena* arena);
  void SharedDtor();
  void InternalSwap(AlgorithmConfigProto* other);
 private:
  friend class ::google::protobuf::internal::AnyMetadata;
  static ::absl::string_view FullMessageName() { return "stream_executor.dnn.AlgorithmConfigProto"; }

 protected:
  explicit AlgorithmConfigProto(::google::protobuf::Arena* arena);
  AlgorithmConfigProto(::google::protobuf::Arena* arena, const AlgorithmConfigProto& from);
  AlgorithmConfigProto(::google::protobuf::Arena* arena, AlgorithmConfigProto&& from) noexcept
      : AlgorithmConfigProto(arena) {
    *this = ::std::move(from);
  }
  const ::google::protobuf::Message::ClassData* GetClassData() const PROTOBUF_FINAL;
  static const ::google::protobuf::Message::ClassDataFull _class_data_;

 public:
  ::google::protobuf::Metadata GetMetadata() const;
  // nested types ----------------------------------------------------

  // accessors -------------------------------------------------------
  enum : int {
    kAlgorithmFieldNumber = 1,
    kAlgorithmNoScratchFieldNumber = 2,
    kScratchSizeFieldNumber = 3,
  };
  // .stream_executor.dnn.AlgorithmProto algorithm = 1;
  bool has_algorithm() const;
  private:
  bool _internal_has_algorithm() const;

  public:
  void clear_algorithm() ;
  const ::stream_executor::dnn::AlgorithmProto& algorithm() const;
  PROTOBUF_NODISCARD ::stream_executor::dnn::AlgorithmProto* release_algorithm();
  ::stream_executor::dnn::AlgorithmProto* mutable_algorithm();
  void set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value);
  void unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value);
  ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm();

  private:
  const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm() const;
  ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm();

  public:
  // .stream_executor.dnn.AlgorithmProto algorithm_no_scratch = 2;
  bool has_algorithm_no_scratch() const;
  private:
  bool _internal_has_algorithm_no_scratch() const;

  public:
  void clear_algorithm_no_scratch() ;
  const ::stream_executor::dnn::AlgorithmProto& algorithm_no_scratch() const;
  PROTOBUF_NODISCARD ::stream_executor::dnn::AlgorithmProto* release_algorithm_no_scratch();
  ::stream_executor::dnn::AlgorithmProto* mutable_algorithm_no_scratch();
  void set_allocated_algorithm_no_scratch(::stream_executor::dnn::AlgorithmProto* value);
  void unsafe_arena_set_allocated_algorithm_no_scratch(::stream_executor::dnn::AlgorithmProto* value);
  ::stream_executor::dnn::AlgorithmProto* unsafe_arena_release_algorithm_no_scratch();

  private:
  const ::stream_executor::dnn::AlgorithmProto& _internal_algorithm_no_scratch() const;
  ::stream_executor::dnn::AlgorithmProto* _internal_mutable_algorithm_no_scratch();

  public:
  // int64 scratch_size = 3;
  bool has_scratch_size() const;
  void clear_scratch_size() ;
  ::int64_t scratch_size() const;
  void set_scratch_size(::int64_t value);

  private:
  ::int64_t _internal_scratch_size() const;
  void _internal_set_scratch_size(::int64_t value);

  public:
  void clear_optional_algorithm();
  OptionalAlgorithmCase optional_algorithm_case() const;
  void clear_optional_algorithm_no_scratch();
  OptionalAlgorithmNoScratchCase optional_algorithm_no_scratch_case() const;
  void clear_optional_scratch_size();
  OptionalScratchSizeCase optional_scratch_size_case() const;
  // @@protoc_insertion_point(class_scope:stream_executor.dnn.AlgorithmConfigProto)
 private:
  class _Internal;
  void set_has_algorithm();
  void set_has_algorithm_no_scratch();
  void set_has_scratch_size();
  inline bool has_optional_algorithm() const;
  inline void clear_has_optional_algorithm();
  inline bool has_optional_algorithm_no_scratch() const;
  inline void clear_has_optional_algorithm_no_scratch();
  inline bool has_optional_scratch_size() const;
  inline void clear_has_optional_scratch_size();
  friend class ::google::protobuf::internal::TcParser;
  static const ::google::protobuf::internal::TcParseTable<
      0, 3, 2,
      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* arena);
    inline explicit Impl_(::google::protobuf::internal::InternalVisibility visibility,
                          ::google::protobuf::Arena* arena, const Impl_& from,
                          const AlgorithmConfigProto& from_msg);
    union OptionalAlgorithmUnion {
      constexpr OptionalAlgorithmUnion() : _constinit_{} {}
      ::google::protobuf::internal::ConstantInitialized _constinit_;
      ::stream_executor::dnn::AlgorithmProto* algorithm_;
    } optional_algorithm_;
    union OptionalAlgorithmNoScratchUnion {
      constexpr OptionalAlgorithmNoScratchUnion() : _constinit_{} {}
      ::google::protobuf::internal::ConstantInitialized _constinit_;
      ::stream_executor::dnn::AlgorithmProto* algorithm_no_scratch_;
    } optional_algorithm_no_scratch_;
    union OptionalScratchSizeUnion {
      constexpr OptionalScratchSizeUnion() : _constinit_{} {}
      ::google::protobuf::internal::ConstantInitialized _constinit_;
      ::int64_t scratch_size_;
    } optional_scratch_size_;
    mutable ::google::protobuf::internal::CachedSize _cached_size_;
    ::uint32_t _oneof_case_[3];
    PROTOBUF_TSAN_DECLARE_MEMBER
  };
  union { Impl_ _impl_; };
  friend struct ::TableStruct_xla_2ftsl_2fprotobuf_2fdnn_2eproto;
};

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




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


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

// TensorDescriptorProto

// repeated int64 dimensions = 1;
inline int TensorDescriptorProto::_internal_dimensions_size() const {
  return _internal_dimensions().size();
}
inline int TensorDescriptorProto::dimensions_size() const {
  return _internal_dimensions_size();
}
inline void TensorDescriptorProto::clear_dimensions() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.dimensions_.Clear();
}
inline ::int64_t TensorDescriptorProto::dimensions(int index) const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.TensorDescriptorProto.dimensions)
  return _internal_dimensions().Get(index);
}
inline void TensorDescriptorProto::set_dimensions(int index, ::int64_t value) {
  _internal_mutable_dimensions()->Set(index, value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.TensorDescriptorProto.dimensions)
}
inline void TensorDescriptorProto::add_dimensions(::int64_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _internal_mutable_dimensions()->Add(value);
  // @@protoc_insertion_point(field_add:stream_executor.dnn.TensorDescriptorProto.dimensions)
}
inline const ::google::protobuf::RepeatedField<::int64_t>& TensorDescriptorProto::dimensions() const
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_list:stream_executor.dnn.TensorDescriptorProto.dimensions)
  return _internal_dimensions();
}
inline ::google::protobuf::RepeatedField<::int64_t>* TensorDescriptorProto::mutable_dimensions()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_mutable_list:stream_executor.dnn.TensorDescriptorProto.dimensions)
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _internal_mutable_dimensions();
}
inline const ::google::protobuf::RepeatedField<::int64_t>&
TensorDescriptorProto::_internal_dimensions() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.dimensions_;
}
inline ::google::protobuf::RepeatedField<::int64_t>* TensorDescriptorProto::_internal_mutable_dimensions() {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return &_impl_.dimensions_;
}

// .stream_executor.dnn.DataType data_type = 2;
inline void TensorDescriptorProto::clear_data_type() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.data_type_ = 0;
}
inline ::stream_executor::dnn::DataType TensorDescriptorProto::data_type() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.TensorDescriptorProto.data_type)
  return _internal_data_type();
}
inline void TensorDescriptorProto::set_data_type(::stream_executor::dnn::DataType value) {
  _internal_set_data_type(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.TensorDescriptorProto.data_type)
}
inline ::stream_executor::dnn::DataType TensorDescriptorProto::_internal_data_type() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return static_cast<::stream_executor::dnn::DataType>(_impl_.data_type_);
}
inline void TensorDescriptorProto::_internal_set_data_type(::stream_executor::dnn::DataType value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.data_type_ = value;
}

// .stream_executor.dnn.DataLayout data_layout = 3;
inline bool TensorDescriptorProto::has_data_layout() const {
  return layout_oneof_case() == kDataLayout;
}
inline void TensorDescriptorProto::set_has_data_layout() {
  _impl_._oneof_case_[0] = kDataLayout;
}
inline void TensorDescriptorProto::clear_data_layout() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (layout_oneof_case() == kDataLayout) {
    _impl_.layout_oneof_.data_layout_ = 0;
    clear_has_layout_oneof();
  }
}
inline ::stream_executor::dnn::DataLayout TensorDescriptorProto::data_layout() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.TensorDescriptorProto.data_layout)
  return _internal_data_layout();
}
inline void TensorDescriptorProto::set_data_layout(::stream_executor::dnn::DataLayout value) {
  if (layout_oneof_case() != kDataLayout) {
    clear_layout_oneof();
    set_has_data_layout();
  }
  _impl_.layout_oneof_.data_layout_ = value;
  // @@protoc_insertion_point(field_set:stream_executor.dnn.TensorDescriptorProto.data_layout)
}
inline ::stream_executor::dnn::DataLayout TensorDescriptorProto::_internal_data_layout() const {
  if (layout_oneof_case() == kDataLayout) {
    return static_cast<::stream_executor::dnn::DataLayout>(_impl_.layout_oneof_.data_layout_);
  }
  return static_cast<::stream_executor::dnn::DataLayout>(0);
}

// .stream_executor.dnn.FilterLayout filter_layout = 4;
inline bool TensorDescriptorProto::has_filter_layout() const {
  return layout_oneof_case() == kFilterLayout;
}
inline void TensorDescriptorProto::set_has_filter_layout() {
  _impl_._oneof_case_[0] = kFilterLayout;
}
inline void TensorDescriptorProto::clear_filter_layout() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (layout_oneof_case() == kFilterLayout) {
    _impl_.layout_oneof_.filter_layout_ = 0;
    clear_has_layout_oneof();
  }
}
inline ::stream_executor::dnn::FilterLayout TensorDescriptorProto::filter_layout() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.TensorDescriptorProto.filter_layout)
  return _internal_filter_layout();
}
inline void TensorDescriptorProto::set_filter_layout(::stream_executor::dnn::FilterLayout value) {
  if (layout_oneof_case() != kFilterLayout) {
    clear_layout_oneof();
    set_has_filter_layout();
  }
  _impl_.layout_oneof_.filter_layout_ = value;
  // @@protoc_insertion_point(field_set:stream_executor.dnn.TensorDescriptorProto.filter_layout)
}
inline ::stream_executor::dnn::FilterLayout TensorDescriptorProto::_internal_filter_layout() const {
  if (layout_oneof_case() == kFilterLayout) {
    return static_cast<::stream_executor::dnn::FilterLayout>(_impl_.layout_oneof_.filter_layout_);
  }
  return static_cast<::stream_executor::dnn::FilterLayout>(0);
}

inline bool TensorDescriptorProto::has_layout_oneof() const {
  return layout_oneof_case() != LAYOUT_ONEOF_NOT_SET;
}
inline void TensorDescriptorProto::clear_has_layout_oneof() {
  _impl_._oneof_case_[0] = LAYOUT_ONEOF_NOT_SET;
}
inline TensorDescriptorProto::LayoutOneofCase TensorDescriptorProto::layout_oneof_case() const {
  return TensorDescriptorProto::LayoutOneofCase(_impl_._oneof_case_[0]);
}
// -------------------------------------------------------------------

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

// AlgorithmProto

// int64 algo_id = 1;
inline void AlgorithmProto::clear_algo_id() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.algo_id_ = ::int64_t{0};
}
inline ::int64_t AlgorithmProto::algo_id() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmProto.algo_id)
  return _internal_algo_id();
}
inline void AlgorithmProto::set_algo_id(::int64_t value) {
  _internal_set_algo_id(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.AlgorithmProto.algo_id)
}
inline ::int64_t AlgorithmProto::_internal_algo_id() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.algo_id_;
}
inline void AlgorithmProto::_internal_set_algo_id(::int64_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.algo_id_ = value;
}

// .stream_executor.dnn.AlgorithmProto.MathType math_type = 2;
inline void AlgorithmProto::clear_math_type() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.math_type_ = 0;
}
inline ::stream_executor::dnn::AlgorithmProto_MathType AlgorithmProto::math_type() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmProto.math_type)
  return _internal_math_type();
}
inline void AlgorithmProto::set_math_type(::stream_executor::dnn::AlgorithmProto_MathType value) {
  _internal_set_math_type(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.AlgorithmProto.math_type)
}
inline ::stream_executor::dnn::AlgorithmProto_MathType AlgorithmProto::_internal_math_type() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return static_cast<::stream_executor::dnn::AlgorithmProto_MathType>(_impl_.math_type_);
}
inline void AlgorithmProto::_internal_set_math_type(::stream_executor::dnn::AlgorithmProto_MathType value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.math_type_ = value;
}

// map<int64, int64> tuning_knobs = 4;
inline int AlgorithmProto::_internal_tuning_knobs_size() const {
  return _internal_tuning_knobs().size();
}
inline int AlgorithmProto::tuning_knobs_size() const {
  return _internal_tuning_knobs_size();
}
inline void AlgorithmProto::clear_tuning_knobs() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.tuning_knobs_.Clear();
}
inline const ::google::protobuf::Map<::int64_t, ::int64_t>& AlgorithmProto::_internal_tuning_knobs() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.tuning_knobs_.GetMap();
}
inline const ::google::protobuf::Map<::int64_t, ::int64_t>& AlgorithmProto::tuning_knobs() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_map:stream_executor.dnn.AlgorithmProto.tuning_knobs)
  return _internal_tuning_knobs();
}
inline ::google::protobuf::Map<::int64_t, ::int64_t>* AlgorithmProto::_internal_mutable_tuning_knobs() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _impl_.tuning_knobs_.MutableMap();
}
inline ::google::protobuf::Map<::int64_t, ::int64_t>* AlgorithmProto::mutable_tuning_knobs() ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_mutable_map:stream_executor.dnn.AlgorithmProto.tuning_knobs)
  return _internal_mutable_tuning_knobs();
}

// bool is_cudnn_frontend = 5 [deprecated = true];
inline void AlgorithmProto::clear_is_cudnn_frontend() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.is_cudnn_frontend_ = false;
}
inline bool AlgorithmProto::is_cudnn_frontend() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmProto.is_cudnn_frontend)
  return _internal_is_cudnn_frontend();
}
inline void AlgorithmProto::set_is_cudnn_frontend(bool value) {
  _internal_set_is_cudnn_frontend(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.AlgorithmProto.is_cudnn_frontend)
}
inline bool AlgorithmProto::_internal_is_cudnn_frontend() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.is_cudnn_frontend_;
}
inline void AlgorithmProto::_internal_set_is_cudnn_frontend(bool value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.is_cudnn_frontend_ = value;
}

// .google.protobuf.UInt64Value workspace_size = 6;
inline bool AlgorithmProto::has_workspace_size() const {
  bool value = (_impl_._has_bits_[0] & 0x00000001u) != 0;
  PROTOBUF_ASSUME(!value || _impl_.workspace_size_ != nullptr);
  return value;
}
inline const ::google::protobuf::UInt64Value& AlgorithmProto::_internal_workspace_size() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  const ::google::protobuf::UInt64Value* p = _impl_.workspace_size_;
  return p != nullptr ? *p : reinterpret_cast<const ::google::protobuf::UInt64Value&>(::google::protobuf::_UInt64Value_default_instance_);
}
inline const ::google::protobuf::UInt64Value& AlgorithmProto::workspace_size() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmProto.workspace_size)
  return _internal_workspace_size();
}
inline void AlgorithmProto::unsafe_arena_set_allocated_workspace_size(::google::protobuf::UInt64Value* value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (GetArena() == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.workspace_size_);
  }
  _impl_.workspace_size_ = reinterpret_cast<::google::protobuf::UInt64Value*>(value);
  if (value != nullptr) {
    _impl_._has_bits_[0] |= 0x00000001u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000001u;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:stream_executor.dnn.AlgorithmProto.workspace_size)
}
inline ::google::protobuf::UInt64Value* AlgorithmProto::release_workspace_size() {
  ::google::protobuf::internal::TSanWrite(&_impl_);

  _impl_._has_bits_[0] &= ~0x00000001u;
  ::google::protobuf::UInt64Value* released = _impl_.workspace_size_;
  _impl_.workspace_size_ = nullptr;
#ifdef PROTOBUF_FORCE_COPY_IN_RELEASE
  auto* old = reinterpret_cast<::google::protobuf::MessageLite*>(released);
  released = ::google::protobuf::internal::DuplicateIfNonNull(released);
  if (GetArena() == nullptr) {
    delete old;
  }
#else   // PROTOBUF_FORCE_COPY_IN_RELEASE
  if (GetArena() != nullptr) {
    released = ::google::protobuf::internal::DuplicateIfNonNull(released);
  }
#endif  // !PROTOBUF_FORCE_COPY_IN_RELEASE
  return released;
}
inline ::google::protobuf::UInt64Value* AlgorithmProto::unsafe_arena_release_workspace_size() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:stream_executor.dnn.AlgorithmProto.workspace_size)

  _impl_._has_bits_[0] &= ~0x00000001u;
  ::google::protobuf::UInt64Value* temp = _impl_.workspace_size_;
  _impl_.workspace_size_ = nullptr;
  return temp;
}
inline ::google::protobuf::UInt64Value* AlgorithmProto::_internal_mutable_workspace_size() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (_impl_.workspace_size_ == nullptr) {
    auto* p = ::google::protobuf::Message::DefaultConstruct<::google::protobuf::UInt64Value>(GetArena());
    _impl_.workspace_size_ = reinterpret_cast<::google::protobuf::UInt64Value*>(p);
  }
  return _impl_.workspace_size_;
}
inline ::google::protobuf::UInt64Value* AlgorithmProto::mutable_workspace_size() ABSL_ATTRIBUTE_LIFETIME_BOUND {
  _impl_._has_bits_[0] |= 0x00000001u;
  ::google::protobuf::UInt64Value* _msg = _internal_mutable_workspace_size();
  // @@protoc_insertion_point(field_mutable:stream_executor.dnn.AlgorithmProto.workspace_size)
  return _msg;
}
inline void AlgorithmProto::set_allocated_workspace_size(::google::protobuf::UInt64Value* value) {
  ::google::protobuf::Arena* message_arena = GetArena();
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (message_arena == nullptr) {
    delete reinterpret_cast<::google::protobuf::MessageLite*>(_impl_.workspace_size_);
  }

  if (value != nullptr) {
    ::google::protobuf::Arena* submessage_arena = reinterpret_cast<::google::protobuf::MessageLite*>(value)->GetArena();
    if (message_arena != submessage_arena) {
      value = ::google::protobuf::internal::GetOwnedMessage(message_arena, value, submessage_arena);
    }
    _impl_._has_bits_[0] |= 0x00000001u;
  } else {
    _impl_._has_bits_[0] &= ~0x00000001u;
  }

  _impl_.workspace_size_ = reinterpret_cast<::google::protobuf::UInt64Value*>(value);
  // @@protoc_insertion_point(field_set_allocated:stream_executor.dnn.AlgorithmProto.workspace_size)
}

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

// AlgorithmConfigProto

// .stream_executor.dnn.AlgorithmProto algorithm = 1;
inline bool AlgorithmConfigProto::has_algorithm() const {
  return optional_algorithm_case() == kAlgorithm;
}
inline bool AlgorithmConfigProto::_internal_has_algorithm() const {
  return optional_algorithm_case() == kAlgorithm;
}
inline void AlgorithmConfigProto::set_has_algorithm() {
  _impl_._oneof_case_[0] = kAlgorithm;
}
inline void AlgorithmConfigProto::clear_algorithm() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (optional_algorithm_case() == kAlgorithm) {
    if (GetArena() == nullptr) {
      delete _impl_.optional_algorithm_.algorithm_;
    } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) {
      ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.optional_algorithm_.algorithm_);
    }
    clear_has_optional_algorithm();
  }
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::release_algorithm() {
  // @@protoc_insertion_point(field_release:stream_executor.dnn.AlgorithmConfigProto.algorithm)
  if (optional_algorithm_case() == kAlgorithm) {
    clear_has_optional_algorithm();
    auto* temp = _impl_.optional_algorithm_.algorithm_;
    if (GetArena() != nullptr) {
      temp = ::google::protobuf::internal::DuplicateIfNonNull(temp);
    }
    _impl_.optional_algorithm_.algorithm_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::stream_executor::dnn::AlgorithmProto& AlgorithmConfigProto::_internal_algorithm() const {
  return optional_algorithm_case() == kAlgorithm ? *_impl_.optional_algorithm_.algorithm_ : reinterpret_cast<::stream_executor::dnn::AlgorithmProto&>(::stream_executor::dnn::_AlgorithmProto_default_instance_);
}
inline const ::stream_executor::dnn::AlgorithmProto& AlgorithmConfigProto::algorithm() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmConfigProto.algorithm)
  return _internal_algorithm();
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::unsafe_arena_release_algorithm() {
  // @@protoc_insertion_point(field_unsafe_arena_release:stream_executor.dnn.AlgorithmConfigProto.algorithm)
  if (optional_algorithm_case() == kAlgorithm) {
    clear_has_optional_algorithm();
    auto* temp = _impl_.optional_algorithm_.algorithm_;
    _impl_.optional_algorithm_.algorithm_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void AlgorithmConfigProto::unsafe_arena_set_allocated_algorithm(::stream_executor::dnn::AlgorithmProto* value) {
  // We rely on the oneof clear method to free the earlier contents
  // of this oneof. We can directly use the pointer we're given to
  // set the new value.
  clear_optional_algorithm();
  if (value) {
    set_has_algorithm();
    _impl_.optional_algorithm_.algorithm_ = value;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:stream_executor.dnn.AlgorithmConfigProto.algorithm)
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::_internal_mutable_algorithm() {
  if (optional_algorithm_case() != kAlgorithm) {
    clear_optional_algorithm();
    set_has_algorithm();
    _impl_.optional_algorithm_.algorithm_ =
        ::google::protobuf::Message::DefaultConstruct<::stream_executor::dnn::AlgorithmProto>(GetArena());
  }
  return _impl_.optional_algorithm_.algorithm_;
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::mutable_algorithm() ABSL_ATTRIBUTE_LIFETIME_BOUND {
  ::stream_executor::dnn::AlgorithmProto* _msg = _internal_mutable_algorithm();
  // @@protoc_insertion_point(field_mutable:stream_executor.dnn.AlgorithmConfigProto.algorithm)
  return _msg;
}

// .stream_executor.dnn.AlgorithmProto algorithm_no_scratch = 2;
inline bool AlgorithmConfigProto::has_algorithm_no_scratch() const {
  return optional_algorithm_no_scratch_case() == kAlgorithmNoScratch;
}
inline bool AlgorithmConfigProto::_internal_has_algorithm_no_scratch() const {
  return optional_algorithm_no_scratch_case() == kAlgorithmNoScratch;
}
inline void AlgorithmConfigProto::set_has_algorithm_no_scratch() {
  _impl_._oneof_case_[1] = kAlgorithmNoScratch;
}
inline void AlgorithmConfigProto::clear_algorithm_no_scratch() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (optional_algorithm_no_scratch_case() == kAlgorithmNoScratch) {
    if (GetArena() == nullptr) {
      delete _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_;
    } else if (::google::protobuf::internal::DebugHardenClearOneofMessageOnArena()) {
      ::google::protobuf::internal::MaybePoisonAfterClear(_impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_);
    }
    clear_has_optional_algorithm_no_scratch();
  }
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::release_algorithm_no_scratch() {
  // @@protoc_insertion_point(field_release:stream_executor.dnn.AlgorithmConfigProto.algorithm_no_scratch)
  if (optional_algorithm_no_scratch_case() == kAlgorithmNoScratch) {
    clear_has_optional_algorithm_no_scratch();
    auto* temp = _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_;
    if (GetArena() != nullptr) {
      temp = ::google::protobuf::internal::DuplicateIfNonNull(temp);
    }
    _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline const ::stream_executor::dnn::AlgorithmProto& AlgorithmConfigProto::_internal_algorithm_no_scratch() const {
  return optional_algorithm_no_scratch_case() == kAlgorithmNoScratch ? *_impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_ : reinterpret_cast<::stream_executor::dnn::AlgorithmProto&>(::stream_executor::dnn::_AlgorithmProto_default_instance_);
}
inline const ::stream_executor::dnn::AlgorithmProto& AlgorithmConfigProto::algorithm_no_scratch() const ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmConfigProto.algorithm_no_scratch)
  return _internal_algorithm_no_scratch();
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::unsafe_arena_release_algorithm_no_scratch() {
  // @@protoc_insertion_point(field_unsafe_arena_release:stream_executor.dnn.AlgorithmConfigProto.algorithm_no_scratch)
  if (optional_algorithm_no_scratch_case() == kAlgorithmNoScratch) {
    clear_has_optional_algorithm_no_scratch();
    auto* temp = _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_;
    _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_ = nullptr;
    return temp;
  } else {
    return nullptr;
  }
}
inline void AlgorithmConfigProto::unsafe_arena_set_allocated_algorithm_no_scratch(::stream_executor::dnn::AlgorithmProto* value) {
  // We rely on the oneof clear method to free the earlier contents
  // of this oneof. We can directly use the pointer we're given to
  // set the new value.
  clear_optional_algorithm_no_scratch();
  if (value) {
    set_has_algorithm_no_scratch();
    _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_ = value;
  }
  // @@protoc_insertion_point(field_unsafe_arena_set_allocated:stream_executor.dnn.AlgorithmConfigProto.algorithm_no_scratch)
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::_internal_mutable_algorithm_no_scratch() {
  if (optional_algorithm_no_scratch_case() != kAlgorithmNoScratch) {
    clear_optional_algorithm_no_scratch();
    set_has_algorithm_no_scratch();
    _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_ =
        ::google::protobuf::Message::DefaultConstruct<::stream_executor::dnn::AlgorithmProto>(GetArena());
  }
  return _impl_.optional_algorithm_no_scratch_.algorithm_no_scratch_;
}
inline ::stream_executor::dnn::AlgorithmProto* AlgorithmConfigProto::mutable_algorithm_no_scratch() ABSL_ATTRIBUTE_LIFETIME_BOUND {
  ::stream_executor::dnn::AlgorithmProto* _msg = _internal_mutable_algorithm_no_scratch();
  // @@protoc_insertion_point(field_mutable:stream_executor.dnn.AlgorithmConfigProto.algorithm_no_scratch)
  return _msg;
}

// int64 scratch_size = 3;
inline bool AlgorithmConfigProto::has_scratch_size() const {
  return optional_scratch_size_case() == kScratchSize;
}
inline void AlgorithmConfigProto::set_has_scratch_size() {
  _impl_._oneof_case_[2] = kScratchSize;
}
inline void AlgorithmConfigProto::clear_scratch_size() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  if (optional_scratch_size_case() == kScratchSize) {
    _impl_.optional_scratch_size_.scratch_size_ = ::int64_t{0};
    clear_has_optional_scratch_size();
  }
}
inline ::int64_t AlgorithmConfigProto::scratch_size() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.AlgorithmConfigProto.scratch_size)
  return _internal_scratch_size();
}
inline void AlgorithmConfigProto::set_scratch_size(::int64_t value) {
  if (optional_scratch_size_case() != kScratchSize) {
    clear_optional_scratch_size();
    set_has_scratch_size();
  }
  _impl_.optional_scratch_size_.scratch_size_ = value;
  // @@protoc_insertion_point(field_set:stream_executor.dnn.AlgorithmConfigProto.scratch_size)
}
inline ::int64_t AlgorithmConfigProto::_internal_scratch_size() const {
  if (optional_scratch_size_case() == kScratchSize) {
    return _impl_.optional_scratch_size_.scratch_size_;
  }
  return ::int64_t{0};
}

inline bool AlgorithmConfigProto::has_optional_algorithm() const {
  return optional_algorithm_case() != OPTIONAL_ALGORITHM_NOT_SET;
}
inline void AlgorithmConfigProto::clear_has_optional_algorithm() {
  _impl_._oneof_case_[0] = OPTIONAL_ALGORITHM_NOT_SET;
}
inline bool AlgorithmConfigProto::has_optional_algorithm_no_scratch() const {
  return optional_algorithm_no_scratch_case() != OPTIONAL_ALGORITHM_NO_SCRATCH_NOT_SET;
}
inline void AlgorithmConfigProto::clear_has_optional_algorithm_no_scratch() {
  _impl_._oneof_case_[1] = OPTIONAL_ALGORITHM_NO_SCRATCH_NOT_SET;
}
inline bool AlgorithmConfigProto::has_optional_scratch_size() const {
  return optional_scratch_size_case() != OPTIONAL_SCRATCH_SIZE_NOT_SET;
}
inline void AlgorithmConfigProto::clear_has_optional_scratch_size() {
  _impl_._oneof_case_[2] = OPTIONAL_SCRATCH_SIZE_NOT_SET;
}
inline AlgorithmConfigProto::OptionalAlgorithmCase AlgorithmConfigProto::optional_algorithm_case() const {
  return AlgorithmConfigProto::OptionalAlgorithmCase(_impl_._oneof_case_[0]);
}
inline AlgorithmConfigProto::OptionalAlgorithmNoScratchCase AlgorithmConfigProto::optional_algorithm_no_scratch_case() const {
  return AlgorithmConfigProto::OptionalAlgorithmNoScratchCase(_impl_._oneof_case_[1]);
}
inline AlgorithmConfigProto::OptionalScratchSizeCase AlgorithmConfigProto::optional_scratch_size_case() const {
  return AlgorithmConfigProto::OptionalScratchSizeCase(_impl_._oneof_case_[2]);
}
// -------------------------------------------------------------------

// ConvolutionDescriptorProto

// repeated int64 paddings = 1;
inline int ConvolutionDescriptorProto::_internal_paddings_size() const {
  return _internal_paddings().size();
}
inline int ConvolutionDescriptorProto::paddings_size() const {
  return _internal_paddings_size();
}
inline void ConvolutionDescriptorProto::clear_paddings() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.paddings_.Clear();
}
inline ::int64_t ConvolutionDescriptorProto::paddings(int index) const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.paddings)
  return _internal_paddings().Get(index);
}
inline void ConvolutionDescriptorProto::set_paddings(int index, ::int64_t value) {
  _internal_mutable_paddings()->Set(index, value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.paddings)
}
inline void ConvolutionDescriptorProto::add_paddings(::int64_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _internal_mutable_paddings()->Add(value);
  // @@protoc_insertion_point(field_add:stream_executor.dnn.ConvolutionDescriptorProto.paddings)
}
inline const ::google::protobuf::RepeatedField<::int64_t>& ConvolutionDescriptorProto::paddings() const
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_list:stream_executor.dnn.ConvolutionDescriptorProto.paddings)
  return _internal_paddings();
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::mutable_paddings()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_mutable_list:stream_executor.dnn.ConvolutionDescriptorProto.paddings)
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _internal_mutable_paddings();
}
inline const ::google::protobuf::RepeatedField<::int64_t>&
ConvolutionDescriptorProto::_internal_paddings() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.paddings_;
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::_internal_mutable_paddings() {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return &_impl_.paddings_;
}

// repeated int64 strides = 2;
inline int ConvolutionDescriptorProto::_internal_strides_size() const {
  return _internal_strides().size();
}
inline int ConvolutionDescriptorProto::strides_size() const {
  return _internal_strides_size();
}
inline void ConvolutionDescriptorProto::clear_strides() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.strides_.Clear();
}
inline ::int64_t ConvolutionDescriptorProto::strides(int index) const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.strides)
  return _internal_strides().Get(index);
}
inline void ConvolutionDescriptorProto::set_strides(int index, ::int64_t value) {
  _internal_mutable_strides()->Set(index, value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.strides)
}
inline void ConvolutionDescriptorProto::add_strides(::int64_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _internal_mutable_strides()->Add(value);
  // @@protoc_insertion_point(field_add:stream_executor.dnn.ConvolutionDescriptorProto.strides)
}
inline const ::google::protobuf::RepeatedField<::int64_t>& ConvolutionDescriptorProto::strides() const
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_list:stream_executor.dnn.ConvolutionDescriptorProto.strides)
  return _internal_strides();
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::mutable_strides()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_mutable_list:stream_executor.dnn.ConvolutionDescriptorProto.strides)
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _internal_mutable_strides();
}
inline const ::google::protobuf::RepeatedField<::int64_t>&
ConvolutionDescriptorProto::_internal_strides() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.strides_;
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::_internal_mutable_strides() {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return &_impl_.strides_;
}

// repeated int64 dilations = 3;
inline int ConvolutionDescriptorProto::_internal_dilations_size() const {
  return _internal_dilations().size();
}
inline int ConvolutionDescriptorProto::dilations_size() const {
  return _internal_dilations_size();
}
inline void ConvolutionDescriptorProto::clear_dilations() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.dilations_.Clear();
}
inline ::int64_t ConvolutionDescriptorProto::dilations(int index) const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.dilations)
  return _internal_dilations().Get(index);
}
inline void ConvolutionDescriptorProto::set_dilations(int index, ::int64_t value) {
  _internal_mutable_dilations()->Set(index, value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.dilations)
}
inline void ConvolutionDescriptorProto::add_dilations(::int64_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _internal_mutable_dilations()->Add(value);
  // @@protoc_insertion_point(field_add:stream_executor.dnn.ConvolutionDescriptorProto.dilations)
}
inline const ::google::protobuf::RepeatedField<::int64_t>& ConvolutionDescriptorProto::dilations() const
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_list:stream_executor.dnn.ConvolutionDescriptorProto.dilations)
  return _internal_dilations();
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::mutable_dilations()
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_mutable_list:stream_executor.dnn.ConvolutionDescriptorProto.dilations)
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _internal_mutable_dilations();
}
inline const ::google::protobuf::RepeatedField<::int64_t>&
ConvolutionDescriptorProto::_internal_dilations() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.dilations_;
}
inline ::google::protobuf::RepeatedField<::int64_t>* ConvolutionDescriptorProto::_internal_mutable_dilations() {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return &_impl_.dilations_;
}

// .stream_executor.dnn.DataType compute_mode = 4;
inline void ConvolutionDescriptorProto::clear_compute_mode() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.compute_mode_ = 0;
}
inline ::stream_executor::dnn::DataType ConvolutionDescriptorProto::compute_mode() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.compute_mode)
  return _internal_compute_mode();
}
inline void ConvolutionDescriptorProto::set_compute_mode(::stream_executor::dnn::DataType value) {
  _internal_set_compute_mode(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.compute_mode)
}
inline ::stream_executor::dnn::DataType ConvolutionDescriptorProto::_internal_compute_mode() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return static_cast<::stream_executor::dnn::DataType>(_impl_.compute_mode_);
}
inline void ConvolutionDescriptorProto::_internal_set_compute_mode(::stream_executor::dnn::DataType value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.compute_mode_ = value;
}

// int32 group_count = 5;
inline void ConvolutionDescriptorProto::clear_group_count() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.group_count_ = 0;
}
inline ::int32_t ConvolutionDescriptorProto::group_count() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.group_count)
  return _internal_group_count();
}
inline void ConvolutionDescriptorProto::set_group_count(::int32_t value) {
  _internal_set_group_count(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.group_count)
}
inline ::int32_t ConvolutionDescriptorProto::_internal_group_count() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.group_count_;
}
inline void ConvolutionDescriptorProto::_internal_set_group_count(::int32_t value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.group_count_ = value;
}

// .stream_executor.dnn.ConvolutionMode convolution_mode = 6;
inline void ConvolutionDescriptorProto::clear_convolution_mode() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.convolution_mode_ = 0;
}
inline ::stream_executor::dnn::ConvolutionMode ConvolutionDescriptorProto::convolution_mode() const {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.convolution_mode)
  return _internal_convolution_mode();
}
inline void ConvolutionDescriptorProto::set_convolution_mode(::stream_executor::dnn::ConvolutionMode value) {
  _internal_set_convolution_mode(value);
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.convolution_mode)
}
inline ::stream_executor::dnn::ConvolutionMode ConvolutionDescriptorProto::_internal_convolution_mode() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return static_cast<::stream_executor::dnn::ConvolutionMode>(_impl_.convolution_mode_);
}
inline void ConvolutionDescriptorProto::_internal_set_convolution_mode(::stream_executor::dnn::ConvolutionMode value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.convolution_mode_ = value;
}

// string name = 7;
inline void ConvolutionDescriptorProto::clear_name() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.name_.ClearToEmpty();
}
inline const std::string& ConvolutionDescriptorProto::name() const
    ABSL_ATTRIBUTE_LIFETIME_BOUND {
  // @@protoc_insertion_point(field_get:stream_executor.dnn.ConvolutionDescriptorProto.name)
  return _internal_name();
}
template <typename Arg_, typename... Args_>
inline PROTOBUF_ALWAYS_INLINE void ConvolutionDescriptorProto::set_name(Arg_&& arg,
                                                     Args_... args) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.name_.Set(static_cast<Arg_&&>(arg), args..., GetArena());
  // @@protoc_insertion_point(field_set:stream_executor.dnn.ConvolutionDescriptorProto.name)
}
inline std::string* ConvolutionDescriptorProto::mutable_name() ABSL_ATTRIBUTE_LIFETIME_BOUND {
  std::string* _s = _internal_mutable_name();
  // @@protoc_insertion_point(field_mutable:stream_executor.dnn.ConvolutionDescriptorProto.name)
  return _s;
}
inline const std::string& ConvolutionDescriptorProto::_internal_name() const {
  ::google::protobuf::internal::TSanRead(&_impl_);
  return _impl_.name_.Get();
}
inline void ConvolutionDescriptorProto::_internal_set_name(const std::string& value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.name_.Set(value, GetArena());
}
inline std::string* ConvolutionDescriptorProto::_internal_mutable_name() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  return _impl_.name_.Mutable( GetArena());
}
inline std::string* ConvolutionDescriptorProto::release_name() {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  // @@protoc_insertion_point(field_release:stream_executor.dnn.ConvolutionDescriptorProto.name)
  return _impl_.name_.Release();
}
inline void ConvolutionDescriptorProto::set_allocated_name(std::string* value) {
  ::google::protobuf::internal::TSanWrite(&_impl_);
  _impl_.name_.SetAllocated(value, GetArena());
  #ifdef PROTOBUF_FORCE_COPY_DEFAULT_STRING
        if (_impl_.name_.IsDefault()) {
          _impl_.name_.Set("", GetArena());
        }
  #endif  // PROTOBUF_FORCE_COPY_DEFAULT_STRING
  // @@protoc_insertion_point(field_set_allocated:stream_executor.dnn.ConvolutionDescriptorProto.name)
}

#ifdef __GNUC__
#pragma GCC diagnostic pop
#endif  // __GNUC__

// @@protoc_insertion_point(namespace_scope)
}  // namespace dnn
}  // namespace stream_executor


namespace google {
namespace protobuf {

template <>
struct is_proto_enum<::stream_executor::dnn::AlgorithmProto_MathType> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::AlgorithmProto_MathType>() {
  return ::stream_executor::dnn::AlgorithmProto_MathType_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::DataType> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::DataType>() {
  return ::stream_executor::dnn::DataType_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::DataLayout> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::DataLayout>() {
  return ::stream_executor::dnn::DataLayout_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::FilterLayout> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::FilterLayout>() {
  return ::stream_executor::dnn::FilterLayout_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::ActivationMode> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::ActivationMode>() {
  return ::stream_executor::dnn::ActivationMode_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::ConvolutionMode> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::ConvolutionMode>() {
  return ::stream_executor::dnn::ConvolutionMode_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::ConvolutionKind> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::ConvolutionKind>() {
  return ::stream_executor::dnn::ConvolutionKind_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::NormKind> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::NormKind>() {
  return ::stream_executor::dnn::NormKind_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::FusedMHAKind> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::FusedMHAKind>() {
  return ::stream_executor::dnn::FusedMHAKind_descriptor();
}
template <>
struct is_proto_enum<::stream_executor::dnn::FMHAMaskKind> : std::true_type {};
template <>
inline const EnumDescriptor* GetEnumDescriptor<::stream_executor::dnn::FMHAMaskKind>() {
  return ::stream_executor::dnn::FMHAMaskKind_descriptor();
}

}  // namespace protobuf
}  // namespace google

// @@protoc_insertion_point(global_scope)

#include "google/protobuf/port_undef.inc"

#endif  // GOOGLE_PROTOBUF_INCLUDED_xla_2ftsl_2fprotobuf_2fdnn_2eproto_2epb_2eh