CUDA内存结构

写在前面

• 本来以为我已经懂了，但是在实践中发现，自己写的代码中同步时延占据了非常大的部分，所以打算再好好整理一下内存模型
• 参考了CUDA编程基础以及该作者的其他文章

GPU结构

• 每个线程在SP（Streaming Processor）中执行，每个SP有自己的寄存器和local memory

• 由多个SP和一块共享内存（Shared Memory）组成一个SM（Streaming Multiprocessor）

• 多个SM和一块全局内存（Global Memory）组成一个GPU

• GPU的内存结构如下图所示：

  graph TD
      A[GPU] -->|多个| B[SM]
      B -->|多个| C[SP（线程处理器）]
      B -->|一块| D[Shared Memory]
      A -->|一块| E[Global Memory]
      C -->|寄存器和local memory| F[Local Memory]

thread, block, grid

• 线程处理器SP对应thread
• 多核处理器SM对应block
• 设备device（一整个GPU）对应grid
• 注意！一般协作的都是块内的线程，也就是一个SM里面的线程
  • 这些线程可以访问同一块共享内存（__shared__）
  • 这些线程可以通过__syncthreads()进行同步
  • 这些线程可以通过atomic函数进行原子操作，比如atomicAdd()
• 不过在CUDA 9.0之后，CUDA支持了整个grid的sync操作：

  grid.sync();

warp

• GPU的每一行由一个控制单元和多个SP组成，这些SP执行的指令是相同的
• 这就是SIMT，这些线程的组合叫做warp
• warp的大小是32个线程
• 严重注意！如果一个warp中的线程执行的指令不一样，那么GPU会将这些线程分开执行，这样会导致性能下降，举例说明：

  __global__ void kernel(int *data) {
      int idx = threadIdx.x;
      if (data[idx] > 0) {  // 部分线程走此分支，另一部分走else
          data[idx] *= 2;
      } else {
          data[idx] += 1;
      }
  }

• ​执行流程：
  • 先执行if (data[idx] > 0)，禁用不满足条件的线程
  • 执行data[idx] *= 2（仅活跃线程）
  • 重新启用所有线程，执行else分支，禁用满足if条件的线程
  • 执行data[idx] += 1（剩余线程）
• 并不会实现if和else的并行！！

TO BE CONTINUED