https://zhuanlan.zhihu.com/p/699754357

好的，我们来对这段代码进行一次非常详尽的、逐行的分析。这段代码是理解CUDA流（Stream）和事件（Event）如何协同工作的绝佳实践范例。

一、 总体目标与设计思想

这段代码的核心目标是演示如何：

1. 创建和使用多个CUDA流（Stream） 来提交不同的GPU任务。
2. 使用 CUDA事件（Event） 在两个不同的流之间建立一个执行依赖关系。具体来说，就是让stream2中的kernel2必须等待stream1中的kernel1执行完毕后才能开始。
3. 在满足上述依赖关系的同时，实现其他独立任务的并发执行。具体来说，就是让stream1中的kernel3与stream2中的kernel2并行运行，以提高GPU的利用率。

代码通过对两块不同的设备内存（d_data1, d_data2）进行操作，并最终打印结果，来验证这个复杂的执行顺序和并发关系是否正确实现。

---

二、 代码逐段详解

1. 头文件和内核定义 (__global__ 函数)

#include <iostream>
#include <cuda_runtime.h>

• iostream: 包含C++的标准输入输出库，主要用于最后使用 std::cout 打印结果。
• cuda_runtime.h: 核心头文件。包含了所有你在代码中看到的 cuda... 函数（如 cudaMalloc, cudaStreamCreate 等）以及相关类型（如 cudaStream_t, dim3）的声明。

__global__ void kernel1(int64_t *data, int64_t repeat) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    for (size_t i = 0; i < repeat; i++) { data[idx] += 1; }
}

__global__ void kernel2(int64_t *data, int64_t repeat) {
    // ...
    for (size_t i = 0; i < repeat; i++) { data[idx] += 2; }
}

__global__ void kernel3(int64_t *data, int64_t repeat) {
    // ...
    for (size_t i = 0; i < repeat; i++) { data[idx] -= 1; }
}

• __global__: 关键字，表示这是一个CUDA内核，即在GPU上由大量线程并行执行的函数。
• kernel1, kernel2, kernel3: 三个简单的计算内核。
• int idx = ...: 这是计算每个线程全局唯一索引的标准公式，确保每个线程处理数组中的一个不同元素。
• for 循环: 这里的循环 (repeat 次) 是为了人为地增加内核的计算量，让它执行得久一点。如果内核瞬间完成，我们就无法观察到并发效果。
• 功能总结:
  • kernel1: 给每个数据元素 +1。
  • kernel2: 给每个数据元素 +2。
  • kernel3: 给每个数据元素 -1。

2. main 函数 - 初始化阶段

const int dataSize = 1024;
const int printSize = 10;
int64_t *h_data = new int64_t[dataSize]; // Host data
int64_t *d_data1, *d_data2; // Device data

• h_data: 在**主机（Host，即CPU）**内存中分配一个数组。h_ 前缀是常见的命名习惯。
• d_data1, d_data2: 声明两个指向**设备（Device，即GPU）**内存的指针。d_ 前缀也是命名习惯。

for (int i = 0; i < dataSize; i++) { h_data[i] = 0; }

• 将主机上的数组 h_data 初始化为全0。

cudaMalloc((void**)&d_data1, dataSize * sizeof(int64_t));
cudaMalloc((void**)&d_data2, dataSize * sizeof(int64_t));

• cudaMalloc: 在GPU上分配内存。这是GPU版的malloc。我们分配了两块独立的内存区域。

cudaMemcpy(d_data1, h_data, ..., cudaMemcpyHostToDevice);
cudaMemcpy(d_data2, h_data, ..., cudaMemcpyHostToDevice);

• cudaMemcpy: 将数据从主机拷贝到设备。
• cudaMemcpyHostToDevice: 明确指定拷贝方向是从CPU到GPU。这里我们将初始为0的h_data分别拷贝到d_data1和d_data2。

3. main 函数 - CUDA核心对象创建

dim3 blockDim(256);
dim3 gridDim((dataSize + blockDim.x - 1) / blockDim.x);

• 定义内核的启动配置。
• blockDim: 指定每个线程块（Block）包含256个线程。
• gridDim: 计算需要多少个线程块才能覆盖所有dataSize个数据。(size + block - 1) / block 是一种确保向上取整的常用技巧。

cudaStream_t stream1, stream2;
cudaEvent_t event1;

• 声明我们将要使用的CUDA对象：两个流和一事件。cudaStream_t 和 cudaEvent_t 只是代表这些对象句柄的类型。

int priorityHigh, priorityLow;
cudaDeviceGetStreamPriorityRange(&priorityLow, &priorityHigh);
cudaStreamCreate(&stream1);
cudaStreamCreateWithPriority(&stream2, cudaStreamDefault, priorityHigh);

• cudaDeviceGetStreamPriorityRange: 查询当前GPU支持的流优先级范围。priorityLow 通常是最高的优先级（数值小），priorityHigh 是最低的（数值大）。
• cudaStreamCreate: 创建一个默认优先级的流 stream1。
• cudaStreamCreateWithPriority: 创建 stream2 并赋予它一个高优先级。这只是一个给GPU调度器的提示，当多个流都准备好执行时，它倾向于先执行高优先级的。注意：这不能保证执行顺序，保证顺序必须用Event。

cudaEventCreate(&event1);

• 创建一个事件对象 event1。

4. main 函数 - 核心逻辑：任务提交与依赖设置

这是整个程序最关键的部分。CPU线程会按顺序执行以下代码，但这些操作在GPU上是异步执行的。

const int64_t repeat = 1000;

• 设置内核中循环的次数。

// Execute kernel1 in stream1
kernel1<<<gridDim, blockDim, 0, stream1>>>(d_data1, repeat);

• 第1步: 将 kernel1 的启动命令放入 stream1 的任务队列。CPU不等待，立即返回。
• stream1 队列状态: [kernel1]

cudaEventRecord(event1, stream1); // Record event1 after kernel1 execution in stream1

• 第2步: 将一个“记录事件”的命令放入 stream1 的队列。
• 含义: “stream1，当你执行完你队列里当前所有的任务（即 kernel1）之后，请把 event1 这个‘旗帜’插上，标记为已完成。”
• stream1 队列状态: [kernel1, record_event1]

// Execute kernel2 in stream2, waiting for event1
cudaStreamWaitEvent(stream2, event1, 0);

• 第3步: 将一个“等待事件”的命令放入 stream2 的任务队列。
• 含义: “stream2，请你停下来，在继续执行你队列里的后续任务之前，必须一直等到 event1 这面‘旗帜’被插上为止。”
• stream2 队列状态: [wait_for_event1]

kernel2<<<gridDim, blockDim, 0, stream2>>>(d_data1, repeat);

• 第4步: 将 kernel2 的启动命令放入 stream2 的队列。它被放在了 wait_for_event1 命令的后面。
• stream2 队列状态: [wait_for_event1, kernel2]

// Execute kernel3 in stream1 on a different array
kernel3<<<gridDim, blockDim, 0, stream1>>>(d_data2, repeat);

• 第5步: 将 kernel3 的启动命令放入 stream1 的队列。它被放在了 record_event1 命令的后面。
• stream1 队列状态: [kernel1, record_event1, kernel3]

5. main 函数 - 同步、验证与清理

// Synchronize streams
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

• 至关重要的一步。到目前为止，CPU已经把所有任务都异步提交了，但GPU可能还在计算。
• cudaStreamSynchronize(stream) 是一个阻塞操作，它会暂停CPU线程，直到指定流中的所有任务（包括计算、记录事件、等待事件等）全部完成。
• 我们在这里同步两个流，是为了确保在进入下一步（从GPU拷贝数据回CPU）之前，所有的GPU计算都已经百分之百完成了。

cudaMemcpy(h_data, d_data1, ..., cudaMemcpyDeviceToHost);
std::cout << "Data after kernel1 and kernel2:" << std::endl;
// ... loop to print h_data ...

• 将 d_data1 的最终结果从GPU拷贝回CPU，并打印前10个元素。
• 预期结果: 初始为0，kernel1 执行后变为1，kernel2 等待 kernel1 完成后再执行，变为 1 + 2 = 3。所以打印结果应为一串 3。

cudaMemcpy(h_data, d_data2, ..., cudaMemcpyDeviceToHost);
std::cout << "Data after kernel3:" << std::endl;
// ... loop to print h_data ...

• 将 d_data2 的最终结果从GPU拷贝回CPU，并打印。
• 预期结果: 初始为0，只有 kernel3 对其操作，结果应为 -1。

// Free device memory and destroy streams and event
cudaFree(d_data1);
cudaFree(d_data2);
delete[] h_data;
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
cudaEventDestroy(event1);

• 资源清理。这是一个好习惯，释放所有分配的资源。
• cudaFree: 释放 cudaMalloc 分配的GPU内存。
• delete[]: 释放 new[] 分配的CPU内存。
• cudaStreamDestroy/cudaEventDestroy:销毁创建的流和事件对象。

---

三、 总结

该代码通过 cudaEventRecord 和 cudaStreamWaitEvent 的精妙组合，成功地构建了一个任务依赖图（DAG），实现了：

1. 串行依赖: kernel2 严格在 kernel1 之后执行。
2. 并行执行: kernel2 和 kernel3 在时间上重叠执行。

最终的打印结果验证了这一复杂调度逻辑的正确性，完美展示了如何利用Stream和Event进行高级CUDA性能优化。

好的，我们来详细分解这两行代码。这是 CUDA 编程中最基础也是最核心的部分之一，它决定了你的内核（Kernel）将以何种方式、由多少个线程来执行。

让我们一步步来理解。

CUDA 的执行模型：Grid、Block、Thread

首先，你需要理解 CUDA 的三层执行模型：

1. Thread (线程): 这是执行的最小单位。成千上万个线程同时执行同一个内核函数。
2. Block (线程块): 线程被组织成“线程块”。一个块内的线程可以相互协作（通过共享内存和同步）。一个块最多可以有1024个线程。
3. Grid (网格): 所有的线程块组合在一起，形成一个“网格”。一次内核启动就是启动一个网格。

当你启动一个内核时，你需要告诉 CUDA：

• Grid 的维度：这个网格有多大？（即包含多少个线程块）
• Block 的维度：每个线程块有多大？（即包含多少个线程）

dim3 类型就是用来定义这些维度的。它是一个可以存储一维、二维或三维尺寸的结构体。

---

dim3 blockDim(256); —— 定义每个线程块的大小

• dim3: 这是一个 CUDA 提供的数据类型，用于表示维度，可以看作是一个包含 x, y, z 三个成员的结构体。
• blockDim: 这是我们给变量起的名字，意为“Block Dimension”（块维度）。
• (256): 这是构造函数的参数。它等价于 dim3 blockDim(256, 1, 1);。

这行代码的含义是：“我决定将我的线程组织成一维的线程块，每个线程块里包含 256 个线程。”

所以，当你启动内核时，GPU 知道每一个工作小组（Block）都有256名工人（Threads）。

为什么是256？
这是一个经验性的选择。线程块的大小通常是32的倍数（因为GPU以32个线程为一组，称为一个Warp，来调度执行）。256（即 8 * 32）是一个非常常见且通常性能不错的选择。其他常见值还有128, 512, 1024等。

---

dim3 gridDim((dataSize + blockDim.x - 1) / blockDim.x); —— 计算需要多少个线程块

这行代码稍微复杂一点，但它的目标很明确：计算需要启动多少个线程块，才能保证有足够的总线程数来处理所有的数据。

让我们把它拆开来看：

• gridDim: 我们给变量起的名字，意为“Grid Dimension”（网格维度）。
• dataSize: 我们需要处理的数据总量。在这个例子中是 1024。
• blockDim.x: 每个线程块中的线程数量。在这里就是我们上面定义的 256。

核心问题: 假设我们有 1024 个数据元素，每个线程处理一个元素。我们每个线程块有 256 个线程。那么，我们需要多少个线程块？

简单的数学:
总线程数 / 每个块的线程数 = 块的数量
1024 / 256 = 4
所以，我们需要4个线程块。

但如果 dataSize 不是 blockDim.x 的整数倍呢？
比如说，dataSize 是 1025。
1025 / 256 = 4.0039...
我们不能启动4.0039个块。我们必须启动5个块才能覆盖所有1025个数据。第5个块虽然只需要处理1个数据，但我们必须启动整个块。多余的线程在内核代码里通常会通过一个判断条件（if (idx < dataSize)）来提前退出，不做任何事。

这就是向上取整 (Ceiling Division) 的需求。

表达式 (dataSize + blockDim.x - 1) / blockDim.x 就是在C++中实现整数向上取整的一个经典技巧。

我们来验证一下：

• 情况1：正好整除

  • dataSize = 1024, blockDim.x = 256
  • (1024 + 256 - 1) / 256
  • = (1279) / 256
  • = 4 (因为是整数除法，小数部分被截断)
  • 结果正确。

• 情况2：不能整除

  • dataSize = 1025, blockDim.x = 256
  • (1025 + 256 - 1) / 256
  • = (1280) / 256
  • = 5
  • 结果正确。

所以，这行代码的含义是：“请计算出为了处理 dataSize 个元素，在每个块有 blockDim.x 个线程的情况下，我总共需要启动的线程块数量（向上取整），并用这个数量来定义我们一维网格的大小。”

总结

放在一起看：

// 我要处理1024个数据
const int dataSize = 1024;

// 决定把工人们（线程）分成小组（块），每组256人。
dim3 blockDim(256);

// 计算为了完成1024个任务，需要多少个这样的小组。
// (1024 + 256 - 1) / 256 = 4。所以需要4个小组。
// 于是，我们的项目（网格）由4个小组构成。
dim3 gridDim((dataSize + blockDim.x - 1) / blockDim.x);

当内核启动时 kernel<<<gridDim, blockDim, ...>>>，GPU就会：

1. 启动一个包含 4 个线程块的网格。
2. 每个线程块里包含 256 个线程。
3. 总共启动的线程数是 4 * 256 = 1024 个。
4. 这1024个线程会并行地执行内核代码，每个线程处理 data 数组中的一个元素。