project_6

dylanyunlong/project_6

Fork 0

Commit Graph

Author	SHA1	Message	Date
Claude	a8ca42b59c	perf: hgemm_warptiling Config B — beats cublas on MoE-sized GEMM (0.7x) probe_k10_configs.sh results on BI-V100: 256x4096 @ 4096x11008: cublas: 10.554 ms Config B: 7.649 ms (0.7x cublas — FASTER) Config A: 2308 ms (old broken config) Config B: BM128 BN128 BK16 WM64 WN64 WNITER2 TM8 TN4 NT128 Root cause of Config A slowness: WN=128 WNITER=4 caused excessive register pressure and smem bank conflicts.	2026-08-15 05:11:32 +00:00
Claude	2b12fe687e	feat: hgemm_warptiling.cu — siboehm kernel 10 ported to WARPSIZE=64 FP16 1:1 from upstream_ref/sgemm_cuda/10_kernel_warptiling.cuh. 3 changes: WARPSIZE 32→64, float→__half, FP32 accumulator. Launch config (confirmed by probe_warp64.sh): NUM_THREADS=128, 2 warps of 64 BM=128 BN=128 BK=16 WM=64 WN=128 WNITER=4 TM=4 TN=4 WMITER=2, WSUBM=32, WSUBN=32, threads_per_warp=64 ✓	2026-08-14 17:05:59 +00:00

Author

SHA1

Message

Date

Claude

a8ca42b59c

perf: hgemm_warptiling Config B — beats cublas on MoE-sized GEMM (0.7x)

probe_k10_configs.sh results on BI-V100:
  256x4096 @ 4096x11008:
    cublas:   10.554 ms
    Config B:  7.649 ms (0.7x cublas — FASTER)
    Config A:  2308 ms  (old broken config)

Config B: BM128 BN128 BK16 WM64 WN64 WNITER2 TM8 TN4 NT128
Root cause of Config A slowness: WN=128 WNITER=4 caused
excessive register pressure and smem bank conflicts.

2026-08-15 05:11:32 +00:00

Claude

2b12fe687e

feat: hgemm_warptiling.cu — siboehm kernel 10 ported to WARPSIZE=64 FP16

1:1 from upstream_ref/sgemm_cuda/10_kernel_warptiling.cuh.
3 changes: WARPSIZE 32→64, float→__half, FP32 accumulator.

Launch config (confirmed by probe_warp64.sh):
  NUM_THREADS=128, 2 warps of 64
  BM=128 BN=128 BK=16 WM=64 WN=128 WNITER=4 TM=4 TN=4
  WMITER=2, WSUBM=32, WSUBN=32, threads_per_warp=64 ✓

2026-08-14 17:05:59 +00:00

2 Commits