Feat/fused silu quant integration

[JackTan25]

No description provided.

[LLLLKKKK]

🤖 AI Code Review — PR #816 — Feat/fused silu quant integration

概述

将 DeepGEMM MoE contiguous executor 中的 SiLU 激活 + FP8 量化从两步（silu_and_mul → sgl_per_token_group_quant_fp8/trt_fp8_quantize_128）融合为单个 Triton kernel silu_and_mul_contiguous_fp8_post_quant，消除中间 bf16 buffer 分配，减少显存占用和 kernel launch 开销。

优点

• 方向正确的性能优化：kernel fusion 消除了 (all_tokens, N//2) 大小的中间 bf16 buffer，对大 batch MoE 场景显存节省可观
• 新 Triton kernel 结构清晰，支持 UE8M0 和标准 FP8 两种量化路径，支持 column-major scale layout 和 TMA alignment
• 代码格式化改进（长行拆分）提升了可读性
• del 语句位置调整正确——在不再需要时尽早释放 tensor

建议改进

P1 - 重要

1. 缺少单元测试
   silu_and_mul_contiguous_fp8_post_quant 是包含 SiLU + FP8 量化 + UE8M0 rounding 的复合 kernel，目前无任何测试覆盖。建议添加测试：

   • 基本正确性：对比 silu_and_mul + sgl_per_token_group_quant_fp8 两步结果
   • UE8M0 路径 vs 非 UE8M0 路径
   • column-major scale layout + TMA alignment
   • 边界 case：all_tokens=1、非 4 对齐的 token 数

2. UE8M0 scale 格式兼容性需确认
   旧代码 sgl_per_token_group_quant_fp8(..., scale_ue8m0=True) 内部会将 scale 打包为 int32 UE8M0 格式。新 kernel 返回 float32 scale（仅做 power-of-2 rounding），注释说 "UE8M0 packing is handled externally"，但 executor 中没有显式 packing 步骤，直接将 float32 scale 传给 m_grouped_fp8_gemm_nt_contiguous。请确认 GEMM 在 disable_ue8m0_cast=False 时能否正确处理 float32 power-of-2 scale。如果 GEMM 期望 packed int32，这里会产生静默数值错误。

P2 - 建议

3. up 变量未显式 cast 到 float32
   kernel 中 gate 被 .to(tl.float32) 但 up 保持 bf16。虽然 Triton 会自动 promote，但显式 cast 更清晰且与同文件 _silu_mul_fp8_quant_deep_gemm_masked kernel 风格一致。建议：

   up = tl.load(input_ptr + up_offset, mask=mask, other=0.0).to(tl.float32)

4. PR description 为空
   对于性能优化 PR，建议补充优化动机、性能基准数据、适用场景说明。

P3 - Nit

5. num_warps=1 硬编码 — 建议添加注释说明选择依据（group_size=128 时合理）。

6. BLOCK=group_size 时 mask 恒为 True — 如果 BLOCK 永远等于 GROUP_SIZE，mask 逻辑冗余。如果设计上允许非 2 的幂 group_size，wrapper 应改为 BLOCK=triton.next_power_of_2(group_size)。

总结

方向正确的 kernel fusion 优化。核心风险在于 UE8M0 scale 格式兼容性（#2）——如果 GEMM 期望 packed int32 而收到 float32，会产生静默数值错误。建议确认兼容性并补充单元测试后合入。

[LLLLKKKK]

🤖 AI Code Review — PR #816 — Feat/fused silu quant integration

概述

将 DeepGEMM MoE contiguous executor 中的 SiLU 激活 + FP8 量化从两步（silu_and_mul + sgl_per_token_group_quant_fp8/trt_fp8_quantize_128）融合为单个 Triton kernel silu_and_mul_contiguous_fp8_post_quant，消除中间 bf16 buffer 分配，减少显存占用和 kernel launch 开销。

优点

• Kernel fusion 方向正确，消除了 (all_tokens, N//2) 的 bf16 中间 buffer，对大 batch 场景显存节省可观
• 新 kernel 同时支持 UE8M0 和普通 float32 scale 两种模式，与现有两条路径对齐
• column-major scale layout + TMA alignment 处理与 DeepGEMM 兼容
• del gateup_output 位置调整合理，尽早释放显存

建议改进

P1 - 重要

1. UE8M0 scale 格式兼容性需确认
   旧代码 UE8M0 路径使用 sgl_per_token_group_quant_fp8(..., scale_ue8m0=True)，该函数内部通过 CUDA kernel 将 scale 打包为 int32 UE8M0 格式（4 个 8-bit exponent 打包到一个 int32）。新 kernel 在 use_ue8m0=True 时仅做 power-of-2 rounding，但仍以 float32 存储 scale，注释说 "UE8M0 packing (int32) is handled externally by the caller"。然而 executor 中没有显式 packing 步骤，直接将 float32 scale 传给 m_grouped_fp8_gemm_nt_contiguous。需要确认 GEMM 在 disable_ue8m0_cast=False 时是否能正确处理 float32 power-of-2 scale（而非已打包的 int32）。如果 GEMM 期望 packed int32，这里会产生静默数值错误。

2. 缺少单元测试
   新增 ~140 行 Triton kernel，逻辑非平凡（SiLU + FP8 quant + UE8M0 rounding），但没有对应测试。建议至少覆盖：

   • 基本正确性：对比 fused kernel 与原始两步流程的输出
   • UE8M0 路径 vs 非 UE8M0 路径
   • column-major scale layout + TMA alignment
   • 边界 case：all_tokens=1、all_tokens 非 4 对齐

P2 - 建议

3. up 变量未显式 cast 到 float32
   kernel 中 gate 被显式 .to(tl.float32) 但 up 保持原始 dtype（bf16）。虽然 Triton 会在乘法时自动 promote，但显式 cast 更清晰且与同文件中 _silu_mul_fp8_quant_deep_gemm_masked kernel 的风格一致。建议统一：

   up = tl.load(input_ptr + up_offset, mask=mask, other=0.0).to(tl.float32)

4. PR description 为空
   没有描述动机、性能数据或设计决策。建议补充性能基准数据（fused vs unfused 延迟对比）和适用场景说明。

P3 - Nit

5. num_warps=1 硬编码 — 对 group_size=128 合理，但建议添加注释说明选择依据。

6. BLOCK=group_size 时 mask 恒为 True — 如果 BLOCK 永远等于 GROUP_SIZE，mask 逻辑冗余。如果设计上允许 BLOCK > GROUP_SIZE，wrapper 应改为 BLOCK=triton.next_power_of_2(group_size)。

总结

方向正确的性能优化。核心风险在于 UE8M0 scale 格式兼容性（问题 #1）——如果 GEMM 期望 packed int32 而收到 float32，会产生静默数值错误。建议确认兼容性并补充单元测试后合入。

[LLLLKKKK]

🤖 AI Code Review — PR #816
Head SHA: c01aa50138f78d495fce398f93055a5d84a7f2be | Verdict: P2

Summary

Introduces a fused SiLU+FP8 quantization Triton kernel for the MoE contiguous DeepGemm executor, eliminating the intermediate bf16 buffer between activation and quantization. Supports both UE8M0 and standard FP8 paths.

Findings

[P2] Kernel uses num_warps=1 unconditionally
With BLOCK=group_size (typically 128), a single warp handles 128 elements (4 per thread). Consider making num_warps adaptive or adding a comment explaining why 1 warp is sufficient.

[P2] No numerical validation test for the fused kernel
The new fused path replaces two separate well-tested operations. A regression test comparing fused vs unfused outputs would catch precision issues.

[Nit] The gate/up layout assumption ([up | gate]) should be documented in the kernel.

Good memory optimization — eliminating the (all_tokens, N//2) bf16 intermediate is meaningful for large MoE models.