4. CUDA编程模型：Kernel启动、Grid-Block-Thread层级、内存访问模式对带宽的影响

好，咱们今天聊点硬核的。CUDA编程模型，说白了就是你怎么指挥GPU干活。我见过不少新手，代码写出来能跑，但带宽利用率惨不忍睹。问题出在哪？多半是对Kernel启动、线程层级和内存访问模式的理解不够深。

我个人习惯，把CUDA编程模型拆成三个层次来看：怎么启动、怎么组织、怎么访存。这三个点搞明白了，带宽优化就有方向了。

4.1 Kernel启动：你的一声令下，GPU怎么接招？

Kernel启动，就是你在CPU端调用一个函数，让GPU去执行。语法很简单，但背后的机制值得琢磨。

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

这里四个参数，我一个个说：

• gridDim：网格维度，决定有多少个Block
• blockDim：线程块维度，决定每个Block有多少个Thread
• sharedMemSize：动态共享内存大小（可选）
• stream：流（可选，用于异步执行）

嗯，这里要注意：Kernel启动是异步的。CPU发完指令就继续往下跑了，不会等GPU算完。我曾经在项目里踩过这个坑——CPU端数据还没准备好，GPU就开始读了，结果全是垃圾值。所以记得加cudaDeviceSynchronize()或者用cudaMemcpy来隐式同步。

4.2 Grid-Block-Thread层级：GPU的军队编制

GPU的线程组织，就像一支军队。你想想看：

• Grid 是整个战役
• Block 是一个作战连队
• Thread 是每个士兵

每个Thread执行相同的代码（Kernel函数），但处理不同的数据。这就是SIMT（单指令多线程）的精髓。

我建议你记住这几个硬件限制：

为什么Warp大小是32？因为GPU的SIMD单元一次处理32个线程。如果你的Block大小不是32的倍数，就会有线程束分化，性能直接打折。我在项目中遇到过，把Block大小从128改成160（32的倍数），带宽提升了15%。

4.3 内存访问模式：带宽的命门

好了，前面都是铺垫。现在说重点——内存访问模式。GPU的显存带宽很高，但前提是你的访问模式要对。

我画了一张图，帮你理解三种典型的内存访问模式：

说白了，GPU最喜欢的是合并访问——连续的线程访问连续的内存地址。这样一次内存事务就能把数据都取回来。反之，如果线程0访问A[0]，线程1访问A[100]，那就得发两次内存请求，带宽直接腰斩。

4.4 实战：如何写出带宽友好的代码

光说不练假把式。咱们看一个实际例子：

// ❌ 坏代码：非合并访问
__global__ void badAccess(float *A, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 注意这里：按列访问，但线程是按行组织的
    float val = A[tid * N + 0];  // 每个线程访问不同行
}

// ✅ 好代码：合并访问
__global__ void goodAccess(float *A, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 连续线程访问连续地址
    float val = A[tid];  // 完美合并
}

为什么会这样？因为GPU的内存控制器一次可以读取128字节（32个float）。如果32个线程访问连续的32个float，一次事务搞定。如果分散访问，就得多次事务。

4.5 共享内存：程序员可控的缓存

除了全局内存，GPU还有共享内存。它就在每个Block内部，速度比全局内存快一个数量级。我建议你把它当成手动管理的L1缓存来用。

// 使用共享内存做矩阵分块
__global__ void sharedMemExample(float *A, float *B, float *C, int N) {
    __shared__ float tileA[16][16];
    __shared__ float tileB[16][16];
    
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    
    float sum = 0.0f;
    for (int k = 0; k < N; k += 16) {
        // 加载到共享内存
        tileA[threadIdx.y][threadIdx.x] = A[row * N + k + threadIdx.x];
        tileB[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();  // 同步！确保所有线程都加载完毕
        
        // 从共享内存计算
        for (int i = 0; i < 16; i++) {
            sum += tileA[threadIdx.y][i] * tileB[i][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

嗯，这里有个关键点：__syncthreads()。它保证Block内所有线程都执行到同一位置。没有它，共享内存里的数据可能还没写完就被读了，结果就是随机错误。我曾经调试过一个bug，就是少了一个__syncthreads()，结果每跑三次就错一次。

4.6 总结：带宽优化的三个层次

好了，咱们捋一捋：

1. Kernel启动：注意异步特性，别在数据没准备好时就启动Kernel
2. 线程组织：Block大小要是32的倍数，避免线程束分化
3. 内存访问：尽量合并访问，多用共享内存做数据复用

我个人习惯，写Kernel之前先画个内存访问模式图。你想想看，如果每个线程的访问地址都能画成一条直线，那带宽利用率肯定高。如果画出来像一团乱麻，那就得重构了。

最后送你一句话：GPU的带宽不是用来浪费的，是用来压榨的。每次写Kernel，都问问自己——我的内存访问模式，GPU喜欢吗？