#0: fix and update test for typecast to compare after round

tenstorrent · Jan 22, 2024 · 865ac2a · 865ac2a
1 parent 05e6b0a
commit 865ac2a
Show file tree

Hide file tree

Showing 8 changed files with 38 additions and 45 deletions.
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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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<const DataType> 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<const DataType> 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<const DataType> 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
@@ -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<const DataType> 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<const DataType> 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<const DataType> output_dtype = std::nullopt);
 
 //unary assign
-inline Tensor assign(const Tensor& input, const MemoryConfig& output_mem_config = operation::DEFAULT_OUTPUT_MEMORY_CONFIG, std::optional<const DataType> 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<const DataType> 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
@@ -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<uint32_t> 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<uint32_t> 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
@@ -43,7 +43,7 @@ namespace tt::tt_metal::detail{
         detail::export_enum<BcastOpDim>(m_tensor);
 
         detail::bind_unary_op<true, true>(m_tensor, "clone", &clone, R"doc(  Returns a new tensor which is a new copy of input tensor ``{0}``.)doc");
-        detail::bind_unary_op<true, true>(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<true, true>(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<false, false, false>(m_tensor, "copy", &copy, R"doc(  Copies the elements from ``{0}`` into ``{1}``. ``{1}`` is modified in place.)doc");
         detail::bind_unary_op<true, true>(m_tensor, "assign", py::overload_cast<const Tensor&, const MemoryConfig&, std::optional<const DataType>>(&assign), R"doc(  Returns a new tensor which is a new copy of input tensor ``{0}``.)doc");