[llvm-branch-commits] [mlir] 9bc39d9 - Initial commit to introduce MLIR MIOpen dialect -> MIOpen C++ translation.
Wen-Heng Chung via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Oct 22 13:20:12 PDT 2020
Author: Wen-Heng (Jack) Chung
Date: 2020-06-05T22:18:19-05:00
New Revision: 9bc39d9e6f9e7046f99dc105c093c78fc8031011
URL: https://github.com/llvm/llvm-project/commit/9bc39d9e6f9e7046f99dc105c093c78fc8031011
DIFF: https://github.com/llvm/llvm-project/commit/9bc39d9e6f9e7046f99dc105c093c78fc8031011.diff
LOG: Initial commit to introduce MLIR MIOpen dialect -> MIOpen C++ translation.
Added:
mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir
Modified:
mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
new file mode 100644
index 000000000000..d3e9b8ee09a2
--- /dev/null
+++ b/mlir/include/mlir/Dialect/MIOpenOps/MIOpenCPP.h
@@ -0,0 +1,40 @@
+//===- MIOpenCPP.h - MLIR to C++ for MIOpen conversion ----------*- C++ -*-===//
+//
+// Part of the MLIR Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file declares the entry point for the MLIR to MIOpen C++ conversion.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_MIOPEN_CPP_H
+#define MLIR_TARGET_MIOPEN_CPP_H
+
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.h"
+#include "mlir/IR/Block.h"
+#include "mlir/IR/Module.h"
+#include "mlir/IR/Value.h"
+
+#include <memory>
+
+namespace llvm {
+class StringRef;
+} // namespace llvm
+
+namespace mlir {
+
+class OwningModuleRef;
+class MLIRContext;
+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);
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_MIOPEN_CPP_H
diff --git a/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt b/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
index ba32051bbb39..b7b0c69febea 100644
--- a/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
+++ b/mlir/lib/Dialect/MIOpenOps/CMakeLists.txt
@@ -7,3 +7,5 @@ add_llvm_library(MLIRMIOpenOps
)
add_dependencies(MLIRMIOpenOps MLIRMIOpenOpsIncGen MLIRStandardOps LLVMSupport)
target_link_libraries(MLIRMIOpenOps LLVMSupport)
+
+add_subdirectory(CppOutput)
diff --git a/mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt b/mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
new file mode 100644
index 000000000000..855985b4b945
--- /dev/null
+++ b/mlir/lib/Dialect/MIOpenOps/CppOutput/CMakeLists.txt
@@ -0,0 +1,12 @@
+add_llvm_library(MLIRMIOpenCpp
+ ConvertToMIOpenCPP.cpp
+
+ ADDITIONAL_HEADER_DIRS
+ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/MIOpenOps
+ )
+target_link_libraries(MLIRMIOpenCpp
+ MLIRIR
+ MLIRMIOpenOps
+ MLIRStandardOps
+ MLIRTranslation)
+
diff --git a/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp b/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
new file mode 100644
index 000000000000..5fe33d695cb3
--- /dev/null
+++ b/mlir/lib/Dialect/MIOpenOps/CppOutput/ConvertToMIOpenCPP.cpp
@@ -0,0 +1,46 @@
+//===- ConvertToMIOpenCPP.cpp - MLIR to MIOpen C++ conversion -------------===//
+//
+// Part of the MLIR Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements a translation between the MLIR MIOpen dialect and C++.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/MIOpenOps/MIOpenCPP.h"
+#include "mlir/Dialect/MIOpenOps/MIOpenOps.h"
+#include "mlir/Dialect/StandardOps/Ops.h"
+
+#include "mlir/Translation.h"
+
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/Support/ToolOutputFile.h"
+
+using namespace mlir;
+
+std::unique_ptr<llvm::StringRef> mlir::translateModuleToMIOpenCPP(ModuleOp m) {
+ // Check constraints:
+ //
+ // The Module should only contain 1 function.
+ // The Function should only contain exactly:
+ // - 0 conv2d op.
+ // - 5 transform ops (1 for filter, 3 for input, 1 for output).
+ // - 1 gridwise gemm op.
+ m.dump();
+
+ return std::make_unique<llvm::StringRef>("Hello World");
+}
+
+static TranslateFromMLIRRegistration
+ toCPP("mlir-to-miopencpp", [](ModuleOp module, llvm::raw_ostream &output) {
+ auto sourceCode = mlir::translateModuleToMIOpenCPP(module);
+ if (!sourceCode)
+ return failure();
+
+ output << *sourceCode;
+ return success();
+ });
diff --git a/mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir b/mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir
new file mode 100644
index 000000000000..4b4bc0031717
--- /dev/null
+++ b/mlir/test/Dialect/MIOpen/CppOutput/miopencpp.mlir
@@ -0,0 +1,145 @@
+// RUN: mlir-translate -mlir-to-miopencpp %s | FileCheck %s
+
+// CHECK: Hello World
+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) {
+ layout = [
+ {
+ dimensions = [0],
+ names = ["gemmK"],
+ transformation = "merge",
+ source_dimensions = [1, 2, 3],
+ source_names = ["c", "y", "x"]
+ },
+ {
+ dimensions = [1],
+ names = ["gemmM"],
+ transformation = "passthrough",
+ source_dimensions = [0],
+ source_names = ["n"]
+ }
+ ]
+ } : memref<?x?x?x?xf32> to memref<?x?xf32>
+
+ // input tensor
+ %input_n_c_hipad_wipad = miopen.transform(%input) {
+ layout = [
+ {
+ dimensions = [0],
+ names = ["n"],
+ transformation = "passthorugh",
+ source_dimensions = [0],
+ source_names = ["n"]
+ },
+ {
+ dimensions = [1],
+ names = ["c"],
+ transformation = "passthorugh",
+ source_dimensions = [1],
+ source_names = ["c"]
+ },
+ {
+ dimensions = [2],
+ names = ["hipad"],
+ transformation = "pad",
+ parameters = [0, 0],
+ source_dimensions = [2],
+ source_names = ["hi"]
+ },
+ {
+ dimensions = [3],
+ names = ["wipad"],
+ transformation = "pad",
+ parameters = [0, 0],
+ source_dimensions = [3],
+ source_names = ["wi"]
+ }
+ ]
+ } : memref<?x?x?x?xf32> to memref<?x?x?x?xf32>
+
+ %input_n_c_y_ho_x_wo = miopen.transform(%input_n_c_hipad_wipad) {
+ layout = [
+ {
+ dimensions = [0],
+ names = ["n"],
+ transformation = "passthrough",
+ source_dimensions = [0],
+ source_names = ["n"]
+ },
+ {
+ dimensions = [1],
+ names = ["c"],
+ transformation = "passthrough",
+ source_dimensions = [1],
+ source_names = ["c"]
+ },
+ {
+ dimensions = [2, 3],
+ names = ["y", "ho"],
+ transformation = "embed",
+ parameters = [2, [1, 1, 0]],
+ source_dimensions = [2],
+ source_names = ["hipad"]
+ },
+ {
+ dimensions = [4, 5],
+ names = ["x", "wo"],
+ transformation = "embed",
+ parameters = [2, [1, 1, 0]],
+ source_dimensions = [2],
+ source_names = ["wipad"]
+ }
+ ]
+ } : memref<?x?x?x?xf32> to memref<?x?x?x?x?x?x?xf32>
+
+ %input_gemmK_gemmN = miopen.transform(%input_n_c_y_ho_x_wo) {
+ layout = [
+ {
+ dimensions = [0],
+ names = ["gemmK"],
+ transformation = "merge",
+ source_dimensions = [1, 2, 4],
+ source_names = ["c", "y", "x"]
+ },
+ {
+ dimensions = [1],
+ names = ["gemmN"],
+ transformation = "merge",
+ source_dimensions = [0, 3, 5],
+ source_names = ["n", "ho", "wo"]
+ }
+ ]
+ } : memref<?x?x?x?x?x?x?xf32> to memref<?x?xf32>
+
+ // output tensor
+ %output_gemmM_gemmN = miopen.transform(%output) {
+ layout = [
+ {
+ dimensions = [0],
+ names = ["gemmM"],
+ transformation = "passthrough",
+ source_dimensions = [1],
+ source_names = ["k"]
+ },
+ {
+ dimensions = [1],
+ names = ["gemmN"],
+ transformation = "merge",
+ source_dimensions = [0, 2, 3],
+ source_names = ["n", "ho", "wo"]
+ }
+ ]
+ } : memref<?x?x?x?xf32> to memref<?x?xf32>
+
+ // apply gridwise GEMM
+ miopen.gridwise_gemm(%filter_gemmK_gemmM, %input_gemmK_gemmN, %output_gemmM_gemmN) {
+ parameters = [
+ // tuning parameters
+ ]
+ } : memref<?x?xf32>,
+ memref<?x?xf32>,
+ memref<?x?xf32>
+
+ return
+}
More information about the llvm-branch-commits
mailing list