diff --git a/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/kernels/moreh_sum_nc.cpp b/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/kernels/moreh_sum_nc.cpp index d2648770e9d..92af5a4f9dc 100644 --- a/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/kernels/moreh_sum_nc.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/kernels/moreh_sum_nc.cpp @@ -34,6 +34,9 @@ void MAIN { } tile_regs_acquire(); + #if defined FP32_DEST_ACC_EN + unpack_reconfig_data_format(cb_in0, cb_add); + #endif add_tiles_init(cb_in0, cb_add); add_tiles(cb_in0, cb_add, first_tile, first_tile, dst0); tile_regs_commit(); @@ -46,6 +49,9 @@ void MAIN { uint32_t cb_out = (last_out) ? (cb_out0) : (cb_intermed0); cb_reserve_back(cb_out, onetile); tile_regs_wait(); + #if defined FP32_DEST_ACC_EN + pack_reconfig_data_format(cb_out); + #endif pack_tile(dst0, cb_out); tile_regs_release(); cb_push_back(cb_out, onetile); diff --git a/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/moreh_sum_nc_impl.cpp b/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/moreh_sum_nc_impl.cpp index 31272c94922..07795d1f861 100644 --- a/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/moreh_sum_nc_impl.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_sum/moreh_sum_nc_impl/moreh_sum_nc_impl.cpp @@ -97,7 +97,7 @@ operation::ProgramWithCallbacks moreh_sum_nc_impl(const Tensor &input, const Ten { {CB::c_in0, in0_t}, // input {CB::c_in1, in1_t}, // zero - {CB::c_intermed0, intermed0_t}, + {CB::c_intermed0, intermed0_t, (fp32_dest_acc_en) ? tt::DataFormat::Float32: cb_data_format}, {CB::c_out0, out0_t}, // output }); ////////////////////////////////////////////////////////////////////////////