From 83d2644c3fd43ae8bd78e06c29e441de44fa9a55 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:22:09 +0000 Subject: [PATCH 01/11] added tests for sum and mean --- .../ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 38 ++++++++++ .../ttml/ttnn_fixed/trivial_ttnn_ops.hpp | 7 ++ tt-train/tests/ttnn_fixed/reduce_ops_test.cpp | 73 +++++++++++++++++++ 3 files changed, 118 insertions(+) create mode 100644 tt-train/tests/ttnn_fixed/reduce_ops_test.cpp diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index 652073c583d..b15f7142ff3 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -7,6 +7,7 @@ #include #include +#include "autograd/auto_context.hpp" #include "core/compute_kernel_config.hpp" #include "core/tt_tensor_utils.hpp" @@ -54,4 +55,41 @@ tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::T return ttnn::multiply(a, inv_b); } +tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { + // auto tensor_shape = t.get_shape(); + // auto shape = core::create_shape({tensor_shape[0], tensor_shape[1], tensor_shape[2], 1}); + // auto* device = &autograd::ctx().get_device(); + // auto mean = core::empty(shape, device, t.memory_config()); + auto res = ttnn::moreh_mean( + t, + dim, + keep_dim, + std::nullopt, + std::nullopt, + std::nullopt, + /* device_compute_kernel_config */ core::ComputeKernelConfig::precise()); + return res; +} +tt::tt_metal ::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { + return ttnn::mean(t, dim, keep_dim, std::nullopt, core::ComputeKernelConfig::precise()); +} + +tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { + // auto tensor_shape = t.get_shape(); + // auto shape = core::create_shape({tensor_shape[0], tensor_shape[1], tensor_shape[2], 1}); + // auto* device = &autograd::ctx().get_device(); + // auto mean = core::empty(shape, device, t.memory_config()); + auto res = ttnn::moreh_sum( + t, + dim, + keep_dim, + std::nullopt, + std::nullopt, + /* device_compute_kernel_config */ core::ComputeKernelConfig::precise()); + return res; +} +tt::tt_metal ::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { + return ttnn::sum(t, dim, keep_dim, std::nullopt, core::ComputeKernelConfig::precise()); +} + } // namespace ttml::ttnn_fixed diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp index 564c985c198..f06b5d2bd5a 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp @@ -5,6 +5,8 @@ #pragma once #include +#include "core/tt_tensor_utils.hpp" + namespace ttml::ttnn_fixed { tt::tt_metal::Tensor sum_over_dim(const tt::tt_metal::Tensor& t, uint32_t dim); @@ -13,4 +15,9 @@ tt::tt_metal::Tensor log_softmax(const tt::tt_metal::Tensor& t, int dim); tt::tt_metal::Tensor softmax(const tt::tt_metal::Tensor& t, int dim); tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::Tensor& b); +tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal ::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); + +tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal ::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); } // namespace ttml::ttnn_fixed diff --git a/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp b/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp new file mode 100644 index 00000000000..c6b26c6aebf --- /dev/null +++ b/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp @@ -0,0 +1,73 @@ +// SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC +// +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include +#include +#include +#include + +#include "autograd/auto_context.hpp" +#include "core/compute_kernel_config.hpp" +#include "core/device.hpp" +#include "core/tt_tensor_utils.hpp" +#include "ttnn_fixed/trivial_ttnn_ops.hpp" + +class ReduceOpTest : public ::testing::Test { +protected: + void SetUp() override { + ttml::autograd::ctx().open_device(); + } + + void TearDown() override { + ttml::autograd::ctx().close_device(); + } +}; + +TEST_F(ReduceOpTest, TestMean) { + xt::random::seed(42); + auto* device = &ttml::autograd::ctx().get_device(); + xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.5, 0.5).reshape({2, 1, 64, 64}); + + auto xtensor_a_tensor = ttml::core::from_xtensor(xtensor_a, device); + + auto ttnn_mean_dim3 = ttml::ttnn_fixed::mean_ttnn(xtensor_a_tensor, 3, true); + auto moreh_mean_dim3 = ttml::ttnn_fixed::mean_moreh(xtensor_a_tensor, 3, true); + + xt::xarray mean_xtensor = xt::mean(xtensor_a, {3}, xt::evaluation_strategy::immediate); + mean_xtensor.reshape({2, 1, 64, 1}); + + auto mean_ttnn = ttml::core::to_xtensor(ttnn_mean_dim3); + auto mean_moreh = ttml::core::to_xtensor(moreh_mean_dim3); + + EXPECT_TRUE(xt::allclose(mean_ttnn, mean_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_ttnn, /*rtol=*/1e-3, /*atol=*/1e-2)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_moreh, /*rtol=*/1e-3, /*atol=*/1e-2)); +} + +TEST_F(ReduceOpTest, TestSum) { + xt::random::seed(42); + auto* device = &ttml::autograd::ctx().get_device(); + xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.1, 0.1).reshape({2, 1, 64, 64}); + + auto xtensor_a_tensor = ttml::core::from_xtensor(xtensor_a, device); + + auto ttnn_sum_dim3 = ttml::ttnn_fixed::sum_ttnn(xtensor_a_tensor, 3, true); + auto moreh_sum_dim3 = ttml::ttnn_fixed::sum_moreh(xtensor_a_tensor, 3, true); + + xt::xarray sum_xtensor = xt::sum(xtensor_a, {3}, xt::evaluation_strategy::immediate); + sum_xtensor.reshape({2, 1, 64, 1}); + + auto sum_ttnn = ttml::core::to_xtensor(ttnn_sum_dim3); + auto sum_moreh = ttml::core::to_xtensor(moreh_sum_dim3); + std::cout << sum_ttnn << std::endl; + std::cout << "------------" << std::endl; + std::cout << sum_xtensor << std::endl; + + EXPECT_TRUE(xt::allclose(sum_ttnn, sum_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); + EXPECT_TRUE(xt::allclose(sum_xtensor, sum_ttnn, /*rtol=*/1e-2, /*atol=*/1e-2)); + EXPECT_TRUE(xt::allclose(sum_xtensor, sum_moreh, /*rtol=*/1e-2, /*atol=*/1e-2)); +} From 28398f6c8d586f3c803c25b48c2354927875f167 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 21:58:34 +0000 Subject: [PATCH 02/11] udated tests --- .../ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 16 +---- tt-train/tests/ttnn_fixed/reduce_ops_test.cpp | 70 +++++++++++++++++-- 2 files changed, 66 insertions(+), 20 deletions(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index b15f7142ff3..fd9cb168b6b 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -14,13 +14,7 @@ namespace ttml::ttnn_fixed { tt::tt_metal::Tensor sum_over_dim(const tt::tt_metal::Tensor& t, uint32_t dim) { - return ttnn::moreh_sum( - t, - /* dim */ dim, - /* keep_dim */ true, - /* output */ std::nullopt, - /* output_mem_config */ std::nullopt, - /*compute_kernel_config */ core::ComputeKernelConfig::precise()); + return sum_ttnn(t, dim, true); } tt::tt_metal::Tensor sum_over_batch(const tt::tt_metal::Tensor& t) { @@ -56,10 +50,6 @@ tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::T } tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { - // auto tensor_shape = t.get_shape(); - // auto shape = core::create_shape({tensor_shape[0], tensor_shape[1], tensor_shape[2], 1}); - // auto* device = &autograd::ctx().get_device(); - // auto mean = core::empty(shape, device, t.memory_config()); auto res = ttnn::moreh_mean( t, dim, @@ -75,10 +65,6 @@ tt::tt_metal ::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool kee } tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { - // auto tensor_shape = t.get_shape(); - // auto shape = core::create_shape({tensor_shape[0], tensor_shape[1], tensor_shape[2], 1}); - // auto* device = &autograd::ctx().get_device(); - // auto mean = core::empty(shape, device, t.memory_config()); auto res = ttnn::moreh_sum( t, dim, diff --git a/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp b/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp index c6b26c6aebf..9c15b982e82 100644 --- a/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp +++ b/tt-train/tests/ttnn_fixed/reduce_ops_test.cpp @@ -27,7 +27,49 @@ class ReduceOpTest : public ::testing::Test { } }; -TEST_F(ReduceOpTest, TestMean) { +TEST_F(ReduceOpTest, TestMeanDim0) { + xt::random::seed(42); + auto* device = &ttml::autograd::ctx().get_device(); + xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.5, 0.5).reshape({2, 1, 64, 64}); + + auto xtensor_a_tensor = ttml::core::from_xtensor(xtensor_a, device); + + auto ttnn_mean_dim0 = ttml::ttnn_fixed::mean_ttnn(xtensor_a_tensor, 0, true); + auto moreh_mean_dim0 = ttml::ttnn_fixed::mean_moreh(xtensor_a_tensor, 0, true); + + xt::xarray mean_xtensor = xt::mean(xtensor_a, {0}, xt::evaluation_strategy::immediate); + mean_xtensor.reshape({1, 1, 64, 64}); + + auto mean_ttnn = ttml::core::to_xtensor(ttnn_mean_dim0); + auto mean_moreh = ttml::core::to_xtensor(moreh_mean_dim0); + + EXPECT_TRUE(xt::allclose(mean_ttnn, mean_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_ttnn, /*rtol=*/1e-3, /*atol=*/1e-2)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_moreh, /*rtol=*/1e-3, /*atol=*/1e-2)); +} + +TEST_F(ReduceOpTest, TestSumDim0) { + xt::random::seed(42); + auto* device = &ttml::autograd::ctx().get_device(); + xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.1, 0.1).reshape({2, 1, 64, 64}); + + auto xtensor_a_tensor = ttml::core::from_xtensor(xtensor_a, device); + + auto ttnn_sum_dim0 = ttml::ttnn_fixed::sum_ttnn(xtensor_a_tensor, 0, true); + auto moreh_sum_dim0 = ttml::ttnn_fixed::sum_moreh(xtensor_a_tensor, 0, true); + + xt::xarray sum_xtensor = xt::sum(xtensor_a, {0}, xt::evaluation_strategy::immediate); + sum_xtensor.reshape({1, 1, 64, 64}); + + auto sum_ttnn = ttml::core::to_xtensor(ttnn_sum_dim0); + auto sum_moreh = ttml::core::to_xtensor(moreh_sum_dim0); + + EXPECT_TRUE(xt::allclose(sum_ttnn, sum_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); + EXPECT_TRUE(xt::allclose(sum_xtensor, sum_ttnn, /*rtol=*/1e-2, /*atol=*/1e-2)); + EXPECT_TRUE(xt::allclose(sum_xtensor, sum_moreh, /*rtol=*/1e-2, /*atol=*/1e-2)); +} + +TEST_F(ReduceOpTest, TestMeanDim3) { xt::random::seed(42); auto* device = &ttml::autograd::ctx().get_device(); xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.5, 0.5).reshape({2, 1, 64, 64}); @@ -48,7 +90,7 @@ TEST_F(ReduceOpTest, TestMean) { EXPECT_TRUE(xt::allclose(mean_xtensor, mean_moreh, /*rtol=*/1e-3, /*atol=*/1e-2)); } -TEST_F(ReduceOpTest, TestSum) { +TEST_F(ReduceOpTest, TestSumDim3) { xt::random::seed(42); auto* device = &ttml::autograd::ctx().get_device(); xt::xarray xtensor_a = xt::random::rand({128 * 64}, -0.1, 0.1).reshape({2, 1, 64, 64}); @@ -63,11 +105,29 @@ TEST_F(ReduceOpTest, TestSum) { auto sum_ttnn = ttml::core::to_xtensor(ttnn_sum_dim3); auto sum_moreh = ttml::core::to_xtensor(moreh_sum_dim3); - std::cout << sum_ttnn << std::endl; - std::cout << "------------" << std::endl; - std::cout << sum_xtensor << std::endl; EXPECT_TRUE(xt::allclose(sum_ttnn, sum_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); EXPECT_TRUE(xt::allclose(sum_xtensor, sum_ttnn, /*rtol=*/1e-2, /*atol=*/1e-2)); EXPECT_TRUE(xt::allclose(sum_xtensor, sum_moreh, /*rtol=*/1e-2, /*atol=*/1e-2)); } + +TEST_F(ReduceOpTest, TestMeanLargeDim3) { + xt::random::seed(42); + auto* device = &ttml::autograd::ctx().get_device(); + xt::xarray xtensor_a = xt::random::rand({1024 * 1024}, -0.5, 0.5).reshape({2, 1, 512, 1024}); + + auto xtensor_a_tensor = ttml::core::from_xtensor(xtensor_a, device); + + auto ttnn_mean_dim3 = ttml::ttnn_fixed::mean_ttnn(xtensor_a_tensor, 3, true); + auto moreh_mean_dim3 = ttml::ttnn_fixed::mean_moreh(xtensor_a_tensor, 3, true); + + xt::xarray mean_xtensor = xt::mean(xtensor_a, {3}, xt::evaluation_strategy::immediate); + mean_xtensor.reshape({2, 1, 512, 1}); + + auto mean_ttnn = ttml::core::to_xtensor(ttnn_mean_dim3); + auto mean_moreh = ttml::core::to_xtensor(moreh_mean_dim3); + + EXPECT_TRUE(xt::allclose(mean_ttnn, mean_moreh, /*rtol=*/1e-4, /*atol=*/1e-3)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_ttnn, /*rtol=*/1e-3, /*atol=*/1e-2)); + EXPECT_TRUE(xt::allclose(mean_xtensor, mean_moreh, /*rtol=*/1e-3, /*atol=*/1e-2)); +} From 14199b2ee11bedabcc14852ae57201e12d5bd4a8 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:00 -0800 Subject: [PATCH 03/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index fd9cb168b6b..7c9446d3331 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -49,7 +49,7 @@ tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::T return ttnn::multiply(a, inv_b); } -tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { +tt::tt_metal::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { auto res = ttnn::moreh_mean( t, dim, From 7718c52cb862880dd01021bab17dceb605107e66 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:06 -0800 Subject: [PATCH 04/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index 7c9446d3331..caccc6a16f4 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -14,7 +14,7 @@ namespace ttml::ttnn_fixed { tt::tt_metal::Tensor sum_over_dim(const tt::tt_metal::Tensor& t, uint32_t dim) { - return sum_ttnn(t, dim, true); + return sum_ttnn(t, dim, /* keepdim */ true); } tt::tt_metal::Tensor sum_over_batch(const tt::tt_metal::Tensor& t) { From 5a1744ee522af56748df4fca1c7a5ad355e71cc0 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:15 -0800 Subject: [PATCH 05/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index caccc6a16f4..0e15de7396a 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -60,7 +60,7 @@ tt::tt_metal::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool kee /* device_compute_kernel_config */ core::ComputeKernelConfig::precise()); return res; } -tt::tt_metal ::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { +tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { return ttnn::mean(t, dim, keep_dim, std::nullopt, core::ComputeKernelConfig::precise()); } From eb4bd397bd01f11ba7b693a5c2beb61b9c96fc92 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:21 -0800 Subject: [PATCH 06/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index 0e15de7396a..8ea69d7cdb6 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -74,7 +74,7 @@ tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool kee /* device_compute_kernel_config */ core::ComputeKernelConfig::precise()); return res; } -tt::tt_metal ::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { +tt::tt_metal::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { return ttnn::sum(t, dim, keep_dim, std::nullopt, core::ComputeKernelConfig::precise()); } From 1fc32cf8b0d7223bd6ebb179ba5a999a073a8a4e Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:26 -0800 Subject: [PATCH 07/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp index f06b5d2bd5a..738d63e486c 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp @@ -16,7 +16,7 @@ tt::tt_metal::Tensor softmax(const tt::tt_metal::Tensor& t, int dim); tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::Tensor& b); tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); -tt::tt_metal ::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal ::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); From 293f265fe5c62df2cf9935ffc95746ee7a9f7ee9 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:31 -0800 Subject: [PATCH 08/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp index 738d63e486c..35975a6d4a6 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp @@ -19,5 +19,5 @@ tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool ke tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); -tt::tt_metal ::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); } // namespace ttml::ttnn_fixed From a87f5299e53c8e21db474674c6967d4016697442 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:37 -0800 Subject: [PATCH 09/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp index 8ea69d7cdb6..ad864fa0b3d 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.cpp @@ -64,7 +64,7 @@ tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep return ttnn::mean(t, dim, keep_dim, std::nullopt, core::ComputeKernelConfig::precise()); } -tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { +tt::tt_metal::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim) { auto res = ttnn::moreh_sum( t, dim, From 9874d0e231e69c9b64a921f7a8deba060361267f Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:44 -0800 Subject: [PATCH 10/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp index 35975a6d4a6..1b1134d84f5 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp @@ -15,7 +15,7 @@ tt::tt_metal::Tensor log_softmax(const tt::tt_metal::Tensor& t, int dim); tt::tt_metal::Tensor softmax(const tt::tt_metal::Tensor& t, int dim); tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::Tensor& b); -tt::tt_metal ::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); From a22376ac757b8e1e126585ec4ea34c069e1727e2 Mon Sep 17 00:00:00 2001 From: Denys Makoviichuk Date: Wed, 18 Dec 2024 19:10:49 -0800 Subject: [PATCH 11/11] Update tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp Co-authored-by: Roman Furko --- tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp index 1b1134d84f5..c8a62d981bc 100644 --- a/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp +++ b/tt-train/sources/ttml/ttnn_fixed/trivial_ttnn_ops.hpp @@ -18,6 +18,6 @@ tt::tt_metal::Tensor divide(const tt::tt_metal::Tensor& a, const tt::tt_metal::T tt::tt_metal::Tensor mean_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal::Tensor mean_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); -tt::tt_metal ::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); +tt::tt_metal::Tensor sum_moreh(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); tt::tt_metal::Tensor sum_ttnn(const tt::tt_metal::Tensor& t, int dim, bool keep_dim); } // namespace ttml::ttnn_fixed