为便于对 Helion 内核进行性能评估与调优，需要一套标准的基准测试工具。原有 baseline 依赖 CUDA ops，在自动调优和测试中不够稳定；改用原生 torch 实现可提高可重复性和易维护性。关联 Issue #32219 和 #32962 分别定义了 Helion 集成框架和候选内核列表。
本 PR 值得精读，尤其是 benchmark_helion_kernels.py 的设计：如何选择基准测试方法（triton.do_bench vs tritonbench 的 cudagraph 实现）、如何组织基线映射。惰性注册的模式也值得在类似场景复用。风险较低，建议合入后通知相关开发者检查自定义脚本是否需要显式调用 import_all_kernels()。
无实质讨论。yushangdi 和 zou3519 直接批准，claude[bot] 自动评论无实际内容。
新增基准测试脚本 scripts/benchmark_helion_kernels.py（+457 行）：
--list 列出可用内核并提示 CUDA 基线映射是否完整。--baseline 选择 autotune（默认，使用 torch.compile 包装的 autotune_baseline_fn）或 cuda（映射至 torch.ops._C.*）。--no-cudagraph 禁用 CUDA graph，并在非 graph 模式下使用 triton.testing.do_bench，在 graph 模式下使用从 tritonbench 复制的 _do_bench_cudagraph_with_cache_clear（避免 L2 缓存不清理导致的性能虚高）。实现惰性注册：修改 vllm/kernels/helion/ops/__init__.py（+27/-5）：
pkgutil 自动遍历导入的逻辑，改为定义 import_all_kernels() 函数，由工具脚本显式调用以按需注册所有内核。vllm/kernels/helion/__init__.py，删除副作用导入 import vllm.kernels.helion.ops，避免运行时自动触发注册。更新自动调优脚本 scripts/autotune_helion_kernels.py（+4 行）：
main() 入口调用 import_all_kernels()，确保所有内核注册后再执行自动调优。为多个 Helion 内核替换 baseline 实现（涉及 silu_and_mul_per_block_quant.py、rms_norm_per_block_quant.py、rms_norm_dynamic_per_token_quant.py、per_token_group_fp8_quant.py、dynamic_per_token_scaled_fp8_quant.py）：
baseline() 从调用 torch.ops._C.* 改为完整的原生 PyTorch 前向计算，包括数据类型判断（int8/fp8）、残差融合、量化缩放等操作，并正确写出到输出 tensor。删除自动导入副作用后，现有正确性测试需 rerun 以确认无回归（PR 说明已通过）。
| 文件 | 模块 | 状态 | 重要度 |
|---|---|---|---|
scripts/benchmark_helion_kernels.py |
基准测试 | added | 8.98 |
vllm/kernels/helion/ops/__init__.py |
内核注册 | modified | 5.55 |
scripts/autotune_helion_kernels.py |
自动调优 | modified | 5.07 |
vllm/kernels/helion/ops/silu_and_mul_per_block_quant.py |
Helion 内核 | modified | 4.91 |
vllm/kernels/helion/ops/rms_norm_per_block_quant.py |
Helion 内核 | modified | 4.86 |
vllm/kernels/helion/ops/rms_norm_dynamic_per_token_quant.py |
Helion 内核 | modified | 4.77 |
scripts/benchmark_helion_kernels.py
dependency-wiring
新增基准测试脚手架，是 PR 的核心交付物。定义了基线映射、参数解析、性能收集与输出逻辑。
# scripts/benchmark_helion_kernels.py（片段）
# 将 Helion 内核名称映射到 torch.ops._C 中的 CUDA op
# 当添加新内核且需要 CUDA 基线时，在此添加条目
CUDA_BASELINE_OPS: dict[str, str] = {
"dynamic_per_token_scaled_fp8_quant": "dynamic_per_token_scaled_fp8_quant",
"fused_qk_norm_rope": "fused_qk_norm_rope",
"per_token_group_fp8_quant": "per_token_group_fp8_quant",
"rms_norm_dynamic_per_token_quant": "rms_norm_dynamic_per_token_quant",
"rms_norm_per_block_quant": "rms_norm_per_block_quant",
"silu_and_mul_per_block_quant": "silu_and_mul_per_block_quant",
"scaled_mm": "cutlass_scaled_mm",
}
# torch.compile 选项，与 vLLM 内部编译方式一致
_TORCH_COMPILE_OPTIONS: dict[str, bool] = {
"enable_auto_functionalized_v2": False,
"size_asserts": False,
"alignment_asserts": False,
"scalar_asserts": False,
"combo_kernels": True,
"benchmark_combo_kernel": True,
}
def make_cuda_baseline(kernel_name: str) -> Callable | None:
"""返回调用 CUDA op 的包装函数，若未映射则返回 None。"""
if kernel_name not in CUDA_BASELINE_OPS:
return None
cuda_op_name = CUDA_BASELINE_OPS[kernel_name]
cuda_op = getattr(torch.ops._C, cuda_op_name, None)
if cuda_op is None:
logger.warning("CUDA op %s not found", cuda_op_name)
return None
def baseline_fn(*args, **kwargs):
return cuda_op(*args, **kwargs)
return baseline_fn
def make_autotune_baseline(kernel) -> Callable | None:
"""返回 autotune_baseline_fn 经过 torch.compile 包装的函数。"""
baseline = getattr(kernel.helion_settings, "autotune_baseline_fn", None)
if baseline is None:
return None
return torch.compile(baseline, **_TORCH_COMPILE_OPTIONS)
vllm/kernels/helion/ops/__init__.py
infrastructure
实现惰性注册机制，从自动导入改为显式函数调用，是此 PR 设计决策的关键文件。
# vllm/kernels/helion/ops/__init__.py
"""
Helion 内核实现包。
导入此包**不会**自动注册内核。运行时代码应导入所需的特定 op 模块:
from vllm.kernels.helion.ops import scaled_mm # noqa: F401
需要完整注册表的工具（如 scripts/autotune_helion_kernels.py）
应显式调用 ``import_all_kernels()``。
"""
import importlib
import pkgutil
def import_all_kernels() -> list[str]:
"""导入所有内核子模块，触发每个模块的 @register_kernel 装饰器。
Returns:
已导入的完整模块名列表。
"""
imported: list[str] = []
for module_info in pkgutil.iter_modules(__path__):
if module_info.ispkg:
continue
module_name = f"{__name__}.{module_info.name}"
importlib.import_module(module_name)
imported.append(module_name)
return imported
vllm/kernels/helion/ops/silu_and_mul_per_block_quant.py
infrastructure
baseline() 从 CUDA op 替换为原生 PyTorch 实现，是 baseline 迁移的代表性文件。
# vllm/kernels/helion/ops/silu_and_mul_per_block_quant.py 中的 baseline 片段
def baseline(
out: torch.Tensor,
input: torch.Tensor,
scales: torch.Tensor,
group_size: int,
scale_ub: torch.Tensor | None = None,
is_scale_transposed: bool = False,
) -> None:
# 改用原生 PyTorch 实现，不再调用 torch.ops._C.silu_and_mul_per_block_quant
num_tokens, intermediate_size = out.shape
groups_per_row = intermediate_size // group_size
quant_dtype = out.dtype
if quant_dtype == torch.int8:
qtype_min, qtype_max = get_int8_min_max()
min_scaling_factor = get_int8_min_scaling_factor()
else: # fp8
qtype_min, qtype_max = get_fp8_min_max()
min_scaling_factor = 1.0 / (qtype_max * 512.0)
# SiLU 激活 → 分组求绝对最大值 → 计算 scale → 量化
act = SiluAndMul.forward_native(input.to(torch.float32))
x_grouped = act.view(num_tokens, groups_per_row, group_size)
s = torch.amax(torch.abs(x_grouped), dim=-1)
if scale_ub is not None:
s = s.clamp(max=scale_ub)
s = (s * (1.0 / qtype_max)).clamp(min=min_scaling_factor)
y = x_grouped / s[:, :, None]
if quant_dtype == torch.int8:
y = y.round()
scales.copy_(s)
out.copy_(
y.clamp(qtype_min, qtype_max).view(num_tokens, intermediate_size).to(out.dtype)
)
当前评论区没有形成足够清晰的争议点或结论，后续有更多讨论时会体现在这里。
import vllm.kernels.helion.ops 时的自动注册，则会因注册失效而崩溃。但该 PR 同时更新了 autotune_helion_kernels.py 显式调用 import_all_kernels()；需要确保所有需注册的场景都做了相应调整。开发者：获得统一的 Helion 内核性能评估工具，可标准化未来内核优化的验证流程。惰性注册减少了非 Helion 场景的无关导入开销。
系统：运行时行为无变化，因基准测试脚本仅作为独立工具使用。自动调优和测试流程依赖新的注册方式，需注意调用顺序。
团队：内核开发流程更加清晰（编写内核 → 注册 → 基准测试 → 自动调优）。需要维护 CUDA_BASELINE_OPS 映射以确保基线完整性。
参与讨论