GPU：合并访存、共享内存与归约

GPU 通过大量并行工作获得吞吐量。一个 kernel 描述许多逻辑线程要执行的程序，但逻辑线程并不意味着每个都有独立物理核心同时运行。理解执行分组、内存层次和同步范围，比只背 API 名字更重要。

grid、block 与 warp

CUDA 使用 grid 包含多个 block，每个 block 包含多个 thread。block 内可以使用共享内存和块级同步，block 之间通常独立调度，不能假设全部同时驻留。NVIDIA 常见 warp 为 32 个线程，同一 warp 的执行路径与活跃 lane 会影响吞吐，但正确同步不应依赖“它们看起来总是一起执行”。

把 N 项映射到一维线程，常见索引为 blockIdx.x * blockDim.x + threadIdx.x。若 N 不是块大小的倍数，末块存在没有对应输入的线程。计算可以让它们使用中性值参与；遇到 block 级 barrier 时，不能让部分线程提前返回而其他线程继续等待。

合并访存到底在比较谁

合并访存关注同一组线程在同一条内存指令上访问的地址。若相邻线程读取相邻 float，地址通常落在少量连续事务内；若相邻线程以很大 stride 读取，可能需要更多事务，浪费传输带宽。对齐、数据类型、架构和缓存都会影响具体事务数量，因此本章不给脱离硬件的固定“几次事务”保证。

常见错误是只检查一个线程自己的循环地址连续，却没有看同一时刻其他 lane 在读哪里。转置中每线程读一行再写一列，可能一边合并、一边分散。共享内存 tile 可帮助在块内重排访问，但还需考虑 bank conflict、同步和占用资源。

块级归约的依赖结构

求和可以先让每个线程取一个输入，再把 block 内的值两两合并。256 项依次经过 128、64、32、16、8、4、2、1 的跨度，做 8 轮。每轮的写入要在下一轮读取前对块内线程可见，因此示例在每轮使用 __syncthreads()。

flowchart LR
  A[256 个输入] --> B[128 个部分和]
  B --> C[64 个部分和]
  C --> D[逐级合并]
  D --> E[每块 1 个部分和]
  E --> F[CPU 或另一个 kernel 合并]

查看流程图文本

flowchart LR
  A[256 个输入] --> B[128 个部分和]
  B --> C[64 个部分和]
  C --> D[逐级合并]
  D --> E[每块 1 个部分和]
  E --> F[CPU 或另一个 kernel 合并]

本站 reduction.cu 固定每块 256 线程，输入 N=100003，故意覆盖非整除尾部。越界线程写零进入共享内存，仍参加所有 barrier。每个 block 的线程 0 写一个独立部分和，最后由 CPU 合并；没有假装 __syncthreads 能同步全 grid。

cpp

// Optional original CUDA example. Not compiled/run on the macOS CPU host.
// 256 threads per block; one input value per thread, padded with zero.
#include <cuda_runtime.h>
#include <algorithm>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#define CUDA_CHECK(call) do { cudaError_t err = (call); if (err != cudaSuccess) { \
    std::fprintf(stderr, "%s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); std::exit(1); } } while (0)

__global__ void block_sum(const float *input, float *partials, int n) {
    __shared__ float scratch[256];
    unsigned lane = threadIdx.x;
    unsigned i = blockIdx.x * blockDim.x + lane;
    scratch[lane] = i < static_cast<unsigned>(n) ? input[i] : 0.0f;
    __syncthreads();
    for (unsigned stride = blockDim.x / 2; stride; stride /= 2) {
        if (lane < stride) scratch[lane] += scratch[lane + stride];
        __syncthreads(); // ALL threads reach each barrier, including inactive lanes.
    }
    if (lane == 0) partials[blockIdx.x] = scratch[0];
}

int main() {
    constexpr int n = 100003; // Tail not a multiple of the block size.
    constexpr int blocks = (n + 255) / 256;
    std::vector<float> input(n), partials(blocks);
    double expected = 0;
    for (int i = 0; i < n; ++i) { input[i] = float(i % 17) / 16.0f; expected += input[i]; }
    float *device_input = nullptr, *device_partials = nullptr;
    CUDA_CHECK(cudaMalloc(reinterpret_cast<void **>(&device_input), input.size() * sizeof(float)));
    CUDA_CHECK(cudaMalloc(reinterpret_cast<void **>(&device_partials), partials.size() * sizeof(float)));
    CUDA_CHECK(cudaMemcpy(device_input, input.data(), input.size() * sizeof(float), cudaMemcpyHostToDevice));
    block_sum<<<blocks, 256>>>(device_input, device_partials, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaMemcpy(partials.data(), device_partials, partials.size() * sizeof(float), cudaMemcpyDeviceToHost));
    double result = 0;
    for (float part : partials) result += part;
    CUDA_CHECK(cudaFree(device_input));
    CUDA_CHECK(cudaFree(device_partials));
    double error = std::fabs(result - expected);
    if (error > 1e-5 * std::max(1.0, std::fabs(expected))) { std::fprintf(stderr, "reduction mismatch\n"); return 1; }
    std::printf("reduction: result=%.9f expected=%.9f error=%.9g\n", result, expected, error);
}

这份实现强调可解释性，不是高性能库替代品。更高效方案可能让一个线程先归约多个输入、用 warp shuffle 降低共享内存流量，再分层合并；这些优化还需正确处理参与掩码、尾部和同步范围。基本规则以 CUDA 12.2.2 编程指南为补充参考，固定此文档版本便于重查。

构建与验证边界

本仓库的 macOS arm64 CPU 环境没有编译或运行这份 CUDA 代码。有支持 NVIDIA CUDA 的环境后，在仓库根目录执行：

make -C labs/parallel gpu
./labs/parallel/reduction

程序检查 CUDA API 返回值、launch 错误和同步错误，再将设备结果与 CPU double 参考比较。输入由 1/16 的倍数组成，使基础例子的求和更易观察；并行浮点加法一般仍可能因顺序变化出现差异，验证时应明确定义绝对/相对误差，而不是要求所有算法逐位相同。

当前代码不计时。未来加性能测量时，要区分 host→device 传输、kernel、device→host 传输和端到端时间；异步 launch 后立即停 CPU 计时器只量到了提交的一部分，不能当作 GPU 完成时间。使用设备事件与必要同步，并记录设备、工具链和输入。

占用率不是最终目标

寄存器与共享内存限制了一个执行单元可同时驻留的线程块。更多驻留工作可帮助隐藏内存等待，但降低每线程资源也可能增加重复计算或寄存器溢出。占用率是诊断指标，最终目标是正确结果的端到端时间。一个高占用率、低数据复用的 kernel 仍可能受带宽限制。

在语言模型中，矩阵乘法、softmax、归一化和注意力的瓶颈不同。融合可以减少中间张量往返设备内存，但会改变寄存器压力、共享内存需求和数值计算顺序。先从 Roofline 判断数据移动，再依据实际 profiler 信息选择优化。

自测：为什么最后一个 block 的越界线程不能直接在第一行 return？

如果其他线程仍会经过块级 barrier，参与线程分歧可能违反同步要求。示例让越界线程提供零并参与同步，保持统一控制流。具体提前退出是否允许要严格遵守使用的同步语义，不能靠偶然运行成功判断。

自测：每个 block 都算完后，block 0 能直接读所有部分和吗？

普通 kernel 内没有因此自动形成全 grid barrier。其他 block 可能尚未执行或写入。常见方法是结束第一个 kernel，再启动第二个归约 kernel；本例通过完成同步后在 CPU 合并。

官方实践对应：CUDA 圆形渲染、CUDA kernels。下一章：实验与实测。

GPU：合并访存、共享内存与归约 ​

grid、block 与 warp ​

合并访存到底在比较谁 ​

块级归约的依赖结构 ​