循环层用户指南
- 循环操作可以并行化，如循环层中所述。我们建议使用 NVIDIA® cuDNN 实现，它会自动执行此操作。
- 当使用标准实现时，尺寸相关参数（小批量大小和隐藏大小）应为：
- 可被 8（对于 FP16）或 4（对于 TF32）整除，以便在 Tensor Cores 上高效运行；请参阅Tensor Core 要求。
- 可被至少 64 且理想情况下为 256 整除，以提高平铺效率；请参阅维度量化效果。
- 大于 128（小批量大小）或 256（隐藏大小），以便受计算速率而非内存带宽限制（NVIDIA V100；阈值因 GPU 而异）；请参阅数学和内存限制。
- 当使用持久实现（仅适用于 FP16 数据）时：
- 隐藏大小应可被 32 整除，以便在 Tensor Cores 上高效运行。通过选择可被更大的 2 的倍数（最多 256）整除的隐藏大小，可以实现更好的平铺效率。
- 小批量大小应可被 8 整除，以便在 Tensor Cores 上高效运行。如果前向传递期间小批量大小超过 96 或后向传递期间超过 32，则可能需要被更大的 2 的幂整除。
- 有关示例，请参阅案例研究：GNMT 的持久性。
- 尝试增加参数以获得更好的效率。例如，将小批量大小加倍通常会导致每个小批量在不到两倍的时间内完成两倍的计算。隐藏大小和小批量大小影响最大。如果序列长度非常小或小批量大小不符合指南，则增加序列长度也可能提高效率。分别参阅输入和隐藏大小和小批量大小和序列长度部分。
循环神经网络 (RNN) 是一种深度神经网络，其中输入数据和先前的隐藏状态都被馈送到网络的层中，从而为网络提供状态和内存。RNN 通常用于基于序列或基于时间的数据。
在训练期间，输入数据以某个小批量大小（小批量中数据序列的数量）和序列长度（每个序列中时间步的数量）馈送到网络。如图图 1所示，RNN 的每一层都由循环单元组成，其数量是该层的隐藏大小。
图 1. 具有两个输入和四个隐藏单元的单个循环层。RNN 通常由多个此类层组成；该层将从隐藏大小为二的层接收其输入。
循环单元有几乎无限数量的类型；许多框架甚至允许您定义自定义循环单元。但是，大多数网络使用一些常见的单元类型，包括长短期记忆 (LSTM) 单元或门控循环单元 (GRU)。这些单元以及较旧且更简单的 ReLU 和 tanh 激活单元都受 cuDNN API 支持。在本指南中，我们重点介绍 LSTM 单元，但使用其他单元类型的网络行为类似。
循环操作不像本指南中讨论的其他一些操作那样容易表示为 GEMM。与每个输入张量独立处理的前馈网络不同，循环网络在以后的时间步长对较早的时间步长具有固有的依赖性。
图 2. 跨 4 个时间步的 3 层 RNN。垂直箭头表示层到层路径中的计算，这取决于先前层的激活。水平箭头表示循环路径中的计算，这取决于先前时间步长的结果。
在不考虑并行化的情况下，循环层的训练可以表示为一组单独的矩阵乘法。每个门（LSTM 单元为四个，GRU 为三个，ReLU 和 Tanh 单元为一个）的每个训练阶段（前向、激活梯度和权重梯度）都等效于一个维度为 1，另两个维度等于隐藏大小的 GEMM。
我们希望并行化这些计算，以最大化 GPU 可用的工作量。当使用 LSTM 单元或 GRU 时，可以连接不同门的 GEMM。GEMM 也可以在小批量中组合，并且如果时间依赖性允许，则可以在序列步骤中组合。下表显示了所得维度示例。
使用 cuDNN，包括cudnnRNNForwardTraining()等函数，这是自动完成的；这些实现也可以通过 PyTorch、TensorFlow 和 MXNet 等框架用于某些类型的循环层。相关技巧可在参数和性能中找到。另一方面，如果 cuDNN 不适用于您的应用程序，则可以手动实现本节中描述的并行类型以提高性能。
| 路径 | 计算阶段 | M | N | K | GEMM 数量 |
|---|---|---|---|---|---|
| 循环 | 前向传播 | 4 x H | B | H | S |
| 激活梯度 | H | B | 4 x H | S | |
| 权重梯度 | H | 4 x H | B x S | 1 | |
| 层到层 | 前向传播 | 4 x H（本层） | B x F | H（上一层） | 1 |
| 激活梯度 | H（本层） | B x F | 4 x H（下一层） | 1 | |
| 权重梯度 | H（上一层） | 4 x H（本层） | B x S | 1 |
此处我们显示了由 LSTM 单元组成的循环层的维度（图 3和图 4），其中连接了四个 GEMM；对于 ReLU/Tanh 单元和 GRU，分别有一个和三个 GEMM。对于循环路径中的前向和激活梯度传递，每个序列步骤都依赖于前一个步骤。我们可以将这些 GEMM 在小批量大小上组合，但不能在不同的序列步骤上组合。
图 3. 循环（a）前向和（b）激活梯度在 LSTM 单元的一层上的等效 GEMM 维度。
相比之下，在层到层路径中，每个序列步骤仅依赖于上一层的输出。因此，我们可以将 GEMM 在小批量大小和多个序列步骤上组合。通常，在整个序列长度上组合效率不高。为每个 GEMM 组合的序列步骤数称为展开因子，并由启发式方法自动选择。
图 4. 层到层（a）前向和（b）激活梯度在 LSTM 单元的单个层对之间的等效 GEMM 维度。
同样，对于两个路径中的权重梯度传递，序列步骤之间不存在依赖性；这些 GEMM 也可以安全地在小批量大小和序列步骤上组合。在此传递期间，GEMM 在整个序列长度上组合。
图 5. （a）循环和（b）层到层权重梯度传递的等效 GEMM 维度，同样使用 LSTM 单元。
在循环和层到层计算中，相同的权重用于计算跨批次和序列的激活。当隐藏大小足够小，以至于权重矩阵可以在本地缓存而不是重复加载时，可以使用具有持久权重的实现。当小批量大小较小时，此持久实现往往优于默认实现；根据您的框架，可以自动选择它。
层的隐藏大小，以及对于层到层操作，相邻层的隐藏大小，对每个 RNN 操作的三个等效 GEMM 维度中的两个有贡献。对于前向和激活梯度传递，这些维度是 M 和 K；对于权重梯度传递，M 和 N。剩余维度由小批量大小和/或序列长度控制。因此，隐藏大小是控制性能的最重要参数之一。
图 6. 对于更大的输入和隐藏节点数量，计算效率更高。权重梯度传递比前向或激活梯度传递执行得更好，因为相应的 GEMM 往往具有更相似的 M 和 N，以及对于循环和层到层路径都足够大的 K。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。
当一个或多个维度较小时，GEMM 的性能往往不佳。对于高达约 256 个单元的隐藏大小，GEMM 很可能受内存限制。还可能发生平铺和 GPU 占用问题。在这些尺寸下，通常可以增加隐藏大小，而对训练持续时间几乎没有影响。如果较小的隐藏大小不可协商，则可能值得在网络的多个层之间实现并行性。这可以在 cuDNN API 中通过使用cudnnSetRNNDescriptor()定义层堆栈来完成；无法自行有效利用 GPU 的层会自动并行计算。
4.2. 小批量大小和序列长度
小批量大小始终至少对一个 GEMM 维度做出贡献，因此较大的小批量往往更有效率。在前向和激活梯度传递期间，“N”维度等于小批量大小（对于循环路径）或小批量大小和序列长度的乘积（对于层到层路径）。这导致了一些有趣的效果：小批量大小往往比序列长度更强烈地影响性能。即使对于较大的序列长度，小批量大小也可能导致性能不佳。
图 7. 对于前向和激活梯度传递，“N”维度取决于小批量，而在层到层计算中，则取决于序列长度。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。
对于权重梯度传递，“K”维度等于循环和层到层路径的小批量大小和序列长度的乘积。因此，对于此维度的相等值，无论小批量大小和序列长度的个别值如何，性能都相似。
图 8. 在权重梯度传递期间，“K”维度对于两个路径都是相同的。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1。
Google 神经机器翻译系统 (GNMT)和许多其他语言处理网络都使用了堆叠的循环层。GNMT 的架构由编码器、解码器和注意力模块组成，编码器和解码器都包含堆叠的 LSTM 层。
在本节前面，我们提到当小批量大小较小时，像这样的循环层的计算速度可能会受到标准实现的数据移动的限制。一种替代方案是持久实现，其中权重矩阵在本地缓存以避免从内存重复加载。在 cuDNN 中，可以使用cudnnRNNAlgo_t选择此实现。一些框架也提供了它，包括 PyTorch，我们在此示例中使用它。持久实现无法使用 PyTorch 手动选择；相反，当预期它能提高性能时，会自动使用它。
关于持久性，有三个主要的限制需要牢记。首先，持久实现仅在使用 FP16 时可用。其次，权重矩阵必须足够小，以便在迭代之间本地缓存。这意味着隐藏大小必须低于阈值（具体值取决于单元类型）。当使用 256 的序列长度和 64 的小批量大小训练 GNMT 时，持久性可以与隐藏大小高达 1024 个单元的 LSTM 层一起使用。确保层足够小以使用持久性可以显着提高性能（图 9）。
图 9. 对于这种类型的循环层，可以选择隐藏大小为 1024 及更低的持久实现。单元过多的层不会从持久性中受益。为进行比较，显示了始终使用标准模式的 cuDNN 执行时间。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1，PyTorch 1.8。
第三，当小批量大小较小时，持久实现通常优于标准实现。图 10a 显示了 LSTM 层在 1024 个单元的训练性能，使用 256 的序列长度，在标准模式和持久模式下跨小批量大小；对于少于 256 个元素的小批量，持久计算明显更快。在此范围内，我们建议选择隐藏大小以启用持久模式（以及所有参数以启用 Tensor Cores，如快速入门清单中所述），只要有可能。对于较大的小批量，标准实现的性能与持久实现相似或更好，则没有必要这样做。PyTorch 在实现之间切换以获得最佳性能，随着小批量增加（图 10b）。
图 10. （a）cuDNN 的训练性能，包括标准实现和持久实现，以及（b）PyTorch 的训练性能，随着小批量大小的变化自动在标准实现和持久实现之间切换。持久实现对于小批量尤其有效。对于较大的小批量，标准实现更快，PyTorch 会选择标准实现。NVIDIA A100-SXM4-80GB，CUDA 11.2，cuDNN 8.1，PyTorch 1.8。
声明
本文档仅供参考，不得视为对产品特定功能、状况或质量的保证。NVIDIA Corporation（“NVIDIA”）对本文档中包含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。对于因使用此类信息或因使用此类信息而可能导致的侵犯第三方专利或其他权利的行为，NVIDIA 概不负责。本文档不构成对开发、发布或交付任何材料（如下定义）、代码或功能的承诺。
NVIDIA 保留在不另行通知的情况下随时对本文档进行更正、修改、增强、改进和任何其他更改的权利。
客户应在下订单前获取最新的相关信息，并应验证此类信息是否为最新且完整。
NVIDIA 产品的销售受 NVIDIA 标准销售条款和条件（在订单确认时提供）的约束，除非 NVIDIA 和客户的授权代表签署的个别销售协议中另有约定（“销售条款”）。NVIDIA 特此明确反对将任何客户通用条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档不直接或间接地形成任何合同义务。
NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或失灵可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对在上述设备或应用中包含和/或使用 NVIDIA 产品不承担任何责任，因此，此类包含和/或使用由客户自行承担风险。
NVIDIA 不作任何陈述或保证，保证基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每种产品的全部参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合并符合客户计划的应用，并为该应用执行必要的测试，以避免应用或产品的默认设置。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同条件和/或要求。NVIDIA 对可能基于或归因于以下原因的任何默认设置、损坏、成本或问题不承担任何责任：（i）以任何违反本文档的方式使用 NVIDIA 产品，或（ii）客户产品设计。
本文档未授予 NVIDIA 专利权、版权或其他 NVIDIA 知识产权下的任何明示或暗示的许可。NVIDIA 发布的有关第三方产品或服务的信息不构成 NVIDIA 授予的使用此类产品或服务的许可，也不构成 NVIDIA 对其的保证或认可。使用此类信息可能需要获得第三方在其专利或其他知识产权下的许可，或者获得 NVIDIA 在其专利或其他知识产权下的许可。
仅当事先获得 NVIDIA 书面批准，在不作更改且完全遵守所有适用的出口法律和法规的情况下复制本文档中的信息，并附带所有相关的条件、限制和声明，才是允许的。
本文档和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”）均按“原样”提供。NVIDIA 对材料不作任何明示、暗示、法定或其他形式的保证，并且明确声明不承担所有关于不侵权、适销性和特定用途适用性的暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论如何造成，也无论责任理论如何）承担责任，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因而遭受任何损害，但 NVIDIA 对本文所述产品的面向客户的累积总责任应根据产品的销售条款进行限制。
Android、Android TV、Google Play 和 Google Play 徽标是 Google, Inc. 的商标。
商标
NVIDIA、NVIDIA 徽标、CUDA、Merlin、RAPIDS、Triton Inference Server、Turing 和 Volta 是 NVIDIA Corporation 在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。