卷积层用户指南 - NVIDIA 文档

摘要

本指南提供提高卷积层性能的技巧。它还详细介绍了参数（包括批量大小、输入和滤波器维度、步幅和空洞率）的影响。

1. 快速入门清单

以下快速入门清单提供了卷积层的具体技巧。

选择可被 8（对于 FP16）或 4（对于 TF32）整除的输入和输出通道数，以便在 Tensor Core 上高效运行。对于大多数 CNN 中的第一个卷积层，其中输入张量由 3 通道图像组成，如果使用步幅为 2，则填充到 4 通道就足够了；请参阅输入和输出通道。
选择可被至少 64 且理想情况下被 256 整除的参数（批量大小、输入和输出通道数），以实现高效平铺并减少开销；请参阅量化效应。
尺寸相关参数（批量大小、输入和输出高度和宽度以及输入和输出通道数）的较大值可以提高并行化程度。与全连接层一样，这提高了运算的效率，但不会减少其绝对持续时间；请参阅卷积参数如何影响性能和子章节。
NVIDIA^® 库提供了一组不同的卷积算法，这些算法具有不同的性能表现，具体取决于卷积的参数。当网络处理的输入大小在每次迭代中相同时，自动调优是一种有效的方法，可确保为网络中的每个卷积选择理想的算法。对于 TensorFlow，默认情况下启用自动调优。对于 PyTorch，通过将 torch.backends.cudnn.benchmark = True 添加到您的代码来启用自动调优。
选择内存中的张量布局以避免转置输入和输出数据。有两种主要的约定，每种约定都以维度的顺序命名：NHWC 和 NCHW。我们建议尽可能使用 NHWC 格式。有关更多详细信息，包括框架支持，请参阅内存中的张量布局：NCHW 与 NHWC。

2. 简介

卷积由输入和滤波器张量的大小以及卷积的行为（例如使用的填充类型）定义。

图 1 说明了定义卷积所需的最小参数集。

图 1. NCHW 输入张量与 KCRS 权重张量的卷积，生成 NKPQ 输出。

在本指南的其余部分，我们将对每个参数使用单字母缩写。

注意

此处的 N 和 K 与通用矩阵乘法 (GEMM) 中的 N 和 K 参数无关；GEMM 维度在此处表示为 M x N x K，以使其与卷积的参数区分开来。GEMM 和 GEMM 性能在NVIDIA 矩阵乘法背景用户指南中进行了解释。

表 1. 定义卷积的参数
参数	张量	含义
`N`	不适用	批量大小
`C`	输入	通道数
`H`		高度
`W`		宽度
`K`	输出	通道数
`P`		高度（通常从其他参数派生）
`Q`		宽度（通常从其他参数派生）
`R`	滤波器	高度
`S`		宽度
`U`		垂直步幅
`V`		水平步幅
`PadH`		垂直维度中的输入填充
`PadW`		水平维度中的输入填充
`DilH`		垂直维度中的空洞率
`DilW`		水平维度中的空洞率

3. 卷积算法

NVIDIA cuDNN 库使用两种主要方法实现卷积：基于隐式 GEMM 和基于变换。

隐式 GEMM 方法是直接卷积的变体，直接对输入权重和激活张量进行运算。或者，可以通过将数据和权重转换为另一个空间、执行更简单的运算（例如，逐点乘法），然后再转换回来来计算卷积。cuDNN 库提供了一些使用 FFT 和 Winograd 变换的卷积实现。

3.1. 使用 cuDNN 选择卷积算法

使用 cuDNN 运行卷积时，例如使用cudnnConvolutionForward()，您可以指定使用哪种通用算法。

NVIDIA cuDNN API 参考提供了用于估计不同算法相对性能的函数。一组以 cudnnGet 为前缀的函数使用一组启发式方法来预测可用算法的相对性能。这些函数评估速度很快；但是，尽管我们不断改进我们的启发式方法，但预测可能并不总是准确的。偶尔可能会发生次优的算法选择，这在不寻常类型的卷积和极端情况下更为常见。

另一组以 cudnnFind 为前缀的函数，测试并报告所有可用算法的性能，以确定给定卷积运算的最有效选项。使用这些函数的好处是选择的算法是最佳选择。但是，由于正在运行实际的性能测试，因此这些函数可能非常耗时且资源密集。

在选择算法后，我们的启发式方法会指定其他底层细节；例如，平铺大小，在量化效应和维度量化效应部分（在矩阵乘法背景用户指南中）中进行了详细讨论。

从 cuDNN 版本 8 开始，NVIDIA cuDNN 后端 API 允许更精确地控制平铺大小和所用算法的其他参数。

4. Tensor Core 使用和性能建议

NVIDIA Tensor Core GPU 使用的主要（不使用变换）执行卷积的方法称为隐式 GEMM。它执行与直接卷积完全相同的数学运算次数，因此在计算上是等效的。

隐式 GEMM 原生地对卷积输入张量进行运算，将计算动态转换为矩阵乘法。重要的是要注意，永远不会在内存中创建相应的矩阵。因此，要计算算术强度，可以使用原始张量大小。

为了说明卷积作为矩阵乘法的概念，让我们首先考虑将卷积滤波器应用于输入数据的单个应用。假设我们将 3x3 卷积应用于 128 通道输入张量。为了计算单个输出值，我们有效地计算两个 1,152 元素向量的点积。一个是滤波器的权重 (3x3x128 = 1,152)。另一个由与权重相乘的数据（激活）值组成，以产生输出。由于并非所有 1,152 个数据值在内存中都是连续的，因此读取原始张量布局，并在动态将其转换为适当的向量形式。为了计算所有输出，我们执行多次点积，这可以看作是矩阵乘法，但由于矩阵是隐式形成的，而不是在内存中创建的，因此这种方法称为隐式 GEMM。但是，要理解卷积的性能，了解这些“虚拟”矩阵的形状和大小可能很有用。

4.1. 内存中的张量布局：NCHW 与 NHWC

卷积通常对四维张量进行运算：由 N 个“图像”组成的批次，这些“图像”具有 C 个通道的 H x W 特征图。

深度学习框架通常在内存中使用 NCHW 和 NHWC 布局（首字母缩略词列出了内存中从最慢到最快变化的维度）。布局选择会影响性能，因为为 Tensor Core 实现的卷积需要 NHWC 布局，并且当输入张量以 NHWC 布局时速度最快。

注意

Tensor Core 仍然可以对 NCHW 布局进行运算，但由于自动转置运算而产生一些开销，如图 2 所示。当输入和输出张量较大或所需计算量较低时（例如当滤波器尺寸较小时），转置开销往往更显着。为了最大限度地提高性能，我们建议使用 NHWC 张量布局。

图 2. 不需要转置的内核 (NHWC) 比需要一个或多个转置的内核 (NCHW) 性能更好。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

在撰写本文时，包括 MXNet（通过 layout 参数，此处的文档页面）、TensorFlow（通过 data_format 参数，此处的文档页面）和 PyTorch（使用 Channels Last 内存格式，此处的文档页面）等框架中提供了 NHWC 布局。除非另有说明，否则本节中的性能示例可以假定使用 NHWC 布局中的输入和输出数据。

4.2. 隐式 GEMM 维度

现在让我们考虑在执行前向卷积、计算激活梯度和计算权重梯度时遇到的矩阵维度。

表 2. 卷积参数到相应 GEMM 参数的转换
计算阶段	"M"	"N"	"K"
前向传播	NPQ	K	CRS
激活梯度	NHW	C	KRS
权重梯度	CRS	K	NPQ

下面显示了“虚拟”矩阵的组成。对于每个pass，都有一个虚拟矩阵，如果显式构造，它将包含比其对应张量更多的值。例如，在前向卷积期间，A 矩阵 (N*P*Q x C*R*S) 由输入激活（维度为 N x H x W x C 的张量）组成。每个单独的输入激活在矩阵中出现 R*S 次，并以必要的偏移量重复，以使该输入值与匹配的 R x S 滤波器通道的重叠值相乘。在计算输入激活梯度时，输出激活梯度会发生类似的 conceptual 扩展，在权重梯度计算期间，输入激活也会发生类似的 conceptual 扩展。

图 3. (a) 前向卷积、(b) 激活梯度计算和 (c) 权重梯度计算的等效 GEMM 的维度

重要的是要重申，这些大小的矩阵不会存储在内存中；它们是帮助解释计算的抽象。“重复”值不是字面上复制的，并且避免了从内存中浪费读取。这在算术强度的计算中直接可见。(前向) 隐式 GEMM 读取两个大小为 NCHW 和 KCRS 的输入张量，并使用每个 NKPQ 输出张量的 CRS 乘法和加法生成一个大小为 NKPQ 的输出张量，总共进行 NKPQCRS 次乘加运算。因此，FP16 中每个元素 2 字节的算术强度为

$算术强度 = \frac{# 运算}{# 字节} = \frac{2 \cdot (N \cdot K \cdot P \cdot Q) \cdot (C \cdot R \cdot S)}{2 \cdot (N \cdot C \cdot H \cdot W + K \cdot C \cdot R \cdot S + N \cdot K \cdot P \cdot Q)} = \frac{N \cdot K \cdot P \cdot Q \cdot C \cdot R \cdot S}{N \cdot C \cdot H \cdot W + K \cdot C \cdot R \cdot S + N \cdot K \cdot P \cdot Q}$

例如，在 256x56x56x64 输入张量上计算 3x3 卷积，生成 256x56x56x128 输出，所有操作均以半精度进行，算术强度为 383.8 FLOPS/字节。

4.3. 量化效应

平铺和波量化效应可能非常显着，尤其对于小问题规模（背景在维度量化效应中）。就像 GEMM 一样，在隐式 GEMM 中，输出矩阵的表示形式被划分为选定大小的平铺，并且这组平铺分布在可用的多处理器上。

我们的测试 GPU 有 108 个 SM。每个 SM 可以并行处理多个线程块，这取决于正在使用的内核；为了获得最佳并行化，隐式 GEMM 应包含 108 个平铺的整数倍。

图 4. 图表显示了滤波器大小为 3x3、输入大小为 16x16、4096 个输入通道和 256 个输出通道的卷积性能。当 N 可被 108 整除时（当创建 216 个平铺的倍数，每个 SM 上并行两个时），观察到最佳性能。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

正如图 4 所示，导致小等效 GEMM 的卷积可能会表现出显着的量化效应。当 N = 54 时，创建 216 个平铺；当 N = 55 时，创建 220 个平铺。前者导致 Tensor Core 高利用率，而后者将需要额外的波来处理剩余的 4 个平铺，从而严重影响性能（图 4 (a)）。一旦卷积足够大，效果就不那么明显了（图 4 (b)）。

值得注意的是，权重梯度量化的行为方式与图 3 中的 GEMM 维度所暗示的不同；出于量化目的，要平铺的矩阵的高度应视为 C（而不是 C*R*S）。这在关于滤波器尺寸的部分末尾进行了更详细的讨论。

4.4. 卷积参数如何影响性能

在本节中，我们将讨论影响性能的趋势。为了简单起见，填充设置为 H = P 和 W = Q，并且步幅和空洞率都等于 1，除非另有说明。

4.4.1. 批量大小、高度和宽度

当将前向卷积表示为 GEMM 时，批量大小、输出高度和输出宽度的乘积 (N*P*Q) 是展开的输入张量（A 矩阵）以及输出（C 矩阵）的“M”维度。

这些参数的单个值对于 GEMM 性能而言并非特别重要；只有最终维度，即乘积 N*P*Q，才具有重要意义。方便的是，在大多数应用程序中，批量大小比有助于输出张量的高度和宽度的参数更容易更改。

通常，效率随着 N*P*Q 的增加而提高，但收益递减。从图 5 中，我们可以看到具有等效 N*P*Q 的点具有大致相当的性能，因为相应的 GEMM 具有相同的维度。

图 5. 同一数据的不同视角。性能随着 N*P*Q 的增加而提高。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

在计算激活梯度时，N*H*W 是等效 GEMM 的“M”维度（与前向卷积中的 N*P*Q 相比）。在滤波器步幅为 1 的情况下，前向卷积和激活梯度计算的性能大致相同。

相反，对于权重梯度计算，N*P*Q 变为累积（GEMM 中的“K”）维度。此维度的性能影响并不那么直接，因为它不会以任何方式影响输出矩阵的平铺。但是，较大的 N*P*Q 值通常会导致更多时间用于乘法和累加元素，而不是 GEMM 计算的设置和拆卸开销，从而提高了整个运算实现的峰值性能的比例。值得注意的是，对于权重梯度，cuDNN 通常也支持在 N*P*Q 维度中进行平铺，因为许多常见层会导致很小的（例如，标准 ResNet 的第一个块中的 64x64）输出矩阵，这些矩阵本身无法提供足够的平铺并行性来保持 GPU 的繁忙状态。

4.4.2. 滤波器尺寸

前向卷积的等效 GEMM 的“K”维度为 C*R*S。如前所述，“K”维度确实会对性能产生影响，并且这种效应对于小 GEMM 最为明显。

当使用 1x1 滤波器时，具有更多输入通道的层在前向卷积中往往表现更好（图 6），因为它最终是 C*R*S 乘积起作用。

图 6. 具有较大滤波器的卷积往往表现更好（C*R*S 很重要）。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

在计算卷积的激活梯度时，K 会影响此维度：“K” = K*R*S。图 6 和图 7 之间存在明显的相似性；通常，前向卷积的 C 相关趋势与激活梯度计算的 K 相关趋势相关，反之亦然。

图 7. 与上图几乎相同的关系，但涉及 K。对于激活梯度，K*R*S 很重要。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

权重梯度计算具有 “M” = C*R*S，因此滤波器尺寸对性能的影响与先前讨论的批量大小、高度和宽度的影响类似；较大的值往往表现更好。但是，在考虑平铺量化时，权重梯度算法与前向和数据梯度算法不同；只有 C 维度，而不是完整的 C*R*S 维度，针对平铺大小进行量化（这意味着，可以忽略滤波器尺寸参数）。例如，当使用 64x64 平铺时，如果 C = 32，则每个平铺的一半（垂直方向）都被浪费了，而与 R 和 S 无关；只有 C 的值才重要。

4.4.3. 输入和输出通道

启用 Tensor Core 的要求取决于使用的 cuDNN 版本。使用 cuDNN v7.6.3 及更高版本，卷积维度将在必要时自动填充，以利用 Tensor Core。早期版本的 cuDNN 更严格：将 Tensor Core 与 NHWC 打包数据一起使用需要 C 和 K 与 TF32 的 4 倍数、FP16 的 8 倍数或 INT8 的 16 倍数对齐。对于 NCHW 打包的 FP16 数据，通道将自动填充为 8 的倍数，以便启用 Tensor Core。但是，将 NCHW 数据与启用 Tensor Core 的内核一起使用会产生一些额外的转置成本，这在内存中的张量布局：NCHW 与 NHWC中进行了讨论。

另一方面，对于这些早期版本的 cuDNN，自动填充不适用于 NHWC 打包数据，因此选择效率较低的后备内核，该内核不使用 Tensor Core。假设 C 和 K 可被 8 整除，则使用 NHWC 数据的卷积确实比使用 NCHW 数据的卷积性能更好。换句话说，如果某个层已与 NCHW 数据一起使用，则会自动进行填充；但是，如果正在使用 NHWC 数据，则选择或填充 C 和 K 为 8 的倍数可以提高性能。

对于 cuDNN v7.6.3 及更高版本，填充是自动的，与数据格式无关。填充会增加一些时间，尽管与启用 Tensor Core 带来的性能提升相比，此成本通常可以忽略不计。值得注意的是，选择 C 和 K 为 FP16 的 8 倍数或与其他数据类型等效的倍数可以获得最佳效率：对于这些情况，不需要填充。

在某些情况下，通道计数很小且不可协商。对于网络中的第一层，通常 C 的值非常小（对于灰度和 RGB 或 YCrCb 图像，分别为 1 或 3）。特殊的卷积实现可满足此需求，特别是对于 C = 4 和步幅为 2 的情况（图 8）。此处显示的数据是使用 cuDNN 8.1 收集的，因此填充是自动的。从 C = 3 到 C = 4 的性能提升不如 7.6.3 之前的版本那么显着，但选择 C = 4 仍然更快，因为没有发生填充。

图 8. C = 4 的专用内核加速了卷积神经网络中的常见第一层（使用 NHWC 数据）。选择 C = 4 或 8 的倍数可获得最佳性能。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

先前在滤波器尺寸（“K” = C*R*S）中讨论了相对于 C 的前向卷积性能，以及相对于 K (“K” = K*R*S) 的激活梯度计算。在批量大小、高度和宽度中提到了 C 对权重更新性能的影响（使用 “M” = C*R*S）。简而言之，较大的值通常会提供更高的效率，但收益递减。

但是，输入和输出通道的数量可能对性能产生更直接的影响；对于前向卷积、激活梯度计算和权重梯度计算，GEMM 维度“N”等于 C 或 K。因此，您对这些参数的选择可能会对性能产生直接影响。

在 K 中变化时，在前向卷积和权重梯度计算中可以看到类似的行为（图 9）。在这两种情况下，“N” = K，对于小的 K 值，会导致强烈的趋势，一旦 K 大于大多数平铺大小，收益就会递减。

图 9. 对于较大的 K，前向卷积和权重梯度计算性能要好得多，直到某个点。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

激活梯度计算中输入通道也存在相同的效果（“N” = C），如图 10 所示。如前所述，C 对激活梯度计算的影响往往与 K 对前向卷积的影响相匹配。

图 10. 激活梯度计算性能随着 C 的增加而提高，但收益递减。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

4.5. 步幅

滤波器步幅（U 和 V）主要通过它们对输入和输出张量维度的影响来影响性能。使用水平和垂直步幅 1，H 和 W 分别与 P 和 Q 大致相等，具体取决于滤波器尺寸和填充。但是，当使用较大的步幅时，输入和输出特征图的大小存在几倍的差异。反过来，这会影响 GEMM 维度和性能。

图 11. 只要输出高度和宽度恒定，前向卷积和权重梯度计算的性能相对不受步幅或输入高度和宽度的变化影响。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

对于相等的 P 和 Q，前向卷积和权重梯度计算表现相似，尽管 H、W、U 和 V 各不相同（图 11）。

图 12. 激活梯度计算性能并非完全由输出维度决定；输入高度和宽度也会产生影响。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

相比之下，激活梯度计算受输入特征图大小的影响更大（图 12）。

4.6. 高性能示例

高性能卷积的示例如图 13 所示。此场景基于卷积层，该卷积层具有大小为 64x64 的输入特征图、大小为 3x3 的滤波器、1024 个输入通道和 1024 个输出通道；每个参数都相当大。使用 NHWC 数据以避免来自额外转置的开销。输入和输出通道数可被 8 整除，因此将启用 Tensor Core。

图 13. 在这里，对于较大的批量大小的卷积，可以看到良好的性能，对于所有三个 pass，都接近 250 TFLOPS。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

在较大的批量大小下，这种情况下的前向卷积以及激活和权重梯度计算的性能约为 250 TFLOPS。

5. 卷积变体

5.1. 空洞卷积

空洞卷积是常规卷积层的变体，它通过在滤波器元素之间插入零来有效地扩展正在应用的滤波器。

空洞率比每对元素之间添加的零数大 1。因此，与滤波器每个通道重叠的整体 2D 区域会增加。

图 14. 3x3 滤波器的空洞；空洞率分别为 (a) 1、(b) 2 和 (c) 3。

$r_{有效} = {空洞率}_{h} \cdot (r - 1) + 1$

$s_{有效} = {空洞率}_{w} \cdot (s - 1) + 1$

空洞率的选择会影响卷积如何表示为虚拟 GEMM，但实际上不会更改该 GEMM 的维度；因此，无论空洞率如何，前向卷积和激活梯度计算的性能都相似（如下所示）。

图 15. 对于非空洞卷积和空洞率为 2 的卷积，前向卷积、激活梯度计算和权重梯度计算的性能非常相似。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。

声明

本文档仅供参考，不应视为对产品的特定功能、条件或质量的保证。NVIDIA Corporation（“NVIDIA”）对本文档中包含的信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因使用此类信息而可能导致的侵犯专利或第三方的其他权利的后果或使用不承担任何责任。本文档不是开发、发布或交付任何材料（如下定义）、代码或功能的承诺。

NVIDIA 保留随时对此文档进行更正、修改、增强、改进和任何其他更改的权利，恕不另行通知。

客户在下订单之前应获取最新的相关信息，并应验证此类信息是最新的和完整的。

英伟达产品根据订单确认时提供的英伟达标准销售条款和条件进行销售，除非英伟达授权代表和客户签署的个别销售协议（“销售条款”）另有约定。英伟达特此明确反对将任何客户通用条款和条件应用于本文件中提及的英伟达产品的购买。本文件不直接或间接地形成任何合同义务。

英伟达产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命维持设备，也不适用于英伟达产品发生故障或失灵时可合理预期会导致人身伤害、死亡或财产或环境损害的应用。英伟达对在此类设备或应用中包含和/或使用英伟达产品不承担任何责任，因此，包含和/或使用此类产品的风险由客户自行承担。

英伟达不声明或保证基于本文档的产品将适用于任何特定用途。英伟达不一定会对每种产品的所有参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合客户计划的应用，并为该应用执行必要的测试，以避免应用或产品出现故障。客户产品设计的缺陷可能会影响英伟达产品的质量和可靠性，并可能导致超出本文档所含内容的附加或不同条件和/或要求。英伟达对任何可能基于或归因于以下原因的任何故障、损坏、成本或问题不承担任何责任：(i) 以违反本文档的任何方式使用英伟达产品或 (ii) 客户产品设计。

本文档未授予任何明示或暗示的许可，以使用任何英伟达专利权、版权或其他英伟达知识产权。英伟达发布的有关第三方产品或服务的信息不构成英伟达授予的使用此类产品或服务的许可，也不构成对此类产品或服务的保证或认可。使用此类信息可能需要获得第三方的专利或其他知识产权许可，或获得英伟达的专利或其他知识产权许可。

只有在事先获得英伟达书面批准的情况下，才可以复制本文档中的信息，并且复制必须未经修改，完全符合所有适用的出口法律和法规，并附带所有相关的条件、限制和声明。

本文档和所有英伟达设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”，单独或统称）均“按原样”提供。英伟达对这些材料不作任何明示、暗示、法定或其他方面的保证，并明确否认所有关于不侵权、适销性和特定用途适用性的默示保证。在法律未禁止的范围内，在任何情况下，英伟达均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害）承担责任，无论其由何种原因引起，也无论其基于何种责任理论，即使英伟达已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，英伟达对客户就本文所述产品承担的总体和累积责任应根据产品的销售条款进行限制。

Google

Android、Android TV、Google Play 和 Google Play 徽标是 Google, Inc. 的商标。

商标

_{NVIDIA、NVIDIA 徽标、CUDA、Merlin、RAPIDS、Triton Inference Server、Turing 和 Volta 是 NVIDIA Corporation 在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。}