feat: pipeline scalar softmax and blackwell mma issue

src/main/scala/radiance/core/TensorCoreBlackwell.scala

@@ -126,7 +126,8 @@ class TensorCoreBlackwell(
   val bIndexReg = RegInit(0.U(log2Ceil(numBFragsPerGroup).W))
   val mGroupReg = RegInit(0.U(log2Ceil(numMGroups).W))
   val substepReg = RegInit(0.U(1.W))
-  val elemReg = RegInit(0.U(log2Ceil(numLanes).W))
+  val issueElemReg = RegInit(0.U(log2Ceil(numLanes).W))
+  val retireElemReg = RegInit(0.U(log2Ceil(numLanes).W))
   val waitCounter = RegInit(0.U(3.W))
 
   val aBuf = Reg(Vec(numAFragsPerSet, UInt(memWidth.W)))
@@ -205,8 +206,8 @@ class TensorCoreBlackwell(
     x((idx + 1) * 16 - 1, idx * 16)
   }
 
-  val elemM = if (numLanes == 4) elemReg(0, 0) else elemReg(1, 0)
-  val elemN = if (numLanes == 4) elemReg(1) else elemReg(2)
+  val elemM = if (numLanes == 4) issueElemReg(0, 0) else issueElemReg(1, 0)
+  val elemN = if (numLanes == 4) issueElemReg(1) else issueElemReg(2)
   dpu.io.in.valid := dpuInValid
   for (k <- 0 until 8) {
     dpu.io.in.bits.a(k) := (
@@ -223,7 +224,7 @@ class TensorCoreBlackwell(
     )
     dpu.io.in.bits.b(k) := Mux(elemN.asBool, halfWord(operandB, 8 + k), halfWord(operandB, k))
   }
-  dpu.io.in.bits.c := cWords(elemReg)
+  dpu.io.in.bits.c := cWords(issueElemReg)
   dpu.io.stall := false.B
   val dpuValid = dpu.io.out.valid
 
@@ -244,7 +245,8 @@ class TensorCoreBlackwell(
     bIndexReg := 0.U
     mGroupReg := 0.U
     substepReg := 0.U
-    elemReg := 0.U
+    issueElemReg := 0.U
+    retireElemReg := 0.U
     switch(io.initiate.bits.op) {
       is(Ops.bwgmma) { state := State.bwLoadAReq }
       is(Ops.tcgen05Cp) { state := State.cpRead }
@@ -313,24 +315,28 @@ class TensorCoreBlackwell(
 
   when(state === State.bwReadCResp) {
     cDataReg := io.tmemC.cRdata
-    elemReg  := 0.U
+    issueElemReg := 0.U
+    retireElemReg := 0.U
     state    := State.bwCompute
   }
 
   when(state === State.bwCompute) {
     dpuInValid := true.B
-    state := State.bwDpuResp
+    when(issueElemReg === (numLanes - 1).U) {
+      state := State.bwDpuResp
+    }.otherwise {
+      issueElemReg := issueElemReg + 1.U
+    }
   }
 
-  when(state === State.bwDpuResp) {
-    when(dpuValid) {
-      mmaDataReg(elemReg) := dpu.io.out.bits.data
-      when(elemReg === (numLanes - 1).U) {
-        state := State.bwWriteCReq
-      }.otherwise {
-        elemReg := elemReg + 1.U
-        state := State.bwCompute
-      }
+  when(dpuValid) {
+    assert(state === State.bwCompute || state === State.bwDpuResp,
+      "BWGMMA DPU response arrived outside the compute states")
+    mmaDataReg(retireElemReg) := dpu.io.out.bits.data
+    when(retireElemReg === (numLanes - 1).U) {
+      state := State.bwWriteCReq
+    }.otherwise {
+      retireElemReg := retireElemReg + 1.U
     }
   }
 

src/main/scala/radiance/tile/RadianceTile.scala

@@ -851,6 +851,9 @@ class RadianceTileModuleImp(outer: RadianceTile)
       core.io.tc_tmem_C_rready := DontCare
       core.io.tc_tmem_C_rdata := DontCare
       core.io.tc_tmem_C_wready := DontCare
+      core.io.sc_tmem_rready := DontCare
+      core.io.sc_tmem_rdata := DontCare
+      core.io.sc_tmem_wready := DontCare
     }
 
     def connectTensorBlackwell = {
@@ -885,59 +888,225 @@ class RadianceTileModuleImp(outer: RadianceTile)
         tcDData.foreach(_ := 0.U)
         tcDTag.foreach(_ := 0.U)
 
-        // TMEM matrix: one shared 2R1W SRAM.  read0 is operand A, read1 is C.
-        // Each warp owns 2KB: A tile and C tile are 1KB each.  The row count
-        // scales with the physical fragment width (16B for 4 lanes, 32B for 8).
+        // TMEM keeps the ISA-visible address space unified while storing the
+        // A and C halves in separate 1R1W arrays.  This avoids duplicating each
+        // bank for two read ports, and still allows common A-read/C-read pairs
+        // to proceed in parallel because they normally hit different arrays.
         val tmemBytesPerWarp = 2048
-        val tmemDepth = outer.numWarps * (tmemBytesPerWarp / outer.tcSmemSize)
-        val tmem = Module(new radiance.memory.TwoReadOneWriteSyncMem(
-          tmemDepth, UInt((outer.tcSmemSize * 8).W)))
+        val tmemFragsPerWarp = tmemBytesPerWarp / outer.tcSmemSize
+        val tmemFragsPerTile = tmemFragsPerWarp / 2
+        val tmemLogicalDepth = outer.numWarps * tmemFragsPerWarp
+        val tmemArrayDepth = outer.numWarps * tmemFragsPerTile
+        val tmemBanks = 4
+        val tmemBankBits = log2Ceil(tmemBanks)
+        val tmemFragAddrBits = log2Ceil(tmemFragsPerWarp)
+        val tmemTileAddrBits = log2Ceil(tmemFragsPerTile)
+        val tmemWarpAddrBits = log2Ceil(outer.numWarps)
+        val tmemPhysAddrBits = log2Ceil(tmemArrayDepth)
+        val tmemBankDepth = tmemArrayDepth / tmemBanks
+        require(isPow2(tmemBanks))
+        require(isPow2(tmemFragsPerWarp))
+        require(tmemFragsPerWarp == tmemFragsPerTile * 2)
+        require(tmemLogicalDepth <= (1 << tmemAddrBits))
+        require(tmemArrayDepth % tmemBanks == 0)
+        require(tmemPhysAddrBits > tmemBankBits)
+        val tmemA = Seq.fill(tmemBanks) {
+          Module(new radiance.memory.TwoPortSyncMem(
+            tmemBankDepth, UInt((outer.tcSmemSize * 8).W)))
+        }
+        val tmemC = Seq.fill(tmemBanks) {
+          Module(new radiance.memory.TwoPortSyncMem(
+            tmemBankDepth, UInt((outer.tcSmemSize * 8).W)))
+        }
 
-        val aReadArb = Module(new RRArbiter(UInt(tmemAddrBits.W), nTC))
-        val cReadArb = Module(new RRArbiter(UInt(tmemAddrBits.W), nTC))
+        class TmemReadReq extends Bundle {
+          val addr = UInt(tmemAddrBits.W)
+          val src = UInt(2.W)
+          val tc = UInt(log2Ceil(nTC max 2).W)
+        }
 
         class TmemWriteReq extends Bundle {
           val addr = UInt(tmemAddrBits.W)
           val data = UInt(tmemDataBits.W)
           val mask = UInt(tmemMaskBits.W)
-        }
-        val cWriteArb = Module(new RRArbiter(new TmemWriteReq, nTC))
-
-        (0 until nTC).foreach { tc =>
-          aReadArb.io.in(tc).valid := core.io.tc_tmem_A_ren(tc)
-          aReadArb.io.in(tc).bits := slice(core.io.tc_tmem_A_raddr, tmemAddrBits, tc)
-          cReadArb.io.in(tc).valid := core.io.tc_tmem_C_ren(tc)
-          cReadArb.io.in(tc).bits := slice(core.io.tc_tmem_C_raddr, tmemAddrBits, tc)
-          cWriteArb.io.in(tc).valid := core.io.tc_tmem_C_wen(tc)
-          cWriteArb.io.in(tc).bits.addr := slice(core.io.tc_tmem_C_waddr, tmemAddrBits, tc)
-          cWriteArb.io.in(tc).bits.data := slice(core.io.tc_tmem_C_wdata, tmemDataBits, tc)
-          cWriteArb.io.in(tc).bits.mask := slice(core.io.tc_tmem_C_mask, tmemMaskBits, tc)
+          val src = UInt(1.W)
+          val tc = UInt(log2Ceil(nTC max 2).W)
         }
 
-        aReadArb.io.out.ready := true.B
-        cReadArb.io.out.ready := true.B
-        cWriteArb.io.out.ready := true.B
+        def tmemIsC(addr: UInt): Bool = addr(tmemTileAddrBits)
+        def tmemPhysAddr(addr: UInt): UInt = {
+          val tileOffset = addr(tmemTileAddrBits - 1, 0)
+          if (tmemWarpAddrBits == 0) {
+            tileOffset
+          } else {
+            Cat(addr(tmemFragAddrBits + tmemWarpAddrBits - 1, tmemFragAddrBits), tileOffset)
+          }
+        }
+        def bank(addr: UInt): UInt = addr(tmemBankBits - 1, 0)
+        def row(addr: UInt): UInt = addr(tmemPhysAddrBits - 1, tmemBankBits)
 
-        tmem.io.ren0 := aReadArb.io.out.fire
-        tmem.io.raddr0 := aReadArb.io.out.bits
-        tmem.io.ren1 := cReadArb.io.out.fire
-        tmem.io.raddr1 := cReadArb.io.out.bits
-        tmem.io.wen := cWriteArb.io.out.fire
-        tmem.io.waddr := cWriteArb.io.out.bits.addr
-        tmem.io.wdata := cWriteArb.io.out.bits.data
-        tmem.io.mask := cWriteArb.io.out.bits.mask
+        val aReady = Wire(Vec(nTC, Bool()))
+        val cReady = Wire(Vec(nTC, Bool()))
+        val wReady = Wire(Vec(nTC, Bool()))
+        val scReadReady = Wire(Bool())
+        val scWriteReady = Wire(Bool())
+        aReady.foreach(_ := false.B)
+        cReady.foreach(_ := false.B)
+        wReady.foreach(_ := false.B)
+        scReadReady := false.B
+        scWriteReady := false.B
 
-        val aReadGrant = RegNext(Mux(aReadArb.io.out.fire, UIntToOH(aReadArb.io.chosen, nTC), 0.U(nTC.W)))
-        val cReadGrant = RegNext(Mux(cReadArb.io.out.fire, UIntToOH(cReadArb.io.chosen, nTC), 0.U(nTC.W)))
-        core.io.tc_tmem_A_rready := VecInit(aReadArb.io.in.map(_.fire)).asUInt
-        core.io.tc_tmem_C_rready := VecInit(cReadArb.io.in.map(_.fire)).asUInt
-        core.io.tc_tmem_C_wready := VecInit(cWriteArb.io.in.map(_.fire)).asUInt
+        val aReadGrant = Wire(Vec(tmemBanks, new TmemReadReq))
+        val cReadGrant = Wire(Vec(tmemBanks, new TmemReadReq))
+        val aReadValid = Wire(Vec(tmemBanks, Bool()))
+        val cReadValid = Wire(Vec(tmemBanks, Bool()))
+        val aWriteGrant = Wire(Vec(tmemBanks, new TmemWriteReq))
+        val cWriteGrant = Wire(Vec(tmemBanks, new TmemWriteReq))
+        val aWriteValid = Wire(Vec(tmemBanks, Bool()))
+        val cWriteValid = Wire(Vec(tmemBanks, Bool()))
+        aReadGrant.foreach(_ := 0.U.asTypeOf(new TmemReadReq))
+        cReadGrant.foreach(_ := 0.U.asTypeOf(new TmemReadReq))
+        aReadValid.foreach(_ := false.B)
+        cReadValid.foreach(_ := false.B)
+        aWriteGrant.foreach(_ := 0.U.asTypeOf(new TmemWriteReq))
+        cWriteGrant.foreach(_ := 0.U.asTypeOf(new TmemWriteReq))
+        aWriteValid.foreach(_ := false.B)
+        cWriteValid.foreach(_ := false.B)
+
+        (0 until tmemBanks).foreach { b =>
+          val readRequests = (0 until nTC).flatMap { tc =>
+            val aAddr = slice(core.io.tc_tmem_A_raddr, tmemAddrBits, tc)
+            val cAddr = slice(core.io.tc_tmem_C_raddr, tmemAddrBits, tc)
+            Seq(
+              (core.io.tc_tmem_A_ren(tc).asBool, aAddr, 0.U(2.W), tc.U),
+              (core.io.tc_tmem_C_ren(tc).asBool, cAddr, 1.U(2.W), tc.U)
+            )
+          } ++ Seq(
+            (core.io.sc_tmem_ren.asBool, core.io.sc_tmem_raddr, 2.U(2.W), 0.U)
+          )
+
+          var aReadUsed = false.B
+          var cReadUsed = false.B
+          readRequests.foreach { case (valid, addr, src, tc) =>
+            val physAddr = tmemPhysAddr(addr)
+            val isC = tmemIsC(addr)
+            val aGrant = valid && !isC && bank(physAddr) === b.U && !aReadUsed
+            val cGrant = valid && isC && bank(physAddr) === b.U && !cReadUsed
+            when(aGrant) {
+              aReadGrant(b).addr := physAddr
+              aReadGrant(b).src := src
+              aReadGrant(b).tc := tc
+            }
+            when(cGrant) {
+              cReadGrant(b).addr := physAddr
+              cReadGrant(b).src := src
+              cReadGrant(b).tc := tc
+            }
+            aReadUsed = aReadUsed || aGrant
+            cReadUsed = cReadUsed || cGrant
+            when(aGrant || cGrant) {
+              when(src === 0.U) { aReady(tc) := true.B }
+              when(src === 1.U) { cReady(tc) := true.B }
+              when(src === 2.U) { scReadReady := true.B }
+            }
+          }
+          aReadValid(b) := aReadUsed
+          cReadValid(b) := cReadUsed
+
+          var aWriteUsed = false.B
+          var cWriteUsed = false.B
+          (0 until nTC).foreach { tc =>
+            val addr = slice(core.io.tc_tmem_C_waddr, tmemAddrBits, tc)
+            val physAddr = tmemPhysAddr(addr)
+            val isC = tmemIsC(addr)
+            val valid = core.io.tc_tmem_C_wen(tc).asBool && bank(physAddr) === b.U
+            val aGrant = valid && !isC && !aWriteUsed
+            val cGrant = valid && isC && !cWriteUsed
+            when(aGrant) {
+              aWriteValid(b) := true.B
+              aWriteGrant(b).addr := physAddr
+              aWriteGrant(b).data := slice(core.io.tc_tmem_C_wdata, tmemDataBits, tc)
+              aWriteGrant(b).mask := slice(core.io.tc_tmem_C_mask, tmemMaskBits, tc)
+              aWriteGrant(b).src := 0.U
+              aWriteGrant(b).tc := tc.U
+              wReady(tc) := true.B
+            }
+            when(cGrant) {
+              cWriteValid(b) := true.B
+              cWriteGrant(b).addr := physAddr
+              cWriteGrant(b).data := slice(core.io.tc_tmem_C_wdata, tmemDataBits, tc)
+              cWriteGrant(b).mask := slice(core.io.tc_tmem_C_mask, tmemMaskBits, tc)
+              cWriteGrant(b).src := 0.U
+              cWriteGrant(b).tc := tc.U
+              wReady(tc) := true.B
+            }
+            aWriteUsed = aWriteUsed || aGrant
+            cWriteUsed = cWriteUsed || cGrant
+          }
+
+          val scWPhysAddr = tmemPhysAddr(core.io.sc_tmem_waddr)
+          val scWIsC = tmemIsC(core.io.sc_tmem_waddr)
+          val scWValid = core.io.sc_tmem_wen.asBool && bank(scWPhysAddr) === b.U
+          val scWAGrant = scWValid && !scWIsC && !aWriteUsed
+          val scWCGrant = scWValid && scWIsC && !cWriteUsed
+          when(scWAGrant) {
+            aWriteValid(b) := true.B
+            aWriteGrant(b).addr := scWPhysAddr
+            aWriteGrant(b).data := core.io.sc_tmem_wdata
+            aWriteGrant(b).mask := core.io.sc_tmem_mask
+            aWriteGrant(b).src := 1.U
+            aWriteGrant(b).tc := 0.U
+            scWriteReady := true.B
+          }
+          when(scWCGrant) {
+            cWriteValid(b) := true.B
+            cWriteGrant(b).addr := scWPhysAddr
+            cWriteGrant(b).data := core.io.sc_tmem_wdata
+            cWriteGrant(b).mask := core.io.sc_tmem_mask
+            cWriteGrant(b).src := 1.U
+            cWriteGrant(b).tc := 0.U
+            scWriteReady := true.B
+          }
+
+          tmemA(b).io.ren := aReadValid(b)
+          tmemA(b).io.raddr := row(aReadGrant(b).addr)
+          tmemA(b).io.wen := aWriteValid(b)
+          tmemA(b).io.waddr := row(aWriteGrant(b).addr)
+          tmemA(b).io.wdata := aWriteGrant(b).data
+          tmemA(b).io.mask := aWriteGrant(b).mask
+          tmemC(b).io.ren := cReadValid(b)
+          tmemC(b).io.raddr := row(cReadGrant(b).addr)
+          tmemC(b).io.wen := cWriteValid(b)
+          tmemC(b).io.waddr := row(cWriteGrant(b).addr)
+          tmemC(b).io.wdata := cWriteGrant(b).data
+          tmemC(b).io.mask := cWriteGrant(b).mask
+        }
+
+        val aReadGrantReg = RegNext(aReadGrant)
+        val cReadGrantReg = RegNext(cReadGrant)
+        val aReadValidReg = RegNext(aReadValid)
+        val cReadValidReg = RegNext(cReadValid)
+        core.io.tc_tmem_A_rready := aReady.asUInt
+        core.io.tc_tmem_C_rready := cReady.asUInt
+        core.io.tc_tmem_C_wready := wReady.asUInt
+        core.io.sc_tmem_rready := scReadReady.asUInt
+        core.io.sc_tmem_wready := scWriteReady.asUInt
         core.io.tc_tmem_A_rdata := VecInit((0 until nTC).map { tc =>
-          Mux(aReadGrant(tc), tmem.io.rdata0, 0.U(tmemDataBits.W))
+          VecInit((0 until tmemBanks).map { b =>
+            Mux(aReadValidReg(b) && aReadGrantReg(b).src === 0.U && aReadGrantReg(b).tc === tc.U, tmemA(b).io.rdata,
+              Mux(cReadValidReg(b) && cReadGrantReg(b).src === 0.U && cReadGrantReg(b).tc === tc.U, tmemC(b).io.rdata, 0.U(tmemDataBits.W)))
+          }).reduce(_ | _)
         }).asUInt
         core.io.tc_tmem_C_rdata := VecInit((0 until nTC).map { tc =>
-          Mux(cReadGrant(tc), tmem.io.rdata1, 0.U(tmemDataBits.W))
+          VecInit((0 until tmemBanks).map { b =>
+            Mux(aReadValidReg(b) && aReadGrantReg(b).src === 1.U && aReadGrantReg(b).tc === tc.U, tmemA(b).io.rdata,
+              Mux(cReadValidReg(b) && cReadGrantReg(b).src === 1.U && cReadGrantReg(b).tc === tc.U, tmemC(b).io.rdata, 0.U(tmemDataBits.W)))
+          }).reduce(_ | _)
         }).asUInt
+        core.io.sc_tmem_rdata := VecInit((0 until tmemBanks).map { b =>
+          Mux(aReadValidReg(b) && aReadGrantReg(b).src === 2.U, tmemA(b).io.rdata,
+            Mux(cReadValidReg(b) && cReadGrantReg(b).src === 2.U, tmemC(b).io.rdata, 0.U(tmemDataBits.W)))
+        }).reduce(_ | _)
 
         // port 2: SMEM B, one TL client per tensor core. RadianceSharedMem arbitrates them.
         (0 until nTC).foreach { tc =>
@@ -1025,6 +1194,9 @@ class RadianceTileModuleImp(outer: RadianceTile)
         core.io.tc_tmem_C_rready := DontCare
         core.io.tc_tmem_C_rdata := DontCare
         core.io.tc_tmem_C_wready := DontCare
+        core.io.sc_tmem_rready := DontCare
+        core.io.sc_tmem_rdata := DontCare
+        core.io.sc_tmem_wready := DontCare
       }
     }
 

src/main/scala/radiance/tile/VortexCore.scala

@@ -120,6 +120,15 @@ class VortexBundle(tile: RadianceTile)(implicit p: Parameters) extends CoreBundl
   val tc_tmem_C_waddr  = Output(UInt((numTensorCores * 9).W))
   val tc_tmem_C_wdata  = Output(UInt((numTensorCores * numLanes * 32).W))
   val tc_tmem_C_mask   = Output(UInt((numTensorCores * numLanes * 4).W))
+  val sc_tmem_ren      = Output(UInt(1.W))
+  val sc_tmem_rready   = Input(UInt(1.W))
+  val sc_tmem_raddr    = Output(UInt(9.W))
+  val sc_tmem_rdata    = Input(UInt((numLanes * 32).W))
+  val sc_tmem_wen      = Output(UInt(1.W))
+  val sc_tmem_wready   = Input(UInt(1.W))
+  val sc_tmem_waddr    = Output(UInt(9.W))
+  val sc_tmem_wdata    = Output(UInt((numLanes * 32).W))
+  val sc_tmem_mask     = Output(UInt((numLanes * 4).W))
 
   // FIXME: hardcoded
   val barrierIdBits = tile.barrierMasterNode.out(0)._2.barrierIdBits
@@ -204,6 +213,7 @@ class Vortex(tile: RadianceTile)(implicit p: Parameters)
   addResource("/vsrc/vortex/hw/rtl/core/VX_dispatch.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_dispatch_unit.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_dispatch_unit_sane.sv")
+  addResource("/vsrc/vortex/hw/rtl/core/VX_tmem_softmax_unit.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_execute.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_fetch.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_gather_unit.sv")
@@ -351,6 +361,7 @@ class Vortex(tile: RadianceTile)(implicit p: Parameters)
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_div.sv")
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_dpi.sv")
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_dsp.sv")
+  addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_exp.sv")
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_fma.sv")
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_ncomp.sv")
   addResource("/vsrc/vortex/hw/rtl/fpu/VX_fpu_rounding.sv")

src/test/scala/radiance/TensorCoreBlackwellTest.scala

@@ -283,6 +283,7 @@ class TensorCoreBlackwellTest extends AnyFlatSpec with ChiselScalatestTester {
 
       var pendingB = Option.empty[(BigInt, BigInt)]
       var sawWriteback = false
+      var cycles = 0
 
       for (_ <- 0 until 20000 if !sawWriteback) {
         // Drive TMEM reads/writes
@@ -306,11 +307,14 @@ class TensorCoreBlackwellTest extends AnyFlatSpec with ChiselScalatestTester {
           } else None
 
           c.clock.step()
+          cycles += 1
           pendingB = nextB
         }
       }
 
       assert(sawWriteback, "BWGMMA did not complete")
+      assert(cycles < 5000,
+        s"BWGMMA took $cycles cycles; fragment elements are not issuing back-to-back")
       c.io.writeback.bits.wid.expect(1.U)
       // Verify all 32 C frags in TMEM
       for (i <- 0 until 32) {