CUDA中的程序并发可以分为两种
- 内核级并发
- 网格级并发
内核级并发是开发程序中经常使用到的，即通过划分block和thread实现同一个内核在GPU上同时并发，将同一个核分别部署到不同的SP上进行同时运行，上述并发方式为我们经常使用到的方式。
然而在内核级并发中，其实 还是存在同步忙等待现象，比如内存copy时，将数据从CPU copy到GPU或者在GPU中的数据结果 copy到CPU中这些过程都需要进行忙等待，以及有时经常遇到两个不相关的内核函数可以同时运行以提升以最大利用率，此时就需要用到CUDA中的网格级并发，网格级并发运行多个内核同时运行在同一个硬件设备中。
流和事件
CUDA中的网格级并发是通过建立不同的流来实现。CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流通过将这些主机代码的操作进行封装，以保持操作的顺序，并且允许这些操作在流中排队，按照先后顺序执行所有的操作。在相同的流中的操作，按照严格的顺序来执行，而在不同流中的操作相互不影响，所以可以通过建立不同的流，来实现网格级并发。
CUDA中的流分为两种隐式和显式流。
隐式流又称为空流，CUDA在内核其他时会默认创建一个流 为隐式流，如果不创建新的流，则所有的操作将封装到默认流中。默认流的对数据传输采用同步，对内核运行采用异步方式。数据传输API一般为cudaMemcpy等，而kernel运行则使用异步方式，这就意味着CUDA 主机端在提交内核之后会立即返回，主机端可以在等待kernel运行期间 做其他操作，即可以将主机和设备端的计算这一段过程重叠并行运行。
显式流可以调用CUDA API进行显式的创建，为一个异步流，即封装的所有操作都可以是异步的，包括内存copy，例如
cudaMemcpyAsync(void *dst, const void *srcm, size_t count, cudaMemcpyKind kind, cudaStream_t stream=0)上述内存copy异步动作必须显式的设置一个流，调用完后即使数据没有copy完毕，也会立即进行返回。
值得注意的一点是， 使用cudaMemcpyAsync()函数时，使用的内存必须是固定主机内存，可以使用以下函数进行申请：
cudaMallocHost(void **ptr, size_t size);
cudaHostAlloc(void **pHost, size_t size, unsigned int flags);在主机虚拟内存中固定分配，可以确保其在CPU内存中的物理位置在应用程序的整个生命周期中保持不变。否则，操作系统可以随时自由改变主机虚拟内存的物理位置。如果在没有固定主机内存的情况下执行一个异步CUDA操作，操作系统可能会在物理层面上移动数组，而CUDA操作运行时将该数组移动到设备中，这样会导致为定义行为。
CUDA相关API
CUDA中有关流的API如下：
创建一个非空流：
__host__ cudaError_t cudaStreamCreate ( cudaStream_t* pStream )此时创建的流为一个异步流。
销毁一个非空流：
__host__ __device__ cudaError_t cudaStreamDestroy ( cudaStream_t stream )CUDA程序中一般典型的模式为：
1：将输入数据从主机移动到设备上。
2：在设备上执行一个内核。
3：将结果从设备移回主机中。
在默认流中，对数据在主机和设备之间移动是通过cudaMemcpy()函数来实现，而该函数是一个同步堵塞函数。而在一个非空流中为一个异步流，可以将主机和设备之间数据移动也改为一个异步操作，这样就能隐藏该步骤的时间消耗。
异步copy函数API为：
__host__ __device__ cudaError_t cudaMemcpyAsync ( void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0 )下面一个例子说明如何实现网格级并发，多个流同时运行：
for (int i= 0; i< nStreams;i++)
{
int offset = i * bytesPerStream;
cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, streams[i]);
kernel<<grid, block,0,streams[i]>>(&d_a[offset]);
cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream,stream[i]);
}
for(int i =0; i< nStream;i++)
{
cudaStreamSynchronize(streams[i]);
}其中cudaStreamSynchronize()函数为强制堵塞主机，一直等到给定流中所有的操作都完成，函数原型为：
__host__ cudaError_t cudaStreamSynchronize ( cudaStream_t stream )
执行的流过程如下：
与传统的串行方式相比，在性能提升上有很大改进。其中数据传输操作虽然分布在不同的流中，但是由于硬件资源PCIe总线限制，无法真正并行执行，但是可以利用kernel异步执行过程中，来进行下一个流的传输，以此来隐藏传输数据的时间消耗，来提高性能。
重叠GPU和CPU执行
CUDA中还提供了另外一种方法用于判断一个异步流中的所有操作是否完毕cudaStreamQuery(),用于查询所有流中的操作是否已经完成，该函数非堵塞某事，函数原型为：
__host__ cudaError_t cudaStreamQuery ( cudaStream_t stream )当流中的操作都完成时，函数返回cudaSuccess，当流中仍然有操作在执行时返回cudaErrorNotReady。利用cudaStreamQuery()函数可以完成来实现GPU和CPU同时执行。
下面以一个内核实现向量与标量加法例子来说明重叠GPU和CPU执行是如何实现的。
内核kernel的实现加法运算如下：
__global__ void kernel(float * g_data,float value)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
g_data[idx] = g_data[idx] + value;
}使用异步操作完成数据copy，计算等操作，如下：
cudaStream_t Stream;
cudaStreamCreate (&Stream);
cudaMemcpyAsync(d_a,h_a,nbytes,cudaMemcpyHostToDevice, Stream);
kernel<<<grid, block,Stream>>>(d_a, value);
cudaMemcpyAsync(h_a, d_a,nbytes, cudaMemcpyDeviceToHost, Stream);异步等待时刻,使用cudaStreamQuery()不断查询是否流中的操作是否已经完毕：
unsigned long int counter = 0;
while（cudaErrorNotReady == cudaStreamQuery()）
{
counter++
}在循环等待过程中，CPU不断执行counter++操作，以实现GPU和CPU同时运算。
参考资料
《CUDA C编程权威指南》
