Author: Wen-Heng (Jack) Chung Date: 2020-06-05T22:18:19-05:00 New Revision: 4600d68a403647ea41eca5c62616d7abd77d9281
URL: https://github.com/llvm/llvm-project/commit/4600d68a403647ea41eca5c62616d7abd77d9281 DIFF: https://github.com/llvm/llvm-project/commit/4600d68a403647ea41eca5c62616d7abd77d9281.diff LOG: Initial commit to add ops into MIOpen dialect. - conv2d - transform - gridwise_gemm add dummy parse / print / verify logic. add dummy test. Added: mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.td mlir/test/Dialect/MIOpen/ops.mlir Modified: mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h mlir/lib/Dialect/MIOpenOps/CMakeLists.txt mlir/lib/Dialect/MIOpenOps/MIOpenOps.cpp Removed: ################################################################################ diff --git a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenOps.h index e0c36c23a854..47341c1637b1 100644 --- 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 new file mode 100644 index 000000000000..1304f16f3b30 --- /dev/null +++ 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 index 474196162792..ba32051bbb39 100644 --- 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 index ced4d25e866f..9408e17b2831 100644 --- 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 new file mode 100644 index 000000000000..a37e54110186 --- /dev/null +++ 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 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
