[llvm-branch-commits] [mlir] 0ac976c - Change emitted names to fixed "mlir".

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


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

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

LOG: Change emitted names to fixed "mlir".

Added: 
    

Modified: 
    mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
    mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
    mlir/test/Dialect/MIOpen/translate.mlir

Removed: 
    


################################################################################
diff  --git a/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp b/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
index cda706c4112c..b0b104bd137d 100644
--- a/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
+++ b/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
@@ -186,13 +186,20 @@ void EmitCppPreamble(llvm::raw_ostream &output, llvm::StringRef layoutStr) {
 // Between Preamble Part 1 and Part 2:
 // #include "gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp"
   output << R"(#include "gridwise_convolution_implicit_gemm_v4r4_)";
-  output << layoutStr << R"(.hpp")";
+
+  // Change to fixed "mlir".
+  //output << layoutStr << R"(.hpp")";
+  output << "mlir" << R"(.hpp")";
+
   output << kCppPreamblePart2;
 // Between Preamble Part 2 and Par 3:
 //    __launch_bounds__(CK_PARAM_TUNABLE_BLOCK_SIZE, 2) void gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw(
-   output << R"(
+  output << R"(
     __launch_bounds__(CK_PARAM_TUNABLE_BLOCK_SIZE, 2) void gridwise_convolution_implicit_gemm_v4r4_)";
-   output << layoutStr;
+  // Change to fixed "mlir".
+  //output << layoutStr;
+  output << "mlir";
+
   output << kCppPreamblePart3;
 }
 
@@ -205,7 +212,11 @@ void EmitCppEpilogue(llvm::raw_ostream &output, llvm::StringRef layoutStr, llvm:
 //    constexpr auto gridwise_conv = GridwiseConvolutionImplicitGemm_v4r4_nchw_kcyx_nkhw
   output << R"(
     constexpr auto gridwise_conv = GridwiseConvolutionImplicitGemm_v4r4_)";
-  output << layoutStr;
+
+  // Change to fixed "mlir".
+  //output << layoutStr;
+  output << "mlir";
+
   output << kCppEpiloguePart1;
 // Between Part1 and Part2:
 //        decltype(in_nchw_desc),
@@ -339,7 +350,11 @@ void EmitHeaderPreamble(llvm::raw_ostream &output, llvm::StringRef layoutStr, ll
   output << kHeaderPreamblePart1;
   output << R"(
 struct GridwiseConvolutionImplicitGemm_v4r4_)";
-  output << layoutStr;
+
+  // Change to fixed "mlir".
+  //output << layoutStr;
+  output << "mlir";
+
   output << kHeaderPreamblePart2;
   output << kHeaderPreamblePart3;
   output << '\n';

diff  --git a/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp b/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
index 27311cb8cfb9..0378a4c113bf 100644
--- a/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
+++ b/mlir/lib/Dialect/MIOpenOps/LowerMIOpenOps.cpp
@@ -208,8 +208,8 @@ struct Conv2DOpRewritePattern : public OpRewritePattern<miopen::Conv2DOp> {
                                                         ArrayAttr::get({
                                                             StringAttr::get("ni", op.getContext()),
                                                             StringAttr::get("ci", op.getContext()),
-                                                            StringAttr::get("hi", op.getContext()),
-                                                            StringAttr::get("wi", op.getContext())
+                                                            StringAttr::get("hipad", op.getContext()),
+                                                            StringAttr::get("wipad", op.getContext())
                                                         }, op.getContext()));
     paddedInputAttrs.push_back(paddedInputOutputLayoutAttr);
     auto paddedInput = rewriter.create<miopen::TransformOp>(op.getLoc(), inputType, op.input(), paddedInputAttrs);

diff  --git a/mlir/test/Dialect/MIOpen/translate.mlir b/mlir/test/Dialect/MIOpen/translate.mlir
index 1139100262e2..a3fedba84cda 100644
--- a/mlir/test/Dialect/MIOpen/translate.mlir
+++ b/mlir/test/Dialect/MIOpen/translate.mlir
@@ -1,8 +1,8 @@
 // RUN: mlir-translate -mlir-to-miopen-cpp %s | FileCheck -check-prefix=MIOPEN-CPP %s
 // RUN: mlir-translate -mlir-to-miopen-hpp %s | FileCheck -check-prefix=MIOPEN-HPP %s
 
-// MIOPEN-CPP:  __launch_bounds__(CK_PARAM_TUNABLE_BLOCK_SIZE, 2) void gridwise_convolution_implicit_gemm_v4r4_kcyx_nicihiwi_nokohowo
-// MIOPEN-HPP: struct GridwiseConvolutionImplicitGemm_v4r4_kcyx_nicihiwi_nokohowo
+// MIOPEN-CPP:  __launch_bounds__(CK_PARAM_TUNABLE_BLOCK_SIZE, 2) void gridwise_convolution_implicit_gemm_v4r4_mlir
+// MIOPEN-HPP: struct GridwiseConvolutionImplicitGemm_v4r4_mlir
 func @miopen_transformed_conv2d(%filter : memref<?x?x?x?xf32>, %input : memref<?x?x?x?xf32>, %output : memref<?x?x?x?xf32>) {
   // filter tensor
   %filter_gemmK_gemmM = miopen.transform(%filter) {
@@ -153,7 +153,7 @@ func @miopen_transformed_conv2d(%filter : memref<?x?x?x?xf32>, %input : memref<?
 // MIOPEN-CPP:    constexpr auto weight_k_c_y_x_desc = make_native_tensor_descriptor(Sequence<k, c, y, x>{}, Sequence<stride_k, stride_c, stride_y, stride_x>{});
 // MIOPEN-CPP:     constexpr auto input_ni_ci_hi_wi_desc = make_native_tensor_descriptor(Sequence<ni, ci, hi, wi>{}, Sequence<stride_ni, stride_ci, stride_hi, stride_wi>{});
 // MIOPEN-CPP:     constexpr auto output_no_ko_ho_wo_desc = make_native_tensor_descriptor(Sequence<no, ko, ho, wo>{}, Sequence<stride_no, stride_ko, stride_ho, stride_wo>{});
-// MIOPEN-CPP:         constexpr auto gridwise_conv = GridwiseConvolutionImplicitGemm_v4r4_kcyx_nicihiwi_nokohowo
+// MIOPEN-CPP:         constexpr auto gridwise_conv = GridwiseConvolutionImplicitGemm_v4r4_mlir
 // MIOPEN-CPP:        decltype(weight_k_c_y_x_desc),
 // MIOPEN-CPP:        decltype(input_ni_ci_hi_wi_desc),
 // MIOPEN-CPP:        decltype(output_no_ko_ho_wo_desc),


        


More information about the llvm-branch-commits mailing list