aclnnHardswishBackwardV2 接口详解CANN ops-nn 中 HardSwish 反向传播算子的两段式调用与 Kernel 实现剖析【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn本文以 CANN ops-nn 仓库中activation/hard_swish_grad_v2模块的官方接口文档为骨架完整讲解aclnnHardswishBackwardV2接口的功能定义、计算公式以及与 V1 接口aclnnHardswishBackward的边界差异、两段式函数原型、参数与错误码约定并配套一份可直接参考的完整 C 调用示例同时结合仓库内的 op_api 实现、AscendC kernel、tiling 定义与测试 golden 代码还原该算子从 Host 侧入参校验、非连续 Tensor 规整到 Device 侧 Compare/Select 分档计算的实际执行路径帮助读者掌握在 NPU 上以 aclnn 两段式接口完成 HardSwish 激活反向梯度计算的完整实战方案。1. 产品支持情况根据接口文档与 模块 READMEaclnnHardswishBackwardV2当前支持情况如下产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I / 500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品Ascend 910√数据类型仅支持 FLOAT16、FLOAT32不支持 BFLOAT16需要注意数据类型支持是区分芯片平台的Ascend 910 平台不支持 BFLOAT16。这一点在 op_api 源码 aclnn_hardswish_backward_v2.cpp 中有明确体现两段式接口内部维护了按平台区分的 dtype 支持列表static const std::initializer_listop::DataType ASCEND910B_DTYPE_SUPPORT_LIST { op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16}; static const std::initializer_listop::DataType ASCEND910_DTYPE_SUPPORT_LIST {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};运行时通过GetDtypeSupportListV2按当前设备选择对应列表再做逐项校验。2. 功能说明与计算公式接口功能aclnnHardswish 前向算子的反向传播完成张量self的梯度计算。计算公式$$ out_{i} gradOutput_{i} \times gradSelf_{i} $$其中gradSelf的计算公式为$$ gradSelf_{i} \begin{cases} 0, self_{i} \le -3, \ self_{i} / 3 0.5, -3 \lt self_{i} \lt 3, \ 1, self_{i} \ge 3 \end{cases} $$2.1 与 aclnnHardswishBackwardV1的边界差异相比于 aclnnHardswishBackward 接口V2 的计算公式只做了细微调整——分档的边界开闭方向不同。对照 V1 公式self 取值V1aclnnHardswishBackwardV2aclnnHardswishBackwardV2self -300self -3-3/3 0.5 -0.5落入中间段0落入左段-3 self 3self/3 0.5self/3 0.5self 33/3 0.5 1.5落入中间段1落入右段self 311V2 在self ±3两个边界点处取的是硬阈值单侧导数值-3 处为 03 处为 1而非中间线性段的延拓值。这一处理与测试参考实现 golden.py 中的 torch 语义完全一致mask_greater x torch.tensor(-3.0, ...) # self -3 才保留线性段 mask_less x torch.tensor(3.0, ...) # self 3 才保留线性段否则取 1因此在需要与主流框架数值行为对齐的场景下应选用 V2 接口。3. 函数原型两段式接口每个算子分为两段式接口必须先调用aclnnHardswishBackwardV2GetWorkspaceSize接口获取计算所需 workspace 大小以及包含算子计算流程的执行器再调用aclnnHardswishBackwardV2接口执行计算。aclnnStatus aclnnHardswishBackwardV2GetWorkspaceSize( const aclTensor* gradOutput, const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)aclnnStatus aclnnHardswishBackwardV2( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)接口声明位于 aclnn_hardswish_backward_v2.h两个函数均以ACLNN_API宏导出供 C/C 工程直接链接调用。3.1 第一段接口 aclnnHardswishBackwardV2GetWorkspaceSize参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续 TensorgradOutputaclTensor*输入表示输入张量公式中的输入 gradOutput-BFLOAT16、FLOAT16、FLOAT32ND1-8√selfaclTensor*输入输入数据。公式中的 self支持空 Tensorshape 需要与 gradOutput、out 相同BFLOAT16、FLOAT16、FLOAT32ND1-8√outaclTensor*输出表示输出张量公式中的 out-BFLOAT16、FLOAT16、FLOAT32ND1-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程-----返回值aclnnStatus返回状态码具体参见aclnn 返回码。第一段接口会完成入参校验出现以下场景时报错返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 gradOutput、self 是空指针ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 的数据类型不在支持的范围之内ACLNN_ERR_PARAM_INVALID161002gradOutput、self 和 out 的数据类型不同ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 的 shape 不同ACLNN_ERR_PARAM_INVALID161002out 和 self 的 shape 不同上述校验在源码 aclnn_hardswish_backward_v2.cpp 的CheckParams函数中按序实现CheckNotNull3Tensor空指针检查含 out→CheckDtypeValid类型检查 →CheckShapeValidshape 检查其中 shape 检查还包含维度数不超过MAX_SUPPORT_DIMS_NUMS即文档中 1~8 维的约束。3.2 第二段接口 aclnnHardswishBackwardV2参数说明参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口获取executor输入op 执行器包含了算子计算流程stream输入指定执行任务的 Stream返回值aclnnStatus返回状态码具体参见aclnn 返回码。4. 调用示例示例代码如下仅供参考具体编译和执行过程请参考编译与运行样例。仓库中对应的完整可编译样例位于 test_aclnn_hard_swish_backward_v2.cpp。#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_hardswish_backward_v2.h #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vectorint64_t shape) { int64_t shape_size 1; for (auto i : shape) { shape_size * i; } return shape_size; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法资源初始化 auto ret aclInit(nullptr); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclInit failed. ERROR: %d\n, ret); return ret); ret aclrtSetDevice(deviceId); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSetDevice failed. ERROR: %d\n, ret); return ret); ret aclrtCreateStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtCreateStream failed. ERROR: %d\n, ret); return ret); return 0; } template typename T int CreateAclTensor(const std::vectorT hostData, const std::vectorint64_t shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMalloc failed. ERROR: %d\n, ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMemcpy failed. ERROR: %d\n, ret); return ret); // 计算连续tensor的strides std::vectorint64_t strides(shape.size(), 1); for (int64_t i shape.size() - 2; i 0; i--) { strides[i] shape[i 1] * strides[i 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1. 固定写法device/stream初始化, 参考acl API手册 // 根据自己的实际device填写deviceId int32_t deviceId 0; aclrtStream stream; auto ret Init(deviceId, stream); CHECK_RET(ret 0, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2. 构造输入与输出 std::vectorint64_t gradOutShape {4, 2}; std::vectorint64_t selfShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* selfDeviceAddr nullptr; void* gradOutDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* self nullptr; aclTensor* gradOut nullptr; aclTensor* out nullptr; std::vectorfloat selfHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat gradOutHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat outHostData {0, 0, 0, 0, 0, 0, 0, 0}; // 创建self aclTensor ret CreateAclTensor(selfHostData, selfShape, selfDeviceAddr, aclDataType::ACL_FLOAT, self); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建gradOut aclTensor ret CreateAclTensor(gradOutHostData, gradOutShape, gradOutDeviceAddr, aclDataType::ACL_FLOAT, gradOut); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建out aclTensor ret CreateAclTensor(outHostData, outShape, outDeviceAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); // 3. 调用CANN算子库API uint64_t workspaceSize 0; aclOpExecutor* executor; // 调用aclnnHardswishBackwardV2第一段接口 ret aclnnHardswishBackwardV2GetWorkspaceSize(gradOut, self, out, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardswishBackwardV2GetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr nullptr; if (workspaceSize 0) { ret aclrtMalloc(workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(allocate workspace failed. ERROR: %d\n, ret); return ret); } // 调用aclnnHardswishBackwardV2第二段接口 ret aclnnHardswishBackwardV2(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardswishBackwardV2 failed. ERROR: %d\n, ret); return ret); // 4. 固定写法同步等待任务执行结束 ret aclrtSynchronizeStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSynchronizeStream failed. ERROR: %d\n, ret); return ret); // 5. 获取输出的值将device侧内存上的结果拷贝至host侧 auto size GetShapeSize(outShape); std::vectorfloat resultData(size, 0); ret aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(copy result from device to host failed. ERROR: %d\n, ret); return ret); for (int64_t i 0; i size; i) { LOG_PRINT(result[%ld] is: %f\n, i, resultData[i]); } // 6. 释放aclTensor aclDestroyTensor(self); aclDestroyTensor(gradOut); aclDestroyTensor(out); // 7. 释放device资源 aclrtFree(selfDeviceAddr); aclrtFree(gradOutDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }按示例数据手算一下预期输出可以加深理解self {0,1,2,3,4,5,6,7}、gradOut同值时gradSelf依次为{0.5, 0.8333…, 1.1667…, 1, 1, 1, 1, 1}注意 self3 时 V2 取 1 而非 1.5out gradOut * gradSelf。5. 源码纵深两段式接口内部的计算图拼装阅读 aclnn_hardswish_backward_v2.cpp 可以看到第一段接口并非只做参数校验它还在 Host 侧用 l0 级 API 把整个计算图拼装到 OpExecutor 中。源码注释给出了完整流程* self other * | | * \ / * Contiguous(workspace_0) Contiguous(workspace_2) * \ / * HardswishBackwardV2(workspace_1) * | * ViewCopy * | * result对应实现逻辑依次为创建执行器CREATE_EXECUTOR()生成持有计算流程的uniqueExecutor参数检查CheckParams空指针 / dtype / shape空 Tensor 快速路径若gradOutput或self为空 Tensor直接workspaceSize 0并返回成功——即空 tensor 在 kernel 中支持非连续 Tensor 规整l0op::Contiguous(gradOutput)与l0op::Contiguous(self)分别将两个输入转为连续布局分别占 workspace_0 / workspace_2这是接口支持非连续 Tensor √的根本原因核心计算l0op::HardSwishGradV2(gradOutputContiguous, selfContiguous, ...)注册 kernel 节点workspace_1结果回写l0op::ViewCopy(hardswishBackwardOpOut, out, ...)将结果拷贝到可能非连续的out上汇总uniqueExecutor-GetWorkspaceSize()汇总各节点 workspaceReleaseTo(executor)将执行器所有权转移给调用方。第二段接口实现极简仅通过CommonOpExecutorRun(workspace, workspaceSize, executor, stream)把 Host 侧拼好的执行图提交到指定 stream 上执行。这也意味着第一段接口的耗时主要是 Host 侧构图真正的 Device 计算发生在第二段。6. 源码纵深Kernel 侧的分档计算与多核流水6.1 Kernel 入口与 TilingKernel 入口在 hard_swish_grad_v2.cpp按芯片核型宏编译分发到不同实现extern C __global__ __aicore__ void hard_swish_grad_v2(GM_ADDR gradOutput, GM_ADDR self, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); #if __CCE_AICORE__ 220 HardSwishGradV2220DTYPE_SELF op; op.Init(gradOutput, self, out, userWs, tilingData); op.Process(); #elif __CCE_AICORE__ 100 HardSwishGradV2100DTYPE_SELF op; op.Init(gradOutput, self, out, userWs, tilingData); op.Process(); #endif }Tiling 数据结构定义于 hard_swish_grad_v2_tiling_def.h由 Host 侧 tilingarch22 tiling / arch35 tiling填充BEGIN_TILING_DATA_DEF(HardSwishGradV2TilingData) TILING_DATA_FIELD_DEF(int64_t, elementNum); // 输入总元素数 TILING_DATA_FIELD_DEF(uint32_t, needCoreNum); // 实际需要的AI Core数 TILING_DATA_FIELD_DEF(uint64_t, ubSize); // 单核UB缓冲大小 TILING_DATA_FIELD_DEF(int64_t, elementNumEachCore); // 每轮每核处理的元素数 END_TILING_DATA_DEF;6.2 分档计算的向量化实现hard_swish_grad_v2_220.h 的Compute函数展示了文档公式在 AscendC VE 指令下的落地方式——用比较掩码 Select实现分段函数全程以 FP32 计算FP16/BF16 输入会先Cast到 FP32// step1: 生成两个比较掩码 CompareScalar(this-maskGreater, this-selfTensorFp32, this-negative, CMPMODE::GT, tmpDataCount); // self -3 CompareScalar(this-maskLessThan, this-selfTensorFp32, this-positive, CMPMODE::LT, tmpDataCount); // self 3 // step2: 先整体算出线性段 self/3 0.5 Muls(this-selfTensorFp32, this-selfTensorFp32, this-oneThird, dataCount); Adds(this-selfTensorFp32, this-selfTensorFp32, this-oneHalf, dataCount); // step3: 用掩码选出分段结果self -3 置 0self 3 置 1 Select(this-selfTensorFp32, this-maskGreater, this-selfTensorFp32, this-zero, SELMODE::VSEL_TENSOR_SCALAR_MODE, dataCount); Select(this-selfTensorFp32, this-maskLessThan, this-selfTensorFp32, this-one, SELMODE::VSEL_TENSOR_SCALAR_MODE, dataCount); // step4: out gradOutput * gradSelf Mul(this-selfTensorFp32, this-gradTensorFp32, this-selfTensorFp32, dataCount);其中oneThird常量取0.33333334fFP32 下 1/3 的近似值与 golden.py 中参考实现使用的0.33333334保持一致这也是测试能以 L1 级精度跨核对齐的关键。掩码条件GT(-3)/LT(3)严格大于/小于与公式中self ≤ -3 取 0、self ≥ 3 取 1的边界语义严格对应再次印证了 V2 接口的边界行为。6.3 多核切分与双缓冲流水线基类 hard_swish_grad_v2_base.h 与 220 实现中的Process展示了 Device 侧的任务切分策略将elementNum按elementNumEachCore切块再由needCoreNumber个 AI Core 轮流领取loopNumber轮尾块remain元素分给最后一个核blockIdx needCoreNumber的核直接返回即只有前 needCoreNum 个核参与计算每个核内采用pingPongFlag双缓冲CopyInAndCastMTE2 搬运 FP32 上转→ComputeVE 计算→CastAndCopyOutV 域下转 MTE3 写回三段配合MTE2_V/V_MTE3/MTE3_MTE2三组硬件事件标志形成搬运-计算重叠的流水线数据类型转换上有细节差异FP16 输出使用RoundMode::CAST_NONEBF16 输出使用RoundMode::CAST_RINT。Host 侧的 shape/dtype 推断逻辑见 hard_swish_grad_v2_infershape.cpp输出 shape 直接拷贝第一个输入gradOutput的 shape输出 dtype 继承输入 dtype这也解释了为何三个张量必须 shape 与 dtype 完全一致。7. 测试验证与精度参考算子测试组织在 tests 目录 下覆盖三个层面测试层文件验证内容Kernel 测试test_hard_swish_grad_v2.cpp直接调用hard_swish_grad_v2kernel与 golden 数据比对aclnn API 测试test_aclnn_hardswish_backward_v2.cpp覆盖两段式接口的入参校验空指针/类型/shape 错误分支与正向调用Host 单测test_hard_swish_grad_v2_infershape.cpp、tiling 测试验证 shape 推断与 tiling 参数生成数值参考实现 golden.py 明确了两点其一参考计算统一提升到 FP32compute_dtype torch.float32 if source_tensor.dtype in (float16, bfloat16)与 kernel内部 FP32 计算、出口转回的策略一致其二精度标准上 kernel 测试采用cross_checkL1 级跨核比对本地比对采用stat_rel_err统计相对误差。数据生成脚本位于 gen_data.py可用于理解测试数据分布如何覆盖self ∈ [-3, 3]的线性段及两侧饱和段。8. 约束说明与使用注意事项确定性计算aclnnHardswishBackwardV2默认确定性实现同一输入多次调用输出一致无需额外配置。shape 约束gradOutput、self、out的维度不超过 8 维且三者 shape 必须完全一致数据格式为 ND。类型约束三张量 dtype 必须一致支持 BFLOAT16 / FLOAT16 / FLOAT32Ascend 910 平台不支持 BFLOAT16。空 Tensor接口支持空 Tensor 输入第一段接口会直接返回workspaceSize 0的成功执行器。workspace 分配示例中if (workspaceSize 0)的判断是必要的——本算子在空 Tensor 或特定条件下可能返回 0用户侧分配 workspace 时应保留该判断释放时同样需按workspaceSize 0条件调用aclrtFree。与 V1 的选型若需要与 PyTorchtorch.nn.Hardswish反向的边界数值±3 处取 0 / 1对齐应选用 V2 接口V1 在self ±3处会落入线性段-0.5 / 1.5。综合以上aclnnHardswishBackwardV2通过标准的两段式 aclnn 接口把入参校验 → 非连续规整 → 分段梯度计算 → 结果回写整条链路封装在 OpExecutor 中调用者只需关注 workspace 分配与 stream 提交结合仓库内的 kernel 源码与 golden 测试可以完整复现并验证其数值行为是 NPU 上实现 HardSwish 反向传播的推荐接口。【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 aclnnHardswishBackwardV2 接口详解：CANN ops-nn 中 HardSwish 反向传播算子的两段式调用与 Kernel 实现剖析
aclnnHardswishBackwardV2 接口详解CANN ops-nn 中 HardSwish 反向传播算子的两段式调用与 Kernel 实现剖析【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn本文以 CANN ops-nn 仓库中activation/hard_swish_grad_v2模块的官方接口文档为骨架完整讲解aclnnHardswishBackwardV2接口的功能定义、计算公式以及与 V1 接口aclnnHardswishBackward的边界差异、两段式函数原型、参数与错误码约定并配套一份可直接参考的完整 C 调用示例同时结合仓库内的 op_api 实现、AscendC kernel、tiling 定义与测试 golden 代码还原该算子从 Host 侧入参校验、非连续 Tensor 规整到 Device 侧 Compare/Select 分档计算的实际执行路径帮助读者掌握在 NPU 上以 aclnn 两段式接口完成 HardSwish 激活反向梯度计算的完整实战方案。1. 产品支持情况根据接口文档与 模块 READMEaclnnHardswishBackwardV2当前支持情况如下产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I / 500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品Ascend 910√数据类型仅支持 FLOAT16、FLOAT32不支持 BFLOAT16需要注意数据类型支持是区分芯片平台的Ascend 910 平台不支持 BFLOAT16。这一点在 op_api 源码 aclnn_hardswish_backward_v2.cpp 中有明确体现两段式接口内部维护了按平台区分的 dtype 支持列表static const std::initializer_listop::DataType ASCEND910B_DTYPE_SUPPORT_LIST { op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16}; static const std::initializer_listop::DataType ASCEND910_DTYPE_SUPPORT_LIST {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};运行时通过GetDtypeSupportListV2按当前设备选择对应列表再做逐项校验。2. 功能说明与计算公式接口功能aclnnHardswish 前向算子的反向传播完成张量self的梯度计算。计算公式$$ out_{i} gradOutput_{i} \times gradSelf_{i} $$其中gradSelf的计算公式为$$ gradSelf_{i} \begin{cases} 0, self_{i} \le -3, \ self_{i} / 3 0.5, -3 \lt self_{i} \lt 3, \ 1, self_{i} \ge 3 \end{cases} $$2.1 与 aclnnHardswishBackwardV1的边界差异相比于 aclnnHardswishBackward 接口V2 的计算公式只做了细微调整——分档的边界开闭方向不同。对照 V1 公式self 取值V1aclnnHardswishBackwardV2aclnnHardswishBackwardV2self -300self -3-3/3 0.5 -0.5落入中间段0落入左段-3 self 3self/3 0.5self/3 0.5self 33/3 0.5 1.5落入中间段1落入右段self 311V2 在self ±3两个边界点处取的是硬阈值单侧导数值-3 处为 03 处为 1而非中间线性段的延拓值。这一处理与测试参考实现 golden.py 中的 torch 语义完全一致mask_greater x torch.tensor(-3.0, ...) # self -3 才保留线性段 mask_less x torch.tensor(3.0, ...) # self 3 才保留线性段否则取 1因此在需要与主流框架数值行为对齐的场景下应选用 V2 接口。3. 函数原型两段式接口每个算子分为两段式接口必须先调用aclnnHardswishBackwardV2GetWorkspaceSize接口获取计算所需 workspace 大小以及包含算子计算流程的执行器再调用aclnnHardswishBackwardV2接口执行计算。aclnnStatus aclnnHardswishBackwardV2GetWorkspaceSize( const aclTensor* gradOutput, const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)aclnnStatus aclnnHardswishBackwardV2( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)接口声明位于 aclnn_hardswish_backward_v2.h两个函数均以ACLNN_API宏导出供 C/C 工程直接链接调用。3.1 第一段接口 aclnnHardswishBackwardV2GetWorkspaceSize参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续 TensorgradOutputaclTensor*输入表示输入张量公式中的输入 gradOutput-BFLOAT16、FLOAT16、FLOAT32ND1-8√selfaclTensor*输入输入数据。公式中的 self支持空 Tensorshape 需要与 gradOutput、out 相同BFLOAT16、FLOAT16、FLOAT32ND1-8√outaclTensor*输出表示输出张量公式中的 out-BFLOAT16、FLOAT16、FLOAT32ND1-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程-----返回值aclnnStatus返回状态码具体参见aclnn 返回码。第一段接口会完成入参校验出现以下场景时报错返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 gradOutput、self 是空指针ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 的数据类型不在支持的范围之内ACLNN_ERR_PARAM_INVALID161002gradOutput、self 和 out 的数据类型不同ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 的 shape 不同ACLNN_ERR_PARAM_INVALID161002out 和 self 的 shape 不同上述校验在源码 aclnn_hardswish_backward_v2.cpp 的CheckParams函数中按序实现CheckNotNull3Tensor空指针检查含 out→CheckDtypeValid类型检查 →CheckShapeValidshape 检查其中 shape 检查还包含维度数不超过MAX_SUPPORT_DIMS_NUMS即文档中 1~8 维的约束。3.2 第二段接口 aclnnHardswishBackwardV2参数说明参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口获取executor输入op 执行器包含了算子计算流程stream输入指定执行任务的 Stream返回值aclnnStatus返回状态码具体参见aclnn 返回码。4. 调用示例示例代码如下仅供参考具体编译和执行过程请参考编译与运行样例。仓库中对应的完整可编译样例位于 test_aclnn_hard_swish_backward_v2.cpp。#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_hardswish_backward_v2.h #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vectorint64_t shape) { int64_t shape_size 1; for (auto i : shape) { shape_size * i; } return shape_size; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法资源初始化 auto ret aclInit(nullptr); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclInit failed. ERROR: %d\n, ret); return ret); ret aclrtSetDevice(deviceId); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSetDevice failed. ERROR: %d\n, ret); return ret); ret aclrtCreateStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtCreateStream failed. ERROR: %d\n, ret); return ret); return 0; } template typename T int CreateAclTensor(const std::vectorT hostData, const std::vectorint64_t shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMalloc failed. ERROR: %d\n, ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMemcpy failed. ERROR: %d\n, ret); return ret); // 计算连续tensor的strides std::vectorint64_t strides(shape.size(), 1); for (int64_t i shape.size() - 2; i 0; i--) { strides[i] shape[i 1] * strides[i 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1. 固定写法device/stream初始化, 参考acl API手册 // 根据自己的实际device填写deviceId int32_t deviceId 0; aclrtStream stream; auto ret Init(deviceId, stream); CHECK_RET(ret 0, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2. 构造输入与输出 std::vectorint64_t gradOutShape {4, 2}; std::vectorint64_t selfShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* selfDeviceAddr nullptr; void* gradOutDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* self nullptr; aclTensor* gradOut nullptr; aclTensor* out nullptr; std::vectorfloat selfHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat gradOutHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat outHostData {0, 0, 0, 0, 0, 0, 0, 0}; // 创建self aclTensor ret CreateAclTensor(selfHostData, selfShape, selfDeviceAddr, aclDataType::ACL_FLOAT, self); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建gradOut aclTensor ret CreateAclTensor(gradOutHostData, gradOutShape, gradOutDeviceAddr, aclDataType::ACL_FLOAT, gradOut); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建out aclTensor ret CreateAclTensor(outHostData, outShape, outDeviceAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); // 3. 调用CANN算子库API uint64_t workspaceSize 0; aclOpExecutor* executor; // 调用aclnnHardswishBackwardV2第一段接口 ret aclnnHardswishBackwardV2GetWorkspaceSize(gradOut, self, out, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardswishBackwardV2GetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr nullptr; if (workspaceSize 0) { ret aclrtMalloc(workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(allocate workspace failed. ERROR: %d\n, ret); return ret); } // 调用aclnnHardswishBackwardV2第二段接口 ret aclnnHardswishBackwardV2(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardswishBackwardV2 failed. ERROR: %d\n, ret); return ret); // 4. 固定写法同步等待任务执行结束 ret aclrtSynchronizeStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSynchronizeStream failed. ERROR: %d\n, ret); return ret); // 5. 获取输出的值将device侧内存上的结果拷贝至host侧 auto size GetShapeSize(outShape); std::vectorfloat resultData(size, 0); ret aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(copy result from device to host failed. ERROR: %d\n, ret); return ret); for (int64_t i 0; i size; i) { LOG_PRINT(result[%ld] is: %f\n, i, resultData[i]); } // 6. 释放aclTensor aclDestroyTensor(self); aclDestroyTensor(gradOut); aclDestroyTensor(out); // 7. 释放device资源 aclrtFree(selfDeviceAddr); aclrtFree(gradOutDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }按示例数据手算一下预期输出可以加深理解self {0,1,2,3,4,5,6,7}、gradOut同值时gradSelf依次为{0.5, 0.8333…, 1.1667…, 1, 1, 1, 1, 1}注意 self3 时 V2 取 1 而非 1.5out gradOut * gradSelf。5. 源码纵深两段式接口内部的计算图拼装阅读 aclnn_hardswish_backward_v2.cpp 可以看到第一段接口并非只做参数校验它还在 Host 侧用 l0 级 API 把整个计算图拼装到 OpExecutor 中。源码注释给出了完整流程* self other * | | * \ / * Contiguous(workspace_0) Contiguous(workspace_2) * \ / * HardswishBackwardV2(workspace_1) * | * ViewCopy * | * result对应实现逻辑依次为创建执行器CREATE_EXECUTOR()生成持有计算流程的uniqueExecutor参数检查CheckParams空指针 / dtype / shape空 Tensor 快速路径若gradOutput或self为空 Tensor直接workspaceSize 0并返回成功——即空 tensor 在 kernel 中支持非连续 Tensor 规整l0op::Contiguous(gradOutput)与l0op::Contiguous(self)分别将两个输入转为连续布局分别占 workspace_0 / workspace_2这是接口支持非连续 Tensor √的根本原因核心计算l0op::HardSwishGradV2(gradOutputContiguous, selfContiguous, ...)注册 kernel 节点workspace_1结果回写l0op::ViewCopy(hardswishBackwardOpOut, out, ...)将结果拷贝到可能非连续的out上汇总uniqueExecutor-GetWorkspaceSize()汇总各节点 workspaceReleaseTo(executor)将执行器所有权转移给调用方。第二段接口实现极简仅通过CommonOpExecutorRun(workspace, workspaceSize, executor, stream)把 Host 侧拼好的执行图提交到指定 stream 上执行。这也意味着第一段接口的耗时主要是 Host 侧构图真正的 Device 计算发生在第二段。6. 源码纵深Kernel 侧的分档计算与多核流水6.1 Kernel 入口与 TilingKernel 入口在 hard_swish_grad_v2.cpp按芯片核型宏编译分发到不同实现extern C __global__ __aicore__ void hard_swish_grad_v2(GM_ADDR gradOutput, GM_ADDR self, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); #if __CCE_AICORE__ 220 HardSwishGradV2220DTYPE_SELF op; op.Init(gradOutput, self, out, userWs, tilingData); op.Process(); #elif __CCE_AICORE__ 100 HardSwishGradV2100DTYPE_SELF op; op.Init(gradOutput, self, out, userWs, tilingData); op.Process(); #endif }Tiling 数据结构定义于 hard_swish_grad_v2_tiling_def.h由 Host 侧 tilingarch22 tiling / arch35 tiling填充BEGIN_TILING_DATA_DEF(HardSwishGradV2TilingData) TILING_DATA_FIELD_DEF(int64_t, elementNum); // 输入总元素数 TILING_DATA_FIELD_DEF(uint32_t, needCoreNum); // 实际需要的AI Core数 TILING_DATA_FIELD_DEF(uint64_t, ubSize); // 单核UB缓冲大小 TILING_DATA_FIELD_DEF(int64_t, elementNumEachCore); // 每轮每核处理的元素数 END_TILING_DATA_DEF;6.2 分档计算的向量化实现hard_swish_grad_v2_220.h 的Compute函数展示了文档公式在 AscendC VE 指令下的落地方式——用比较掩码 Select实现分段函数全程以 FP32 计算FP16/BF16 输入会先Cast到 FP32// step1: 生成两个比较掩码 CompareScalar(this-maskGreater, this-selfTensorFp32, this-negative, CMPMODE::GT, tmpDataCount); // self -3 CompareScalar(this-maskLessThan, this-selfTensorFp32, this-positive, CMPMODE::LT, tmpDataCount); // self 3 // step2: 先整体算出线性段 self/3 0.5 Muls(this-selfTensorFp32, this-selfTensorFp32, this-oneThird, dataCount); Adds(this-selfTensorFp32, this-selfTensorFp32, this-oneHalf, dataCount); // step3: 用掩码选出分段结果self -3 置 0self 3 置 1 Select(this-selfTensorFp32, this-maskGreater, this-selfTensorFp32, this-zero, SELMODE::VSEL_TENSOR_SCALAR_MODE, dataCount); Select(this-selfTensorFp32, this-maskLessThan, this-selfTensorFp32, this-one, SELMODE::VSEL_TENSOR_SCALAR_MODE, dataCount); // step4: out gradOutput * gradSelf Mul(this-selfTensorFp32, this-gradTensorFp32, this-selfTensorFp32, dataCount);其中oneThird常量取0.33333334fFP32 下 1/3 的近似值与 golden.py 中参考实现使用的0.33333334保持一致这也是测试能以 L1 级精度跨核对齐的关键。掩码条件GT(-3)/LT(3)严格大于/小于与公式中self ≤ -3 取 0、self ≥ 3 取 1的边界语义严格对应再次印证了 V2 接口的边界行为。6.3 多核切分与双缓冲流水线基类 hard_swish_grad_v2_base.h 与 220 实现中的Process展示了 Device 侧的任务切分策略将elementNum按elementNumEachCore切块再由needCoreNumber个 AI Core 轮流领取loopNumber轮尾块remain元素分给最后一个核blockIdx needCoreNumber的核直接返回即只有前 needCoreNum 个核参与计算每个核内采用pingPongFlag双缓冲CopyInAndCastMTE2 搬运 FP32 上转→ComputeVE 计算→CastAndCopyOutV 域下转 MTE3 写回三段配合MTE2_V/V_MTE3/MTE3_MTE2三组硬件事件标志形成搬运-计算重叠的流水线数据类型转换上有细节差异FP16 输出使用RoundMode::CAST_NONEBF16 输出使用RoundMode::CAST_RINT。Host 侧的 shape/dtype 推断逻辑见 hard_swish_grad_v2_infershape.cpp输出 shape 直接拷贝第一个输入gradOutput的 shape输出 dtype 继承输入 dtype这也解释了为何三个张量必须 shape 与 dtype 完全一致。7. 测试验证与精度参考算子测试组织在 tests 目录 下覆盖三个层面测试层文件验证内容Kernel 测试test_hard_swish_grad_v2.cpp直接调用hard_swish_grad_v2kernel与 golden 数据比对aclnn API 测试test_aclnn_hardswish_backward_v2.cpp覆盖两段式接口的入参校验空指针/类型/shape 错误分支与正向调用Host 单测test_hard_swish_grad_v2_infershape.cpp、tiling 测试验证 shape 推断与 tiling 参数生成数值参考实现 golden.py 明确了两点其一参考计算统一提升到 FP32compute_dtype torch.float32 if source_tensor.dtype in (float16, bfloat16)与 kernel内部 FP32 计算、出口转回的策略一致其二精度标准上 kernel 测试采用cross_checkL1 级跨核比对本地比对采用stat_rel_err统计相对误差。数据生成脚本位于 gen_data.py可用于理解测试数据分布如何覆盖self ∈ [-3, 3]的线性段及两侧饱和段。8. 约束说明与使用注意事项确定性计算aclnnHardswishBackwardV2默认确定性实现同一输入多次调用输出一致无需额外配置。shape 约束gradOutput、self、out的维度不超过 8 维且三者 shape 必须完全一致数据格式为 ND。类型约束三张量 dtype 必须一致支持 BFLOAT16 / FLOAT16 / FLOAT32Ascend 910 平台不支持 BFLOAT16。空 Tensor接口支持空 Tensor 输入第一段接口会直接返回workspaceSize 0的成功执行器。workspace 分配示例中if (workspaceSize 0)的判断是必要的——本算子在空 Tensor 或特定条件下可能返回 0用户侧分配 workspace 时应保留该判断释放时同样需按workspaceSize 0条件调用aclrtFree。与 V1 的选型若需要与 PyTorchtorch.nn.Hardswish反向的边界数值±3 处取 0 / 1对齐应选用 V2 接口V1 在self ±3处会落入线性段-0.5 / 1.5。综合以上aclnnHardswishBackwardV2通过标准的两段式 aclnn 接口把入参校验 → 非连续规整 → 分段梯度计算 → 结果回写整条链路封装在 OpExecutor 中调用者只需关注 workspace 分配与 stream 提交结合仓库内的 kernel 源码与 golden 测试可以完整复现并验证其数值行为是 NPU 上实现 HardSwish 反向传播的推荐接口。【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考