CUDA C 编程权威指南 Grossman 第6章 流和并发

6.1 流和事件概述

CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流能够封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的所有操作之后执行，并且可以查询排队操作的状态。这些操作包括在主机与设备间进行数据传输，内核启动一级大多数由主机发起但由设备处理的其他命令。

流中操作的执行相对于主机总是异步的。

我们的的任务是使用CUDA的API来确保一个异步操作在运行结果被使用之前可以完成。在同一个CUDA流中的操作有严格的执行顺序，而在不同的CUDA流中的操作在执行顺序上不受限制。

使用多个流同时启动多个内核，可以实现网格级并发。

在同一时间内将流中排队的操作与其他有用的操作一起执行，可以隐藏执行那些操作的开销。

在许多情况下，执行内核比传输数据耗时更多。在这些情况下，可以完全隐藏CPU和GPU之间的通信延迟。通过将内核执行和数据传输调度到不同的流中，这些操作可以重叠，程序运行的总时间将被缩短。流在CUDA的API调用粒度上可实现流水线和双缓冲技术。

CUDA的API函数一般可以分为同步或异步。具有同步行为的函数会阻塞主机端线程，直到它们完成。具有异步行为的函数被调用后，会立即将控制权返还给主机。

异步函数和流是在CUDA中构建网格级并发的两个基本支柱。

从软件的角度来看，CUDA操作在不同的流中并发运行；而从硬件上来看，不一定总是如此。根据PCIe总线争用或每个SM资源的可用性，完成不同的CUDA流可能任然需要互相等待。

6.1.1 CUDA流

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

隐式声明的流(空流、默认流)；

显式声明的流；

如果没有显示地指定一个流，那么内核启动和数据传输将默认使用空流。

非空流可以被显式地创建和管理。如果想要重叠不同的CUDA操作，必须使用非空流。

基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发;

        重叠主机计算和设备计算；

        重叠主机计算和设备间的数据传输；

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

        并发设备计算；

cudaMemcpy();

kernel<<<>>>;

cudaMemcpy();

要理解一个CUDA程序，应该从设备和主机两个角度去考虑。

从设备的角度来看，以上3个操作都被发布到默认的流中，并且按发布顺序执行。设备不知道其他被执行的主机操作。

从主机的角度来看，每个数据传输都是同步的，在等待它们完成时，都将强制空闲主机时间。

内核启动是异步的，所以无论内核是否完成，主机的应用程序几乎都立即恢复执行。这种内核启动的默认异步情况使它可以直接重叠设备和主机计算。

数据传输也可以被异步发布，但是必须显式地设置一个CUDA流来装载它们。CUDA提供了cudaMmcpyAsync函数。

cudaStreamCreat创建了一个可以显式管理的非空流。之后，就可以被当做流参数共异步API函数来使用。

在使用异步函数API时，常见的疑惑在于它们可能从先前启动的异步操作中返回错误代码。

当执行异步数据传输时，必须使用固定主机内存。可以用cudaMallocHost或者cudaHostAlloc。

在主机虚拟内存中固定分配，可以确保其在CPU内存中的物理位置在应用程序的整个生命周期保持不变。否则，操作系统可以随时自由改变主机虚拟内存的物理位置。

1.  
   // 声明
2.  
   cudaStream_t stream;
3.  
   // 初始化
4.  
   cudaStreamCreat(&stream);
5.  
   // 释放
6.  
   cudaStreamDestroy(stream);

在一个流中，当cudaStreamDestroy函数被调用时，如果流中仍有未完成的工作，cudaStreamDestory函数将立即返回，当流中所有的工作都已完成时，与流相关的资源将自动释放。

因为所有CUDA流操作都是异步的，所以CUDA的API提供了两个函数来检查流中所有操作是否都已经完成。

cudaStreamSynchronize强制阻塞主机，直到在给定流中所有操作都完成了。

cudaStreamQuery会检查流中所有操作是否都已经完成，但在它们完成前不会阻塞主机。当所有操作都完成时cudaStreamQuery函数会返回cudaSuccess，当一个或多个操作扔在执行或等待执行时返回cudaErrorNotReady。

1.  
   for (int i = 0; i < nStreams; i )
2.  
   {
3.  
   int offset = i * bytesPerStream;
4.  
   cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, streams[i]);
5.  
   kernel<<<grid, block, 0, streams[i]>>>(&d_a[offset]);
6.  
   cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
7.  
   }
8.  
    
9.  
   for (int i = 0; i < nStreams; i )
10.  
    cudaStreamSynchronize(streams[i]);

在图中来看，数据传输虽然分布在不同的流中，但是并没有并发执行。这是由于一个共享资源导致的：PCIe总线。具有双工PCIe总线的设备可以重叠两个数据传输，但它们必须在不同的流中以及不同的方向中。

并发内核的数目的最大数量是依赖设备而确定的。

6.1.2 流调度

6.1.2.1 虚假的依赖关系

虽然Fermi GPU支持16路并发，即多达16个网格同时执行，但是所有的流最终任务值被多路复用到单一的硬件工作队列中。这种单一的流水线可能会导致虚假的依赖关系。在工作队列中，一个被阻塞的操作会将队列中该操作的所有操作都阻塞，即使它们属于不同的流。

6.1.2.2 Hyper-Q技术

Kepler的Hyper-Q技术，使用多个工作队列，从而减少了虚假的依赖关系。Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的连接，允许多个CPU线程或进程在单一GPU上同时启动工作。

Hyper-Q使用32个硬件工作队列，每个流分配一个工作队列。如果创建的流超过32个，多个流将共享一个硬件工作队列。这样做的结果是可实现全流级并发，并且其具有最小的虚假流间依赖关系。

6.1.3 流的优先级

对于计算能力3.5以上的设备，可以给流分配优先级。使用cudaStreamCreatWithPriority()创建了一个具有指定整数优先级的流，并在流中返回一个句柄。这个优先级是与流中的工作调度相关的。高优先级流的网格可以优先占有低优先级流已经执行的任务。

流优先级不会影响数据传输操作，只对计算内核有影响。

如果指定的优先级超过设备定义的范围，它会被自动限制为定义范围内的最低值或最高值。

可以用cudaDeviceGetStreamPriorityRange()查询优先级的允许范围。

按照惯例，一个较低的整数值表示更高的优先级。如果当前设备不支持流优先级，cudaDeviceStreamGetPriorityRange()将0返回给这两个参数。

6.1.4 CUDA事件

CUDA中事件本质是CUDA流中的标记，它与该流内操作流中特定点相关联。

可以使用事件来执行以下两个基本任务：

        同步流的执行；

        监控设备的进展；

CUDA的API提供了在流中任意点插入事件以及查询事件完成的函数。

只有当一个给定CUDA流中先前的所有操作都执行结束后，记录在该流中的事件才会起作用。在默认流中指定的事件，适用于CUDA流中先前的所有操作。

6.1.4.1 创建和销毁

1.  
   // 声明
2.  
   cudaEvent_t event;
3.  
   // 创建
4.  
   cudaEventCreate();
5.  
   // 销毁
6.  
   cudaEventDestory();

当cudaEventDestroy()函数被调用时，如果事件尚未起作用，则调用立即返回，当事件被标记完成时自动释放与该事件相关的资源。

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

事件在流中标记了一个点。它们可以用来检查正在执行的流操作是否已经到达了给定点。

当事件从工作队列中取出来时，这个操作的唯一作用就是通过主机端标记来指示完成状态。

一个事件使用如下函数排队进入CUDA流中，cudaEventRecord()。

已经排队进入CUDA流中的事件可用于等待或测试在指定流中先前操作的完成情况。

等待一个事件会阻塞主机线程的调用，它可以用cudaEventSynchronize()来执行。

对于流来说，cudaEventSynchronize函数类似于cudaStreamSynchronize函数，但cudaEventSynchronize函数允许主机等待流执行的中间点。

可以使用如下代码测试一个事件是否可以不用阻塞主机应用程序来完成：cudaEventQuery()。

cudaEventElapsedTime()函数用来记录两个事件标记的CUDA操作的运行时间，返回启动和停止之间的运行时间。

事件的启动和停止不必在同一个CUDA流中。请注意，如果在非空流中记录启动事件或停止事件时，返回的时间可能比预期的要大。这是因为cudaEventRecord函数是异步的，并且不能保证计算的延迟正好处于两个事件之间。

1.  
   cudaEvent_t start, stop;
2.  
   cudaEventCreat(start);
3.  
   cudaEventCreat(stop);
4.  
    
5.  
   cudaEventRecord(start);
6.  
   kernel<<<grid, block>>>(args);
7.  
   cudaEventRecord(stop);
8.  
   cudaEventSynchronize(stop);
9.  
    
10.  
    float time;
11.  
    cudaEventElapsedTime(&time, start, stop);
12.  
     
13.  
    cudaEventDestroy(start);
14.  
    cudaEventDestroy(stop);

6.1.5 流同步

在非默认流中，所有的操作对于主机线程都是非阻塞的，因此会遇到需要在同一个流中运行主机和运算操作同步的情况。

从主机的角度出发，CUDA操作可以分为两大类：

        内存相关操作；

        内核启动；

对于主机来说，内核启动总是异步的。许多内存操作本质上是同步的，但是CUDA运行时也为内存操作提供了异步函数。

有两种类型的流：

        异步流(非空流)；

        同步流(空流);

在主机上非空流是一个异步流，其上所有操作都不阻塞主机执行。另一方面，被隐式声明的空流是主机上的同步流。

大多数添加到空流上的操作都会导致主机在先前所有的操作上阻塞，主要的异常是内核启动。

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

        阻塞流；

        非阻塞流；

虽然非空流在主机上是非阻塞的，但是非空流内的操作可以被空流中的操作所阻塞。如果一个非空流是阻塞流，则空流可以阻塞该非空流中的操作。如果一个非空流是非阻塞的，则它不会阻塞空流中的操作。

6.1.5.1 阻塞流和非阻塞流

使用cudaStreamCreate()函数创建的流是阻塞流，这意味着在这些流中操作执行可以被阻塞，一直等到空流中先前的操作执行结束。

空流是隐式流，在相同的CUDA上下文中，它和其他所有的阻塞流同步。

一般情况下，当操作被发布到空流中，在该操作被执行之前，CUDA上下文会等待所有先前的操作发布到所有的阻塞流中。此外，任何发布到阻塞流中的操作，会被挂起等待，直到空流中先前的操作执行结束才开始执行。

1.  
   kernel<<<1, 1, 0, stream_1>>>();
2.  
   kerne2<<<1, 1>>>();
3.  
   kerne3<<<1, 1, 0, stream_2>>>();

直到kernel_1执行结束，kernel_2才会在GPU上执行，kernel‘_2执行结束后，kernel_3才开始执行。

CUDA运行时提供了一个定制函数，它是关于空流的非空流行为，代码如下;

cudaStreamCreatWithFlags()。cudaStreamNonBlocking标记使得非空流对于空流的阻塞行为失效。在kernel1-3案例中，如果stream1,2都是用NonBlocking标记创造的话，那么所有核函数执行都不会被阻塞。

6.1.5.2 隐式同步

CUDA包含两种类型的主机-设备同步：显式和隐式。在前面已经介绍了许多执行显式同步的函数，如cudaDeviceSynchronize()，’cudaStreamSynchronize(), cudaEventSynchronize()函数被主机显式调用，使得在设备上任务执行和主机线程同步。在应用程序的逻辑点，可以手动插入显式同步调用。

前文介绍的隐式同步的例子，如cudaMemcpy函数，这是由于主机的应用程序在数据传输完成 之前会被阻塞。

理解隐式同步很重要，因为无意中调用隐式同步主机和设备的函数，可能会导致意想不到的性能下降。

带有隐式同步行为的运行时函数可能会导致不必要的阻塞，这种阻塞通常发生在设备层面上。许多与内存相关的操作意味着当前设备上所有先前的操作都有阻塞。例如：

        锁页主机内存分配；

        设备内存分配；

        同一设备上两个地址之间的内存复制；

        一级缓存/共享内存配置的修改；

6.1.5.3 显式同步

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

        同步设备；cudaDeviceSynchronize()使主机线程等待直到所有和当前设备相关的计算和通信完成。应尽量少用此函数。

        同步流；cudaStreamSynchronize()阻塞主机线程直到流中所有的操作完成为止，使用cudaStreamQuery函数可以完成非阻塞测试。

        同步流中事件；cudaEventSynchronize();CUDA事件也可用于细粒度阻塞和同步。

        使用事件跨流同步；

使用cudaStreamWaitEvent函数提供了一个使用CUDA事件引入流间依赖关系比较灵活的方法。

在流中执行任何排队的操作之前，并且在cudaStreamWaitEvent调用之后，cudaStreamWaitEvent函数能使指定流等待指定事件。该事件可能与同一个流相关，也可以与不同的流相关。

6.1.5.4 可配置事件

CUDA运行时提供了一种方式来定制事件的行为和性能，代码如下：cudaEventCreateWithFlags()

cudaEventBlockingSync指定使用cudaEventSynchronize函数同步事件会阻塞调用的线程，调用的线程在另一个将要休眠的线程或进程上运行，而不是放弃核心，直到事件满足依赖关系。

cudaEventDisableTimeing表明创建的事件只能用来进行同步，不需要基类时序数据。除去时间戳花费的总开销，提高了调用cudaStreamWaitEvent和cudaEventQuery函数调用的性能。

cudaEventInterprocess表明创建的事件可能被用作进程间事件。

6.2 并发内核执行

第一个实例演示了如何使用多个流并发运行多个核函数。这个简单 的例子将介绍并发内核执行的几个基本问题：

        使用深度优先或广度优先方法的调度工作；

        调整硬件工作队列；

        在Kepler设备或Fermi设备上避免虚假的依赖关系；

        检查默认流的阻塞行为；

        在非默认流之间添加依赖关系；

        检查资源使用是如何影响并发的；

6.2.1 非空流中的并发内核

1.  
   __global__ void kernel_1()
2.  
   {
3.  
   double sum = 0;
4.  
   for (int i = 0; i < N; i )
5.  
   sum = tan(0.1) * tan(0.1);
6.  
   }
7.  
    
8.  
   cudaStream_t* streams = (cudaStream_t*)malloc(n_streams * sizeof(cudaStream_t));
9.  
   for (int i = 0; i < n_streams; i )
10.  
    cudaStreamCreate(&streams[i]);
11.  
     
12.  
    dim3 grid, block;
13.  
     
14.  
    cudaEvent_t start, stop;
15.  
    cudaEventCreat(&start);
16.  
    cudaEventCreat(&stop);
17.  
     
18.  
    cudaEventRecord(start);
19.  
    for (int i = 0; i < n_streams; i )
20.  
    {
21.  
    kernel_1<<<grid, block, 0, streams[i]>>>();
22.  
    kernel_2<<<grid, block, 0, streams[i]>>>();
23.  
    kernel_3<<<grid, block, 0, streams[i]>>>();
24.  
    kernel_4<<<grid, block, 0, streams[i]>>>();
25.  
    }
26.  
     
27.  
    cudaEventRecord(stop);
28.  
     
29.  
    cudaEventSynchronize(stop);
30.  
    cudaEventElapsedTime(&elapsed_time, start, stop);

6.2.2 Fermi GPU上的虚假的依赖关系

因为每个流的第一个任务不依赖之前的任何任务，并且有可用的SM，所以它可以立即执行。之后，调度流i 1的第二个任务，然而它对第一个任务的依赖却阻止它被执行，这就会导致任务执行再次被阻塞。

这种虚假的依赖关系是由主机调度内核的顺序引起的。

利用深度优先方法得到的工作队列中的任务如图所示。由于所有流被多路复用到一个硬件工作队列中，所以前面的流就连续阻塞了后面的流。

采用广度优先顺序可以确保工作队列中相邻的任务来自于不同的流。因此，任何相邻的任务对之间都不存在虚假的依赖关系，从而实现了并发内核执行。

6.2.3 使用openmp调度操作

本节将使用多个主机线程将操作调度到多个流中，并使用一个线程来管理每一个流。

在使用OpenMP的同时使用CUDA,不仅可以提高便携性和生产效率，而且还可以提高主机代码的性能。

1.  
   omp_set_num_threads(n_streams)
2.  
   #pragma omp parallel
3.  
   {
4.  
   int i = omp_get_thread_num();
5.  
   kernel_1<<<grid, block, 0, streams[i]>>>();
6.  
   kernel_2<<<grid, block, 0, streams[i]>>>();
7.  
   kernel_3<<<grid, block, 0, streams[i]>>>();
8.  
   kernel_4<<<grid, block, 0, streams[i]>>>();
9.  
   }

将OpenMP创建的主机线程与CUDA流间一对一映射。

在一般情况下，如果每个流在内核执行之前，期间或之后有额外的工作待完成，那么它可以包含在同一个OpenMP并行区域里，并且跨流和线程进行重叠。

6.2.4 用环境变量调整流的行为

虽然，Kepler设备支持的硬件工作队列的最大数量是32。然而，默认情况下并发硬件连接的数量被限制为8，由于每个连接都需要额外的内存和资源，所以设置默认为8，减少了不必要的全部32个工作队列的用意程序的资源消耗。

可以使用CUDA_DEVICE_MAX_CONNECTIONS来调整并行硬件连接的数量。

对于bash和bourne Shell而言，代码如下;

        export CUDA_DEVICE_MAX_CONNECTIONS = 32

对于C Shell而言

        setenv CUDA_DEVICE_MAX_CONNECTIONS = 32

当流的数量超过硬件连接的数量，多个流将共享一个连接。当多个流共享相同的硬件工作队列时，可能会产生虚假的依赖关系。

6.2.5 GPU资源的并发限制

有限的内核资源可以抑制应用程序中可能出现的内核并发的数量。

通常，会创建成百货数千个线程。有了这么多线程，可用的硬件资源可能会成为并发的主要限制因素，因为它们阻止启动符合条件的内核。

6.2.6 默认流的阻塞行为

如果某个内核在默认流中被启动，所以再非空流上所有之后的操作都会被阻塞，知道默认流中的操作被完成。

6.2.7 创建流间依赖关系

在复杂的应用程序中，引入流间依赖关系式很有用的，它可以在一个流中阻塞操作直到另一个流中的操作完成。事件可以用来添加流间依赖关系。

1.  
   cudaEvent_t* kernelEvent = (cudaEvent_t*)malloc(n_streams * sizeof(cudaEvent_t));
2.  
   for (int i = 0; i < n_streams; i )
3.  
   cudaEventCreatWithFalgs(&kernelEvent[i], cudaEventDisableTiming);
4.  
    
5.  
   for (int i = 0; i < n_streams; i )
6.  
   {
7.  
   kernel_1<<<grid, block, 0, streams[i]>>>();
8.  
   kernel_2<<<grid, block, 0, streams[i]>>>();
9.  
   kernel_3<<<grid, block, 0, streams[i]>>>();
10.  
    kernel_4<<<grid, block, 0, streams[i]>>>();
11.  
     
12.  
    cudaEventRecord(kernelEvent[i], streams[i]);
13.  
    cudaStreamWaitEvent(streams[n_streams - 1], kernelEvent[i], 0);
14.  
    }

从结果来看，最后一个流，在其他所有流完成后才能开始启动工作。

6.3 重叠内核执行和数据传输

由于Fermi和Kepler设备有两个复制引擎：一个用于数据传输到设备，另一个用于从设备将数据提取出来。因此，最多重叠两个数据传输，并且只有当它们的方向不同且被调度到不同的流时才能这样做。否则，所有的数据传输都是串行的。

在应用程序中，还需要检测数据传输和内核执行之间的关系，从而区分以下两种情况：

        如果一个内核使用数据A，那么对A的数据传输必须要安排在内核启动前，且必须位于同一流中；

        如果一个内核完全不使用数据A，那么内核执行和数据传输就可以位于不同的流中。

6.3.1 使用深度优先的调度重叠

为了在向量的加法中实现重叠，需要将输入和输出数据集划分成子集，并将来自一个子集的通信与来自于其他子集的计算进行重叠。

为了在向量加法中实现重叠，需要将输入和输出数据集划分为子集，并将来自于一个子集的通信与来自于其他子集的计算进行重叠。

要重叠数据传输和内核执行，必须使用异步复制函数。因为异步复制函数需要固定的主机内存，所以首先需要使用cudaHostAlloc函数，在固定主机内存上修改主机数组的分配：

1.  
   cudaHostAlloc((void**)&gpuRef, nBytes, cudaHostAllocDefault);
2.  
   cudaHostAlloc((void**)&hostRef, nBytes, cudaHostAllocDefault);
3.  
    
4.  
   int iEle = nElem / NSTREAMS;
5.  
   for (int i = 0; i < NSTREAMS; i )
6.  
   {
7.  
   int offset = i * iEle;
8.  
   cudaMemcpyAsync(&d_A[offset], &h_A[offset], iBytes, cudaMemcpyHostToDevice, streams[i]);
9.  
   cudaMemcpyAsync(&d_B[offset], &h_B[offset], iBytes, cudaMemcpyHostToDevice, streams[i]);
10.  
    sumArray<<<grid, block, 0, streams[i]>>>(&d_A[offset], &d_B[offset], d_C[offset], iEle);
11.  
     
12.  
    }

由于这些内存复制和内核启动对于主机而言是异步的，因此全部的工作负载都可以毫无阻塞的在流之间分配。通过将数据传输和该数据上的计算放在同一个流中，输入向量，内核计算以及输出向量之间的依赖关系可以被保持。

此时出现了3中重叠：

        不同流中的内核互相重叠；

        内核和其他流中的数据传输重叠；

        在不同流以及不同方向的数据传输的互相重叠；

也出现了2种阻塞：

        内核被同一流中先前的数据传输所阻塞；

        从主机到设备的数据传输被同一方向上先前的数据传输所阻塞；

网格管理单元(Grid Management Unit)

        Kepler引入了一个新的网格管理和调度控制系统，即网格管理单元；

        GMU可以暂停新网格的调度，使网格排队等待且暂停网格直到它们准备好执行，这样就使运行时变得非常灵活强大，动态并行就是一个很好的例子。

         在Feimi设备上，网格直接从流队列中被传到CUDA工作分配器。

        GMU创建了多个硬件工作队列，从而减少或消除了虚假的依赖关系。

6.3.2 使用广度优先调度重叠

1.  
   for (int i = 0; i < NSTREAMS; i )
2.  
   {
3.  
   int offset = i * iEle;
4.  
   cudaMemcpyAsync(&d_A[offset], &h_A[offset], iBytes, cudaMemcpyHostToDevice, streams[i]);
5.  
   cudaMemcpyAsync(&d_B[offset], &h_B[offset], iBytes, cudaMemcpyHostToDevice, streams[i]);
6.  
    
7.  
   }
8.  
    
9.  
   for (int i = 0; i < NSTREAMS; i )
10.  
    {
11.  
    int offset = i * iEle;
12.  
    sumArray<<<grid, block, 0, streams[i]>>>(&d_A[offset], &d_B[offset], d_C[offset], iEle);
13.  
    }
14.  
     
15.  
    for (int i = 0; i < NSTREAMS; i )
16.  
    {
17.  
    int offset = i * iEle;
18.  
    cudaMemcpyAsync(&d_C[offset], &h_C[offset], iBytes, cudaMemcpyDeviceToHost, streams[i]);
19.  
     
20.  
    }

在Kepler设备上，与深度优先的方法相比，广度优先没有明显差异。在Fermi设备上，因为虚假依赖，在整体性能方面，使用广度优先方法不如深度优先方法。由于主机到设备复制队列上的争用导致虚假依赖关系，在主机到设备的完成之气那，将阻止所有内核启动。

6.4 重叠CPU和GPU执行

本节示例主要包括两个部分：

        内核被调度到默认流中；

        等待GPU内核时执行主机计算；

1.  
   cudaMemcpyAsync(&d_A, &h_A, iBytes, cudaMemcpyHostToDevice);
2.  
   cudaMemcpyAsync(&d_B, &h_B, iBytes, cudaMemcpyHostToDevice);
3.  
   sumArray<<<grid, block>>>(&d_A, &d_B, d_C, iEle);
4.  
   cudaMemcpyAsync(&d_C, &h_C, iBytes, cudaMemcpyDeviceToHost);

GPU和CPU执行重叠是比较简单的，因为所有的内核启动在默认情况下都是异步的。

6.5 流回调

流回调是另一种到CUDA流中排列等待的操作。一旦流回调之前所有的流操作全部完成，被流回调指定的主机端函数就会被CUDA运行时所调用。此函数由应用程序提供，并允许任何主机端逻辑插入到CUDA流中。

流回调是另一种GPU和CPU同步机制。

流回调函数是由应用程序提供的一个主机函数，并在流中使用一下API函数注册：

cudaStreamAddCallback。此函数为提供的流添加了一个回调函数。在流中所有先前排队的操作完成后，回调函数才能在主机上执行。

每使用cudaStreamAddCallback一次，只执行一次回调，并阻塞队列中排在其后面的工作，直到回调函数完成。

对于回调函数的两个限制：

        从回调函数中不可以调用CUDA的API函数；

        从回调函数中不可以执行同步；

一般来说，对互相关联或其他CUDA操作相关的回调顺序做任何假设都是有风险的，可能会导致代码的不稳定。

1.  
    
2.  
   void CUDART_CB my_callback(cudaStream_t stream, cudaError_t status, void* data)
3.  
   printf("call back from stram %d\n", *((int*) data));
4.  
    
5.  
   for (int i = 0; i < n_streams; i )
6.  
   {
7.  
   kernel_1<<<grid, block, 0, streams[i]>>>();
8.  
   kernel_2<<<grid, block, 0, streams[i]>>>();
9.  
   kernel_3<<<grid, block, 0, streams[i]>>>();
10.  
    kernel_4<<<grid, block, 0, streams[i]>>>();
11.  
     
12.  
    cudaStreamAddCallback(streams[i], my_callback, (void*)(stream_ids i), 0);
13.  
    }

6.6 总结

流允许高级CUDA操作在独立的流中排队执行，CUDA流可以实现粗粒度并发。

因为CUDA支持异步操作和大多数版本的运行时函数，所以它可以在多个CUAD流中调度计算和通信。

从概念上说，如果CUDA操作之间存在依赖，则它们必须在同一个流中被调度。没有依赖关系的操作可以在任何流中被调度。

在CUDA中，通常有3种不同类型的重叠方案来隐藏计算和通信延迟：

        在设备上重叠多个并发的内核；

        重叠带有传入或传出设备数据传输的CUDA内核；

        重叠CPU与GPU执行；

需要注意一下问题：

        平衡内核资源需求和并发资源需求。在设备上一次启动过多的计算任务，可能会导致内核串行，这使得硬件资源的工作块变得可用。

        如果可能，避免使用默认流执行异步操作。放在默认流中的操作可能会阻塞其他非默认CUDA流的进展。

        在Fermi设备上，从深度优先和广度优先两方面考虑主机的调度。这个选择可以通过消除共享硬件工作队列上的虚假依赖关系，显著影响其性能。

        要注意隐式同步的函数，并且充分利用它们和异步函数来避免性能的降低。

这篇好文章是转载于：学新通技术网