是如何减速深度学习模型的训练和推理环节的 汽车长翅膀 GPU

作者 |Lucas de Lima Nogueira

编译|岳扬

Image by the author with the assistance of AI (​ ​​ ​)

现如今，当咱们提及深度学习时，人们人造而然地会联想到经过 GPU 来增强其性能。

GPU（图形解决器，Graphical Processing Units）后来是为了减速图像（images）及 2D、3D 图形（graphics）的渲染而生。但仰仗其弱小的并行运算才干，GPU 的运行范围迅速拓展，已裁减至深度学习（deep learning）等运行畛域。

GPU 在深度学习模型中的运行始于 2000 年代中前期，2012 年AlexNet的横空入世更是将这种趋向推向高潮。 AlexNet，这款由Alex Krizhevsky、Ilya Sutskever 和 Geoffrey Hinton 共同设计、研发的卷积神经网络，在 2012 年的 ImageNet Large Scale Visual Recognition Challenge (ILSVRC) 上一举成名。这一胜利具有里程碑式的意义，它不只证明了深度神经网络在图像分类畛域（image classification）的出色性能，同时也彰显了经常使用 GPU 训练大型模型的有效性。

在这一技术打破之后，GPU 在深度学习模型中的运行愈发宽泛，PyTorch 和 TensorFlow 等框架应运而生。

如今，咱们只有在 PyTorch 中轻敲 .to("cuda")，即可将数据传递给 GPU，从而减速模型的训练。但在通常中，深度学习算法终究是如何奇妙地利用 GPU 算力的呢？让咱们一探求竟吧！

深度学习的外围架构，如神经网络、CNNs、RNNs 和 transformer，其实质都围绕着矩阵加法（matrix addition）、矩阵乘法（matrix multiplication）以及对矩阵运行函数（applying a function a matrix）等基本数学操作开展。因此，优化这些外围运算，便是优化深度学习模型性能的关键所在。

那么，让咱们从最基础的场景说起。构想一下，你要求对两个向量口头相加操作 C = A + B。

可以用 C 言语繁难成功这一配置：

不难发现，传统上，计算机需逐个访问向量中的各个元素（elements），在每次迭代中按顺序对每对元素启动加法运算。但有一点要求留意，各对元素间的加法操作互不影响，即恣意一对元素的加法不依赖于其它任何一对。那么，若咱们能同时口头这些数学运算，成功一切元素对（pairs of elements）的并行相加，成果会如何呢？

间接做法是借助 CPU 的多线程配置，并行口头一切数学运算。但在深度学习畛域，咱们要求解决的向量规模渺小，往往蕴含数百万个元素。通常状况下，普通 CPU 只能同时解决十几条线程。此时，GPU 的长处便凸显进去！目前的干流 GPU 能够同时运转数百万个线程，极大地提高了解决大规模向量中数学运算的效率。

01 GPU vs. CPU comparison

虽然从单次运算（single operation）的解决速度来看，CPU 或者略胜 GPU 一筹，但 GPU 的长处在于其出色的并行解决才干。究其根源，这一状况源于两者设计初衷的差异。 CPU 的设计并重于高效口头繁多序列的操作（即线程（thread）），但一次性仅能同时解决几十个；相比之下，GPU 的设计指标是成功数百万个线程的并行运算，虽有所就义单个线程的运算速度，却在全体并行性能上成功了质的飞跃。

打个比如，你可以将 CPU 视作一辆炫酷的法拉利（Ferrari）跑车，而 GPU 则似乎一辆宽阔的公交车。倘若你的义务仅仅是运送一位乘客，毫无不懂，法拉利（CPU）是最佳选用。但是，如若的运输需求是运送多位乘客，即使法拉利（CPU）单程速度占优，公交车（GPU）却能一次性容纳所有乘客，其群体运输效率远超法拉利屡次独自接送的效率。由此可见， CPU 更适于解决延续性的繁多义务，而 GPU 则在并行解决少量义务时展现出色的效劳。

Image by the author with the assistance of AI (​ ​​ ​)

为了成功更出色的并行计算才干，GPU 在设计上偏差于将更多晶体管资源（transistors）投入到数据解决中，而非数据缓存（data caching）和流控机制（flow contro），这与 CPU 的设计思绪天壤之别。CPU 为了优化繁多线程的口头效率和复杂指令集的解决，特地划拨了少量的晶体管来增强这些方面的性能。

下图活泼地描述了 CPU 与 GPU 在芯片资源调配上的清楚差异。

Image by the author with inspiration from CUDA C++ Programming Guide

CPU 装备了高性能内核（powerful cores）与更为精妙的缓存内存架构（cache memory architecture）（消耗了相当多的晶体管资源），这种设计打算能够极大地优化顺序义务的口头速度。而图形解决器（GPU）则着重于内核（cores）数量，以成功更高的并行解决才干。

如今曾经引见完这些基础常识，那么在实践运行中，咱们应如何有效应用并行计算的长处呢？

02 Introduction to CUDA

当咱们着手构建深度学习模型时，很或者会偏差于驳回诸如 PyTorch 或 TensorFlow 这类广受欢迎的 Python 开发库。虽然如此，一个不争的理想是，这些库的外围代码都是 C/C++ 代码。另外，正如咱们先前所提及的，应用 GPU 放慢数据的解决速度往往是一种干流优化打算。此时，CUDA 的关键作用便凸显进去！CUDA 是一致计算设施架构（Compute Unified Device Architecture）的缩写，是英伟达（NVIDIA）为使 GPU 能够在通用计算畛域大放光荣而精心打造的平台。与DirectX被游戏引擎用于图形运算（graphical computation）不同，CUDA 使开发人员能够将英伟达（NVIDIA）的 GPU 计算才干集成到通用软件中，而不只仅局限于图形渲染。

为了成功这一指标，CUDA 推出了一款基于 C/C++ 的繁难接口（CUDA C/C++），协助开发者调用 GPU 虚构指令集（virtual intruction se）及口头特定操作（specific operations）（如在 CPU 与 GPU 间传输数据）。

在继续深化技术细节之前，咱们有必要廓清几个 CUDA 编程的基础概念和专业术语：

因此，在一份经常使用 CUDA 撰写的基本代码（basic code）中，程序主体在 host (CPU) 上口头，随后将数据传递给 device (GPU) ，并调用 kernels (functions) 在 device (GPU) 上并行运转。这些 kernels 由多条线程同时口头。运算成功后，结果再从 device (GPU) 回传至 host (CPU) 。

话说回来，让咱们再次聚焦于两组向量相加这个详细疑问：

借助 CUDA C/C++，编程人员能够创立一种被称为 kernels 的 C/C++ 函数；一旦这些 kernels 被调用， N 个不同的 CUDA 线程会并行口头 N 次。

若想定义这类 kernel，可运用 ​ ​__global__​ ​​ 关键字作为申明限定符（declaration specifier），而若欲设定口头该 kernel 的详细 CUDA 线程数目，则需驳回 ​​ 来成功：

每个 CUDA 线程在口头 kernel 时，都会被赋予一个唯一无二的线程 ID，即 threadIdx，它可以经过 kernel 中的预设变量失掉。上述示例代码将两个长度（size）均为 N 的向量 A 和 B 相加，并将结果保留到向量 C 中。值得咱们留意的是，相较于循环逐次解决成对加法的传统串行方式，CUDA 的长处在于其能够并行应用 N 个线程，一次性性成功所有加法运算。

不过，在运转上述这段代码前，咱们还需对其启动一次性修正。 切记，kernel 函数的运转环境是 device (GPU) ，这象征着一切相关数据均须驻留于 device 的内存之中。 要到达这一要求，可以借助 CUDA 提供的以下内置函数：

间接将变量 A、B 和 C 传入 kernel 的做法并不实用于本状况，咱们应当经常使用指针。在 CUDA 编程环境下，host 数组（比如示例中的 A、B 和 C）无法间接用于 kernel 启动（<<<...>>>）。鉴于 CUDA kernels 的上班空间为 device 的内存（device memory），故需向 kernel 提供 device 指针（device pointers）（d_A、d_B 和 d_C），以确保其能在 device 的内存上运转。

除此之外，咱们还需经过调用 cudaMalloc 函数在 device 上划分外存空间，并运用 cudaMemcpy 成功 host 和 device 之间的数据传输。

至此，咱们可在代码中成功向量 A 和 B 的初始化，并在程序开头处清算 CUDA 内存（cuda memory）。

另外，调用 kernel 后，务必拔出 ​ ​cudaDeviceSynchronize();​ ​ 这一行代码。该函数的作用在于协调 host 线程与 device 间的同步，确保 host 线程在继续口头前，device 已成功一切先前提交的 CUDA 操作。

此外， CUDA 的失误检测机制雷同无法或缺，这种检测机制能协助咱们及时发现并修正 GPU 上潜在的程序缺点（bugs）。倘若疏忽此环节，device 线程（CPU）将继续运转，而 CUDA 相关的缺点排查则将变得意外辣手，很难识别与 CUDA 相关的失误。

上方是这两种技术的详细成功方式：

要编译和运转 CUDA 代码，首先要求确保系统中已装有 CUDA 工具包（CUDA toolkit）。紧接着，经常使用 nvcc —— NVIDIA CUDA 编译器成功相关代码编译上班。

但是，的代码尚存优化空间。在前述示例中，咱们解决的向量规模仅为 N = 1000，这一数值偏小，难以充沛展现 GPU 弱小的并行解决才干。特地是在深度学习场景下，咱们时常要应答含有数以百万计参数的巨型向量。但是，倘若尝试将 N 的数值设为 500000，并驳回 <<<1, 500000>>> 的方式运转 kernel ，上述代码便会抛出失误。因此，为了完善代码，使之能顺利口头此类大规模运算，咱们亟需掌握 CUDA 编程中的外围思念 —— 线程层级结构（Thread hierarchy）。

03 Thread hierarchy（线程层级结构）

调用 kernel 函数时，驳回的是 <<<number_of_blocks, threads_per_block>>> 这种格局（notation）。因此，在上述示例中，咱们是以单个线程块的方式，启动了 N 个 CUDA 线程。但是，每个线程块所能容纳的线程数量都有限度，这是由于一切处于同一线程块内的线程，都被要求共存于同一流式多解决器外围（streaming multiprocessor core），并共同经常使用该外围的内存资源。

欲查问这一限度数量的详细数值，可经过以下代码成功：

就作者经常使用的 GPU 而言，其繁多线程块最多能承载 1024 个线程。因此，为了有效解决示例中提及的巨型向量（massive vector），咱们必定部署更多线程块，以成功更大规模的线程并发口头。 同时，这些线程块被精心规划成网格状结构（grids），如下图所展现：

​ ​​ ​ (CC BY-SA 3.0)

如今，咱们可以经过以下路径失掉线程 ID：

于是，该代码脚本更新为：

04 性能对比剖析

下表展现了在解决不同大小向量的加法运算时，CPU 与 GPU 的计算性能对比状况。

Image by the author

显而易见，GPU 的解决效劳长处，唯有在解决大规模向量时方能得以凸显。此外，切勿漠视一件事， 此处的期间对比仅仅考量了 kernel/function 的口头耗时，而未将 host 和 device 间数据传输所需的期间归入思考范围 。虽然在大少数状况下，数据传输的期间开支微无余道，但就咱们目前仅口头繁难加法运算（simple addition operation）的情景而言，这局部期间消耗却显得相对可观。因此，咱们应当铭刻， GPU 的计算性能，仅在面对那些既高度依赖计算才干又适宜大规模并行解决的义务时，才干得以酣畅淋漓地展现。

05 多维线程解决（Multidimensional threads）

如今，咱们曾经知道如何优化繁难数组操作（simple array operation）的性能了。但是，在解决深度学习模型时，必定要解决矩阵和张量运算（matrix and tensor operations）。在前文的示例中，咱们仅经常使用了内含 N 个线程的一维线程块（one-dimensional blocks）。但是，口头多维线程块（multidimensional thread blocks）（最高允许三维）雷同也是齐全可行的。因此，为了繁难起见，当咱们要求解决矩阵运算时，可运转一个由 N x M 个线程组成的线程块。还可以经过 row = threadIdx.x 来确定矩阵的行索引，而 col = threadIdx.y 则可用来失掉列索引。此外，为了简化操作，还可以经常使用 dim3 变量类型定义 number_of_blocks 和 threads_per_block。

下文的示例代码展现了如何成功两个矩阵的相加运算。

此外，咱们还可以将此示例进一步拓展，成功对多个线程块的解决：

此外，咱们也可以用雷同的思绪将这个示例裁减到三维运算（3-dimensional operations）操作的解决。

上文曾经引见了解决多维数据（multidimensional>​__device__​​​ 申明限定符（declaration specifier）来成功。这种限定符定义了可由 device （GPU）间接调用的函数（functions）。因此，这些函数仅能在 ​ ​__global__​ ​​ 或其余 ​ ​__device__​ ​ 函数中被调用。上方这个示例展现了如何对一个向量启动 sigmoid 运算（这是深度学习模型中极端经常出现的一种运算方式）。

至此，咱们曾经掌握了 CUDA 编程的外围概念，如今可以着手构建 CUDA kernels 了。关于深度学习模型而言，其实质就是一系列触及矩阵（matrix）与张量（tensor）的运算操作，包括但不限于求和（sum）、乘法（multiplication）、卷积（convolution）以及归一化（normalization ）等。举个例子，一个基础的矩阵乘法算法，可以经过以下方式成功并行化：

咱们可以留意到，在 GPU 版本的矩阵乘法算法中，循环次数清楚缩小，从而清楚优化了运算解决速度。上方这张图表直观地展现了 N x N 矩阵乘法在 CPU 与 GPU 上的性能对比状况：

Image by the author

咱们会发现， 随着矩阵大小（matrix size）的增大，GPU 在解决矩阵乘法运算时的性能优化幅度更大。

接上去，让咱们聚焦于一个基础的神经网络模型，其外围运算通常体现为 y = σ(Wx + b)，如下图所示：

Image by the author

上述运算关键触及矩阵乘法（matrix multiplication）、矩阵加法（matrix addition）以及对数组施加函数变换（applying a function to an array）。如若你已掌握这些并行化解决技术，象征着你如今齐全具有了从零构建、并在 GPU 上构建神经网络的才干！

06 Conclusion

本文咱们讨论了经过 GPU processing（译者注：经常使用 GPU启动数据解决和计算。）优化深度学习模型效劳的入门概念。不过，有一点还要求指出， 本文所引见的内容仅仅是皮毛，面前还暗藏着很多很多更深档次的物品。PyTorch 和 Tensorflow 等框架成功了诸多初级性能优化技术，涵盖了 optimized memory access、batched operations 等复杂概念（其底层应用了基于 CUDA 的 cuBLAS 和 cuDNN 等库）。 但愿这篇文章能够让各位读者对经常使用 .to("cuda") 方法，在 GPU 上构建、运转深度学习模型时的底层原理，有个初步的了解。

Thanks so much for reading!