流（Streams）

        应用程序通过流来管理并行。一个流是一个顺次执行的命令序列。不同的流之间并行执行，没有固定的执行顺序。

1、流的创建与销毁

        定义一个流的过程通常包括：创建一个流对象，然后指定它为内核启动或者主机设备间数据传输的流参数。下面的一段代码创建了两个流并且在锁页内存中分配了一块float类型的数组hostPtr：

这部分代码中有一点需要注意：为了并行化数据拷贝和内核执行，主机端内存必须分配为锁页（page-locked）内存。

要销毁一个流需要调用函数cudaStreamDestroy()

1. int *host_a, *host_b, *host_c;  
2. int *dev_a, *dev_b, *dev_c;  
3.   
4. cudaError_t cudaStatus;  
5. cudaStatus = cudaMalloc((void **)&dev_a, N * sizeof(int));  
6. if (cudaStatus != cudaSuccess)  
7. {  
8.     printf("cudaMalloc dev_a failed!\n");  
9. }  
10.   
11. cudaStatus = cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);  
12. if (cudaStatus != cudaSuccess)  
13. {  
14.     printf("cudaHostAlloc host_a failed!\n");  
15. }  
16. for (int i = 0; i < FULL_DATA_SIZE; i += N)  
17. {  
18.     cudaStatus = cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);  
19.     if (cudaStatus != cudaSuccess)  
20.     {  
21.         printf("cudaMemcpyAsync a failed!\n");  
22.     }  
23.   
24.     cudaStatus = cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);  
25.     if (cudaStatus != cudaSuccess)  
26.     {  
27.         printf("cudaMemcpyAsync b failed!\n");  
28.     }  
29.   
30.     kernel << > >(dev_a, dev_b, dev_c);  
31.   
32.     cudaStatus = cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);  
33.     if (cudaStatus != cudaSuccess)  
34.     {  
35.         printf("cudaMemcpyAsync c failed!\n");  
36.     }  
37. }  
38. 这段代码中并没有使用cudaMemcpy()，而是通过cudaMemcpyAsync()在GPU与主机之间复制数据。函数差异虽小，但却很重要。cudaMemcpy()的行为类似于C库函数memcpy()。尤其是，这个函数将以同步方式执行，也就是说，当函数返回时，复制操作已经完成。
    异步函数的行为与同步函数相反，在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动或完成。我们能够保证的是复制操作肯定会在下一个被放入流中的操作启动之前执行。任何传递给cudaMemcpyAsync()的主机内存指针都必须已经通过cudaHostAlloc()分配好内存，也就是，只能以异步方式对页锁定内存进行复制操作。
    注意，在核函数调用的尖括号中有一个流参数stream，此时核函数调用将是异步的。从技术上来说，当循环迭代完一次时，有可能不会启动任何内存复制或核函数执行。但能够确保的是，第一次放入流中的复制操作将在第二次复制操作之前执行，第二个复制操作将在核函数启动之前执行完成。这意味着，代码中for循环的完成不保证流的完成，每个流中的任务都可能处于等待状态。
    5） 当for循环结束时，队列中应该包含了许多等待GPU执行的工作。如果想要确保GPU执行完了计算与内存复制等操作，那么就需要将GPU与主机同步。也就是说，主机在继续执行之前，要首先等待GPU执行完成。可以调用cudaStreamSynchronize()并指定想要等待的流：
    [cpp] view plain copy

    1. cudaStatus = cudaStreamSynchronize(stream);  

    6） 当程序执行到stream与主机同步之后的代码时，所用计算与复制操作都已完成。此时需要释放缓冲区，并销毁对GPU操作进行排队的流：
    [cpp] view plain copy

    1. cudaFreeHost(host_a);  
    2. cudaFreeHost(host_b);  
    3. cudaFreeHost(host_c);  
    4. cudaFree(dev_a);  
    5. cudaFree(dev_b);  
    6. cudaFree(dev_c);  
    7. cudaStreamDestroy(stream);  

    至此，单个流的使用已经讲完。