Quartz 4

GPU的进程线程调度问题

9 min read

GPU 的进程与线程

进程与线程进程/线程通信、同步、死锁、生产者 - 消费者模型等，都是基于传统 CPU 的冯·诺依曼架构和多核并行模型设计的。

而现代 GPU（图形处理器） 的架构、编程模型和执行方式与 CPU 有根本性差异，因此这些问题在 GPU 上的体现方式完全不同。

🔍 一、CPU Vs GPU 核心差异（背景知识）

维度CPUGPU
设计目标低延迟、复杂控制流高吞吐、大规模并行
核心数少（4-64）多（数千 CUDA 核心 / 流处理器）
线程模型重量级线程（OS 线程）轻量级线程（Warp/Wavefront）
内存模型共享虚拟内存，缓存层次复杂分层内存（全局、共享、寄存器等）
执行方式MIMD（多指令多数据）SIMT（单指令多线程）
同步粒度线程/进程级线程块（Block）、网格（Grid）级

✅ 所以，不能简单地将 CPU 的并发模型套用到 GPU 上。

🧩 二、GPU 编程模型简述（以 NVIDIA CUDA 为例）

  • Kernel：运行在 GPU 上的函数，由成千上万个线程并行执行。
  • Thread：最细粒度的执行单元。
  • Block：一组线程（如 256 或 512），共享 Shared Memory 和 Barrier Synchronization。
  • Grid：多个 Block 的集合，可跨 SM（流式多处理器）调度。
__global__ void my_kernel(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 每个线程处理一个元素
    data[idx] *= 2;
}

🚫 三、传统“问题”在 GPU 上的体现与差异

1. 进程间通信（IPC）？

  • ❌ 不适用：GPU 本身不运行“进程”。
  • ✅ 替代形式：
    • CPU 与 GPU 之间通信：通过 PCIe 总线传输数据（cudaMemcpy）
    • GPU 与 GPU 通信（多 GPU）：通过 NVLink 或 PCIe + cudaMemcpyPeer
    • 通信瓶颈：CPU-GPU 数据传输是主要性能瓶颈，远高于 GPU 内部计算。

💡 优化策略：尽量减少主机（Host）与设备（Device）之间的数据拷贝，尽可能在 GPU 上完成整个计算流程。

2. 线程间通信？

  • ❌ 不能像 CPU 线程那样通过共享变量 + 锁通信。
  • ✅ GPU 线程通信方式：
    • 共享内存（Shared Memory）：Block 内线程通过 __shared__ 变量共享数据（类似 CPU 的“共享变量”）。
    • 全局内存（Global Memory）：所有线程可读写，但延迟高。
    • 寄存器和本地内存：线程私有。

⚠️ 但不能使用互斥锁或阻塞队列！因为：

  • 锁会导致Warp 内线程分歧（Divergence）
  • 阻塞操作会挂起整个 Warp（32 个线程），造成资源浪费

3. 线程同步？

  • ✅ 支持，但粒度有限：
    • __syncthreads()：Block 内所有线程同步（屏障）
    • __threadfence()：内存栅栏，保证内存可见性
    • 不支持跨 Block 同步（CUDA 中无全局屏障）

❗ 重要限制：不能在分支中调用 __syncthreads()，否则行为未定义（Warp 分歧）。

__global__ void add(float* a, float* b, float* c) {
    int idx = threadIdx.x;
    float temp = a[idx] + b[idx];
    __syncthreads();  // 确保所有线程完成计算
    c[idx] = temp;
}

✅ 同步用于协调 Shared Memory 的读写顺序，而非实现“生产者 - 消费者”。

4. 生产者 - 消费者模型？

  • ❌ 不能直接实现传统的阻塞队列。
  • ✅ 替代思路：
    • 批处理（Batch Processing）：一次性将所有数据放入 GPU，一次性处理完。
    • 流水线（Pipeline）：多个 Kernel 串联，前一个 Kernel 输出作为下一个输入（无需显式队列）。
    • 异步流（Streams）：使用多个 CUDA Stream 实现重叠计算与传输。
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 重叠：计算和数据传输并行
kernel1<<<grid, block, 0, stream1>>>(d_data1);
kernel2<<<grid, block, 0, stream2>>>(d_data2);
cudaMemcpyAsync(h_result1, d_data1, size, cudaMemcpyDeviceToHost, stream1);

✅ 用 “数据流”和“异步执行” 替代“阻塞队列”。

5. 死锁？

  • ❌ 传统死锁（如锁循环等待）在 GPU Kernel 中几乎不可能发生。
  • ✅ 但可能出现 “逻辑死锁”或“执行停滞”：
    • 无限循环：某个线程陷入 while(1)，导致整个 Warp 卡住。
    • Warp 分歧 + 同步：部分线程执行 __syncthreads()，部分不执行，导致未执行的线程永远等不到同步点。
    • 资源耗尽：如共享内存不足，Kernel 启动失败。

💡 GPU 死锁通常是编程错误导致的执行停滞，而非资源竞争。

🌐 四、现代 GPU 编程中的“通信”与“同步”新范式

传统 CPU 概念GPU 等价或替代方案
进程间通信Host-Device 数据传输（cudaMemcpy）
线程间通信Shared Memory + Global Memory
线程同步__syncthreads()（Block 内）、内存栅栏
阻塞队列批处理 + 异步流（Streams）
生产者 - 消费者流水线 Kernel + 双缓冲（Double Buffering）
死锁预防避免无限循环、Warp 分歧、正确使用同步

🛠 五、实际应用：GPU 中的“生产者 - 消费者”思想（间接体现）

虽然不能直接实现阻塞队列，但可以用批处理和流水线模拟：

场景：实时图像处理

CPU（摄像头） 
    ↓ (生产图像帧)
GPU Kernel 1: 图像预处理 
    ↓ (中间结果)
GPU Kernel 2: 特征提取 
    ↓ (结果)
CPU（显示或存储）
  • 使用 两个 CUDA Stream 和 双缓冲：
    • Stream A 处理 Frame 1
    • Stream B 处理 Frame 2
    • 重叠计算与数据传输
  • 实现类似生产者 - 消费者的流水线效果，但无阻塞、无锁。

🧠 六、总结：GPU 并发的本质

传统 CPU 并发现代 GPU 并发
控制流驱动：复杂逻辑、分支多数据流驱动：简单逻辑、大规模并行
线程少，上下文切换开销大线程多，Warp 调度零开销
共享内存 + 锁同步分层内存 + 显式同步（仅 Block 内）
通信靠 IPC、管道、消息队列通信靠内存读写、异步流、主机交互
易发生死锁、竞态死锁罕见，但需避免执行停滞

✅ 最终结论

  1. GPU 不使用传统的“进程/线程通信”模型，而是基于大规模数据并行和SIMT 执行模型。
  2. 同步仅限于线程块内，跨块同步需通过主机协调。
  3. “生产者 - 消费者”通过批处理、流水线、异步流实现，而非阻塞队列。
  4. 死锁概念弱化，但需避免无限循环和同步错误。
  5. 性能瓶颈在数据传输，而非计算或锁竞争。

💡 编程哲学转变：

  • CPU：“如何协调少量线程”
  • GPU：“如何让数千线程高效并行”

掌握这些差异，才能真正发挥 GPU 的强大算力，写出高性能的并行程序。

Graph View

Backlinks