CUDA 编程简介

CUDA 编程简介

经常有人问我：“为什么深度学习模型都运行在 GPU 上？”

简而言之——这些深度模型在 GPU 上运行得更快。

接着问题又来了，“谁什么运行在 GPU 上会更快？”

那就要从 GPU 的工作原理和 CUDA 异构并行计算讲起了。

今天，我会用尽量通俗的语言为大家讲解 GPU 的工作原理和 CUDA 异构编程。

GPU 计算简史

GPU 最初是为了图形加速、二维和三维图形实时渲染而设计的。然而，由于 GPU 能够执行多个并行操作，其用途已经扩展到包括深度学习在内的其他应用领域。

GPU 在深度学习模型中的应用始于 2005 年左右，并在 2012 年随着 AlexNet 的出现而变得非常流行。AlexNet 是 Alex Krizhevsky、Ilya Sutskever 和 Geoffrey Hinton 设计的卷积神经网络，它在 2012 年赢得了 ImageNet 大规模视觉识别挑战（ILSVRC）。这一胜利是一个重要的里程碑，它证明了深度神经网络在图像分类中的有效性，以及在训练大型模型时使用 GPU 的优势。

继这一突破之后，使用 GPU 进行深度学习模型的训练和推理变得日益流行，这也促成了像 PyTorch 和 TensorFlow 这样的框架的创建。

如今，在 PyTorch 中，我们只需写上 .to("cuda") 就可以将数据发送到 GPU，并加速训练速度。但实际上，这背后隐藏了大量的优化与开发工作，下面我会一一为大家打开这个黑盒！

为什么 GPU 会更快

深度学习架构，如神经网络、CNN、RNN 和 Transformer，基本上是通过数学运算构建的，例如矩阵加法、矩阵乘法以及对矩阵应用函数。因此，如果我们找到一种优化这些操作的方法，我们就能提升深度学习模型的性能。

GPU 强大的并行处理能力使其非常适合执行这些大规模的数学运算。每个 GPU 包含数百到数千个处理核心，可以同时处理多个操作。与传统的 CPU 相比，CPU 通常只有较少的核心并侧重于执行连续任务。通过在 GPU 上执行并行计算，深度学习训练过程可以显著加速，从而大幅缩短模型训练时间并提高效率。

举个简单的例子，假设你想要将两个向量 $A$ 和 $B$ 相加，得到一个新的向量 $C$，即 $C=A+B$。
$$C=A+B=\left[\begin{array}{c}{a}_1\\ a_2\\ a_3\\ ⋮\\ a_n\end{array}\right]+\left[\begin{array}{c}{b}_1\\ b_2\\ b_3\\ ⋮\\ b_n\end{array}\right]=\left[\begin{array}{c}{a}_1+b_1\\ a_2+b2\\ a_3+b3\\ ⋮\\ a_n+b_n\end{array}\right]$$
在 C 语言中实现两个向量相加的代码如下：

void AddTwoVectors(flaot A[], float B[], float C[]) {
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
}

正如你所看到的，计算机必须遍历向量，在每次迭代中依次将每对元素相加。但这些操作是彼此独立的，第 i 对元素相加并不依赖于任何其他元素。如果我们能够同时执行这些操作，以并行方式将每对元素相加，那么理论上速度岂不是可以快 n 倍！

当然，一个简单的方法是使用 CPU 的多线程来并行执行所有计算。然而，当涉及到深度学习模型时，我们需要处理的是包含数百万元素的巨大向量。而常规 CPU 同时只能处理大约十几个线程。这就是 GPU 发挥用武之地的时候了！现代 GPU 可以同时运行数百万个线程，从而大大提高巨大向量上进行数学运算的性能。

GPU vs. CPU

尽管对于单个操作，CPU 的计算速度可能比 GPU 更快，但 GPU 的优势在于其并行计算能力。根本原因在于 CPU 和 GPU 的设计目标不同。CPU 设计用于尽可能快速地执行一系列操作（线程），且同一时间只能执行几十个线程；而 GPU 则设计为可以并行执行数百万个线程（虽然牺牲了单个线程的速度）。这种设计差异使得 GPU 在处理需要大规模并行处理的任务时更为有效。

NVIDIA 的这个视频很好地展示了 CPU 和 GPU 工作方式的差异。

NVIDIA 现场形象展示 CPU 和 GPU 工作原理上的区别

举个例子，我们可以将 CPU 想象成一辆法拉利，而 GPU 则是一辆公交车。如果你的任务是运送一个人，法拉利（CPU）是更好的选择。然而，如果你需要运送很多人，尽管法拉利（CPU）每趟的速度更快，但公交车（GPU）可以一次性将所有人运送到目的地，完成整体运送的速度会比法拉利往返几次快得多。因此，CPU 更适合处理顺序操作，而 GPU 更适合处理并行操作。这就是为什么在需要大量并行处理的计算任务中，使用 GPU 通常更为高效。

Ferrari vs. Bus （该图由 AI 生成）

为了提供更高的并行计算能力，GPU 设计为将更多的晶体管用于数据处理，而不是数据缓存和流程控制，这与 CPU 不同，CPU 将相当一部分晶体管用于缓存和控制，从而优化单线程性能和执行复杂指令。这种设计使得 GPU 在处理大规模数据并行计算时更加高效，而 CPU 则更擅长处理需要快速执行复杂任务和优化任务流程的场景。

下图展示了 CPU 与 GPU 在芯片资源分配上的差异。

CPU 拥有强大的核心以及更复杂的缓存架构（为此分配了大量的晶体管）。这种设计使得 CPU 在处理顺序操作时更加迅速。另一方面，GPU 优先配置大量的核心，以实现更高级别的并行性。这种核心数量的优势使得 GPU 在执行同时需要大量计算的任务时表现出色。

这就好比 CPU 是一个数学教授，而 GPU 是一百万个小学生。如果让它们解一个偏微分方程，再多小学生也自然无法胜任这种复杂任务。而如果是将两个 $1000\times 1000$ 的矩阵相加，一个老教授绝对干不过一群小学生，因为小学生可以每人认领一对元素将其相加，一百万个小学生同时算加法，瞬间就能完成两个大矩阵相加。

Professor vs. Pupils （该图由 AI 生成）

现在我们已经理解了这些基本概念，那么在实际应用中我们如何利用 GPU 的并行计算能力呢？这时，我们的主角 CUDA 就要粉墨登场了。

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 编写的基本代码中，程序在 host（CPU）上运行，将数据发送到 device（GPU），并启动 kernel（函数）在 device（GPU）上执行。这些 kernel 由多个线程并行执行。执行完成后，结果从 device（GPU）传回 host（CPU）。

让我们回到两个向量相加这个问题上来。之前的 C 语言代码如下：

#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);
    ...
}

我们可以看到，上面代码使用了一个 for 循环，迭代 $N$ 次，每次将一对数据现价。而在CUDA C/C++中，程序员可以定义称为 kernel 的C/C++函数，当调用这些函数时，它们将由 $N$ 个不同的 CUDA线程并行执行 $N$ 次。

要定义一个核函数，可以使用 __global__ 声明符，同时使用 <<<...>>> 符号来指定执行此 kernel 函数的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() {
    ...
    // N 个线程同时调用 kernel
    AddTwoVectors<<<1, N>>>(A, B, C);
    ...
}

每个线程执行一个 kernel 函数，并通过内建变量获得一个唯一的线程ID threadIdx。上述代码将两个大小为 $N$ 的向量 A 和 B 相加，并将结果存储到向量 C 中。你会注意到，CUDA 允许我们使用 $N$ 个并行线程同时执行所有这些操作，而不是使用循环来逐一执行每对加法，因此代码中不会出现 for 等循环结构。

但在我们可以运行这段代码之前，我们还需要做另一项修改。重要的是要记住，kernel 函数在 device（GPU）内运行。因此，所有的数据都需要存储在 device（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 的 device 指针

    // 在 device 上为向量 A, B, C 分配内存
    cudaMalloc((void **)&d_A, N * sizeof(float));
    cudaMalloc((void **)&d_B, N * sizeof(float));
    cudaMalloc((void **)&d_C, N * sizeof(float));

    // 从 host 中拷贝向量 A, B, C 到 device
    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);
    
    // 从 device 拷贝向量 C 到 host
    cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
}

在 CUDA 编程中，我们不能直接传递变量 A, B, C 给 kernel ，我们需要使用指针。CUDA kernel 操作 device 内存，因此需要将 device 指针（d_A, d_B, d_C）传递给 kernel 以便其进行操作。

此外，我们需要使用 cudaMalloc 在 device 上分配内存，并使用 cudaMemcpy 在 host 和 device 之间复制数据。

现在，我们可以添加向量 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 的 device 指针

    // 在 device 上为向量 A, B, C 分配内存
    cudaMalloc((void **)&d_A, N * sizeof(float));
    cudaMalloc((void **)&d_B, N * sizeof(float));
    cudaMalloc((void **)&d_C, N * sizeof(float));

    // 从 host 中拷贝向量 A, B, C 到 device
    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);
    
    // 从 device 拷贝向量 C 到 host
    cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

    // 释放 device 内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
}

此外，我们需要在调用 kernel 后添加 cudaDeviceSynchronize(); 这是一个用于 host 线程与 device 同步的函数。当调用该函数时，host 线程会等待，直到 device 上所有之前发出的 CUDA 命令完成后才继续执行。

还有一点很重要，我们需要添加一些 CUDA 错误检查，这样我们可以识别 GPU 上的错误。如果我们不添加这些检查，代码将继续执行 host 线程（CPU），这会使得很难识别与 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 的 device 指针

    // 在 device 上为向量 A, B, C 分配内存
    cudaMalloc((void **)&d_A, N * sizeof(float));
    cudaMalloc((void **)&d_B, N * sizeof(float));
    cudaMalloc((void **)&d_C, N * sizeof(float));

    // 从 host 中拷贝向量 A, B, C 到 device
    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);

    // 错误检查
    cudaError_t error = cudaGetLastError();
    if(error != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(error));
        exit(-1);
    }
    
    // 等待所有 CUDA 线程执行完毕
    cudaDeviceSynchronize();
    
    // 从 device 拷贝向量 C 到 host
    cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

    // 释放 device 内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
}

要编译和运行 CUDA 代码，你需要确保你的系统上安装了 CUDA 工具包。然后，你可以使用 NVIDIA CUDA 编译器 nvcc 来编译代码。如果你的机器上没有 GPU，你可以使用 Google Colab。你只需要在运行时→笔记设置中选择 GPU，然后将代码保存在一个 example.cu 文件中并运行：

%%shell
nvcc example.cu -o compiled_example # 编译
./compiled_example # 运行

# 你还可以使用 bug 检测工具 sanitizer 来运行代码
compute-sanitizer --tool memcheck ./compiled_example 

上面的代码虽然可以编译运行，但还有很大的优化空间。示例中 N = 1000，这个数字太小了，无法充分展示 GPU 的并行计算能力。此外，在处理深度学习问题时，我们通常要处理包含数百万参数的庞大向量。然而，如果我们将 N 设置为 N = 500000，并将运行 kernel 配置为 <<<1, 500000>>> 时，就会报错。因此，为了改进代码我们首先需要理解 CUDA 编程的一个重要概念：线程层次结构（Thread Hierarchy）。

线程层次结构

在 CUDA 编程中，调用 kernel 函数使用的符号是 <<<number_of_blocks, threads_per_block>>>。因此，在我们上面的例子中，我们运行了 1 个块（block），其中包含 N 个 CUDA 线程。然而，每个块能支持的线程数量是有限的。这是因为块内的每个线程都必须位于同一个流式多处理器核上，并且必须共享该核的内存资源。

你可以使用以下代码片段来获取这个限制：

int device;
cudaDeviceProp props;
cudaGetDevice(&device);
cudaGetDeviceProperties(&props, device);
printf("每个块的最大线程数: %d\n", props.maxThreadsPerBlock);

目前 Colab 所支持的 GPU 中，一个线程块最多可以包含 1024 个线程。因此，为了处理庞大向量，我们需要更多的块来执行更多的线程。此外，块被组织成网格，如下图所示：

现在，可以使用以下方式获取线程 ID：

int i = blockIdx.x * blockDim.x + threadIdx.x;

于是，我们的代码可以改造为：

#include <stdio.h>

// 定义 kernel
__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 的 device 指针

    // 在 device 上为向量 A, B, C 分配内存
    cudaMalloc((void **)&d_A, N * sizeof(float));
    cudaMalloc((void **)&d_B, N * sizeof(float));
    cudaMalloc((void **)&d_C, N * sizeof(float));

    // 从 host 中拷贝向量 A, B, C 到 device
    cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);

    // 多个块，每个块 threads_per_block 个线程调用 kernel
    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 error: %s\n", cudaGetErrorString(error));
        exit(-1);
    }

    // 等待所有 CUDA 线程执行完毕
    cudaDeviceSynchronize();

    // 从 device 拷贝向量 C 到 host
    cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

    // 释放 device 内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
}

性能比较

下面是对不同大小的两个向量进行相加操作，在 CPU 和 GPU 上性能的比较。

如上图所见，GPU 的优势只有在向量尺寸 N 很大时才变得明显。此外，需要牢记的是，这个时间比较仅考虑了 kernel 函数的执行时间。没有考虑 host 和 device 之间拷贝数据的时间，尽管在大多数情况下拷贝数据的时间可能不显著，但在我们的案例中相对来说是相当可观的，因为我们只执行了一个简单的加法操作。因此，重要的是要记住，GPU 计算只有在处理高度计算密集型且高度并行化的计算任务时，才能展现其优势。

多维线程

上面的例子处理的是简单的一维向量。但是在深度学习模型中，我们处理的更多是矩阵和张量。在之前的例子中，我们只使用了一维的块，每个块包含 N 个线程。其实，也可以执行多维线程块（最多 3 维）。因此，为了方便，如果你需要进行矩阵运算，你可以运行一个 $N\times M$ 线程的线程块。在这种情况下，你可以获取矩阵行列索引，行为 row = threadIdx.x，列为 col = threadIdx.y。此外，为了方便，你可以使用 dim3 类型的变量来定义number_of_blocks 和 threads_per_block。

下面的例子展示了如何将两个矩阵相加：

#include <stdio.h>

// 定义 kernel
__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 个线程同时调用 kernel
    dim3 threads_per_block(N, N);
    AddTwoMatrices<<<1, threads_per_block>>>(A, B, C);
    ...
}

当然我们也可以将其扩展为多个块：

#include <stdio.h>

// 定义 kernel
__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() {
    ...
    // number_of_blocks 个块，每个块 NxN 个线程同时调用 kernel
    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 函数

现在你已经知道如何操作多维数据了，接下来需要学习另一个重要且简单的概念：如何在 kernel 中调用函数。基本上，这是通过使用 __device__ 声明符来实现的。这定义了可以直接被 device（GPU）调用的函数。因此，它们只能从 __globle__ 或其他 __device__ 函数中被调用。下面的例子展示了如何对一个向量应用 Sigmoid 操作（这是深度学习模型中非常常见的操作）。

#include <math.h>

// Sigmoid 函数
__device__ float sigmoid(float x) {
    return 1 / (1 + expf(-x));
}

// 定义 kernel 对两个向量应用 sigmoid 函数
__global__ void sigmoidActivation(float input[], float output[]) {
    int i = threadIdx.x;
    output[i] = sigmoid(input[i]);
   
}

更多案例

现在你已经了解了 CUDA 编程的基本概念，接下来你就可以开始创建 CUDA 应用了。在深度学习模型中，基本上是一系列对矩阵或张量的运算，如加法、乘法、卷积、归一化等。我们看一下矩阵乘法如何并行化计算：
$$C_{m\times p}=A_{m\times n}\times B_{n\times p}$$

// 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 版本中我们使用了更少的循环，这导致操作处理更快。下面是两个版本的代码在 $N\times N$ 矩阵乘法上的性能比较：

如上图所见，随着矩阵大小的增加，GPU 处理的性能提升在矩阵乘法操作中更为显著。

我们还可以在 GPU 上实现神经网络。考虑一个基本的神经网络，其可以表示为 $y=\sigma (Wx+b)$ 操作，如下所示：

这些操作主要包括矩阵乘法、矩阵加法以及对数组应用函数，你只需要将上面学到的内容加以组合，就可以实现运行在 GPU 上的神经网络了！

总结

在这篇文章中，我介绍了关于 GPU 并行计算的入门概念。然而这些内容只是基础，要想真正掌握 GPU 异构编程还有很多需要学习的内容。像 PyTorch 和 Tensorflow 这样的库实现了多种优化技术，涉及到更多更复杂的概念，如优化的内存访问、批处理等（它们利用基于 CUDA 的库，如 cuBLAS 和 cuDNN）。然而，我希望这篇文章能帮助你了解当你编写 .to("cuda") 并在 GPU 上执行深度学习模型时幕后发生了什么。

在未来的文章中，我将尝试介绍更复杂的关于 CUDA 编程的概念。请在评论中告诉我你的想法，或者你希望我下次讲什么！非常感谢你的阅读！