Reference

【CUDA 基础】6.1 流和事件概述

CUDA —— Stream and Event

CUDA 7 Stream流简化并发性

2.7.CUDA流

CUDA 官方文档

计算与传输重叠

CPU 与 GPU 间交互时涉及两个引擎：内存复制引擎和核函数执行引擎，内存复制引擎负责 CPU 与 GPU 间的数据传输，核函数执行引擎负责 CPU 向 GPU 部署核函数任务

这两个引擎可以理解为两个独立的并发任务队列，CPU 将任务添加到不同的队列中，GPU 驱动程序负责执行队列中的任务

由于这两个队列是相互独立且并发的，因此在 GPU 执行核函数时，可以同时进行 CPU 与 GPU 之间的数据传输，即计算与传输重叠（Overlap）

使用 cudaGetDeviceProperties() 函数可以查看 GPU 的设备信息，包括 CUDA 版本号、内存大小、最大线程数目等，通过其中的 deviceOverlap 属性可以判断该 GPU 是否支持计算与传输重叠

cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (prop.deviceOverlap) {
    printf("GPU 支持计算与传输重叠，可以使用流进行加速\n");
} else {
    printf("GPU 不支持计算与传输重叠, 无法使用流进行加速\n" );
}

CUDA 流

在程序中实现计算与重叠功能，需要使用 CUDA 流（CUDA Stream），其可以理解为一系列异步 GPU 操作，这些操作的执行顺序是按照 host 端代码中的顺序在 GPU 上执行的

流能封装一系列的异步操作，且保持这些操作在流中排队，使得在前面所有操作启动之后再启动后续的操作

一个流中的不同操作有着严格的顺序，但不同流之间没有任何限制，多个流同时启动多个内核，就形成了网格级别的并行

如图所示，灰色部分代表进行内存复制，蓝色部分代表执行核函数，当流 stream0 执行核函数时，流 stream1 能够进行内存复制，当流 stream1 执行核函数时，流 stream 0 和流 stream2 能够进行内存复制，这就将内存复制的执行时间尽可能的压缩，从而提高了程序执行效率

流的类型

CUDA 流可分为以下两种：

空流：隐式声明的流，即默认流
非空流：显式声明的流

如果没有声明一个流，那么所有 CUDA 操作都是在默认的空流中完成的，例如分配 device 端显存、host 端向 device 端传递数据、启动核函数等

但空流由于是隐式声明的，没有默认名，无法进行管理，因此若想控制流，就需要使用非空流

需要注意的是，在同一个流内不能实现计算与传输重叠，这是因为在同一个流中的计算需要的数据一般都是依赖于传输的数据，如果传输没有完成就开始计算，那么计算访问的数据就是错误的，所以 CUDA 中只能重叠不同流中的传输与计算

也就是说，如果没有显式声明非空流只使用空流的话，那么是无法实现计算与传输重叠的

流的使用

基本思路

利用流实现计算与传输重叠的基本思路为：

利用 cudaStreamCreate() 函数创建多个流
在每个流上利用 cudaMemcpyAsync() 函数将 host 端数据异步传输到 device 端中
在每个流上执行核函数
在每个流上利用 cudaMemcpyAsync() 函数将 device 端数据异步传输到 host 端中
利用 cudaStreamSynchronize() 流同步函数对多个流进行同步
利用 cudaStreamDestroy() 销毁创建的流

流的创建

CUDA 中封装了非空流类型 cudaStream_t，使用如下代码可以声明一个流

1	cudaStream_t stream;

在声明一个流后，这个流是无法使用的，需要使用 cudaStreamCreate() 来为其分配资源，函数原型如下：

1	__host__ cudaError_t cudaStreamCreate(cudaStream_t* pStream)

在非空流中需要启动核函数时，可在 <<<grid, block>>> 中附加的非空流的启动设置

1	kernel_fun<<<grid, block, sharedMemSize, stream>>>(...);

需要注意的是，host 端虚拟内存中分配的数据在物理内存中是随时可能被移动的，那么在执行异步数据传输时，必须要保证 host 端的内存是固定的，在整个生存周期中位置不变，否则如果操作系统移动了数据的物理地址，那么 device 可能会回到之前的物理地址取数据，导致出现未定义的错误

这就要求在 host 端分配内存时需要使用 cudaMallocHost() 来分配锁定内存

流的资源传递

CUDA 中使用 cudaMemcpyAsync() 函数将流的资源内存异步的复制到目标内存中，前三个参数与 cudaMemcpy() 函数类似，第四个参数指定在哪个流中进行资源传递

1	__host__ __device__ cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0)

流的同步

由于 CUDA 流的执行是异步的，那么就需要同步函数在必要的时候执行同步操作

1 2	__host__ cudaError_t cudaStreamSynchronize(cudaStream_t stream); __host__ cudaError_t cudaStreamQuery(cudaStream_t stream);

cudaStreamSynchronize() 会强制 host 端阻塞等待，直到流中所有操作执行完成；cudaStreamQuery() 会检查流中的操作是否全部完成，不会阻塞 host 端，如果流中所有操作都执行完成，那么会返回 cudaSuccess，否则返回 cudaErrorNotReady

流的销毁

cudaStreamDestroy() 用于回收非空流的资源，函数原型如下：

1	__host__ __device__ cudaError_t cudaStreamDestroy(cudaStream_t stream)

由于流和主机端是异步的，在使用 cudaStreamDestroy() 回收非空流的资源的时，很有可能流还在执行，这时候，该函数会正常执行，但不会立刻停止流，而是等待流执行完成后，立刻回收该流中的资源

流同步

同步流与异步流

流分成阻塞流和非阻塞流，在非空流中所有操作都是非阻塞的，所以流启动以后，host 端还要完成自己的任务，这时就需要在某些时刻去同步 host 端与 device 端流之间的进度，或者同步流和流之间的进度

对于流的两种类型来说，有：

空流：同步流，其中部分操作会造成阻塞，令 host 端等待操作完成
非空流：异步流，通常不会阻塞 host 端

阻塞流与非阻塞流

虽然非空流都是异步操作，不存在阻塞 host 端的情况，但有时可能会被空流中的操作阻塞，因此对于非空流，有：

阻塞流：当一个非空流被声明为阻塞流，会被空流阻塞
非阻塞流：当一个非空流被声明为非阻塞流时，对空流的阻塞行为失效

空流不需要任何显式声明，因此其是阻塞的，跟所有阻塞流同步，而使用 cudaStreamCreate() 创建流时，创建出来的流是阻塞流

使用 cudaStreamCreateWithFlags() 可以显式的创建阻塞流或非阻塞流，函数原型如下：

1	__host__ __device__ cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags)

其中，第二个参数就是选择要创建的流是阻塞的还是非阻塞的

// 默认为阻塞流
flags = cudaStreamDefault
// 非阻塞流，对空流的阻塞行为失效
flags = cudaStreamNonBlocking

实例

下述代码给出了多个非空流中调度 CUDA 操作的实例

#include <stdio.h>
#define N 5
#define STREAM_N 5
#define PI acos(-1.0)

__global__ void kernel(double *x, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride) {
        x[i] = pow(PI,i);
    }
}

int main() {
    cudaDeviceProp prop;
    int whichDevice;
    cudaGetDevice(&whichDevice);
    cudaGetDeviceProperties(&prop, whichDevice);
    if (prop.deviceOverlap) {
        printf("GPU 支持计算与传输重叠，可以使用流进行加速\n");
    } else {
        printf("GPU 不支持计算与传输重叠, 无法使用流进行加速\n" );
    }

    // 声明非空流
    cudaStream_t streams[STREAM_N];

    double *h_data[STREAM_N];
    double *d_data[STREAM_N];

    for (int i = 0; i < STREAM_N; i++) {
        // 为非空流分配资源
        cudaStreamCreate(&streams[i]);
        // 申请host锁定内存
        cudaMallocHost(&h_data[i], sizeof(double) * N);
        // 申请device内存
        cudaMalloc(&d_data[i], sizeof(double) * N);
    }

    for (int i = 0; i < STREAM_N; i++) {
        cudaMemcpyAsync(d_data[i], h_data[i], sizeof(double) * N, cudaMemcpyHostToDevice, streams[i]);
        kernel<<<1, 64, 0, streams[i]>>>(d_data[i], N);
        cudaMemcpyAsync(h_data[i], d_data[i], sizeof(double) * N, cudaMemcpyDeviceToHost, streams[i]);
    }

    // 流同步函数
    for (int i = 0; i < STREAM_N; i++) {
        cudaStreamSynchronize(streams[i]);
    }

    for (int i = 0; i < STREAM_N; i++) {
        for (int j = 0; j < N; j++) {
            printf("%lf ", h_data[i][j]);
        }
        printf("\n");
    }

    for (int i = 0; i < STREAM_N; i++) {
        // 释放host锁定内存
        cudaFreeHost(h_data[i]);
        // 释放device显存
        cudaFree(d_data[i]);
    }

    return 0;
}

Alex_McAvoy

NVIDIA CUDA2023春训营（七）CUDA 流