创作不易，方便的话点点关注，谢谢
文章结尾有最新热度的文章，感兴趣的可以去看看。
本文是经过严格查阅相关权威文献和资料，形成的专业的可靠的内容。全文数据都有据可依，可回溯。特别申明：数据和资料已获得授权。本文内容，不涉及任何偏颇观点，用中立态度客观事实描述事情本身
文章有点长(6660字阅读时长：15分)，期望您能坚持看完，并有所收获。
如今，当我们谈到深度学习时，通常会将其实施与利用 GPU 来提高性能联系起来。图形处理器（GPU）最初设计用于加速图像、二维和三维图形的渲染。然而，由于 GPU 能够执行许多并行操作，其用途已扩展到深度学习等应用领域。
在深度学习模型中使用 GPU 大约始于 2000 年代中后期，2012 年 AlexNet 的出现使其变得非常流行。AlexNet 是由 Alex Krizhevsky、Ilya Sutskever 和 Geoffrey Hinton 设计的卷积神经网络，在 2012 年的 ImageNet 大规模视觉识别挑战赛（ILSVRC）中获胜。这一胜利具有里程碑式的意义，因为它证明了深度神经网络在图像分类方面的有效性，以及使用 GPU 训练大型模型的有效性。
在这一突破之后，将 GPU 用于深度学习模型变得越来越流行，这也促进了 PyTorch 和 TensorFlow 等框架的诞生。如今，我们只需在 PyTorch 中写入 .to("cuda")，就能将数据发送到 GPU，并期望加速训练。但在实践中，深度学习算法如何利用 GPU 的计算性能呢？让我们一探究竟！
神经网络、CNN、RNN 和变换器等深度学习架构基本上是通过矩阵加法、矩阵乘法和矩阵函数应用等数学运算构建的。因此，如果我们能找到优化这些运算的方法，就能提高深度学习模型的性能。
那么，让我们从简单的开始。想象一下，你想把两个向量 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 为此目的分配了很大一部分晶体管，以优化单线程性能和复杂指令的执行。
下图说明了 CPU 与 GPU 的芯片资源分配情况。
CPU 拥有强大的内核和更复杂的高速缓冲存储器架构（为此分配了大量的晶体管）。这种设计可以更快地处理顺序操作。另一方面，GPU 优先考虑拥有大量内核，以实现更高水平的并行性。
既然我们已经了解了这些基本概念，那么在实践中如何利用这种并行计算能力呢？
CUDA 简介
当你运行一些深度学习模型时，你可能会选择使用一些流行的 Python 库，如 PyTorch 或 TensorFlow。不过，众所周知，这些库的核心代码都是 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: refers to the CPU and its memory; 主机：指 CPU 及其内存； device: refers to the GPU and its memory; 设备：指 GPU 及其内存； kernel: refers to a function that is executed on the device (GPU); 内核：指在设备（GPU）上执行的功能；
因此，在使用 CUDA 编写的基本代码中，程序在主机（CPU）上运行，向设备（GPU）发送数据，并启动内核（函数）在设备（GPU）上执行。这些内核由多个线程并行执行。执行结束后，结果将从设备（GPU）传回主机（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++ 中，程序员可以定义称为内核的 C/C++ 函数，调用这些函数时，N 个不同的 CUDA 线程会并行执行 N 次。
要定义内核，可以使用 global 声明指定符，并使用 <<...>>> 符号指定执行该内核的 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); ...}
每个线程都会执行内核，并获得一个唯一的线程 ID threadIdx，可通过内置变量在内核中访问。上面的代码将大小为 N 的两个向量 A 和 B 相加，并将结果存储到向量 C 中。正如你所注意到的，CUDA 允许我们使用 N 个并行线程同时执行所有这些操作，而不是使用一个循环来顺序执行每个成对加法。
不过，在运行这段代码之前，我们还需要进行另一项修改。必须记住，内核函数是在设备（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）。CUDA 内核在设备内存上运行，因此需要将设备指针（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__ voidAddTwoVectors(float A[], float B[], float C[]) {int i = threadIdx.x; C[i] = A[i] + B[i];}intmain() {int N = 1000; // Size of the vectorsfloat A[N], B[N], C[N]; // Arrays for vectors A, B, and C// Initialize vectors A and Bfor (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。只需在运行时 → 笔记本设置中选择 GPU，然后将代码保存到 example.cu 文件中并运行即可：
%%shellnvcc example.cu -o compiled_example # compile./compiled_example # run# you can also run the code with bug detection sanitizercompute-sanitizer --tool memcheck ./compiled_example
不过，我们的代码仍未完全优化。上面的例子使用了一个大小为 N = 1000 的向量。但是，这只是一个很小的数字，并不能充分展示 GPU 的并行化能力。此外，在处理深度学习问题时，我们经常要处理数百万个参数的海量向量。但是，如果我们尝试设置 N = 500000，并使用上述示例中的<<<1, 500000>>运行内核，就会出现错误。因此，要改进代码并执行此类操作，我们首先需要了解 CUDA 编程的一个重要概念：线程层次结构。
线程结构
内核函数的调用使用 <<>> 符号完成。因此，在上面的例子中，我们用 N 个 CUDA 线程运行 1 个块。不过，每个区块可支持的线程数量都有限制。这是因为区块内的每个线程都必须位于同一个流式多处理器内核上，并且必须共享该内核的内存资源。 您可以使用下面的代码片段获得这一限制：
int device;cudaDeviceProp props;cudaGetDevice(&device);cudaGetDeviceProperties(&props, device);printf("Maximum threads per block: %d\n", props.maxThreadsPerBlock);
在当前的 Colab GPU 上，一个线程块最多可包含 1024 个线程。因此，我们需要更多的区块来执行更多的线程，以便处理示例中的大量向量。此外，如下图所示，线程块被组织成网格： 现在，可以使用
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 vectorsint 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); // 1024float A[N], B[N], C[N]; // Arrays for vectors A, B, and C// Initialize vectors A and Bfor (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 CcudaMalloc((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 devicecudaMemcpy(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 blockint 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 executedcudaDeviceSynchronize();// Copy vector C from device to hostcudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);}
性能比较
下面比较了 CPU 和 GPU 在不同向量大小的情况下对这一添加两个向量操作的计算结果。
此外，请记住，这种时间比较只考虑了内核/函数的执行。虽然在大多数情况下这一时间并不重要，但在我们的案例中却相对可观，因为我们只是执行了一个简单的加法操作。因此，重要的是要记住，只有在处理高度并行化的计算密集型计算时，GPU 计算才能显示出其优势。
多线程
好了，现在我们知道如何提高简单数组操作的性能了。但在处理深度学习模型时，我们需要处理矩阵和张量操作。在前面的例子中，我们只使用了 N 个线程的一维块。不过，也可以执行多维线程块（最多 3 维）。因此，为了方便起见，如果需要运行矩阵操作，可以运行一个由 NxM 个线程组成的线程块。在这种情况下，可以通过 row = threadIdx.x, col = threadIdx.y 来获取矩阵行列索引。此外，为了方便起见，还可以使用 dim3 变量类型来定义块数（number_of_block）和线程数（threads_per_block）。
下面的示例说明了如何添加两个矩阵。
#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 声明指定符来实现。这定义了设备（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 versionvoid 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 上运行！
结论
在这篇文章中，我们介绍了有关 GPU 处理的入门概念，以提高深度学习模型的性能。不过，还需要指出的是，您所看到的概念只是基础知识，要学的东西还有很多。PyTorch 和 Tensorflow 等库实现的优化技术涉及其他更复杂的概念，如优化内存访问、分批操作等（它们利用了基于 CUDA 构建的库，如 cuBLAS 和 cuDNN）。不过，我希望这篇文章能帮助你理清在 GPU 上编写 .to("cuda") 和执行深度学习模型时的幕后工作。
在今后的文章中，我将尝试介绍有关 CUDA 编程的更多复杂概念。请在评论中告诉我您的想法或您希望我接下来写些什么！感谢您的阅读。
点个“在看”不失联
最新热门文章推荐：
国外C++程序员分享：2024/2025年C++是否还值得学习？
国外Python程序员分享：如何用Python构建一个多代理AI应用
国外CUDA程序员分享：2024年GPU编程CUDA C++（从环境安装到进阶技巧）
国外Python程序员分享：2024年NumPy高性能计算库（高级技巧）
外国人眼中的卢湖川：从大连理工到全球舞台，他的科研成果震撼世界！
外国人眼中的张祥雨：交大90后男神博士，3年看1800篇论文，还入选福布斯精英榜
参考文献： 《图片来源网络》 《CUDA Programming Guide — NVIDIA CUDA Programming documentation.》 《CUDA Documentation — NVIDIA complete CUDA documentation.》 《CUDA Neural Network training implementation — Pure CUDA C++ implementation of a neural network training.》 《CUDA LLM training implementation — Training implementation of LLM with pure CUDA C.》
本文使用 文章同步助手 同步
