流:
应用程序通过流管理上述并发操作。 流是按顺序执行的一系列命令(可能由不同的主机线程发出)。 另一方面,不同的流可以相对于彼此或同时地不按顺序执行它们的命令; 这种行为是不能保证的,因此不应该依赖于正确性(例如,内核间通信是未定义的)
流的创造和销毁:
通过创建流对象并将其指定为内核启动序列和主机< - >设备内存副本的流参数来定义流。 以下代码示例创建两个流,并在pagelocked内存中分配一个浮点数组hostPtr:
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
这些流中的每一个都由以下代码示例定义为从主机到设备的一个内存副本,一个内核启动以及一个从设备到主机的内存副本的序列:
for (int i = 0; i < 2; ++i) {
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
MyKernel << <100, 512, 0, stream[i] >> >
(outputDevPtr + i * size, inputDevPtr + i * size, size);
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
size, cudaMemcpyDeviceToHost, stream[i]);
}
每个流将其输入数组hostPtr的部分复制到设备内存中的数组inputDevPtr,通过调用MyKernel()来处理设备上的inputDevPtr,并将结果outputDevPtr复制回hostPtr的同一部分。 重叠行为根据设备的能力描述在这个示例中流如何重叠。 请注意,hostPtr必须指向页面锁定的主机内存,以便发生任何重叠。
流通过调用cudaStreamDestroy()来释放:
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);
如果设备在调用cudaStreamDestroy()时仍然在流中工作,则函数将立即返回,一旦设备完成了流中的所有工作,与流关联的资源将自动释放。
默认流:
内核启动和主机< - >设备内存副本,不指定任何流参数,或者等同于将流参数设置为零,发送到默认流。 因此他们是按顺序执行的。
对于使用--default-stream per-thread编译标志(或者在包含CUDA头文件(cuda.h和cuda_runtime.h)之前定义了CUDA_API_PER_THREAD_DEFAULT_STREAM宏)的代码,默认流是常规流,每个主机线程 有它自己的默认流
对于使用--default-stream传统编译标志编译的代码,默认流是一个称为NULL流的特殊流,每个设备都有一个用于所有主机线程的NULL流。 NULL流是特殊的,因为它会导致隐式同步,如隐式同步中所述。
对于编译时未指定--default-stream编译标志的代码, - default-stream legacy被认为是默认值
显式同步:
有多种方法可以明确地使流彼此同步:
cudaDeviceSynchronize()等待,直到所有主机线程的所有流中的所有前面的命令都完成;
cudaStreamSynchronize()将流作为参数,并等待,直到给定流中的所有前面的命令都完成为止。 它可用于将主机与特定的流同步,从而允许其他流在设备上继续执行。
cudaStreamWaitEvent()将一个流和一个事件作为参数(请参阅Events来描述事件),并在调用cudaStreamWaitEvent()之后,将所有添加到给定流的命令延迟执行直到给定事件完成。 流可以为0,在这种情况下,在调用cudaStreamWaitEvent()等待事件之后,所有命令添加到任何流;
cudaStreamQuery()为应用程序提供了一种方法来知道流中的所有前面的命令是否完成;
为避免不必要的减速,所有这些同步功能通常最适合用于计时目的,或隔离发生故障的启动或内存拷贝。