如何在CUDA程序中简化内核和数据副本的并发

异构计算是指高效地使用系统中的所有处理器，包括 CPU 和 GPU 。为此，应用程序必须在多个处理器上并发执行函数。 CUDA 应用程序通过在 streams 中执行异步命令来管理并发性，这些命令是按顺序执行的。不同的流可以并发地执行它们的命令，也可以彼此无序地执行它们的命令。[见帖子[See the post 如何在 CUDA C / C ++中实现数据传输的重叠 ]

在不指定流的情况下执行异步 CUDA 命令时，运行时使用默认流。在 CUDA 7 之前，默认流是一个特殊流，它隐式地与设备上的所有其他流同步。

CUDA 7 引入了大量强大的新功能 ，包括一个新的选项，可以为每个主机线程使用独立的默认流，这避免了传统默认流的序列化。在这篇文章中，我将向您展示如何在 CUDA 程序中简化实现内核和数据副本之间的并发。
 

CUDA 中的异步命令

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

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

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

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

默认流

　　在并发性对性能不重要的情况下，默认流很有用。在 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 。

多流示例

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

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 ：使用新的每线程默认流选项的多流示例，它支持完全并发执行。

多线程示例

让我们看另一个例子，该示例旨在演示新的默认流行为如何使多线程应用程序更容易实现执行并发。下面的例子创建了八个 POSIX 线程，每个线程在默认流上调用我们的内核，然后同步默认流。（我们需要在本例中进行同步，以确保探查器在程序退出之前获得内核开始和结束时间戳。）

#include 
#include 

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

void *launch_kernel(void *dummy)
{
    float *data;
    cudaMalloc(&data, N * sizeof(float));

    kernel<<<1, 64>>>(data, N);

    cudaStreamSynchronize(0);

    return NULL;
}

int main()
{
    const int num_threads = 8;

    pthread_t threads[num_threads];

    for (int i = 0; i < num_threads; i++) {
        if (pthread_create(&threads[i], NULL, launch_kernel, 0)) {
            fprintf(stderr, "Error creating threadn");
            return 1;
        }
    }

    for (int i = 0; i < num_threads; i++) {
        if(pthread_join(threads[i], NULL)) {
            fprintf(stderr, "Error joining threadn");
            return 2;
        }
    }

    cudaDeviceReset();

    return 0;
}

首先，让我们编译时不使用任何选项来测试遗留的默认流行为。

nvcc ./pthread_test.cu -o pthreads_legacy

当我们在 nvvp 中运行它时，我们看到一个流，默认流，所有内核启动都序列化，如图 3 所示。

图 3 ：一个具有遗留默认流行为的多线程示例：所有八个线程都被序列化。

让我们用新的 per-thread default stream 选项编译它。

nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread

图 4 显示，对于每个线程的默认流，每个线程都会自动创建一个新的流，它们不会同步，因此所有八个线程的内核都会并发运行。

图 4 ：每个线程默认流的多线程示例：所有八个线程的内核同时运行。

更多提示

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

　　记住：对于每线程的默认流，每个线程中的默认流的行为与常规流相同，只要同步和并发就可以了。对于传统的默认流，这是不正确的。

　　--default-stream 选项是按编译单元应用的，因此请确保将其应用于所有需要它的 nvcc 命令行。

　　cudaDeviceSynchronize（） 继续同步设备上的所有内容，甚至使用新的每线程默认流选项。如果您只想同步单个流，请使用 cudaStreamSynchronize（cudaStream_t stream） ，如我们的第二个示例所示。

　　从 CUDA 7 开始，您还可以使用句柄 cudaStreamPerThread 显式地访问每线程的默认流，也可以使用句柄 cudaStreamLegacy 访问旧的默认流。请注意， cudaStreamLegacy 仍然隐式地与每个线程的默认流同步，如果您碰巧在一个程序中混合使用它们。

　　您可以通过将 cudaStreamCreate（） 标志传递给 cudaStreamCreate（） 来创建不与传统默认流同步的 非阻塞流 。

　　关于作者

　　Mark Harris 是 NVIDIA 杰出的工程师，致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： GPGPU （图形处理单元上的通用计算）。

　　审核编辑：郭婷