《CUDA编程》11.CUDA流

本章将介绍CUDA流

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

1. 核函数计算与数据传输之间的并行
2. 主机计算与数据传输之间的并行
3. 不同数据传输之间的并行
4. 核函数计算与主机计算之间的并行
5. 不同核函数之间的并行

为了实现上述所说的并行，需要合理的使用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流中，按照顺序依次执行，即：

1. 主机发出命令执行cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
2. 等命令1执行完毕之后，执行命令cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
3. 等命令2执行完毕后，主机发出命令执行sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N); ，注意：在发出调用核函数的命令之后，主机不会等待该命令执行完毕，因为此时是设备在执行操作，所以主机紧接着会发出下一个命令
4. 然而，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。