2022-06-04 22:10:00 +08:00
|
|
|
#include "triton/Dialect/Triton/IR/Dialect.h"
|
|
|
|
#include "triton/Dialect/TritonGPU/IR/Dialect.h"
|
|
|
|
|
2022-06-06 21:03:58 +08:00
|
|
|
#include "triton/Dialect/Triton/Transforms/Passes.h"
|
|
|
|
#include "triton/Dialect/TritonGPU/Transforms/Passes.h"
|
|
|
|
|
2022-06-07 19:33:51 +08:00
|
|
|
#include "triton/Conversion/Passes.h"
|
|
|
|
|
2022-06-04 22:10:00 +08:00
|
|
|
#include "mlir/IR/Dialect.h"
|
|
|
|
#include "mlir/InitAllPasses.h"
|
|
|
|
#include "mlir/Support/MlirOptMain.h"
|
|
|
|
|
2022-07-26 17:25:03 -07:00
|
|
|
namespace mlir {
|
|
|
|
namespace test {
|
2022-07-19 13:38:48 -07:00
|
|
|
void registerTestAlignmentPass();
|
2022-08-04 11:22:49 -07:00
|
|
|
void registerTestAllocationPass();
|
|
|
|
} // namespace test
|
2022-07-26 17:25:03 -07:00
|
|
|
} // namespace mlir
|
2022-06-04 22:10:00 +08:00
|
|
|
|
|
|
|
int main(int argc, char **argv) {
|
|
|
|
mlir::registerAllPasses();
|
2022-06-06 21:03:58 +08:00
|
|
|
mlir::registerTritonPasses();
|
|
|
|
mlir::registerTritonGPUPasses();
|
2022-07-19 13:38:48 -07:00
|
|
|
mlir::test::registerTestAlignmentPass();
|
2022-08-04 11:22:49 -07:00
|
|
|
mlir::test::registerTestAllocationPass();
|
2022-06-07 19:33:51 +08:00
|
|
|
mlir::triton::registerConvertTritonToTritonGPUPass();
|
2022-08-04 10:15:45 +08:00
|
|
|
mlir::triton::registerConvertTritonGPUToLLVMPass();
|
2022-06-04 22:10:00 +08:00
|
|
|
|
2022-06-06 21:03:58 +08:00
|
|
|
// TODO: register Triton & TritonGPU passes
|
2022-06-04 22:10:00 +08:00
|
|
|
mlir::DialectRegistry registry;
|
2022-07-26 17:25:03 -07:00
|
|
|
registry
|
|
|
|
.insert<mlir::triton::TritonDialect, mlir::triton::gpu::TritonGPUDialect,
|
|
|
|
mlir::arith::ArithmeticDialect, mlir::StandardOpsDialect,
|
|
|
|
mlir::scf::SCFDialect>();
|
2022-06-04 22:10:00 +08:00
|
|
|
|
2022-07-26 17:25:03 -07:00
|
|
|
return mlir::asMainReturnCode(mlir::MlirOptMain(
|
|
|
|
argc, argv, "Triton (GPU) optimizer driver\n", registry));
|
2022-06-04 22:10:00 +08:00
|
|
|
}
|