/* Autogenerated by mlir-tblgen; don't manually edit */

#ifdef GEN_PASS_DECL
// Generate declarations for all passes.
#define GEN_PASS_DECL_CONVERTPURECALLOPSPASS
#define GEN_PASS_DECL_ERASEDEADFUNCTIONSPASS
#define GEN_PASS_DECL_EXPANDFLOATOPSPASS
#define GEN_PASS_DECL_EXPANDINTEGERPOWERPASS
#define GEN_PASS_DECL_FLATTENTENSORSPASS
#define GEN_PASS_DECL_LOWERTENSORSPASS
#define GEN_PASS_DECL_LOWERXLAINTRINSICLIBPASS
#define GEN_PASS_DECL_LOWERXLALOOPSTOSCFPASS
#define GEN_PASS_DECL_LOWERXLATOSCFPASS
#define GEN_PASS_DECL_MERGEPOINTERSTOSAMESLICEPASS
#define GEN_PASS_DECL_PROPAGATESLICEINDICESPASS
#define GEN_PASS_DECL_SAFEINTEGERARITHMETICPASS
#define GEN_PASS_DECL_SIMPLIFYAFFINEPASS
#define GEN_PASS_DECL_SIMPLIFYARITHPASS
#define GEN_PASS_DECL_UNSWITCHLOOPSPASS
#define GEN_PASS_DECL_VECTORIZELOADSANDSTORESPASS
#undef GEN_PASS_DECL
#endif // GEN_PASS_DECL

//===----------------------------------------------------------------------===//
// ConvertPureCallOpsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTPURECALLOPSPASS
#undef GEN_PASS_DECL_CONVERTPURECALLOPSPASS
#endif // GEN_PASS_DECL_CONVERTPURECALLOPSPASS
#ifdef GEN_PASS_DEF_CONVERTPURECALLOPSPASS
namespace impl {

template <typename DerivedT>
class ConvertPureCallOpsPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = ConvertPureCallOpsPassBase;

  ConvertPureCallOpsPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  ConvertPureCallOpsPassBase(const ConvertPureCallOpsPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  ConvertPureCallOpsPassBase& operator=(const ConvertPureCallOpsPassBase &) = delete;
  ConvertPureCallOpsPassBase(ConvertPureCallOpsPassBase &&) = delete;
  ConvertPureCallOpsPassBase& operator=(ConvertPureCallOpsPassBase &&) = delete;
  ~ConvertPureCallOpsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-convert-pure-call-ops");
  }
  ::llvm::StringRef getArgument() const override { return "xla-convert-pure-call-ops"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Converts xla.pure_call to func.call)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("ConvertPureCallOpsPass");
  }
  ::llvm::StringRef getName() const override { return "ConvertPureCallOpsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<xla::XlaDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(ConvertPureCallOpsPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTPURECALLOPSPASS
#endif // GEN_PASS_DEF_CONVERTPURECALLOPSPASS

//===----------------------------------------------------------------------===//
// EraseDeadFunctionsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ERASEDEADFUNCTIONSPASS
#undef GEN_PASS_DECL_ERASEDEADFUNCTIONSPASS
#endif // GEN_PASS_DECL_ERASEDEADFUNCTIONSPASS
#ifdef GEN_PASS_DEF_ERASEDEADFUNCTIONSPASS
namespace impl {

template <typename DerivedT>
class EraseDeadFunctionsPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = EraseDeadFunctionsPassBase;

  EraseDeadFunctionsPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  EraseDeadFunctionsPassBase(const EraseDeadFunctionsPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  EraseDeadFunctionsPassBase& operator=(const EraseDeadFunctionsPassBase &) = delete;
  EraseDeadFunctionsPassBase(EraseDeadFunctionsPassBase &&) = delete;
  EraseDeadFunctionsPassBase& operator=(EraseDeadFunctionsPassBase &&) = delete;
  ~EraseDeadFunctionsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-erase-dead-functions");
  }
  ::llvm::StringRef getArgument() const override { return "xla-erase-dead-functions"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Deletes unused functions)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("EraseDeadFunctionsPass");
  }
  ::llvm::StringRef getName() const override { return "EraseDeadFunctionsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<xla::XlaDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(EraseDeadFunctionsPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_ERASEDEADFUNCTIONSPASS
#endif // GEN_PASS_DEF_ERASEDEADFUNCTIONSPASS

//===----------------------------------------------------------------------===//
// ExpandFloatOpsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_EXPANDFLOATOPSPASS
struct ExpandFloatOpsPassOptions {
  bool approximate_tanh_ = true;
};
#undef GEN_PASS_DECL_EXPANDFLOATOPSPASS
#endif // GEN_PASS_DECL_EXPANDFLOATOPSPASS
#ifdef GEN_PASS_DEF_EXPANDFLOATOPSPASS
namespace impl {

template <typename DerivedT>
class ExpandFloatOpsPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = ExpandFloatOpsPassBase;

  ExpandFloatOpsPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  ExpandFloatOpsPassBase(const ExpandFloatOpsPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  ExpandFloatOpsPassBase& operator=(const ExpandFloatOpsPassBase &) = delete;
  ExpandFloatOpsPassBase(ExpandFloatOpsPassBase &&) = delete;
  ExpandFloatOpsPassBase& operator=(ExpandFloatOpsPassBase &&) = delete;
  ~ExpandFloatOpsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-expand-float-ops");
  }
  ::llvm::StringRef getArgument() const override { return "xla-expand-float-ops"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Expands float ops that are not natively supported.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("ExpandFloatOpsPass");
  }
  ::llvm::StringRef getName() const override { return "ExpandFloatOpsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::arith::ArithDialect>();
    registry.insert<mlir::math::MathDialect>();
    registry.insert<mlir::mhlo::MhloDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(ExpandFloatOpsPassBase<DerivedT>)

  ExpandFloatOpsPassBase(ExpandFloatOpsPassOptions options) : ExpandFloatOpsPassBase() {
    approximate_tanh_ = std::move(options.approximate_tanh_);
  }
protected:
  ::mlir::Pass::Option<bool> approximate_tanh_{*this, "approximate_tanh", ::llvm::cl::desc(R"PO(Use built-in MLIR approximation for tanh.)PO"), ::llvm::cl::init(true)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_EXPANDFLOATOPSPASS
#endif // GEN_PASS_DEF_EXPANDFLOATOPSPASS

//===----------------------------------------------------------------------===//
// ExpandIntegerPowerPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_EXPANDINTEGERPOWERPASS
std::unique_ptr<::mlir::Pass> createExpandIntegerPowerPass();
#undef GEN_PASS_DECL_EXPANDINTEGERPOWERPASS
#endif // GEN_PASS_DECL_EXPANDINTEGERPOWERPASS
#ifdef GEN_PASS_DEF_EXPANDINTEGERPOWERPASS

namespace impl {
  std::unique_ptr<::mlir::Pass> createExpandIntegerPowerPass();
} // namespace impl
namespace impl {

template <typename DerivedT>
class ExpandIntegerPowerPassBase : public ::mlir::OperationPass<> {
public:
  using Base = ExpandIntegerPowerPassBase;

  ExpandIntegerPowerPassBase() : ::mlir::OperationPass<>(::mlir::TypeID::get<DerivedT>()) {}
  ExpandIntegerPowerPassBase(const ExpandIntegerPowerPassBase &other) : ::mlir::OperationPass<>(other) {}
  ExpandIntegerPowerPassBase& operator=(const ExpandIntegerPowerPassBase &) = delete;
  ExpandIntegerPowerPassBase(ExpandIntegerPowerPassBase &&) = delete;
  ExpandIntegerPowerPassBase& operator=(ExpandIntegerPowerPassBase &&) = delete;
  ~ExpandIntegerPowerPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-expand-integer-power");
  }
  ::llvm::StringRef getArgument() const override { return "xla-expand-integer-power"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Expands mlir::math::IPowIop to Exponentiation by squaring.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("ExpandIntegerPowerPass");
  }
  ::llvm::StringRef getName() const override { return "ExpandIntegerPowerPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::arith::ArithDialect>();
    registry.insert<mlir::math::MathDialect>();
    registry.insert<mlir::scf::SCFDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(ExpandIntegerPowerPassBase<DerivedT>)

protected:
private:

  friend std::unique_ptr<::mlir::Pass> createExpandIntegerPowerPass() {
    return std::make_unique<DerivedT>();
  }
};
} // namespace impl

std::unique_ptr<::mlir::Pass> createExpandIntegerPowerPass() {
  return impl::createExpandIntegerPowerPass();
}
#undef GEN_PASS_DEF_EXPANDINTEGERPOWERPASS
#endif // GEN_PASS_DEF_EXPANDINTEGERPOWERPASS

//===----------------------------------------------------------------------===//
// FlattenTensorsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_FLATTENTENSORSPASS
#undef GEN_PASS_DECL_FLATTENTENSORSPASS
#endif // GEN_PASS_DECL_FLATTENTENSORSPASS
#ifdef GEN_PASS_DEF_FLATTENTENSORSPASS
namespace impl {

template <typename DerivedT>
class FlattenTensorsPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = FlattenTensorsPassBase;

  FlattenTensorsPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  FlattenTensorsPassBase(const FlattenTensorsPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  FlattenTensorsPassBase& operator=(const FlattenTensorsPassBase &) = delete;
  FlattenTensorsPassBase(FlattenTensorsPassBase &&) = delete;
  FlattenTensorsPassBase& operator=(FlattenTensorsPassBase &&) = delete;
  ~FlattenTensorsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-flatten-tensors");
  }
  ::llvm::StringRef getArgument() const override { return "xla-flatten-tensors"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Flatten tensors.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("FlattenTensorsPass");
  }
  ::llvm::StringRef getName() const override { return "FlattenTensorsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<mlir::tensor::TensorDialect>();
    registry.insert<xla::gpu::XlaGpuDialect>();
    registry.insert<xla::XlaDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(FlattenTensorsPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_FLATTENTENSORSPASS
#endif // GEN_PASS_DEF_FLATTENTENSORSPASS

//===----------------------------------------------------------------------===//
// LowerTensorsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERTENSORSPASS
struct LowerTensorsPassOptions {
  std::string gpu_device_info_;
  std::string target_type_ = "gpu";
};
#undef GEN_PASS_DECL_LOWERTENSORSPASS
#endif // GEN_PASS_DECL_LOWERTENSORSPASS
#ifdef GEN_PASS_DEF_LOWERTENSORSPASS
namespace impl {

template <typename DerivedT>
class LowerTensorsPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = LowerTensorsPassBase;

  LowerTensorsPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  LowerTensorsPassBase(const LowerTensorsPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  LowerTensorsPassBase& operator=(const LowerTensorsPassBase &) = delete;
  LowerTensorsPassBase(LowerTensorsPassBase &&) = delete;
  LowerTensorsPassBase& operator=(LowerTensorsPassBase &&) = delete;
  ~LowerTensorsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-lower-tensors");
  }
  ::llvm::StringRef getArgument() const override { return "xla-lower-tensors"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Lowers tensors to llvm pointers and loads/stores.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("LowerTensorsPass");
  }
  ::llvm::StringRef getName() const override { return "LowerTensorsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::LLVM::LLVMDialect>();
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<mlir::gpu::GPUDialect>();
    registry.insert<mlir::scf::SCFDialect>();
    registry.insert<mlir::tensor::TensorDialect>();
    registry.insert<xla::gpu::XlaGpuDialect>();
    registry.insert<xla::XlaDialect>();
    registry.insert<mlir::vector::VectorDialect>();
    registry.insert<mlir::ROCDL::ROCDLDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LowerTensorsPassBase<DerivedT>)

  LowerTensorsPassBase(LowerTensorsPassOptions options) : LowerTensorsPassBase() {
    gpu_device_info_ = std::move(options.gpu_device_info_);
    target_type_ = std::move(options.target_type_);
  }
protected:
  ::mlir::Pass::Option<std::string> gpu_device_info_{*this, "gpu_device_info", ::llvm::cl::desc(R"PO(Serialized stream_executor::GPUDeviceInfo proto.)PO")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc(R"PO(Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty.)PO"), ::llvm::cl::init("gpu")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERTENSORSPASS
#endif // GEN_PASS_DEF_LOWERTENSORSPASS

//===----------------------------------------------------------------------===//
// LowerXlaIntrinsicLibPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERXLAINTRINSICLIBPASS
#undef GEN_PASS_DECL_LOWERXLAINTRINSICLIBPASS
#endif // GEN_PASS_DECL_LOWERXLAINTRINSICLIBPASS
#ifdef GEN_PASS_DEF_LOWERXLAINTRINSICLIBPASS
namespace impl {

template <typename DerivedT>
class LowerXlaIntrinsicLibPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = LowerXlaIntrinsicLibPassBase;

  LowerXlaIntrinsicLibPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  LowerXlaIntrinsicLibPassBase(const LowerXlaIntrinsicLibPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  LowerXlaIntrinsicLibPassBase& operator=(const LowerXlaIntrinsicLibPassBase &) = delete;
  LowerXlaIntrinsicLibPassBase(LowerXlaIntrinsicLibPassBase &&) = delete;
  LowerXlaIntrinsicLibPassBase& operator=(LowerXlaIntrinsicLibPassBase &&) = delete;
  ~LowerXlaIntrinsicLibPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-lower-xla-intrinsic-lib");
  }
  ::llvm::StringRef getArgument() const override { return "xla-lower-xla-intrinsic-lib"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Lowers XLA math function calls to xla.* intrinsic lib calls.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("LowerXlaIntrinsicLibPass");
  }
  ::llvm::StringRef getName() const override { return "LowerXlaIntrinsicLibPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<mlir::vector::VectorDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LowerXlaIntrinsicLibPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERXLAINTRINSICLIBPASS
#endif // GEN_PASS_DEF_LOWERXLAINTRINSICLIBPASS

//===----------------------------------------------------------------------===//
// LowerXlaLoopsToScfPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERXLALOOPSTOSCFPASS
#undef GEN_PASS_DECL_LOWERXLALOOPSTOSCFPASS
#endif // GEN_PASS_DECL_LOWERXLALOOPSTOSCFPASS
#ifdef GEN_PASS_DEF_LOWERXLALOOPSTOSCFPASS
namespace impl {

template <typename DerivedT>
class LowerXlaLoopsToScfPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = LowerXlaLoopsToScfPassBase;

  LowerXlaLoopsToScfPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  LowerXlaLoopsToScfPassBase(const LowerXlaLoopsToScfPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  LowerXlaLoopsToScfPassBase& operator=(const LowerXlaLoopsToScfPassBase &) = delete;
  LowerXlaLoopsToScfPassBase(LowerXlaLoopsToScfPassBase &&) = delete;
  LowerXlaLoopsToScfPassBase& operator=(LowerXlaLoopsToScfPassBase &&) = delete;
  ~LowerXlaLoopsToScfPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-lower-xla-loops-to-scf");
  }
  ::llvm::StringRef getArgument() const override { return "xla-lower-xla-loops-to-scf"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Lowers xla.loop to SCF.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("LowerXlaLoopsToScfPass");
  }
  ::llvm::StringRef getName() const override { return "LowerXlaLoopsToScfPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::scf::SCFDialect>();
    registry.insert<mlir::tensor::TensorDialect>();
    registry.insert<xla::gpu::XlaGpuDialect>();
    registry.insert<xla::XlaDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LowerXlaLoopsToScfPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERXLALOOPSTOSCFPASS
#endif // GEN_PASS_DEF_LOWERXLALOOPSTOSCFPASS

//===----------------------------------------------------------------------===//
// LowerXlaToScfPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERXLATOSCFPASS
struct LowerXlaToScfPassOptions {
  int64_t warp_size = 32;
};
#undef GEN_PASS_DECL_LOWERXLATOSCFPASS
#endif // GEN_PASS_DECL_LOWERXLATOSCFPASS
#ifdef GEN_PASS_DEF_LOWERXLATOSCFPASS
namespace impl {

template <typename DerivedT>
class LowerXlaToScfPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = LowerXlaToScfPassBase;

  LowerXlaToScfPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  LowerXlaToScfPassBase(const LowerXlaToScfPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  LowerXlaToScfPassBase& operator=(const LowerXlaToScfPassBase &) = delete;
  LowerXlaToScfPassBase(LowerXlaToScfPassBase &&) = delete;
  LowerXlaToScfPassBase& operator=(LowerXlaToScfPassBase &&) = delete;
  ~LowerXlaToScfPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-lower-xla-to-scf");
  }
  ::llvm::StringRef getArgument() const override { return "xla-lower-xla-to-scf"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Lowers xla to SCF.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("LowerXlaToScfPass");
  }
  ::llvm::StringRef getName() const override { return "LowerXlaToScfPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::gpu::GPUDialect>();
    registry.insert<mlir::LLVM::LLVMDialect>();
    registry.insert<mlir::scf::SCFDialect>();
    registry.insert<mlir::tensor::TensorDialect>();
    registry.insert<xla::gpu::XlaGpuDialect>();
    registry.insert<xla::XlaDialect>();
    registry.insert<mlir::vector::VectorDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LowerXlaToScfPassBase<DerivedT>)

  LowerXlaToScfPassBase(LowerXlaToScfPassOptions options) : LowerXlaToScfPassBase() {
    warp_size = std::move(options.warp_size);
  }
protected:
  ::mlir::Pass::Option<int64_t> warp_size{*this, "warp_size", ::llvm::cl::desc(R"PO(Warp size.)PO"), ::llvm::cl::init(32)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERXLATOSCFPASS
#endif // GEN_PASS_DEF_LOWERXLATOSCFPASS

//===----------------------------------------------------------------------===//
// MergePointersToSameSlicePass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_MERGEPOINTERSTOSAMESLICEPASS
#undef GEN_PASS_DECL_MERGEPOINTERSTOSAMESLICEPASS
#endif // GEN_PASS_DECL_MERGEPOINTERSTOSAMESLICEPASS
#ifdef GEN_PASS_DEF_MERGEPOINTERSTOSAMESLICEPASS
namespace impl {

template <typename DerivedT>
class MergePointersToSameSlicePassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = MergePointersToSameSlicePassBase;

  MergePointersToSameSlicePassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  MergePointersToSameSlicePassBase(const MergePointersToSameSlicePassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  MergePointersToSameSlicePassBase& operator=(const MergePointersToSameSlicePassBase &) = delete;
  MergePointersToSameSlicePassBase(MergePointersToSameSlicePassBase &&) = delete;
  MergePointersToSameSlicePassBase& operator=(MergePointersToSameSlicePassBase &&) = delete;
  ~MergePointersToSameSlicePassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-merge-pointers");
  }
  ::llvm::StringRef getArgument() const override { return "xla-merge-pointers"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Merges pointers that share slices.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("MergePointersToSameSlicePass");
  }
  ::llvm::StringRef getName() const override { return "MergePointersToSameSlicePass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(MergePointersToSameSlicePassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_MERGEPOINTERSTOSAMESLICEPASS
#endif // GEN_PASS_DEF_MERGEPOINTERSTOSAMESLICEPASS

//===----------------------------------------------------------------------===//
// PropagateSliceIndicesPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_PROPAGATESLICEINDICESPASS
#undef GEN_PASS_DECL_PROPAGATESLICEINDICESPASS
#endif // GEN_PASS_DECL_PROPAGATESLICEINDICESPASS
#ifdef GEN_PASS_DEF_PROPAGATESLICEINDICESPASS
namespace impl {

template <typename DerivedT>
class PropagateSliceIndicesPassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = PropagateSliceIndicesPassBase;

  PropagateSliceIndicesPassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  PropagateSliceIndicesPassBase(const PropagateSliceIndicesPassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  PropagateSliceIndicesPassBase& operator=(const PropagateSliceIndicesPassBase &) = delete;
  PropagateSliceIndicesPassBase(PropagateSliceIndicesPassBase &&) = delete;
  PropagateSliceIndicesPassBase& operator=(PropagateSliceIndicesPassBase &&) = delete;
  ~PropagateSliceIndicesPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-propagate-slice-indices");
  }
  ::llvm::StringRef getArgument() const override { return "xla-propagate-slice-indices"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Propagates slice indices from the entry function to all callees.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("PropagateSliceIndicesPass");
  }
  ::llvm::StringRef getName() const override { return "PropagateSliceIndicesPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(PropagateSliceIndicesPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_PROPAGATESLICEINDICESPASS
#endif // GEN_PASS_DEF_PROPAGATESLICEINDICESPASS

//===----------------------------------------------------------------------===//
// SafeIntegerArithmeticPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_SAFEINTEGERARITHMETICPASS
std::unique_ptr<::mlir::Pass> createSafeIntegerArithmeticPass();
#undef GEN_PASS_DECL_SAFEINTEGERARITHMETICPASS
#endif // GEN_PASS_DECL_SAFEINTEGERARITHMETICPASS
#ifdef GEN_PASS_DEF_SAFEINTEGERARITHMETICPASS

namespace impl {
  std::unique_ptr<::mlir::Pass> createSafeIntegerArithmeticPass();
} // namespace impl
namespace impl {

template <typename DerivedT>
class SafeIntegerArithmeticPassBase : public ::mlir::OperationPass<> {
public:
  using Base = SafeIntegerArithmeticPassBase;

  SafeIntegerArithmeticPassBase() : ::mlir::OperationPass<>(::mlir::TypeID::get<DerivedT>()) {}
  SafeIntegerArithmeticPassBase(const SafeIntegerArithmeticPassBase &other) : ::mlir::OperationPass<>(other) {}
  SafeIntegerArithmeticPassBase& operator=(const SafeIntegerArithmeticPassBase &) = delete;
  SafeIntegerArithmeticPassBase(SafeIntegerArithmeticPassBase &&) = delete;
  SafeIntegerArithmeticPassBase& operator=(SafeIntegerArithmeticPassBase &&) = delete;
  ~SafeIntegerArithmeticPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-safe-integer-arithmetic");
  }
  ::llvm::StringRef getArgument() const override { return "xla-safe-integer-arithmetic"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Make integer arithmetics safe.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("SafeIntegerArithmeticPass");
  }
  ::llvm::StringRef getName() const override { return "SafeIntegerArithmeticPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(SafeIntegerArithmeticPassBase<DerivedT>)

protected:
private:

  friend std::unique_ptr<::mlir::Pass> createSafeIntegerArithmeticPass() {
    return std::make_unique<DerivedT>();
  }
};
} // namespace impl

std::unique_ptr<::mlir::Pass> createSafeIntegerArithmeticPass() {
  return impl::createSafeIntegerArithmeticPass();
}
#undef GEN_PASS_DEF_SAFEINTEGERARITHMETICPASS
#endif // GEN_PASS_DEF_SAFEINTEGERARITHMETICPASS

//===----------------------------------------------------------------------===//
// SimplifyAffinePass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_SIMPLIFYAFFINEPASS
#undef GEN_PASS_DECL_SIMPLIFYAFFINEPASS
#endif // GEN_PASS_DECL_SIMPLIFYAFFINEPASS
#ifdef GEN_PASS_DEF_SIMPLIFYAFFINEPASS
namespace impl {

template <typename DerivedT>
class SimplifyAffinePassBase : public ::mlir::OperationPass<mlir::ModuleOp> {
public:
  using Base = SimplifyAffinePassBase;

  SimplifyAffinePassBase() : ::mlir::OperationPass<mlir::ModuleOp>(::mlir::TypeID::get<DerivedT>()) {}
  SimplifyAffinePassBase(const SimplifyAffinePassBase &other) : ::mlir::OperationPass<mlir::ModuleOp>(other) {}
  SimplifyAffinePassBase& operator=(const SimplifyAffinePassBase &) = delete;
  SimplifyAffinePassBase(SimplifyAffinePassBase &&) = delete;
  SimplifyAffinePassBase& operator=(SimplifyAffinePassBase &&) = delete;
  ~SimplifyAffinePassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-simplify-affine");
  }
  ::llvm::StringRef getArgument() const override { return "xla-simplify-affine"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Simplifies affine.apply using XLA's range-aware simplifier.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("SimplifyAffinePass");
  }
  ::llvm::StringRef getName() const override { return "SimplifyAffinePass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::affine::AffineDialect>();
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<mlir::scf::SCFDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(SimplifyAffinePassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_SIMPLIFYAFFINEPASS
#endif // GEN_PASS_DEF_SIMPLIFYAFFINEPASS

//===----------------------------------------------------------------------===//
// SimplifyArithPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_SIMPLIFYARITHPASS
struct SimplifyArithPassOptions {
  bool fast_min_max_ = false;
};
#undef GEN_PASS_DECL_SIMPLIFYARITHPASS
#endif // GEN_PASS_DECL_SIMPLIFYARITHPASS
#ifdef GEN_PASS_DEF_SIMPLIFYARITHPASS
namespace impl {

template <typename DerivedT>
class SimplifyArithPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = SimplifyArithPassBase;

  SimplifyArithPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  SimplifyArithPassBase(const SimplifyArithPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  SimplifyArithPassBase& operator=(const SimplifyArithPassBase &) = delete;
  SimplifyArithPassBase(SimplifyArithPassBase &&) = delete;
  SimplifyArithPassBase& operator=(SimplifyArithPassBase &&) = delete;
  ~SimplifyArithPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-simplify-arith");
  }
  ::llvm::StringRef getArgument() const override { return "xla-simplify-arith"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Simplifies arith using XLA's range-aware simplifier.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("SimplifyArithPass");
  }
  ::llvm::StringRef getName() const override { return "SimplifyArithPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::arith::ArithDialect>();
    registry.insert<mlir::func::FuncDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(SimplifyArithPassBase<DerivedT>)

  SimplifyArithPassBase(SimplifyArithPassOptions options) : SimplifyArithPassBase() {
    fast_min_max_ = std::move(options.fast_min_max_);
  }
protected:
  ::mlir::Pass::Option<bool> fast_min_max_{*this, "fast_min_max", ::llvm::cl::desc(R"PO(Use fast min/max (does not propagate nan).)PO"), ::llvm::cl::init(false)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_SIMPLIFYARITHPASS
#endif // GEN_PASS_DEF_SIMPLIFYARITHPASS

//===----------------------------------------------------------------------===//
// UnswitchLoopsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_UNSWITCHLOOPSPASS
#undef GEN_PASS_DECL_UNSWITCHLOOPSPASS
#endif // GEN_PASS_DECL_UNSWITCHLOOPSPASS
#ifdef GEN_PASS_DEF_UNSWITCHLOOPSPASS
namespace impl {

template <typename DerivedT>
class UnswitchLoopsPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = UnswitchLoopsPassBase;

  UnswitchLoopsPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  UnswitchLoopsPassBase(const UnswitchLoopsPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  UnswitchLoopsPassBase& operator=(const UnswitchLoopsPassBase &) = delete;
  UnswitchLoopsPassBase(UnswitchLoopsPassBase &&) = delete;
  UnswitchLoopsPassBase& operator=(UnswitchLoopsPassBase &&) = delete;
  ~UnswitchLoopsPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-unswitch-loops");
  }
  ::llvm::StringRef getArgument() const override { return "xla-unswitch-loops"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Swaps scf.if and scf.for.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("UnswitchLoopsPass");
  }
  ::llvm::StringRef getName() const override { return "UnswitchLoopsPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::func::FuncDialect>();
    registry.insert<mlir::scf::SCFDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(UnswitchLoopsPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_UNSWITCHLOOPSPASS
#endif // GEN_PASS_DEF_UNSWITCHLOOPSPASS

//===----------------------------------------------------------------------===//
// VectorizeLoadsAndStoresPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_VECTORIZELOADSANDSTORESPASS
struct VectorizeLoadsAndStoresPassOptions {
  std::string gpu_device_info_;
  std::string target_type_ = "gpu";
};
#undef GEN_PASS_DECL_VECTORIZELOADSANDSTORESPASS
#endif // GEN_PASS_DECL_VECTORIZELOADSANDSTORESPASS
#ifdef GEN_PASS_DEF_VECTORIZELOADSANDSTORESPASS
namespace impl {

template <typename DerivedT>
class VectorizeLoadsAndStoresPassBase : public ::mlir::OperationPass<mlir::func::FuncOp> {
public:
  using Base = VectorizeLoadsAndStoresPassBase;

  VectorizeLoadsAndStoresPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  VectorizeLoadsAndStoresPassBase(const VectorizeLoadsAndStoresPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(other) {}
  VectorizeLoadsAndStoresPassBase& operator=(const VectorizeLoadsAndStoresPassBase &) = delete;
  VectorizeLoadsAndStoresPassBase(VectorizeLoadsAndStoresPassBase &&) = delete;
  VectorizeLoadsAndStoresPassBase& operator=(VectorizeLoadsAndStoresPassBase &&) = delete;
  ~VectorizeLoadsAndStoresPassBase() = default;

  /// Returns the command-line argument attached to this pass.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xla-vectorize-loads-stores");
  }
  ::llvm::StringRef getArgument() const override { return "xla-vectorize-loads-stores"; }

  ::llvm::StringRef getDescription() const override { return R"PD(Vectorizes loads and stores.)PD"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("VectorizeLoadsAndStoresPass");
  }
  ::llvm::StringRef getName() const override { return "VectorizeLoadsAndStoresPass"; }

  /// Support isa/dyn_cast functionality for the derived pass class.
  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  /// A clone method to create a copy of this pass.
  std::unique_ptr<::mlir::Pass> clonePass() const override {
    return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
  }

  /// Return the dialect that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::vector::VectorDialect>();
  }

  /// Explicitly declare the TypeID for this class. We declare an explicit private
  /// instantiation because Pass classes should only be visible by the current
  /// library.
  MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(VectorizeLoadsAndStoresPassBase<DerivedT>)

  VectorizeLoadsAndStoresPassBase(VectorizeLoadsAndStoresPassOptions options) : VectorizeLoadsAndStoresPassBase() {
    gpu_device_info_ = std::move(options.gpu_device_info_);
    target_type_ = std::move(options.target_type_);
  }
protected:
  ::mlir::Pass::Option<std::string> gpu_device_info_{*this, "gpu_device_info", ::llvm::cl::desc(R"PO(Serialized stream_executor::GPUDeviceInfo proto.)PO")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc(R"PO(Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty.)PO"), ::llvm::cl::init("gpu")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_VECTORIZELOADSANDSTORESPASS
#endif // GEN_PASS_DEF_VECTORIZELOADSANDSTORESPASS
#ifdef GEN_PASS_REGISTRATION
// Generate registrations for all passes.
#define GEN_PASS_REGISTRATION_CONVERTPURECALLOPSPASS
#define GEN_PASS_REGISTRATION_ERASEDEADFUNCTIONSPASS
#define GEN_PASS_REGISTRATION_EXPANDFLOATOPSPASS
#define GEN_PASS_REGISTRATION_EXPANDINTEGERPOWERPASS
#define GEN_PASS_REGISTRATION_FLATTENTENSORSPASS
#define GEN_PASS_REGISTRATION_LOWERTENSORSPASS
#define GEN_PASS_REGISTRATION_LOWERXLAINTRINSICLIBPASS
#define GEN_PASS_REGISTRATION_LOWERXLALOOPSTOSCFPASS
#define GEN_PASS_REGISTRATION_LOWERXLATOSCFPASS
#define GEN_PASS_REGISTRATION_MERGEPOINTERSTOSAMESLICEPASS
#define GEN_PASS_REGISTRATION_PROPAGATESLICEINDICESPASS
#define GEN_PASS_REGISTRATION_SAFEINTEGERARITHMETICPASS
#define GEN_PASS_REGISTRATION_SIMPLIFYAFFINEPASS
#define GEN_PASS_REGISTRATION_SIMPLIFYARITHPASS
#define GEN_PASS_REGISTRATION_UNSWITCHLOOPSPASS
#define GEN_PASS_REGISTRATION_VECTORIZELOADSANDSTORESPASS
#endif // GEN_PASS_REGISTRATION

//===----------------------------------------------------------------------===//
// ConvertPureCallOpsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_CONVERTPURECALLOPSPASS

inline void registerConvertPureCallOpsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateConvertPureCallOpsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerConvertPureCallOpsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateConvertPureCallOpsPass();
  });
}

#undef GEN_PASS_REGISTRATION_CONVERTPURECALLOPSPASS
#endif // GEN_PASS_REGISTRATION_CONVERTPURECALLOPSPASS

//===----------------------------------------------------------------------===//
// EraseDeadFunctionsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_ERASEDEADFUNCTIONSPASS

inline void registerEraseDeadFunctionsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateEraseDeadFunctionsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerEraseDeadFunctionsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateEraseDeadFunctionsPass();
  });
}

#undef GEN_PASS_REGISTRATION_ERASEDEADFUNCTIONSPASS
#endif // GEN_PASS_REGISTRATION_ERASEDEADFUNCTIONSPASS

//===----------------------------------------------------------------------===//
// ExpandFloatOpsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_EXPANDFLOATOPSPASS

inline void registerExpandFloatOpsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateExpandFloatOpsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerExpandFloatOpsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateExpandFloatOpsPass();
  });
}

#undef GEN_PASS_REGISTRATION_EXPANDFLOATOPSPASS
#endif // GEN_PASS_REGISTRATION_EXPANDFLOATOPSPASS

//===----------------------------------------------------------------------===//
// ExpandIntegerPowerPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_EXPANDINTEGERPOWERPASS

inline void registerExpandIntegerPowerPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return createExpandIntegerPowerPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerExpandIntegerPowerPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return createExpandIntegerPowerPass();
  });
}

#undef GEN_PASS_REGISTRATION_EXPANDINTEGERPOWERPASS
#endif // GEN_PASS_REGISTRATION_EXPANDINTEGERPOWERPASS

//===----------------------------------------------------------------------===//
// FlattenTensorsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_FLATTENTENSORSPASS

inline void registerFlattenTensorsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateFlattenTensorsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerFlattenTensorsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateFlattenTensorsPass();
  });
}

#undef GEN_PASS_REGISTRATION_FLATTENTENSORSPASS
#endif // GEN_PASS_REGISTRATION_FLATTENTENSORSPASS

//===----------------------------------------------------------------------===//
// LowerTensorsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_LOWERTENSORSPASS

inline void registerLowerTensorsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerTensorsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerLowerTensorsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerTensorsPass();
  });
}

#undef GEN_PASS_REGISTRATION_LOWERTENSORSPASS
#endif // GEN_PASS_REGISTRATION_LOWERTENSORSPASS

//===----------------------------------------------------------------------===//
// LowerXlaIntrinsicLibPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_LOWERXLAINTRINSICLIBPASS

inline void registerLowerXlaIntrinsicLibPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaIntrinsicLibPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerLowerXlaIntrinsicLibPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaIntrinsicLibPass();
  });
}

#undef GEN_PASS_REGISTRATION_LOWERXLAINTRINSICLIBPASS
#endif // GEN_PASS_REGISTRATION_LOWERXLAINTRINSICLIBPASS

//===----------------------------------------------------------------------===//
// LowerXlaLoopsToScfPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_LOWERXLALOOPSTOSCFPASS

inline void registerLowerXlaLoopsToScfPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaLoopsToScfPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerLowerXlaLoopsToScfPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaLoopsToScfPass();
  });
}

#undef GEN_PASS_REGISTRATION_LOWERXLALOOPSTOSCFPASS
#endif // GEN_PASS_REGISTRATION_LOWERXLALOOPSTOSCFPASS

//===----------------------------------------------------------------------===//
// LowerXlaToScfPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_LOWERXLATOSCFPASS

inline void registerLowerXlaToScfPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaToScfPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerLowerXlaToScfPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateLowerXlaToScfPass();
  });
}

#undef GEN_PASS_REGISTRATION_LOWERXLATOSCFPASS
#endif // GEN_PASS_REGISTRATION_LOWERXLATOSCFPASS

//===----------------------------------------------------------------------===//
// MergePointersToSameSlicePass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_MERGEPOINTERSTOSAMESLICEPASS

inline void registerMergePointersToSameSlicePass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateMergePointersToSameSlicePass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerMergePointersToSameSlicePassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateMergePointersToSameSlicePass();
  });
}

#undef GEN_PASS_REGISTRATION_MERGEPOINTERSTOSAMESLICEPASS
#endif // GEN_PASS_REGISTRATION_MERGEPOINTERSTOSAMESLICEPASS

//===----------------------------------------------------------------------===//
// PropagateSliceIndicesPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_PROPAGATESLICEINDICESPASS

inline void registerPropagateSliceIndicesPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreatePropagateSliceIndicesPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerPropagateSliceIndicesPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreatePropagateSliceIndicesPass();
  });
}

#undef GEN_PASS_REGISTRATION_PROPAGATESLICEINDICESPASS
#endif // GEN_PASS_REGISTRATION_PROPAGATESLICEINDICESPASS

//===----------------------------------------------------------------------===//
// SafeIntegerArithmeticPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_SAFEINTEGERARITHMETICPASS

inline void registerSafeIntegerArithmeticPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return createSafeIntegerArithmeticPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerSafeIntegerArithmeticPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return createSafeIntegerArithmeticPass();
  });
}

#undef GEN_PASS_REGISTRATION_SAFEINTEGERARITHMETICPASS
#endif // GEN_PASS_REGISTRATION_SAFEINTEGERARITHMETICPASS

//===----------------------------------------------------------------------===//
// SimplifyAffinePass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_SIMPLIFYAFFINEPASS

inline void registerSimplifyAffinePass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateSimplifyAffinePass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerSimplifyAffinePassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateSimplifyAffinePass();
  });
}

#undef GEN_PASS_REGISTRATION_SIMPLIFYAFFINEPASS
#endif // GEN_PASS_REGISTRATION_SIMPLIFYAFFINEPASS

//===----------------------------------------------------------------------===//
// SimplifyArithPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_SIMPLIFYARITHPASS

inline void registerSimplifyArithPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateSimplifyArithPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerSimplifyArithPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateSimplifyArithPass();
  });
}

#undef GEN_PASS_REGISTRATION_SIMPLIFYARITHPASS
#endif // GEN_PASS_REGISTRATION_SIMPLIFYARITHPASS

//===----------------------------------------------------------------------===//
// UnswitchLoopsPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_UNSWITCHLOOPSPASS

inline void registerUnswitchLoopsPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateUnswitchLoopsPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerUnswitchLoopsPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateUnswitchLoopsPass();
  });
}

#undef GEN_PASS_REGISTRATION_UNSWITCHLOOPSPASS
#endif // GEN_PASS_REGISTRATION_UNSWITCHLOOPSPASS

//===----------------------------------------------------------------------===//
// VectorizeLoadsAndStoresPass Registration
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_REGISTRATION_VECTORIZELOADSANDSTORESPASS

inline void registerVectorizeLoadsAndStoresPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateVectorizeLoadsAndStoresPass();
  });
}

// Old registration code, kept for temporary backwards compatibility.
inline void registerVectorizeLoadsAndStoresPassPass() {
  ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
    return CreateVectorizeLoadsAndStoresPass();
  });
}

#undef GEN_PASS_REGISTRATION_VECTORIZELOADSANDSTORESPASS
#endif // GEN_PASS_REGISTRATION_VECTORIZELOADSANDSTORESPASS
#ifdef GEN_PASS_REGISTRATION

//===----------------------------------------------------------------------===//
// Transforms Registration
//===----------------------------------------------------------------------===//

inline void registerTransformsPasses() {
  registerConvertPureCallOpsPass();
  registerEraseDeadFunctionsPass();
  registerExpandFloatOpsPass();
  registerExpandIntegerPowerPass();
  registerFlattenTensorsPass();
  registerLowerTensorsPass();
  registerLowerXlaIntrinsicLibPass();
  registerLowerXlaLoopsToScfPass();
  registerLowerXlaToScfPass();
  registerMergePointersToSameSlicePass();
  registerPropagateSliceIndicesPass();
  registerSafeIntegerArithmeticPass();
  registerSimplifyAffinePass();
  registerSimplifyArithPass();
  registerUnswitchLoopsPass();
  registerVectorizeLoadsAndStoresPass();
}
#undef GEN_PASS_REGISTRATION
#endif // GEN_PASS_REGISTRATION
#ifdef GEN_PASS_CLASSES
#error "GEN_PASS_CLASSES is deprecated; use per-pass macros"
#undef GEN_PASS_CLASSES
#endif // GEN_PASS_CLASSES