Flash-DWC：在GPU上实现深度卷积的高效计算

《ACM Transactions on Architecture and Code Optimization》：Flash-DWC: Making Depthwise Convolution Compute-Efficient on GPUs

　　

摘要

深度卷积（DWC）是高效卷积神经网络（CNN）中的关键算子。近期模型越来越多地采用大尺寸滤波器DWC来捕获长程依赖关系，实现了与视觉Transformer相媲美的精度。然而，在GPU上，DWC的实际效率远低于其低FLOPs所暗示的水平，可能成为运行时瓶颈，原因在于算术强度低、与Tensor Core不匹配以及内存访问效率低下。为解决这些问题，我们提出了Flash-DWC，一个使DWC计算高效的GPU库。Flash-DWC引入了一种基于GEMV的数据流协同设计，联合优化数据复用、访问顺序、布局和任务分块。该设计提高了算术强度，同时保持了访问连续性和128字节合并内存访问。Flash-DWC进一步将DWC的GEMV转化为紧凑的高密度GEMM，运行在Tensor Core上，实现高效的混合精度加速。此外，矩阵压缩、双路径分拆GEMM流水线和基于寄存器的输入复用进一步释放片上资源、隐藏内存延迟并提高占用率。我们将这些技术扩展到前向和反向传播，以实现高效的端到端训练。与最快的cuDNN DWC算法相比，Flash-DWC的基于GEMV的CUDA Core内核和基于GEMM的Tensor Core内核分别实现了5.09倍和9.01倍的平均加速。在CNN训练中，Flash-DWC获得了1.03×至14.18倍的加速，并在ImageNet-1K上与PyTorch的收敛性相匹配。通过使DWC计算高效，Flash-DWC促进了更大、更具表达力滤波器的使用。

AI摘要

AI生成摘要（实验性）

本摘要由自动化工具生成，未经论文作者撰写或审核。此摘要旨在支持文献发现，帮助读者评估相关性，并协助来自相邻研究领域的读者理解本工作。本摘要作为作者提供的摘要的补充，作者摘要仍是论文的主要摘要。全文仍是权威版本。点击此处了解更多信息。点击此处评论本摘要的准确性、清晰度和实用性。这样做将有助于改进和未来的重新生成版本。

1 引言

对于卷积神经网络（CNN）而言，在保持表征能力的同时降低计算成本一直是一个长期目标。为此，许多高效CNN [12, 13, 39, 40] 用深度卷积（DWC）替换了标准卷积，大幅降低FLOPs的同时保持了空间感受野。自2022年以来，新一波的架构 [15, 25, 26, 29, 41, 44, 46, 47, 48] 进一步利用大滤波器DWC（例如7×7、31×31）来捕获长程依赖关系，在更低的FLOPs下实现了与视觉Transformer [9, 22, 23, 43] 相当甚至更优的精度。除ImageNet分类之外，大滤波器DWC在目标检测和语义分割等下游任务上提供了更大的增益 [2, 7, 8, 17, 20]，突显了其作为可迁移骨干网络的价值。

然而，DWC的低FLOPs并不能保证在GPU上获得成比例的加速。随着DWC FLOPs随滤波器大小的二次方增长，这一问题变得更加突出。我们的PyTorch性能分析显示，对于评估的滤波器大于7×7的工作负载 [7, 20, 25, 41]，DWC算子可以占用每批次训练时间的一半以上。这种运行时成本限制了更大、更具表达力滤波器的实际应用。近期工作 [8, 17, 20] 通过使用分解或稀疏滤波器来缓解此问题，但此类近似往往会牺牲表征能力，而密集滤波器通常仍然更准确 [2]。

这种效率低下源于DWC与GPU架构之间的根本性不匹配：（i）算术强度低：每个输入元素参与的计算量远低于标准卷积，导致计算单元未被充分利用；（ii）Tensor Core不匹配：DWC本质上作为一系列小型GEMV的集合运行，而非大型GEMM，使其难以利用高吞吐量的矩阵单元；（iii）内存访问效率低：DWC仅在空间维度上进行累加，破坏了通道方向的连续性，降低了内存合并访问能力。因此，PyTorch [36] 和cuDNN [4] 等主流库仅利用了GPU计算容量的一小部分来执行DWC。先前的优化部分缓解了这些问题，但局限性仍然存在。基于CUDA Core的方法 [27, 28, 32, 34, 37] 在混合精度吞吐量方面受到限制；基于GEMM或Tensor Core的方法 [32, 34, 38] 往往会引入冗余或稀疏模式。此外，一些方法仅针对较窄的设置（如小滤波器）进行了调优，或未能充分利用软件流水线和内存合并访问。

我们没有重新设计模型，而是提出一个问题：DWC本身能否在GPU上实现计算高效？我们的核心洞察是，加速DWC的关键不在于减少FLOPs，而在于隐藏内存延迟以充分饱和计算单元。为此，我们提出了Flash-DWC，一个以统一方式解决DWC瓶颈的GPU库：（i）基于GEMV的数据流协同设计联合优化数据复用、访问顺序、布局和任务分块，从而提高算术强度的同时保持访问连续性和128字节合并访问；（ii）紧凑变换将DWC的GEMV映射为高密度GEMM，实现高效的Tensor Core利用。结合资源高效的核函数优化，这些技术将DWC从内存瓶颈转变为计算瓶颈，并支持完整的训练流水线。Flash-DWC集成于PyTorch，源代码可在我们的GitHub仓库中获得。

我们在RTX 4090、A100 80GB和RTX 5880 Ada GPU上评估了Flash-DWC。与最快的cuDNN DWC算法相比，基于GEMV的CUDA Core内核实现了5.09倍的平均加速；基于GEMM的Tensor Core内核进一步实现了9.01倍的平均加速（相对于cuDNN），以及在CUDA Core内核基础上2.01倍的加速。在端到端训练中，Flash-DWC获得了1.03×至14.18倍的加速，同时在ImageNet-1K [6] 上匹配PyTorch的收敛性。

我们的贡献总结如下：

——基于GEMV的数据流协同设计。我们利用DWC GEMV中的滤波器复用以及跨重叠补丁的宽度优先输入复用。这将算术强度从朴素DWC的1提高到CUDA Core上的超过15，在GEMV到GEMM变换后超过40，同时保持访问连续性和内存合并访问。

——紧凑的GEMV到GEMM变换。我们以56.25%的密度将DWC GEMV映射为GEMM，是朴素填充的4.5倍。该变换支持多通道、大滤波器工作负载，并与我们的GEMV协同设计无缝集成，在Tensor Core上实现混合精度加速。

——资源高效的核函数优化。对于Tensor Core内核，双路径分拆GEMM流水线配合在线矩阵压缩，将50%的计算与数据加载重叠，同时将共享内存使用量大幅削减以提高占用率和L1容量。对于CUDA Core内核，基于寄存器的输入复用同样最小化共享内存使用量以释放片上内存。

——优化的训练流水线。针对滤波器梯度计算中低并行度和不规则输出形状的问题，我们通过累加轴并行化和灵活分块来解决，同时混合累加加速混合精度执行。核函数融合减少了小滤波器的数据移动。

2 背景与设计原则

为了使我们的优化有立足之地，我们首先形式化描述标准和深度卷积，分析DWC的瓶颈，然后确立我们的初步设计选择。

2.1 标准和深度卷积

为了形式化描述二维卷积，我们定义以下符号：N表示批大小；IH、IW和IC分别表示输入高度、宽度和通道数；OH、OW和OC分别表示输出高度、宽度和通道数；FH和FW分别表示滤波器高度和宽度。

标准二维卷积遵循全对全模式。它将输入特征图 X∈R^(N×IH×IW×IC) 与滤波器 W∈R^(OC×FH×FW×IC) 卷积以生成输出特征图 Y∈R^(N×OH×OW×OC)。在所有 N×OH×OW 个批量-空间位置上，每个输入补丁与每个 OC 滤波器执行一次 FH×FW×IC 内积。这允许通过Im2col自然地映射到通用矩阵乘法（GEMM），其中一个矩阵包含所有展开的输入补丁，另一个包含所有展开的滤波器 [4]。

深度二维卷积（DWC）遵循全对一模式。它将 X 与滤波器 W∈R^(OC×FH×FW) 卷积以生成 Y。每个输入通道都有其专用的滤波器对应输出通道，因此 IC等于OC。在所有批量-空间位置上，每个输入补丁与其专用滤波器执行一次 FH×FW 内积。因此，每个通道的操作自然地映射到通用矩阵-向量乘法（GEMV），如图1所示，其中矩阵包含该通道中所有展开的输入补丁，向量是展开的滤波器。

图1. 深度卷积。左侧展示了一个三通道的DWC：每个通道的操作可以通过Im2col转换为GEMV。右侧详细说明了通道0的Im2col变换：朴素地将每个展开的滤波器片段填充为8×4矩阵以用于Tensor Core GEMM会产生至少78%的稀疏度；虚线框表示对应前4个滤波器元素的GEMM。

2.2 深度卷积的性能瓶颈

我们从三个方面分析DWC的性能瓶颈。

算术强度。标准卷积（GEMM）在 O(n2) 数据上执行 O(n3) 计算，将每个矩阵行在所有列方向上复用。相比之下，DWC（GEMV）在 O(n2) 数据上执行 O(n2) 计算，每个矩阵行仅使用一次。这种数据复用的根本差异导致了算术强度的巨大差距。具体来说，DWC和标准卷积的理论上限分别为 η_D = (2×N×OH×OW×OC×FH×FW) / (N×IH×IW×IC + OC×FH×FW) 和 η_S = (2×N×OH×OW×OC×FH×FW×IC) / (N×IH×IW×IC + OC×FH×FW×IC)。这些简化为 η_D ≈ (1/IC)×η_S，因为分辨率和批大小总是很大的。由于CNN中IC通常很大（例如64到1024），DWC的算术强度比标准卷积低数个数量级。这量化了DWC为何仍然受内存限制且无法饱和计算单元。

Tensor Core不匹配。在现代GPU上，Tensor Core是用于混合精度GEMM的专用单元，提供比通用CUDA Core高得多的峰值吞吐量。Tensor Core要求GEMM维度为 m≥8、n≥8 和 k≥4。标准卷积通过其具有大通道数的GEMM结构自然地满足这些约束。然而，DWC作为一系列GEMV的集合运行，每个GEMV对于Tensor Core的GEMM维度来说太小。如图1所示，朴素地将每个展开的滤波器片段填充到8×4满足最小m8n8k4 GEMM形状，但浪费了7/8的矩阵列。这带来了大量不必要的计算、冗余寄存器、共享内存使用量和潜在的warp发散。因此，朴素Tensor Core映射的DWC相比CUDA Core无法获得有意义的加速。这促成了我们在第4节中提出的高效GEMV到GEMM变换。

内存访问。在标准卷积中，累加发生在 FH×FW×IC 维度上，这些维度完全连续且可以合并访问。例如，在NHWC布局中，IC始终是32的倍数，因此可以沿IC轴以128字节对W和X进行向量化加载，最大化内存事务粒度。相比之下，DWC仅在 FH×FW 上累加。在NHWC格式中，空间位置与IC步幅交错，破坏了访问连续性。此外，在NCHW和NHWC布局中，W的维度顺序为 OC×FH×FW，其中FH和FW通常为奇数且因此不对齐到128字节边界。因此，DWC不太能充分利用内存带宽。

2.3 设计原则

为了克服这些瓶颈，我们的目标是隐藏内存延迟并饱和计算单元，遵循内存感知核函数中的通用原则 [5, 10]。这转化为两个方向：（i）最大化片上数据复用以最小化片外内存流量；（ii）通过高效布局和访问模式最大化带宽。为实现这些目标，我们首先选择卷积方法和张量布局。

直接卷积方法。这指的是直接计算DWC，而不是使用FFT或Winograd域变换。虽然FFT [31] 和Winograd [16] 可以减少FLOPs，但它们引入了额外的变换、中间存储、数据移动和有限精度舍入。这种权衡对DWC不利：减少的计算提供边际收益，而额外的数据移动加剧了主要的内存瓶颈并可能抵消FLOP节省，最终降低吞吐量。直接卷积保持了最直接的数据流，允许我们将片上资源（如寄存器、共享内存）专门用于最大化数据复用和带宽。

统一的通道最后布局。我们将滤波器从常规布局 W∈R^(OC×FH×FW) 重排为 W∈R^(FH×FW×OC)；我们将输入和输出特征图存储为 X∈R^(N×IH×IW×IC) 和 Y∈R^(N×OH×OW×OC)。对于DWC，IC=OC，因此所有张量都将通道轴放在最后。这使得滤波器和特征图都能沿通道轴进行合并访问。由于X和Y保持NHWC布局，Flash-DWC与主流库的相邻操作（如归一化和逐点卷积）自然互操作。

3 基于GEMV的深度卷积数据流协同设计

基于直接卷积和统一的通道最后布局，我们首先在CUDA Core上以FP32精度优化DWC的原始GEMV形式，然后再扩展到混合精度Tensor Core。这种分离是刻意的：它隔离了GEMV优化的贡献，并为后续的GEMV到GEMM变换奠定基础。

我们的设计以数据复用为中心，但其有效性依赖于布局、任务分块和访问顺序：（i）任务分块创建复用机会，并利用统一的通道最后布局实现128字节访问；（ii）滤波器复用利用跨空间位置共享的滤波器；（iii）宽度优先输入复用保持跨重叠补丁的访问连续性。

图2展示了我们将朴素DWC1的算术强度从1提升到超过15的过程。图2. 基于GEMV的Flash-DWC内核的算术强度（η，操作数/元素）。(a) 每一步优化后强度（η）的逐步增长。(b) 最终强度与滤波器宽度FW的关系。

3.1 工作负载分块

我们将输出张量Y∈R^(N×OH×OW×OC) 分块为BN个连续通道和BM个批次-空间位置，总计产生(OC/BN)×(NOH×OW/BM)个块。由于OC与NOH×OW的乘积比其各自因子变化更小，这种分块方式在不同层之间保持稳定的块数。

我们配置每个块使用4个warp，以饱和每个流式多处理器（SM）的4个warp调度器。在warp内部，我们将32个线程排列为8×4的模式：(i) 8个线程使用128位指令执行128字节合并加载，读取32个连续通道，最大化内存事务粒度；(ii) 4个线程获取相邻的输入补丁，创建输入复用的重叠区域。为了最大化warp之间的输入补丁重叠，我们将BN设为32，恰好是每个warp的通道数，因此所有4个warp完全跨越批次-空间维度。

3.2 通过GEMV实现滤波器复用

我们利用DWC的GEMV结构来复用滤波器数据。如图3(a)所示，每个线程被分配跨4个通道的TM个连续空间位置的补丁。在每个通道内，所有TM个位置共享相同的滤波器，因此加载一个滤波器元素即可在相应的输入元素上执行TM次融合乘加（FMA）操作。这将算术强度提升到2TM/(TM+1)（>1），并随TM增长。过大的TM会减少活跃线程数，损害延迟隐藏效果，因此我们经验性地将TM设为8，得到算术强度≈1.78。

图3. GEMV中单个通道内的滤波器复用。(a) 在一个线程中，一个滤波器元素在TM个不同空间位置的内积中被复用；xi,j(t)表示第t个内积中滤波器偏移(i,j)处的输入元素。(b) 在一个块中，512/BN个线程协作将一个滤波器片段加载到SMEM中，然后复用共享片段来计算各自的TM个内积。

为了进一步提升滤波器复用，我们采用块级协作。如图3(b)所示，对于每个通道，所有512/BN个线程共享相同的滤波器。因此，它们协作将一个滤波器片段加载到共享内存（SMEM）中，从而将许多线程级GEMV合并为一个更大的块级GEMV。算术强度随后提升到2TM/(TM+BN/512)。当BN=32且TM=8时，这大约为1.98，接近理论最大值2.3。

3.3 跨重叠补丁的输入复用

虽然滤波器复用已达到最大化，但输入复用仍然是主要瓶颈。与滤波器不同，滤波器占用空间小且更易于缓存，输入补丁随空间维度和批次大小增长，因此更难保留在片上内存中。幸运的是，相邻的输入补丁存在相当程度的重叠，特别是在步长为1的卷积中。现代CNN通常仅在阶段边界处使用相对少量的步长为2的卷积进行下采样，且往往使用较小的滤波器。在每个阶段内，反复出现的步长为1的DWC块保持空间分辨率不变，在同一尺度上细化特征，并通常使用较大的滤波器来扩展感受场。因此，步长为1的DWC主导了许多CNN工作负载，使得输入重叠复用具有广泛适用性。

在每个通道内，我们评估了图4中展示的两项复用方法。(i) 矩形复用同时利用高度和宽度方向的重叠，提供更高的理论复用率。然而，它从间隔为IC×IW的不连续行中加载元素，且操作模式在各迭代间变化。跨高度轴迭代复用数据需要O(TM×FW)的额外存储来保留多行数据，或者需要重排维度的遍历顺序，这会破坏内存连续性。(ii) 宽度优先复用虽然仅限于宽度轴的重叠，但保持完全连续的内存访问且对缓存友好。在初始批量加载后，后续迭代以统一方式操作，持有的数据简单地移位并复用于下一次迭代。

图4. 单个通道内的矩形和宽度优先复用方法。此示例展示了3×3滤波器的第0步到第8步的累加过程。蓝色表示当前迭代加载的元素；灰色表示从上一迭代复用的元素；黄色标记本迭代持有但未使用的元素。

基于我们的配置TM=8，宽度优先复用加载Nw=FH(FW+8?1)个输入元素，而矩形复用（2×4）加载Nr=(FH+2?1)(FW+4?1)个元素，并且需要至少多25%的存储空间用于数据复用。对于FH,FW≥3，Nw/Nr的比率最多为1.25，意味着宽度优先复用最多产生25%的额外加载量，同时保持了简单性和访问连续性。因此我们选择宽度优先复用。为实现它，我们将任意2D滤波器∈R^(FH×FW)分解为FH个1D单元∈R^(FW)，用每个单元执行1D DWC，然后累加结果。这种方法可以推广到高维DWC，将其视为一系列1D DWC的序列。

我们使用寄存器而非SMEM来实现宽度优先复用。寄存器提供最低延迟的操作数访问并直接为计算流水线供数，而SMEM有约30个周期的延迟且与L1缓存共享资源。最小化SMEM使用可保留L1容量并允许每个SM运行更多活跃块，从而提升占用率。由于每个线程维护私有的输入复用窗口，使用SMEM暂存会引入额外的加载/存储指令和地址计算开销。由于每个线程每通道处理TM个相邻的输入补丁，我们分配4TM个寄存器用于累加，另分配4TM个寄存器缓存输入元素。这种寄存器分配恰好满足每个GEMV的需求，并与cuBLAS GEMM内核的寄存器占用量相匹配，维持高占用率。

3.4 基于GEMV的内核工作流程

我们将上述设计整合为一个连贯的工作流程，如图1算法所示。主循环在两个关键函数Prefetch和GEMV之间交替。

**Prefetch（预取）。** 每个块加载一个FW×BN的滤波器块和一个BM×BN的输入块，其中每个线程加载TM×4个输入元素，跨越TM个空间位置的4个通道。滤波器加载使用cp.async.cg PTX指令，该指令发出带有显式提交/等待控制的异步全局到SMEM拷贝。这避免了寄存器中转暂存，实现了更细粒度的软件流水线，性能比标准加载提升约8%。对滤波器SMEM数组W进行双缓冲，将Prefetch与GEMV重叠执行，再获得约5%的性能提升。为避免银行冲突，我们将W的最后一个维度填充128位。

**GEMV。** 每个分解后的1D DWC以FW次累加迭代的GEMV形式执行，这些迭代通过C++模板在编译时展开。从第二次迭代开始，每个线程从其寄存器缓存X中复用4(TM?1)个元素并仅加载4个新的输入元素，同样每次迭代执行TM次FMA。结合滤波器复用，算术强度变为2TM/(1+(TM?1)/FW+BN/512)，随FW增长，如图2(b)所示。当BN=32且TM=8时，随着FW→+∞，强度趋近于15.06。

对于CUDA-Core路径，当FW≤51时，缓存滤波器数据使用的SMEM少于可用SMEM的10%，将大部分片上内存留给L1缓存。Nsight测量表明，当FW≤9时L1缓存命中率可超过80%，且对于FW≤51保持高于50%。由于L1带宽比全局内存高10倍以上，实际算术强度显著高于仅从全局内存访问次数估算的值。

4 基于GEMM的混合精度深度卷积

虽然基于GEMV的优化显著提升了FP32吞吐量，但性能最终受限于CUDA-Core峰值。为实现更高效的高混合精度DWC，我们利用Tensor Core，这需要根本性的转变：将DWC原生的GEMV转换为GEMM。

基于第3节中我们的基于GEMV的设计，我们将其适配为FP16/FP32混合精度，这是DNN训练中的标准实践。该扩展依赖于三项关键创新：(i) 一种56.25%稠密度的GEMV到GEMM变换，专为DWC的多通道、大滤波器设置量身定制；(ii) 通过矩阵压缩和宽度优先输入复用最小化GEMM冗余，大幅削减SMEM使用并提升算术强度；(iii) 一种双路径分割GEMM流水线，将50%的计算与数据加载重叠，同时保持复用模式。图5展示了我们将算术强度从1提升到超过40的过程。

图5. 基于GEMM的Flash-DWC的算术强度（?，操作数/元素）。(a) 每一步优化后强度（?）的逐步增长。(b) 最终强度与滤波器宽度FW的关系：Flash-DWC表示基于GEMV的CUDA-Core内核，Flash-DWC-T1、-T2和-T4是基于GEMM的Tensor-Core内核。

4.1 面向Tensor Core

FP16/FP32混合精度以FP16存储张量并执行乘法以提高效率，同时以FP32累加以保证数值稳定性。对于DWC，FP32累加至关重要：由于精度和动态范围有限，纯FP16累加可能导致不可接受的精度损失。当张量从均匀分布U(?1,1)中采样时，FP16累加DWC的平均绝对相对误差（MARE）相对于FP64真值超过1%，有时甚至超过10%；而FP16/FP32保持MARE在10?3量级，满足典型训练需求。该结果表明FP32累加能有效控制数值误差；剩下的挑战是在GPU上高效执行混合精度DWC。

然而，CUDA Core上的FP16/FP32 DWC面临硬件限制。大多数GPU架构（如Ampere、Ada、Hopper）不在CUDA Core上提供原生混合精度FMA，因此每次FMA都需要显式的FP16到FP32转换。这将吞吐量限制在FP32 CUDA-Core上限，低于某些GPU上更高的FP16峰值；同时浪费了计算周期在数据类型转换上，使得对于大滤波器，FP16/FP32 DWC比FP32 DWC更慢，如图6所示。

图6. Flash-DWC内核的吞吐量比较。滤波器尺寸（FH=FW，x轴）从3到51。吞吐量（y轴）单位为TFLOPS。子图标题指示GPU和NHWC格式的特征图维度。Flash-DWC表示基于GEMV的CUDA-Core内核；Flash-DWC-T1、-T2和-T4表示基于GEMM的Tensor-Core内核。

为了克服CUDA-Core的限制，我们转向Tensor Core，它原生支持更高吞吐量的FP16/FP32 GEMM。这还避免了FMA与辅助任务之间的资源竞争：Tensor Core处理FMA，而辅助任务（如地址计算、条件分支）仍在CUDA Core上执行。此后，我们简称FP16/FP32混合精度为FP16，并将Tensor-Core内核记为Flash-DWC-T。如图6所示，Flash-DWC-T显著优于FP16和FP32 CUDA-Core内核。

4.2 从GEMV转换为GEMM

如第2.2节所述，朴素地将DWC的GEMV填充为GEMM会浪费78%的矩阵列。类似的挑战也出现在模板计算中，最近的研究[3, 19]将模板GEMV转换为超过50%稠密度的GEMM。然而，这些方法不适合DWC，因为底层张量形状存在根本差异。模板计算操作单个通道，使用适中的滤波器尺寸，而DWC涉及数百或数千个通道（例如1024、2048），并且可以使用高达31×31或更大的滤波器[2, 7, 20]。

这些差异影响了滤波器和输入两种变换。

**(i) 滤波器变换：** 模板滤波器较小且为单通道，因此变换后的滤波器矩阵可以完全驻留在常量内存缓存中[3, 19]，甚至寄存器中[21]。尽管变换本身复杂且引入了矩阵冗余，但其开销可以忽略不计：规模小，在内核启动前执行一次，且结果矩阵在多次启动间全局共享。相比之下，在DWC中，即使是3×3的多通道滤波器也可能超过64 KB的常量内存限制。此外，由于每个块处理128字节的通道段以最大化内存事务粒度，大于8×8的滤波器会溢出每个SM的8 KB常量缓存。因此，DWC必须将滤波器分割为片段，并将变换结果缓存到更大的SMEM中，许多近期GPU每个SM可提供128 KB或更多。

**(ii) 输入变换：** 模板方法通过按不规则空间偏移收集输入元素来构建输入矩阵。这是可行的，因为所有访问的元素位于同一通道且在内存中相邻；偏移量还可以预计算为完全适配缓存的查找表。然而，在DWC中，大量通道在空间位置之间引入了较长的内存步长，使得这种不规则的、基于表驱动的收集效率低下。为了最大化带宽，DWC要求在通道轴上进行合并加载并在宽度轴上进行连续访问。

DWC的实用GEMV到GEMM变换必须满足三项要求：(i) 将滤波器和输入补丁分割为片段，即时变换它们，并将结果以最少冗余缓存到SMEM中，以释放片上资源；(ii) 将片段变换与GEMM计算重叠，保持变换简洁以减少GEMM间延迟；(iii) 与我们的宽度优先输入复用对齐，以增强算术强度并保持访问连续性。

我们的变换将一批M个GEMV打包为一个GEMM。每个GEMV操作TM个相邻的输入补丁，且所有GEMV共享一个一维滤波器片段∈R^(Ffrag)。图7展示了该变换：

图7. DWC的GEMV到GEMM变换。16个空间相邻的GEMV，每个在8×9矩阵和长度为9的向量之间执行，被打包为适合Tensor Core的m16n8k16 GEMM。

在此，GEMV[i] 表示打包批次中的第 i 个 GEMV。(i) 滤波器变换：我们将滤波器片段展开为滤波器矩阵 ∈R^k×TM，其中第 i 列将片段放置在索引 i 处，并在其余位置填充零至长度 K，其中 K=TM+Ffrag?1。(ii) 输入变换：对于每个 GEMV，我们将其 TM 个输入补丁压缩为长度 K 的向量，其中 K 是不同输入元素的总数。将所有 M 个这样的向量按行堆叠，得到输入矩阵 ∈RM×K。在生成的 GEMM 中，输入矩阵是稠密的，而滤波器矩阵的密度为 Ffrag/K = 1?(TM?1)/(Ffrag+TM?1)。对于 FP16/FP32 GEMM，Tensor Core 支持形状 m8n8k4、m16n8k8 和 m16n8k16。我们丢弃 m8n8k4，因为其算术强度较低。由于较大的 Ffrag 可获得更高的密度，我们选择 m16n8k16 并设置 {M, TM, Ffrag} 为 {16, 8, 9}，实现 56.25% 的密度（4096 次 FMA 中有 2304 次非零）。该密度是朴素填充的 4.5 倍，与模板变换 [3, 19] 的密度相当。保持第 3.3 节中的宽度优先分解，我们将任何 2D 滤波器分割为 FH × ?FW/9? 个宽度为 9 的片段，并累加其 1D DWC 结果。

4.3 GEMM 冗余最小化

使用与第 3.1 节相同的分块方法，我们将工作负载划分为 OCBNNOHOWBM 的块，每块 4 个 warp。由于 FP16 相比 FP32 将数据量减半，我们将每线程处理的通道数加倍（从 4 增至 8），每 warp 处理的通道数加倍（从 32 增至 64），以维持 128 字节的内存事务。相应地，我们设置 BN=64 和 BM=128，这最大化了输入补丁的重叠，同时匹配 m16n8k16 GEMM 的 16×8 瓦片大小。将多个 GEMV 打包为 GEMM 将算术强度提升至 4096·0.5625/(256+128)=6，但仍然低于我们 CUDA Core 内核的强度。此外，在 SMEM 中存储完整的滤波器和输入矩阵需要 256BN 和 4BNBM 字节，这达到了 48 KB 的静态分配限制，可能降低占用率和 L1 缓存容量。为解决这两个问题，我们通过两项优化来减少 GEMM 冗余：(i) 压缩 SMEM 中的输入和滤波器矩阵，并在寄存器中即时重建它们以用于 Tensor Core GEMM；(ii) 采用宽度优先的输入重用以最小化从全局内存加载矩阵数据。

在滤波器矩阵 ∈R16×8 中，9 元素的滤波器片段在 8 列中重复，如图 8(b) 所示。因此我们将其存储为非重复格式：在片段两侧各填充 7 个零，得到 23 元素的向量。要重建第 i 列，块只需从索引 i 开始提取 16 个连续元素，无需 warp 分歧。这将滤波器矩阵的存储成本降低到仅 46BN 字节，为原始值的 18%，将总 SMEM 使用降至 34.9 KB。我们将填充后的滤波器片段初始化为零，并让块中的线程协作将滤波器数据加载到 SMEM 中，以最大化数据重用。这将算术强度提升至 4096·0.5625/(256+FW/γ) > 8.69，其中 γ=?FW/9? 是每个滤波器行的宽度为 9 的片段数。

图 8. 共享内存中的矩阵压缩。(a) 输入矩阵：当 θ=4 且 OW≥32 时，四个连续行共享 3×8 个重叠元素，因此 16×16 的输入矩阵被压缩为五个 4×8 缓冲区。(b) 滤波器矩阵：存储为 23 元素向量以消除列冗余。

在输入矩阵 ∈R16×16 中，当 OW≥16 时，相邻行沿宽度轴重叠 8 个元素，如图 8(a) 所示。因此，当 OW≥8θ 时，我们将 θ 个连续行从 16θ 个元素压缩为 8(θ+1) 个元素，存储在 (θ+1) 个宽度为 8 的输入 SMEM 缓冲区中。这将算术强度提升至 4096·0.5625/(128(1+θ?1)+FW/γ)。我们将结果得到的 θ=1、2、4 的 Tensor Core 内核分别称为 Flash-DWC-T1、-T2 和 -T4。T2 和 T4 实现的算术强度分别为 11.96 和 14.35，同时将 SMEM 使用降至 26.9 KB 和 22.9 KB。在填充 SMEM 以避免 bank 冲突后，Flash-DWC-T 在每个 SM 上保留了超过 70% 的片上内存用于并发块和 L1 缓存。

我们在 SMEM 中实现宽度优先的输入重用以适配 Tensor Core GEMM。由于每个 GEMM 处理宽度为 9 的滤波器片段和输入矩阵的 16 元素行，我们可以为下一个 GEMM 迭代重用每行 7 个输入元素。为在避免 warp 分歧的同时最大化重用，前 (θ?1) 个缓冲区从下一个缓冲区重用 7 个元素并仅从 X 加载 1 个新元素，而最后一个缓冲区加载 8 个新元素。这将算术强度提升至 4096·0.5625/(128θ?1+112γ?1+16+FW/γ)，该值随 FW 增大而增大，如图 5(b) 所示。对于 T4，当 FW≥3 时强度超过 13.6，当 FW≥11 时超过 20.4。当 FW→+∞ 时，T2 和 T4 的强度分别趋近于 25.88 和 40.42，远高于 CUDA Core 内核的 15.06 上限。

4.4 基于 GEMM 的内核工作流

我们将基于 GEMV 的设计和 GEMV-to-GEMM 变换整合为 Flash-DWC-T 的连贯工作流，如算法 2 所示。主循环（第 4-11 行）首先沿高度轴迭代。每次高度轴迭代包含 ?FW/9? 次宽度轴迭代，这些迭代通过 C++ 模板在编译时展开。由于每个块处理 BN 个通道，每次迭代执行一批 BN 个 m16n8k16 GEMM，每个通道一个。

软件流水线。我们的 GEMV-to-GEMM 变换通过偏移映射和数据加载实现。为将数据加载与 GEMM 计算重叠，我们将每个 BN 批次的 m16n8k16 GEMM 拆分为两个 m16n8k8 子单元，称为 GEMM1 和 GEMM2，如图 9 所示。这种分解形成了双路径拆分 GEMM 流水线，由 Prefetch1→GEMM1 和 Prefetch2→GEMM2 组成。它可以被视为双缓冲的广义形式。然而，与分块 GEMM 内核中的标准两阶段乒乓缓冲不同，这两条路径是不对称的，并使用部分重叠的 SMEM 区域。具体而言，GEMM1 使用第一个 θ 个输入 SMEM 缓冲区和填充后滤波器片段的前 8 个非零元素，而 GEMM2 使用最后 θ 个缓冲区和最后 8 个非零元素。

图 9. 双路径拆分 GEMM 流水线。虚线框显示 BN 批次中一个代表性的 GEMM。每个 m16n8k16 GEMM 对应一个通道，并被拆分为两个 m16n8k8 子单元。箭头表示数据依赖；迭代间的同步通过重叠 SMEM 上的数据重用来强制实现。

在每次迭代中，Prefetch1 和 Prefetch2 作为两组异步加载发出。Prefetch1 完成后 GEMM1 开始，同时 Prefetch2 并行进行。GEMM2 在 GEMM1 和 Prefetch2 都完成后开始。由于这两个子 GEMM 单元共享重叠的 SMEM 区域，下一次迭代的 Prefetch1 必须等待前一次 GEMM2 释放这些区域。即便如此，该流水线在单个块内将 50% 的 GEMM 计算与数据加载重叠。此外，低 SMEM 使用使多个块能在一个 SM 上并发执行，进一步在块之间重叠数据加载与 GEMM 计算。

数据加载。对于输入元素，我们使用 cp.async.ca PTX 指令，将数据缓存在 L1 中重用 SM 内的输入补丁重叠。滤波器元素没有空间重叠，因此我们用 cp.async.cg 加载它们，仅在 L2 中缓存，从而为输入元素保留 L1 容量。在第一次宽度轴迭代之前，我们将初始大批量输入元素加载到 SMEM 中（第 5 行）。在每次宽度轴迭代中：Prefetch1 重用输入 SMEM 缓冲区来准备 GEMM1 所需的 (TM?1/TM=7/8) 输入元素，并从全局内存加载剩余的 (1/TM=1/8)。同时，Prefetch2 加载 GEMM2 所需的最后一个输入 SMEM 缓冲区。与其在输入 SMEM 缓冲区之间复制数据，我们循环重新编号缓冲区并相应更新其访问索引，如图 10 所示。

图 10. 用于输入重用的循环缓冲区重新编号，无需显式数据复制。蓝色箭头表示移位缓冲区以进行数据重用。在每次迭代中，编号 0–7 表示每个缓冲区内的访问索引；新加载和重用的元素分别用红色和黑色字体表示。单元格颜色跟踪每个元素最初被加载到的缓冲区，重用的元素在迭代中保留其颜色。

输入特征图上的零填充隐式处理。配合谓词掩码，cp.async 跳过零填充区域并向 SMEM 写入零，从而避免 warp 分歧。为节省寄存器并增强活跃线程数，我们将每批 32 个谓词掩码压缩为一个 32 位整数，每个空间位置使用一个位。由于 FW 和填充大小是常数，宽度轴谓词掩码预先计算并在所有高度轴迭代中重用。此外，我们扫描每个块以检测所有线程在高度轴上是否共享相同的填充零。如果是，我们裁剪高度轴迭代以排除这些共享零，使用 FHs(≥0) 和 FHe(
gemm 计算。gemm1 和 gemm2 各自计算一批 bn 个通道级 m16n8k8 gemm。每个 warp 在 tensor core 上执行 bn/4 条 mma.sync.aligned.m16n8k8 ptx 指令。为避免 bank 冲突，我们将 smem 数组 w 和 x 的最后一个维度填充 128 位 [2]。主循环结束后，我们通过 smem 排列寄存器矩阵 v（第 12 行），并将结果 bm×bn dwc 瓦片以 128 字节事务存储到全局内存中。

为什么不用稀疏 tensor core？尽管稀疏 tensor core（例如在 sparstencil [19] 中使用）具有密集 tensor core 两倍的理论吞吐量，但它们不适合 dwc。首先，其支持的最小 gemm 形状 m16n8k16 将迫使我们放弃双路径拆分 gemm 流水线，该流水线隐藏了 50% 的 gemm 计算。其次，即使我们重新组织数据流以满足 2:4 稀疏性，gemm 运行时间最多减少 43.75%，这还不计额外开销和数据访问连续性的损失。由于 dwc 本质上是内存受限而非计算受限，切换到稀疏 tensor core 将以牺牲内存延迟隐藏为代价换取微薄的计算增益，最终降低整体性能。

5 将 flash-dwc 扩展到 cnn 训练

基于第 3 节和第 4 节中优化的内核，我们将 flash-dwc 扩展到完整的 cnn 训练，并将其集成到 pytorch 2.10 [36] 中作为即插即用库，其中 flashdwc.ops 用于操作调用，flashdwc.nn 用于模型构建。

5.1 反向传播

前向卷积（fc）直接由 flash-dwc 内核支持。训练还需要两个反向算子：(i) 反向数据卷积（bdc），将输出梯度 ?y 与转置滤波器 wt 卷积以计算输入梯度 ?x；(ii) 反向滤波器卷积（bfc），将输入特征图 x 与 ?y 卷积以计算滤波器梯度 ?w。由于 bdc 在旋转滤波器后具有相同的 dwc 结构，它可以重用 flash-dwc 内核。主要挑战在于 bfc：其小输出张量导致并行度不足，且不匹配原始内核的分块方法。图 11 展示了 flash-dwc 如何支持所有三个阶段。

图 11. flash-dwc 用于 cnn 训练。fc 和 bdc 利用 tensor core 和 cuda core 内核。对于 bfc，我们通过将 ?y 分块为 z 个子 bfc 并以原子方式累加其结果来增加并行度。如图所示，每个产生 11×11 输出的子 bfc 产生 6 个块：mega kernel 通过组合两个 tm=2 和 9 的内核来匹配 fw=11；β=4 最小化 fh 轴填充。

bdc 实现。我们将 w 旋转 180° 得到 wt，然后在 wt 和 ?y 上执行 flash-dwc 内核。由于旋转开销可忽略不计，bdc 的性能与 fc 类似，同样受益于 tensor core。

累加轴并行化。bfc 的输出 ?w ∈ r^{fh×fw×oc} 仍然较小，且不随批次大小或空间分辨率扩展。这与 fc 和 bdc 不同，后者的输出较大，自然提供充足的并行度。如果我们仅通过分块 ?w 来并行化 bfc，产生的块数太少，无法充分利用硬件。我们通过沿累加轴并行化来解决此问题。具体而言，我们将 ?y 拆分为 z 个瓦片，每个瓦片作为独立的子 bfc 处理。这使并行度提高 z 倍，但引入了通过 fp32 原子加法合并子 bfc 结果的额外开销。为平衡这两个效应，我们根据 dwc 参数和 sm 数量动态决定 z。如果累加长度不能被 z 整除，我们将余数附加到最后一个瓦片，并调整 z 以最小化此余数以实现均衡的工作负载。我们进一步通过一次性的性能分析步骤来细化 z，测量性能并选择最佳值。

灵活分块。bfc 的不规则输出形状不匹配原始 flash-dwc 内核的分块方法。在 bm=128 时，这些内核将空间维度分块为 16×8、8×16 或 4×32。这些形状非常适合特征图，后者通常较大且是 8 的倍数（例如 112×112、64×64）。然而，bfc 的输出 ?w 具有奇数大小的空间维度，较小或不对齐于这些瓦片，导致大量零填充和不必要的计算。 作为起始和结束位置（第 4 行）。这在未引入 warp 分歧的情况下略微减少了零填充开销。 gemm 计算。gemm1 和 gemm2 各自计算一批 bn 个通道级 m16n8k8 gemm。每个 warp 在 tensor core 上执行 bn 4 条 mma.sync.aligned.m16n8k8 ptx 指令。为避免 bank 冲突，我们将 smem 数组 w 和 x 的最后一个维度填充 128 位 [2]。主循环结束后，我们通过 smem 排列寄存器矩阵 v（第 12 行），并将结果 bm×bn dwc 瓦片以 128 字节事务存储到全局内存中。 为什么不用稀疏 tensor core？尽管稀疏 tensor core（例如在 sparstencil [19] 中使用）具有密集 tensor core 两倍的理论吞吐量，但它们不适合 dwc。首先，其支持的最小 gemm 形状 m16n8k16 将迫使我们放弃双路径拆分 gemm 流水线，该流水线隐藏了 50% 的 gemm 计算。其次，即使我们重新组织数据流以满足 2:4 稀疏性，gemm 运行时间最多减少 43.75%，这还不计额外开销和数据访问连续性的损失。由于 dwc 本质上是内存受限而非计算受限，切换到稀疏 tensor core 将以牺牲内存延迟隐藏为代价换取微薄的计算增益，最终降低整体性能。 5 将 flash-dwc 扩展到 cnn 训练 基于第 3 节和第 4 节中优化的内核，我们将 flash-dwc 扩展到完整的 cnn 训练，并将其集成到 pytorch 2.10 [36] 中作为即插即用库，其中 flashdwc.ops 用于操作调用，flashdwc.nn 用于模型构建。 5.1 反向传播 前向卷积（fc）直接由 flash-dwc 内核支持。训练还需要两个反向算子：(i) 反向数据卷积（bdc），将输出梯度 ?y 与转置滤波器 wt 卷积以计算输入梯度 ?x；(ii) 反向滤波器卷积（bfc），将输入特征图 x 与 ?y 卷积以计算滤波器梯度 ?w。由于 bdc 在旋转滤波器后具有相同的 dwc 结构，它可以重用 flash-dwc 内核。主要挑战在于 bfc：其小输出张量导致并行度不足，且不匹配原始内核的分块方法。图 11 展示了 flash-dwc 如何支持所有三个阶段。 图 11. flash-dwc 用于 cnn 训练。fc 和 bdc 利用 tensor core 和 cuda core 内核。对于 bfc，我们通过将 ?y 分块为 z 个子 bfc 并以原子方式累加其结果来增加并行度。如图所示，每个产生 11×11 输出的子 bfc 产生 6 个块：mega kernel 通过组合两个 tm=2 和 9 的内核来匹配 fw=11；β=4 最小化 fh 轴填充。 bdc 实现。我们将 w 旋转 180° 得到 wt，然后在 wt 和 ?y 上执行 flash-dwc 内核。由于旋转开销可忽略不计，bdc 的性能与 fc 类似，同样受益于 tensor core。 累加轴并行化。bfc 的输出 ?w ∈ r^{fh×fw×oc} 仍然较小，且不随批次大小或空间分辨率扩展。这与 fc 和 bdc 不同，后者的输出较大，自然提供充足的并行度。如果我们仅通过分块 ?w 来并行化 bfc，产生的块数太少，无法充分利用硬件。我们通过沿累加轴并行化来解决此问题。具体而言，我们将 ?y 拆分为 z 个瓦片，每个瓦片作为独立的子 bfc 处理。这使并行度提高 z 倍，但引入了通过 fp32 原子加法合并子 bfc 结果的额外开销。为平衡这两个效应，我们根据 dwc 参数和 sm 数量动态决定 z。如果累加长度不能被 z 整除，我们将余数附加到最后一个瓦片，并调整 z 以最小化此余数以实现均衡的工作负载。我们进一步通过一次性的性能分析步骤来细化 z，测量性能并选择最佳值。 灵活分块。bfc 的不规则输出形状不匹配原始 flash-dwc 内核的分块方法。在 bm=128 时，这些内核将空间维度分块为 16×8、8×16 或 4×32。这些形状非常适合特征图，后者通常较大且是 8 的倍数（例如 112×112、64×64）。然而，bfc 的输出 ?w>
gemm 计算。gemm1 和 gemm2 各自计算一批 bn 个通道级 m16n8k8 gemm。每个 warp 在 tensor core 上执行 bn/4 条 mma.sync.aligned.m16n8k8 ptx 指令。为避免 bank 冲突，我们将 smem 数组 w 和 x 的最后一个维度填充 128 位 [2]。主循环结束后，我们通过 smem 排列寄存器矩阵 v（第 12 行），并将结果 bm×bn dwc 瓦片以 128 字节事务存储到全局内存中。

为什么不用稀疏 tensor core？尽管稀疏 tensor core（例如在 sparstencil [19] 中使用）具有密集 tensor core 两倍的理论吞吐量，但它们不适合 dwc。首先，其支持的最小 gemm 形状 m16n8k16 将迫使我们放弃双路径拆分 gemm 流水线，该流水线隐藏了 50% 的 gemm 计算。其次，即使我们重新组织数据流以满足 2:4 稀疏性，gemm 运行时间最多减少 43.75%，这还不计额外开销和数据访问连续性的损失。由于 dwc 本质上是内存受限而非计算受限，切换到稀疏 tensor core 将以牺牲内存延迟隐藏为代价换取微薄的计算增益，最终降低整体性能。

5 将 flash-dwc 扩展到 cnn 训练

基于第 3 节和第 4 节中优化的内核，我们将 flash-dwc 扩展到完整的 cnn 训练，并将其集成到 pytorch 2.10 [36] 中作为即插即用库，其中 flashdwc.ops 用于操作调用，flashdwc.nn 用于模型构建。

5.1 反向传播

前向卷积（fc）直接由 flash-dwc 内核支持。训练还需要两个反向算子：(i) 反向数据卷积（bdc），将输出梯度 ?y 与转置滤波器 wt 卷积以计算输入梯度 ?x；(ii) 反向滤波器卷积（bfc），将输入特征图 x 与 ?y 卷积以计算滤波器梯度 ?w。由于 bdc 在旋转滤波器后具有相同的 dwc 结构，它可以重用 flash-dwc 内核。主要挑战在于 bfc：其小输出张量导致并行度不足，且不匹配原始内核的分块方法。图 11 展示了 flash-dwc 如何支持所有三个阶段。

图 11. flash-dwc 用于 cnn 训练。fc 和 bdc 利用 tensor core 和 cuda core 内核。对于 bfc，我们通过将 ?y 分块为 z 个子 bfc 并以原子方式累加其结果来增加并行度。如图所示，每个产生 11×11 输出的子 bfc 产生 6 个块：mega kernel 通过组合两个 tm=2 和 9 的内核来匹配 fw=11；β=4 最小化 fh 轴填充。

bdc 实现。我们将 w 旋转 180° 得到 wt，然后在 wt 和 ?y 上执行 flash-dwc 内核。由于旋转开销可忽略不计，bdc 的性能与 fc 类似，同样受益于 tensor core。

累加轴并行化。bfc 的输出 ?w ∈ r^{fh×fw×oc} 仍然较小，且不随批次大小或空间分辨率扩展。这与 fc 和 bdc 不同，后者的输出较大，自然提供充足的并行度。如果我们仅通过分块 ?w 来并行化 bfc，产生的块数太少，无法充分利用硬件。我们通过沿累加轴并行化来解决此问题。具体而言，我们将 ?y 拆分为 z 个瓦片，每个瓦片作为独立的子 bfc 处理。这使并行度提高 z 倍，但引入了通过 fp32 原子加法合并子 bfc 结果的额外开销。为平衡这两个效应，我们根据 dwc 参数和 sm 数量动态决定 z。如果累加长度不能被 z 整除，我们将余数附加到最后一个瓦片，并调整 z 以最小化此余数以实现均衡的工作负载。我们进一步通过一次性的性能分析步骤来细化 z，测量性能并选择最佳值。

灵活分块。bfc 的不规则输出形状不匹配原始 flash-dwc 内核的分块方法。在 bm=128 时，这些内核将空间维度分块为 16×8、8×16 或 4×32。这些形状非常适合特征图，后者通常较大且是 8 的倍数（例如 112×112、64×64）。然而，bfc 的输出 ?w 具有奇数大小的空间维度，较小或不对齐于这些瓦片，导致大量零填充和不必要的计算。>

Flash-DWC-T 将 BM 固定为 128 以支持 m16n8k16 GEMM，但其宽度为 9 的碎片常与 BFC 累加瓦片发生对齐偏差。相比之下，我们的 CUDA Core 内核基于 GEMV，不受 Tensor Core 形状约束。因此，我们对其进行灵活分块适配，以最小化零填充，更好地匹配 BFC 形状。我们调整 TM——即每个线程处理的空间位置数，取值从 1 到 9——并将两个具有不同 TM 值的核心融合为一个大内核，以无缝匹配任意 FW。由于累加长度 NOHOW 通常是 8 的倍数，我们以长度 8 为单位处理，并将每个线程块组织为 β×8 个线程：(i) 沿通道轴，8 个线程执行 128 字节的合并加载；(ii) 沿 FH 轴，β 个线程复用补丁重叠部分，协作加载一个长度 8 的单位以执行块级 GEMV。我们从 β∈{4,8} 中选择，因为两者都能整除 8，并保持总线程数为 32 的倍数，从而避免非活跃的 warp 通道。在两者之间，我们选择能最小化 FH 轴上零填充的一个；如果两者零填充相同，则取较大值以最大化补丁重叠。混合累加。它作为可选优化用于加速混合精度 BFC：每个子 BFC 以 FP16 累加，最终结果以 FP32 合并。这避免了显式数据类型转换，减少了寄存器占用，并可以利用某些 GPU 上更高的 FP16 CUDA Core 吞吐量。其精度仍然可接受，原因有二。(i) BFC 的长累加长度可以拆分为许多（例如 ≥128）短子单元。以 FP16 累加每个子单元可限制误差增长，而以 FP32 合并结果则受益于其更大的尾数位数和动态范围，基本避免了精度损失和数值溢出。(ii) 与 FC 和 BDC 不同——后者的对称权重初始化会最大化正负抵消——BFC 的输入 X 和 ?Y 由于激活函数和损失函数的非线性，围绕零值的对称性较弱。这减少了 FP16 舍入误差。第 6.3 节显示混合累加与全 FP32 累加具有相似的收敛性。5.2 内核融合 DWC 通常紧接在批归一化（BN）和激活函数（Act）之后，形成一个常见模式：BN+Act+DWC。对于计算单元未完全饱和的小滤波器，我们提供了可选的将 BN 和 Act 融合到 Flash-DWC 内核中的方案，以利用空闲周期并减少数据移动。我们在 CUDA Core 内核上实现此融合，这些内核将输入数据缓存在寄存器中，因此非常适合在 DWC 之前进行即时 BN 和 Act 计算。尽管存在由瓦片重叠引起的重算，但额外计算负载仍低于计算受限阈值，因此对于小滤波器，融合内核的延迟几乎与对应的未融合 DWC 内核相同。6 评估 我们首先评估 CUDA Core 内核，以验证我们的 GEMV 数据流协同设计，并分析从内存受限向计算受限操作的转变。然后评估基于 GEMM 的 Tensor Core 内核，以展示 GEMV 到 GEMM 的转换如何进一步突破 CUDA Core 的性能上限。最后，我们在 ImageNet-1K 上对 Flash-DWC 进行端到端 CNN 训练评估，衡量加速比和收敛性。6.1 FP32 CUDA Core 内核评估 6.1.1 设置。我们在 RTX 4090 和 A100 80GB 上进行评估，涵盖 Ada-Lovelace GeForce 和 Ampere 数据中心 GPU。我们将 Flash-DWC 与 DWC 基线和计算受限参考进行比较：(i) cuDNN 9.18：我们同时启用 cuDNN 的 Tensor Core 和 CUDA Core 算法。CuGEMM-DWC 指其直接算法和 GEMM 算法中最快者；CuFFT-DWC 指其 FFT 和 Winograd 算法中最快者。(ii) Inductor-DWC：由 PyTorch 2.10 torch.compile 的 Inductor 后端和 max-autotune 生成。(iii) cuBLAS-GEMM：它不是 DWC 实现，而是作为计算受限参考，衡量 Flash-DWC 接近 CUDA Core 性能上限的程度。其维度设为 m=n=k 且 2m3≈2NOHOWOCFHFW，使其总 FLOPs 与 DWC 匹配。DWC 参数遵循常见 CNN 配置。(i) 滤波器从小型的 3×3 [12, 13, 39] 到超大型的 51×51 [2, 20]，涵盖 7×7 [25, 41, 44]、9×9 [8, 41] 和 31×31 [7] 等代表性尺寸。(ii) 特征图按 NHWC 顺序进行系统性缩放：128×128×128×128、128×64×64×256 和 128×32×32×1024。这些尺寸足以饱和 GPU 并行度，并在测试中保持可比的负载。6.1.2 结果与缩放分析 图 12 中所有吞吐量值均在预热运行后经过 500 次运行取平均。图 13 展示了随着滤波器尺寸增大算术强度增加时，CUDA Core Flash-DWC 内核的吞吐量及其相对于 cuBLAS-GEMM 的吞吐量比值如何缩放。图 12. FP32 吞吐量。滤波器尺寸（FH=FW，x 轴）范围从 3 到 51。吞吐量（y 轴）单位为 TFLOPS。子图标题列出 GPU 类型和 NHWC 格式的特征图维度。图 13. FP32 CUDA Core Flash-DWC 内核随算术强度的缩放。滤波器尺寸（x 轴，FH=FW）范围从 3 到 51。左 y 轴显示以 TFLOPS 为单位的吞吐量。右 y 轴显示两个归一化指标：算术强度比（强度相对于其渐近极限），以及 Flash-DWC 吞吐量相对于计算受限 cuBLAS-GEMM 的比值。所有值均跨特征图维度取平均。Flash-DWC 在滤波器尺寸增大时表现出从内存受限向计算受限操作的明显转变。这一趋势符合我们的算术强度模型 2TM1+(TM?1)/FW+BN/512，并经硬件测量验证。对于小滤波器（FW≤9），内核仍处于内存受限状态。在此区间，(TM?1)/FW 项占主导，因此算术强度和吞吐量都随 FW 快速增加。因此，运行时间的增长远慢于 O(FHFW) 的 FLOP 数量：将滤波器从 3×3 增大到 7×7，运行时间仅延长 2.2 倍，尽管 FLOPs 增加了约 5.44 倍。这里 A100 优于 RTX 4090，因为其内存带宽更高（约 2 TB/s 对比约 1 TB/s），证实内存流量是主要性能瓶颈。对于较大滤波器，算术强度曲线逐渐趋于平缓，而吞吐量和相对于 cuBLAS-GEMM 的吞吐量比值却持续上升。这意味着 Flash-DWC 不仅绝对速度更快，而且逐步接近计算受限行为。在两种 GPU 上，当 FW≥13 时，Flash-DWC 达到 cuBLAS-GEMM 吞吐量的 65%；在 A100 上当 FW≥15 时超过 90%，表明更高的算术强度将 DWC 推向 CUDA Core 性能上限。屋顶线交叉点出现在 FW≥11 处，此时在 FLOPs 更高的 RTX 4090 上 Flash-DWC 的吞吐量（82.6 TFLOPS）超过在带宽更高的 A100 上的吞吐量（19.5 TFLOPS），表明瓶颈已从内存带宽转移到计算能力。这一转变带来了相对于所有 DWC 基线的巨大加速。Flash-DWC 比最快的 cuDNN DWC 算法快 1.08×–8.48×（平均 5.09×），比 CuGEMM-DWC 快 1.08×–8.72×（平均 5.98×），比 CuFFT-DWC 快 1.73×–83.05×（平均 13.90×），比 Inductor-DWC 快 2.65×–11.51×（平均 5.65×）。对于汇总加速比，平均值指算术平均。对应的几何平均加速比分别为 4.67×、5.62×、9.48× 和 5.42×。RTX 4090 上的加速比更大，因为其 CUDA Core 峰值更高：例如，相对于 CuGEMM-DWC，Flash-DWC 在 RTX 4090 上实现 6.34× 平均加速，而 A100 上为 5.61×。Flash-DWC 还受益于较小的特征图，其中更高的缓存命中率改善了吞吐量。6.2 FP16 Tensor Core 内核评估 6.2.1 设置。我们使用与第 6.1 节相同的基线、DWC 参数集和 GPU，区别在于所有 DWC 基线现在操作于 FP16 张量。对于 cuDNN，我们报告其 FP16 和 FP32 累加模式中最快结果。FP32 cuBLAS-GEMM 仍作为计算受限参考，因为 CUDA Core 上的 FP16/FP32 DWC 需要显式 FP16 到 FP32 转换，因此受限于 FP32 吞吐量而非标称 FP16 峰值，如第 4.1 节所述。在 RTX 4090 上，这一选择进一步合理，因为其 FP16 和 FP32 CUDA Core 峰值相同。我们使用 cuBLAS-GEMM 来评估 Flash-DWC-T 是否超越 CUDA Core 性能上限。6.2.2 结果与缩放分析 图 14 中所有吞吐量值均在预热运行后经过 500 次运行取平均。图 15 展示了随着滤波器尺寸增大算术强度增加时，Flash-DWC-T4 的吞吐量及其相对于 cuBLAS-GEMM 的吞吐量比值如何缩放。图 14. FP16 吞吐量。滤波器尺寸（x 轴，FH=FW）范围从 3 到 51。吞吐量（y 轴）单位为 TFLOPS。子图标题列出 GPU 类型和 NHWC 格式的特征图维度。所有 DWC 算法操作于 FP16 张量；cuBLAS-GEMM 为 FP32，作为 CUDA Core 性能上限的参考。图 15. FP16 Flash-DWC-T4 随算术强度的缩放。滤波器尺寸（x 轴，FH=FW）范围从 3 到 51。左 y 轴显示以 TFLOPS 为单位的吞吐量。右 y 轴显示两个归一化指标：算术强度比（定义为强度相对于其渐近极限），以及 Flash-DWC-T4 吞吐量相对于 CUDA Core cuBLAS-GEMM 的比值。超过 100% 的值表明 Tensor Core 加速超过了 CUDA Core 计算上限。所有值均跨特征图维度取平均。与对应的 CUDA Core 版本相比，Flash-DWC-T 实现 0.96×–4.35× 加速（平均 2.01×）。唯一的次线性情况出现在 3×3 滤波器上。更普遍地，两种内核（CUDA Core 和 Tensor Core）在小滤波器（FW≤7）时仍处于内存受限状态，因此表现出相似的吞吐量。然而，随着滤波器尺寸增大算术强度增长，Flash-DWC-T 提升更快，因为 Tensor Core 提供更高速吞吐量并避免了显式 FP16 到 FP32 转换。相比之下，在 A100 上，CUDA Core 内核的吞吐量在大滤波器时甚至略有下降，因为一旦 CUDA Core 饱和，数据类型转换会与 FMA 争夺计算周期。如图 15 所示，Flash-DWC-T 吞吐量遵循其算术强度模型 4096?0.5625128θ?1+112γ?1+16+FW/γ，并与其相对于 cuBLAS-GEMM 的吞吐量比值紧密吻合。吞吐量还以 9 为步长递增，反映了先宽度分解为宽度 9 的碎片。在三种 Tensor Core 配置中，Flash-DWC-T4 实现最高吞吐量，其次是 T2，然后是 T1。这一顺序与它们的输入矩阵压缩比一致：更大的 θ 产生更高算术强度和更低 SMEM 使用。综合来看，这些结果表明算术强度仍然是 Tensor Core 上性能的主要预测因素，正如在 CUDA Core 内核中一样。对于中等大小的滤波器，Flash-DWC-T 在两种 GPU 上甚至超过了 cuBLAS-GEMM 吞吐量。一旦 Flash-DWC 在 CUDA Core 上大幅缓解了 DWC 的内存瓶颈，主导限制不再是内存流量，而是 CUDA Core 计算吞吐量。Tensor Core 加速提升了这一上限，使 DWC 吞吐量能够超越 CUDA Core GEMM 参考继续扩展。从这个意义上说，超越 cuBLAS-GEMM 反映了从内存带宽到计算吞吐量的转变，并说明为什么 Tensor Core 对于进一步加速是必要的。这一效应在 A100 上更为显著，因为其 Tensor Core 与 CUDA Core 之间的峰值吞吐量差距更大。相应地，Flash-DWC-T 在该 GPU 上相对于 cuBLAS-GEMM 及其 CUDA Core 对应版本均取得更大优势。总体而言，Flash-DWC-T 在 A100 上达到 cuBLAS-GEMM 吞吐量的 0.36×–2.68×（平均 1.67×），而在 RTX 4090 上为 0.1×–1.71×（平均 0.99×）。Flash-DWC-T 相对于所有 DWC 基线也取得巨大收益。重要的是，它在整个滤波器尺寸谱上均实现了正加速，而非以牺牲小滤波器性能为代价来提升大滤波器性能。它比最快的 cuDNN DWC 算法快 1.84×–16.38×（平均 9.01×），比 CuGEMM-DWC 快 1.84×–16.38×（平均 10.58×），比 CuFFT-DWC 快 3.29×–102.63×（平均 23.35×），比 Inductor-DWC 快 2.88×–17.73×（平均 10.46×）。对应的几何平均加速比分别为 8.21×、9.83×、16.78× 和 9.95×。Flash-DWC-T 还受益于较小的特征图，因为更高的缓存命中率。6.2.3 跨 GPU 比较 尽管 A100 的 FP16/FP32 Tensor Core 峰值（312 TFLOPS）高于 RTX 4090（165 TFLOPS），但 Flash-DWC-T 在 RTX 4090 上实现了更高的绝对吞吐量：平均/峰值为 50.6/90.3 TFLOPS，而 A100 上为 30.3/49.5 TFLOPS 平均/峰值。考虑到我们的 GEMV 到 GEMM 转换的 56.25% 密度后，密度归一化吞吐量在 RTX 4090 上达到平均/峰值 90.0/160.5 TFLOPS，而 A100 上为 53.9/88.0 TFLOPS，RTX 4090 的峰值接近其标称 FP16/FP32 Tensor Core 峰值。这一跨 GPU 差距不能完全由 Tensor Core 峰值或 DRAM 带宽单独解释。为了更好地理解这一行为，我们使用 Nsight Compute 对 Flash-DWC-T 进行了剖析。在剖析的案例中，RTX 4090 始终实现了比 A100 更高的 L2 命中率。对于 128×64×64×256 特征图，9×9、17×17、31×31 和 51×51 滤波器的 RTX 4090/A100 L2 命中率分别为 90.5/82.3%、95.4/89.1%、97.9/93.5% 和 99.0/96.2%。这表明 RTX 4090 更大的 L2 缓存（72 MB 对比 A100 的 40 MB）有助于捕获 Flash-DWC-T 暴露的输入复用。在宽度轴迭代中，Flash-DWC-T 执行初始批量加载，随后进行许多小的增量加载。初始批量加载可以受益于 A100 的更高带宽（约 2 TB/s 对比 RTX 4090 的约 1 TB/s），这与其在小滤波器（FW≤5）上的优势一致。随着滤波器尺寸增大，由于 Flash-DWC-T 的先宽度访问模式和工作负载分布，增量加载变得更加缓存驻留，减少了片外带宽压力。在此区间，RTX 4090 受益于其更大的 L2 缓存并实现更高吞吐量。其他架构因素也可能对差距有所贡献。

诸如谓词评估、地址计算和循环控制等辅助任务在 Tensor Cores 之外执行，而异步复制和共享内存行为会影响数据移动与 Tensor Core 计算之间的重叠。这种跨 GPU 的行为并非 Flash-DWC-T 独有；cuDNN 和 Inductor 的 DWC 实现在 RTX 4090 上的吞吐量也高于 A100，比率平均为 1.52×，峰值达到 2.06×。结合 L2 缓存命中率计数器，这一一致趋势支持我们的解释：由于 DWC 的跨通道复用有限且空间访问模式精细，其性能对缓存层次结构和流水线行为的敏感度高于对 Tensor Core 峰值或 DRAM 带宽的敏感度。

6.3 CNN 训练评估
我们选择训练作为端到端评估场景，因为它不仅涉及推理中使用的正向 DWC 算子，还涉及反向 DWC（BDC）、反向 FCD（BFC）、优化器交互以及收敛保持。该设置评估 Flash-DWC 的效率，同时验证其在完整训练流水线中的安全可用性。

我们选取四个代表性模型族 [7, 20, 25, 41]，涵盖具有中等大小稠密、大型稠密以及分解大型滤波器的 DWC 层。对于所有模型，我们将 PyTorch 的原生 DWC 层替换为 Flash-DWC 实现，同时保持 Flash-DWC 与 PyTorch 评估的模型架构完全一致。对于降采样阶段，当特征图尺寸缩小至名义滤波器尺寸以下时，我们将滤波器裁剪为（特征图尺寸 ?1），因为此类超大滤波器在 ImageNet 识别上通常带来很少的精度提升，却会增加 FLOPs。Flash-DWC 和 PyTorch 使用完全相同的训练设置，包括默认的 Adam 优化器 [14]、10?? 的学习率，以及用于 FP16 混合精度的损失缩放 [33]。我们在两款 Ada-Lovelace 架构的 GPU 上进行实验：RTX 4090（24 GB）和 RTX 5880 Ada（工作站 GPU，48 GB）。如第 6.1 节所示，Ada-Lovelace 架构为评估的 DWC 实现提供了更高的吞吐量，使其非常适合 DWC 工作负载。

6.3.1 速度
我们在 FP32 和 FP16 下测量端到端的每批次训练时间，先进行 50 个批次的预热，然后对 100 个批次取平均值。每次计时迭代涵盖完整的训练步骤，包括数据加载与预处理、主机到设备的传输、正向与反向传播以及优化器更新。输入图像被调整为 256×256，这是一种常见分辨率。为了最大化 GPU 利用率，对于每个模型，我们从 {16, 24, 32, 48, 64, 128, 256} 中选择能够放入内存中的最大批次大小，并对 Flash-DWC 和 PyTorch 使用相同的批次大小。表 1(a) 报告了由此得出的加速比。

表 1. Flash-DWC 相对于 PyTorch 的端到端训练加速比。HA 表示针对 BFC 的混合累加，↑ 表示其在端到端训练中的增益。

Flash-DWC 在 FP32 下实现 1.44×–6.27× 的加速，在 FP16 下实现 1.03×–11.23× 的加速，Tensor Core 加速在 FP16 下带来了大幅更高的峰值加速比。加速比在同一架构的基础版本与放大版本之间保持稳定，表明 Flash-DWC 能够很好地随通道数扩展。由于标准卷积和全连接层的 FLOPs 随通道数呈二次增长，端到端加速比在放大版本中略微减小。混合累加将核函数吞吐量提高了最多约 30%，但受限于 Amdahl 定律，其端到端增益仅为 0.1%–4.6%，这表明 Flash-DWC 已消除了 DWC 瓶颈。

端到端加速比随滤波器尺寸增大而增加。对于 7×7 滤波器的 ConvNeXt，加速比为 1.03×–1.70×；将滤波器扩展到 9×9 后，加速比提升至 1.94×–2.75×。7×7 ConvNeXt 的 FP16 加速比低于 FP32 加速比，因为 PyTorch 在 FP16 下会调度专用核函数，导致改进空间更小。对于使用大量大通道（1024、1536）9×9 滤波器 DWC 的 ConvMixer，Flash-DWC 实现了 3.64×–6.11× 的显著加速；将滤波器扩大到 11×11 后，加速比提升至 4.55×–7.19×。对于具有 31×31 和 51×51 超大滤波器的模型，Flash-DWC 为 RepLKNet 实现了 4.50×–8.64× 的加速，为 SLaK 实现了 3.38×–11.23× 的加速。在 SLaK 中，步幅分解3降低了 Flash-DWC 的算术强度和运行时间占比，因此加速比较低：条纹版本的加速比为 3.38×–4.20×，而稠密版本的加速比为 5.65×–11.23×。核函数融合可以进一步加速小滤波器或中等滤波器场景，但出于公平比较的目的，此处未启用。

Flash-DWC 大幅削弱了更大稠密滤波器带来的运行时惩罚。对于 ConvNeXt 和 ConvMixer，将滤波器从 7×7/9×9 增加到 9×9/11×11，尽管 DWC FLOPs 分别增加了约 65% 和约 49%，端到端训练时间仅分别增加了约 1.3% 和约 15%。类似地，稠密 SLaK 的端到端训练时间仅比其条纹版本增加约 8%，尽管其 DWC FLOPs 大数倍。

我们进一步将输入分辨率扩大到 512×512，这与分割和对象检测等密集预测任务相关。如表 1(b) 所示，Flash-DWC 实现了 1.04×–14.18× 的相似加速比，证明其跨输入尺寸的鲁棒性。

6.3.2 收敛性
我们首先验证 Flash-DWC 在单层上不会引入可测量的精度退化，使用从 U(?1,1) 中抽取的随机张量。在 FP32 下，Flash-DWC 的输出和梯度与 PyTorch 的输出和梯度几乎相同，通常具有零 MARE。在 FP16 下，Flash-DWC 相对于 FP64 真值的 MARE 约为 10?3 至 10??，与 PyTorch 自身的 FP16 实现相当。

我们在 FP16 混合精度下训练代表性模型 [7, 25, 41]，这是一个比 FP32 更具挑战性的设置。输入图像调整为 256×256，批次大小为 64。如图 16 所示，Flash-DWC 的收敛性与 PyTorch 相似，而混合累加的收敛性与全 FP32 累加相当。这些结果证实 Flash-DWC 在提供大幅加速的同时保留了训练精度。

图 16. CNN 训练损失。子图标题列出了模型及其最大滤波器尺寸。HA 表示针对 BFC 的混合累加。

7 可扩展性
Flash-DWC 可扩展至高级 GPU 架构特性和低位数据类型。两个方向都需要核函数级别的重新调优，但不会改变 Flash-DWC 的主算法。

高级 GPU 架构特性可进一步提升 Flash-DWC 的性能。（i）Warp 特化可以为消费者 warp 提供更大的寄存器预算，为生产者 warp 提供较小的寄存器预算，使消费者能够执行更大的子 GEMM，实现更高的数据复用。（ii）对于生产者，Flash-DWC 具有规则的预取模式，可以视为一系列连续的 128 字节段的跨步序列。因此，许多细粒度的 cp.async 指令可以由较少的多维 cp.async.bulk.tensor 传输替代。这可以将全局内存加载和部分地址生成工作卸载到 Tensor Memory Accelerator（TMA）上，实现更高的带宽和更好的重叠。（iii）在集群内，分布式共享内存（DSMEM）可能实现跨 CTA（块）的输入补丁重叠和共享滤波器片段的复用。（iv）对于消费者，多个 warp 级别的 MMA 指令可以重新组织以利用更大的 warpgroup 级别的 WGMMA 指令， potentially 提高 Tensor Core 利用率。

扩大 Tensor Core GEMM 粒度可能需要重新组织工作负载分布和 CTA 形状；它还可能增加累加器寄存器压力并降低占用率。这些开销可以通过 warp 特化、基于 TMA 的异步传输和 WGMMA 执行部分抵消。DSMEM 的收益将取决于能否在 GEMM 计算背后隐藏集群级别的 DSMEM 访问。

低位数据类型，如 FP8、FP6、FP4 和微缩放格式，通常以混合精度方式使用：操作数以较低精度存储和相乘，而累加保持较高精度，如 FP32 或 FP16，具体取决于精度需求。Flash-DWC-T 的数据复用模式、GEMV 到 GEMM 转换、矩阵压缩和双路径 split-GEMM 流水线在很大程度上是数据类型无关的。因此，只要 Tensor Cores 支持目标混合精度格式，Flash-DWC 就可以扩展到低位数据类型。

低位数据类型可以减少内存流量并在 Tensor Cores 上实现更高的峰值吞吐量，两者都可以改善 Flash-DWC-T 的性能。支持它们需要核函数级别的重新调优，包括针对 128 字节内存访问调整数据布局、选择合适的 Tensor Core 指令以及处理量化/反量化。这些修改可能影响软件流水线、内存访问行为和辅助 CUDA Core 开销，因此其对性能和精度的影响应针对特定部署进行验证。

8 相关工作
CNN 中的深度卷积。深度卷积（DWC）已成为现代 CNN 的关键组件。早期的轻量级网络 [12, 13, 39, 40] 采用小滤波器 DWC 以减少 FLOPs 同时保持精度。最近，大型滤波器 DWC 被采用以扩大感受野并建模长程依赖关系，实现了与视觉 Transformer 竞争的精度 [9, 22, 23, 43]。例如，ConvNeXt [25, 44] 使用 7×7 DWC，RepLKNet [7] 扩展到 31×31，SLaK [20] 进一步将条纹卷积推进到 51×51。PeLK [2] 表明，尽管 FLOPs 更高，稠密 DWC 仍然优于条纹变体，从而推动了参数高效的稠密设计。其他工作 [8, 15, 17, 26, 29, 41, 46, 47, 48] 也利用大型滤波器 DWC 来捕获全局上下文。尽管有这些架构上的进步，DWC 在 GPU 上仍然是内存受限的，限制了其实际训练效率，尤其是对于大滤波器。

DWC 优化。先前的工作已在 GPU 上优化 DWC。大多数 [4, 27, 28, 32, 34, 36, 37] 在 CUDA Cores 上运行，从而限制了性能，并且在混合精度下产生数据类型转换开销。传统布局，尤其是 NCHW [4, 28, 32, 36, 37]，阻碍了 DWC 的合并内存访问。[27, 28] 提出了有效的方法来复用输入补丁重叠，但它们仍然是面向 CUDA Core 的，并且主要针对小滤波器（例如 ≤5×5）。虽然 cuDNN [34] 可以在 Tensor Cores 上执行 DWC，但其 Tensor Core 性能通常低于 CUDA Core 性能。对角线重分解 [38] 将 DWC 转换为标准卷积，可以被高性能 BLAS 库加速；然而，其高度稀疏的转换滤波器引入了大量不必要的计算，使其比最近版本的 cuDNN 更慢。[37] 将深度卷积和逐点卷积融合以减少数据移动；它对小批次推理有效，但对于较大的批次大小（例如 16），由于重新计算和控制复杂性，效率降低。MegEngine [32] 为大滤波器 DWC 提供了两种方法：其直接卷积方法使用 SMEM 进行数据复用，但在 CUDA Cores 上运行且缺乏软件流水线；其批量 GEMM 方法将 FLOPs 增加了 IH×IW/(FH×FW) 倍，在大型特征图或小批次大小时可能效率较低。与优化 DWC 本身正交的是，硬件感知剪枝技术 [35] 算法上减少模型大小和 FLOPs。

FPGA 工作 [1, 18, 24, 30, 42] 设计专用数据流和核心以提高硬件利用率和片上数据复用。面向 CPU 的工作 [11, 49] 引入 ARM 特定的寄存器/缓存优化和核函数融合。[45] 通过基于扫描的分块和层融合来优化 CNN 推理加速器上的 DWC，减少 DRAM 访问。这些非 GPU 工作表明，DWC 加速主要依赖于数据局部性和硬件利用率，而不仅仅是 FLOPs。虽然它们的平台特定机制不能直接应用于 GPU，但我们共享这种由局部性驱动的观点，并研究其在 GPU 上的实现。

Tensor Coes 用于不规则计算。Tensor Coes 专为稠密 GEMM 设计，限制了其适用于稀疏或不规则计算的能力。最近的工作如 ConvStencil [3] 和 SparStencil [19] 将模板计算转换为具有超过 50% 结构化密度的 GEMM。然而，这些方法在小型单通道滤波器上运行，无法扩展到 DWC 的多通道大滤波器设置。

我们的工作。Flash-DWC 与先前的工作有三个方面的不同。（i）它为原生 DWC 引入了基于 GEMV 的数据流协同设计，联合优化数据复用、访问顺序、布局和工作负载分块，以提高算术强度，同时保持访问连续性和内存合并。（ii）它引入了 DWC 特定的 GEMV 到 GEMM 转换，将 DWC 映射到紧凑的 56.25% 稠密的 Tensor Core GEMM，从而实现高效的混合精度执行而无严重冗余。（iii）它将这些技术与资源高效的核函数优化相结合，以支持完整训练，包括反向传播和 PyTorch 集成。

9 结论
我们提出了 Flash-DWC，一个使 DWC 计算高效的 GPU 库。通过结合基于 GEMV 的数据流协同设计、紧凑的 GEMV 到 GEMM 转换以及资源高效的核函数优化，Flash-DWC 将 DWC 从内存受限操作转向计算受限操作，并将这种加速扩展到端到端的 CNN 训练。实验表明，Flash-DWC 在核函数级别评估和端到端训练中都实现了显著的加速。总体而言，Flash-DWC 使大滤波器 DWC 对于高效且精确的视觉应用更加实用。

致谢
作者感谢匿名审稿人的建设性评论，这些评论极大地改进了本稿件。我们向 Qingcai Jiang 表达诚挚的感谢，他的深刻讨论帮助塑造了这一研究方向。