From b7db8c5afd9fb452be082c6ccc197225f08c292d Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Tue, 9 Jul 2024 06:49:24 +0000 Subject: [PATCH 01/11] #10071: Merge floor_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 -- .../backward_ops/test_backward_floor.py | 3 ++- .../tt_dnn/op_library/backward/backward_ops.cpp | 10 ---------- .../tt_dnn/op_library/backward/backward_ops.hpp | 3 --- .../tt_lib_bindings_tensor_backward_ops.cpp | 17 ----------------- .../unary_backward/device/unary_backward_op.cpp | 2 ++ .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 5 +++++ 9 files changed, 11 insertions(+), 33 deletions(-) diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index eb913f593a3..7291df8e70c 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -193,6 +193,7 @@ Pointwise Unary ttnn/elu_bw ttnn/celu_bw ttnn/rpow_bw + ttnn/floor_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index b4b942d6133..54ab9221b0f 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -928,8 +928,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.repeat_bw -.. autofunction:: tt_lib.tensor.floor_bw - .. autofunction:: tt_lib.tensor.round_bw .. autofunction:: tt_lib.tensor.unary_div_no_nan_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py index b9fd8c12135..8d71f822865 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py @@ -5,6 +5,7 @@ import torch import pytest import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range @@ -22,7 +23,7 @@ def test_bw_floor(input_shapes, device): pyt_y = torch.floor(in_data) - tt_output_tensor_on_device = tt_lib.tensor.floor_bw(grad_tensor) + tt_output_tensor_on_device = ttnn.floor_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index a60a9883150..329761eb192 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -1748,16 +1748,6 @@ std::vector repeat_bw( return operation::decorate_as_composite(__func__, _repeat_bw)(grad, input, shape, output_mem_config); } -std::vector _floor_bw(const Tensor& grad, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor t_zero = zeros_like(grad, output_mem_config); - grad_tensor.emplace_back(t_zero); - return grad_tensor; -} -std::vector floor_bw(const Tensor& grad, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _floor_bw)(grad, output_mem_config); -} - std::vector _round_bw(const Tensor& grad, const MemoryConfig& output_mem_config) { std::vector grad_tensor; Tensor t_zero = zeros_like(grad, output_mem_config); diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 18a5e4b99d8..d8c69403085 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -422,9 +422,6 @@ std::vector complex_sub_bw( std::vector repeat_bw( const Tensor& grad, const Tensor& input, const Shape& shape, const MemoryConfig& output_mem_config); -std::vector floor_bw( - const Tensor& grad, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector round_bw( const Tensor& grad, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index cffa0bde7d6..a4cbf453f9d 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -1153,23 +1153,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - - m_tensor.def("floor_bw", &tt::tt_metal::floor_bw, - py::arg("grad").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Returns an tensor of zeros like ``grad`` tensor - - Input tensor must have BFLOAT16 data type. - - Output tensor will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - - m_tensor.def("round_bw", &tt::tt_metal::round_bw, py::arg("grad").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Returns an tensor of zeros like ``grad`` tensor diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 775252685fa..89fce03079e 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -388,6 +388,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _relu_bw; case UnaryBackwardOpType::LOGIT_BW: return _logit_bw; + case UnaryBackwardOpType::FLOOR_BW: + return _floor_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index c59943916f8..ca74b38c485 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -92,6 +92,7 @@ constexpr auto leaky_relu_bw = ttnn::register_operation>("ttnn::elu_bw"); constexpr auto celu_bw = ttnn::register_operation>("ttnn::celu_bw"); constexpr auto rpow_bw = ttnn::register_operation>("ttnn::rpow_bw"); +constexpr auto floor_bw = ttnn::register_operation>("ttnn::floor_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index f1a2e12c1a7..9f7caca571d 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -262,6 +262,11 @@ void py_module(py::module& module) { module, ttnn::logit_bw, R"doc(Performs backward operations for logit on :attr:`input_tensor` or attr:`input_tensor_a` with given :attr:`grad_tensor`.)doc"); + + detail::bind_unary_backward( + module, + ttnn::floor_bw, + R"doc(Performs backward operations for floor on :attr:`input_tensor` or attr:`input_tensor_a`, attr:`input_tensor_b` with given :attr:`grad_tensor`.)doc"); detail::bind_unary_backward( module, From accb1ca33d57ea6933a80775b4a69624d4a07ca7 Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Tue, 9 Jul 2024 07:21:55 +0000 Subject: [PATCH 02/11] #10071: Merge round_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 -- .../backward_ops/test_backward_round.py | 3 ++- .../tt_dnn/op_library/backward/backward_ops.cpp | 10 ---------- .../tt_dnn/op_library/backward/backward_ops.hpp | 3 --- .../tt_lib_bindings_tensor_backward_ops.cpp | 15 --------------- .../unary_backward/device/unary_backward_op.cpp | 17 ++++++++++++++++- .../unary_backward/device/unary_backward_op.hpp | 2 ++ .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 6 ++++++ 10 files changed, 28 insertions(+), 32 deletions(-) diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index 7291df8e70c..ffcb12cfee9 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -194,6 +194,7 @@ Pointwise Unary ttnn/celu_bw ttnn/rpow_bw ttnn/floor_bw + ttnn/round_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 54ab9221b0f..1f9c3f0a6b6 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -928,8 +928,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.repeat_bw -.. autofunction:: tt_lib.tensor.round_bw - .. autofunction:: tt_lib.tensor.unary_div_no_nan_bw Loss Functions diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py index cc9b7109a9c..3d5f9122359 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py @@ -5,6 +5,7 @@ import torch import pytest import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range @@ -21,7 +22,7 @@ def test_bw_round(input_shapes, device): in_data, input_tensor = data_gen_with_range(input_shapes, -200, 201, device, required_grad=True) pyt_y = torch.round(in_data) - tt_output_tensor_on_device = tt_lib.tensor.round_bw(grad_tensor) + tt_output_tensor_on_device = ttnn.round_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index 329761eb192..42591a78ace 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -1748,16 +1748,6 @@ std::vector repeat_bw( return operation::decorate_as_composite(__func__, _repeat_bw)(grad, input, shape, output_mem_config); } -std::vector _round_bw(const Tensor& grad, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor t_zero = zeros_like(grad, output_mem_config); - grad_tensor.emplace_back(t_zero); - return grad_tensor; -} -std::vector round_bw(const Tensor& grad, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _round_bw)(grad, output_mem_config); -} - std::vector _unary_div_no_nan_bw( const Tensor& grad, const Tensor& input, float scalar, const MemoryConfig& output_mem_config) { std::vector grad_tensor; diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index d8c69403085..2e37d1ce0ad 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -422,9 +422,6 @@ std::vector complex_sub_bw( std::vector repeat_bw( const Tensor& grad, const Tensor& input, const Shape& shape, const MemoryConfig& output_mem_config); -std::vector round_bw( - const Tensor& grad, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector unary_div_no_nan_bw( const Tensor& grad, const Tensor& input, diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index a4cbf453f9d..50208801e9c 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -1153,21 +1153,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("round_bw", &tt::tt_metal::round_bw, - py::arg("grad").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Returns an tensor of zeros like ``grad`` tensor - - Input tensor must have BFLOAT16 data type. - - Output tensor will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient Tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - m_tensor.def("unary_div_no_nan_bw", &tt::tt_metal::unary_div_no_nan_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("scalar") = 1.0f, py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward operations for division with given ``grad`` and ``scalar`` with no nan. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 89fce03079e..bc312c32f6f 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -264,7 +264,6 @@ std::vector _logit_bw(const Tensor& grad, const Tensor& input, const Mem return grad_tensor; } - std::vector _hardshrink_bw( const Tensor& grad, const Tensor& input_tensor, float lambd, const MemoryConfig& output_mem_config) { std::vector grad_tensor; @@ -358,6 +357,20 @@ std::vector _rpow_bw( } +std::vector _floor_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor t_zero = tt::tt_metal::zeros_like(grad, output_mem_config); + grad_tensor.emplace_back(t_zero); + return grad_tensor; +} + +std::vector _round_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor t_zero = tt::tt_metal::zeros_like(grad, output_mem_config); + grad_tensor.emplace_back(t_zero); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -390,6 +403,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _logit_bw; case UnaryBackwardOpType::FLOOR_BW: return _floor_bw; + case UnaryBackwardOpType::ROUND_BW: + return _round_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 410fb4bd098..4fb7556cc23 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -40,6 +40,8 @@ enum class UnaryBackwardOpType { ELU_BW, CELU_BW, RPOW_BW, + FLOOR_BW, + ROUND_BW, }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index ca74b38c485..ef34bd3d231 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -93,6 +93,7 @@ constexpr auto elu_bw = ttnn::register_operation>("ttnn::celu_bw"); constexpr auto rpow_bw = ttnn::register_operation>("ttnn::rpow_bw"); constexpr auto floor_bw = ttnn::register_operation>("ttnn::floor_bw"); +constexpr auto round_bw = ttnn::register_operation>("ttnn::round_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 9f7caca571d..23cc5687638 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -302,6 +302,12 @@ void py_module(py::module& module) { module, ttnn::rpow_bw, R"doc(Performs backward operations for rpow on :attr:`input_tensor`, :attr:`exponent` with given :attr:`grad_tensor`.)doc"); + + detail::bind_unary_backward( + module, + ttnn::round_bw, + R"doc(Performs backward operations for round on :attr:`input_tensor` or attr:`input_tensor_a`, attr:`input_tensor_b` with given :attr:`grad_tensor`.)doc"); + } } // namespace binary_backward From 7aa63dbcce5f42dfd81ec5a4ea2a40340249800c Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Tue, 9 Jul 2024 09:49:12 +0000 Subject: [PATCH 03/11] #10071: Merge log_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 +- docs/source/ttnn/ttnn/ttnn/floor_bw.rst | 6 ++++++ docs/source/ttnn/ttnn/ttnn/log_bw.rst | 6 ++++++ docs/source/ttnn/ttnn/ttnn/round_bw.rst | 6 ++++++ .../backward_ops/test_backward_floor.py | 1 - .../backward_ops/test_backward_log.py | 6 +++--- .../backward_ops/test_backward_round.py | 1 - .../op_library/backward/backward_ops.cpp | 20 ------------------- .../op_library/backward/backward_ops.hpp | 2 +- .../tt_lib_bindings_tensor_backward_ops.cpp | 16 --------------- .../device/unary_backward_op.cpp | 19 ++++++++++++++++++ .../device/unary_backward_op.hpp | 1 + .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 9 +++++++-- 15 files changed, 52 insertions(+), 45 deletions(-) create mode 100644 docs/source/ttnn/ttnn/ttnn/floor_bw.rst create mode 100644 docs/source/ttnn/ttnn/ttnn/log_bw.rst create mode 100644 docs/source/ttnn/ttnn/ttnn/round_bw.rst diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index ffcb12cfee9..2fa51831698 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -195,6 +195,7 @@ Pointwise Unary ttnn/rpow_bw ttnn/floor_bw ttnn/round_bw + ttnn/log_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 1f9c3f0a6b6..c72c5199fa2 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -814,7 +814,7 @@ Backward Operations .. autofunction:: tt_lib.tensor.fill_bw -.. autofunction:: tt_lib.tensor.log_bw +.. autofunction:: tt_lib.tensor.unary_sub_bw .. autofunction:: tt_lib.tensor.abs_bw diff --git a/docs/source/ttnn/ttnn/ttnn/floor_bw.rst b/docs/source/ttnn/ttnn/ttnn/floor_bw.rst new file mode 100644 index 00000000000..54e31216ba6 --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/floor_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.floor_bw: + + ttnn.floor_bw + ############## + + .. autofunction:: ttnn.floor_bw diff --git a/docs/source/ttnn/ttnn/ttnn/log_bw.rst b/docs/source/ttnn/ttnn/ttnn/log_bw.rst new file mode 100644 index 00000000000..027646a942e --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/log_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.log_bw: + + ttnn.log_bw + ############ + + .. autofunction:: ttnn.log_bw diff --git a/docs/source/ttnn/ttnn/ttnn/round_bw.rst b/docs/source/ttnn/ttnn/ttnn/round_bw.rst new file mode 100644 index 00000000000..8dbf21e706a --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/round_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.round_bw: + + ttnn.round_bw + ############## + + .. autofunction:: ttnn.round_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py index 8d71f822865..c9efe6c5e86 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py @@ -4,7 +4,6 @@ import torch import pytest -import tt_lib import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py index c404ff025c4..3418ea67dfa 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py @@ -4,7 +4,7 @@ import torch import pytest -import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import ( data_gen_with_val, compare_pcc, @@ -23,7 +23,7 @@ def test_bw_log_0(input_shapes, device): in_data, input_tensor = data_gen_with_val(input_shapes, device, True, val=0) grad_data, grad_tensor = data_gen_with_range(input_shapes, -1, 1, device) - tt_output_tensor_on_device = tt_lib.tensor.log_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.log_bw(grad_tensor, input_tensor) in_data.retain_grad() @@ -47,7 +47,7 @@ def test_bw_log_0(input_shapes, device): def test_bw_log(input_shapes, device): in_data, input_tensor = data_gen_with_range(input_shapes, -100, 100, device, True) grad_data, grad_tensor = data_gen_with_range(input_shapes, -100, 100, device) - tt_output_tensor_on_device = tt_lib.tensor.log_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.log_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py index 3d5f9122359..baabffc0009 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py @@ -4,7 +4,6 @@ import torch import pytest -import tt_lib import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index 42591a78ace..c4cc6f28821 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -280,26 +280,6 @@ std::vector ne_bw(const Tensor& grad, const MemoryConfig& output_mem_con return operation::decorate_as_composite(__func__, _ne_bw)(grad, output_mem_config); } -std::vector _log_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor grad_a = ttnn::multiply(grad, ttnn::reciprocal(input, output_mem_config), std::nullopt, output_mem_config); - Tensor t_inf = full_like(input, std::numeric_limits::infinity(), output_mem_config); - Tensor t_nan = full_like(input, std::nanf(""), output_mem_config); - grad_tensor.emplace_back(where( - ttnn::eqz(input, output_mem_config), - where( - ttnn::eqz(grad, output_mem_config), - t_nan, - ttnn::multiply(t_inf, ttnn::sign(grad, output_mem_config), std::nullopt, output_mem_config), - output_mem_config), - grad_a, - output_mem_config)); - return grad_tensor; -} -std::vector log_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _log_bw)(grad, input, output_mem_config); -} - std::vector _abs_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; Tensor result = ttnn::multiply(grad, ttnn::sign(input, output_mem_config), std::nullopt, output_mem_config); diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 2e37d1ce0ad..76b150b1f36 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -101,7 +101,7 @@ std::vector> tanh_bw( std::vector fill_bw( const Tensor& grad, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector log_bw( +std::vector unary_sub_bw( const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index 50208801e9c..f661e4e6841 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -328,22 +328,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("log_bw", &tt::tt_metal::log_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Performs backward operations for logarithm of ``input`` tensors with given ``grad``. - - Input tensors must have BFLOAT16 data type. - - Output tensors will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor add is applied to", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - m_tensor.def("abs_bw", &tt::tt_metal::abs_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward operations for abs of ``input`` tensors with given ``grad``. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index bc312c32f6f..3a21990e5ac 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -371,6 +371,23 @@ std::vector _round_bw(const Tensor& grad, const Tensor& input, const Mem return grad_tensor; } +std::vector _log_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor grad_a = ttnn::multiply(grad, ttnn::reciprocal(input, output_mem_config), std::nullopt, output_mem_config); + Tensor t_inf = tt::tt_metal::full_like(input, std::numeric_limits::infinity(), output_mem_config); + Tensor t_nan = tt::tt_metal::full_like(input, std::nanf(""), output_mem_config); + grad_tensor.emplace_back(where( + ttnn::eqz(input, output_mem_config), + where( + ttnn::eqz(grad, output_mem_config), + t_nan, + ttnn::multiply(t_inf, ttnn::sign(grad, output_mem_config), std::nullopt, output_mem_config), + output_mem_config), + grad_a, + output_mem_config)); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -405,6 +422,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _floor_bw; case UnaryBackwardOpType::ROUND_BW: return _round_bw; + case UnaryBackwardOpType::LOG_BW: + return _log_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 4fb7556cc23..5377b3b5fe9 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -42,6 +42,7 @@ enum class UnaryBackwardOpType { RPOW_BW, FLOOR_BW, ROUND_BW, + LOG_BW, }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index ef34bd3d231..8da1cd7c67c 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -94,6 +94,7 @@ constexpr auto celu_bw = ttnn::register_operation>("ttnn::rpow_bw"); constexpr auto floor_bw = ttnn::register_operation>("ttnn::floor_bw"); constexpr auto round_bw = ttnn::register_operation>("ttnn::round_bw"); +constexpr auto log_bw = ttnn::register_operation>("ttnn::log_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 23cc5687638..6975ee4c5cf 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -266,7 +266,7 @@ void py_module(py::module& module) { detail::bind_unary_backward( module, ttnn::floor_bw, - R"doc(Performs backward operations for floor on :attr:`input_tensor` or attr:`input_tensor_a`, attr:`input_tensor_b` with given :attr:`grad_tensor`.)doc"); + R"doc(Performs backward operations for floor on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); detail::bind_unary_backward( module, @@ -306,7 +306,12 @@ void py_module(py::module& module) { detail::bind_unary_backward( module, ttnn::round_bw, - R"doc(Performs backward operations for round on :attr:`input_tensor` or attr:`input_tensor_a`, attr:`input_tensor_b` with given :attr:`grad_tensor`.)doc"); + R"doc(Performs backward operations for round on :attr:`input_tensor` with given :attr:`grad_tensor`.)doc"); + + detail::bind_unary_backward( + module, + ttnn::log_bw, + R"doc(Performs backward operations for logarithm on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); } From bc8f79300668545754f630e1e7d3fd4af9e4c307 Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Tue, 9 Jul 2024 17:37:51 +0000 Subject: [PATCH 04/11] #10071: Merge relu6_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 +- docs/source/ttnn/ttnn/ttnn/relu6_bw.rst | 6 ++++ .../backward_ops/test_backward_relu6.py | 4 +-- .../op_library/backward/backward_ops.cpp | 36 ++++++++----------- .../op_library/backward/backward_ops.hpp | 3 +- .../tt_lib_bindings_tensor_backward_ops.cpp | 14 ++++---- .../device/unary_backward_op.cpp | 25 +++++++++++++ .../device/unary_backward_op.hpp | 1 + .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 5 +++ 11 files changed, 66 insertions(+), 32 deletions(-) create mode 100644 docs/source/ttnn/ttnn/ttnn/relu6_bw.rst diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index 2fa51831698..7e7e74621ce 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -196,6 +196,7 @@ Pointwise Unary ttnn/floor_bw ttnn/round_bw ttnn/log_bw + ttnn/relu6_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index c72c5199fa2..322eeffc8a4 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -880,7 +880,7 @@ Backward Operations .. autofunction:: tt_lib.tensor.reciprocal_bw -.. autofunction:: tt_lib.tensor.relu6_bw +.. autofunction:: tt_lib.tensor.rpow_bw .. autofunction:: tt_lib.tensor.silu_bw diff --git a/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst b/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst new file mode 100644 index 00000000000..1a61f8b3d70 --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.relu6_bw: + + ttnn.relu6_bw + ############## + + .. autofunction:: ttnn.relu6_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py index 2ee155ac401..15aff1a6855 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py @@ -4,7 +4,7 @@ import torch import pytest -import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range @@ -22,7 +22,7 @@ def test_bw_relu6(input_shapes, device): pyt_y = torch.nn.functional.relu6(in_data) - tt_output_tensor_on_device = tt_lib.tensor.relu6_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.relu6_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index c4cc6f28821..821cf9c5472 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -970,33 +970,27 @@ std::vector reciprocal_bw(const Tensor& grad, const Tensor& input, const return operation::decorate_as_composite(__func__, _reciprocal_bw)(grad, input, output_mem_config); } -std::vector _relu6_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { +std::vector _rpow_bw( + const Tensor& grad, const Tensor& input, float exponent, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor zero_tensor = zeros_like(input, output_mem_config); - Tensor one_tensor = ones_like(input, output_mem_config); - Tensor six_tensor = full_like(input, 6, output_mem_config); - Tensor grad_result = - where(ttnn::le(input, zero_tensor, std::nullopt, output_mem_config), zero_tensor, six_tensor, output_mem_config); - grad_result = where( - ttnn::logical_and( - ttnn::gtz(input, output_mem_config), - ttnn::lt(input, six_tensor, std::nullopt, output_mem_config), - std::nullopt, - output_mem_config), - grad, - grad_result, - output_mem_config); - grad_result = - where(ttnn::ge(input, six_tensor, std::nullopt, output_mem_config), zero_tensor, grad_result, output_mem_config); - + float t_nan = std::nanf(""); + Tensor grad_result = zeros_like(input, output_mem_config); + if (exponent != 0.0) { + grad_result = + ttnn::multiply(grad, + ttnn::multiply(pow(input, exponent - 1, output_mem_config), exponent, std::nullopt, output_mem_config), + std::nullopt, + output_mem_config); + grad_result = where(ttnn::ltz(input, output_mem_config), t_nan, grad_result, output_mem_config); + } grad_tensor.emplace_back(grad_result); return grad_tensor; } -std::vector relu6_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _relu6_bw)(grad, input, output_mem_config); +std::vector rpow_bw( + const Tensor& grad, const Tensor& input, float exponent, const MemoryConfig& output_mem_config) { + return operation::decorate_as_composite(__func__, _rpow_bw)(grad, input, exponent, output_mem_config); } - // Silu // result: grad * sigmoid_result * (1 + input * (1 - sigmoid_result)) std::vector _silu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 76b150b1f36..50740012285 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -284,9 +284,10 @@ std::vector reciprocal_bw( const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector relu6_bw( +std::vector rpow_bw( const Tensor& grad, const Tensor& input, + float exponent, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); std::vector silu_bw( diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index f661e4e6841..2842e6235e3 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -807,22 +807,22 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("relu6_bw", &tt::tt_metal::relu6_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Returns an tensor of backward operation of relu6 for ``input`` tensor and ``grad`` tensor. + m_tensor.def("rpow_bw", &tt::tt_metal::rpow_bw, + py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("exponent").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( + Performs backward operations for rpow for the ``input`` and ``exponent`` with given ``grad`` Input tensors must have BFLOAT16 data type. - Output tensors will have BFLOAT16 data type. + Output tensor will have BFLOAT16 data type. .. csv-table:: :header: "Argument", "Description", "Data type", "Valid range", "Required" "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor relu6 is applied to", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" + "input", "Tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" + "exponent", "exponent", "float", ">0.0", "Yes" "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - + )doc"); m_tensor.def("silu_bw", &tt::tt_metal::silu_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 3a21990e5ac..183c6343038 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -388,6 +388,29 @@ std::vector _log_bw(const Tensor& grad, const Tensor& input, const Memor return grad_tensor; } +std::vector _relu6_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor zero_tensor = tt::tt_metal::zeros_like(input, output_mem_config); + Tensor one_tensor = tt::tt_metal::ones_like(input, output_mem_config); + Tensor six_tensor = tt::tt_metal::full_like(input, 6, output_mem_config); + Tensor grad_result = + where(ttnn::le(input, zero_tensor, std::nullopt, output_mem_config), zero_tensor, six_tensor, output_mem_config); + grad_result = where( + ttnn::logical_and( + ttnn::gtz(input, output_mem_config), + ttnn::lt(input, six_tensor, std::nullopt, output_mem_config), + std::nullopt, + output_mem_config), + grad, + grad_result, + output_mem_config); + grad_result = + where(ttnn::ge(input, six_tensor, std::nullopt, output_mem_config), zero_tensor, grad_result, output_mem_config); + + grad_tensor.emplace_back(grad_result); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -424,6 +447,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _round_bw; case UnaryBackwardOpType::LOG_BW: return _log_bw; + case UnaryBackwardOpType::RELU6_BW: + return _relu6_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 5377b3b5fe9..5cac6f62779 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -43,6 +43,7 @@ enum class UnaryBackwardOpType { FLOOR_BW, ROUND_BW, LOG_BW, + RELU6_BW }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index 8da1cd7c67c..6b75af8a8c3 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -95,6 +95,7 @@ constexpr auto rpow_bw = ttnn::register_operation>("ttnn::floor_bw"); constexpr auto round_bw = ttnn::register_operation>("ttnn::round_bw"); constexpr auto log_bw = ttnn::register_operation>("ttnn::log_bw"); +constexpr auto relu6_bw = ttnn::register_operation>("ttnn::relu6_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 6975ee4c5cf..44ef0602ab6 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -313,6 +313,11 @@ void py_module(py::module& module) { ttnn::log_bw, R"doc(Performs backward operations for logarithm on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + detail::bind_unary_backward( + module, + ttnn::relu6_bw, + R"doc(Performs backward operations for relu6 on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + } } // namespace binary_backward From 0baac6e0ce7ec083532575b4fe73fee7fa5d0ccb Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Tue, 9 Jul 2024 18:01:33 +0000 Subject: [PATCH 05/11] #10071: Merge abs_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 -- docs/source/ttnn/ttnn/ttnn/abs_bw.rst | 6 ++++++ .../backward_ops/test_backward_abs.py | 4 ++-- .../op_library/backward/backward_ops.cpp | 18 ++++++++++++++---- .../op_library/backward/backward_ops.hpp | 5 ----- .../tt_lib_bindings_tensor_backward_ops.cpp | 16 ---------------- .../device/unary_backward_op.cpp | 9 +++++++++ .../device/unary_backward_op.hpp | 3 ++- .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 5 +++++ 11 files changed, 40 insertions(+), 30 deletions(-) create mode 100644 docs/source/ttnn/ttnn/ttnn/abs_bw.rst diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index 7e7e74621ce..f29b309ef75 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -197,6 +197,7 @@ Pointwise Unary ttnn/round_bw ttnn/log_bw ttnn/relu6_bw + ttnn/abs_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 322eeffc8a4..5acd9c68ced 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -816,8 +816,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.unary_sub_bw -.. autofunction:: tt_lib.tensor.abs_bw - .. autofunction:: tt_lib.tensor.complex_abs_bw .. autofunction:: tt_lib.tensor.lt_bw diff --git a/docs/source/ttnn/ttnn/ttnn/abs_bw.rst b/docs/source/ttnn/ttnn/ttnn/abs_bw.rst new file mode 100644 index 00000000000..1f39180f834 --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/abs_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.abs_bw: + + ttnn.abs_bw + ############ + + .. autofunction:: ttnn.abs_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py index 83e419f7c65..02b5481f7ec 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py @@ -4,7 +4,7 @@ import torch import pytest -import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc @@ -22,7 +22,7 @@ def test_bw_abs(input_shapes, device): pyt_y = torch.abs(in_data) - tt_output_tensor_on_device = tt_lib.tensor.abs_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.abs_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index 821cf9c5472..caee9737bf7 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -280,14 +280,24 @@ std::vector ne_bw(const Tensor& grad, const MemoryConfig& output_mem_con return operation::decorate_as_composite(__func__, _ne_bw)(grad, output_mem_config); } -std::vector _abs_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { +std::vector _rsqrt_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor result = ttnn::multiply(grad, ttnn::sign(input, output_mem_config), std::nullopt, output_mem_config); + Tensor rsqrt_result = power(rsqrt(input, true, output_mem_config), 3, output_mem_config); + Tensor result = mul_unary(ttnn::multiply(grad, rsqrt_result, std::nullopt, output_mem_config), -0.5, output_mem_config); + float t_inf = std::numeric_limits::infinity(); + result = where(eqz(input, output_mem_config), t_inf, result, output_mem_config); + float t_nan = std::nanf(""); + result = where(ltz(input, output_mem_config), t_nan, result, output_mem_config); + result = where( + ttnn::logical_and(eqz(input, output_mem_config), eqz(grad, output_mem_config), std::nullopt, output_mem_config), + t_nan, + result, + output_mem_config); grad_tensor.emplace_back(result); return grad_tensor; } -std::vector abs_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _abs_bw)(grad, input, output_mem_config); +std::vector rsqrt_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + return operation::decorate_as_composite(__func__, _rsqrt_bw)(grad, input, output_mem_config); } // bw(expm1) = grad * expm1(input) + 1 diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 50740012285..8d6606065d2 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -111,11 +111,6 @@ std::vector binary_le_bw( const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector abs_bw( - const Tensor& grad, - const Tensor& input, - const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector complex_abs_bw( const Tensor& grad, const Tensor& input, diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index 2842e6235e3..0b39023b0be 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -328,22 +328,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("abs_bw", &tt::tt_metal::abs_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Performs backward operations for abs of ``input`` tensors with given ``grad``. - - Input tensors must have BFLOAT16 data type. - - Output tensor will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor add is applied to", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - m_tensor.def("complex_abs_bw", py::overload_cast(&complex_abs_bw), py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward operations for abs of complex ``input`` tensor with given ``grad``. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 183c6343038..2b95d936e86 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -411,6 +411,13 @@ std::vector _relu6_bw(const Tensor& grad, const Tensor& input, const Mem return grad_tensor; } +std::vector _abs_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor result = ttnn::multiply(grad, ttnn::sign(input, output_mem_config), std::nullopt, output_mem_config); + grad_tensor.emplace_back(result); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -449,6 +456,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _log_bw; case UnaryBackwardOpType::RELU6_BW: return _relu6_bw; + case UnaryBackwardOpType::ABS_BW: + return _abs_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 5cac6f62779..3a938c7f1c9 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -43,7 +43,8 @@ enum class UnaryBackwardOpType { FLOOR_BW, ROUND_BW, LOG_BW, - RELU6_BW + RELU6_BW, + ABS_BW, }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index 6b75af8a8c3..c523de91a4c 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -96,6 +96,7 @@ constexpr auto floor_bw = ttnn::register_operation>("ttnn::round_bw"); constexpr auto log_bw = ttnn::register_operation>("ttnn::log_bw"); constexpr auto relu6_bw = ttnn::register_operation>("ttnn::relu6_bw"); +constexpr auto abs_bw = ttnn::register_operation>("ttnn::abs_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 44ef0602ab6..7144c851be8 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -318,6 +318,11 @@ void py_module(py::module& module) { ttnn::relu6_bw, R"doc(Performs backward operations for relu6 on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + detail::bind_unary_backward( + module, + ttnn::abs_bw, + R"doc(Performs backward operations for abs on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + } } // namespace binary_backward From 3481d61b584bb293d171fa9c4f73f6f8bcb8ba7f Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Wed, 10 Jul 2024 05:44:26 +0000 Subject: [PATCH 06/11] #10071: Merge silu_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 -- docs/source/ttnn/ttnn/ttnn/silu_bw.rst | 6 +++++ .../backward_ops/test_backward_silu.py | 4 ++-- .../op_library/backward/backward_ops.cpp | 22 ------------------- .../op_library/backward/backward_ops.hpp | 5 ----- .../tt_lib_bindings_tensor_backward_ops.cpp | 17 -------------- .../device/unary_backward_op.cpp | 19 ++++++++++++++++ .../device/unary_backward_op.hpp | 1 + .../eltwise/unary_backward/unary_backward.hpp | 1 + .../unary_backward/unary_backward_pybind.hpp | 5 +++++ 11 files changed, 35 insertions(+), 48 deletions(-) create mode 100644 docs/source/ttnn/ttnn/ttnn/silu_bw.rst diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index f29b309ef75..2d866a4c7fb 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -198,6 +198,7 @@ Pointwise Unary ttnn/log_bw ttnn/relu6_bw ttnn/abs_bw + ttnn/silu_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 5acd9c68ced..7bbd91c7546 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -880,8 +880,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.rpow_bw -.. autofunction:: tt_lib.tensor.silu_bw - .. autofunction:: tt_lib.tensor.selu_bw .. autofunction:: tt_lib.tensor.square_bw diff --git a/docs/source/ttnn/ttnn/ttnn/silu_bw.rst b/docs/source/ttnn/ttnn/ttnn/silu_bw.rst new file mode 100644 index 00000000000..16c8ea8efe4 --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/silu_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.silu_bw: + + ttnn.silu_bw + ############# + + .. autofunction:: ttnn.silu_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py index 455c87ce191..c86af349dbe 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py @@ -4,7 +4,7 @@ import torch import pytest -import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc @@ -22,7 +22,7 @@ def test_bw_silu(input_shapes, device): pyt_y = torch.nn.functional.silu(in_data) - tt_output_tensor_on_device = tt_lib.tensor.silu_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.silu_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index caee9737bf7..fde73df19c9 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -1001,28 +1001,6 @@ std::vector rpow_bw( return operation::decorate_as_composite(__func__, _rpow_bw)(grad, input, exponent, output_mem_config); } -// Silu -// result: grad * sigmoid_result * (1 + input * (1 - sigmoid_result)) -std::vector _silu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor grad_sigmoid = ttnn::multiply(grad, ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config); - Tensor add_sub = ttnn::add( - ttnn::multiply(ttnn::subtract(ttnn::full_like(input, 1.0f) , ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config), - input, - std::nullopt, - output_mem_config), - 1.0f, - std::nullopt, - output_mem_config); - Tensor grad_result = ttnn::multiply(grad_sigmoid, add_sub, std::nullopt, output_mem_config); - - grad_tensor.emplace_back(grad_result); - return grad_tensor; -} -std::vector silu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _silu_bw)(grad, input, output_mem_config); -} - // Selu // result: torch.where(input > 0, grad * lambd, grad * lambd * alpha * torch.exp(input)) std::vector _selu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 8d6606065d2..c1d60d32ea4 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -285,11 +285,6 @@ std::vector rpow_bw( float exponent, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector silu_bw( - const Tensor& grad, - const Tensor& input, - const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector selu_bw( const Tensor& grad, const Tensor& input, diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index 0b39023b0be..ea2dc4907a1 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -808,23 +808,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("silu_bw", &tt::tt_metal::silu_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Performs backward operations for silu sin of ``input`` tensors with given ``grad``. - - Input tensors must have BFLOAT16 data type. - - Output tensors will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor silu_bw is applied to", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - - m_tensor.def("selu_bw", &tt::tt_metal::selu_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward operations for selu sin of ``input`` tensors with given ``grad``. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 2b95d936e86..252d61eeb02 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -418,6 +418,23 @@ std::vector _abs_bw(const Tensor& grad, const Tensor& input, const Memor return grad_tensor; } +// Silu +// result: grad * sigmoid_result * (1 + input * (1 - sigmoid_result)) +std::vector _silu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor grad_sigmoid = ttnn::multiply(grad, sigmoid(input, output_mem_config), std::nullopt, output_mem_config); + Tensor add_sub = add1( + ttnn::multiply(sub_unary(1.0f, sigmoid(input, output_mem_config), output_mem_config), + input, + std::nullopt, + output_mem_config), + output_mem_config); + Tensor grad_result = ttnn::multiply(grad_sigmoid, add_sub, std::nullopt, output_mem_config); + + grad_tensor.emplace_back(grad_result); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -458,6 +475,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _relu6_bw; case UnaryBackwardOpType::ABS_BW: return _abs_bw; + case UnaryBackwardOpType::SILU_BW: + return _silu_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 3a938c7f1c9..7b6fb78fc21 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -45,6 +45,7 @@ enum class UnaryBackwardOpType { LOG_BW, RELU6_BW, ABS_BW, + SILU_BW }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index c523de91a4c..f42d61b4197 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -97,6 +97,7 @@ constexpr auto round_bw = ttnn::register_operation>("ttnn::log_bw"); constexpr auto relu6_bw = ttnn::register_operation>("ttnn::relu6_bw"); constexpr auto abs_bw = ttnn::register_operation>("ttnn::abs_bw"); +constexpr auto silu_bw = ttnn::register_operation>("ttnn::silu_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 7144c851be8..047d0eb56b9 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -323,6 +323,11 @@ void py_module(py::module& module) { ttnn::abs_bw, R"doc(Performs backward operations for abs on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + detail::bind_unary_backward( + module, + ttnn::silu_bw, + R"doc(Performs backward operations for silu on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + } } // namespace binary_backward From f049411b5e2db1ab5ec55f7277c0806a71f8c44d Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Wed, 10 Jul 2024 07:21:12 +0000 Subject: [PATCH 07/11] #10071: Merge selu_bw to TTNN --- docs/source/ttnn/ttnn/api.rst | 1 + docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 2 -- docs/source/ttnn/ttnn/ttnn/selu_bw.rst | 6 +++++ .../backward_ops/test_backward_selu.py | 4 +-- .../op_library/backward/backward_ops.cpp | 21 --------------- .../op_library/backward/backward_ops.hpp | 5 ---- .../tt_lib_bindings_tensor_backward_ops.cpp | 16 ----------- .../device/unary_backward_op.cpp | 27 ++++++++++++++++--- .../device/unary_backward_op.hpp | 3 ++- .../eltwise/unary_backward/unary_backward.hpp | 2 +- .../unary_backward/unary_backward_pybind.hpp | 5 ++++ 11 files changed, 41 insertions(+), 51 deletions(-) create mode 100644 docs/source/ttnn/ttnn/ttnn/selu_bw.rst diff --git a/docs/source/ttnn/ttnn/api.rst b/docs/source/ttnn/ttnn/api.rst index 2d866a4c7fb..a04456f3c6d 100644 --- a/docs/source/ttnn/ttnn/api.rst +++ b/docs/source/ttnn/ttnn/api.rst @@ -199,6 +199,7 @@ Pointwise Unary ttnn/relu6_bw ttnn/abs_bw ttnn/silu_bw + ttnn/selu_bw Pointwise Binary ================ diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 7bbd91c7546..8e59eb3a369 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -880,8 +880,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.rpow_bw -.. autofunction:: tt_lib.tensor.selu_bw - .. autofunction:: tt_lib.tensor.square_bw .. autofunction:: tt_lib.tensor.tanhshrink_bw diff --git a/docs/source/ttnn/ttnn/ttnn/selu_bw.rst b/docs/source/ttnn/ttnn/ttnn/selu_bw.rst new file mode 100644 index 00000000000..591459bf6ec --- /dev/null +++ b/docs/source/ttnn/ttnn/ttnn/selu_bw.rst @@ -0,0 +1,6 @@ +.. _ttnn.selu_bw: + + ttnn.selu_bw + ############# + + .. autofunction:: ttnn.selu_bw diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py index b834db1fbe1..4218ffa1524 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py +++ b/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py @@ -4,7 +4,7 @@ import torch import pytest -import tt_lib +import ttnn from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc @@ -22,7 +22,7 @@ def test_bw_selu(input_shapes, device): pyt_y = torch.nn.functional.selu(in_data) - tt_output_tensor_on_device = tt_lib.tensor.selu_bw(grad_tensor, input_tensor) + tt_output_tensor_on_device = ttnn.selu_bw(grad_tensor, input_tensor) in_data.retain_grad() diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index fde73df19c9..8f17edc8f92 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -1001,27 +1001,6 @@ std::vector rpow_bw( return operation::decorate_as_composite(__func__, _rpow_bw)(grad, input, exponent, output_mem_config); } -// Selu -// result: torch.where(input > 0, grad * lambd, grad * lambd * alpha * torch.exp(input)) -std::vector _selu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor grad_lambd = ttnn::multiply(grad, 1.0507f, std::nullopt, output_mem_config); - Tensor grad_result = where( - ttnn::gtz(input, output_mem_config), - grad_lambd, - ttnn::multiply(ttnn::multiply(grad_lambd, 1.673260f, std::nullopt, output_mem_config), - ttnn::exp(input, false, output_mem_config), - std::nullopt, - output_mem_config), - output_mem_config); - grad_tensor.emplace_back(grad_result); - return grad_tensor; -} -std::vector selu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _selu_bw)(grad, input, output_mem_config); -} - - // Autoformat support Tensor change_layout_to_tile(const Tensor& temp, const MemoryConfig& output_mem_config) { auto formatted_input_tensor = temp; diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index c1d60d32ea4..1e4b649a92c 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -285,11 +285,6 @@ std::vector rpow_bw( float exponent, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector selu_bw( - const Tensor& grad, - const Tensor& input, - const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector square_bw( const Tensor& grad, const Tensor& input, diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index ea2dc4907a1..5c7864f1a1b 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -808,22 +808,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("selu_bw", &tt::tt_metal::selu_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Performs backward operations for selu sin of ``input`` tensors with given ``grad``. - - Input tensors must have BFLOAT16 data type. - - Output tensors will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor selu_bw is applied to", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - m_tensor.def("square_bw", &tt::tt_metal::square_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward square operations on ``input`` tensors with given ``grad``. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index 252d61eeb02..c498663e660 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -422,12 +422,14 @@ std::vector _abs_bw(const Tensor& grad, const Tensor& input, const Memor // result: grad * sigmoid_result * (1 + input * (1 - sigmoid_result)) std::vector _silu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor grad_sigmoid = ttnn::multiply(grad, sigmoid(input, output_mem_config), std::nullopt, output_mem_config); - Tensor add_sub = add1( - ttnn::multiply(sub_unary(1.0f, sigmoid(input, output_mem_config), output_mem_config), + Tensor grad_sigmoid = ttnn::multiply(grad, ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config); + Tensor add_sub = ttnn::add( + ttnn::multiply(ttnn::subtract(tt::tt_metal::full_like(input, 1.0f) , ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config), input, std::nullopt, output_mem_config), + 1.0f, + std::nullopt, output_mem_config); Tensor grad_result = ttnn::multiply(grad_sigmoid, add_sub, std::nullopt, output_mem_config); @@ -435,6 +437,23 @@ std::vector _silu_bw(const Tensor& grad, const Tensor& input, const Memo return grad_tensor; } +// Selu +// result: torch.where(input > 0, grad * lambd, grad * lambd * alpha * torch.exp(input)) +std::vector _selu_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { + std::vector grad_tensor; + Tensor grad_lambd = ttnn::multiply(grad, 1.0507f, std::nullopt, output_mem_config); + Tensor grad_result = where( + ttnn::gtz(input, output_mem_config), + grad_lambd, + ttnn::multiply(ttnn::multiply(grad_lambd, 1.673260f, std::nullopt, output_mem_config), + ttnn::exp(input, false, output_mem_config), + std::nullopt, + output_mem_config), + output_mem_config); + grad_tensor.emplace_back(grad_result); + return grad_tensor; +} + std::function(const Tensor&, const Tensor&, const MemoryConfig&)> UnaryBackwardFunction::get_function_type1(UnaryBackwardOpType OpType){ switch (OpType) { case UnaryBackwardOpType::ASSIGN_BW: @@ -477,6 +496,8 @@ std::function(const Tensor&, const Tensor&, const Memo return _abs_bw; case UnaryBackwardOpType::SILU_BW: return _silu_bw; + case UnaryBackwardOpType::SELU_BW: + return _selu_bw; default: TT_ASSERT(false && "Undefined op type"); return 0; diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp index 7b6fb78fc21..4851ff39af3 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.hpp @@ -45,7 +45,8 @@ enum class UnaryBackwardOpType { LOG_BW, RELU6_BW, ABS_BW, - SILU_BW + SILU_BW, + SELU_BW, }; struct UnaryBackwardFunction{ diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp index f42d61b4197..2a5bf896106 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward.hpp @@ -98,6 +98,6 @@ constexpr auto log_bw = ttnn::register_operation>("ttnn::relu6_bw"); constexpr auto abs_bw = ttnn::register_operation>("ttnn::abs_bw"); constexpr auto silu_bw = ttnn::register_operation>("ttnn::silu_bw"); - +constexpr auto selu_bw = ttnn::register_operation>("ttnn::selu_bw"); } // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp index 047d0eb56b9..efd7eb2b5b4 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/unary_backward_pybind.hpp @@ -328,6 +328,11 @@ void py_module(py::module& module) { ttnn::silu_bw, R"doc(Performs backward operations for silu on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + detail::bind_unary_backward( + module, + ttnn::selu_bw, + R"doc(Performs backward operations for selu on :attr:`input_tensor` with given :attr:`grad_tensor`)doc"); + } } // namespace binary_backward From 873523bac47798352b401ba8ec8dc01f079ef743 Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Wed, 10 Jul 2024 07:30:54 +0000 Subject: [PATCH 08/11] #10071: Fix rebase errors --- .../eltwise/unary_backward/device/unary_backward_op.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index c498663e660..c9a05b8566e 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -144,7 +144,6 @@ std::vector _log_sigmoid_bw(const Tensor& grad, const Tensor& input, con return grad_tensor; } - std::vector _fill_zero_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; Tensor result = tt::tt_metal::zeros_like(grad, output_mem_config); From 297e916de8e3222867f326092d92e49f9c11d74b Mon Sep 17 00:00:00 2001 From: VirdhatchaniKN Date: Wed, 10 Jul 2024 09:11:15 +0000 Subject: [PATCH 09/11] #10071: Update files --- docs/source/ttnn/ttnn/ttnn/abs_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/floor_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/log_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/relu6_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/round_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/selu_bw.rst | 4 ++-- docs/source/ttnn/ttnn/ttnn/silu_bw.rst | 4 ++-- 7 files changed, 14 insertions(+), 14 deletions(-) diff --git a/docs/source/ttnn/ttnn/ttnn/abs_bw.rst b/docs/source/ttnn/ttnn/ttnn/abs_bw.rst index 1f39180f834..dc397a7c134 100644 --- a/docs/source/ttnn/ttnn/ttnn/abs_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/abs_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.abs_bw: - ttnn.abs_bw - ############ +ttnn.abs_bw +########### .. autofunction:: ttnn.abs_bw diff --git a/docs/source/ttnn/ttnn/ttnn/floor_bw.rst b/docs/source/ttnn/ttnn/ttnn/floor_bw.rst index 54e31216ba6..6baffddfda2 100644 --- a/docs/source/ttnn/ttnn/ttnn/floor_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/floor_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.floor_bw: - ttnn.floor_bw - ############## +ttnn.floor_bw +############# .. autofunction:: ttnn.floor_bw diff --git a/docs/source/ttnn/ttnn/ttnn/log_bw.rst b/docs/source/ttnn/ttnn/ttnn/log_bw.rst index 027646a942e..5f6d4ffe8f2 100644 --- a/docs/source/ttnn/ttnn/ttnn/log_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/log_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.log_bw: - ttnn.log_bw - ############ +ttnn.log_bw +########### .. autofunction:: ttnn.log_bw diff --git a/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst b/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst index 1a61f8b3d70..fc9ea28206d 100644 --- a/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/relu6_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.relu6_bw: - ttnn.relu6_bw - ############## +ttnn.relu6_bw +############# .. autofunction:: ttnn.relu6_bw diff --git a/docs/source/ttnn/ttnn/ttnn/round_bw.rst b/docs/source/ttnn/ttnn/ttnn/round_bw.rst index 8dbf21e706a..d38af90674c 100644 --- a/docs/source/ttnn/ttnn/ttnn/round_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/round_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.round_bw: - ttnn.round_bw - ############## +ttnn.round_bw +############# .. autofunction:: ttnn.round_bw diff --git a/docs/source/ttnn/ttnn/ttnn/selu_bw.rst b/docs/source/ttnn/ttnn/ttnn/selu_bw.rst index 591459bf6ec..e3679d997bd 100644 --- a/docs/source/ttnn/ttnn/ttnn/selu_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/selu_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.selu_bw: - ttnn.selu_bw - ############# +ttnn.selu_bw +############ .. autofunction:: ttnn.selu_bw diff --git a/docs/source/ttnn/ttnn/ttnn/silu_bw.rst b/docs/source/ttnn/ttnn/ttnn/silu_bw.rst index 16c8ea8efe4..2f18db474fc 100644 --- a/docs/source/ttnn/ttnn/ttnn/silu_bw.rst +++ b/docs/source/ttnn/ttnn/ttnn/silu_bw.rst @@ -1,6 +1,6 @@ .. _ttnn.silu_bw: - ttnn.silu_bw - ############# +ttnn.silu_bw +############ .. autofunction:: ttnn.silu_bw From 0112a8fc2b61330e929daace2f2887f2252522a7 Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Wed, 10 Jul 2024 15:03:34 +0000 Subject: [PATCH 10/11] #10071: Move test files to TTNN --- .../unit_tests/operations/backward}/test_backward_abs.py | 2 +- .../unit_tests/operations/backward}/test_backward_floor.py | 2 +- .../unit_tests/operations/backward}/test_backward_log.py | 2 +- .../unit_tests/operations/backward}/test_backward_relu6.py | 2 +- .../unit_tests/operations/backward}/test_backward_round.py | 2 +- .../unit_tests/operations/backward}/test_backward_selu.py | 2 +- .../unit_tests/operations/backward}/test_backward_silu.py | 2 +- 7 files changed, 7 insertions(+), 7 deletions(-) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_abs.py (87%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_floor.py (87%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_log.py (94%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_relu6.py (87%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_round.py (87%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_selu.py (87%) rename tests/{tt_eager/python_api_testing/unit_testing/backward_ops => ttnn/unit_tests/operations/backward}/test_backward_silu.py (87%) diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py b/tests/ttnn/unit_tests/operations/backward/test_backward_abs.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_abs.py index 02b5481f7ec..4da39c4a436 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_abs.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_abs.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py b/tests/ttnn/unit_tests/operations/backward/test_backward_floor.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_floor.py index c9efe6c5e86..bedf8ea2fef 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_floor.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_floor.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py b/tests/ttnn/unit_tests/operations/backward/test_backward_log.py similarity index 94% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_log.py index 3418ea67dfa..1fcec1f3958 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_log.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_log.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import ( +from tests.ttnn.unit_tests.operations.backward.utility_funcs import ( data_gen_with_val, compare_pcc, data_gen_with_range, diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py b/tests/ttnn/unit_tests/operations/backward/test_backward_relu6.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_relu6.py index 15aff1a6855..21e72e32a28 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_relu6.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_relu6.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py b/tests/ttnn/unit_tests/operations/backward/test_backward_round.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_round.py index baabffc0009..0ea2865a08c 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_round.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_round.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import compare_pcc, data_gen_with_range +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py b/tests/ttnn/unit_tests/operations/backward/test_backward_selu.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_selu.py index 4218ffa1524..7adc3c09ee1 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_selu.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_selu.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py b/tests/ttnn/unit_tests/operations/backward/test_backward_silu.py similarity index 87% rename from tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py rename to tests/ttnn/unit_tests/operations/backward/test_backward_silu.py index c86af349dbe..577a929682b 100644 --- a/tests/tt_eager/python_api_testing/unit_testing/backward_ops/test_backward_silu.py +++ b/tests/ttnn/unit_tests/operations/backward/test_backward_silu.py @@ -5,7 +5,7 @@ import torch import pytest import ttnn -from tests.tt_eager.python_api_testing.unit_testing.backward_ops.utility_funcs import data_gen_with_range, compare_pcc +from tests.ttnn.unit_tests.operations.backward.utility_funcs import data_gen_with_range, compare_pcc @pytest.mark.parametrize( From 75d99cbbd51965b323579da998474e8568453d5b Mon Sep 17 00:00:00 2001 From: mouliraj-mcw Date: Thu, 11 Jul 2024 14:59:38 +0000 Subject: [PATCH 11/11] #10071: Update CPP files --- docs/source/ttnn/ttnn/dependencies/tt_lib.rst | 4 -- .../op_library/backward/backward_ops.cpp | 41 ------------------- .../op_library/backward/backward_ops.hpp | 10 ----- .../tt_lib_bindings_tensor_backward_ops.cpp | 17 -------- .../device/unary_backward_op.cpp | 16 ++++---- 5 files changed, 8 insertions(+), 80 deletions(-) diff --git a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst index 8e59eb3a369..8f71224774a 100644 --- a/docs/source/ttnn/ttnn/dependencies/tt_lib.rst +++ b/docs/source/ttnn/ttnn/dependencies/tt_lib.rst @@ -814,8 +814,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.fill_bw -.. autofunction:: tt_lib.tensor.unary_sub_bw - .. autofunction:: tt_lib.tensor.complex_abs_bw .. autofunction:: tt_lib.tensor.lt_bw @@ -878,8 +876,6 @@ Backward Operations .. autofunction:: tt_lib.tensor.reciprocal_bw -.. autofunction:: tt_lib.tensor.rpow_bw - .. autofunction:: tt_lib.tensor.square_bw .. autofunction:: tt_lib.tensor.tanhshrink_bw diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp index 8f17edc8f92..7f878cc5bf8 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.cpp @@ -280,26 +280,6 @@ std::vector ne_bw(const Tensor& grad, const MemoryConfig& output_mem_con return operation::decorate_as_composite(__func__, _ne_bw)(grad, output_mem_config); } -std::vector _rsqrt_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - Tensor rsqrt_result = power(rsqrt(input, true, output_mem_config), 3, output_mem_config); - Tensor result = mul_unary(ttnn::multiply(grad, rsqrt_result, std::nullopt, output_mem_config), -0.5, output_mem_config); - float t_inf = std::numeric_limits::infinity(); - result = where(eqz(input, output_mem_config), t_inf, result, output_mem_config); - float t_nan = std::nanf(""); - result = where(ltz(input, output_mem_config), t_nan, result, output_mem_config); - result = where( - ttnn::logical_and(eqz(input, output_mem_config), eqz(grad, output_mem_config), std::nullopt, output_mem_config), - t_nan, - result, - output_mem_config); - grad_tensor.emplace_back(result); - return grad_tensor; -} -std::vector rsqrt_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _rsqrt_bw)(grad, input, output_mem_config); -} - // bw(expm1) = grad * expm1(input) + 1 std::vector _expm1_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; @@ -980,27 +960,6 @@ std::vector reciprocal_bw(const Tensor& grad, const Tensor& input, const return operation::decorate_as_composite(__func__, _reciprocal_bw)(grad, input, output_mem_config); } -std::vector _rpow_bw( - const Tensor& grad, const Tensor& input, float exponent, const MemoryConfig& output_mem_config) { - std::vector grad_tensor; - float t_nan = std::nanf(""); - Tensor grad_result = zeros_like(input, output_mem_config); - if (exponent != 0.0) { - grad_result = - ttnn::multiply(grad, - ttnn::multiply(pow(input, exponent - 1, output_mem_config), exponent, std::nullopt, output_mem_config), - std::nullopt, - output_mem_config); - grad_result = where(ttnn::ltz(input, output_mem_config), t_nan, grad_result, output_mem_config); - } - grad_tensor.emplace_back(grad_result); - return grad_tensor; -} -std::vector rpow_bw( - const Tensor& grad, const Tensor& input, float exponent, const MemoryConfig& output_mem_config) { - return operation::decorate_as_composite(__func__, _rpow_bw)(grad, input, exponent, output_mem_config); -} - // Autoformat support Tensor change_layout_to_tile(const Tensor& temp, const MemoryConfig& output_mem_config) { auto formatted_input_tensor = temp; diff --git a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp index 1e4b649a92c..905f3ad2adf 100644 --- a/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp +++ b/tt_eager/tt_dnn/op_library/backward/backward_ops.hpp @@ -101,10 +101,6 @@ std::vector> tanh_bw( std::vector fill_bw( const Tensor& grad, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector unary_sub_bw( - const Tensor& grad, - const Tensor& input, - const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); std::vector binary_le_bw( const Tensor& grad, @@ -279,12 +275,6 @@ std::vector reciprocal_bw( const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); -std::vector rpow_bw( - const Tensor& grad, - const Tensor& input, - float exponent, - const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG); - std::vector square_bw( const Tensor& grad, const Tensor& input, diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp index 5c7864f1a1b..64ab793d044 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_backward_ops.cpp @@ -791,23 +791,6 @@ namespace tt::tt_metal::detail{ "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" )doc"); - m_tensor.def("rpow_bw", &tt::tt_metal::rpow_bw, - py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("exponent").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( - Performs backward operations for rpow for the ``input`` and ``exponent`` with given ``grad`` - - Input tensors must have BFLOAT16 data type. - - Output tensor will have BFLOAT16 data type. - - .. csv-table:: - :header: "Argument", "Description", "Data type", "Valid range", "Required" - - "grad", "Gradient tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "input", "Tensor", "Tensor", "Tensor of shape [W, Z, Y, X]", "Yes" - "exponent", "exponent", "float", ">0.0", "Yes" - "output_mem_config", "Layout of tensor in TT Accelerator device memory banks", "MemoryConfig", "Default is interleaved in DRAM", "No" - )doc"); - m_tensor.def("square_bw", &tt::tt_metal::square_bw, py::arg("grad").noconvert(), py::arg("input").noconvert(), py::arg("output_mem_config").noconvert() = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, R"doc( Performs backward square operations on ``input`` tensors with given ``grad``. diff --git a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp index c9a05b8566e..bbcffb294c3 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/unary_backward/device/unary_backward_op.cpp @@ -358,14 +358,14 @@ std::vector _rpow_bw( std::vector _floor_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor t_zero = tt::tt_metal::zeros_like(grad, output_mem_config); + Tensor t_zero = ttnn::operations::creation::zeros_like(grad); grad_tensor.emplace_back(t_zero); return grad_tensor; } std::vector _round_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor t_zero = tt::tt_metal::zeros_like(grad, output_mem_config); + Tensor t_zero = ttnn::operations::creation::zeros_like(grad); grad_tensor.emplace_back(t_zero); return grad_tensor; } @@ -373,8 +373,8 @@ std::vector _round_bw(const Tensor& grad, const Tensor& input, const Mem std::vector _log_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; Tensor grad_a = ttnn::multiply(grad, ttnn::reciprocal(input, output_mem_config), std::nullopt, output_mem_config); - Tensor t_inf = tt::tt_metal::full_like(input, std::numeric_limits::infinity(), output_mem_config); - Tensor t_nan = tt::tt_metal::full_like(input, std::nanf(""), output_mem_config); + Tensor t_inf = ttnn::operations::creation::full_like(input, std::numeric_limits::infinity()); + Tensor t_nan = ttnn::operations::creation::full_like(input, std::nanf("")); grad_tensor.emplace_back(where( ttnn::eqz(input, output_mem_config), where( @@ -389,9 +389,9 @@ std::vector _log_bw(const Tensor& grad, const Tensor& input, const Memor std::vector _relu6_bw(const Tensor& grad, const Tensor& input, const MemoryConfig& output_mem_config) { std::vector grad_tensor; - Tensor zero_tensor = tt::tt_metal::zeros_like(input, output_mem_config); - Tensor one_tensor = tt::tt_metal::ones_like(input, output_mem_config); - Tensor six_tensor = tt::tt_metal::full_like(input, 6, output_mem_config); + Tensor zero_tensor = ttnn::operations::creation::zeros_like(input); + Tensor one_tensor = ttnn::operations::creation::ones_like(input); + Tensor six_tensor = ttnn::operations::creation::full_like(input, 6); Tensor grad_result = where(ttnn::le(input, zero_tensor, std::nullopt, output_mem_config), zero_tensor, six_tensor, output_mem_config); grad_result = where( @@ -423,7 +423,7 @@ std::vector _silu_bw(const Tensor& grad, const Tensor& input, const Memo std::vector grad_tensor; Tensor grad_sigmoid = ttnn::multiply(grad, ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config); Tensor add_sub = ttnn::add( - ttnn::multiply(ttnn::subtract(tt::tt_metal::full_like(input, 1.0f) , ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config), + ttnn::multiply(ttnn::subtract(ttnn::operations::creation::full_like(input, 1.0f) , ttnn::sigmoid(input, output_mem_config), std::nullopt, output_mem_config), input, std::nullopt, output_mem_config),