CANN ops-cv 图像算子 ExtractImagePatches滑动窗口图像块提取的原理、参数与实现解析【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cvExtractImagePatches 是 CANN ops-cv 算子库中面向图像处理与目标检测场景的算子它从 4D 图像张量中按滑动窗口提取局部图像块patch并把每个 patch 展平拼接到通道维语义上兼容 TensorFlow 的tf.extract_image_patches。本文以 image/extract_image_patches/README.md 为主线结合仓库中的算子定义、形状推导、tiling 与 SIMT kernel 源码完整讲解其功能语义、参数约束、输出形状推导、NPU 上的实现机制与图模式调用方式帮助你快速理解并在自己的模型中正确使用该算子。一、产品支持情况从算子注册配置extract_image_patches_def.cpp可以看到该算子为ascend950平台注册了 AICore 配置与 README 中的产品支持矩阵一致产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品×二、功能说明与计算公式算子功能从 4D 输入图像中按滑动窗口提取图像块patch并将每个 patch 展平拼接到通道维。它属于 im2col 类的数据搬运算子没有浮点计算常用于卷积替代、局部特征聚合等场景。对于 NHWC 格式输入x形状为[N, H, W, C]输出y形状为[N, out_h, out_w, C * kH * kW]计算公式为$$ y[n, i, j, c \cdot kH \cdot kW p \cdot kW q] x[n, i \cdot stride_h p \cdot rate_h - pad_h, j \cdot stride_w q \cdot rate_w - pad_w, c] $$其中kH ksizes[H]kW ksizes[W]patch 高/宽stride_h strides[H]stride_w strides[W]滑动步长rate_h rates[H]rate_w rates[W]扩张率dilationpad_h/pad_w由 padding 模式SAME/VALID决定的 padding 量采样点越界时VALID 模式对应输出位置不存在SAME 模式补零。输出通道维的排列顺序为(p * kW q) * C c即先按 patch 内行p、列q展开再叠加输入通道c。这一顺序在 kernel 源码extract_image_patches_simt.h中通过坐标反推显式实现先由oc反解pq oc / C、c oc % C再由pq反解p pq / kW、q pq % kW。三、参数说明算子共 1 个输入、1 个输出和 4 个属性定义见 extract_image_patches_def.cpp参数名输入/输出/属性描述数据类型数据格式x输入4D 图像输入支持 NHWC 与 NCHW 两种 origin format。公式中的 x。FLOAT16、FLOAT、BF16、INT8、UINT8NDy输出patch 展平到通道维的输出format 与输入 origin format 一致。公式中的 y。FLOAT16、FLOAT、BF16、INT8、UINT8NDksizes属性patch 大小长度 4 的 ListIntN/C 维必须为 1。ListInt-strides属性滑动步长长度 4 的 ListIntN/C 维必须为 1。ListInt-rates属性扩张率dilation长度 4 的 ListIntN/C 维必须为 1。ListInt-padding属性padding 模式取值为 SAME 或 VALID。String-从源码看dtype 与 format 是成对校验的x与y各声明了 10 组(dtype, format)组合即 5 种 dtype × NHWC/NCHW 两种 format且输入输出同 dtype、同 format。此外算子配置了DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)与DynamicCompileStaticFlag(true)支持动态 shape 与动态编译。四、约束说明输入必须为 4D 张量NHWC 或 NCHW 格式。ksizes/strides/rates 长度必须为 4且 N/C 维必须为 1。strides 的 H/W 维必须大于 0rates 的 H/W 维必须大于等于 1。padding 取值仅支持 SAME 或 VALID。输入输出 dtype 相同无类型提升。纯数据搬运算子所有特殊值NaN/Inf/0/-0原样透传SAME padding 越界补零值为 0.0。这些约束在源码中有完整落地详见 extract_image_patches_infershape.cpp三个 ListInt 属性长度必须为 4否则报invalid listsizeksizes 的[0]与通道维下标必须为 1strides 的空间维必须 0rates 的空间维必须 1padding 非 SAME 且非 VALID 时报invalid value输入 dtype 只接受 FLOAT16、FLOAT、BF16、INT8、UINT8 五种tiling 侧 extract_image_patches_tiling.cpp 再次校验。另外在 format 上还有一个实现层面的约束tiling 解析输入时要求 origin NHWC 不能配 storage NCHWextract_image_patches_tiling.cpp即 NHWC 输入必须以 NHWC 存储而 origin NCHW 的输入允许以 NHWC 存储此时属性列表会自动做[N,C,H,W] - [N,H,W,C]的重排见GetAttrByStorageFormat。五、输出形状推导源码级验证输出形状由 InferShape4ExtractImagePatches 推导其规则与 README 描述完全一致有效卷积核尺寸考虑扩张effFilterH (kH - 1) * rateH 1effFilterW (kW - 1) * rateW 1VALID 模式outH (inH - effFilterH strideH) / strideH整除若结果为负会被钳制为 0SAME 模式outH (inH strideH - 1) / strideH向上取整输出通道outC inC * kH * kW输出 dtype 与输入 dtype 相同由 InferDataTypeExtractImagePatches 直接透传。形状推导还处理了两类边界情况输入为 UnknownRankrank 未知时输出同样置为 UnknownRank输入某维为 -1动态 shape时对应的输出维保持 -1由推导逻辑原样透传。在 SAME 模式下tiling 侧还会计算实际 padding 量padTotalH max((outH - 1) * strideH effFilterH - inH, 0)padHTop padTotalH / 2W 方向同理ComputeOutputParams。注意 padding 采用上/左优先的分配方式整数除法向下取整与 TensorFlow 的 SAME 语义一致。六、NPU 实现原理tiling 分块与 SIMT kernel1. Tiling输出元素级并行与多核切分该算子是纯数据搬运算子采用输出元素级并行策略每个输出元素由输入的一个元素搬运而来SAME 越界位置补零。Host 侧 tilingextract_image_patches_tiling.cpp负责多核切分perCoreElements ceil(totalElements / coreNum)并做两轮调整低于PER_CORE_MIN_ELEMENTS 1024时抬升到 1024再向上对齐到 32BLOCK_ALIGN最后按needCoreNum ceil(totalElements / perCoreElements)决定实际启动核数索引位宽自适应输出总元素数不超过INT32_MAX时使用 32 位索引EXTRACT_IMAGE_PATCHES_TPL_IDX_32BIT否则回退 64 位索引同时按存储格式设置formatModeNHWC/NCHW两者共同组成 tiling keyTilingData 填充包括输出坐标反推除数divNhc/divWc用于 NHWCdivNcHw/divHw用于 NCHW、outC/outW/kW、输入空间维inH/inW、滑动参数stride/rate/pad、输入各维 GM 物理步长以及每核元素数perCoreElements/lastCoreElements等结构定义见 extract_image_patches_tiling_data.hWorkspace申请 1 块系统 workspaceGetLibApiWorkSpaceSize。2. KernelUB 传参与无除法坐标反解kernel 入口extract_image_patches.cpp根据模板参数formatMode0NHWC1NCHW与idxBit064 位132 位实例化。两个 VFvector function路径的实现要点extract_image_patches_simt.h所有 tiling 标量参数先写入 UBFillUbParams/FillUintDivParams再以__ubuf__指针传给 VF坐标反解全部使用magic number shift的免除法整数除法Simt::UintDiv5 组除数NHWC 的divNhc/divWc/outC/inC/kWNCHW 的divNcHw/divHw/outW/inC/kW在发射前由GetUintDivMagicAndShift预计算每个线程按threadIdx.x blockDim.x步长遍历本核输出区间先反解出(n, i, j, oc)再反解oc得到(p, q, c)随后计算输入采样位置inHPos i*strideH p*rateH - padHTop、inWPos j*strideW q*rateW - padWLeft采样点落在[0, inH) × [0, inW)内则直接从 GM 搬入outputGm[loopIdx] inputGm[inOffset]否则写 0SAME 补零索引类型为 32 位时单 block 发射 1024 线程64 位时 512 线程THREADS常量Process入口对totalElements 0直接返回。NHWC 与 NCHW 两条路径共享同一个 TilingData 结构体仅按存储格式取用不同字段输出物理布局分别为[N, out_h, out_w, out_c]与[N, out_c, out_h, out_w]。七、调用说明调用方式调用样例说明图模式调用test_geir_extract_image_patches参见 算子调用 完成算子编译和验证。图模式示例的核心片段CreateOppInGraph给出了一个最小可用配置auto add1 op::ExtractImagePatches(extract1); add1.set_attr_ksizes({1, 2, 2, 1}); // patch 2x2 add1.set_attr_strides({1, 1, 1, 1}); // 步长 1 add1.set_attr_rates({1, 1, 1, 1}); // 无扩张 add1.set_attr_padding(VALID); std::vectorint64_t xShape {1, 3, 3, 1}; // 输入 [N,H,W,C] std::vectorint64_t yShape {1, 2, 2, 4}; // 输出 [1,2,2,4] [N,2,2,C*kH*kW]即输入[1,3,3,1]、VALID 模式下 2×2 patch、步长 1 时输出为[1, 2, 2, 4]outC 1*2*2 4与公式一致。示例随后通过 GE 会话完成构图、dump、运行与输入输出 bin 文件落盘可配合 算子调用 文档中的编译验证流程运行。八、正确性验证与 TensorFlow 的交叉对照仓库在 tests/assets/golden.py 中提供了完整的黄金数据golden生成逻辑直接以 TensorFlow 为参考实现做交叉验证优先调用tf.extract_image_patches旧版本回退到tf.image.extract_patches自动识别输入 formatNCHW或NHWCNCHW 时先转置为 NHWC 计算后再转回[N, C, H, W]布局对属性做与算子约束一致的校验长度 4、N/C 维为 1、spatial 维正值等精度标准FLOAT16/FLOAT/BF16 使用cross_checkL1 级别容差INT8/UINT8 使用binary_equal逐字节精确比对——这从侧面印证了 README 中纯数据搬运算子、无类型提升、特殊值原样透传的语义。此外tests/ut/op_host/test_extract_image_patches_infershape.cpp 与 tests/ut/op_host/arch35/test_extract_image_patches_tiling.cpp 分别覆盖了形状推导与 tiling 计算两个 Host 侧环节的单元测试可作为实现行为的补充依据。九、框架适配extract_image_patches_tf_plugin.cpp 中注册了 TensorFlow 框架的算子映射OriginOpType(ExtractImagePatches)自动映射到本算子的输入输出与属性并将输入输出的 origin format 与存储 format 统一设置为FORMAT_NHWC。这意味着来自 TensorFlow 的ExtractImagePatches节点在 CANN 图上会以 NHWC 布局落盘与 kernel 的 NHWC 路径直接对接。小结ExtractImagePatches 是 ops-cv 中一个语义简单但实现颇具代表性的数据搬运算子语义上严格对齐 TensorFlow支持 NHWC/NCHW 双布局与 SAME/VALID 双 padding支持动态 shape实现上以输出元素级并行 SIMT VF 为核心通过 magic number 除法、UB 标量透传、多核分块与 32/64 位索引自适应来逼近数据搬运的极限效率。理解该算子也等于理解了 ops-cv 中一批 im2col 类算子的通用实现套路。实际使用时请牢记4D 输入、属性长度 4 且 N/C 维为 1、padding 仅限 SAME/VALIDdtype 五选一输出outC inC * kH * kW且顺序为(p*kWq)*Cc。【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cv创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN ops-cv 图像算子 ExtractImagePatches：滑动窗口图像块提取的原理、参数与实现解析
CANN ops-cv 图像算子 ExtractImagePatches滑动窗口图像块提取的原理、参数与实现解析【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cvExtractImagePatches 是 CANN ops-cv 算子库中面向图像处理与目标检测场景的算子它从 4D 图像张量中按滑动窗口提取局部图像块patch并把每个 patch 展平拼接到通道维语义上兼容 TensorFlow 的tf.extract_image_patches。本文以 image/extract_image_patches/README.md 为主线结合仓库中的算子定义、形状推导、tiling 与 SIMT kernel 源码完整讲解其功能语义、参数约束、输出形状推导、NPU 上的实现机制与图模式调用方式帮助你快速理解并在自己的模型中正确使用该算子。一、产品支持情况从算子注册配置extract_image_patches_def.cpp可以看到该算子为ascend950平台注册了 AICore 配置与 README 中的产品支持矩阵一致产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品×二、功能说明与计算公式算子功能从 4D 输入图像中按滑动窗口提取图像块patch并将每个 patch 展平拼接到通道维。它属于 im2col 类的数据搬运算子没有浮点计算常用于卷积替代、局部特征聚合等场景。对于 NHWC 格式输入x形状为[N, H, W, C]输出y形状为[N, out_h, out_w, C * kH * kW]计算公式为$$ y[n, i, j, c \cdot kH \cdot kW p \cdot kW q] x[n, i \cdot stride_h p \cdot rate_h - pad_h, j \cdot stride_w q \cdot rate_w - pad_w, c] $$其中kH ksizes[H]kW ksizes[W]patch 高/宽stride_h strides[H]stride_w strides[W]滑动步长rate_h rates[H]rate_w rates[W]扩张率dilationpad_h/pad_w由 padding 模式SAME/VALID决定的 padding 量采样点越界时VALID 模式对应输出位置不存在SAME 模式补零。输出通道维的排列顺序为(p * kW q) * C c即先按 patch 内行p、列q展开再叠加输入通道c。这一顺序在 kernel 源码extract_image_patches_simt.h中通过坐标反推显式实现先由oc反解pq oc / C、c oc % C再由pq反解p pq / kW、q pq % kW。三、参数说明算子共 1 个输入、1 个输出和 4 个属性定义见 extract_image_patches_def.cpp参数名输入/输出/属性描述数据类型数据格式x输入4D 图像输入支持 NHWC 与 NCHW 两种 origin format。公式中的 x。FLOAT16、FLOAT、BF16、INT8、UINT8NDy输出patch 展平到通道维的输出format 与输入 origin format 一致。公式中的 y。FLOAT16、FLOAT、BF16、INT8、UINT8NDksizes属性patch 大小长度 4 的 ListIntN/C 维必须为 1。ListInt-strides属性滑动步长长度 4 的 ListIntN/C 维必须为 1。ListInt-rates属性扩张率dilation长度 4 的 ListIntN/C 维必须为 1。ListInt-padding属性padding 模式取值为 SAME 或 VALID。String-从源码看dtype 与 format 是成对校验的x与y各声明了 10 组(dtype, format)组合即 5 种 dtype × NHWC/NCHW 两种 format且输入输出同 dtype、同 format。此外算子配置了DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)与DynamicCompileStaticFlag(true)支持动态 shape 与动态编译。四、约束说明输入必须为 4D 张量NHWC 或 NCHW 格式。ksizes/strides/rates 长度必须为 4且 N/C 维必须为 1。strides 的 H/W 维必须大于 0rates 的 H/W 维必须大于等于 1。padding 取值仅支持 SAME 或 VALID。输入输出 dtype 相同无类型提升。纯数据搬运算子所有特殊值NaN/Inf/0/-0原样透传SAME padding 越界补零值为 0.0。这些约束在源码中有完整落地详见 extract_image_patches_infershape.cpp三个 ListInt 属性长度必须为 4否则报invalid listsizeksizes 的[0]与通道维下标必须为 1strides 的空间维必须 0rates 的空间维必须 1padding 非 SAME 且非 VALID 时报invalid value输入 dtype 只接受 FLOAT16、FLOAT、BF16、INT8、UINT8 五种tiling 侧 extract_image_patches_tiling.cpp 再次校验。另外在 format 上还有一个实现层面的约束tiling 解析输入时要求 origin NHWC 不能配 storage NCHWextract_image_patches_tiling.cpp即 NHWC 输入必须以 NHWC 存储而 origin NCHW 的输入允许以 NHWC 存储此时属性列表会自动做[N,C,H,W] - [N,H,W,C]的重排见GetAttrByStorageFormat。五、输出形状推导源码级验证输出形状由 InferShape4ExtractImagePatches 推导其规则与 README 描述完全一致有效卷积核尺寸考虑扩张effFilterH (kH - 1) * rateH 1effFilterW (kW - 1) * rateW 1VALID 模式outH (inH - effFilterH strideH) / strideH整除若结果为负会被钳制为 0SAME 模式outH (inH strideH - 1) / strideH向上取整输出通道outC inC * kH * kW输出 dtype 与输入 dtype 相同由 InferDataTypeExtractImagePatches 直接透传。形状推导还处理了两类边界情况输入为 UnknownRankrank 未知时输出同样置为 UnknownRank输入某维为 -1动态 shape时对应的输出维保持 -1由推导逻辑原样透传。在 SAME 模式下tiling 侧还会计算实际 padding 量padTotalH max((outH - 1) * strideH effFilterH - inH, 0)padHTop padTotalH / 2W 方向同理ComputeOutputParams。注意 padding 采用上/左优先的分配方式整数除法向下取整与 TensorFlow 的 SAME 语义一致。六、NPU 实现原理tiling 分块与 SIMT kernel1. Tiling输出元素级并行与多核切分该算子是纯数据搬运算子采用输出元素级并行策略每个输出元素由输入的一个元素搬运而来SAME 越界位置补零。Host 侧 tilingextract_image_patches_tiling.cpp负责多核切分perCoreElements ceil(totalElements / coreNum)并做两轮调整低于PER_CORE_MIN_ELEMENTS 1024时抬升到 1024再向上对齐到 32BLOCK_ALIGN最后按needCoreNum ceil(totalElements / perCoreElements)决定实际启动核数索引位宽自适应输出总元素数不超过INT32_MAX时使用 32 位索引EXTRACT_IMAGE_PATCHES_TPL_IDX_32BIT否则回退 64 位索引同时按存储格式设置formatModeNHWC/NCHW两者共同组成 tiling keyTilingData 填充包括输出坐标反推除数divNhc/divWc用于 NHWCdivNcHw/divHw用于 NCHW、outC/outW/kW、输入空间维inH/inW、滑动参数stride/rate/pad、输入各维 GM 物理步长以及每核元素数perCoreElements/lastCoreElements等结构定义见 extract_image_patches_tiling_data.hWorkspace申请 1 块系统 workspaceGetLibApiWorkSpaceSize。2. KernelUB 传参与无除法坐标反解kernel 入口extract_image_patches.cpp根据模板参数formatMode0NHWC1NCHW与idxBit064 位132 位实例化。两个 VFvector function路径的实现要点extract_image_patches_simt.h所有 tiling 标量参数先写入 UBFillUbParams/FillUintDivParams再以__ubuf__指针传给 VF坐标反解全部使用magic number shift的免除法整数除法Simt::UintDiv5 组除数NHWC 的divNhc/divWc/outC/inC/kWNCHW 的divNcHw/divHw/outW/inC/kW在发射前由GetUintDivMagicAndShift预计算每个线程按threadIdx.x blockDim.x步长遍历本核输出区间先反解出(n, i, j, oc)再反解oc得到(p, q, c)随后计算输入采样位置inHPos i*strideH p*rateH - padHTop、inWPos j*strideW q*rateW - padWLeft采样点落在[0, inH) × [0, inW)内则直接从 GM 搬入outputGm[loopIdx] inputGm[inOffset]否则写 0SAME 补零索引类型为 32 位时单 block 发射 1024 线程64 位时 512 线程THREADS常量Process入口对totalElements 0直接返回。NHWC 与 NCHW 两条路径共享同一个 TilingData 结构体仅按存储格式取用不同字段输出物理布局分别为[N, out_h, out_w, out_c]与[N, out_c, out_h, out_w]。七、调用说明调用方式调用样例说明图模式调用test_geir_extract_image_patches参见 算子调用 完成算子编译和验证。图模式示例的核心片段CreateOppInGraph给出了一个最小可用配置auto add1 op::ExtractImagePatches(extract1); add1.set_attr_ksizes({1, 2, 2, 1}); // patch 2x2 add1.set_attr_strides({1, 1, 1, 1}); // 步长 1 add1.set_attr_rates({1, 1, 1, 1}); // 无扩张 add1.set_attr_padding(VALID); std::vectorint64_t xShape {1, 3, 3, 1}; // 输入 [N,H,W,C] std::vectorint64_t yShape {1, 2, 2, 4}; // 输出 [1,2,2,4] [N,2,2,C*kH*kW]即输入[1,3,3,1]、VALID 模式下 2×2 patch、步长 1 时输出为[1, 2, 2, 4]outC 1*2*2 4与公式一致。示例随后通过 GE 会话完成构图、dump、运行与输入输出 bin 文件落盘可配合 算子调用 文档中的编译验证流程运行。八、正确性验证与 TensorFlow 的交叉对照仓库在 tests/assets/golden.py 中提供了完整的黄金数据golden生成逻辑直接以 TensorFlow 为参考实现做交叉验证优先调用tf.extract_image_patches旧版本回退到tf.image.extract_patches自动识别输入 formatNCHW或NHWCNCHW 时先转置为 NHWC 计算后再转回[N, C, H, W]布局对属性做与算子约束一致的校验长度 4、N/C 维为 1、spatial 维正值等精度标准FLOAT16/FLOAT/BF16 使用cross_checkL1 级别容差INT8/UINT8 使用binary_equal逐字节精确比对——这从侧面印证了 README 中纯数据搬运算子、无类型提升、特殊值原样透传的语义。此外tests/ut/op_host/test_extract_image_patches_infershape.cpp 与 tests/ut/op_host/arch35/test_extract_image_patches_tiling.cpp 分别覆盖了形状推导与 tiling 计算两个 Host 侧环节的单元测试可作为实现行为的补充依据。九、框架适配extract_image_patches_tf_plugin.cpp 中注册了 TensorFlow 框架的算子映射OriginOpType(ExtractImagePatches)自动映射到本算子的输入输出与属性并将输入输出的 origin format 与存储 format 统一设置为FORMAT_NHWC。这意味着来自 TensorFlow 的ExtractImagePatches节点在 CANN 图上会以 NHWC 布局落盘与 kernel 的 NHWC 路径直接对接。小结ExtractImagePatches 是 ops-cv 中一个语义简单但实现颇具代表性的数据搬运算子语义上严格对齐 TensorFlow支持 NHWC/NCHW 双布局与 SAME/VALID 双 padding支持动态 shape实现上以输出元素级并行 SIMT VF 为核心通过 magic number 除法、UB 标量透传、多核分块与 32/64 位索引自适应来逼近数据搬运的极限效率。理解该算子也等于理解了 ops-cv 中一批 im2col 类算子的通用实现套路。实际使用时请牢记4D 输入、属性长度 4 且 N/C 维为 1、padding 仅限 SAME/VALIDdtype 五选一输出outC inC * kH * kW且顺序为(p*kWq)*Cc。【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cv创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考