diff --git a/tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_copy.py b/tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_copy.py index 8f4d9bb727b..fd90ebbe4da 100644 --- a/tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_copy.py +++ b/tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_copy.py @@ -25,12 +25,12 @@ @pytest.mark.parametrize( "input_shapes", [ + [[1, 1, 1, 30]], # Single core + [[1, 1, 300, 380]], # multi core + [[1, 3, 320, 380]], # multi core [[1, 1, 32, 32]], # Single core - # [[1, 1, 300, 380]], # multi core - # [[1, 3, 320, 380]], # multi core - # [[1, 1, 32, 32]], # Single core - # [[1, 1, 320, 384]], # Multi core - # [[1, 3, 320, 384]], # Multi core + [[1, 1, 320, 384]], # Multi core + [[1, 3, 320, 384]], # Multi core ], ) @pytest.mark.parametrize( diff --git a/tests/tt_eager/python_api_testing/sweep_tests/pytorch_ops.py b/tests/tt_eager/python_api_testing/sweep_tests/pytorch_ops.py index 570571b22dd..a4d865056bc 100644 --- a/tests/tt_eager/python_api_testing/sweep_tests/pytorch_ops.py +++ b/tests/tt_eager/python_api_testing/sweep_tests/pytorch_ops.py @@ -35,7 +35,7 @@ def clone(x, *args, **kwargs): def typecast(x, *args, **kwargs): - return x.to(torch.float32) + return x.to(torch.float32).round() def move(x, *args, **kwargs): diff --git a/tests/tt_eager/python_api_testing/sweep_tests/run_pytorch_ci_tests.py b/tests/tt_eager/python_api_testing/sweep_tests/run_pytorch_ci_tests.py index 64d7868f0fb..6caa0d442d3 100644 --- a/tests/tt_eager/python_api_testing/sweep_tests/run_pytorch_ci_tests.py +++ b/tests/tt_eager/python_api_testing/sweep_tests/run_pytorch_ci_tests.py @@ -28,7 +28,7 @@ def run_single_pytorch_test( env="", plot_func=None, ): - assert test_name in op_map.keys() + assert test_name in op_map default_env_dict = {} # Get env variables from CLI diff --git a/tests/tt_eager/python_api_testing/sweep_tests/tt_lib_ops.py b/tests/tt_eager/python_api_testing/sweep_tests/tt_lib_ops.py index 9fcc8f3d916..fbf6717bb3e 100644 --- a/tests/tt_eager/python_api_testing/sweep_tests/tt_lib_ops.py +++ b/tests/tt_eager/python_api_testing/sweep_tests/tt_lib_ops.py @@ -112,7 +112,7 @@ def typecast( t0 = setup_tt_tensor(x, device, layout[0], input_mem_config[0], dtype[0]) t1 = ttl.tensor.typecast(t0, output_mem_config=output_mem_config) - return tt2torch_tensor(t1) + return tt2torch_tensor(t1).round() @setup_host_and_device diff --git a/tt_eager/tt_dnn/op_library/copy/copy_op.cpp b/tt_eager/tt_dnn/op_library/copy/copy_op.cpp index 3dbfebab470..00c4fb35fb8 100644 --- a/tt_eager/tt_dnn/op_library/copy/copy_op.cpp +++ b/tt_eager/tt_dnn/op_library/copy/copy_op.cpp @@ -78,6 +78,30 @@ tt::stl::reflection::Attributes Copy::attributes() const { }; } +Tensor copy(const Tensor& src_tensor, const Tensor& dst_tensor) { + operation::run(Copy{dst_tensor.memory_config(), dst_tensor.dtype()}, {src_tensor, dst_tensor}); + return dst_tensor; +} + +Tensor clone(const Tensor& input, const MemoryConfig& output_mem_config, std::optional output_dtype) { + return operation::run(Copy{output_mem_config, output_dtype.value_or(input.dtype())}, {input}).at(0); +} + +Tensor typecast(const Tensor& input_tensor, const MemoryConfig& output_mem_config, std::optional output_dtype ) { + return operation::run(Copy{output_mem_config, output_dtype.value_or(input_tensor.dtype())}, {input_tensor}).at(0); +} + +//unary assign +Tensor assign(const Tensor& input, const MemoryConfig& output_mem_config , std::optional output_dtype ) { + return operation::run(Copy{output_mem_config, output_dtype.value_or(input.dtype())}, {input}).at(0); +} + +// binary assign +Tensor assign(const Tensor& input_a, const Tensor& input_b) { + operation::run(Copy{input_b.memory_config(), input_b.dtype()}, {input_a, input_b}); + return input_b; +} + } // namespace tt_metal } // namespace tt diff --git a/tt_eager/tt_dnn/op_library/copy/copy_op.hpp b/tt_eager/tt_dnn/op_library/copy/copy_op.hpp index f4f994d21f3..11381c58890 100644 --- a/tt_eager/tt_dnn/op_library/copy/copy_op.hpp +++ b/tt_eager/tt_dnn/op_library/copy/copy_op.hpp @@ -37,29 +37,17 @@ struct Copy { operation::ProgramWithCallbacks copy_multi_core(const Tensor &input, const Tensor &output, bool backwards = false); operation::ProgramWithCallbacks copy_single_core(const Tensor &input, const Tensor &output, bool backwards = false); -inline Tensor copy(const Tensor& src_tensor, const Tensor& dst_tensor) { - operation::run(Copy{dst_tensor.memory_config(), dst_tensor.dtype()}, {src_tensor, dst_tensor}); - return dst_tensor; -} +Tensor copy(const Tensor& src_tensor, const Tensor& dst_tensor); -inline Tensor clone(const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, std::optional output_dtype = std::nullopt) { - return operation::run(Copy{output_mem_config, output_dtype.value_or(input.dtype())}, {input}).at(0); -} +Tensor clone(const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, std::optional output_dtype = std::nullopt); -inline Tensor typecast(const Tensor& input_tensor, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG) { - return operation::run(Copy{output_mem_config}, {input_tensor}).at(0); -} +Tensor typecast(const Tensor& input_tensor, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, std::optional output_dtype = std::nullopt); //unary assign -inline Tensor assign(const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, std::optional output_dtype = std::nullopt) { - return operation::run(Copy{output_mem_config, output_dtype.value_or(input.dtype())}, {input}).at(0); -} +Tensor assign(const Tensor& input, const MemoryConfig& output_mem_config, std::optional output_dtype = std::nullopt); // binary assign -inline Tensor assign(const Tensor& input_a, const Tensor& input_b) { - operation::run(Copy{input_b.memory_config(), input_b.dtype()}, {input_a, input_b}); - return input_b; -} +Tensor assign(const Tensor& input_a, const Tensor& input_b); } // namespace tt_metal diff --git a/tt_eager/tt_dnn/op_library/copy/single_core/copy_op_single_core.cpp b/tt_eager/tt_dnn/op_library/copy/single_core/copy_op_single_core.cpp index 1cb93fc98d2..744bcd70464 100644 --- a/tt_eager/tt_dnn/op_library/copy/single_core/copy_op_single_core.cpp +++ b/tt_eager/tt_dnn/op_library/copy/single_core/copy_op_single_core.cpp @@ -102,7 +102,6 @@ operation::ProgramWithCallbacks copy_single_core(const Tensor &input, const Tens core, tt_metal::WriterDataMovementConfig{.compile_args = writer_compile_time_args, .defines = kernel_defines}); -<<<<<<< HEAD if (convert_dtype) { vector compute_kernel_args = { num_units @@ -114,24 +113,6 @@ operation::ProgramWithCallbacks copy_single_core(const Tensor &input, const Tens tt_metal::ComputeConfig{.compile_args=compute_kernel_args} ); } -======= - bool fp32_dest_acc_en = false; - bool math_approx_mode = false; - vector compute_kernel_args = { - uint(num_units) // per_core_tile_cnt - }; - - auto eltwise_unary_kernel = tt_metal::CreateKernel( - program, - "tt_eager/tt_dnn/kernels/compute/eltwise_copy.cpp", - core, - tt_metal::ComputeConfig{ - .math_fidelity = MathFidelity::HiFi4, - .fp32_dest_acc_en = fp32_dest_acc_en, - .math_approx_mode = math_approx_mode, - .compile_args = compute_kernel_args - }); ->>>>>>> 8f2e54a4b... #0: Add support for typecast if (tilized) { SetRuntimeArgs( diff --git a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_dm_ops.cpp b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_dm_ops.cpp index 7f19e53b9f0..9568f2595a3 100644 --- a/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_dm_ops.cpp +++ b/tt_eager/tt_lib/csrc/tt_lib_bindings_tensor_dm_ops.cpp @@ -43,7 +43,7 @@ namespace tt::tt_metal::detail{ detail::export_enum(m_tensor); detail::bind_unary_op(m_tensor, "clone", &clone, R"doc( Returns a new tensor which is a new copy of input tensor ``{0}``.)doc"); - detail::bind_unary_op(m_tensor, "typecast", &typecast, R"doc( Returns a new tensor which is a new copy of input tensor with new datatype``{0}``.)doc"); + detail::bind_unary_op(m_tensor, "typecast", &typecast, R"doc( Returns a new tensor which is a typecast of input tensor with new datatype``{0}``.)doc"); detail::bind_binary_op(m_tensor, "copy", ©, R"doc( Copies the elements from ``{0}`` into ``{1}``. ``{1}`` is modified in place.)doc"); detail::bind_unary_op(m_tensor, "assign", py::overload_cast>(&assign), R"doc( Returns a new tensor which is a new copy of input tensor ``{0}``.)doc");