第二章 CUDA编程基础回顾：GPU架构、线程层次、内存模型与核函数编写

在深入cuBLAS、cuFFT、cuRAND等CUDA数学库之前，必须牢固掌握CUDA编程的核心概念。数学库本质上是高度优化的核函数集合，其性能上限由GPU硬件架构决定，其调用方式则受线程层次与内存模型的约束。本章将系统回顾GPU的硬件架构、线程组织方式、内存体系以及核函数的编写与编译流程，为后续理解数学库的内部机制与性能特征奠定基础。

上图展示了CUDA编程模型的核心要素及其相互关系。硬件架构决定了执行资源，线程层次定义了并行粒度，内存模型规定了数据访问路径，而核函数则是将这些要素串联起来的执行入口。数学库（cuBLAS、cuFFT、cuRAND）正是构建在这一基础之上的高度优化实现。

2.1 GPU硬件架构核心原理

理解GPU架构是编写高效CUDA代码的前提。现代NVIDIA GPU（如Ampere、Hopper、Ada Lovelace架构）采用SIMT（单指令多线程）执行模型，其基本执行单元是流式多处理器（SM）。每个SM包含多个CUDA核心（FP32/FP64单元）、特殊功能单元（SFU）、加载/存储单元（LD/ST）、共享内存、寄存器堆以及线程调度器。

关键架构特征包括：

• SM内部并行：一个SM可同时驻留多个线程块，线程块内的线程通过共享内存和同步机制协作。
• Warp调度：线程以32个为一组（称为Warp）被调度执行。Warp是硬件调度的最小单位，同一Warp内的线程执行相同的指令（可能因分支而分叉）。
• 内存层级：每个线程拥有私有寄存器；每个线程块拥有共享内存（SM内低延迟）；所有线程共享全局内存（显存，通过L2缓存访问）。
• 并发能力：GPU通过大量线程的并发执行来隐藏内存访问延迟。当某些Warp等待内存数据时，调度器立即切换到其他就绪的Warp。

对于数学库而言，cuBLAS的矩阵乘法、cuFFT的蝶形运算、cuRAND的随机数生成，其性能高度依赖于对SM资源（寄存器、共享内存、Warp调度）的充分利用。例如，cuBLAS的GEMM（通用矩阵乘法）内核会使用共享内存进行数据分块（Tiling），以减少全局内存访问次数。

2.2 线程层次与索引映射

CUDA采用三层线程层次：Grid（网格）→ Block（线程块）→ Thread（线程）。核函数启动时，需要指定Grid和Block的维度（最多三维）。这种层次结构直接映射到硬件：Grid对应整个GPU，Block对应SM，Thread对应CUDA核心。

线程索引通过内置变量获取：threadIdx（块内索引）、blockIdx（块索引）、blockDim（块维度）、gridDim（网格维度）。全局线程ID的计算公式为：

// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维情况（常用于矩阵）
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col;

在调用数学库时，虽然不需要直接编写这些索引计算，但理解其原理有助于：

• 预估数据规模与线程配置的匹配关系（例如，cuBLAS会根据矩阵大小自动选择最优的Block配置）。
• 调试自定义核函数时定位线程执行逻辑。
• 理解性能瓶颈——当Block尺寸过小（如64线程）时，SM占用率不足；过大（如1024线程）时，寄存器压力增加。

层次	硬件对应	典型尺寸	同步范围
Thread	CUDA核心	1	无
Warp	调度器	32线程	隐式同步
Block	SM	128-512线程	__syncthreads()
Grid	GPU	数百至数千Block	核函数结束

2.3 内存模型与数据搬运策略

CUDA内存模型是性能优化的核心战场。不同内存的访问速度差异可达两个数量级，数学库的性能优化很大程度上就是内存访问优化。

内存类型及特点：

• 寄存器：每个线程私有，访问速度最快（~1周期），但容量有限（每线程最多255个）。
• 共享内存：每个Block共享，访问延迟约20-30周期，容量通常为48KB-228KB（可配置）。用于线程间数据交换和缓存。
• 全局内存：所有线程可见，容量大（数GB），但延迟高达400-800周期。必须通过合并访问（Coalesced Access）提高带宽利用率。
• 常量/纹理内存：有专门缓存，适合只读数据和特定访问模式。
• 主机端内存：可分页内存（Pageable）与页锁定内存（Pinned），后者支持DMA直接传输，速度更快。

在数学库调用中，数据搬运是不可避免的开销。例如，使用cuBLAS进行矩阵乘法时，需要将矩阵从主机拷贝到设备（cudaMemcpy），计算完成后再拷贝回主机。优化策略包括：

1. 使用页锁定内存（cudaMallocHost）以加速主机-设备传输。
2. 使用异步拷贝（cudaMemcpyAsync）与计算重叠。
3. 在设备端直接生成数据（如cuRAND在设备端生成随机数），避免不必要的主机-设备往返。

2.4 核函数编写与编译流程

核函数（Kernel）是运行在GPU上的函数，使用__global__修饰符声明。其编写需遵循以下规则：

• 返回类型必须为void。
• 不能使用可变参数（如printf在设备端受限）。
• 不能调用主机端函数（除非使用__device__修饰）。
• 通过<<<grid, block>>>语法启动。

一个典型的核函数示例：

__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int n = 1024;
size_t bytes = n * sizeof(float);
float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, bytes);
cudaMalloc(&d_B, bytes);
cudaMalloc(&d_C, bytes);
cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

int threads = 256;
int blocks = (n + threads - 1) / threads;
vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, n);

cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

编译流程使用NVCC编译器：

nvcc -arch=sm_80 -O3 -o vector_add vector_add.cu

NVCC会将.cu文件中的主机代码与设备代码分离，设备代码编译为PTX（并行线程执行）或SASS（流式汇编），主机代码则通过C++编译器处理。编译选项-arch指定计算能力（如sm_80对应Ampere架构），-O3启用优化。

2.5 案例拆解：从零开始优化一个向量加法

为巩固上述概念，我们通过一个实际案例展示从朴素实现到性能优化的完整过程。假设需要计算长度为1000万的浮点向量加法。

初始实现：使用默认配置（256线程/块），直接全局内存访问。

__global__ void vectorAddNaive(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

性能分析：使用nvprof或ncu分析，发现内存吞吐量仅为理论峰值的40%。原因是未使用向量化访问（float4）和网格-块配置不合理。

优化步骤：

1. 向量化访问：使用float4类型，一次处理4个元素，减少指令数。
2. 网格-块调整：将Block设为256，Grid设为ceil(n/(256*4))，确保每个线程处理4个连续元素。
3. 使用页锁定内存：主机端使用cudaMallocHost分配，减少拷贝时间。

__global__ void vectorAddVec4(const float4 *A, const float4 *B, float4 *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i].x = A[i].x + B[i].x;
        C[i].y = A[i].y + B[i].y;
        C[i].z = A[i].z + B[i].z;
        C[i].w = A[i].w + B[i].w;
    }
}

结果对比：

版本	耗时（ms）	有效带宽（GB/s）	说明
朴素版	3.2	12.5	未优化，内存访问未合并
向量化版	1.1	36.4	float4访问，带宽提升2.9倍
向量化+页锁定	0.9	44.4	拷贝时间减少20%

复盘经验：向量化访问是内存密集型核函数最有效的优化手段之一。cuBLAS等数学库内部大量使用float4/double2等向量类型，并采用分块策略（Tiling）将数据加载到共享内存，以最大化数据复用率。

2.6 常见误区与总结

总结：本章回顾了CUDA编程的核心基础。硬件架构决定了资源上限，线程层次提供了并行组织方式，内存模型是性能优化的关键战场，核函数编写与编译流程则是将算法映射到GPU的桥梁。对于后续的数学库学习，需要重点理解：cuBLAS如何利用共享内存和寄存器优化矩阵运算；cuFFT如何通过Warp级协作实现蝶形变换；cuRAND如何利用GPU的并行随机数生成算法。这些库的内部实现都遵循本章所述的基本原则。