第3章 CUDA核心语法

3. CUDA核心语法：核函数(Kernel)编写、线程层次结构(Block/Grid)、内存模型(Global/Shared/Local)

好，咱们直接进入正题。CUDA编程，说白了就是让GPU帮你干活。但GPU不是CPU，你不能直接扔个for循环给它。你得写一种特殊的函数——核函数（Kernel）。这东西是GPU并行计算的灵魂。

我个人习惯把核函数想象成一个“分身术”。你写一份代码，GPU会复制出成千上万个线程，每个线程都跑这份代码。但每个线程处理的数据不一样。嗯，这就是并行的本质。

3.1 核函数长什么样？

先看个最简单的例子。把两个数组加起来：

// CPU上的写法
void vecAdd(float* A, float* B, float* C, int N) {
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
}

// GPU上的核函数
__global__ void vecAddKernel(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

看到区别了吗？关键就两点：

__global__ 修饰符——告诉编译器这是核函数，跑在GPU上
threadIdx.x、blockIdx.x、blockDim.x——这三个变量帮你算出当前线程该处理哪个数据

我在项目中遇到过不少新手，上来就写 for (int i = 0; i < N; i++)。兄弟，千万别！GPU里每个线程只干一件事，循环是CPU的思维模式。

小技巧：核函数里尽量别用递归。GPU的栈空间很有限，递归容易爆栈。我吃过这个亏，调试了一整天。

3.2 线程层次结构：Block和Grid

为什么要有Block和Grid？你想想看，GPU里有几千个核心，但你要处理的数据可能是几百万甚至上亿。怎么组织这些线程？

CUDA用了两层结构：

Grid（网格）—— 整个任务，包含所有Block
Block（线程块）—— 一组线程，可以协作
Thread（线程）—— 最小的执行单元

调用核函数时，你得指定这两层的大小：

// 定义线程块大小（每个块256个线程）
int threadsPerBlock = 256;
// 计算需要多少个块
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

// 启动核函数
vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里有个坑，我曾经踩过：Block里的线程数不能超过1024（老架构是512）。你设个2048，程序直接崩，连个错误提示都不给。嗯，CUDA的错误处理就是这么“优雅”。

注意：Block内的线程数最好是32的倍数。为什么？因为GPU的warp（线程束）大小就是32。不是32的倍数，会有线程束浪费。我一般用128或256，性能比较稳。

线程索引的计算公式，我建议你死记硬背：

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

这公式其实很简单：blockIdx.x * blockDim.x 是当前块之前的所有线程数，加上 threadIdx.x 就是全局唯一ID。

3.3 内存模型：Global、Shared、Local

GPU的内存不像CPU那样“大一统”。它分好几层，速度差异巨大。我画个图帮你理解：

这张图你得刻在脑子里。我刚开始做CUDA时，总把数据放Global Memory里，结果性能惨不忍睹。后来才明白：能放Shared就别放Global，能放Register就别放Shared。

3.4 三种内存的使用场景

内存类型	作用域	生命周期	典型用途
Global	所有线程 + Host	程序运行期间	存储原始数据、最终结果
Shared	单个Block内	Block生命周期	线程间数据共享、缓存
Local/Register	单个线程	线程生命周期	临时变量、中间结果

举个例子，矩阵乘法。如果直接从Global读数据，每个线程都要读很多次，带宽根本不够。我一般这么做：

__global__ void matMulShared(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    // 分块加载到Shared Memory
    for (int t = 0; t < N / TILE_SIZE; t++) {
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        
        __syncthreads();  // 同步！确保所有线程都加载完了
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();  // 同步！防止覆盖数据
    }
    
    C[row * N + col] = sum;
}

    关键点：注意 __syncthreads() 这个同步函数。没有它，线程A可能还在读数据，线程B已经把数据覆盖了。我曾经因为少写一个同步，结果算出来的结果每次都不一样，排查了整整两天。

3.5 避坑指南

最后，分享几个我踩过的坑：

Bank Conflict——Shared Memory有32个bank，如果同一warp里的线程访问同一个bank的不同地址，就会冲突。解决办法：加padding或者调整访问模式。
寄存器溢出——每个线程的寄存器有限（一般是255个）。变量太多会被挤到Local Memory里，速度骤降。我一般会限制每个线程的变量数，不够就拆成多个核函数。
内存对齐——Global Memory访问最好128字节对齐。不对齐的话，带宽利用率直接腰斩。我习惯用 cudaMallocPitch 来分配二维数组。

嗯，这些内容够你消化一阵子了。记住：写CUDA代码，脑子里时刻想着“数据在哪里”、“数据怎么流动”。把数据放在对的地方，性能自然就上来了。

我的习惯：写核函数之前，先在纸上画出数据流图。哪个数据放Shared，哪个放Global，线程怎么协作。画清楚了再写代码，事半功倍。

📚 课程目录

3. CUDA核心语法：核函数(Kernel)编写、线程层次结构(Block/Grid)、内存模型(Global/Shared/Local)

3.1 核函数长什么样？

3.2 线程层次结构：Block和Grid

3.3 内存模型：Global、Shared、Local

3.4 三种内存的使用场景

3.5 避坑指南