Richard Yan
c114a7a4ab
new gemm kernel
2024-11-08 20:55:27 -08:00
Hansung Kim
367fa927f8
sgemm_impl: Fix default FP_SIZE to 16
2024-10-29 14:56:18 -07:00
Hansung Kim
e1b0fc3944
generate_matrix.py: Rand [0,1); also save non-swizzled row-major B
2024-10-29 14:55:32 -07:00
Hansung Kim
24064dc7a2
sgemm_impl: Do proper addr gen and store for wgmma
2024-10-29 01:31:55 -07:00
Hansung Kim
bd7a8e39b9
sgemm_impl: Split out smem addr gen to functions
...
so that the addr gen code can also be used for wgmma.
2024-10-29 01:30:48 -07:00
Hansung Kim
ae98ae6e93
sgemm_tcore: Fix DMA smem addresses, add markers
...
Take into account that DMA writes B tiles starting from the end of the
quartile.
2024-10-28 17:26:07 -07:00
Hansung Kim
b4dadfaf61
Merge remote-tracking branch 'origin/kernels' into kernels-hopper
2024-10-28 14:25:18 -07:00
Richard Yan
fd1c9f4729
update gemmini dma kernel
2024-10-28 13:47:13 -07:00
Hansung Kim
d0421426be
sgemm_tcore: Hardcode CISC spadQuartile addresses
2024-10-28 12:49:22 -07:00
Hansung Kim
e55c8b480e
sgemm_impl: Comment out GEMMINI_DMA code in single_tile
...
This is already done in the higher-level thread_block_gemm function, and
flash also has explicit DMA sync code. Also having this executed twice
sometimes triggers vx_bar movement into a branch which we really want to
avoid.
2024-10-28 12:47:49 -07:00
Hansung Kim
36eb50060f
sgemm_impl: Add skeleton wgmma routine for single_tile
2024-10-28 12:47:20 -07:00
Hansung Kim
e8a943e893
Push golden sgemm_tcore kernels
2024-10-27 19:40:08 -07:00
Richard Yan
379d863456
Merge branch 'kernels' of https://github.com/hansungk/vortex-private into kernels
2024-10-24 17:31:01 -07:00
Richard Yan
fb928b5cda
new unaligned access kernel, update idle kernel
2024-10-24 17:28:59 -07:00
Richard Yan
8cc0c3bae4
fp16 no dma kernel
2024-10-24 17:12:34 -07:00
Richard Yan
04a64dee7a
fp16 dma kernel
2024-10-24 17:12:02 -07:00
Hansung Kim
6417a625b1
sgemm_impl: Add tiling params for hopper tensor core
2024-10-23 19:50:18 -07:00
Hansung Kim
68cd6455fe
sgemm_impl: Add mmio reconverge barrier to avoid slip-off; switch to FP32
2024-10-02 15:17:44 -07:00
Hansung Kim
34d0956cd5
tensor: Attempt row-major mapping for C store (WIP)
...
Doesn't work because 1x2 jagged mapping is required to achieve
throughput for storing the bigger C matrix (2x4, vs. 2x2 in A).
2024-10-02 15:14:55 -07:00
Hansung Kim
3490294626
generate_matrix.py: switch to fp16 rand, generate row-major A
2024-10-02 11:01:23 -07:00
Hansung Kim
db2789bf23
Add asm label for cisc compute
2024-10-02 10:59:14 -07:00
Hansung Kim
6f6ee5616f
Add convergent attribute to vx_barrier
...
Note this attribute is only supported by Clang, so this will only be
applied to the kernel binary but not runtime.
2024-10-02 10:57:45 -07:00
Hansung Kim
221d5f75c2
flash: Optimize smem alloc for tcore for 8banks
...
Divide into first half & last half for warpgroup 0 & 1, and
allocate Q/K and P/V in different banks for parallel acccess.
2024-09-19 21:31:39 -07:00
Hansung Kim
d0ef06cec1
flash: Complete Q_IS_K_MAJOR code for GEMM II
2024-09-19 20:36:03 -07:00
Hansung Kim
b9cafd6372
idle: unused const
2024-09-18 18:10:29 -07:00
Hansung Kim
be15cffbf3
flash: Revert to gemmini config, remove DEBUG and unnecessary checks
2024-09-12 14:25:33 -07:00
Hansung Kim
b5916f3f07
flash: Fix hardcoded barrier for tcore; move tcore-specific flags
2024-09-11 22:08:44 -07:00
Hansung Kim
d69707f686
flash: Enable GEMM II fence; Pull 1st KV move out of the loop
2024-09-11 19:24:06 -07:00
Hansung Kim
18cf0e73cd
flash: Add early return for warp-indivisible row iter
2024-09-11 00:56:09 -07:00
Hansung Kim
068d48534e
flash: Swap S1/S0 to avoid GEMM II - softmax bank conflict
...
+ remove spurrious fences to better overlap GEMM I and DMA
2024-09-11 00:55:36 -07:00
Hansung Kim
ba66d2c2bd
sgemm_impl: barrier dumb dumb
2024-09-11 00:01:56 -07:00
Hansung Kim
dc746272fb
flash: Conditionally enable GEMM II fence code, fix tile_k for DEBUG
2024-09-10 22:53:35 -07:00
Hansung Kim
28b2eaec8f
sgemm_gemmini_dma: Fix tile size to (128,64,128)
2024-09-10 18:29:40 -07:00
Hansung Kim
2152c80ffd
sgemm_impl: Add missing reconvergence barrier after mmio
2024-09-10 18:05:01 -07:00
Hansung Kim
ccddd0bcc9
sgemm_impl: Remove unused FLEXIBLE_LAYOUT
2024-09-10 15:54:17 -07:00
Hansung Kim
90e03894fc
flash: Add flag in SMEM for dependency check on O
...
TODO: results unverified.
Stalls O rescale until GEMM II finishes.
2024-09-10 13:42:47 -07:00
Hansung Kim
88760596cb
flash: Remove bogus mvout to SMEM code
2024-09-09 17:18:59 -07:00
Hansung Kim
a17edac875
flash: Fix barrier stall with DEBUG
...
Verified for up to P_expected on 2nd iter; O_before_PV is partially
correct
2024-09-09 17:02:05 -07:00
Hansung Kim
b652e25945
flash: Warp-specialize between warp 0 and 1-7
...
Finishes without stalls; No dependency check between O rescale and
GEMM-II.
2024-09-09 16:42:30 -07:00
Hansung Kim
d31c8ffd7d
flash: Fix grid size to hw cluster size
...
Verified fast config, minus the barrier stall at the end.
2024-09-09 15:44:03 -07:00
Hansung Kim
829af5d429
flash: Comment out mvout to smem
...
Verified up to O_before_PV; still stalls without DEBUG
2024-09-09 15:21:49 -07:00
Hansung Kim
ecc800964a
flash: Change smem alloc for less bank conflicts; noskip stc
2024-09-09 13:47:18 -07:00
Hansung Kim
1f51f7f9d4
sgemm_impl: Mark threadblock_barrier convergent
...
Thank you Chris Lattner
2024-09-08 22:49:38 -07:00
Hansung Kim
714b9f501e
flash: Restructure to do delayed fence for better concurrency
...
Verified up to O_before_PV of 2nd iteration; O_after_PV needs preload
fix.
FIXME: Stalls at barrier without DEBUG set.
2024-09-08 22:16:17 -07:00
Hansung Kim
6911843a82
flash: Remove unnecessary dmem preload, fix rowmax/rowsum dependency
2024-09-08 21:11:59 -07:00
Hansung Kim
a4dd45bc1b
flash: Replace CISC with RISC
...
spadQuartile in hw does not match spad addresses in kernel; match them
later for optimization.
2024-09-08 20:52:28 -07:00
Hansung Kim
6547e92757
flash: Load Q to both quartiles; preload O for acc
2024-09-08 19:47:55 -07:00
Hansung Kim
8efa6868ea
flash: Restructure for full software pipelining
...
Verified up to P and O before PV; need to fix iteration for V load.
2024-09-08 18:45:32 -07:00
Hansung Kim
cdb8377b62
flash: Do GEMM II in Gemmini; verify 1st iteration
2024-09-08 16:09:06 -07:00
Hansung Kim
3f50ac57ee
flash: use 12bit dma interface
2024-09-08 15:29:56 -07:00
