[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