diff --git a/hw/rtl/core/VX_tensor_blackwell_core.sv b/hw/rtl/core/VX_tensor_blackwell_core.sv index 30ee6c4a..52db3e37 100644 --- a/hw/rtl/core/VX_tensor_blackwell_core.sv +++ b/hw/rtl/core/VX_tensor_blackwell_core.sv @@ -78,7 +78,9 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #( wire [`NW_WIDTH-1:0] writeback_wid; wire [4:0] writeback_rd; logic writeback_ready; - wire [`NUM_THREADS-1:0][`XLEN-1:0] writeback_data; + wire [7:0][`XLEN-1:0] tensor_writeback_data; + wire [`NUM_THREADS-1:0][`XLEN-1:0] writeback_data = + tensor_writeback_data[`NUM_THREADS-1:0]; wire metadata_valid = !metadata_queue_empty; wire bwgmma = metadata_valid && @@ -135,14 +137,14 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #( .io_writeback_bits_last(writeback_last), .io_writeback_bits_wid(writeback_wid), .io_writeback_bits_rd(writeback_rd), - .io_writeback_bits_data_0(writeback_data[0]), - .io_writeback_bits_data_1(writeback_data[1]), - .io_writeback_bits_data_2(writeback_data[2]), - .io_writeback_bits_data_3(writeback_data[3]), - .io_writeback_bits_data_4(writeback_data[4]), - .io_writeback_bits_data_5(writeback_data[5]), - .io_writeback_bits_data_6(writeback_data[6]), - .io_writeback_bits_data_7(writeback_data[7]), + .io_writeback_bits_data_0(tensor_writeback_data[0]), + .io_writeback_bits_data_1(tensor_writeback_data[1]), + .io_writeback_bits_data_2(tensor_writeback_data[2]), + .io_writeback_bits_data_3(tensor_writeback_data[3]), + .io_writeback_bits_data_4(tensor_writeback_data[4]), + .io_writeback_bits_data_5(tensor_writeback_data[5]), + .io_writeback_bits_data_6(tensor_writeback_data[6]), + .io_writeback_bits_data_7(tensor_writeback_data[7]), .io_respA_ready(tmem_if.rsp_ready), .io_respA_valid(tmem_if.rsp_valid),