vLLM 对不同 GPU SM 架构的适配机制深度剖析
·AI
vLLM 对不同 GPU SM 架构的适配机制深度剖析
前几篇 vLLM 文章分别剖析了 PagedAttention、fused_moe kernel、CUTLASS 集成和分布式推理。但有一个贯穿所有这些组件的横切关注点一直没有展开：vLLM 如何在一套代码中适配从 V100 到 B200 跨越七年的 GPU 架构？
这不是一个简单的 if-else。从 SM70（Volta）到 SM121（DGX Spark），每一代 GPU 引入了不同的硬件特性——Tensor Core 指令集不同、内存层级不同、支持的数值精度不同。vLLM 需要在每个计算路径上做出正确的后端选择，才能在每一代硬件上都跑出接近最优的性能。
这篇文章直接从 vLLM 源码出发，系统梳理 Attention、GEMM、MoE、量化四个维度的 SM 适配机制。
一、GPU 架构代际速览
先建立硬件背景。从 2017 年到 2025 年，NVIDIA 的数据中心 GPU 经历了以下代际：
| SM 版本 | 架构名 | 代表 GPU | 年份 | 关键特性 |
|---|---|---|---|---|
| SM70 | Volta | V100 | 2017 | 首代 Tensor Core (FP16)，HBM2 |
| SM75 | Turing | T4, RTX 2080 | 2018 | INT8/INT4 Tensor Core，混合精度推理 |
| SM80 | Ampere | A100 | 2020 | BF16, TF32, 2:4 稀疏，cp.async, mma.sync |
| SM89 | Ada Lovelace | L40S, RTX 4090 | 2022 | FP8 (E4M3/E5M2) Tensor Core |
| SM90 | Hopper | H100, H200 | 2023 | WGMMA, TMA, Thread Block Cluster, FP8 native |
| SM100 | Blackwell DC | B100, B200, GB200 | 2025 | 5th Gen TC, NVFP4, MXFP8, 2x FP8 吞吐 |
| SM120 | Blackwell RTX | RTX 5090 | 2025 | 消费级 Blackwell, NVFP4 |
| SM121 | Blackwell Spark | GB10 (DGX Spark) | 2025 | ARM SoC, 桌面级 Blackwell |
每一代的关键差异直接影响 vLLM 的 kernel 选择：
- SM70 没有 INT8 Tensor Core → 不支持 AWQ/GPTQ Marlin kernel
- SM80 引入 BF16 → SM75 及以下只能用 FP16
- SM89 引入 FP8 → 原生 FP8 量化在 SM89+ 才可用
- SM90 引入 TMA/WGMMA → CUTLASS 3.x kernel 只在 SM90+ 运行
- SM100 引入 NVFP4 → 4-bit 量化的硬件加速
- SM100 和 SM120 是不同的架构 → SM100 的 kernel 不能直接跑在 SM120 上
二、Attention 后端选择
2.1 选择逻辑源码
Attention 后端的选择是 vLLM 中最复杂的 SM 适配逻辑之一。核心函数在 vllm/platforms/cuda.py：
# vllm/platforms/cuda.py
def _get_backend_priorities(
use_mla: bool,
device_capability: DeviceCapability,
num_heads: int | None = None,
kv_cache_dtype: CacheDType | None = None,
) -> list[AttentionBackendEnum]:
if use_mla:
if device_capability.major == 10: # Blackwell 数据中心
# Sparse MLA 优先级取决于 KV cache 精度和 head 数量
if kv_cache_dtype.startswith("fp8"):
sparse_backends = [FLASHINFER_MLA_SPARSE, FLASHMLA_SPARSE]
elif num_heads <= 16:
sparse_backends = [FLASHINFER_MLA_SPARSE, FLASHMLA_SPARSE]
else:
sparse_backends = [FLASHMLA_SPARSE, FLASHINFER_MLA_SPARSE]
return [
FLASHINFER_MLA, # FlashInfer MLA (SM100 only)
CUTLASS_MLA, # CUTLASS MLA (SM100 only)
FLASH_ATTN_MLA, # FlashAttn MLA (SM90 only)
FLASHMLA, # FlashMLA (SM89/SM100)
TRITON_MLA, # Triton MLA (any SM)
*sparse_backends,
]
else: # 非 Blackwell (Hopper 及以下)
return [
FLASH_ATTN_MLA, # SM90
FLASHMLA, # SM89+
FLASHINFER_MLA, # SM100 only, 在非 Blackwell 上会被跳过
TRITON_MLA, # fallback
FLASHMLA_SPARSE,
]
else: # 标准 Attention
if device_capability.major == 10: # Blackwell
return [FLASHINFER, FLASH_ATTN, TRITON_ATTN, FLEX_ATTENTION]
else:
return [FLASH_ATTN, FLASHINFER, TRITON_ATTN, FLEX_ATTENTION]
vLLM 会按优先级逐个尝试后端，调用每个后端的 supports_compute_capability() 检查是否兼容当前 GPU。
2.2 各后端的 SM 兼容范围
# FlashAttention: 需要 SM80+ (Ampere 及以上)
class FlashAttentionBackend:
@classmethod
def supports_compute_capability(cls, capability):
return capability >= DeviceCapability(8, 0)
# FlashInfer: SM75 到 SM121
class FlashInferBackend:
@classmethod
def supports_compute_capability(cls, capability):
return (capability >= DeviceCapability(7, 5) and
capability <= DeviceCapability(12, 1))
# Triton Attention: 无 SM 限制 (universal fallback)
class TritonAttentionBackend:
@classmethod
def supports_compute_capability(cls, capability):
return True
MLA 后端的 SM 限制更为严格：
| MLA 后端 | SM 要求 | 说明 |
|---|---|---|
| FlashInfer MLA | SM100 only | Blackwell 数据中心专用 |
| CUTLASS MLA | SM100 only | Blackwell CUTLASS kernel |
| FlashAttn MLA | SM90 only | Hopper 专用 |
| FlashMLA | SM89 或 SM100 | Ada Lovelace / Blackwell |
| FlashInfer MLA Sparse | SM100 only | 稀疏注意力 |
| FlashMLA Sparse | SM89 或 SM100 | 稀疏注意力 |
| Triton MLA | 任意 SM | Universal fallback |
2.3 FlashAttention 版本选择
在 SM80+ 的范围内，FlashAttention 还有版本选择逻辑：
# vllm/v1/attention/backends/fa_utils.py
# SM90 (Hopper): 优先 FA3
# SM100 (Blackwell): 优先 FA4
# 其他 SM80+: 回退到 FA2
# 注意：FA3 无法在 Blackwell 上运行（显式 fallback）
# FA4 有 TMEM 容量限制：head_size 必须 <=128 或恰好 192
2.4 SM70 的困境
V100（SM70）是 vLLM 支持的最低端 GPU，但面临多重限制：
- FlashAttention 要求 SM80+ → 不可用
- FlashInfer 要求 SM75+ → 不可用
- 只剩 Triton Attention 作为 fallback → 性能远低于 FlashAttention
- 不支持 BF16 → 只能用 FP16
- 没有 INT8 Tensor Core → 量化选择受限
这就是为什么社区分支 1Cat-vLLM 需要额外的适配工作（见第八节）。
三、GEMM 后端选择
3.1 CUTLASS SM 版本分发
CUTLASS 量化 GEMM 的 SM 分发是 vLLM 中最清晰的架构适配模式：
// csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu
void cutlass_scaled_mm(torch::Tensor& c, ...) {
int32_t version_num = get_sm_version_num();
if (version_num >= 120) cutlass_scaled_mm_sm120(...); // RTX Blackwell
else if (version_num >= 100) cutlass_scaled_mm_sm100(...); // Blackwell DC
else if (version_num >= 90) cutlass_scaled_mm_sm90(...); // Hopper
else if (version_num == 89) cutlass_scaled_mm_sm89(...); // Ada Lovelace
else if (version_num >= 80) cutlass_scaled_mm_sm80(...); // Ampere
else if (version_num >= 75) cutlass_scaled_mm_sm75(...); // Turing
// SM70 及以下: 不支持 CUTLASS 量化 GEMM
}
每个 SM 版本使用不同的 CUTLASS 代际和特性：
| SM | CUTLASS 版本 | Tensor Core 指令 | 特有调度策略 |
|---|---|---|---|
| SM75 | 2.x | WMMA (INT8) | 基础流水线 |
| SM80 | 2.x | mma.sync (INT8/FP16) |
cp.async, 多级流水线, SplitK |
| SM89 | 2.x | mma.sync (FP8) |
同 SM80 + FP8 支持 |
| SM90 | 3.x | WGMMA (FP8/INT8) | TMA, Warp Specialization, Pingpong/Cooperative |
| SM100 | 3.x | 5th Gen TC | KernelScheduleAuto, 更大 Cluster |
| SM120 | 3.x | RTX Blackwell TC | 独立编译目标，与 SM100 不共享 kernel |
3.2 FP8 支持的能力检查
// csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu
bool cutlass_scaled_mm_supports_fp8(int64_t capability) {
// SM90 (Hopper): CUDA 12.0+
if (capability >= 90) return CUDA_VERSION >= 12000;
// SM89 (Ada): CUDA 12.4+ (需要更新的驱动)
if (capability >= 89) return CUDA_VERSION >= 12040;
// SM80 及以下: 不支持 FP8
return false;
}
bool cutlass_scaled_mm_supports_block_fp8(int64_t capability) {
// SM100 (Blackwell): CUDA 12.8+
if (capability >= 100) return CUDA_VERSION >= 12080;
// SM90 (Hopper): CUDA 12.0+
if (capability >= 90) return CUDA_VERSION >= 12000;
return false;
}
bool cutlass_scaled_mm_supports_fp4(int64_t capability) {
// NVFP4 只在 SM100+ 且 CUDA 12.8+ 才支持
return capability >= 100 && CUDA_VERSION >= 12080;
}
3.3 默认 GEMM 路径
对于未量化的 BF16/FP16 密集 GEMM，vLLM 始终使用 cuBLAS（通过 torch.mm），不涉及 SM 版本分发——cuBLAS 内部自动选择最优 kernel。
CUTLASS 只用于量化 GEMM（需要 fused dequantize/scale 的场景）。
四、MoE Kernel 选择
MoE 的后端选择比 Attention 和 GEMM 更复杂，因为涉及多种量化格式和多种硬件后端的交叉：
4.1 FP8 MoE 后端优先级
# vllm/model_executor/layers/fused_moe/oracle/fp8.py
# 全局优先级（简化）：
# AITER > FlashInfer+TRTLLM > FlashInfer+CUTLASS > DeepGEMM
# > vLLM CUTLASS > Triton > Marlin > Batched 变体
# Hopper (SM90) 特殊规则：
if sm90 and block_fp8:
if tensor_parallel:
prefer Triton # Triton fused_moe 在 TP 场景更快
if expert_parallel:
prefer FlashInfer_CUTLASS # EP 需要 batched 格式
4.2 各 MoE 后端的 SM 要求
| MoE 后端 | SM 要求 | 量化格式 | 说明 |
|---|---|---|---|
| Triton fused_moe | 任意 | BF16/FP8 | Universal，但 FP8 性能次优 |
| vLLM CUTLASS MoE | SM90+ | FP8 | Grouped GEMM |
| FlashInfer CUTLASS MoE | SM90/100/110/120 | FP8/NVFP4/MXFP4 | 多平台支持 |
| FlashInfer TRTLLM MoE | SM100 only | FP8/BF16/NVFP4 | Blackwell DC 专用 |
| FlashInfer CuteDSL MoE | SM100 only | - | Blackwell DC 专用 |
| DeepGEMM MoE | SM90+ | FP8 block-wise | N > 512 场景 |
| TRTLLM FP8 MoE | SM100 only | FP8 | TensorRT-LLM kernel |
| TRTLLM NVFP4 MoE | SM100 only | NVFP4 | TensorRT-LLM kernel |
| Batched DeepGEMM MXFP8 | SM100 only | MXFP8 | EP 场景 |
| Marlin MoE | SM75+ | WNA16 | Weight-only 量化 fallback |
4.3 Blackwell 上的 MoE 特殊逻辑
# vllm/model_executor/layers/fused_moe/triton_cutlass_moe.py
class TritonOrCutlassExperts(FallbackExperts):
"""Blackwell 上的 MoE 后端选择"""
def _select_experts_impl(self, hidden_states, w1, w2):
if self.is_sm100 and hidden_states.shape[0] <= 8:
# Blackwell 小 batch：CUTLASS grouped GEMM launch 开销大
# 回退到 Triton（单 kernel, 更低 launch overhead）
return self.fallback_experts # TritonExperts
else:
return self.experts # CutlassExpertsFp8
4.4 NVFP4 和 MXFP8 MoE
这两种格式是 Blackwell 的新特性：
# NVFP4 MoE 优先级：
# FlashInfer+TRTLLM (SM100) > FlashInfer+CuteDSL (SM100)
# > FlashInfer+CUTLASS (SM100/110/120) > vLLM CUTLASS (SM100/110/120) > Marlin
# MXFP8 MoE：仅支持 FlashInfer+TRTLLM (SM100)
# MXFP4 MoE 路径更复杂：
# SM100: FlashInfer+TRTLLM / FlashInfer+CUTLASS
# SM90: FlashInfer+CUTLASS (BF16 mode)
# SM90-109: Triton (但 SM110 有 bug, SM120 需要 Triton 修复)
# Fallback: Marlin
五、量化支持矩阵
5.1 get_min_capability() 汇总
vLLM 的每个量化方案都通过 get_min_capability() 声明最低 SM 要求：
# 直接从源码提取的最低 SM 版本要求
class GGUFConfig: get_min_capability = 60 # Pascal+
class GPTQConfig: get_min_capability = 60 # Pascal+
class BitsAndBytesConfig: get_min_capability = 70 # Volta+
class AWQConfig: get_min_capability = 75 # Turing+
class AWQMarlinConfig: get_min_capability = 75 # Turing+
class GPTQMarlinConfig: get_min_capability = 75 # Turing+
class Fp8Config: get_min_capability = 75 # Turing+ (Marlin fallback on <SM89)
class ModelOptFp8: get_min_capability = 89 # Ada Lovelace+
class ModelOptNvfp4: get_min_capability = 75 # Turing+ (Marlin fallback)
class FBGEMMFp8Config: get_min_capability = 80 # Ampere+
class ExpertsInt8Config: get_min_capability = 80 # Ampere+
class MxfpConfig: get_min_capability = 80 # Ampere+ (hardware accel SM90+)
class MxFp8Config: get_min_capability = 100 # Blackwell+
5.2 FP8 的双路径逻辑
FP8 量化的 get_min_capability 返回 75（Turing），但这不意味着 Turing 有 FP8 Tensor Core。实际逻辑是：
# vllm/model_executor/layers/quantization/fp8.py
class Fp8Config:
def get_min_capability(cls) -> int:
return 75 # Turing 可以用 Marlin 做 weight-only FP8
def __init__(self, ...):
self.use_marlin = (
not current_platform.has_device_capability(89)
or envs.VLLM_TEST_FORCE_FP8_MARLIN
)
# SM75-88: 用 Marlin kernel (weight-only, dequant on the fly)
# SM89+: 用原生 FP8 CUTLASS kernel (hardware accelerated)
这是一种「优雅降级」模式： 同一个量化格式（FP8），在不同 SM 上选择不同的执行路径。SM89+ 有硬件 FP8 Tensor Core，用 CUTLASS FP8 kernel 获得最佳性能。SM75-88 没有 FP8 硬件支持，但可以用 Marlin kernel 在 FP16 Tensor Core 上做 weight-only FP8 dequantize。
5.3 完整 SM × 功能支持矩阵
SM70 SM75 SM80 SM89 SM90 SM100 SM120 SM121
V100 T4 A100 L40S H100 B200 5090 Spark
───── ───── ───── ───── ───── ───── ───── ─────
数据类型支持
FP32 ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
FP16 ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
BF16 ✗ ✗ ✓ ✓ ✓ ✓ ✓ ✓
FP8 (E4M3) ✗ ✗ ✗ ✓ ✓ ✓ ✓ ✓
NVFP4 ✗ ✗ ✗ ✗ ✗ ✓ ✓ ✓
MXFP8 ✗ ✗ ✗ ✗ ✗ ✓ ✗(?) ✗(?)
Attention 后端
FlashAttention ✗ ✗ ✓FA2 ✓FA2 ✓FA3 ✓FA4 ✓FA2 ✓FA2
FlashInfer ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓
Triton Attn ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
FlexAttention ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
MLA 后端
FlashAttn MLA ✗ ✗ ✗ ✗ ✓ ✗ ✗ ✗
FlashMLA ✗ ✗ ✗ ✓ ✗ ✓ ✗ ✗
CUTLASS MLA ✗ ✗ ✗ ✗ ✗ ✓ ✗ ✗
FlashInfer MLA ✗ ✗ ✗ ✗ ✗ ✓ ✗ ✗
Triton MLA ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
CUTLASS GEMM
INT8 W8A8 ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓
FP8 W8A8 ✗ ✗ ✗ ✓ ✓ ✓ ✓ ✓
Block-wise FP8 ✗ ✗ ✗ ✗ ✓ ✓ ✓(?) ✗(?)
W4A8 混合精度 ✗ ✗ ✗ ✗ ✓ ✓ ✓(?) ✗(?)
NVFP4 GEMM ✗ ✗ ✗ ✗ ✗ ✓ ✓ ✗(?)
2:4 Sparse ✗ ✗ ✗ ✗ ✓ ✓ ✓(?) ✗(?)
Grouped GEMM ✗ ✗ ✗ ✗ ✓ ✓ ✓(?) ✗(?)
MoE Kernel
Triton fused_moe ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
CUTLASS MoE FP8 ✗ ✗ ✗ ✗ ✓ ✓ ✓ ✗(?)
TRTLLM MoE FP8 ✗ ✗ ✗ ✗ ✗ ✓ ✗ ✗
TRTLLM MoE NVFP4 ✗ ✗ ✗ ✗ ✗ ✓ ✗ ✗
DeepGEMM MoE ✗ ✗ ✗ ✗ ✓ ✓(?) ✗ ✗
Marlin MoE ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓
量化方案
GGUF ✗ ✗ ✓ ✓ ✓ ✓ ✓ ✓
GPTQ ✗ ✗ ✓ ✓ ✓ ✓ ✓ ✓
BitsAndBytes ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓
AWQ ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓
GPTQ-Marlin ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓
FP8 (Marlin路径) ✗ ✓ ✓ ✓ ─ ─ ─ ─
FP8 (native) ✗ ✗ ✗ ✓ ✓ ✓ ✓ ✓
FBGEMM FP8 ✗ ✗ ✓ ✓ ✓ ✓ ✓ ✓
ModelOpt FP8 ✗ ✗ ✗ ✓ ✓ ✓ ✓ ✓
ModelOpt MXFP8 ✗ ✗ ✗ ✗ ✗ ✓ ✗(?) ✗(?)
NVFP4 量化 ✗ ✗ ✗ ✗ ✗ ✓ ✓ ✓(?)
✓ = 支持 ✗ = 不支持 ✓(?) = 可能支持但尚未完全验证 ─ = 不适用
六、Platform 抽象层
6.1 DeviceCapability
vLLM 通过 DeviceCapability NamedTuple 封装 SM 版本：
# vllm/platforms/interface.py
class DeviceCapability(NamedTuple):
major: int # SM major version (e.g., 9 for Hopper)
minor: int # SM minor version (e.g., 0 for SM90)
def to_int(self) -> int:
"""SM90 → 90, SM100 → 100, SM121 → 121"""
return self.major * 10 + self.minor
# 支持比较运算：
# DeviceCapability(9, 0) >= DeviceCapability(8, 0) → True
6.2 Capability 检查 API
# vllm/platforms/cuda.py
class CudaPlatformBase(Platform):
@classmethod
def has_device_capability(cls, N: int) -> bool:
"""检查 SM >= N（例如 has_device_capability(89) 检查是否是 Ada+）"""
cap = cls.get_device_capability()
return cap.to_int() >= N
@classmethod
def is_device_capability(cls, N: int) -> bool:
"""精确匹配（例如 is_device_capability(90) 只匹配 Hopper）"""
cap = cls.get_device_capability()
return cap.to_int() == N
@classmethod
def is_device_capability_family(cls, N: int) -> bool:
"""检查 SM 家族（例如 is_device_capability_family(100)
匹配 SM100/SM103，但不匹配 SM120）"""
cap = cls.get_device_capability()
return cap.major * 10 == N
@classmethod
def supports_fp8(cls) -> bool:
return cls.has_device_capability(89)
@property
def supported_dtypes(self) -> list[torch.dtype]:
if self.has_device_capability(80):
return [torch.bfloat16, torch.float16, torch.float32]
if self.has_device_capability(60):
return [torch.float16, torch.float32] # 无 BF16
return [torch.float32]
is_device_capability_family 的重要性： SM100（B200 数据中心）和 SM120（RTX 5090 消费级）虽然都是 Blackwell，但 不是同一个 SM 家族。很多 TRTLLM kernel 用 is_device_capability_family(100) 检查，意味着它们只能在 B100/B200 上运行，不能在 RTX 5090 上运行。这是一个容易踩的坑。
七、Blackwell (SM100/SM120) 的最新适配
7.1 SM100 vs SM120 的关键差异
SM100（数据中心 Blackwell）和 SM120（消费级 RTX Blackwell）是两条独立的编译和适配路径：
# CMakeLists.txt 中的独立编译开关：
ENABLE_SCALED_MM_SM100 # SM100 CUTLASS scaled_mm
ENABLE_SCALED_MM_SM120 # SM120 CUTLASS scaled_mm (独立编译)
ENABLE_NVFP4_SM100 # SM100 NVFP4 kernel
ENABLE_NVFP4_SM120 # SM120 NVFP4 kernel (独立编译)
ENABLE_CUTLASS_MOE_SM100 # SM100 MoE grouped GEMM
ENABLE_CUTLASS_MOE_SM120 # SM120 MoE grouped GEMM (独立编译)
SM100 的 kernel 不能直接在 SM120 上运行——它们使用不同的 CUTLASS 架构目标，需要单独编译和实例化。这是因为 SM100 和 SM120 的 Tensor Core 微架构有差异（SM100 是完整的数据中心规格，SM120 是面积和功耗优化的消费级版本）。
7.2 SM100 独占功能
以下 kernel 只在 SM100 上可用，SM120 无法使用：
| 功能 | 说明 |
|---|---|
| CUTLASS MLA | DeepSeek MLA 注意力 |
| FlashInfer MLA / MLA Sparse | Blackwell MLA kernel |
| FlashAttention 4 | 利用 TMEM 的下一代 FA |
| TRTLLM FP8 MoE | TRT-LLM kernel |
| TRTLLM NVFP4 MoE | TRT-LLM kernel |
| FlashInfer CuteDSL MoE | CuteDSL kernel |
| MXFP8 量化 | Microscale FP8 |
| Batched DeepGEMM MXFP8 | EP 场景 |
7.3 SM120/SM121 的适配进展
SM120（RTX 5090）和 SM121（DGX Spark GB10）的适配仍在积极开发中：
- CUTLASS FP8 scaled_mm：已有 SM120 专用路径
- NVFP4 GEMM：已有 SM120 专用 kernel (
cutlass_scaled_fp4_mm_sm120a) - FlashInfer：支持 SM75 到 SM121 的完整范围
- FlashInfer CUTLASS MoE：显式支持 SM120 家族
- CUTLASS MoE FP4：支持 SM100/SM110/SM120 家族
已知问题：
- Issue #31085：SM120 NVFP4 MoE 支持
- Issue #36821：SM121 DGX Spark/Acer GN100 在 aarch64 上的支持
八、社区适配：1Cat-vLLM 与 V100 支持
8.1 上游 vLLM 在 V100 上的痛点
虽然 vLLM 的 get_min_capability() 最低到 SM60（GPTQ/GGUF），但实际上 V100 (SM70) 上的可用功能非常有限：
V100 (SM70) 的可用路径：
Attention: Triton Attention only (FlashAttention 需要 SM80)
GEMM: cuBLAS FP16 only (无 CUTLASS 量化 GEMM, 无 BF16)
量化: BitsAndBytes, GPTQ (非 Marlin), GGUF
MoE: Triton fused_moe (FP16 only)
AWQ: ✗ (upstream 需要 SM75)
FP8: ✗
Marlin: ✗ (需要 SM75)
对于仍在使用 V100 集群的用户（尤其是学术界和小团队），这意味着很多量化模型无法部署，性能也远低于理论水平。
8.2 1Cat-vLLM 的解决方案
1Cat-vLLM 是一个面向 V100/SM70 的实验性 vLLM 分支，核心改动包括：
① SM70 AWQ Kernel。 上游 AWQ 的 CUDA kernel 使用 INT8 Tensor Core（SM75+）。1Cat-vLLM 集成了 lmdeploy TurboMind 的 SM70 WMMA kernel——使用 Volta 的 WMMA 指令（FP16 Tensor Core）做 weight-only INT4 dequantize + FP16 GEMM：
上游 AWQ (SM75+):
INT4 weights → INT8 dequant → INT8 Tensor Core GEMM → FP16 output
1Cat-vLLM AWQ (SM70):
INT4 weights → FP16 dequant → WMMA FP16 GEMM → FP16 output
使用 lmdeploy TurboMind SM70 WMMA kernel
② MLA Prefill 的 SDPA Fallback。 DeepSeek V3 等 MLA 模型的 prefill 阶段需要 FlashAttention（SM80+）。在 V100 上，1Cat-vLLM 提供了基于 PyTorch 原生 SDPA（Scaled Dot Product Attention）的 fallback 路径：
MLA Prefill 后端优先级 (1Cat-vLLM):
FlashInfer MLA → SM100 only
cuDNN → SM80+ (有时不稳定)
SDPA fallback → SM70+ (PyTorch native, 最慢但兼容性最好)
FlashAttention → SM80+
// _need_sdpa_prefill() 检查 GPU 能力是否低于 SM80
// 如果 cap[0] < 8 (即 SM70)，则使用 SDPA 作为 prefill 后端
③ Triton Attention 修复。 V100 的 Triton 编译有一些特有的 bug（如 shared memory 限制不同），1Cat-vLLM 提供了针对性的运行时修复。
④ 推荐的运行时配置：
# 1Cat-vLLM V100 推荐配置
VLLM_DISABLE_PYNCCL=1 python -m vllm.entrypoints.openai.api_server \
--model Qwen/Qwen2.5-7B-Instruct-AWQ \
--attention-backend TRITON_ATTN \ # V100 只能用 Triton
--disable-custom-all-reduce \ # V100 不支持自定义 AllReduce
--gpu-memory-utilization 0.80 \ # V100 16GB 需要保守一些
--tensor-parallel-size 4 # 4×V100 16GB
8.3 向下兼容的工程挑战
1Cat-vLLM 的存在揭示了一个深层问题：开源 LLM 推理框架面临的「向下兼容税」。
每一代 GPU 的新特性（FP8、TMA、NVFP4）都让 kernel 更快更高效，但也让代码更复杂：
支持 N 代 GPU 的代码复杂度：
每新增一代 GPU：
+ 新的 CUTLASS 编译目标
+ 新的 Attention 后端适配
+ 新的量化格式 kernel
+ 新的 MoE 后端
+ 新的 CI/CD 测试矩阵
当前 vLLM 支持的 SM 范围：SM70 - SM121
需要维护的编译变体：~7 个 CUTLASS 架构目标
需要测试的 GPU 类型：V100, T4, A100, L40S, H100, B200, RTX 5090...
上游 vLLM 的策略是 渐进式放弃旧架构：
- 核心功能（cuBLAS GEMM、Triton Attention）保持广泛兼容
- 新的高性能路径（CUTLASS FP8、MLA、NVFP4）只在新架构上实现
- 社区分支（如 1Cat-vLLM）负责旧架构的深度适配
这是一种合理的工程取舍：让 80% 的用户（使用 Ampere+ GPU）获得最佳体验，同时不阻止 20% 的用户（使用旧 GPU）通过社区方案继续使用。
九、Kernel 选择的运行时流程
把所有维度串起来，看一个完整的 kernel 选择流程：
模型加载时：
┌─────────────────────────────────────────────────────┐
│ 1. 检测 GPU: get_device_capability() │
│ → DeviceCapability(9, 0) = Hopper H100 │
│ │
│ 2. 检查量化配置: Fp8Config │
│ → min_capability = 75 ✓ │
│ → has_device_capability(89) = True │
│ → use_marlin = False (使用 native FP8) │
│ │
│ 3. 选择 Attention 后端: │
│ 优先级: FlashAttn > FlashInfer > Triton │
│ FlashAttn.supports_compute_capability(SM90) = True │
│ → 选择 FlashAttention (FA3 for Hopper) │
│ │
│ 4. 选择 GEMM 后端: │
│ FP8 Linear → cutlass_scaled_mm │
│ → SM90: cutlass_scaled_mm_sm90() │
│ → CUTLASS 3.x, WGMMA, TMA │
│ │
│ 5. 如果是 MoE 模型: │
│ FP8 MoE → oracle/fp8.py │
│ → SM90 + block FP8 + TP: prefer Triton fused_moe │
└─────────────────────────────────────────────────────┘
推理时：
┌─────────────────────────────────────────────────────┐
│ Linear 层: │
│ FP8 input → CUTLASS SM90 FP8 kernel │
│ TileShape<128,128,128>, Pingpong, ClusterShape<2,1>│
│ │
│ Attention 层: │
│ Q,K,V → FlashAttention 3 (Hopper TMA) │
│ PagedAttention KV Cache access │
│ │
│ MoE 层 (if applicable): │
│ Triton fused_moe kernel │
│ Block-wise FP8 dequant inside GEMM loop │
└─────────────────────────────────────────────────────┘
十、总结
vLLM 的 SM 适配机制可以归纳为几个设计原则：
① 分层分发，逐级 fallback。 每个计算路径（Attention、GEMM、MoE、量化）独立做 SM 版本检查和后端选择。高性能 kernel 有严格的 SM 要求，但总有一个 universal fallback（Triton Attention、cuBLAS GEMM、Triton fused_moe）保底。
② 编译期分离，运行时分发。 CUTLASS kernel 按 SM 版本分开编译（sm75、sm80、sm89、sm90、sm100、sm120 各自独立的编译单元），运行时通过 get_sm_version_num() 选择正确的 kernel。这避免了编译时间爆炸，也让每个 kernel 可以用对应架构的最优指令。
③ 优雅降级而非硬性拒绝。 FP8 量化在 SM75 上不是报错，而是回退到 Marlin kernel 做 weight-only dequant。MLA 注意力在没有 FlashAttention 的 GPU 上回退到 Triton MLA。用户不需要知道底层选了哪个 kernel——vLLM 自动选择当前硬件上的最优路径。
④ SM100 和 SM120 是两个世界。 这是一个反直觉但重要的事实：数据中心 Blackwell (B200) 和消费级 Blackwell (RTX 5090) 不共享 kernel。SM100 独占的功能（CUTLASS MLA、TRTLLM MoE、MXFP8）不会在 SM120 上运行。
⑤ 社区驱动的长尾兼容。 上游 vLLM 聚焦主流硬件的性能优化（A100/H100/B200），旧架构的深度适配留给社区分支。1Cat-vLLM 为 V100 集成了 WMMA kernel 和 SDPA fallback，这种模式让 vLLM 既能保持主线代码的简洁，又能覆盖长尾用户。
从更大的视角看，vLLM 的 SM 适配机制是 GPU 软件生态碎片化的一个缩影。每一代新硬件都带来新的指令集和新的最优实践，但旧硬件不会立即退役。在「追求极致性能」和「保持广泛兼容」之间取得平衡，是每个 GPU 软件框架都面临的永恒挑战。
参考资源
- vLLM 源码: platforms/cuda.py
- vLLM 源码: attention backends
- vLLM 源码: CUTLASS SM dispatch
- vLLM 源码: MoE oracle
- 1Cat-vLLM: V100/SM70 实验分支
- NVIDIA CUDA Compute Capability
- CUTLASS: CUDA Templates for Linear Algebra
- FlashInfer: Efficient Attention Engine (MLSys 2025)
- vLLM + NVIDIA InferenceMAX on Blackwell
- vLLM SM120 NVFP4 MoE Support (Issue #31085)
- vLLM SM121 DGX Spark Support (Issue #36821)