-
Notifications
You must be signed in to change notification settings - Fork 13
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Add CMake target for EmitC TTNN dylibs (#1241)
* Add cmake target for emitc ttnn dylibs * tiny adjustments to c++ examples
- Loading branch information
1 parent
6dd5eba
commit edc964b
Showing
6 changed files
with
133 additions
and
25 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,16 +1,36 @@ | ||
| # WIP | ||
| ## Table of contents | ||
|
|
||
| - [TTNN Standalone](#ttnn-standalone) | ||
| - [Usage](#usage) | ||
| - [TTNN Dylib](#ttnn-dylib) | ||
|
|
||
| ## TTNN Standalone | ||
|
|
||
| TTNN Standalone is a post-compile tuning tool. | ||
|
|
||
| Third party ML models (PyTorch, Jax, ONNX, ...) can be compiled to a set of TTNN library calls in C++. This generated code can then be manually fine-tuned outside of the compiler environment. TTNN Standalone tool offers all the scaffolding needed to run the C++ code on device (build & run scripts). | ||
|
|
||
| ### Usage | ||
|
|
||
| Following script compiles and runs the ttnn standalone | ||
| ```bash | ||
| cd tools/ttnn-standalone | ||
| # Compile a model to C++ code | ||
| ./build/bin/ttmlir-opt --ttir-load-system-desc --ttir-implicit-device --ttir-layout --convert-ttir-to-ttnn --ttnn-decompose-layouts --ttnn-deallocate --convert-ttnn-to-emitc test/ttmlir/Silicon/TTNN/emitc/simple_add.mlir | ./build/bin/ttmlir-translate --mlir-to-cpp -allow-unregistered-dialect | ||
|
|
||
| # Copy paste the generated function into `ttnn-standalone.cpp`. | ||
|
|
||
| # Adapt the `main()` function in `ttnn-standalone.cpp` to feed tensors needed for the model | ||
|
|
||
| # Run the following script from within this folder (`tools/ttnn-standalone`) to compile and run the ttnn standalone: | ||
|
|
||
| ./run | ||
| ``` | ||
|
|
||
| Note: if you receive this error | ||
| ```bash | ||
| -bash: ./run: Permission denied | ||
| ``` | ||
| you may run | ||
| ```bash | ||
| chmod +x run | ||
| ``` | ||
| running `chmod +x run` will allow the execution of the script. | ||
|
|
||
| ## TTNN Dylib | ||
|
|
||
| Similarly to the Standalone, this tool offers the ability to compile third party ML models, but to dylibs. Initial intent for compiled dylibs is to be used in testing infrastructure, but sky's the limit :) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,45 @@ | ||
| // SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC | ||
| // | ||
| // SPDX-License-Identifier: Apache-2.0 | ||
|
|
||
| #include "ttnn-dylib.hpp" | ||
|
|
||
| // Forward function example | ||
| // | ||
| std::vector<ttnn::Tensor> forward(std::vector<ttnn::Tensor> inputs) { | ||
| ttnn::Tensor v1 = inputs[0]; | ||
| ttnn::Tensor v2 = inputs[1]; | ||
| ttnn::Device *v3 = ttnn::DeviceGetter::getInstance(); | ||
| ttnn::MemoryConfig v4 = ttnn::MemoryConfig( | ||
| ttnn::TensorMemoryLayout::INTERLEAVED, ttnn::BufferType::DRAM); | ||
| ttnn::Tensor v5 = ttnn::to_device(v1, v3, v4); | ||
| ttnn::Tensor v6 = | ||
| ttnn::to_layout(v5, ttnn::Layout::TILE, std::nullopt, std::nullopt, | ||
| static_cast<::ttnn::Device *>(nullptr)); | ||
| ttnn::deallocate(v5, false); | ||
| ttnn::MemoryConfig v7 = ttnn::MemoryConfig( | ||
| ttnn::TensorMemoryLayout::INTERLEAVED, ttnn::BufferType::DRAM); | ||
| ttnn::Tensor v8 = ttnn::to_device(v2, v3, v7); | ||
| ttnn::Tensor v9 = | ||
| ttnn::to_layout(v8, ttnn::Layout::TILE, std::nullopt, std::nullopt, | ||
| static_cast<::ttnn::Device *>(nullptr)); | ||
| ttnn::deallocate(v8, false); | ||
| ttnn::Shape v10 = ttnn::Shape(tt::tt_metal::LegacyShape({ | ||
| 32, | ||
| 32, | ||
| })); | ||
| ttnn::MemoryConfig v11 = ttnn::MemoryConfig( | ||
| ttnn::TensorMemoryLayout::INTERLEAVED, ttnn::BufferType::DRAM); | ||
| ttnn::Tensor v12 = | ||
| ttnn::empty(v10, ttnn::DataType::BFLOAT16, ttnn::Layout::TILE, v3, v11); | ||
| ttnn::Tensor v13 = ttnn::add(v6, v9, std::nullopt, std::nullopt, v12); | ||
| ttnn::deallocate(v9, false); | ||
| ttnn::deallocate(v6, false); | ||
| ttnn::Tensor v14 = ttnn::from_device(v13); | ||
| ttnn::deallocate(v12, false); | ||
| ttnn::Tensor v15 = | ||
| ttnn::to_layout(v14, ttnn::Layout::ROW_MAJOR, std::nullopt, std::nullopt, | ||
| static_cast<::ttnn::Device *>(nullptr)); | ||
| ttnn::deallocate(v14, false); | ||
| return std::vector<ttnn::Tensor>{v15}; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,7 @@ | ||
| // SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC | ||
| // | ||
| // SPDX-License-Identifier: Apache-2.0 | ||
|
|
||
| #include "ttnn-precompiled.hpp" | ||
|
|
||
| std::vector<ttnn::Tensor> forward(std::vector<ttnn::Tensor> inputs); |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters