WEBKT

CUDA 动态并行中的同步机制：cudaDeviceSynchronize, cudaStreamSynchronize, __syncthreads 深度解析

175 0 0 0

为什么需要同步？

CUDA 中的同步机制概览

1. __syncthreads()：线程块内同步

2. cudaDeviceSynchronize()：设备同步

3. cudaStreamSynchronize()：流同步

4. 总结和选择

5. 动态并行与同步

6. 总结

7. 进阶阅读

你好，我是老码农。今天我们来聊聊 CUDA 编程中一个非常重要的概念：同步 (Synchronization)。特别是针对动态并行 (Dynamic Parallelism) 的场景，理解并正确使用同步机制是编写高性能 CUDA 代码的关键。我们将深入探讨 cudaDeviceSynchronize(), cudaStreamSynchronize(), 和 __syncthreads() 这几个核心同步函数，分析它们之间的区别，以及各自适用的场景。

为什么需要同步？

在理解具体的同步函数之前，我们先来思考一个问题：为什么我们需要同步？

CUDA 是一种并行计算平台，它允许我们利用 GPU 的强大计算能力来加速各种应用。GPU 的核心数量众多，可以同时执行大量的线程。这些线程被组织成线程块 (Thread Block)，线程块又被组织成网格 (Grid)。

当我们在 CUDA 中编写程序时，我们通常会将计算任务分解成多个可以并行执行的子任务，然后将这些子任务分配给 GPU 上的线程。然而，由于线程的执行顺序是不确定的，并且不同线程之间可能需要共享数据或依赖关系，如果不进行适当的同步，就会出现数据竞争 (Data Race) 和其他并发问题，最终导致程序结果错误。

总结来说，同步的主要目的是：

  • 保证数据一致性： 确保多个线程对共享数据的访问是有序的，避免数据竞争。
  • 协调线程执行： 控制线程的执行顺序，满足线程之间的依赖关系。
  • 避免资源冲突： 确保线程之间不会同时访问或修改相同的资源，例如显存。

CUDA 中的同步机制概览

CUDA 提供了多种同步机制，它们可以分为两大类：

  1. 设备端同步 (Device-side Synchronization): 发生在 GPU 内部，用于线程块内部或不同线程块之间的同步。
  2. 主机端同步 (Host-side Synchronization): 发生在 CPU (主机) 和 GPU (设备) 之间，用于控制 CPU 和 GPU 之间的交互。

我们今天主要关注的是设备端同步和主机端同步中的几个关键函数：

  • __syncthreads()：线程块内同步，属于设备端同步。用于线程块内部所有线程的同步。
  • cudaDeviceSynchronize()：设备同步，属于主机端同步。用于阻塞 CPU，直到 GPU 上所有任务完成。
  • cudaStreamSynchronize()：流同步，属于主机端同步。用于阻塞 CPU，直到指定流上的所有任务完成。

接下来，我们将详细介绍这三个函数，并结合例子来理解它们的使用方法和适用场景。

1. __syncthreads()：线程块内同步

__syncthreads() 是 CUDA 中最基本的同步原语之一，它用于线程块 (Thread Block) 内部的所有线程之间的同步。调用 __syncthreads() 的线程会等待，直到线程块内的所有其他线程都执行到 __syncthreads() 处。也就是说，当一个线程执行到 __syncthreads() 时，它会停下来等待，直到该线程块内的所有其他线程都执行到 __syncthreads()，然后所有线程才能继续执行后续代码。

作用:

  • 确保线程块内的数据一致性。例如，当一个线程需要读取其他线程写入的数据时，可以使用 __syncthreads() 来确保所有写操作都已经完成。
  • 协调线程块内的线程执行顺序。

使用场景：

  • 共享内存的读写： 在使用共享内存进行线程间通信时，通常需要在读写操作之间插入 __syncthreads()。
  • 归约 (Reduction) 操作： 在进行归约操作时，需要使用 __syncthreads() 来同步中间结果。
  • 分而治之 (Divide and Conquer) 算法： 在并行算法中，可能需要将一个任务分解成多个子任务，并使用 __syncthreads() 来同步子任务的执行。

代码示例：

#include <iostream>
#include <cuda_runtime.h>
 
__global__ void exampleKernel(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
 
    if (idx < n) {
        // 模拟一些计算
        data[idx] = idx * 2;
 
        // 线程块内同步，确保所有线程都完成数据写入
        __syncthreads();
 
        // 读取其他线程写入的数据
        int sum = 0;
        for (int i = 0; i < blockDim.x; ++i) {
            sum += data[blockIdx.x * blockDim.x + i];
        }
 
        // 只有第一个线程打印结果
        if (threadIdx.x == 0) {
            printf("Block %d: Sum = %d\n", blockIdx.x, sum);
        }
    }
}
 
int main() {
    int n = 16;
    int *h_data, *d_data;
 
    // 分配主机内存
    h_data = (int *)malloc(n * sizeof(int));
 
    // 分配设备内存
    cudaMalloc((void **)&d_data, n * sizeof(int));
 
    // 初始化主机数据
    for (int i = 0; i < n; ++i) {
        h_data[i] = 0;
    }
 
    // 设置 grid 和 block 的维度
    dim3 blockDim(4);
    dim3 gridDim((n + blockDim.x - 1) / blockDim.x);
 
    // 启动 kernel
    exampleKernel<<<gridDim, blockDim>>>(d_data, n);
 
    // 将设备数据拷贝回主机
    cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
 
    // 打印结果
    for (int i = 0; i < n; ++i) {
        printf("h_data[%d] = %d\n", i, h_data[i]);
    }
 
    // 释放内存
    free(h_data);
    cudaFree(d_data);
 
    return 0;
}

代码解释：

  1. 数据初始化: 在主机端分配和初始化数据，包括主机端数组 h_data 和设备端数组 d_data。
  2. Kernel 调用: 调用 exampleKernel，设置 gridDim 和 blockDim。gridDim 定义了启动的线程块数量，blockDim 定义了每个线程块中的线程数量。在本例中，blockDim 为 4，这意味着每个线程块有 4 个线程，gridDim 根据数据大小计算得到，保证所有数据都被处理。
  3. Kernel 逻辑:
    • 每个线程计算自己的索引 idx。
    • 在 if (idx < n) 内部，每个线程写入数据 data[idx] = idx * 2;。
    • __syncthreads(): 关键在于 __syncthreads() 的使用，它确保在所有线程读取数据之前，所有线程都完成了写入操作。__syncthreads() 使得线程块内的所有线程同步，确保了数据的一致性。
    • 每个线程读取线程块内其他线程写入的数据，计算 sum。
    • 只有 threadIdx.x == 0 的线程打印 sum 的结果。
  4. 数据拷贝与释放: 将设备端的数据拷贝回主机端，并释放主机和设备的内存。

注意事项：

  • __syncthreads() 只能在同一个线程块 (Thread Block) 内使用。不能跨线程块同步。
  • __syncthreads() 会导致性能开销。过度使用会导致性能下降，因此需要谨慎使用。
  • 线程块内的所有线程必须同时执行到 __syncthreads()，否则会导致程序死锁 (Deadlock)。
  • __syncthreads() 是一个 barrier，所有线程必须到达这个点才能继续执行。

2. cudaDeviceSynchronize()：设备同步

cudaDeviceSynchronize() 是一个主机端 (Host-side) 的同步函数，它的作用是阻塞 CPU (主机)，直到 GPU 上所有的 CUDA 任务都完成。这些 CUDA 任务包括所有流 (Stream) 中的所有 kernel 调用、内存拷贝操作等。换句话说，cudaDeviceSynchronize() 会等待 GPU 完成它所接收到的所有工作，才会继续执行 CPU 上的代码。

作用：

  • 确保主机和设备之间的同步，保证 GPU 上的操作已经完成，CPU 才能访问 GPU 上的数据或者进行下一步操作。
  • 可以用来测量 CUDA kernel 的执行时间。在 kernel 调用前后分别调用 cudaDeviceSynchronize()，然后计算时间差。

使用场景：

  • 需要在 CPU 上访问 GPU 计算结果时： 在从 GPU 拷贝数据回 CPU 之前，需要调用 cudaDeviceSynchronize() 来确保 GPU 已经完成计算。
  • 需要在 GPU 上执行一系列依赖任务时： 确保前面的任务完成后，再启动后面的任务。
  • 测量 kernel 执行时间： 在 kernel 调用前后使用 cudaDeviceSynchronize()。

代码示例：

#include <iostream>
#include <cuda_runtime.h>
#include <chrono>
 
__global__ void simpleKernel(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = idx * 2;
    }
}
 
int main() {
    int n = 1024 * 1024;
    int *h_data, *d_data;
 
    // 分配主机内存
    h_data = (int *)malloc(n * sizeof(int));
 
    // 分配设备内存
    cudaMalloc((void **)&d_data, n * sizeof(int));
 
    // 设置 grid 和 block 的维度
    dim3 blockDim(256);
    dim3 gridDim((n + blockDim.x - 1) / blockDim.x);
 
    // 启动 kernel
    auto start = std::chrono::high_resolution_clock::now();
    simpleKernel<<<gridDim, blockDim>>>(d_data, n);
 
    // 同步设备，等待 kernel 完成
    cudaDeviceSynchronize();
    auto end = std::chrono::high_resolution_clock::now();
 
    // 计算 kernel 执行时间
    auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - start);
    std::cout << "Kernel execution time: " << duration.count() / 1000.0 << " ms" << std::endl;
 
    // 将设备数据拷贝回主机
    cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
 
    // 释放内存
    free(h_data);
    cudaFree(d_data);
 
    return 0;
}

代码解释：

  1. 时间测量： 在启动 kernel 前后使用 std::chrono 库来记录时间。
  2. Kernel 调用: 调用 simpleKernel，然后立即执行 cudaDeviceSynchronize()。
  3. cudaDeviceSynchronize(): 它会阻塞 CPU，直到 GPU 上所有的 CUDA 任务（在本例中，就是 simpleKernel 的执行）都完成。这确保了在测量 kernel 执行时间时，不会受到 CPU 并发执行的影响。
  4. 时间计算: 计算 kernel 执行时间，并打印出来。
  5. 数据拷贝与释放: 在 cudaDeviceSynchronize() 之后，可以安全地将设备端的数据拷贝回主机端，并释放内存。

注意事项：

  • cudaDeviceSynchronize() 会导致 CPU 阻塞，降低程序并发性。在不需要同步的情况下，应尽量避免使用它，以免影响程序性能。
  • cudaDeviceSynchronize() 会等待 GPU 上所有任务完成，包括所有流 (Stream) 上的任务。如果你的 CUDA 程序使用了多个流，并且这些流之间没有依赖关系，那么使用 cudaStreamSynchronize() 可能会更有效率。

3. cudaStreamSynchronize()：流同步

cudaStreamSynchronize() 也是一个主机端 (Host-side) 的同步函数，但它的作用范围比 cudaDeviceSynchronize() 更小。cudaStreamSynchronize() 阻塞 CPU，直到指定流 (Stream) 上的所有 CUDA 任务都完成。CUDA 流是一种抽象，用于表示 GPU 上的一系列操作。你可以将 kernel 调用、内存拷贝等操作添加到特定的流中。默认情况下，CUDA 使用一个默认流 (也称为 NULL 流)。

作用：

  • 确保指定流上的操作已经完成，CPU 才能访问流中的数据或者进行下一步操作。
  • 可以用于管理多个流之间的依赖关系。

使用场景：

  • 多流并行： 在使用多个流来提高程序并发性时，需要使用 cudaStreamSynchronize() 来同步不同流之间的操作。例如，一个流用于计算，另一个流用于数据传输，你可以使用 cudaStreamSynchronize() 来确保数据传输完成后，再使用计算结果。
  • 管理流之间的依赖关系： 确保一个流上的操作完成后，才能启动另一个流上的操作。
  • 测量特定流的执行时间： 类似于 cudaDeviceSynchronize()，你可以在流上启动 kernel 前后分别调用 cudaStreamSynchronize() 来测量该流的执行时间。

代码示例：

#include <iostream>
#include <cuda_runtime.h>
#include <chrono>
 
__global__ void kernel1(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = idx * 1;
    }
}
 
__global__ void kernel2(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2;
    }
}
 
int main() {
    int n = 1024 * 1024;
    int *h_data, *d_data;
    cudaStream_t stream1, stream2;
 
    // 分配主机内存
    h_data = (int *)malloc(n * sizeof(int));
 
    // 分配设备内存
    cudaMalloc((void **)&d_data, n * sizeof(int));
 
    // 创建流
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
 
    // 设置 grid 和 block 的维度
    dim3 blockDim(256);
    dim3 gridDim((n + blockDim.x - 1) / blockDim.x);
 
    // 启动 kernel1 在 stream1 上
    kernel1<<<gridDim, blockDim, 0, stream1>>>(d_data, n);
 
    // 启动 kernel2 在 stream2 上
    kernel2<<<gridDim, blockDim, 0, stream2>>>(d_data, n);
 
    // 将数据从设备拷贝回主机，使用 stream1
    cudaMemcpyAsync(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost, stream1);
 
    // 同步 stream1
    cudaStreamSynchronize(stream1);
 
    // 打印结果
    std::cout << "First 10 elements of the result after stream1 sync:" << std::endl;
    for (int i = 0; i < 10; ++i) {
        std::cout << h_data[i] << " ";
    }
    std::cout << std::endl;
 
    // 同步 stream2, 确保 kernel2 完成
    cudaStreamSynchronize(stream2);
 
    // 将数据从设备拷贝回主机，使用 stream2
    cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
 
    // 打印结果
    std::cout << "First 10 elements of the result after stream2 sync:" << std::endl;
    for (int i = 0; i < 10; ++i) {
        std::cout << h_data[i] << " ";
    }
    std::cout << std::endl;
 
    // 释放资源
    cudaStreamDestroy(stream1);
    cudaStreamDestroy(stream2);
    free(h_data);
    cudaFree(d_data);
 
    return 0;
}

代码解释：

  1. 流的创建： 使用 cudaStreamCreate() 创建了两个流 stream1 和 stream2。
  2. Kernel 调用： 分别使用 kernel1 和 kernel2，并在调用时指定了不同的流。kernel1 在 stream1 上执行，kernel2 在 stream2 上执行。注意 Kernel 调用的第三个参数是共享内存大小（设为 0 表示使用默认值），第四个参数指定了流。
  3. 异步拷贝： 使用 cudaMemcpyAsync() 将数据从设备拷贝回主机，并且指定了使用 stream1。这意味着拷贝操作也会被添加到 stream1 中。
  4. cudaStreamSynchronize(): 调用 cudaStreamSynchronize(stream1) 来阻塞 CPU，直到 stream1 上所有的操作 (包括 kernel1 的执行和数据拷贝) 都完成。这样可以确保在访问数据之前，kernel1 已经完成了计算。
  5. 打印结果: 打印结果。
  6. 同步 stream2: 使用 cudaStreamSynchronize(stream2) 来阻塞 CPU，直到 stream2 上所有的操作 (包括 kernel2 的执行) 都完成。
  7. 数据拷贝: 将设备端的数据拷贝回主机端。
  8. 打印结果: 打印结果。
  9. 资源释放: 销毁流，释放内存。

注意事项：

  • cudaStreamSynchronize() 只会阻塞 CPU，直到指定流上的所有操作完成。对于其他流上的操作，它不会产生影响。
  • 使用多个流可以提高程序的并发性，但也会增加程序的复杂性。需要仔细管理流之间的依赖关系，避免数据竞争和死锁。
  • 流之间可以通过 cudaStreamWaitEvent() 和 cudaEventRecord() 来实现同步，创建流之间的依赖关系。

4. 总结和选择

同步函数 作用范围 阻塞对象 适用场景 影响并发性 复杂程度
__syncthreads() 线程块内部 线程 共享内存读写，归约操作，分而治之算法 无 低
cudaDeviceSynchronize() 设备所有任务 CPU 在主机端访问 GPU 计算结果，测量 kernel 执行时间，确保依赖任务顺序 高 低
cudaStreamSynchronize() 指定流上的所有任务 CPU 多流并行，管理流之间的依赖关系，测量特定流的执行时间 中 中

如何选择？

  • 线程块内同步： 如果需要在线程块内部进行同步，例如访问共享内存，那么使用 __syncthreads()。
  • 全局设备同步： 如果需要在主机端等待 GPU 上所有的任务完成，例如在拷贝数据回 CPU 之前，使用 cudaDeviceSynchronize()。
  • 流同步： 如果你的 CUDA 程序使用了多个流，需要同步特定的流，或者管理流之间的依赖关系，那么使用 cudaStreamSynchronize()。在可能的情况下，优先使用 cudaStreamSynchronize()，因为它对整体并发性的影响比 cudaDeviceSynchronize() 要小。

5. 动态并行与同步

CUDA 动态并行 (Dynamic Parallelism) 允许 GPU 上的 kernel 启动其他 kernel。这为编写更灵活和更复杂的 GPU 程序提供了可能。在使用动态并行时，同步机制变得尤为重要，因为你需要确保子 kernel 在正确的时间执行，并且能够访问父 kernel 生成的数据。

动态并行场景下的同步考虑：

  • 子 kernel 的启动依赖： 父 kernel 启动子 kernel 的时候，可能需要确保父 kernel 已经完成某些计算，或者已经将数据准备好。这时，可以使用 cudaDeviceSynchronize() 或 cudaStreamSynchronize() 来同步。
  • 共享数据： 父 kernel 和子 kernel 之间可能需要共享数据。在这种情况下，需要使用 cudaDeviceSynchronize() 或 cudaStreamSynchronize() 来确保数据一致性。在某些情况下，可以使用统一寻址 (Unified Addressing) 来简化数据共享。
  • 嵌套同步： 子 kernel 内部可能也需要使用 __syncthreads() 进行线程块内同步。这种嵌套的同步机制需要仔细设计，避免死锁。__syncthreads() 只能在同一个线程块内使用，不能跨线程块或者 kernel。cudaDeviceSynchronize() 和 cudaStreamSynchronize() 可以用于跨 kernel 的同步。

动态并行示例 (简略)：

#include <iostream>
#include <cuda_runtime.h>
 
__global__ void childKernel(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2;
    }
}
 
__global__ void parentKernel(int *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        // 启动子 kernel
        childKernel<<<1, 256>>>(data, n);
    }
}
 
int main() {
    int n = 1024;
    int *h_data, *d_data;
 
    // 分配主机内存
    h_data = (int *)malloc(n * sizeof(int));
 
    // 分配设备内存
    cudaMalloc((void **)&d_data, n * sizeof(int));
 
    // 初始化数据
    for (int i = 0; i < n; ++i) {
        h_data[i] = i + 1;
    }
 
    // 将数据拷贝到设备
    cudaMemcpy(d_data, h_data, n * sizeof(int), cudaMemcpyHostToDevice);
 
    // 启动父 kernel
    parentKernel<<<1, 256>>>(d_data, n);
 
    // 同步，等待所有 kernel 完成
    cudaDeviceSynchronize();
 
    // 将数据拷贝回主机
    cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
 
    // 打印结果
    for (int i = 0; i < 10; ++i) {
        std::cout << h_data[i] << " ";
    }
    std::cout << std::endl;
 
    // 释放内存
    free(h_data);
    cudaFree(d_data);
 
    return 0;
}

代码解释：

  1. 父 kernel 启动子 kernel: parentKernel 启动了 childKernel。这展示了动态并行的基本用法。
  2. 同步： cudaDeviceSynchronize() 用于确保父 kernel 和子 kernel 都执行完毕，然后再将结果拷贝回主机。如果去掉 cudaDeviceSynchronize()，结果可能会不正确，因为 CPU 可能会在子 kernel 还没执行完的时候就去读取数据。

注意： 这个例子非常简单，只展示了动态并行的基本流程。在实际的 CUDA 程序中，动态并行的使用场景会更加复杂，需要根据具体的应用场景来选择合适的同步机制。

6. 总结

CUDA 同步是编写高效、正确的 GPU 程序的关键。本文介绍了三个核心的同步函数：__syncthreads(), cudaDeviceSynchronize(), 和 cudaStreamSynchronize()。我们详细分析了它们的作用、使用场景和注意事项，并通过代码示例进行了演示。特别地，我们还讨论了在动态并行场景下如何使用同步机制。

记住，选择正确的同步机制取决于你的具体需求。__syncthreads() 用于线程块内同步，cudaDeviceSynchronize() 用于主机和设备之间的全局同步，而 cudaStreamSynchronize() 提供了更细粒度的流同步。理解这些函数的区别，并根据实际情况选择合适的同步方式，可以帮助你编写出更高效、更可靠的 CUDA 程序。

希望这篇文章对你有所帮助! 如果你还有其他问题，欢迎在评论区留言。

7. 进阶阅读

为了更深入地理解 CUDA 同步，建议你继续学习以下内容：

  • CUDA Streams (CUDA 流): 深入了解 CUDA 流的概念和用法。
  • CUDA Events (CUDA 事件): 学习如何使用 CUDA 事件来实现更灵活的同步机制。
  • CUDA Graphs (CUDA 图): 了解 CUDA 图的概念，它提供了一种更高级的方式来管理 CUDA 任务的依赖关系。
  • 统一寻址 (Unified Addressing): 了解统一寻址的概念，它可以简化父 kernel 和子 kernel 之间的数据共享。
  • CUDA 编程最佳实践： 学习 CUDA 编程的最佳实践，例如如何避免死锁，如何优化性能。

祝你在 CUDA 编程的道路上越走越远！

老码农 CUDAGPU并行计算同步__syncthreadscudaDeviceSynchronizecudaStreamSynchronize

评论点评

打赏赞助
sponsor

感谢您的支持让我们更好的前行

分享

QRcode

https://www.webkt.com/article/8113