CUDA的流和并发

YYRAN_ZZU

已于 2024-12-07 20:56:45 修改

阅读量1.4k

点赞数 30

CC 4.0 BY-SA版权

文章标签： linux gpu算力 c++

于 2024-12-07 19:34:53 首次发布

本文链接：https://blog.youkuaiyun.com/YYRAN_ZZU/article/details/144163833

流和并发

1. 流和事件概述
2. 并发内核执行
3. 重叠内核执行和数据传输
- 3.1 使用深度优先调度重叠例子
- 3.2 使用广度优先调度重建例子
4. 重叠GPU和CPU执行
5. 流回调

前面从编程模型，执行模型和内存模型的角度进行介绍，本节将研究网格级的并发。

1. 流和事件概述

CUDA流是一系列异步的CUDA操作(按照主机代码确定的顺序在设备上执行)，流封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的所有操作之后执行，并且可以查询排队操作的状态。从软件的角度开看，CUDA操作在不同的流中并发运行，从硬件上看，不一定总是如此。根据PCIe总线争用或每个SM资源的可用性，完成不同的CUDA流可能仍然需要互相等待。

1.1 CUDA流

所有的CUDA操作(包括内核和数据传输)都在一个流中显示或隐式的运行。流分为两种类型

隐式声明的流（空流）
显示声明的流(非空流）

如果没有显式的指定一个流，那么内核启动和数据传输将默认使用空流。另一方面，非空流可以被显示地创建和管理，如果想重叠不同的CUDA操作，必须使用非空流。基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发：

重叠主机计算和设备计算
重叠主机计算和主机与设备间的数据传输
重叠主机与设备间的数据传输和设备计算
并发设备计算

对于函数cudaMemcpy的异步版本：
cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0);
注意附加的流标识作为第五个参数，默认情况下，标识符被设置为默认流。
使用下面的函数创建一个非空流：
cudaError_t cudaStreamCreate(cudaStream_t *pStream);
当执行异步数据传输时，必须使用固定主机内存。可以使用cudaMallocHost函数或cudaHostAlloc函数分配固定内存
可以使用以下代码释放流中的资源：
cudaError_t cudaStreamDestroy(cudaStream_t stream);
在上一个流中，当cudaStreamDestroy函数被调用时，如果流中仍有为完成的工作，cudaStreamDestroy函数将立即返回，当流中所有工作都已完成时，与流相关的自愿将被自动释放。
因为所有的CUDA流操作都是异步的，所有CUDA的API提供两个函数来检查流中所有操作是否已经完成：
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
cudaStreamSynchronize 强制阻塞主机，直到给定流中所有操作都完成了
cudaStreamQuery会检查流中所有操作是否都已经完成，但是它们完成前不会阻塞主机

1.2 流调度

概念上讲，所有的流可以同时运行。但是，当将流映射到物理硬件时并不是总是这样的。

1.2.1 虚假的依赖关系

比如支持多路并发，可以多个网格同时执行，但是所有的流最终时被多路复用到单一的硬件工作队列中。

1.2.2 Hyper-Q技术

使用多个硬件队列，可以减少虚假的依赖关系。在不改变任何现有代码的情况下看到显著的性能提升。

1.3 流的优先级

使用下面函数可以创建一个具有特定优先级的流：
cudaError_t cudaStreamCreateWithPriority( cudaStream_t *stream, unsigned int flags, int priority );
高优先级流网格队列可以优先占有低优先级流已经执行的工作。对于一个给定的设备，可以使用以下的函数查询优先级的允许范围：
cudaError_t cudaDeviceGetStreamPriorityRange( int *leastPriority, int *greatestPriority );
如果当前的设备不支持流优先级，上面的函数将0返回给两个参数

1.4 CUDA 事件

CUDA中事件本质上时CUDA流中的标记，它与该流内操作中特定点相关联。使用事件来执行以下两个基本任务：

同步流的执行
监控设备的进展

CUDA的API提供了在流中任意点插入事件以及查询事件完成的函数。只有当一个给定CUDA流中先前的所有操作都执行结束后，记录在该流内的事件才会起作用。在默认流中指定实际爱你，适用于CUDA流中先前所有的操作。

1.4.1 创建和销毁

一个事件声明如下：
cudaEvent_t event;
一旦被声明，事件可以使用如下代码进行创建：
cudaError_t cudaEventCreate(cudaEvent_t *event);
销毁一个事件：
cudaError_t cudaEventDestroy(cudaEvent_t event);
当cudaEventDestroy函数被调用时，如果事件尚未起作用，则调用立即返回，当事件被标记完成时，自动释放与该事件相关的资源。

1.4.2 记录事件和计算运行时间

事件在流执行中标记了一个点，用来检查正在执行的流操作是否已经到了给定点。一个事件使用如下函数排队进入CUDA流：
cudaError_t cudaEventRecord( cudaEvent_t event, cudaStream_t stream = 0 );
等待一个事件会阻塞主机线程的调用，可以用下面的函数来执行：
cudaError_t cudaEventSynchronize(cudaEvent_t event);
使用如下函数测试一个事件是否可以不用阻塞主机应用程序来完成：
cudaError_t cudaEventQuery(cudaEvent_t event);
下面的函数来计算被两个事件标记的CUDA操作的运行时间：
cudaError_t cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end);
下面的示例代码演示如何将事件用于时间设备操作：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void kernel_example() {
   
   
    // GPU 上的代码
}

int main() {
   
   
    cudaEvent_t start, stop;
    float elapsedTime;
    cudaError_t err;
    // 创建事件
    err = cudaEventCreate(&start);
    if (err != cudaSuccess) {
   
   
        printf("Failed to create start event: %s\n", cudaGetErrorString(err));
        return -1;
    }
    err = cudaEventCreate(&stop);
    if (err != cudaSuccess) {
   
   
        printf("Failed to create stop event: %s\n", cudaGetErrorString(err));
        cudaEventDestroy(start);
        return -1;
    }
    // 记录开始事件
    err = cudaEventRecord(start, 0);
    if (err != cudaSuccess) {
   
   
        printf("Failed to record start event: %s\n", cudaGetErrorString(err));
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
        return -1;
    }
    // 启动 GPU Kernel
    kernel_example<<<1, 1>>>();
    // 注意：对于默认流，通常需要使用 cudaDeviceSynchronize() 来确保 Kernel 执行完成。
    // 但对于非默认流，可能需要使用其他同步机制或确保流中的事件已正确设置。
    cudaDeviceSynchronize(); // 确保 Kernel 执行完成（对于默认流是必需的）
    // 记录结束事件
    err = cudaEventRecord(stop, 0);
    if (err != cudaSuccess) {
   
   
        printf("Failed to record stop event: %s\n", cudaGetErrorString(err));
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
        return -1;
    }
    // 等待结束事件完成（对于某些情况可能是可选的，具体取决于程序逻辑）
    // err = cudaEventSynchronize(stop); // 如果前面已经使用了 cudaDeviceSynchronize()，这里可以省略
    // if (err != cudaSuccess) {
   
   
    //     printf("Failed to synchronize stop event: %s\n", cudaGetErrorString(err));
    //     cudaEventDestroy(start);
    //     cudaEventDestroy(stop);
    //     return -1;
    // }
    // 计算并打印执行时间
    err = cudaEventElapsedTime(&elapsedTime, start, stop);
    if (err != cudaSuccess) {
   
   
        printf("Failed to get elapsed time: %s\n", cudaGetErrorString(err));
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
        return -1;
    }
    printf("Kernel execution time: %.2f ms\n", elapsedTime);
    // 销毁事件
    err = cudaEventDestroy(start);
    if (err != cudaSuccess) {
   
   
        printf("Failed to destroy start event: %s\n", cudaGetErrorString(err));
    }
    err = cudaEventDestroy(stop);
    if (err != cudaSuccess) {
   
   
        printf("Failed to destroy stop event: %s\n", cudaGetErrorString(err));
    }
    return 0;
}

1.5 流同步

在非默认流中，所有的操作对于主机线程都是非阻塞的，因此会遇到需要在同一个流中运行主机和运算操作同步的情况。
从主机的角度来说，CUDA操作可以分为两大类：

内存相关操作
内核启动

对于主机来说，内核启动总是异步的。CUDA运行时也为内存操作的执行提供了异步函数：

异步流（非空流）
同步流（空流/默认流）

非空流可进一步分为以下两种类型：

阻塞流
非阻塞流

下面的部分中，将介绍如何使用阻塞流和非阻塞流

1.5.1 阻塞流和非阻塞流

CUDA运行时提供了一个定制函数，关于空流的非空流行为，代码如下：
cudaError_t cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags);
flags：用于设置流行为的标志位，可以指定不同的流类型。常用的标志位包括：

cudaStreamDefault：CUDA 流的默认创建标记。
cudaStreamNonBlocking：创建非阻塞流。在通过该标记创建的 CUDA 流中运行的工作可以与流 0（NULL 流，即默认流）中的工作同时运行，并且该流不应与默认流执行隐式同步。

1.5.2 隐式同步

CUDA包括两种类型的主机-设备同步：显式和隐式。前面介绍的执行显式同步函数，如cudaDeviceSynchronize,cudaSteamSynchronize以及cudaEventSynchronize函数属于被主机显式调用，使得在设备上任务执行和主机线程同步。
例如调用cudaMemcpy函数，可以隐式同步设备和主机，由于主机的应用程序在数据传输完成前会被阻塞。由于此函数的主要目的不是同步，因此其同步的产生是隐式的，
隐式同步行为的运行时函数可能会导致不必要的阻塞，这种阻塞通常发生在设备层面。例如：

锁页主机内存分配
设备内存分配
设备内存初始化
同一设备上两个地址之间的内存复制
一级缓存/共享内存配置的修改

1.5.3 显式同步

CUDA运行时在网格级支持显式同步CUDA程序的几种方法：

同步设备
同步流
同步流中的事件
使用事件跨流同步

下述函数可以阻塞一个主机线程直到设备完成所有先前的任务：
cudaError_t cudaDeviceSynchronize(void);
使用cudaStreamSynchronize函数可以阻塞主机线程直到流中所有操作完成为止，使用cudaStreamQuery函数可以完成非阻塞测试
使用cudaEventSynchronize函数和cudaEventQuery函数，CUDA事件也可以用于细粒度阻塞和同步
此外cudaStreamWaitEvent函数提供了一个使用CUDA事件引入流间依赖关系比较灵活的方法

1.5.4 可配置事件

CUDA运行时提供了一种方式来定制事件的行为和性能：
cudaError_t cudaEventCreateWithFlags( cudaEvent_t *event, unsigned int flags );
cudaEventDefault：默认事件，适合大多数用途。
cudaEventBlockingSync：阻塞同步事件，使用 cudaEventSynchronize 等待此事件完成时，会导致 CPU 等待 GPU 完成操作。
cudaEventDisableTiming：禁用计时的事件，创建的事件将不能用于测量时间（例如，不能使用 cudaEventElapsedTime）。
cudaEventInterprocess：进程间事件，允许不同进程间的事件同步（在支持该功能的系统上）

2. 并发内核执行

2.1 非空流的并发内核

使用NVIDIA的可视化性能分析器nvvp可视化并发核函数执行。

#include "../common/common.h"
#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>

/*
 * This example demonstrates submitting work to a CUDA stream in depth-first
 * order. Work submission in depth-first order may introduce false-dependencies
 * between unrelated tasks in different CUDA streams, limiting the parallelism
 * of a CUDA application. kernel_1, kernel_2, kernel_3, and kernel_4 simply
 * implement identical, dummy computation. Separate kernels are used to make the
 * scheduling of these kernels simpler to visualize in the Visual Profiler.
 */

#define N 300000
#define NSTREAM 4

__global__ void kernel_1()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_2()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_3()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_4()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

int main(int argc, char **argv)
{
   
   
    int n_streams = NSTREAM;
    int isize = 1;
    int iblock = 1;
    int bigcase = 0;

    // get argument from command line
    if (argc > 1) n_streams = atoi(argv[1]);

    if (argc > 2) bigcase = atoi(argv[2]);

    float elapsed_time;

    // set up max connectioin
    char* iname = "CUDA_DEVICE_MAX_CONNECTIONS";
    setenv (iname, "32", 1);
    char *ivalue =  getenv (iname);
    printf ("%s = %s\n", iname, ivalue);

    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("> Using Device %d: %s with num_streams=%d\n", dev, deviceProp.name,
           n_streams);
    CHECK(cudaSetDevice(dev));

    // check if device support hyper-q
    if (deviceProp.major < 3 || (deviceProp.major == 3 && deviceProp.minor < 5))
    {
   
   
        if (deviceProp.concurrentKernels == 0)
        {
   
   
            printf("> GPU does not support concurrent kernel execution (SM 3.5 "
                    "or higher required)\n");
            printf("> CUDA kernel runs will be serialized\n");
        }
        else
        {
   
   
            printf("> GPU does not support HyperQ\n");
            printf("> CUDA kernel runs will have limited concurrency\n");
        }
    }

    printf("> Compute Capability %d.%d hardware with %d multi-processors\n",
           deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);

    // Allocate and initialize an array of stream handles
    cudaStream_t *streams = (cudaStream_t *) malloc(n_streams * sizeof(
                                cudaStream_t));

    for (int i = 0 ; i < n_streams ; i++)
    {
   
   
        CHECK(cudaStreamCreate(&(streams[i])));
    }

    // run kernel with more threads
    if (bigcase == 1)
    {
   
   
        iblock = 512;
        isize = 1 << 12;
    }

    // set up execution configuration
    dim3 block (iblock);
    dim3 grid  (isize / iblock);
    printf("> grid %d block %d\n", grid.x, block.x);

    // creat events
    cudaEvent_t start, stop;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    // record start event
    CHECK(cudaEventRecord(start, 0));
    // dispatch job with depth first ordering
    for (int i = 0; i < n_streams; i++)
    {
   
   
        kernel_1<<<grid, block, 0, streams[i]>>>();
        kernel_2<<<grid, block, 0, streams[i]>>>();
        kernel_3<<<grid, block, 0, streams[i]>>>();
        kernel_4<<<grid, block, 0, streams[i]>>>();
    }
    // record stop event
    CHECK(cudaEventRecord(stop, 0));
    CHECK(cudaEventSynchronize(stop));

    // calculate elapsed time
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    printf("Measured time for parallel execution = %.3fs\n",
           elapsed_time / 1000.0f);
    // release all stream
    for (int i = 0 ; i < n_streams ; i++)
    {
   
   
        CHECK(cudaStreamDestroy(streams[i]));
    }
    free(streams);
    // destroy events
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));
    // reset device
    CHECK(cudaDeviceReset());
    return 0;
}

2.2 GPU上的虚假依赖关系

在不支持Hyper-Q的GPU上，为避免虚假的依赖关系，用广度优先的方法从主机中调度工作：

#include "../common/common.h"
#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>
/*
 * This example demonstrates submitting work to a CUDA stream in breadth-first
 * order. Work submission in breadth-first order prevents false-dependencies
 * from reducing the parallelism of an application. kernel_1, kernel_2,
 * kernel_3, and kernel_4 simply implement identical, dummy computation.
 * Separate kernels are used to make the scheduling of these kernels simpler to
 * visualize in the Visual Profiler.
 */
#define N 300000
#define NSTREAM 4

__global__ void kernel_1()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_2()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_3()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}

__global__ void kernel_4()
{
   
   
    double sum = 0.0;

    for(int i = 0; i < N; i++)
    {
   
   
        sum = sum + tan(0.1) * tan(0.1);
    }
}
int main(int argc, char **argv)
{
   
   
    int n_streams = NSTREAM;
    int isize = 1;
    int iblock = 1;
    int bigcase = 0;
    // get argument from command line
    if (argc > 1) n_streams = atoi(argv[1]);
    if (argc > 2) bigcase = atoi(argv[2]);
    float elapsed_time;
    // set up max connectioin
    char * iname = "CUDA_DEVICE_MAX_CONNECTIONS";
    setenv (iname, "32", 1);
    char *ivalue =  getenv (iname);
    printf ("%s = %s\n", iname, ivalue);
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("