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), _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
