合并受阻
变更摘要
本 PR 为 exercises 目录新增一套基于 Ascend C 的 Leakyrelu 自定义算子工程：Host 侧完成算子注册、形状推导与多核/UB 切分 Tiling，并通过 LeakyreluTilingData 把切分参数与 negative_slope 属性下发到 Kernel；Kernel 侧按 tilingKey 分派 half/float 两种模板实例，采用双缓冲流水搬运与正负区间分段计算。同时补齐了构建脚本、aclnn 调用示例、精度比对样例数据以及 Host/Kernel 的 UT 框架。
主要改动
- 算子定义与形状推导: 新增
Leakyrelu算子注册（op_host/leakyrelu_def.cpp），输入x、输出y支持DT_FLOAT与DT_FLOAT16、FORMAT_ND，并声明Float类型属性negative_slope及ascend910b配置；leakyrelu_infershape.cpp中以输入 shape 直接推导输出 shape。 - Tiling 与属性下发:
op_host/leakyrelu_tiling.cpp读取平台coreNum/ubSize，按 32B 对齐与MIN_SPLIT_THRESHOLD做多核切分、按输入/输出双缓冲加 float 中间缓冲估算 UB 容量，向LeakyreluTilingData写入totalNum、blockFactor、ubFactor、negativeSlope、slopeTag，并依据 dtype 设置 tilingKey（half→LEAKYRELU_TPL_SCH_MODE_0、float→LEAKYRELU_TPL_SCH_MODE_1）；GetNegativeSlope在属性缺失时回退默认值。 - Kernel 实现: 新增模板类
NsLeakyrelu::Leakyrelu<T>（op_kernel/leakyrelu.h），使用TQue双缓冲DataCopy搬运、按ubLength_分块循环，计算统一在 float 域完成（Relu/Sub/Muls/Add），half分支经Cast升/降精度，尾块按对齐长度搬入但只写回有效元素；op_kernel/leakyrelu.cpp为__global__ __aicore__入口，用if constexpr按schMode实例化 half/float 两种 kernel。 - TilingData 与模板参数定义: 新增公共结构
LeakyreluTilingData（含LEAKYRELU_DEFAULT_SLOPE与LEAKYRELU_SLOPE_TAG标记，用于兼容未下发属性的场景）与leakyrelu_tiling_key.h中的ASCENDC_TPL_ARGS_DECL/ASCENDC_TPL_SEL模板声明，供 Host tiling 与 Kernel 共用。 - 构建、示例与测试配套: 新增
build.sh（支持-j、--make_clean、-u、-e，含源码指纹校验与产物检查）、examples/run.sh配合test_aclnn_leakyrelu.cpp/sample_support.h走 aclnn 接口调用、samples/case.json+reference.py+sample.py生成与校验 golden 数据，以及tests/ut下的TilingContextFaker/InfershapeContextFaker上下文构造器、ExecuteTestCase断言工具、Host tiling 用例与 CPU 侧 kernel UT。
代码审查
审查结论
- [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/examples/sample_support.h:44fix: Tensor 构造函数在 aclrtMalloc 之后抛异常时泄漏 device 显存 - [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/op_host/leakyrelu_infershape.cpp:25fix: InferShape 在输入 shape 为空时仍返回 GRAPH_SUCCESS，输出 shape 未被设置 - [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/op_kernel/leakyrelu.h:165fix: 尾块使用基础 DataCopy 搬运非 32B 对齐长度，可能越界写输出张量 - [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/op_kernel/leakyrelu.h:117fix: CopyIn 按 AlignUp 长度搬入，末核会越过输入张量末尾读取 - [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/op_kernel/leakyrelu.h:67fix: slopeTag 门控导致属性未下发时静默使用默认 slope - [建议]
2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/personal_Code_demo111111_11/code/exercises/tests/ut/op_host/CMakeLists.txt:4fix: SOC_VERSION 默认值大小写与 socToUpper 表键不匹配导致 SoC 版本被解析为空
| 类型 | 数量 |
|---|---|
| 🔴 阻塞 | 0 |
| 🟡 建议 | 6 |
💬 仅评论
/compile
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
⚠️ This PR does not yet meet the following requirements:lgtm (requires ≥ 2 person(s) per module)、approve (requires ≥ 1 person(s) per module)
Module Approval Details
| module | lgtm status | approve status |
|---|---|---|
| repo-cann/cann-outreach | ❌ (0/2)(You can also ask: 王艳秋, jxlang, Carolina_yuan, shaoyf, yanhf) | ❌ (0/1)(You can also ask: shaoyf, 王艳秋, 傅涛, jxlang, luzx66) |
💡 Tip:
- Committer can comment
/approveor/lgtm- Commenting
/approveimplies both code review (lgtm) and intent to merge (approve)
CLA Signature Pass
Code_demo111111, thanks for your pull request. All authors of the commits have signed the CLA. 👍
CANN-robotCANN-robot🟡 Medium Priority
变更行（leakyrelu.h:67-68）：negativeSlope_ = (tilingData->slopeTag == LEAKYRELU_SLOPE_TAG) ? tilingData->negativeSlope : LEAKYRELU_DEFAULT_SLOPE;
- 触发条件：host tiling 只要没有把
slopeTag填为LEAKYRELU_SLOPE_TAG（字段顺序/结构体版本不一致、属性解析遗漏、新增字段未同步初始化等），kernel 就会静默回退到硬编码 0.01。
🟡 Medium Priority
变更行（leakyrelu.h:117-118）：const int64_t copyNum = AlignUp(currentNum, alignElems_); DataCopy(xLocal, inputGMX[offset], static_cast<int32_t>(copyNum));
- 触发条件：totalNum 不是
alignElems_（float 8 / half 16）的整数倍时，最后一轮 tail 的 copyNum 会向上取整到 32B。 - 失败模式：MTE2 会读到输入张量末尾之后最多 31B（half 为 30B）未分配的 GM，越界读；在开启越界检查的环境（仿真/UT/部分硬件校验）会报 AIC Error 或读到脏数据参与计算（虽然结果不写回，但可能污染生成 NaN/Inf 而影响 Add 的整块向量计算）。
建议：对非对齐尾块使用 DataCopyPad 按 currentNum*sizeof(T) 实际字节数搬入，或确保 host tiling 下发的 totalNum 已 32B 对齐，避免越界读。
🟡 Medium Priority
变更行：leakyrelu_infershape.cpp 第 25-32 行，if (in_shape != nullptr) { *output_shape = *in_shape; } 之后无条件 return ge::GRAPH_SUCCESS;。
契约/行为：Leakyrelu 只有 1 个必选输入，InferShape 的正常契约是「输出 shape = 输入 shape」；当 context->GetInputShape(0) 返回 nullptr（输入未找到、可选输入/子图切分边界、部分动态 shape 场景）时，本实现既不设置 *output_shape，也不报错，直接返回成功。
失效模式：框架把未被写入（rank/dims 处于初始/未知状态）的输出 Shape 当作合法结果继续下传，后续算子按错误的 rank/维度做 tiling 与内存分配，可能计算出错误的输出尺寸或触发越界读写；GetOutputShape(i) 的 nullptr 已经返回 GRAPH_FAILED，而更关键的输入侧缺失却被静默放过，属于错误处理不一致。
|
30 | + if (in_shape == nullptr) { |
|
31 | + return ge::GRAPH_FAILED; |
| 30
| - } |
|
32 | + } |
|
33 | + *output_shape = *in_shape; |
🟡 Medium Priority
变更行（leakyrelu.h:165）：DataCopy(outputGMY[offset], yLocal, static_cast<int32_t>(currentNum));
- 契约：基础 DataCopy（UB→GM，MTE3）要求搬运数据量为 32B 的整数倍；
currentNum在尾块由blockLength_ - progress * ubLength_得到，当 totalNum 不是alignElems_（float 8 个元素 / half 16 个元素）的整数倍时必然非对齐。 - 失败模式：MTE3 按整块搬运时会把 32B 块尾部的冗余数据写到
outputGMY[offset+currentNum]之后的 GM，踩坏输出张量尾部或相邻 tensor 数据；在开启越界检测时（kernel UT / 仿真）直接报 AIC Error。注释“只写回有效元素，避免越界污染输出张量尾部”的意图正确，但基础 DataCopy 无法保证，非对齐尾块必须用 DataCopyPad。
建议：尾块改 DataCopyPad 按实际字节数写回，或由 host tiling 保证 totalNum 32B 对齐后再用 DataCopy；同时 CopyOut 与 CopyIn 的对齐策略需保持一致。
🟡 Medium Priority
变更行（sample_support.h:44-53）：check(aclrtMalloc(&device, ...)) 成功分配 device 显存后，后续两步仍可能抛异常——check(aclrtMemcpy(...))（第 46 行）与 if (!tensor) throw（第 52-53 行）。
失败模式：C++ 规定构造函数抛出异常时该对象的析构函数不会执行（成员 host/file 会被析构，但 ~Tensor() 不会），因此已分配的 device GM 显存无人释放；同理 host 之外的 aclTensor 无需释放，但 device 一定会泄漏。
触发条件：aclrtMemcpy 返回非 0（如设备侧内存不足、非法拷贝参数）或 aclCreateTensor 返回 nullptr（strides/shape 非法、内存不足）。本 helper 被每个 Tensor 复用，一旦出现多次失败会累积泄漏。
|
53 | + check(aclrtMalloc(&device, std::max(size_t(1), size), ACL_MEM_MALLOC_HUGE_FIRST)); |
|
54 | + try { |
|
55 | + if (size) |
|
56 | + check(aclrtMemcpy(device, size, host.data(), size, ACL_MEMCPY_HOST_TO_DEVICE)); |
|
57 | + std::vector<int64_t> strides(shape.size(), 1); |
|
58 | + for (size_t i = shape.size(); i > 1; --i) |
|
59 | + strides[i - 2] = strides[i - 1] * shape[i - 1]; |
|
60 | + tensor = aclCreateTensor(shape.data(), shape.size(), type, strides.data(), 0, ACL_FORMAT_ND, |
|
61 | + shape.data(), shape.size(), device); |
|
62 | + if (!tensor) |
| 53
| - throw std::runtime_error("tensor creation failed"); |
|
63 | + throw std::runtime_error("tensor creation failed"); |
|
64 | + } catch (...) { |
|
65 | + aclrtFree(device); |
|
66 | + device = nullptr; |
|
67 | + throw; |
|
68 | + } |
🟡 Medium Priority
变更行：op_host/CMakeLists.txt:4-6 设置 SOC_VERSION "Ascend910B"，第 39-41 行以 BUILD_SOC_VERSION=${SOC_VERSION} 下发宏，tiling_case_executor.cpp:125 用 STR(BUILD_SOC_VERSION) 得到字符串 "Ascend910B"。
影响链路：tiling_case_executor.cpp:118-124 的 socToUpper 表键全部是小写（"ascend910b"→"Ascend910B"、"ascend910_93" 等），第 126-127 行用 socToUpper[buildSocVersion] 查 "Ascend910B" 必然 miss；由于用的是 map::operator[]，miss 时不是报错而是插入并返回空串，于是 buildSocVersion 变成 ""。
失败模式：run.sh 不传 -DSOC_VERSION，因此恒走该默认值。随后（1）第 144 行 socToArch[""] 再插入空值，socversions = {{"NpuArch",""},{"Short_SoC_version",""}} 被 SetPlatformRes("version", ...) 写入平台信息；（2）compileInfoString 结尾拼成 "socVersion":""。任何依赖 SoC 版本/NpuArch 做分支的 tiling（如按架构选择 UB 切分或特性开关）会取到空字符串走默认/错误分支，而测试仍会“通过”，把平台配置错误静默带过，属于 UT 侧算出的期望值不可信的可靠性问题。修复方向：默认值改为小写 "ascend910b"（或用 socToUpper.find + 大小写无关比较，miss 时直接报错）。
建议：将默认 SOC_VERSION 改为与 tiling_case_executor.cpp 中 socToUpper 键一致的小写形式（"ascend910b"），或在该处统一转小写后再查表；同时确认 BUILD_SOC_VERSION 宏最终能命中 socToUpper 与 socToArch 两张表。
|
6 | + if(NOT DEFINED SOC_VERSION) |
|
7 | + set(SOC_VERSION "ascend910b") |
| 6
8 | endif() |
个人开发者
团队信息
成员
算子: exercises
本算子基于Ascend C多核架构实现LeakyReLU激活计算：Host侧完成输入切分与斜率属性解析，经TilingData传递给Kernel；Kernel侧各核独立处理数据分片，采用双缓冲流水搬运，正负区间分段计算（正数直通、负数乘以斜率），支持half与float两种精度。