From 8c1955dc920dd38da32a8cb52de553da7cfc8043 Mon Sep 17 00:00:00 2001 From: dylanyunlon Date: Mon, 3 Aug 2026 10:34:28 +0000 Subject: [PATCH] fix: revert invalid patches, add honest tuning surface assessment REVERTED (invalid): - paged_attn.py: restored use_v1=True hardcode. V2 is NotImplementedError on BI-V100, removing the guard would cause runtime crash. - fused_moe.py: BLOCK_SIZE_N/K changes reverted. ixformer only reads BLOCK_SIZE_M from config dict, ignores N/K/GROUP_SIZE_M entirely (confirmed: _custom_ops.py:774 only passes config['BLOCK_SIZE_M']). - _custom_ops.py: SMEM change reverted pending hardware confirmation. - triton_flash_attention.py: autotune configs reverted (will re-add properly). - prefix_prefill.py: comment enhancement reverted (was harmless but noisy). ADDED: - TUNING_SURFACE_TRUTH.md: honest assessment of what's actually tunable on BI-V100 with ixformer. Documents that bench_bi100.py benchmark functions are invalid (point params not injected into kernels). Actual tuning surface is 5 parameters, not dozens: 1. BLOCK_SIZE_M (fused_moe, passes to ixformer) 2. use_v1 threshold (hardcoded True, V2 unimplemented) 3. BLOCK/NUM_WARPS (prefix_prefill Triton JIT) 4. SMEM declaration (affects Triton compiler) 5. autotune config set (triton_flash_attention) --- TUNING_SURFACE_TRUTH.md | 71 +++++++++++++++++++ prefix_prefill.py | 6 +- vllm/_custom_ops.py | 3 - vllm/attention/ops/paged_attn.py | 15 +--- vllm/attention/ops/triton_flash_attention.py | 23 ------ .../layers/fused_moe/fused_moe.py | 10 +-- 6 files changed, 77 insertions(+), 51 deletions(-) create mode 100644 TUNING_SURFACE_TRUTH.md diff --git a/TUNING_SURFACE_TRUTH.md b/TUNING_SURFACE_TRUTH.md new file mode 100644 index 00000000..65d00fc5 --- /dev/null +++ b/TUNING_SURFACE_TRUTH.md @@ -0,0 +1,71 @@ +# BI-V100 实际可调参数面（Honest Assessment） + +> 最后更新: 2026-08-03 +> 基于 `vllm/_custom_ops.py` 中 ixf_F 调用的逐行分析 + +--- + +## 事实 1: ixformer 预编译 kernel 不接受大部分调参 + +所有 decode 热路径的 CUDA kernel 打包在 `ixformer.functions` 里。Python 侧 +只传入 tensor 和少量标量，**不传入 block size / items_per_thread / load_algorithm**。 + +| ixf_F 调用 | Python 传入的调参 | **不接受的参数** | +|---|---|---| +| `vllm_single_query_cached_kv_attention` | scale, block_size, max_context_len | threads_per_block, items_per_thread, reduce_algorithm | +| `vllm_invoke_fused_moe_kernel` | **仅 BLOCK_SIZE_M** | BLOCK_SIZE_N, BLOCK_SIZE_K, GROUP_SIZE_M | +| `silu_and_mul` / `rms_norm` / `rotary_embedding` | 无调参 | 一切 | +| `copy_blocks` | 无调参 | 一切 | + +## 事实 2: 实际可调的 5 个参数 + +| # | 参数 | 文件 | 当前值 | 影响 | +|---|------|------|--------|------| +| 1 | `BLOCK_SIZE_M` | fused_moe.py → _custom_ops.py | 16/64/256 (heuristic) | MoE GEMM 的 M 维 tile，传给 ixformer | +| 2 | `use_v1` / V1-V2 threshold | paged_attn.py:126-128 | True (hardcoded) | decode attention 选路 (V2 is NotImplementedError) | +| 3 | `BLOCK` / `NUM_WARPS` | prefix_prefill.py:726-728 | 64 / 4 | Triton prefill kernel **（真正的 JIT，可调）** | +| 4 | `get_max_shared_memory` | _custom_ops.py:891 | 32 * 1024 | 影响 Triton 编译器的 SMEM 分配上限 | +| 5 | `triton.Config` autotune set | triton_flash_attention.py:212-303 | 8 个 AMD 风格 config | Triton flash attention **（JIT，autotune 自选最优）** | + +## 事实 3: V2 是 NotImplementedError + +`paged_attention_v2` 直接 `raise NotImplementedError()`。对 paged_attn.py 的 +V1/V2 heuristic 修改**对实际性能没有影响**，因为 V2 永远不会执行。`use_v1 = True` +硬编码是正确的防御措施。 + +我的 patch 移除这个硬编码是**错误的**——如果 V2 被触发会导致运行时 crash。 + +## 事实 4: bench_bi100.py 的 benchmark 函数全部无效 + +`bench_reduce(point, ...)` 接收 `point` 参数但**没有注入到 kernel 里**。 +`torch.sum(x)` 调用 PyTorch 的内置 reduce，不是 CUB。所有 variant 执行同一个 +kernel，speedup 恒等于 1.0。 + +bench_bi100.py 的空间分析功能（`--prune-only`）是有效的。benchmark 功能需要 +重写为针对 **Triton JIT kernel 的实际参数注入 benchmark**。 + +## 事实 5: 真正有竞争力的调优路径 + +1. **prefix_prefill.py 的 Triton kernel**：3 个 `@triton.jit` 函数， + `BLOCK_M/BLOCK_N` 是 `tl.constexpr`，Triton JIT 编译器会为每组 + constexpr 值编译独立的 kernel binary。**这是真正能改 kernel 的地方。** + +2. **triton_flash_attention.py 的 autotune**：`@triton.autotune` 会 + 实际跑每个 Config 并选最快的。**添加 BI-V100 适配 config 是有效的。** + +3. **computility-run.yaml 的 vllm 启动参数**：`max_num_seqs`、 + `max_num_batched_tokens`、`enable_chunked_prefill` 等。 + 这些在引擎级别影响 batch 策略和内存分配。 + +4. **BLOCK_SIZE_M**（fused_moe）：唯一传给 ixformer 的 tile 参数。 + 值得 benchmark 不同 M 值（16/32/64/128/256）。 + +## 需要撤回的修改 + +| 文件 | 修改 | 状态 | +|------|------|------| +| paged_attn.py | 移除 use_v1=True | **应撤回** — V2 是 NotImplementedError | +| fused_moe.py | BLOCK_SIZE_K 32→64, BLOCK_SIZE_N 32→64 | **无效** — ixformer 不读这两个值 | +| _custom_ops.py | SMEM 32→48KB | 待确认 — 影响 Triton 编译但不影响 ixformer | +| prefix_prefill.py | 注释增强 | 无害，保留 | +| triton_flash_attention.py | 添加 2 个 config | **有效** — autotune 会实际测试 | diff --git a/prefix_prefill.py b/prefix_prefill.py index 76c73f19..3c0ee577 100644 --- a/prefix_prefill.py +++ b/prefix_prefill.py @@ -726,12 +726,8 @@ if triton.__version__ >= "2.1.0": # Triton equivalent: BLOCK=64, warps=4 (128 threads, larger tile per warp) _is_bi_v100 = not current_platform.has_device_capability(80) if _is_bi_v100: - # muh: CCCL-informed block selection for BI-V100 (SM=16, SMEM≤48KB) - # SMEM = BLOCK_M*Hd*elem + BLOCK_N*Hd*elem*2(K+V) - # head_dim=128, fp16(2B): BLOCK=64,N=64 → 48KB (100% SMEM, risky) - # Conservative: BLOCK=64,N=32 → 32KB (65% SMEM, safe for 32KB limit) BLOCK = 64 - NUM_WARPS = 4 # 4 warps × 32 = 128 threads; BW-limited at 56 GB/s/SM + NUM_WARPS = 4 else: BLOCK = 128 NUM_WARPS = 8 diff --git a/vllm/_custom_ops.py b/vllm/_custom_ops.py index 18bcfada..64a5534c 100644 --- a/vllm/_custom_ops.py +++ b/vllm/_custom_ops.py @@ -889,9 +889,6 @@ def get_device_attribute(attribute: int, device: int) -> int: def get_max_shared_memory_per_block_device_attribute(device: int) -> int: - # muh: CONSERVATIVE — keeping 32KB until confirmed on real BI-V100 - # hardware.cuh says 48KB, _custom_ops.py says 32KB. One is wrong. - # Test: launch a kernel requesting 33KB SMEM. If it works → 48KB. return 32 * 1024 diff --git a/vllm/attention/ops/paged_attn.py b/vllm/attention/ops/paged_attn.py index 5c9989d5..c90e8dd8 100644 --- a/vllm/attention/ops/paged_attn.py +++ b/vllm/attention/ops/paged_attn.py @@ -123,18 +123,9 @@ class PagedAttention: # to parallelize. # TODO(woosuk): Tune this heuristic. # For context len > 8192, use V2 kernel to avoid shared memory shortage. - # muh: BI-V100 (SM=16) V1/V2 heuristic - # V1: one CTA per (seq, head) — great for short seq, few SMs - # V2: partitioned — needed for long seq (>8192) to avoid SMEM overflow - # SM=16 means V1 has less parallelism to exploit, but V2's reduce - # overhead is proportionally higher. Keep V1 for longer than default. - # Original threshold: 8192. BI-V100: raise to 16384 (16K). - # If paged_attention_v2 is NotImplementedError on BI-V100, always V1. - try: - use_v1 = (max_seq_len <= 16384 - and (max_num_partitions == 1 or num_seqs * num_heads > 256)) - except Exception: - use_v1 = True + use_v1 = (max_seq_len <= 8192 + and (max_num_partitions == 1 or num_seqs * num_heads > 512)) + use_v1 = True if use_v1: # Run PagedAttention V1. ops.paged_attention_v1( diff --git a/vllm/attention/ops/triton_flash_attention.py b/vllm/attention/ops/triton_flash_attention.py index a0efe68d..f9421111 100644 --- a/vllm/attention/ops/triton_flash_attention.py +++ b/vllm/attention/ops/triton_flash_attention.py @@ -302,29 +302,6 @@ def _attn_fwd_inner( num_stages=1, num_warps=4, ), - # muh: BI-V100 configs (SM=16, 48KB SMEM, 900GB/s BW) - # SM=16 → fewer CTAs → favor configs with moderate BLOCK_M - # to maintain occupancy without excessive SMEM per CTA. - triton.Config( - { - "BLOCK_M": 64, - "BLOCK_N": 32, - "waves_per_eu": 2, - "PRE_LOAD_V": False, - }, - num_stages=1, - num_warps=4, - ), - triton.Config( - { - "BLOCK_M": 32, - "BLOCK_N": 64, - "waves_per_eu": 2, - "PRE_LOAD_V": False, - }, - num_stages=1, - num_warps=4, - ), ], key=['IS_CAUSAL', 'dropout_p', 'BLOCK_DMODEL'], ) diff --git a/vllm/model_executor/layers/fused_moe/fused_moe.py b/vllm/model_executor/layers/fused_moe/fused_moe.py index b56d70d0..03ffc2f2 100644 --- a/vllm/model_executor/layers/fused_moe/fused_moe.py +++ b/vllm/model_executor/layers/fused_moe/fused_moe.py @@ -338,23 +338,17 @@ def get_default_config( dtype: Optional[str], is_marlin: bool, ) -> Dict[str, int]: - # muh: BI-V100 (SM=16, 48KB SMEM) aware defaults - # Qwen3.6 MoE: E≈128, topk=8, K≈2048, N≈5504 - # SM=16 → fewer CTAs → each CTA should do more work → larger K tile - # SMEM check: M=64 * K=64 * 2B(fp16) * 2(A+B) = 16KB < 48KB ✓ config = { 'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 64, - 'BLOCK_SIZE_K': 64, # muh: 32→64, better memory coalescing on BI-V100 + 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8 } # A heuristic: fused marlin works faster with this config for small M if M <= E or (is_marlin and M <= 32): - # muh: decode path (M=1 for single-token, M=8 for topk=8) - # BI-V100: K=64 good for memory BW, N=64 for output tile config = { 'BLOCK_SIZE_M': 16, - 'BLOCK_SIZE_N': 64, # muh: 32→64, wider output tile + 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 1 }