wenyi-li 大规模 Mean 归约优化（行二次切分）
任务特征
- 数据尺寸：(1000, 8192)，非reduce轴中等，reduce轴较大
优化：行二次切分
pid = tl.program_id(0)
for m_start in range(0, BLOCK_SIZE_M, SUB_BLOCK_SIZE_M):
m_offsets = pid * BLOCK_SIZE_M + m_start + tl.arange(0, SUB_BLOCK_SIZE_M)
目的：
- 每个kernel计算多行（BLOCK_SIZE_M），减少总线程块数量
- Kernel内对行进行二次切分（SUB_BLOCK_SIZE_M），避免超出硬件缓存
Autotune 配置
# （AI core=40）
# 1. grid<40 -> 28.64 us
triton.Config({'BLOCK_SIZE_M': 50, 'SUB_BLOCK_SIZE_M': 25, 'BLOCK_SIZE_N': 512})
# 2. grid=40，SUB切分含尾块 -> 16.54 us
triton.Config({'BLOCK_SIZE_M': 25, 'SUB_BLOCK_SIZE_M': 4, 'BLOCK_SIZE_N': 4096})
# 3. grid=40，SUB切分不含尾块 -> 16.00 us 最优
triton.Config({'BLOCK_SIZE_M': 25, 'SUB_BLOCK_SIZE_M': 25, 'BLOCK_SIZE_N': 512})
# 4. grid>40，且非核数整数倍 -> 25.86 us
triton.Config({'BLOCK_SIZE_M': 20, 'SUB_BLOCK_SIZE_M': 20, 'BLOCK_SIZE_N': 512})
总结
- grid等于核数，SUB切分不含尾块时性能最优
- 尾块计算会降低性能
- grid超出核数且非核数整数倍时，各核计算任务不均匀，性能较差