1. 默认流
当程序对性能要求不高时，可使用默认流。
1.1. cuda7之前的默认流行为
每个设备具有一个默认流，不同主机线程调用默认流时会产生隐式同步。默认流只有在所有常规流运行完成后方可运行。默认流与常规流存在同步。
1.2. cuda7的默认流行为
每个主机线程具有一个默认流。默认流被视为常规流，不会与其他流发生同步。
1.2.1. 如何开启每线程默认流
要在 nvcc 7 及更高版本中启用每线程默认流，您可以在包含 CUDA 头（ cuda.h 或 cuda_runtime.h ）之前，使用 nvcc 命令行选项 CUDA 或 #define 编译 CUDA_API_PER_THREAD_DEFAULT_STREAM 预处理器宏。需要注意的是：当代码由 nvcc 编译时，不能使用 #define CUDA_API_PER_THREAD_DEFAULT_STREAM 在. cu 文件中启用此行为，因为 nvcc 在翻译单元的顶部隐式包含了 cuda_runtime.h 。
下图为开普勒GPU上，添加 --default-stream per-thread 选项到nvcc前后的nvvp运行状态图。未添加的情况下，default流导致常规流阻塞。
const int N = 1 << 20;
__global__ void kernel(float *x, int n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
x[i] = sqrt(pow(3.14159,i));
}
}
int main()
{
const int num_streams = 8;
cudaStream_t streams[num_streams];
float *data[num_streams];
for (int i = 0; i < num_streams; i++) {
cudaStreamCreate(&streams[i]);
cudaMalloc(&data[i], N * sizeof(float));
// launch one worker kernel per stream
kernel<<<1, 64, 0, streams[i]>>>(data[i], N);
// launch a dummy kernel on the default stream
kernel<<<1, 1>>>(0, 0);
}
cudaDeviceReset();
return 0;
}
1.3. 默认流的开发提示
- 记住：对于每线程的默认流，每个线程中的默认流的行为与常规流相同，只要同步和并发就可以了。对于传统的默认流，这是不正确的。
- --default-stream 选项是按编译单元应用的，因此请确保将其应用于所有需要它的 nvcc 命令行。
- cudaDeviceSynchronize() 继续同步设备上的所有内容，甚至使用新的每线程默认流选项。如果您只想同步单个流，请使用 cudaStreamSynchronize(cudaStream_t stream) ，如我们的第二个示例所示。
- 从 CUDA 7 开始，您还可以使用句柄 cudaStreamPerThread 显式地访问每线程的默认流，也可以使用句柄 cudaStreamLegacy 访问旧的默认流。请注意， cudaStreamLegacy 仍然隐式地与每个线程的默认流同步，如果您碰巧在一个程序中混合使用它们。
- 您可以通过将 cudaStreamCreate() 标志传递给 cudaStreamCreate() 来创建不与传统默认流同步的 非阻塞流 。 流的异步
2. 流的异步
2.1. 同步流的方式
阻塞CPU线程直到GPU完成工作。多数情况下严重损害主机线程性能。
2.1.1. 暴力流同步
cudaDeviceSynchronize();
2.1.2. 其他流同步
cudaStreamWaitEvent(event);
cudaStreamSynchronize(stream);
cudaStreamQuery(stream);
cudaEventSynchronize(event);
cudaEventQuery(event);
2.2. 异步流拷贝内存
不会增加访存效率
