人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载导读本文以 CANN ops-nn 仓库中 aclnnMaxPoolV3.md 为核心系统讲解aclnnMaxPoolV3二维最大池化算子的功能定义、输出尺寸计算、两段式接口原型与全部参数语义并结合experimental/pooling/max_pool_v3目录下的算子定义、InferShape、Tiling 与 Kernel 源码以及 ST/UT 测试深入剖析该算子在 Atlas A2 训练/推理系列产品Ascend910B 及后续同代 SoC上的实现原理。读完本文你将掌握aclnnMaxPoolV3GetWorkspaceSizeaclnnMaxPoolV3的正确调用流程、各参数取值范围与约束、返回码错误语义并能对照源码理解其在 NPU 上的 shape 推导、负载均衡切分与确定性计算机制。一、功能概述2D 最大池化的 v3 实现aclnnMaxPoolV3是 CANN ops-nn 提供的**二维最大池化Max Pooling 2D**ACLNN 接口对 4 维 NCHW 输入张量在 H、W 两个空间维度上执行滑动窗口取最大值操作支持自定义ksize窗口大小、strides步长、pads填充和ceil_mode输出尺寸取整模式。目录 experimental/pooling/max_pool_v3 中的v3后缀是实现区分标识对外暴露的 ACLNN 接口名称为aclnnMaxPoolV3与仓库中其他 max pool 实现如max_pool_v2、max_pool_with_argmax_v3等相区分。1.1 计算公式以 NCHW 格式为例H 维度索引为 2W 维度索引为 3每个输出位置取窗口内最大值$$ y[n, c, h_o, w_o] \max_{i0}^{kH-1} \max_{j0}^{kW-1} x[n, c, h_o \cdot sH i - padT, w_o \cdot sW j - padL] $$窗口超出输入边界的元素视为 $-\infty$不参与最大值计算。这也与仓库 ST 测试的 PyTorch 参考实现一致——在 executor_aclnnMaxPoolV3.py 中非对称 padding 场景会先对输入执行torch.nn.functional.pad(..., valuefloat(-inf))再用padding0的max_pool2d计算用 $-inf$ 填充边界保证池化语义一致。1.2 输出尺寸计算CALCULATED padding 模式输出尺寸由输入尺寸、padding、kernel size、stride 与ceil_mode共同决定$$ H_{out} \left\lfloor \frac{H_{in} padT padB - kH (ceil_mode; ?; sH - 1 : 0)}{sH} \right\rfloor 1 $$$$ W_{out} \left\lfloor \frac{W_{in} padL padR - kW (ceil_mode; ?; sW - 1 : 0)}{sW} \right\rfloor 1 $$当ceil_mode true时如果最后一个窗口起始位置超出H_in padT或 W 侧超出W_in padL会额外减少一个输出避免最后一个窗口完全落在填充区域之外。该逻辑在 Host 侧有精确的源码对应InferShape 与 Tiling 共用 max_pool_v3_util.h 中的CalculateUpdateDim函数其实现为int64_t outputSize DivRtn(dim_size padL padR - ksize (ceil_mode ? stride - 1 : 0), stride) 1; if (ceil_mode) { if ((outputSize - 1) * stride dim_size padL) { --outputSize; } }注意其中的DivRtn是向下取整除法floor division见 max_pool_v3_util.h当被除数为负时向负无穷方向取整这与 C/C 的截断除法行为不同是保证尺寸公式在边界场景下正确性的关键细节。二、产品支持情况产品是否支持Atlas A2 训练系列产品 / Atlas A2 推理系列产品√数据类型的支持随 SoC 不同而略有差异Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持 FLOAT16、FLOAT32、BFLOAT16其中BFLOAT16 仅 Ascend910B 及后续同代 SoC 支持文档原文以term标注说明。该差异在源码中得到双重印证在 ACLNN 接口层aclnn_max_pool_v3.cpp 定义了两份数据类型支持列表ASCEND910_DTYPE_SUPPORT_LIST仅DT_FLOAT、DT_FLOAT16与ASCEND910B_DTYPE_SUPPORT_LIST增加DT_BF16并按实际 SoC 选取在 Tiling 层max_pool_v3_tiling.cpp 会在socVersion既不是ASCEND910B也不是ASCEND310B且输入为 BF16 时报错BF16 not supported on this SoC version.从编译期/运行期双保险杜绝 BF16 在不受支持硬件上运行。三、函数原型与两段式调用流程aclnnMaxPoolV3遵循 CANN 标准的两段式接口设计必须先调用第一段aclnnMaxPoolV3GetWorkspaceSize获取执行器executor和 workspace 大小再调用第二段aclnnMaxPoolV3执行计算。3.1 第一段接口aclnnStatus aclnnMaxPoolV3GetWorkspaceSize( const aclTensor* x, const aclIntArray* ksize, const aclIntArray* strides, const aclIntArray* pads, const aclScalar* ceilMode, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)3.2 第二段接口aclnnStatus aclnnMaxPoolV3( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)从接口实现看aclnn_max_pool_v3.cpp 的第一段接口完成以下工作调用CheckParams做入参校验空指针、数据类型、维度检查对空 Tensor 做短路处理当x或out为空时直接返回workspaceSize 0无需真正构建计算图对非连续输入调用l0op::Contiguous做连续化再通过ReshapeSelfValueGetActivation处理 shape读取ceilModeToInt64() ! 0视为 true并调用l0op::MaxPoolV3构建算子执行图通过GetWorkspaceSize()获取实际需要的 workspace 大小并释放 executor。第二段接口则直接调用CommonOpExecutorRun(workspace, workspaceSize, executor, stream)完成 Device 侧任务下发。四、aclnnMaxPoolV3GetWorkspaceSize 参数详解参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续TensorxaclTensor*输入4 维输入张量NCHW 格式支持空 Tensor不支持 broadcastBFLOAT16仅 Ascend910B 及后续同代 SoC 支持、FLOAT16、FLOAT32ND4√ksizeaclIntArray*输入池化窗口大小长度为 4格式 [N, C, H, W]不可为空H、W 维度值必须大于 0int64_t[]-[4]-stridesaclIntArray*输入池化步长长度为 4格式 [N, C, H, W]不可为空H、W 维度值必须大于 0int64_t[]-[4]-padsaclIntArray*输入填充大小长度为 4格式 [pad_top, pad_bottom, pad_left, pad_right]允许为空等同于全 0 填充默认值为 {0, 0, 0, 0}int64_t[]-[4]-ceilModeaclScalar*输入ceil 模式开关非 0 表示使用 ceil 模式计算输出尺寸允许为空等同于 false默认值为 false(0)int64_t---outaclTensor*输出池化后的输出张量支持空 Tensor数据类型必须与 x 一致BFLOAT16仅 Ascend910B 及后续同代 SoC 支持、FLOAT16、FLOAT32ND4√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小-----executoraclOpExecutor**输出返回 op 执行器包含算子计算流程-----4.1 参数语义的源码级说明x 与 out 的 4 维 NCHW 约束算子图定义 max_pool_v3_proto.h 将输入限定为TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16})算子定义 max_pool_v3_def.cpp 进一步声明输入输出均要求FORMAT_ND并开启AutoContiguous这与不支持 broadcast、隐式类型提升的约束对应。ksize / strides 的空间维度必须大于 0在 InferShape 与 Tiling 共用的 max_pool_v3_util.h 中ValidateSpatialDims会校验ksize[h,w]与strides[h,w]均大于 0否则返回GRAPH_FAILED并记录错误日志。pads 与 ceil_mode 均为可选属性算子定义max_pool_v3_def.cpp中pads默认{0, 0, 0, 0}、ceil_mode默认falseInferShape 与 Tiling 在属性缺失时都会回退到默认值DEFAULT_PADS与false见 max_pool_v3_infershape.cpp 和 max_pool_v3_tiling.cpp。pads 数组的语义布局[pad_top, pad_bottom, pad_left, pad_right]代码中以PAD_TOP 0、PAD_BOTTOM 1、PAD_LEFT 2、PAD_RIGHT 3四个常量索引见 max_pool_v3_util.h。注意它与 PyTorchmax_pool2d的对称 padding 语义不同PyTorch 无法直接表达非对称 padding这正是 ST 参考实现中需要先F.pad再池化的原因。五、返回值与错误码两段接口均返回aclnnStatus状态码具体取值参见aclnn返回码。第一段接口aclnnMaxPoolV3GetWorkspaceSize会完成入参校验以下场景直接报错返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 x、ksize、strides 或 out 是空指针ACLNN_ERR_PARAM_INVALID161002x 或 out 的数据类型不在支持范围内ACLNN_ERR_PARAM_INVALID161002x 和 out 的数据类型不一致ACLNN_ERR_PARAM_INVALID161002x 或 out 的维度大于 8对照 aclnn_max_pool_v3.cpp 的CheckParams实现可见先做CheckNotNull2Tensor(x, out)与ksize/strides的空指针检查对应 161001再通过CheckDtypeValidActivation按 SoC 选择的支持列表校验类型、用OP_CHECK_MAX_DIM校验维度上限MAX_SUPPORT_DIMS_NUMS对应 161002。六、aclnnMaxPoolV3 参数详解参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口 aclnnMaxPoolV3GetWorkspaceSize 获取executor输入op 执行器包含算子计算流程stream输入指定执行任务的 Stream使用范式伪代码// 1. 第一段获取 workspace 大小与执行器 aclnnStatus ret aclnnMaxPoolV3GetWorkspaceSize(x, ksize, strides, pads, ceilMode, out, workspaceSize, executor); // 2. 申请 Device 侧 workspace void* workspace aclrtMalloc(workspaceSize, ...); // 3. 第二段执行计算 ret aclnnMaxPoolV3(workspace, workspaceSize, executor, stream); // 4. 释放资源workspace、executor、tensor 等七、约束说明汇总默认确定性实现算子保证相同输入在多次运行时结果一致相关背景可参考确定性计算说明。不支持 broadcast。不支持隐式类型提升输入输出类型必须严格一致。输入张量必须为 4 维NCHW 格式。ksize和strides的 H、W 维度值必须大于 0。BFLOAT16 仅在 Ascend910B 及后续同代 SoC 上支持。支持非连续 Tensor接口层会内部执行连续化l0op::Contiguous与ViewCopy相关机制参见非连续 Tensor 说明。八、源码级实现原理从图定义到 Kernel 下发aclnnMaxPoolV3的完整调用链跨越四个层次均在 experimental/pooling/max_pool_v3 目录下可查8.1 图 IR 定义op_graphmax_pool_v3_proto.h 通过REG_OP(MaxPoolV3)注册图算子输入xFLOAT16/FLOAT/BF16、输出y、必选属性ksize/stridesListInt、可选属性pads默认{0,0,0,0}与ceil_mode默认 false。8.2 算子定义与 InferShapeop_hostmax_pool_v3_def.cpp 注册OpDef声明输入输出数据类型与 ND 格式声明四个属性及默认值并为 AICore 配置ascend910b平台开启DynamicCompileStaticFlag(true)、DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)与PrecisionReduceFlag(true)说明该算子支持动态 shape 并允许精度降级优化。max_pool_v3_infershape.cpp 实现 shape 与 dtype 推导空间维度为 4 时直接通过CalculateUpdateDim计算 H/W 输出维度否则输出未知维度InferDtypeMaxPoolV3简单地将输出类型设置为输入类型对应不支持隐式类型提升约束。8.3 Tiling 策略op_hostmax_pool_v3_tiling.cpp 是性能调度的核心可以观察到三个关键设计输出元素级负载均衡ComputeCoreDistribution按big core / small core方案把n*c*hOut*wOut个输出元素分给最多coreNum个核前formerNum个核分到ceil(total/cores)个元素其余核分到floor(total/cores)个元素见 max_pool_v3_tiling.cpp单元素 Tile 双缓冲每个输出元素是一个独立 tiletileDataNum 1Tiling 数据结构中共享常量MAX_POOL_V3_BUFFER_NUM 2表示输入队列采用双缓冲见 max_pool_v3_tiling_data.hUB 容量与 BF16 平台校验ValidatePoolingOutput会按数据类型换算elementsPerBlockfloat32 每 32B block 2 个元素float16/bf16 每 block 4 个元素校验kH*kW窗口能否装入 UB见 max_pool_v3_tiling.cpp。Tiling 结果通过MaxPoolV3TilingData结构体max_pool_v3_tiling_data.h在 Host 与 Device 之间传递字段覆盖每核元素分布、NCHW 各维度尺寸、池化参数kH/kW/sH/sW/padT/padL以及派生量inHW/outHW。8.4 Kernel 执行op_kernelmax_pool_v3.cpp 是 AscendC 内核入口__global__ __aicore__函数按KERNEL_TYPE_AIV_ONLY类型任务运行通过REGISTER_TILING_DEFAULT注册并解析 Tiling 数据然后实例化KernelMaxPoolV3DTYPE_X将每核元素分布、shape 与池化参数传入op.Init(...)后调用op.Process()完成窗口内取最大值与写回。8.5 测试覆盖testsST 测试executor_aclnnMaxPoolV3.py 提供 NPU/CPU 双后端的 PyTorch 参考实现基于max_pool2d非对称 padding 时用 $-inf$ 显式 pad配合 all_aclnnMaxPoolV3.json 定义的用例数据驱动测试UT 测试test_aclnn_max_pool_v3.cpp 覆盖 op_api 接口层test_max_pool_v3_infershape.cpp 覆盖 Host 侧 shape 推导。九、编译与运行该算子位于experimental目录属于实验特性编译与运行需显式携带--experimental参数命令来自 max_pool_v3 README# 编译算子包Ascend910B SoC bash build.sh --pkg --socascend910b --experimental --opsmax_pool_v3 # 运行示例eager 模式 cust 自定义算子 bash build.sh --run_example max_pool_v3 eager cust --experimental更多工程级编译细节可参考仓库根目录的 CMakeLists.txt 与 docs/QUICKSTART.md。十、总结aclnnMaxPoolV3是 CANN ops-nn 面向 Atlas A2 训练/推理系列产品提供的 2D 最大池化 ACLNN 接口核心特性可归纳为标准两段式接口GetWorkspaceSize完成校验、连续化与执行图构建aclnnMaxPoolV3负责实际下发执行完整参数语义ksize/strides按 [N,C,H,W] 布局且 H/W 必须大于 0pads按 [top,bottom,left,right] 布局且可缺省ceil_mode控制输出尺寸的取整模式严格类型与格式约束仅支持 NCHW ND 格式的 FLOAT16/FLOAT32/BFLOAT16BF16 限 Ascend910B 及同代 SoC不支持 broadcast 与隐式类型提升支持非连续 Tensor源码全链路可查从op_graph图定义、op_host的 InferShape/Tiling 到op_kernel的 AscendC 实现与 ST/UT 测试均有完整落地CalculateUpdateDim的向下取整除法、big/small core 负载均衡与 UB 容量校验是实现正确性与性能的关键细节。如需进一步深入推荐继续阅读 aclnnMaxPoolV3.md 原始文档、两段式接口说明 以及 aclnn 返回码说明。赞分享人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载相关推荐CANN ops-nn 算子 aclnnHardsigmoidBackward 接口全解析两段式调用、参数约束与源码实现CANN ops nn 算子 aclnnHardsigmoidBackward 接口全解析两段式调用、参数约束与源码实现 本文基于 CANN ops nn 开人工智能算子库深度学习CANNAscendCANN ops-nn IndexFillD 算子完全指南aclnn 两段式接口调用、参数语义与 NPU 实现原理CANN ops nn IndexFillD 算子完全指南aclnn 两段式接口调用、参数语义与 NPU 实现原理 本篇技术指南围绕 CANN ops nn人工智能算子库深度学习CANNAscendCANN ops-nn GatherNd 算子全解析计算语义、aclnnGatherNd 两段式接口调用与 NPU 源码实现CANN ops nn GatherNd 算子全解析计算语义、aclnnGatherNd 两段式接口调用与 NPU 源码实现 导读 本文以 index/gat人工智能算子库深度学习CANNAscend创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考