wu-arch/vortex
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fix Blackwell tensor retirement accounting

Browse Source
This commit is contained in:
Zhongdi LUO
2026-07-16 07:27:11 +00:00
parent a0b08f4c81
commit 3a8ff9490a
1 changed files with 5 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
hw/rtl/core/VX_tensor_blackwell_core.sv
View File

@@ -201,7 +201,9 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
if (sync_launch_pending) begin if (sync_launch_pending) begin
if (sync_writeback_done) begin if (sync_writeback_done) begin
commit_select_tensor = 1'b1; 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; commit_sync_done = commit_if.ready;
end else if (sync_no_wb_done) begin end else if (sync_no_wb_done) begin
commit_select_tensor = 1'b0; commit_select_tensor = 1'b0;
@@ -217,7 +219,7 @@ module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
writeback_ready = 1'b0; writeback_ready = 1'b0;
commit_select_tensor = 1'b0; commit_select_tensor = 1'b0;
sync_done = initiate_valid && initiate_ready; 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 end else if (tcgen05_cp || tcgen05_ld || tcgen05_st || tcgen05_cb) begin
writeback_ready = 1'b0; writeback_ready = 1'b0;
commit_select_tensor = 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; commit_if.data.eop = sync_launch_is_ld ? 1'b1 : writeback_last;
end else begin end else begin
commit_if.valid = (sync_no_wb_done && metadata_valid) || 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 && (metadata_valid && !sync_launch_pending && !tensor_launch_op &&
(!tensor_wait_op || tensor_idle)); (!tensor_wait_op || tensor_idle));
commit_if.data.uuid = execute_if_data_uuid; commit_if.data.uuid = execute_if_data_uuid;