diff --git a/ttnn/cpp/ttnn/operations/data_movement/reshape_view/reshape.cpp b/ttnn/cpp/ttnn/operations/data_movement/reshape_view/reshape.cpp index 8fb90cbc761..904fdf88a1f 100644 --- a/ttnn/cpp/ttnn/operations/data_movement/reshape_view/reshape.cpp +++ b/ttnn/cpp/ttnn/operations/data_movement/reshape_view/reshape.cpp @@ -52,36 +52,6 @@ ttnn::Tensor convert_tile_to_rm( (tensor.get_dtype() == DataType::BFLOAT8_B) ? ttnn::typecast(new_tensor, tensor.get_dtype()) : new_tensor; return new_tensor; } -ttnn::Tensor host_reshape(const ttnn::Tensor& tensor, const ttnn::Shape& shape) { - //This function is due to embedding issue 15558, once the issue is fixed we want to delete it - tt::log_warning("host_reshape is deprecated and will be removed in the near future"); - if (!ttnn::has_storage_type_of(tensor, ttnn::StorageType::DEVICE)) { - return tensor.reshape(shape); - } - auto tensor_shape = tensor.shape(); - auto layout = tensor.layout(); - auto device = tensor.device(); - auto memory_config = tensor.memory_config(); - auto host_tensor = tensor.cpu(); - auto rm_tensor = ttnn::to_layout(host_tensor, ttnn::ROW_MAJOR_LAYOUT, std::nullopt, std::nullopt, (Device*)nullptr); - - if (tensor_shape.has_tile_padding()) { - ttnn::Tensor slice_input; - auto host_tensor_4d = unsqueeze_to_4D(rm_tensor); - auto tensor_shape_4d = host_tensor_4d.shape(); - ttnn::SmallVector begins({0, 0, 0, 0}); - ttnn::SmallVector ends( - {tensor_shape_4d[0], tensor_shape_4d[1], tensor_shape_4d[2], tensor_shape_4d[3]}); - ttnn::SmallVector step({1, 1, 1, 1}); - host_tensor_4d = ttnn::slice(host_tensor_4d, begins, ends, step, std::nullopt); - host_tensor = squeeze_from_4D(host_tensor_4d, tensor_shape.rank()); - } - auto host_reshape_tensor = rm_tensor.reshape(shape); - auto final_layout_tensor = - ttnn::to_layout(host_reshape_tensor, layout, std::nullopt, std::nullopt, (Device*)nullptr); - auto device_tensor = ttnn::data_transfer_to_device(final_layout_tensor, device, memory_config); - return device_tensor; -} //Wrapper to turn the ND-> MD problem into 3D->3D for tiled and 2D->2D for Row Major @@ -399,7 +369,7 @@ ttnn::Tensor ReshapeViewOperation::invoke( return tensor.reshape(shape); } //This is a completely incorrect test but it is due to issue 15558 - return detail::host_reshape(tensor, shape); + TT_FATAL(false, "Attempting to reshape between two shapes with different volumes"); } // Catch-all // Do the reshape in row-major diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/binary_ng_program_factory.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/binary_ng_program_factory.cpp index 6c710e3d193..bce5cfdc824 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/binary_ng_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/binary_ng_program_factory.cpp @@ -302,14 +302,14 @@ BinaryNgDeviceOperation::ProgramFactory::cached_program_t BinaryNgDeviceOperatio // How many tiles to store per input CB (double buffer) constexpr uint32_t num_tiles_per_cb = 2; auto [a_cb, a_cb_handle] = - create_cb(tt::CB::c_in0, program, all_device_cores, a_single_tile_size, num_tiles_per_cb, a_data_format); + create_cb(tt::CBIndex::c_0, program, all_device_cores, a_single_tile_size, num_tiles_per_cb, a_data_format); auto [c_cb, c_cb_handle] = - create_cb(tt::CB::c_out0, program, all_device_cores, c_single_tile_size, num_tiles_per_cb, c_data_format); + create_cb(tt::CBIndex::c_2, program, all_device_cores, c_single_tile_size, num_tiles_per_cb, c_data_format); // If b is a scalar, we only need one tile in the CB uint32_t b_num_tiles_per_cb = b_buffer != nullptr ? num_tiles_per_cb : 1; auto [b_cb, b_cb_handle] = - create_cb(tt::CB::c_in1, program, all_device_cores, b_single_tile_size, b_num_tiles_per_cb, b_data_format); + create_cb(tt::CBIndex::c_1, program, all_device_cores, b_single_tile_size, b_num_tiles_per_cb, b_data_format); auto a_is_dram = static_cast(a_buffer->buffer_type() == tt_metal::BufferType::DRAM); bool b_is_dram = false; diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary.cpp index fdfdd2f650f..dff89bfd613 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary.cpp @@ -41,9 +41,9 @@ void MAIN { return; } - constexpr auto cb_in0 = tt::CB::c_in0; - constexpr auto cb_in1 = tt::CB::c_in1; - constexpr auto cb_out0 = tt::CB::c_out0; + constexpr auto cb_in0 = tt::CBIndex::c_0; + constexpr auto cb_in1 = tt::CBIndex::c_1; + constexpr auto cb_out0 = tt::CBIndex::c_2; #if BCAST_INPUT auto cb_bcast = cb_in1; diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_no_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_no_bcast.cpp index 6275e672c49..1e3703fa4b8 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_no_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_no_bcast.cpp @@ -10,9 +10,9 @@ namespace NAMESPACE { void MAIN { uint32_t num_tiles = get_arg_val(0); - constexpr auto cb_in0 = tt::CB::c_in0; - constexpr auto cb_in1 = tt::CB::c_in1; - constexpr auto cb_out0 = tt::CB::c_out0; + constexpr auto cb_in0 = tt::CBIndex::c_0; + constexpr auto cb_in1 = tt::CBIndex::c_1; + constexpr auto cb_out0 = tt::CBIndex::c_2; binary_op_init_common(cb_in0, cb_in1, cb_out0); add_tiles_init(); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_scalar.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_scalar.cpp index 4e57b6aba57..2b377fb52ca 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_scalar.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/compute/eltwise_binary_scalar.cpp @@ -10,9 +10,9 @@ namespace NAMESPACE { void MAIN { uint32_t num_tiles = get_arg_val(0); - constexpr auto cb_in0 = tt::CB::c_in0; - constexpr auto cb_in1 = tt::CB::c_in1; - constexpr auto cb_out0 = tt::CB::c_out0; + constexpr auto cb_in0 = tt::CBIndex::c_0; + constexpr auto cb_in1 = tt::CBIndex::c_1; + constexpr auto cb_out0 = tt::CBIndex::c_2; binary_op_init_common(cb_in0, cb_in1, cb_out0); add_tiles_init(); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_col_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_col_bcast.cpp index aa98c9575d9..cf0a1e4e13d 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_col_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_col_bcast.cpp @@ -21,7 +21,7 @@ void kernel_main() { constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; - constexpr auto cb_id_src = tt::CB::c_in0; + constexpr auto cb_id_src = tt::CBIndex::c_0; constexpr uint32_t onetile = 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_no_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_no_bcast.cpp index 36cbaaf3b0b..2379206448a 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_no_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_no_bcast.cpp @@ -18,7 +18,7 @@ void kernel_main() { constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; - constexpr auto cb_id_src = tt::CB::c_in0; + constexpr auto cb_id_src = tt::CBIndex::c_0; constexpr uint32_t onetile = 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_row_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_row_bcast.cpp index c5eadbe61d7..7278e6b3510 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_row_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_row_bcast.cpp @@ -21,7 +21,7 @@ void kernel_main() { constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; - constexpr auto cb_id_src = tt::CB::c_in0; + constexpr auto cb_id_src = tt::CBIndex::c_0; constexpr uint32_t onetile = 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_scalar_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_scalar_bcast.cpp index 6d23f3ac2f2..d1f0922d648 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_scalar_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/reader_interleaved_scalar_bcast.cpp @@ -20,7 +20,7 @@ void kernel_main() { constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; - constexpr auto cb_id_src = tt::CB::c_in0; + constexpr auto cb_id_src = tt::CBIndex::c_0; constexpr uint32_t onetile = 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_col_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_col_bcast.cpp index 4c8d2337e17..f17b23684ff 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_col_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_col_bcast.cpp @@ -22,7 +22,7 @@ void kernel_main() { constexpr uint32_t onetile = 1; - constexpr auto cb_id_src = tt::CB::c_in1; + constexpr auto cb_id_src = tt::CBIndex::c_1; constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); const DataFormat src_data_format = get_dataformat(cb_id_src); @@ -30,7 +30,7 @@ void kernel_main() { const InterleavedAddrGenFast src = { .bank_base_address = src_addr, .page_size = src_tile_bytes, .data_format = src_data_format}; - constexpr auto cb_id_dst = tt::CB::c_out0; + constexpr auto cb_id_dst = tt::CBIndex::c_2; constexpr bool dst_is_dram = get_compile_time_arg_val(1) == 1; const uint32_t dst_tile_bytes = get_tile_size(cb_id_dst); const DataFormat dst_data_format = get_dataformat(cb_id_dst); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_no_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_no_bcast.cpp index 0ac36437c03..4ef2f809e8e 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_no_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_no_bcast.cpp @@ -20,7 +20,7 @@ void kernel_main() { constexpr uint32_t onetile = 1; - constexpr auto cb_id_src = tt::CB::c_in1; + constexpr auto cb_id_src = tt::CBIndex::c_1; constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); const DataFormat src_data_format = get_dataformat(cb_id_src); @@ -28,7 +28,7 @@ void kernel_main() { const InterleavedAddrGenFast src = { .bank_base_address = src_addr, .page_size = src_tile_bytes, .data_format = src_data_format}; - constexpr auto cb_id_dst = tt::CB::c_out0; + constexpr auto cb_id_dst = tt::CBIndex::c_2; constexpr bool dst_is_dram = get_compile_time_arg_val(1) == 1; const uint32_t dst_tile_bytes = get_tile_size(cb_id_dst); const DataFormat dst_data_format = get_dataformat(cb_id_dst); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_row_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_row_bcast.cpp index f02dc62537b..65ff8e60f69 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_row_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_row_bcast.cpp @@ -22,7 +22,7 @@ void kernel_main() { constexpr uint32_t onetile = 1; - constexpr auto cb_id_src = tt::CB::c_in1; + constexpr auto cb_id_src = tt::CBIndex::c_1; constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); const DataFormat src_data_format = get_dataformat(cb_id_src); @@ -30,7 +30,7 @@ void kernel_main() { const InterleavedAddrGenFast src = { .bank_base_address = src_addr, .page_size = src_tile_bytes, .data_format = src_data_format}; - constexpr auto cb_id_dst = tt::CB::c_out0; + constexpr auto cb_id_dst = tt::CBIndex::c_2; constexpr bool dst_is_dram = get_compile_time_arg_val(1) == 1; const uint32_t dst_tile_bytes = get_tile_size(cb_id_dst); const DataFormat dst_data_format = get_dataformat(cb_id_dst); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar.cpp index 4aa94dac0ab..452d1dafaa7 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar.cpp @@ -21,8 +21,8 @@ void kernel_main() { constexpr bool dst_is_dram = get_compile_time_arg_val(1) == 1; - constexpr auto cb_id_src = tt::CB::c_in1; - constexpr auto cb_id_dst = tt::CB::c_out0; + constexpr auto cb_id_src = tt::CBIndex::c_1; + constexpr auto cb_id_dst = tt::CBIndex::c_2; constexpr uint32_t onetile = 1; const uint32_t dst_tile_bytes = get_tile_size(cb_id_dst); diff --git a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar_bcast.cpp b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar_bcast.cpp index d0e866c7c50..18915373d25 100644 --- a/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar_bcast.cpp +++ b/ttnn/cpp/ttnn/operations/eltwise/binary_ng/device/kernels/dataflow/writer_interleaved_scalar_bcast.cpp @@ -20,7 +20,7 @@ void kernel_main() { constexpr uint32_t onetile = 1; - constexpr auto cb_id_src = tt::CB::c_in1; + constexpr auto cb_id_src = tt::CBIndex::c_1; constexpr bool src_is_dram = get_compile_time_arg_val(0) == 1; const uint32_t src_tile_bytes = get_tile_size(cb_id_src); const DataFormat src_data_format = get_dataformat(cb_id_src); @@ -28,7 +28,7 @@ void kernel_main() { const InterleavedAddrGenFast src = { .bank_base_address = src_addr, .page_size = src_tile_bytes, .data_format = src_data_format}; - constexpr auto cb_id_dst = tt::CB::c_out0; + constexpr auto cb_id_dst = tt::CBIndex::c_2; constexpr bool dst_is_dram = get_compile_time_arg_val(1) == 1; const uint32_t dst_tile_bytes = get_tile_size(cb_id_dst); const DataFormat dst_data_format = get_dataformat(cb_id_dst);