ATVC Pool模板实战用TILE_LAYOUT与TILE_PADDING实现自定义Edge边缘检测算子【免费下载链接】atvcATVCAscend C Templates for Vector Compute是为基于Ascend C开发的典型Vector算子封装的一系列模板头文件的集合可帮助用户快速开发典型Vector算子。项目地址: https://gitcode.com/cann/atvcATVCAscend C Templates for Vector Compute的样例集examples/edge/README.md通过一个典型的非逐元素算子——自定义 Edge 边缘检测算子演示了如何利用 ATVC 的 Pool 模板把元素结果依赖周围相邻元素这类计算从手写 GM/UB 搬运与分块逻辑中解放出来。读完本文你将理解 Pool 模板Tile 分块 Padding 邻域的设计原理掌握TILE_LAYOUT、TILE_PADDING两个编译态参数的约束与取值逻辑并能够照 edge.cpp 的 Kernel 直调方式完成自定义邻域算子的开发、编译与精度验证。样例概述与适用范围本样例的目标是利用 ATVC 实现自定义 Edge 单算子并完成功能验证其关键属性如下与 README 保持一致算子功能自定义 Edge 计算功能——每个元素的结果由其周围相邻元素通过自定义计算得到使用的 ATVC 模板PoolATVC::Kernel::PoolOpTemplate调用方式Kernel 直调方式直接启动核函数支持的 AI 处理器型号Ascend 910C、Ascend 910B。Edge 算子的数学公式定义如下输入为二维数组例如 x [ x0, x1, x2, ... x3, x4, x5, ... x6, x7, x8, ...] y4 min(abs(((x2 x5 x8) - (x0 x3 x6)) / 3), 255) 以此类推其他元素的计算结果。即中心元素x4的结果等于右列三元素之和减去左列三元素之和除以 3 后取绝对值并与 255 取最小值。这是一个典型的水平方向 3x3 邻域算子类似 Sobel 水平梯度天然依赖左、右、上、下相邻元素——这正是选择 Pool 模板的原因Pool 的分块机制可以为每个 Tile 自动带出上下左右的 Padding 邻域数据使读一个块、算一块的模型直接适配邻域计算。算子规格如下项目说明算子类型OpTypeEdge输入 namex输入 width / height1023 / 2517输入 data type / formatfloat / ND输出 namez输出 width / height1023 / 2517输出 data type / formatfloat / ND核函数名EdgeCustom规格限制当前模板只支持 2 维 shape 按 16 元素个数对齐、TILE_LAYOUT{16, 16}、TILE_PADDING{8, 8, 1, 1}的场景注意规格限制一栏的含义TILE_LAYOUT{16, 16}指基本块宽 16、高 16 个元素TILE_PADDING{8, 8, 1, 1}的四个分量对应left/right/up/down定义见 pool_common.h 中的PoolTilePadding即左右各扩展 8 个元素、上下各扩展 1 个元素恰好覆盖公式所需的左右相邻列与上下相邻行。输入宽高 1023、2517 均满足按 16 元素个数对齐实际为16N 剩余的尾块裁剪场景由模板自动处理。为什么 Edge 算子要选 Pool 模板从源码结构看Pool 模板 PoolOpTemplate 的核心能力是二维 Tile 分块 Padding 邻域搬运 尾块裁剪它把用户代码限制在一个固定形状的 UB 基本块内分块调度host 端将整图按TILE_LAYOUT切分为 tile 网格kernel 端每个 AI Core 负责其中若干 tilecurCoreTileCnt_通过UpdateCurTileOffset()逐块推进见 pool_op_template.h邻域搬运CopyIn每个 tile 从 GM 拷入 UB 时DataCopyPad会额外带上TILE_PADDING指定的上下左右 Padding 区域见 CopyInAllTensors。对 Edge 算子来说这就是免费获得的 3x3 邻域tile 左右各多 8 列、上下各多 1 行尾块裁剪CalcCurTileScope靠近图像右边界/下边界的 tile 会按需收缩 padding避免越界读取见 CalcCurTileScope因此 1023×2517 这样不能被 tile 整除的形状可以直接运行。也就是说用户只需编写一个基本块内的计算逻辑Compute仿函数分块、搬运、边界裁剪、多核分配全部由模板承担。编译态参数TILE_LAYOUT 与 TILE_PADDING在 edge.cpp 中计算仿函数Edge2C3ComputeFunc以静态常量形式声明了两个模板级参数template typename Traits struct Edge2C3ComputeFunc { static constexpr ATVC::Layout2Dim TILE_LAYOUT{16, 16}; // 基本块宽高 宽需要32B对齐 未裁剪前的 static constexpr ATVC::PoolTilePadding TILE_PADDING{ 8, 8, 1, 1}; // tile块上下左右padding的设置left/right需要32B对齐 未裁剪前基础值 // ... };两者分别决定了 UB 中基本块的形状与邻域扩展范围取值时必须满足对齐约束TILE_LAYOUT{16, 16}基本块为 16×16 个 float 元素每行 16×4B 64B满足 UB 拷贝宽度 32B 对齐要求代码注释中宽需要32B对齐即指TILE_LAYOUT.width × sizeof(T)须为 32 的倍数TILE_PADDING{8, 8, 1, 1}left8、right8、up1、down1。左右 padding 各 8 元素 × 4B 32B满足注释中left/right 需要 32B 对齐的约束上下 padding 各 1 行恰好覆盖公式中±width的上下相邻行引用。由此每个 AI Core 处理的 UB 基本块尺寸为basicTensorCnt_ (TILE_LAYOUT.width TILE_PADDING.left TILE_PADDING.right) × (TILE_LAYOUT.height TILE_PADDING.up TILE_PADDING.down) (16 8 8) × (16 1 1) 32 × 18 576 个元素该值由模板在 CalcCurCoreStart 中计算basicTensorCnt_用户仿函数里同样可用c.GetSize()获取本例中输出张量大小为 576 个元素即calcSize。host 端 tiling 由 CalcPoolTiling 完成它根据totalLayout与tileLayout计算 tile 网格数并分配 AI Coreuint32_t tileNum ((totalH tileH - 1) / tileH) * ((totalW tileW - 1) / tileW); param.tilingData.blockNum tileNum; if (tileNum compileInfo.vectorCoreNum) { param.tilingData.blockNum compileInfo.vectorCoreNum; // tile数超过核数时按核数截断 } param.tilingData.numPerBlock tileNum / param.tilingData.blockNum; // 每核平均处理的tile数 param.tilingData.tailBlockCnt tileNum % param.tilingData.blockNum; // 需要多跑一轮的尾核数对本例 1023×2517、tile 16×16横向ceil(1023/16)65列、纵向ceil(2517/16)158行共65×15810270个 tile远多于 AI Core 数量因此blockNum被截断为vectorCoreNum每个核通过numPerBlock个 tile尾核再多 1 个遍历全图。PoolTilingData各字段含义可参考 pool_common.h。核函数Edge2C3ComputeFunc 的计算逻辑自定义计算仿函数通过重载operator()接收模板派发的三个LocalTensor输入a、输出c均为 float来自 UB 中的基本块与临时张量tempint32_t用于索引运算。完整签名见 edge.cpp其计算可以拆成两个阶段理解阶段一用 Gather 组装左右相邻列算出右列三元素和 − 左列三元素和static constexpr uint32_t TENSOR_WIDTH TILE_PADDING.left TILE_LAYOUT.width TILE_PADDING.right; // 基本块内存布局以3x3局部窗口为例: // 0 1 2 // 3 4 5 // 6 7 8 AscendC::CreateVecIndexU(temp, (int32_t)2, calcSize); // temp 2,4,6,8,...步长2的行内偏移 AscendC::MulsU(temp, temp, sizeT, calcSize); // 元素个数 - 字节偏移 AscendC::LocalTensoruint32_t tempRef temp.template ReinterpretCastuint32_t(); AscendC::Gather(c, a, tempRef, 0, calcSize - 2); // c[i] a[i2]右列元素 AscendC::Sub(a, c, a, calcSize); // a[i] - a[i2]先做 a[i] - 右列 AscendC::AddsU(temp, temp, (sizeT * -3), calcSize); // 字节偏移 -12即指向左列 AscendC::Relu(temp, temp, 1); // 边界保护负索引置0裁剪左padding AscendC::Gather(c, a, tempRef, 0, calcSize - 2); // 再累加上左列这里有两个值得注意的工程设计索引复用CreateVecIndex以步长 2 生成 2、4、6、8… 的偏移序列一次Gather即可取出每个 3 行窗口中的右列偏移 2 的 1、4、7 位置把 3x3 窗口内的列访问转化为对基本块的一维批量操作避免在标量循环里逐个寻址Relu做越界裁剪在图像最左列左 padding 区域已被CalcCurTileScope收缩为 0i-2会算成负字节偏移。源码中Relu(temp, temp, 1)将这些负索引截断为 0对应 edge.cpp 处保证最左列输出为 0——因为公式中右列和 − 左列和在无左邻域时按模板语义取 0这也是样例精度校验只覆盖内部区域的原因之一。阶段二对窗口中心列做min(abs((sum_right - sum_left)/3), 255)AscendC::Add(a[TENSOR_WIDTH], c, c[TENSOR_WIDTH * 2], calcSize - TENSOR_WIDTH * 2); // 上窗口中心列 下窗口中心列先累加到 c AscendC::Add(c, a, c, calcSize); // 加上本窗口中心列即 a[i]已含 ±左右列 AscendC::Muls(c, c, 1 / 3.0f, calcSize); AscendC::Abs(c, c, calcSize); AscendC::Mins(c, c, 255.0f, calcSize);TENSOR_WIDTH 8 16 8 32是带 padding 的一行元素数因此a[TENSOR_WIDTH]恰好指向下一窗口的中心列。整段计算全程使用 Ascend C 的向量 APIGather/Sub/Add/Muls/Relu/Abs/Mins每个算子间用PipeBarrierPIPE_V()同步与 Pool 模板的CopyIn → Compute → CopyOut流水线见 Process衔接。需要说明的是README 给出的公式以中心元素 y4 为例由于上下 padding 为 1 行Add的两步跨窗口累加实现的是上、中、下三个窗口中心列求和与公式中x2 x5 x8、x0 x3 x6的行向展开一致。Kernel 入口与 Host 端调用流程EdgeCustom是一个__global__ __aicore__核函数仅 4 行核心代码edge.cppstatic constexpr ATVC::Layout2Dim totalLayout{1023, 2517}; // 原图宽高 template class Traits, const auto totalLayout __global__ __aicore__ void EdgeCustom(GM_ADDR a, GM_ADDR c, ATVC::PoolParam param) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); // 将Edge2C3ComputeFunc仿函数作为模板参数传入实例化PoolOpTemplate模板类 auto op ATVC::Kernel::PoolOpTemplateEdge2C3ComputeFuncTraits, totalLayout(); op.Run(a, c, param); // 按输入、输出、param的顺序传入Run函数实现GM-GM的数据计算 }三个要点KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)声明该核函数仅运行在 AIVVector核上totalLayout以模板参数传入分块、边界裁剪逻辑在编译期即可展开无需运行期传 shapeRun(a, c, param)是模板统一入口PoolOpTemplate内部完成参数校验CheckPoolParam、本核起始 tile 计算、UB 队列初始化Init输入/输出各nBufferNum2个 576 元素的基本块缓冲以及分块主循环。Host 端main()的完整调用链edge.cpp生成1023×2517的随机输入取值 [1, 100)并用 CPU 参考实现计算 golden内部 1 元素边框不参与校验InitializeACL初始化 acl 上下文与 streamexample_common.h调用ATVC::Host::CalcPoolTilingPoolOpTraits(totalLayout, Edge2C3ComputeFuncPoolOpTraits::TILE_LAYOUT, param)计算 tiling 参数并取param.tilingData.blockNum作为启动的 AI Core 数量aclrtMalloc申请 device 内存并拷入输入然后 Kernel 直调EdgeCustomPoolOpTraits, totalLayoutblockNum, nullptr, stream(xDevice, zDevice, param);同步后将输出拷回 host用IsClose绝对/相对容差 1e-3见 example_common.h逐元素比对全部通过则打印Accuracy verification passed.清理 device/host 内存并反初始化 acl。此外edge.cpp中保留了ATVC_DEBUG_MODE 2时的 profiling 分支额外重复运行 19 次以稳定采样配合下文运行脚本的 profiling 模式使用。编译与运行在 ATVC 代码仓根目录下执行来自 README 的算子运行章节cd ./examples bash run_examples.sh edgerun_examples.sh 的工作方式compile_operator校验bisheng编译器可用随 CANN 安装包提供需先设置好环境变量默认以 NPU 模式编译等价命令形如bisheng -x asc --npu-archdav-2201 edge.cpp -o edge \ -I ${ATVC_HOME_DIR}/include -I ${CURRENT_DIR}/common \ -ltiling_api -lplatform -lm -ldl \ -L${ASCEND_INSTALL_PATH}/lib64其中--npu-archdav-2201对应 Ascend 910 系列芯片-I include指向 ATVC 模板头文件目录编译成功后执行./edge二进制退出码为 0 时打印Sample edge passed!可选调试参数--run-modebash run_examples.sh edge --run-modedebug_print附加-DATVC_DEBUG_MODE1开启模板内的DebugPrintf打印分块偏移、padding、tiling 参数等见PrintPoolParam/CalcCurTileScope中的日志bash run_examples.sh edge --run-modeprofiling附加-DATVC_DEBUG_MODE2并用msprof --ai-coreon --ascendclon --runtime-apion --task-timeon采集性能数据。两点说明-I include与-I examples/common使单文件工程即可复用pool/host、pool/kernel全部模板与example_common.h中的 acl 初始化/比对工具运行前提是已安装对应 CANN 软件包并导出ASCEND_HOME_PATH脚本会依次回退到$HOME/Ascend/ascend-toolkit/latest与/usr/local/Ascend/ascend-toolkit/latest。最后需要注意 README 的明确声明当前PoolOpTemplate暂不支持 ATVC 调试调优功能如 tanh_grad 样例中展示的 Tiling 超参调优相关功能待后续补充。因此本样例中CalcPoolTiling内部使用默认的PoolTilingHyperParam如singleCoreBaseLine512、nBufferNum2见 pool_host.h暂不能像其他模板那样通过超参进行性能调优TILE_LAYOUT/TILE_PADDING仍可在满足 32B 对齐约束的前提下自行调整。总结从本样例提炼的 Pool 模板使用模式步骤本样例对应实现依据定义OpTraits输入/输出/临时张量类型OpTraitsOpInputsfloat, OpOutputsfloat, OpTempsint32_tedge.cpp编写 Compute 仿函数声明TILE_LAYOUT、TILE_PADDING并实现块内计算Edge2C3ComputeFunc::operator()edge.cpp定义__global__核函数实例化PoolOpTemplate并调用RunEdgeCustomedge.cppHost 端计算 tiling 并取blockNumATVC::Host::CalcPoolTilingPoolOpTraitspool_host.hKernel 直调blockNum, nullptr, stream并做精度比对main()edge.cpp一键编译运行bash run_examples.sh edge [--run-modedebug_print\|profiling]run_examples.sh对后续开发者的启示当你的算子属于结果依赖固定邻域窗口的类别边缘检测、方向梯度、局部滤波等且邻域宽度可由 padding 覆盖时Pool 模板可以把最繁琐的 tile 调度、越界裁剪和多核数据划分全部模板化你只需保证 UB 基本块TILE_LAYOUT TILE_PADDING足够装下邻域窗口并编写块内的向量计算即可。【免费下载链接】atvcATVCAscend C Templates for Vector Compute是为基于Ascend C开发的典型Vector算子封装的一系列模板头文件的集合可帮助用户快速开发典型Vector算子。项目地址: https://gitcode.com/cann/atvc创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考