图片由作者使用AI（https://copilot.microsoft.com/images/create）协助创建
如今，当我们谈到深度学习时，通常会将其实现与利用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）和变压器主要通过数学运算如矩阵加法、矩阵乘法和对矩阵应用函数来构建。因此，如果我们找到优化这些运算的方法，就可以提高深度学习模型的性能。
所以让我们从简单开始。想象你想将两个向量 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];
}
}正如你可以注意到的，计算机必须遍历向量，在每次迭代中顺序地将每对元素相加。但这些操作是相互独立的。第 ith 对元素的加法并不依赖于任何其他对元素。那么，如果我们能够并发地执行这些操作，同时并行地将所有元素对相加会怎么样呢？
一个直接的方法是使用CPU多线程来并行运行所有的计算。然而，当涉及到深度学习模型时，我们处理的是包含数百万个元素的巨大向量。普通的CPU一次只能处理大约十几个线程。这时，GPU就发挥作用了！现代的GPU可以同时运行数百万个线程，从而增强对这些大规模向量的数学运算性能。
GPU 与 CPU 对比尽管对于单个操作来说，CPU 的计算可能比 GPU 更快，但 GPU 的优势在于其并行化能力。原因在于它们的设计目标不同。CPU 被设计为尽可能快速地执行一系列操作（线程）（并且一次只能同时执行几十个线程），而 GPU 被设计为并行执行数百万个线程（同时牺牲了单个线程的速度）。
参见下方视频：
为了说明这一点，想象一下CPU就像一辆法拉利，而GPU就像一辆公交车。如果你的任务是移动一个人，法拉利（CPU）是更好的选择。然而，如果你要移动几个人，即使法拉利（CPU）每次行程更快，公交车（GPU）可以一次性运送所有人，比法拉利多次往返运送所有人更快。因此，CPU更适合处理顺序操作，而GPU则更适合处理并行操作。
图片由作者使用AI（https://copilot.microsoft.com/images/create）协助创建
为了提供更高的并行能力，GPU设计分配了更多的晶体管用于数据处理，而不是数据缓存和流控制，这与CPU不同，CPU为了优化单线程性能和复杂指令的执行，会分配大量的晶体管用于数据缓存和流控制。
下图展示了CPU与GPU之间的芯片资源分配情况。
图片由作者创作，灵感来源于CUDA C++编程指南
CPU 具有强大的核心和更复杂的缓存内存架构（为此分配了大量的晶体管）。这种设计使得处理顺序操作的速度更快。另一方面，GPU 优先考虑拥有大量的核心以实现更高的并行性。
现在我们理解了这些基本概念，如何在实践中利用这些并行计算能力呢？
CUDA 介绍当你运行一些深度学习模型时，你可能会选择使用一些流行的Python库，如PyTorch或TensorFlow。然而，众所周知，这些库的核心是用C/C++代码实现的。此外，正如我们之前提到的，你可能会使用GPU来加速处理。这时CUDA就派上用场了！CUDA代表 计算统一架构，它是NVIDIA为在其GPU上进行通用计算而开发的一个平台。因此，虽然DirectX被游戏引擎用来处理图形计算，CUDA使开发人员能够将NVIDIA的GPU计算能力集成到他们的通用软件应用程序中，而不仅仅局限于图形渲染。
为了实现这一点，CUDA 提供了一个简单的 C/C++ 基础接口（CUDA C/C++），该接口允许访问 GPU 的虚拟指令集和特定操作（例如在 CPU 和 GPU 之间移动数据）。
在我们继续之前，让我们先了解一些基本的CUDA编程概念和术语：
- host : 指的是CPU及其内存；
- device : 指的是GPU及其内存；
- kernel : 指的是在设备（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>
// 内核定义
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
...
// 使用 N 个线程调用内核
AddTwoVectors<<<1, N>>>(A, B, C);
...
}每个线程执行内核并被赋予一个唯一的线程ID threadIdx，该ID可以通过内建变量在内核中访问。上述代码将两个大小为N的向量A和B相加，并将结果存储到向量C中。如你所见，我们没有使用循环来顺序执行每对元素的加法操作，而是使用CUDA并行使用N个线程同时执行所有这些操作。
但在我们运行这段代码之前，我们需要进行另一项修改。重要的是要记住，kernel 函数在设备（GPU）中运行。因此，它所有的数据都需要存储在设备内存中。你可以通过使用以下 CUDA 内置函数来实现这一点：
#include <stdio.h>
// Kernel 定义
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // 向量的大小
float A[N], B[N], C[N]; // 向量 A、B 和 C 的数组
...
float *d_A, *d_B, *d_C; // 向量 A、B 和 C 的设备指针
// 为向量 A、B 和 C 在设备上分配内存
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// 将向量 A 和 B 从主机复制到设备
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// 使用 N 个线程调用 Kernel
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// 将向量 C 从设备复制到主机
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
}而不是直接将变量 A、B 和 C 传递给 kernel，我们需要使用指针。在 CUDA 编程中，你不能直接在内核启动 (<<<...>>>) 中使用 host 数组（如示例中的 A、B 和 C）。CUDA 内核操作设备内存，因此你需要将设备指针（如 d_A、d_B 和 d_C）传递给内核以便其进行操作。
除此之外，我们需要使用 cudaMalloc 在 设备 上分配内存，并使用 cudaMemcpy 在 主机 和 设备 之间复制数据。
现在我们可以添加向量A和B的初始化，并在代码的末尾刷新cuda内存。
#include <stdio.h>
// Kernel 定义
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // 向量的大小
float A[N], B[N], C[N]; // 向量A、B和C的数组
// 初始化向量A和B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // 向量A、B和C的设备指针
// 为向量A、B和C在设备上分配内存
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// 将向量A和B从主机复制到设备
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// 使用N个线程调用Kernel
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// 将向量C从设备复制到主机
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// 释放设备内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}此外，我们还需要在调用内核之后添加 cudaDeviceSynchronize();。这是一个用于将主机线程与设备同步的函数。当调用此函数时，主机线程将等待所有先前在设备上发出的CUDA命令完成后再继续执行。
除此之外，还需要添加一些CUDA错误检查，以便我们可以识别GPU上的错误。如果我们不添加这些检查，代码将继续执行_主机_线程（CPU），这将使得难以识别与CUDA相关的错误。
以下两种技术的实现：
#include <stdio.h>
// 内核定义
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // 向量的大小
float A[N], B[N], C[N]; // 用于向量 A、B 和 C 的数组
// 初始化向量 A 和 B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // 向量 A、B 和 C 的设备指针
// 为向量 A、B 和 C 在设备上分配内存
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// 将向量 A 和 B 从主机复制到设备
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// 使用 N 个线程调用内核
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// 检查错误
cudaError_t error = cudaGetLastError();
if(error != cudaSuccess) {
printf("CUDA 错误: %s\n", cudaGetErrorString(error));
exit(-1);
}
// 等待所有 CUDA 线程执行完毕
cudaDeviceSynchronize();
// 将向量 C 从设备复制到主机
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// 释放设备内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}为了编译和运行CUDA代码，你需要确保CUDA工具包已安装在你的系统上。然后，你可以使用nvcc（NVIDIA CUDA编译器）来编译代码。如果你的机器上没有GPU，你可以使用Google Colab。你需要在Runtime → Notebook设置中选择一个GPU，然后将代码保存到一个example.cu文件中并运行：
%%shell
nvcc example.cu -o compiled_example # 编译
./compiled_example # 运行
# 你也可以使用内存检查工具运行代码
compute-sanitizer --tool memcheck ./compiled_example 然而，我们的代码还没有完全优化。上面的例子使用了一个大小为 N = 1000 的向量。但是，这是一个较小的数字，无法充分展示 GPU 的并行化能力。此外，在处理深度学习问题时，我们通常会处理包含数百万参数的巨大向量。然而，如果我们尝试设置，例如 N = 500000，并使用上面的例子以 <<<1, 500000>>> 运行内核，将会抛出错误。因此，为了改进代码并执行此类操作，我们首先需要理解 CUDA 编程中的一个重要概念：线程层次结构。
内核函数的调用使用符号 <<<number_of_blocks, threads_per_block>>>。因此，在我们上面的例子中，我们运行了1个块，包含N个CUDA线程。然而，每个块对其支持的线程数量有限制。这是因为块内的每个线程都必须位于同一个流多处理器核心上，并且必须共享该核心的内存资源。
您可以使用以下代码片段获取此限制：
int 设备;
cudaDeviceProp 属性;
cudaGetDevice(&设备);
cudaGetDeviceProperties(&属性, 设备);
printf("每个块的最大线程数: %d\n", 属性.maxThreadsPerBlock);在当前的 Colab GPU 上，一个线程块可能包含多达 1024 个线程。因此，我们需要更多的块来执行更多的线程，以便处理示例中的大规模向量。此外，块被组织成网格，如下面所示：
https://handwiki.org/wiki/index.php?curid=1157670 (CC BY-SA 3.0)
现在，可以通过以下方式访问线程ID：
int i = blockIdx.x * blockDim.x + threadIdx.x;所以，我们的脚本变成了：
#include <stdio.h>
// 内核定义
__global__ void AddTwoVectors(float A[], float B[], float C[], int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) // 为了避免超出数组范围
C[i] = A[i] + B[i];
}
int main() {
int N = 500000; // 向量的大小
int threads_per_block;
int device;
cudaDeviceProp props;
cudaGetDevice(&device);
cudaGetDeviceProperties(&props, device);
threads_per_block = props.maxThreadsPerBlock;
printf("每个块的最大线程数: %d\n", threads_per_block); // 1024
float A[N], B[N], C[N]; // 用于向量A、B和C的数组
// 初始化向量A和B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // 向量A、B和C的设备指针
// 为向量A、B和C在设备上分配内存
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// 将向量A和B从主机复制到设备
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// 使用多个块和每个块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);
// 检查错误
cudaError_t error = cudaGetLastError();
if (error != cudaSuccess) {
printf("CUDA错误: %s\n", cudaGetErrorString(error));
exit(-1);
}
// 等待所有CUDA线程执行完毕
cudaDeviceSynchronize();
// 将向量C从设备复制到主机
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// 释放设备内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}下面比较了对于不同向量大小，CPU和GPU执行两个向量加法操作的计算情况。
图片由作者提供
如您所见，GPU 处理的优势仅在较大的向量大小 N 时才变得明显。另外，请记住，这次比较的时间仅考虑内核/函数的执行时间。没有考虑在 host 和 device 之间复制数据所需的时间，尽管在大多数情况下这可能并不显著，但在我们的情况下，由于我们只进行简单的加法操作，这相对是比较重要的。因此，重要的是要记住，GPU 计算仅在处理高度计算密集型且高度并行化的计算时才表现出其优势。
多维线程好的，现在我们知道如何提高简单数组操作的性能。但在处理深度学习模型时，我们需要处理矩阵和张量操作。在我们之前的例子中，我们只使用了一维的N个线程的块。然而，也可以执行多维线程块（最多3维）。因此，为了方便起见，如果你需要运行矩阵操作，你可以运行NxM个线程的线程块。在这种情况下，你可以通过 row = threadIdx.x, col = threadIdx.y 获取矩阵的行和列索引。此外，为了方便起见，你可以使用 dim3 变量类型来定义 number_of_blocks 和 threads_per_block。
下面的例子说明了如何添加两个矩阵。
#include <stdio.h>
// 核函数定义
__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() {
...
// 使用1个NxN线程块调用核函数
dim3 threads_per_block(N, N);
AddTwoMatrices<<<1, threads_per_block>>>(A, B, C);
...
}您也可以扩展此示例以处理多个块：
#include <stdio.h>
// 内核定义
__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() {
...
// 使用NxN线程的1个块调用内核
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 函数
__device__ float sigmoid(float x) {
return 1 / (1 + expf(-x));
}
// 对向量应用 Sigmoid 函数的内核定义
__global__ void sigmoidActivation(float input[], float output[]) {
int i = threadIdx.x;
output[i] = sigmoid(input[i]);
}所以，现在你已经知道了CUDA编程的基本重要概念，你可以开始创建CUDA内核了。在深度学习模型的情况下，它们基本上是一系列矩阵和张量操作，如求和、乘法、卷积、归一化等。例如，一个简单的矩阵乘法算法可以并行化如下：
// GPU 版本
__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版本
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上的神经网络！
结论在这篇文章中，我们介绍了关于GPU处理以增强深度学习模型性能的一些基础概念。然而，也需要注意的是，你所看到的概念只是基础知识，还有很多需要学习的内容。像PyTorch和Tensorflow这样的库实现了包括优化内存访问、批处理操作等更复杂概念在内的优化技术（它们利用了基于CUDA构建的库，如cuBLAS和cuDNN）。不过，希望这篇文章能帮助你理解当你编写.to("cuda")并执行GPU上的深度学习模型时背后发生了什么。
在未来的文章中，我将尝试介绍更多关于CUDA编程的复杂概念。请在评论中告诉我你的想法，或者告诉我你希望我在下一篇文章中写些什么！非常感谢你的阅读！😊
进一步阅读CUDA 编程指南 — NVIDIA CUDA 编程文档。
CUDA 文档 — NVIDIA 完整的 CUDA 文档。
CUDA神经网络训练实现 — 纯CUDA C++实现的神经网络训练。
CUDA LLM 训练实现— 使用纯 CUDA C 实现的 LLM 训练。
共同學習，寫下你的評論
評論加載中...
作者其他優質文章
