#ifdef GEN_PASS_DECL
#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
#ifdef GEN_PASS_DECL_GPUASYNCREGIONPASS
#undef GEN_PASS_DECL_GPUASYNCREGIONPASS
#endif
#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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuAsyncRegionPass");
}
::llvm::StringRef getName() const override { return "GpuAsyncRegionPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<async::AsyncDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuAsyncRegionPassBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_GPUASYNCREGIONPASS
#endif
#ifdef GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#undef GEN_PASS_DECL_GPUDECOMPOSEMEMREFSPASS
#endif
#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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuDecomposeMemrefsPass");
}
::llvm::StringRef getName() const override { return "GpuDecomposeMemrefsPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
registry.insert<mlir::memref::MemRefDialect>();
registry.insert<mlir::affine::AffineDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuDecomposeMemrefsPassBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_GPUDECOMPOSEMEMREFSPASS
#endif
#ifdef GEN_PASS_DECL_GPUELIMINATEBARRIERS
std::unique_ptr<::mlir::Pass> createGpuEliminateBarriers();
#undef GEN_PASS_DECL_GPUELIMINATEBARRIERS
#endif
#ifdef GEN_PASS_DEF_GPUELIMINATEBARRIERS
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuEliminateBarriers();
}
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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuEliminateBarriers");
}
::llvm::StringRef getName() const override { return "GpuEliminateBarriers"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
registry.insert<mlir::memref::MemRefDialect>();
registry.insert<mlir::scf::SCFDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuEliminateBarriersBase<DerivedT>)
protected:
private:
friend std::unique_ptr<::mlir::Pass> createGpuEliminateBarriers() {
return std::make_unique<DerivedT>();
}
};
}
std::unique_ptr<::mlir::Pass> createGpuEliminateBarriers() {
return impl::createGpuEliminateBarriers();
}
#undef GEN_PASS_DEF_GPUELIMINATEBARRIERS
#endif
#ifdef GEN_PASS_DECL_GPUKERNELOUTLINING
#undef GEN_PASS_DECL_GPUKERNELOUTLINING
#endif
#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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuKernelOutlining");
}
::llvm::StringRef getName() const override { return "GpuKernelOutlining"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::DLTIDialect>();
registry.insert<cf::ControlFlowDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuKernelOutliningBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_GPUKERNELOUTLINING
#endif
#ifdef GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#undef GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONS
#endif
#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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuLaunchSinkIndexComputations");
}
::llvm::StringRef getName() const override { return "GpuLaunchSinkIndexComputations"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuLaunchSinkIndexComputationsBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS
#endif
#ifdef GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#undef GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#endif
#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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuMapParallelLoopsPass");
}
::llvm::StringRef getName() const override { return "GpuMapParallelLoopsPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(GpuMapParallelLoopsPassBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_GPUMAPPARALLELLOOPSPASS
#endif
#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
#ifdef GEN_PASS_DEF_GPUMODULETOBINARYPASS
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuModuleToBinaryPass();
}
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuModuleToBinaryPass(GpuModuleToBinaryPassOptions options);
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuModuleToBinaryPass");
}
::llvm::StringRef getName() const override { return "GpuModuleToBinaryPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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));
}
};
}
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
#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
#ifdef GEN_PASS_DEF_GPUNVVMATTACHTARGET
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuNVVMAttachTarget();
}
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuNVVMAttachTarget(GpuNVVMAttachTargetOptions options);
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuNVVMAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuNVVMAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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));
}
};
}
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
#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
#ifdef GEN_PASS_DEF_GPUROCDLATTACHTARGET
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuROCDLAttachTarget();
}
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuROCDLAttachTarget(GpuROCDLAttachTargetOptions options);
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuROCDLAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuROCDLAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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));
}
};
}
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
#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
#ifdef GEN_PASS_DEF_GPUSPIRVATTACHTARGET
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuSPIRVAttachTarget();
}
namespace impl {
std::unique_ptr<::mlir::Pass> createGpuSPIRVAttachTarget(GpuSPIRVAttachTargetOptions options);
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuSPIRVAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuSPIRVAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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));
}
};
}
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
#ifdef GEN_PASS_REGISTRATION
inline void registerGpuAsyncRegionPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuAsyncRegionPass();
});
}
inline void registerGpuAsyncRegionPassPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuAsyncRegionPass();
});
}
inline void registerGpuDecomposeMemrefsPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuDecomposeMemrefsPass();
});
}
inline void registerGpuDecomposeMemrefsPassPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuDecomposeMemrefsPass();
});
}
inline void registerGpuEliminateBarriers() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuEliminateBarriers();
});
}
inline void registerGpuEliminateBarriersPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuEliminateBarriers();
});
}
inline void registerGpuKernelOutlining() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuKernelOutliningPass();
});
}
inline void registerGpuKernelOutliningPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuKernelOutliningPass();
});
}
inline void registerGpuLaunchSinkIndexComputations() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuLauchSinkIndexComputationsPass();
});
}
inline void registerGpuLaunchSinkIndexComputationsPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuLauchSinkIndexComputationsPass();
});
}
inline void registerGpuMapParallelLoopsPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuMapParallelLoopsPass();
});
}
inline void registerGpuMapParallelLoopsPassPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::createGpuMapParallelLoopsPass();
});
}
inline void registerGpuModuleToBinaryPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuModuleToBinaryPass();
});
}
inline void registerGpuModuleToBinaryPassPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuModuleToBinaryPass();
});
}
inline void registerGpuNVVMAttachTarget() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuNVVMAttachTarget();
});
}
inline void registerGpuNVVMAttachTargetPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuNVVMAttachTarget();
});
}
inline void registerGpuROCDLAttachTarget() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuROCDLAttachTarget();
});
}
inline void registerGpuROCDLAttachTargetPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuROCDLAttachTarget();
});
}
inline void registerGpuSPIRVAttachTarget() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuSPIRVAttachTarget();
});
}
inline void registerGpuSPIRVAttachTargetPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return createGpuSPIRVAttachTarget();
});
}
inline void registerGPUPasses() {
registerGpuAsyncRegionPass();
registerGpuDecomposeMemrefsPass();
registerGpuEliminateBarriers();
registerGpuKernelOutlining();
registerGpuLaunchSinkIndexComputations();
registerGpuMapParallelLoopsPass();
registerGpuModuleToBinaryPass();
registerGpuNVVMAttachTarget();
registerGpuROCDLAttachTarget();
registerGpuSPIRVAttachTarget();
}
#undef GEN_PASS_REGISTRATION
#endif
#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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuAsyncRegionPass");
}
::llvm::StringRef getName() const override { return "GpuAsyncRegionPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<async::AsyncDialect>();
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuDecomposeMemrefsPass");
}
::llvm::StringRef getName() const override { return "GpuDecomposeMemrefsPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
registry.insert<mlir::memref::MemRefDialect>();
registry.insert<mlir::affine::AffineDialect>();
}
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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuEliminateBarriers");
}
::llvm::StringRef getName() const override { return "GpuEliminateBarriers"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
registry.insert<mlir::memref::MemRefDialect>();
registry.insert<mlir::scf::SCFDialect>();
}
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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuKernelOutlining");
}
::llvm::StringRef getName() const override { return "GpuKernelOutlining"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::DLTIDialect>();
registry.insert<cf::ControlFlowDialect>();
}
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;
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"; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuLaunchSinkIndexComputations");
}
::llvm::StringRef getName() const override { return "GpuLaunchSinkIndexComputations"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuMapParallelLoopsPass");
}
::llvm::StringRef getName() const override { return "GpuMapParallelLoopsPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
registry.insert<mlir::gpu::GPUDialect>();
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuModuleToBinaryPass");
}
::llvm::StringRef getName() const override { return "GpuModuleToBinaryPass"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuNVVMAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuNVVMAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuROCDLAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuROCDLAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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;
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."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("GpuSPIRVAttachTarget");
}
::llvm::StringRef getName() const override { return "GpuSPIRVAttachTarget"; }
static bool classof(const ::mlir::Pass *pass) {
return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
}
std::unique_ptr<::mlir::Pass> clonePass() const override {
return std::make_unique<DerivedT>(*static_cast<const DerivedT *>(this));
}
void getDependentDialects(::mlir::DialectRegistry ®istry) const override {
}
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