GPU 架构与并行计算：从 SIMT 到 GPGPU 编程

GPU 早已不再只是图形渲染的专用处理器。从深度学习训练到科学计算，从加密货币挖矿到实时视频编码，通用 GPU 计算（GPGPU） 正在重塑高性能计算的版图。理解 GPU 的底层架构——SIMT 执行模型、Warp 调度、显存层次——是写出高效并行程序的前提。本文从硬件执行机制出发，对比 Vulkan Compute Shader 与 CUDA 两大编程接口，带你构建完整的 GPGPU 知识体系。

一、SIMT 执行模型：GPU 并行的核心抽象

SIMT（Single Instruction, Multiple Threads，单指令多线程） 是 NVIDIA 提出的执行模型，它是对 SIMD（单指令多数据）的进化。SIMD 中一条指令同时操作多个数据通道；而 SIMT 中，每条指令被多个独立线程执行，每个线程拥有独立的寄存器状态、程序计数器和执行路径。

SIMT 的关键特性在于：

• 线程级并行：每个线程可以独立分支、独立寻址，逻辑上等同于一个独立的执行流。
• 硬件成组调度：虽然逻辑独立，但硬件将线程组织成固定大小的组（Warp / Wavefront）进行调度。
• 隐式同步：Warp 内线程默认按 lock-step 执行同一条指令，分支发散时硬件自动串行化处理路径。

与 CPU 的 SMT（同步多线程）不同，SIMT 不是为了掩盖单核心的指令延迟，而是用海量线程掩盖全局内存延迟。一个现代 GPU 可以同时容纳数万个活跃线程，当某组线程等待显存加载时，调度器立即切换到另一组就绪线程，几乎没有空闲周期。

二、Warp、Wavefront 与 Thread Block

2.1 Warp（NVIDIA）与 Wavefront（AMD）

Warp 是 NVIDIA GPU 调度的基本单位，固定包含 32 个线程。AMD GPU 的对应概念称为 Wavefront，大小为 64 线程。Warp 内所有线程共享同一条指令流，但被分配不同的数据（通过线程 ID 区分）。

Warp 的执行特点：

2.2 Thread Block（线程块 / Work Group）

线程块是程序员显式组织的线程集合，具有以下属性：

• 同块内线程可以同步：通过 __syncthreads()（CUDA）或 barrier()（Vulkan/GLSL）实现。
• 共享内存可见：同一块内线程可以访问高带宽的 Shared Memory（也称为 Local Data Share）。
• 独立执行：不同线程块之间没有任何执行顺序保证，可以并行也可以串行，取决于硬件资源。
• 最大线程数限制：NVIDIA 通常为 1024 线程/块；AMD 为 1024（RDNA）或更大（CDNA）。

// CUDA: 启动一个 Grid，包含多个 Block，每个 Block 包含 256 个线程
__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

// 主机端启动
vecAdd<<< (n + 255) / 256, 256 >>>(d_a, d_b, d_c, n);

三、显存层次：从寄存器到全局显存

GPU 的显存层次设计是其高性能的关键。理解各级延迟、容量和带宽，是优化并行程序的基础。

优化原则：尽量让数据访问停留在靠近计算单元的位置。寄存器最快但容量极小；Shared Memory 是手动管理的"程序员可控缓存"，常用作线程协作的数据交换区；全局显存访问必须高度合并（Coalesced），否则带宽利用率会断崖式下跌。

四、三大性能陷阱：Divergence、Occupancy 与 Coalescing

4.1 Warp Divergence（Warp 分支发散）

当 Warp 内线程进入不同的代码分支时，硬件无法同时执行两条路径，只能串行化执行：先执行 if 分支中活跃的线程，再执行 else 分支。未走某条路径的线程被硬件屏蔽（predicated off），但仍在占用执行槽位。

// 坏示例：Warp 内线程随机分支，导致严重 Divergence
__global__ void badBranch(int* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        if (data[tid] % 2 == 0) {
            data[tid] = data[tid] * 2;   //  Warp 内部分线程走这里
        } else {
            data[tid] = data[tid] + 1;   //  另一部分走这里
        }
    }
}

// 好示例：确保 Warp 级对齐，整个 Warp 同时走同一分支
// 例如让相邻 32 个数据的奇偶性一致，或以 Warp 大小对齐的重排

4.2 Occupancy（占用率）

Occupancy 定义为：活跃 Warp 数 / SM 支持的最大 Warp 数。高 Occupancy 意味着有更多 Warp 可供调度，从而更好地隐藏延迟。影响因素：

• 寄存器用量：每线程使用寄存器越多，SM 能并发的线程越少。
• Shared Memory 用量：每块使用 Shared Memory 越多，SM 上能同时驻留的块越少。
• 块内线程数：块太小（如低于 64）可能导致块级资源分配粒度浪费。

通常目标是将 Occupancy 提升到 50% 以上，但并非越高越好——有时候通过寄存器换算法效率（如展开循环减少指令数），即使 Occupancy 降低，实际性能反而提升。

4.3 Memory Coalescing（内存合并访问）

当 Warp 内线程访问全局显存时，硬件会将多个线程的请求合并成尽可能少的缓存行事务。理想情况下，线程 tid 访问地址 base + tid * sizeof(T)，Warp 的 32 个线程访问一段连续 128 字节，合并为一次 128 字节事务。

如果线程访问地址散乱（如随机索引），则每个线程可能触发独立的缓存行加载，有效带宽暴跌至峰值的 1/10 甚至更低。

// 好的合并访问：连续、对齐
__global__ void coalescedRead(const float* input, float* output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        output[tid] = input[tid];   // tid 连续 → 地址连续 → 完美合并
    }
}

// 坏的跨步访问：高维转置常见陷阱
__global__ void stridedRead(const float* input, float* output, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    // 线程按列增长，但按行读取：跨步 = width * sizeof(float)
    output[row * width + col] = input[col * width + row];
}
// 优化方式：Shared Memory tiling（见矩阵乘法章节）

五、Compute Shader 编程模型

Vulkan 的 Compute Shader 是跨平台的 GPGPU 入口，不依赖 CUDA 生态，可在 Windows、Linux、Android 和 macOS（通过 MoltenVK）上运行。

5.1 执行层级

5.2 Vulkan Compute Shader 示例：向量加法

// shader.comp
#version 450
layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in;

layout(set = 0, binding = 0) readonly buffer InputA { float data[]; } inputA;
layout(set = 0, binding = 1) readonly buffer InputB { float data[]; } inputB;
layout(set = 0, binding = 2) writeonly buffer OutputC { float data[]; } outputC;

void main() {
    uint gid = gl_GlobalInvocationID.x;
    outputC.data[gid] = inputA.data[gid] + inputB.data[gid];
}

//主机端 C++（Vulkan 伪代码，省略错误检查和管线细节）
void dispatchCompute(VkCommandBuffer cmd, uint32_t n) {
    // Compute pipeline 已绑定，DescriptorSet 已更新
    uint32_t groupCountX = (n + 255) / 256;
    vkCmdDispatch(cmd, groupCountX, 1, 1);
}

5.3 Shared Memory + Barrier 的协作模式

// shader.comp - Shared Memory 归约求和（每 Work Group 对 512 个元素求和）
#version 450
layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;

layout(set = 0, binding = 0) readonly  buffer Input  { float data[]; } inputData;
layout(set = 0, binding = 1) writeonly buffer Output { float data[]; } outputData;

shared float sharedData[512];

void main() {
    uint lid = gl_LocalInvocationID.x;
    uint gid = gl_GlobalInvocationID.x;

    sharedData[lid] = inputData.data[gid];
    barrier();                    // 确保所有线程写完成
    memoryBarrierShared();

    // 树形归约
    for (uint stride = 256; stride > 0; stride >>= 1) {
        if (lid < stride) {
            sharedData[lid] += sharedData[lid + stride];
        }
        barrier();
        memoryBarrierShared();
    }

    if (lid == 0) {
        outputData.data[gl_WorkGroupID.x] = sharedData[0];
    }
}

六、GPGPU 经典算法

6.1 并行归约（Reduction）

归约是将数组聚合为单一值（如求和、求最大、求最小）的操作。高效实现在于：

1. 每个 Block 用 Shared Memory 做本地归约，避免全局内存的反复读写。
2. 多轮 Kernel 启动解决跨 Block 依赖，或借助 Atomic 操作 / 二次归约 Kernel。
3. 展开循环（循环展开）减少同步次数和指令开销。

// CUDA 归约 Kernel（简化版，假设线程数等于数组长度）
__global__ void reduceSum(const float* input, float* output, int n) {
    extern __shared__ float sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? input[i] : 0.0f;
    __syncthreads();

    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) output[blockIdx.x] = sdata[0];
}

6.2 并行前缀和（Prefix Sum / Scan）

前缀和计算输出数组 output[i] = input[0] + ... + input[i]。GPU 上的高效实现通常采用 Blelloch Scan（Work-Efficient Scan）：

• Up-Sweep（归约）阶段：树形求和，自底向上。
• Down-Sweep（扫描）阶段：自顶向下传播前缀值。

前缀和是并行算法中的基石操作，广泛应用于流压缩、排序（Radix Sort）、稀疏矩阵运算和光线追踪中的活跃光线压缩。

6.3 矩阵乘法（Tiled Matrix Multiply）

Naive 全局内存矩阵乘会触及显存带宽瓶颈。Tiling 技术通过 Shared Memory 缓存子矩阵，将全局显存访问从 O(N^3) 降至 O(N^3 / TILE_SIZE)：

// CUDA Tiled Matrix Multiply（简化示意，TILE_WIDTH = 16）
__global__ void matMulTiled(const float* A, const float* B, float* C,
                            int N, int TILE_WIDTH) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];

    int row = blockIdx.y * TILE_WIDTH + threadIdx.y;
    int col = blockIdx.x * TILE_WIDTH + threadIdx.x;
    float sum = 0.0f;

    for (int m = 0; m < (N + TILE_WIDTH - 1) / TILE_WIDTH; ++m) {
        As[threadIdx.y][threadIdx.x] =
            (row < N && m * TILE_WIDTH + threadIdx.x < N)
                ? A[row * N + m * TILE_WIDTH + threadIdx.x] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] =
            (col < N && m * TILE_WIDTH + threadIdx.y < N)
                ? B[(m * TILE_WIDTH + threadIdx.y) * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE_WIDTH; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }

    if (row < N && col < N) C[row * N + col] = sum;
}

七、Vulkan Compute vs CUDA 对比

决策建议：如果你的目标平台锁定 NVIDIA 且依赖 cuBLAS/cuDNN 等库，CUDA 是不二之选。如果你需要跨平台部署（如游戏引擎同时面向 PC 和主机），或需要 Compute 与 Graphics 无缝混合，Vulkan Compute Shader 更具战略价值。

八、GPU vs CPU：架构思维的根本差异

核心的思维转换是：CPU 优化 latency，GPU 优化 throughput。不要把 CPU 的递归、复杂分支、指针追逐搬到 GPU 上——那不是它的战场。

九、FAQ

Q1：SIMD 和 SIMT 的本质区别是什么？

SIMD 中数据通道（lane）没有独立状态，一条指令操作向量寄存器上的多个数据；SIMT 中每个线程有独立的寄存器文件和状态，可以独立执行不同指令，只是被硬件打包到 Warp 中调度。SIMT 的编程模型更像多线程程序，而 SIMD 更偏向数据向量处理。

Q2：Warp Divergence 一定会导致性能减半吗？

不一定。如果 Warp 内线程分为两条路径，则两条路径串行执行，最坏情况下接近减半。但如果 Divergence 发生在少数 Warp 中，整体影响有限。优化时应尽量保证 Warp 级对齐：例如让数组维度按 32 对齐，或重排数据使相邻元素具有相同的分支走向。

Q3：Occupancy 100% 就是最好的吗？

不是。100% Occupancy 说明 SM 上驻满了 Warp，但如果 Kernel 本身受限于指令吞吐或全局显存带宽，再增加活跃 Warp 也无济于事。有时候减少寄存器压力和 Shared Memory 使用量，反而能提升 Occupancy 并改善延迟隐藏效果——这是一个需要 profile 指导的权衡。

Q4：Compute Shader 可以替代 CUDA 用于深度学习吗？

理论上可以，但目前不推荐。Vulkan Compute 缺乏成熟的 BLAS 和 DNN 加速库实现，而 cuDNN 和 TensorRT 经过多年优化，性能差距在数量级上。除非你有强烈的跨平台需求且愿意自行实现算子融合，否则训练场景仍以 CUDA 为主。

Q5：为什么 Shared Memory 比 L1 Cache 更快可控？

Shared Memory 和 L1 在物理上通常共享同一块 SRAM，但 Shared Memory 是显式地址管理：程序员精确控制哪些数据存入、何时写入、何时读取，避免了缓存替换策略的不确定性。在矩阵乘等协作算法中，Tiling 到 Shared Memory 是获得接近峰值算力的关键。

Q6：Mobile GPU（如 Mali、Adreno）的 Warp 大小是否与桌面不同？

是的。Arm Mali GPU 的 Warp（称为 Quad / Thread group）通常为 8 或 16 线程；Qualcomm Adreno 使用 64 线程的 Wave。这不仅影响 Divergence 的粒度，也影响你选择的 Block / Work Group 大小。移动端优化应遵循厂商的 best practice guide，通常 Work Group 取 64 或 128 的倍数。

十、总结

GPU 的并行计算威力建立在 SIMT 执行模型、Warp 级调度 和 严格分层的显存体系 之上。要写出高效的 GPGPU 程序，必须时刻关注：

1. Warp Divergence：保证 Warp 内线程走同一路径。
2. Memory Coalescing：让 Warp 的访存尽量合并为连续事务。
3. 显存层次利用：活用 Shared Memory 做 Tile 缓存和线程协作。
4. Occupancy 与寄存器/共享内存的平衡：用 Nsight / RenderDoc 实际 Profile，不要盲信理论。

CUDA 提供了最成熟的生态和调试工具，适合 NVIDIA 独占场景；Vulkan Compute Shader 以高复杂度和跨平台能力为代价，换来了与图形管线的原生互通。两者并不是替代关系，而是针对不同部署策略的工具选择。掌握它们的底层共性——Warp、Barrier、Tiling、Latency Hiding——就能在 GPGPU 的世界中游刃有余。

推荐阅读

• NVIDIA CUDA C Programming Guide（官方文档）
• “Programming Massively Parallel Processors”（David Kirk & Wen-mei Hwu）
• Vulkan Specification - Chapter 31: Compute Shaders