CUDA Stream流并发性

目录

1、CUDA 中的异步命令

2、默认流

3、Multistream多流示例

---

异构计算是指高效地使用系统中的所有处理器，包括 CPU 和 GPU 。为此，应用程序必须在多个处理器上并发执行函数。 CUDA 应用程序通过在 streams 中执行异步命令来管理并发性，这些命令是按顺序执行的。不同的流可以并发地执行它们的命令，也可以彼此无序地执行它们的命令。在不指定流的情况下执行异步 CUDA 命令时，运行时使用默认流。在 CUDA 7 之前，默认流是一个特殊流，它隐式地与设备上的所有其他流同步。CUDA7引入了大量强大的新功能，包括一个新的选项，可以为每个主机线程使用独立的默认流，这避免了传统默认流的序列化。本文将展示如何在 CUDA 程序中简化实现内核和数据副本之间的并发。

1、CUDA 中的异步命令

如 CUDA C 编程指南所述，异步命令在设备完成请求的任务之前将控制权返回给调用主机线程（它们是非阻塞的）。这些命令是：

• 内核启动；
• 存储器在两个地址之间复制到同一设备存储器；
• 从主机到设备的 64kb 或更少内存块的内存拷贝；
• 由后缀为 Async 的函数执行的内存复制；
• 内存设置函数调用。

为内核启动或主机设备内存复制指定流是可选的；可以调用 CUDA 命令而不指定流（或通过将 stream 参数设置为零）。下面两行代码都在默认流上启动内核。

  kernel<<< blocks, threads, bytes >>>();    // default stream
  kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0

2、默认流

在并发性对性能不重要的情况下，默认流很有用。在 CUDA 7 之前，每个设备都有一个用于所有主机线程的默认流，这会导致隐式同步。正如 CUDA C 编程指南中的“隐式同步”一节所述，如果主机线程向它们之间的默认流发出任何 CUDA 命令，来自不同流的两个命令就不能并发运行。CUDA 7 引入了一个新选项， 每线程默认流 ，它有两个效果。首先，它为每个主机线程提供自己的默认流。这意味着不同主机线程向默认流发出的命令可以并发运行。其次，这些默认流是常规流。这意味着默认流中的命令可以与非默认流中的命令同时运行。要在 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 。

3、Multistream多流示例

看一个小例子。下面的代码简单地在八个流上启动一个简单内核的八个副本。只为每个网格启动一个线程块，这样就有足够的资源同时运行多个线程块。作为遗留默认流如何导致序列化的示例，在默认流上添加不起作用的虚拟内核启动

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;
}

首先让检查遗留行为，通过不带选项的编译：

nvcc ./stream_test.cu -o stream_legacy

可以在 NVIDIA visualprofiler （ nvvp ）中运行该程序，以获得显示所有流和内核启动的时间轴。图 1 显示了 Macbook Pro 上生成的内核时间线，该 Macbook Pro 带有 NVIDIA GeForce GT 750M （一台开普勒 GPU ）。可以看到默认流上虚拟内核的非常小的条，以及它们如何导致所有其他流序列化。

一个简单的多流示例在将任何交错内核发送到默认流时不会实现并发

现在尝试新的每线程默认流。

nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread

图 2 显示了来自 nvvp 的结果。在这里可以看到九个流之间的完全并发：默认流（在本例中映射到流 14 ）和创建的其他八个流。请注意，虚拟内核运行得如此之快，以至于很难看到在这个图像中默认流上有八个调用。

图 2 ：使用新的每线程默认流选项的多流示例，它支持完全并发执行

在为并发进行编程时，还需要记住以下几点。

• 记住：对于每线程的默认流，每个线程中的默认流的行为与常规流相同，只要同步和并发就可以了。对于传统的默认流，这是不正确的。
• --default-stream 选项是按编译单元应用的，确保将其应用于所有需要它的 nvcc 命令行。
• cudaDeviceSynchronize() 继续同步设备上的所有内容，甚至使用新的每线程默认流选项。如果只想同步单个流，请使用 cudaStreamSynchronize(cudaStream_t stream) ，如的第二个示例所示。
• 从 CUDA 7 开始，还可以使用句柄 cudaStreamPerThread 显式地访问每线程的默认流，也可以使用句柄 cudaStreamLegacy 访问旧的默认流。请注意， cudaStreamLegacy 仍然隐式地与每个线程的默认流同步，如果碰巧在一个程序中混合使用它们。
• 可以通过将 cudaStreamCreate() 标志传递给 cudaStreamCreate() 来创建不与传统默认流同步的非堵塞流 。

参考文献：

CUDA 7 Stream流简化并发性 - 吴建明wujianming - 博客园