2、GPU计算基础：CPU与GPU架构对比、并行计算概念、GPU的线程层次（线程、块、网格）、内存层次概述

在进入CUDA编程之前，必须先建立对GPU硬件架构和编程模型的正确认知。很多初学者直接上手写kernel，却因为不理解线程组织方式和内存访问模式，导致程序运行缓慢甚至结果错误。本章将从CPU与GPU的架构差异出发，解释为什么GPU适合大规模并行计算，然后深入剖析CUDA的线程层次（线程、块、网格）和内存层次，为后续的编程实践打下坚实的理论基础。

2.1 CPU与GPU架构对比：从“少而强”到“多而专”

核心原理：两种不同的设计哲学

CPU（中央处理器）和GPU（图形处理器）的架构差异源于它们的设计目标不同。CPU面向的是延迟优化——它需要快速处理单个复杂的指令流，因此拥有强大的控制单元（分支预测、乱序执行）、大容量的缓存（L1/L2/L3）以及较少的算术逻辑单元（ALU）。相比之下，GPU面向的是吞吐量优化——它需要同时处理成千上万个简单的计算任务，因此将更多的晶体管用于ALU，而控制单元和缓存相对简化。

以NVIDIA A100 GPU为例，它拥有108个流式多处理器（SM），每个SM包含64个FP32 CUDA核心，总计6912个CUDA核心。而一颗高端服务器CPU（如AMD EPYC 7763）通常只有64个物理核心。这种数量级的差异意味着GPU在并行计算场景下具有巨大的吞吐量优势，但在单线程性能上远不及CPU。

判断方法：如何选择CPU还是GPU

在实际项目中，判断一个任务是否适合GPU加速，可以从以下三个维度评估：

• 数据并行性：任务是否可以分解为大量相互独立的子任务？例如，图像每个像素的颜色变换、矩阵每个元素的运算，都是天然并行的。
• 计算密度：算术运算与内存访问的比例是否足够高？如果任务主要是内存拷贝（memory-bound），GPU的高带宽优势可能被延迟掩盖。
• 分支复杂度：任务是否包含大量分支判断？GPU的SIMT（单指令多线程）执行模型要求同一线程束（warp）内的线程执行相同指令，分支发散会导致部分线程等待，降低效率。

案例拆解：向量加法的CPU与GPU实现对比

假设我们需要对两个长度为100万的浮点数组进行逐元素相加。在CPU上，我们通常使用多线程（如OpenMP）将数组分段，每个线程处理一段连续数据。在GPU上，我们为每个元素分配一个线程，一次性启动100万个线程并行执行。以下是两种方式的伪代码对比：

// CPU 多线程方式（伪代码）
for (int t = 0; t < num_threads; t++) {
    int start = t * chunk_size;
    int end = start + chunk_size;
    for (int i = start; i < end; i++) {
        c[i] = a[i] + b[i];
    }
}

// GPU CUDA 方式（伪代码）
__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}
// 启动配置：vecAdd<<<ceil(n/256), 256>>>(a, b, c, n);

在这个案例中，GPU版本通过blockIdx.x、blockDim.x和threadIdx.x三个内置变量计算出全局线程索引，每个线程只处理一个元素。这种“一个线程对应一个数据元素”的模式是CUDA编程中最基本的并行模式。

常见误区与总结

2.2 并行计算概念：从任务并行到数据并行

核心原理：并行计算的两种基本模式

并行计算分为任务并行（Task Parallelism）和数据并行（Data Parallelism）。任务并行是将不同的任务分配给不同的处理单元，每个任务执行不同的指令；数据并行则是将同一操作应用于不同的数据元素，所有处理单元执行相同的指令。GPU的SIMT架构天然适合数据并行，因为它的硬件设计就是围绕“大量线程执行相同指令”这一场景优化的。

在CUDA中，数据并行通常表现为SPMD（单程序多数据）模式：所有线程执行同一个kernel函数，但每个线程处理的数据不同。线程通过内置变量（threadIdx、blockIdx）来区分自己负责的数据范围。

操作流程：设计一个并行算法的步骤

1. 分解：将问题分解为可并行执行的子任务，确定每个子任务的数据范围。
2. 映射：将子任务映射到线程层次上，决定每个线程处理多少数据、线程块如何组织。
3. 通信：确定线程之间是否需要交换数据，如果需要，选择合适的内存空间（共享内存或全局内存）和同步机制。
4. 聚合：将并行结果合并为最终输出，注意避免数据竞争。

案例拆解：矩阵乘法的并行化设计

矩阵乘法C = A × B是并行计算的经典案例。假设A是M×K矩阵，B是K×N矩阵，C是M×N矩阵。最直接的并行策略是：每个线程计算C中的一个元素。此时，线程需要读取A的一行和B的一列，计算内积。但这种方式存在大量重复的内存访问——A的每一行被N个线程重复读取，B的每一列被M个线程重复读取。

更高效的策略是使用分块矩阵乘法（Tiled Matrix Multiplication）：将矩阵划分为多个子块，每个线程块负责计算一个子块，并将子块数据加载到共享内存中，减少全局内存访问次数。这种优化将在后续章节详细展开。

常见误区与总结

2.3 GPU的线程层次：线程、块、网格

核心原理：三层组织结构的含义

CUDA的线程层次是理解GPU编程的关键。它由三个层级组成：

• 线程（Thread）：最小的执行单元，每个线程有独立的程序计数器、寄存器和局部变量。
• 线程块（Block）：一组线程的集合，块内线程可以通过共享内存通信，并通过__syncthreads()同步。同一块内的线程可以协作完成一个子任务。
• 网格（Grid）：一组线程块的集合，网格内的块之间不能直接通信，只能通过全局内存间接交互。

这种分层结构的设计初衷是：可扩展性。GPU硬件不知道网格有多大，它只需要将线程块调度到可用的SM上即可。无论网格包含10个块还是10000个块，硬件都能自动分配，程序无需修改。

判断方法：如何确定线程块大小

线程块大小（blockDim）的选择直接影响性能。以下是常用的判断依据：

考虑因素	建议	原因
线程束大小	块大小取32的倍数（如128、256）	GPU以32个线程为一组（warp）调度，非倍数会导致最后一个warp不满，浪费计算资源
寄存器压力	块大小不宜过大	每个SM的寄存器文件有限，块越大，每个线程可用的寄存器越少，可能导致寄存器溢出到本地内存
共享内存容量	块大小与共享内存用量匹配	共享内存是块内共享的，块越大，每个线程可用的共享内存越少
硬件限制	块大小不超过1024	CUDA规范规定每个块最多1024个线程

案例拆解：图像模糊处理的线程映射

假设我们需要对一张1920×1080的灰度图像进行3×3均值模糊。每个输出像素需要读取周围9个像素的值。我们可以将图像划分为多个块，每个块处理一个16×16的像素区域。此时，块内线程需要访问相邻像素，存在边界数据重叠。解决方案是：每个线程块在共享内存中加载一个18×18的区域（包含边界），然后块内线程从共享内存读取数据，避免重复访问全局内存。

在这个案例中，线程块的大小为16×16=256个线程，网格大小为(1920/16)×(1080/16)=120×68个块。每个线程负责计算一个输出像素。这种“块内协作+块间独立”的模式是GPU编程的典型范式。

常见误区与总结

2.4 内存层次概述：从寄存器到全局内存

核心原理：不同内存的访问速度与作用域

CUDA的内存层次与线程层次紧密对应。从快到慢、从私有到共享，主要包含以下几种内存：

内存类型	作用域	生命周期	访问速度	典型用途
寄存器（Register）	单个线程	线程执行期间	最快（~1周期）	局部变量、中间计算结果
共享内存（Shared Memory）	线程块内	块执行期间	快（~20-30周期）	块内数据共享、数据重用
全局内存（Global Memory）	所有线程	程序运行期间	慢（~400-600周期）	主数据存储、CPU与GPU通信
常量内存（Constant Memory）	所有线程	程序运行期间	快（缓存命中时）	只读常量、查找表
纹理内存（Texture Memory）	所有线程	程序运行期间	快（缓存命中时）	图像处理、空间局部性访问
本地内存（Local Memory）	单个线程	线程执行期间	慢（实际在全局内存中）	寄存器溢出、数组索引访问

操作流程：内存选择的决策路径

1. 判断数据是否只属于单个线程：如果是，优先使用寄存器；如果寄存器不够用，编译器会自动溢出到本地内存（应尽量避免）。
2. 判断数据是否需要在块内共享：如果是，使用共享内存，并注意同步问题。
3. 判断数据是否只读且所有线程访问相同地址：如果是，使用常量内存或纹理内存，利用缓存加速。
4. 其他情况：使用全局内存，但要注意访问模式——尽量让同一warp内的线程访问连续地址（合并访问），以提高带宽利用率。

案例拆解：归约（Reduction）操作的内存优化

归约操作（如求数组所有元素的和）是内存优化的经典案例。朴素实现是每个线程读取一个元素，然后通过原子操作累加到全局变量。但原子操作会导致严重的竞争冲突。更优的方案是：

1. 每个线程块将部分数据加载到共享内存中。
2. 块内线程通过树形归约（两两相加）在共享内存中完成局部求和。
3. 每个块将局部结果写入全局内存。
4. 最后在CPU端（或单个块）汇总所有块的局部结果。

这种方案将全局内存的原子操作次数从N次减少到块数（如1000次），大幅提升了性能。关键点在于：共享内存的访问速度远快于全局内存，且避免了原子操作的竞争开销。

常见误区与总结