写第一个 CUDA kernel：索引、同步与启动配置

前两篇建立了执行模型和内存层次的直觉。这一篇动手：把一个 CUDA kernel 的骨架拆开——线程怎么找到自己负责的数据、什么时候需要同步、启动时该选多大的 block。这些是写任何算子都绕不开的基本功。

一、kernel 的解剖：向量加法

最小的有用 kernel 是逐元素操作。向量加法 c = a + b：

extern "C" __global__
void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

几个要点：

• __global__ 标记这是一个从 host 启动、在 device 执行的 kernel，返回类型必须是 void。
• 每个线程算自己那一个元素。全局索引 i 由 blockIdx.x * blockDim.x + threadIdx.x 得到——block 偏移加上 block 内偏移，这是一维场景的标准写法。
• if (i < n) 是边界检查。线程总数通常向上取整到 block 的整数倍，会多出一些线程，必须挡住越界访问，否则就是非法内存访问。
• const float* __restrict__ 提示编译器指针不重叠，有时能帮助优化；这里省略了 __restrict__，访存优化篇 会用到。

启动时用三尖括号配置 grid 和 block：

int threads = 256;
int blocks = (n + threads - 1) / threads;   // 向上取整覆盖所有元素
vadd<<<blocks, threads>>>(d_a, d_b, d_c, n);

二、grid-stride loop：让 kernel 与数据规模解耦

“一个线程一个元素”在 \(n\) 很大时会启动巨量 block。更通用、更稳健的写法是 grid-stride loop：启动一个适配 GPU 规模的固定 grid，每个线程用步长 gridDim.x * blockDim.x 跨着处理多个元素。

extern "C" __global__
void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    for (; i < n; i += stride) c[i] = a[i] + b[i];
}

好处：launch 配置不再被 \(n\) 绑死，可以按 GPU 的 SM 数量选一个让占用饱和的 grid；同一份 kernel 处理任意大小输入；调试时甚至能用单 block 跑。NVIDIA 推荐把它作为默认写法。本系列后续的流式 kernel 多用这种结构。

三、__syncthreads：block 内同步

逐元素操作不需要线程间通信。但一旦线程要共享中间结果（比如先把数据搬进 shared memory 再协作计算），就需要同步。__syncthreads() 是 block 内的栅栏：它保证 block 内所有线程都到达这一点、且之前的 shared/global 写入对彼此可见，才继续往下走。

一个典型模式（后面归约、GEMM 都会用）：

__shared__ float tile[256];
tile[threadIdx.x] = in[i];   // 每个线程往 shared 写一份
__syncthreads();             // 等所有线程写完
// 现在可以安全地读别的线程写的 tile[...]

三个常见错误：

1. 在分支里调用 __syncthreads() 且不是所有线程都进入该分支——会死锁，因为栅栏要求 block 内全员到达。
2. 忘记同步就读别人写的 shared 数据——读到旧值，结果随机错误。
3. 以为它能跨 block 同步——它只管一个 block。跨 block 的同步要靠 kernel 边界（启动新 kernel）或 cooperative groups 的 grid 同步。

注意 __syncthreads() 只同步线程，不替代 warp 内的内存可见性细节；warp 级原语和 __syncwarp() 在 reduction 篇 展开。

四、launch 配置：block 该选多大

block 大小必须是 32 的倍数（warp 粒度），上限 1024。怎么选？用 grid-stride 版本的向量加法实测不同 block 大小（grid 固定为 \(38 \times 16\)，\(n=2^{24}\)，CUDA event 取中位数）：

结论很清楚：block=32 明显偏慢（约低 9%），因为每 block 只有一个 warp，SM 上能并存的 warp 受 block 数上限（每 SM 16 个）压制，占用不足、延迟隐藏差；block ≥ 64 之后基本持平。实用默认值是 128、256 或 512，避免用 32 这种过小的 block。这个最优区间对访存密集 kernel 普遍适用；计算密集或重 shared memory 的 kernel 要结合 occupancy 单独权衡，见 第 06 篇。

五、错误检查：别让错误静默

CUDA 的 API 和 kernel 启动失败常常不抛异常，而是返回错误码。kernel 内的非法访问可能要到下次同步才暴露。生产代码必须检查：

#define CUDA_CHECK(call) do {                                  \
    cudaError_t _e = (call);                                   \
    if (_e != cudaSuccess) {                                   \
        fprintf(stderr, "CUDA error %s at %s:%d\n",            \
                cudaGetErrorString(_e), __FILE__, __LINE__);   \
        abort();                                               \
    }                                                          \
} while (0)

vadd<<<blocks, threads>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaGetLastError());      // 捕获启动配置错误（如 block 太大）
CUDA_CHECK(cudaDeviceSynchronize()); // 捕获 kernel 执行期错误

cudaGetLastError() 抓启动参数错误，cudaDeviceSynchronize() 等 kernel 跑完并抓执行期错误（如越界）。开发期还应配合 compute-sanitizer 做越界和竞态检查，见 调试篇。这里用的是 runtime API（cudaError_t / cudaSuccess）的教学写法；本系列的实测脚本走 driver API（cuda-python + NVRTC），对应的错误检查是把 CUresult 与 CU_SUCCESS 比较、用 cuGetErrorString 取信息，思路一致，只是符号不同。

六、完整流程：host 侧五步

把一个 kernel 真正跑起来，host 侧的标准流程是：

flowchart LR
  A["1. cudaMalloc<br/>分配 device 内存"] --> B["2. cudaMemcpy H2D<br/>拷入输入"]
  B --> C["3. kernel&lt;&lt;&lt;grid,block&gt;&gt;&gt;<br/>启动计算"]
  C --> D["4. cudaMemcpy D2H<br/>拷回结果"]
  D --> E["5. cudaFree<br/>释放"]

注意第 2、4 步的 host↔︎device 拷贝走 PCIe，往往比 kernel 本身慢得多。真实工作负载里要尽量让数据留在 device、多个 kernel 串起来跑，避免来回搬。把多个操作融合成一个 kernel 减少中间数据落地，是 kernel fusion 篇 的主题。

七、小结与下一步

• kernel 用 blockIdx*blockDim+threadIdx 算全局索引，必须做边界检查；grid-stride loop 让 kernel 与数据规模解耦，是推荐默认写法。
• __syncthreads() 是 block 内栅栏，用于共享中间结果；不能放在部分线程才进入的分支里，也不能跨 block。
• block 大小避免用 32，128–512 是安全默认；过小的 block 因占用不足实测慢约 9%。
• 必须检查 CUDA 错误，否则失败会静默。

会写 kernel 之后，下一篇进入第一个真正的性能主题——访存优化：合并访问、bank conflict 与对齐，看访问模式如何决定你能拿到多少带宽。