已合并
【fix】: Asin/Acos/Remainder修改为Simt接口 #2045
【fix】: Asin/Acos/Remainder修改为Simt接口 #2045
已合并
gaoxin创建于 15 天前
gaoxin成员
15 天前

Pull Request

描述

将 Autofuse V2 的 Asin/Acos 与 Remainder(int32) 算子执行路径切换为 SIMT 接口实现,与既有 Sin/Cos/Expm1 的 SIMT 机制对齐:SIMT kernel(__simt_vf__ + LAUNCH_BOUND(1024),线程按 threadIdx.x 步进 blockDim.x 逐元素计算)通过 AscendC::Simt::VF_CALL 启动,并配套 40 KiB SIMT dcache 预留。外部 ASCIR 类型、Python API、ATT 注册、tmp buffer 计算和 ApiCall 调用 ABI 均保持不变。

问题

  1. Asin/Acos:此前直接调用 AscendC 高阶 API(adv_api/math/asin.h/acos.h),性能比单算子差。
  2. Remainder(int32):原实现基于 AscendC::Reg:: 向量指令,性能比单算子差。

修改方案

  1. 新增 asin.h/acos.h regbase:AsinSimtCompute/AcosSimtCompute(调用 Simt::Asin/Simt::Acos)+ AsinExtend/AcosExtend wrapper;CMake 生成 raw string 并在 ascendc_reg_base_api_register.cpp 注册;Codegen 类 GetApiName() 切换为 AsinExtend/AcosExtend,LoadApiHeaderFiles() 加载 asin_reg_base.h/acos_reg_base.h,IncludeApiHeaderFiles() 增加 SIMT 官方头;GetConversionDtype() 增加 FP16→FLOAT 转换(与 BF16→FLOAT 一致)。
  2. remainder.h int32 路径替换为 SIMT:新增 RemainderIntSimtCompute(参考 ops-math 仓 FloorModInt_1:% 取余 + 符号差异修正 rem + src2,保证 floor 语义),通过 AscendC::Simt::VF_CALL 启动;删除 AscendC::Reg:: 向量实现 RemainderIntImplVF;Codegen IncludeApiHeaderFiles() 增加 simt_api/cpp/kernel_simt_intf.h。float 路径保持不变。
  3. 图信息补全(ascgraph_info_complete.cpp):CompleteElewiseApiInfo 为 Asin/Acos 及 int32 输入的 Remainder 设置 40 KiB SIMT dcache 预留;kOpTypeToComputeType 补全 Asin/Acos/Remainder 的 kComputeElewise 条目。
  4. 测试断言增强:asin/acos backend generator 断言 AsinExtend/AcosExtend、*SimtCompute、SIMT 头文件和 tiling dcache(40960);remainder int32 generator 同步增加 RemainderIntSimtCompute 与 dcache 断言。

变更类型

关联的Issue

如何测试

  1. cmake --build build --target aihac_codegen -j 8 编译通过。
  2. 运行 asin_bf16_test_codegen_v2 / acos_bf16_test_codegen_v2:断言 Extend/SimtCompute/SIMT 头/tiling dcache 全部通过。
  3. 运行 remainder_int32_store_test_e2e_v2:ICPU(ascend950 模拟)3/3 通过,覆盖对齐/非对齐/较大 shape,输入含正负除数符号修正场景,golden 独立计算 floor 语义。
  4. 回归:optimize_ut 651/651、test_main 878/878(0 failed)、regbase UT RemainderFloor_Test 通过;clang-format(v18.1.8 pre-commit 与 v21.1.4 双版本)检查通过。

核对清单

其他信息

likedislike
Pull Request已成功合入, 合并人@CANN-robot
(感谢 gaoxin 的贡献)
Ggaoxin成员
15 天前 创建了 pull request,commit 39c6a02a
atomgit-bot
atomgit-bot
15 天前 评论:

变更摘要

本 PR 将 Asin/Acos 算子从原实现修改为 SIMT 接口(参考既有 Sin/Cos 实现):新增 asin.h/acos.h regbase 头文件并注册 raw string,Codegen 侧切换 GetApiName 与 Load/Include 头文件,GetConversionDtype 增加 FP16→FLOAT 映射,同时在节点补全阶段为 Asin/Acos 设置 40 KiB SIMT dcache 预留并补全 ComputeType,最后在 bf16 backend E2E 测试中新增 kernel 文本断言。该变更保持 ASCIR 类型、Python API、ATT 注册、tmp buffer 计算和 UnaryApiTmpCall 调用 ABI 不变。

主要改动

  • 新增 SIMT regbase 头文件并注册: 新增 autofuse/v35/ascendc/api_regbase/asin.h、acos.h,定义 SIMT 核函数 AsinSimtCompute/AcosSimtCompute 及封装 AscendC::Simt::VF_CALL 调用的 AsinExtend/AcosExtend(线程数 ASIN_THREAD_NUM/ACOS_THREAD_NUM = 1024),并在 ascendc_reg_base_api_register.cpp 中将 asin_reg_base.h/acos_reg_base.h 加入 api_to_file 映射完成注册。
  • Codegen 切换接口与头文件: AsinAscIrCodegenImplV2/AcosAscIrCodegenImplV2 的 GetApiName 由 "Asin"/"Acos" 改为 "AsinExtend"/"AcosExtend",新增 LoadApiHeaderFiles 返回 asin_reg_base.h/acos_reg_base.h,并在 IncludeApiHeaderFiles 中补充 simt_api/cpp/kernel_simt_intf.h 与 simt_api/math_functions.h。
  • float16 转 float 处理: 两个算子的 GetConversionDtype 转换映射在原有 {DT_BF16, DT_FLOAT} 基础上新增 {DT_FLOAT16, DT_FLOAT},使 float16 输入统一转换为 float 处理。
  • 节点 SIMT dcache 与 ComputeType 补全: CompleteElewiseApiInfo 将 Asin/Acos 纳入设置 kSimtDcacheSize(40 KiB)的条件,并在 kOpTypeToComputeType 映射表中为 Asin::Type/Acos::Type 补充 kComputeElewise。
  • E2E 测试断言增强: asin_bf16_backend_generate.cpp 与 acos_bf16_backend_generator.cpp 新增对生成的 kernel 中 AsinExtend/AcosExtend、AsinSimtCompute/AcosSimtCompute、simt_api/cpp/kernel_simt_intf.h 以及 tiling 中 dcache 大小(40 * 1024)的文本断言。
likedislike
不准确?
atomgit-bot
atomgit-bot
15 天前 评论:

代码审查

✅ 未发现问题

likedislike
不准确?
CANN-robotCANN-robot成员
15 天前 添加了label:cann-cla/yes
CANN-robot
CANN-robot成员
15 天前 评论:

Thanks for your pull-request.
The full list of commands accepted by me can be found at here.
You can get sig-info at here.
You can self-configure the PR merge rules for this repository. For more details, please refer to Here.
For more, you also can visit HICANN.


PR Approval Progress

✅ Congratulations! All modules have met the lgtm and approve requirements.

Module Approval Details

module lgtm status approve status
repo-cann/graph-autofusion ✅ yangyongqiang0606, zhang_shengjie, liyuewei (3/2) ✅ yangyongqiang0606 (1/1)

💡 Tip:

  • Committer can comment /approve or /lgtm
  • Commenting /approve implies both code review (lgtm) and intent to merge (approve)

CLA Signature Pass

gcw_V3YyYBt1, thanks for your pull request. All authors of the commits have signed the CLA. 👍

likedislike
CANN-robotCANN-robot成员
15 天前 将niuyuhu,yangyongqiang0606,liyuewei,xchu42,zhang_shengjie,depeng1994,wqtshg_wt,peiyang,shengnan666,XuebinYang,sjtulxh,zhanj,zhujingjing,xuyafei,wangxiaotian995设为评审人
此处折叠了6条事件消息 查看更多
CANN-robotCANN-robot成员
15 天前 添加了label:cann-cla/yes
gaoxin成员
15 天前 评论:

compile

likedislike
Ggaoxin成员
15 天前 预合并成功(commit_id: 4027603168d2d9ec4b2e5daa2b8407e21cd11d2a)
CANN-robot
CANN-robot成员
15 天前 评论:
🚀 CI 流水线已启动
📋 执行详情: 点击查看流水线
likedislike
CANN-robotCANN-robot成员
15 天前 添加了label:ci-pipeline-running
此处折叠了9条事件消息 查看更多
CANN-robotCANN-robot成员
14 天前 删除了label:ci-pipeline-passedstat/needs-squash
CANN-robot
CANN-robot成员
14 天前 评论:

Notification

This pull request has been changed(code update) or closed, so removes the following label(s): ci-pipeline-passed, stat/needs-squash.

likedislike
CANN-robotCANN-robot成员
14 天前 添加了label:cann-cla/yes
gaoxin成员
14 天前 评论:

compile

likedislike
Ggaoxin成员
14 天前 预合并成功(commit_id: 1d5247734c05e1a79e17561359e17a6b25d62dcc)
CANN-robot
CANN-robot成员
14 天前 评论:
🚀 CI 流水线已启动
📋 执行详情: 点击查看流水线
likedislike
CANN-robotCANN-robot成员
14 天前 添加了label:ci-pipeline-running
此处折叠了6条事件消息 查看更多
Ggaoxin成员
14 天前 修改了pull request 的描述
liyuewei
liyuewei成员
14 天前 评论:

/lgtm

likedislike
zhang_shengjie成员
14 天前 评论:

/lgtm

likedislike
CANN-robotCANN-robot成员
14 天前 添加了label:lgtm
yangyongqiang
yangyongqiang成员
14 天前 评论:

/approve

likedislike
CANN-robotCANN-robot成员
14 天前 添加了label:approved
CANN-robotCANN-robot成员
14 天前 合入了pull request