#8632: Update fp32 dest acc support in moreh_sum_nc

tenstorrent · Jun 1, 2024 · bfa47be · bfa47be
1 parent 80965ba
commit bfa47be
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 1 deletion.
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
@@ -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
@@ -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
         });
     ////////////////////////////////////////////////////////////////////////////