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