CUDA 编程基础

CUDA 编程模型

CUDA 程序在由 CPU 和 GPU 组成的异构系统中运行。CPU 及其内存称为 Host（主机），GPU 及其内存称为 Device（设备）。程序从 Host 代码开始执行；Host 通过 CUDA 运行时 API 分配设备内存、传输数据并启动 GPU 函数。由 Host 启动、在 Device 上执行的函数称为 Kernel（核函数）。一次 Kernel 启动会让多个 GPU 线程运行同一函数，线程索引决定各自处理的数据。

图 1 展示使用 cudaMalloc 和 cudaMemcpy 的显式数据传输流程。Kernel 启动后，Host 线程继续执行，GPU 异步处理计算。Host 读取结果前，通过同步操作等待 Kernel 完成。

CUDA 线程组织

CUDA C++ 的启动配置给出 Grid（网格） 和 Block（线程块） 的尺寸；Grid 包含多个 Block，Block 包含多个 Thread（线程）。Block 内的线程在执行时进一步组成 Warp。因此，Grid、Block、Thread 是程序指定的组织层次，而 Warp 是理解这些线程如何执行所需的分组概念。

Grid：一次 Kernel 的线程块集合

一次 Kernel 启动产生一个 Grid。Grid 可以是一维、二维或三维，尺寸由启动配置的第一个参数指定，例如 kernel<<<gridSize, blockSize>>>(...) 中的 gridSize。Grid 的 Block 数量可以超过 GPU 同时能够执行的数量，其余 Block 排队等待硬件调度。blockIdx 标识 Block 在 Grid 中的位置，gridDim 给出 Grid 在各维的 Block 数。

Block：可协作的线程组

Block 的尺寸由启动配置的第二个参数指定；同一个 Grid 中的 Block 具有相同尺寸。线程通过 threadIdx 获取块内索引，通过 blockDim 获取块的尺寸。同一 Block 的线程可使用共享内存交换数据，并用 __syncthreads() 做块内同步。GPU 根据资源情况安排不同 Block 的执行顺序；__syncthreads() 的同步范围是当前 Block。

一维数组通常使用 i = blockIdx.x × blockDim.x + threadIdx.x 计算全局下标。若每个 Block 有 256 个线程，处理 1000 个元素需要 ⌈1000 / 256⌉ = 4 个 Block；启动的 1024 个线程中，最后 24 个线程必须在读写前检查 i < 1000。

Warp：SIMT 执行分组

以本页后面的向量加法为例，vectorAdd<<<4, 256>>>(...) 表示启动 4 个 Block，每个 Block 有 256 个线程。Grid 和 Block 的尺寸由这两个参数指定，Warp 则由 SM 根据 Block 中的线程编号自动组成。每个 Warp 固定有 32 个 lane，lane 是线程在 Warp 内的编号，范围为 0～31。对于一维 Block，编号 0～31 的线程属于第一个 Warp，32～63 属于第二个，以此类推。

SIMT（Single-Instruction, Multiple-Threads，单指令多线程）描述一个 Warp 中的线程执行同一份 Kernel 代码，并分别处理各自的数据。每个线程保留自己的寄存器状态。

同一 Warp 的线程也可以根据条件进入不同分支。例如，对于一维 Block，条件 (threadIdx.x & 1) == 0 会让相邻线程分别进入奇、偶两条路径。GPU 分别处理这些路径，每次由当前路径对应的 lane 执行指令。这种 Warp 分支发散可能降低执行单元利用率。详见 官方 Warps and SIMT 说明。

GPU 硬件执行与软件映射

GPU 包含多个 SM（Streaming Multiprocessor，流式多处理器）。每个 SM 具有寄存器文件、执行单元及片上数据缓存等资源；设备级 L2 缓存和全局内存由多个 SM 使用。一个 Block 的线程在同一个 SM 上执行；SM 的资源允许时可以同时驻留多个 Block。设备根据各 SM 的可用资源动态分配 Block。

SM 的 Warp 调度器从已就绪的 Warp 中选择下一条要执行的指令。当一个 Warp 等待数据时，其他已就绪的 Warp 可以继续执行。调度机制见 官方 Hardware Multithreading 说明。

表中的 Grid、Block、Warp 和 Thread 描述程序及执行层次，SM、寄存器文件、执行单元和缓存描述硬件资源。可同时驻留的 Block 与 Warp 数量受每个线程使用的寄存器、每个 Block 使用的共享内存、线程数及设备上限共同约束；具体上限随 GPU 架构变化。

内存空间与数据传输

CUDA 内存空间的主要属性是可见范围、生命周期和物理位置。寄存器与局部内存对单个线程私有，共享内存供同一 Block 的线程使用，全局内存可由设备上的多个 Block 访问。局部内存的可见范围是单个线程，其数据物理存放在设备内存中。

L1 和 L2 缓存由硬件管理；共享内存由程序显式使用，块内线程交换数据时需遵守同步要求。全局内存中的数组可由 cudaMalloc 分配，并通过 cudaMemcpy 在 Host 与 Device 之间复制。相邻线程访问相邻地址通常有利于合并全局内存请求；合并效果取决于访问尺寸与地址分布，详见 官方合并访存说明。

向量加法 Kernel

向量加法的输入是长度为 n 的数组 a、b，输出满足 c[i] = a[i] + b[i]。各元素的计算相互独立，因此一个线程可以负责一个输出下标。Kernel 用全局索引确定元素位置，并在访问设备内存前检查边界。

__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];
}

完整的 vector_add.cu 包含 Host 端数组初始化、设备内存分配与复制、Kernel 启动、错误检查、结果验证及资源释放。其启动配置为 threadsPerBlock = 256、blocks = (n + 255) / 256。启动后，cudaGetLastError() 检查启动错误；cudaDeviceSynchronize() 等待执行完成并报告执行期错误。相关行为见 CUDA C++ 官方示例。

具备 NVIDIA GPU 与 CUDA Toolkit 的环境可在下载文件所在目录执行以下命令。示例使用可精确表示的小整数构造输入，因此使用直接相等比较；一般浮点计算应根据误差要求设置容差。

nvcc -O2 -std=c++17 vector_add.cu -o vector_add
./vector_add
# 预期输出：OK: 1000 elements, 4 blocks, 256 threads/block

算子优化的常见方向

CUDA 算子优化围绕访存、计算和调度三类开销展开。向量加法每个元素读取两个 float、写入一个 float，并执行一次加法。对于较大的输入，访存通常是首先需要检查的开销；可以从全局内存访问是否合并、带宽使用情况入手，再选择优化手段。

全局内存访问与数据复用

同一 Warp 的相邻线程访问相邻地址，通常有助于将全局内存请求合并为较少的内存事务。共享内存适合多个线程反复使用同一批数据的算子，例如矩阵乘法中的分块计算：线程先从全局内存加载数据，再由同一 Block 的线程在共享内存中复用。向量加法的线程各自读取对应元素并写回结果，直接采用连续的全局内存访问即可。相关机制见 CUDA Programming Guide · Memory Access Patterns。

共享内存 Bank Conflict

共享内存分成多个可并行访问的 Bank。同一 Warp 的线程若在一次访问中请求同一 Bank 的不同地址，硬件需要把这些请求拆开处理，形成 Bank Conflict（存储体冲突）。例如，按列读取某些二维共享内存数组会让相邻线程访问相隔固定步长的地址，可能集中到同一 Bank；调整数组步长或填充一列可改变映射。参见 官方 Shared Memory Access Patterns。

向量化访存与指令数量

当单个线程连续处理多个元素时，可以使用 float4 等向量类型组织数据。在地址满足对齐要求且编译器生成向量指令时，一条 128 bit 加载指令可读取 4 个 float，减少加载指令的发射次数。可以查看生成的指令并测量 Kernel 时间，评估这种写法在目标算子上的效果。参见 NVIDIA 向量化访存说明。

Block 大小、活跃 Warp 数量和 Warp 内分支发散也会影响执行效率。合适的 Block 配置可以让 SM 保持足够多的活跃 Warp，在部分 Warp 等待数据时继续发射其他 Warp 的指令；实际配置可结合分析工具与运行时间评估。

性能测试工具概览

工具的观察范围不同：CUDA Event 适合测量指定 Stream 中的一段 GPU 时间；Nsight Systems（nsys）给出 CPU、CUDA API、Kernel 与内存复制的时间线，用于判断启动、同步和数据传输的耗时；Nsight Compute（ncu）分析单个 Kernel 的启动配置、内存访问和 Occupancy 等指标。入门时先确认结果正确，再确定计时范围，最后根据时间线或 Kernel 指标定位问题。

nsys profile --trace=cuda -o vector_add_timeline ./vector_add
ncu ./vector_add
ncu --list-sections
ncu --section MemoryWorkloadAnalysis ./vector_add

nsys profile 用于观察整体执行时间线；ncu 默认收集基础 Kernel 指标，--section 可针对特定问题收集更多数据。使用 ncu --list-sections 查询当前环境可收集的章节。Windows 下将命令中的 ./vector_add 换为 .\vector_add.exe。示例的 1000 个元素用于检查执行流程；研究带宽时应扩大输入并记录设备与测试条件。完整计时方法、工具指标解释与访存映射实测见CUDA 性能测试与调优；命令细节见 Nsight Systems User Guide 和 Nsight Compute CLI Guide。

结论

Grid 与 Block 规定程序的线程组织，Warp 描述 Block 内线程的 SIMT 执行分组，SM 提供实际执行资源。向量加法用线程索引和边界判断连接这些层次。算子优化需要分别检查全局内存事务、共享内存 Bank Conflict、指令数量、并行度和分支行为，再借助 CUDA Event、Nsight Systems 与 Nsight Compute 验证瓶颈和收益。

参考资料

• NVIDIA CUDA Programming Guide · Programming Model
• NVIDIA CUDA Programming Guide · Writing SIMT Kernels
• NVIDIA CUDA Programming Guide · Intro to CUDA C++
• CUDA C++ Best Practices Guide、NVIDIA Vectorized Memory Access
• Nsight Systems User Guide、Nsight Compute CLI Guide