GPU 内存架构

GPU 内存架

GPU Memories

Global Memory

  • 一个在GPU上的主要内存。
  • 在GPU上存在大量Global Memory。所有的运行中的线程都可从Global Memory中读取信息。
  • 所有的地址都是Byte Addressable。
  • 读写速度最慢。

Local Mmeory

  • 另一个在GPU上的主要内存。
  • 当没有足够的寄存器时，每个线程会调用Local Memory (被称作Register Spilling)。
  • Addresssable，被用于不被常量indexed数组。
  • 每个线程有独立的 Local Memory。

Caches

  • L1， L2 Cache作用于Global Memory和Local Memory。
  • L1 和 Shared Memory共享地址，并且Configable它们的具体分配。
  • 所有的 Global Memory 读写都通过L2 Cache，包括CPU的读写请求。

Shared Memory

  • 速度仅次于寄存器（registers），在同一个Block内的线程中被共享。

Constant Memory

  • 只读并拥有自己的Cache。

Texture Memory

  • 同样拥有自己的Cache。（Read-Only to GPU， Sets Up by CPU)
  • 有一些额外的关于Addressing的奇技淫巧，方便一些如interpolation的计算。
    Memory Access Speed

Streams Multiprocessors(SM)

Warp 和 Active Block 及 线程数量。

  • Warp是Scheduler执行的基本单元。每个Warp拥有32个线程
  • 一个重要的提高GPU效率的方法为Latency Hidden，即当某些“远程”资源（例如贴图纹理）的读取没有到位时，可通过Context Switch来运行别的线程Warp使GPU占有率保持在一定数值。当资源到位时，可转换会之前的Warp继续运行。
  • 高占有率（High Occupancy）需要足够的Warp数量:
    • 反例1， 每个block里thread数量太少：当一个Block只有一个Warp（32 threads）， 由于每个SM最多支持16个active blocks，所以也只有16个active warp。当前硬件支持64-128个active Warp。
    • 反例2， 每个block里thread数量太多：由于register总数量限制，则active blocks的数量将被迫减少，同业也导致active warp的数量减少。
    • 一般选用 128 - 512 threads per block

__syncthreads()

通过__syncthreads来确保每一段cuda代码所加载的资源完全已经准备好