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

#ifdef GEN_PASS_DECL
// Generate declarations for all passes.
#define GEN_PASS_DECL_ADDLOOPUNROLLFLAGSPASS
#define GEN_PASS_DECL_ADDREDUCTIONFASTMATHFLAGSPASS
#define GEN_PASS_DECL_EXPANDFLOATOPSPASS
#define GEN_PASS_DECL_LOWERTOLLVMPASS
#define GEN_PASS_DECL_LOWERXLASHAREDPASS
#define GEN_PASS_DECL_PEELWORKGROUPLOOPPASS
#undef GEN_PASS_DECL
#endif // GEN_PASS_DECL

//===----------------------------------------------------------------------===//
// AddLoopUnrollFlagsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ADDLOOPUNROLLFLAGSPASS
struct AddLoopUnrollFlagsPassOptions {
  int32_t max_nested_bits_ = 256;
};
#undef GEN_PASS_DECL_ADDLOOPUNROLLFLAGSPASS
#endif // GEN_PASS_DECL_ADDLOOPUNROLLFLAGSPASS
#ifdef GEN_PASS_DEF_ADDLOOPUNROLLFLAGSPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Add loop unroll flags to large nested loops."; }

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

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

  AddLoopUnrollFlagsPassBase(AddLoopUnrollFlagsPassOptions options) : AddLoopUnrollFlagsPassBase() {
    max_nested_bits_ = std::move(options.max_nested_bits_);
  }
protected:
  ::mlir::Pass::Option<int32_t> max_nested_bits_{*this, "max_nested_bits", ::llvm::cl::desc("The maximum number of bits accessed in a nested loop before disabling unrolling"), ::llvm::cl::init(256)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_ADDLOOPUNROLLFLAGSPASS
#endif // GEN_PASS_DEF_ADDLOOPUNROLLFLAGSPASS

//===----------------------------------------------------------------------===//
// AddReductionFastMathFlagsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ADDREDUCTIONFASTMATHFLAGSPASS
#undef GEN_PASS_DECL_ADDREDUCTIONFASTMATHFLAGSPASS
#endif // GEN_PASS_DECL_ADDREDUCTIONFASTMATHFLAGSPASS
#ifdef GEN_PASS_DEF_ADDREDUCTIONFASTMATHFLAGSPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Add fast math flags to reduction functions."; }

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

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

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

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

//===----------------------------------------------------------------------===//
// 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::func::FuncOp> {
public:
  using Base = ExpandFloatOpsPassBase;

  ExpandFloatOpsPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  ExpandFloatOpsPassBase(const ExpandFloatOpsPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(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-cpu-expand-float-ops");
  }
  ::llvm::StringRef getArgument() const override { return "xla-cpu-expand-float-ops"; }

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

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

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

//===----------------------------------------------------------------------===//
// LowerToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERTOLLVMPASS
struct LowerToLLVMPassOptions {
  int32_t prefer_vector_width_ = 256;
};
std::unique_ptr<::mlir::Pass> createLowerToLLVMPass();
std::unique_ptr<::mlir::Pass> createLowerToLLVMPass(LowerToLLVMPassOptions options);
#undef GEN_PASS_DECL_LOWERTOLLVMPASS
#endif // GEN_PASS_DECL_LOWERTOLLVMPASS
#ifdef GEN_PASS_DEF_LOWERTOLLVMPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createLowerToLLVMPass(LowerToLLVMPassOptions options);
} // namespace impl
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-cpu-lower-to-llvm");
  }
  ::llvm::StringRef getArgument() const override { return "xla-cpu-lower-to-llvm"; }

  ::llvm::StringRef getDescription() const override { return "Lowering from tensors + xla_cpu 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::tensor::TensorDialect>();
    registry.insert<xla::cpu::XlaCpuDialect>();
    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(LowerToLLVMPassBase<DerivedT>)

  LowerToLLVMPassBase(LowerToLLVMPassOptions options) : LowerToLLVMPassBase() {
    prefer_vector_width_ = std::move(options.prefer_vector_width_);
  }
protected:
  ::mlir::Pass::Option<int32_t> prefer_vector_width_{*this, "prefer_vector_width", ::llvm::cl::desc("prefer-vector-width value to set on entry function"), ::llvm::cl::init(256)};
private:

  friend std::unique_ptr<::mlir::Pass> createLowerToLLVMPass() {
    return std::make_unique<DerivedT>();
  }

  friend std::unique_ptr<::mlir::Pass> createLowerToLLVMPass(LowerToLLVMPassOptions options) {
    return std::make_unique<DerivedT>(std::move(options));
  }
};
} // namespace impl

std::unique_ptr<::mlir::Pass> createLowerToLLVMPass() {
  return impl::createLowerToLLVMPass();
}

std::unique_ptr<::mlir::Pass> createLowerToLLVMPass(LowerToLLVMPassOptions options) {
  return impl::createLowerToLLVMPass(std::move(options));
}
#undef GEN_PASS_DEF_LOWERTOLLVMPASS
#endif // GEN_PASS_DEF_LOWERTOLLVMPASS

//===----------------------------------------------------------------------===//
// LowerXlaSharedPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_LOWERXLASHAREDPASS
#undef GEN_PASS_DECL_LOWERXLASHAREDPASS
#endif // GEN_PASS_DECL_LOWERXLASHAREDPASS
#ifdef GEN_PASS_DEF_LOWERXLASHAREDPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lowers XLA shared dialect ops to XLA cpu compatable ops."; }

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

  /// 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<xla::XlaDialect>();
    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(LowerXlaSharedPassBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// PeelWorkgroupLoopPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_PEELWORKGROUPLOOPPASS
#undef GEN_PASS_DECL_PEELWORKGROUPLOOPPASS
#endif // GEN_PASS_DECL_PEELWORKGROUPLOOPPASS
#ifdef GEN_PASS_DEF_PEELWORKGROUPLOOPPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Split the loop into two, creating one with no constraints."; }

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

  /// 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<xla::XlaDialect>();
    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(PeelWorkgroupLoopPassBase<DerivedT>)

protected:
private:
};
} // namespace impl
#undef GEN_PASS_DEF_PEELWORKGROUPLOOPPASS
#endif // GEN_PASS_DEF_PEELWORKGROUPLOOPPASS
#ifdef GEN_PASS_REGISTRATION

//===----------------------------------------------------------------------===//
// AddLoopUnrollFlagsPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// AddReductionFastMathFlagsPass Registration
//===----------------------------------------------------------------------===//

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

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

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

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

//===----------------------------------------------------------------------===//
// LowerXlaSharedPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// PeelWorkgroupLoopPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// XlaCpuTransforms Registration
//===----------------------------------------------------------------------===//

inline void registerXlaCpuTransformsPasses() {
  registerAddLoopUnrollFlagsPass();
  registerAddReductionFastMathFlagsPass();
  registerExpandFloatOpsPass();
  registerLowerToLLVMPass();
  registerLowerXlaSharedPass();
  registerPeelWorkgroupLoopPass();
}
#undef GEN_PASS_REGISTRATION
#endif // GEN_PASS_REGISTRATION
// Deprecated. Please use the new per-pass macros.
#ifdef GEN_PASS_CLASSES

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

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

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

  ::llvm::StringRef getDescription() const override { return "Add loop unroll flags to large nested loops."; }

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

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

protected:
  ::mlir::Pass::Option<int32_t> max_nested_bits_{*this, "max_nested_bits", ::llvm::cl::desc("The maximum number of bits accessed in a nested loop before disabling unrolling"), ::llvm::cl::init(256)};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Add fast math flags to reduction functions."; }

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

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

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

protected:
};

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

  ExpandFloatOpsPassBase() : ::mlir::OperationPass<mlir::func::FuncOp>(::mlir::TypeID::get<DerivedT>()) {}
  ExpandFloatOpsPassBase(const ExpandFloatOpsPassBase &other) : ::mlir::OperationPass<mlir::func::FuncOp>(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-cpu-expand-float-ops");
  }
  ::llvm::StringRef getArgument() const override { return "xla-cpu-expand-float-ops"; }

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

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

protected:
};

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-cpu-lower-to-llvm");
  }
  ::llvm::StringRef getArgument() const override { return "xla-cpu-lower-to-llvm"; }

  ::llvm::StringRef getDescription() const override { return "Lowering from tensors + xla_cpu 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::tensor::TensorDialect>();
    registry.insert<xla::cpu::XlaCpuDialect>();
    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(LowerToLLVMPassBase<DerivedT>)

protected:
  ::mlir::Pass::Option<int32_t> prefer_vector_width_{*this, "prefer_vector_width", ::llvm::cl::desc("prefer-vector-width value to set on entry function"), ::llvm::cl::init(256)};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lowers XLA shared dialect ops to XLA cpu compatable ops."; }

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

  /// 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<xla::XlaDialect>();
    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(LowerXlaSharedPassBase<DerivedT>)

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Split the loop into two, creating one with no constraints."; }

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

  /// 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<xla::XlaDialect>();
    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(PeelWorkgroupLoopPassBase<DerivedT>)

protected:
};
#undef GEN_PASS_CLASSES
#endif // GEN_PASS_CLASSES