【CUDA】Stream and Event

cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaDeviceSynchronize();
// ... ... host计算
cudaMemcpy(..., cudaMemcpyDeviceToHost);

从 device 角度看，所有者三个操作都是使用的默认stream，并且按照代码从上到下的顺序依次执行，device 本身是不知道其他的 host 操作怎样执行的

从 host 角度来看，数据传输都是同步的并且会一直等待，直到操作完成。不过不同于数据传输，kernel 的 launch 是异步的，host 立刻就能重新得到控制权，不用管 kernel 是否执行完毕，从而进行下一步动作。很明显，这种异步行为有助于重叠device和host之间的运算时间

数据传输也是可以异步执行的，使用时必须显示的声明一个 stream 来分派执行

cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);

注意新增加的最后一个参数。在 host issue 了这个函数给 device 执行后，控制权可以立刻返还给 host。上面代码使用了默认 stream，若要创建一个新的 stream 则使用下面的API：

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

注意：使用该函数的一个比较常见的错误，或者说容易引起混乱的地方是，这个函数返回的 error code 可能是上一次调用异步函数产生的

当执行一次异步数据传输时，必须使用pinned（或者non-pageable）memory

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

通过在将该内存 pin 到 host 的虚拟内存上，就可以将该 memory 的物理位置强制分配到CPU内存中以便使之在整个程序生命周期中保持不变。否则的话，操作系统可能会在任意时刻改变该 host 端的虚拟内存对应的物理地址

若异步数据传输函数没有使用 pinned host memory，操作系统就可能将数据从一块物理空间移动到另一块物理空间（因为是异步的，CPU在执行其他的动作就可能影响这块数据），而此时 cuda runtime 正在执行数据的传输，这会导致不确定的行为

执行 kernel 时要想设置 stream 只需加一个stream参数即可

kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);

// 非默认的stream声明
cudaStream_t stream;
// 初始化
cudaStreamCreate(&stream);
// 资源释放
cudaError_t cudaStreamDestroy(cudaStream_t stream);

当执行资源释放的时候，若仍然有 stream 的工作未完成，那么虽然该函数仍然会立刻返回，但相关的工作做完后，这些资源才会自动释放

由于所有 stram 的执行都是异步的，就需要一些API在必要时进行同步：

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

第一个会强制 host 阻塞等待，直至 stream 中所有操作完成为止；第二个会检查 stream 中的操作是否全部完成，即使有操作没完成也不会阻塞 host。若所有操作都完成了，则返回 cudaSuccess，否则返回 cudaErrorNotReady

for (int i = 0; i < nStreams; i++) {
    int offset = i * bytesPerStream;
    cudaMemcpyAsync(&d_a[offset], &a[offset], bytePerStream, streams[i]);
    kernel<<<grid, block, 0, streams[i]>>>(&d_a[offset]);
    cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
}

for (int i = 0; i < nStreams; i++)
    cudaStreamSynchronize(streams[i]);

上述代码中使用了三个stream，数据传输和kernel运算都被分配在了这几个并发的stre