`ifdef EXT_T_ENABLE
`include "VX_fpu_define.vh"
`include "VX_trace.vh"

module VX_tensor_blackwell_core_block import VX_gpu_pkg::*; #(
    parameter ISW,
    parameter FP16
) (
    input clk,
    input reset,

    VX_execute_if.slave execute_if,
    VX_tc_rf_if.master  regfile_if,
    VX_tc_bus_if.master tmem_if,
    // tmem_C is now a direct SRAM port (no TileLink)
    output logic                            tmem_C_wen,
    output logic                            tmem_C_ren,
    output logic [8:0]                      tmem_C_waddr,
    output logic [8:0]                      tmem_C_raddr,
    output logic [`NUM_THREADS*`XLEN-1:0]   tmem_C_wdata,
    output logic [`NUM_THREADS*`XLEN/8-1:0] tmem_C_mask,
    input  logic [`NUM_THREADS*`XLEN-1:0]   tmem_C_rdata,
    VX_tc_bus_if.master smem_B_if,
    VX_commit_if.master commit_if
);
    localparam NUM_LANES = `NUM_THREADS;
    localparam METADATA_QUEUE_DEPTH = 2;

    wire [`UUID_WIDTH-1:0]          execute_if_data_uuid;
    wire [`NW_WIDTH-1:0]            execute_if_data_wid;
    wire [NUM_LANES-1:0]            execute_if_data_tmask;
    wire [`INST_ALU_BITS-1:0]       execute_if_data_op_type;
    wire [`XLEN-1:0]                execute_if_data_PC;
    wire                            execute_if_data_wb;
    wire [`NR_BITS-1:0]             execute_if_data_rd;
    wire [NUM_LANES-1:0][`XLEN-1:0] execute_if_data_rs1;
    wire [NUM_LANES-1:0][`XLEN-1:0] execute_if_data_rs2;
    wire [NUM_LANES-1:0][`XLEN-1:0] execute_if_data_rs3;

    wire metadata_queue_full;
    wire metadata_queue_empty;
    assign execute_if.ready = !metadata_queue_full;

    logic metadata_deq;

    wire operand_enq_fire = execute_if.valid && execute_if.ready;
    wire enq = operand_enq_fire;
    wire deq = metadata_deq;

    localparam DATAW = `UUID_WIDTH + `NW_WIDTH + `NUM_THREADS + `INST_ALU_BITS + `XLEN + 1 +
                       `NR_BITS + (NUM_LANES * `XLEN) + (NUM_LANES * `XLEN) + (NUM_LANES * `XLEN);
    VX_fifo_queue #(
        .DATAW(DATAW),
        .DEPTH(METADATA_QUEUE_DEPTH)
    ) pending_uops (
        .clk(clk),
        .reset(reset),
        .push(enq),
        .pop(deq),
        .data_in({execute_if.data.uuid, execute_if.data.wid,
                  execute_if.data.tmask, execute_if.data.op_type, execute_if.data.PC,
                  execute_if.data.wb, execute_if.data.rd,
                  execute_if.data.rs1_data, execute_if.data.rs2_data, execute_if.data.rs3_data}),
        .data_out({execute_if_data_uuid, execute_if_data_wid,
                   execute_if_data_tmask, execute_if_data_op_type, execute_if_data_PC,
                   execute_if_data_wb, execute_if_data_rd,
                   execute_if_data_rs1, execute_if_data_rs2, execute_if_data_rs3}),
        .empty(metadata_queue_empty),
        `UNUSED_PIN(alm_empty),
        .full(metadata_queue_full),
        `UNUSED_PIN(alm_full),
        `UNUSED_PIN(size)
    );

    wire initiate_ready;
    wire writeback_valid;
    wire writeback_last;
    wire [`NW_WIDTH-1:0] writeback_wid;
    wire [4:0] writeback_rd;
    logic writeback_ready;
    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 &&
                  (execute_if_data_op_type == `INST_TENSOR_BWGMMA);
    wire bwgmma_wait = metadata_valid &&
                       (execute_if_data_op_type == `INST_TENSOR_BWGMMA_WAIT);
    wire tcgen05_cp = metadata_valid &&
                      (execute_if_data_op_type == `INST_TENSOR_TCGEN05_CP);
    wire tcgen05_cp_wait = metadata_valid &&
                           (execute_if_data_op_type == `INST_TENSOR_TCGEN05_CP_WAIT);
    wire tcgen05_ld = metadata_valid &&
                      (execute_if_data_op_type == `INST_TENSOR_TCGEN05_LD);
    wire tcgen05_st = metadata_valid &&
                      (execute_if_data_op_type == `INST_TENSOR_TCGEN05_ST);
    wire tcgen05_cb = metadata_valid &&
                      (execute_if_data_op_type == `INST_TENSOR_TCGEN05_CB);
    wire tensor_launch_op = bwgmma || tcgen05_cp || tcgen05_ld || tcgen05_st || tcgen05_cb;
    wire tensor_wait_op = bwgmma_wait || tcgen05_cp_wait;
    wire [2:0] initiate_op =
        bwgmma    ? 3'd0 :
        bwgmma_wait ? 3'd1 :
        tcgen05_cp ? 3'd2 :
        tcgen05_cp_wait ? 3'd3 :
        tcgen05_ld ? 3'd4 :
        tcgen05_st ? 3'd5 :
        tcgen05_cb ? 3'd6 :
        3'd0;
    wire initiate_valid = metadata_valid && tensor_launch_op && !sync_launch_pending && commit_if.ready;
    wire [`NW_WIDTH-1:0] initiate_wid = execute_if_data_wid;
    wire [4:0] initiate_rd = execute_if_data_rd[4:0];
    wire [`XLEN-1:0] initiate_addr_a = execute_if_data_rs1[0];
    wire [`XLEN-1:0] initiate_addr_b = execute_if_data_rs2[0];
    wire [`XLEN-1:0] initiate_addr_c = execute_if_data_rs3[0];

    logic sync_launch_pending;
    logic sync_launch_is_ld;
    logic sync_launch_is_bwgmma;

    TensorCoreBlackwell tensor_blackwell_core (
        .clock(clk),
        .reset(reset),

        .io_initiate_ready(initiate_ready),
        .io_initiate_valid(initiate_valid),
        .io_initiate_bits_op(initiate_op),
        .io_initiate_bits_wid(initiate_wid),
        .io_initiate_bits_rd(initiate_rd),
        .io_initiate_bits_addressA(initiate_addr_a),
        .io_initiate_bits_addressB(initiate_addr_b),
        .io_initiate_bits_addressC(initiate_addr_c),

        .io_writeback_ready(writeback_ready),
        .io_writeback_valid(writeback_valid),
        .io_writeback_bits_last(writeback_last),
        .io_writeback_bits_wid(writeback_wid),
        .io_writeback_bits_rd(writeback_rd),
        .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),
        .io_respA_bits_source(tmem_if.rsp_data.tag),
        .io_respA_bits_data(tmem_if.rsp_data.data),
        .io_respB_ready(smem_B_if.rsp_ready),
        .io_respB_valid(smem_B_if.rsp_valid),
        .io_respB_bits_source(smem_B_if.rsp_data.tag),
        .io_respB_bits_data(smem_B_if.rsp_data.data),
        .io_respC(regfile_if.rsp_data.data),

        .io_reqA_ready(tmem_if.req_ready),
        .io_reqA_valid(tmem_if.req_valid),
        .io_reqA_bits_rw(tmem_if.req_data.rw),
        .io_reqA_bits_byteen(tmem_if.req_data.byteen),
        .io_reqA_bits_source(tmem_if.req_data.tag),
        .io_reqA_bits_address(tmem_if.req_data.addr),
        .io_reqA_bits_data(tmem_if.req_data.data),
        .io_tmemC_wen(tmem_C_wen),
        .io_tmemC_ren(tmem_C_ren),
        .io_tmemC_waddr(tmem_C_waddr),
        .io_tmemC_raddr(tmem_C_raddr),
        .io_tmemC_wdata(tmem_C_wdata),
        .io_tmemC_mask(tmem_C_mask),
        .io_tmemC_rdata(tmem_C_rdata),
        .io_reqB_ready(smem_B_if.req_ready),
        .io_reqB_valid(smem_B_if.req_valid),
        .io_reqB_bits_rw(smem_B_if.req_data.rw),
        .io_reqB_bits_byteen(smem_B_if.req_data.byteen),
        .io_reqB_bits_source(smem_B_if.req_data.tag),
        .io_reqB_bits_address(smem_B_if.req_data.addr),
        .io_reqB_bits_data(smem_B_if.req_data.data),
        .io_reqC_valid(regfile_if.req_valid),
        .io_reqC_bits(regfile_if.req_data.rs[4:0])
    );

    assign regfile_if.req_data.rs[5] = 1'b1;
    assign regfile_if.req_data.wis = '0;

    logic commit_select_tensor;
    logic commit_sync_done;
    logic sync_done;
    wire sync_writeback_done = sync_launch_pending && (sync_launch_is_ld || sync_launch_is_bwgmma) && writeback_valid;
    wire sync_no_wb_done = sync_launch_pending && !(sync_launch_is_ld || sync_launch_is_bwgmma) && initiate_ready;
    wire tensor_idle = initiate_ready && !writeback_valid && !sync_launch_pending;

    always @(*) begin
        metadata_deq = 1'b0;
        commit_select_tensor = 1'b0;
        commit_sync_done = 1'b0;
        sync_done = 1'b0;
        writeback_ready = commit_if.ready && sync_launch_pending && (sync_launch_is_ld || sync_launch_is_bwgmma);

        if (sync_launch_pending) begin
            if (sync_writeback_done) begin
                commit_select_tensor = 1'b1;
                metadata_deq = commit_if.ready;
                commit_sync_done = commit_if.ready;
            end else if (sync_no_wb_done) begin
                commit_select_tensor = 1'b0;
                metadata_deq = metadata_valid && commit_if.ready;
                commit_sync_done = metadata_valid && commit_if.ready;
            end
        end else if (metadata_valid) begin
            if (tensor_wait_op) begin
                writeback_ready = 1'b0;
                commit_select_tensor = 1'b0;
                metadata_deq = metadata_valid && commit_if.ready && tensor_idle;
            end else if (bwgmma) begin
                writeback_ready = 1'b0;
                commit_select_tensor = 1'b0;
                sync_done = initiate_valid && initiate_ready;
                metadata_deq = 1'b0;
            end else if (tcgen05_cp || tcgen05_ld || tcgen05_st || tcgen05_cb) begin
                writeback_ready = 1'b0;
                commit_select_tensor = 1'b0;
                sync_done = initiate_valid && initiate_ready;
                // suppress commit until sync_launch_pending is set next cycle
                metadata_deq = 1'b0;
            end else begin
                commit_select_tensor = 1'b0;
                metadata_deq = metadata_valid && commit_if.ready;
            end
        end else begin
            commit_select_tensor = 1'b1;
        end

        if (commit_select_tensor) begin
            commit_if.valid       = sync_writeback_done;
            commit_if.data.uuid   = sync_launch_is_bwgmma ? '0 : execute_if_data_uuid;
            commit_if.data.wid    = sync_launch_is_bwgmma ? writeback_wid : execute_if_data_wid;
            commit_if.data.tmask  = sync_launch_is_bwgmma ? {NUM_LANES{1'b1}} : execute_if_data_tmask;
            commit_if.data.PC     = sync_launch_is_bwgmma ? '0 : execute_if_data_PC;
            commit_if.data.wb     = sync_launch_is_ld;
            commit_if.data.rd     = sync_launch_is_ld ? execute_if_data_rd : '0;
            commit_if.data.data   = writeback_data;
            commit_if.data.tensor = sync_launch_is_bwgmma;
            commit_if.data.pid    = 1'b0;
            commit_if.data.sop    = 1'b1;
            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 && !tensor_launch_op &&
                                     (!tensor_wait_op || tensor_idle));
            commit_if.data.uuid   = execute_if_data_uuid;
            commit_if.data.wid    = execute_if_data_wid;
            commit_if.data.tmask  = execute_if_data_tmask;
            commit_if.data.PC     = execute_if_data_PC;
            commit_if.data.wb     = execute_if_data_wb;
            commit_if.data.rd     = execute_if_data_rd;
            commit_if.data.data   = '0;
            commit_if.data.tensor = 1'b0;
            commit_if.data.pid    = 1'b0;
            commit_if.data.sop    = 1'b1;
            commit_if.data.eop    = 1'b1;
        end
    end

    always @(posedge clk) begin
        if (reset) begin
            sync_launch_pending <= 1'b0;
            sync_launch_is_ld <= 1'b0;
            sync_launch_is_bwgmma <= 1'b0;
        end else begin
            if (sync_done) begin
                sync_launch_pending <= 1'b1;
                sync_launch_is_ld <= tcgen05_ld;
                sync_launch_is_bwgmma <= bwgmma;
            end else if (commit_sync_done) begin
                sync_launch_pending <= 1'b0;
                sync_launch_is_ld <= 1'b0;
                sync_launch_is_bwgmma <= 1'b0;
            end
        end
    end

`ifdef SIMULATION
`ifdef DBG_TRACE_CORE_PIPELINE_VCS
    always @(posedge clk) begin
        if (!reset && ($time > `TRACE_STARTTIME)) begin
            if (operand_enq_fire) begin
                `TRACE(2, ("%d: tensor-blackwell-enq: wid=%0d, PC=0x%0h, op=0x%0h, uuid=%0d, queue_full=%b\n",
                    $time, execute_if.data.wid, execute_if.data.PC, execute_if.data.op_type,
                    execute_if.data.uuid, metadata_queue_full));
            end
            if (metadata_valid) begin
                `TRACE(2, ("%d: tensor-blackwell-meta: wid=%0d, PC=0x%0h, op=0x%0h, deq=%b, commit_valid=%b, commit_ready=%b, initiate_valid=%b, initiate_ready=%b, writeback_valid=%b, writeback_ready=%b, sync_pending=%b, sync_is_ld=%b, sync_done=%b, commit_sync_done=%b, queue_full=%b\n",
                    $time, execute_if_data_wid, execute_if_data_PC, execute_if_data_op_type,
                    metadata_deq, commit_if.valid, commit_if.ready, initiate_valid, initiate_ready,
                    writeback_valid, writeback_ready, sync_launch_pending, sync_launch_is_ld,
                    sync_done, commit_sync_done, metadata_queue_full));
            end
            if (writeback_valid || initiate_valid) begin
                `TRACE(2, ("%d: tensor-blackwell-handshake: initiate=%b/%b, op=%0d, addrA=0x%0h, addrB=0x%0h, writeback=%b/%b, wb_wid=%0d, wb_rd=%0d, wb_last=%b\n",
                    $time, initiate_valid, initiate_ready, initiate_op, initiate_addr_a, initiate_addr_b,
                    writeback_valid, writeback_ready, writeback_wid, writeback_rd, writeback_last));
            end
            if (tmem_if.req_valid || tmem_C_wen || tmem_C_ren || smem_B_if.req_valid || tmem_if.rsp_valid || smem_B_if.rsp_valid) begin
                `TRACE(2, ("%d: tensor-blackwell-mem: tmemA_req=%b/%b tmemA_rw=%b tmemA_addr=0x%0h tmemA_rsp=%b/%b, tmemC_wen=%b tmemC_ren=%b tmemC_waddr=0x%0h tmemC_raddr=0x%0h, smemB_req=%b/%b smemB_rw=%b smemB_addr=0x%0h smemB_rsp=%b/%b, initiate_ready=%b\n",
                    $time, tmem_if.req_valid, tmem_if.req_ready, tmem_if.req_data.rw, tmem_if.req_data.addr, tmem_if.rsp_valid, tmem_if.rsp_ready,
                    tmem_C_wen, tmem_C_ren, tmem_C_waddr, tmem_C_raddr,
                    smem_B_if.req_valid, smem_B_if.req_ready, smem_B_if.req_data.rw, smem_B_if.req_data.addr, smem_B_if.rsp_valid, smem_B_if.rsp_ready,
                    initiate_ready));
            end
        end
    end
`endif
`endif
endmodule

`endif
