From 3a8ff9490a04f2104ec3fd0a12c5cef049b49688 Mon Sep 17 00:00:00 2001 From: Zhongdi LUO Date: Thu, 16 Jul 2026 07:27:11 +0000 Subject: [PATCH] Fix Blackwell tensor retirement accounting --- hw/rtl/core/VX_tensor_blackwell_core.sv | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/hw/rtl/core/VX_tensor_blackwell_core.sv b/hw/rtl/core/VX_tensor_blackwell_core.sv index 52db3e37..0d9f21b5 100644 --- a/hw/rtl/core/VX_tensor_blackwell_core.sv +++ b/hw/rtl/core/VX_tensor_blackwell_core.sv @@ -201,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; @@ -217,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; @@ -247,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;