fix: configure tensor DPU precision by architecture

.gitmodules

@@ -1,3 +1,3 @@
 [submodule "src/main/resources/vsrc/vortex"]
 	path = src/main/resources/vsrc/vortex
-	url = https://git.nudt.space/wu-arch/vortex.git
+	url = https://github.com/hansungk/vortex.git

radiance.mk

@@ -17,14 +17,14 @@ RADIANCE_VSRC_DIR = $(base_dir)/generators/radiance/src/main/resources/vsrc
 ifeq ($(shell echo $(CONFIG) | grep -E "SynConfig$$"),$(CONFIG))
     EXTRA_SIM_PREPROC_DEFINES += +define+SYNTHESIS +define+NDEBUG +define+DPI_DISABLE
 endif
-ifeq ($(shell echo $(CONFIG) | grep -E "FP16Config$$"),$(CONFIG))
+ifeq ($(shell echo $(CONFIG) | grep -E "(FP16|Volta|Ampere)Config$$"),$(CONFIG))
-    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=8
+    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=1 +define+TENSOR_DPU_FP16
 endif
 ifeq ($(shell echo $(CONFIG) | grep -E "HopperConfig$$"),$(CONFIG))
-    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=4 +define+EXT_T_HOPPER
+    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=1 +define+EXT_T_HOPPER +define+TENSOR_DPU_FP16
 endif
 ifeq ($(shell echo $(CONFIG) | grep -E "BlackwellConfig$$"),$(CONFIG))
-    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=1 +define+NUM_WARPS=4 +define+NUM_THREADS=4 +define+NUM_TENSOR_WARPS=2 +define+EXT_T_BLACKWELL
+    EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=1 +define+EXT_T_BLACKWELL +define+TENSOR_DPU_FP16
 endif
 ifeq ($(shell echo $(CONFIG) | grep -E "FlashConfig$$"),$(CONFIG))
     EXTRA_SIM_PREPROC_DEFINES += +define+NUM_CORES=4

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

@@ -14,7 +14,8 @@ class TensorCoreBlackwell(
     val numFPRegs: Int = 32
 ) extends Module {
   require(half, "Blackwell MMA currently supports FP16 inputs only")
-  require(numLanes == 8, "Blackwell MMA currently assumes 8 lanes")
+  require(numLanes == 4 || numLanes == 8,
+    s"Blackwell MMA currently supports 4 or 8 lanes, got ${numLanes}")
 
   val numWarpBits = log2Ceil(numWarps)
   val sourceWidth = log2Ceil(numSourceIds)
@@ -26,11 +27,16 @@ class TensorCoreBlackwell(
   val fragOffsetBits = log2Ceil(memWidth / 8)
 
   val numSets = 4
-  val numAFragsPerSet = 8
   val numBGroups = 4
-  val numBFragsPerGroup = 2
+  val numSubsteps = 2
-  val numMGroups = 4
+  val mElemsPerFrag = if (numLanes == 4) 2 else 4
-  val numCFrags = 32
+  val numMGroups = 16 / mElemsPerFrag
+  val numAFragsPerMGroup = 2
+  val numAFragsPerSet = numMGroups * numAFragsPerMGroup
+  val numBFragsPerSubstep = if (numLanes == 4) 2 else 1
+  val numBFragsPerGroup = numSubsteps * numBFragsPerSubstep
+  val numBFragsPerSet = numBGroups * numBFragsPerGroup
+  val numCFrags = numBGroups * numMGroups * numSubsteps
 
   object Ops {
     val bwgmma :: bwgmmaWait :: tcgen05Cp :: tcgen05CpWait :: tcgen05Ld :: tcgen05St :: tcgen05Cb :: Nil = Enum(7)
@@ -57,21 +63,13 @@ class TensorCoreBlackwell(
 
   // Direct SRAM port for TMEM (no TileLink overhead)
   class TmemSramPort extends Bundle {
-    val aRen    = Output(Bool())
+    val wen   = Output(Bool())
-    val aRready = Input(Bool())
+    val ren   = Output(Bool())
-    val aRaddr  = Output(UInt(log2Ceil(numWarps * numCFrags * 2).W))
+    val waddr = Output(UInt(log2Ceil(numWarps * numCFrags * 2).W))
-    val aRdata  = Input(UInt(memWidth.W))
+    val raddr = Output(UInt(log2Ceil(numWarps * numCFrags * 2).W))
-
+    val wdata = Output(UInt(memWidth.W))
-    val cRen    = Output(Bool())
+    val mask  = Output(UInt(maskWidth.W))
-    val cRready = Input(Bool())
+    val rdata = Input(UInt(memWidth.W))
-    val cRaddr  = Output(UInt(log2Ceil(numWarps * numCFrags * 2).W))
-    val cRdata  = Input(UInt(memWidth.W))
-
-    val cWen    = Output(Bool())
-    val cWready = Input(Bool())
-    val cWaddr  = Output(UInt(log2Ceil(numWarps * numCFrags * 2).W))
-    val cWdata  = Output(UInt(memWidth.W))
-    val cMask   = Output(UInt(maskWidth.W))
   }
 
   val io = IO(new Bundle {
@@ -102,7 +100,7 @@ class TensorCoreBlackwell(
     val idle, bwLoadAReq, bwLoadAResp, bwLoadBReq, bwLoadBResp,
         bwReadCReq, bwReadCResp, bwCompute, bwDpuResp, bwWriteCReq,
         bwWriteCWait, bwDone, cpRead, cpWrite, ldReq, stReq, stWrite, waitWb,
-        cbRead, cbCapture, cbWrite = Value
+        cbRead, cbWrite = Value
   }
   val state = RegInit(State.idle)
 
@@ -136,10 +134,11 @@ class TensorCoreBlackwell(
     base + (fragIndex << fragOffsetBits).asUInt
   }
 
-  val aFragIndex = (setReg << 3) + aIndexReg
+  val aFragIndex = (setReg * numAFragsPerSet.U) + aIndexReg
-  val bFragIndex = (setReg << 3) + (bGroupReg << 1) + bIndexReg
+  val bFragIndex =
-  val stepIndex = Cat(bGroupReg, mGroupReg)
+    (setReg * numBFragsPerSet.U) + (bGroupReg * numBFragsPerGroup.U) + bIndexReg
-  val cFragIndex = (stepIndex << 1) + substepReg
+  val cFragIndex =
+    (((bGroupReg * numMGroups.U) + mGroupReg) * numSubsteps.U) + substepReg
   val aReqAddress = byteAddress(addrAReg, aFragIndex)
   val bReqAddress = byteAddress(addrBReg, bFragIndex)
   val cReqAddress = byteAddress(addrCReg, cFragIndex)
@@ -155,14 +154,12 @@ class TensorCoreBlackwell(
   io.reqA <> reqA
   io.reqB <> reqB
 
-  io.tmemC.aRen   := false.B
+  io.tmemC.wen   := false.B
-  io.tmemC.aRaddr := 0.U
+  io.tmemC.ren   := false.B
-  io.tmemC.cRen   := false.B
+  io.tmemC.waddr := 0.U
-  io.tmemC.cRaddr := 0.U
+  io.tmemC.raddr := 0.U
-  io.tmemC.cWen   := false.B
+  io.tmemC.wdata := 0.U
-  io.tmemC.cWaddr := 0.U
+  io.tmemC.mask  := 0.U
-  io.tmemC.cWdata := 0.U
-  io.tmemC.cMask  := 0.U
 
   val wbValid = RegInit(false.B)
   val wbData = Reg(Vec(numLanes, UInt(laneWidth.W)))
@@ -181,7 +178,12 @@ class TensorCoreBlackwell(
   io.initiate.ready := state === State.idle && !wbValid
 
   val operandA = Cat(aBuf((mGroupReg << 1) + 1.U), aBuf(mGroupReg << 1))
-  val operandB = bBuf(substepReg)
+  val operandB =
+    if (numLanes == 4) {
+      Cat(bBuf((substepReg << 1) + 1.U), bBuf(substepReg << 1))
+    } else {
+      bBuf(substepReg)
+    }
   val cWords = cDataReg.asTypeOf(Vec(numLanes, UInt(laneWidth.W)))
   val dpuInValid = WireDefault(false.B)
   val dpu = Module(new TensorDotProductUnit(
@@ -193,16 +195,22 @@ class TensorCoreBlackwell(
     x((idx + 1) * 16 - 1, idx * 16)
   }
 
-  val elemM = elemReg(1, 0)
+  val elemM = if (numLanes == 4) elemReg(0, 0) else elemReg(1, 0)
-  val elemN = elemReg(2)
+  val elemN = if (numLanes == 4) elemReg(1) else elemReg(2)
   dpu.io.in.valid := dpuInValid
   for (k <- 0 until 8) {
-    dpu.io.in.bits.a(k) := MuxLookup(elemM, halfWord(operandA, k))(Seq(
+    dpu.io.in.bits.a(k) := (
-      0.U -> halfWord(operandA, k),
+      if (numLanes == 4) {
-      1.U -> halfWord(operandA, 8 + k),
+        Mux(elemM.asBool, halfWord(operandA, 8 + k), halfWord(operandA, k))
-      2.U -> halfWord(operandA, 16 + k),
+      } else {
-      3.U -> halfWord(operandA, 24 + k)
+        MuxLookup(elemM, halfWord(operandA, k))(Seq(
-    ))
+          0.U -> halfWord(operandA, k),
+          1.U -> halfWord(operandA, 8 + k),
+          2.U -> halfWord(operandA, 16 + k),
+          3.U -> halfWord(operandA, 24 + k)
+        ))
+      }
+    )
     dpu.io.in.bits.b(k) := Mux(elemN.asBool, halfWord(operandB, 8 + k), halfWord(operandB, k))
   }
   dpu.io.in.bits.c := cWords(elemReg)
@@ -239,15 +247,13 @@ class TensorCoreBlackwell(
   }
 
   when(state === State.bwLoadAReq) {
-    io.tmemC.aRen   := true.B
+    io.tmemC.ren   := true.B
-    io.tmemC.aRaddr := tmemABase + aFragIndex
+    io.tmemC.raddr := tmemABase + aFragIndex
-    when(io.tmemC.aRready) {
+    state := State.bwLoadAResp
-      state := State.bwLoadAResp
-    }
   }
 
   when(state === State.bwLoadAResp) {
-    aBuf(aIndexReg) := io.tmemC.aRdata
+    aBuf(aIndexReg) := io.tmemC.rdata
     when(aIndexReg === (numAFragsPerSet - 1).U) {
       bGroupReg := 0.U
       bIndexReg := 0.U
@@ -286,15 +292,13 @@ class TensorCoreBlackwell(
   }
 
   when(state === State.bwReadCReq) {
-    io.tmemC.cRen   := true.B
+    io.tmemC.ren   := true.B
-    io.tmemC.cRaddr := tmemCBase + cFragIndex
+    io.tmemC.raddr := tmemCBase + cFragIndex
-    when(io.tmemC.cRready) {
+    state := State.bwReadCResp
-      state := State.bwReadCResp
-    }
   }
 
   when(state === State.bwReadCResp) {
-    cDataReg := io.tmemC.cRdata
+    cDataReg := io.tmemC.rdata
     elemReg  := 0.U
     state    := State.bwCompute
   }
@@ -317,36 +321,34 @@ class TensorCoreBlackwell(
   }
 
   when(state === State.bwWriteCReq) {
-    io.tmemC.cWen   := true.B
+    io.tmemC.wen   := true.B
-    io.tmemC.cWaddr := tmemCBase + cFragIndex
+    io.tmemC.waddr := tmemCBase + cFragIndex
-    io.tmemC.cWdata := mmaDataReg.asUInt
+    io.tmemC.wdata := mmaDataReg.asUInt
-    io.tmemC.cMask  := Fill(maskWidth, 1.U(1.W))
+    io.tmemC.mask  := Fill(maskWidth, 1.U(1.W))
-    when(io.tmemC.cWready) {
+    when(substepReg === 0.U) {
-      when(substepReg === 0.U) {
+      substepReg := 1.U
-        substepReg := 1.U
+      state := State.bwReadCReq
-        state := State.bwReadCReq
+    }.elsewhen(mGroupReg =/= (numMGroups - 1).U) {
-      }.elsewhen(mGroupReg =/= (numMGroups - 1).U) {
+      substepReg := 0.U
-        substepReg := 0.U
+      mGroupReg := mGroupReg + 1.U
-        mGroupReg := mGroupReg + 1.U
+      state := State.bwReadCReq
-        state := State.bwReadCReq
+    }.elsewhen(bGroupReg =/= (numBGroups - 1).U) {
-      }.elsewhen(bGroupReg =/= (numBGroups - 1).U) {
+      substepReg := 0.U
-        substepReg := 0.U
+      mGroupReg := 0.U
-        mGroupReg := 0.U
+      bGroupReg := bGroupReg + 1.U
-        bGroupReg := bGroupReg + 1.U
+      bIndexReg := 0.U
-        bIndexReg := 0.U
+      state := State.bwLoadBReq
-        state := State.bwLoadBReq
+    }.elsewhen(setReg =/= (numSets - 1).U) {
-      }.elsewhen(setReg =/= (numSets - 1).U) {
+      substepReg := 0.U
-        substepReg := 0.U
+      mGroupReg := 0.U
-        mGroupReg := 0.U
+      bGroupReg := 0.U
-        bGroupReg := 0.U
+      bIndexReg := 0.U
-        bIndexReg := 0.U
+      setReg := setReg + 1.U
-        setReg := setReg + 1.U
+      aIndexReg := 0.U
-        aIndexReg := 0.U
+      state := State.bwLoadAReq
-        state := State.bwLoadAReq
+    }.otherwise {
-      }.otherwise {
+      waitCounter := 7.U
-        waitCounter := 7.U
+      state := State.bwWriteCWait
-        state := State.bwWriteCWait
-      }
     }
   }
 
@@ -377,26 +379,24 @@ class TensorCoreBlackwell(
   }
 
   when(state === State.cpWrite) {
-    io.respA.ready := io.tmemC.cWready
+    io.respA.ready := true.B
-    io.tmemC.cWen   := io.respA.valid
-    io.tmemC.cWaddr := (addrAReg >> fragOffsetBits.U).asUInt
-    io.tmemC.cWdata := io.respA.bits.data
-    io.tmemC.cMask  := Fill(maskWidth, 1.U(1.W))
     when(io.respA.fire) {
+      io.tmemC.wen   := true.B
+      io.tmemC.waddr := (addrAReg >> fragOffsetBits.U).asUInt
+      io.tmemC.wdata := io.respA.bits.data
+      io.tmemC.mask  := Fill(maskWidth, 1.U(1.W))
       state := State.idle
     }
   }
 
   when(state === State.ldReq) {
-    io.tmemC.cRen   := true.B
+    io.tmemC.ren   := true.B
-    io.tmemC.cRaddr := (addrAReg >> fragOffsetBits.U).asUInt
+    io.tmemC.raddr := (addrAReg >> fragOffsetBits.U).asUInt
-    when(io.tmemC.cRready) {
+    state := State.waitWb
-      state := State.waitWb
-    }
   }
 
   when(state === State.waitWb && opReg === Ops.tcgen05Ld) {
-    wbData  := io.tmemC.cRdata.asTypeOf(Vec(numLanes, UInt(laneWidth.W)))
+    wbData  := io.tmemC.rdata.asTypeOf(Vec(numLanes, UInt(laneWidth.W)))
     wbValid := true.B
     state   := State.idle
   }
@@ -407,25 +407,16 @@ class TensorCoreBlackwell(
   }
 
   when(state === State.stWrite) {
-    io.tmemC.cWen   := true.B
+    io.tmemC.wen   := true.B
-    io.tmemC.cWaddr := (addrAReg >> fragOffsetBits.U).asUInt
+    io.tmemC.waddr := (addrAReg >> fragOffsetBits.U).asUInt
-    io.tmemC.cWdata := io.respC
+    io.tmemC.wdata := io.respC
-    io.tmemC.cMask  := Fill(maskWidth, 1.U(1.W))
+    io.tmemC.mask  := Fill(maskWidth, 1.U(1.W))
-    when(io.tmemC.cWready) {
+    state := State.idle
-      state := State.idle
-    }
   }
 
   when(state === State.cbRead) {
-    io.tmemC.cRen   := true.B
+    io.tmemC.ren   := true.B
-    io.tmemC.cRaddr := (addrAReg >> fragOffsetBits.U).asUInt
+    io.tmemC.raddr := (addrAReg >> fragOffsetBits.U).asUInt
-    when(io.tmemC.cRready) {
-      state := State.cbCapture
-    }
-  }
-
-  when(state === State.cbCapture) {
-    cDataReg := io.tmemC.cRdata
     state := State.cbWrite
   }
 
@@ -435,7 +426,7 @@ class TensorCoreBlackwell(
     reqA.bits.byteen := Fill(maskWidth, 1.U(1.W))
     reqA.bits.address := addrBReg
     reqA.bits.source := sourceCounter
-    reqA.bits.data := cDataReg
+    reqA.bits.data := io.tmemC.rdata
     when(reqA.fire) {
       bumpSource()
       state := State.waitWb

src/main/scala/radiance/core/TensorCoreDecoupled.scala

@@ -47,8 +47,16 @@ class TensorCoreDecoupled(
     val numSourceIds: Int = 16,
     val numFPRegs: Int = 32
 ) extends Module {
+  require(numLanes == 4 || numLanes == 8,
+    s"Hopper tensor core supports 4 or 8 lanes, got ${numLanes}")
   val tilingParams =
-    if (half) TensorTilingParams.fp16 else TensorTilingParams.fp32
+    if (half && numLanes == 4) {
+      TensorTilingParams(m = 16, n = 16, k = 32, mc = 4, nc = 2, kc = 8)
+    } else if (half) {
+      TensorTilingParams.fp16
+    } else {
+      TensorTilingParams.fp32
+    }
   val numWarpBits = log2Ceil(numWarps)
   val wordSize = if (half) 2 else 4
   val wordSizeInBits = wordSize * 8/*bits*/
@@ -127,7 +135,8 @@ class TensorCoreDecoupled(
   // or [0,n/2), where 2 is the stride can be read in a single request size.
   require(tilingParams.m == tilingParams.n,
           "currently only supports square SMEM tile")
-  val numIndices = tilingParams.m / 2/*FIXME:hardcoded?*/
+  val fragmentBytes = memWidth / 8
+  val numIndices = (tilingParams.m * tilingParams.kc * wordSize) / fragmentBytes
   val indexBits = log2Ceil(numIndices)
   val lastIndex = (1 << indexBits) - 1
 
@@ -345,8 +354,10 @@ class TensorCoreDecoupled(
 
   // serialize every two B responses into one full 4x4 B tile
   // FIXME: do the same for A
+  val numBFragmentsPerTile =
+    (tilingParams.nc * tilingParams.kc * wordSize) / fragmentBytes
   val fullB = Module(new FillBuffer(
-    chiselTypeOf(respQueueB.bits.data), 2/*substeps*/
+    chiselTypeOf(respQueueB.bits.data), numBFragmentsPerTile
   ))
   fullB.io.enq.valid := respQueueB.valid
   fullB.io.enq.bits := respQueueB.bits.data
@@ -524,10 +535,13 @@ class TensorCoreDecoupled(
   // select the correct 4x4 tile from A operand buffer
   val numTilesM = tilingParams.m / tilingParams.mc
   val numTilesMBits = log2Ceil(numTilesM)
+  val numAFragmentsPerTile =
+    (tilingParams.mc * tilingParams.kc * wordSize) / fragmentBytes
   def selectOperandA(buf: Vec[UInt]): UInt = {
     require(buf.length == numIndices)
     val stepM = stepCompute & ((1 << numTilesMBits) - 1).U
-    Cat(buf((stepM << 1) + 1.U), buf(stepM << 1))
+    val base = stepM * numAFragmentsPerTile.U
+    Cat((0 until numAFragmentsPerTile).reverse.map(i => buf(base + i.U)))
   }
   val operandA = selectOperandA(fullABuf.io.deq.bits.data)
   val operandATag = fullABuf.io.deq.bits.tag

src/main/scala/radiance/subsystem/Configs.scala

@@ -51,7 +51,6 @@ class WithRadianceCores(
   tensorCoreFP16: Boolean,
   tensorCoreDecoupled: Boolean,
   tensorCoreBlackwell: Boolean,
-  numTensorWarps: Int,
   startupAddress: BigInt,
   useVxCache: Boolean
 ) extends Config((site, _, up) => {
@@ -64,7 +63,6 @@ class WithRadianceCores(
         tensorCoreFP16 = tensorCoreFP16,
         tensorCoreDecoupled = tensorCoreDecoupled,
         tensorCoreBlackwell = tensorCoreBlackwell,
-        numTensorWarps = numTensorWarps,
         startupAddress = startupAddress
       ),
       btb = None,
@@ -103,7 +101,6 @@ class WithRadianceCores(
   def this(n: Int, location: HierarchicalLocation = InSubsystem,
     tensorCoreFP16: Boolean = false, tensorCoreDecoupled: Boolean = false,
     tensorCoreBlackwell: Boolean = false,
-    numTensorWarps: Int = 4,
     startupAddress: BigInt = BigInt("10100", 16),
     useVxCache: Boolean = false)
   = this(n, location, RocketCrossingParams(
@@ -113,7 +110,7 @@ class WithRadianceCores(
       case InSubsystem => CBUS
       case InCluster(clusterId) => CCBUS(clusterId)
     }
-  ), tensorCoreFP16, tensorCoreDecoupled, tensorCoreBlackwell, numTensorWarps, startupAddress, useVxCache)
+  ), tensorCoreFP16, tensorCoreDecoupled, tensorCoreBlackwell, startupAddress, useVxCache)
 }
 
 class WithBlackwellTensorCore(location: HierarchicalLocation = InSubsystem) extends Config((site, _, up) => {

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

@@ -102,7 +102,6 @@ case class VortexCoreParams(
   tensorCoreFP16: Boolean = false, // FP16 if true, FP32 if false
   tensorCoreDecoupled: Boolean = false, // hopper-style SMEM operand decoupling
   tensorCoreBlackwell: Boolean = false, // blackwell-style TMEM + SMEM tensor core
-  numTensorWarps: Int = 4,
   startupAddress: BigInt = BigInt("10100", 16), // initial warp PC programmed through startup DCRs
   debugROB: Boolean = false, // if enabled, uses a C++ debug ROB to generate trace-with-wdata
   haveCease: Boolean = true, // non-standard CEASE instruction
@@ -211,9 +210,7 @@ class RadianceTile private (
     case Some(false) => 1
     case None => 1
   }
-  // Must match VX_gpu_pkg.sv: ICACHE_TAG_WIDTH = domain + UUID + wid.
+  val imemTagWidth = UUID_WIDTH + NW_WIDTH
-  val imemDomainWidth = 1
-  val imemTagWidth = imemDomainWidth + UUID_WIDTH + NW_WIDTH
 
   require(numWarps >= numLsuLanes,
         s"Vortex core requires numWarps (${numWarps}) >= numLsuLanes (${numLsuLanes})")
@@ -288,17 +285,9 @@ class RadianceTile private (
     )
   }
 
-  val tcSmemSize = 32
+  val tcSmemSize = numLsuLanes * 4
-  val numTensorWarps = radianceParams.core.numTensorWarps
-  val numScalarWarps = numWarps - numTensorWarps
-  require(numTensorWarps > 0 && numTensorWarps < numWarps,
-        s"Wu requires 0 < numTensorWarps (${numTensorWarps}) < numWarps (${numWarps})")
-  val numTensorCores = if (radianceParams.core.tensorCoreBlackwell) numTensorWarps else 1
-  if (radianceParams.core.tensorCoreBlackwell) {
-    require(numTensorCores == numTensorWarps, "Wu Blackwell binding requires one Tensor Core per Tensor warp")
-  }
   val tensorUsesAsyncMem = radianceParams.core.tensorCoreDecoupled || radianceParams.core.tensorCoreBlackwell
-  val tcSmemNodeCount = if (radianceParams.core.tensorCoreDecoupled) 2 else if (radianceParams.core.tensorCoreBlackwell) numTensorCores else 0
+  val tcSmemNodeCount = if (radianceParams.core.tensorCoreDecoupled) 2 else if (radianceParams.core.tensorCoreBlackwell) 1 else 0
   val tcSmemNodes = Seq.tabulate(tcSmemNodeCount) { i =>
     TLClientNode(Seq(TLMasterPortParameters.v2(
       masters = Seq(TLMasterParameters.v2(
@@ -315,21 +304,19 @@ class RadianceTile private (
   }
 
   // For Blackwell, tcSmemNodes accesses SMEM (bwgmma B operand)
-  // tcGmemNodes provide global memory access for cp (global→tmem) and cb (tmem→global)
+  // tcGmemNode provides global memory access for cp (global→tmem) and cb (tmem→global)
-  val tcGmemNodes = if (radianceParams.core.tensorCoreBlackwell) {
+  val tcGmemNode = if (radianceParams.core.tensorCoreBlackwell) Some(TLClientNode(Seq(
-    Seq.tabulate(numTensorCores) { i =>
+    TLMasterPortParameters.v2(masters = Seq(TLMasterParameters.v2(
-      TLClientNode(Seq(TLMasterPortParameters.v2(masters = Seq(TLMasterParameters.v2(
+      name = s"rad_tc_gmem_${radianceParams.coreId}",
-        name = s"rad_tc_gmem_${radianceParams.coreId}_$i",
+      sourceId = IdRange(0, 1 << dmemSourceWidth),
-        sourceId = IdRange(0, 1 << dmemSourceWidth),
+      supports = TLSlaveToMasterTransferSizes(
-        supports = TLSlaveToMasterTransferSizes(
+        probe = TransferSizes(1, tcSmemSize),
-          probe = TransferSizes(1, tcSmemSize),
+        get = TransferSizes(1, tcSmemSize),
-          get = TransferSizes(1, tcSmemSize),
+        putFull = TransferSizes(1, tcSmemSize),
-          putFull = TransferSizes(1, tcSmemSize),
+      ),
-        ),
+      requestFifo = true
-        requestFifo = true
+    )))
-      )))))
+  ))) else None
-    }
-  } else Seq.empty
 
   // combine outgoing per-lane dmemNode into 1 idenity node
   //
@@ -419,7 +406,7 @@ class RadianceTile private (
     // imemNodes.foreach { tlMasterXbar.node := TLWidthWidget(4) := _ }
     tlMasterXbar.node :=* AddressOrNode(base) :=* icacheNode
     tlMasterXbar.node :=* AddressOrNode(base) :=* dcacheNode
-    tcGmemNodes.foreach { n => tlMasterXbar.node := AddressOrNode(base) := n }
+    tcGmemNode.foreach { n => tlMasterXbar.node := AddressOrNode(base) := n }
   }
 
   /* below are copied from rocket */
@@ -777,13 +764,14 @@ class RadianceTileModuleImp(outer: RadianceTile)
     }
 
     def connectTensor = {
+      core.io.tc_tmem_C_rdata := DontCare
       if (outer.radianceParams.core.tensorCoreDecoupled) {
         val tcb0 = new {
           val addr = core.io.tc_a_bits_address(31, 0)
           val tag = core.io.tc_a_bits_tag(outer.tensorTagWidth - 1, 0)
           val write = core.io.tc_a_bits_write(0)
-          val mask = core.io.tc_a_bits_mask(31, 0)
+          val mask = core.io.tc_a_bits_mask(outer.tcSmemSize - 1, 0)
-          val data = core.io.tc_a_bits_data(255, 0)
+          val data = core.io.tc_a_bits_data(outer.tcSmemSize * 8 - 1, 0)
           val aValid = core.io.tc_a_valid(0)
           val dReady = core.io.tc_d_ready(0)
         }
@@ -791,8 +779,8 @@ class RadianceTileModuleImp(outer: RadianceTile)
           val addr = core.io.tc_a_bits_address(63, 32)
           val tag = core.io.tc_a_bits_tag(4 + outer.tensorTagWidth - 1, 4)
           val write = core.io.tc_a_bits_write(1)
-          val mask = core.io.tc_a_bits_mask(63, 32)
+          val mask = core.io.tc_a_bits_mask(2 * outer.tcSmemSize - 1, outer.tcSmemSize)
-          val data = core.io.tc_a_bits_data(511, 256)
+          val data = core.io.tc_a_bits_data(2 * outer.tcSmemSize * 8 - 1, outer.tcSmemSize * 8)
           val aValid = core.io.tc_a_valid(1)
           val dReady = core.io.tc_d_ready(1)
         }
@@ -802,8 +790,8 @@ class RadianceTileModuleImp(outer: RadianceTile)
           val adapter = Module(
             new VortexTLAdapter(
               outer.smemSourceWidth,
-              new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = 32 * 8),
+              new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
-              new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = 32 * 8),
+              new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
               client
             )
           )
@@ -813,7 +801,7 @@ class RadianceTileModuleImp(outer: RadianceTile)
           adapter.io.inReq.valid := bundle.aValid
           adapter.io.inReq.bits.address := bundle.addr
           adapter.io.inReq.bits.source := bundle.tag
-          adapter.io.inReq.bits.size := 5.U // 256 bits
+          adapter.io.inReq.bits.size := log2Ceil(outer.tcSmemSize).U
           adapter.io.inReq.bits.opcode := Mux(bundle.write.asBool, TLMessages.PutFullData, TLMessages.Get)
           adapter.io.inReq.bits.mask := bundle.mask
           adapter.io.inReq.bits.data := bundle.data
@@ -825,7 +813,7 @@ class RadianceTileModuleImp(outer: RadianceTile)
         }
         core.io.tc_a_ready := Cat(0.U(1.W), adapters.last.io.inReq.ready, adapters.head.io.inReq.ready)
         core.io.tc_d_valid := Cat(0.U(1.W), adapters.last.io.inResp.valid, adapters.head.io.inResp.valid)
-        core.io.tc_d_bits_data := Cat(0.U((32 * 8).W), adapters.last.io.inResp.bits.data, adapters.head.io.inResp.bits.data)
+        core.io.tc_d_bits_data := Cat(0.U((outer.tcSmemSize * 8).W), adapters.last.io.inResp.bits.data, adapters.head.io.inResp.bits.data)
         core.io.tc_d_bits_tag := Cat(0.U(outer.tensorTagWidth.W), adapters.last.io.inResp.bits.source, adapters.head.io.inResp.bits.source)
         require(core.io.tc_d_bits_data.widthOption.get == adapters.head.io.inResp.bits.data.widthOption.get * 3)
         require(core.io.tc_d_bits_tag.widthOption.get == adapters.head.io.inResp.bits.source.widthOption.get * 3)
@@ -835,160 +823,87 @@ class RadianceTileModuleImp(outer: RadianceTile)
         core.io.tc_d_bits_data := DontCare
         core.io.tc_d_bits_tag := DontCare
       }
-      core.io.tc_tmem_A_rready := DontCare
-      core.io.tc_tmem_A_rdata := DontCare
-      core.io.tc_tmem_C_rready := DontCare
-      core.io.tc_tmem_C_rdata := DontCare
-      core.io.tc_tmem_C_wready := DontCare
     }
 
     def connectTensorBlackwell = {
       if (outer.radianceParams.core.tensorCoreBlackwell) {
         require(outer.tcSmemNodes.nonEmpty)
-        require(outer.tcSmemNodes.length == outer.numTensorCores)
-        require(outer.tcGmemNodes.length == outer.numTensorCores)
 
-        val nTC = outer.numTensorCores
+        // TMEM matrix: direct SRAM (no TileLink), connected via VortexCore IO.
-        val tcPorts = 3
+        // Each warp owns 2KB; row count scales with the lane-dependent fragment width.
-        val tcDataBits = outer.tcSmemSize * 8
+        val tmemBytesPerWarp = 2048
-        val tmemAddrBits = 9
+        val tmemDepth = outer.numWarps * (tmemBytesPerWarp / outer.tcSmemSize)
-        val tmemDataBits = outer.numLsuLanes * 32
-        val tmemMaskBits = outer.numLsuLanes * 4
-
-        def slice(u: UInt, width: Int, idx: Int): UInt = u(width * (idx + 1) - 1, width * idx)
-        def port(tc: Int, p: Int): Int = tc * tcPorts + p
-
-        val tcAReady = Wire(Vec(nTC * tcPorts, Bool()))
-        val tcDValid = Wire(Vec(nTC * tcPorts, Bool()))
-        val tcDData = Wire(Vec(nTC * tcPorts, UInt(tcDataBits.W)))
-        val tcDTag = Wire(Vec(nTC * tcPorts, UInt(outer.tensorTagWidth.W)))
-        tcAReady.foreach(_ := false.B)
-        tcDValid.foreach(_ := false.B)
-        tcDData.foreach(_ := 0.U)
-        tcDTag.foreach(_ := 0.U)
-
-        // TMEM matrix: one shared 2R1W SRAM.  read0 is operand A, read1 is C.
-        // Each warp needs 2 tiles (A + C), each tile = 32 frags × 32B = 1KB
-        val tmemDepth = outer.numWarps * outer.tcSmemSize * 2  // numWarps × 64 rows
         val tmem = Module(new radiance.memory.TwoReadOneWriteSyncMem(
           tmemDepth, UInt((outer.tcSmemSize * 8).W)))
+        tmem.io.ren0   := core.io.tc_tmem_C_ren
+        tmem.io.raddr0 := core.io.tc_tmem_C_raddr
+        core.io.tc_tmem_C_rdata := tmem.io.rdata0
+        tmem.io.ren1   := false.B
+        tmem.io.raddr1 := 0.U
+        tmem.io.wen    := core.io.tc_tmem_C_wen
+        tmem.io.waddr  := core.io.tc_tmem_C_waddr
+        tmem.io.wdata  := core.io.tc_tmem_C_wdata
+        tmem.io.mask   := core.io.tc_tmem_C_mask
 
-        val aReadArb = Module(new RRArbiter(UInt(tmemAddrBits.W), nTC))
+        // smem_B (port 2): Global Memory via TileLink
-        val cReadArb = Module(new RRArbiter(UInt(tmemAddrBits.W), nTC))
+        val smemBBundle = new {
-
+          val addr   = core.io.tc_a_bits_address(95, 64)
-        class TmemWriteReq extends Bundle {
+          val tag    = core.io.tc_a_bits_tag(8 + outer.tensorTagWidth - 1, 8)
-          val addr = UInt(tmemAddrBits.W)
+          val write  = core.io.tc_a_bits_write(2)
-          val data = UInt(tmemDataBits.W)
+          val mask   = core.io.tc_a_bits_mask(3 * outer.tcSmemSize - 1, 2 * outer.tcSmemSize)
-          val mask = UInt(tmemMaskBits.W)
+          val data   = core.io.tc_a_bits_data(3 * outer.tcSmemSize * 8 - 1, 2 * outer.tcSmemSize * 8)
+          val aValid = core.io.tc_a_valid(2)
+          val dReady = core.io.tc_d_ready(2)
         }
-        val cWriteArb = Module(new RRArbiter(new TmemWriteReq, nTC))
+        val client = outer.tcSmemNodes.head.out.head
+        val adapter = Module(new VortexTLAdapter(
+          outer.smemSourceWidth,
+          new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
+          new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
+          client
+        ))
+        adapter.io.inReq.bits <> DontCare
+        adapter.io.inReq.valid        := smemBBundle.aValid
+        adapter.io.inReq.bits.address := smemBBundle.addr
+        adapter.io.inReq.bits.source  := smemBBundle.tag
+        adapter.io.inReq.bits.size    := log2Ceil(outer.tcSmemSize).U
+        adapter.io.inReq.bits.opcode  := Mux(smemBBundle.write.asBool, TLMessages.PutFullData, TLMessages.Get)
+        adapter.io.inReq.bits.mask    := smemBBundle.mask
+        adapter.io.inReq.bits.data    := smemBBundle.data
+        adapter.io.inResp.ready       := smemBBundle.dReady
+        client._1.a <> adapter.io.outReq
+        adapter.io.outResp <> client._1.d
 
-        (0 until nTC).foreach { tc =>
+        // port 0: global memory (cp/cb)
-          aReadArb.io.in(tc).valid := core.io.tc_tmem_A_ren(tc)
+        val gmemClient = outer.tcGmemNode.get.out.head
-          aReadArb.io.in(tc).bits := slice(core.io.tc_tmem_A_raddr, tmemAddrBits, tc)
+        val gmemAdapter = Module(new VortexTLAdapter(
-          cReadArb.io.in(tc).valid := core.io.tc_tmem_C_ren(tc)
+          outer.dmemSourceWidth,
-          cReadArb.io.in(tc).bits := slice(core.io.tc_tmem_C_raddr, tmemAddrBits, tc)
+          new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
-          cWriteArb.io.in(tc).valid := core.io.tc_tmem_C_wen(tc)
+          new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = outer.tcSmemSize * 8),
-          cWriteArb.io.in(tc).bits.addr := slice(core.io.tc_tmem_C_waddr, tmemAddrBits, tc)
+          gmemClient
-          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)
+        gmemAdapter.io.inReq.bits <> DontCare
-        }
+        gmemAdapter.io.inReq.valid        := core.io.tc_a_valid(0)
+        gmemAdapter.io.inReq.bits.address := core.io.tc_a_bits_address(31, 0)
+        gmemAdapter.io.inReq.bits.source  := core.io.tc_a_bits_tag(outer.tensorTagWidth - 1, 0)
+        gmemAdapter.io.inReq.bits.size    := log2Ceil(outer.tcSmemSize).U
+        gmemAdapter.io.inReq.bits.opcode  := Mux(core.io.tc_a_bits_write(0).asBool, TLMessages.PutFullData, TLMessages.Get)
+        gmemAdapter.io.inReq.bits.mask    := core.io.tc_a_bits_mask(outer.tcSmemSize - 1, 0)
+        gmemAdapter.io.inReq.bits.data    := core.io.tc_a_bits_data(outer.tcSmemSize * 8 - 1, 0)
+        gmemAdapter.io.inResp.ready       := core.io.tc_d_ready(0)
+        gmemClient._1.a <> gmemAdapter.io.outReq
+        gmemAdapter.io.outResp <> gmemClient._1.d
 
-        aReadArb.io.out.ready := true.B
+        core.io.tc_a_ready    := Cat(adapter.io.inReq.ready,  0.U(1.W), gmemAdapter.io.inReq.ready)
-        cReadArb.io.out.ready := true.B
+        core.io.tc_d_valid    := Cat(adapter.io.inResp.valid, 0.U(1.W), gmemAdapter.io.inResp.valid)
-        cWriteArb.io.out.ready := true.B
+        core.io.tc_d_bits_data := Cat(adapter.io.inResp.bits.data, 0.U((outer.tcSmemSize * 8).W), gmemAdapter.io.inResp.bits.data)
-
+        core.io.tc_d_bits_tag  := Cat(adapter.io.inResp.bits.source, 0.U(outer.tensorTagWidth.W), gmemAdapter.io.inResp.bits.source)
-        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 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
-        core.io.tc_tmem_A_rdata := VecInit((0 until nTC).map { tc =>
-          Mux(aReadGrant(tc), tmem.io.rdata0, 0.U(tmemDataBits.W))
-        }).asUInt
-        core.io.tc_tmem_C_rdata := VecInit((0 until nTC).map { tc =>
-          Mux(cReadGrant(tc), tmem.io.rdata1, 0.U(tmemDataBits.W))
-        }).asUInt
-
-        // port 2: SMEM B, one TL client per tensor core. RadianceSharedMem arbitrates them.
-        (0 until nTC).foreach { tc =>
-          val p2 = port(tc, 2)
-          val client = outer.tcSmemNodes(tc).out.head
-          val adapter = Module(new VortexTLAdapter(
-            outer.smemSourceWidth,
-            new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = tcDataBits),
-            new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = tcDataBits),
-            client
-          ))
-          adapter.io.inReq.bits <> DontCare
-          adapter.io.inReq.valid := core.io.tc_a_valid(p2)
-          adapter.io.inReq.bits.address := slice(core.io.tc_a_bits_address, 32, p2)
-          adapter.io.inReq.bits.source := slice(core.io.tc_a_bits_tag, outer.tensorTagWidth, p2)
-          adapter.io.inReq.bits.size := 5.U
-          adapter.io.inReq.bits.opcode := Mux(core.io.tc_a_bits_write(p2).asBool, TLMessages.PutFullData, TLMessages.Get)
-          adapter.io.inReq.bits.mask := slice(core.io.tc_a_bits_mask, 32, p2)
-          adapter.io.inReq.bits.data := slice(core.io.tc_a_bits_data, tcDataBits, p2)
-          adapter.io.inResp.ready := core.io.tc_d_ready(p2)
-          client._1.a <> adapter.io.outReq
-          adapter.io.outResp <> client._1.d
-
-          tcAReady(p2) := adapter.io.inReq.ready
-          tcDValid(p2) := adapter.io.inResp.valid
-          tcDData(p2) := adapter.io.inResp.bits.data
-          tcDTag(p2) := adapter.io.inResp.bits.source
-        }
-
-        // port 0: global memory (cp/cb), one TL client per tensor core.
-        (0 until nTC).foreach { tc =>
-          val p0 = port(tc, 0)
-          val gmemClient = outer.tcGmemNodes(tc).out.head
-          val gmemAdapter = Module(new VortexTLAdapter(
-            outer.dmemSourceWidth,
-            new VortexBundleA(tagWidth = outer.tensorTagWidth, dataWidth = tcDataBits),
-            new VortexBundleD(tagWidth = outer.tensorTagWidth, dataWidth = tcDataBits),
-            gmemClient
-          ))
-          gmemAdapter.io.inReq.bits <> DontCare
-          gmemAdapter.io.inReq.valid := core.io.tc_a_valid(p0)
-          gmemAdapter.io.inReq.bits.address := slice(core.io.tc_a_bits_address, 32, p0)
-          gmemAdapter.io.inReq.bits.source := slice(core.io.tc_a_bits_tag, outer.tensorTagWidth, p0)
-          gmemAdapter.io.inReq.bits.size := 5.U
-          gmemAdapter.io.inReq.bits.opcode := Mux(core.io.tc_a_bits_write(p0).asBool, TLMessages.PutFullData, TLMessages.Get)
-          gmemAdapter.io.inReq.bits.mask := slice(core.io.tc_a_bits_mask, 32, p0)
-          gmemAdapter.io.inReq.bits.data := slice(core.io.tc_a_bits_data, tcDataBits, p0)
-          gmemAdapter.io.inResp.ready := core.io.tc_d_ready(p0)
-          gmemClient._1.a <> gmemAdapter.io.outReq
-          gmemAdapter.io.outResp <> gmemClient._1.d
-
-          tcAReady(p0) := gmemAdapter.io.inReq.ready
-          tcDValid(p0) := gmemAdapter.io.inResp.valid
-          tcDData(p0) := gmemAdapter.io.inResp.bits.data
-          tcDTag(p0) := gmemAdapter.io.inResp.bits.source
-        }
-
-        core.io.tc_a_ready := tcAReady.asUInt
-        core.io.tc_d_valid := tcDValid.asUInt
-        core.io.tc_d_bits_data := tcDData.asUInt
-        core.io.tc_d_bits_tag := tcDTag.asUInt
       } else {
         core.io.tc_a_ready    := false.B
         core.io.tc_d_valid    := false.B
         core.io.tc_d_bits_data := DontCare
         core.io.tc_d_bits_tag  := DontCare
-        core.io.tc_tmem_A_rready := DontCare
-        core.io.tc_tmem_A_rdata := DontCare
-        core.io.tc_tmem_C_rready := DontCare
         core.io.tc_tmem_C_rdata := DontCare
-        core.io.tc_tmem_C_wready := DontCare
       }
     }
 
@@ -1093,11 +1008,7 @@ class RadianceTileModuleImp(outer: RadianceTile)
     tensor.io.reqA.ready := false.B
     tensor.io.reqB.ready := false.B
     tensor.io.writeback.ready := false.B
-    tensor.io.tmemC.aRready := false.B
+    tensor.io.tmemC.rdata := DontCare
-    tensor.io.tmemC.aRdata := DontCare
-    tensor.io.tmemC.cRready := false.B
-    tensor.io.tmemC.cRdata := DontCare
-    tensor.io.tmemC.cWready := false.B
     dontTouch(tensor.io)
   } else {
     if (outer.radianceParams.core.tensorCoreFP16) {

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

@@ -90,36 +90,28 @@ class VortexBundle(tile: RadianceTile)(implicit p: Parameters) extends CoreBundl
   val smem_d_bits_data = Input(UInt((tile.numLsuLanes * 32).W))
   val smem_d_ready = Output(UInt((tile.numLsuLanes * 1).W))
 
-  val numTensorCores = if (tile.radianceParams.core.tensorCoreBlackwell) tile.numTensorCores else 1
   val tcPortCount = 3
-  val tcFlatPortCount = tcPortCount * numTensorCores
+  val tc_a_valid = Output(UInt(tcPortCount.W))
-  val tc_a_valid = Output(UInt(tcFlatPortCount.W))
+  val tc_a_bits_write = Output(UInt(tcPortCount.W))
-  val tc_a_bits_write = Output(UInt(tcFlatPortCount.W))
+  val tc_a_bits_address = Output(UInt((tcPortCount * 32).W))
-  val tc_a_bits_address = Output(UInt((tcFlatPortCount * 32).W))
+  val tc_a_bits_tag = Output(UInt((tcPortCount * 4).W))
-  val tc_a_bits_tag = Output(UInt((tcFlatPortCount * 4).W))
+  val tc_a_bits_mask = Output(UInt((tcPortCount * tile.numLsuLanes * 4).W))
-  val tc_a_bits_mask = Output(UInt((tcFlatPortCount * 32).W))
+  val tc_a_bits_data = Output(UInt((tcPortCount * tile.numLsuLanes * 32).W))
-  val tc_a_bits_data = Output(UInt((tcFlatPortCount * 32 * 8).W))
+  val tc_a_ready = Input(UInt(tcPortCount.W))
-  val tc_a_ready = Input(UInt(tcFlatPortCount.W))
+  val tc_d_valid = Input(UInt(tcPortCount.W))
-  val tc_d_valid = Input(UInt(tcFlatPortCount.W))
+  val tc_d_bits_data = Input(UInt((tcPortCount * tile.numLsuLanes * 32).W))
-  val tc_d_bits_data = Input(UInt((tcFlatPortCount * 32 * 8).W))
+  val tc_d_bits_tag = Input(UInt((tcPortCount * 4).W))
-  val tc_d_bits_tag = Input(UInt((tcFlatPortCount * 4).W))
+  val tc_d_ready = Output(UInt(tcPortCount.W))
-  val tc_d_ready = Output(UInt(tcFlatPortCount.W))
 
-  // Direct SRAM ports for shared TMEM (bypasses TileLink)
+  // Direct SRAM port for TMEM C (bypasses TileLink)
   val numLanes = tile.numLsuLanes
-  val tc_tmem_A_ren    = Output(UInt(numTensorCores.W))
+  val tc_tmem_C_wen   = Output(Bool())
-  val tc_tmem_A_rready = Input(UInt(numTensorCores.W))
+  val tc_tmem_C_ren   = Output(Bool())
-  val tc_tmem_A_raddr  = Output(UInt((numTensorCores * 9).W))
+  val tc_tmem_C_waddr = Output(UInt(9.W))
-  val tc_tmem_A_rdata  = Input(UInt((numTensorCores * numLanes * 32).W))
+  val tc_tmem_C_raddr = Output(UInt(9.W))
-  val tc_tmem_C_ren    = Output(UInt(numTensorCores.W))
+  val tc_tmem_C_wdata = Output(UInt((numLanes * 32).W))
-  val tc_tmem_C_rready = Input(UInt(numTensorCores.W))
+  val tc_tmem_C_mask  = Output(UInt((numLanes * 4).W))
-  val tc_tmem_C_raddr  = Output(UInt((numTensorCores * 9).W))
+  val tc_tmem_C_rdata = Input(UInt((numLanes * 32).W))
-  val tc_tmem_C_rdata  = Input(UInt((numTensorCores * numLanes * 32).W))
-  val tc_tmem_C_wen    = Output(UInt(numTensorCores.W))
-  val tc_tmem_C_wready = Input(UInt(numTensorCores.W))
-  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))
 
   // FIXME: hardcoded
   val barrierIdBits = tile.barrierMasterNode.out(0)._2.barrierIdBits
@@ -156,7 +148,7 @@ class Vortex(tile: RadianceTile)(implicit p: Parameters)
         "TENSOR_FP16" -> (if (tile.radianceParams.core.tensorCoreFP16) 1 else 0),
         "STARTUP_ADDR" -> tile.radianceParams.core.startupAddress,
         "NUM_THREADS" -> tile.numLsuLanes,
-        "NUM_TENSOR_CORES" -> (if (tile.radianceParams.core.tensorCoreBlackwell) tile.numTensorCores else 1)
+        "TC_DATA_WIDTH" -> (tile.numLsuLanes * 32)
       )
     )
     with HasBlackBoxResource with HasBlackBoxPath {
@@ -220,7 +212,6 @@ class Vortex(tile: RadianceTile)(implicit p: Parameters)
   addResource("/vsrc/vortex/hw/rtl/core/VX_scoreboard.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_sfu_unit.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_smem_unit.sv")
-  addResource("/vsrc/vortex/hw/rtl/core/VX_tensor_ctrl_unit.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_split_join.sv")
   addResource("/vsrc/vortex/hw/rtl/core/VX_trace.vh")
   addResource("/vsrc/vortex/hw/rtl/core/VX_wctl_unit.sv")