第一章：CUDA编程基础回顾

各位同学好，我是你们的老朋友。今天咱们聊聊CUDA编程里最基础、也最绕不开的几个概念。说实话，我见过不少新手一上来就追各种花里胡哨的优化技巧，结果连Kernel Launch的配置都写不对——这就像盖楼不打地基，迟早要塌。

这一章，我会带着大家把Grid-Block-Thread映射、各种内存的使用、同步和原子操作这些基本功再过一遍。嗯，都是些老生常谈的东西，但相信我，里面有不少坑是我自己踩过的。

1.1 Kernel Launch：你写的第一个CUDA函数

Kernel Launch，说白了就是让GPU干活的那条指令。你写一个函数，前面加个__global__，然后调用的时候用三括号语法<<<grid, block>>>。就这么简单？其实不然。

// 一个简单的向量加法Kernel
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 调用方式
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

我个人习惯把blocksPerGrid和threadsPerBlock分开定义，这样调试的时候一目了然。你在项目里可能会遇到一个问题：threadsPerBlock设多大合适？我建议一般取128到512之间，具体要看你的Kernel里用了多少寄存器。

1.2 Grid-Block-Thread映射：搞清楚谁在干活

你想想看，GPU里有成千上万个线程在跑。怎么知道哪个线程处理哪个数据？靠的就是这个映射关系。

• threadIdx：线程在Block内的索引，可以是1D、2D或3D
• blockIdx：Block在Grid内的索引
• blockDim：Block的维度大小
• gridDim：Grid的维度大小

最常见的映射公式：

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

这个公式我闭着眼睛都能写出来。但要注意，如果你的数据是2D的，比如图像处理，那就得用2D的映射：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col;

为什么会这样？因为GPU的线程调度是以Warp（32个线程）为单位的。如果你用1D映射处理2D数据，访存效率会大打折扣。我在做矩阵转置的时候吃过这个亏，后来改成2D映射，性能直接翻倍。

1.3 内存三兄弟：Global、Local、Shared

CUDA里的内存类型，说白了就三种：Global Memory（全局内存）、Local Memory（本地内存）、Shared Memory（共享内存）。它们各有各的脾气。

Global Memory是最大的，但也是最慢的。我建议你尽量少直接读写Global Memory，能复用数据就复用。Shared Memory是块宝，同一个Block里的线程可以共享数据，速度比Global快两个数量级。

// 使用Shared Memory做矩阵分块乘法
__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;
    
    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;
}

这段代码里，__shared__声明了Shared Memory数组。__syncthreads()是同步屏障，确保所有线程都加载完数据再开始计算。我刚开始写这个的时候，忘了加第二个__syncthreads()，结果数据被覆盖了，算出来的结果全是错的——嗯，这种错误排查起来特别痛苦。

1.4 同步与原子操作：别让线程打架

多线程编程最头疼的问题就是数据竞争。CUDA里提供了两种机制来解决：同步（Synchronization）和原子操作（Atomic Operations）。

同步：__syncthreads()

这个函数是Block级别的同步屏障。所有线程必须都执行到这一行，才能继续往下走。我建议你在以下场景必须使用：

• Shared Memory读写之间
• 一个线程写数据，另一个线程读同一个数据
• 分块算法中，不同阶段之间

但要注意，__syncthreads()不能用在if-else分支里。如果一个线程进了if，另一个线程进了else，那就会死锁。我曾经在代码里写过这样的逻辑，结果程序卡死了，排查了半天才发现是同步位置不对。

原子操作：atomicAdd、atomicCAS等

原子操作用于处理全局变量的并发更新。比如你要统计数组里有多少个正数：

__global__ void countPositive(float* data, int* count, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N && data[idx] > 0) {
        atomicAdd(count, 1);
    }
}

这里atomicAdd保证了对count的累加是线程安全的。但原子操作是有性能代价的——它会让多个线程串行化。我建议你尽量少用原子操作，能用并行规约（Parallel Reduction）解决的问题，就别用原子操作。

1.5 本章知识体系

下面这张图是我自己画的，把这一章的核心逻辑串起来了。你仔细看看，应该能对CUDA编程基础有个整体的认识。

这张图把Kernel Launch、线程映射、内存类型、同步原子操作串在了一起。你写代码的时候，脑子里要有这张图——每个线程在哪个Block里，访问哪块内存，需不需要同步，心里都得有数。

好了，这一章的内容就到这里。CUDA编程基础虽然看起来简单，但真正用好用对，还是需要大量实践的。我建议你回去把上面的代码示例都跑一遍，改改参数，看看性能变化。嗯，动手才是最好的学习方式。