Pull Request已成功合入, 合并人@CANN-robot
(感谢 繁依Fanyi 的贡献)变更摘要
本PR实现了Challenge01-DivCustom逐元素除法算子,补全了完整的Ascend C设备端kernel算子与Host端ACL调用逻辑。设备端通过KernelDiv类使用双缓冲队列(inQueueX/inQueueY/outQueueZ)实现CopyIn→Compute→CopyOut流水线,在8个AI Core上分块处理16384个float32元素的逐元素除法;Host端通过kernel_div函数完成ACL初始化、设备内存申请、数据搬运、kernel启动与结果回拷,并显式指定KERNEL_TYPE_AIV_ONLY任务类型。
主要改动
- 新增
div_custom.asc设备端算子实现:定义了KernelDiv类,使用双缓冲队列与分块策略(BLOCK_DIM=8、TILE_NUM=8、BUFFER_NUM=2)实现完整的CopyIn→Compute→CopyOut流水线,通过AscendC::Div执行逐元素除法 - 新增
div_custom全局kernel入口函数:通过__global__ __aicore__声明kernel,接收DivCustomTilingData分块参数,内部实例化KernelDiv并调用Init/Process,显式指定KERNEL_TYPE_AIV_ONLY任务类型 - 新增
kernel_divHost端调用函数:完成ACL运行时初始化、设备内存分配(aclrtMalloc)、Host到Device数据搬运、kernel启动与流同步、Device到Host结果回拷以及资源释放的完整生命周期管理 - 新增测试验证与构建脚本:
main函数构造16384个元素的测试向量(valueX=1.2f,valueY=2.3f)并通过VerifyResult逐元素比对验证精度;run.sh提供一键构建与运行流程


代码审查
审查总结
| 文件 | 审查结果 |
|---|---|
CMakeLists.txt |
✅ 无问题(与模板一致) |
run.sh |
✅ 无问题(与模板一致,&& 链正确,变量已引号) |
div_custom.asc |
⚠️ P3 ×1(VerifyResult 长度校验不足) |
各优先级问题数: P0: 0, P1: 0, P2: 0, P3: 1
总体风险判断:低风险。 算子核心实现(Tiling 切分、双缓冲流水线、CopyIn → Compute → CopyOut 数据搬运与计算、ACL Host 端内存管理与 Kernel 启动)逻辑正确,与题目要求的数据规模(16384 元素 / 8 个 AI Core)完全匹配。唯一发现的是 VerifyResult 中 std::equal 的三参数用法未校验输出向量长度,在当前测试用例下不会触发,但属于可改进的健壮性问题。代码已通过人工验证,可以合入。
⚠️ 已识别出整体风险,但无法提取行内评论,请参考整体评估。


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
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/cann-outreach | ✅ 傅涛, shaoyf (2/2) | ✅ shaoyf, 傅涛 (2/1) |
💡 Tip:
- Committer can comment
/approveor/lgtm- Commenting
/approveimplies both code review (lgtm) and intent to merge (approve)


变更描述 / Description
补全 Challenge01-DivCustom 逐元素除法算子,实现完整的 Ascend C 设备端流水线与 Host 端调用:
改动类型 / Change Type
关联 Issue / Related Issues
测试信息 / Testing
bash run.sh[Success] Case accuracy is verification passed.自验证截图