Remove unused miopen.gridwise_gemm_ex operation and its tests.

llvm · Sep 9, 2020 · a8851d0 · a8851d0
1 parent 017ab45
commit a8851d0
Show file tree

Hide file tree

Showing 4 changed files with 1 addition and 61 deletions.
diff --git a/mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td b/mlir/include/mlir/Dialect/MIOpen/MIOpenOps.td
@@ -92,23 +92,6 @@ def MIOpen_GridwiseGemmOp :
   }];
 }
 
-// gridwise_gemm_ex would eventually be replaced with gridwise_gemm
-// keep it as a separate op for now.
-// the difference with gridwise_gemm is gridwise_gemm_ex carries tuning
-// parameters. On the other hand, gridwise_gemm only carries convolution
-// configuration.
-def MIOpen_GridwiseGemmExOp:
-    MIOpen_Op<"gridwise_gemm_ex">,
-    Arguments<(ins MemRefRankOf<[F32], [2]>,
-                   MemRefRankOf<[F32], [2]>,
-                   MemRefRankOf<[F32], [2]>)> {
-  let summary = "Gridwise GEMM extended";
-  let description = [{
-    The `miopen.gridwise_gemm_ex` op computes gridwise GEMM, with tuning
-    parameters supplied in the attributes.
-  }];
-}
-
 // Memory allocation on GPU memory hierachy.
 def MIOpen_GpuAllocOp:
     MIOpen_Op<"alloc">,

diff --git a/mlir/lib/Dialect/MIOpen/MIOpenOps.cpp b/mlir/lib/Dialect/MIOpen/MIOpenOps.cpp
@@ -164,30 +164,6 @@ static LogicalResult verify(GridwiseGemmOp op) {
   return success();
 }
 
-//===----------------------------------------------------------------------===//
-// GridwiseGemmExOp
-//===----------------------------------------------------------------------===//
-
-static ParseResult parseGridwiseGemmExOp(OpAsmParser &parser, OperationState &result) {
-  SmallVector<OpAsmParser::OperandType, 3> ops;
-  SmallVector<Type, 3> types;
-  return failure(
-      parser.parseOperandList(ops, OpAsmParser::Delimiter::Paren) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonTypeList(types) ||
-      parser.resolveOperands(ops, types, parser.getNameLoc(), result.operands));
-}
-
-static void print(OpAsmPrinter &p, GridwiseGemmExOp op) {
-  p << op.getOperationName() << "(" << op.getOperands() << ")";
-  p.printOptionalAttrDict(op.getAttrs());
-  p << " : " << op.getOperandTypes();
-}
-
-static LogicalResult verify(GridwiseGemmExOp op) {
-  return success();
-}
-
 //===----------------------------------------------------------------------===//
 // GpuAllocOp
 //===----------------------------------------------------------------------===//

diff --git a/mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir b/mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_level.mlir
@@ -7,8 +7,7 @@
 // RUN: mlir-opt -miopen-lowering %s | FileCheck %s
 
 func @miopen_gridwise_gemm(%matrix_a : memref<?x?xf32>, %matrix_b : memref<?x?xf32>, %matrix_c : memref<?x?xf32>) {
-  // TBD: replace with gridwise_gemm once the lowering logic is complete.
-  miopen.gridwise_gemm_ex(%matrix_a, %matrix_b, %matrix_c) {
+  miopen.gridwise_gemm(%matrix_a, %matrix_b, %matrix_c) {
     block_size = 256,
 
     m_per_block = 128,

diff --git a/mlir/test/Dialect/MIOpen/ops_2.mlir b/mlir/test/Dialect/MIOpen/ops_2.mlir
@@ -2,24 +2,6 @@
 // RUN: mlir-opt %s | mlir-opt | FileCheck %s
 // Run: mlir-opt -mlir-print-op-generic %s | mlir-opt | FileCheck %s
 
-func @miopen_gridwise_gemm_ex(%A : memref<?x?xf32>, %B : memref<?x?xf32>, %C : memref<?x?xf32>) {
-  miopen.gridwise_gemm_ex(%A, %B, %C) {
-    filter_layout = ["k", "c", "y", "x"],
-    filter_dimension = [1, 2, 3, 4],
-    input_layout = ["n", "c", "hi", "wi"],
-    input_dimension = [5, 6, 7, 8],
-    output_layout = ["n", "k", "ho", "wo"],
-    output_dimension = [9, 10, 11, 12],
-    strides = [1, 1],
-    dilations = [1, 1],
-    padding = [0, 0]
-  } : memref<?x?xf32>, memref<?x?xf32>, memref<?x?xf32>
-  return
-}
-
-// CHECK-LABEL: func @miopen_gridwise_gemm_ex
-//  CHECK-NEXT: miopen.gridwise_gemm_ex
-
 func @miopen_alloc() {
   // allocation on global.
   %buffer_global = miopen.alloc() : memref<1024xi8>