企业官网建设流程全解析

CANN ops-cv 图像算子 ExtractImagePatches：滑动窗口图像块提取的原理、参数与实现解析

【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cv

ExtractImagePatches 是 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 中的产品支持矩阵一致：

二、功能说明与计算公式

算子功能：从 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]：扩张率（dilation）；
• pad_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：

从源码看，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 维必须大于 0，rates 的 H/W 维必须大于等于 1。
• padding 取值仅支持 "SAME" 或 "VALID"。
• 输入输出 dtype 相同，无类型提升。
• 纯数据搬运算子，所有特殊值（NaN/Inf/+0/-0）原样透传，SAME padding 越界补零值为 +0.0。

这些约束在源码中有完整落地，详见 extract_image_patches_infershape.cpp：

• 三个 ListInt 属性长度必须为 4，否则报invalid listsize；
• ksizes 的[0]与通道维下标必须为 1；
• strides 的空间维必须> 0，rates 的空间维必须>= 1；
• padding 非 "SAME" 且非 "VALID" 时报invalid value；
• 输入 dtype 只接受 FLOAT16、FLOAT、BF16、INT8、UINT8 五种（tiling 侧 extract_image_patches_tiling.cpp 再次校验）。

另外在 format 上还有一个实现层面的约束：tiling 解析输入时要求 origin NHWC 不能配 storage NCHW（extract_image_patches_tiling.cpp），即 NHWC 输入必须以 NHWC 存储；而 origin NCHW 的输入允许以 NHWC 存储（此时属性列表会自动做[N,C,H,W] -> [N,H,W,C]的重排，见GetAttrByStorageFormat）。

五、输出形状推导（源码级验证）

输出形状由 InferShape4ExtractImagePatches 推导，其规则与 README 描述完全一致：

• 有效卷积核尺寸（考虑扩张）：effFilterH = (kH - 1) * rateH + 1，effFilterW = (kW - 1) * rateW + 1；
• VALID 模式：outH = (inH - effFilterH + strideH) / strideH（整除），若结果为负会被钳制为 0；
• SAME 模式：outH = (inH + strideH - 1) / strideH（向上取整）；
• 输出通道：outC = inC * kH * kW；
• 输出 dtype 与输入 dtype 相同（由 InferDataTypeExtractImagePatches 直接透传）。

形状推导还处理了两类边界情况：

• 输入为 UnknownRank（rank 未知）时，输出同样置为 UnknownRank；
• 输入某维为 -1（动态 shape）时，对应的输出维保持 -1，由推导逻辑原样透传。

在 SAME 模式下，tiling 侧还会计算实际 padding 量：padTotalH = max((outH - 1) * strideH + effFilterH - inH, 0)，padHTop = padTotalH / 2，W 方向同理（ComputeOutputParams）。注意 padding 采用"上/左优先"的分配方式（整数除法向下取整），与 TensorFlow 的 SAME 语义一致。

六、NPU 实现原理：tiling 分块与 SIMT kernel

1. Tiling：输出元素级并行与多核切分

该算子是纯数据搬运算子，采用输出元素级并行策略，每个输出元素由输入的一个元素搬运而来（SAME 越界位置补零）。Host 侧 tiling（extract_image_patches_tiling.cpp）负责：

• 多核切分：perCoreElements = ceil(totalElements / coreNum)，并做两轮调整：低于PER_CORE_MIN_ELEMENTS = 1024时抬升到 1024，再向上对齐到 32（BLOCK_ALIGN），最后按needCoreNum = ceil(totalElements / perCoreElements)决定实际启动核数；
• 索引位宽自适应：输出总元素数不超过INT32_MAX时使用 32 位索引（EXTRACT_IMAGE_PATCHES_TPL_IDX_32BIT），否则回退 64 位索引；同时按存储格式设置formatMode（NHWC/NCHW），两者共同组成 tiling key；
• TilingData 填充：包括输出坐标反推除数（divNhc/divWc用于 NHWC，divNcHw/divHw用于 NCHW）、outC/outW/kW、输入空间维inH/inW、滑动参数stride/rate/pad、输入各维 GM 物理步长，以及每核元素数perCoreElements/lastCoreElements等，结构定义见 extract_image_patches_tiling_data.h；
• Workspace：申请 1 块系统 workspace（GetLibApiWorkSpaceSize）。

2. Kernel：UB 传参与无除法坐标反解

kernel 入口（extract_image_patches.cpp）根据模板参数formatMode（0=NHWC，1=NCHW）与idxBit（0=64 位，1=32 位）实例化。两个 VF（vector function）路径的实现要点（extract_image_patches_simt.h）：

• 所有 tiling 标量参数先写入 UB（FillUbParams/FillUintDivParams），再以__ubuf__指针传给 VF；
• 坐标反解全部使用magic number + shift的免除法整数除法（Simt::UintDiv），5 组除数（NHWC 的divNhc/divWc/outC/inC/kW，NCHW 的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]），否则写 0（SAME 补零）；
• 索引类型为 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]。

七、调用说明

图模式示例的核心片段（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::vector<int64_t> xShape = {1, 3, 3, 1}; // 输入 [N,H,W,C] std::vector<int64_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；
• 自动识别输入 format（NCHW或NHWC），NCHW 时先转置为 NHWC 计算后再转回[N, C, H, W]布局；
• 对属性做与算子约束一致的校验（长度 4、N/C 维为 1、spatial 维正值等）；
• 精度标准：FLOAT16/FLOAT/BF16 使用cross_check（L1 级别容差），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/VALID，dtype 五选一，输出outC = inC * kH * kW且顺序为(p*kW+q)*C+c。

【免费下载链接】ops-cv本项目是CANN提供的图像处理、目标检测相关的算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-cv