[Performance]: [V100 / sm_70] Volta build enablement + fused-MoE BLOCK_K performance finding

[KumphanartDansiri]

Proposal to improve performance

I selected this as a performance discussion rather than a bug report because this is a measured Volta/V100 enablement and tuning note, with a prepared PR path if it is useful to Aphrodite.

I searched the open issues and did not find an exact duplicate. This is specifically about CUDA V100 / sm_70 enablement and a fused-MoE decode performance cliff on Volta, not ROCm/CPU install, FlashInfer, FP8 KV cache, or eager-mode TP behavior.

Summary:

1. V100 / sm_70 build enablement

   • I have a small 4-commit patch series that re-enables Tesla V100 (sm_70) for Aphrodite 0.21.
   • The series re-adds 7.0 to the CUDA >=12.8 arch list, excludes sm80-only EXL3 kernels, guards Marlin op registrations that trap on Volta, and adds V100 smoke/build tooling.
   • I also used a CUDA 12.6 / cu126 Dockerfile path.
   • Validated on 8x V100-SXM2-32GB.

2. FP16/BF16 fused-MoE BLOCK_K performance issue on Volta

   • The decode/small-M path chooses BLOCK_SIZE_K=128.
   • On V100 this causes severe register spill / memory traffic.
   • Changing small-M Volta decode to BLOCK_SIZE_K=64 restores expected MoE behavior.
   • This affects FP16/BF16 MoE decode, not dense models and not my custom FP8 path.

3. Possible follow-up: compressed-tensors FP8 on sm_70

   • Aphrodite's compressed_tensors_w8a16_fp8 path appears sm75-gated.
   • I have a working out-of-tree V100 FP8 W8A16 plugin in a separate repo, but that part is not PR-ready for Aphrodite yet.
   • I would treat this as future work after the build + MoE path.

If broad V100 support is useful to Aphrodite, I can prepare this as one or more PRs. My suggested split would be:

• PR 1: V100 / sm_70 build re-enable and smoke tooling.
• PR 2: FP16/BF16 fused-MoE Volta small-M heuristic + V100 config JSONs.
• Later discussion/PR: compressed-tensors FP8 fallback for sm_70.

If this direction is welcome, I can split the work into small PRs and run Aphrodite's formatter (sh ./formatting.sh) before review.

Report of performance regression

Environment:

• Hardware: 8x NVIDIA Tesla V100-SXM2-32GB
• Architecture: Volta / sm_70
• Stack tested: Aphrodite 0.21 source build
• Toolchain: CUDA 12.6 / cu126 torch path
• Serving mode: cudagraph decode

Build validation:

• With the local sm_70 re-enable series, Aphrodite 0.21 builds and serves modern models on V100.
• Example smoke result: gemma-4-31B around 29 tok/s on V100, roughly in line with my vLLM 0.19 V100 path.

MoE performance finding:

On Volta, the fused-MoE decode/small-M path using BLOCK_SIZE_K=128 is much slower than expected. The practical symptom is that sparse MoE can become slower than a dense model of similar class, which is backwards for decode.

Measured in my vLLM-side repro on the same V100 box:

Model	Stock FP16 MoE	Patched/tuned MoE	Improvement
Qwen3.6-35B-A3B, C1	~15.6 tok/s	~65.9 tok/s	~4.2x
Qwen3.6-35B-A3B, C8 aggregate	~24.9 tok/s	~173.9 tok/s	~7.0x aggregate
Gemma-4-26B-A4B, C1	~10.9 tok/s	~43.7 tok/s	~4.0x
Gemma-4-26B-A4B, C8 aggregate	~28.3 tok/s	~155.9 tok/s	~5.5x aggregate

Root cause summary:

• The first suspect was num_stages, because V100 lacks cp.async.
• Sweeping num_stages did not explain the slowdown.
• The real lever was BLOCK_SIZE_K.
• Small-M decode with BLOCK_SIZE_K=128 spills badly on Volta.
• BLOCK_SIZE_K=64 removes the pathology.
• Output remained bit-identical in the tested A/B; this is a speed/config fix, not an approximation.

A possible Aphrodite heuristic shape:

if has_device_capability(70) and not has_device_capability(80) and M <= 64:
    if M <= 4:
        return {
            "BLOCK_SIZE_M": 16,
            "BLOCK_SIZE_N": 32,
            "BLOCK_SIZE_K": 64,
            "GROUP_SIZE_M": 1,
            "num_warps": 4,
            "num_stages": 2,
        }
    return {
        "BLOCK_SIZE_M": 16,
        "BLOCK_SIZE_N": 128,
        "BLOCK_SIZE_K": 64,
        "GROUP_SIZE_M": 1,
        "num_warps": 8,
        "num_stages": 2,
    }
Scope/caveats:
This is for FP16/BF16 MoE decode on Volta.
It does not claim V100 beats modern GPUs.
It does not touch dense models.
It does not depend on my FP8 plugin.
The goal is to remove a Volta-specific default/config pathology and make Aphrodite's broad-architecture story stronger.
Reference write-up / evidence repo:
https://github.com/KumphanartDansiri/v100-vllm-2026
Chapter 2 documents the MoE BLOCK_K finding.
The prepared Aphrodite packet is under upstream_feedback/aphrodite/.

### Misc discussion on performance

The main reason I am filing this as a discussion first is that I am not sure whether Aphrodite currently wants to carry V100 / `sm_70` support as an upstream path.

From my side, the work is prepared and measured, but I do not want to open a PR if the project direction has changed. If maintainers are interested, I can provide the patch branch / split PRs.

My preferred upstream shape would be:

1. First discuss whether V100 / `sm_70` is in scope.
2. If yes, send the build re-enable PR.
3. Then send the FP16/BF16 MoE small-M heuristic + tuned config JSONs.
4. Keep compressed-tensors FP8 for a later discussion, because that needs more cleanup before it is PR-ready.

I am happy to adjust the format if this belongs as a feature request or PR instead of a performance discussion.

### Your current environment (if you think it is necessary)

Hardware:
- 8x NVIDIA Tesla V100-SXM2-32GB
- Volta / sm_70
- NVLink box

Software / build context:
- Aphrodite 0.21 source build experiment
- CUDA 12.6 / cu126 torch path
- Python editable/source-build workflow
- Serving mode tested with cudagraph decode

Related reference implementation / measurements:
- vLLM 0.19 + 0.21 source builds on the same V100 box
- Public write-up: https://github.com/KumphanartDansiri/v100-vllm-2026
- Companion V100 FP8/plugin repo: https://github.com/KumphanartDansiri/vllm-fp8-w8a16-sm70

Notes:
- I am not submitting a patch in this issue yet.
- I can run Aphrodite formatting (`sh ./formatting.sh`) before opening a PR.
- The local V100/Aphrodite patch series is prepared but not included in this issue; I can provide it as a branch or PR if maintainers want it.