GPU架构与并行计算基础

在开始CUDA编程之前，必须先理解GPU的硬件设计哲学。CPU与GPU的差异并非简单的“谁更快”，而是源于不同的设计目标：CPU面向延迟优化，GPU面向吞吐量优化。本章将深入拆解NVIDIA GPU的核心硬件单元（SM、SP、内存层次），并与CPU进行系统对比，最后建立线程、块、网格这一CUDA编程模型中最基本的抽象概念。理解这些底层机制，是后续编写高效内核函数的基石。

2.1 GPU硬件架构：从SM到SP的层级拆解

核心原理：GPU的并行计算单元

NVIDIA GPU的硬件设计围绕流式多处理器（Streaming Multiprocessor, SM）展开。SM是GPU的基本计算单元，一个GPU包含多个SM（例如A100有108个SM，RTX 3090有82个SM）。每个SM内部又包含多个流处理器（Streaming Processor, SP），SP即CUDA核心，是执行浮点运算的最小硬件单元。以Ampere架构为例，每个SM包含128个CUDA核心、4个Tensor Core、4个纹理单元和4个调度器。

SM内部的关键组件包括：

• CUDA核心（SP）：执行整数和浮点运算，是并行计算的主力。
• 调度器（Warp Scheduler）：每个SM有4个调度器，每个调度器每时钟周期可发射一条指令到32个线程组成的Warp。
• 寄存器文件（Register File）：每个SM拥有256KB寄存器，分为多个32位寄存器，供所有线程共享。
• 共享内存（Shared Memory）：每个SM有192KB（可配置），供块内线程高速通信。
• L1缓存：与共享内存共用物理空间，可配置为不同比例。

GPU的线程调度以Warp（32个线程）为最小单位。当一个Warp中的线程执行同一指令时，称为“锁步”执行；若出现分支发散，则Warp会串行执行各分支路径，导致性能下降。

判断方法：如何识别GPU的硬件规格

在编写CUDA代码前，必须了解目标GPU的硬件参数。使用以下方法获取信息：

1. 运行 nvidia-smi 查看GPU型号、显存大小、驱动版本。
2. 在CUDA代码中调用 cudaGetDeviceProperties() 获取SM数量、每个SM的最大线程数、共享内存大小等。
3. 查阅官方规格表：例如A100的SM数为108，每个SM最大线程数为2048，因此最大并发线程数为108×2048=221,184。

案例拆解：从硬件规格推导内核配置

假设我们使用RTX 3090（82个SM，每个SM最大2048线程），需要处理100万个元素的向量加法。初始配置为：

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; // 约3907块

此时GPU可同时驻留的最大线程数为 82×2048=167,936。由于3907个块中，每个块256线程，总线程数为1,000,192，远超驻留上限，因此GPU会分批次调度块执行。实际执行时，每个SM同时驻留约2048/256=8个块，82个SM共驻留656个块，其余块排队等待。

若改用 threadsPerBlock=512，则每个SM驻留4个块，块数减少为1954，调度开销降低，但可能因块数不足导致SM空闲。因此，块大小选择需平衡资源利用率和调度开销。

常见误区与总结

总结：SM是硬件资源的核心，SP是执行单元，Warp是调度单位。理解SM的资源上限（寄存器、共享内存、线程数）是优化CUDA程序的第一步。

2.2 CPU与GPU对比：设计哲学与适用场景

核心原理：两种不同的计算范式

CPU与GPU的差异源于设计目标不同。CPU面向延迟优化，需要快速响应单个复杂任务；GPU面向吞吐量优化，需要同时处理海量简单任务。具体对比见下表：

CPU的每个核心都拥有独立的ALU、控制单元和缓存，能高效处理分支密集的代码；GPU的众多核心共享控制单元和缓存，通过大规模并行掩盖内存延迟。

判断方法：何时使用GPU加速

并非所有程序都适合GPU。判断标准如下：

1. 数据并行性：任务能否分解为大量独立的小任务？例如向量运算、矩阵乘法、像素处理。
2. 计算密度：算术运算与内存访问的比值是否足够高？若程序是内存密集型（如链表遍历），GPU的带宽优势可能被延迟掩盖。
3. 分支复杂度：若代码包含大量if-else且分支发散严重，GPU的Warp执行效率会大幅下降。

案例拆解：矩阵乘法的CPU与GPU实现对比

以512×512矩阵乘法为例，CPU单线程实现需要约512³=1.34亿次乘加运算。在3GHz CPU上，若每次运算需4个时钟周期，则耗时约0.18秒。GPU实现时，将输出矩阵的每个元素分配给一个线程，共262,144个线程，每个线程执行一次点积。在A100上，理论峰值算力为19.5 TFLOPS，实际可达约0.01秒，提速约18倍。

但若矩阵规模为32×32（1024个元素），CPU耗时约0.1毫秒，GPU因内核启动开销（约5-10微秒）和数据传输开销（PCIe 3.0约12GB/s），总耗时可能超过0.5毫秒，反而更慢。因此，GPU加速需要“足够大的问题规模”来摊销固定开销。

常见误区与总结

总结：CPU擅长复杂控制流和低延迟任务，GPU擅长大规模并行计算。选择硬件时需评估数据规模、并行度、分支复杂度和传输开销。

2.3 并行计算基本概念：线程、块、网格

核心原理：CUDA的线程层次模型

CUDA将并行任务组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。这种层次设计是为了匹配GPU的硬件调度机制：

• 线程（Thread）：最小的执行单元，每个线程执行相同的内核函数，但通过线程ID访问不同的数据。
• 线程块（Block）：一组线程的集合，块内线程可通过共享内存通信，并通过__syncthreads()同步。块被调度到单个SM上执行。
• 网格（Grid）：所有线程块的集合，对应一次内核启动。网格中的块可以并行或串行执行（取决于SM数量）。

线程和块可以使用一维、二维或三维索引，便于处理图像（二维）、体数据（三维）等结构。例如，处理1024×1024图像时，可定义dim3 block(16,16)，dim3 grid(64,64)，每个线程对应一个像素。

操作流程：定义和启动内核

编写CUDA内核的完整流程如下：

1. 使用 __global__ 关键字声明内核函数（在GPU上执行，由CPU调用）。
2. 在主机代码中定义网格和块维度：dim3 grid(blocksPerGrid); dim3 block(threadsPerBlock);
3. 调用内核：kernel<<<grid, block>>>(args);
4. 在内核中使用 blockIdx.x、threadIdx.x、blockDim.x 计算全局线程ID。

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}
// 主机端调用
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

全局线程ID的计算公式为 blockIdx.x * blockDim.x + threadIdx.x，这是CUDA编程中最基础的索引模式。

案例拆解：图像灰度化中的线程映射

假设有一张1920×1080的RGB图像，需要转换为灰度图。每个像素的灰度值计算公式为 gray = 0.299R + 0.587G + 0.114B。我们定义二维线程块：

dim3 block(32, 8);  // 每个块256线程
dim3 grid((1920 + 31)/32, (1080 + 7)/8);  // grid(60, 135)
__global__ void rgbToGray(unsigned char *rgb, unsigned char *gray, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        int idx = (y * width + x) * 3;
        gray[y * width + x] = 0.299f * rgb[idx] + 0.587f * rgb[idx+1] + 0.114f * rgb[idx+2];
    }
}

该映射中，每个线程处理一个像素，块内线程共享x和y方向的连续性，有利于内存访问合并（coalescing）。若使用一维映射，则需手动计算x和y坐标，且内存访问模式可能不连续。

常见误区与总结

总结：线程、块、网格是CUDA编程的核心抽象。理解三者与硬件（SP、SM、GPU）的映射关系，是编写高效内核的前提。实际编程中，需根据数据规模和硬件资源合理选择块大小和网格维度。