XXX WIP tuning miopen.xdlops_gemm_v2 lowering logic.

llvm · Sep 11, 2020 · f8b0090 · f8b0090
1 parent 12e4d12
commit f8b0090
Show file tree

Hide file tree

Showing 2 changed files with 13 additions and 5 deletions.
diff --git a/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h b/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
@@ -5835,7 +5835,9 @@ struct BlockwiseGemmV2RewritePattern
       xdlopsGemmV2Op0.setAttr("m", op.getAttr("m"));
       xdlopsGemmV2Op0.setAttr("n", op.getAttr("n"));
       xdlopsGemmV2Op0.setAttr("k", op.getAttr("k"));
-      xdlopsGemmV2Op0.setAttr("m_per_wave", op.getAttr("m_per_wave"));
+      // TBD. hard-coded as 64 for now.
+      xdlopsGemmV2Op0.setAttr("m_per_wave", b.getI32IntegerAttr(64));
+      //xdlopsGemmV2Op0.setAttr("m_per_wave", op.getAttr("m_per_wave"));
       xdlopsGemmV2Op0.setAttr("n_per_wave", op.getAttr("n_per_wave"));
       xdlopsGemmV2Op0.setAttr("coord_transforms",
                               op.getAttr("coord_transforms"));
@@ -5853,7 +5855,9 @@ struct BlockwiseGemmV2RewritePattern
       xdlopsGemmV2Op1.setAttr("m", op.getAttr("m"));
       xdlopsGemmV2Op1.setAttr("n", op.getAttr("n"));
       xdlopsGemmV2Op1.setAttr("k", op.getAttr("k"));
-      xdlopsGemmV2Op1.setAttr("m_per_wave", op.getAttr("m_per_wave"));
+      // TBD. hard-coded as 64 for now.
+      xdlopsGemmV2Op1.setAttr("m_per_wave", b.getI32IntegerAttr(64));
+      //xdlopsGemmV2Op1.setAttr("m_per_wave", op.getAttr("m_per_wave"));
       xdlopsGemmV2Op1.setAttr("n_per_wave", op.getAttr("n_per_wave"));
       xdlopsGemmV2Op1.setAttr("coord_transforms",
                               op.getAttr("coord_transforms"));
@@ -5878,7 +5882,9 @@ struct BlockwiseGemmV2RewritePattern
       xdlopsGemmV2Op0.setAttr("m", op.getAttr("m"));
       xdlopsGemmV2Op0.setAttr("n", op.getAttr("n"));
       xdlopsGemmV2Op0.setAttr("k", op.getAttr("k"));
-      xdlopsGemmV2Op0.setAttr("m_per_wave", op.getAttr("m_per_wave"));
+      // TBD. hard-coded as 64 for now.
+      xdlopsGemmV2Op0.setAttr("m_per_wave", b.getI32IntegerAttr(64));
+      //xdlopsGemmV2Op0.setAttr("m_per_wave", op.getAttr("m_per_wave"));
       xdlopsGemmV2Op0.setAttr("n_per_wave", op.getAttr("n_per_wave"));
       xdlopsGemmV2Op0.setAttr("coord_transforms",
                               op.getAttr("coord_transforms"));
@@ -5896,7 +5902,9 @@ struct BlockwiseGemmV2RewritePattern
       xdlopsGemmV2Op1.setAttr("m", op.getAttr("m"));
       xdlopsGemmV2Op1.setAttr("n", op.getAttr("n"));
       xdlopsGemmV2Op1.setAttr("k", op.getAttr("k"));
-      xdlopsGemmV2Op1.setAttr("m_per_wave", op.getAttr("m_per_wave"));
+      // TBD. hard-coded as 64 for now.
+      xdlopsGemmV2Op1.setAttr("m_per_wave", b.getI32IntegerAttr(64));
+      //xdlopsGemmV2Op1.setAttr("m_per_wave", op.getAttr("m_per_wave"));
       xdlopsGemmV2Op1.setAttr("n_per_wave", op.getAttr("n_per_wave"));
       xdlopsGemmV2Op1.setAttr("coord_transforms",
                               op.getAttr("coord_transforms"));

diff --git a/mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir b/mlir/test/Dialect/MIOpen/lowering_xdlops_gemm_v2.mlir
@@ -1,5 +1,5 @@
 // XFAIL: *
-// RUN: mlir-opt -miopen-lowering-step4 %s | FileCheck %s
+// RUN: mlir-opt -miopen-lowering-step5 %s | FileCheck %s
 
 func @miopen_xdlops_gemm_v2_two_results(%matrixA : memref<12288xf32, 3>, %matrixB : memref<12288xf32, 3>) -> (vector<32xf32>, vector<32xf32>) {
   %c0 = constant 0 : index