CUDA101 - 03 内存模型与规约优化

2026/03/22

• TODO：
  • Q1：没太听懂warp那里，warp内部是不用原子累加的吗，warp之间是atomadd？
    • A1：明白了，之前是完全串行，现在是warp内部并行（但是这样为什么warp内部不会竞争呢？），warp之间串行，所以速度变成原来的32x了；
    • A2：并非明白，warp内部是只有一个线程在干活，它模拟了另外32个线程也在干活的情况，但是之前的atomicADD的数量是thread的数量（全串行），但现在每一个warp才会有一个atomicADD，它的数量变成了原来的1 / #(warp)，每个warp的计算长度是32（串行），但比原来的要强。
  • Q2：intra-block怎么就利用率更低了？
    • A2：哦，我好像明白了，他之前是一个warp内并行，现在变成一个block内部并行了（是吗？）
• __syncthreads()可能会出现死锁，因此必须每个线程都能跑到对应的位置才可以。
• 用__shfl_down_sync()的逻辑是，warp内部天然同步。
  • 不会出现一个warp内部有的线程已经结束了循环，有的还没有的情况，因为不满足循环条件的thread只会被标记成inactive，但是还会参与这个轮次。
• 这节课讲了下面这几个东西：
  • warp reduce
  • block reduce
  • multi-block reduce（two-pass）
  • inter-block reduce（基于cooperative groups做块间规约）
• __align__是一种对齐约束（表示这块内存的起始地址要满足按 sizeof(T) 字节对齐），unsigned char表示申请一块原始字节数组（一块没有类型的内存buffer）。
  • 所以这句话的意思是：在共享内存上申请一块动态大小的原始字节数组shared_mem_raw，它的内存地址要按照sizeof(T)对齐。

      extern __shared__ __align__(sizeof(T)) unsigned char shared_mem_raw[];
    

  • 为什么不直接写成T呢？
    • 问题1：类型固定，这样写意味着这块shared memory只能存T，但是现实中很多kernel需要一块shared memory，要分成多个区域，存不同类型的数据；
    • 问题2：不好做手动内存布局，比如"[float buffer][int buffer][warp buffer]"就不好做。
  • 所以这样写，先申请一块裸内存（raw memory），再自己控制布局：

      T* smem = reinterpret_cast<T*>(buffer);
    

• 为了能让cooperative kernel运行，需要确定以下几件事：
  • 限制grid size，保证所有 block 能同时驻留在 GPU 上，用cudaOccupancyMaxActiveBlocksPerMultiprocessor;
  • 检查硬件支持，用cudaDevAttrCooperativeLaunch；
  • 用特殊API发射kernel：cudaLaunchCooperativeKernel。
• 课后问题的答案：https://chatgpt.com/s/t_69c9204e85188191a521df0bbfbd139e