已关闭
[Requirement|需求建议]: 新增SIMD与SIMT混合编程Softmax性能调优样例 #1648
liuyi424创建于  9月9日关闭于  16 天前
liuyi424成员
9月9日 创建

Background(背景信息)

当前 examples/05_simd_simt_hybrid/02_best_practices 下的混编样例覆盖了访存合并、gridDim配置、MTE任务队列、整数快速除法等场景,缺少行内规约类算子的调优样例。row-wise softmax是典型的行内规约算子:每行内部需要求max和sum两类规约,行与行之间彼此独立,且实际网络中大量出现最后一维不超过512的形状。

本需求新增 simd_simt_softmax 样例,展示行内规约场景下从朴素实现到编译期模板特化的完整调优路径:

  • Case 0:SIMT VF直接读GM基线,一个Warp处理一行,行内max/sum通过Warp规约完成;
  • Case 1:MTE搬入UB中转,双份tile缓冲区轮换,MTE2搬入、SIMT VF计算与MTE3搬出流水重叠;
  • Case 2:按cols上界(64/128/256/384/512)生成编译期bucket模板kernel,定长数组中间值驻留寄存器,Host侧按shape动态选择线程数。

Origin(信息来源)

无

Benefit / Necessity (价值/作用)

  1. 补齐混编样例在行内规约场景的空白,与已有访存、执行配置、指令优化类样例形成互补,覆盖“一个Warp处理一行”的典型线程组织方式。
  2. 展示编译期维度特化(bucket模板)与运行时参数的权衡:规约维度较小时把维度上界固化到编译期,定长数组中间值可驻留寄存器,避免重复读输入和重复计算exp。
  3. 展示Host侧按shape动态选择线程数与SIMT VF __launch_bounds__ 的配合,说明最大线程数与单线程可用寄存器数量的关系。
  4. 样例含中英文README、上板实测性能数据(Case 0/1/2 Task Duration分别为16.516/14.670/10.491μs)与msOpProf性能调试说明,便于用户复现和学习。

Design(设计方案)

  1. 样例目录为 examples/05_simd_simt_hybrid/02_best_practices/simd_simt_softmax,包含:

    • softmax_gm.h:Case 0,SIMT直接读GM基线实现;
    • softmax_ub.h:Case 1,UB中转与MTE双缓冲实现;
    • softmax_bucket.h:Case 2,编译期bucket模板主路径;
    • softmax_host.asc:Host侧运行入口,参照reduce_sum样例风格实现(CHECK_ACL宏、单一职责函数分解);
    • scripts/gen_data.py、scripts/verify_result.py:数据生成与精度校验,精度阈值atol=rtol=1e-3。
  2. 三个Case通过CMake编译选项 SCENARIO_NUM(0/1/2,默认2)选择,未选中的Case不参与编译。

  3. kernel主体为 __global__ __vector__ 核函数:MTE流水负责GM与UB之间的数据搬运,asc_vf_call 调起SIMT VF完成行内计算,事件机制管理双缓冲生命周期。

  4. Host侧根据数据量动态计算线程块数量(AIV核数与rows取小)与bucket路径线程数(512/1024/2048按shape选择,cols大于256时考虑寄存器压力收敛为1024)。

  5. 同步更新 05_simd_simt_hybrid 上层README与 02_best_practices 样例索引(中英文)。

likedislike
Lliuyi424成员
9月9日 添加了label:requirement
liuyi424成员
9月9日 评论:

该需求已通过 MR 实现:https://gitcode.com/cann/asc-devkit/merge_requests/5793 ,合入后本issue可关闭。

likedislike
Lliuyi424成员
9月9日 将 liuyi424 设为负责人
shaonaiteshaonaite成员
9月10日 issue类型由 任务 改变为 需求
Lliuyi424成员
16 天前 issue状态由 待办的 改变为 已确认
Lliuyi424成员
16 天前 issue状态由 已确认 改变为 已解决
Lliuyi424成员
16 天前 关闭了 issue
CANN-robotCANN-robot成员
16 天前 添加了label:resolved