#ifdef GEN_PASS_DECL
#define GEN_PASS_DECL_OPTIMIZESHAREDMEMORY
#undef GEN_PASS_DECL
#endif
#ifdef GEN_PASS_DECL_OPTIMIZESHAREDMEMORY
#undef GEN_PASS_DECL_OPTIMIZESHAREDMEMORY
#endif
#ifdef GEN_PASS_DEF_OPTIMIZESHAREDMEMORY
namespace impl {
template <typename DerivedT>
class OptimizeSharedMemoryBase : public ::mlir::OperationPass<> {
public:
using Base = OptimizeSharedMemoryBase;
OptimizeSharedMemoryBase() : ::mlir::OperationPass<>(::mlir::TypeID::get<DerivedT>()) {}
OptimizeSharedMemoryBase(const OptimizeSharedMemoryBase &other) : ::mlir::OperationPass<>(other) {}
OptimizeSharedMemoryBase& operator=(const OptimizeSharedMemoryBase &) = delete;
OptimizeSharedMemoryBase(OptimizeSharedMemoryBase &&) = delete;
OptimizeSharedMemoryBase& operator=(OptimizeSharedMemoryBase &&) = delete;
~OptimizeSharedMemoryBase() = default;
static constexpr ::llvm::StringLiteral getArgumentName() {
return ::llvm::StringLiteral("nvgpu-optimize-shared-memory");
}
::llvm::StringRef getArgument() const override { return "nvgpu-optimize-shared-memory"; }
::llvm::StringRef getDescription() const override { return "Optimizes accesses to shard memory memrefs in order to reduce bank conflicts."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("OptimizeSharedMemory");
}
::llvm::StringRef getName() const override { return "OptimizeSharedMemory"; }
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<memref::MemRefDialect>();
registry.insert<vector::VectorDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(OptimizeSharedMemoryBase<DerivedT>)
protected:
private:
};
}
#undef GEN_PASS_DEF_OPTIMIZESHAREDMEMORY
#endif
#ifdef GEN_PASS_REGISTRATION
inline void registerOptimizeSharedMemory() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::nvgpu::createOptimizeSharedMemoryPass();
});
}
inline void registerOptimizeSharedMemoryPass() {
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
return mlir::nvgpu::createOptimizeSharedMemoryPass();
});
}
inline void registerNVGPUPasses() {
registerOptimizeSharedMemory();
}
#undef GEN_PASS_REGISTRATION
#endif
#ifdef GEN_PASS_CLASSES
template <typename DerivedT>
class OptimizeSharedMemoryBase : public ::mlir::OperationPass<> {
public:
using Base = OptimizeSharedMemoryBase;
OptimizeSharedMemoryBase() : ::mlir::OperationPass<>(::mlir::TypeID::get<DerivedT>()) {}
OptimizeSharedMemoryBase(const OptimizeSharedMemoryBase &other) : ::mlir::OperationPass<>(other) {}
OptimizeSharedMemoryBase& operator=(const OptimizeSharedMemoryBase &) = delete;
OptimizeSharedMemoryBase(OptimizeSharedMemoryBase &&) = delete;
OptimizeSharedMemoryBase& operator=(OptimizeSharedMemoryBase &&) = delete;
~OptimizeSharedMemoryBase() = default;
static constexpr ::llvm::StringLiteral getArgumentName() {
return ::llvm::StringLiteral("nvgpu-optimize-shared-memory");
}
::llvm::StringRef getArgument() const override { return "nvgpu-optimize-shared-memory"; }
::llvm::StringRef getDescription() const override { return "Optimizes accesses to shard memory memrefs in order to reduce bank conflicts."; }
static constexpr ::llvm::StringLiteral getPassName() {
return ::llvm::StringLiteral("OptimizeSharedMemory");
}
::llvm::StringRef getName() const override { return "OptimizeSharedMemory"; }
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<memref::MemRefDialect>();
registry.insert<vector::VectorDialect>();
}
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(OptimizeSharedMemoryBase<DerivedT>)
protected:
};
#undef GEN_PASS_CLASSES
#endif