深入GPU编程：从硬件架构到内核优化

第一部分：揭秘GPU硬件与执行模型
要编写高效的GPU程序，首先必须理解其运行的舞台——GPU硬件本身，以及将软件指令映射到硬件上的执行模型。这部分将建立起连接抽象编程与物理现实的桥梁。

1.1 从宏观到微观：GPU的组织结构
从宏观上看，一块GPU并非一个单一的处理器，而是一个由众多更小的处理单元组成的并行计算巨舰。这些核心处理单元被称为流式多处理器（Streaming Multiprocessors, SMs）。以NVIDIA H100 GPU为例，它集成了多达144个SM。

每个SM可以被看作一个功能完备的“迷你处理器”，是GPU计算能力的基本单位。SM内部又包含了更微观的组件：

处理核心（Processing Cores）：例如，H100的每个SM拥有128个FP32（32位浮点）核心，这些是执行具体数学运算（如加法、乘法）的物理单元。
调度器（Schedulers）：比如Warp Scheduler，负责管理和调度指令的执行。
寄存器文件（Register File）：高速的片上存储，供线程存放临时变量。
片上内存（On-chip Memory）：包括L1缓存和程序员可控的共享内存（Shared Memory）。
这种架构设计体现了GPU与CPU之间根本性的哲学差异。CPU的设计目标是低延迟（Latency-Oriented），它将大量的芯片面积用于复杂的控制逻辑、分支预测单元和巨大的缓存，以确保单个任务能以最快的速度完成。而GPU的设计目标是高吞吐（Throughput-Oriented），它将绝大部分芯片面积用于构建成百上千个简单的算术逻辑单元（ALU，即处理核心），同时配备相对简单的控制逻辑和较小的缓存。这种设计使得GPU能够同时处理海量的并行任务，通过并行计算的规模来弥补单个任务的执行延迟，从而在整体上实现惊人的计算吞吐量。

1.2 CUDA编程模型的抽象
为了驾驭GPU强大的并行能力，NVIDIA推出了CUDA（Compute Unified Device Architecture）平台。CUDA提供了一个编程模型，它将复杂的硬件抽象为程序员易于理解的层次结构：Thread（线程）、Block（线程块）和Grid（网格）。

线程（Thread）：是最基本的执行单位。程序员编写的kernel（内核函数）代码实际上是从单个线程的视角来描述的。每个线程都执行同样的代码，但通过其唯一的ID（threadIdx）来处理不同的数据。
线程块（Thread Block）：由一组线程组成（在现代GPU上最多可达1024个线程）。一个线程块中的所有线程会被调度到同一个SM上执行。块内的线程可以通过高速的**共享内存（Shared Memory）**进行数据交换和协作，并且可以通过__syncthreads()指令进行同步。
网格（Grid）：由一次内核启动所创建的所有线程块组成。一个网格内的所有线程块是相互独立的，它们可以被GPU的全局调度器以任意顺序、并行地分发到GPU上所有可用的SM上执行。
这种抽象模型与硬件的映射关系非常清晰：当程序员启动一个内核时，定义了一个Grid。GPU的硬件调度器接收这个Grid，并将其中的Thread Block分配给各个SM。一旦一个Block被分配给一个SM，它就会在该SM上执行直至完成，不会中途迁移到其他SM。

让我们通过一个经典的“向量加法”例子来具体感受这个模型。

// CUDA C++ Kernel for Vector Addition
// __global__ 关键字表示这是一个在GPU上执行的内核函数
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
// 计算当前线程的全局唯一ID
// blockIdx.x: 当前线程块在网格中的X维度索引
// blockDim.x: 每个线程块在X维度的大小（线程数量）
// threadIdx.x: 当前线程在线程块中的X维度索引
int i = blockIdx.x * blockDim.x + threadIdx.x;

// 防止越界访问
if (i < numElements)
{
C[i] = A[i] + B[i];
}
}

int main()
{
const int numElements = 50000;
const size_t size = numElements * sizeof(float);

// 1. 在主机（CPU）上分配内存
float *h_A = (float *)malloc(size);
float *h_B = (float *)malloc(size);
float *h_C = (float *)malloc(size);

// 初始化主机上的数据…

// 2. 在设备（GPU）上分配内存
float *d_A = NULL;
cudaMalloc((void **)&d_A, size);
float *d_B = NULL;
cudaMalloc((void **)&d_B, size);
float *d_C = NULL;
cudaMalloc((void **)&d_C, size);

// 3. 将数据从主机拷贝到设备
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

// 4. 定义内核启动配置
int threadsPerBlock = 256;
// 计算需要的线程块数量，确保覆盖所有元素
int blocksPerGrid = (numElements + threadsPerBlock – 1) / threadsPerBlock;

// 5. 启动内核 <<<网格大小, 线程块大小>>>
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

// 6. 将结果从设备拷贝回主机
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// 7. 释放设备和主机内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);

return 0;
}

在这个例子中，我们定义了一个vectorAdd内核。每个线程通过blockIdx.x * blockDim.x + threadIdx.x计算出一个全局唯一的索引i，并负责计算结果向量C中第i个元素的和。在主函数main中，我们遵循了典型的CUDA程序流程：在主机（CPU）和设备（GPU）上分配内存 (malloc vs. cudaMalloc)，通过cudaMemcpy在两者之间传输数据，使用<<<…>>>语法配置网格和线程块大小并启动内核，最后将结果传回并释放所有资源。

1.3 执行的核心：Warp
虽然CUDA编程模型以Thread为基本单位，但在硬件层面，真正的调度和执行单位是Warp（线程束）。一个Warp由32个连续的线程组成。当一个Thread Block被分配给一个SM后，SM会将其中的线程划分为多个Warp。例如，一个包含256个线程的Block会被划分为256/32=8个Warp。

所有Warp的执行都遵循SIMT（Single Instruction, Multiple Threads，单指令多线程）模型。这意味着在一个时钟周期内，一个Warp中的所有32个线程同时执行同一条指令，但每个线程都在自己的私有数据（例如，自己的寄存器）上进行操作。Thread是程序员的逻辑视角，而Warp是硬件的物理现实。理解这一点是从入门到精通的关键一步，因为几乎所有的性能优化都与Warp的行为密切相关。

GPU实现高吞吐的秘诀在于延迟隐藏（Latency Hiding）。当一个Warp执行一条指令，比如从全局内存读取数据时，数据返回需要数百个时钟周期，此时该Warp会进入“等待”或“停滞”（stall）状态。如果此时无事可做，宝贵的计算核心就会被闲置。然而，SM的Warp Scheduler（线程束调度器）被设计用来解决这个问题。在一个SM上，通常会同时驻留多个Warp。当一个Warp停滞时，调度器会立即、几乎零开销地切换到另一个已经准备就绪（其所需数据或指令已到位）的Warp并执行它的指令。通过在大量等待的Warp之间快速切换，调度器确保了计算核心总有活干，从而将内存访问等长延迟操作的负面影响“隐藏”起来，维持了极高的计算单元利用率。

这种设计解释了为什么我们通常需要在GPU上启动远超其物理核心数量的线程——目的就是为了给每个SM提供足够多的Warp，形成一个庞大的“待选池”，让Warp Scheduler总有选择的余地。为了最大化资源利用率，我们应尽量使线程块的大小成为Warp大小（32）的整数倍。如果线程块大小不是32的倍数，最后一个Warp将是部分填充的，但它仍然会占用一个完整Warp的硬件资源（如调度槽位），造成浪费。

第二部分：驾驭GPU内存层次结构
如果说Warp是GPU执行的心跳，那么内存访问就是其性能的命脉。GPU的计算速度远超其从主内存（全局内存）获取数据的速度。因此，能否高效地管理和利用GPU的多层级内存，是决定一个CUDA内核性能好坏的最关键因素。

2.1 内存金字塔
GPU的内存系统是一个金字塔结构的层次体系，越顶层的内存速度越快、但容量越小、访问范围也越局限。

寄存器（Registers）：位于金字塔顶端，是GPU上最快的内存。每个线程都拥有一组私有的寄存器，用于存储局部变量和计算的中间结果。寄存器的访问延迟极低，几乎可以瞬时完成。其分配由编译器自动管理，程序员无法直接通过地址访问。但需要注意的是，每个SM的寄存器文件总大小是固定的，如果单个线程使用的寄存器过多，就会限制能够同时驻留在该SM上的线程（和Warp）数量，从而影响Occupancy（占用率）。
共享内存/L1缓存（Shared Memory / L1 Cache）：位于SM内部的片上高速缓存。它的延迟远低于全局内存，带宽高。共享内存的作用域是线程块，块内的所有线程都可以访问同一块共享内存，这使其成为线程间高效通信和协作的理想工具。与寄存器不同，共享内存是程序员可编程的缓存，需要通过__shared__关键字在内核中显式声明和管理。
全局内存（Global Memory）：即GPU显存（DRAM），是容量最大（通常为GB级别）但速度最慢的内存。它对所有Grid中的所有线程以及主机CPU都可见，是主机与设备之间数据交换的主要场所。所有输入数据和最终输出结果通常都存放在这里。对全局内存的访问是性能优化的核心和难点。
只读内存（Read-Only Memory）：包括常量内存（Constant Memory）和纹理内存（Texture Memory）。它们也驻留在设备DRAM中，但在SM上有专门的只读缓存。常量内存适合广播（将同一个值高效地发送给Warp中的所有线程）场景，而纹理内存则为具有空间局部性的访问模式（如图像处理中的2D访问）提供了优化。
下表总结了主要内存类型的特性，以便于比较和决策。

特性 (Characteristic) 寄存器 (Registers) 共享内存 (Shared Memory) 全局内存 (Global Memory)
作用域 (Scope) 线程私有 (Per-thread) 块内共享 (Per-block) 全局/应用可见 (Per-grid/application)
延迟 (Latency) 极低 (~1 cycle) 低 (~tens of cycles) 高 (~hundreds of cycles)
带宽 (Bandwidth) 极高 (TB/s) 高 (TB/s) 中等 (~1.5 TB/s on A100)
容量 (Size) 小 (KB per SM) 小 (KB per SM) 大 (GB per device)
程序员控制 (Programmer Control) 间接 (编译器管理) 直接 (__shared__) 直接 (cudaMalloc)
主要用途 (Primary Use Case) 局部变量、临时值 线程间通信、可编程缓存 主要数据存储
2.2 性能命脉：全局内存合并访问
全局内存访问是CUDA程序中最常见的性能瓶颈。虽然GPU的内存总线很宽，但只有当访问模式正确时，才能充分利用其带宽。这个正确的模式就是内存合并（Memory Coalescing）。

内存合并指的是，当一个Warp中的32个线程同时访问一块连续且对齐的全局内存时，硬件能够将这32个独立的、微小的内存请求“合并”成一个或少数几个大的内存事务来完成。

理想模式：最理想的访问模式是，Warp中的第k个线程（thread_k）访问内存地址address + k * sizeof(type)。例如，当一个Warp处理float类型（4字节）数组时，如果32个线程分别访问array到array，它们访问的是一个连续的128字节内存块。现代GPU硬件可以非常高效地处理这种请求，一次性获取所有数据。

非合并访问的代价：

跨步访问（Strided Access）：如果线程访问内存时存在固定的间隔，例如thread_k访问array[k * 2]，那么Warp访问的内存地址将是0, 2, 4,…，不再连续。硬件为了满足这些分散的请求，必须发起多次内存事务，其中获取的许多数据都是不需要的，从而导致大量带宽被浪费。
非对齐访问（Misaligned Access）：即使访问是连续的，但如果Warp请求的内存块的起始地址没有与硬件内存段（如32字节、64字节或128字节）的边界对齐，也可能需要额外的内存事务来获取数据，降低效率。
让我们以一个二维矩阵（以行主序存储）的访问为例。假设一个Warp的32个线程要读取数据：

合并访问（Coalesced）：如果32个线程同时读取矩阵的同一行中的32个连续元素，它们的内存访问是连续的，可以实现合并。
非合并访问（Uncoalesced）：如果32个线程同时读取矩阵的同一列中的32个元素，由于数据是按行存储的，这些元素的内存地址会相隔一个行宽（width），访问模式变成了跨步访问，性能会急剧下降。
因此，在设计内核时，必须精心安排数据在内存中的布局和线程的访问方式，以确保Warp内的线程访问是连续的。

2.3 程序员的缓存：共享内存的高级应用
共享内存为程序员提供了对抗高昂全局内存延迟的强大武器。它的主要应用场景有两个：一是作为块内线程间通信的媒介，二是被用作一个由程序员手动管理的缓存，以减少对全局内存的重复访问。

**分块（Tiling）**是利用共享内存作为缓存的最经典、最有效的技术之一，尤其适用于矩阵乘法这类具有高度数据复用性的计算。

朴素矩阵乘法：一个简单的实现是，让每个线程负责计算输出矩阵C中的一个元素。为此，该线程需要循环读取输入矩阵A的一整行和B的一整列。这意味着对A和B中的每个元素，都会被从全局内存中重复读取多次，造成了巨大的带宽浪费。

分块矩阵乘法：

加载分块（Tile Load）：将计算任务分解。每个线程块负责计算输出矩阵C的一个小子块（Tile）。首先，线程块中的所有线程协同，从全局内存中加载计算该子块所需的A和B的对应子块到共享内存中。这一步的加载操作必须设计成合并访问，以保证效率。
同步（Synchronization）：块内所有线程调用__syncthreads()进行同步，确保共享内存中的数据块已全部加载完毕，可以被所有线程安全访问。
块内计算（On-chip Computation）：线程从共享内存中读取数据，进行子矩阵的乘法运算，并将中间结果累加在自己的私有寄存器中。由于共享内存的访问速度比全局内存快几个数量级，这一步非常高效。
循环迭代：重复步骤1-3，加载A和B的下一个分块，继续累加计算结果，直到遍历完所有需要的分块。
写回结果：所有计算完成后，每个线程将最终累加在寄存器中的结果一次性写回到全局内存的目标位置。
通过这种方式，原本需要对全局内存进行N次访问的数据，现在只需要访问一次。数据被加载到高速的共享内存后被反复利用，极大地降低了对全局内存带宽的压力。这种优化策略的本质，是将一个受限于内存带宽的计算问题，转化为一个受限于片上计算延迟的问题，从而让SM中强大的计算单元能够摆脱内存的束缚，全力投入运算。

需要注意的是，共享内存也被划分为32个存储体（Banks）。如果一个Warp中的多个线程同时访问同一个Bank中的数据（除了广播情况），就会发生Bank冲突，导致访问被串行化，从而损失性能。因此在设计共享内存数据布局时，需要考虑访问模式以避免或减少Bank冲突，有时甚至需要通过填充（padding）数组来错开访问。