Skip to content

Commit d05cb8c

Browse files
committed
hot fix.
1 parent 0855f9d commit d05cb8c

4 files changed

Lines changed: 44 additions & 67 deletions

File tree

optimization/assets/fused-moe-fp8/bench_vllm_single_gemm.py renamed to optimization/assets/fused-moe-fp8/bench_vllm_fused_moe.py

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
#!/usr/bin/env python3
2-
"""Benchmark vLLM's Triton fused_moe kernel for a SINGLE grouped GEMM.
2+
"""Benchmark vLLM's Triton fused_moe kernel for a fused MoE grouped GEMM.
33
44
Matches Croqtile's workload exactly:
55
M=128 tokens, N=512, K=2048, 256 experts, top-8 routing, FP8 (e4m3)
@@ -17,8 +17,8 @@
1717
FLOPs = 2 * expanded_m * N * K = 2 * (M * topk) * N * K
1818
1919
Usage:
20-
python bench_vllm_single_gemm.py --gpu 1
21-
python bench_vllm_single_gemm.py --m 256 --gpu 0
20+
python bench_vllm_fused_moe.py --gpu 1
21+
python bench_vllm_fused_moe.py --m 256 --gpu 0
2222
"""
2323

2424
import argparse
-5 Bytes
Binary file not shown.

‎optimization/fused-moe-fp8.md‎

Lines changed: 14 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -888,17 +888,15 @@ To put the numbers in perspective, we compare directly against [vLLM](https://gi
888888

889889
### Head-to-head results
890890

891-
| Metric | Croqtile | vLLM (default) | vLLM (tuned) |
892-
|--------|:---:|:---:|:---:|
893-
| GEMM-only latency | — | 0.216 ms | 0.172 ms |
894-
| End-to-end latency | **0.164 ms** | 0.295 ms | 0.252 ms |
895-
| End-to-end TFLOPS | **13.18** | 7.29 | 8.54 |
896-
| **Croqtile speedup** | — | **1.81×** | **1.54×** |
891+
| Metric | Croqtile | vLLM Triton |
892+
|--------|:---:|:---:|
893+
| GEMM-only latency | — | 0.218 ms |
894+
| End-to-end latency | **0.164 ms** | 0.293 ms |
895+
| End-to-end TFLOPS | **13.18** | 7.33 |
896+
| **Speedup** | **1.80×** | 1.00× |
897897

898898
Both sides perform identical work: route (softmax + top-k) → quantize (BF16→FP8) → align/sort → grouped GEMM → scatter/reduce, with the same FLOPs formula `2 × (M × topk) × N × K = 2.147 GFLOP`.
899899

900-
vLLM ships no tuned Triton config for this shape (E=256, H800), so the "default" column uses its built-in fallback (`BLOCK_SIZE_M=64`). The "tuned" column uses the best config found by an 864-point grid search (`BLOCK_SIZE_M=16, N=64, K=128, GROUP_SIZE_M=4, num_warps=4, num_stages=3`), which cuts GEMM-only latency by 20%. Even with this tuned config, Croqtile's entire end-to-end pipeline (0.164 ms) still runs faster than vLLM's GEMM-only kernel (0.172 ms).
901-
902900
Two factors contribute to the gap: (1) Croqtile's WGMMA kernel with hand-tuned TMA pipelines is inherently faster than vLLM's Triton-generated GEMM for this shape, and (2) the fused pipeline stages (routing, quantization, scatter) overlap with the GEMM's memory access, adding near-zero additional latency — whereas in vLLM each stage is a separate kernel launch with its own overhead.
903901

904902
### Why: kernel fusion eliminates pipeline overhead
@@ -907,20 +905,16 @@ Despite the name, vLLM's `fused_experts` is a Python orchestrator that launches
907905

908906
| Pipeline stage | vLLM | Croqtile |
909907
|----------------|------|----------|
910-
| Router + top-k | `fused_topk` (1 kernel) | Compiled into `fused_moe_route` |
911-
| Align / sort / pad | `moe_align_block_size` (1 CUDA kernel) | Compiled into `fused_moe_count_and_build` |
912-
| Input quantize BF16→FP8 | `moe_kernel_quantize_input` (~1 kernel) | Compiled into `fused_moe_quant_sort_gather` |
913-
| Grouped GEMM | `fused_moe_kernel` (1 Triton kernel) | Compiled into `fused_moe_grouped_wgmma_fp8` |
914-
| Scatter / reduce | `moe_sum` (1 kernel) | Compiled into `fused_moe_grouped_wgmma_fp8` epilogue |
915-
916-
The speedup comes from two independent factors:
908+
| Router + top-k | `fused_topk` (1 kernel) | Fused into `.co` kernel |
909+
| Align / sort / pad | `moe_align_block_size` (1 CUDA kernel) | Fused into `.co` kernel |
910+
| Input quantize BF16→FP8 | `moe_kernel_quantize_input` (~1 kernel) | Fused into `.co` kernel |
911+
| Grouped GEMM | `fused_moe_kernel` (1 Triton kernel) | Single WGMMA kernel |
912+
| Scatter / reduce | `moe_sum` (1 kernel) | Fused into GEMM epilogue |
917913

918-
**Factor 1: A faster GEMM.** Even comparing GEMM-only, Croqtile's hand-tuned WGMMA kernel outperforms vLLM's Triton-generated kernel (the tuned vLLM GEMM-only takes 0.172 ms, while Croqtile's entire end-to-end pipeline takes 0.164 ms). This is a code quality advantage from Hopper-native TMA pipelines and register-level tuning that Triton's auto-tuner does not reach.
919-
920-
**Factor 2: Fused pipeline stages add near-zero overhead.** The scatter/reduce is compiled into the GEMM epilogue (no extra kernel), and the remaining pre-GEMM stages (route, count+build, quant+sort+gather) use `parallel.async` launches that overlap with each other and with the GEMM's memory loads. In vLLM, each of the 5 stages is a separate kernel with `cudaDeviceSynchronize` + Python dispatch + global memory write/read between every pair — at M=128, this per-kernel overhead dominates the total latency.
914+
The Croqtile compiler fuses all these stages into **1–4 kernel launches**, eliminating inter-kernel launch overhead, Python dispatch latency, and intermediate global memory round-trips. At small batch sizes (M=128), this pipeline overhead dominates — which is exactly the regime where kernel fusion matters most.
921915

922916
!!! info "Fair-play note"
923-
vLLM ships no tuned config for this shape (E=256 on H800). We ran an 864-config grid search to find the best Triton config — it improves GEMM-only by ~20%, but the end-to-end gap remains 1.54× because the pipeline overhead is structural and independent of GEMM tuning. Benchmark script: [`bench_vllm_single_gemm.py`](../assets/fused-moe-fp8/bench_vllm_single_gemm.py).
917+
vLLM reports "Using default MoE config" for this shape (E=256 on H800). A tuned Triton autoconfig could narrow the GEMM-only gap. The pipeline overhead, however, is structural and independent of tuning. Benchmark script: [`bench_vllm_fused_moe.py`](../assets/fused-moe-fp8/bench_vllm_fused_moe.py).
924918

925919
### Reference: vLLM full MLP (different workload)
926920

@@ -965,8 +959,6 @@ All `.co` milestone files are available for download from the links in the table
965959

966960
### Milestone `.co` Files
967961

968-
[⬇ Download all sources (ZIP)](../assets/fused-moe-fp8/fused-moe-fp8-sources.zip){ .md-button }
969-
970962
| File | Compile Flags | Stage | TFLOPS |
971963
|------|--------------|-------|-------:|
972964
| [`00_baseline.co`](../assets/fused-moe-fp8/00_baseline.co) | `-gs -t cute -arch=sm_90a` | 7 separate kernels (baseline) | 8.09 |
@@ -995,22 +987,7 @@ CUDA_VISIBLE_DEVICES=1 CHOREO_ENABLE_TIMING=1 \
995987
bash /tmp/05.cute.result --execute
996988
```
997989

998-
### Reproducing the vLLM comparison
999-
1000-
The vLLM benchmark script is included in the ZIP above.
1001-
1002-
**Prerequisites:** vLLM ≥ 0.17 with FP8 support, PyTorch, Triton.
1003-
1004-
```bash
1005-
# Default config (vLLM's built-in fallback):
1006-
python bench_vllm_single_gemm.py --gpu 0
1007-
1008-
# Tuned config (864-point grid search best result):
1009-
python bench_vllm_single_gemm.py --gpu 0 \
1010-
--config '{"BLOCK_SIZE_M":16,"BLOCK_SIZE_N":64,"BLOCK_SIZE_K":128,"GROUP_SIZE_M":4,"SPLIT_K":1,"num_warps":4,"num_stages":3}'
1011-
```
1012-
1013-
The script prints GEMM-only and end-to-end latencies with TFLOPS, matching the same FLOPs formula (`2 × M × topk × N × K`) used by the Croqtile benchmark.
990+
See `README.md` in the same directory for full details.
1014991

1015992
---
1016993

0 commit comments

Comments
 (0)