Rebase to LLVM/MLIR tip.

mlir/include/mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToStd.h → mlir/include/mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h

@@ -19,16 +19,12 @@
 
 namespace mlir {
 class LLVMTypeConverter;
-class FuncOp;
-class ModuleOp;
-template <typename T> class OpPassBase;
-
-namespace miopen {
+class TypeConverter;
+class Pass;
 
 /// Create a pass to convert MIOpen operations to std operations.
-std::unique_ptr<OpPassBase<ModuleOp>> createLowerMIOpenOpsToStdPass();
+std::unique_ptr<Pass> createLowerMIOpenOpsToGPUPass();
 
-} // namespace miopen
 } // namespace mlir
 
 #endif // MLIR_DIALECT_MIOPEN_CONVERT_MIOPEN_OPS_TO_LLVM_H

mlir/include/mlir/Conversion/Passes.td

@@ -292,4 +292,14 @@ def ConvertVectorToLLVM : Pass<"convert-vector-to-llvm", "ModuleOp"> {
   let constructor = "mlir::createConvertVectorToLLVMPass()";
 }
 
+//===----------------------------------------------------------------------===//
+// MIOpenToGPU
+//===----------------------------------------------------------------------===//
+
+def ConvertMIOpenToGPU : Pass<"convert-miopen-to-gpu", "ModuleOp"> {
+  let summary = "Lower the operations from the MIOpen dialect into the GPU "
+                "dialect";
+  let constructor = "mlir::createLowerMIOpenOpsToGPUPass()";
+}
+
 #endif // MLIR_CONVERSION_PASSES

mlir/include/mlir/Dialect/CMakeLists.txt

@@ -4,7 +4,7 @@ add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
 add_subdirectory(LoopOps)
-add_subdirectory(MIOpenOps)
+add_subdirectory(MIOpen)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)
 add_subdirectory(SCF)

mlir/include/mlir/Dialect/MIOpen/CMakeLists.txt

@@ -0,0 +1,8 @@
+add_mlir_dialect(MIOpenOps miopen)
+add_mlir_doc(MIOpenOps -gen-dialect-doc MIOpenDialect Dialects/)
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls)
+add_public_tablegen_target(MLIRMIOpenPassIncGen)
+
+add_mlir_doc(Passes -gen-pass-doc MIOpenPasses ./)

mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h → mlir/include/mlir/Dialect/MIOpen/MIOpenOps.h

@@ -19,21 +19,16 @@
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/IR/StandardTypes.h"
 #include "mlir/IR/Types.h"
-#include "mlir/Transforms/LoopLikeInterface.h"
 
 namespace mlir {
 namespace miopen {
 
 enum ConvOpType { Conv2DOpType, Conv2DBwdDataOpType, Conv2DBwdWeightOpType };
 
-class MIOpenOpsDialect : public Dialect {
-public:
-  MIOpenOpsDialect(MLIRContext *context);
-  static StringRef getDialectNamespace() { return "miopen"; }
-};
+#include "mlir/Dialect/MIOpen/MIOpenOpsDialect.h.inc"
 
 #define GET_OP_CLASSES
-#include "mlir/Dialect/MIOpenOps/MIOpenOps.h.inc"
+#include "mlir/Dialect/MIOpen/MIOpenOps.h.inc"
 
 } // end namespace miopen
 } // end namespace mlir

mlir/include/mlir/Dialect/MIOpenOps/Passes.h → mlir/include/mlir/Dialect/MIOpen/Passes.h

@@ -17,29 +17,27 @@
 #include "llvm/ADT/ArrayRef.h"
 
 namespace mlir {
-class FuncOp;
-class ModuleOp;
-template <typename T> class OpPassBase;
+class Pass;
 
 namespace miopen {
 
 /// Create a pass to convert MIOpen conv2d operations to transform and
 /// gridwise_gemm operations.
-std::unique_ptr<OpPassBase<ModuleOp>> createLowerMIOpenOpsStep1Pass();
+std::unique_ptr<Pass> createLowerMIOpenOpsStep1Pass();
 
 /// Create a pass to convert MIOpen gridwise_gemm operations to blockwise
 /// operations.
-std::unique_ptr<OpPassBase<ModuleOp>> createLowerMIOpenOpsStep2Pass();
+std::unique_ptr<Pass> createLowerMIOpenOpsStep2Pass();
 
 /// Create a pass to convert MIOpen blockwise operations to threadwise
 /// operations.
-std::unique_ptr<OpPassBase<ModuleOp>> createLowerMIOpenOpsStep3Pass();
+std::unique_ptr<Pass> createLowerMIOpenOpsStep3Pass();
 
 /// Create a pass to convert transform operations to affine maps.
-std::unique_ptr<OpPassBase<FuncOp>> createAffineTransformPass();
+std::unique_ptr<Pass> createAffineTransformPass();
 
 /// Create a pass to affix tuning parameters to gridwise gemm ops.
-std::unique_ptr<OpPassBase<FuncOp>> createAffixTuningParametersPass();
+std::unique_ptr<Pass> createAffixTuningParametersPass();
 
 } // namespace miopen
 } // namespace mlir

mlir/include/mlir/Dialect/MIOpen/Passes.td

@@ -0,0 +1,39 @@
+//===-- Passes.td - MIOpen pass definition file ------------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_MIOPEN_PASSES
+#define MLIR_DIALECT_MIOPEN_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def MIOpenOpsStep1Pass : Pass<"miopen-lowering", "ModuleOp"> {
+  let summary = "expand convolution into coordinate transformations and gridwise gemm";
+  let constructor = "mlir::miopen::createLowerMIOpenOpsStep1Pass()";
+}
+
+def MIOpenOpsAffineTransformPass : FunctionPass<"miopen-affine-transform"> {
+  let summary = "build affine maps for miopen.transform ops";
+  let constructor = "mlir::miopen::createAffineTransformPass()";
+}
+
+def MIOpenOpsAffixTuningParametersPass : FunctionPass<"miopen-affix-params"> {
+  let summary = "populate tuning parameters for miopen.gridwise_gemm ops";
+  let constructor = "mlir::miopen::createAffixTuningParametersPass()";
+}
+
+def MIOpenOpsStep2Pass : Pass<"miopen-lowering-step2", "ModuleOp"> {
+  let summary = "expand gridwise gemm into blockwise copy, blockwise gemm, and threadwise copy";
+  let constructor = "mlir::miopen::createLowerMIOpenOpsStep2Pass()";
+}
+
+def MIOpenOpsStep3Pass : Pass<"miopen-lowering-step3", "ModuleOp"> {
+  let summary = "expand blockwise copy into threadwise copy, blockwise gemm into threadwise gemm";
+  let constructor = "mlir::miopen::createLowerMIOpenOpsStep3Pass()";
+}
+
+#endif // MLIR_DIALECT_MIOPEN_PASSES

mlir/include/mlir/InitAllDialects.h

@@ -22,6 +22,8 @@
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
+#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/MIOpen/MIOpenOps.h"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
 #include "mlir/Dialect/Quant/QuantOps.h"
 #include "mlir/Dialect/SCF/SCF.h"
@@ -54,6 +56,7 @@ inline void registerAllDialects() {
     registerDialect<ROCDL::ROCDLDialect>();
     registerDialect<SDBMDialect>();
     registerDialect<shape::ShapeDialect>();
+    registerDialect<miopen::MIOpenDialect>();
     return true;
   }();
   (void)init_once;

mlir/include/mlir/InitAllPasses.h

@@ -23,6 +23,9 @@
 #include "mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h"
 #include "mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h"
 #include "mlir/Conversion/LinalgToStandard/LinalgToStandard.h"
+#include "mlir/Conversion/LoopToStandard/ConvertLoopToStandard.h"
+#include "mlir/Conversion/LoopsToGPU/LoopsToGPUPass.h"
+#include "mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h"
 #include "mlir/Conversion/SCFToGPU/SCFToGPUPass.h"
 #include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
 #include "mlir/Conversion/ShapeToStandard/ShapeToStandard.h"
@@ -34,6 +37,8 @@
 #include "mlir/Dialect/GPU/Passes.h"
 #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
 #include "mlir/Dialect/Linalg/Passes.h"
+#include "mlir/Dialect/LoopOps/Passes.h"
+#include "mlir/Dialect/MIOpen/Passes.h"
 #include "mlir/Dialect/Quant/Passes.h"
 #include "mlir/Dialect/SCF/Passes.h"
 #include "mlir/Dialect/SPIRV/Passes.h"
@@ -94,6 +99,10 @@ inline void registerAllPasses() {
   // Standard
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/StandardOps/Transforms/Passes.h.inc"
+
+  // MIOpen
+#define GEN_PASS_REGISTRATION
+#include "mlir/Dialect/MIOpen/Passes.h.inc"
 }
 
 } // namespace mlir

mlir/include/mlir/InitAllTranslations.h

@@ -23,6 +23,7 @@ void registerToSPIRVTranslation();
 void registerToNVVMIRTranslation();
 void registerToROCDLIRTranslation();
 void registerAVX512ToLLVMIRTranslation();
+void registerToMIOpenCPPTranslation();
 
 // This function should be called before creating any MLIRContext if one
 // expects all the possible translations to be made available to the context
@@ -36,6 +37,7 @@ inline void registerAllTranslations() {
     registerToNVVMIRTranslation();
     registerToROCDLIRTranslation();
     registerAVX512ToLLVMIRTranslation();
+    registerToMIOpenCPPTranslation();
     return true;
   }();
   (void)init_once;

mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h → mlir/include/mlir/Target/MIOpenCPP.h

@@ -13,7 +13,7 @@
 #ifndef MLIR_TARGET_MIOPEN_CPP_H
 #define MLIR_TARGET_MIOPEN_CPP_H
 
-#include "mlir/Dialect/MIOpenOps/MIOpenOps.h"
+#include "mlir/Dialect/MIOpen/MIOpenOps.h"
 #include "mlir/IR/Block.h"
 #include "mlir/IR/Module.h"
 #include "mlir/IR/Value.h"
@@ -149,13 +149,13 @@ class TunableParametersBase {
     yin >> params;
   }
   int operator[](llvm::StringRef str) {
-    if (params.find(str) != params.end()) {
-      return params[str];
+    if (params.find(str.str()) != params.end()) {
+      return params[str.str()];
     }
     return 0;
   }
   void setValue(llvm::StringRef str, int value) {
-    params[str] = value;
+    params[str.str()] = value;
   }
 protected:
   std::map<std::string, int> params;

mlir/lib/Conversion/CMakeLists.txt

@@ -7,6 +7,9 @@ add_subdirectory(GPUToSPIRV)
 add_subdirectory(GPUToVulkan)
 add_subdirectory(LinalgToLLVM)
 add_subdirectory(LinalgToSPIRV)
+add_subdirectory(LoopsToGPU)
+add_subdirectory(LoopToStandard)
+add_subdirectory(MIOpenToGPU)
 add_subdirectory(LinalgToStandard)
 add_subdirectory(SCFToGPU)
 add_subdirectory(SCFToStandard)

mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt

@@ -0,0 +1,23 @@
+add_mlir_conversion_library(MLIRMIOpenToGPU
+  MIOpenToGPU.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Conversion/MIOpenToGPU
+
+  DEPENDS
+  MLIRConversionPassIncGen
+)
+target_link_libraries(MLIRMIOpenToGPU
+  PUBLIC
+  MLIRGPU
+  MLIRLLVMIR
+  MLIRIR
+  MLIRPass
+  MLIRStandardOps
+  MLIRStandardToLLVM
+  MLIRSupport
+  MLIRTransforms
+  MLIRVector
+  LLVMCore
+  LLVMSupport
+  )

mlir/lib/Dialect/MIOpenOps/Transform/ConvertMIOpenOpsToStd.cpp → mlir/lib/Conversion/MIOpenToGPU/MIOpenToGPU.cpp

@@ -1,4 +1,4 @@
-//===- ConvertMIOpenOpsToStd.cpp - MLIR MIOpen ops lowering passes ---------------===//
+//===- MIOpenToGPU.cpp - MLIR MIOpen ops lowering passes ---------------===//
 //
 // Copyright 2020 The MLIR Authors.
 //
@@ -19,14 +19,16 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/MIOpenOps/ConvertMIOpenOpsToStd.h"
+
+#include "mlir/Conversion/MIOpenToGPU/MIOpenToGPU.h"
+#include "../PassDetail.h"
 
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/MIOpenOps/MIOpenOps.h"
-#include "mlir/Dialect/MIOpenOps/Passes.h"
-#include "mlir/Dialect/StandardOps/Ops.h"
+#include "mlir/Dialect/MIOpen/MIOpenOps.h"
+#include "mlir/Dialect/MIOpen/Passes.h"
+#include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/Attributes.h"
@@ -49,13 +51,13 @@
 using namespace mlir;
 
 namespace {
-struct LowerMIOpenOpsToStdPass : public ModulePass<LowerMIOpenOpsToStdPass> {
-  void runOnModule() override;
+struct LowerMIOpenOpsToGPUPass : public ConvertMIOpenToGPUBase<LowerMIOpenOpsToGPUPass> {
+  void runOnOperation() override;
 };
 } // end anonymous namespace
 
-void LowerMIOpenOpsToStdPass::runOnModule() {
-  auto m = getModule();
+void LowerMIOpenOpsToGPUPass::runOnOperation() {
+  auto m = getOperation();
 
   for (auto func : m.getOps<FuncOp>()) {
     LLVMTypeConverter converter(&getContext());
@@ -130,14 +132,15 @@ void LowerMIOpenOpsToStdPass::runOnModule() {
     func.walk([&](miopen::LdsBarrierOp op) {
       OpBuilder b(op.getContext());
       auto loc = op.getLoc();
-      if (!getModule().lookupSymbol<FuncOp>("lds_barrier")) {
+      auto module = op.getParentOfType<ModuleOp>();
+      if (!module.lookupSymbol<FuncOp>("lds_barrier")) {
         auto funcType = b.getFunctionType({}, {});
 
         StringRef funcName = "lds_barrier";
-        b.setInsertionPoint(getModule().getBody(), getModule().getBody()->begin());
+        b.setInsertionPoint(module.getBody(), module.getBody()->begin());
         auto func = b.create<FuncOp>(loc, funcName, funcType, ArrayRef<NamedAttribute>{});
       }
-      auto barrierFunc = getModule().lookupSymbol<FuncOp>("lds_barrier");
+      auto barrierFunc = module.lookupSymbol<FuncOp>("lds_barrier");
       b.setInsertionPoint(op);
       b.create<CallOp>(loc, ArrayRef<Type>{},
                        b.getSymbolRefAttr(barrierFunc),
@@ -147,10 +150,6 @@ void LowerMIOpenOpsToStdPass::runOnModule() {
   }
 }
 
-std::unique_ptr<OpPassBase<ModuleOp>> mlir::miopen::createLowerMIOpenOpsToStdPass() {
-  return std::make_unique<LowerMIOpenOpsToStdPass>();
+std::unique_ptr<Pass> mlir::createLowerMIOpenOpsToGPUPass() {
+  return std::make_unique<LowerMIOpenOpsToGPUPass>();
 }
-
-static PassRegistration<LowerMIOpenOpsToStdPass>
-    lowerMIOpenOpsToStdPass("miopen-lowering-step4",
-                       "Lower MIOpen ops to std dialect.");

mlir/lib/Dialect/CMakeLists.txt

@@ -4,7 +4,7 @@ add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
 add_subdirectory(LoopOps)
-add_subdirectory(MIOpenOps)
+add_subdirectory(MIOpen)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)
 add_subdirectory(SCF)

mlir/lib/Dialect/MIOpen/CMakeLists.txt

@@ -0,0 +1,19 @@
+file(GLOB globbed *.c *.cpp)
+add_mlir_dialect_library(MLIRMIOpenOps
+  ${globbed}
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/MIOpen
+
+  DEPENDS
+  MLIRMIOpenOpsIncGen
+  )
+target_link_libraries(MLIRMIOpenOps
+  PUBLIC
+  MLIRIR
+  MLIRStandardOps
+  LLVMSupport
+  )
+
+add_subdirectory(Driver)
+add_subdirectory(Transforms)

mlir/lib/Dialect/MIOpenOps/Driver/CMakeLists.txt → mlir/lib/Dialect/MIOpen/Driver/CMakeLists.txt

@@ -12,5 +12,9 @@ add_llvm_tool(mlir-miopen-driver
   mlir-miopen-driver.cpp
 )
 llvm_update_compile_flags(mlir-miopen-driver)
-whole_archive_link(mlir-miopen-driver ${LIBS})
-target_link_libraries(mlir-miopen-driver PRIVATE MLIRIR ${LIBS} LLVMSupport)
+get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
+target_link_libraries(mlir-miopen-driver PRIVATE
+  ${dialect_libs}
+  MLIRIR
+  ${LIBS}
+  LLVMSupport)