#16373: Move conv specific ops from tensor_utils to conv2d.

Signed-off-by: Nilaykumar Patel <nkpatel@tenstorrent.com>

tests/scripts/run_tt_eager.py

@@ -34,7 +34,7 @@
     TestEntry("tt_eager/tests/ops/test_bcast_op", "ops/test_bcast_op"),
     TestEntry("tt_eager/tests/ops/test_transpose_op", "ops/test_transpose_op"),
     TestEntry("tt_eager/tests/ops/test_sliding_window_ops", "ops/test_sliding_window_ops"),
-    TestEntry("tt_eager/tests/ops/test_tensor_utils", "ops/test_tensor_utils"),
+    TestEntry("tt_eager/tests/ops/test_conv_prepare_weights_and_biases", "ops/test_conv_prepare_weights_and_biases"),
     TestEntry("tt_eager/tests/ops/test_bmm_op", "ops/test_bmm_op"),
     void_for_bh(void_for_whb0(TestEntry("tt_eager/tests/ops/test_eltwise_unary_op", "ops/test_eltwise_unary_op"))),
     void_for_whb0(

tests/tt_eager/CMakeLists.txt

@@ -20,7 +20,7 @@ set(TT_EAGER_TESTS_OPS
     ops/test_sfpu.cpp
     ops/test_sliding_window_ops.cpp
     ops/test_fold_op.cpp
-    ops/test_tensor_utils.cpp
+    ops/test_conv_prepare_weights_and_biases.cpp
 )
 
 set(TT_EAGER_TESTS_TENSORS

tests/tt_eager/ops/test_tensor_utils.cpp → tests/tt_eager/ops/test_conv_prepare_weights_and_biases.cpp

@@ -5,7 +5,6 @@
 #include "common/assert.hpp"
 #include "common/bfloat16.hpp"
 #include "ttnn/cpp/ttnn/tensor/host_buffer/functions.hpp"
-#include "ttnn/cpp/ttnn/tensor/tensor_utils.hpp"
 #include "ttnn/cpp/ttnn/tensor/types.hpp"
 #include "ttnn/tensor/host_buffer/functions.hpp"
 #include "ttnn/tensor/host_buffer/types.hpp"
@@ -14,7 +13,7 @@
 #include "ttnn/operations/creation.hpp"
 #include "ttnn/operations/functions.hpp"
 #include "ttnn/tensor/types.hpp"
-#include "ttnn/tensor/tensor_utils.hpp"
+#include "ttnn/cpp/ttnn/operations/conv/conv2d/prepare_conv2d_weights.hpp"
 
 static std::vector<std::vector<bfloat16>> ref_weight_in = {
     {
@@ -452,8 +451,8 @@ static void test_convert_conv_weight_tensor_to_tiled_layout_block_sharded() {
         for (auto j = 0; j < input_buffer.size(); j++) {
             input_buffer[j] = ref_weight_in[i][j];
         }
-        auto output_tensor =
-            convert_conv_weight_tensor_to_tiled_layout_block_sharded(input_tensor, shards[i], DataType::BFLOAT16);
+        auto output_tensor = ttnn::operations::conv::convert_conv_weight_tensor_to_tiled_layout_block_sharded(
+            input_tensor, shards[i], DataType::BFLOAT16);
         auto out_buffer = owned_buffer::get_as<bfloat16>(output_tensor);
 
         TT_FATAL(compare_out_with_ref(out_buffer, ref_weight_out[i]) == 0, "Error");
@@ -465,8 +464,8 @@ static void test_convert_conv_bias_tensor_to_tiled_layout_block_sharded() {
     for (auto i = 0; i < bias_tensor_shape.size(); i++) {
         auto input_tensor = ttnn::random::random(bias_tensor_shape[i], DataType::BFLOAT16).to(Layout::ROW_MAJOR).cpu();
         auto input_buffer = owned_buffer::get_as<bfloat16>(input_tensor);
-        auto output_tensor =
-            convert_conv_bias_tensor_to_tiled_layout_block_sharded(input_tensor, shards[i], DataType::BFLOAT16);
+        auto output_tensor = ttnn::operations::conv::convert_conv_bias_tensor_to_tiled_layout_block_sharded(
+            input_tensor, shards[i], DataType::BFLOAT16);
         auto out_buffer = owned_buffer::get_as<bfloat16>(output_tensor);
         /* Expected output should be same as input buffer except some padding*/
         TT_FATAL(compare_out_with_ref(out_buffer, input_buffer) == 0, "Error");

ttnn/cpp/ttnn/operations/conv/conv2d/conv2d_utils.cpp

@@ -72,35 +72,6 @@ uint32_t find_closest_largest_divisor_with_num_padding(uint32_t num1, uint32_t n
     return divisor;
 }
 
-// Converts convolution weights to tilized 2d matrix layout.
-// Returns a new tensor with layout=Tile
-Tensor convert_conv_weight_tensor_to_tiled_layout(
-    const Tensor& conv_weight_tensor,
-    uint32_t in1_block_h,
-    uint32_t in1_block_w,
-    std::optional<DataType> output_dtype) {
-    return tt::tt_metal::convert_conv_weight_tensor_to_tiled_layout(
-        std::move(conv_weight_tensor), in1_block_h, in1_block_w, output_dtype);
-}
-
-// Converts convolution weights to tilized 2d matrix layout with special block height padding
-// Returns a new tensor with layout=Tile
-Tensor convert_conv_weight_tensor_to_special_padding_tiled_layout(
-    const Tensor& conv_weight_tensor,
-    uint32_t in1_block_h,
-    uint32_t in1_block_w,
-    std::optional<DataType> output_dtype) {
-    return tt::tt_metal::convert_conv_weight_tensor_to_special_padding_tiled_layout(
-        std::move(conv_weight_tensor), in1_block_h, in1_block_w, output_dtype);
-}
-
-// Converts convolution weights to grouped layout with padded zeros
-Tensor convert_conv_weight_tensor_to_grouped_layout(
-    const Tensor& conv_weight_tensor, uint32_t num_groups, DataType output_dtype) {
-    return tt::tt_metal::convert_conv_weight_tensor_to_grouped_layout(
-        std::move(conv_weight_tensor), num_groups, output_dtype);
-}
-
 ParallelConfig determine_parallel_config_non_tile_mul_width(
     const TensorMemoryLayout shard_layout,
     uint32_t batch_size,

ttnn/cpp/ttnn/operations/conv/conv2d/conv2d_utils.hpp

@@ -171,26 +171,6 @@ shard_or_reshard_tensor_if_required(
     bool auto_shard,
     bool is_non_tile_mul_width = false);
 
-// Converts convolution weights to tilized 2d matrix layout.
-// Returns a new tensor with layout=Tile
-Tensor convert_conv_weight_tensor_to_tiled_layout(
-    const Tensor& conv_weight_tensor,
-    uint32_t in1_block_h,
-    uint32_t in1_block_w,
-    std::optional<DataType> output_dtype = std::nullopt);
-
-// Converts convolution weights to tilized 2d matrix layout with special block height padding
-// Returns a new tensor with layout=Tile
-Tensor convert_conv_weight_tensor_to_special_padding_tiled_layout(
-    const Tensor& conv_weight_tensor,
-    uint32_t in1_block_h,
-    uint32_t in1_block_w,
-    std::optional<DataType> output_dtype = std::nullopt);
-
-// Converts convolution weights to grouped layout with padded zeros
-Tensor convert_conv_weight_tensor_to_grouped_layout(
-    const Tensor& conv_weight_tensor, uint32_t num_groups, DataType output_dtype);
-
 std::ostream& operator<<(std::ostream& os, const Conv2dConfig& config);
 
 }  // namespace operations::conv