[llvm-branch-commits] [mlir] e587a8a - Inital commit to add MIOpen Conv2D to Transform and GridwiseGemm transform pass.

Wen-Heng Chung via llvm-branch-commits llvm-branch-commits at lists.llvm.org
Thu Oct 22 13:20:23 PDT 2020


Author: Wen-Heng (Jack) Chung
Date: 2020-06-05T22:18:20-05:00
New Revision: e587a8a23c9812610dd1e79a3e1211e1f4d8aba5

URL: https://github.com/llvm/llvm-project/commit/e587a8a23c9812610dd1e79a3e1211e1f4d8aba5
DIFF: https://github.com/llvm/llvm-project/commit/e587a8a23c9812610dd1e79a3e1211e1f4d8aba5.diff

LOG: Inital commit to add MIOpen Conv2D to Transform and GridwiseGemm transform pass.

Added: 
    mlir/include/mlir/Dialect/MIOpenOps/Passes.h
    mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
    mlir/test/Dialect/MIOpen/lowering.mlir

Modified: 
    mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
index 8ffd66647f3f..1f531d9176ab 100644
--- a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
+++ b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
@@ -37,9 +37,9 @@ class MIOpen_Op<string mnemonic, list<OpTrait> traits = []> :
 
 def MIOpen_Conv2DOp :
     MIOpen_Op<"conv2d">,
-    Arguments<(ins MemRefRankOf<[F32], [4]>,
-                   MemRefRankOf<[F32], [4]>,
-                   MemRefRankOf<[F32], [4]>)> {
+    Arguments<(ins MemRefRankOf<[F32], [4]>:$filter,
+                   MemRefRankOf<[F32], [4]>:$input,
+                   MemRefRankOf<[F32], [4]>:$output)> {
   let summary = "2D convolution";
   let description = [{
     The `miopen.conv2d` op computes 2D convolution.

diff  --git a/mlir/include/mlir/Dialect/MIOpenOps/Passes.h b/mlir/include/mlir/Dialect/MIOpenOps/Passes.h
new file mode 100644
index 000000000000..6752b71c5598
--- /dev/null
+++ b/mlir/include/mlir/Dialect/MIOpenOps/Passes.h
@@ -0,0 +1,33 @@
+//===- Passes.h - Linalg pass entry points ----------------------*- C++ -*-===//
+//
+// Part of the MLIR 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This header file defines prototypes that expose pass constructors.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_MIOPEN_PASSES_H_
+#define MLIR_DIALECT_MIOPEN_PASSES_H_
+
+#include "mlir/Support/LLVM.h"
+#include "llvm/ADT/ArrayRef.h"
+
+namespace mlir {
+class FuncOp;
+class ModuleOp;
+template <typename T> class OpPassBase;
+
+namespace miopen {
+
+/// Create a pass to convert MIOpen conv2d operations to transform and
+/// gridwise_gemm operations.
+std::unique_ptr<OpPassBase<ModuleOp>> createLowerMIOpenOpsPass();
+
+} // namespace miopen
+} // namespace mlir
+
+#endif // MLIR_DIALECT_MIOPEN_PASSES_H_

diff  --git a/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp b/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
new file mode 100644
index 000000000000..2a00ed675122
--- /dev/null
+++ b/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
@@ -0,0 +1,82 @@
+//===- LowerMIOpenOps.cpp - MLIR MIOpen ops lowering passes ---------------===//
+//
+// Copyright 2020 The MLIR Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//   http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+// =============================================================================
+//
+// This pass converts miopen.conv2d into miopen.transform and
+// miopen.gridwise_gemm.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.h"
+#include "mlir/Dialect/MIOpenOps/Passes.h"
+#include "mlir/Dialect/StandardOps/Ops.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Module.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/IR/StandardTypes.h"
+#include "mlir/IR/Types.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Pass/PassManager.h"
+#include "mlir/Pass/PassRegistry.h"
+#include "mlir/Transforms/DialectConversion.h"
+#include "mlir/Transforms/Passes.h"
+#include "mlir/Support/LogicalResult.h"
+
+using namespace mlir;
+
+struct Conv2DOpRewritePattern : public OpRewritePattern<miopen::Conv2DOp> {
+  using OpRewritePattern<miopen::Conv2DOp>::OpRewritePattern;
+
+  PatternMatchResult
+  matchAndRewrite(miopen::Conv2DOp op, PatternRewriter &rewriter) const override {
+    rewriter.create<miopen::TransformOp>(op.getLoc(), op.filter().getType(), op.filter());
+
+    rewriter.create<miopen::TransformOp>(op.getLoc(), op.input().getType(), op.input());
+    rewriter.create<miopen::TransformOp>(op.getLoc(), op.input().getType(), op.input());
+    rewriter.create<miopen::TransformOp>(op.getLoc(), op.input().getType(), op.input());
+
+    rewriter.create<miopen::TransformOp>(op.getLoc(), op.output().getType(), op.output());
+
+    //rewriter.create<miopen::GridwiseGemmOp>(op.getLoc(), op.filter(), op.input(), op.output());
+
+    // Finally, erase the original Conv2D op.
+    op.erase();
+
+    return matchSuccess();
+  }
+};
+
+namespace {
+struct LowerMIOpenOpsPass : public ModulePass<LowerMIOpenOpsPass> {
+  void runOnModule() override;
+};
+} // end anonymous namespace
+
+void LowerMIOpenOpsPass::runOnModule() {
+  OwningRewritePatternList patterns;
+  patterns.insert<Conv2DOpRewritePattern>(&getContext());
+  applyPatternsGreedily(getModule(), patterns);
+}
+
+std::unique_ptr<OpPassBase<ModuleOp>> mlir::miopen::createLowerMIOpenOpsPass() {
+  return std::make_unique<LowerMIOpenOpsPass>();
+}
+
+static PassRegistration<LowerMIOpenOpsPass>
+    lowerMIOpenOpsPass("miopen-lowering",
+                       "Lower MIOpen conv2d into transform and gridwise_gemm.");

diff  --git a/mlir/test/Dialect/MIOpen/lowering.mlir b/mlir/test/Dialect/MIOpen/lowering.mlir
new file mode 100644
index 000000000000..e7734cef5a29
--- /dev/null
+++ b/mlir/test/Dialect/MIOpen/lowering.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-opt -miopen-lowering %s | FileCheck %s
+
+func @miopen_conv2d(%filter : memref<?x?x?x?xf32>, %input : memref<?x?x?x?xf32>, %output : memref<?x?x?x?xf32>) {
+  miopen.conv2d(%filter, %input, %output) {
+    filter_layout = ["k", "c", "y", "x"],
+    input_layout = ["n", "c", "hi", "wi"],
+    output_layout = ["n", "k", "ho", "wo"],
+    dilations = [1, 1],
+    strides = [1, 1],
+    padding = [0, 0]
+  } : memref<?x?x?x?xf32>, memref<?x?x?x?xf32>, memref<?x?x?x?xf32>
+  return
+}
+// CHECK-LABEL: func @miopen_conv2d
+//  CHECK-NOT: miopen.conv2d
+//  CHECK-NEXT: miopen.transform
+//  CHECK-NEXT: miopen.transform
+//  CHECK-NEXT: miopen.transform
+//  CHECK-NEXT: miopen.transform
+//  CHECK-NEXT: miopen.transform
+//  TBD-CHECK-NEXT: miopen.gridwise_gemm


        


More information about the llvm-branch-commits mailing list