Streams

• https://devblogs.nvidia.com/how-overlap-data-transfers-cuda-cc/

• https://github.com/NVIDIA-developer-blog/code-samples/blob/master/series/cuda-cpp/overlap-data-transfers/async.cu

• http://on-demand.gputechconf.com/gtc/2015/presentation/S5530-Stephen-Jones.pdf

• http://on-demand.gputechconf.com/gtc/2015/video/S5530.html

  S4158-cuda-streams-best-practices-common-pitfalls.pdf

Non-default

cudaStream_t stream1;
auto result = cudaStreamCreate(&stream1);
auto result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1);
increment<<<1, N, 0, stream1>>>(d_a);
auto result = cudaStreamDestroy(stream1);

Synchronization with streams

• cudaDeviceSynchroniza(): 阻塞 host，等待所有 device 上的操作完成，通常会极大损害性能
• cudaStreamSynchronize(stream) : 阻塞 host，等待所有特定 stream 上的操作完成
• cudaStreamQuery(stream)：不阻塞，测试特定 stream 上的操作是否全部完成
• cudaEventSynchroniza(event) , cudaEventQuery(event)
• cudaStreamWaitEvent(event)

Overlapping Kernel Execution and Data Transfers

通过 deviceQuery 可查询到 concurrent copy and execution，如图：

需要满足下列条件才能 overlapping

1. "concurrent copy and execution" is capable
2. execution and transfer both occur in different, non-default streams