feat(moe): support BF16 activations in CuTe DSL NVFP4 MoE - #4048
Merged
aleozlx merged 118 commits intoAug 7, 2026
Merged
Conversation
Contributor
|
Important Review skippedNo new commits to review since the last review. ⚙️ Run configurationConfiguration used: defaults Review profile: CHILL Plan: Pro Plus Run ID: You can disable this status message by setting the Use the checkbox below for a quick retry:
📝 WalkthroughWalkthroughAdds CuTe DSL W4A16 MoE execution with NVFP4 weight decoding, mode-aware APIs, activation-pack plumbing, stream-aware permutation dispatch, expanded tests, and distributed EP/TP profiling benchmarks. Existing W4A4 paths gain explicit quantization-mode and per-token-scale handling. ChangesW4A16 CuTe DSL MoE
Estimated code review effort: 5 (Critical) | ~120 minutes Possibly related PRs
Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request introduces support for BF16 activations with online NVFP4 weight dequantization and a deterministic two-stage finalize mode (non-fused) in the CuTe DSL MoE implementation for Blackwell GPUs. It updates the underlying kernels, bindings, host preparation, and tests to support these new execution paths. The review feedback recommends adding data type validation for moe_output in the BF16 activation path to ensure it is torch.bfloat16 and prevent potential runtime issues.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
Contributor
Author
E2M1/E4M3 to BF16 compatibility experimentI evaluated whether the CUDA 13.2-only direct conversion sequence could be replaced by the older SM120 W4A16 bit-manipulation design without losing performance. Designs comparedThe current path decodes each packed pair with: I tested two older-toolkit-compatible alternatives:
Scale preprocessing and rebasing were cached and occurred outside the measured kernel execution. CorrectnessThe final packed bit-manipulation implementation passed the existing BF16 activation numerical coverage: CUDA_VISIBLE_DEVICES=0 \
FLASHINFER_DISABLE_VERSION_CHECK=1 \
python3 -m pytest -q \
tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeBf16Activation::test_numerical_accuracyResult: PerformanceEnvironment:
Command: CUDA_VISIBLE_DEVICES=0 \
FLASHINFER_DISABLE_VERSION_CHECK=1 \
python3 benchmarks/bench_moe_deepseek.py \
--warmup 10 \
--iters 100 \
--cute-dsl-only \
--use-bf16-activation
The packed bit path reaches 0.841x geomean performance, corresponding to an 18.9% geomean latency regression. Even after matching the SM120 packing strategy, the extra integer permutations, masks, and shifts remain visible in this transform-heavy kernel. DecisionThis is not iso-performance, so I reverted the experiment and kept the CUDA 13.2 direct E2M1/E4M3-to-BF16 instructions. The PR source remains unchanged by this experiment. Supporting an older toolkit through the SM120-style decode would impose a material runtime cost rather than being a compatibility-only implementation choice. |
Contributor
Author
|
Follow-up on the final W4A16 transform-loop cleanup:
The CUDA 13.0 rerun used |
Contributor
Author
|
Follow-up on the sparse W4A16 topology audit: The retained design uses three route-density tiers: N=32 through 16 routed rows/expert, N=64 through 32 rows/expert, and N=128 above that. Sparse and medium DeepSeek GEMM1 shapes use independent 1-SM B200, CUDA 13.2, CUPTI enabled, PDL enabled, 10 warmups, 100 iterations, three-run medians:
The mixed-load alternatives were not retained:
Validation on the final code: CUDA_VISIBLE_DEVICES=0 pytest -q \
tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeBf16Activation
# 8 passed
pre-commit run --all-files
# passed |
Contributor
Author
|
CUDA 13.0 results have been refreshed on the current PR head ( Environment: B200, PyTorch 2.11.0+cu130, CUDA 13.0, CuTe DSL compiler 13.1, CUPTI 13.0.85, PDL enabled, CUDA Graph enabled, 10 warmups, and 100 measured iterations. Each value is the median of three independent runs; all nine runs reported The W4A4 baseline starts from BF16 and includes per-token NVFP4 4over6 MSE quantization, with E4M3 max 256, FP4 quantization fast math disabled, 4over6 error fast math disabled, and deterministic two-stage finalize. W4A16 consumes BF16 directly without activation quantization.
Compatibility-path correctness was rerun across the full focused matrix: The PR body now uses these current post-finalize-refactor CUDA 13.0 results instead of the stale pre-refactor headline table. |
ziang-and
force-pushed
the
agent-cutedsl-moe-bf16-activation
branch
from
July 20, 2026 22:16
9c5f115 to
6dd1ed7
Compare
Contributor
Author
W4A8 prototype: design and performance resultI preserved the complete W4A8 experiment in one commit and then reverted it:
The implementation was numerically functional, and the TMA work materially improved the first version. However, online NVFP4-to-MXFP8 weight conversion still made the final W4A8 strategy substantially slower than W4A4 plus activation quantization. I therefore do not think this implementation should ship in its current form. Prototype contract
Activation and weight pipelinesThe first version followed the W4A4 activation path: no standalone permutation, four The optimized tactic added a separate route permutation and scale-aware permutation. It wrote activation scales directly into the tensor-core 128x4 swizzled layout, then loaded activation data and scales with TMA. This reduced the activation producer to one warp and left the remaining load warps idle. The final version also loaded packed NVFP4 data and E4M3 scales with TMA on one shared transaction barrier, removing per-fragment global scale-index calculations from the transform warps. Four transform warps received 224 registers each for NVFP4 decode, FP32 scaling, and MXFP8 requantization. Optimization experimentsSpeedup is baseline latency divided by experiment latency; values above 1.0 are improvements. These are complete 1-4096-token EP8 and TP8 sweeps.
Other configurations evaluated during this work included raster-M versus raster-N scheduling, 224 versus 240 transform registers, CP-async/TMA GEMM1-GEMM2 mixtures, and larger candidate tactic sets. They were numerically valid but did not improve the full sweep enough to retain. The final autotuner was deliberately reduced to the three tactics above. The 2-CTA path was made strictly numerically correct before performance evaluation. Correct block-scaled MMA required the NVFP4 E4M3 scale fragments to be available to both peer CTAs. The working design replicated the required input/scale state and split packed weight data and weight scales into separate TMA pipelines/barriers. That extra synchronization and pipeline traffic reduced both EP and TP performance by roughly 7-9%. Removing only the 2-CTA tactic did not recover performance; restoring the 1-CTA combined data/scale TMA transaction did. The 2-CTA implementation was therefore removed from the squashed prototype. ValidationThe retained 1-CTA prototype passed: W4A8 coverage included functional and wrapper APIs, PDL on/off, fused/deterministic finalize, EP1 and EP8 local-expert routing, SwiGLU and ReLU2, in-place weight-scale updates, and every retained autotune tactic. The isolated strict 2-CTA case, the full W4A8 class, and the combined W4A8/W4A16 selection also passed before the 2-CTA performance result was rejected. End-to-end benchmarkEnvironment: 8x B200, PyTorch 2.13.0+cu132, CUDA 13.2, CuTe DSL compiler 13.3, CUPTI available, and deterministic finalize. Each row is the median of 100 maximum-rank CUDA-event samples after 10 warmups, with an L2 flush before each sample. EP8 includes real dispatch and combine. TP8 starts from the same distributed logical token set and includes activation all-gather and output all-reduce. W4A4 starts from BF16 and includes CuTe DSL per-token NVFP4 4over6 quantization. W4A8 starts from BF16 and includes CuTe DSL MXFP8 activation quantization. The benchmark command was: FLASHINFER_NVFP4_4OVER6=1 \
FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 \
FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE \
FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 \
FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 \
torchrun --standalone --nproc_per_node=8 \
benchmarks/bench_cute_dsl_moe_distributed.py \
--num-gpus 8 \
--use-per-token-activation \
--no-fused-finalizeAll slowdown columns below use W4A4 EP8 as the strategy baseline, matching the existing distributed comparison.
For a same-topology comparison, W4A8/W4A4 is 2.092x geomean under EP8 and 1.804x under TP8. The W4A16 controls in the same run were 0.913x EP8 and 0.989x TP8 relative to the W4A4 EP8 baseline. Nsight Systems breakdownThe profiler used the same numerical contract and default 32/4096-token cases: FLASHINFER_NVFP4_4OVER6=1 \
FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 \
FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE \
FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 \
FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 \
python3 benchmarks/bench_cute_dsl_moe_distributed.py \
--mode profile_nsys \
--num-gpus 8 \
--use-per-token-activation \
--no-fused-finalizeValues are mean staged GPU-kernel time per rank per captured iteration in milliseconds, obtained by dividing each Nsight aggregate by 8 ranks x 10 iterations. They are diagnostic kernel-time sums, not application wall-clock spans, so concurrent communication is not subtracted. 32 global tokens:
4096 global tokens:
The breakdown makes the limiting factor clear. W4A8 GEMM1 plus GEMM2 account for 71-78% of its staged time in three of the four profiles and 68% in TP8 at 4096 tokens. At 4096 EP8, removing W4A4's two activation-quantization stages saves only 0.073 ms, while W4A8's two online-conversion GEMMs add 1.255 ms relative to W4A4's GEMMs. Communication and standalone quantization are not the main gap; the SIMT NVFP4 -> BF16 -> MXFP8 conversion in both GEMM pipelines is. The TMA work reduced data movement and scale-addressing overhead substantially, but it did not remove that conversion cost. A future W4A8 design would need a materially different weight-conversion contract or substantially better overlap between conversion and MMA, rather than additional launch-configuration tuning around this kernel structure. |
Contributor
Author
W4A16 FP32 weight-scale placement ablationI tested the alternative contract of fully dequantizing each NVFP4 weight to BF16 in the transform/prologue instead of applying the per-expert FP32 scale to the FP32 accumulator in the epilogue. The two matched implementations were:
The ablation was an internal compile-time switch, was included in the compiled-kernel cache key, and independently ran the existing W4A16 autotuner. It did not change the public API, routing, communication, tactics, PDL, activation fusion, or finalize algorithm. Setup
FLASHINFER_DISABLE_VERSION_CHECK=1 \
FLASHINFER_NVFP4_4OVER6=1 \
FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 \
FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE \
FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 \
FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 \
torchrun --standalone --nproc-per-node=8 \
benchmarks/bench_cute_dsl_moe_distributed.py \
--warmup 10 --iters 100 \
--use-per-token-activation --no-fused-finalizeThe prologue run adds
The repeated W4A4 control was stable between the two runs: EP8 prologue-run/baseline geomean was 0.999x and TP8 was 1.001x, so the W4A16 difference is not explained by run-to-run drift. A focused W4A16 route-tile numerical test also passed with the prologue contract ( ConclusionThe prologue contract is neutral only at the smallest one- or two-token points. Across the complete sweep it is 23.6% slower for EP8 and 16.8% slower for TP8 by geomean, with the largest regressions in the 32-512 token region. Moving the scale into the transform path applies FP32 multiply-and-round work to every decoded weight element, while the current design applies one FP32 scale to each already-accumulated output element. This increases pressure on the online-dequantization pipeline, which is already the sensitive part of W4A16. Based on this ablation, we should retain the current FP32-accumulator epilogue scaling contract. |
ziang-and
force-pushed
the
agent-cutedsl-moe-bf16-activation
branch
from
July 23, 2026 23:40
417b07b to
60d18d2
Compare
Contributor
Author
|
@coderabbitai review |
Contributor
✅ Action performedReview finished.
|
zianglih
marked this pull request as ready for review
July 24, 2026 09:18
zianglih
requested review from
Anerudhan,
IwakuraRein,
aleozlx,
jiahanc,
nv-yunzheq,
samuellees and
yzh119
as code owners
July 24, 2026 09:18
Contributor
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
zianglih
requested review from
Aneureka,
bkryu,
cyx-6,
feih-nv,
saltyminty,
sricketts,
yongwww and
yyihuang
as code owners
July 24, 2026 09:18
5 tasks
5 tasks
This was referenced Aug 7, 2026
aleozlx
added a commit
to aleozlx/flashinfer
that referenced
this pull request
Aug 7, 2026
Resolves conflicts with flashinfer-ai#4048 (c9172e6), which added `use_fused_finalize: bool = True` to ExecutionConfig -- the same knob this branch relocates -- and with the CuteDslNvfp4Runner rework that added the NVFP4/W4A16 variant split and an enable_pdl hoist. Resolution: * ExecutionConfig keeps only enable_pdl / tune_max_num_tokens; both finalize knobs now live on MoEFinalizeConfig, which is the point of this branch and of issue flashinfer-ai#4325. * MoEFinalizeConfig.use_fused_finalize adopts main's semantics exactly: `bool = True` rather than this branch's original `Optional[bool] = None` sentinel. main now threads the value explicitly into both inner runners, so the "None means backend default" indirection bought nothing; a plain bool keeps behaviour identical to main. * Both call sites in CuteDslNvfp4Runner (NVFP4 and W4A16 branches) read config.finalize.use_fused_finalize; main's enable_pdl hoist and variant split are kept as-is. Verified: repr round-trip for both configs, ExecutionConfig no longer exposes either finalize knob, use_fused_finalize still defaults True. pre-commit passes. No GPU available here, so no MoE kernel test was run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
aleozlx
added a commit
that referenced
this pull request
Aug 7, 2026
## 📌 Description The rewrite in #4027 changed `tma_load_2d` in `csrc/fused_moe/monomoe/src/ptx_utils.h` from `.shared::cluster` to `.shared::cta`. The `.shared::cta` form requires PTX ISA 8.6 / CUDA 12.8+, breaking the `aot-build-import (cu126)` CI job with `ptxas: State space incorrect for instruction 'cp.async.bulk.tensor'`. Fix: revert to `.shared::cluster` and remove the erroneous `.tile` qualifier (`.tile` is only valid with `.shared::cta` and was never correct on the cluster form). This matches the SM90 fallback in CUTLASS `cute/arch/copy_sm90_tma.hpp`, which uses `shared::cluster` without `.tile` for all pre-SM120 targets. Restores CUDA 12.0+ compatibility. No performance impact: monomoe launches with cluster size 1, so `.shared::cluster` and `.shared::cta` are semantically identical for this kernel. The `.tile` qualifier is a PTX syntax disambiguator only; hardware behavior is unchanged. The first commit added a CUDA >= 12.8 version gate as an initial workaround; the second commit replaces it with this root-cause fix and reverts the gate. ## 🔍 Related Issues Reported in #4048 (comment): `ptxas monomoe_binding.ptx: State space incorrect for instruction 'cp.async.bulk.tensor'` on cu126 AOT build. ## 🚀 Pull Request Checklist - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [ ] `aot-build-import (x64, cu126)` and `aot-build-import (arm64, cu126)` pass in CI - [ ] `aot-build-import (x64/arm64, cu128/cu129/cu130)` continue to pass ## Reviewer Notes The PTX change is a one-liner: `shared::cta.global.tile` → `shared::cluster.global` (dropping `.tile`). CUTLASS's `cute/arch/copy_sm90_tma.hpp` uses exactly this form for the SM90 path. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved compatibility for tensor loading across supported PTX and CUDA versions. * Updated tensor addressing behavior for single-CTA kernel launches. * Improved support across compatible hardware and toolchains. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
5 tasks
This was referenced Aug 10, 2026
Merged
IwakuraRein
pushed a commit
that referenced
this pull request
Aug 16, 2026
## 📌 Description @HumansAnd - Extend CuTe DSL W4A16 fused MoE with tanh-approximate GeGLU and SiTU. - Reuse the activation math and validation introduced for W4A4 in #4009: - GeGLU-tanh: `up * GELU(gate, approximate="tanh")`. - SiTU: `ActivationType.Swiglu` plus `situ_beta`, with optional `situ_linear_beta` for the up-branch clamp. - Apply each expert's W1 alpha before activation, matching the existing W4A4 contract. - Include both SiTU parameters in autotuner and compiled-kernel cache identities. - Keep activation in GEMM1; GEMM2 routing-weight reduction and fused finalize are unchanged. - Extend existing Cartesian tests across W4A4 per-tensor, W4A4 per-token, and W4A16 modes; no new standalone test functions are added. ## 🔍 Related Issues - #4009 - #4048 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see the [pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). B200 source-only validation with `lmsysorg/sglang:nightly-dev-cu13-20260806-ae5f8c94`: - `pre-commit run --all-files` — passed. - GeGLU/SiTU functional and cache-key matrix — 11 passed. - W4A16 GeGLU/SiTU wrapper-autotune rows — 2 passed, 6 deselected. - `moe_utils is_aot: False`; editable-install data links resolved inside the synced source checkout. ## Reviewer Notes - SiTU intentionally remains `ActivationType.Swiglu` plus its beta parameters, matching #4009. - Scope follows #4009's functional and wrapper APIs; the unified `CuteDslNvfp4Runner` activation configuration is unchanged. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added support for SiLU-in-the-unit (SiTU) parameters in W4A16 mixture-of-experts execution. * Added GeGLU-Tanh activation support for applicable gated operations. * Improved tuning and cache handling for activation and quantization configurations. * **Bug Fixes** * Corrected activation parameter handling across wrapper and functional execution paths. * **Tests** * Expanded accuracy and autotuning coverage across W4A4 and W4A16 modes. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
This was referenced Aug 17, 2026
1 task
IwakuraRein
pushed a commit
that referenced
this pull request
Aug 18, 2026
## 📌 Description @HumansAnd Fixes W4A16 CuTe DSL MoE scratch-workspace retention exposed by the companion SGLang Nemotron-3 integration. - **Root cause:** each `CuteDslFusedMoEW4A16Runner` retained its sort, permuted-hidden, and intermediate BF16 tensors in `_workspace_cache`. The key includes the route tile, so an autotuned runner could retain six large arenas; model serving then retained a selected arena through every layer-owned wrapper. These are live tensor references, so `torch.cuda.empty_cache()` cannot release them. - **Observed geometry:** at `M=16384`, `top_k=22`, 128 local experts, hidden 1024, intermediate 2688, and route tile 128, one runner retains 735.75 MiB of hidden scratch plus 1,931.34 MiB of intermediate scratch and sort buffers. The reproduced 1.89 GiB OOM is the intermediate allocation. - **Fix:** make W4A16 scratch invocation-local, matching the established W4A4 wrapper path. PyTorch's caching allocator reuses eager allocations after the current-stream work completes, while CUDA graph-private pools own captured storage and replay addresses. - **API scope:** no shared/module-global arena, no wrapper sharing, and no new stream or graph serialization contract. The top-level wrapper, kernel inputs, tactics, outputs, and numerical contract are unchanged; the lower-level runner-only `launch_w4a16_moe(..., workspace_cache=...)` keyword is removed. - **Validation history:** two temporary behavioral tests at [`678867c3`](zianglih@678867c) verified eager workspace lifetime and two-call CUDA-graph replay. [`c5ea45e4`](zianglih@c5ea45e) removes them from the final diff for a smaller PR while preserving them in branch history. - **Consumer integration:** [sgl-project/sglang#35120](sgl-project/sglang#35120) depends on this fix and validates online-quantized Nemotron W4A16 accuracy plus serialized Qwen3 W4A16 disk-reload invariance. ## 🔍 Related Issues - W4A16 support was introduced in [#4048](#4048); Nemotron SiTU support was added in [#4394](#4394). - Required by [sgl-project/sglang#35120](sgl-project/sglang#35120). ## Validation Environment: - Image: `lmsysorg/sglang:nightly-dev-cu13-20260817-d91c3682` - Image digest: `sha256:fa8774dd128600a09fd6d46670b06fb69a55dac8a3881e50ccf0916a45eb39af` - Hardware: 8x NVIDIA B300 (SM103), driver `590.48.01` - PyTorch: `2.13.0+cu130`; CUDA: `13.0`; NCCL: `2.29.7` - FlashInfer base: `nightly-v0.6.18-20260817` / `e77a4a0d276367895c3b50a642fd8f326c03fb72` - FlashInfer validation commit: `678867c392be6d76f8004e8533c379eb0c222e0f`; final head: `c5ea45e4422e4c8be8059d79b27581971891132a` - CuTe DSL: `nvidia-cutlass-dsl==4.7.0`; Nsight Systems: `2026.4.1` ### Workspace lifetime and CUDA graph validation The following ran at implementation commit `678867c3`; the final head changes only by removing these temporary tests. ```bash CUDA_VISIBLE_DEVICES=0 \ PYTHONPATH=/hai-workspace/flashinfer-runtime-w4a16-fix:/hai-workspace/flashinfer-w4a16-workspace \ python3 -m pytest -vv -s \ /hai-workspace/flashinfer-w4a16-workspace/tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeW4A16::test_wrapper_releases_eager_workspace_with_live_wrappers \ /hai-workspace/flashinfer-w4a16-workspace/tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeW4A16::test_wrapper_cuda_graph_owns_invocation_workspace ``` ```text # Targeted workspace-lifetime and CUDA-graph regression test_wrapper_releases_eager_workspace_with_live_wrappers PASSED test_wrapper_cuda_graph_owns_invocation_workspace PASSED 2 passed, 70 warnings in 918.73s # Full TestCuteDslMoeW4A16 class (14 existing + 2 temporary tests) 16 passed, 598 warnings in 1133.16s ``` All reported warnings are CuTe DSL or dependency deprecation warnings. ### SGLang Nemotron-3 end-to-end The companion SGLang PR validates online expert-weight quantization with the default 16,384-token prefill-graph capacity and `--mem-fraction-static 0.7`. ```text # Online FP8 checkpoint with online NVFP4 expert-weight quantization Capture target prefill CUDA graph end. elapsed=205.47 s, mem usage=4.74 GB, avail mem=56.87 GB. max_total_num_tokens=30259081, chunked_prefill_size=16384, max_prefill_tokens=16384, available_gpu_mem=54.84 GB Total latency: 127.126 s Score: 0.975 Output throughput: 631.254 token/s avg_spec_accept_length=3.1296052121306137 Ran 1 test in 1890.684s OK ``` The serialized ModelOpt Qwen3 test also completed two same-checkpoint disk reloads across weight/CUDA-graph release and resume while preserving deterministic decode text, token IDs, and logprobs: ```text # Serialized nvidia/Qwen3-30B-A3B-NVFP4 W4A16 disk reload Ran 1 test in 1205.499s OK ``` ### PR #4048 distributed benchmark compatibility This compatibility rerun executed at final PR head `c5ea45e4422e4c8be8059d79b27581971891132a`. `benchmarks/bench_cute_dsl_moe_distributed.py` is byte-identical between #4048's merged head and this head (`sha256:f61bd31757758bc661ddc94551d04c8a762a27612e0fad10c74c1c3240cc87fa`). The imported runtime copy is also byte-identical to the final head for both changed source files: - `flashinfer/fused_moe/cute_dsl/blackwell/moe_w4a16.py`: `sha256:cd11e0c1c37653b0a58403de04fa31ea20ae75a6ebab1b8330e7091e32d729f7` - `flashinfer/fused_moe/cute_dsl/tuner.py`: `sha256:c2dbbad8f2643f1fcd7f0677ec759d150910f6c44418f261cebd5dcba4c2fd39` Each of the three complete runs uses 10 warmups and 100 cold-L2, maximum-rank CUDA-event samples per cell. The HAI pod additionally needs `PET_LOCAL_ADDR=127.0.0.1` because its hostname is not resolvable in pod-local DNS. `--allreduce-backend auto` selected FlashInfer MNNVL, confirmed by a separate one-token `--verbose` diagnostic in the same environment. ```bash set -euo pipefail export FLASHINFER_SOURCE=/hai-workspace/flashinfer-w4a16-workspace export FLASHINFER_RUNTIME=/hai-workspace/flashinfer-runtime-w4a16-fix export RESULTS_DIR=/hai-workspace/sglang-cutedsl-w4a16-results/pr4048-distributed-c5ea-20260817b export PYTHONPATH="${FLASHINFER_RUNTIME}:${FLASHINFER_SOURCE}" export PET_LOCAL_ADDR=127.0.0.1 export FLASHINFER_DISABLE_VERSION_CHECK=1 export FLASHINFER_NVFP4_4OVER6=1 export FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 export FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE export FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 export FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 mkdir -p "${RESULTS_DIR}" for run in 1 2 3; do torchrun --standalone --nproc-per-node=8 \ "${FLASHINFER_SOURCE}/benchmarks/bench_cute_dsl_moe_distributed.py" \ --warmup 10 --iters 100 \ --use-per-token-activation --no-fused-finalize \ 2>&1 | tee "${RESULTS_DIR}/end-to-end-run-${run}.log" done ``` All three runs emitted the expected 52 rows: 13 token counts x EP8/TP8 x W4A4/W4A16. <details> <summary>Run 1: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.168512 DISTRIBUTED_CSV,ep,w4a16,1,8,0.944432 DISTRIBUTED_CSV,ep,w4a4,2,8,1.226576 DISTRIBUTED_CSV,ep,w4a16,2,8,0.957840 DISTRIBUTED_CSV,ep,w4a4,4,8,1.222032 DISTRIBUTED_CSV,ep,w4a16,4,8,0.916528 DISTRIBUTED_CSV,ep,w4a4,8,8,1.184704 DISTRIBUTED_CSV,ep,w4a16,8,8,0.963344 DISTRIBUTED_CSV,ep,w4a4,16,8,1.224800 DISTRIBUTED_CSV,ep,w4a16,16,8,0.951408 DISTRIBUTED_CSV,ep,w4a4,32,8,1.234592 DISTRIBUTED_CSV,ep,w4a16,32,8,0.967936 DISTRIBUTED_CSV,ep,w4a4,64,8,1.239088 DISTRIBUTED_CSV,ep,w4a16,64,8,0.958080 DISTRIBUTED_CSV,ep,w4a4,128,8,1.227296 DISTRIBUTED_CSV,ep,w4a16,128,8,0.943760 DISTRIBUTED_CSV,ep,w4a4,256,8,1.228256 DISTRIBUTED_CSV,ep,w4a16,256,8,0.955056 DISTRIBUTED_CSV,ep,w4a4,512,8,1.220816 DISTRIBUTED_CSV,ep,w4a16,512,8,0.949472 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.217200 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.950672 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.226176 DISTRIBUTED_CSV,ep,w4a16,2048,8,1.002928 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.270112 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.163104 DISTRIBUTED_CSV,tp,w4a4,1,8,1.246640 DISTRIBUTED_CSV,tp,w4a16,1,8,0.929936 DISTRIBUTED_CSV,tp,w4a4,2,8,1.193216 DISTRIBUTED_CSV,tp,w4a16,2,8,0.905024 DISTRIBUTED_CSV,tp,w4a4,4,8,1.191968 DISTRIBUTED_CSV,tp,w4a16,4,8,0.914336 DISTRIBUTED_CSV,tp,w4a4,8,8,1.232912 DISTRIBUTED_CSV,tp,w4a16,8,8,1.005168 DISTRIBUTED_CSV,tp,w4a4,16,8,1.275024 DISTRIBUTED_CSV,tp,w4a16,16,8,0.988672 DISTRIBUTED_CSV,tp,w4a4,32,8,1.302352 DISTRIBUTED_CSV,tp,w4a16,32,8,0.991168 DISTRIBUTED_CSV,tp,w4a4,64,8,1.287776 DISTRIBUTED_CSV,tp,w4a16,64,8,1.007712 DISTRIBUTED_CSV,tp,w4a4,128,8,1.288544 DISTRIBUTED_CSV,tp,w4a16,128,8,0.976032 DISTRIBUTED_CSV,tp,w4a4,256,8,1.294512 DISTRIBUTED_CSV,tp,w4a16,256,8,1.014432 DISTRIBUTED_CSV,tp,w4a4,512,8,1.325776 DISTRIBUTED_CSV,tp,w4a16,512,8,1.021552 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.333536 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.087632 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.386016 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.226896 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.534096 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.496576 ``` </details> <details> <summary>Run 2: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.198144 DISTRIBUTED_CSV,ep,w4a16,1,8,0.936384 DISTRIBUTED_CSV,ep,w4a4,2,8,1.208656 DISTRIBUTED_CSV,ep,w4a16,2,8,0.914576 DISTRIBUTED_CSV,ep,w4a4,4,8,1.227664 DISTRIBUTED_CSV,ep,w4a16,4,8,0.951744 DISTRIBUTED_CSV,ep,w4a4,8,8,1.208256 DISTRIBUTED_CSV,ep,w4a16,8,8,1.068048 DISTRIBUTED_CSV,ep,w4a4,16,8,1.231184 DISTRIBUTED_CSV,ep,w4a16,16,8,0.946000 DISTRIBUTED_CSV,ep,w4a4,32,8,1.234944 DISTRIBUTED_CSV,ep,w4a16,32,8,0.953472 DISTRIBUTED_CSV,ep,w4a4,64,8,1.382608 DISTRIBUTED_CSV,ep,w4a16,64,8,0.955488 DISTRIBUTED_CSV,ep,w4a4,128,8,1.217792 DISTRIBUTED_CSV,ep,w4a16,128,8,0.942240 DISTRIBUTED_CSV,ep,w4a4,256,8,1.223472 DISTRIBUTED_CSV,ep,w4a16,256,8,0.939152 DISTRIBUTED_CSV,ep,w4a4,512,8,1.224704 DISTRIBUTED_CSV,ep,w4a16,512,8,0.953024 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.231840 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.955184 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.230240 DISTRIBUTED_CSV,ep,w4a16,2048,8,0.984656 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.246976 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.152400 DISTRIBUTED_CSV,tp,w4a4,1,8,1.205248 DISTRIBUTED_CSV,tp,w4a16,1,8,0.952688 DISTRIBUTED_CSV,tp,w4a4,2,8,1.221104 DISTRIBUTED_CSV,tp,w4a16,2,8,0.927984 DISTRIBUTED_CSV,tp,w4a4,4,8,1.222752 DISTRIBUTED_CSV,tp,w4a16,4,8,0.930240 DISTRIBUTED_CSV,tp,w4a4,8,8,1.264928 DISTRIBUTED_CSV,tp,w4a16,8,8,1.002752 DISTRIBUTED_CSV,tp,w4a4,16,8,1.280784 DISTRIBUTED_CSV,tp,w4a16,16,8,0.994400 DISTRIBUTED_CSV,tp,w4a4,32,8,1.280864 DISTRIBUTED_CSV,tp,w4a16,32,8,1.002256 DISTRIBUTED_CSV,tp,w4a4,64,8,1.291056 DISTRIBUTED_CSV,tp,w4a16,64,8,1.005040 DISTRIBUTED_CSV,tp,w4a4,128,8,1.281040 DISTRIBUTED_CSV,tp,w4a16,128,8,1.001472 DISTRIBUTED_CSV,tp,w4a4,256,8,1.286368 DISTRIBUTED_CSV,tp,w4a16,256,8,0.997664 DISTRIBUTED_CSV,tp,w4a4,512,8,1.297824 DISTRIBUTED_CSV,tp,w4a16,512,8,1.012064 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.319184 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.068624 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.358832 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.234288 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.534160 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.490592 ``` </details> <details> <summary>Run 3: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.154336 DISTRIBUTED_CSV,ep,w4a16,1,8,0.908048 DISTRIBUTED_CSV,ep,w4a4,2,8,1.202480 DISTRIBUTED_CSV,ep,w4a16,2,8,0.942320 DISTRIBUTED_CSV,ep,w4a4,4,8,1.217984 DISTRIBUTED_CSV,ep,w4a16,4,8,0.925856 DISTRIBUTED_CSV,ep,w4a4,8,8,1.219968 DISTRIBUTED_CSV,ep,w4a16,8,8,0.925536 DISTRIBUTED_CSV,ep,w4a4,16,8,1.211344 DISTRIBUTED_CSV,ep,w4a16,16,8,0.944256 DISTRIBUTED_CSV,ep,w4a4,32,8,1.227168 DISTRIBUTED_CSV,ep,w4a16,32,8,0.924864 DISTRIBUTED_CSV,ep,w4a4,64,8,1.209744 DISTRIBUTED_CSV,ep,w4a16,64,8,0.942160 DISTRIBUTED_CSV,ep,w4a4,128,8,1.222128 DISTRIBUTED_CSV,ep,w4a16,128,8,0.950096 DISTRIBUTED_CSV,ep,w4a4,256,8,1.228112 DISTRIBUTED_CSV,ep,w4a16,256,8,0.955936 DISTRIBUTED_CSV,ep,w4a4,512,8,1.238400 DISTRIBUTED_CSV,ep,w4a16,512,8,0.946848 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.222480 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.950592 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.247856 DISTRIBUTED_CSV,ep,w4a16,2048,8,0.995744 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.271680 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.173120 DISTRIBUTED_CSV,tp,w4a4,1,8,1.248944 DISTRIBUTED_CSV,tp,w4a16,1,8,0.931824 DISTRIBUTED_CSV,tp,w4a4,2,8,1.218480 DISTRIBUTED_CSV,tp,w4a16,2,8,0.910384 DISTRIBUTED_CSV,tp,w4a4,4,8,1.222864 DISTRIBUTED_CSV,tp,w4a16,4,8,0.919216 DISTRIBUTED_CSV,tp,w4a4,8,8,1.287648 DISTRIBUTED_CSV,tp,w4a16,8,8,0.991392 DISTRIBUTED_CSV,tp,w4a4,16,8,1.270656 DISTRIBUTED_CSV,tp,w4a16,16,8,0.994000 DISTRIBUTED_CSV,tp,w4a4,32,8,1.291312 DISTRIBUTED_CSV,tp,w4a16,32,8,0.996016 DISTRIBUTED_CSV,tp,w4a4,64,8,1.303648 DISTRIBUTED_CSV,tp,w4a16,64,8,0.994432 DISTRIBUTED_CSV,tp,w4a4,128,8,1.294864 DISTRIBUTED_CSV,tp,w4a16,128,8,1.004208 DISTRIBUTED_CSV,tp,w4a4,256,8,1.309296 DISTRIBUTED_CSV,tp,w4a16,256,8,1.022560 DISTRIBUTED_CSV,tp,w4a4,512,8,1.310064 DISTRIBUTED_CSV,tp,w4a16,512,8,1.026016 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.332736 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.082720 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.383648 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.229296 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.546464 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.511744 ``` </details> The table below takes the median across the three complete runs for each cell. `speedup = median(W4A4 EP8) / median(strategy)`, so values above 1 favor the strategy. | Tokens | W4A4 EP8 ms | W4A4 TP8 ms | TP speedup | W4A16 EP8 ms | EP speedup | W4A16 TP8 ms | TP speedup | |---:|---:|---:|---:|---:|---:|---:|---:| | 1 | 1.169 | 1.247 | 0.937x | 0.936 | 1.248x | 0.932 | 1.254x | | 2 | 1.209 | 1.218 | 0.992x | 0.942 | 1.283x | 0.910 | 1.328x | | 4 | 1.222 | 1.223 | 0.999x | 0.926 | 1.320x | 0.919 | 1.329x | | 8 | 1.208 | 1.265 | 0.955x | 0.963 | 1.254x | 1.003 | 1.205x | | 16 | 1.225 | 1.275 | 0.961x | 0.946 | 1.295x | 0.994 | 1.232x | | 32 | 1.235 | 1.291 | 0.956x | 0.953 | 1.295x | 0.996 | 1.240x | | 64 | 1.239 | 1.291 | 0.960x | 0.955 | 1.297x | 1.005 | 1.233x | | 128 | 1.222 | 1.289 | 0.948x | 0.944 | 1.295x | 1.001 | 1.220x | | 256 | 1.228 | 1.295 | 0.949x | 0.955 | 1.286x | 1.014 | 1.211x | | 512 | 1.225 | 1.310 | 0.935x | 0.949 | 1.290x | 1.022 | 1.199x | | 1024 | 1.222 | 1.333 | 0.917x | 0.951 | 1.286x | 1.083 | 1.129x | | 2048 | 1.230 | 1.384 | 0.889x | 0.996 | 1.235x | 1.229 | 1.001x | | 4096 | 1.270 | 1.534 | 0.828x | 1.163 | 1.092x | 1.497 | 0.849x | Geometric-mean speedups versus W4A4 EP8 are `0.940x` for W4A4 TP8, `1.266x` for W4A16 EP8, and `1.179x` for W4A16 TP8. Using the same exported environment above, the Nsight Systems driver launches its own eight-rank workers for all eight cases: two token counts x EP8/TP8 x W4A4/W4A16. ```bash python3 "${FLASHINFER_SOURCE}/benchmarks/bench_cute_dsl_moe_distributed.py" \ --mode profile_nsys --warmup 10 --profile-iters 20 \ --use-per-token-activation --no-fused-finalize \ --nsys-output-dir "${RESULTS_DIR}/nsys" \ 2>&1 | tee "${RESULTS_DIR}/nsys.log" ``` <details> <summary>32-token cases: four raw Nsight Systems stage tables</summary> ```text Nsight Systems kernel breakdown: EP8 W4A4, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.35% | 0.559 | 160 | 3.491 | deepseek_v3_topk_kernel dispatch | 1 | 0.15% | 0.240 | 160 | 1.502 | moeA2APrepareDispatchKernel dispatch | 2 | 19.77% | 31.380 | 160 | 196.127 | moeA2ADispatchKernel dispatch | 3 | 11.53% | 18.301 | 160 | 114.379 | moeA2ASanitizeExpertIdsKernel activation prep/quant | 1 | 0.11% | 0.169 | 160 | 1.059 | vectorized_elementwise_kernel activation prep/quant | 2 | 0.47% | 0.750 | 160 | 4.688 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.39% | 0.619 | 160 | 3.871 | routingIndicesClusterKernel local MoE | 2 | 7.29% | 11.573 | 160 | 72.334 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.12% | 0.197 | 160 | 1.230 | vectorized_elementwise_kernel local MoE | 4 | 0.81% | 1.282 | 160 | 8.013 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 3.28% | 5.204 | 160 | 32.524 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 0.88% | 1.400 | 160 | 8.750 | moeUnpermuteKernel combine | 1 | 0.27% | 0.429 | 160 | 2.682 | moeA2APrepareCombineKernel combine | 2 | 54.58% | 86.649 | 160 | 541.557 | moeA2ACombineKernel Nsight Systems kernel breakdown: EP8 W4A16, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.66% | 0.557 | 160 | 3.483 | deepseek_v3_topk_kernel dispatch | 1 | 0.29% | 0.240 | 160 | 1.498 | moeA2APrepareDispatchKernel dispatch | 2 | 21.84% | 18.310 | 160 | 114.439 | moeA2ADispatchKernel dispatch | 3 | 4.86% | 4.071 | 160 | 25.444 | moeA2ASanitizeExpertIdsKernel local MoE | 1 | 0.76% | 0.639 | 160 | 3.995 | routingIndicesClusterKernel local MoE | 2 | 0.78% | 0.656 | 160 | 4.099 | moePermuteKernel local MoE | 3 | 12.99% | 10.896 | 160 | 68.099 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 7.05% | 5.910 | 160 | 36.940 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 1.34% | 1.125 | 160 | 7.032 | moeUnpermuteKernel combine | 1 | 0.47% | 0.390 | 160 | 2.440 | moeA2APrepareCombineKernel combine | 2 | 48.96% | 41.055 | 160 | 256.594 | moeA2ACombineKernel Nsight Systems kernel breakdown: TP8 W4A4, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 13.30% | 35.050 | 160 | 219.062 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.21% | 0.566 | 160 | 3.535 | deepseek_v3_topk_kernel activation prep/quant | 1 | 0.06% | 0.153 | 160 | 0.954 | vectorized_elementwise_kernel activation prep/quant | 2 | 0.27% | 0.706 | 160 | 4.411 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.23% | 0.600 | 160 | 3.747 | routingIndicesClusterKernel local MoE | 2 | 4.46% | 11.747 | 160 | 73.420 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.08% | 0.201 | 160 | 1.257 | vectorized_elementwise_kernel local MoE | 4 | 2.18% | 5.746 | 160 | 35.910 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 2.35% | 6.188 | 160 | 38.673 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 1.09% | 2.865 | 160 | 17.906 | moeUnpermuteKernel all-reduce | 1 | 75.78% | 199.723 | 160 | 1248.270 | twoshotAllreduceKernel Nsight Systems kernel breakdown: TP8 W4A16, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 18.23% | 17.083 | 160 | 106.768 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.61% | 0.568 | 160 | 3.547 | deepseek_v3_topk_kernel local MoE | 1 | 0.63% | 0.591 | 160 | 3.692 | routingIndicesClusterKernel local MoE | 2 | 0.87% | 0.815 | 160 | 5.093 | moePermuteKernel local MoE | 3 | 11.64% | 10.909 | 160 | 68.178 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 7.32% | 6.863 | 160 | 42.897 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 2.21% | 2.067 | 160 | 12.920 | moeUnpermuteKernel all-reduce | 1 | 58.50% | 54.823 | 160 | 342.642 | twoshotAllreduceKernel ``` </details> <details> <summary>4,096-token cases: four raw Nsight Systems stage tables</summary> ```text Nsight Systems kernel breakdown: EP8 W4A4, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.39% | 0.663 | 160 | 4.143 | deepseek_v3_topk_kernel dispatch | 1 | 0.15% | 0.247 | 160 | 1.546 | moeA2APrepareDispatchKernel dispatch | 2 | 18.26% | 30.738 | 160 | 192.110 | moeA2ADispatchKernel dispatch | 3 | 9.38% | 15.786 | 160 | 98.664 | moeA2ASanitizeExpertIdsKernel activation prep/quant | 1 | 0.15% | 0.259 | 160 | 1.620 | vectorized_elementwise_kernel activation prep/quant | 2 | 2.04% | 3.435 | 160 | 21.470 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.63% | 1.060 | 160 | 6.623 | routingIndicesClusterKernel local MoE | 2 | 10.34% | 17.407 | 160 | 108.791 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.23% | 0.385 | 160 | 2.405 | vectorized_elementwise_kernel local MoE | 4 | 5.91% | 9.940 | 160 | 62.123 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 5.83% | 9.808 | 160 | 61.299 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 3.10% | 5.219 | 160 | 32.616 | moeUnpermuteKernel combine | 1 | 0.93% | 1.571 | 160 | 9.817 | moeA2APrepareCombineKernel combine | 2 | 42.66% | 71.811 | 160 | 448.816 | moeA2ACombineKernel Nsight Systems kernel breakdown: EP8 W4A16, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.45% | 0.658 | 160 | 4.112 | deepseek_v3_topk_kernel dispatch | 1 | 0.16% | 0.240 | 160 | 1.500 | moeA2APrepareDispatchKernel dispatch | 2 | 13.16% | 19.373 | 160 | 121.083 | moeA2ADispatchKernel dispatch | 3 | 2.70% | 3.982 | 160 | 24.887 | moeA2ASanitizeExpertIdsKernel local MoE | 1 | 0.71% | 1.039 | 160 | 6.494 | routingIndicesClusterKernel local MoE | 2 | 2.20% | 3.241 | 160 | 20.257 | moePermuteKernel local MoE | 3 | 22.48% | 33.100 | 160 | 206.873 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 14.91% | 21.957 | 160 | 137.229 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 3.47% | 5.106 | 160 | 31.915 | moeUnpermuteKernel combine | 1 | 1.44% | 2.125 | 160 | 13.279 | moeA2APrepareCombineKernel combine | 2 | 38.32% | 56.418 | 160 | 352.613 | moeA2ACombineKernel Nsight Systems kernel breakdown: TP8 W4A4, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 20.71% | 49.212 | 160 | 307.574 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.68% | 1.627 | 160 | 10.170 | deepseek_v3_topk_kernel activation prep/quant | 1 | 0.10% | 0.245 | 160 | 1.534 | vectorized_elementwise_kernel activation prep/quant | 2 | 2.01% | 4.774 | 160 | 29.837 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.53% | 1.258 | 160 | 7.864 | routingIndicesClusterKernel local MoE | 2 | 8.26% | 19.630 | 160 | 122.688 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.09% | 0.215 | 160 | 1.345 | vectorized_elementwise_kernel local MoE | 4 | 4.90% | 11.640 | 160 | 72.748 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 9.40% | 22.334 | 160 | 139.585 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 7.93% | 18.841 | 160 | 117.757 | moeUnpermuteKernel all-reduce | 1 | 45.37% | 107.793 | 160 | 673.706 | twoshotAllreduceKernel Nsight Systems kernel breakdown: TP8 W4A16, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 14.75% | 31.958 | 160 | 199.739 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.75% | 1.632 | 160 | 10.197 | deepseek_v3_topk_kernel local MoE | 1 | 0.57% | 1.243 | 160 | 7.769 | routingIndicesClusterKernel local MoE | 2 | 9.78% | 21.189 | 160 | 132.433 | moePermuteKernel local MoE | 3 | 16.70% | 36.189 | 160 | 226.178 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 18.50% | 40.095 | 160 | 250.596 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 8.52% | 18.467 | 160 | 115.418 | moeUnpermuteKernel all-reduce | 1 | 30.42% | 65.919 | 160 | 411.993 | twoshotAllreduceKernel ``` </details> The raw `total (ms)` columns above aggregate all eight ranks and 20 captured iterations. Dividing each case's sum by `8 * 20 = 160` gives: | Tokens | Mode | W4A4 staged ms | W4A16 staged ms | W4A4 / W4A16 | |---:|:---:|---:|---:|---:| | 32 | EP8 | 0.992 | 0.524 | 1.893x | | 32 | TP8 | 1.647 | 0.586 | 2.812x | | 4096 | EP8 | 1.052 | 0.920 | 1.143x | | 4096 | TP8 | 1.485 | 1.354 | 1.096x | Limitations: - This is a current-head compatibility run, not a base-vs-head performance A/B. It does not by itself prove that invocation-local allocation has zero latency overhead. - It ran on B300/SM103 with CUDA 13.0; #4048's published tables used B200 with different CUDA/runtime revisions, so absolute timings are not directly comparable. - Two end-to-end cells had isolated run ranges around 14% (EP8 W4A16 at 8 tokens and EP8 W4A4 at 64 tokens); use the reported three-run medians and avoid sub-percent conclusions. - Nsight staged sums are diagnostic kernel-time aggregates, not wall-clock latency, and do not subtract overlapping kernels. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes Please focus on the workspace lifetime and CUDA-graph ownership contract. This intentionally follows W4A4's invocation-local scratch pattern instead of adding a process-global or cross-wrapper workspace cache. The final diff contains no new test code; the validation-only behavioral tests and their passing results remain reviewable at `678867c3`. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved W4A16 mixture-of-experts operation reliability by allocating required working memory for each invocation. * Improved handling of workloads with varying sizes by removing reliance on previously cached capacity. * **Refactor** * Simplified workspace management while preserving existing operation behavior. * Streamlined the W4A16 execution interface for more predictable invocation behavior. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
5 tasks
jefby
pushed a commit
to jefby/flashinfer
that referenced
this pull request
Aug 19, 2026
## 📌 Description @HumansAnd - Extend CuTe DSL W4A16 fused MoE with tanh-approximate GeGLU and SiTU. - Reuse the activation math and validation introduced for W4A4 in flashinfer-ai#4009: - GeGLU-tanh: `up * GELU(gate, approximate="tanh")`. - SiTU: `ActivationType.Swiglu` plus `situ_beta`, with optional `situ_linear_beta` for the up-branch clamp. - Apply each expert's W1 alpha before activation, matching the existing W4A4 contract. - Include both SiTU parameters in autotuner and compiled-kernel cache identities. - Keep activation in GEMM1; GEMM2 routing-weight reduction and fused finalize are unchanged. - Extend existing Cartesian tests across W4A4 per-tensor, W4A4 per-token, and W4A16 modes; no new standalone test functions are added. ## 🔍 Related Issues - flashinfer-ai#4009 - flashinfer-ai#4048 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see the [pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). B200 source-only validation with `lmsysorg/sglang:nightly-dev-cu13-20260806-ae5f8c94`: - `pre-commit run --all-files` — passed. - GeGLU/SiTU functional and cache-key matrix — 11 passed. - W4A16 GeGLU/SiTU wrapper-autotune rows — 2 passed, 6 deselected. - `moe_utils is_aot: False`; editable-install data links resolved inside the synced source checkout. ## Reviewer Notes - SiTU intentionally remains `ActivationType.Swiglu` plus its beta parameters, matching flashinfer-ai#4009. - Scope follows flashinfer-ai#4009's functional and wrapper APIs; the unified `CuteDslNvfp4Runner` activation configuration is unchanged. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added support for SiLU-in-the-unit (SiTU) parameters in W4A16 mixture-of-experts execution. * Added GeGLU-Tanh activation support for applicable gated operations. * Improved tuning and cache handling for activation and quantization configurations. * **Bug Fixes** * Corrected activation parameter handling across wrapper and functional execution paths. * **Tests** * Expanded accuracy and autotuning coverage across W4A4 and W4A16 modes. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
jefby
pushed a commit
to jefby/flashinfer
that referenced
this pull request
Aug 19, 2026
## 📌 Description @HumansAnd Fixes W4A16 CuTe DSL MoE scratch-workspace retention exposed by the companion SGLang Nemotron-3 integration. - **Root cause:** each `CuteDslFusedMoEW4A16Runner` retained its sort, permuted-hidden, and intermediate BF16 tensors in `_workspace_cache`. The key includes the route tile, so an autotuned runner could retain six large arenas; model serving then retained a selected arena through every layer-owned wrapper. These are live tensor references, so `torch.cuda.empty_cache()` cannot release them. - **Observed geometry:** at `M=16384`, `top_k=22`, 128 local experts, hidden 1024, intermediate 2688, and route tile 128, one runner retains 735.75 MiB of hidden scratch plus 1,931.34 MiB of intermediate scratch and sort buffers. The reproduced 1.89 GiB OOM is the intermediate allocation. - **Fix:** make W4A16 scratch invocation-local, matching the established W4A4 wrapper path. PyTorch's caching allocator reuses eager allocations after the current-stream work completes, while CUDA graph-private pools own captured storage and replay addresses. - **API scope:** no shared/module-global arena, no wrapper sharing, and no new stream or graph serialization contract. The top-level wrapper, kernel inputs, tactics, outputs, and numerical contract are unchanged; the lower-level runner-only `launch_w4a16_moe(..., workspace_cache=...)` keyword is removed. - **Validation history:** two temporary behavioral tests at [`678867c3`](zianglih@678867c) verified eager workspace lifetime and two-call CUDA-graph replay. [`c5ea45e4`](zianglih@c5ea45e) removes them from the final diff for a smaller PR while preserving them in branch history. - **Consumer integration:** [sgl-project/sglang#35120](sgl-project/sglang#35120) depends on this fix and validates online-quantized Nemotron W4A16 accuracy plus serialized Qwen3 W4A16 disk-reload invariance. ## 🔍 Related Issues - W4A16 support was introduced in [flashinfer-ai#4048](flashinfer-ai#4048); Nemotron SiTU support was added in [flashinfer-ai#4394](flashinfer-ai#4394). - Required by [sgl-project/sglang#35120](sgl-project/sglang#35120). ## Validation Environment: - Image: `lmsysorg/sglang:nightly-dev-cu13-20260817-d91c3682` - Image digest: `sha256:fa8774dd128600a09fd6d46670b06fb69a55dac8a3881e50ccf0916a45eb39af` - Hardware: 8x NVIDIA B300 (SM103), driver `590.48.01` - PyTorch: `2.13.0+cu130`; CUDA: `13.0`; NCCL: `2.29.7` - FlashInfer base: `nightly-v0.6.18-20260817` / `e77a4a0d276367895c3b50a642fd8f326c03fb72` - FlashInfer validation commit: `678867c392be6d76f8004e8533c379eb0c222e0f`; final head: `c5ea45e4422e4c8be8059d79b27581971891132a` - CuTe DSL: `nvidia-cutlass-dsl==4.7.0`; Nsight Systems: `2026.4.1` ### Workspace lifetime and CUDA graph validation The following ran at implementation commit `678867c3`; the final head changes only by removing these temporary tests. ```bash CUDA_VISIBLE_DEVICES=0 \ PYTHONPATH=/hai-workspace/flashinfer-runtime-w4a16-fix:/hai-workspace/flashinfer-w4a16-workspace \ python3 -m pytest -vv -s \ /hai-workspace/flashinfer-w4a16-workspace/tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeW4A16::test_wrapper_releases_eager_workspace_with_live_wrappers \ /hai-workspace/flashinfer-w4a16-workspace/tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoeW4A16::test_wrapper_cuda_graph_owns_invocation_workspace ``` ```text # Targeted workspace-lifetime and CUDA-graph regression test_wrapper_releases_eager_workspace_with_live_wrappers PASSED test_wrapper_cuda_graph_owns_invocation_workspace PASSED 2 passed, 70 warnings in 918.73s # Full TestCuteDslMoeW4A16 class (14 existing + 2 temporary tests) 16 passed, 598 warnings in 1133.16s ``` All reported warnings are CuTe DSL or dependency deprecation warnings. ### SGLang Nemotron-3 end-to-end The companion SGLang PR validates online expert-weight quantization with the default 16,384-token prefill-graph capacity and `--mem-fraction-static 0.7`. ```text # Online FP8 checkpoint with online NVFP4 expert-weight quantization Capture target prefill CUDA graph end. elapsed=205.47 s, mem usage=4.74 GB, avail mem=56.87 GB. max_total_num_tokens=30259081, chunked_prefill_size=16384, max_prefill_tokens=16384, available_gpu_mem=54.84 GB Total latency: 127.126 s Score: 0.975 Output throughput: 631.254 token/s avg_spec_accept_length=3.1296052121306137 Ran 1 test in 1890.684s OK ``` The serialized ModelOpt Qwen3 test also completed two same-checkpoint disk reloads across weight/CUDA-graph release and resume while preserving deterministic decode text, token IDs, and logprobs: ```text # Serialized nvidia/Qwen3-30B-A3B-NVFP4 W4A16 disk reload Ran 1 test in 1205.499s OK ``` ### PR flashinfer-ai#4048 distributed benchmark compatibility This compatibility rerun executed at final PR head `c5ea45e4422e4c8be8059d79b27581971891132a`. `benchmarks/bench_cute_dsl_moe_distributed.py` is byte-identical between flashinfer-ai#4048's merged head and this head (`sha256:f61bd31757758bc661ddc94551d04c8a762a27612e0fad10c74c1c3240cc87fa`). The imported runtime copy is also byte-identical to the final head for both changed source files: - `flashinfer/fused_moe/cute_dsl/blackwell/moe_w4a16.py`: `sha256:cd11e0c1c37653b0a58403de04fa31ea20ae75a6ebab1b8330e7091e32d729f7` - `flashinfer/fused_moe/cute_dsl/tuner.py`: `sha256:c2dbbad8f2643f1fcd7f0677ec759d150910f6c44418f261cebd5dcba4c2fd39` Each of the three complete runs uses 10 warmups and 100 cold-L2, maximum-rank CUDA-event samples per cell. The HAI pod additionally needs `PET_LOCAL_ADDR=127.0.0.1` because its hostname is not resolvable in pod-local DNS. `--allreduce-backend auto` selected FlashInfer MNNVL, confirmed by a separate one-token `--verbose` diagnostic in the same environment. ```bash set -euo pipefail export FLASHINFER_SOURCE=/hai-workspace/flashinfer-w4a16-workspace export FLASHINFER_RUNTIME=/hai-workspace/flashinfer-runtime-w4a16-fix export RESULTS_DIR=/hai-workspace/sglang-cutedsl-w4a16-results/pr4048-distributed-c5ea-20260817b export PYTHONPATH="${FLASHINFER_RUNTIME}:${FLASHINFER_SOURCE}" export PET_LOCAL_ADDR=127.0.0.1 export FLASHINFER_DISABLE_VERSION_CHECK=1 export FLASHINFER_NVFP4_4OVER6=1 export FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 export FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE export FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 export FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 mkdir -p "${RESULTS_DIR}" for run in 1 2 3; do torchrun --standalone --nproc-per-node=8 \ "${FLASHINFER_SOURCE}/benchmarks/bench_cute_dsl_moe_distributed.py" \ --warmup 10 --iters 100 \ --use-per-token-activation --no-fused-finalize \ 2>&1 | tee "${RESULTS_DIR}/end-to-end-run-${run}.log" done ``` All three runs emitted the expected 52 rows: 13 token counts x EP8/TP8 x W4A4/W4A16. <details> <summary>Run 1: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.168512 DISTRIBUTED_CSV,ep,w4a16,1,8,0.944432 DISTRIBUTED_CSV,ep,w4a4,2,8,1.226576 DISTRIBUTED_CSV,ep,w4a16,2,8,0.957840 DISTRIBUTED_CSV,ep,w4a4,4,8,1.222032 DISTRIBUTED_CSV,ep,w4a16,4,8,0.916528 DISTRIBUTED_CSV,ep,w4a4,8,8,1.184704 DISTRIBUTED_CSV,ep,w4a16,8,8,0.963344 DISTRIBUTED_CSV,ep,w4a4,16,8,1.224800 DISTRIBUTED_CSV,ep,w4a16,16,8,0.951408 DISTRIBUTED_CSV,ep,w4a4,32,8,1.234592 DISTRIBUTED_CSV,ep,w4a16,32,8,0.967936 DISTRIBUTED_CSV,ep,w4a4,64,8,1.239088 DISTRIBUTED_CSV,ep,w4a16,64,8,0.958080 DISTRIBUTED_CSV,ep,w4a4,128,8,1.227296 DISTRIBUTED_CSV,ep,w4a16,128,8,0.943760 DISTRIBUTED_CSV,ep,w4a4,256,8,1.228256 DISTRIBUTED_CSV,ep,w4a16,256,8,0.955056 DISTRIBUTED_CSV,ep,w4a4,512,8,1.220816 DISTRIBUTED_CSV,ep,w4a16,512,8,0.949472 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.217200 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.950672 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.226176 DISTRIBUTED_CSV,ep,w4a16,2048,8,1.002928 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.270112 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.163104 DISTRIBUTED_CSV,tp,w4a4,1,8,1.246640 DISTRIBUTED_CSV,tp,w4a16,1,8,0.929936 DISTRIBUTED_CSV,tp,w4a4,2,8,1.193216 DISTRIBUTED_CSV,tp,w4a16,2,8,0.905024 DISTRIBUTED_CSV,tp,w4a4,4,8,1.191968 DISTRIBUTED_CSV,tp,w4a16,4,8,0.914336 DISTRIBUTED_CSV,tp,w4a4,8,8,1.232912 DISTRIBUTED_CSV,tp,w4a16,8,8,1.005168 DISTRIBUTED_CSV,tp,w4a4,16,8,1.275024 DISTRIBUTED_CSV,tp,w4a16,16,8,0.988672 DISTRIBUTED_CSV,tp,w4a4,32,8,1.302352 DISTRIBUTED_CSV,tp,w4a16,32,8,0.991168 DISTRIBUTED_CSV,tp,w4a4,64,8,1.287776 DISTRIBUTED_CSV,tp,w4a16,64,8,1.007712 DISTRIBUTED_CSV,tp,w4a4,128,8,1.288544 DISTRIBUTED_CSV,tp,w4a16,128,8,0.976032 DISTRIBUTED_CSV,tp,w4a4,256,8,1.294512 DISTRIBUTED_CSV,tp,w4a16,256,8,1.014432 DISTRIBUTED_CSV,tp,w4a4,512,8,1.325776 DISTRIBUTED_CSV,tp,w4a16,512,8,1.021552 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.333536 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.087632 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.386016 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.226896 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.534096 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.496576 ``` </details> <details> <summary>Run 2: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.198144 DISTRIBUTED_CSV,ep,w4a16,1,8,0.936384 DISTRIBUTED_CSV,ep,w4a4,2,8,1.208656 DISTRIBUTED_CSV,ep,w4a16,2,8,0.914576 DISTRIBUTED_CSV,ep,w4a4,4,8,1.227664 DISTRIBUTED_CSV,ep,w4a16,4,8,0.951744 DISTRIBUTED_CSV,ep,w4a4,8,8,1.208256 DISTRIBUTED_CSV,ep,w4a16,8,8,1.068048 DISTRIBUTED_CSV,ep,w4a4,16,8,1.231184 DISTRIBUTED_CSV,ep,w4a16,16,8,0.946000 DISTRIBUTED_CSV,ep,w4a4,32,8,1.234944 DISTRIBUTED_CSV,ep,w4a16,32,8,0.953472 DISTRIBUTED_CSV,ep,w4a4,64,8,1.382608 DISTRIBUTED_CSV,ep,w4a16,64,8,0.955488 DISTRIBUTED_CSV,ep,w4a4,128,8,1.217792 DISTRIBUTED_CSV,ep,w4a16,128,8,0.942240 DISTRIBUTED_CSV,ep,w4a4,256,8,1.223472 DISTRIBUTED_CSV,ep,w4a16,256,8,0.939152 DISTRIBUTED_CSV,ep,w4a4,512,8,1.224704 DISTRIBUTED_CSV,ep,w4a16,512,8,0.953024 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.231840 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.955184 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.230240 DISTRIBUTED_CSV,ep,w4a16,2048,8,0.984656 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.246976 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.152400 DISTRIBUTED_CSV,tp,w4a4,1,8,1.205248 DISTRIBUTED_CSV,tp,w4a16,1,8,0.952688 DISTRIBUTED_CSV,tp,w4a4,2,8,1.221104 DISTRIBUTED_CSV,tp,w4a16,2,8,0.927984 DISTRIBUTED_CSV,tp,w4a4,4,8,1.222752 DISTRIBUTED_CSV,tp,w4a16,4,8,0.930240 DISTRIBUTED_CSV,tp,w4a4,8,8,1.264928 DISTRIBUTED_CSV,tp,w4a16,8,8,1.002752 DISTRIBUTED_CSV,tp,w4a4,16,8,1.280784 DISTRIBUTED_CSV,tp,w4a16,16,8,0.994400 DISTRIBUTED_CSV,tp,w4a4,32,8,1.280864 DISTRIBUTED_CSV,tp,w4a16,32,8,1.002256 DISTRIBUTED_CSV,tp,w4a4,64,8,1.291056 DISTRIBUTED_CSV,tp,w4a16,64,8,1.005040 DISTRIBUTED_CSV,tp,w4a4,128,8,1.281040 DISTRIBUTED_CSV,tp,w4a16,128,8,1.001472 DISTRIBUTED_CSV,tp,w4a4,256,8,1.286368 DISTRIBUTED_CSV,tp,w4a16,256,8,0.997664 DISTRIBUTED_CSV,tp,w4a4,512,8,1.297824 DISTRIBUTED_CSV,tp,w4a16,512,8,1.012064 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.319184 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.068624 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.358832 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.234288 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.534160 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.490592 ``` </details> <details> <summary>Run 3: all 52 raw DISTRIBUTED_CSV rows</summary> ```text DISTRIBUTED_CSV,ep,w4a4,1,8,1.154336 DISTRIBUTED_CSV,ep,w4a16,1,8,0.908048 DISTRIBUTED_CSV,ep,w4a4,2,8,1.202480 DISTRIBUTED_CSV,ep,w4a16,2,8,0.942320 DISTRIBUTED_CSV,ep,w4a4,4,8,1.217984 DISTRIBUTED_CSV,ep,w4a16,4,8,0.925856 DISTRIBUTED_CSV,ep,w4a4,8,8,1.219968 DISTRIBUTED_CSV,ep,w4a16,8,8,0.925536 DISTRIBUTED_CSV,ep,w4a4,16,8,1.211344 DISTRIBUTED_CSV,ep,w4a16,16,8,0.944256 DISTRIBUTED_CSV,ep,w4a4,32,8,1.227168 DISTRIBUTED_CSV,ep,w4a16,32,8,0.924864 DISTRIBUTED_CSV,ep,w4a4,64,8,1.209744 DISTRIBUTED_CSV,ep,w4a16,64,8,0.942160 DISTRIBUTED_CSV,ep,w4a4,128,8,1.222128 DISTRIBUTED_CSV,ep,w4a16,128,8,0.950096 DISTRIBUTED_CSV,ep,w4a4,256,8,1.228112 DISTRIBUTED_CSV,ep,w4a16,256,8,0.955936 DISTRIBUTED_CSV,ep,w4a4,512,8,1.238400 DISTRIBUTED_CSV,ep,w4a16,512,8,0.946848 DISTRIBUTED_CSV,ep,w4a4,1024,8,1.222480 DISTRIBUTED_CSV,ep,w4a16,1024,8,0.950592 DISTRIBUTED_CSV,ep,w4a4,2048,8,1.247856 DISTRIBUTED_CSV,ep,w4a16,2048,8,0.995744 DISTRIBUTED_CSV,ep,w4a4,4096,8,1.271680 DISTRIBUTED_CSV,ep,w4a16,4096,8,1.173120 DISTRIBUTED_CSV,tp,w4a4,1,8,1.248944 DISTRIBUTED_CSV,tp,w4a16,1,8,0.931824 DISTRIBUTED_CSV,tp,w4a4,2,8,1.218480 DISTRIBUTED_CSV,tp,w4a16,2,8,0.910384 DISTRIBUTED_CSV,tp,w4a4,4,8,1.222864 DISTRIBUTED_CSV,tp,w4a16,4,8,0.919216 DISTRIBUTED_CSV,tp,w4a4,8,8,1.287648 DISTRIBUTED_CSV,tp,w4a16,8,8,0.991392 DISTRIBUTED_CSV,tp,w4a4,16,8,1.270656 DISTRIBUTED_CSV,tp,w4a16,16,8,0.994000 DISTRIBUTED_CSV,tp,w4a4,32,8,1.291312 DISTRIBUTED_CSV,tp,w4a16,32,8,0.996016 DISTRIBUTED_CSV,tp,w4a4,64,8,1.303648 DISTRIBUTED_CSV,tp,w4a16,64,8,0.994432 DISTRIBUTED_CSV,tp,w4a4,128,8,1.294864 DISTRIBUTED_CSV,tp,w4a16,128,8,1.004208 DISTRIBUTED_CSV,tp,w4a4,256,8,1.309296 DISTRIBUTED_CSV,tp,w4a16,256,8,1.022560 DISTRIBUTED_CSV,tp,w4a4,512,8,1.310064 DISTRIBUTED_CSV,tp,w4a16,512,8,1.026016 DISTRIBUTED_CSV,tp,w4a4,1024,8,1.332736 DISTRIBUTED_CSV,tp,w4a16,1024,8,1.082720 DISTRIBUTED_CSV,tp,w4a4,2048,8,1.383648 DISTRIBUTED_CSV,tp,w4a16,2048,8,1.229296 DISTRIBUTED_CSV,tp,w4a4,4096,8,1.546464 DISTRIBUTED_CSV,tp,w4a16,4096,8,1.511744 ``` </details> The table below takes the median across the three complete runs for each cell. `speedup = median(W4A4 EP8) / median(strategy)`, so values above 1 favor the strategy. | Tokens | W4A4 EP8 ms | W4A4 TP8 ms | TP speedup | W4A16 EP8 ms | EP speedup | W4A16 TP8 ms | TP speedup | |---:|---:|---:|---:|---:|---:|---:|---:| | 1 | 1.169 | 1.247 | 0.937x | 0.936 | 1.248x | 0.932 | 1.254x | | 2 | 1.209 | 1.218 | 0.992x | 0.942 | 1.283x | 0.910 | 1.328x | | 4 | 1.222 | 1.223 | 0.999x | 0.926 | 1.320x | 0.919 | 1.329x | | 8 | 1.208 | 1.265 | 0.955x | 0.963 | 1.254x | 1.003 | 1.205x | | 16 | 1.225 | 1.275 | 0.961x | 0.946 | 1.295x | 0.994 | 1.232x | | 32 | 1.235 | 1.291 | 0.956x | 0.953 | 1.295x | 0.996 | 1.240x | | 64 | 1.239 | 1.291 | 0.960x | 0.955 | 1.297x | 1.005 | 1.233x | | 128 | 1.222 | 1.289 | 0.948x | 0.944 | 1.295x | 1.001 | 1.220x | | 256 | 1.228 | 1.295 | 0.949x | 0.955 | 1.286x | 1.014 | 1.211x | | 512 | 1.225 | 1.310 | 0.935x | 0.949 | 1.290x | 1.022 | 1.199x | | 1024 | 1.222 | 1.333 | 0.917x | 0.951 | 1.286x | 1.083 | 1.129x | | 2048 | 1.230 | 1.384 | 0.889x | 0.996 | 1.235x | 1.229 | 1.001x | | 4096 | 1.270 | 1.534 | 0.828x | 1.163 | 1.092x | 1.497 | 0.849x | Geometric-mean speedups versus W4A4 EP8 are `0.940x` for W4A4 TP8, `1.266x` for W4A16 EP8, and `1.179x` for W4A16 TP8. Using the same exported environment above, the Nsight Systems driver launches its own eight-rank workers for all eight cases: two token counts x EP8/TP8 x W4A4/W4A16. ```bash python3 "${FLASHINFER_SOURCE}/benchmarks/bench_cute_dsl_moe_distributed.py" \ --mode profile_nsys --warmup 10 --profile-iters 20 \ --use-per-token-activation --no-fused-finalize \ --nsys-output-dir "${RESULTS_DIR}/nsys" \ 2>&1 | tee "${RESULTS_DIR}/nsys.log" ``` <details> <summary>32-token cases: four raw Nsight Systems stage tables</summary> ```text Nsight Systems kernel breakdown: EP8 W4A4, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.35% | 0.559 | 160 | 3.491 | deepseek_v3_topk_kernel dispatch | 1 | 0.15% | 0.240 | 160 | 1.502 | moeA2APrepareDispatchKernel dispatch | 2 | 19.77% | 31.380 | 160 | 196.127 | moeA2ADispatchKernel dispatch | 3 | 11.53% | 18.301 | 160 | 114.379 | moeA2ASanitizeExpertIdsKernel activation prep/quant | 1 | 0.11% | 0.169 | 160 | 1.059 | vectorized_elementwise_kernel activation prep/quant | 2 | 0.47% | 0.750 | 160 | 4.688 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.39% | 0.619 | 160 | 3.871 | routingIndicesClusterKernel local MoE | 2 | 7.29% | 11.573 | 160 | 72.334 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.12% | 0.197 | 160 | 1.230 | vectorized_elementwise_kernel local MoE | 4 | 0.81% | 1.282 | 160 | 8.013 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 3.28% | 5.204 | 160 | 32.524 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 0.88% | 1.400 | 160 | 8.750 | moeUnpermuteKernel combine | 1 | 0.27% | 0.429 | 160 | 2.682 | moeA2APrepareCombineKernel combine | 2 | 54.58% | 86.649 | 160 | 541.557 | moeA2ACombineKernel Nsight Systems kernel breakdown: EP8 W4A16, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.66% | 0.557 | 160 | 3.483 | deepseek_v3_topk_kernel dispatch | 1 | 0.29% | 0.240 | 160 | 1.498 | moeA2APrepareDispatchKernel dispatch | 2 | 21.84% | 18.310 | 160 | 114.439 | moeA2ADispatchKernel dispatch | 3 | 4.86% | 4.071 | 160 | 25.444 | moeA2ASanitizeExpertIdsKernel local MoE | 1 | 0.76% | 0.639 | 160 | 3.995 | routingIndicesClusterKernel local MoE | 2 | 0.78% | 0.656 | 160 | 4.099 | moePermuteKernel local MoE | 3 | 12.99% | 10.896 | 160 | 68.099 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 7.05% | 5.910 | 160 | 36.940 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 1.34% | 1.125 | 160 | 7.032 | moeUnpermuteKernel combine | 1 | 0.47% | 0.390 | 160 | 2.440 | moeA2APrepareCombineKernel combine | 2 | 48.96% | 41.055 | 160 | 256.594 | moeA2ACombineKernel Nsight Systems kernel breakdown: TP8 W4A4, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 13.30% | 35.050 | 160 | 219.062 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.21% | 0.566 | 160 | 3.535 | deepseek_v3_topk_kernel activation prep/quant | 1 | 0.06% | 0.153 | 160 | 0.954 | vectorized_elementwise_kernel activation prep/quant | 2 | 0.27% | 0.706 | 160 | 4.411 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.23% | 0.600 | 160 | 3.747 | routingIndicesClusterKernel local MoE | 2 | 4.46% | 11.747 | 160 | 73.420 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.08% | 0.201 | 160 | 1.257 | vectorized_elementwise_kernel local MoE | 4 | 2.18% | 5.746 | 160 | 35.910 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 2.35% | 6.188 | 160 | 38.673 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 1.09% | 2.865 | 160 | 17.906 | moeUnpermuteKernel all-reduce | 1 | 75.78% | 199.723 | 160 | 1248.270 | twoshotAllreduceKernel Nsight Systems kernel breakdown: TP8 W4A16, 32 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 18.23% | 17.083 | 160 | 106.768 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.61% | 0.568 | 160 | 3.547 | deepseek_v3_topk_kernel local MoE | 1 | 0.63% | 0.591 | 160 | 3.692 | routingIndicesClusterKernel local MoE | 2 | 0.87% | 0.815 | 160 | 5.093 | moePermuteKernel local MoE | 3 | 11.64% | 10.909 | 160 | 68.178 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 7.32% | 6.863 | 160 | 42.897 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 2.21% | 2.067 | 160 | 12.920 | moeUnpermuteKernel all-reduce | 1 | 58.50% | 54.823 | 160 | 342.642 | twoshotAllreduceKernel ``` </details> <details> <summary>4,096-token cases: four raw Nsight Systems stage tables</summary> ```text Nsight Systems kernel breakdown: EP8 W4A4, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.39% | 0.663 | 160 | 4.143 | deepseek_v3_topk_kernel dispatch | 1 | 0.15% | 0.247 | 160 | 1.546 | moeA2APrepareDispatchKernel dispatch | 2 | 18.26% | 30.738 | 160 | 192.110 | moeA2ADispatchKernel dispatch | 3 | 9.38% | 15.786 | 160 | 98.664 | moeA2ASanitizeExpertIdsKernel activation prep/quant | 1 | 0.15% | 0.259 | 160 | 1.620 | vectorized_elementwise_kernel activation prep/quant | 2 | 2.04% | 3.435 | 160 | 21.470 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.63% | 1.060 | 160 | 6.623 | routingIndicesClusterKernel local MoE | 2 | 10.34% | 17.407 | 160 | 108.791 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.23% | 0.385 | 160 | 2.405 | vectorized_elementwise_kernel local MoE | 4 | 5.91% | 9.940 | 160 | 62.123 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 5.83% | 9.808 | 160 | 61.299 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 3.10% | 5.219 | 160 | 32.616 | moeUnpermuteKernel combine | 1 | 0.93% | 1.571 | 160 | 9.817 | moeA2APrepareCombineKernel combine | 2 | 42.66% | 71.811 | 160 | 448.816 | moeA2ACombineKernel Nsight Systems kernel breakdown: EP8 W4A16, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel routing | 1 | 0.45% | 0.658 | 160 | 4.112 | deepseek_v3_topk_kernel dispatch | 1 | 0.16% | 0.240 | 160 | 1.500 | moeA2APrepareDispatchKernel dispatch | 2 | 13.16% | 19.373 | 160 | 121.083 | moeA2ADispatchKernel dispatch | 3 | 2.70% | 3.982 | 160 | 24.887 | moeA2ASanitizeExpertIdsKernel local MoE | 1 | 0.71% | 1.039 | 160 | 6.494 | routingIndicesClusterKernel local MoE | 2 | 2.20% | 3.241 | 160 | 20.257 | moePermuteKernel local MoE | 3 | 22.48% | 33.100 | 160 | 206.873 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 14.91% | 21.957 | 160 | 137.229 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 3.47% | 5.106 | 160 | 31.915 | moeUnpermuteKernel combine | 1 | 1.44% | 2.125 | 160 | 13.279 | moeA2APrepareCombineKernel combine | 2 | 38.32% | 56.418 | 160 | 352.613 | moeA2ACombineKernel Nsight Systems kernel breakdown: TP8 W4A4, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 20.71% | 49.212 | 160 | 307.574 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.68% | 1.627 | 160 | 10.170 | deepseek_v3_topk_kernel activation prep/quant | 1 | 0.10% | 0.245 | 160 | 1.534 | vectorized_elementwise_kernel activation prep/quant | 2 | 2.01% | 4.774 | 160 | 29.837 | kernel_cutlass_kernel_flashinf_1... local MoE | 1 | 0.53% | 1.258 | 160 | 7.864 | routingIndicesClusterKernel local MoE | 2 | 8.26% | 19.630 | 160 | 122.688 | kernel_cutlass_kernel_flashinf_2... local MoE | 3 | 0.09% | 0.215 | 160 | 1.345 | vectorized_elementwise_kernel local MoE | 4 | 4.90% | 11.640 | 160 | 72.748 | kernel_cutlass_kernel_flashinf_3... local MoE | 5 | 9.40% | 22.334 | 160 | 139.585 | kernel_cutlass_kernel_flashinf_4... local MoE | 6 | 7.93% | 18.841 | 160 | 117.757 | moeUnpermuteKernel all-reduce | 1 | 45.37% | 107.793 | 160 | 673.706 | twoshotAllreduceKernel Nsight Systems kernel breakdown: TP8 W4A16, 4096 global tokens stage | launch | GPU time | total (ms) | instances | avg (us) | kernel all-gather | 1 | 14.75% | 31.958 | 160 | 199.739 | ncclDevKernel_AllGather_RING_LL routing | 1 | 0.75% | 1.632 | 160 | 10.197 | deepseek_v3_topk_kernel local MoE | 1 | 0.57% | 1.243 | 160 | 7.769 | routingIndicesClusterKernel local MoE | 2 | 9.78% | 21.189 | 160 | 132.433 | moePermuteKernel local MoE | 3 | 16.70% | 36.189 | 160 | 226.178 | kernel_cutlass_kernel_flashinf_1... local MoE | 4 | 18.50% | 40.095 | 160 | 250.596 | kernel_cutlass_kernel_flashinf_1... local MoE | 5 | 8.52% | 18.467 | 160 | 115.418 | moeUnpermuteKernel all-reduce | 1 | 30.42% | 65.919 | 160 | 411.993 | twoshotAllreduceKernel ``` </details> The raw `total (ms)` columns above aggregate all eight ranks and 20 captured iterations. Dividing each case's sum by `8 * 20 = 160` gives: | Tokens | Mode | W4A4 staged ms | W4A16 staged ms | W4A4 / W4A16 | |---:|:---:|---:|---:|---:| | 32 | EP8 | 0.992 | 0.524 | 1.893x | | 32 | TP8 | 1.647 | 0.586 | 2.812x | | 4096 | EP8 | 1.052 | 0.920 | 1.143x | | 4096 | TP8 | 1.485 | 1.354 | 1.096x | Limitations: - This is a current-head compatibility run, not a base-vs-head performance A/B. It does not by itself prove that invocation-local allocation has zero latency overhead. - It ran on B300/SM103 with CUDA 13.0; flashinfer-ai#4048's published tables used B200 with different CUDA/runtime revisions, so absolute timings are not directly comparable. - Two end-to-end cells had isolated run ranges around 14% (EP8 W4A16 at 8 tokens and EP8 W4A4 at 64 tokens); use the reported three-run medians and avoid sub-percent conclusions. - Nsight staged sums are diagnostic kernel-time aggregates, not wall-clock latency, and do not subtract overlapping kernels. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes Please focus on the workspace lifetime and CUDA-graph ownership contract. This intentionally follows W4A4's invocation-local scratch pattern instead of adding a process-global or cross-wrapper workspace cache. The final diff contains no new test code; the validation-only behavioral tests and their passing results remain reviewable at `678867c3`. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved W4A16 mixture-of-experts operation reliability by allocating required working memory for each invocation. * Improved handling of workloads with varying sizes by removing reliance on previously cached capacity. * **Refactor** * Simplified workspace management while preserving existing operation behavior. * Streamlined the W4A16 execution interface for more predictable invocation behavior. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
5 tasks
This was referenced Aug 30, 2026
11 tasks
aleozlx
pushed a commit
that referenced
this pull request
Sep 10, 2026
<!-- .github/pull_request_template.md -->
## 📌 Description
@humansand
Compare CuTe DSL W4A16/W4A4 MoE against pure TRTLLM BF16 using the same
DeepSeek-V3 source inputs. Extend `benchmarks/bench_moe_deepseek.py`
with BF16 and NVFP4 precision modes in the existing `bench_trtllm`
helper.
- **BF16 contract:** BF16 activations and expert weights; public
`TrtllmBf16Config.prepare_weights` applies gated row reorder and
BlockMajorK layout outside timing, with no quantization.
- **Integration:** `TRTLLM NVFP4` / `TRTLLM BF16` output and
`trtllm-nvfp4` / `trtllm-bf16` selectors for benchmarking/profiling;
shared routing, autotuning and measurement setup with explicit precision
branches. Both precisions reuse the same cached routed realization per
input row.
- **Finalize:** `--no-fused-finalize` applies to CuTe. TRTLLM retains
native finalize; FP4 activation flags do not affect pure BF16.
- **Scope:** benchmark script only; no kernel or distributed-script
changes. Three configurations below: inference per-tensor, inference
per-token, and the deterministic RL configuration from #4048.
### Source and environment
- **Measured source:** `1ac2233d43df747af455ed4a8406d7439d686d9f`
(FlashInfer 0.6.18), used for all nine sweeps. **Implementation head:**
`2079fa00c4a2a0fb2e95bb284978a703653a4830`. The nine sweeps precede the
shared-helper/name/CLI-validation updates. Separate diagnostic
measurements below identify their own source commits; they do not
replace the original sweeps. Paired refactor correctness validation is
linked in Tests.
- **Hardware:** retained C2 devbox, one B300 SXM6 AC
(`CUDA_VISIBLE_DEVICES=0`, SM103, 148 SMs, 267.7 GiB) on an eight-B300
host; driver 590.48.01 and image CUDA compatibility library 595.58.03.
- **Image:** `nvcr.io/nvidia/pytorch:26.05-py3`; digest
`nvcr.io/nvidia/pytorch@sha256:222d8b18e671be5c3ef91cb41727a2572a0b23f59ded6c39f373a96946f6f2ba`.
- **Software:** CUDA toolkit 13.2 (`nvcc V13.2.78`), PyTorch
`2.12.0a0+5aff3928d8.nv26.05`, Python 3.12.3, CuTe DSL `[cu13]==4.7.0`,
`cupti-python==13.2.0`, `nvidia-cuda-cupti==13.2.86`,
`apache-tvm-ffi==0.1.13.post3`; `nvidia-cuda-nvdisasm==13.3.73` is the
DSL disassembler dependency, while compiler/runtime remain CUDA 13.2.
- **Provenance:** editable metadata still reports initial commit
`e713467f`; the measured Git checkout is `1ac2233d`. Their difference
only shares routed benchmark inputs across TRTLLM precisions. Warm
compilation caches are retained.
### Reproduction and configuration
Run inside the image above. Setup the exact measured source; the
optional EP-transport build is disabled because no communication is
performed.
```sh
set -euo pipefail
mkdir -p /workspace/flashinfer-moe-bf16-benchmark/artifacts
cd /workspace/flashinfer-moe-bf16-benchmark
git clone --branch zianglih/bench-deepseek-moe-bf16 https://github.com/zianglih/flashinfer.git flashinfer
cd flashinfer
git checkout --detach 1ac2233d43df747af455ed4a8406d7439d686d9f
git submodule update --init --recursive
export BUILD_NVEP=0 FLASHINFER_BUILD_NO_PIP=1 MAX_JOBS=8
python3 -m pip install 'setuptools>=77' 'apache-tvm-ffi>=0.1.6,!=0.1.8,!=0.1.8.post0,<0.2'
python3 -m pip install 'nvidia-cutlass-dsl[cu13]==4.7.0' 'cupti-python==13.2.0' 'nvidia-cuda-cupti==13.2.86' 'nvidia-cuda-nvdisasm==13.3.73' ninja einops pytest pytest-timeout
python3 -m pip install --no-build-isolation --no-deps -e .
python3 -m flashinfer.collect_env > ../artifacts/environment.txt
python3 -m pip freeze > ../artifacts/pip-freeze.txt
nvcc --version >> ../artifacts/environment.txt
nvidia-smi --query-gpu=name,uuid,driver_version,memory.total --format=csv >> ../artifacts/environment.txt
```
Run each command below from that checkout. All commands use logits
routing, wrapper API, TP1, initial activation quantization included,
CUDA graphs, CUPTI and autotuning. The existing order was per-token 1,
RL 1, per-token 2, RL 2, per-token 3, RL 3. Only the added per-tensor
variant was subsequently run three times consecutively. Every repetition
starts a fresh Python process.
**Inference per-tensor — `Per-token activation: False`:** omit
`--use-per-token-activation`; use the existing global-scale path, 4over6
disabled, FP4 quantization fast math enabled, CuTe fused atomic
finalize. The default CLI includes CUTLASS NVFP4 in this configuration.
```sh
env \
-u FLASHINFER_NVFP4_4OVER6 \
-u FLASHINFER_NVFP4_4OVER6_E4M3_USE_256 \
-u FLASHINFER_NVFP4_4OVER6_ERR_MODE \
-u FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH \
-u FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH \
CUDA_VISIBLE_DEVICES=0 FLASHINFER_DISABLE_VERSION_CHECK=1 MAX_JOBS=8 \
python3 benchmarks/bench_moe_deepseek.py \
--num-tokens 1,2,4,8,16,32,64,128,256,512,1024,2048,4096 \
--ep 8 --warmup 10 --iters 100 --routing-input-mode logits \
--include-activation-quant
```
**Inference per-token — `Per-token activation: True`:** per-token
activation scaling, 4over6 disabled, FP4 quantization fast math enabled,
CuTe fused atomic finalize. CUTLASS is omitted because it does not
consume per-token scales.
```sh
env \
-u FLASHINFER_NVFP4_4OVER6 \
-u FLASHINFER_NVFP4_4OVER6_E4M3_USE_256 \
-u FLASHINFER_NVFP4_4OVER6_ERR_MODE \
-u FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH \
-u FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH \
CUDA_VISIBLE_DEVICES=0 FLASHINFER_DISABLE_VERSION_CHECK=1 MAX_JOBS=8 \
python3 benchmarks/bench_moe_deepseek.py \
--num-tokens 1,2,4,8,16,32,64,128,256,512,1024,2048,4096 \
--ep 8 --warmup 10 --iters 100 --routing-input-mode logits \
--include-activation-quant --use-per-token-activation
```
**Deterministic RL — `Per-token activation: True`:** per-token 4over6
MSE, E4M3 max 256, error fast math enabled, FP4 quantization fast math
disabled, CuTe two-stage finalize. This is the requested configuration
name, not a full-model determinism claim; CUTLASS is omitted.
```sh
env \
-u FLASHINFER_NVFP4_4OVER6 \
-u FLASHINFER_NVFP4_4OVER6_E4M3_USE_256 \
-u FLASHINFER_NVFP4_4OVER6_ERR_MODE \
-u FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH \
-u FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH \
CUDA_VISIBLE_DEVICES=0 FLASHINFER_DISABLE_VERSION_CHECK=1 MAX_JOBS=8 \
FLASHINFER_NVFP4_4OVER6=1 \
FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 \
FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE \
FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 \
FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1 \
python3 benchmarks/bench_moe_deepseek.py \
--num-tokens 1,2,4,8,16,32,64,128,256,512,1024,2048,4096 \
--ep 8 --warmup 10 --iters 100 --routing-input-mode logits \
--include-activation-quant --use-per-token-activation --no-fused-finalize
```
- **Timing:** 10 warmups, 100 iterations/case, cold L2; each run reports
median iteration latency. Routing, expert computation, finalize and
initial FP4 activation quantization are timed. W4A16/BF16 consume BF16
directly. Weight preparation, compilation and autotuning are outside
measurement.
- **Workload:** H=7168, I=2048, E=256, 32 local experts (EP8
simulation), top-k=8, 8 groups/top-4 groups, routed scale=2.5, bias
scale=0.01, seed 42, tokens 1–4096 in powers of two. All arms share
source BF16 inputs and logits within each row. TFLOPS uses expected
uniform local work fraction 32/256; printed routing statistics show
actual local routes.
- **Capture:** an outer wrapper serializes returned `BenchResult`
records after measurement; original CLI/timing functions are unchanged.
All nine processes exited 0. All 507 captured latencies (6×13×4 +
3×13×5) were checked against the printed values before aggregation.
### Complete raw performance results
All nine complete result tables follow. Backend names and table spacing
are normalized for display; decorative divider lines are omitted.
Original logs and captures remain unchanged, as do every result value
and column. Identical printed configuration headers are shown once per
three-run group. Raw latency prints to 0.001 ms; raw speedups use
unrounded values. The existing six per-token/RL measurements remain
unchanged; only the three per-tensor runs were added.
**Inference per-tensor — printed configuration for all three runs:**
```text
DeepSeek-V3 MoE Performance Benchmark
GPU: NVIDIA B300 SXM6 AC
CuteDSL API: Wrapper
Per-token activation: False
Initial activation quantization: True
CuteDSL modes: W4A4 and W4A16; baselines: TRTLLM NVFP4 and TRTLLM BF16
Tensor parallelism simulation: TP=1
CUDA profiler capture: False
CuteDSL finalize: atomic fused
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
DeepSeek-V3 MoE Benchmark: CuteDSL W4A4/W4A16 vs CUTLASS vs TRTLLM NVFP4 / TRTLLM BF16 (EP=8, TP=1)
Model: hidden=7168, intermediate=2048, experts=256, top_k=8
EP Config: 32 local experts (simulating 8-way parallelism)
TP Config: intermediate size 2048 (simulating 1-way parallelism)
CUDA Graph: enabled, CUPTI: enabled
Routing bias scale: 0.01 (larger values tend to create expert imbalance)
Timed initial activation quantization for FP4-activation backends: included; W4A16 and TRTLLM BF16 consume BF16 directly
CuteDSL finalize: atomic fused
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
```
<details>
<summary>Inference per-tensor run 1 — 13 token counts, 5
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | CUTLASS | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs CUTLASS | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.031 2.8 | 0.040 2.2 | 0.048 1.8 | 0.026 3.4 | 0.051 1.7 | 1.53x 1.21x | 0.82x 0.65x | 1.63x 1.29x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.040 4.4 | 0.046 3.9 | 0.058 3.1 | 0.036 4.9 | 0.079 2.2 | 1.43x 1.26x | 0.89x 0.79x | 1.95x 1.72x | TRTLLM NVFP4 | 4 | 0/ 1/ 0.00
4 | 0.057 6.2 | 0.063 5.6 | 0.070 5.0 | 0.051 6.9 | 0.117 3.0 | 1.24x 1.12x | 0.89x 0.81x | 2.05x 1.85x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.073 9.7 | 0.079 8.9 | 0.086 8.2 | 0.064 11.0 | 0.157 4.5 | 1.18x 1.10x | 0.88x 0.81x | 2.15x 1.99x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.083 17.0 | 0.085 16.5 | 0.097 14.5 | 0.076 18.7 | 0.192 7.3 | 1.17x 1.14x | 0.91x 0.89x | 2.32x 2.25x | TRTLLM NVFP4 | 13 | 0/ 2/ 0.00
32 | 0.113 24.8 | 0.119 23.8 | 0.127 22.1 | 0.106 26.5 | 0.285 9.9 | 1.12x 1.07x | 0.94x 0.90x | 2.52x 2.41x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.146 38.5 | 0.153 36.9 | 0.159 35.6 | 0.139 40.6 | 0.388 14.5 | 1.08x 1.04x | 0.95x 0.91x | 2.65x 2.54x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.150 75.3 | 0.158 71.5 | 0.164 68.8 | 0.144 78.3 | 0.399 28.3 | 1.10x 1.04x | 0.96x 0.91x | 2.66x 2.53x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.158 142.8 | 0.164 137.2 | 0.173 130.0 | 0.154 146.7 | 0.427 52.8 | 1.10x 1.06x | 0.97x 0.94x | 2.71x 2.60x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.160 282.2 | 0.169 267.2 | 0.179 251.4 | 0.164 275.3 | 0.435 103.7 | 1.12x 1.06x | 1.03x 0.97x | 2.72x 2.58x | W4A4 | 32 | 7/ 23/ 14.00
1024 | 0.165 545.5 | 0.189 476.0 | 0.197 458.2 | 0.295 305.4 | 0.444 203.3 | 1.19x 1.04x | 1.79x 1.56x | 2.68x 2.34x | W4A4 | 32 | 17/ 42/ 28.50
2048 | 0.177 1018.3 | 0.233 773.7 | 0.223 810.3 | 0.299 602.6 | 0.475 379.4 | 1.26x 0.95x | 1.69x 1.28x | 2.68x 2.04x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.219 1646.7 | 0.363 993.0 | 0.281 1282.8 | 0.309 1167.4 | 0.529 682.5 | 1.28x 0.77x | 1.41x 0.85x | 2.41x 1.45x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Inference per-tensor run 2 — 13 token counts, 5
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | CUTLASS | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs CUTLASS | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.031 2.8 | 0.039 2.2 | 0.048 1.8 | 0.026 3.4 | 0.051 1.7 | 1.53x 1.21x | 0.82x 0.65x | 1.63x 1.29x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.040 4.4 | 0.046 3.8 | 0.058 3.1 | 0.036 4.9 | 0.079 2.2 | 1.43x 1.26x | 0.89x 0.79x | 1.95x 1.72x | TRTLLM NVFP4 | 4 | 0/ 1/ 0.00
4 | 0.057 6.2 | 0.063 5.6 | 0.070 5.0 | 0.051 6.9 | 0.117 3.0 | 1.24x 1.12x | 0.89x 0.81x | 2.05x 1.85x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.073 9.7 | 0.079 9.0 | 0.086 8.2 | 0.064 11.0 | 0.157 4.5 | 1.18x 1.10x | 0.88x 0.81x | 2.15x 1.99x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.083 17.0 | 0.088 16.0 | 0.097 14.5 | 0.076 18.6 | 0.192 7.3 | 1.17x 1.10x | 0.91x 0.86x | 2.32x 2.17x | TRTLLM NVFP4 | 13 | 0/ 2/ 0.00
32 | 0.114 24.8 | 0.118 23.8 | 0.127 22.1 | 0.106 26.6 | 0.285 9.9 | 1.12x 1.08x | 0.93x 0.90x | 2.51x 2.41x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.147 38.4 | 0.151 37.4 | 0.159 35.5 | 0.140 40.3 | 0.388 14.5 | 1.08x 1.05x | 0.95x 0.93x | 2.64x 2.57x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.150 75.3 | 0.158 71.3 | 0.164 68.8 | 0.144 78.2 | 0.399 28.2 | 1.09x 1.04x | 0.96x 0.91x | 2.67x 2.53x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.158 143.0 | 0.164 137.1 | 0.173 130.2 | 0.154 146.5 | 0.428 52.7 | 1.10x 1.05x | 0.98x 0.94x | 2.71x 2.60x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.160 281.9 | 0.169 267.3 | 0.179 251.3 | 0.164 275.4 | 0.435 103.8 | 1.12x 1.06x | 1.02x 0.97x | 2.72x 2.58x | W4A4 | 32 | 7/ 23/ 14.00
1024 | 0.165 546.4 | 0.190 475.7 | 0.197 458.4 | 0.295 305.4 | 0.443 203.6 | 1.19x 1.04x | 1.79x 1.56x | 2.68x 2.34x | W4A4 | 32 | 17/ 42/ 28.50
2048 | 0.177 1018.2 | 0.233 772.9 | 0.223 810.6 | 0.299 603.6 | 0.476 379.1 | 1.26x 0.95x | 1.69x 1.28x | 2.69x 2.04x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.217 1661.9 | 0.363 994.3 | 0.281 1283.9 | 0.309 1166.6 | 0.530 680.7 | 1.29x 0.77x | 1.42x 0.85x | 2.44x 1.46x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Inference per-tensor run 3 — 13 token counts, 5
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | CUTLASS | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs CUTLASS | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.031 2.8 | 0.040 2.2 | 0.048 1.8 | 0.026 3.4 | 0.051 1.7 | 1.53x 1.21x | 0.82x 0.65x | 1.64x 1.29x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.043 4.1 | 0.046 3.9 | 0.058 3.1 | 0.036 4.9 | 0.078 2.2 | 1.34x 1.26x | 0.83x 0.79x | 1.82x 1.71x | TRTLLM NVFP4 | 4 | 0/ 1/ 0.00
4 | 0.057 6.2 | 0.063 5.6 | 0.071 5.0 | 0.051 6.9 | 0.117 3.0 | 1.24x 1.12x | 0.90x 0.81x | 2.05x 1.85x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.073 9.7 | 0.079 9.0 | 0.086 8.2 | 0.064 11.0 | 0.156 4.5 | 1.18x 1.10x | 0.88x 0.81x | 2.15x 1.99x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.083 17.0 | 0.085 16.5 | 0.097 14.5 | 0.076 18.6 | 0.189 7.5 | 1.17x 1.14x | 0.91x 0.89x | 2.28x 2.22x | TRTLLM NVFP4 | 13 | 0/ 2/ 0.00
32 | 0.114 24.8 | 0.118 23.8 | 0.127 22.1 | 0.106 26.6 | 0.286 9.9 | 1.12x 1.08x | 0.93x 0.89x | 2.52x 2.41x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.147 38.5 | 0.151 37.3 | 0.158 35.6 | 0.140 40.3 | 0.387 14.6 | 1.08x 1.05x | 0.95x 0.93x | 2.64x 2.56x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.150 75.4 | 0.156 72.1 | 0.164 68.9 | 0.144 78.2 | 0.399 28.3 | 1.09x 1.05x | 0.96x 0.92x | 2.67x 2.55x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.158 142.7 | 0.165 136.7 | 0.173 130.0 | 0.154 146.9 | 0.428 52.7 | 1.10x 1.05x | 0.97x 0.93x | 2.70x 2.59x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.160 282.0 | 0.169 267.2 | 0.179 251.4 | 0.164 275.1 | 0.435 103.7 | 1.12x 1.06x | 1.03x 0.97x | 2.72x 2.58x | W4A4 | 32 | 7/ 23/ 14.00
1024 | 0.165 546.2 | 0.190 475.9 | 0.195 463.0 | 0.296 304.9 | 0.443 203.5 | 1.18x 1.03x | 1.79x 1.56x | 2.68x 2.34x | W4A4 | 32 | 17/ 42/ 28.50
2048 | 0.177 1019.1 | 0.233 774.0 | 0.223 810.6 | 0.299 603.0 | 0.477 378.4 | 1.26x 0.95x | 1.69x 1.28x | 2.69x 2.05x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.217 1661.4 | 0.363 993.5 | 0.280 1288.5 | 0.309 1168.4 | 0.529 681.8 | 1.29x 0.77x | 1.42x 0.85x | 2.44x 1.46x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
**Inference per-token — printed configuration for all three runs:**
```text
DeepSeek-V3 MoE Performance Benchmark
GPU: NVIDIA B300 SXM6 AC
CuteDSL API: Wrapper
Per-token activation: True
Initial activation quantization: True
CuteDSL modes: W4A4 and W4A16; baselines: TRTLLM NVFP4 and TRTLLM BF16
Tensor parallelism simulation: TP=1
CUDA profiler capture: False
CuteDSL finalize: atomic fused
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
DeepSeek-V3 MoE Benchmark: CuteDSL W4A4/W4A16 vs TRTLLM NVFP4 / TRTLLM BF16 (EP=8, TP=1)
Model: hidden=7168, intermediate=2048, experts=256, top_k=8
EP Config: 32 local experts (simulating 8-way parallelism)
TP Config: intermediate size 2048 (simulating 1-way parallelism)
CUDA Graph: enabled, CUPTI: enabled
Routing bias scale: 0.01 (larger values tend to create expert imbalance)
Timed initial activation quantization for FP4-activation backends: included; W4A16 and TRTLLM BF16 consume BF16 directly
CuteDSL finalize: atomic fused
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
CUTLASS omitted: it does not consume the per-token activation scale.
```
<details>
<summary>Inference per-token run 1 — 13 token counts, 4
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.042 2.1 | 0.040 2.2 | 0.034 2.6 | 0.053 1.7 | 0.82x 0.87x | 1.26x 1.34x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.053 3.3 | 0.046 3.9 | 0.046 3.8 | 0.079 2.2 | 0.87x 1.01x | 1.49x 1.72x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.072 4.9 | 0.062 5.7 | 0.060 5.9 | 0.116 3.0 | 0.83x 0.97x | 1.61x 1.87x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.088 8.0 | 0.078 9.0 | 0.074 9.5 | 0.157 4.5 | 0.84x 0.95x | 1.79x 2.01x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.097 14.6 | 0.084 16.8 | 0.086 16.3 | 0.190 7.4 | 0.89x 1.03x | 1.97x 2.27x | W4A16 | 13 | 0/ 2/ 0.00
32 | 0.132 21.4 | 0.120 23.5 | 0.116 24.3 | 0.287 9.8 | 0.88x 0.97x | 2.18x 2.39x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.165 34.2 | 0.153 36.8 | 0.148 38.0 | 0.386 14.6 | 0.90x 0.97x | 2.34x 2.52x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.170 66.5 | 0.157 71.8 | 0.154 73.2 | 0.398 28.3 | 0.91x 0.98x | 2.34x 2.53x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.178 126.8 | 0.164 137.3 | 0.164 137.9 | 0.427 52.8 | 0.92x 1.00x | 2.40x 2.60x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.182 247.3 | 0.169 267.6 | 0.174 259.3 | 0.436 103.5 | 0.95x 1.03x | 2.39x 2.58x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.192 470.5 | 0.190 475.4 | 0.316 285.6 | 0.443 203.8 | 1.65x 1.66x | 2.31x 2.33x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.213 848.8 | 0.233 774.5 | 0.326 553.7 | 0.477 378.2 | 1.53x 1.40x | 2.24x 2.05x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.268 1346.3 | 0.362 997.3 | 0.338 1066.3 | 0.529 681.8 | 1.26x 0.94x | 1.97x 1.46x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Inference per-token run 2 — 13 token counts, 4
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.048 1.8 | 0.039 2.2 | 0.035 2.5 | 0.051 1.7 | 0.73x 0.88x | 1.06x 1.28x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.053 3.3 | 0.046 3.9 | 0.046 3.8 | 0.079 2.2 | 0.87x 1.01x | 1.49x 1.73x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.072 4.9 | 0.062 5.7 | 0.060 5.9 | 0.117 3.0 | 0.83x 0.97x | 1.62x 1.90x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.088 8.0 | 0.078 9.0 | 0.074 9.5 | 0.156 4.5 | 0.84x 0.95x | 1.78x 2.00x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.097 14.6 | 0.087 16.3 | 0.087 16.3 | 0.190 7.4 | 0.90x 1.00x | 1.96x 2.19x | W4A16 | 13 | 0/ 2/ 0.00
32 | 0.132 21.4 | 0.120 23.5 | 0.116 24.3 | 0.287 9.8 | 0.88x 0.97x | 2.18x 2.39x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.165 34.2 | 0.153 36.8 | 0.149 37.9 | 0.387 14.6 | 0.90x 0.97x | 2.34x 2.53x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.170 66.4 | 0.157 71.7 | 0.154 73.1 | 0.398 28.3 | 0.91x 0.98x | 2.34x 2.53x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.178 126.7 | 0.164 137.2 | 0.164 137.8 | 0.427 52.8 | 0.92x 1.00x | 2.40x 2.60x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.183 247.1 | 0.169 267.4 | 0.174 259.2 | 0.436 103.5 | 0.95x 1.03x | 2.39x 2.58x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.191 471.0 | 0.190 475.1 | 0.317 284.7 | 0.442 203.9 | 1.65x 1.67x | 2.31x 2.33x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.213 848.6 | 0.233 774.4 | 0.326 553.2 | 0.477 378.4 | 1.53x 1.40x | 2.24x 2.05x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.268 1345.8 | 0.363 994.6 | 0.339 1065.5 | 0.528 683.5 | 1.26x 0.93x | 1.97x 1.46x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Inference per-token run 3 — 13 token counts, 4
backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.048 1.8 | 0.040 2.2 | 0.034 2.6 | 0.051 1.7 | 0.72x 0.87x | 1.06x 1.28x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.053 3.3 | 0.046 3.9 | 0.046 3.8 | 0.079 2.2 | 0.87x 1.01x | 1.49x 1.73x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.072 4.9 | 0.064 5.5 | 0.060 5.9 | 0.116 3.0 | 0.83x 0.94x | 1.61x 1.83x | TRTLLM NVFP4 | 7 | 0/ 1/ 0.00
8 | 0.088 8.0 | 0.078 9.0 | 0.074 9.5 | 0.156 4.5 | 0.84x 0.95x | 1.78x 2.00x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.097 14.6 | 0.087 16.3 | 0.086 16.4 | 0.190 7.4 | 0.89x 0.99x | 1.97x 2.20x | TRTLLM NVFP4 | 13 | 0/ 2/ 0.00
32 | 0.132 21.4 | 0.120 23.5 | 0.116 24.3 | 0.287 9.8 | 0.88x 0.97x | 2.18x 2.39x | TRTLLM NVFP4 | 21 | 0/ 3/ 1.00
64 | 0.165 34.1 | 0.153 36.8 | 0.149 37.8 | 0.387 14.6 | 0.90x 0.97x | 2.34x 2.53x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.170 66.4 | 0.157 71.8 | 0.154 73.2 | 0.400 28.2 | 0.91x 0.98x | 2.36x 2.55x | TRTLLM NVFP4 | 30 | 0/ 8/ 3.00
256 | 0.178 126.9 | 0.164 137.2 | 0.163 138.0 | 0.427 52.8 | 0.92x 0.99x | 2.40x 2.60x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.182 247.7 | 0.168 267.7 | 0.174 258.9 | 0.436 103.4 | 0.96x 1.03x | 2.39x 2.59x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.192 470.6 | 0.190 475.5 | 0.316 285.5 | 0.443 203.8 | 1.65x 1.67x | 2.31x 2.33x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.212 848.9 | 0.233 774.4 | 0.326 554.0 | 0.477 378.1 | 1.53x 1.40x | 2.25x 2.05x | W4A4 | 32 | 38/ 74/ 57.50
4096 | 0.268 1346.0 | 0.362 997.1 | 0.338 1066.0 | 0.529 682.3 | 1.26x 0.94x | 1.97x 1.46x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
**Deterministic RL — printed configuration for all three runs:**
```text
DeepSeek-V3 MoE Performance Benchmark
GPU: NVIDIA B300 SXM6 AC
CuteDSL API: Wrapper
Per-token activation: True
Initial activation quantization: True
CuteDSL modes: W4A4 and W4A16; baselines: TRTLLM NVFP4 and TRTLLM BF16
Tensor parallelism simulation: TP=1
CUDA profiler capture: False
CuteDSL finalize: deterministic two-stage
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
DeepSeek-V3 MoE Benchmark: CuteDSL W4A4/W4A16 vs TRTLLM NVFP4 / TRTLLM BF16 (EP=8, TP=1)
Model: hidden=7168, intermediate=2048, experts=256, top_k=8
EP Config: 32 local experts (simulating 8-way parallelism)
TP Config: intermediate size 2048 (simulating 1-way parallelism)
CUDA Graph: enabled, CUPTI: enabled
Routing bias scale: 0.01 (larger values tend to create expert imbalance)
Timed initial activation quantization for FP4-activation backends: included; W4A16 and TRTLLM BF16 consume BF16 directly
CuteDSL finalize: deterministic two-stage
TRTLLM NVFP4 / TRTLLM BF16 finalize: native (unaffected by --no-fused-finalize).
CUTLASS omitted: it does not consume the per-token activation scale.
```
<details>
<summary>Deterministic RL run 1 — 13 token counts, 4 backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.048 1.8 | 0.039 2.2 | 0.036 2.5 | 0.051 1.7 | 0.74x 0.91x | 1.05x 1.29x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.059 3.0 | 0.046 3.8 | 0.047 3.7 | 0.078 2.3 | 0.80x 1.03x | 1.32x 1.70x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.078 4.5 | 0.061 5.8 | 0.061 5.7 | 0.117 3.0 | 0.79x 1.01x | 1.51x 1.93x | W4A16 | 7 | 0/ 1/ 0.00
8 | 0.093 7.6 | 0.076 9.3 | 0.076 9.3 | 0.157 4.5 | 0.82x 1.00x | 1.69x 2.07x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.104 13.6 | 0.084 16.9 | 0.087 16.3 | 0.191 7.4 | 0.83x 1.04x | 1.83x 2.28x | W4A16 | 13 | 0/ 2/ 0.00
32 | 0.138 20.4 | 0.117 24.1 | 0.117 24.0 | 0.286 9.9 | 0.85x 1.01x | 2.07x 2.45x | W4A16 | 21 | 0/ 3/ 1.00
64 | 0.175 32.3 | 0.150 37.5 | 0.150 37.5 | 0.386 14.6 | 0.86x 1.00x | 2.21x 2.57x | W4A16 | 29 | 0/ 5/ 2.00
128 | 0.180 62.8 | 0.155 72.9 | 0.156 72.1 | 0.398 28.3 | 0.87x 1.01x | 2.21x 2.57x | W4A16 | 30 | 0/ 8/ 3.00
256 | 0.188 119.6 | 0.167 134.7 | 0.166 136.0 | 0.427 52.8 | 0.88x 0.99x | 2.26x 2.55x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.194 232.1 | 0.167 270.1 | 0.177 254.2 | 0.435 103.6 | 0.91x 1.06x | 2.24x 2.61x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.208 433.7 | 0.185 487.7 | 0.319 282.7 | 0.443 203.7 | 1.53x 1.73x | 2.13x 2.39x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.238 758.6 | 0.224 807.1 | 0.334 540.8 | 0.477 378.0 | 1.40x 1.49x | 2.01x 2.14x | W4A16 | 32 | 38/ 74/ 57.50
4096 | 0.311 1160.4 | 0.328 1098.6 | 0.350 1030.2 | 0.529 681.9 | 1.13x 1.07x | 1.70x 1.61x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Deterministic RL run 2 — 13 token counts, 4 backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.048 1.8 | 0.039 2.2 | 0.036 2.5 | 0.053 1.7 | 0.74x 0.91x | 1.10x 1.35x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.059 3.0 | 0.046 3.8 | 0.047 3.7 | 0.079 2.2 | 0.80x 1.03x | 1.34x 1.71x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.077 4.5 | 0.060 5.8 | 0.061 5.7 | 0.116 3.0 | 0.79x 1.02x | 1.50x 1.92x | W4A16 | 7 | 0/ 1/ 0.00
8 | 0.093 7.6 | 0.075 9.4 | 0.075 9.4 | 0.160 4.4 | 0.81x 1.00x | 1.72x 2.13x | TRTLLM NVFP4 | 10 | 0/ 2/ 0.00
16 | 0.104 13.5 | 0.083 16.9 | 0.087 16.3 | 0.190 7.4 | 0.83x 1.04x | 1.83x 2.28x | W4A16 | 13 | 0/ 2/ 0.00
32 | 0.138 20.4 | 0.115 24.5 | 0.117 24.1 | 0.287 9.8 | 0.85x 1.02x | 2.08x 2.49x | W4A16 | 21 | 0/ 3/ 1.00
64 | 0.175 32.3 | 0.150 37.5 | 0.150 37.6 | 0.386 14.6 | 0.86x 1.00x | 2.21x 2.57x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.180 62.8 | 0.153 73.5 | 0.156 72.1 | 0.400 28.2 | 0.87x 1.02x | 2.23x 2.61x | W4A16 | 30 | 0/ 8/ 3.00
256 | 0.189 119.6 | 0.167 134.7 | 0.165 136.3 | 0.427 52.8 | 0.88x 0.99x | 2.26x 2.55x | TRTLLM NVFP4 | 32 | 1/ 11/ 7.00
512 | 0.194 232.2 | 0.167 270.2 | 0.177 254.4 | 0.436 103.5 | 0.91x 1.06x | 2.24x 2.61x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.208 433.7 | 0.185 487.9 | 0.320 282.2 | 0.443 203.5 | 1.54x 1.73x | 2.13x 2.40x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.238 759.1 | 0.224 807.1 | 0.333 542.0 | 0.477 378.2 | 1.40x 1.49x | 2.01x 2.13x | W4A16 | 32 | 38/ 74/ 57.50
4096 | 0.311 1160.6 | 0.328 1098.4 | 0.350 1029.4 | 0.528 682.8 | 1.13x 1.07x | 1.70x 1.61x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
<details>
<summary>Deterministic RL run 3 — 13 token counts, 4 backends</summary>
```text
Tokens | CuteDSL W4A4 | CuteDSL W4A16 | TRTLLM NVFP4 | TRTLLM BF16 | Speedup vs TRTLLM NVFP4 | Speedup vs TRTLLM BF16 | Winner | Active | Stats
| ms TFLOPS | ms TFLOPS | ms TFLOPS | ms TFLOPS | W4A4 W4A16 | W4A4 W4A16 | | experts | min/max/median
1 | 0.048 1.8 | 0.040 2.2 | 0.036 2.4 | 0.053 1.7 | 0.75x 0.90x | 1.09x 1.31x | TRTLLM NVFP4 | 2 | 0/ 1/ 0.00
2 | 0.059 3.0 | 0.046 3.8 | 0.047 3.7 | 0.079 2.2 | 0.81x 1.03x | 1.34x 1.71x | W4A16 | 4 | 0/ 1/ 0.00
4 | 0.077 4.6 | 0.060 5.8 | 0.061 5.7 | 0.117 3.0 | 0.79x 1.01x | 1.51x 1.93x | W4A16 | 7 | 0/ 1/ 0.00
8 | 0.093 7.6 | 0.075 9.4 | 0.075 9.4 | 0.157 4.5 | 0.81x 1.00x | 1.68x 2.09x | W4A16 | 10 | 0/ 2/ 0.00
16 | 0.104 13.6 | 0.083 16.9 | 0.087 16.3 | 0.190 7.4 | 0.84x 1.04x | 1.83x 2.28x | W4A16 | 13 | 0/ 2/ 0.00
32 | 0.139 20.3 | 0.117 24.1 | 0.117 24.1 | 0.287 9.8 | 0.84x 1.00x | 2.07x 2.46x | W4A16 | 21 | 0/ 3/ 1.00
64 | 0.175 32.3 | 0.152 37.0 | 0.150 37.5 | 0.386 14.6 | 0.86x 0.99x | 2.21x 2.54x | TRTLLM NVFP4 | 29 | 0/ 5/ 2.00
128 | 0.180 62.8 | 0.155 72.8 | 0.156 72.2 | 0.398 28.4 | 0.87x 1.01x | 2.21x 2.57x | W4A16 | 30 | 0/ 8/ 3.00
256 | 0.189 119.5 | 0.162 139.1 | 0.166 136.0 | 0.427 52.8 | 0.88x 1.02x | 2.26x 2.63x | W4A16 | 32 | 1/ 11/ 7.00
512 | 0.194 232.3 | 0.167 270.1 | 0.177 254.3 | 0.436 103.3 | 0.91x 1.06x | 2.25x 2.61x | W4A16 | 32 | 7/ 23/ 14.00
1024 | 0.208 433.5 | 0.185 487.5 | 0.322 280.3 | 0.443 203.4 | 1.55x 1.74x | 2.13x 2.40x | W4A16 | 32 | 17/ 42/ 28.50
2048 | 0.238 759.5 | 0.223 807.2 | 0.333 542.3 | 0.477 378.1 | 1.40x 1.49x | 2.01x 2.14x | W4A16 | 32 | 38/ 74/ 57.50
4096 | 0.311 1159.4 | 0.329 1097.9 | 0.351 1029.3 | 0.528 683.4 | 1.13x 1.07x | 1.70x 1.61x | W4A4 | 32 | 81/147/ 117.00
Speedup > 1.0 means that CuTe DSL mode is faster than the comparison backend
```
</details>
### Derived performance summary
Latencies are **medians of three fresh-process runs**, calculated from
unrounded captures and displayed in milliseconds to three decimal
places. Each **W4A16 speedup** is the named TRTLLM baseline's median
latency divided by the CuTe DSL W4A16 median latency at that token
count. These are ratios of medians; values above 1 mean W4A16 is faster;
values below 1 mean W4A16 is slower. Both NVFP4 and pure BF16
comparisons are explicit.
**Inference per-tensor — median of three runs**
| Tokens | CuTe W4A4 (ms) | CuTe W4A16 (ms) | CUTLASS NVFP4 (ms) |
TRTLLM NVFP4 (ms) | TRTLLM BF16 (ms) | W4A16 speedup over TRTLLM NVFP4 |
W4A16 speedup over TRTLLM BF16 |
| ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: |
| 1 | 0.031 | 0.040 | 0.048 | 0.026 | 0.051 | 0.65x | 1.29x |
| 2 | 0.040 | 0.046 | 0.058 | 0.036 | 0.079 | 0.79x | 1.72x |
| 4 | 0.057 | 0.063 | 0.070 | 0.051 | 0.117 | 0.81x | 1.85x |
| 8 | 0.073 | 0.079 | 0.086 | 0.064 | 0.157 | 0.81x | 1.99x |
| 16 | 0.083 | 0.085 | 0.097 | 0.076 | 0.192 | 0.89x | 2.25x |
| 32 | 0.114 | 0.118 | 0.127 | 0.106 | 0.285 | 0.89x | 2.41x |
| 64 | 0.147 | 0.151 | 0.159 | 0.140 | 0.388 | 0.92x | 2.56x |
| 128 | 0.150 | 0.158 | 0.164 | 0.144 | 0.399 | 0.91x | 2.53x |
| 256 | 0.158 | 0.164 | 0.173 | 0.154 | 0.428 | 0.93x | 2.60x |
| 512 | 0.160 | 0.169 | 0.179 | 0.164 | 0.435 | 0.97x | 2.58x |
| 1024 | 0.165 | 0.190 | 0.197 | 0.295 | 0.443 | 1.56x | 2.34x |
| 2048 | 0.177 | 0.233 | 0.223 | 0.299 | 0.476 | 1.28x | 2.04x |
| 4096 | 0.217 | 0.363 | 0.281 | 0.309 | 0.529 | 0.85x | 1.46x |
**Inference per-token — median of three runs**
| Tokens | CuTe W4A4 (ms) | CuTe W4A16 (ms) | TRTLLM NVFP4 (ms) | TRTLLM
BF16 (ms) | W4A16 speedup over TRTLLM NVFP4 | W4A16 speedup over TRTLLM
BF16 |
| ---: | ---: | ---: | ---: | ---: | ---: | ---: |
| 1 | 0.048 | 0.040 | 0.034 | 0.051 | 0.87x | 1.28x |
| 2 | 0.053 | 0.046 | 0.046 | 0.079 | 1.01x | 1.73x |
| 4 | 0.072 | 0.062 | 0.060 | 0.116 | 0.97x | 1.88x |
| 8 | 0.088 | 0.078 | 0.074 | 0.156 | 0.95x | 2.00x |
| 16 | 0.097 | 0.087 | 0.086 | 0.190 | 1.00x | 2.20x |
| 32 | 0.132 | 0.120 | 0.116 | 0.287 | 0.97x | 2.39x |
| 64 | 0.165 | 0.153 | 0.149 | 0.387 | 0.97x | 2.52x |
| 128 | 0.170 | 0.157 | 0.154 | 0.398 | 0.98x | 2.53x |
| 256 | 0.178 | 0.164 | 0.164 | 0.427 | 1.00x | 2.60x |
| 512 | 0.182 | 0.169 | 0.174 | 0.436 | 1.03x | 2.58x |
| 1024 | 0.192 | 0.190 | 0.316 | 0.443 | 1.66x | 2.33x |
| 2048 | 0.213 | 0.233 | 0.326 | 0.477 | 1.40x | 2.05x |
| 4096 | 0.268 | 0.362 | 0.338 | 0.529 | 0.94x | 1.46x |
**Deterministic RL — median of three runs**
| Tokens | CuTe W4A4 (ms) | CuTe W4A16 (ms) | TRTLLM NVFP4 (ms) | TRTLLM
BF16 (ms) | W4A16 speedup over TRTLLM NVFP4 | W4A16 speedup over TRTLLM
BF16 |
| ---: | ---: | ---: | ---: | ---: | ---: | ---: |
| 1 | 0.048 | 0.039 | 0.036 | 0.053 | 0.91x | 1.34x |
| 2 | 0.059 | 0.046 | 0.047 | 0.079 | 1.03x | 1.71x |
| 4 | 0.077 | 0.060 | 0.061 | 0.117 | 1.01x | 1.93x |
| 8 | 0.093 | 0.075 | 0.075 | 0.157 | 1.00x | 2.09x |
| 16 | 0.104 | 0.083 | 0.087 | 0.190 | 1.04x | 2.28x |
| 32 | 0.138 | 0.117 | 0.117 | 0.287 | 1.00x | 2.46x |
| 64 | 0.175 | 0.150 | 0.150 | 0.386 | 1.00x | 2.57x |
| 128 | 0.180 | 0.155 | 0.156 | 0.398 | 1.01x | 2.57x |
| 256 | 0.189 | 0.167 | 0.166 | 0.427 | 0.99x | 2.55x |
| 512 | 0.194 | 0.167 | 0.177 | 0.436 | 1.06x | 2.61x |
| 1024 | 0.208 | 0.185 | 0.320 | 0.443 | 1.73x | 2.40x |
| 2048 | 0.238 | 0.224 | 0.333 | 0.477 | 1.49x | 2.13x |
| 4096 | 0.311 | 0.328 | 0.350 | 0.528 | 1.07x | 1.61x |
- **Inference per-tensor:** W4A16 speedup 0.65–1.56x over TRTLLM NVFP4
(2/13 faster); 1.29–2.60x over TRTLLM BF16 (13/13 faster).
- **Inference per-token:** W4A16 speedup 0.87–1.66x over TRTLLM NVFP4
(4/13 faster); 1.28–2.60x over TRTLLM BF16 (13/13 faster).
- **Deterministic RL:** W4A16 speedup 0.91–1.73x over TRTLLM NVFP4 (9/13
faster); 1.34–2.61x over TRTLLM BF16 (13/13 faster).
### BF16 performance investigation
**Finding:** the W4A16 advantage is present in GPU expert computation.
TRTLLM BF16 tuning is active; exhaustive native-tactic checks on the
actual inputs do not close the gap. Preparation, tuning, L2 flushing and
host dispatch are outside the reported CUPTI GPU span. Independent
single-call graph events reproduce the ordering. These diagnostics
supplement the unchanged nine sweeps.
**Provenance:** profiles, actual-input tactic sweep and event controls
use `2079fa00c4a2a0fb2e95bb284978a703653a4830` on the same retained
B300/CUDA 13.2 image above. Nsight Systems 2026.2.1, GPU 0; tactic
diagnostics GPU 1, serialized with GPU 0 work. Initial cache/default
audit uses `a3fce87d5beec29775e9bc3917df99ba6bc4db9c`. Naming commit
`d0caa842` and profiling-only commit `2079fa00` are separate; normal
benchmark functions/timing remain unchanged from the refactor validation
source.
**Autotuning:** 2,852 successful native candidate profiles across 21
token buckets, zero failures; all ten measured warmup/capture dispatches
used nonfallback cache hits with tuning disabled. The selected tactic
reaches the native launcher. At 32 tokens, all 352 actual-input
candidates give a best 284.5635 µs versus 287.3155 µs for the selected
baseline (about 1%, single sweep). At 4096, the selected `(128,0)` is
also the exhaustive eight-candidate winner at 529.3520 µs. Default `-1`
is ~15% slower there. The v1 tuner's synthetic routing activates 11/30
local experts at 32/4096, versus 21/32 on actual inputs; shorter
synthetic probe times therefore are not the final benchmark latency. The
exhaustive actual-input sweep checks this selection concern directly.
TRTLLM BF16 tuned/default controls (µs; each cell is a 100-iteration
median; 3 alternating pairs):
```csv
tokens,tuned_tactic,repeat,tuned_us,default_us
32,8:92,1,286.9315,287.1400
32,8:92,2,287.1560,287.0915
32,8:92,3,287.0440,287.0440
4096,128:0,1,528.9665,609.1280
4096,128:0,2,530.0870,610.4075
4096,128:0,3,531.5110,611.6720
```
All actual-input native candidates (median of 100 iterations, µs; one
sample each). Each CSV row lists latencies in ascending contiguous
config-ID order; `(tileN, config)` is the native tactic. This includes
every 352/8 enumerated candidate plus default.
```text
tokens=32; default(-1)=286.8350
tile=8; config=0..143
295.8595,286.0840,297.8440,285.7475,299.2360,285.5240,295.9080,297.4760,286.0515,299.2035,286.1795,285.4120,297.0920,287.0435,298.6920,286.5315,300.6760,286.2120,296.9640,298.5000,286.8360,300.1955,286.5155,286.3720,295.2840,285.3795,296.8685,285.3000,299.1075,284.6115,295.2520,296.9475,285.4115,299.1080,284.8520,284.5635,296.2920,286.4840,298.1160,286.3075,300.0350,285.9240,296.1640,298.5960,286.8840,300.0355,286.2435,285.6995,297.1560,286.8360,298.8200,286.7560,300.6760,286.3080,296.5795,298.5475,286.8035,300.8355,286.7235,286.1800,296.7555,286.9640,298.3240,287.0435,300.3080,286.3885,296.9475,297.9720,287.0115,300.1160,286.6915,286.4990,296.2600,285.9395,297.4600,285.8280,299.6360,285.6360,295.9720,297.3800,286.1155,299.5075,285.6360,285.3955,297.2680,287.2360,298.4360,286.8200,300.6600,286.3555,297.1070,298.7720,286.9630,301.0115,286.7235,286.4840,295.3160,285.2355,296.7560,285.2200,299.2040,284.9795,295.3315,296.9640,285.4600,299.2995,285.0440,284.5800,296.3880,286.6440,297.8755,286.3395,299.7320,285.7800,296.1955,297.9715,287.1240,300.1160,286.1960,285.9715,297.1395,287.0595,299.0280,286.9150,300.9635,286.3720,297.1560,298.4520,286.9155,300.9960,286.6755,286.5800,296.9000,287.0920,298.2125,287.1080,300.1795,286.5635,296.5800,298.5800,287.0280,300.3085,286.8520,286.2920
tile=16; config=0..143
292.6915,294.5480,285.8280,296.8680,285.6355,285.1880,292.8355,285.6040,294.7075,296.8840,285.6200,285.3160,293.3160,295.2365,286.3875,297.7965,286.3235,286.0680,293.0440,286.6600,295.0760,297.4760,286.4200,286.2920,291.9400,293.9885,285.0910,296.1645,285.0600,284.7080,292.0835,284.9640,293.8280,296.4360,285.1070,284.7230,292.6760,294.9800,285.9240,297.1880,286.2595,285.3320,292.7240,285.9880,294.7080,296.8680,286.0040,285.5075,293.7630,295.8605,287.0440,298.0835,287.0920,286.5960,293.5720,286.7880,295.6520,297.4760,287.0120,286.1800,293.2840,295.4600,286.8670,297.9560,286.6115,286.2115,293.0435,286.3720,295.3805,297.2840,286.9325,286.0675,292.2435,294.4520,285.6040,296.4360,285.8760,285.1560,292.7245,285.5240,294.7560,296.9000,285.7320,285.3005,293.2845,295.3000,286.3245,297.5245,286.4360,286.5000,293.1080,286.4680,295.2520,297.5875,286.7240,286.1480,291.8760,294.1960,285.1080,296.0360,284.9640,284.8200,291.6040,285.1235,294.0360,296.2920,285.3160,284.7555,292.6440,294.8200,285.8920,296.9640,286.0520,285.1880,292.7560,285.9550,294.5955,297.2520,286.0840,285.7315,293.7000,295.8920,286.8520,298.2760,286.9000,286.9320,293.7320,286.8520,295.3640,297.7485,287.1720,286.6755,293.3965,295.4440,286.3725,297.3800,286.9160,286.1795,293.5725,286.5480,295.1400,297.5710,286.5950,286.2915
tile=32; config=0..63
292.4360,286.0520,295.0600,285.7955,292.2285,295.1080,285.5555,286.2440,294.3400,287.7800,297.5880,287.7320,294.1960,297.2200,287.8120,287.8920,293.0760,286.9160,296.5960,286.7720,292.8675,296.4680,286.8035,286.7240,293.6680,287.7165,297.1715,287.5720,293.5565,296.9160,287.6040,287.4920,292.1795,286.2110,295.2680,285.7475,292.1635,295.4600,286.0200,286.0675,294.1480,288.0195,297.5405,287.7955,294.1480,297.2200,287.6845,288.0680,293.2840,286.9160,296.1320,286.7075,293.0920,296.3240,286.9320,286.9800,293.8600,287.7960,296.7075,287.4290,293.5720,296.9160,287.3000,287.6835
tokens=4096; default(-1)=609.7365
tile=64; config=0..3
611.4960,611.9765,610.2970,611.5760
tile=128; config=0..3
529.3520,529.5590,530.8710,530.3430
```
**Nsight Systems complete device breakdown (µs):** 20 cold-L2, one-call
graph replays per case after tuning and three uncaptured replays. Every
non-dash activity cell is the mean of 20 instances; both temporary-zero
kernels are shown separately. All target kernels/memsets are accounted
for inside synchronized `moe` NVTX ranges; each contains one
`cudaGraphLaunch`. The 20 separately labeled L2 flushes are excluded.
Kernel sums are diagnostic, not latency: [programmatic dependent
launch](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html)
allows dependent kernel intervals to overlap. GPU span measures the
first activity start through last activity end; uncovered gaps use the
union of intervals. Inference profiles cover both inference activation
variants because those flags do not change W4A16 or TRTLLM BF16; RL is
separately profiled.
**Inference:**
| Activity / metric | CuTe W4A16 / 32 | TRTLLM BF16 / 32 | CuTe W4A16 /
4096 | TRTLLM BF16 / 4096 |
| --- | ---: | ---: | ---: | ---: |
| Top-k routing | 3.150 | 3.346 | 9.635 | 9.214 |
| Temporary zero 1 | 0.587 | — | 0.672 | — |
| Temporary zero 2 | 0.628 | — | 0.702 | — |
| Route map | 3.709 | 3.171 | 6.832 | 6.339 |
| BF16 permute | 4.648 | — | 22.986 | — |
| GEMM1 + SwiGLU | 67.812 | 189.739 | 164.956 | 331.377 |
| Output zero | 0.906 | — | 9.722 | — |
| GEMM2 | 36.946 | 89.677 | 147.033 | 150.361 |
| Finalize | — | 5.779 | — | 38.018 |
| GPU activity sum | 118.386 | 291.712 | 362.538 | 535.309 |
| GPU span | 118.522 | 285.562 | 359.142 | 527.871 |
| Uncovered GPU gaps | 0.668 | 0.000 | 0.644 | 0.000 |
| Graph launch CPU API (excluded) | 9.867 | 8.062 | 7.405 | 8.048 |
| L2 flush (excluded) | 67.826 | 67.925 | 67.882 | 67.810 |
**Deterministic RL:**
| Activity / metric | CuTe W4A16 / 32 | TRTLLM BF16 / 32 | CuTe W4A16 /
4096 | TRTLLM BF16 / 4096 |
| --- | ---: | ---: | ---: | ---: |
| Top-k routing | 3.094 | 3.283 | 9.622 | 9.283 |
| Temporary zero 1 | 0.592 | — | 0.677 | — |
| Temporary zero 2 | 0.590 | — | 0.694 | — |
| Route map | 3.578 | 3.014 | 7.237 | 6.586 |
| BF16 permute | 4.526 | — | 23.088 | — |
| GEMM1 + SwiGLU | 68.505 | 188.811 | 165.900 | 331.810 |
| Output zero | — | — | — | — |
| GEMM2 | 43.343 | 89.194 | 110.439 | 150.190 |
| Finalize | 5.880 | 5.755 | 31.789 | 38.095 |
| GPU activity sum | 130.108 | 290.058 | 349.446 | 535.964 |
| GPU span | 114.262 | 284.682 | 326.137 | 528.484 |
| Uncovered GPU gaps | 0.365 | 0.000 | 0.329 | 0.000 |
| Graph launch CPU API (excluded) | 7.163 | 6.141 | 10.132 | 8.223 |
| L2 flush (excluded) | 67.928 | 67.817 | 67.880 | 67.842 |
Kernel mapping: top-k=`deepseek_v3_topk_kernel` / `routingMainKernel`;
route map=`routingIndicesClusterKernel` / `routingIndicesCoopKernel`;
temporary zeros=`vectorized_elementwise_kernel`;
permute/unpermute=`moePermuteKernel` / `moeUnpermuteKernel`; CuTe
GEMMs=`Sm100W4A16GroupedGemmKernel` in launch order; TRTLLM BF16
GEMMs=`bmm_Bfloat16_Bfloat16Bfloat16...` (FC1 has `swiGlu`); native
finalize=`finalizeKernel` / `finalizeKernelVecLoad`. CuTe inference
GEMM2 includes atomic finalize; RL uses the separate unpermute row. At
32 tokens both GEMMs explain the gap; at 4096 the largest difference is
GEMM1, with native finalize also contributing. TRTLLM BF16 has no
uncovered device gaps in these traces.
**Independent timer control:** same tuned callable, one call per graph,
100 samples after 10 warmups; zero a 2×L2 byte buffer and synchronize
before each start event, then `start.record(); graph.replay();
end.record(); end.synchronize()`. CUPTI separately measures the same
prepared callable. Every case passed with finite outputs. Table contains
every case's median in µs. Event intervals add roughly 4–6 µs but retain
the gap; host submission cannot explain the BF16 slowdown. Warm-cache
results are sequential diagnostics, not an isolated cache speedup claim.
Ordinary `--no-cupti` is not this control: it uses ten-call graphs and
rotates only input kwargs, leaving closure-held weights unrotated.
| Config | Tokens | Backend | Cold CUPTI | Cold events | Warm events |
| --- | ---: | --- | ---: | ---: | ---: |
| inference | 1 | CuTe W4A16 | 38.9440 | 42.9440 | 39.8240 |
| inference | 1 | TRTLLM BF16 | 50.7525 | 56.8640 | 46.5280 |
| inference | 32 | CuTe W4A16 | 118.9450 | 124.4480 | 113.4080 |
| inference | 32 | TRTLLM BF16 | 286.0665 | 292.3360 | 276.1920 |
| inference | 4096 | CuTe W4A16 | 361.9550 | 366.7200 | 378.1440 |
| inference | 4096 | TRTLLM BF16 | 528.9640 | 533.6480 | 543.2800 |
| rl | 1 | CuTe W4A16 | 40.1280 | 45.8080 | 42.1760 |
| rl | 1 | TRTLLM BF16 | 53.0885 | 58.8960 | 48.0480 |
| rl | 32 | CuTe W4A16 | 115.0410 | 120.3360 | 112.3520 |
| rl | 32 | TRTLLM BF16 | 286.5785 | 292.5920 | 276.2240 |
| rl | 4096 | CuTe W4A16 | 327.4590 | 333.0560 | 355.2640 |
| rl | 4096 | TRTLLM BF16 | 528.6275 | 533.4560 | 542.9120 |
### Nsight Compute: fixed normal-tuned tactics
- **Scope:** inference per-token, EP8/TP1, B300 GPU 0, NCU 2026.1.1,
source `2079fa00`. Tactics were recorded during normal tuning before NCU
injection, then fixed by a diagnostic helper; all captured calls had
tuning disabled. Each report contains exactly FC1 (with SwiGLU) and FC2,
nine replay passes each. Node replay flushes caches and isolates PDL
kernels: these durations are diagnostics, not additive end-to-end
latency.
- **Tactics:** TRTLLM BF16: `(8, 37)` at 32; `(128, 1)` at 4096. CuteDSL
W4A16 uses the same tactic for both GEMMs: `((128, 8, 256), (2, 1),
True)` at 32; `((256, 128, 256), (2, 1), True)` at 4096.
Raw counters below: R/W = DRAM bytes; D/T/S = DRAM throughput /
tensor-pipe activity / SM throughput (% of sustained elapsed peak);
LD/ST = local-memory sectors; warps = active warps per active SM cycle;
A/T = achieved/theoretical occupancy (%).
| Tokens | Backend | GEMM | R bytes | W bytes | ns | D/T/S % | LD/ST |
Warps | Occupancy A/T % | Registers/thread | Shared bytes/block |
| ---: | --- | --- | ---: | ---: | ---: | --- | --- | ---: | --- | ---:
| ---: |
| 32 | TRTLLM BF16 | FC1 | 1233656832 | 3507200 | 176896 |
91.17/2.42/25.37 | 0/0 | 11.33 | 17.70/18.75 | 168 | 181904 |
| 32 | TRTLLM BF16 | FC2 | 616784896 | 4100096 | 93760 |
86.35/2.27/23.87 | 0/0 | 7.95 | 12.42/12.50 | 255 | 216720 |
| 4096 | TRTLLM BF16 | FC1 | 2064285184 | 17972224 | 323040 |
84.03/40.83/45.78 | 0/0 | 10.84 | 16.94/18.75 | 168 | 187032 |
| 4096 | TRTLLM BF16 | FC2 | 960271104 | 50072064 | 152864 |
86.17/44.00/48.40 | 0/0 | 6.98 | 10.91/12.50 | 255 | 219800 |
| 32 | CuteDSL W4A16 | FC1 | 349258752 | 3499264 | 65632 |
70.11/6.61/53.35 | 0/0 | 15.86 | 24.77/25.00 | 128 | 215040 |
| 32 | CuteDSL W4A16 | FC2 | 174409728 | 3022592 | 37504 |
61.78/5.74/49.50 | 0/0 | 15.61 | 24.39/25.00 | 128 | 206880 |
| 4096 | CuteDSL W4A16 | FC1 | 637545728 | 13989120 | 165696 |
51.26/87.20/87.81 | 0/0 | 15.84 | 24.75/25.00 | 128 | 223232 |
| 4096 | CuteDSL W4A16 | FC2 | 322100992 | 7225344 | 148096 |
28.99/44.90/46.44 | 0/0 | 15.06 | 23.53/25.00 | 128 | 215168 |
Remaining launch counters (all eight kernels): occupancy limits = 32
blocks, 1 block by registers, 1 by shared memory; warp limit = 5/8
blocks for BF16 FC1/FC2 and 4 for W4A16. Cluster dimensions = `(2,1,1)`,
except BF16 at 32 uses `(1,1,1)`. `occupancy_cluster_pct` = 3.12%;
`occupancy_cluster_gpu_pct` = 0.59/0.39% for BF16 FC1/FC2 and 0.78% for
both W4A16 kernels (launch estimates, not measured GPU activity). Full
raw CSV/report metadata is retained.
**Interpretation:** at 4096, BF16 still reaches 84–86% DRAM throughput
with 41–44% tensor activity; W4A16 FC1 reaches 87% tensor activity. BF16
reads 3.025 GB versus W4A16's 0.960 GB. EP8 leaves 3715 local
assignments across 32 experts (116/expert); 4096 is not each expert's
GEMM M. No local load/store traffic was measured in any kernel.
Registers and shared memory both limit residency to one block/SM; low
occupancy alone does not establish an avoidable defect. These counters
support a bandwidth explanation and do not demonstrate a spilling or
launch-overhead bug.
Public CLI equivalent on `2079fa00`: TRTLLM BF16 at 32 passed normal
cache save then fresh-process `--no-autotune` load (both exit 0, target
cache hit, no retuning, unchanged cache hash). The measured counters
above used the helper-frozen tactics, not this smoke cache. Run in the
same inference environment with NVFP4 overrides unset. Require a
target-operation `Config cache hit`; missing/incompatible caches can
otherwise fall back.
```bash
export CUDA_VISIBLE_DEVICES=0 FLASHINFER_DISABLE_VERSION_CHECK=1 MAX_JOBS=8 FLASHINFER_LOGGING_LEVEL=info
m=(dram__bytes_{read,write}.sum {dram__throughput,sm__throughput,sm__pipe_tensor_cycles_active}.avg.pct_of_peak_sustained_elapsed gpu__time_duration.sum l1tex__t_sectors_pipe_lsu_mem_local_op_{ld,st}.sum sm__warps_active.avg.{per_cycle_active,pct_of_peak_sustained_active} sm__maximum_warps_per_active_cycle_pct launch__{registers_per_thread,shared_mem_per_block,occupancy_limit_{blocks,registers,shared_mem,warps},cluster_dim_{x,y,z},occupancy_cluster_pct,occupancy_cluster_gpu_pct})
for n in 32 4096; do
for b in cute-dsl-w4a16 trtllm-bf16; do
a=(benchmarks/bench_moe_deepseek.py --num-tokens "$n" --ep 8 --tp 1 --routing-input-mode logits --use-per-token-activation --include-activation-quant --profile-cuda --profile-backend "$b" --profile-iters 1 --cache "$b-$n.json")
python "${a[@]}"
test -s "$b-$n.json" || exit 1
ncu --profile-from-start off --replay-mode kernel --graph-profiling node --nvtx --nvtx-include moe/ --kernel-name-base demangled --rename-kernels off --kernel-name 'regex:.*(Sm100W4A16GroupedGemmKernel|bmm_Bfloat16_Bfloat16Bfloat16_).*' --launch-count 2 --cache-control all --clock-control none --metrics "$(IFS=,; echo "${m[*]}")" --csv --page raw --print-units base --export "$b-$n" python "${a[@]}" --no-autotune
done
done
```
**Why 4096 remains bandwidth-sensitive:** the actual EP8 routing has
3715/32 = 116.09 rows per expert. Useful BF16 GEMM FLOPs divided by
unique weight bytes give 116.09 FLOP/byte; including minimum activation
traffic gives about 108–112. The advertised dense BF16 rate is 36
PFLOPS/node ÷ 2 ÷ 8 = 2.25 PFLOPS/GPU ([NVIDIA HGX
specifications](https://www.nvidia.com/en-us/data-center/hgx/)); with [8
TB/s
HBM](https://docs.nvidia.com/enterprise-reference-architectures/hgx-ai-factory/latest/components.html),
the ideal crossover is about 281 FLOP/byte. These are analytical peak
bounds, supported here by the measured DRAM counters. Tile128 pads to 40
tiles/5120 slots (72.56% useful rows). The 3.025 GB measured BF16 reads
are close to 2.819 GB unique weights and below 3.523 GB if every token
tile reread them. Compression moves W4A16 GEMM1 toward compute
saturation; equal tensor-core input precision does not give equal memory
traffic.
**Reproduce the eight Nsight traces** from the implementation checkout
and image above. Output files remain on the retained devbox; no
compilation caches are removed.
```bash
git checkout 2079fa00c4a2a0fb2e95bb284978a703653a4830
export CUDA_VISIBLE_DEVICES=0 FLASHINFER_DISABLE_VERSION_CHECK=1 MAX_JOBS=8
for mode in inference rl; do
unset FLASHINFER_NVFP4_4OVER6 FLASHINFER_NVFP4_4OVER6_E4M3_USE_256
unset FLASHINFER_NVFP4_4OVER6_ERR_MODE FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH
unset FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH
extra=()
if [[ $mode == rl ]]; then
export FLASHINFER_NVFP4_4OVER6=1 FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1
export FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1
export FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH=1
extra=(--no-fused-finalize)
fi
for n in 32 4096; do
for backend in cute-dsl-w4a16 trtllm-bf16; do
nsys profile --sample=none --cpuctxsw=none --trace=cuda,nvtx \
--cuda-graph-trace=node --capture-range=cudaProfilerApi \
--capture-range-end=stop -o "$mode-$backend-t$n" \
python3 benchmarks/bench_moe_deepseek.py --num-tokens "$n" --ep 8 \
--warmup 10 --iters 100 --use-per-token-activation --include-activation-quant \
--profile-cuda --profile-backend "$backend" --profile-iters 20 "${extra[@]}"
nsys export --type=sqlite -o "$mode-$backend-t$n.sqlite" "$mode-$backend-t$n.nsys-rep"
done
done
done
```
Aggregate all GPU activities contained by each synchronized `moe` NVTX
range, including target memsets; verify 20 ranges and one graph launch
per range. Report per-kernel `sum(end-start)/20`, per-range GPU
`max(end)-min(start)`, and CPU graph-launch duration separately. Exclude
`l2_flush`; do not sum kernel durations as wall-clock latency. These are
single-GPU local-expert diagnostics, not distributed or full-model
speedups.
### Interpretation and limits
- Single-GPU EP8 shard shapes: no dispatch/combine, all-gather,
all-reduce, network or full-model runtime is timed.
- Both TRTLLM arms use native finalize in every configuration. Native
logits routing stores selected weights in BF16; CuTe route weights are
FP32. This is not a numerical-parity claim between BF16/quantized paths
or proof of RL-training determinism.
- New per-tensor runs also remeasure the unaffected W4A16/BF16 paths.
Differences across separate sweeps are not attributed solely to
activation scaling. Coverage is B300/SM103 and the stated shapes; other
architectures and distributed execution were not tested.
## 🔍 Related Issues
Related implementation and benchmark configurations: #4048.
## 🚀 Pull Request Checklist
Thank you for contributing to FlashInfer! Before we review your pull
request, please make sure the following items are complete.
### ✅ Pre-commit Checks
- [x] I have installed `pre-commit` by running `pip install pre-commit`
(or used your preferred method).
- [x] I have installed the hooks with `pre-commit install`.
- [ ] I have run the hooks manually with `pre-commit run --all-files`
and fixed any reported issues.
> If you are unsure about how to set up `pre-commit`, see [the
pre-commit documentation](https://pre-commit.com/).
## 🧪 Tests
- [ ] Tests have been added or updated as needed.
- [ ] All tests are passing (`unittest`, etc.).
- **Original measured source:** seven existing BF16 GPU tests passed
(310 deselected); eight independent reference cases passed with
unchanged `rtol=atol=0.03`, three exact eager repeats and three exact
CUDA-graph replays per case. BF16 CLI/profiler smoke passed. [Complete
original test output, helper and
commands](https://github.com/flashinfer-ai/flashinfer/pull/4985#issuecomment-5549485234).
- **Shared-helper refactor `a3fce87d`:** 43 paired GPU cases matched
measured source `1ac2233d` for public entry points, bound kernel
arguments, finalized output bytes and timing inputs; renamed selectors
and NVFP4 profiler smoke passed; an excluded profiling backend is
rejected with exit 2. Autotuning was disabled for these bounded
equivalence checks. [Complete refactor validation and
reproduction](https://github.com/flashinfer-ai/flashinfer/pull/4985#issuecomment-5549433238).
- **Later head `2079fa00`:** printed-name fixtures and external-profiler
NVTX ranges are validated separately; the five benchmark functions and
non-profile timing branch remain AST-identical to the 43-case refactor
source. Those 43 cases are attributed to the refactor commit above.
- **Static scope:** file-scoped pre-commit hooks, Ruff, formatting,
Python compilation and `git diff --check` passed. Mypy is skipped by the
configured hook because this benchmark is outside its `flashinfer/`
scope. No upstream test files were changed; all-files hooks and the full
repository test suite were not run. These correctness checks do not
replace the nine performance sweeps.
## 🔬 Experimental Track
<!-- Only for PRs submitted under the experimental policy
(CONTRIBUTING.md → "Experimental APIs and Backends").
Leave this section untouched for normal PRs. -->
- [ ] This PR is **experimental**: it adds or changes code under
`flashinfer/experimental/` and/or an `@flashinfer_experimental_api`.
Tracking issue: #
- [ ] The tracking issue names an owner, the reason for the experimental
path, and a graduation plan with a target release.
- [ ] Core changes are limited to a thin entry point (signature, shared
validation, feature-gate check, backend selection, handoff).
- [ ] Tests live in `tests/experimental/` and were validated on the
intended hardware; a runnable example is included.
- [ ] Nothing is registered in `flashinfer/aot.py`, and no experimental
backend is reachable from `backend="auto"` without
`FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1`. (Calling an
`@flashinfer_experimental_api` or naming a backend explicitly is itself
the opt-in and needs no environment variable.)
- [ ] **Test scope declared below.** The experimental CI lane runs
exactly these targets, so keep them as narrow as the change allows.
<!-- Required for experimental PRs. Replace the commented lines below
with your targets.
Do not delete the fence or change its `experimental-tests` tag — the
experimental-track
watcher reads it verbatim to decide which targets to ask CI for. -->
```experimental-tests
# One target per line: a directory or a file. (A pytest ::selector is not
# supported -- the sharding runner cannot consume one.) Must be under
# tests/experimental/ and must exist. Delete these comment lines and add yours, e.g.
#
# tests/experimental/test_my_backend.py
# tests/experimental/my_backend/
#
# Declaring the whole tree (tests/experimental/) is allowed but means every
# experimental PR pays for every other feature's tests, in every matrix cell.
```
## Reviewer Notes
Please review the precision-specific TRTLLM weight/scale arguments,
shared routing and timing boundary, explicit backend/finalize labels,
and the three configuration definitions. All nine raw tables and both
W4A16 baseline speedup columns are above; complete validation output is
linked in Tests. Unchecked all-files/full-suite items reflect the stated
scope.
<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit
* **New Features**
* Unified TensorRT-LLM benchmarking for NVFP4 and BF16 precision modes.
* Added support for logits and pre-routed execution scenarios.
* Added the `trtllm-bf16` backend to benchmark selection, validation,
tables, speedup reporting, and command-line help.
* Improved consistency for routed workload comparisons across
TensorRT-LLM precision modes.
* **Bug Fixes**
* Profile-backend selection now rejects values not included in the
selected backend list.
* Benchmark status and winner labels now clearly distinguish NVFP4 and
BF16 results.
<!-- end of auto-generated comment: release notes by coderabbit.ai -->
This was referenced Sep 13, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
📌 Description
@HumansAnd
Adds an explicit
quant_mode="w4a16"path to the SM100 CuTe DSL NVFP4 MoE APIs.Design walkthrough
quant_mode="w4a16"mirrors the existing SM12x CuTe DSL MoE quant-mode dispatch instead of inferring the path from tensor dtype. W4A4/NVFP4 remains the default; W4A16 requires BF16x,x_sf=None,fc2_input_scale=None, and no per-token scale. The unified runner uses the existingMoEConfig.quant.variant,RoutingInputMode.PackedPrecomputed,TunableRunner,use_fused_finalize, andenable_pdlcontracts.moe_sort-> BF16moe_permute-> W4A16 GEMM1 -> W4A16 GEMM2 -> optionalmoe_unpermute. Sorting reuses the existing DeepSeek routing implementation; permutation and deterministic combine reuse the repository'smoePermuteKernelandmoeUnpermuteKernel. Fused finalize also reusesmoe_output_memset_inplacebefore GEMM2's in-kernel atomic reduction.use_fused_finalizecontract andTruedefault are preserved. Fused mode applies routing weights and atomically reduces in GEMM2; deterministic mode writes expanded route rows and lets sharedmoe_unpermuteapply top-k scales and combine them.moe_sortproduces the same expert/tile maps and EP-local metadata consumed by existing CuTe DSL MoE paths. Preallocated sort/workspace buffers preserve CUDA-graph addresses. Serving-owned 6D weight-scale views are read directly, so in-place updates remain visible.TunableRunner/AutoTuner, hybrid token buckets, cold-L2 profiling, per-shapecan_implementfiltering, and fixed fallback. The tactic key is(gemm1_tactic, gemm2_tactic); each member is((mma_m, route_tile, 256), (cluster_m, 1), True)for(mma_tiler_mnk, cluster_shape_mn, raster_along_m), and both shareroute_tilein{8, 16, 32, 64, 128, 192}. Unlike W4A4's Cartesian product, W4A16 uses curated 1-CTA, 2-CTA, and mixed GEMM pairs because grouped scheduling fixes cluster N to 1 and pipeline storage excludes some tile pairs.51b6e718records a W4A8 prototype and3e1283a8removes it after noncompetitive results, so W4A8 is intentionally outside this PR. It can reuse thequant_modeinterface later, potentially with PTX 9.4'smul.rn.e4m3x4.e2m1x4.e4m3x4.satfinite fp8x4, fp4x4, scale_x4;, a packed E2M1x4 x E4M3x4 to E4M3x4 operation onsm_100a.Performance
Workloads
Both isolated comparisons use B200, CUDA graphs, CUPTI timing, and an EP8 shape simulation with 32 local experts. Initial activation quantization is timed for W4A4 and TRTLLM; W4A16 consumes BF16 directly.
Speedup is TRTLLM latency divided by W4A16 latency; values above 1 favor W4A16. TRTLLM falls back to its default tactic for 1, 2, and 4 tokens because those shapes are outside its tuning buckets.
CUDA 13.2 W4A16 NSYS & NCU
GEMM1:
GEMM1 is mostly BF16 tensor core bound. PDL gives a 20us~30us overlap, showed by both NSYS and NCU.
GEMM2:
Pipeline utilization is lower maybe due to the early PDL overlap.
CUDA 13.0 W4A16 vs TRTLLM
Current head
b4c44feaon B200 with PyTorch 2.11.0+cu130.Inference
Deterministic RL
CUDA 13.2 W4A16 vs TRTLLM
Current head
b4c44feaon B200 with PyTorch 2.13.0+cu132.Inference
Deterministic RL
CUDA 13.0 end-to-end multi-GPU RL workload
Current head
b4c44feaon 8x B200, PyTorch 2.11.0+cu130, CuTe DSL compiler 13.1, CUPTI 13.0.85, and NCCL 2.30.7.CUDA 13.0 Nsight Systems RL breakdown
Nsight Systems 2026.1.3 captured 20 iterations for each EP8/TP8 and W4A4/W4A16 case. Values are aggregate staged kernel time divided by 8 ranks x 20 iterations; concurrent kernels are not removed.
32 global tokens
4096 global tokens
CUDA 13.2 end-to-end multi-GPU RL workload
Current head
b2442579on 8x B200, PyTorch 2.13.0+cu132, CuTe DSL 4.6.1, CUPTI 13.2.75, and NCCL 2.29.7.W4A4 EP8 is the baseline and includes per-token NVFP4 4over6 MSE quantization plus deterministic finalize; W4A16 starts from BF16. EP includes dispatch/combine, and TP includes all-gather/all-reduce. Each cell is the median of three runs with 10 warmups and 100 maximum-rank CUDA-event samples; speedup is W4A4 EP8 latency divided by strategy latency.
CUDA 13.2 Nsight Systems RL breakdown
Nsight Systems 2025.6.3 captured 20 iterations for each EP8/TP8 and W4A4/W4A16 case. Values are aggregate staged kernel time divided by 8 ranks x 20 iterations; concurrent kernels are not removed.
32 global tokens
4096 global tokens
🔍 Related Issues
🚀 Pull Request Checklist
Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete.
✅ Pre-commit Checks
pre-commitby runningpip install pre-commit(or used your preferred method).pre-commit install.pre-commit run --all-filesand fixed any reported issues.🧪 Tests
unittest, etc.).Coverage includes functional and wrapper APIs, W4A4/W4A16, SwiGLU/ReLU2, PDL on/off, both finalize modes, EP1/EP8, tactic boundaries, CUDA graphs, in-place weight-scale updates, and the EP compute bridge.
Reviewer Notes
autotune(True)before capture; capture then reuses the selected tactic and runner-owned workspace.Summary by CodeRabbit
New Features
Improvements
Tests