CUDA本可以更快：从零开始学习CUDA

实用 CUDA

CUDA是由NVIDIA开发的并行计算平台和编程模型，用于在图形处理单元上进行通用计算。我们可以用它来加速昂贵的计算，将负载分配到多个处理器上。例如，在一些 zk-SNARKs 中，我们必须计算一个 multiscalar multiplication，这涉及到对椭圆曲线上的大量点进行求和（例如，100,000,000），∑akPk∑akPk，其中PkPk是曲线上的点，而akak是正整数。我们还可以使用 CUDA 来解决其他任务高度可并行化的问题，例如求解微分方程、执行快速傅里叶变换、排序元素等。

CUDA 也广泛应用于机器学习算法，尤其是在涉及深度学习的算法中。它也常用于游戏引擎、图像处理和科学目的的模拟中。

在 GPU 加速的应用程序中，工作负载的顺序部分在 CPU 上运行，而处理大型数据块则在数千个 GPU 核心上并行运行。GPU 经过优化以运行此类工作！总体理念是不同的核心并行独立地运行同一组指令（SIMT 或 Single Instruction, Multiple Thread 模型）。

可以在 这里 找到 CUDA 编程模型的精彩介绍。

在这篇文章中，我们将专注于 CUDA 代码，使用 google colab 来展示和运行示例。但在我们开始编写代码之前，我们需要了解一些构建块。

构建块

通过 CUDA，我们可以并行运行多个线程来处理数据。这些线程被分组到不同的处理单元中，并具有自己的数据共享和同步原语。

逻辑处理单元

应用程序最基本的构建块是线程。然后将线程分组到 Warps 中，然后再分组到 Blocks 中，最后包含在 Grid 中。

根据我们的算法，可以忽略 Warps，或者像我们稍后将看到的那样，使用 Warps 来进一步优化我们的应用程序。

在撰写本文时，每个 Warp 有 32 个线程，每个 Block 有 1024 个线程或 32 个 Warps。

物理处理单元和内存

Blocks 在 Streaming Multiprocessors 中运行。每个流多处理器有 8 个 CUDA Cores*。这些核心也可以称为 Shaders 或 Streaming Processors。

一个繁忙的多处理器执行一个 Warp，它们的指令并行运行。由于 Warp 线程在同一个多处理器中运行，它们可以通过 Registers 快速交换信息。这很有用，因为在尽可能多的线程中运行我们的应用程序后，提高性能的方法是减少内存访问。

现在我们已经介绍了寄存器，我们可以问的下一个问题是我们如何在 Warps 之间以及 Blocks 之间共享信息？让我们向上遍历内存层次结构。

每个 Streaming Multiprocessor 都有一个 SRAM。其大小取决于显卡。例如，在 V100 中，它是 128 KiB，在 A100 中是 192 KiB。

这个 SRAM 有双重目的。首先，它以对程序员透明的方式用作 L1 cache。第二个用途是作为 Shared Memory。此共享内存使程序员能够以快速的方式在 Block 内共享数据。

由于 SRAM 具有两个功能，因此 CUDA 允许程序员定义有多少 SRAM 可以用作 L1 缓存，有多少用作共享内存。

最后，我们有 Global Memory。该内存是我们在显卡规范中看到的 GPU 内存，也是使用 cudaAlloc() 分配的内存。全局内存允许我们在线程块之间无缝地共享数据。

正如硬件中经常发生的那样，当我们移动到更大的内存时，操作会变得更加昂贵。

图片来自 Cuda Refresher - Nvidia blog

* Nvidia 还为其 Tensor Cores GPU 发布了一种新型核心，称为 Tensor Cores。这些核心可以以混合精度运行浮点数的小矩阵乘法作为本机操作，以进一步优化机器学习算法

CUDA 编程

开始 - 简单数组加法

我们将从并行化一些基本操作并仅使用全局内存开始。让我们开始编写一个添加两个数组的程序。

在我们开始之前，我们需要执行一些标准程序：

• 对于要在设备中使用的每个函数或内核，我们需要定义以下内容：
  • 我们使用多少个 Blocks？
  • 我们在每个 Block 中包含多少个线程？
  • Blocks 如何索引？
• 我们需要分配数据并将其复制到设备

为内核选择最佳参数是其自身的主题，但请务必记住，每个 Block 的线程数应为每个 Warp 的线程数（32）的倍数。

最后，我们需要决定如何索引 Blocks。我们可以将它们设置为作为 1、2 或 3 维数组访问。然后，我们在一个典型的数组、一个矩阵或一个立方体之间进行选择。

这只是设备的索引，无关紧要。但对于程序员来说，选择一些与要解决的问题相关的东西很有帮助。如果我们添加数组，则一个维度是合适的；如果我们正在处理图像，则 2 是最佳选择；并且，如果我们正在处理 3D 模型，则使用 3 维矩阵是有意义的。

在我们的例子中，我们将定义以下维度：

dim3 threadsPerBlock(128);
dim3 numBlocks(1024*1024);

如果我们想要一个二维数组，我们可以这样做

dim3 threadsPerBlock(128);
dim3 numBlocks(1024,1024);

现在，我们还需要在我们的设备中分配一些内存并复制我们想要添加的数组。

假设我们想要添加两个字节数组 array1 和 array2，大小为 AMOUNT_OF_ELEMENTS。然后我们可以为两个数组和一个结果保留字节：

char* array1_in_device;
char* array2_in_device;
char* result_in_device;

cudaMalloc(&array1_in_device, AMOUNT_OF_ELEMENTS);
cudaMalloc(&array2_in_device, AMOUNT_OF_ELEMENTS);
cudaMalloc(&result_in_device, AMOUNT_OF_ELEMENTS);

cudaMemcpy(array1_in_device, array1, AMOUNT_OF_ELEMENTS, cudaMemcpyHostToDevice);
cudaMemcpy(array2_in_device, array2, AMOUNT_OF_ELEMENTS, cudaMemcpyHostToDevice);

请注意，如果我们在 CUDA 中相加后不需要原始数组，我们不需要将结果存储在不同的位置。并且通常只使用一个 malloc，然后使用数据的地址对指针进行索引。但由于这是第一个程序，我们将使其尽可能简单。

现在，让我们专注于算法。

解决此问题的简单非 CUDA 代码如下所示：

for(int i = 0; i < MAX_ELEMENTS; i++)
    solution_array[i] = a[i] + b[i]

如果我们假设每个索引都有一个核心，我们可以删除 for 并让每个线程计算一个加法。情况并非总是如此，并且会产生不够灵活的解决方案。然后，我们将需要使用 strides。

Strides 无非是 for 循环中的步骤，用于在线程之间分配负载。例如，如果我们有一个 stride 为 4，则线程 0 将处理元素 0 3 7 11 ...，线程 1 将处理元素 1 4 8 12 ...，依此类推。

我们可以使用 CUDA 原语使我们的算法足够灵活以处理不同大小的数组和 Blocks，而不是将 stride 固定为一个数字。然后，我们使用 CUDA 的算法将变为：

__global__ void sum_arrays(char* array1, char* array2, char* result){

    uint globalThreadID = blockIdx.x*blockDim.x+threadIdx.x;
    uint stride = gridDim.x*blockDim.x;

    for (int i = globalThreadID; i < AMOUNT_OF_ELEMENTS; i += stride){
        result[i] = array1[i] + array2[i];
    }
}

这里的 __global__ 表示它是在设备上运行的可以从主机调用的函数。

blockIdx 是 Block 的 id，blockDim 是 Block 中元素的数量。ThreadIdx 是 Block 中线程的 id。请注意，通过执行

uint globalThreadID = blockIdx.x*blockDim.x+threadIdx.x;

我们获得了一个唯一的 ThreadID，独立于 Block，这对于分配工作很有用。

该 stride 定义为我们必须均匀分配工作的线程数。

最后，要从主机调用此函数，我们使用以下命令：

sum_arrays<<<numBlocks, threadsPerBlock>>>(
    array1_in_device, array2_in_device, result_in_device
);

完整的代码可以通过复制以下 google colab 来读取和运行。我们还添加了一些矩阵添加的示例，以显示具有更多维度时的索引如何工作。

Host - Device 并行

让我们继续以不同的方式使用相同的 sum_arrays() 函数，并检查另一种情况。

假设我们从主机在设备中调用我们的函数；之后，我们编写 CPU 的操作。在这种情况下会发生什么？代码是否运行，或者它是否等待设备？

为了回答第一个问题，让我们采取一些措施。

我们将编写一个程序，该程序对一个小数组进行大量工作，然后分两部分检索数据。我们还将测量调用函数以检索两部分所需的时间。

由于代码有点长，我们将把它留在我们之前使用的同一个 google colab 中，所以请随意复制并自己运行。

那会发生什么？

我们可以看到函数调用几乎不花费任何时间，并且第二个块的 memcpy 也很快。在两个函数的中间，第一个 memcpy 花费了大部分时间，几乎是第二个 memcpy 的 1000 倍！然而，操作是相同的。这是怎么回事？

答案是内核与主机并发运行，程序只有在需要数据时才会被阻止。memcpy 没有花费那么多时间，但它是第一个需要结果的函数调用，因此它必须等待设备完成。

为了使其更明显，我们将使用另一个原语：

cudaDeviceSynchronize();

使用此函数，所有时间都花在等待设备上，并且两个 memcpy 花费的时间相同。

并且知道我们可以在 GPU 和 CPU 中同时运行代码，我们可以进一步优化我们的密集型应用程序。

多流代码

知道当我们启动一个内核并尝试在本地运行代码时会发生什么，我们可以问以下问题：如果我们同时启动多个内核会发生什么？它们也可以并行运行吗？内存传输呢？

让我们尝试回答这些问题。

内核和 memcpy 函数在其流中按顺序运行。在我们之前看到的例子中，没有明确提到流，因此使用了默认流。

但是，我们可以创建更多可以使用的流，使用 cudaStreamCreate，然后将内核分配给新的流。

让我们看一个包含两个内核的示例：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
foo<<<blocks,threads,0,stream1>>>();
foo<<<blocks,threads,0,stream2>>>();
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

这样，如果一个内核不足以充分利用设备，我们可以用许多其他可以并行运行的任务来填充它。如果上一个示例的两个内核都使用了设备 50% 的功能，那么我们将拥有完全占用率。

由于我们有许多内核在运行，因此最好使用 memcpy 的异步版本，以便在数据传入时立即开始移动数据。

例如：

cudaMemcpyAsync(&results, &results_in_kernel1, AMOUNT_OF_ELEMENTS, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(&results, &results_in_kernel2, AMOUNT_OF_ELEMENTS, cudaMemcpyDeviceToHost, stream2);

假设计算花费大量时间在设备和主机之间传输数据。由于 GPU 支持同时传输和计算，因此异步内存传输可以与内核执行并行完成。

如果你想要更多这方面的例子，我们在 colab 中编写了一个完整的例子

同步和事件

我们如何同步工作？

在主机代码中，我们可以使用不同的同步级别。从多到少同步，我们可以使用的一些 API 调用是：

• 我们可以使用 cudaDeviceSynchronize() 同步所有内容，这将阻止主机，直到所有发出的 CUDA 调用完成；
• 我们可以使用 cudaStreamSynchronize(stream) 同步关于特定流的内容，这将阻止主机，直到 stream 中所有发出的 CUDA 调用完成；
• 或者我们可以使用 events 更有选择地同步主机或设备。

事件

CUDA 事件提供了一种机制来指示流中何时发生了操作。它们有助于分析和同步。

事件具有布尔状态：“Occurred”（这是默认状态）或“Not Occurred”。

管理事件和同步流

创建、删除和排队事件的最常见方式是：

• cudaEventCreate(&event) 创建一个 event；
• cudaEventDestroy(&event) 销毁一个 event；
• cudaEventRecord(&event, stream)
  • 将 event 状态设置为“Not Occurred”，
  • 将 event 排队到 stream 中，并且
  • 当 event 到达 stream 的前面时， event 状态设置为发生。

我们如何确保在继续执行之前发生了某些事件？

• 如果 event 已经发生，则 cudaEventQuery(event) 返回 CUDA_SUCCESS。
• cudaEventSynchronize(event) 阻止主机，直到 event 发生。
• cudaStreamWaitEvent(stream, event)
  • 在此调用之后，阻止 stream 上的所有启动，直到 event 发生
  • 不会阻止主机。

总结

CUDA 允许我们通过在 GPU 之间分配负载来加速昂贵的计算。为了充分利用这些功能，我们需要重新思考我们如何进行计算，寻找可以轻松并行化的算法（例如快速傅里叶变换）。在这篇文章中，我们回顾了 CUDA 的基础知识，什么是线程和 Warps，以及如何管理和同步事件。GPU 可以提供工具来缩短 zk-SNARK 中的证明和验证时间，从而为许多令人兴奋的应用程序打开大门。在以后的文章中，我们将介绍更多 CUDA 的高级主题。

• 原文链接： blog.lambdaclass.com/cud...
• 登链社区 AI 助手，为大家转译优秀英文文章，如有翻译不通的地方，还请包涵～