为什么深度学习模型在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编程中的一个关键概念：线程的层级结构。

线程层次结构

核心函数的调用是通过<<>>这样的标记来完成的。比如，在我们之前的例子中，我们执行了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 上运行的您自己的神经网络！

总结

本文我们探讨了提升深度学习模型性能的GPU处理基础知识。PyTorch和TensorFlow等库应用了包含优化内存访问、批量处理等更高级概念的优化技术（它们使用了在CUDA基础上构建的库，比如cuBLAS和cuDNN）。希望本文能够帮助你理解当你执行.to("cuda")并利用GPU运行深度学习模型时，背后所发生的机制。