llvm/tools/mlir/include/mlir/Conversion/Passes.h.inc

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

#ifdef GEN_PASS_DECL
// Generate declarations for all passes.
#define GEN_PASS_DECL_ARITHTOAMDGPUCONVERSIONPASS
#define GEN_PASS_DECL_ARITHTOARMSMECONVERSIONPASS
#define GEN_PASS_DECL_ARITHTOLLVMCONVERSIONPASS
#define GEN_PASS_DECL_CONVERTAMDGPUTOROCDL
#define GEN_PASS_DECL_CONVERTAFFINEFORTOGPU
#define GEN_PASS_DECL_CONVERTAFFINETOSTANDARD
#define GEN_PASS_DECL_CONVERTARITHTOEMITC
#define GEN_PASS_DECL_CONVERTARITHTOSPIRV
#define GEN_PASS_DECL_CONVERTARMNEON2DTOINTR
#define GEN_PASS_DECL_CONVERTARMSMETOLLVM
#define GEN_PASS_DECL_CONVERTARMSMETOSCF
#define GEN_PASS_DECL_CONVERTASYNCTOLLVMPASS
#define GEN_PASS_DECL_CONVERTBUFFERIZATIONTOMEMREF
#define GEN_PASS_DECL_CONVERTCOMPLEXTOLLVMPASS
#define GEN_PASS_DECL_CONVERTCOMPLEXTOLIBM
#define GEN_PASS_DECL_CONVERTCOMPLEXTOSPIRVPASS
#define GEN_PASS_DECL_CONVERTCOMPLEXTOSTANDARD
#define GEN_PASS_DECL_CONVERTCONTROLFLOWTOLLVMPASS
#define GEN_PASS_DECL_CONVERTCONTROLFLOWTOSPIRV
#define GEN_PASS_DECL_CONVERTFUNCTOEMITC
#define GEN_PASS_DECL_CONVERTFUNCTOLLVMPASS
#define GEN_PASS_DECL_CONVERTFUNCTOSPIRV
#define GEN_PASS_DECL_CONVERTGPUTOSPIRV
#define GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
#define GEN_PASS_DECL_CONVERTGPUOPSTOLLVMSPVOPS
#define GEN_PASS_DECL_CONVERTGPUOPSTONVVMOPS
#define GEN_PASS_DECL_CONVERTGPUOPSTOROCDLOPS
#define GEN_PASS_DECL_CONVERTINDEXTOLLVMPASS
#define GEN_PASS_DECL_CONVERTINDEXTOSPIRVPASS
#define GEN_PASS_DECL_CONVERTLINALGTOSTANDARD
#define GEN_PASS_DECL_CONVERTMATHTOFUNCS
#define GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
#define GEN_PASS_DECL_CONVERTMATHTOLIBM
#define GEN_PASS_DECL_CONVERTMATHTOROCDL
#define GEN_PASS_DECL_CONVERTMATHTOSPIRV
#define GEN_PASS_DECL_CONVERTMEMREFTOEMITC
#define GEN_PASS_DECL_CONVERTMEMREFTOSPIRV
#define GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS
#define GEN_PASS_DECL_CONVERTNVVMTOLLVMPASS
#define GEN_PASS_DECL_CONVERTOPENACCTOSCF
#define GEN_PASS_DECL_CONVERTOPENMPTOLLVMPASS
#define GEN_PASS_DECL_CONVERTPDLTOPDLINTERP
#define GEN_PASS_DECL_CONVERTPARALLELLOOPTOGPU
#define GEN_PASS_DECL_CONVERTSCFTOOPENMPPASS
#define GEN_PASS_DECL_CONVERTSPIRVTOLLVMPASS
#define GEN_PASS_DECL_CONVERTSHAPECONSTRAINTS
#define GEN_PASS_DECL_CONVERTSHAPETOSTANDARD
#define GEN_PASS_DECL_CONVERTTENSORTOLINALG
#define GEN_PASS_DECL_CONVERTTENSORTOSPIRV
#define GEN_PASS_DECL_CONVERTTOLLVMPASS
#define GEN_PASS_DECL_CONVERTTOSPIRVPASS
#define GEN_PASS_DECL_CONVERTVECTORTOARMSME
#define GEN_PASS_DECL_CONVERTVECTORTOGPU
#define GEN_PASS_DECL_CONVERTVECTORTOLLVMPASS
#define GEN_PASS_DECL_CONVERTVECTORTOSCF
#define GEN_PASS_DECL_CONVERTVECTORTOSPIRV
#define GEN_PASS_DECL_CONVERTVECTORTOXEGPU
#define GEN_PASS_DECL_CONVERTVULKANLAUNCHFUNCTOVULKANCALLSPASS
#define GEN_PASS_DECL_FINALIZEMEMREFTOLLVMCONVERSIONPASS
#define GEN_PASS_DECL_GPUTOLLVMCONVERSIONPASS
#define GEN_PASS_DECL_LIFTCONTROLFLOWTOSCFPASS
#define GEN_PASS_DECL_LOWERHOSTCODETOLLVMPASS
#define GEN_PASS_DECL_MAPMEMREFSTORAGECLASS
#define GEN_PASS_DECL_RECONCILEUNREALIZEDCASTS
#define GEN_PASS_DECL_SCFTOCONTROLFLOW
#define GEN_PASS_DECL_SCFTOEMITC
#define GEN_PASS_DECL_SCFTOSPIRV
#define GEN_PASS_DECL_SETLLVMMODULEDATALAYOUTPASS
#define GEN_PASS_DECL_TOSATOARITH
#define GEN_PASS_DECL_TOSATOLINALG
#define GEN_PASS_DECL_TOSATOLINALGNAMED
#define GEN_PASS_DECL_TOSATOMLPROGRAM
#define GEN_PASS_DECL_TOSATOSCF
#define GEN_PASS_DECL_TOSATOTENSOR
#define GEN_PASS_DECL_UBTOLLVMCONVERSIONPASS
#define GEN_PASS_DECL_UBTOSPIRVCONVERSIONPASS
#undef GEN_PASS_DECL
#endif // GEN_PASS_DECL

//===----------------------------------------------------------------------===//
// ArithToAMDGPUConversionPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ARITHTOAMDGPUCONVERSIONPASS
struct ArithToAMDGPUConversionPassOptions {
  std::string chipset = "gfx000";
  bool saturateFP8Truncf = false;
  bool allowPackedF16Rtz = false;
};
std::unique_ptr<::mlir::Pass> createArithToAMDGPUConversionPass();
std::unique_ptr<::mlir::Pass> createArithToAMDGPUConversionPass(ArithToAMDGPUConversionPassOptions options);
#undef GEN_PASS_DECL_ARITHTOAMDGPUCONVERSIONPASS
#endif // GEN_PASS_DECL_ARITHTOAMDGPUCONVERSIONPASS
#ifdef GEN_PASS_DEF_ARITHTOAMDGPUCONVERSIONPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createArithToAMDGPUConversionPass(ArithToAMDGPUConversionPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arith operations to AMDGPU-specific implementations"; }

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

  /// 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<amdgpu::AMDGPUDialect>();
    registry.insert<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(ArithToAMDGPUConversionPassBase<DerivedT>)

  ArithToAMDGPUConversionPassBase(ArithToAMDGPUConversionPassOptions options) : ArithToAMDGPUConversionPassBase() {
    chipset = std::move(options.chipset);
    saturateFP8Truncf = std::move(options.saturateFP8Truncf);
    allowPackedF16Rtz = std::move(options.allowPackedF16Rtz);
  }
protected:
  ::mlir::Pass::Option<std::string> chipset{*this, "chipset", ::llvm::cl::desc("Chipset that these operations will run on"), ::llvm::cl::init("gfx000")};
  ::mlir::Pass::Option<bool> saturateFP8Truncf{*this, "saturate-fp8-truncf", ::llvm::cl::desc("Use saturating truncation for 8-bit float types"), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> allowPackedF16Rtz{*this, "allow-packed-f16-round-to-zero", ::llvm::cl::desc("Whether we should allow f32->f16 packed round-to-zero conversion"), ::llvm::cl::init(false)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createArithToAMDGPUConversionPass(ArithToAMDGPUConversionPassOptions options) {
  return impl::createArithToAMDGPUConversionPass(std::move(options));
}
#undef GEN_PASS_DEF_ARITHTOAMDGPUCONVERSIONPASS
#endif // GEN_PASS_DEF_ARITHTOAMDGPUCONVERSIONPASS

//===----------------------------------------------------------------------===//
// ArithToArmSMEConversionPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ARITHTOARMSMECONVERSIONPASS
std::unique_ptr<::mlir::Pass> createArithToArmSMEConversionPass();
#undef GEN_PASS_DECL_ARITHTOARMSMECONVERSIONPASS
#endif // GEN_PASS_DECL_ARITHTOARMSMECONVERSIONPASS
#ifdef GEN_PASS_DEF_ARITHTOARMSMECONVERSIONPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arith dialect to ArmSME dialect"; }

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ArithToLLVMConversionPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_ARITHTOLLVMCONVERSIONPASS
struct ArithToLLVMConversionPassOptions {
  unsigned indexBitwidth = 0;
};
std::unique_ptr<::mlir::Pass> createArithToLLVMConversionPass();
std::unique_ptr<::mlir::Pass> createArithToLLVMConversionPass(ArithToLLVMConversionPassOptions options);
#undef GEN_PASS_DECL_ARITHTOLLVMCONVERSIONPASS
#endif // GEN_PASS_DECL_ARITHTOLLVMCONVERSIONPASS
#ifdef GEN_PASS_DEF_ARITHTOLLVMCONVERSIONPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createArithToLLVMConversionPass(ArithToLLVMConversionPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arith dialect to LLVM dialect"; }

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

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

  ArithToLLVMConversionPassBase(ArithToLLVMConversionPassOptions options) : ArithToLLVMConversionPassBase() {
    indexBitwidth = std::move(options.indexBitwidth);
  }
protected:
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createArithToLLVMConversionPass(ArithToLLVMConversionPassOptions options) {
  return impl::createArithToLLVMConversionPass(std::move(options));
}
#undef GEN_PASS_DEF_ARITHTOLLVMCONVERSIONPASS
#endif // GEN_PASS_DEF_ARITHTOLLVMCONVERSIONPASS

//===----------------------------------------------------------------------===//
// ConvertAMDGPUToROCDL
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTAMDGPUTOROCDL
struct ConvertAMDGPUToROCDLOptions {
  std::string chipset = "gfx000";
};
#undef GEN_PASS_DECL_CONVERTAMDGPUTOROCDL
#endif // GEN_PASS_DECL_CONVERTAMDGPUTOROCDL
#ifdef GEN_PASS_DEF_CONVERTAMDGPUTOROCDL
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert AMDGPU dialect to ROCDL dialect"; }

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

  /// 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<LLVM::LLVMDialect>();
    registry.insert<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(ConvertAMDGPUToROCDLBase<DerivedT>)

  ConvertAMDGPUToROCDLBase(ConvertAMDGPUToROCDLOptions options) : ConvertAMDGPUToROCDLBase() {
    chipset = std::move(options.chipset);
  }
protected:
  ::mlir::Pass::Option<std::string> chipset{*this, "chipset", ::llvm::cl::desc("Chipset that these operations will run on"), ::llvm::cl::init("gfx000")};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTAMDGPUTOROCDL
#endif // GEN_PASS_DEF_CONVERTAMDGPUTOROCDL

//===----------------------------------------------------------------------===//
// ConvertAffineForToGPU
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTAFFINEFORTOGPU
struct ConvertAffineForToGPUOptions {
  unsigned numBlockDims = 1u;
  unsigned numThreadDims = 1u;
};
#undef GEN_PASS_DECL_CONVERTAFFINEFORTOGPU
#endif // GEN_PASS_DECL_CONVERTAFFINEFORTOGPU
#ifdef GEN_PASS_DEF_CONVERTAFFINEFORTOGPU
namespace impl {

template <typename DerivedT>
class ConvertAffineForToGPUBase : public ::mlir::InterfacePass<FunctionOpInterface> {
public:
  using Base = ConvertAffineForToGPUBase;

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

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

  ::llvm::StringRef getDescription() const override { return "Convert top-level AffineFor Ops to GPU kernels"; }

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

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

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

  ConvertAffineForToGPUBase(ConvertAffineForToGPUOptions options) : ConvertAffineForToGPUBase() {
    numBlockDims = std::move(options.numBlockDims);
    numThreadDims = std::move(options.numThreadDims);
  }
protected:
  ::mlir::Pass::Option<unsigned> numBlockDims{*this, "gpu-block-dims", ::llvm::cl::desc("Number of GPU block dimensions for mapping"), ::llvm::cl::init(1u)};
  ::mlir::Pass::Option<unsigned> numThreadDims{*this, "gpu-thread-dims", ::llvm::cl::desc("Number of GPU thread dimensions for mapping"), ::llvm::cl::init(1u)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTAFFINEFORTOGPU
#endif // GEN_PASS_DEF_CONVERTAFFINEFORTOGPU

//===----------------------------------------------------------------------===//
// ConvertAffineToStandard
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTAFFINETOSTANDARD
#undef GEN_PASS_DECL_CONVERTAFFINETOSTANDARD
#endif // GEN_PASS_DECL_CONVERTAFFINETOSTANDARD
#ifdef GEN_PASS_DEF_CONVERTAFFINETOSTANDARD
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lower Affine operations to a combination of Standard and SCF operations"; }

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

  /// 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<memref::MemRefDialect>();
    registry.insert<scf::SCFDialect>();
    registry.insert<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(ConvertAffineToStandardBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// ConvertArithToEmitC
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTARITHTOEMITC
std::unique_ptr<::mlir::Pass> createConvertArithToEmitC();
#undef GEN_PASS_DECL_CONVERTARITHTOEMITC
#endif // GEN_PASS_DECL_CONVERTARITHTOEMITC
#ifdef GEN_PASS_DEF_CONVERTARITHTOEMITC

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arith dialect to EmitC dialect"; }

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertArithToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTARITHTOSPIRV
struct ConvertArithToSPIRVOptions {
  bool emulateLT32BitScalarTypes = true;
};
#undef GEN_PASS_DECL_CONVERTARITHTOSPIRV
#endif // GEN_PASS_DECL_CONVERTARITHTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTARITHTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arith dialect to SPIR-V dialect"; }

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

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

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

  ConvertArithToSPIRVBase(ConvertArithToSPIRVOptions options) : ConvertArithToSPIRVBase() {
    emulateLT32BitScalarTypes = std::move(options.emulateLT32BitScalarTypes);
  }
protected:
  ::mlir::Pass::Option<bool> emulateLT32BitScalarTypes{*this, "emulate-lt-32-bit-scalar-types", ::llvm::cl::desc("Emulate narrower scalar types with 32-bit ones if not supported by the target"), ::llvm::cl::init(true)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTARITHTOSPIRV
#endif // GEN_PASS_DEF_CONVERTARITHTOSPIRV

//===----------------------------------------------------------------------===//
// ConvertArmNeon2dToIntr
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTARMNEON2DTOINTR
#undef GEN_PASS_DECL_CONVERTARMNEON2DTOINTR
#endif // GEN_PASS_DECL_CONVERTARMNEON2DTOINTR
#ifdef GEN_PASS_DEF_CONVERTARMNEON2DTOINTR
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Arm NEON structured ops to intrinsics"; }

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

  /// 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<arm_neon::ArmNeonDialect>();
    registry.insert<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(ConvertArmNeon2dToIntrBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// ConvertArmSMEToLLVM
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTARMSMETOLLVM
struct ConvertArmSMEToLLVMOptions {
  bool dumpTileLiveRanges = false;
};
#undef GEN_PASS_DECL_CONVERTARMSMETOLLVM
#endif // GEN_PASS_DECL_CONVERTARMSMETOLLVM
#ifdef GEN_PASS_DEF_CONVERTARMSMETOLLVM
namespace impl {

template <typename DerivedT>
class ConvertArmSMEToLLVMBase : public ::mlir::InterfacePass<FunctionOpInterface> {
public:
  using Base = ConvertArmSMEToLLVMBase;

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

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

  ::llvm::StringRef getDescription() const override { return "Lower the operations from the ArmSME dialect into the LLVM dialect"; }

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

  /// 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<arm_sme::ArmSMEDialect>();
    registry.insert<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(ConvertArmSMEToLLVMBase<DerivedT>)

  ConvertArmSMEToLLVMBase(ConvertArmSMEToLLVMOptions options) : ConvertArmSMEToLLVMBase() {
    dumpTileLiveRanges = std::move(options.dumpTileLiveRanges);
  }
protected:
  ::mlir::Pass::Option<bool> dumpTileLiveRanges{*this, "dump-tile-live-ranges", ::llvm::cl::desc("Dump the live ranges of SME tiles (for debugging)"), ::llvm::cl::init(false)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTARMSMETOLLVM
#endif // GEN_PASS_DEF_CONVERTARMSMETOLLVM

//===----------------------------------------------------------------------===//
// ConvertArmSMEToSCF
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTARMSMETOSCF
#undef GEN_PASS_DECL_CONVERTARMSMETOSCF
#endif // GEN_PASS_DECL_CONVERTARMSMETOSCF
#ifdef GEN_PASS_DEF_CONVERTARMSMETOSCF
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lower the operations from the ArmSME dialect into the SCF dialect"; }

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

  /// 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<scf::SCFDialect>();
    registry.insert<arith::ArithDialect>();
    registry.insert<vector::VectorDialect>();
    registry.insert<arm_sme::ArmSMEDialect>();
  }

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

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

//===----------------------------------------------------------------------===//
// ConvertAsyncToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTASYNCTOLLVMPASS
std::unique_ptr<::mlir::Pass> createConvertAsyncToLLVMPass();
#undef GEN_PASS_DECL_CONVERTASYNCTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTASYNCTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTASYNCTOLLVMPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert the operations from the async dialect into the LLVM dialect"; }

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

  /// 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<arith::ArithDialect>();
    registry.insert<async::AsyncDialect>();
    registry.insert<LLVM::LLVMDialect>();
    registry.insert<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(ConvertAsyncToLLVMPassBase<DerivedT>)

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertBufferizationToMemRef
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTBUFFERIZATIONTOMEMREF
#undef GEN_PASS_DECL_CONVERTBUFFERIZATIONTOMEMREF
#endif // GEN_PASS_DECL_CONVERTBUFFERIZATIONTOMEMREF
#ifdef GEN_PASS_DEF_CONVERTBUFFERIZATIONTOMEMREF
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert operations from the Bufferization dialect to the MemRef dialect"; }

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

  /// 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<arith::ArithDialect>();
    registry.insert<memref::MemRefDialect>();
    registry.insert<scf::SCFDialect>();
    registry.insert<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(ConvertBufferizationToMemRefBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// ConvertComplexToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCOMPLEXTOLLVMPASS
std::unique_ptr<::mlir::Pass> createConvertComplexToLLVMPass();
#undef GEN_PASS_DECL_CONVERTCOMPLEXTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTCOMPLEXTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTCOMPLEXTOLLVMPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Complex dialect to LLVM dialect"; }

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertComplexToLibm
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCOMPLEXTOLIBM
#undef GEN_PASS_DECL_CONVERTCOMPLEXTOLIBM
#endif // GEN_PASS_DECL_CONVERTCOMPLEXTOLIBM
#ifdef GEN_PASS_DEF_CONVERTCOMPLEXTOLIBM
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Complex dialect to libm calls"; }

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

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

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

//===----------------------------------------------------------------------===//
// ConvertComplexToSPIRVPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCOMPLEXTOSPIRVPASS
std::unique_ptr<::mlir::Pass> createConvertComplexToSPIRVPass();
#undef GEN_PASS_DECL_CONVERTCOMPLEXTOSPIRVPASS
#endif // GEN_PASS_DECL_CONVERTCOMPLEXTOSPIRVPASS
#ifdef GEN_PASS_DEF_CONVERTCOMPLEXTOSPIRVPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Complex dialect to SPIRV dialect"; }

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertComplexToStandard
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCOMPLEXTOSTANDARD
#undef GEN_PASS_DECL_CONVERTCOMPLEXTOSTANDARD
#endif // GEN_PASS_DECL_CONVERTCOMPLEXTOSTANDARD
#ifdef GEN_PASS_DEF_CONVERTCOMPLEXTOSTANDARD
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Complex dialect to standard dialect"; }

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

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

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

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

//===----------------------------------------------------------------------===//
// ConvertControlFlowToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCONTROLFLOWTOLLVMPASS
struct ConvertControlFlowToLLVMPassOptions {
  unsigned indexBitwidth = 0;
};
std::unique_ptr<::mlir::Pass> createConvertControlFlowToLLVMPass();
std::unique_ptr<::mlir::Pass> createConvertControlFlowToLLVMPass(ConvertControlFlowToLLVMPassOptions options);
#undef GEN_PASS_DECL_CONVERTCONTROLFLOWTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTCONTROLFLOWTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTCONTROLFLOWTOLLVMPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertControlFlowToLLVMPass(ConvertControlFlowToLLVMPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert ControlFlow operations to the LLVM dialect"; }

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

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

  ConvertControlFlowToLLVMPassBase(ConvertControlFlowToLLVMPassOptions options) : ConvertControlFlowToLLVMPassBase() {
    indexBitwidth = std::move(options.indexBitwidth);
  }
protected:
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertControlFlowToLLVMPass(ConvertControlFlowToLLVMPassOptions options) {
  return impl::createConvertControlFlowToLLVMPass(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTCONTROLFLOWTOLLVMPASS
#endif // GEN_PASS_DEF_CONVERTCONTROLFLOWTOLLVMPASS

//===----------------------------------------------------------------------===//
// ConvertControlFlowToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTCONTROLFLOWTOSPIRV
struct ConvertControlFlowToSPIRVOptions {
  bool emulateLT32BitScalarTypes = true;
};
#undef GEN_PASS_DECL_CONVERTCONTROLFLOWTOSPIRV
#endif // GEN_PASS_DECL_CONVERTCONTROLFLOWTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTCONTROLFLOWTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert ControlFlow dialect to SPIR-V dialect"; }

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

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

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

  ConvertControlFlowToSPIRVBase(ConvertControlFlowToSPIRVOptions options) : ConvertControlFlowToSPIRVBase() {
    emulateLT32BitScalarTypes = std::move(options.emulateLT32BitScalarTypes);
  }
protected:
  ::mlir::Pass::Option<bool> emulateLT32BitScalarTypes{*this, "emulate-lt-32-bit-scalar-types", ::llvm::cl::desc("Emulate narrower scalar types with 32-bit ones if not supported by the target"), ::llvm::cl::init(true)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTCONTROLFLOWTOSPIRV
#endif // GEN_PASS_DEF_CONVERTCONTROLFLOWTOSPIRV

//===----------------------------------------------------------------------===//
// ConvertFuncToEmitC
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTFUNCTOEMITC
std::unique_ptr<::mlir::Pass> createConvertFuncToEmitC();
#undef GEN_PASS_DECL_CONVERTFUNCTOEMITC
#endif // GEN_PASS_DECL_CONVERTFUNCTOEMITC
#ifdef GEN_PASS_DEF_CONVERTFUNCTOEMITC

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Func dialect to EmitC dialect"; }

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertFuncToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTFUNCTOLLVMPASS
struct ConvertFuncToLLVMPassOptions {
  bool useBarePtrCallConv = false;
  unsigned indexBitwidth = 0;
};
std::unique_ptr<::mlir::Pass> createConvertFuncToLLVMPass();
std::unique_ptr<::mlir::Pass> createConvertFuncToLLVMPass(ConvertFuncToLLVMPassOptions options);
#undef GEN_PASS_DECL_CONVERTFUNCTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTFUNCTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTFUNCTOLLVMPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertFuncToLLVMPass(ConvertFuncToLLVMPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert from the Func dialect to the LLVM dialect"; }

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

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

  ConvertFuncToLLVMPassBase(ConvertFuncToLLVMPassOptions options) : ConvertFuncToLLVMPassBase() {
    useBarePtrCallConv = std::move(options.useBarePtrCallConv);
    indexBitwidth = std::move(options.indexBitwidth);
  }
protected:
  ::mlir::Pass::Option<bool> useBarePtrCallConv{*this, "use-bare-ptr-memref-call-conv", ::llvm::cl::desc("Replace FuncOp's MemRef arguments with bare pointers to the MemRef element types"), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertFuncToLLVMPass(ConvertFuncToLLVMPassOptions options) {
  return impl::createConvertFuncToLLVMPass(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTFUNCTOLLVMPASS
#endif // GEN_PASS_DEF_CONVERTFUNCTOLLVMPASS

//===----------------------------------------------------------------------===//
// ConvertFuncToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTFUNCTOSPIRV
struct ConvertFuncToSPIRVOptions {
  bool emulateLT32BitScalarTypes = true;
};
#undef GEN_PASS_DECL_CONVERTFUNCTOSPIRV
#endif // GEN_PASS_DECL_CONVERTFUNCTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTFUNCTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Func dialect to SPIR-V dialect"; }

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

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

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

  ConvertFuncToSPIRVBase(ConvertFuncToSPIRVOptions options) : ConvertFuncToSPIRVBase() {
    emulateLT32BitScalarTypes = std::move(options.emulateLT32BitScalarTypes);
  }
protected:
  ::mlir::Pass::Option<bool> emulateLT32BitScalarTypes{*this, "emulate-lt-32-bit-scalar-types", ::llvm::cl::desc("Emulate narrower scalar types with 32-bit ones if not supported by the target"), ::llvm::cl::init(true)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTFUNCTOSPIRV
#endif // GEN_PASS_DEF_CONVERTFUNCTOSPIRV

//===----------------------------------------------------------------------===//
// ConvertGPUToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTGPUTOSPIRV
struct ConvertGPUToSPIRVOptions {
  bool use64bitIndex = false;
};
#undef GEN_PASS_DECL_CONVERTGPUTOSPIRV
#endif // GEN_PASS_DECL_CONVERTGPUTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTGPUTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert GPU dialect to SPIR-V dialect"; }

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

  /// 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<func::FuncDialect>();
    registry.insert<spirv::SPIRVDialect>();
  }

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

  ConvertGPUToSPIRVBase(ConvertGPUToSPIRVOptions options) : ConvertGPUToSPIRVBase() {
    use64bitIndex = std::move(options.use64bitIndex);
  }
protected:
  ::mlir::Pass::Option<bool> use64bitIndex{*this, "use-64bit-index", ::llvm::cl::desc("Use 64-bit integers to convert index types"), ::llvm::cl::init(false)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTGPUTOSPIRV
#endif // GEN_PASS_DEF_CONVERTGPUTOSPIRV

//===----------------------------------------------------------------------===//
// ConvertGpuLaunchFuncToVulkanLaunchFunc
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
#undef GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
#endif // GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
#ifdef GEN_PASS_DEF_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert gpu.launch_func to vulkanLaunch external call"; }

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

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

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

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

//===----------------------------------------------------------------------===//
// ConvertGpuOpsToLLVMSPVOps
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTGPUOPSTOLLVMSPVOPS
struct ConvertGpuOpsToLLVMSPVOpsOptions {
  unsigned indexBitwidth = 0;
};
std::unique_ptr<::mlir::Pass> createConvertGpuOpsToLLVMSPVOps();
std::unique_ptr<::mlir::Pass> createConvertGpuOpsToLLVMSPVOps(ConvertGpuOpsToLLVMSPVOpsOptions options);
#undef GEN_PASS_DECL_CONVERTGPUOPSTOLLVMSPVOPS
#endif // GEN_PASS_DECL_CONVERTGPUOPSTOLLVMSPVOPS
#ifdef GEN_PASS_DEF_CONVERTGPUOPSTOLLVMSPVOPS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertGpuOpsToLLVMSPVOps(ConvertGpuOpsToLLVMSPVOpsOptions options);
} // namespace impl
namespace impl {

template <typename DerivedT>
class ConvertGpuOpsToLLVMSPVOpsBase : public ::mlir::OperationPass<gpu::GPUModuleOp> {
public:
  using Base = ConvertGpuOpsToLLVMSPVOpsBase;

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

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

  ::llvm::StringRef getDescription() const override { return "Generate LLVM operations to be ingested by a SPIR-V backend for gpu operations"; }

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

  /// 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<LLVM::LLVMDialect>();
    registry.insert<spirv::SPIRVDialect>();
  }

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

  ConvertGpuOpsToLLVMSPVOpsBase(ConvertGpuOpsToLLVMSPVOpsOptions options) : ConvertGpuOpsToLLVMSPVOpsBase() {
    indexBitwidth = std::move(options.indexBitwidth);
  }
protected:
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertGpuOpsToLLVMSPVOps(ConvertGpuOpsToLLVMSPVOpsOptions options) {
  return impl::createConvertGpuOpsToLLVMSPVOps(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTGPUOPSTOLLVMSPVOPS
#endif // GEN_PASS_DEF_CONVERTGPUOPSTOLLVMSPVOPS

//===----------------------------------------------------------------------===//
// ConvertGpuOpsToNVVMOps
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTGPUOPSTONVVMOPS
struct ConvertGpuOpsToNVVMOpsOptions {
  unsigned indexBitwidth = 0;
  bool hasRedux = false;
  bool useBarePtrCallConv = false;
};
std::unique_ptr<::mlir::Pass> createConvertGpuOpsToNVVMOps();
std::unique_ptr<::mlir::Pass> createConvertGpuOpsToNVVMOps(ConvertGpuOpsToNVVMOpsOptions options);
#undef GEN_PASS_DECL_CONVERTGPUOPSTONVVMOPS
#endif // GEN_PASS_DECL_CONVERTGPUOPSTONVVMOPS
#ifdef GEN_PASS_DEF_CONVERTGPUOPSTONVVMOPS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertGpuOpsToNVVMOps(ConvertGpuOpsToNVVMOpsOptions options);
} // namespace impl
namespace impl {

template <typename DerivedT>
class ConvertGpuOpsToNVVMOpsBase : public ::mlir::OperationPass<gpu::GPUModuleOp> {
public:
  using Base = ConvertGpuOpsToNVVMOpsBase;

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

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

  ::llvm::StringRef getDescription() const override { return "Generate NVVM operations for gpu operations"; }

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

  /// 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<cf::ControlFlowDialect>();
    registry.insert<memref::MemRefDialect>();
    registry.insert<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(ConvertGpuOpsToNVVMOpsBase<DerivedT>)

  ConvertGpuOpsToNVVMOpsBase(ConvertGpuOpsToNVVMOpsOptions options) : ConvertGpuOpsToNVVMOpsBase() {
    indexBitwidth = std::move(options.indexBitwidth);
    hasRedux = std::move(options.hasRedux);
    useBarePtrCallConv = std::move(options.useBarePtrCallConv);
  }
protected:
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
  ::mlir::Pass::Option<bool> hasRedux{*this, "has-redux", ::llvm::cl::desc("Target gpu supports redux"), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> useBarePtrCallConv{*this, "use-bare-ptr-memref-call-conv", ::llvm::cl::desc("Replace memref arguments in GPU functions with bare pointers. All memrefs must have static shape."), ::llvm::cl::init(false)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertGpuOpsToNVVMOps(ConvertGpuOpsToNVVMOpsOptions options) {
  return impl::createConvertGpuOpsToNVVMOps(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTGPUOPSTONVVMOPS
#endif // GEN_PASS_DEF_CONVERTGPUOPSTONVVMOPS

//===----------------------------------------------------------------------===//
// ConvertGpuOpsToROCDLOps
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTGPUOPSTOROCDLOPS
struct ConvertGpuOpsToROCDLOpsOptions {
  std::string chipset = "gfx000";
  unsigned indexBitwidth = 0;
  bool useBarePtrCallConv = false;
  ::mlir::gpu::amd::Runtime runtime = ::mlir::gpu::amd::Runtime::Unknown;
};
#undef GEN_PASS_DECL_CONVERTGPUOPSTOROCDLOPS
#endif // GEN_PASS_DECL_CONVERTGPUOPSTOROCDLOPS
#ifdef GEN_PASS_DEF_CONVERTGPUOPSTOROCDLOPS
namespace impl {

template <typename DerivedT>
class ConvertGpuOpsToROCDLOpsBase : public ::mlir::OperationPass<gpu::GPUModuleOp> {
public:
  using Base = ConvertGpuOpsToROCDLOpsBase;

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

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

  ::llvm::StringRef getDescription() const override { return "Generate ROCDL operations for gpu operations"; }

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

  /// 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<ROCDL::ROCDLDialect>();
    registry.insert<cf::ControlFlowDialect>();
    registry.insert<memref::MemRefDialect>();
  }

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

  ConvertGpuOpsToROCDLOpsBase(ConvertGpuOpsToROCDLOpsOptions options) : ConvertGpuOpsToROCDLOpsBase() {
    chipset = std::move(options.chipset);
    indexBitwidth = std::move(options.indexBitwidth);
    useBarePtrCallConv = std::move(options.useBarePtrCallConv);
    runtime = std::move(options.runtime);
  }
protected:
  ::mlir::Pass::Option<std::string> chipset{*this, "chipset", ::llvm::cl::desc("Chipset that these operations will run on"), ::llvm::cl::init("gfx000")};
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
  ::mlir::Pass::Option<bool> useBarePtrCallConv{*this, "use-bare-ptr-memref-call-conv", ::llvm::cl::desc("Replace memref arguments in GPU functions with bare pointers.All memrefs must have static shape"), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<::mlir::gpu::amd::Runtime> runtime{*this, "runtime", ::llvm::cl::desc("Runtime code will be run on (default is Unknown, can also use HIP or OpenCl)"), ::llvm::cl::init(::mlir::gpu::amd::Runtime::Unknown), ::llvm::cl::values(
            clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"),
            clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"),
            clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL")
          )};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTGPUOPSTOROCDLOPS
#endif // GEN_PASS_DEF_CONVERTGPUOPSTOROCDLOPS

//===----------------------------------------------------------------------===//
// ConvertIndexToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTINDEXTOLLVMPASS
struct ConvertIndexToLLVMPassOptions {
  unsigned indexBitwidth = 0;
};
std::unique_ptr<::mlir::Pass> createConvertIndexToLLVMPass();
std::unique_ptr<::mlir::Pass> createConvertIndexToLLVMPass(ConvertIndexToLLVMPassOptions options);
#undef GEN_PASS_DECL_CONVERTINDEXTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTINDEXTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTINDEXTOLLVMPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertIndexToLLVMPass(ConvertIndexToLLVMPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lower the `index` dialect to the `llvm` dialect."; }

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

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

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

  ConvertIndexToLLVMPassBase(ConvertIndexToLLVMPassOptions options) : ConvertIndexToLLVMPassBase() {
    indexBitwidth = std::move(options.indexBitwidth);
  }
protected:
  ::mlir::Pass::Option<unsigned> indexBitwidth{*this, "index-bitwidth", ::llvm::cl::desc("Bitwidth of the index type, 0 to use size of machine word"), ::llvm::cl::init(0)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertIndexToLLVMPass(ConvertIndexToLLVMPassOptions options) {
  return impl::createConvertIndexToLLVMPass(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTINDEXTOLLVMPASS
#endif // GEN_PASS_DEF_CONVERTINDEXTOLLVMPASS

//===----------------------------------------------------------------------===//
// ConvertIndexToSPIRVPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTINDEXTOSPIRVPASS
struct ConvertIndexToSPIRVPassOptions {
  bool use64bitIndex = false;
};
std::unique_ptr<::mlir::Pass> createConvertIndexToSPIRVPass();
std::unique_ptr<::mlir::Pass> createConvertIndexToSPIRVPass(ConvertIndexToSPIRVPassOptions options);
#undef GEN_PASS_DECL_CONVERTINDEXTOSPIRVPASS
#endif // GEN_PASS_DECL_CONVERTINDEXTOSPIRVPASS
#ifdef GEN_PASS_DEF_CONVERTINDEXTOSPIRVPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertIndexToSPIRVPass(ConvertIndexToSPIRVPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Lower the `index` dialect to the `spirv` dialect."; }

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

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

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

  ConvertIndexToSPIRVPassBase(ConvertIndexToSPIRVPassOptions options) : ConvertIndexToSPIRVPassBase() {
    use64bitIndex = std::move(options.use64bitIndex);
  }
protected:
  ::mlir::Pass::Option<bool> use64bitIndex{*this, "use-64bit-index", ::llvm::cl::desc("Use 64-bit integers to convert index types"), ::llvm::cl::init(false)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertIndexToSPIRVPass(ConvertIndexToSPIRVPassOptions options) {
  return impl::createConvertIndexToSPIRVPass(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTINDEXTOSPIRVPASS
#endif // GEN_PASS_DEF_CONVERTINDEXTOSPIRVPASS

//===----------------------------------------------------------------------===//
// ConvertLinalgToStandard
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTLINALGTOSTANDARD
#undef GEN_PASS_DECL_CONVERTLINALGTOSTANDARD
#endif // GEN_PASS_DECL_CONVERTLINALGTOSTANDARD
#ifdef GEN_PASS_DEF_CONVERTLINALGTOSTANDARD
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert the operations from the linalg dialect into the Standard dialect"; }

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

  /// 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<func::FuncDialect>();
    registry.insert<memref::MemRefDialect>();
  }

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

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

//===----------------------------------------------------------------------===//
// ConvertMathToFuncs
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMATHTOFUNCS
struct ConvertMathToFuncsOptions {
  unsigned minWidthOfFPowIExponent = 1;
  bool convertCtlz = false;
};
std::unique_ptr<::mlir::Pass> createConvertMathToFuncs();
std::unique_ptr<::mlir::Pass> createConvertMathToFuncs(ConvertMathToFuncsOptions options);
#undef GEN_PASS_DECL_CONVERTMATHTOFUNCS
#endif // GEN_PASS_DECL_CONVERTMATHTOFUNCS
#ifdef GEN_PASS_DEF_CONVERTMATHTOFUNCS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertMathToFuncs(ConvertMathToFuncsOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Math operations to calls of outlined implementations."; }

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

  /// 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<arith::ArithDialect>();
    registry.insert<cf::ControlFlowDialect>();
    registry.insert<func::FuncDialect>();
    registry.insert<scf::SCFDialect>();
    registry.insert<vector::VectorDialect>();
    registry.insert<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(ConvertMathToFuncsBase<DerivedT>)

  ConvertMathToFuncsBase(ConvertMathToFuncsOptions options) : ConvertMathToFuncsBase() {
    minWidthOfFPowIExponent = std::move(options.minWidthOfFPowIExponent);
    convertCtlz = std::move(options.convertCtlz);
  }
protected:
  ::mlir::Pass::Option<unsigned> minWidthOfFPowIExponent{*this, "min-width-of-fpowi-exponent", ::llvm::cl::desc("Convert FPowI only if the width of its exponent's integer type is greater than or equal to this value"), ::llvm::cl::init(1)};
  ::mlir::Pass::Option<bool> convertCtlz{*this, "convert-ctlz", ::llvm::cl::desc("Convert math.ctlz to a software implementation. Enable for targets that do not natively support ctlz."), ::llvm::cl::init(false)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertMathToFuncs(ConvertMathToFuncsOptions options) {
  return impl::createConvertMathToFuncs(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTMATHTOFUNCS
#endif // GEN_PASS_DEF_CONVERTMATHTOFUNCS

//===----------------------------------------------------------------------===//
// ConvertMathToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
struct ConvertMathToLLVMPassOptions {
  bool approximateLog1p = true;
};
std::unique_ptr<::mlir::Pass> createConvertMathToLLVMPass();
std::unique_ptr<::mlir::Pass> createConvertMathToLLVMPass(ConvertMathToLLVMPassOptions options);
#undef GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTMATHTOLLVMPASS

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

namespace impl {
  std::unique_ptr<::mlir::Pass> createConvertMathToLLVMPass(ConvertMathToLLVMPassOptions options);
} // namespace impl
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Math dialect to LLVM dialect"; }

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

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

  ConvertMathToLLVMPassBase(ConvertMathToLLVMPassOptions options) : ConvertMathToLLVMPassBase() {
    approximateLog1p = std::move(options.approximateLog1p);
  }
protected:
  ::mlir::Pass::Option<bool> approximateLog1p{*this, "approximate-log1p", ::llvm::cl::desc("Enable approximation of Log1p."), ::llvm::cl::init(true)};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createConvertMathToLLVMPass(ConvertMathToLLVMPassOptions options) {
  return impl::createConvertMathToLLVMPass(std::move(options));
}
#undef GEN_PASS_DEF_CONVERTMATHTOLLVMPASS
#endif // GEN_PASS_DEF_CONVERTMATHTOLLVMPASS

//===----------------------------------------------------------------------===//
// ConvertMathToLibm
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMATHTOLIBM
#undef GEN_PASS_DECL_CONVERTMATHTOLIBM
#endif // GEN_PASS_DECL_CONVERTMATHTOLIBM
#ifdef GEN_PASS_DEF_CONVERTMATHTOLIBM
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Math dialect to libm calls"; }

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

  /// 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<arith::ArithDialect>();
    registry.insert<func::FuncDialect>();
    registry.insert<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(ConvertMathToLibmBase<DerivedT>)

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

//===----------------------------------------------------------------------===//
// ConvertMathToROCDL
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMATHTOROCDL
std::unique_ptr<::mlir::Pass> createConvertMathToROCDL();
#undef GEN_PASS_DECL_CONVERTMATHTOROCDL
#endif // GEN_PASS_DECL_CONVERTMATHTOROCDL
#ifdef GEN_PASS_DEF_CONVERTMATHTOROCDL

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Math dialect to ROCDL library calls"; }

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

  /// 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<arith::ArithDialect>();
    registry.insert<func::FuncDialect>();
    registry.insert<ROCDL::ROCDLDialect>();
    registry.insert<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(ConvertMathToROCDLBase<DerivedT>)

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertMathToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMATHTOSPIRV
#undef GEN_PASS_DECL_CONVERTMATHTOSPIRV
#endif // GEN_PASS_DECL_CONVERTMATHTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTMATHTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert Math dialect to SPIR-V dialect"; }

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

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

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

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

//===----------------------------------------------------------------------===//
// ConvertMemRefToEmitC
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMEMREFTOEMITC
std::unique_ptr<::mlir::Pass> createConvertMemRefToEmitC();
#undef GEN_PASS_DECL_CONVERTMEMREFTOEMITC
#endif // GEN_PASS_DECL_CONVERTMEMREFTOEMITC
#ifdef GEN_PASS_DEF_CONVERTMEMREFTOEMITC

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert MemRef dialect to EmitC dialect"; }

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertMemRefToSPIRV
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTMEMREFTOSPIRV
struct ConvertMemRefToSPIRVOptions {
  int boolNumBits = 8;
  bool use64bitIndex = false;
};
#undef GEN_PASS_DECL_CONVERTMEMREFTOSPIRV
#endif // GEN_PASS_DECL_CONVERTMEMREFTOSPIRV
#ifdef GEN_PASS_DEF_CONVERTMEMREFTOSPIRV
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert MemRef dialect to SPIR-V dialect"; }

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

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

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

  ConvertMemRefToSPIRVBase(ConvertMemRefToSPIRVOptions options) : ConvertMemRefToSPIRVBase() {
    boolNumBits = std::move(options.boolNumBits);
    use64bitIndex = std::move(options.use64bitIndex);
  }
protected:
  ::mlir::Pass::Option<int> boolNumBits{*this, "bool-num-bits", ::llvm::cl::desc("The number of bits to store a boolean value"), ::llvm::cl::init(8)};
  ::mlir::Pass::Option<bool> use64bitIndex{*this, "use-64bit-index", ::llvm::cl::desc("Use 64-bit integers to convert index types"), ::llvm::cl::init(false)};
private:
};
} // namespace impl
#undef GEN_PASS_DEF_CONVERTMEMREFTOSPIRV
#endif // GEN_PASS_DEF_CONVERTMEMREFTOSPIRV

//===----------------------------------------------------------------------===//
// ConvertNVGPUToNVVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS
std::unique_ptr<::mlir::Pass> createConvertNVGPUToNVVMPass();
#undef GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS
#endif // GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS
#ifdef GEN_PASS_DEF_CONVERTNVGPUTONVVMPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert NVGPU dialect to NVVM dialect"; }

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertNVVMToLLVMPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTNVVMTOLLVMPASS
std::unique_ptr<::mlir::Pass> createConvertNVVMToLLVMPass();
#undef GEN_PASS_DECL_CONVERTNVVMTOLLVMPASS
#endif // GEN_PASS_DECL_CONVERTNVVMTOLLVMPASS
#ifdef GEN_PASS_DEF_CONVERTNVVMTOLLVMPASS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert NVVM to PTX with Inline Assembly in LLVM dialect"; }

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// ConvertOpenACCToSCF
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_CONVERTOPENACCTOSCF
#undef GEN_PASS_DECL_CONVERTOPENACCTOSCF
#endif // GEN_PASS_DECL_CONVERTOPENACCTOSCF
#ifdef GEN_PASS_DEF_CONVERTOPENACCTOSCF
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Convert the OpenACC ops to OpenACC with SCF dialect"; }

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

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