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 的打算性能呢？让我们来理解一下吧！
神经网络、CNN、RNN 和 transformer 等深度学习架构基本上是利用数学运算构建的，例如矩阵加法、矩阵乘法和运用函数矩阵。因此，如果我们找到一种方法来优化这些操作，我们可以提高深度学习模型的性能。
以是，让我们从大略开始。假设您想添加两个向量 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 可以同时运行数百万个线程，从而增强了这些数学运算在海量向量上的性能。
只管 CPU 打算在单个操作中可能比 GPU 更快，但 GPU 的上风在于其并行化功能。这样做的缘故原由是它们的设计目标不同。虽然 CPU 被设计为尽可能快地实行一系列操作（线程），并且只能同时实行数十个操作（线程），但 GPU 被设计为并行实行数百万个操作（同时捐躯单个线程的速率）。
举例来说，想象一下 CPU 就像法拉利，而 GPU 就像巴士。如果您的任务是移动一个人，法拉利 （CPU） 是更好的选择。但是，如果您要运送几个人，纵然法拉利 （CPU） 每次旅行速率更快，巴士 （GPU） 也可以一次性运送所有人，一次运送所有人的速率比法拉利多次行驶的速率更快。因此，CPU 更适宜处理顺序操作，而 GPU 则更适宜并行操作。
为了供应更高的并行功能，GPU 设计为数据处理分配了比数据缓存和流量掌握更多的晶体管，这与为此目的分配大量晶体管的 CPU 不同，以优化单线程性能和繁芜的指令实行。
下图解释了 CPU 与 GPU 的芯片资源分布。
Image by the author with inspiration from CUDA C++ Programming Guide
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>// 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，该 ID 可通过内置变量在内核内访问。上面的代码将两个大小为 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__ 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。您只须要在“运行时”→“条记本设置”中选择一个 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 编程的一个主要观点：线程层次构造。
线程层次构造内核函数的调用是利用符号 <<>> 完成的。因此，在上面的示例中，我们运行 1 个具有 N 个 CUDA 线程的块。但是，每个块可以支持的线程数都有限定。发生这种情形的缘故原由是，块中的每个线程都须要位于同一个流式多处理器内核上，并且必须共享该内核的内存资源。
您可以利用以下代码片段获取此限定：
int device;cudaDeviceProp props;cudaGetDevice(&device);cudaGetDeviceProperties(&props, device);printf("Maximum threads per block: %d\n", props.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>// 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 的打算，为不同的矢量大小添加了两个矢量运算。
Image by the author
正如人们所看到的，GPU 处理的上风只有在矢量大小为 N 的情形下才会显现出来。其余，请记住，此韶光比较仅考虑内核/函数的实行。它没有考虑到在主机和设备之间复制数据的韶光，只管在大多数情形下可能并不主要，但在我们的情形下，它相对较大，由于我们只实行一个大略的加法操作。因此，主要的是要记住，GPU 打算仅在处理高度并行化的高度打算密集型打算时才显示出其上风。
多维线程好了，现在我们知道如何提高大略数组操作的性能了。但是在处理深度学习模型时，我们须要处理矩阵和张量运算。在前面的示例中，我们只利用了具有 N 个线程的一维模块。但是，也可以实行多维线程块（最多 3 维）。因此，为了方便起见，如果须要运行矩阵操作，可以运行NxM线程的线程块。在本例中，您可以获取矩阵行列索引，如 row = threadIdx.x， col = threadIdx.y。此外，为方便起见，您可以利用 dim3 变量类型来定义number_of_blocks和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); ...}
您还可以将此示例扩展为利用相同想法处理 3 维操作。
现在您已经知道如何操作多维数据，还有另一个主要且大略的观点须要学习：如何在内核中调用函数。基本上，这只需利用 __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 版本上，我们的循环更少，从而加快了操作的处理速率。以下是 NxN 矩阵乘法的 CPU 和 GPU 性能比较：
Image by the author
正如您可能不雅观察到的，随着矩阵大小的增加，矩阵乘法运算的 GPU 处理性能提升乃至更高。
现在，考虑一个基本的神经网络，它紧张涉及 y = σ（Wx + b） 运算，如下所示：
Image by the author
这些操作紧张包括矩阵乘法、矩阵加法和将函数运用于数组，所有这些操作您都已经熟习并行化技能。因此，您现在能够从头开始实现自己的在 GPU 上运行的神经网络！
在这篇文章中，我们先容了有关 GPU 处理的先容性观点，以增强深度学习模型的性能。但是，同样主要的是要提到，您所看到的观点只是根本知识，还有很多东西须要学习。像 PyTorch 和 Tensorflow 这样的库实现了涉及其他更繁芜观点的优化技能，例如优化的内存访问、批处理操作等（它们利用基于 CUDA 构建的库，例如 cuBLAS 和 cuDNN）。但是，我希望这篇文章有助于澄清您在 GPU 上编写 .to（“cuda”） 并实行深度学习模型时幕后发生的事情。
