Merge pull request #29 from jerryyin/miopen-dialect-libMLIRMIOpen

Adding libMLIRMIOpen.a for cpp generation API
llvm · Aug 7, 2020 · ccad4d9 · ccad4d9
2 parents 3388282 + e100ab7
commit ccad4d9
Show file tree

Hide file tree

Showing 16 changed files with 555 additions and 228 deletions.
diff --git a/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h b/mlir/include/mlir/Dialect/MIOpen/LowerMIOpenOps.h
@@ -10,6 +10,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#ifndef MLIR_DIALECT_MIOPEN_LOWERMIOPENOPS_H
+#define MLIR_DIALECT_MIOPEN_LOWERMIOPENOPS_H
+
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
@@ -915,42 +918,6 @@ struct Conv2DRewritePattern : public OpRewritePattern<T> {
   }
 };
 
-// High level convolution operation always have
-// [filter, input, output]
-// as the convolution argument. The only difference between different
-// hight level convolution operations is the argument sequence. For
-// simplicity, we always arrange the first two arguments to be input
-// and the last argument to be output
-template <>
-const ArgumentFields Conv2DRewritePattern<miopen::Conv2DOp>::fields = {
-    {0, 1, 2},
-    {"KM", "KN", "MN"},
-};
-template <>
-const miopen::ConvOpType Conv2DRewritePattern<miopen::Conv2DOp>::convOpType =
-    miopen::ConvOpType::Conv2DOpType;
-
-template <>
-const ArgumentFields Conv2DRewritePattern<miopen::Conv2DBwdDataOp>::fields = {
-    {0, 2, 1},
-    {"KM", "MN", "KN"},
-};
-
-template <>
-const miopen::ConvOpType Conv2DRewritePattern<miopen::Conv2DBwdDataOp>::convOpType =
-miopen::ConvOpType::Conv2DBwdDataOpType;
-
-template <>
-const ArgumentFields Conv2DRewritePattern<miopen::Conv2DBwdWeightOp>::fields = {
-	{2, 1, 0},
-	{"MN", "KN", "KM"},
-};
-
-template <>
-const miopen::ConvOpType Conv2DRewritePattern<miopen::Conv2DBwdWeightOp>::convOpType =
-miopen::ConvOpType::Conv2DBwdWeightOpType;
-
-
 // Explicitly instantiate the template to operation type
 template struct Conv2DRewritePattern<miopen::Conv2DOp>;
 template struct Conv2DRewritePattern<miopen::Conv2DBwdDataOp>;
@@ -3004,3 +2971,4 @@ struct TransformRewritePattern : public OpRewritePattern<miopen::TransformOp> {
     return success();
   }
 };
+#endif // MLIR_DIALECT_MIOPEN_LOWERMIOPENOPS_H
diff --git a/mlir/include/mlir/Target/MIOpenCPP.h b/mlir/include/mlir/Target/MIOpenCPP.h
@@ -27,17 +27,17 @@ class ModuleOp;
 /// Convert the given MLIR module into MIOpen C++ . In case of error, report it
 /// to the error handler registered with the MLIR context, if any (obtained from
 /// the MLIR module), and return `nullptr`.
-std::unique_ptr<llvm::StringRef> translateModuleToMIOpenCpp(ModuleOp m);
+void translateModuleToMIOpenCpp(ModuleOp m, std::string &source);
 
 /// Convert the given MLIR module into MIOpen C++ Header. In case of error, report it
 /// to the error handler registered with the MLIR context, if any (obtained from
 /// the MLIR module), and return `nullptr`.
-std::unique_ptr<llvm::StringRef> translateModuleToMIOpenHeader(ModuleOp m);
+void translateModuleToMIOpenHeader(ModuleOp m, std::string &header);
 
 /// Convert the given MLIR module into MIOpen C++ compilation flags. In case of error, report it
 /// to the error handler registered with the MLIR context, if any (obtained from
 /// the MLIR module), and return `nullptr`.
-std::unique_ptr<llvm::StringRef> translateModuleToMIOpenCFlags(ModuleOp m);
+void translateModuleToMIOpenCFlags(ModuleOp m, std::string &cflags);
 
 /// Convert the given MLIR module into MIOpen C++ . In case of error, report it
 /// to the error handler registered with the MLIR context, if any (obtained from

diff --git a/mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt b/mlir/lib/Conversion/MIOpenToGPU/CMakeLists.txt
@@ -10,6 +10,7 @@ add_mlir_conversion_library(MLIRMIOpenToGPU
 target_link_libraries(MLIRMIOpenToGPU
   PUBLIC
   MLIRAffineToStandard
+  MLIRMIOpenTransforms
   MLIRGPU
   MLIRLLVMIR
   MLIRIR

diff --git a/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp b/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
@@ -71,6 +71,43 @@ struct LowerMIOpenOpsStep5Pass
 };
 } // end anonymous namespace
 
+// High level convolution operation always have
+// [filter, input, output]
+// as the convolution argument. The only difference between different
+// hight level convolution operations is the argument sequence. For
+// simplicity, we always arrange the first two arguments to be input
+// and the last argument to be output
+template <>
+const ArgumentFields Conv2DRewritePattern<miopen::Conv2DOp>::fields = {
+    {0, 1, 2},
+    {"KM", "KN", "MN"},
+};
+template <>
+const miopen::ConvOpType Conv2DRewritePattern<miopen::Conv2DOp>::convOpType =
+    miopen::ConvOpType::Conv2DOpType;
+
+template <>
+const ArgumentFields Conv2DRewritePattern<miopen::Conv2DBwdDataOp>::fields = {
+    {0, 2, 1},
+    {"KM", "MN", "KN"},
+};
+
+template <>
+const miopen::ConvOpType
+    Conv2DRewritePattern<miopen::Conv2DBwdDataOp>::convOpType =
+        miopen::ConvOpType::Conv2DBwdDataOpType;
+
+template <>
+const ArgumentFields Conv2DRewritePattern<miopen::Conv2DBwdWeightOp>::fields = {
+    {2, 1, 0},
+    {"MN", "KN", "KM"},
+};
+
+template <>
+const miopen::ConvOpType
+    Conv2DRewritePattern<miopen::Conv2DBwdWeightOp>::convOpType =
+        miopen::ConvOpType::Conv2DBwdWeightOpType;
+
 void LowerMIOpenOpsStep1Pass::runOnOperation() {
   OwningRewritePatternList patterns;
   patterns.insert<Conv2DRewritePattern<miopen::Conv2DOp>>(&getContext());

diff --git a/mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp b/mlir/lib/Target/CppOutput/ConvertToMIOpenCPP.cpp
@@ -40,36 +40,33 @@ cl::opt<bool> IsPopulateTunableParameters("populate-tunable-parameters-to-yaml-f
 namespace mlir {
 void registerToMIOpenCPPTranslation() {
   // non-XDLOPS kernel generation.
-  TranslateFromMLIRRegistration
-      toCpp("mlir-to-miopen-cpp", [](ModuleOp module, llvm::raw_ostream &output) {
-        auto sourceCode = mlir::translateModuleToMIOpenCpp(module);
-        if (!sourceCode)
-          return failure();
-
-        output << *sourceCode;
+  TranslateFromMLIRRegistration toCpp(
+      "mlir-to-miopen-cpp", [](ModuleOp module, llvm::raw_ostream &output) {
+        std::string source;
+        mlir::translateModuleToMIOpenCpp(module, source);
+
+        output << source;
         return success();
       });
-
-  TranslateFromMLIRRegistration
-      toHeader("mlir-to-miopen-hpp", [](ModuleOp module, llvm::raw_ostream &output) {
-        auto sourceCode = mlir::translateModuleToMIOpenHeader(module);
-        if (!sourceCode)
-          return failure();
-
-        output << *sourceCode;
+
+  TranslateFromMLIRRegistration toHeader(
+      "mlir-to-miopen-hpp", [](ModuleOp module, llvm::raw_ostream &output) {
+        std::string header;
+        mlir::translateModuleToMIOpenHeader(module, header);
+
+        output << header;
         return success();
       });
-
-  TranslateFromMLIRRegistration
-      toCFlags("mlir-to-miopen-cflags", [](ModuleOp module, llvm::raw_ostream &output) {
-        auto sourceCode = mlir::translateModuleToMIOpenCFlags(module);
-        if (!sourceCode)
-          return failure();
-
-        output << *sourceCode;
+
+  TranslateFromMLIRRegistration toCFlags(
+      "mlir-to-miopen-cflags", [](ModuleOp module, llvm::raw_ostream &output) {
+        std::string cflags;
+        mlir::translateModuleToMIOpenCFlags(module, cflags);
+
+        output << cflags;
         return success();
       });
-  
+
   // XDLOPS kernel generation.
   TranslateFromMLIRRegistration
       toCppXDLOPS("mlir-to-miopen-cpp-xdlops", [](ModuleOp module, llvm::raw_ostream &output) {

diff --git a/mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp b/mlir/lib/Target/CppOutput/gridwise_convolution_implicit_gemm_v4r4.cpp
@@ -28,9 +28,6 @@
 using namespace mlir;
 
 namespace {
-// result string to keep C++ source / header / flags emission.
-std::string resultStr;
-
 static constexpr StringLiteral kVarArgName[3] = {"p_wei_global", "p_in_global",
                                                  "p_out_global"};
 
@@ -684,8 +681,8 @@ static void ObtainModuleInfo(ModuleOp &m,
 
 } // namespace
 
-std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenHeader(ModuleOp m) {
-  llvm::raw_string_ostream output(resultStr);
+void mlir::translateModuleToMIOpenHeader(ModuleOp m, std::string &header) {
+  llvm::raw_string_ostream output(header);
 
   // Enumerate FuncOp instances inside the ModuleOp.
   for (auto f : m.getOps<FuncOp>()) {
@@ -865,11 +862,10 @@ std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenHeader(ModuleOp m)
   }
 
   output.flush();
-  return std::make_unique<llvm::StringRef>(resultStr);
 }
 
-std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenCpp(ModuleOp m) {
-  llvm::raw_string_ostream output(resultStr);
+void mlir::translateModuleToMIOpenCpp(ModuleOp m, std::string &source) {
+  llvm::raw_string_ostream output(source);
 
   // Enumerate FuncOp instances inside the ModuleOp.
   for (auto f : m.getOps<FuncOp>()) {
@@ -912,11 +908,10 @@ std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenCpp(ModuleOp m) {
   }
 
   output.flush();
-  return std::make_unique<llvm::StringRef>(resultStr);
 }
 
-std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenCFlags(ModuleOp m) {
-  llvm::raw_string_ostream output(resultStr);
+void mlir::translateModuleToMIOpenCFlags(ModuleOp m, std::string &cflags) {
+  llvm::raw_string_ostream output(cflags);
 
   for (auto f : m.getOps<FuncOp>()) {
     output << f.getName() << "\n";
@@ -1049,5 +1044,4 @@ std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenCFlags(ModuleOp m)
   }
 
   output.flush();
-  return std::make_unique<llvm::StringRef>(resultStr);
 }
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
@@ -81,6 +81,7 @@ endif()
 if(MLIR_MIOPEN_DRIVER_ENABLED)
   list(APPEND MLIR_TEST_DEPENDS
     mlir-miopen-driver
+    mlir-miopen-lib-test
     mlir-translate
     opt
     llc

diff --git a/mlir/test/mlir-miopen-lib/populate_bwd.mlir b/mlir/test/mlir-miopen-lib/populate_bwd.mlir
@@ -0,0 +1,7 @@
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option cflags | FileCheck %s --check-prefix=CFLAGS
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option source | FileCheck %s --check-prefix=SOURCE
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option header | FileCheck %s --check-prefix=HEADER
+
+// CFLAGS: miopen_conv2d_bwd_data_kcyx_nchw_nkhw
+// SOURCE: void gridwise_convolution_backward_data_implicit_gemm_v1r1_mlir
+// HEADER: struct GridwiseConvolutionBackwardDataImplicitGemm_v1r1_mlir 
diff --git a/mlir/test/mlir-miopen-lib/populate_bww.mlir b/mlir/test/mlir-miopen-lib/populate_bww.mlir
@@ -0,0 +1,7 @@
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option cflags | FileCheck %s --check-prefix=CFLAGS
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option source | FileCheck %s --check-prefix=SOURCE
+// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option header | FileCheck %s --check-prefix=HEADER
+
+// CFLAGS: miopen_conv2d_bwd_weight_kcyx_nchw_nkhw
+// SOURCE: void gridwise_convolution_backward_weight_implicit_gemm_v4r4_mlir
+// HEADER: struct GridwiseConvolutionBackwardWeightImplicitGemm_v4r4_mlir 
diff --git a/mlir/test/mlir-miopen-lib/populate_fw.mlir b/mlir/test/mlir-miopen-lib/populate_fw.mlir
@@ -0,0 +1,7 @@
+// RUN: mlir-miopen-lib-test --args " --operation conv2d --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option cflags | FileCheck %s --check-prefix=CFLAGS
+// RUN: mlir-miopen-lib-test --args " --operation conv2d --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option source | FileCheck %s --check-prefix=SOURCE
+// RUN: mlir-miopen-lib-test --args " --operation conv2d --fil_layout kcyx --in_layout nchw --out_layout nkhw --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0" --option header | FileCheck %s --check-prefix=HEADER
+
+// CFLAGS: miopen_conv2d_kcyx_nchw_nkhw
+// SOURCE: void gridwise_convolution_implicit_gemm_v4r4_mlir
+// HEADER: struct GridwiseConvolutionImplicitGemm_v4r4_mlir 
diff --git a/mlir/tools/mlir-miopen-driver/CMakeLists.txt b/mlir/tools/mlir-miopen-driver/CMakeLists.txt
@@ -19,15 +19,81 @@ set(LIBS
   MLIRTransforms
   MLIRSupport
   MLIRIR
+  MLIRTargetMIOpenCppTranslation
 )
 
 add_llvm_executable(mlir-miopen-driver
+  PARTIAL_SOURCES_INTENDED
+
   mlir-miopen-driver.cpp
 
   DEPENDS
   ${LIBS}
 )
+
 llvm_update_compile_flags(mlir-miopen-driver)
 target_link_libraries(mlir-miopen-driver PRIVATE ${LIBS})
-
 mlir_check_link_libraries(mlir-miopen-driver)
+
+
+  llvm_add_library(MLIRMIOpenThin
+  PARTIAL_SOURCES_INTENDED
+
+    mlir-miopen-lib.cpp
+
+    LINK_LIBS
+    ${LIBS}
+    )
+
+add_llvm_executable(mlir-miopen-lib-test
+  PARTIAL_SOURCES_INTENDED
+
+  mlir-miopen-lib-test.cpp
+
+  DEPENDS
+  MLIRMIOpenThin
+  ${LIBS}
+)
+
+llvm_update_compile_flags(mlir-miopen-lib-test)
+target_link_libraries(mlir-miopen-lib-test PRIVATE MLIRMIOpenThin ${LIBS})
+mlir_check_link_libraries(mlir-miopen-lib-test)
+
+# Static library target, enabled only when building static libs
+if( NOT BUILD_SHARED_LIBS )
+  function(combine_archives output_archive)
+    set(mri_file ${CMAKE_CURRENT_BINARY_DIR}/${output_archive}.mri)
+    set(full_output_path ${LLVM_LIBRARY_DIR}/lib${output_archive}.a)
+    set(output_archive_dummy_file ./${output_archive}.dummy.cpp)
+    set(install_path /opt/rocm/lib/lib${output_archive}.a)
+
+    # Step one: construct mri file.
+    add_custom_command(OUTPUT ${output_archive_dummy_file}
+                       COMMAND if [ -f ${output_archive}.mri ]\; then rm ${output_archive}.mri\; fi
+                       COMMAND touch ${output_archive}.mri
+                       COMMAND echo "create ${full_output_path}" >> ${output_archive}.mri
+                       COMMAND for archive in ${LLVM_LIBRARY_DIR}/*.a\; 
+                       do echo "addlib $$archive" >> ${output_archive}.mri \; done
+                       COMMAND echo "save" >> ${output_archive}.mri
+                       COMMAND echo "end" >> ${output_archive}.mri
+                       COMMAND touch ${output_archive_dummy_file}
+                       DEPENDS MLIRMIOpenThin)
+
+    # Step two: use mri file to generate the fat library.
+    llvm_add_library(${output_archive}
+      PARTIAL_SOURCES_INTENDED
+      STATIC ${output_archive_dummy_file})
+    add_custom_command(TARGET ${output_archive}
+                       POST_BUILD
+                       COMMAND ${CMAKE_AR} -M < ${mri_file}
+                       COMMAND ${CMAKE_COMMAND} -E copy ${full_output_path} ${install_path}
+                       DEPENDS ${output_archive_dummy_file})
+  endfunction(combine_archives)
+
+  combine_archives(MLIRMIOpen)
+
+  add_custom_target(libMLIRMIOpen ALL
+    DEPENDS
+    MLIRMIOpen
+    )
+endif()