From b2d2f87c062bcdddefcf0b77ebd5ac76d018af27 Mon Sep 17 00:00:00 2001 From: Marko Rakita Date: Wed, 7 Aug 2024 12:51:39 +0200 Subject: [PATCH 01/11] Add stablehlo compile --- .gitignore | 1 + env/CMakeLists.txt | 2 +- third_party/CMakeLists.txt | 67 ++-------------------------------- third_party/stablehlo.cmake | 21 +++++++++++ third_party/tt-metal.cmake | 73 +++++++++++++++++++++++++++++++++++++ 5 files changed, 100 insertions(+), 64 deletions(-) create mode 100644 third_party/stablehlo.cmake create mode 100644 third_party/tt-metal.cmake diff --git a/.gitignore b/.gitignore index 75f4c3980..f18af65bb 100644 --- a/.gitignore +++ b/.gitignore @@ -1,5 +1,6 @@ .local build +third_party/stablehlo third_party/tt-metal .DS_STORE .vscode/* diff --git a/env/CMakeLists.txt b/env/CMakeLists.txt index f19b60ee4..01e956fd1 100644 --- a/env/CMakeLists.txt +++ b/env/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.20.0) project(ttmlir-toolchain LANGUAGES CXX C) set(FLATBUFFERS_VERSION "fb9afbafc7dfe226b9db54d4923bfb8839635274") -set(LLVM_PROJECT_VERSION "9ddfe62f5c11e3f65f444209f514029ded2d58b9") +set(LLVM_PROJECT_VERSION "a6d7828f4c50c1ec7b0b5f61fe59d7a768175dcc") include(ExternalProject) diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index 9d6584738..034792c8f 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -1,65 +1,6 @@ -include(ExternalProject) +option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" OFF) -if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") - set(ARCH_NAME "grayskull") - set(ARCH_EXTRA_DIR "grayskull") -elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") - set(ARCH_NAME "wormhole") - set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") -elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") - set(ARCH_NAME "blackhole") - set(ARCH_EXTRA_DIR "blackhole") -else() - message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") +include(tt-metal.cmake) +if (TTMLIR_ENABLE_STABLEHLO) +include(stablehlo.cmake) endif() - -set(TTMETAL_INCLUDE_DIRS - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f - PARENT_SCOPE -) - -set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) -set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) - -set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) -set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) - -ExternalProject_Add( - tt-metal - PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal - CMAKE_GENERATOR Ninja - CMAKE_ARGS - -DCMAKE_BUILD_TYPE=Release - -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} - -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - -DENABLE_TRACY=${TT_RUNTIME_ENABLE_PERF_TRACE} - GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git - GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 - GIT_PROGRESS ON - BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} -) - -set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) - -list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) -list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) - -foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) - add_library(${lib_name} SHARED IMPORTED GLOBAL) - set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) - add_dependencies(${lib_name} tt-metal) -endforeach() diff --git a/third_party/stablehlo.cmake b/third_party/stablehlo.cmake new file mode 100644 index 000000000..dce19e6f2 --- /dev/null +++ b/third_party/stablehlo.cmake @@ -0,0 +1,21 @@ +include(ExternalProject) + +set(STABLE_HLO_VERSION "v1.0.0") + +ExternalProject_Add( + stablehlo + PREFIX ${TTMLIR_SOURCE_DIR}/third_party/stablehlo + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + -DCMAKE_INSTALL_PREFIX=${TTMLIR_TOOLCHAIN_DIR} + # TODO: Figure out if we need this (or benefit from it) when it comes to linking stablehlo. + #-DSTABLEHLO_BUILD_EMBEDDED=ON + -DMLIR_DIR=${TTMLIR_TOOLCHAIN_DIR}/lib/cmake/mlir + GIT_REPOSITORY https://github.com/openxla/stablehlo + GIT_TAG ${STABLE_HLO_VERSION} + GIT_PROGRESS ON +) \ No newline at end of file diff --git a/third_party/tt-metal.cmake b/third_party/tt-metal.cmake new file mode 100644 index 000000000..7d9493899 --- /dev/null +++ b/third_party/tt-metal.cmake @@ -0,0 +1,73 @@ +include(ExternalProject) + +if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") + set(ARCH_NAME "grayskull") + set(ARCH_EXTRA_DIR "grayskull") +elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") + set(ARCH_NAME "wormhole") + set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") +elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") + set(ARCH_NAME "blackhole") + set(ARCH_EXTRA_DIR "blackhole") +else() + message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") +endif() + +if (TT_RUNTIME_ENABLE_PERF_TRACE) + add_compile_definitions(TRACY_ENABLE) + set(ENV{ENABLE_TRACY} "1") +else() + set(ENV{ENABLE_TRACY} "0") +endif() + +set(TTMETAL_INCLUDE_DIRS + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager + PARENT_SCOPE +) + +set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) +set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) +set(TTEAGER_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_eager.so) + +set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) +set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) +set(TTEAGER_LIBRARY_PATH ${TTEAGER_LIBRARY_PATH} PARENT_SCOPE) + + +ExternalProject_Add( + tt-metal + PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + -DTRACY_ENABLE=${TT_RUNTIME_ENABLE_PERF_TRACE} + GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git + GIT_TAG v0.49.0 + GIT_PROGRESS ON + BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} ${TTEAGER_LIBRARY_PATH} +) + +set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) + +list(APPEND library_names TTNN_LIBRARY TTEAGER_LIBRARY TTMETAL_LIBRARY) +list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} ${TTEAGER_LIBRARY_PATH}) + +foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) + add_library(${lib_name} SHARED IMPORTED GLOBAL) + set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) + add_dependencies(${lib_name} tt-metal) +endforeach() From cb4d067d3b2014540bfcfb8bb7ba9746d933a0ba Mon Sep 17 00:00:00 2001 From: Marko Rakita Date: Wed, 7 Aug 2024 15:04:34 +0200 Subject: [PATCH 02/11] Implement StableHLO to TTIR AddOp conversion --- include/ttmlir/Conversion/Passes.h | 1 + include/ttmlir/Conversion/Passes.td | 7 ++ .../StableHLOToTTIR/StableHLOToTTIR.h | 17 +++ lib/Conversion/CMakeLists.txt | 2 + lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 12 ++ .../StableHLOToTTIR/StableHLOToTTIR.cpp | 104 ++++++++++++++++++ 6 files changed, 143 insertions(+) create mode 100644 include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h create mode 100644 lib/Conversion/StableHLOToTTIR/CMakeLists.txt create mode 100644 lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp diff --git a/include/ttmlir/Conversion/Passes.h b/include/ttmlir/Conversion/Passes.h index bd4ee2753..2e4bfb82e 100644 --- a/include/ttmlir/Conversion/Passes.h +++ b/include/ttmlir/Conversion/Passes.h @@ -5,6 +5,7 @@ #ifndef TTMLIR_CONVERSION_PASSES_H #define TTMLIR_CONVERSION_PASSES_H +#include "ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h" #include "ttmlir/Conversion/TTIRToTTNN/TTIRToTTNN.h" #include "ttmlir/Conversion/TTNNToEmitC/TTNNToEmitC.h" #include "ttmlir/Conversion/TosaToTTIR/TosaToTTIR.h" diff --git a/include/ttmlir/Conversion/Passes.td b/include/ttmlir/Conversion/Passes.td index 42f2b267d..2b7d04162 100644 --- a/include/ttmlir/Conversion/Passes.td +++ b/include/ttmlir/Conversion/Passes.td @@ -7,6 +7,13 @@ include "mlir/Pass/PassBase.td" +def ConvertStableHLOToTTIR : Pass<"convert-stablehlo-to-ttir", "::mlir::ModuleOp"> { + let summary = "Convert StableHLO dialect to TTIR dialect."; + let constructor = "createConvertStableHLOToTTIRPass()"; + // TODO(mrakita): Probably will need to add some include here for StableHLO dialect. + let dependentDialects = ["mlir::stablehlo::StablehloDialect", "mlir::tt::ttir::TTIRDialect"]; +} + def ConvertTosaToTTIR : Pass<"convert-tosa-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert TOSA dialect to TTIR dialect."; let constructor = "createConvertTosaToTTIRPass()"; diff --git a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h new file mode 100644 index 000000000..985c5cd0c --- /dev/null +++ b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h @@ -0,0 +1,17 @@ +// SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC +// +// SPDX-License-Identifier: Apache-2.0 + +#ifndef TTMLIR_CONVERSION_StableHLOToTTIR_StableHLOToTTIR_H +#define TTMLIR_CONVERSION_StableHLOToTTIR_StableHLOToTTIR_H + +#include "mlir/IR/BuiltinOps.h" +#include "mlir/Pass/Pass.h" + +namespace mlir::tt { + +std::unique_ptr> createConvertStableHLOToTTIRPass(); + +} // namespace mlir::tt + +#endif \ No newline at end of file diff --git a/lib/Conversion/CMakeLists.txt b/lib/Conversion/CMakeLists.txt index c46bd1e68..f4d003b9f 100644 --- a/lib/Conversion/CMakeLists.txt +++ b/lib/Conversion/CMakeLists.txt @@ -1,3 +1,4 @@ +add_subdirectory(StableHLOToTTIR) add_subdirectory(TosaToTTIR) add_subdirectory(TTNNToEmitC) add_subdirectory(TTIRToTTNN) @@ -5,6 +6,7 @@ add_subdirectory(TTIRToTTNN) add_library(TTMLIRConversions INTERFACE) target_link_libraries(TTMLIRConversions INTERFACE + TTMLIRStableHLOToTTIR TTMLIRTosaToTTIR TTMLIRTTNNToEmitC TTMLIRTTIRToTTNN diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt new file mode 100644 index 000000000..7459dacba --- /dev/null +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -0,0 +1,12 @@ +add_mlir_library(TTMLIRStableHLOToTTIR + StableHLOToTTIR.cpp + + ADDITIONAL_HEADER_DIRS + ${PROJECT_SOURCE_DIR}/include/ttmlir/Conversion/StableHLOToTTIR + + DEPENDS + TTMLIRConversionPassIncGen + + LINK_LIBS PUBLIC + MLIR +) diff --git a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp new file mode 100644 index 000000000..8a1016c8e --- /dev/null +++ b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp @@ -0,0 +1,104 @@ +// SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC +// +// SPDX-License-Identifier: Apache-2.0 + +#include "ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h" +#include "ttmlir/Dialect/TT/IR/TT.h" +#include "ttmlir/Dialect/TT/IR/TTOpsTypes.h" +#include "ttmlir/Dialect/TTIR/IR/TTIR.h" +#include "ttmlir/Dialect/TTIR/IR/TTIROps.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "stablehlo/dialect/StablehloOps.h" + +using namespace mlir; +using namespace tt; + +namespace mlir::tt::ttir { + +#define GEN_PASS_DEF_CONVERTSTABLEHLOTOTTIR +#include "ttmlir/Conversion/Passes.h.inc" + +} // namespace mlir::tt::ttir + +namespace { + +template +class StableHLOToTTIROpConversionPattern : public OpConversionPattern { + using OpConversionPattern::OpConversionPattern; + +public: + LogicalResult + matchAndRewrite(SrcOp srcOp, Adaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + auto outputType = mlir::cast(srcOp.getResult().getType()); + auto outputTensor = rewriter.create( + srcOp.getLoc(), outputType.getShape(), outputType.getElementType()); + rewriter.replaceOpWithNewOp( + srcOp, TypeRange(outputTensor.getType()), adaptor.getOperands(), + ValueRange(outputTensor), + rewriter.getArrayAttr( + SmallVector(adaptor.getOperands().size() + 1, + rewriter.getAttr( + OperandConstraint::AnyDeviceTile)))); + return success(); + } +}; + +struct ConvertStableHLOToTTIRPass + : public ttir::impl::ConvertStableHLOToTTIRBase { + void runOnOperation() override { + mlir::ConversionTarget target(getContext()); + + target.addIllegalDialect(); + + target.addLegalDialect(); + target.addLegalOp(); + target.addLegalOp(); + target.addLegalOp(); + target.addLegalOp(); + + // For now keep the same type assuming StableHLO ops operate on builtin tensor. + TypeConverter typeConverter; + typeConverter.addConversion([](Type type) { + assert(isa(type) && + "only ranked tensor type supported"); + return type; + }); + RewritePatternSet patterns(&getContext()); + + // Add conversion patterns. + patterns + .add>( + typeConverter, &getContext()); + + // Apply conversion. + if (failed( + applyFullConversion(getOperation(), target, std::move(patterns)))) { + signalPassFailure(); + return; + } + } +}; + +} // namespace + +namespace mlir::tt { + +std::unique_ptr> createConvertStableHLOToTTIRPass() { + return std::make_unique(); +} + +} // namespace mlir::tt From 0e00099a8484530148da0edf9a1c37f34e80e723 Mon Sep 17 00:00:00 2001 From: Marko Rakita Date: Fri, 9 Aug 2024 15:15:50 +0200 Subject: [PATCH 03/11] Attempt linking stablehlo libraries --- .../Conversion/StableHLOToTTIR/StableHLOToTTIR.h | 2 +- lib/CMakeLists.txt | 4 ++++ lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 6 ++++++ lib/RegisterAll.cpp | 4 ++++ third_party/stablehlo.cmake | 15 ++++++++++++++- 5 files changed, 29 insertions(+), 2 deletions(-) diff --git a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h index 985c5cd0c..318d8d762 100644 --- a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h +++ b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h @@ -14,4 +14,4 @@ std::unique_ptr> createConvertStableHLOToTTIRPass(); } // namespace mlir::tt -#endif \ No newline at end of file +#endif diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index 4126bc41f..a7d653197 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -1,3 +1,7 @@ +# StableHLO includes +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) + add_subdirectory(CAPI) add_subdirectory(Conversion) add_subdirectory(Dialect) diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt index 7459dacba..90e300155 100644 --- a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -1,12 +1,18 @@ +get_property(STABLEHLO_LIBS GLOBAL PROPERTY STABLEHLO_LIBS) + add_mlir_library(TTMLIRStableHLOToTTIR StableHLOToTTIR.cpp ADDITIONAL_HEADER_DIRS ${PROJECT_SOURCE_DIR}/include/ttmlir/Conversion/StableHLOToTTIR + # StableHLO includes + ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo + ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build DEPENDS TTMLIRConversionPassIncGen LINK_LIBS PUBLIC MLIR + ${STABLEHLO_LIBS} ) diff --git a/lib/RegisterAll.cpp b/lib/RegisterAll.cpp index f4d25e300..3eff42e25 100644 --- a/lib/RegisterAll.cpp +++ b/lib/RegisterAll.cpp @@ -18,6 +18,8 @@ #include "ttmlir/Dialect/TTNN/Pipelines/Passes.h" #include "ttmlir/Dialect/TTNN/Transforms/Passes.h" +#include "stablehlo/dialect/Register.h" + void mlir::tt::registerAllDialects(mlir::DialectRegistry ®istry) { registry .insert(); + + mlir::stablehlo::registerAllDialects(registry); } void mlir::tt::registerAllPasses() { diff --git a/third_party/stablehlo.cmake b/third_party/stablehlo.cmake index dce19e6f2..fac2fd144 100644 --- a/third_party/stablehlo.cmake +++ b/third_party/stablehlo.cmake @@ -18,4 +18,17 @@ ExternalProject_Add( GIT_REPOSITORY https://github.com/openxla/stablehlo GIT_TAG ${STABLE_HLO_VERSION} GIT_PROGRESS ON -) \ No newline at end of file +) + +file(GLOB TTMLIR_STABLEHLO_LIBRARIES "${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib/*.a") +foreach(TTMLIR_STABLEHLO_LIBRARY ${TTMLIR_STABLEHLO_LIBRARIES}) + get_filename_component(lib_name ${TTMLIR_STABLEHLO_LIBRARY} NAME_WE) + string(REPLACE "lib" "" lib_name ${lib_name}) # Remove the "lib" prefix if it exists + message(STATUS "Adding TTMLIR library: ${lib_name}") + add_library(${lib_name} SHARED IMPORTED GLOBAL) + set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${TTMLIR_STABLEHLO_LIBRARY}) + add_dependencies(${lib_name} stablehlo) + list(APPEND STABLEHLO_LIBRARIES_LIST ${lib_name}) +endforeach() + +set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") \ No newline at end of file From 25867121e2a022688c391cb17c32db2628bff407 Mon Sep 17 00:00:00 2001 From: Marko Rakita Date: Fri, 9 Aug 2024 16:33:31 +0200 Subject: [PATCH 04/11] Add StableHLO to TTIR AddOp test --- test/ttmlir/Dialect/TTIR/stablehlo_to_ttir_addop.mlir | 10 ++++++++++ 1 file changed, 10 insertions(+) create mode 100644 test/ttmlir/Dialect/TTIR/stablehlo_to_ttir_addop.mlir diff --git a/test/ttmlir/Dialect/TTIR/stablehlo_to_ttir_addop.mlir b/test/ttmlir/Dialect/TTIR/stablehlo_to_ttir_addop.mlir new file mode 100644 index 000000000..1b22bff89 --- /dev/null +++ b/test/ttmlir/Dialect/TTIR/stablehlo_to_ttir_addop.mlir @@ -0,0 +1,10 @@ +// RUN: ttmlir-opt --convert-stablehlo-to-ttir %s | FileCheck %s +#any_device = #tt.operand_constraint +module @jit_eltwise_add attributes {} { + func.func public @test_add(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> { + %0 = stablehlo.add %arg0, %arg1 : tensor<13x21x3xf32> + // CHECK: %[[C:.*]] = tensor.empty[[C:.*]] + // CHECK: %[[C:.*]] = "ttir.add"[[C:.*]] + return %0 : tensor<13x21x3xf32> + } +} From 26f18853dc757dec0cb32ec9b71edfc9cfff3b12 Mon Sep 17 00:00:00 2001 From: uazizTT Date: Fri, 9 Aug 2024 14:48:34 -0400 Subject: [PATCH 05/11] Enable stablehlo in docker build. --- .github/workflows/docker-build.yml | 1 + CMakeLists.txt | 1 + third_party/CMakeLists.txt | 4 ++-- 3 files changed, 4 insertions(+), 2 deletions(-) diff --git a/.github/workflows/docker-build.yml b/.github/workflows/docker-build.yml index fb1724c97..8bb1930b9 100644 --- a/.github/workflows/docker-build.yml +++ b/.github/workflows/docker-build.yml @@ -55,6 +55,7 @@ jobs: -DCMAKE_BUILD_TYPE=${{ matrix.build.build_type }} \ -DTTMLIR_ENABLE_RUNTIME=${{ matrix.build.enable_runtime }} \ -DTTMLIR_ENABLE_RUNTIME_TESTS=${{ matrix.build.enable_runtime }} \ + -DTTMLIR_ENABLE_STABLEHLO=ON \ -S ${{ steps.strings.outputs.work-dir }} - name: Build diff --git a/CMakeLists.txt b/CMakeLists.txt index a2778a0a0..6432a01c5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,6 +7,7 @@ endif() option(TT_RUNTIME_ENABLE_PERF_TRACE "Enable performance mode" OFF) option(TTMLIR_ENABLE_RUNTIME "Enable runtime" OFF) +option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" OFF) set(CMAKE_BUILD_WITH_INSTALL_NAME_DIR ON) diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index 034792c8f..34be20d71 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -1,6 +1,6 @@ -option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" OFF) +include(ExternalProject) include(tt-metal.cmake) if (TTMLIR_ENABLE_STABLEHLO) -include(stablehlo.cmake) + include(stablehlo.cmake) endif() From ac2f6be69724711ae4de2fbfe0cf8bd5ac327cc2 Mon Sep 17 00:00:00 2001 From: uazizTT Date: Fri, 9 Aug 2024 15:26:30 -0400 Subject: [PATCH 06/11] restore changes to CMakeLists. --- third_party/tt-metal.cmake | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/third_party/tt-metal.cmake b/third_party/tt-metal.cmake index 7d9493899..bb1ff8729 100644 --- a/third_party/tt-metal.cmake +++ b/third_party/tt-metal.cmake @@ -22,6 +22,7 @@ endif() set(TTMETAL_INCLUDE_DIRS ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd @@ -31,19 +32,17 @@ set(TTMETAL_INCLUDE_DIRS ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f PARENT_SCOPE ) set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) -set(TTEAGER_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_eager.so) set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) -set(TTEAGER_LIBRARY_PATH ${TTEAGER_LIBRARY_PATH} PARENT_SCOPE) - ExternalProject_Add( tt-metal @@ -56,15 +55,15 @@ ExternalProject_Add( -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} -DTRACY_ENABLE=${TT_RUNTIME_ENABLE_PERF_TRACE} GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git - GIT_TAG v0.49.0 + GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 GIT_PROGRESS ON - BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} ${TTEAGER_LIBRARY_PATH} + BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} ) set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) -list(APPEND library_names TTNN_LIBRARY TTEAGER_LIBRARY TTMETAL_LIBRARY) -list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} ${TTEAGER_LIBRARY_PATH}) +list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) +list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) add_library(${lib_name} SHARED IMPORTED GLOBAL) From d9dd55706c05025a52bfff42b1e7477baf68de9a Mon Sep 17 00:00:00 2001 From: uazizTT Date: Mon, 12 Aug 2024 12:49:03 -0400 Subject: [PATCH 07/11] Sevaral fixes to resolve compilation and linking issues. --- CMakeLists.txt | 3 +- include/CMakeLists.txt | 2 ++ include/ttmlir/Conversion/Passes.h | 1 + lib/CMakeLists.txt | 8 ++++- lib/Conversion/CMakeLists.txt | 5 +++ lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 6 +++- lib/RegisterAll.cpp | 3 +- lib/SharedLib/CMakeLists.txt | 4 ++- runtime/lib/binary.cpp | 1 + third_party/stablehlo.cmake | 7 ++-- tools/CMakeLists.txt | 5 +++ tools/ttmlir-opt/CMakeLists.txt | 36 +++++++++++++++++-- tools/ttmlir-opt/ttmlir-opt.cpp | 3 ++ 13 files changed, 73 insertions(+), 11 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6432a01c5..23d3e0128 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,7 +7,7 @@ endif() option(TT_RUNTIME_ENABLE_PERF_TRACE "Enable performance mode" OFF) option(TTMLIR_ENABLE_RUNTIME "Enable runtime" OFF) -option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" OFF) +option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" ON) set(CMAKE_BUILD_WITH_INSTALL_NAME_DIR ON) @@ -44,6 +44,7 @@ include_directories(SYSTEM ${LLVM_INCLUDE_DIRS}) include_directories(SYSTEM ${MLIR_INCLUDE_DIRS}) include_directories(${TTMLIR_SOURCE_DIR}/include) include_directories(${TTMLIR_BINARY_DIR}/include) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/flatbuffers/include) link_directories(${LLVM_BUILD_LIBRARY_DIR}) add_definitions(${LLVM_DEFINITIONS}) include(TTMLIRPythonSitePackages) diff --git a/include/CMakeLists.txt b/include/CMakeLists.txt index 2b51d0d34..6e82d9606 100644 --- a/include/CMakeLists.txt +++ b/include/CMakeLists.txt @@ -1 +1,3 @@ add_subdirectory(ttmlir) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) diff --git a/include/ttmlir/Conversion/Passes.h b/include/ttmlir/Conversion/Passes.h index 2e4bfb82e..db1af667f 100644 --- a/include/ttmlir/Conversion/Passes.h +++ b/include/ttmlir/Conversion/Passes.h @@ -6,6 +6,7 @@ #define TTMLIR_CONVERSION_PASSES_H #include "ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h" + #include "ttmlir/Conversion/TTIRToTTNN/TTIRToTTNN.h" #include "ttmlir/Conversion/TTNNToEmitC/TTNNToEmitC.h" #include "ttmlir/Conversion/TosaToTTIR/TosaToTTIR.h" diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index a7d653197..874cad3c3 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -1,4 +1,3 @@ -# StableHLO includes include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) @@ -14,6 +13,10 @@ if (TTMLIR_ENABLE_RUNTIME) endif() add_mlir_library(TTMLIRStatic STATIC RegisterAll.cpp + + DEPENDS + stablehlo + LINK_LIBS PUBLIC MLIR MLIRTTDialect @@ -27,4 +30,7 @@ add_mlir_library(TTMLIRStatic STATIC RegisterAll.cpp MLIRTTMetalDialect MLIRTTMetalTransforms MLIRTTNNPipelines + TTMLIRStableHLOToTTIR ) + +add_dependencies(TTMLIRStatic stablehlo) diff --git a/lib/Conversion/CMakeLists.txt b/lib/Conversion/CMakeLists.txt index f4d003b9f..4ec7f3882 100644 --- a/lib/Conversion/CMakeLists.txt +++ b/lib/Conversion/CMakeLists.txt @@ -1,3 +1,6 @@ +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) + add_subdirectory(StableHLOToTTIR) add_subdirectory(TosaToTTIR) add_subdirectory(TTNNToEmitC) @@ -5,6 +8,8 @@ add_subdirectory(TTIRToTTNN) add_library(TTMLIRConversions INTERFACE) +add_dependencies(TTMLIRConversions stablehlo) + target_link_libraries(TTMLIRConversions INTERFACE TTMLIRStableHLOToTTIR TTMLIRTosaToTTIR diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt index 90e300155..30d5b7f7c 100644 --- a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -1,3 +1,6 @@ +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) + get_property(STABLEHLO_LIBS GLOBAL PROPERTY STABLEHLO_LIBS) add_mlir_library(TTMLIRStableHLOToTTIR @@ -11,8 +14,9 @@ add_mlir_library(TTMLIRStableHLOToTTIR DEPENDS TTMLIRConversionPassIncGen + stablehlo LINK_LIBS PUBLIC - MLIR + MLIRTransformUtils ${STABLEHLO_LIBS} ) diff --git a/lib/RegisterAll.cpp b/lib/RegisterAll.cpp index 3eff42e25..e6acbb54e 100644 --- a/lib/RegisterAll.cpp +++ b/lib/RegisterAll.cpp @@ -30,13 +30,12 @@ void mlir::tt::registerAllDialects(mlir::DialectRegistry ®istry) { mlir::scf::SCFDialect, mlir::cf::ControlFlowDialect, mlir::tosa::TosaDialect, mlir::vector::VectorDialect, mlir::emitc::EmitCDialect>(); - mlir::stablehlo::registerAllDialects(registry); } void mlir::tt::registerAllPasses() { // Register all dialect conversion passes - // + mlir::tt::registerTTMLIRConversionPasses(); mlir::tt::ttir::registerPasses(); diff --git a/lib/SharedLib/CMakeLists.txt b/lib/SharedLib/CMakeLists.txt index 81eeb5ae5..853ef49f3 100644 --- a/lib/SharedLib/CMakeLists.txt +++ b/lib/SharedLib/CMakeLists.txt @@ -18,6 +18,7 @@ set(TTMLIR_LIBS MLIRTTIRAnalysis MLIRTTNNPipelines TTMLIRTTNNToEmitC + TTMLIRStableHLOToTTIR ) # We supply empty.cpp because CMake does not allow creating a library without sources. @@ -27,6 +28,7 @@ add_dependencies(TTMLIR ${TTMLIR_LIBS} ${TTNN_RUNTIME_LIBS} ${TTNN_LIBS} + ${STABLEHLO_LIBS} ) target_link_libraries(TTMLIR PRIVATE @@ -41,7 +43,7 @@ target_link_libraries(TTMLIR PRIVATE -Wl,--no-whole-archive ${TTNN_LIBS} - + ${STABLEHLO_LIBS} flatbuffers ) diff --git a/runtime/lib/binary.cpp b/runtime/lib/binary.cpp index db84a8e23..498c245b3 100644 --- a/runtime/lib/binary.cpp +++ b/runtime/lib/binary.cpp @@ -4,6 +4,7 @@ #include +#pragma clang diagnostic ignored "-Wcovered-switch-default" #include "flatbuffers/idl.h" #include "tt/runtime/types.h" diff --git a/third_party/stablehlo.cmake b/third_party/stablehlo.cmake index fac2fd144..2ec106661 100644 --- a/third_party/stablehlo.cmake +++ b/third_party/stablehlo.cmake @@ -4,14 +4,14 @@ set(STABLE_HLO_VERSION "v1.0.0") ExternalProject_Add( stablehlo - PREFIX ${TTMLIR_SOURCE_DIR}/third_party/stablehlo + PREFIX ${PROJECT_SOURCE_DIR}/third_party/stablehlo CMAKE_GENERATOR Ninja CMAKE_ARGS -DCMAKE_BUILD_TYPE=Release -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} -DCMAKE_INSTALL_PREFIX=${TTMLIR_TOOLCHAIN_DIR} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} # TODO: Figure out if we need this (or benefit from it) when it comes to linking stablehlo. #-DSTABLEHLO_BUILD_EMBEDDED=ON -DMLIR_DIR=${TTMLIR_TOOLCHAIN_DIR}/lib/cmake/mlir @@ -31,4 +31,5 @@ foreach(TTMLIR_STABLEHLO_LIBRARY ${TTMLIR_STABLEHLO_LIBRARIES}) list(APPEND STABLEHLO_LIBRARIES_LIST ${lib_name}) endforeach() -set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") \ No newline at end of file +set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") + diff --git a/tools/CMakeLists.txt b/tools/CMakeLists.txt index 3899b3775..3717682b5 100644 --- a/tools/CMakeLists.txt +++ b/tools/CMakeLists.txt @@ -1,2 +1,7 @@ +include(ExternalProject) + +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) +include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) + add_subdirectory(ttmlir-opt) add_subdirectory(ttmlir-translate) diff --git a/tools/ttmlir-opt/CMakeLists.txt b/tools/ttmlir-opt/CMakeLists.txt index ef8810f9e..9c62121dd 100644 --- a/tools/ttmlir-opt/CMakeLists.txt +++ b/tools/ttmlir-opt/CMakeLists.txt @@ -1,9 +1,41 @@ + +include(CMakePrintHelpers) + +set(stablehlo_libs +CheckOps +ChloCAPI +ChloOps +InterpreterOps +StablehloAssemblyFormat +StablehloBase +StablehloBroadcastUtils +StablehloCAPI +StablehloLinalgTransforms +StablehloOps +StablehloPasses +StablehloPortableApi +StablehloRegister +StablehloSerialization +StablehloTOSATransforms +StablehloTestUtils +StablehloTypeInference +Version +VhloCAPI +VhloOps +VhloTypes) + get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) -set(LIBS ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp TTMLIRStatic) +set(LIBS ${stablehlo_libs} ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp TTMLIRStatic) + add_llvm_executable(ttmlir-opt ttmlir-opt.cpp) +add_dependencies(ttmlir-opt stablehlo) llvm_update_compile_flags(ttmlir-opt) +target_link_directories(ttmlir-opt PRIVATE ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib) + target_link_libraries(ttmlir-opt PRIVATE ${LIBS}) -mlir_check_all_link_libraries(ttmlir-opt) +mlir_check_all_link_libraries(ttmlir-opt ${LIBS}) + +set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") diff --git a/tools/ttmlir-opt/ttmlir-opt.cpp b/tools/ttmlir-opt/ttmlir-opt.cpp index d4f65dc2e..68a711432 100644 --- a/tools/ttmlir-opt/ttmlir-opt.cpp +++ b/tools/ttmlir-opt/ttmlir-opt.cpp @@ -7,6 +7,8 @@ #include "mlir/InitAllPasses.h" #include "mlir/Support/FileUtilities.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" +#include "stablehlo/dialect/Register.h" + #include "ttmlir/RegisterAll.h" @@ -16,6 +18,7 @@ int main(int argc, char **argv) { mlir::DialectRegistry registry; mlir::tt::registerAllDialects(registry); + mlir::stablehlo::registerAllDialects(registry); return mlir::asMainReturnCode( mlir::MlirOptMain(argc, argv, "ttmlir optimizer driver\n", registry)); From bc02c9bbcbe91682d6ed432cd1fa017834b297ae Mon Sep 17 00:00:00 2001 From: uazizTT Date: Tue, 13 Aug 2024 13:21:16 +0000 Subject: [PATCH 08/11] Enclose StableHLO features to be only enabled when TTMLIR_ENABLE_STABLEHLO flag is set. --- include/ttmlir/Conversion/Passes.h | 2 + include/ttmlir/Conversion/Passes.td | 3 ++ .../StableHLOToTTIR/StableHLOToTTIR.h | 4 +- lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 3 +- .../StableHLOToTTIR/StableHLOToTTIR.cpp | 4 +- lib/RegisterAll.cpp | 6 ++- lib/SharedLib/CMakeLists.txt | 2 + tools/ttmlir-opt/CMakeLists.txt | 45 ++++++++++++++++--- tools/ttmlir-opt/ttmlir-opt.cpp | 6 ++- 9 files changed, 61 insertions(+), 14 deletions(-) diff --git a/include/ttmlir/Conversion/Passes.h b/include/ttmlir/Conversion/Passes.h index db1af667f..cf7958ba1 100644 --- a/include/ttmlir/Conversion/Passes.h +++ b/include/ttmlir/Conversion/Passes.h @@ -5,7 +5,9 @@ #ifndef TTMLIR_CONVERSION_PASSES_H #define TTMLIR_CONVERSION_PASSES_H +#ifdef TTMLIR_ENABLE_STABLEHLO #include "ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h" +#endif #include "ttmlir/Conversion/TTIRToTTNN/TTIRToTTNN.h" #include "ttmlir/Conversion/TTNNToEmitC/TTNNToEmitC.h" diff --git a/include/ttmlir/Conversion/Passes.td b/include/ttmlir/Conversion/Passes.td index 2b7d04162..f0e0961f6 100644 --- a/include/ttmlir/Conversion/Passes.td +++ b/include/ttmlir/Conversion/Passes.td @@ -7,12 +7,15 @@ include "mlir/Pass/PassBase.td" + +#ifdef TTMLIR_ENABLE_STABLEHLO def ConvertStableHLOToTTIR : Pass<"convert-stablehlo-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert StableHLO dialect to TTIR dialect."; let constructor = "createConvertStableHLOToTTIRPass()"; // TODO(mrakita): Probably will need to add some include here for StableHLO dialect. let dependentDialects = ["mlir::stablehlo::StablehloDialect", "mlir::tt::ttir::TTIRDialect"]; } +#endif def ConvertTosaToTTIR : Pass<"convert-tosa-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert TOSA dialect to TTIR dialect."; diff --git a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h index 318d8d762..8fe0a9bb5 100644 --- a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h +++ b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h @@ -9,9 +9,9 @@ #include "mlir/Pass/Pass.h" namespace mlir::tt { - +#ifdef TTMLIR_ENABLE_STABLEHLO std::unique_ptr> createConvertStableHLOToTTIRPass(); - +#endif } // namespace mlir::tt #endif diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt index 30d5b7f7c..72051a5df 100644 --- a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -2,7 +2,7 @@ include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) get_property(STABLEHLO_LIBS GLOBAL PROPERTY STABLEHLO_LIBS) - +if (TTMLIR_ENABLE_STABLEHLO) add_mlir_library(TTMLIRStableHLOToTTIR StableHLOToTTIR.cpp @@ -20,3 +20,4 @@ add_mlir_library(TTMLIRStableHLOToTTIR MLIRTransformUtils ${STABLEHLO_LIBS} ) +endif () diff --git a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp index 8a1016c8e..17d8761fe 100644 --- a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp +++ b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp @@ -19,9 +19,8 @@ #include #include #include - +#ifdef TTMLIR_ENABLE_STABLEHLO #include "stablehlo/dialect/StablehloOps.h" - using namespace mlir; using namespace tt; @@ -102,3 +101,4 @@ std::unique_ptr> createConvertStableHLOToTTIRPass() { } } // namespace mlir::tt +#endif diff --git a/lib/RegisterAll.cpp b/lib/RegisterAll.cpp index e6acbb54e..2ef90dcf3 100644 --- a/lib/RegisterAll.cpp +++ b/lib/RegisterAll.cpp @@ -17,9 +17,9 @@ #include "ttmlir/Dialect/TTNN/IR/TTNN.h" #include "ttmlir/Dialect/TTNN/Pipelines/Passes.h" #include "ttmlir/Dialect/TTNN/Transforms/Passes.h" - +#ifdef TTMLIR_ENABLE_STABLEHLO #include "stablehlo/dialect/Register.h" - +#endif void mlir::tt::registerAllDialects(mlir::DialectRegistry ®istry) { registry .insert(); +#ifdef TTMLIR_ENABLE_STABLEHLO mlir::stablehlo::registerAllDialects(registry); +#endif } void mlir::tt::registerAllPasses() { diff --git a/lib/SharedLib/CMakeLists.txt b/lib/SharedLib/CMakeLists.txt index 853ef49f3..23294c4bd 100644 --- a/lib/SharedLib/CMakeLists.txt +++ b/lib/SharedLib/CMakeLists.txt @@ -18,7 +18,9 @@ set(TTMLIR_LIBS MLIRTTIRAnalysis MLIRTTNNPipelines TTMLIRTTNNToEmitC +if (TTMLIR_ENABLE_STABLEHLO) TTMLIRStableHLOToTTIR +endif () ) # We supply empty.cpp because CMake does not allow creating a library without sources. diff --git a/tools/ttmlir-opt/CMakeLists.txt b/tools/ttmlir-opt/CMakeLists.txt index 9c62121dd..149c3f714 100644 --- a/tools/ttmlir-opt/CMakeLists.txt +++ b/tools/ttmlir-opt/CMakeLists.txt @@ -1,6 +1,3 @@ - -include(CMakePrintHelpers) - set(stablehlo_libs CheckOps ChloCAPI @@ -14,6 +11,21 @@ StablehloLinalgTransforms StablehloOps StablehloPasses StablehloPortableApi +StablehloReferenceApi +StablehloReferenceAxes +StablehloReferenceConfiguration +StablehloReferenceElement +StablehloReferenceErrors +StablehloReferenceIndex +StablehloReferenceNumPy +StablehloReferenceOps +StablehloReferenceProcess +StablehloReferenceProcessGrid +StablehloReferenceScope +StablehloReferenceTensor +StablehloReferenceToken +StablehloReferenceTypes +StablehloReferenceValue StablehloRegister StablehloSerialization StablehloTOSATransforms @@ -26,13 +38,36 @@ VhloTypes) get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) -set(LIBS ${stablehlo_libs} ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp TTMLIRStatic) +set(LIBS TTMLIRStableHLOToTTIR TTMLIRCAPI ${stablehlo_libs} ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp +MLIRTTIRAnalysis +TTMLIRTTIRToTTNN +MLIRTTNNTransforms +TTMLIRTosaToTTIR +TTMLIRTTNNToEmitC +MLIRTTMetalTransforms +MLIRTTIRTransforms +MLIRTTMetalDialect +MLIRTTKernelDialect +MLIRTTIRDialect +MLIRTTNNDialect +MLIRTTNNPipelines +MLIRTTDialect +TTNNTargetFlatbuffer +TTMLIRCAPI +TTMLIRStableHLOToTTIR +TTMLIRStatic) add_llvm_executable(ttmlir-opt ttmlir-opt.cpp) add_dependencies(ttmlir-opt stablehlo) llvm_update_compile_flags(ttmlir-opt) -target_link_directories(ttmlir-opt PRIVATE ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib) +target_link_directories(ttmlir-opt PUBLIC ${TTMLIR_TOOLCHAIN_DIR}/lib) +target_link_directories(ttmlir-opt PUBLIC ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib) + +ADD_CUSTOM_COMMAND(TARGET ttmlir-opt + PRE_BUILD + COMMAND cp ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib/*.a ${TTMLIR_TOOLCHAIN_DIR}/lib + DEPENDS stablehlo) target_link_libraries(ttmlir-opt PRIVATE ${LIBS}) diff --git a/tools/ttmlir-opt/ttmlir-opt.cpp b/tools/ttmlir-opt/ttmlir-opt.cpp index 68a711432..34e5f3cad 100644 --- a/tools/ttmlir-opt/ttmlir-opt.cpp +++ b/tools/ttmlir-opt/ttmlir-opt.cpp @@ -7,8 +7,9 @@ #include "mlir/InitAllPasses.h" #include "mlir/Support/FileUtilities.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" +#ifdef TTMLIR_ENABLE_STABLEHLO #include "stablehlo/dialect/Register.h" - +#endif #include "ttmlir/RegisterAll.h" @@ -18,8 +19,9 @@ int main(int argc, char **argv) { mlir::DialectRegistry registry; mlir::tt::registerAllDialects(registry); +#ifdef TTMLIR_ENABLE_STABLEHLO mlir::stablehlo::registerAllDialects(registry); - +#endif return mlir::asMainReturnCode( mlir::MlirOptMain(argc, argv, "ttmlir optimizer driver\n", registry)); } From d5d51f6d2c288821a4c40931153097edb844c41a Mon Sep 17 00:00:00 2001 From: uazizTT Date: Wed, 14 Aug 2024 20:06:54 -0400 Subject: [PATCH 09/11] Move StableHLO to env dir and resolve build issues. --- CMakeLists.txt | 6 ++ env/CMakeLists.txt | 22 +++++- include/ttmlir/Conversion/Passes.td | 3 - .../StableHLOToTTIR/StableHLOToTTIR.h | 8 ++ .../Dialect/TTKernel/IR/TTKernelOpsTypes.td | 4 +- lib/CMakeLists.txt | 42 +++++----- lib/Conversion/CMakeLists.txt | 24 ++++-- lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 10 +-- .../StableHLOToTTIR/StableHLOToTTIR.cpp | 3 +- lib/RegisterAll.cpp | 4 - lib/SharedLib/CMakeLists.txt | 7 +- third_party/CMakeLists.txt | 72 ++++++++++++++++- third_party/stablehlo.cmake | 35 --------- third_party/tt-metal.cmake | 72 ----------------- tools/ttmlir-opt/CMakeLists.txt | 77 +++---------------- 15 files changed, 161 insertions(+), 228 deletions(-) delete mode 100644 third_party/stablehlo.cmake delete mode 100644 third_party/tt-metal.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 23d3e0128..d84888e97 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -9,6 +9,10 @@ option(TT_RUNTIME_ENABLE_PERF_TRACE "Enable performance mode" OFF) option(TTMLIR_ENABLE_RUNTIME "Enable runtime" OFF) option(TTMLIR_ENABLE_STABLEHLO "Enable StableHLO support" ON) +if (TTMLIR_ENABLE_STABLEHLO) + add_compile_definitions(TTMLIR_ENABLE_STABLEHLO) +endif() + set(CMAKE_BUILD_WITH_INSTALL_NAME_DIR ON) set(CMAKE_CXX_STANDARD 17 CACHE STRING "C++ standard to conform to") @@ -42,6 +46,8 @@ set(TTMLIR_BINARY_DIR ${PROJECT_BINARY_DIR}) set(LLVM_LIT_TOOLS_DIR "${TTMLIR_TOOLCHAIN_DIR}/src/llvm-project/llvm/utils/lit") include_directories(SYSTEM ${LLVM_INCLUDE_DIRS}) include_directories(SYSTEM ${MLIR_INCLUDE_DIRS}) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) include_directories(${TTMLIR_SOURCE_DIR}/include) include_directories(${TTMLIR_BINARY_DIR}/include) include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/flatbuffers/include) diff --git a/env/CMakeLists.txt b/env/CMakeLists.txt index 01e956fd1..a8d1607cd 100644 --- a/env/CMakeLists.txt +++ b/env/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.20.0) project(ttmlir-toolchain LANGUAGES CXX C) set(FLATBUFFERS_VERSION "fb9afbafc7dfe226b9db54d4923bfb8839635274") -set(LLVM_PROJECT_VERSION "a6d7828f4c50c1ec7b0b5f61fe59d7a768175dcc") +set(LLVM_PROJECT_VERSION "9ddfe62f5c11e3f65f444209f514029ded2d58b9") include(ExternalProject) @@ -67,4 +67,24 @@ ExternalProject_Add( DEPENDS python-venv ) +set(STABLE_HLO_VERSION "v1.5.0") + +ExternalProject_Add( + stablehlo + PREFIX ${TTMLIR_TOOLCHAIN_DIR} + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_INSTALL_PREFIX=${TTMLIR_TOOLCHAIN_DIR} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + # TODO: Figure out if we need this (or benefit from it) when it comes to linking stablehlo. + #-DSTABLEHLO_BUILD_EMBEDDED=ON + -DMLIR_DIR=${TTMLIR_TOOLCHAIN_DIR}/lib/cmake/mlir + GIT_REPOSITORY https://github.com/openxla/stablehlo + GIT_TAG ${STABLE_HLO_VERSION} + GIT_PROGRESS ON +) + add_custom_target(llvm-lit ALL COMMAND cp ${TTMLIR_TOOLCHAIN_DIR}/src/llvm-project-build/bin/llvm-lit ${TTMLIR_TOOLCHAIN_DIR}/bin/llvm-lit DEPENDS llvm-project) diff --git a/include/ttmlir/Conversion/Passes.td b/include/ttmlir/Conversion/Passes.td index f0e0961f6..2b7d04162 100644 --- a/include/ttmlir/Conversion/Passes.td +++ b/include/ttmlir/Conversion/Passes.td @@ -7,15 +7,12 @@ include "mlir/Pass/PassBase.td" - -#ifdef TTMLIR_ENABLE_STABLEHLO def ConvertStableHLOToTTIR : Pass<"convert-stablehlo-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert StableHLO dialect to TTIR dialect."; let constructor = "createConvertStableHLOToTTIRPass()"; // TODO(mrakita): Probably will need to add some include here for StableHLO dialect. let dependentDialects = ["mlir::stablehlo::StablehloDialect", "mlir::tt::ttir::TTIRDialect"]; } -#endif def ConvertTosaToTTIR : Pass<"convert-tosa-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert TOSA dialect to TTIR dialect."; diff --git a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h index 8fe0a9bb5..1df166900 100644 --- a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h +++ b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h @@ -9,9 +9,17 @@ #include "mlir/Pass/Pass.h" namespace mlir::tt { + #ifdef TTMLIR_ENABLE_STABLEHLO std::unique_ptr> createConvertStableHLOToTTIRPass(); +#else +// define an empty function in case conversion definition does not exist. +std::unique_ptr> createConvertStableHLOToTTIRPass() +{ +return std::unique_ptr>(nullptr); +} #endif + } // namespace mlir::tt #endif diff --git a/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td b/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td index 1fce627d6..3e257a9fb 100644 --- a/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td +++ b/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td @@ -31,8 +31,8 @@ def TTKernel_CB : TTKernel_Type<"CB", "cb"> { static CBType get(::mlir::MLIRContext *context, uint64_t address, uint64_t port, - MemRefType memref) { - uint64_t numBuffers = 1; + MemRefType memref, + uint64_t numBuffers = 1) { uint64_t pageSize = 0; if (::mlir::isa<::mlir::tt::TileType>(memref.getElementType())) { pageSize = ::mlir::cast<::mlir::tt::TileType>(memref.getElementType()).getSizeBytes(); diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index 874cad3c3..e0bf4f0df 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -1,5 +1,5 @@ -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) add_subdirectory(CAPI) add_subdirectory(Conversion) @@ -8,29 +8,33 @@ add_subdirectory(Target) # Shared library will include runtime code # so we only build it if runtime is enabled -if (TTMLIR_ENABLE_RUNTIME) - add_subdirectory(SharedLib) + +set(ttmlir_dependencies "") +set(link_libs +MLIR; +MLIRTTDialect; +MLIRTTIRDialect; +MLIRTTIRTransforms; +TTMLIRConversions; +MLIRTTIRAnalysis; +MLIRTTNNDialect; +MLIRTTNNTransforms; +MLIRTTKernelDialect; +MLIRTTMetalDialect; +MLIRTTMetalTransforms; +MLIRTTNNPipelines +) + +if (TTMLIR_ENABLE_STABLEHLO) + list(APPEND link_libs TTMLIRStableHLOToTTIR) endif() add_mlir_library(TTMLIRStatic STATIC RegisterAll.cpp DEPENDS - stablehlo + ${ttmlir_dependencies} LINK_LIBS PUBLIC - MLIR - MLIRTTDialect - MLIRTTIRDialect - MLIRTTIRTransforms - TTMLIRConversions - MLIRTTIRAnalysis - MLIRTTNNDialect - MLIRTTNNTransforms - MLIRTTKernelDialect - MLIRTTMetalDialect - MLIRTTMetalTransforms - MLIRTTNNPipelines - TTMLIRStableHLOToTTIR + ${link_libs} ) -add_dependencies(TTMLIRStatic stablehlo) diff --git a/lib/Conversion/CMakeLists.txt b/lib/Conversion/CMakeLists.txt index 4ec7f3882..0a8298967 100644 --- a/lib/Conversion/CMakeLists.txt +++ b/lib/Conversion/CMakeLists.txt @@ -1,18 +1,26 @@ -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) +add_library(TTMLIRConversions INTERFACE) +if (TTMLIR_ENABLE_STABLEHLO) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) add_subdirectory(StableHLOToTTIR) +add_dependencies(TTMLIRConversions stablehlo) +endif() + add_subdirectory(TosaToTTIR) add_subdirectory(TTNNToEmitC) add_subdirectory(TTIRToTTNN) -add_library(TTMLIRConversions INTERFACE) +set(link_libs +TTMLIRTosaToTTIR; +TTMLIRTTNNToEmitC; +TTMLIRTTIRToTTNN +) -add_dependencies(TTMLIRConversions stablehlo) +if (TTMLIR_ENABLE_STABLEHLO) +list(APPEND link_libs TTMLIRStableHLOToTTIR) +endif() target_link_libraries(TTMLIRConversions INTERFACE - TTMLIRStableHLOToTTIR - TTMLIRTosaToTTIR - TTMLIRTTNNToEmitC - TTMLIRTTIRToTTNN + ${link_libs} ) diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt index 72051a5df..8cd88059d 100644 --- a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -1,23 +1,17 @@ -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) get_property(STABLEHLO_LIBS GLOBAL PROPERTY STABLEHLO_LIBS) -if (TTMLIR_ENABLE_STABLEHLO) add_mlir_library(TTMLIRStableHLOToTTIR StableHLOToTTIR.cpp ADDITIONAL_HEADER_DIRS ${PROJECT_SOURCE_DIR}/include/ttmlir/Conversion/StableHLOToTTIR - # StableHLO includes - ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo - ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build DEPENDS TTMLIRConversionPassIncGen - stablehlo LINK_LIBS PUBLIC MLIRTransformUtils ${STABLEHLO_LIBS} ) -endif () diff --git a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp index 17d8761fe..3fe52ab47 100644 --- a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp +++ b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIR.cpp @@ -19,7 +19,7 @@ #include #include #include -#ifdef TTMLIR_ENABLE_STABLEHLO + #include "stablehlo/dialect/StablehloOps.h" using namespace mlir; using namespace tt; @@ -101,4 +101,3 @@ std::unique_ptr> createConvertStableHLOToTTIRPass() { } } // namespace mlir::tt -#endif diff --git a/lib/RegisterAll.cpp b/lib/RegisterAll.cpp index 2ef90dcf3..38025d6ef 100644 --- a/lib/RegisterAll.cpp +++ b/lib/RegisterAll.cpp @@ -17,9 +17,7 @@ #include "ttmlir/Dialect/TTNN/IR/TTNN.h" #include "ttmlir/Dialect/TTNN/Pipelines/Passes.h" #include "ttmlir/Dialect/TTNN/Transforms/Passes.h" -#ifdef TTMLIR_ENABLE_STABLEHLO #include "stablehlo/dialect/Register.h" -#endif void mlir::tt::registerAllDialects(mlir::DialectRegistry ®istry) { registry .insert(); -#ifdef TTMLIR_ENABLE_STABLEHLO mlir::stablehlo::registerAllDialects(registry); -#endif } void mlir::tt::registerAllPasses() { diff --git a/lib/SharedLib/CMakeLists.txt b/lib/SharedLib/CMakeLists.txt index 23294c4bd..ceaf708c4 100644 --- a/lib/SharedLib/CMakeLists.txt +++ b/lib/SharedLib/CMakeLists.txt @@ -18,11 +18,12 @@ set(TTMLIR_LIBS MLIRTTIRAnalysis MLIRTTNNPipelines TTMLIRTTNNToEmitC -if (TTMLIR_ENABLE_STABLEHLO) - TTMLIRStableHLOToTTIR -endif () ) +if (TTMLIR_ENABLE_STABLEHLO) +list(APPEND TTMLIR_LIBS TTMLIRStableHLOToTTIR) +endif() + # We supply empty.cpp because CMake does not allow creating a library without sources. add_library(TTMLIR SHARED empty.cpp) diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index 34be20d71..bb1ff8729 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -1,6 +1,72 @@ include(ExternalProject) -include(tt-metal.cmake) -if (TTMLIR_ENABLE_STABLEHLO) - include(stablehlo.cmake) +if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") + set(ARCH_NAME "grayskull") + set(ARCH_EXTRA_DIR "grayskull") +elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") + set(ARCH_NAME "wormhole") + set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") +elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") + set(ARCH_NAME "blackhole") + set(ARCH_EXTRA_DIR "blackhole") +else() + message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") endif() + +if (TT_RUNTIME_ENABLE_PERF_TRACE) + add_compile_definitions(TRACY_ENABLE) + set(ENV{ENABLE_TRACY} "1") +else() + set(ENV{ENABLE_TRACY} "0") +endif() + +set(TTMETAL_INCLUDE_DIRS + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f + PARENT_SCOPE +) + +set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) +set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) + +set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) +set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) + +ExternalProject_Add( + tt-metal + PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + -DTRACY_ENABLE=${TT_RUNTIME_ENABLE_PERF_TRACE} + GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git + GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 + GIT_PROGRESS ON + BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} +) + +set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) + +list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) +list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) + +foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) + add_library(${lib_name} SHARED IMPORTED GLOBAL) + set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) + add_dependencies(${lib_name} tt-metal) +endforeach() diff --git a/third_party/stablehlo.cmake b/third_party/stablehlo.cmake deleted file mode 100644 index 2ec106661..000000000 --- a/third_party/stablehlo.cmake +++ /dev/null @@ -1,35 +0,0 @@ -include(ExternalProject) - -set(STABLE_HLO_VERSION "v1.0.0") - -ExternalProject_Add( - stablehlo - PREFIX ${PROJECT_SOURCE_DIR}/third_party/stablehlo - CMAKE_GENERATOR Ninja - CMAKE_ARGS - -DCMAKE_BUILD_TYPE=Release - -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} - -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_INSTALL_PREFIX=${TTMLIR_TOOLCHAIN_DIR} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - # TODO: Figure out if we need this (or benefit from it) when it comes to linking stablehlo. - #-DSTABLEHLO_BUILD_EMBEDDED=ON - -DMLIR_DIR=${TTMLIR_TOOLCHAIN_DIR}/lib/cmake/mlir - GIT_REPOSITORY https://github.com/openxla/stablehlo - GIT_TAG ${STABLE_HLO_VERSION} - GIT_PROGRESS ON -) - -file(GLOB TTMLIR_STABLEHLO_LIBRARIES "${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib/*.a") -foreach(TTMLIR_STABLEHLO_LIBRARY ${TTMLIR_STABLEHLO_LIBRARIES}) - get_filename_component(lib_name ${TTMLIR_STABLEHLO_LIBRARY} NAME_WE) - string(REPLACE "lib" "" lib_name ${lib_name}) # Remove the "lib" prefix if it exists - message(STATUS "Adding TTMLIR library: ${lib_name}") - add_library(${lib_name} SHARED IMPORTED GLOBAL) - set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${TTMLIR_STABLEHLO_LIBRARY}) - add_dependencies(${lib_name} stablehlo) - list(APPEND STABLEHLO_LIBRARIES_LIST ${lib_name}) -endforeach() - -set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") - diff --git a/third_party/tt-metal.cmake b/third_party/tt-metal.cmake deleted file mode 100644 index bb1ff8729..000000000 --- a/third_party/tt-metal.cmake +++ /dev/null @@ -1,72 +0,0 @@ -include(ExternalProject) - -if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") - set(ARCH_NAME "grayskull") - set(ARCH_EXTRA_DIR "grayskull") -elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") - set(ARCH_NAME "wormhole") - set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") -elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") - set(ARCH_NAME "blackhole") - set(ARCH_EXTRA_DIR "blackhole") -else() - message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") -endif() - -if (TT_RUNTIME_ENABLE_PERF_TRACE) - add_compile_definitions(TRACY_ENABLE) - set(ENV{ENABLE_TRACY} "1") -else() - set(ENV{ENABLE_TRACY} "0") -endif() - -set(TTMETAL_INCLUDE_DIRS - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f - PARENT_SCOPE -) - -set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) -set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) - -set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) -set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) - -ExternalProject_Add( - tt-metal - PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal - CMAKE_GENERATOR Ninja - CMAKE_ARGS - -DCMAKE_BUILD_TYPE=Release - -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} - -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - -DTRACY_ENABLE=${TT_RUNTIME_ENABLE_PERF_TRACE} - GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git - GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 - GIT_PROGRESS ON - BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} -) - -set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) - -list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) -list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) - -foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) - add_library(${lib_name} SHARED IMPORTED GLOBAL) - set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) - add_dependencies(${lib_name} tt-metal) -endforeach() diff --git a/tools/ttmlir-opt/CMakeLists.txt b/tools/ttmlir-opt/CMakeLists.txt index 149c3f714..7e3284bd5 100644 --- a/tools/ttmlir-opt/CMakeLists.txt +++ b/tools/ttmlir-opt/CMakeLists.txt @@ -1,76 +1,17 @@ -set(stablehlo_libs -CheckOps -ChloCAPI -ChloOps -InterpreterOps -StablehloAssemblyFormat -StablehloBase -StablehloBroadcastUtils -StablehloCAPI -StablehloLinalgTransforms -StablehloOps -StablehloPasses -StablehloPortableApi -StablehloReferenceApi -StablehloReferenceAxes -StablehloReferenceConfiguration -StablehloReferenceElement -StablehloReferenceErrors -StablehloReferenceIndex -StablehloReferenceNumPy -StablehloReferenceOps -StablehloReferenceProcess -StablehloReferenceProcessGrid -StablehloReferenceScope -StablehloReferenceTensor -StablehloReferenceToken -StablehloReferenceTypes -StablehloReferenceValue -StablehloRegister -StablehloSerialization -StablehloTOSATransforms -StablehloTestUtils -StablehloTypeInference -Version -VhloCAPI -VhloOps -VhloTypes) - get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) -set(LIBS TTMLIRStableHLOToTTIR TTMLIRCAPI ${stablehlo_libs} ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp -MLIRTTIRAnalysis -TTMLIRTTIRToTTNN -MLIRTTNNTransforms -TTMLIRTosaToTTIR -TTMLIRTTNNToEmitC -MLIRTTMetalTransforms -MLIRTTIRTransforms -MLIRTTMetalDialect -MLIRTTKernelDialect -MLIRTTIRDialect -MLIRTTNNDialect -MLIRTTNNPipelines -MLIRTTDialect -TTNNTargetFlatbuffer -TTMLIRCAPI -TTMLIRStableHLOToTTIR -TTMLIRStatic) +set(LIBS ${dialect_libs} ${conversion_libs} StablehloBase TTNNTargetFlatbuffer MLIRTTIRDialect TTMLIRConversions MLIRTransformUtils MLIROptLib MLIRTargetCpp TTMLIRStatic CheckOps ChloCAPI ChloOps InterpreterOps StablehloAssemblyFormat StablehloBase StablehloBroadcastUtils StablehloCAPI StablehloLinalgTransforms StablehloOps StablehloPasses StablehloPortableApi StablehloRegister StablehloSerialization StablehloTOSATransforms StablehloTestUtils StablehloTypeInference Version VhloCAPI VhloOps VhloTypes) -add_llvm_executable(ttmlir-opt ttmlir-opt.cpp) -add_dependencies(ttmlir-opt stablehlo) +if (TTMLIR_ENABLE_STABLEHLO) + list(APPEND LIBS TTMLIRStableHLOToTTIR) +endif() -llvm_update_compile_flags(ttmlir-opt) -target_link_directories(ttmlir-opt PUBLIC ${TTMLIR_TOOLCHAIN_DIR}/lib) -target_link_directories(ttmlir-opt PUBLIC ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib) +add_llvm_executable(ttmlir-opt ttmlir-opt.cpp) -ADD_CUSTOM_COMMAND(TARGET ttmlir-opt - PRE_BUILD - COMMAND cp ${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build/lib/*.a ${TTMLIR_TOOLCHAIN_DIR}/lib - DEPENDS stablehlo) +target_link_directories(ttmlir-opt PRIVATE ${TTMLIR_PROJECT_DIR}/build/lib) +target_link_libraries(ttmlir-opt PRIVATE ${LIBS} libStablehloBase.a) +target_link_directories(ttmlir-opt PRIVATE ${TTMLIR_TOOLCHAIN_DIR}/lib) -target_link_libraries(ttmlir-opt PRIVATE ${LIBS}) +llvm_update_compile_flags(ttmlir-opt) mlir_check_all_link_libraries(ttmlir-opt ${LIBS}) - -set_property(GLOBAL PROPERTY STABLEHLO_LIBS "${STABLEHLO_LIBRARIES_LIST}") From cb90754fcf3dd04b55390b1bf7d95b0580009e1c Mon Sep 17 00:00:00 2001 From: uazizTT Date: Mon, 19 Aug 2024 15:56:08 -0400 Subject: [PATCH 10/11] Switch stablehlo to embedded build. --- .gitmodules | 3 ++ CMakeLists.txt | 10 +++-- env/CMakeLists.txt | 20 --------- include/CMakeLists.txt | 2 - include/ttmlir/Conversion/CMakeLists.txt | 8 ++++ include/ttmlir/Conversion/Passes.h | 1 - include/ttmlir/Conversion/Passes.td | 4 +- .../StableHLOToTTIR/StableHLOToTTIR.h | 9 +--- .../Dialect/TTKernel/IR/TTKernelOpsTypes.td | 4 +- lib/CMakeLists.txt | 42 +++++++------------ lib/Conversion/CMakeLists.txt | 12 +++--- lib/Conversion/StableHLOToTTIR/CMakeLists.txt | 7 ++-- lib/RegisterAll.cpp | 9 +++- lib/SharedLib/CMakeLists.txt | 7 +--- runtime/lib/binary.cpp | 1 - test/ttmlir/ttmlir-opt.mlir | 17 +++++++- third_party/CMakeLists.txt | 9 +--- tools/CMakeLists.txt | 5 --- tools/ttmlir-opt/CMakeLists.txt | 14 ++----- tools/ttmlir-opt/ttmlir-opt.cpp | 7 +--- 20 files changed, 77 insertions(+), 114 deletions(-) create mode 100644 .gitmodules diff --git a/.gitmodules b/.gitmodules new file mode 100644 index 000000000..7c1905a0b --- /dev/null +++ b/.gitmodules @@ -0,0 +1,3 @@ +[submodule "third_party/stablehlo"] + path = third_party/stablehlo + url = https://github.com/openxla/stablehlo diff --git a/CMakeLists.txt b/CMakeLists.txt index d84888e97..169067144 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -44,17 +44,21 @@ set(TTMLIR_TOOLCHAIN_DIR $ENV{TTMLIR_TOOLCHAIN_DIR}) set(TTMLIR_SOURCE_DIR ${PROJECT_SOURCE_DIR}) set(TTMLIR_BINARY_DIR ${PROJECT_BINARY_DIR}) set(LLVM_LIT_TOOLS_DIR "${TTMLIR_TOOLCHAIN_DIR}/src/llvm-project/llvm/utils/lit") +include_directories(${PROJECT_SOURCE_DIR}/include) include_directories(SYSTEM ${LLVM_INCLUDE_DIRS}) include_directories(SYSTEM ${MLIR_INCLUDE_DIRS}) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) include_directories(${TTMLIR_SOURCE_DIR}/include) include_directories(${TTMLIR_BINARY_DIR}/include) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/flatbuffers/include) link_directories(${LLVM_BUILD_LIBRARY_DIR}) add_definitions(${LLVM_DEFINITIONS}) include(TTMLIRPythonSitePackages) +if (TTMLIR_ENABLE_STABLEHLO) + set(STABLEHLO_BUILD_EMBEDDED ON) + add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/third_party/stablehlo ${CMAKE_CURRENT_BINARY_DIR}/stablehlo EXCLUDE_FROM_ALL) + include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/stablehlo) +endif() + add_subdirectory(third_party) add_subdirectory(include) add_subdirectory(lib) diff --git a/env/CMakeLists.txt b/env/CMakeLists.txt index a8d1607cd..f19b60ee4 100644 --- a/env/CMakeLists.txt +++ b/env/CMakeLists.txt @@ -67,24 +67,4 @@ ExternalProject_Add( DEPENDS python-venv ) -set(STABLE_HLO_VERSION "v1.5.0") - -ExternalProject_Add( - stablehlo - PREFIX ${TTMLIR_TOOLCHAIN_DIR} - CMAKE_GENERATOR Ninja - CMAKE_ARGS - -DCMAKE_BUILD_TYPE=Release - -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} - -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_INSTALL_PREFIX=${TTMLIR_TOOLCHAIN_DIR} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - # TODO: Figure out if we need this (or benefit from it) when it comes to linking stablehlo. - #-DSTABLEHLO_BUILD_EMBEDDED=ON - -DMLIR_DIR=${TTMLIR_TOOLCHAIN_DIR}/lib/cmake/mlir - GIT_REPOSITORY https://github.com/openxla/stablehlo - GIT_TAG ${STABLE_HLO_VERSION} - GIT_PROGRESS ON -) - add_custom_target(llvm-lit ALL COMMAND cp ${TTMLIR_TOOLCHAIN_DIR}/src/llvm-project-build/bin/llvm-lit ${TTMLIR_TOOLCHAIN_DIR}/bin/llvm-lit DEPENDS llvm-project) diff --git a/include/CMakeLists.txt b/include/CMakeLists.txt index 6e82d9606..2b51d0d34 100644 --- a/include/CMakeLists.txt +++ b/include/CMakeLists.txt @@ -1,3 +1 @@ add_subdirectory(ttmlir) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) diff --git a/include/ttmlir/Conversion/CMakeLists.txt b/include/ttmlir/Conversion/CMakeLists.txt index 120097e29..891fa5608 100644 --- a/include/ttmlir/Conversion/CMakeLists.txt +++ b/include/ttmlir/Conversion/CMakeLists.txt @@ -1,3 +1,11 @@ +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) +include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) +include_directories(${TTMLIR_SOURCE_DIR}/include) + set(LLVM_TARGET_DEFINITIONS Passes.td) +if (TTMLIR_ENABLE_STABLEHLO) +mlir_tablegen(Passes.h.inc -gen-pass-decls -name TTMLIRConversion -DTTMLIR_ENABLE_STABLEHLO) +else() mlir_tablegen(Passes.h.inc -gen-pass-decls -name TTMLIRConversion) +endif() add_public_tablegen_target(TTMLIRConversionPassIncGen) diff --git a/include/ttmlir/Conversion/Passes.h b/include/ttmlir/Conversion/Passes.h index cf7958ba1..b12e9ebb8 100644 --- a/include/ttmlir/Conversion/Passes.h +++ b/include/ttmlir/Conversion/Passes.h @@ -8,7 +8,6 @@ #ifdef TTMLIR_ENABLE_STABLEHLO #include "ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h" #endif - #include "ttmlir/Conversion/TTIRToTTNN/TTIRToTTNN.h" #include "ttmlir/Conversion/TTNNToEmitC/TTNNToEmitC.h" #include "ttmlir/Conversion/TosaToTTIR/TosaToTTIR.h" diff --git a/include/ttmlir/Conversion/Passes.td b/include/ttmlir/Conversion/Passes.td index 2b7d04162..4925e61cc 100644 --- a/include/ttmlir/Conversion/Passes.td +++ b/include/ttmlir/Conversion/Passes.td @@ -7,12 +7,14 @@ include "mlir/Pass/PassBase.td" +#ifdef TTMLIR_ENABLE_STABLEHLO def ConvertStableHLOToTTIR : Pass<"convert-stablehlo-to-ttir", "::mlir::ModuleOp"> { - let summary = "Convert StableHLO dialect to TTIR dialect."; +let summary = "Convert StableHLO dialect to TTIR dialect."; let constructor = "createConvertStableHLOToTTIRPass()"; // TODO(mrakita): Probably will need to add some include here for StableHLO dialect. let dependentDialects = ["mlir::stablehlo::StablehloDialect", "mlir::tt::ttir::TTIRDialect"]; } +#endif def ConvertTosaToTTIR : Pass<"convert-tosa-to-ttir", "::mlir::ModuleOp"> { let summary = "Convert TOSA dialect to TTIR dialect."; diff --git a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h index 1df166900..24fa4a30c 100644 --- a/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h +++ b/include/ttmlir/Conversion/StableHLOToTTIR/StableHLOToTTIR.h @@ -5,19 +5,14 @@ #ifndef TTMLIR_CONVERSION_StableHLOToTTIR_StableHLOToTTIR_H #define TTMLIR_CONVERSION_StableHLOToTTIR_StableHLOToTTIR_H + #include "mlir/IR/BuiltinOps.h" #include "mlir/Pass/Pass.h" namespace mlir::tt { #ifdef TTMLIR_ENABLE_STABLEHLO -std::unique_ptr> createConvertStableHLOToTTIRPass(); -#else -// define an empty function in case conversion definition does not exist. -std::unique_ptr> createConvertStableHLOToTTIRPass() -{ -return std::unique_ptr>(nullptr); -} + std::unique_ptr> createConvertStableHLOToTTIRPass(); #endif } // namespace mlir::tt diff --git a/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td b/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td index 3e257a9fb..1fce627d6 100644 --- a/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td +++ b/include/ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.td @@ -31,8 +31,8 @@ def TTKernel_CB : TTKernel_Type<"CB", "cb"> { static CBType get(::mlir::MLIRContext *context, uint64_t address, uint64_t port, - MemRefType memref, - uint64_t numBuffers = 1) { + MemRefType memref) { + uint64_t numBuffers = 1; uint64_t pageSize = 0; if (::mlir::isa<::mlir::tt::TileType>(memref.getElementType())) { pageSize = ::mlir::cast<::mlir::tt::TileType>(memref.getElementType()).getSizeBytes(); diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index e0bf4f0df..4126bc41f 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -1,6 +1,3 @@ -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) - add_subdirectory(CAPI) add_subdirectory(Conversion) add_subdirectory(Dialect) @@ -8,33 +5,22 @@ add_subdirectory(Target) # Shared library will include runtime code # so we only build it if runtime is enabled - -set(ttmlir_dependencies "") -set(link_libs -MLIR; -MLIRTTDialect; -MLIRTTIRDialect; -MLIRTTIRTransforms; -TTMLIRConversions; -MLIRTTIRAnalysis; -MLIRTTNNDialect; -MLIRTTNNTransforms; -MLIRTTKernelDialect; -MLIRTTMetalDialect; -MLIRTTMetalTransforms; -MLIRTTNNPipelines -) - -if (TTMLIR_ENABLE_STABLEHLO) - list(APPEND link_libs TTMLIRStableHLOToTTIR) +if (TTMLIR_ENABLE_RUNTIME) + add_subdirectory(SharedLib) endif() add_mlir_library(TTMLIRStatic STATIC RegisterAll.cpp - - DEPENDS - ${ttmlir_dependencies} - LINK_LIBS PUBLIC - ${link_libs} + MLIR + MLIRTTDialect + MLIRTTIRDialect + MLIRTTIRTransforms + TTMLIRConversions + MLIRTTIRAnalysis + MLIRTTNNDialect + MLIRTTNNTransforms + MLIRTTKernelDialect + MLIRTTMetalDialect + MLIRTTMetalTransforms + MLIRTTNNPipelines ) - diff --git a/lib/Conversion/CMakeLists.txt b/lib/Conversion/CMakeLists.txt index 0a8298967..3a4c6523d 100644 --- a/lib/Conversion/CMakeLists.txt +++ b/lib/Conversion/CMakeLists.txt @@ -1,15 +1,13 @@ add_library(TTMLIRConversions INTERFACE) +add_subdirectory(TosaToTTIR) +add_subdirectory(TTNNToEmitC) +add_subdirectory(TTIRToTTNN) if (TTMLIR_ENABLE_STABLEHLO) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) -include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) add_subdirectory(StableHLOToTTIR) -add_dependencies(TTMLIRConversions stablehlo) endif() -add_subdirectory(TosaToTTIR) -add_subdirectory(TTNNToEmitC) -add_subdirectory(TTIRToTTNN) +include_directories(${TTMLIR_SOURCE_DIR}/include) set(link_libs TTMLIRTosaToTTIR; @@ -22,5 +20,5 @@ list(APPEND link_libs TTMLIRStableHLOToTTIR) endif() target_link_libraries(TTMLIRConversions INTERFACE - ${link_libs} +${link_libs} ) diff --git a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt index 8cd88059d..18281ccf3 100644 --- a/lib/Conversion/StableHLOToTTIR/CMakeLists.txt +++ b/lib/Conversion/StableHLOToTTIR/CMakeLists.txt @@ -1,7 +1,9 @@ include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo) include_directories(${TTMLIR_TOOLCHAIN_DIR}/src/stablehlo-build) +include_directories(${TTMLIR_SOURCE_DIR}/include) +include_directories(${PROJECT_SOURCE_DIR}/include) +include_directories(${PROJECT_SOURCE_DIR}/build/stablehlo) -get_property(STABLEHLO_LIBS GLOBAL PROPERTY STABLEHLO_LIBS) add_mlir_library(TTMLIRStableHLOToTTIR StableHLOToTTIR.cpp @@ -12,6 +14,5 @@ add_mlir_library(TTMLIRStableHLOToTTIR TTMLIRConversionPassIncGen LINK_LIBS PUBLIC - MLIRTransformUtils - ${STABLEHLO_LIBS} + MLIR ) diff --git a/lib/RegisterAll.cpp b/lib/RegisterAll.cpp index 38025d6ef..98faeaa84 100644 --- a/lib/RegisterAll.cpp +++ b/lib/RegisterAll.cpp @@ -17,7 +17,10 @@ #include "ttmlir/Dialect/TTNN/IR/TTNN.h" #include "ttmlir/Dialect/TTNN/Pipelines/Passes.h" #include "ttmlir/Dialect/TTNN/Transforms/Passes.h" +#ifdef TTMLIR_ENABLE_STABLEHLO #include "stablehlo/dialect/Register.h" +#endif + void mlir::tt::registerAllDialects(mlir::DialectRegistry ®istry) { registry .insert(); - mlir::stablehlo::registerAllDialects(registry); +#if TTMLIR_ENABLE_STABLEHLO + mlir::stablehlo::registerAllDialects(registry); +#endif } void mlir::tt::registerAllPasses() { // Register all dialect conversion passes - + // mlir::tt::registerTTMLIRConversionPasses(); mlir::tt::ttir::registerPasses(); diff --git a/lib/SharedLib/CMakeLists.txt b/lib/SharedLib/CMakeLists.txt index ceaf708c4..81eeb5ae5 100644 --- a/lib/SharedLib/CMakeLists.txt +++ b/lib/SharedLib/CMakeLists.txt @@ -20,10 +20,6 @@ set(TTMLIR_LIBS TTMLIRTTNNToEmitC ) -if (TTMLIR_ENABLE_STABLEHLO) -list(APPEND TTMLIR_LIBS TTMLIRStableHLOToTTIR) -endif() - # We supply empty.cpp because CMake does not allow creating a library without sources. add_library(TTMLIR SHARED empty.cpp) @@ -31,7 +27,6 @@ add_dependencies(TTMLIR ${TTMLIR_LIBS} ${TTNN_RUNTIME_LIBS} ${TTNN_LIBS} - ${STABLEHLO_LIBS} ) target_link_libraries(TTMLIR PRIVATE @@ -46,7 +41,7 @@ target_link_libraries(TTMLIR PRIVATE -Wl,--no-whole-archive ${TTNN_LIBS} - ${STABLEHLO_LIBS} + flatbuffers ) diff --git a/runtime/lib/binary.cpp b/runtime/lib/binary.cpp index 498c245b3..db84a8e23 100644 --- a/runtime/lib/binary.cpp +++ b/runtime/lib/binary.cpp @@ -4,7 +4,6 @@ #include -#pragma clang diagnostic ignored "-Wcovered-switch-default" #include "flatbuffers/idl.h" #include "tt/runtime/types.h" diff --git a/test/ttmlir/ttmlir-opt.mlir b/test/ttmlir/ttmlir-opt.mlir index eddf4d1d8..5a0c871f5 100644 --- a/test/ttmlir/ttmlir-opt.mlir +++ b/test/ttmlir/ttmlir-opt.mlir @@ -1,3 +1,18 @@ // RUN: ttmlir-opt --show-dialects | FileCheck %s // CHECK: Available Dialects: -// CHECK-SAME: arith,builtin,cf,emitc,func,linalg,ml_program,scf,tensor,tosa,tt,ttir,ttkernel,ttmetal,ttnn,vector +// CHECK: arith +// CHECK: builtin +// CHECK: cf +// CHECK: emitc +// CHECK: func +// CHECK: linalg +// CHECK: ml_program +// CHECK: scf +// CHECK: tensor +// CHECK: tosa +// CHECK: tt +// CHECK: ttir +// CHECK: ttkernel +// CHECK: ttmetal +// CHECK: ttnn +// CHECK: vector diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index bb1ff8729..9d6584738 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -13,13 +13,6 @@ else() message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") endif() -if (TT_RUNTIME_ENABLE_PERF_TRACE) - add_compile_definitions(TRACY_ENABLE) - set(ENV{ENABLE_TRACY} "1") -else() - set(ENV{ENABLE_TRACY} "0") -endif() - set(TTMETAL_INCLUDE_DIRS ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated @@ -53,7 +46,7 @@ ExternalProject_Add( -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - -DTRACY_ENABLE=${TT_RUNTIME_ENABLE_PERF_TRACE} + -DENABLE_TRACY=${TT_RUNTIME_ENABLE_PERF_TRACE} GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 GIT_PROGRESS ON diff --git a/tools/CMakeLists.txt b/tools/CMakeLists.txt index 3717682b5..3899b3775 100644 --- a/tools/CMakeLists.txt +++ b/tools/CMakeLists.txt @@ -1,7 +1,2 @@ -include(ExternalProject) - -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo) -include_directories(${PROJECT_SOURCE_DIR}/third_party/stablehlo/src/stablehlo-build) - add_subdirectory(ttmlir-opt) add_subdirectory(ttmlir-translate) diff --git a/tools/ttmlir-opt/CMakeLists.txt b/tools/ttmlir-opt/CMakeLists.txt index 7e3284bd5..ef8810f9e 100644 --- a/tools/ttmlir-opt/CMakeLists.txt +++ b/tools/ttmlir-opt/CMakeLists.txt @@ -1,17 +1,9 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) -set(LIBS ${dialect_libs} ${conversion_libs} StablehloBase TTNNTargetFlatbuffer MLIRTTIRDialect TTMLIRConversions MLIRTransformUtils MLIROptLib MLIRTargetCpp TTMLIRStatic CheckOps ChloCAPI ChloOps InterpreterOps StablehloAssemblyFormat StablehloBase StablehloBroadcastUtils StablehloCAPI StablehloLinalgTransforms StablehloOps StablehloPasses StablehloPortableApi StablehloRegister StablehloSerialization StablehloTOSATransforms StablehloTestUtils StablehloTypeInference Version VhloCAPI VhloOps VhloTypes) - -if (TTMLIR_ENABLE_STABLEHLO) - list(APPEND LIBS TTMLIRStableHLOToTTIR) -endif() - +set(LIBS ${dialect_libs} ${conversion_libs} MLIROptLib MLIRTargetCpp TTMLIRStatic) add_llvm_executable(ttmlir-opt ttmlir-opt.cpp) -target_link_directories(ttmlir-opt PRIVATE ${TTMLIR_PROJECT_DIR}/build/lib) -target_link_libraries(ttmlir-opt PRIVATE ${LIBS} libStablehloBase.a) -target_link_directories(ttmlir-opt PRIVATE ${TTMLIR_TOOLCHAIN_DIR}/lib) - llvm_update_compile_flags(ttmlir-opt) +target_link_libraries(ttmlir-opt PRIVATE ${LIBS}) -mlir_check_all_link_libraries(ttmlir-opt ${LIBS}) +mlir_check_all_link_libraries(ttmlir-opt) diff --git a/tools/ttmlir-opt/ttmlir-opt.cpp b/tools/ttmlir-opt/ttmlir-opt.cpp index 34e5f3cad..d4f65dc2e 100644 --- a/tools/ttmlir-opt/ttmlir-opt.cpp +++ b/tools/ttmlir-opt/ttmlir-opt.cpp @@ -7,9 +7,6 @@ #include "mlir/InitAllPasses.h" #include "mlir/Support/FileUtilities.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" -#ifdef TTMLIR_ENABLE_STABLEHLO -#include "stablehlo/dialect/Register.h" -#endif #include "ttmlir/RegisterAll.h" @@ -19,9 +16,7 @@ int main(int argc, char **argv) { mlir::DialectRegistry registry; mlir::tt::registerAllDialects(registry); -#ifdef TTMLIR_ENABLE_STABLEHLO - mlir::stablehlo::registerAllDialects(registry); -#endif + return mlir::asMainReturnCode( mlir::MlirOptMain(argc, argv, "ttmlir optimizer driver\n", registry)); } From 28a64828b2f6fbfa00d1edfc4c683b7ba8485b9b Mon Sep 17 00:00:00 2001 From: uazizTT Date: Mon, 19 Aug 2024 16:57:40 -0400 Subject: [PATCH 11/11] Change git config. --- .github/workflows/docker-build.yml | 1 - .gitmodules | 3 -- third_party/CMakeLists.txt | 65 ++---------------------------- third_party/stablehlo.cmake | 16 ++++++++ third_party/tt-metal.cmake | 65 ++++++++++++++++++++++++++++++ 5 files changed, 84 insertions(+), 66 deletions(-) delete mode 100644 .gitmodules create mode 100644 third_party/stablehlo.cmake create mode 100644 third_party/tt-metal.cmake diff --git a/.github/workflows/docker-build.yml b/.github/workflows/docker-build.yml index 8bb1930b9..fb1724c97 100644 --- a/.github/workflows/docker-build.yml +++ b/.github/workflows/docker-build.yml @@ -55,7 +55,6 @@ jobs: -DCMAKE_BUILD_TYPE=${{ matrix.build.build_type }} \ -DTTMLIR_ENABLE_RUNTIME=${{ matrix.build.enable_runtime }} \ -DTTMLIR_ENABLE_RUNTIME_TESTS=${{ matrix.build.enable_runtime }} \ - -DTTMLIR_ENABLE_STABLEHLO=ON \ -S ${{ steps.strings.outputs.work-dir }} - name: Build diff --git a/.gitmodules b/.gitmodules deleted file mode 100644 index 7c1905a0b..000000000 --- a/.gitmodules +++ /dev/null @@ -1,3 +0,0 @@ -[submodule "third_party/stablehlo"] - path = third_party/stablehlo - url = https://github.com/openxla/stablehlo diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index 9d6584738..03a78615a 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -1,65 +1,6 @@ include(ExternalProject) -if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") - set(ARCH_NAME "grayskull") - set(ARCH_EXTRA_DIR "grayskull") -elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") - set(ARCH_NAME "wormhole") - set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") -elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") - set(ARCH_NAME "blackhole") - set(ARCH_EXTRA_DIR "blackhole") -else() - message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") +include (tt-metal.cmake) +if (TTMLIR_ENABLE_STABLEHLO) + include(stablehlo.cmake) endif() - -set(TTMETAL_INCLUDE_DIRS - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager - ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f - PARENT_SCOPE -) - -set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) -set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) - -set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) -set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) -set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) - -ExternalProject_Add( - tt-metal - PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal - CMAKE_GENERATOR Ninja - CMAKE_ARGS - -DCMAKE_BUILD_TYPE=Release - -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} - -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} - -DENABLE_TRACY=${TT_RUNTIME_ENABLE_PERF_TRACE} - GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git - GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 - GIT_PROGRESS ON - BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} -) - -set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) - -list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) -list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) - -foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) - add_library(${lib_name} SHARED IMPORTED GLOBAL) - set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) - add_dependencies(${lib_name} tt-metal) -endforeach() diff --git a/third_party/stablehlo.cmake b/third_party/stablehlo.cmake new file mode 100644 index 000000000..a6479c84b --- /dev/null +++ b/third_party/stablehlo.cmake @@ -0,0 +1,16 @@ +include(ExternalProject) + +ExternalProject_Add( + tt-metal + PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + -DENABLE_TRACY=${TT_RUNTIME_ENABLE_PERF_TRACE} + GIT_REPOSITORY https://github.com/openxla/stablehlo.git + GIT_TAG v1.5.0 + GIT_PROGRESS ON +) diff --git a/third_party/tt-metal.cmake b/third_party/tt-metal.cmake new file mode 100644 index 000000000..9d6584738 --- /dev/null +++ b/third_party/tt-metal.cmake @@ -0,0 +1,65 @@ +include(ExternalProject) + +if ("$ENV{ARCH_NAME}" STREQUAL "grayskull") + set(ARCH_NAME "grayskull") + set(ARCH_EXTRA_DIR "grayskull") +elseif ("$ENV{ARCH_NAME}" STREQUAL "wormhole_b0") + set(ARCH_NAME "wormhole") + set(ARCH_EXTRA_DIR "wormhole/wormhole_b0_defines") +elseif ("$ENV{ARCH_NAME}" STREQUAL "blackhole") + set(ARCH_NAME "blackhole") + set(ARCH_EXTRA_DIR "blackhole") +else() + message(FATAL_ERROR "Unsupported ARCH_NAME: $ENV{ARCH_NAME}") +endif() + +set(TTMETAL_INCLUDE_DIRS + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/deprecated + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/fmt + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f + PARENT_SCOPE +) + +set(TTMETAL_LIBRARY_DIR ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/lib) +set(TTNN_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/_ttnn.so) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_DIR}/libtt_metal.so) + +set(TTMETAL_LIBRARY_DIR ${TTMETAL_LIBRARY_DIR} PARENT_SCOPE) +set(TTNN_LIBRARY_PATH ${TTNN_LIBRARY_PATH} PARENT_SCOPE) +set(TTMETAL_LIBRARY_PATH ${TTMETAL_LIBRARY_PATH} PARENT_SCOPE) + +ExternalProject_Add( + tt-metal + PREFIX ${TTMLIR_SOURCE_DIR}/third_party/tt-metal + CMAKE_GENERATOR Ninja + CMAKE_ARGS + -DCMAKE_BUILD_TYPE=Release + -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} + -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} + -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER} + -DENABLE_TRACY=${TT_RUNTIME_ENABLE_PERF_TRACE} + GIT_REPOSITORY https://github.com/tenstorrent/tt-metal.git + GIT_TAG f6a2e5cb2b857bf4c72401bea68adf98c25bbe47 + GIT_PROGRESS ON + BUILD_BYPRODUCTS ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH} +) + +set_target_properties(tt-metal PROPERTIES EXCLUDE_FROM_ALL TRUE) + +list(APPEND library_names TTNN_LIBRARY TTMETAL_LIBRARY) +list(APPEND library_paths ${TTNN_LIBRARY_PATH} ${TTMETAL_LIBRARY_PATH}) + +foreach(lib_name lib_path IN ZIP_LISTS library_names library_paths) + add_library(${lib_name} SHARED IMPORTED GLOBAL) + set_target_properties(${lib_name} PROPERTIES EXCLUDE_FROM_ALL TRUE IMPORTED_LOCATION ${lib_path}) + add_dependencies(${lib_name} tt-metal) +endforeach()