为什么大模型在GPU上运行更快？
引言
当前，提到深度学习，我们很自然地会想到利用GPU来提升运算效率。GPU最初是为了加速图像渲染和2D、3D图形处理而设计的。但它们强大的并行处理能力，使得它们在深度学习等更广泛的领域中也发挥了重要作用。
深度学习模型开始采用GPU是在2000年代中期到晚期，到了2012年，随着AlexNet的诞生，这种使用变得极为普遍。AlexNet是由Alex Krizhevsky、Ilya Sutskever和Geoffrey Hinton共同设计的卷积神经网络，它在2012年的ImageNet大规模视觉识别挑战赛（ILSVRC）中获胜。这一胜利不仅证明了深度神经网络在图像分类上的巨大潜力，也展示了使用GPU进行大型模型训练的优势。
自那以后，使用GPU进行深度学习模型训练变得日益流行，这也催生了PyTorch和TensorFlow等框架的诞生。如今，我们只需在PyTorch中简单地写上.to("cuda")，即可将数据传输至GPU，期待训练过程能够更快。但深度学习算法是如何在实际中利用GPU的计算能力的呢？让我们一探究竟。
深度学习架构，如神经网络、卷积神经网络（CNNs）、循环神经网络（RNNs）和变换器（transformers），本质上是通过矩阵加法、矩阵乘法以及对矩阵应用函数等数学运算构建的。如果我们能够优化这些运算，就能提升深度学习模型的效率。
让我们从基础开始。设想你需要将两个向量A和B相加得到向量C，即C = A + B。
在 C 中的一个简单实现是：
`void AddTwoVectors(flaot A[], float B[], float C[]) {` `for (int i = 0; i < N; i++) {` `C[i] = A[i] + B[i];` `}``}`
你可能会注意到，计算机需要逐个遍历向量中的元素，每次迭代都依次将一对元素相加。这些加法操作是独立进行的，即对第i个元素对的加法并不依赖于其他任何元素对。那么，如果我们能够同时进行这些操作，一次性并行地完成所有元素对的加法，又会如何呢？
一种简单的解决方案是利用CPU的多线程功能，来并行处理所有的计算任务。但是，在处理深度学习模型时，我们面对的是包含数百万元素的大型向量。一般CPU能够同时处理的线程数量大约只有十几个。
这时，GPU的优势就显现出来了！现代GPU能够同时执行数百万的线程，极大地提升了对这些庞大向量进行数学运算的效率。
GPU 与 CPU 比较
虽然CPU在单个操作的速度上可能超过GPU，但GPU的真正优势在于其强大的并行处理功能。这背后的原因在于两者设计初衷的差异。CPU的设计宗旨是尽可能快速地完成一系列操作序列，它能够同时处理的线程数量有限，大约只有几十个；相比之下，GPU的设计宗旨是为了能够同时执行数百万条线程，即便这意味着牺牲了单个线程的执行速度。
举个例子，我们可以把CPU比作一辆法拉利跑车，而GPU则相当于一辆大巴。如果你只需要运送一个人，那么法拉利（CPU）无疑是更佳的选择。但如果你的任务是运送一群人，尽管法拉利（CPU）每次运送的速度更快，但大巴（GPU）却能够一次性将所有人送达，这样一次性完成运输的速度，要比法拉利多次往返运送要快得多。所以，CPU更适合执行顺序串行操作，而GPU则更擅长处理并行操作。
为了实现更强的并行处理功能，GPU在设计时将更多的晶体管资源用于执行数据处理任务，而不是像CPU那样，将大量晶体管用于数据缓存和流程控制，这样做是为了提升单线程的处理速度和复杂指令的执行效率。
下面的图表展示了CPU和GPU在芯片资源分配上的差异。
CPU配备了功能强大的核心和更为复杂的缓存内存结构（为此投入了大量的晶体管资源）。这样的设计让CPU在处理顺序任务时更为迅速。而GPU则侧重于拥有众多核心，以此来达到更高的并行处理水平。
既然我们已经掌握了这些基础概念，那么在实际应用中，我们该如何发挥这些并行计算的优势呢？
CUDA简介
当您启动某个深度学习模型时，您可能会倾向于选择像PyTorch或TensorFlow这样的流行Python库。但这些库的底层实际上是在运行C/C++代码，这是众所周知的事实。此外，正如我们之前所讨论的，您可能会利用GPU来提升处理速度。这就引入了CUDA的概念！CUDA，即Compute Unified Architecture，是NVIDIA为其GPU开发的一个平台，用于执行通用计算任务。因此，DirectX被游戏引擎用于图形计算，而CUDA则允许开发者将NVIDIA的GPU计算能力整合到他们的应用程序中，不仅限于图形渲染。
为了实现这一点，CUDA提供了一个基于C/C++的简洁接口（CUDA C/C++），它能够访问GPU的虚拟指令集和一些特定操作，比如在CPU和GPU之间传输数据。
在我们深入之前，先来理解一些基本的CUDA编程概念和术语：
-
host：指CPU及其内存；
-
device：指GPU及其内存；
-
kernel：指在设备（GPU）上执行的函数；
在用CUDA编写的简单代码中，程序在host（CPU）上运行，将数据发送至device（GPU），并启动kernel（函数）在device（GPU）上执行。这些kernel由多个线程并行执行。执行完毕后，结果会从device（GPU）传回host（CPU）。
现在，让我们回到添加两个向量的问题上：
`#include <stdio.h>`` ``void AddTwoVectors(flaot A[], float B[], float C[]) {` `for (int i = 0; i < N; i++) {` `C[i] = A[i] + B[i];` `}``}`` ``int main() {` `...` `AddTwoVectors(A, B, C);` `...``}`
在CUDA C/C++编程环境中，开发者能够创建被称为kernels的C/C++函数，这些函数一旦被触发，就能由N个不同的CUDA线程同时执行N次。
定义一个kernel时，我们用__global__关键字来声明，而执行这个kernel的CUDA线程数量可以通过特殊的<<<...>>>标记来设置：
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoVectors(float A[], float B[], float C[]) {` `int i = threadIdx.x;` `C[i] = A[i] + B[i];``}`` ``int main() {` `...` `// Kernel invocation with N threads` `AddTwoVectors<<<1, N>>>(A, B, C);` `...``}`
每个执行核心（thread）在运行核心函数（kernel）时，都会被分配一个独一无二的核心标识符 threadIdx，这个标识符可以在核心函数内部通过内建变量来获取。上述代码实现了两个大小为N的向量A和B的相加操作，并将相加结果存放到向量C中。你会注意到，与传统的顺序循环处理每一对元素相加的方式不同，CUDA技术允许我们通过并行使用N个核心来同时完成所有这些操作。
但在我们实际运行这段代码之前，还需要进行一些调整。需要牢记的是，核心函数是在设备（GPU）上执行的。这意味着它使用的所有数据都应当存储在GPU的内存中。我们可以通过调用CUDA提供的一系列内建函数来完成这一数据的迁移：
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoVectors(float A[], float B[], float C[]) {` `int i = threadIdx.x;` `C[i] = A[i] + B[i];``}`` ``int main() {`` ` `int N = 1000; // Size of the vectors` `float A[N], B[N], C[N]; // Arrays for vectors A, B, and C`` ` `...`` ` `float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C`` ` `// Allocate memory on the device for vectors A, B, and C` `cudaMalloc((void **)&d_A, N * sizeof(float));` `cudaMalloc((void **)&d_B, N * sizeof(float));` `cudaMalloc((void **)&d_C, N * sizeof(float));`` ` `// Copy vectors A and B from host to device` `cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);` `cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);`` ` `// Kernel invocation with N threads` `AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);`` ` `// Copy vector C from device to host` `cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);`` ``}`
我们不能将变量A、B和C直接传入核心函数，而应该使用指针。在CUDA编程中，你无法在核心函数调用（标记为<<<...>>>）中直接使用主机上的数组（比如示例中的A、B和C）。核心函数是在设备内存中运行的，因此你需要将设备指针（d_A、d_B和d_C）传入核心函数，以便它能够进行操作。
除此之外，我们还需要通过调用cudaMalloc函数在设备上分配内存，并利用cudaMemcpy函数在主机内存和设备内存之间传输数据。
现在，我们可以在代码的最后添加向量A和B的初始化步骤，并在结束时刷新CUDA内存。
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoVectors(float A[], float B[], float C[]) {` `int i = threadIdx.x;` `C[i] = A[i] + B[i];``}`` ``int main() {`` ` `int N = 1000; // Size of the vectors` `float A[N], B[N], C[N]; // Arrays for vectors A, B, and C`` ` `// Initialize vectors A and B` `for (int i = 0; i < N; ++i) {` `A[i] = 1;` `B[i] = 3;` `}`` ` `float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C`` ` `// Allocate memory on the device for vectors A, B, and C` `cudaMalloc((void **)&d_A, N * sizeof(float));` `cudaMalloc((void **)&d_B, N * sizeof(float));` `cudaMalloc((void **)&d_C, N * sizeof(float));`` ` `// Copy vectors A and B from host to device` `cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);` `cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);`` ` `// Kernel invocation with N threads` `AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);`` ` `// Copy vector C from device to host` `cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);`` ` `// Free device memory` `cudaFree(d_A);` `cudaFree(d_B);` `cudaFree(d_C);``}`
此外，我们在核心函数调用之后，需要加入 `cudaDeviceSynchronize();` 这个调用。这个函数的作用是确保主机线程与设备之间的同步。调用此函数后，主机线程会暂停，直到设备上所有先前发出的CUDA命令都执行完毕才会继续。
此外，重要的是要加入一些CUDA错误检查机制，以便我们能够发现GPU上的错误。如果我们忽略了这些检查，代码会持续执行主机线程（即CPU的线程），这将使得发现与CUDA相关的错误变得困难。
以下是这两种技术的实现方法：
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoVectors(float A[], float B[], float C[]) {` `int i = threadIdx.x;` `C[i] = A[i] + B[i];``}`` ``int main() {`` ` `int N = 1000; // Size of the vectors` `float A[N], B[N], C[N]; // Arrays for vectors A, B, and C`` ` `// Initialize vectors A and B` `for (int i = 0; i < N; ++i) {` `A[i] = 1;` `B[i] = 3;` `}`` ` `float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C`` ` `// Allocate memory on the device for vectors A, B, and C` `cudaMalloc((void **)&d_A, N * sizeof(float));` `cudaMalloc((void **)&d_B, N * sizeof(float));` `cudaMalloc((void **)&d_C, N * sizeof(float));`` ` `// Copy vectors A and B from host to device` `cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);` `cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);`` ` `// Kernel invocation with N threads` `AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);`` ` `// Check for error` `cudaError_t error = cudaGetLastError();` `if(error != cudaSuccess) {` `printf("CUDA error: %s\n", cudaGetErrorString(error));` `exit(-1);` `}`` ` `// Waits untill all CUDA threads are executed` `cudaDeviceSynchronize();`` ` `// Copy vector C from device to host` `cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);`` ` `// Free device memory` `cudaFree(d_A);` `cudaFree(d_B);` `cudaFree(d_C);``}`
为了编译和执行CUDA程序，首先得保证你的计算机上已经安装了CUDA工具集。接着，你可以利用NVIDIA的CUDA编译器nvcc来编译你的代码。如果你的计算机不具备GPU，你可以考虑使用Google Colab平台。你只需在“Runtime”菜单下的“Notebook settings”选项中选择相应的GPU，然后将你的代码保存为example.cu文件，并执行它。
`%%shell``nvcc example.cu -o compiled_example # compile``./compiled_example # run`` ``# you can also run the code with bug detection sanitizer``compute-sanitizer --tool memcheck ./compiled_example`
不过，我们的代码优化还有待提升。例如，上述示例中的向量大小仅为N = 1000。这个数值偏小，不足以完全体现GPU的并行处理优势。在深度学习问题中，我们经常要处理包含数百万参数的大型向量。如果我们尝试将N设置为500000，并像之前的例子那样以<<<1, 500000>>>的方式调用核心函数，会遇到错误。因此，为了优化代码并执行这样的操作，我们首先需要理解CUDA编程中的一个关键概念：线程的层级结构。
线程层次结构
核心函数的调用是通过<<<number_of_blocks, threads_per_block>>>这样的标记来完成的。比如，在我们之前的例子中，我们执行了1个包含N个CUDA线程的区块。但是，每个区块支持的线程数是有上限的。这是因为区块内的所有线程都需要位于同一个流式多处理器核心上，并且需要共享该核心的内存资源。
你可以通过以下代码片段来查询这个上限值：
`int device;``cudaDeviceProp props;``cudaGetDevice(&device);``cudaGetDeviceProperties(&props, device);``printf("Maximum threads per block: %d\n", props.maxThreadsPerBlock);`
在Colab平台的当前GPU配置中，单个线程块最多可以包含1024个线程。因此，为了在示例中处理大型向量，我们需要更多的线程块来执行更多的线程。同时，这些线程块被进一步组织成更大的结构——网格，就像下面展示的那样：
现在，可以使用以下方式访问线程 ID：
int i = blockIdx.x * blockDim.x + threadIdx.x;
所以，我们的脚本变成：
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoVectors(float A[], float B[], float C[], int N) {` `int i = blockIdx.x * blockDim.x + threadIdx.x;` `if (i < N) // To avoid exceeding array limit` `C[i] = A[i] + B[i];``}`` ``int main() {` `int N = 500000; // Size of the vectors` `int threads_per_block;` `int device;` `cudaDeviceProp props;` `cudaGetDevice(&device);` `cudaGetDeviceProperties(&props, device);` `threads_per_block = props.maxThreadsPerBlock;` `printf("Maximum threads per block: %d\n", threads_per_block); // 1024`` ` `float A[N], B[N], C[N]; // Arrays for vectors A, B, and C`` ` `// Initialize vectors A and B` `for (int i = 0; i < N; ++i) {` `A[i] = 1;` `B[i] = 3;` `}`` ` `float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C`` ` `// Allocate memory on the device for vectors A, B, and C` `cudaMalloc((void **)&d_A, N * sizeof(float));` `cudaMalloc((void **)&d_B, N * sizeof(float));` `cudaMalloc((void **)&d_C, N * sizeof(float));`` ` `// Copy vectors A and B from host to device` `cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);` `cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);`` ` `// Kernel invocation with multiple blocks and threads_per_block threads per block` `int number_of_blocks = (N + threads_per_block - 1) / threads_per_block;` `AddTwoVectors<<<number_of_blocks, threads_per_block>>>(d_A, d_B, d_C, N);`` ` `// Check for error` `cudaError_t error = cudaGetLastError();` `if (error != cudaSuccess) {` `printf("CUDA error: %s\n", cudaGetErrorString(error));` `exit(-1);` `}`` ` `// Wait until all CUDA threads are executed` `cudaDeviceSynchronize();`` ` `// Copy vector C from device to host` `cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);`` ` `// Free device memory` `cudaFree(d_A);` `cudaFree(d_B);` `cudaFree(d_C);`` ``}`
性能对比
下面对不同向量大小的两个向量相加运算的 CPU 和 GPU 计算进行了比较。
显而易见，GPU处理的性能优势在处理大规模向量N时才会明显体现出来。此外，需要记住的是，这里的时间比较仅针对核心函数的执行时间，并未包括在主机和设备间传输数据所需的时间。虽然在大多数情况下，数据传输时间可能并不显著，但在我们只进行简单加法操作的情况下，这部分时间却相对较长。因此，我们必须意识到，GPU在处理那些既计算密集又高度可并行化的计算任务时，才能真正发挥其性能优势。
多维线程
明白了，我们现在掌握了如何提升基本数组操作效率的方法。但在深度学习模型的实践中，我们更多地需要处理矩阵和张量的操作。回顾我们之前的示例，我们仅使用了一维区块，每个区块包含N个线程。实际上，我们可以执行更高维度的区块（最多可至三维）。因此，如果你需要进行矩阵运算，可以方便地设置一个NxM的线程区块。在这种情况下，可以通过row = threadIdx.x和col = threadIdx.y来获取矩阵的行和列索引。此外，为了简化操作，可以使用dim3数据类型来指定区块的数量和每个区块中的线程数。
以下示例展示了如何实现两个矩阵的相加操作。
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoMatrices(float A[N][N], float B[N][N], float C[N][N]) {` `int i = threadIdx.x;` `int j = threadIdx.y;` `C[i][j] = A[i][j] + B[i][j];``}`` ``int main() {` `...` `// Kernel invocation with 1 block of NxN threads` `dim3 threads_per_block(N, N);` `AddTwoMatrices<<<1, threads_per_block>>>(A, B, C);` `...``}`
您还可以扩展此示例以处理多个块：
`#include <stdio.h>`` ``// Kernel definition``__global__ void AddTwoMatrices(float A[N][N],float B[N][N],float C[N][N]) {` `int i = blockIdx.x * blockDim.x + threadIdx.x;` `int j = blockIdx.y * blockDim.y + threadIdx.y;` `if (i < N && j < N) {` `C[i][j] = A[i][j] + B[i][j];` `}``}`` ``int main() {` `...` `// Kernel invocation with 1 block of NxN threads` `dim3 threads_per_block(32, 32);` `dim3 number_of_blocks((N + threads_per_block.x - 1) ∕ threads_per_block.x, (N + threads_per_block.y - 1) ∕ threads_per_block.y);` `AddTwoMatrices<<<number_of_blocks, threads_per_block>>>(A, B, C);` `...``}`
您可以按照这个示例的思路，进一步扩展到处理三维数据的操作。
既然您已经掌握了多维数据的操作方式，接下来要学习另一个既重要又简单的概念：在核心函数内部如何调用函数。这通常是通过__device__关键字来实现的。使用__device__关键字定义的函数可以直接在设备（即GPU）上调用。这意味着，这些函数只能在__global__核心函数或其他__device__函数中被调用。以下示例展示了如何在向量上应用sigmoid函数——这是深度学习模型中非常普遍的一种操作。
`#include <math.h>`` ``// Sigmoid function``__device__ float sigmoid(float x) {` `return 1 / (1 + expf(-x));``}`` ``// Kernel definition for applying sigmoid function to a vector``__global__ void sigmoidActivation(float input[], float output[]) {` `int i = threadIdx.x;` `output[i] = sigmoid(input[i]);``}`
明白了CUDA编程的基础关键概念后，您就可以着手编写CUDA核心函数了。对于深度学习模型，它们通常包含一系列矩阵和张量操作，比如求和、乘法、卷积、归一化等操作。以矩阵乘法为例，一个简单的算法可以通过以下方式实现并行处理：
`// GPU version`` ``__global__ void matMul(float A[M][N], float B[N][P], float C[M][P]) {` `int row = blockIdx.x * blockDim.x + threadIdx.x;` `int col = blockIdx.y * blockDim.y + threadIdx.y;`` ` `if (row < M && col < P) {` `float C_value = 0;` `for (int i = 0; i < N; i++) {` `C_value += A[row][i] * B[i][col];` `}` `C[row][col] = C_value;` `}``}`
现在将其与下面两个矩阵乘法的普通 CPU 实现进行比较：
`// CPU version`` ``void matMul(float A[M][N], float B[N][P], float C[M][P]) {` `for (int row = 0; row < M; row++) {` `for (int col = 0; col < P; col++) {` `float C_value = 0;` `for (int i = 0; i < N; i++) {` `C_value += A[row][i] * B[i][col];` `}` `C[row][col] = C_value;` `}` `}``}`
您可以注意到，在 GPU 版本上，我们的循环更少，从而可以更快地处理操作。下面是CPU和GPU在NxN矩阵乘法上的性能比较：
正如您所观察到的，随着矩阵大小的增加，矩阵乘法运算的 GPU 处理性能提升甚至更高。
现在，考虑一个基本的神经网络，它主要涉及 y = σ(Wx + b) 操作，如下所示：
这些操作主要包括矩阵乘法、矩阵加法以及将函数应用于数组，所有这些操作您都已经熟悉了并行化技术。因此，您现在能够从头开始实现在 GPU 上运行的您自己的神经网络！
总结
本文[1]我们探讨了提升深度学习模型性能的GPU处理基础知识。PyTorch和TensorFlow等库应用了包含优化内存访问、批量处理等更高级概念的优化技术（它们使用了在CUDA基础上构建的库，比如cuBLAS和cuDNN）。希望本文能够帮助你理解当你执行.to("cuda")并利用GPU运行深度学习模型时，背后所发生的机制。
Reference
[1]
更多每日开发小技巧
尽在****未闻 Code Telegram Channel !
END
未闻 Code·知识星球开放啦！
一对一答疑爬虫相关问题
职业生涯咨询
面试经验分享
每周直播分享
......
未闻 Code·知识星球期待与你相见~
一二线大厂在职员工
十多年码龄的编程老鸟
国内外高校在读学生
中小学刚刚入门的新人
在“未闻 Code技术交流群”等你来！
入群方式：添加微信“mekingname”，备注“粉丝群”（谢绝广告党，非诚勿扰！)
