Fix Blackwell tensor retirement accounting

hw/rtl/VX_core_wrapper.sv

@@ -82,7 +82,7 @@ module Vortex import VX_gpu_pkg::*; #(
     output [2:0]   tc_a_bits_write,
     output [95:0]  tc_a_bits_address,
     output [3 * TC_TAG_WIDTH - 1:0] tc_a_bits_tag,
-    output [3 * 32 - 1:0] tc_a_bits_mask,
+    output [3 * (TC_DATA_WIDTH / 8) - 1:0] tc_a_bits_mask,
     output [3 * TC_DATA_WIDTH - 1:0] tc_a_bits_data,
     output [2:0]   tc_d_ready,
     input  [2:0]   tc_d_valid,

hw/rtl/VX_platform.vh

@@ -33,9 +33,9 @@
 
 `ifdef SYNTHESIS
 `define NUM_BARRIERS 8
-`define NUM_CORES 4
-`define NUM_THREADS 8
-`define NUM_WARPS 8
+`define NUM_CORES 1
+`define NUM_THREADS 4
+`define NUM_WARPS 4
 
 `define FPU_FPNEW
 // `define FIRESIM

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),
@@ -199,7 +201,9 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
         if (sync_launch_pending) begin
             if (sync_writeback_done) begin
                 commit_select_tensor = 1'b1;
-                metadata_deq = commit_if.ready;
+                // BWGMMA retires when it launches. Its later tensor writeback
+                // is a ghost commit and must not dequeue the next uop.
+                metadata_deq = commit_if.ready && !sync_launch_is_bwgmma;
                 commit_sync_done = commit_if.ready;
             end else if (sync_no_wb_done) begin
                 commit_select_tensor = 1'b0;
@@ -215,7 +219,7 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
                 writeback_ready = 1'b0;
                 commit_select_tensor = 1'b0;
                 sync_done = initiate_valid && initiate_ready;
-                metadata_deq = 1'b0;
+                metadata_deq = sync_done;
             end else if (tcgen05_cp || tcgen05_ld || tcgen05_st || tcgen05_cb) begin
                 writeback_ready = 1'b0;
                 commit_select_tensor = 1'b0;
@@ -245,6 +249,7 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
             commit_if.data.eop    = sync_launch_is_ld ? 1'b1 : writeback_last;
         end else begin
             commit_if.valid       = (sync_no_wb_done && metadata_valid) ||
+                                    (metadata_valid && !sync_launch_pending && bwgmma && initiate_ready) ||
                                     (metadata_valid && !sync_launch_pending && !tensor_launch_op &&
                                      (!tensor_wait_op || tensor_idle));
             commit_if.data.uuid   = execute_if_data_uuid;