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

@@ -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 会实际测试 |

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

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
 
 

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
+        use_v1 = (max_seq_len <= 8192
-        # V1: one CTA per (seq, head) — great for short seq, few SMs
+                  and (max_num_partitions == 1 or num_seqs * num_heads > 512))
-        # V2: partitioned — needed for long seq (>8192) to avoid SMEM overflow
+        use_v1 = True
-        # 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
         if use_v1:
             # Run PagedAttention V1.
             ops.paged_attention_v1(

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'],
 )

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
         }