Initial commit to add ops into MIOpen dialect.

- conv2d - transform - gridwise_gemm add dummy parse / print / verify logic. add dummy test.
llvm · Jun 6, 2020 · 4600d68 · 4600d68
1 parent 125143f
commit 4600d68
Show file tree

Hide file tree

Showing 5 changed files with 127 additions and 9 deletions.
diff --git a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h
@@ -28,8 +28,8 @@ class MIOpenOpsDialect : public Dialect {
   static StringRef getDialectNamespace() { return "miopen"; }
 };
 
-//#define GET_OP_CLASSES
-//#include "mlir/Dialect/MIOpenOps/MIOpenOps.h.inc"
+#define GET_OP_CLASSES
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.h.inc"
 
 } // end namespace miopen
 } // end namespace mlir

diff --git a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td
@@ -0,0 +1,44 @@
+//===- MIOpenOps.td - MIOpen operation definitions ---------*- tablegen -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines MLIR MIOpen operations.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MIOPEN_OPS
+#define MIOPEN_OPS
+
+include "mlir/IR/OpBase.td"
+//include "mlir/Transforms/LoopLikeInterface.td"
+
+def MIOpen_Dialect : Dialect {
+  let name = "miopen";
+  let cppNamespace = "";
+}
+
+// Base class for MIOpen dialect ops.
+class MIOpen_Op<string mnemonic, list<OpTrait> traits = []> :
+    Op<MIOpen_Dialect, mnemonic, traits> {
+  // For every standard op, there needs to be a:
+  //   * void print(OpAsmPrinter &p, ${C++ class of Op} op)
+  //   * LogicalResult verify(${C++ class of Op} op)
+  //   * ParseResult parse${C++ class of Op}(OpAsmParser &parser,
+  //                                         OperationState &result)
+  // functions.
+  let printer = [{ return ::print(p, *this); }];
+  let verifier = [{ return ::verify(*this); }];
+  let parser = [{ return ::parse$cppClass(parser, result); }];
+}
+
+def MIOpen_Conv2DOp : MIOpen_Op<"conv2d">;
+
+def MIOpen_TransformOp : MIOpen_Op<"transform">;
+
+def MIOpen_GridwiseGemmOp : MIOpen_Op<"gridwise_gemm">;
+
+#endif // MIOPEN_OPS
diff --git a/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt b/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
@@ -5,5 +5,5 @@ add_llvm_library(MLIRMIOpenOps
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/MIOpenOps
   )
-add_dependencies(MLIRMIOpenOps MLIRStandardOps LLVMSupport)
+add_dependencies(MLIRMIOpenOps MLIRMIOpenOpsIncGen MLIRStandardOps LLVMSupport)
 target_link_libraries(MLIRMIOpenOps LLVMSupport)
diff --git a/mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp b/mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp
@@ -36,17 +36,65 @@ namespace {
 
 MIOpenOpsDialect::MIOpenOpsDialect(MLIRContext *context)
     : Dialect(getDialectNamespace(), context) {
-//  addOperations<
-//#define GET_OP_LIST
-//#include "mlir/Dialect/MIOpenOps/MIOpenOps.cpp.inc"
-//      >();
+  addOperations<
+#define GET_OP_LIST
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.cpp.inc"
+      >();
 
   //addInterfaces<LoopSideEffectsInterface>();
 }
 
+//===----------------------------------------------------------------------===//
+// Conv2DOp
+//===----------------------------------------------------------------------===//
+
+static ParseResult parseConv2DOp(OpAsmParser &parser, OperationState &result) {
+  return success();
+}
+
+static void print(OpAsmPrinter &p, Conv2DOp op) {
+  p << Conv2DOp::getOperationName();
+}
+
+static LogicalResult verify(Conv2DOp op) {
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
+// TransformOp
+//===----------------------------------------------------------------------===//
+
+static ParseResult parseTransformOp(OpAsmParser &parser, OperationState &result) {
+  return success();
+}
+
+static void print(OpAsmPrinter &p, TransformOp op) {
+  p << TransformOp::getOperationName();
+}
+
+static LogicalResult verify(TransformOp op) {
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
+// GridwiseGemmOp
+//===----------------------------------------------------------------------===//
+
+static ParseResult parseGridwiseGemmOp(OpAsmParser &parser, OperationState &result) {
+  return success();
+}
+
+static void print(OpAsmPrinter &p, GridwiseGemmOp op) {
+  p << GridwiseGemmOp::getOperationName();
+}
+
+static LogicalResult verify(GridwiseGemmOp op) {
+  return success();
+}
+
 //===----------------------------------------------------------------------===//
 // TableGen'd op method definitions
 //===----------------------------------------------------------------------===//
 
-//#define GET_OP_CLASSES
-//#include "mlir/Dialect/MIOpenOps/MIOpenOps.cpp.inc"
+#define GET_OP_CLASSES
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.cpp.inc"
diff --git a/mlir/test/Dialect/MIOpen/ops.mlir b/mlir/test/Dialect/MIOpen/ops.mlir
@@ -0,0 +1,26 @@
+// RUN: mlir-opt %s | FileCheck %s
+// RUN: mlir-opt %s | mlir-opt | FileCheck %s
+// Run: mlir-opt -mlir-print-op-generic %s | mlir-opt | 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
+  return
+}
+// CHECK-LABEL: func @miopen_conv2d
+//  CHECK-NEXT: miopen.conv2d
+
+func @miopen_transform(%memref : memref<?x?x?x?xf32>) {
+  miopen.transform
+  return
+}
+
+// CHECK-LABEL: func @miopen_transform
+//  CHECK-NEXT: miopen.transform
+
+func @miopen_gridwise_gemm(%A : memref<?x?xf32>, %B : memref<?x?xf32>, %C : memref<?x?xf32>) {
+  miopen.gridwise_gemm
+  return
+}
+
+// CHECK-LABEL: func @miopen_gridwise_gemm
+//  CHECK-NEXT: miopen.gridwise_gemm