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

Compare commits

merge into: wu-arch:nvidia-style-arch
Branches Tags
wu-arch:wu-architecture wu-arch:pre-wu-blackwell wu-arch:nvidia-style-arch wu-arch:wu-tmem-bank-scalar-softmax wu-arch:wu-blackwell-fp8-e4m3 wu-arch:wu-blackwell-4lane-tc wu-arch:master
...
pull from: wu-arch:pre-wu-blackwell
Branches Tags
wu-arch:pre-wu-blackwell wu-arch:nvidia-style-arch wu-arch:wu-tmem-bank-scalar-softmax wu-arch:wu-blackwell-fp8-e4m3 wu-arch:wu-blackwell-4lane-tc wu-arch:wu-architecture wu-arch:master

2 Commits
nvidia-sty ... pre-wu-bla

Author SHA1 Message Date
Zhongdi LUO
3a8ff9490a Fix Blackwell tensor retirement accounting 2026-07-16 07:27:11 +00:00
Zhongdi LUO
a0b08f4c81 fix: isolate unused Blackwell writeback lanes 2026-07-13 10:27:08 +00:00
1 changed files with 16 additions and 11 deletions
Expand all files Collapse all files

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

@@ -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;