《CUDA编程》11.CUDA流

本章将介绍CUDA流

CUDA程序的并行层次主要有两个：一个是核函数内部的并行，一个是核函数外部的并行，核函数外部的并行主要指：

核函数计算与数据传输之间的并行
主机计算与数据传输之间的并行
不同数据传输之间的并行
核函数计算与主机计算之间的并行
不同核函数之间的并行

为了实现上述所说的并行，需要合理的使用CUDA流

1 CUDA流

定义：指由主机发出的在一个设备中执行的CUDA操作序列

一个CUDA流中各个操作的次序是由主机控制，并按照主机发布的次序执行。

来自两个不同CUDA流中的操作不一定按照某个次序执行，有可能交错或者并发执行

①默认流 / 空流： 默认流是指当没有显式指定流时，CUDA的API调用所使用的流。
②非默认流 / 非空流： 在CUDA编程中显式创建并使用的流

一个CUDA流由类型cudaStream_t的变量表示，可由如下CUDA运行时API产生：

cudaError_t cudaStreamCreate(cudaStream_t*);

该函数的输入参数是 cudaStream_t 类型的指针，返回一个错误代号。

CUDA 流可由如下 CUDA 运行时 API 函数销毁：

cudaError_t cudaStreamDestroy(cudaStream_t);

该函数的输入参数是 cudaStream_t 类型的变量，返回一个错误代号。

下面展示一个CUDA流的定义、产生和销毁：

cudaStream_t stream_1;
cudaStreamCreate(&stream_1);
cudaStreamDestroy(stream_1);

为了实现不同流之间的并发，主机在向某个 CUDA 流中发布一系列命令之后必须马上获得程序的控制权，不用等待该 CUDA 流中的命令在设备中执行完毕。为了检查一个 CUDA 流中的所有操作是否都在设备中执行完毕，CUDA 运行时 API 提供了如下两个函数：

cudaError_t cudaStreamSynchronize(cudaStream_t stream);强制阻塞主机，直到 CUDA 流 stream 中的所有操作都执行完毕。
cudaError_t cudaStreamQuery(cudaStream_t stream);不会阻塞主机，只是检查 CUDA 流 stream 中的所有操作是否都执行完毕

若是成功返回 cudaSuccess，否则返回 cudaErrorNotReady。

2 在默认流中重叠主机和设备计算

虽然在一个默认的CUDA流中的所有操作都是顺序执行的，但可以通过一些方法在默认流中重叠主机和设备的计算，我们通过下面遗一串代码来理解：

cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice); 
sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N); 
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);

以上的4中操作将在默认的CUDA流中，按照顺序依次执行，即：

主机发出命令执行cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
等命令1执行完毕之后，执行命令cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
等命令2执行完毕后，主机发出命令执行sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N); ，注意：在发出调用核函数的命令之后，主机不会等待该命令执行完毕，因为此时是设备在执行操作，所以主机紧接着会发出下一个命令
然而，cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);不会被立即执行，因为这是默认流中的 CUDA 操作，必须等待前一个 CUDA 操作（即核函数的调用）执行完毕才会开始执行。

如果我们能让主机调用核函数之后，同时去进行一些计算，就能提升主机的利用率，这也就是在默认流中重叠主机和设备计算

下面代码展示了一种做法：

#include <cuda_runtime.h>
#include <iostream>#define N 1024 * 1024
#define M N * sizeof(float)
#define THREADS_PER_BLOCK 256// CUDA 核函数：设备上执行数组相加
__global__ void sum(const float* x, const float* y, float* z, int n) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < n) {z[idx] = x[idx] + y[idx];}
}int main() {// 定义主机和设备内存float* h_x, * h_y, * h_z;float* d_x, * d_y, * d_z;cudaMallocHost((void**)&h_x, M); // 主机内存：分页锁定内存cudaMallocHost((void**)&h_y, M);cudaMallocHost((void**)&h_z, M);cudaMalloc((void**)&d_x, M);     // 设备内存cudaMalloc((void**)&d_y, M);cudaMalloc((void**)&d_z, M);// 初始化主机数据for (int i = 0; i < N; ++i) {h_x[i] = static_cast<float>(i);h_y[i] = static_cast<float>(i * 2);}// 异步将数据从主机传输到设备cudaMemcpyAsync(d_x, h_x, M, cudaMemcpyHostToDevice, cudaStreamDefault);cudaMemcpyAsync(d_y, h_y, M, cudaMemcpyHostToDevice, cudaStreamDefault);// 启动核函数计算int grid_size = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;sum << <grid_size, THREADS_PER_BLOCK >> > (d_x, d_y, d_z, N);// 主机端进行其他计算（在等待核函数完成的同时）float host_computation_result = 0.0f;for (int i = 0; i < N; i += 100) {host_computation_result += h_x[i] * h_y[i]; // 示例主机计算}std::cout << "Host computation result: " << host_computation_result << std::endl;// 异步从设备传输数据到主机cudaMemcpyAsync(h_z, d_z, M, cudaMemcpyDeviceToHost, cudaStreamDefault);// 同步等待设备所有任务完成cudaDeviceSynchronize();// 验证结果bool success = true;for (int i = 0; i < N; ++i) {if (h_z[i] != h_x[i] + h_y[i]) {success = false;break;}}if (success) {std::cout << "Array addition completed successfully!" << std::endl;}else {std::cout << "Error in array addition!" << std::endl;}// 释放内存cudaFree(d_x);cudaFree(d_y);cudaFree(d_z);cudaFreeHost(h_x);cudaFreeHost(h_y);cudaFreeHost(h_z);return 0;
}

运行结果：
在这里插入图片描述
通过这种方式，主机能够在 GPU 进行核函数计算时进行自身的运算，以实现主机和设备计算的重叠。

3 在非默认流中重叠多个核函数的执行

虽然在一个默认流中就可以实现主机计算和设备计算的并行，但是要实现多个核函数之间的并行必须使用多个 CUDA 流。

我们这里仅讨论使用多个非默认流的情况，使用非默认流时，核函数的执行配置中必须包含一个流对象。一个名为 my_kernel 的核函数可以用以下方法调用：

my_kernel<<<N_grid, N_block, N_shared, stream_id>>>(函数参数);

stream_id 是 CUDA 流的编号，说明核函数在编号为stream_id 的 CUDA 流中执行，而且使用了 N_shared 字节的动态共享内存，注意：如果使用非空流，但不想使用共享内存，则应将N_shared设置为0，不能忽略不写

my_kernel<<<N_grid, N_block, 0, stream_id>>>(函数参数);

3.1 重叠多个核函数的例子

#include <cuda_runtime.h>
#include <iostream>#define N 1024 * 1024
#define THREADS_PER_BLOCK 256// 核函数1：对数组每个元素加1
__global__ void kernelAddOne(float* data, int n) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < n) {data[idx] += 1.0f;}
}// 核函数2：对数组每个元素乘2
__global__ void kernelMultiplyTwo(float* data, int n) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < n) {data[idx] *= 2.0f;}
}int main() {// 定义主机和设备内存float* h_data, * d_data1, * d_data2;// 分配主机内存并初始化size_t size = N * sizeof(float);h_data = (float*)malloc(size);for (int i = 0; i < N; ++i) {h_data[i] = static_cast<float>(i);}// 分配设备内存cudaMalloc((void**)&d_data1, size);cudaMalloc((void**)&d_data2, size);// 将数据从主机传输到设备cudaMemcpy(d_data1, h_data, size, cudaMemcpyHostToDevice);cudaMemcpy(d_data2, h_data, size, cudaMemcpyHostToDevice);// 创建两个流cudaStream_t stream1, stream2;cudaStreamCreate(&stream1);cudaStreamCreate(&stream2);// 计算网格大小int gridSize = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;// 在不同流中启动核函数kernelAddOne << <gridSize, THREADS_PER_BLOCK, 0, stream1 >> > (d_data1, N);kernelMultiplyTwo << <gridSize, THREADS_PER_BLOCK, 0, stream2 >> > (d_data2, N);// 异步从设备传输数据到主机cudaMemcpyAsync(h_data, d_data1, size, cudaMemcpyDeviceToHost, stream1);cudaMemcpyAsync(h_data, d_data2, size, cudaMemcpyDeviceToHost, stream2);// 等待所有流完成cudaStreamSynchronize(stream1);cudaStreamSynchronize(stream2);// 清理cudaFree(d_data1);cudaFree(d_data2);cudaStreamDestroy(stream1);cudaStreamDestroy(stream2);free(h_data);std::cout << "Kernels executed in parallel streams!" << std::endl;return 0;
}

在 stream1 中调用 kernelAddOne，在 stream2 中调用 kernelMultiplyTwo。因为它们在不同流中，因此可以并行执行。

利用 CUDA 流并发多个核函数可以提升 GPU 硬件的利用率，减少闲置的 SM，从而从整体上获得性能提升。

4 在非默认流中重叠核函数的执行与数据传递

在上述代码中，我们发现“把数据从设备复制到主机”的代码使用的是cudaMemcpyAsync而不是以前使用的cudaMemcpy，前者便是后者的异步版本。

异步传输由 GPU 中的 DMA（direct memory access）直接实现，不需要主机参与。如果用同步的数据传输函数，主机在向一个流发出数据传输的命令后，将无法立刻获得控制权，必须等待数据传输完毕。

cudaMemcpyAsync 只比 cudaMemcpy 多一个参数。该函数的最后一个参数就是所在流的变量，异步传输函数的原型为：

cudaError_t cudaMemcpyAsync ( 
void *dst, 
const void *src, 
size_t count, 
enum cudaMemcpyKind kind, 
cudaStream_t stream 
);

在使用异步的数据传输函数时，需要将主机内存定义为不可分页内存（non-pageable memory）或者固定内存（pinned memory）
如果将可分页内存传给 cudaMemcpyAsync 函数，则会导同步传输。

4.1 不可分页内存

可分页内存（Pageable Memory）：默认情况下，主机（CPU）分配的内存都是“可分页内存“
不可分页内存（Pinned Memory）或“分页锁定内存”：固定在物理内存中的，CUDA 使用 cudaMallocHost 函数和cudaHostAlloc来分配这种内存。
cudaError_t cudaMallocHost(void** ptr, size_t size); cudaError_t
cudaHostAlloc(void** ptr, size_t size, size_t flags);
若函数 cudaHostAlloc 的第三个参数取默认值 cudaHostAllocDefault，则以上两个函数完全等价。

由以上函数分配的主机内存必须由如下函数释放：

cudaError_t cudaFreeHost(void* ptr);

如果不小心用了 free 函数释放不可分页主机内存，会出现运行错误。

4.2 示例分析

如果仅使用一个 CUDA 流（如默认流），那么以上 3 个操作在设备中一定是顺序的：
在这里插入图片描述
如果简单地将以上 3 个 CUDA 操作放入 3 个不同的流，相比仅使用一个 CUDA 流的情形依然不能得到加速，因为以上 3 个操作在逻辑上是有先后次序的。如果使用 3 个流，其执行流程可以理解如下：
要利用多个流提升性能，就必须创造出在逻辑上可以并发执行的 CUDA 操作。一个方法是将以上 3 个 CUDA 操作都分成若干等份，然后在每个流中发布一个 CUDA 操作序列。例如，使用两个流时，我们将以上 3 个 CUDA 操作都分成两等份。在理想情况下，它们的执行流程可以如下：
在这里插入图片描述
注意，这里的每个 CUDA 操作所处理的数据量只有使用一个 CUDA 流时的一半

如果 H2D、KER、和 D2H 这 3 个 CUDA 操作的执行时间都相同，那么就能有效地隐藏一个 CUDA 流中两个 CUDA 操作的执行时间，使得总的执行效率相比使用单个 CUDA 流的情形提升到 6/4 = 1:5 倍。

我们可以类似地分析使用更多流的情形。例如，当使用 4 个流并将每个流中的 CUDA 操作所处理的数据量变为最初的 1/4 时，在理想的情况下可以得到如下执行流程：
在这里插入图片描述
此时，总的执行效率相比使用单个 CUDA 流的情形提升到 12/6 = 2 倍。不难理解，随着流的数目的增加，在理想情况下能得到的加速比将趋近于 3。