已关闭
[Requirement|需求建议]: 新增SIMD与SIMT混合编程Softmax性能调优样例 #1648
liuyi424创建于 9月9日关闭于 16 天前
liuyi424
9月9日 评论:
9月9日 评论:
该需求已通过 MR 实现:https://gitcode.com/cann/asc-devkit/merge_requests/5793 ,合入后本issue可关闭。


9月10日 issue类型由 任务 改变为 需求
16 天前 添加了label:resolved
该需求已通过 MR 实现:https://gitcode.com/cann/asc-devkit/merge_requests/5793 ,合入后本issue可关闭。


Background(背景信息)
当前
examples/05_simd_simt_hybrid/02_best_practices下的混编样例覆盖了访存合并、gridDim配置、MTE任务队列、整数快速除法等场景,缺少行内规约类算子的调优样例。row-wise softmax是典型的行内规约算子:每行内部需要求max和sum两类规约,行与行之间彼此独立,且实际网络中大量出现最后一维不超过512的形状。本需求新增
simd_simt_softmax样例,展示行内规约场景下从朴素实现到编译期模板特化的完整调优路径:Origin(信息来源)
无
Benefit / Necessity (价值/作用)
__launch_bounds__的配合,说明最大线程数与单线程可用寄存器数量的关系。Design(设计方案)
样例目录为
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。三个Case通过CMake编译选项
SCENARIO_NUM(0/1/2,默认2)选择,未选中的Case不参与编译。kernel主体为
__global__ __vector__核函数:MTE流水负责GM与UB之间的数据搬运,asc_vf_call调起SIMT VF完成行内计算,事件机制管理双缓冲生命周期。Host侧根据数据量动态计算线程块数量(AIV核数与rows取小)与bucket路径线程数(512/1024/2048按shape选择,cols大于256时考虑寄存器压力收敛为1024)。
同步更新
05_simd_simt_hybrid上层README与02_best_practices样例索引(中英文)。