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

Compare commits

base: wu-arch:wu-architecture
Branches Tags
wu-arch:wu-architecture wu-arch:pre-wu-blackwell wu-arch:wu-blackwell-fp8-e4m3 wu-arch:wu-tmem-bank-scalar-softmax wu-arch:nvidia-style-arch wu-arch:feature/scalar-tmem-softmax-vector wu-arch:wu-blackwell-4lane-tc wu-arch:asplos-ae
..
compare: wu-arch:wu-blackwell-fp8-e4m3
Branches Tags
wu-arch:pre-wu-blackwell wu-arch:wu-blackwell-fp8-e4m3 wu-arch:wu-tmem-bank-scalar-softmax wu-arch:nvidia-style-arch wu-arch:feature/scalar-tmem-softmax-vector wu-arch:wu-blackwell-4lane-tc wu-arch:wu-architecture wu-arch:asplos-ae

11 Commits
wu-archite ... wu-blackwe

Author SHA1 Message Date
Zhongdi LUO
80bc54de0c feat: use systolic array for blackwell bwgmma 2026-07-15 07:26:30 +00:00
Zhongdi LUO
b5f716fa3c chore: update radiance bwgmma fix 2026-07-15 07:10:51 +00:00
Zhongdi LUO
5d9f3633bf config: reduce Radiance shared memory to 32 KiB 2026-07-13 08:58:31 +00:00
Zhongdi LUO
a7b886d1f3 feat: add blackwell synthesis configuration 2026-07-12 23:59:39 +00:00
Zhongdi LUO
8fe292015c chore: update radiance pipelined softmax 2026-07-12 02:11:22 +00:00
Zhongdi LUO
5c88e9b67f chore: update radiance fp8 blackwell pointer 2026-07-03 08:41:05 +00:00
Zhongdi LUO
55b7fa2007 chore: update radiance fexp support 2026-07-02 07:26:14 +00:00
Zhongdi LUO
bdca34b2cc Update Radiance for scalar TMEM support 2026-06-24 06:25:39 +00:00
Zhongdi LUO
b0a79c26ee Update radiance pre-TMEM-bank baseline 2026-06-21 08:20:52 +00:00
Zhongdi LUO
d764de5115 Update radiance Blackwell SMEM alignment fix 2026-05-27 08:44:49 +00:00
Zhongdi LUO
3cb58af336 Update radiance for 4-lane Blackwell tensor core 2026-05-27 05:55:17 +00:00
2 changed files with 13 additions and 8 deletions
Expand all files Collapse all files

19
generators/chipyard/src/main/scala/config/RadianceConfigs.scala
View File

@@ -27,6 +27,7 @@ class VirgoConfig extends RadianceClusterConfig
class VirgoFP16Config extends RadianceFP16ClusterConfig
class VirgoHopperConfig extends Radiance4CFP16ClusterConfig
class VirgoBlackwellConfig extends RadianceBlackwellClusterConfig
class VirgoBlackwellSynConfig extends RadianceBlackwellClusterSynConfig
class VirgoFlashConfig extends RadianceClusterConfig
class VirgoSynConfig extends RadianceClusterSynConfig
class VirgoFP16SynConfig extends RadianceFP16ClusterSynConfig
@@ -57,7 +58,7 @@ class RadianceBaseConfig extends Config(
class RadianceFP16ClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 16, accSizeInKB = 32, tileSize = (8, 4, 8), dataType = RadianceGemminiDataType.FP16) ++
new radiance.subsystem.WithRadianceCores(8, location = InCluster(0), tensorCoreFP16 = true, tensorCoreDecoupled = false, useVxCache = false) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 128 << 10, numBanks = 4, numWords = 16) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 4, numWords = 16) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
new radiance.subsystem.WithVortexL1Banks(nBanks = 8) ++
new radiance.subsystem.WithRadianceCluster(0) ++
@@ -66,7 +67,7 @@ class RadianceFP16ClusterConfig extends Config(
class Radiance8B8WFP16ClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 16, accSizeInKB = 32, tileSize = (8, 4, 8), dataType = RadianceGemminiDataType.FP16) ++
new radiance.subsystem.WithRadianceCores(8, location = InCluster(0), tensorCoreFP16 = true, tensorCoreDecoupled = false, useVxCache = false) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 128 << 10, numBanks = 8, numWords = 8) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 8, numWords = 8) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
new radiance.subsystem.WithVortexL1Banks(nBanks = 8) ++
new radiance.subsystem.WithRadianceCluster(0) ++
@@ -75,12 +76,12 @@ class Radiance8B8WFP16ClusterConfig extends Config(
class Radiance4CFP16ClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 16, accSizeInKB = 32, tileSize = (8, 4, 8), dataType = RadianceGemminiDataType.FP16) ++
new radiance.subsystem.WithRadianceCores(4, location = InCluster(0), tensorCoreFP16 = true, tensorCoreDecoupled = true, useVxCache = false) ++
// new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 128 << 10, numBanks = 4, numWords = 16,
// new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 4, numWords = 16,
// memType = radiance.subsystem.TwoReadOneWrite,
// serializeUnaligned = radiance.subsystem.CoreSerialized) ++
// NOTE: Hopper Tensor Core does not work with 16-word config due to the
// address alignment requirement
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 128 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
new radiance.subsystem.WithVortexL1Banks(nBanks = 8) ++
new radiance.subsystem.WithRadianceCluster(0) ++
@@ -89,7 +90,7 @@ class Radiance4CFP16ClusterConfig extends Config(
class RadianceBlackwellClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 16, accSizeInKB = 32, tileSize = (8, 4, 8), dataType = RadianceGemminiDataType.FP16) ++
new radiance.subsystem.WithRadianceCores(1, location = InCluster(0), tensorCoreFP16 = true, tensorCoreDecoupled = false, tensorCoreBlackwell = true, numTensorWarps = 2, startupAddress = BigInt("80000000", 16), useVxCache = false) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 128 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
new radiance.subsystem.WithVortexL1Banks(nBanks = 8) ++
new radiance.subsystem.WithRadianceCluster(0) ++
@@ -101,7 +102,7 @@ class RadianceClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 8, accSizeInKB = 16, tileSize = 8) ++
new radiance.subsystem.WithRadianceCores(4, location = InCluster(0), tensorCoreFP16 = false, tensorCoreDecoupled = false, useVxCache = false) ++
// new radiance.subsystem.WithRadianceFrameBuffer(x"ff018000", 16, 0x8000, x"ff011000", "fb0") ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 256 << 10/*KBytes*/, numBanks = 8, numWords = 8,
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10/*KBytes*/, numBanks = 8, numWords = 8,
// memType = radiance.subsystem.TwoReadOneWrite,
serializeUnaligned = radiance.subsystem.CoreSerialized) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
@@ -134,7 +135,7 @@ class RadianceBigLittleClusterConfig extends Config(
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 4, accSizeInKB = 16, tileSize = 16) ++
new radiance.subsystem.WithRadianceGemmini(location = InCluster(0), dim = 8, accSizeInKB = 16, tileSize = 8) ++
new radiance.subsystem.WithRadianceCores(2, location = InCluster(0), useVxCache = false) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 64 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithRadianceSharedMem(address = x"ff000000", size = 32 << 10, numBanks = 4, numWords = 8) ++
new radiance.subsystem.WithCoalescer(nNewSrcIds = 16) ++
new radiance.subsystem.WithVortexL1Banks(nBanks = 8)++
new radiance.subsystem.WithRadianceCluster(0) ++
@@ -152,6 +153,10 @@ class Radiance4CFP16ClusterSynConfig extends Config(
new radiance.subsystem.WithRadianceSimParams(false) ++
new Radiance4CFP16ClusterConfig)
class RadianceBlackwellClusterSynConfig extends Config(
new radiance.subsystem.WithRadianceSimParams(false) ++
new RadianceBlackwellClusterConfig)
class RadianceBigLittleClusterSynConfig extends Config(
new radiance.subsystem.WithRadianceSimParams(false) ++
new RadianceBigLittleClusterConfig)

2
generators/radiance

Submodule generators/radiance updated: 126523c5d2...7b769e9e84