2. CUDA编程模型核心

要写出第一个真正能跑的CUDA核函数，必须先建立对CUDA编程模型的整体认知。很多初学者把CUDA当成“C语言加几个关键字”，结果核函数能编译却算错，或者能算对却慢得离谱。根本原因在于：CUDA不是顺序执行模型，而是线程层次、内存层次、执行模型与SIMT架构四者耦合的并行模型。理解这四层关系，向量加法才不是“抄来的模板”，而是可推导、可验证、可优化的起点。

线程层次：从网格到线程的三级映射

CUDA把一次核函数启动组织成三级结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。核函数以线程为单位执行，但线程不是孤立运行的，它通过内置变量知道自己“是谁、在哪”。

• gridDim：网格中线程块的数量，可以是1D/2D/3D。
• blockDim：每个线程块中线程的数量，同样支持多维。
• blockIdx：当前线程块在网格中的索引。
• threadIdx：当前线程在线程块内的索引。

全局线程编号由这两级索引组合得到。以1D为例：int i = blockIdx.x * blockDim.x + threadIdx.x;。这个公式是向量加法的“地址生成器”，也是后面所有索引推导的基础。线程块大小通常取128、256、512等32的倍数，因为硬件以warp（32线程）为调度单位，非32倍数会造成最后一个warp部分线程闲置。

内存层次：谁快谁慢，谁共享谁私有

CUDA的内存层次直接决定性能。从线程视角看，内存分为：

• 寄存器（Register）：线程私有，最快，数量有限。
• 共享内存（Shared Memory）：线程块内共享，片上，速度快，需显式声明__shared__。
• 全局内存（Global Memory）：所有线程可见，容量大但延迟高，是向量加法中数组A、B、C的存放位置。
• 常量内存、纹理内存：特殊用途，读多写少场景使用。

向量加法中，A、B、C三个数组都在全局内存。核函数每个线程读取A[i]、B[i]，计算后写回C[i]。这里的关键是合并访问（coalesced access）：相邻线程访问相邻地址，硬件能把32个线程的访问合并成少量内存事务。如果索引写错，比如用threadIdx.x * blockDim.x + blockIdx.x，相邻线程访问的地址会跨步，带宽利用率骤降。

执行模型与SIMT：为什么分支会拖慢速度

CUDA采用SIMT（Single Instruction, Multiple Threads）架构。硬件把32个线程组成一个warp，warp内所有线程在同一周期执行同一条指令，但操作各自的数据。这带来两个重要后果：

• 分支发散（Branch Divergence）：如果warp内线程走不同分支，硬件会串行执行各分支路径，只有对应线程活跃，其余线程屏蔽。向量加法没有分支，所以不受影响，但这是后续复杂核函数必须警惕的。
• 延迟隐藏（Latency Hiding）：全局内存访问延迟数百周期，GPU靠大量warp切换来隐藏延迟。因此占用率（Occupancy）——即SM上活跃warp数与最大支持warp数之比——是重要指标。线程块太小或寄存器用量太高都会降低占用率。

下面这张图把线程层次、内存层次与SIMT执行关系放在一起，帮助建立整体映射：

第一个向量加法核函数

假设有两个长度为N的浮点数组A和B，要计算C = A + B。核函数如下：

__global__ void vecAdd(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];
    }
}

这里有几个必须解释清楚的点。第一，__global__表示这是核函数，在GPU上执行、由CPU调用。第二，if (i < N)是边界检查，因为N通常不是blockDim.x的整数倍，最后一个线程块中会有线程越界。第三，参数中的指针是设备指针，必须由cudaMalloc分配。

主机端启动代码：

int N = 100000;
size_t bytes = N * sizeof(float);
float *h_A = (float*)malloc(bytes);
float *h_B = (float*)malloc(bytes);
float *h_C = (float*)malloc(bytes);
// 初始化 h_A、h_B ...

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 blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

注意gridSize的计算用了向上取整，保证覆盖全部N个元素。如果N=100000，blockSize=256，则gridSize=391，总线程数=391×256=100096，多出的96个线程由if (i < N)拦截。

正确性验证：不能只看“跑通了”

核函数执行完不报错，不代表结果正确。必须做逐元素比对。最直接的方法是在CPU上算一遍参考结果，然后与GPU结果比较：

bool ok = true;
for (int i = 0; i < N; i++) {
    float ref = h_A[i] + h_B[i];
    if (fabs(h_C[i] - ref) > 1e-5) {
        printf("Mismatch at %d: GPU=%f, CPU=%f\n", i, h_C[i], ref);
        ok = false;
        break;
    }
}
printf(ok ? "PASS\n" : "FAIL\n");

浮点比较不能用==，因为GPU和CPU的浮点运算顺序、精度可能略有差异，通常用绝对误差或相对误差阈值。对于向量加法这种单次加法，误差极小，1e-5足够严格。

常见错误与排查方向

第一个核函数最容易踩的坑有三类。第一类是索引错误：把blockIdx.x * blockDim.x + threadIdx.x写成threadIdx.x * blockDim.x + blockIdx.x，导致相邻线程访问跨步地址，结果错乱且性能差。第二类是边界遗漏：去掉if (i < N)后，越界线程会读写非法内存，可能不报错但污染其他数据。第三类是内存拷贝方向错误：把cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost写反，导致设备指针拿到垃圾数据。

排查时建议先用小规模N（如N=1024）验证逻辑，再放大到实际规模。小规模下可以打印每个线程的i值，确认索引映射正确。确认逻辑无误后，再关注性能：用nvprof或Nsight Compute查看内存吞吐、占用率和warp执行效率。

从“能跑”到“跑得好”

向量加法本身计算量极小，性能瓶颈几乎全在内存带宽。一个合格的实现应该做到：相邻线程访问相邻地址（合并访问）、线程块大小为32的倍数、边界检查不引入额外分支发散（实际上边界检查只影响最后一个块，可接受）。如果进一步优化，可以考虑使用float4向量化加载，让每个线程一次处理4个元素，减少指令数并提高内存事务效率。但那是下一步的事——先把基础版本写对、验对，再谈优化。

本节的核心不是记住一个模板，而是理解：线程层次决定了“谁算哪个元素”，内存层次决定了“数据放在哪、怎么访问”，SIMT决定了“为什么这样写才快”。这三者共同支撑起向量加法的正确性与性能。后续更复杂的核函数，都是在这个骨架上生长出来的。