3. CUDA编程模型：Kernel启动、线程层次结构与内存模型

好，咱们今天聊聊CUDA编程模型的核心三件套。说实话，我刚接触CUDA那会儿，最头疼的就是搞不清Grid、Block、Thread到底怎么分工。后来调了上百个kernel，才慢慢摸到门道。

你想想看，GPU里有几千个核心，怎么让它们听话地干活？答案就在这三个层次里。

3.1 Kernel启动：从CPU到GPU的“发令枪”

Kernel启动，说白了就是CPU告诉GPU：“嘿，该干活了！”

语法很简单，但背后门道不少：

// 典型的Kernel启动
myKernel<<<gridDim, blockDim, sharedMem, stream>>>(args);

四个参数，我一个个说：

• gridDim：网格维度，决定有多少个block
• blockDim：块维度，决定每个block里有多少thread
• sharedMem：每个block的动态共享内存大小（字节），默认0
• stream：流ID，默认0表示默认流

我个人习惯，在调试阶段会用cudaDeviceSynchronize()强制同步，确保万无一失。但线上代码里，我会尽量用事件(event)来做细粒度同步，性能会好很多。

3.2 线程层次结构：Grid/Block/Thread

这个结构，我建议你把它想象成一个军队：

• Grid（网格） = 整个战役
• Block（块） = 每个连队
• Thread（线程） = 每个士兵

每个线程都有自己唯一的ID，通过内置变量获取：

int tid = threadIdx.x + blockIdx.x * blockDim.x;

嗯，这里要注意：维度可以是1D、2D、3D。处理图像时我常用2D，处理体数据时用3D。

我在项目中遇到过一个问题：把blockDim设成了(32, 32)，也就是1024个线程。结果共享内存不够用了...后来改成(16, 16)，256个线程，性能反而提升了。为什么？因为每个block占用的资源少了，SM能同时跑更多block。

3.3 内存模型：GPU的“五级存储体系”

GPU的内存，说白了就是“距离越近越快，但容量越小”。我画了张图，帮你理清关系：

这张图我画了好几次才满意。你看，从下往上，速度越来越快，但容量越来越小。实际编程时，你得在这两者之间找平衡。

3.4 内存访问模式：合并访问是关键

全局内存有个特性：合并访问。什么意思？就是相邻线程访问相邻地址时，硬件会把多次访问合并成一次事务。

举个例子，两种访问方式对比：

// 好的方式：合并访问
float val = data[threadIdx.x];  // 相邻线程访问相邻地址

// 坏的方式：非合并访问
float val = data[threadIdx.x * N];  // 跨步访问，性能暴跌

我曾经在一个项目里，就因为没注意这个，性能差了5倍。排查了两天才发现是访问模式的问题...从那以后，我写任何kernel都会先画个内存访问图。

3.5 共享内存：手动管理的“L1缓存”

共享内存是GPU编程的精华。它比全局内存快10倍以上，但需要手动管理。

典型用法：

__global__ void sharedMemExample(float *in, float *out, int N) {
    __shared__ float sdata[256];  // 声明共享内存
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    // 1. 从全局内存加载到共享内存
    sdata[tid] = in[idx];
    __syncthreads();  // 同步！确保所有线程加载完成
    
    // 2. 在共享内存上计算（比如归约）
    for (int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }
    
    // 3. 写回全局内存
    if (tid == 0) out[blockIdx.x] = sdata[0];
}

3.6 内存模型总结

嗯，这张表我建议你存下来。每次写kernel前看一眼，问问自己：数据放哪里最合适？

最后说一句：CUDA编程模型看着复杂，但核心就三件事——怎么启动、怎么组织线程、怎么管理数据。把这三点搞透了，剩下的就是熟练工了。

好了，这一章就到这。记住：纸上得来终觉浅，绝知此事要躬行。赶紧打开你的CUDA环境，动手试试吧。