4. CUDA基础语法：执行配置、内置变量与数据类型

好，咱们进入正题。CUDA编程里最基础也最绕不开的，就是那对尖括号 <<< >>> 和几个内置变量。很多新手一上来就被这些符号搞晕了，觉得像天书。其实说白了，它们就是告诉GPU「谁干活、干多少活、怎么干」的一套指令。

我个人习惯把CUDA程序想象成一个工厂流水线。CPU是厂长，负责调度；GPU是车间，里面有很多工人（线程）。那 <<<grid, block>>> 就是厂长在喊：「张三，你带10个人去搬那堆货！」—— grid 是张三这个小组长，block 是那10个人。

4.1 执行配置：<<<grid, block>>> 到底在配置什么？

先看一个最典型的调用：

kernel_name<<<1, 256>>>(d_a, d_b, d_c);

这行代码的意思是：启动1个线程块（block），每个块里有256个线程。总共就是 1 × 256 = 256 个线程在干活。

你可能会问：「为什么是256？不是128、512？」嗯，这个问题我当年也纠结过。我在项目中遇到过，选256是因为它刚好是32的倍数——GPU的warp（线程束）是32个线程一组，256正好是8个warp，调度起来最舒服。选128也行，但256在大多数场景下能更好地隐藏内存延迟。

执行配置的完整形式其实是：

<<<Dg, Db, Ns, S>>>

• Dg：grid的维度，可以是 dim3 类型，比如 dim3(2, 3, 1)
• Db：block的维度，比如 dim3(16, 16, 1)
• Ns：每个block动态分配的共享内存大小（字节），默认0
• S：指定在哪个CUDA流上执行，默认0（默认流）

大多数时候我们只用前两个参数。但第三个参数Ns，我建议你记住它——我在做矩阵乘法优化时，就靠它给每个block分配共享内存，性能直接翻倍。

4.2 内置变量：threadIdx, blockIdx, blockDim, gridDim

这四个变量是CUDA给每个线程发的「身份证」。你想想看，几千上万个线程同时跑，每个线程怎么知道自己该处理哪部分数据？就靠它们。

它们都是 dim3 类型，有 .x、.y、.z 三个分量。一维时只用 .x。

最常用的一个公式，用来计算全局线程ID：

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

这个公式我闭着眼睛都能写出来。它把block的索引和线程的索引拼成一个全局唯一的ID。比如grid有4个block，每个block有256个线程，那第2个block的第100个线程的全局ID就是：1×256 + 100 = 356。

二维或三维的情况稍微复杂一点，但原理一样：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;  // 二维转一维索引

我曾经在图像处理项目里用过二维grid，每个线程处理一个像素。用 threadIdx.x 和 threadIdx.y 分别对应像素的列和行，代码写起来特别直观。

4.3 数据类型：host与device的「语言」差异

CUDA里数据类型分两套：一套在CPU（host）上用，一套在GPU（device）上用。虽然大部分基础类型（int, float, double）两边通用，但有几个特殊类型你得记住。

4.3.1 基本数据类型

• int, float, double：两边通用，但 double 在GPU上计算速度比 float 慢很多（有些老架构甚至慢8倍）。我建议非必要不用double。
• char, short：适合存储小数值，节省显存。
• unsigned int：索引计算时常用，避免负数。

4.3.2 内置向量类型

CUDA提供了 int2, float3, double4 等向量类型。它们有 .x, .y, .z, .w 分量。

float4 pixel = make_float4(0.5f, 0.3f, 0.8f, 1.0f);
float r = pixel.x;  // 红色通道

用向量类型的好处是：一次内存读取可以拿到多个值，带宽利用率更高。我在做颜色空间转换时，用 uchar4 一次读入RGBA四个通道，比分别读快了不少。

4.3.3 特殊类型：dim3 和 cudaError_t

• dim3：用于定义grid和block的维度。默认值是1，所以 dim3(256) 等价于 dim3(256, 1, 1)。
• cudaError_t：CUDA API的返回值类型。我强烈建议每次调用CUDA函数后都检查返回值，否则出错了你都不知道。

cudaError_t err = cudaMalloc(&d_ptr, size);
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
    return -1;
}

4.4 知识体系总览

下面这张图把本章的核心逻辑串起来了。你可以看到，从执行配置到内置变量，再到数据类型，它们共同构成了CUDA kernel的「骨架」。

嗯，这张图把本章内容串起来了。执行配置决定了「有多少人干活」，内置变量决定了「每个人干哪块活」，数据类型决定了「活干得精不精细」。三者缺一不可。

4.5 一个完整的例子

最后，给你看一个完整的向量加法kernel，把今天讲的东西都用上：

__global__ void vecAdd(float *a, float *b, float *c, int n) {
    // 计算全局线程ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查：防止越界
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

// 调用时
int n = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

注意那个 if (tid < n) 的边界检查。我刚开始写CUDA时经常忘了加，结果数组最后几个元素被乱写，调试到崩溃。后来养成了习惯：只要数据长度不是block大小的整数倍，就一定要加边界检查。