From a8ca42b59cc0e4c4a4d2d78d1347b72207fc26e9 Mon Sep 17 00:00:00 2001 From: Claude Date: Sat, 15 Aug 2026 05:11:32 +0000 Subject: [PATCH] =?UTF-8?q?perf:=20hgemm=5Fwarptiling=20Config=20B=20?= =?UTF-8?q?=E2=80=94=20beats=20cublas=20on=20MoE-sized=20GEMM=20(0.7x)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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. --- ex_engine/xllm_kernels/cuda/hgemm_warptiling.cu | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/ex_engine/xllm_kernels/cuda/hgemm_warptiling.cu b/ex_engine/xllm_kernels/cuda/hgemm_warptiling.cu index d38f2abe..6272af97 100644 --- a/ex_engine/xllm_kernels/cuda/hgemm_warptiling.cu +++ b/ex_engine/xllm_kernels/cuda/hgemm_warptiling.cu @@ -179,14 +179,17 @@ void launch_hgemm_warptiling( __half* C, cudaStream_t stream) { - // Config for BI-V100 (warp_size=64, 128KB smem, 16 SMs): + // Config B — best on BI-V100 (beats cublas 0.7x on 256x4096@4096x11008): + // probe_k10_configs.sh confirmed: 7.6ms vs cublas 10.5ms // 128 threads = 2 warps of 64 - // probe confirmed: NUM_WARPS=2, WMITER=2, threads_per_warp=64 ✓ + // WMITER = (64*64)/(64*8*4*2) = 4096/4096 = 1 + // WSUBM = 64/1 = 64, WSUBN = 64/2 = 32 + // threads_per_warp = (64/8)*(32/4) = 8*8 = 64 ✓ constexpr int NUM_THREADS = 128; constexpr int BM = 128, BN = 128, BK = 16; - constexpr int WM = 64, WN = 128; - constexpr int WNITER = 4; - constexpr int TM = 4, TN = 4; + constexpr int WM = 64, WN = 64; + constexpr int WNITER = 2; + constexpr int TM = 8, TN = 4; dim3 grid(CEIL_DIV(N, BN), CEIL_DIV(M, BM)); dim3 block(NUM_THREADS);