Passes.h.inc | Explore in Territory

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

#ifdef GEN_PASS_DECL
// Generate declarations for all passes.
#define GEN_PASS_DECL_GPUASYNCREGIONPASS
#define GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#define GEN_PASS_DECL_GPUELIMINATEBARRIERS
#define GEN_PASS_DECL_GPUKERNELOUTLINING
#define GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#define GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#define GEN_PASS_DECL_GPUMODULETOBINARYPASS
#define GEN_PASS_DECL_GPUNVVMATTACHTARGET
#define GEN_PASS_DECL_GPUROCDLATTACHTARGET
#define GEN_PASS_DECL_GPUSPIRVATTACHTARGET
#undef GEN_PASS_DECL
#endif // GEN_PASS_DECL

//===----------------------------------------------------------------------===//
// GpuAsyncRegionPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUASYNCREGIONPASS
#undef GEN_PASS_DECL_GPUASYNCREGIONPASS
#endif // GEN_PASS_DECL_GPUASYNCREGIONPASS
#ifdef GEN_PASS_DEF_GPUASYNCREGIONPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Make GPU ops async"; }

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

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

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

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

//===----------------------------------------------------------------------===//
// GpuDecomposeMemrefsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#undef GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#endif // GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#ifdef GEN_PASS_DEF_GPUDECOMPOSEMEMREFSPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Decomposes memref index computation into explicit ops."; }

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

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

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

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

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

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

//===----------------------------------------------------------------------===//
// GpuEliminateBarriers
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUELIMINATEBARRIERS
std::unique_ptr<::mlir::Pass> createGpuEliminateBarriers();
#undef GEN_PASS_DECL_GPUELIMINATEBARRIERS
#endif // GEN_PASS_DECL_GPUELIMINATEBARRIERS
#ifdef GEN_PASS_DEF_GPUELIMINATEBARRIERS

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Erase unnecessary barriers"; }

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

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

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

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

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

protected:
private:

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

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

//===----------------------------------------------------------------------===//
// GpuKernelOutlining
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUKERNELOUTLINING
#undef GEN_PASS_DECL_GPUKERNELOUTLINING
#endif // GEN_PASS_DECL_GPUKERNELOUTLINING
#ifdef GEN_PASS_DEF_GPUKERNELOUTLINING
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Outline gpu.launch bodies to kernel functions"; }

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

  /// 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::DLTIDialect>();
    registry.insert<cf::ControlFlowDialect>();
  }

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

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

//===----------------------------------------------------------------------===//
// GpuLaunchSinkIndexComputations
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#undef GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#endif // GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#ifdef GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Sink index computations into gpu.launch body"; }

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

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

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

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

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

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

//===----------------------------------------------------------------------===//
// GpuMapParallelLoopsPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#undef GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#endif // GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#ifdef GEN_PASS_DEF_GPUMAPPARALLELLOOPSPASS
namespace impl {

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

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

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

  ::llvm::StringRef getDescription() const override { return "Greedily maps loops to GPU hardware dimensions."; }

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

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

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

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

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

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

//===----------------------------------------------------------------------===//
// GpuModuleToBinaryPass
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUMODULETOBINARYPASS
struct GpuModuleToBinaryPassOptions { … };
std::unique_ptr<::mlir::Pass> createGpuModuleToBinaryPass();
std::unique_ptr<::mlir::Pass> createGpuModuleToBinaryPass(GpuModuleToBinaryPassOptions options);
#undef GEN_PASS_DECL_GPUMODULETOBINARYPASS
#endif // GEN_PASS_DECL_GPUMODULETOBINARYPASS
#ifdef GEN_PASS_DEF_GPUMODULETOBINARYPASS

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

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Transforms a GPU module into a GPU binary."; }

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

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

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

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

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

  GpuModuleToBinaryPassBase(GpuModuleToBinaryPassOptions options) : GpuModuleToBinaryPassBase() {
    toolkitPath = std::move(options.toolkitPath);
    linkFiles = std::move(options.linkFiles);
    cmdOptions = std::move(options.cmdOptions);
    compilationTarget = std::move(options.compilationTarget);
  }
protected:
  ::mlir::Pass::Option<std::string> toolkitPath{*this, "toolkit", ::llvm::cl::desc("Toolkit path."), ::llvm::cl::init("")};
  ::mlir::Pass::ListOption<std::string> linkFiles{*this, "l", ::llvm::cl::desc("Extra files to link to.")};
  ::mlir::Pass::Option<std::string> cmdOptions{*this, "opts", ::llvm::cl::desc("Command line options to pass to the tools."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> compilationTarget{*this, "format", ::llvm::cl::desc("The target representation of the compilation process."), ::llvm::cl::init("fatbin")};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createGpuModuleToBinaryPass(GpuModuleToBinaryPassOptions options) {
  return impl::createGpuModuleToBinaryPass(std::move(options));
}
#undef GEN_PASS_DEF_GPUMODULETOBINARYPASS
#endif // GEN_PASS_DEF_GPUMODULETOBINARYPASS

//===----------------------------------------------------------------------===//
// GpuNVVMAttachTarget
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUNVVMATTACHTARGET
struct GpuNVVMAttachTargetOptions { … };
std::unique_ptr<::mlir::Pass> createGpuNVVMAttachTarget();
std::unique_ptr<::mlir::Pass> createGpuNVVMAttachTarget(GpuNVVMAttachTargetOptions options);
#undef GEN_PASS_DECL_GPUNVVMATTACHTARGET
#endif // GEN_PASS_DECL_GPUNVVMATTACHTARGET
#ifdef GEN_PASS_DEF_GPUNVVMATTACHTARGET

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

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches an NVVM target attribute to a GPU Module."; }

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

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

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

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

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

  GpuNVVMAttachTargetBase(GpuNVVMAttachTargetOptions options) : GpuNVVMAttachTargetBase() {
    moduleMatcher = std::move(options.moduleMatcher);
    triple = std::move(options.triple);
    chip = std::move(options.chip);
    features = std::move(options.features);
    optLevel = std::move(options.optLevel);
    fastFlag = std::move(options.fastFlag);
    ftzFlag = std::move(options.ftzFlag);
    linkLibs = std::move(options.linkLibs);
  }
protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> triple{*this, "triple", ::llvm::cl::desc("Target triple."), ::llvm::cl::init("nvptx64-nvidia-cuda")};
  ::mlir::Pass::Option<std::string> chip{*this, "chip", ::llvm::cl::desc("Target chip."), ::llvm::cl::init("sm_50")};
  ::mlir::Pass::Option<std::string> features{*this, "features", ::llvm::cl::desc("Target features."), ::llvm::cl::init("+ptx60")};
  ::mlir::Pass::Option<unsigned> optLevel{*this, "O", ::llvm::cl::desc("Optimization level."), ::llvm::cl::init(2)};
  ::mlir::Pass::Option<bool> fastFlag{*this, "fast", ::llvm::cl::desc("Enable fast math mode."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> ftzFlag{*this, "ftz", ::llvm::cl::desc("Enable flush to zero for denormals."), ::llvm::cl::init(false)};
  ::mlir::Pass::ListOption<std::string> linkLibs{*this, "l", ::llvm::cl::desc("Extra bitcode libraries paths to link to.")};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createGpuNVVMAttachTarget(GpuNVVMAttachTargetOptions options) {
  return impl::createGpuNVVMAttachTarget(std::move(options));
}
#undef GEN_PASS_DEF_GPUNVVMATTACHTARGET
#endif // GEN_PASS_DEF_GPUNVVMATTACHTARGET

//===----------------------------------------------------------------------===//
// GpuROCDLAttachTarget
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUROCDLATTACHTARGET
struct GpuROCDLAttachTargetOptions { … };
std::unique_ptr<::mlir::Pass> createGpuROCDLAttachTarget();
std::unique_ptr<::mlir::Pass> createGpuROCDLAttachTarget(GpuROCDLAttachTargetOptions options);
#undef GEN_PASS_DECL_GPUROCDLATTACHTARGET
#endif // GEN_PASS_DECL_GPUROCDLATTACHTARGET
#ifdef GEN_PASS_DEF_GPUROCDLATTACHTARGET

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

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches a ROCDL target attribute to a GPU Module."; }

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

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

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

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

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

  GpuROCDLAttachTargetBase(GpuROCDLAttachTargetOptions options) : GpuROCDLAttachTargetBase() {
    moduleMatcher = std::move(options.moduleMatcher);
    triple = std::move(options.triple);
    chip = std::move(options.chip);
    features = std::move(options.features);
    abiVersion = std::move(options.abiVersion);
    optLevel = std::move(options.optLevel);
    wave64Flag = std::move(options.wave64Flag);
    fastFlag = std::move(options.fastFlag);
    dazFlag = std::move(options.dazFlag);
    finiteOnlyFlag = std::move(options.finiteOnlyFlag);
    unsafeMathFlag = std::move(options.unsafeMathFlag);
    correctSqrtFlag = std::move(options.correctSqrtFlag);
    linkLibs = std::move(options.linkLibs);
  }
protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> triple{*this, "triple", ::llvm::cl::desc("Target triple."), ::llvm::cl::init("amdgcn-amd-amdhsa")};
  ::mlir::Pass::Option<std::string> chip{*this, "chip", ::llvm::cl::desc("Target chip."), ::llvm::cl::init("gfx900")};
  ::mlir::Pass::Option<std::string> features{*this, "features", ::llvm::cl::desc("Target features."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> abiVersion{*this, "abi", ::llvm::cl::desc("ABI version."), ::llvm::cl::init("500")};
  ::mlir::Pass::Option<unsigned> optLevel{*this, "O", ::llvm::cl::desc("Optimization level."), ::llvm::cl::init(2)};
  ::mlir::Pass::Option<bool> wave64Flag{*this, "wave64", ::llvm::cl::desc("Use Wave64 mode."), ::llvm::cl::init(true)};
  ::mlir::Pass::Option<bool> fastFlag{*this, "fast", ::llvm::cl::desc("Enable fast relaxed math opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> dazFlag{*this, "daz", ::llvm::cl::desc("Enable denormals are zero opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> finiteOnlyFlag{*this, "finite-only", ::llvm::cl::desc("Enable finite only opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> unsafeMathFlag{*this, "unsafe-math", ::llvm::cl::desc("Enable unsafe math opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> correctSqrtFlag{*this, "correct-sqrt", ::llvm::cl::desc("Enable correct rounded sqrt."), ::llvm::cl::init(true)};
  ::mlir::Pass::ListOption<std::string> linkLibs{*this, "l", ::llvm::cl::desc("Extra bitcode libraries paths to link to.")};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createGpuROCDLAttachTarget(GpuROCDLAttachTargetOptions options) {
  return impl::createGpuROCDLAttachTarget(std::move(options));
}
#undef GEN_PASS_DEF_GPUROCDLATTACHTARGET
#endif // GEN_PASS_DEF_GPUROCDLATTACHTARGET

//===----------------------------------------------------------------------===//
// GpuSPIRVAttachTarget
//===----------------------------------------------------------------------===//
#ifdef GEN_PASS_DECL_GPUSPIRVATTACHTARGET
struct GpuSPIRVAttachTargetOptions { … };
std::unique_ptr<::mlir::Pass> createGpuSPIRVAttachTarget();
std::unique_ptr<::mlir::Pass> createGpuSPIRVAttachTarget(GpuSPIRVAttachTargetOptions options);
#undef GEN_PASS_DECL_GPUSPIRVATTACHTARGET
#endif // GEN_PASS_DECL_GPUSPIRVATTACHTARGET
#ifdef GEN_PASS_DEF_GPUSPIRVATTACHTARGET

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

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

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches an SPIR-V target attribute to a GPU Module."; }

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

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

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

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

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

  GpuSPIRVAttachTargetBase(GpuSPIRVAttachTargetOptions options) : GpuSPIRVAttachTargetBase() {
    moduleMatcher = std::move(options.moduleMatcher);
    spirvVersion = std::move(options.spirvVersion);
    spirvCapabilities = std::move(options.spirvCapabilities);
    spirvExtensions = std::move(options.spirvExtensions);
    clientApi = std::move(options.clientApi);
    deviceVendor = std::move(options.deviceVendor);
    deviceType = std::move(options.deviceType);
    deviceId = std::move(options.deviceId);
  }
protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> spirvVersion{*this, "ver", ::llvm::cl::desc("SPIR-V Version."), ::llvm::cl::init("v1.0")};
  ::mlir::Pass::ListOption<std::string> spirvCapabilities{*this, "caps", ::llvm::cl::desc("List of supported SPIR-V Capabilities")};
  ::mlir::Pass::ListOption<std::string> spirvExtensions{*this, "exts", ::llvm::cl::desc("List of supported SPIR-V Extensions")};
  ::mlir::Pass::Option<std::string> clientApi{*this, "client_api", ::llvm::cl::desc("Client API"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<std::string> deviceVendor{*this, "vendor", ::llvm::cl::desc("Device Vendor"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<std::string> deviceType{*this, "device_type", ::llvm::cl::desc("Device Type"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<uint32_t> deviceId{*this, "device_id", ::llvm::cl::desc("Device ID")};
private:

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

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

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

std::unique_ptr<::mlir::Pass> createGpuSPIRVAttachTarget(GpuSPIRVAttachTargetOptions options) {
  return impl::createGpuSPIRVAttachTarget(std::move(options));
}
#undef GEN_PASS_DEF_GPUSPIRVATTACHTARGET
#endif // GEN_PASS_DEF_GPUSPIRVATTACHTARGET
#ifdef GEN_PASS_REGISTRATION

//===----------------------------------------------------------------------===//
// GpuAsyncRegionPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuDecomposeMemrefsPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuEliminateBarriers Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuKernelOutlining Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuLaunchSinkIndexComputations Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuMapParallelLoopsPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuModuleToBinaryPass Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuNVVMAttachTarget Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuROCDLAttachTarget Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GpuSPIRVAttachTarget Registration
//===----------------------------------------------------------------------===//

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

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

//===----------------------------------------------------------------------===//
// GPU Registration
//===----------------------------------------------------------------------===//

inline void registerGPUPasses() {
  registerGpuAsyncRegionPass();
  registerGpuDecomposeMemrefsPass();
  registerGpuEliminateBarriers();
  registerGpuKernelOutlining();
  registerGpuLaunchSinkIndexComputations();
  registerGpuMapParallelLoopsPass();
  registerGpuModuleToBinaryPass();
  registerGpuNVVMAttachTarget();
  registerGpuROCDLAttachTarget();
  registerGpuSPIRVAttachTarget();
}
#undef GEN_PASS_REGISTRATION
#endif // GEN_PASS_REGISTRATION
// Deprecated. Please use the new per-pass macros.
#ifdef GEN_PASS_CLASSES

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

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

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

  ::llvm::StringRef getDescription() const override { return "Make GPU ops async"; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<async::AsyncDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Decomposes memref index computation into explicit ops."; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::gpu::GPUDialect>();
    registry.insert<mlir::memref::MemRefDialect>();
    registry.insert<mlir::affine::AffineDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Erase unnecessary barriers"; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::gpu::GPUDialect>();
    registry.insert<mlir::memref::MemRefDialect>();
    registry.insert<mlir::scf::SCFDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Outline gpu.launch bodies to kernel functions"; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::DLTIDialect>();
    registry.insert<cf::ControlFlowDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Sink index computations into gpu.launch body"; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::gpu::GPUDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Greedily maps loops to GPU hardware dimensions."; }

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

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

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

  /// Register the dialects that must be loaded in the context before this pass.
  void getDependentDialects(::mlir::DialectRegistry &registry) const override {
    registry.insert<mlir::gpu::GPUDialect>();
  }

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

protected:
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Transforms a GPU module into a GPU binary."; }

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

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

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

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

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

protected:
  ::mlir::Pass::Option<std::string> toolkitPath{*this, "toolkit", ::llvm::cl::desc("Toolkit path."), ::llvm::cl::init("")};
  ::mlir::Pass::ListOption<std::string> linkFiles{*this, "l", ::llvm::cl::desc("Extra files to link to.")};
  ::mlir::Pass::Option<std::string> cmdOptions{*this, "opts", ::llvm::cl::desc("Command line options to pass to the tools."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> compilationTarget{*this, "format", ::llvm::cl::desc("The target representation of the compilation process."), ::llvm::cl::init("fatbin")};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches an NVVM target attribute to a GPU Module."; }

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

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

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

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

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

protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> triple{*this, "triple", ::llvm::cl::desc("Target triple."), ::llvm::cl::init("nvptx64-nvidia-cuda")};
  ::mlir::Pass::Option<std::string> chip{*this, "chip", ::llvm::cl::desc("Target chip."), ::llvm::cl::init("sm_50")};
  ::mlir::Pass::Option<std::string> features{*this, "features", ::llvm::cl::desc("Target features."), ::llvm::cl::init("+ptx60")};
  ::mlir::Pass::Option<unsigned> optLevel{*this, "O", ::llvm::cl::desc("Optimization level."), ::llvm::cl::init(2)};
  ::mlir::Pass::Option<bool> fastFlag{*this, "fast", ::llvm::cl::desc("Enable fast math mode."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> ftzFlag{*this, "ftz", ::llvm::cl::desc("Enable flush to zero for denormals."), ::llvm::cl::init(false)};
  ::mlir::Pass::ListOption<std::string> linkLibs{*this, "l", ::llvm::cl::desc("Extra bitcode libraries paths to link to.")};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches a ROCDL target attribute to a GPU Module."; }

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

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

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

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

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

protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> triple{*this, "triple", ::llvm::cl::desc("Target triple."), ::llvm::cl::init("amdgcn-amd-amdhsa")};
  ::mlir::Pass::Option<std::string> chip{*this, "chip", ::llvm::cl::desc("Target chip."), ::llvm::cl::init("gfx900")};
  ::mlir::Pass::Option<std::string> features{*this, "features", ::llvm::cl::desc("Target features."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> abiVersion{*this, "abi", ::llvm::cl::desc("ABI version."), ::llvm::cl::init("500")};
  ::mlir::Pass::Option<unsigned> optLevel{*this, "O", ::llvm::cl::desc("Optimization level."), ::llvm::cl::init(2)};
  ::mlir::Pass::Option<bool> wave64Flag{*this, "wave64", ::llvm::cl::desc("Use Wave64 mode."), ::llvm::cl::init(true)};
  ::mlir::Pass::Option<bool> fastFlag{*this, "fast", ::llvm::cl::desc("Enable fast relaxed math opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> dazFlag{*this, "daz", ::llvm::cl::desc("Enable denormals are zero opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> finiteOnlyFlag{*this, "finite-only", ::llvm::cl::desc("Enable finite only opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> unsafeMathFlag{*this, "unsafe-math", ::llvm::cl::desc("Enable unsafe math opt."), ::llvm::cl::init(false)};
  ::mlir::Pass::Option<bool> correctSqrtFlag{*this, "correct-sqrt", ::llvm::cl::desc("Enable correct rounded sqrt."), ::llvm::cl::init(true)};
  ::mlir::Pass::ListOption<std::string> linkLibs{*this, "l", ::llvm::cl::desc("Extra bitcode libraries paths to link to.")};
};

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

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

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

  ::llvm::StringRef getDescription() const override { return "Attaches an SPIR-V target attribute to a GPU Module."; }

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

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

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

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

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

protected:
  ::mlir::Pass::Option<std::string> moduleMatcher{*this, "module", ::llvm::cl::desc("Regex used to identify the modules to attach the target to."), ::llvm::cl::init("")};
  ::mlir::Pass::Option<std::string> spirvVersion{*this, "ver", ::llvm::cl::desc("SPIR-V Version."), ::llvm::cl::init("v1.0")};
  ::mlir::Pass::ListOption<std::string> spirvCapabilities{*this, "caps", ::llvm::cl::desc("List of supported SPIR-V Capabilities")};
  ::mlir::Pass::ListOption<std::string> spirvExtensions{*this, "exts", ::llvm::cl::desc("List of supported SPIR-V Extensions")};
  ::mlir::Pass::Option<std::string> clientApi{*this, "client_api", ::llvm::cl::desc("Client API"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<std::string> deviceVendor{*this, "vendor", ::llvm::cl::desc("Device Vendor"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<std::string> deviceType{*this, "device_type", ::llvm::cl::desc("Device Type"), ::llvm::cl::init("Unknown")};
  ::mlir::Pass::Option<uint32_t> deviceId{*this, "device_id", ::llvm::cl::desc("Device ID")};
};
#undef GEN_PASS_CLASSES
#endif // GEN_PASS_CLASSES
llvm/tools/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h.inc
