/* 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_FLATTENTENSORSPASS
#define GEN_PASS_DECL_LOWERTENSORSPASS
#define GEN_PASS_DECL_LOWERTOLLVMPASS
#define GEN_PASS_DECL_LOWERXLALOOPSTOSCFPASS
#define GEN_PASS_DECL_LOWERXLAMATHLIBPASS
#define GEN_PASS_DECL_LOWERXLATOSCFPASS
#define GEN_PASS_DECL_MERGEPOINTERSTOSAMESLICEPASS
#define GEN_PASS_DECL_PROPAGATEALIASSCOPESPASS
#define GEN_PASS_DECL_PROPAGATESLICEINDICESPASS
#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 "Converts xla.pure_call to func.call"; }

  /// 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 "Deletes unused functions"; }

  /// 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
#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 "Expands float ops that are not natively supported."; }

  /// 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>)

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

//===----------------------------------------------------------------------===//
// 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 "Flatten tensors."; }

  /// 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 "Lowers tensors to llvm pointers and loads/stores."; }

  /// 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("Serialized stream_executor::GPUDeviceInfo proto.")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc("Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty."), ::llvm::cl::init("gpu")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERTENSORSPASS
#endif // GEN_PASS_DEF_LOWERTENSORSPASS

//===----------------------------------------------------------------------===//
// LowerToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERTOLLVMPASS
struct LowerToLLVMPassOptions {
  std::string gpu_device_info_;
  std::string target_type_ = "gpu";
};
#undef GEN_PASS_DECL_LOWERTOLLVMPASS
#endif // GEN_PASS_DECL_LOWERTOLLVMPASS
#ifdef GEN_PASS_DEF_LOWERTOLLVMPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lowers to LLVM."; }

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

  /// 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::LLVM::LLVMDialect>();
    registry.insert<mlir::NVVM::NVVMDialect>();
  }

  /// 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(LowerToLLVMPassBase<DerivedT>)

  LowerToLLVMPassBase(LowerToLLVMPassOptions options) : LowerToLLVMPassBase() {
    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("Serialized stream_executor::GPUDeviceInfo proto.")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc("Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty."), ::llvm::cl::init("gpu")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_LOWERTOLLVMPASS
#endif // GEN_PASS_DEF_LOWERTOLLVMPASS

//===----------------------------------------------------------------------===//
// 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 "Lowers xla.loop to SCF."; }

  /// 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

//===----------------------------------------------------------------------===//
// LowerXlaMathLibPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERXLAMATHLIBPASS
#undef GEN_PASS_DECL_LOWERXLAMATHLIBPASS
#endif // GEN_PASS_DECL_LOWERXLAMATHLIBPASS
#ifdef GEN_PASS_DEF_LOWERXLAMATHLIBPASS
namespace impl {

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

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

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

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

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

  /// 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(LowerXlaMathLibPassBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// 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 "Lowers xla to SCF."; }

  /// 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("Warp size."), ::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 "Merges pointers that share slices."; }

  /// 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

//===----------------------------------------------------------------------===//
// PropagateAliasScopesPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_PROPAGATEALIASSCOPESPASS
std::unique_ptr<::mlir::Pass> createPropagateAliasScopesPass();
#undef GEN_PASS_DECL_PROPAGATEALIASSCOPESPASS
#endif // GEN_PASS_DECL_PROPAGATEALIASSCOPESPASS
#ifdef GEN_PASS_DEF_PROPAGATEALIASSCOPESPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Propagates alias scopes from function args to tensor modifications."; }

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

  /// 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::LLVM::LLVMDialect>();
  }

  /// 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(PropagateAliasScopesPassBase<DerivedT>)

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// 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 "Propagates slice indices from the entry function to all callees."; }

  /// 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

//===----------------------------------------------------------------------===//
// 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 "Simplifies affine.apply using XLA's range-aware simplifier."; }

  /// 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 "Simplifies arith using XLA's range-aware simplifier."; }

  /// 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("Use fast min/max (does not propagate nan)."), ::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 "Swaps scf.if and scf.for."; }

  /// 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 "Vectorizes loads and stores."; }

  /// 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("Serialized stream_executor::GPUDeviceInfo proto.")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc("Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty."), ::llvm::cl::init("gpu")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_VECTORIZELOADSANDSTORESPASS
#endif // GEN_PASS_DEF_VECTORIZELOADSANDSTORESPASS
#ifdef GEN_PASS_REGISTRATION

//===----------------------------------------------------------------------===//
// ConvertPureCallOpsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// EraseDeadFunctionsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// ExpandFloatOpsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// FlattenTensorsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// LowerTensorsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// LowerToLLVMPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// LowerXlaLoopsToScfPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// LowerXlaMathLibPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// LowerXlaToScfPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// MergePointersToSameSlicePass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// PropagateAliasScopesPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// PropagateSliceIndicesPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// SimplifyAffinePass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// SimplifyArithPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// UnswitchLoopsPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

//===----------------------------------------------------------------------===//
// VectorizeLoadsAndStoresPass Registration
//===----------------------------------------------------------------------===//

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();
  });
}

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

inline void registerTransformsPasses() {
  registerConvertPureCallOpsPass();
  registerEraseDeadFunctionsPass();
  registerExpandFloatOpsPass();
  registerFlattenTensorsPass();
  registerLowerTensorsPass();
  registerLowerToLLVMPass();
  registerLowerXlaLoopsToScfPass();
  registerLowerXlaMathLibPass();
  registerLowerXlaToScfPass();
  registerMergePointersToSameSlicePass();
  registerPropagateAliasScopesPass();
  registerPropagateSliceIndicesPass();
  registerSimplifyAffinePass();
  registerSimplifyArithPass();
  registerUnswitchLoopsPass();
  registerVectorizeLoadsAndStoresPass();
}
#undef GEN_PASS_REGISTRATION
#endif // GEN_PASS_REGISTRATION
// Deprecated. Please use the new per-pass macros.
#ifdef GEN_PASS_CLASSES

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 "Converts xla.pure_call to func.call"; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Deletes unused functions"; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Expands float ops that are not natively supported."; }

  /// 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));
  }

  /// Register the dialects 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>)

protected:
};

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 "Flatten tensors."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Lowers tensors to llvm pointers and loads/stores."; }

  /// 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));
  }

  /// Register the dialects 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>)

protected:
  ::mlir::Pass::Option<std::string> gpu_device_info_{*this, "gpu_device_info", ::llvm::cl::desc("Serialized stream_executor::GPUDeviceInfo proto.")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc("Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty."), ::llvm::cl::init("gpu")};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lowers to LLVM."; }

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

  /// 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));
  }

  /// Register the dialects 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::LLVM::LLVMDialect>();
    registry.insert<mlir::NVVM::NVVMDialect>();
  }

  /// 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(LowerToLLVMPassBase<DerivedT>)

protected:
  ::mlir::Pass::Option<std::string> gpu_device_info_{*this, "gpu_device_info", ::llvm::cl::desc("Serialized stream_executor::GPUDeviceInfo proto.")};
  ::mlir::Pass::Option<std::string> target_type_{*this, "target_type", ::llvm::cl::desc("Whether the pass targets a 'cpu' or 'gpu'. If 'cpu', gpu_device_info_ must be empty."), ::llvm::cl::init("gpu")};
};

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 "Lowers xla.loop to SCF."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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

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

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

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

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

  /// 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));
  }

  /// Register the dialects 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(LowerXlaMathLibPassBase<DerivedT>)

protected:
};

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 "Lowers xla to SCF."; }

  /// 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));
  }

  /// Register the dialects 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>)

protected:
  ::mlir::Pass::Option<int64_t> warp_size{*this, "warp_size", ::llvm::cl::desc("Warp size."), ::llvm::cl::init(32)};
};

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 "Merges pointers that share slices."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Propagates alias scopes from function args to tensor modifications."; }

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

  /// 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));
  }

  /// Register the dialects 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::LLVM::LLVMDialect>();
  }

  /// 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(PropagateAliasScopesPassBase<DerivedT>)

protected:
};

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 "Propagates slice indices from the entry function to all callees."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Simplifies affine.apply using XLA's range-aware simplifier."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Simplifies arith using XLA's range-aware simplifier."; }

  /// 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));
  }

  /// Register the dialects 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>)

protected:
  ::mlir::Pass::Option<bool> fast_min_max_{*this, "fast_min_max", ::llvm::cl::desc("Use fast min/max (does not propagate nan)."), ::llvm::cl::init(false)};
};

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 "Swaps scf.if and scf.for."; }

  /// 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));
  }

  /// Register the dialects 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:
};

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 "Vectorizes loads and stores."; }

  /// 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));
  }

  /// Register the dialects 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>)

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