10.4 使用单个CUDA流

稍后将会看到，仅当使用多个流时才能显现出流的真正威力，但我们首先通过在应用程序中使用单个流来说明流的用法。假设有一个CUDA C核函数，该函数带有两个输入数据缓冲区，a和b。核函数将对这些缓冲区中相应位置上的值执行某种计算，并将生成的结果保存到输出缓冲区c。矢量加法示例就采用了类似的计算模式，但在这个示例中，我们将计算a中三个值和b中三个值的平均值：

include"../common/book.h"   
#define N (1024\*1024)   
#define FULL_DATA_SIZE  $(N^{*}20)$    
.global_void kernel(int \*a,int \*b,int \*c）{ intidx  $=$  threadIdx.x  $^+$  blockIdx.x \*blockDim.x; if(idx< N）{ int idx1  $= (\mathrm{idx} + 1)\% 256$  · int idx2  $= (\mathrm{idx} + 2)\% 256$  float as  $= (\mathrm{a[idx]} + \mathrm{a[idx1]} + \mathrm{a[idx2]}) / 3.0f;$  float bs  $= (\mathrm{b[idx]} + \mathrm{b[idx1]} + \mathrm{b[idx2]}) / 3.0f;$  c[idx]  $= (\mathrm{as} + \mathrm{bs}) / 2$  1

这个核函数不是很重要，因此不需要花太多时间来理解该函数执行的操作。可以把它看成是一个抽象函数，因为这个示例中重要的是函数main()中与流相关的代码。

int main(void) {  
   udaDeviceProp prop;  
    int whichDevice;  
    HANDLE_ERROR(udaGetDevice(&whichDevice));  
    HANDLE_ERROR(udaGetDeviceProperties(&prop, whichDevice));  
if (!prop_deviceOverlap) {  
    printf("Device will not handle overlaps, so no "speed up from streams\n");  
    return 0;  
}

我们做的第一件事情就是选择一个支持设备重叠（Device Overlap）功能的设备。支持设备重叠功能的GPU能够在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。正如前面提到的，我们将使用多个流来实现这种计算与数据传输的重叠，但首先来看看如何创建和使用一个流。与其他需要测量性能提升（或者降低）的示例一样，首先创建和启动一

个事件计时器：

cudaEvent_t start, stop; float elapsedTime;

//启动计时器

HANDLE_ERROR(udaEventCreate(&start));  
HANDLE_ERROR(udaEventCreate(&stop));  
HANDLE_ERROR(udaEventRecord(start, 0));

在启动计时器后，接下来创建在应用程序中使用的流：

//初始化流

CUDAStream_t stream;   
HANDLE_ERROR( CUDAStreamCreate(&stream));

这就是创建流时需要的全部工作，并没有太多值得注意的地方，接下来是数据分配操作。

int *host_a, *host_b, *host_c;  
int *dev_a, *dev_b, *dev_c;

//在GPU上分配内存

HANDLE_ERROR(udaMalloc(void**)&dev_a, N * sizeof(int) );  
HANDLE_ERROR(udaMalloc(void**)&dev_b, N * sizeof(int) );  
HANDLE_ERROR(udaMalloc(void**)&dev_c, N * sizeof(int) );

//分配由流使用的页锁定内存

HANDLE_ERROR(udaHostAlloc(void**)&host_a, FULL_DATA_SIZE * sizeof(int),udaHostAllocDefault) );   
HANDLE_ERROR(udaHostAlloc(void**)&host_b, FULL_DATA_SIZE * sizeof(int),udaHostAllocDefault) ;   
HANDLE_ERROR(udaHostAlloc(void**)&host_c, FULL_DATA_SIZE * sizeof(int),udaHostAllocDefault) ;   
for (int i=0; i<FULL_DATA_SIZE; i++) { host_a[i] = rand(); host_b[i] = rand(); }

我们在GPU和主机上分别分配好了输入内存和输出内存。注意，由于程序将使用主机上的固定内存，因此调用cudaHostAlloc()来执行内存分配操作。使用固定内存的原因并不只在于使

复制操作执行得更快，还存在另外一个好处。稍后将更详细地进行分析，我们将使用一种新的CUDAMemcpy()函数，并且在这个新函数中需要页锁定主机内存。在分配完输入内存后，调用C的库函数rand()并用随机整数填充主机内存。

在创建了流和计时事件，并且分配了设备内存和主机内存后，就准备好了执行一些计算。通常，我们会将这个阶段一带而过，只是将两个输入缓冲区复制到GPU，启动核函数，然后将输出缓冲区复制回主机。我们将再次沿用这种模式，只是进行了一些小修改。

首先，我们不将输入缓冲区整体都复制到GPU，而是将输入缓冲区划分为更小的块，并在每个块上执行一个包含三个步骤的过程。我们将一部分输入缓冲区复制到GPU，在这部分缓冲区上运行核函数，然后将输出缓冲区中的这部分结果复制回主机。想象一下需要使用这种方法的一种情形：GPU的内存远少于主机内存，由于整个缓冲区无法一次性填充到GPU，因此需要分块进行计算。执行“分块”计算的代码如下所示：

//在整体数据上循环，每个数据块的大小为N

for (int i=0; i<FULL_DATA_SIZE; i+=N) { // 将锁定内存以异步方式复制到设备上  
    HANDLE_ERROR(udaMemcpyAsync( dev_a, host_a+i, N * sizeof(int),udaMemcpyHostToDevice, stream));  
    HANDLE_ERROR(udaMemcpyAsync( dev_b, host_b+i, N * sizeof(int),udaMemcpyHostToDevice, stream));  
    kernel<<N/256,256,0,stream>>>(dev_a, dev_b, dev_c);  
// 将数据从设备复制到锁定内存  
    HANDLE_ERROR(udaMemcpyAsync( host_c+i, dev_c, N * sizeof(int),udaMemcpyDeviceToHost, stream));

你将注意到，在前面的代码段中有两个不同之处。首先，代码没有使用熟悉的CUDAMemcpy()，而是通过一个新函数CUDAMemcpyAsync()在GPU与主机之间复制数据。这些函数之间的差异虽然很小，但却很重要。CUDAMemcpy()的行为类似于C库函数memcpy()。尤其是，这个函数将以同步方式执行，这意味着，当函数返回时，复制操作就已经完成，并且在输出缓冲区中包含了复制进去的内容。

异步函数的行为与同步函数相反，通过名字cusamemcpyAsync()就可以知道。在调用cusamemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动，更无法保证它是否已经结束。我们能够得到的保证是，复制操作肯定会当下一个被放入流中的操作之前执行。任何传递给cusamemcpyAsync()的主机内存指针都必须已经通过cusahostAlloc()分配好内存。也就是，你只能以异步方式对页锁定内存进行复制操作。

注意，在核函数调用的尖括号中还可以带有一个流参数。此时核函数调用将是异步的，就像之前与GPU之间的内存复制操作一样。从技术上来说，当循环迭代完一次时，有可能不会启动任何内存复制或核函数执行。前面提到过，我们能够确保的是，第一次放入流中的复制操作将在第二次复制操作之前执行。此外，第二个复制操作将在核函数启动之前完成，而核函数将在第三次复制操作开始之前完成。正如在本章前面提到的，流就像一个有序的工作队列，GPU从该队列中依次取出工作并执行。

当for()循环结束时，在队列中应该包含了许多等待GPU执行的工作。如果想要确保GPU执行完了计算和内存复制等操作，那么就需要将GPU与主机同步。也就是说，主机在继续执行之前，要首先等待GPU执行完成。可以调用CUDAStreamSynchronize()并指定想要等待的流：

//将计算结果从页锁定内存复制到主机内存

HANDLE_ERROR(udaStreamSynchronize( stream ) );

当程序执行到stream与主机同步之后的代码时，所有的计算和复制操作都已经完成，因此可以停止计时器，收集性能数据，并释放输入缓冲区和输出缓冲区。

HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop))；   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop））;   
printf（“Time taken:  $\text{日} 3 . 1 f$  ms\n”， elapsedTime）;

//释放流和内存

HANDLE_ERROR(udaFreeHost(host_a)）；  
HANDLE_ERROR(udaFreeHost(host_b))；  
HANDLE_ERROR(udaFreeHost(host_c))；  
HANDLE_ERROR(udaFree(dev_a)）；  
HANDLE_ERROR(udaFree(dev_b)）；  
HANDLE_ERROR(udaFree(dev_c)）；

最后，在退出应用程序之前，记得销毁对GPU操作进行排队的流。

HANDLE_ERROR(udaStreamDestroy( stream ) );

return 0;

坦白地说，这个示例并没有充分说明流的强大功能。当然，如果当主机正在执行一些工作时，GPU也正忙于处理填充到流的工作，那么即使使用单个流也有助于应用程序速度的提升。但即使不需要在主机上做太多的工作，我们仍然可以通过使用流来加速应用程序，在下一节中，我们将看到如何实现这个目的。

10.4_使用单个CUDA流

10.4 使用单个CUDA流