4. CUDA编程模型核心：Kernel函数定义与调用、线程层次结构、内置变量

好，咱们今天聊点硬核的。CUDA编程模型的核心，说白了就是三件事：怎么写Kernel、怎么组织线程、怎么用内置变量。这三件事搞明白了，你写CUDA程序就算入门了。

我记得刚接触CUDA那会儿，最让我困惑的就是——这玩意儿到底怎么并行起来的？CPU上我们写个for循环，一个个跑。GPU上呢？你得把任务拆成成千上万个线程，让它们同时干活。这个思维转变，是CUDA编程的第一道坎。

4.1 Kernel函数：GPU上跑的“轻量级任务”

Kernel函数，就是跑在GPU上的函数。定义它的时候，有个关键修饰符——__global__。你看到这个，就知道它是给GPU用的。

来看个最简单的例子：

// Kernel定义
__global__ void vecAdd(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];
    }
}

// 调用方式
int main() {
    // ... 数据准备和内存分配 ...
    vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, N);
    // ... 结果回收 ...
}

这里有个坑，我踩过好几次——Kernel函数里别用全局变量。所有数据都得通过参数传进去，或者用常量内存、纹理内存这些特殊方式。为什么？因为GPU的线程模型决定了，每个线程看到的“全局”是不一样的。

4.2 线程层次结构：Grid、Block、Thread

CUDA的线程组织方式，是个三层结构：Grid → Block → Thread。我画个图帮你理解：

这个结构图你看懂了吗？一个Grid包含多个Block，每个Block包含多个Thread。为什么搞这么复杂？说白了，是为了适配GPU的硬件架构。

Block里的线程可以共享显存（Shared Memory），还能同步（__syncthreads()）。但不同Block之间的线程，老死不相往来——没法直接通信，也没法同步。这个设计是有道理的：Block映射到SM（流多处理器），SM内部通信快，SM之间通信慢。硬件决定了软件模型。

4.3 内置变量：threadIdx、blockIdx、blockDim、gridDim

CUDA提供了几个内置变量，帮你定位当前线程在全局中的位置。这些变量在Kernel函数里可以直接用，不需要声明。

变量名	类型	含义	取值范围
threadIdx	dim3	线程在Block内的索引	0 到 blockDim - 1
blockIdx	dim3	Block在Grid内的索引	0 到 gridDim - 1
blockDim	dim3	Block的维度（包含多少个线程）	由调用时指定
gridDim	dim3	Grid的维度（包含多少个Block）	由调用时指定

这些变量都是 dim3 类型，有 .x、.y、.z 三个分量。你可以用一维、二维、三维来组织线程。我个人最常用的是一维，处理数组、向量这些线性数据时特别方便。

计算全局索引的公式，你得刻在脑子里：

// 一维情况
int tid = threadIdx.x + blockIdx.x * blockDim.x;

// 二维情况
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
int idx = iy * width + ix;  // 线性化

// 三维情况
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
int iz = threadIdx.z + blockIdx.z * blockDim.z;
int idx = iz * height * width + iy * width + ix;  // 线性化

4.4 线程配置：怎么选Block大小和Grid大小

这是个经典问题。Block大小选多少？Grid大小怎么定？我直接给经验值：

• Block大小：通常是32的倍数（因为warp大小是32）。128、256、512都是常见选择。
• Grid大小：根据数据量来算。比如你有N个元素，Block大小是256，那Grid大小就是 ceil(N / 256)。
• 最大限制：每个Block最多1024个线程（计算能力2.0以上）。Grid的每个维度最多2^31-1个Block。

为什么Block大小要选32的倍数？因为GPU执行指令时，是以32个线程为一组（叫warp）来调度的。Block大小是32的倍数，能保证warp被充分利用，不会出现“半空”的warp浪费计算资源。

// 一个典型的线程配置
int N = 1000000;
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;  // 向上取整

vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, N);

这里有个细节——边界检查。因为 numBlocks * blockSize 可能大于N，所以Kernel函数里一定要加 if (i < N) 的判断。不然线程会访问越界内存，轻则数据错误，重则程序崩溃。

嗯，这些就是CUDA编程模型的核心骨架。你把这些搞明白了，后面学共享内存、原子操作、流并行这些高级特性时，就会觉得顺理成章。代码写多了，你会发现——CUDA其实没那么神秘，它就是一套用线程换时间的编程范式。

---