Rebase to tip of LLVM / MLIR on 5/18/2020.

- Loop dialect is now renamed as SCF.
llvm · Jun 6, 2020 · cdea5aa · cdea5aa
1 parent 014b84f
commit cdea5aa
Show file tree

Hide file tree

Showing 10 changed files with 29 additions and 36 deletions.
diff --git a/mlir/include/mlir/Dialect/CMakeLists.txt b/mlir/include/mlir/Dialect/CMakeLists.txt
@@ -3,7 +3,6 @@ add_subdirectory(AVX512)
 add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
-add_subdirectory(LoopOps)
 add_subdirectory(MIOpen)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)

diff --git a/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h b/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
@@ -12,9 +12,9 @@
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
 #include "mlir/Dialect/MIOpen/MIOpenOps.h"
 #include "mlir/Dialect/MIOpen/Passes.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/IR/AffineMap.h"
@@ -1631,8 +1631,8 @@ struct GridwiseGemmRewritePattern : public OpRewritePattern<miopen::GridwiseGemm
     auto loopIterationConstantIndexOp =
         b.create<ConstantIndexOp>(loc, loopIteration);
     auto loopOp =
-        b.create<loop::ForOp>(loc, zeroConstantIndexOp,
-                              loopIterationConstantIndexOp, oneConstantIndexOp);
+        b.create<scf::ForOp>(loc, zeroConstantIndexOp,
+                             loopIterationConstantIndexOp, oneConstantIndexOp);
 
     // inside the loop.
     auto lb = OpBuilder::atBlockTerminator(loopOp.getBody());
@@ -1855,8 +1855,8 @@ struct BlockwiseGemmRewritePattern : public OpRewritePattern<miopen::BlockwiseGe
     auto loopIterationConstantIndexOp =
         b.create<ConstantIndexOp>(loc, loopIteration);
     auto loopOp =
-        b.create<loop::ForOp>(loc, zeroConstantIndexOp,
-                              loopIterationConstantIndexOp, oneConstantIndexOp);
+        b.create<scf::ForOp>(loc, zeroConstantIndexOp,
+                             loopIterationConstantIndexOp, oneConstantIndexOp);
 
     // inside the main loop.
     auto lb = OpBuilder::atBlockTerminator(loopOp.getBody());
@@ -1868,7 +1868,7 @@ struct BlockwiseGemmRewritePattern : public OpRewritePattern<miopen::BlockwiseGe
     auto loopReadMatrixAIteration = MRepeat;
     auto loopReadMatrixAIterationConstantIndexOp =
         lb.create<ConstantIndexOp>(loc, loopReadMatrixAIteration);
-    auto loopReadMatrixAOp = lb.create<loop::ForOp>(
+    auto loopReadMatrixAOp = lb.create<scf::ForOp>(
         loc, zeroConstantIndexOp, loopReadMatrixAIterationConstantIndexOp,
         oneConstantIndexOp);
 
@@ -1909,7 +1909,7 @@ struct BlockwiseGemmRewritePattern : public OpRewritePattern<miopen::BlockwiseGe
     auto loopReadMatrixBIteration = NRepeat;
     auto loopReadMatrixBIterationConstantIndexOp =
         lb.create<ConstantIndexOp>(loc, loopReadMatrixBIteration);
-    auto loopReadMatrixBOp = lb.create<loop::ForOp>(
+    auto loopReadMatrixBOp = lb.create<scf::ForOp>(
         loc, zeroConstantIndexOp, loopReadMatrixBIterationConstantIndexOp,
         oneConstantIndexOp);
 
@@ -2093,7 +2093,7 @@ struct FillRewritePattern : public OpRewritePattern<miopen::FillOp> {
     if (inputShape.size() == 1) {
       // Rank 1 loop.
       auto loopIteration = b.create<ConstantIndexOp>(loc, inputShape[0]);
-      auto loopOp = b.create<loop::ForOp>(loc, zero, loopIteration, one);
+      auto loopOp = b.create<scf::ForOp>(loc, zero, loopIteration, one);
 
       // inside loop.
       auto lb = OpBuilder::atBlockTerminator(loopOp.getBody());
@@ -2105,7 +2105,7 @@ struct FillRewritePattern : public OpRewritePattern<miopen::FillOp> {
     } else if (inputShape.size() == 2) {
       // Rank 2 loop.
       auto loop0Iteration = b.create<ConstantIndexOp>(loc, inputShape[0]);
-      auto loop0Op = b.create<loop::ForOp>(loc, zero, loop0Iteration, one);
+      auto loop0Op = b.create<scf::ForOp>(loc, zero, loop0Iteration, one);
 
       // inside outer loop.
       auto l0b = OpBuilder::atBlockTerminator(loop0Op.getBody());
@@ -2114,7 +2114,7 @@ struct FillRewritePattern : public OpRewritePattern<miopen::FillOp> {
         auto iter0 = b.create<ConstantIndexOp>(loc, i);
 
         auto loop1Iteration = b.create<ConstantIndexOp>(loc, inputShape[1]);
-        auto loop1Op = l0b.create<loop::ForOp>(loc, zero, loop1Iteration, one);
+        auto loop1Op = l0b.create<scf::ForOp>(loc, zero, loop1Iteration, one);
 
         // inside inner loop.
         auto l1b = OpBuilder::atBlockTerminator(loop1Op.getBody());
@@ -2367,18 +2367,18 @@ struct ThreadwiseCopyRewritePattern
 
       // outer loop.
       auto outerLoopOp =
-          b.create<loop::ForOp>(loc, zeroConstantIndexOp,
-                                NSliceRowConstantIndexOp, oneConstantIndexOp);
+          b.create<scf::ForOp>(loc, zeroConstantIndexOp,
+                               NSliceRowConstantIndexOp, oneConstantIndexOp);
 
       // inside the outer loop.
       auto lob = OpBuilder::atBlockTerminator(outerLoopOp.getBody());
       auto ivo = outerLoopOp.getInductionVar();
       auto ivo_i32 = lob.create<IndexCastOp>(loc, ivo, b.getIntegerType(32));
 
       // inner loop
-      auto innerLoopOp = lob.create<loop::ForOp>(loc, zeroConstantIndexOp,
-                                                 NSliceColConstantIndexOp,
-                                                 DataPerAccessConstantIndexOp);
+      auto innerLoopOp = lob.create<scf::ForOp>(loc, zeroConstantIndexOp,
+                                                NSliceColConstantIndexOp,
+                                                DataPerAccessConstantIndexOp);
 
       // inside the inner loop.
       auto lib = OpBuilder::atBlockTerminator(innerLoopOp.getBody());
@@ -2510,15 +2510,15 @@ struct ThreadwiseCopyRewritePattern
       // llvm::errs() << "\n";
 
       // Emit loops for vector loads / stores.
-      SmallVector<loop::ForOp, 2> loopOps;
+      SmallVector<scf::ForOp, 2> loopOps;
       SmallVector<OpBuilder, 2> loopBuilders;
       SmallVector<Value, 2> loopIVs;
       SmallVector<Value, 2> loopIV_i32s;
       for (unsigned iter = 0; iter < dimAccessOrder.size(); ++iter) {
         auto dim = dimAccessOrder[iter].template cast<IntegerAttr>().getInt();
         auto loopBuilder = (iter == 0) ? b : loopBuilders[iter - 1];
 
-        auto loopOp = loopBuilder.create<loop::ForOp>(
+        auto loopOp = loopBuilder.create<scf::ForOp>(
             loc, zeroConstantIndexOp, loopBounds[dim], oneConstantIndexOp);
         loopOps.push_back(loopOp);
         auto loopOpBuilder = OpBuilder::atBlockTerminator(loopOp.getBody());

diff --git a/mlir/include/mlir/InitAllDialects.h b/mlir/include/mlir/InitAllDialects.h
@@ -22,7 +22,6 @@
 #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"

diff --git a/mlir/include/mlir/InitAllPasses.h b/mlir/include/mlir/InitAllPasses.h
@@ -23,8 +23,6 @@
 #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"
@@ -37,7 +35,6 @@
 #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"

diff --git a/mlir/lib/Conversion/CMakeLists.txt b/mlir/lib/Conversion/CMakeLists.txt
@@ -7,8 +7,6 @@ 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)

diff --git a/mlir/lib/Dialect/CMakeLists.txt b/mlir/lib/Dialect/CMakeLists.txt
@@ -3,7 +3,6 @@ add_subdirectory(AVX512)
 add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
-add_subdirectory(LoopOps)
 add_subdirectory(MIOpen)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)

diff --git a/mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt b/mlir/lib/Dialect/MIOpen/Transforms/CMakeLists.txt
@@ -8,19 +8,18 @@ add_mlir_dialect_library(MLIRMIOpenTransforms
 
   DEPENDS
   MLIRMIOpenPassIncGen
+  MLIRSCFToStandard
   MLIRStandardOps
-  MLIRLoopToStandard
 )
 
 target_link_libraries(MLIRMIOpenTransforms
   PUBLIC
   MLIRAffineToStandard
   MLIRIR
-  MLIRLoopOps
-  MLIRLoopToStandard
   MLIRLLVMIR
   MLIRMIOpenOps
   MLIRPass
+  MLIRSCFToStandard
   MLIRStandardOps
   MLIRSupport
   MLIRTransformUtils

diff --git a/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp b/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
@@ -24,7 +24,7 @@
 #include "mlir/Dialect/MIOpen/LowerMIOpenOps.h"
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
-#include "mlir/Conversion/LoopToStandard/ConvertLoopToStandard.h"
+#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
 #include "mlir/Dialect/MIOpen/MIOpenOps.h"
 #include "mlir/Dialect/MIOpen/Passes.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"

diff --git a/mlir/lib/Target/CppOutput/CMakeLists.txt b/mlir/lib/Target/CppOutput/CMakeLists.txt
@@ -5,8 +5,14 @@ add_mlir_translation_library(MLIRTargetMIOpenCppTranslation
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target
-  )
-target_link_libraries(MLIRTargetMIOpenCppTranslation
+
+  DEPENDS
+  intrinsics_gen
+
+  LINK_COMPONENTS
+  Core
+
+  LINK_LIBS PUBLIC
   LLVMSupport
   MLIRIR
   MLIRMIOpenOps

diff --git a/mlir/lib/Transforms/CMakeLists.txt b/mlir/lib/Transforms/CMakeLists.txt
@@ -31,12 +31,8 @@ add_mlir_library(MLIRTransforms
   MLIRAffineOps
   MLIRAnalysis
   MLIRLoopLikeInterface
-<<<<<<< HEAD
-  MLIRSCF
-=======
-  MLIRLoopOps
   MLIRMIOpenOps
->>>>>>> 60114561518... Rebase to LLVM/MLIR tip.
+  MLIRSCF
   MLIRPass
   MLIRTransformUtils
   MLIRVector