3. CUDA编程模型与内核启动

CUDA 编程模型的核心思想是“主机端负责调度，设备端负责并行计算”。一个典型的 CUDA 程序由主机代码（Host）和设备代码（Device）组成：主机端负责分配显存、拷贝数据、配置并启动内核；设备端则由成千上万个线程并行执行核函数。理解核函数定义、线程索引计算、网格与块维度配置、启动参数、流与事件、错误处理以及同步与异步机制，是编写正确且高效 CUDA 算子的基础。

核函数定义与线程索引

核函数使用 __global__ 限定符声明，它在设备端执行、由主机端调用，返回值必须为 void。核函数内部通过内置变量 threadIdx、blockIdx、blockDim、gridDim 计算全局线程索引。以三维线程块为例，全局索引计算如下：

__global__ void addKernel(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

这里 blockIdx.x 是当前线程块在网格中的索引，blockDim.x 是每个块的线程数，threadIdx.x 是线程在块内的索引。三者组合得到全局唯一索引 idx。边界判断 if (idx < n) 必不可少，因为网格总线程数通常大于实际数据量，越界访问会导致未定义行为甚至程序崩溃。

对于二维或三维数据，索引计算需要扩展到对应维度。例如处理 width × height 的图像时，常用二维块配置：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int offset = y * width + x;
if (x < width && y < height) {
    // 处理像素 (x, y)
}

这种“行优先”的线性化方式与 C/C++ 多维数组内存布局一致，是 CUDA 中最常见的索引约定。需要特别注意的是，threadIdx 和 blockIdx 都是无符号整数类型，在计算偏移量时若涉及减法或负数，应先转换为有符号类型，避免下溢。

网格与块维度配置

网格（Grid）和线程块（Block）的维度配置直接决定并行粒度和硬件资源利用率。每个线程块最多 1024 个线程（计算能力 2.x 及以上），块内线程以 32 个为一组构成一个 warp，warp 是 SM 调度的基本单位。因此，块大小通常设置为 32 的整数倍，如 128、256、512，以充分利用硬件。

网格大小由数据量和块大小共同决定，常用向上取整公式：

int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
addKernel<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

例如 n = 1000、blockSize = 256 时，gridSize = 4，总线程数为 1024，其中最后 24 个线程因 idx >= n 被边界判断过滤。块大小的选择需要权衡：块太小会导致块数量过多、调度开销上升；块太大则可能因寄存器或共享内存限制导致占用率下降。一般建议从 128 或 256 开始调优。

启动配置参数与流、事件

内核启动语法为 kernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args...)。其中 sharedMemBytes 指定动态共享内存大小（默认 0），stream 指定执行流（默认 0 表示默认流）。使用非默认流可以实现内核与数据传输的重叠：

cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(d_a, d_b, d_c, n);
cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
cudaStreamDestroy(stream);

事件（Event）用于测量内核耗时或实现流间同步。典型用法是记录起始和结束事件，再通过 cudaEventElapsedTime 获取毫秒级时间差。事件也可配合 cudaStreamWaitEvent 实现跨流依赖，避免不必要的全局同步。

错误处理与同步机制

CUDA 的异步特性意味着内核启动后主机端不会立即知道设备端是否出错。因此每次内核启动后应检查错误：

kernel<<<grid, block>>>(args...);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}

对于调试阶段，可在启动后调用 cudaDeviceSynchronize() 强制等待设备完成，从而捕获运行时错误。但生产代码中应避免频繁全局同步，改用流同步或事件同步以保持异步流水线。

同步与异步的边界需要明确：cudaMemcpy 是同步的，会阻塞主机直到拷贝完成；cudaMemcpyAsync 是异步的，需要配合流同步才能保证数据就绪。内核启动本身是异步的，主机端调用后立即返回。理解这些语义，才能正确编排“拷贝—计算—拷贝”的执行顺序，避免数据竞争和结果错误。

实践中的配置与排错要点

在实际算子开发中，建议遵循以下检查顺序：第一，确认 cudaMalloc 和 cudaMemcpy 的返回值均为 cudaSuccess；第二，内核启动后立即调用 cudaGetLastError 捕获启动配置错误（如块大小超过 1024）；第三，调试阶段使用 cudaDeviceSynchronize 定位运行时错误；第四，性能测试时改用事件计时，避免同步干扰流水线。

网格与块维度的选择没有唯一最优解，需要结合数据规模、寄存器压力、共享内存用量和 SM 占用率综合调优。一个实用的起点是：块大小取 256，网格大小按数据量向上取整，然后通过调整块大小（128/256/512）和每线程处理元素数（1/2/4）来观察性能变化。对于访存密集型算子，增大每线程处理元素数有助于提高指令级并行和访存效率；对于计算密集型算子，则更应关注寄存器使用和占用率。

同步与异步的合理编排是 CUDA 性能优化的关键。默认流（stream 0）具有隐式同步语义，会阻塞其他流的执行；使用非默认流可以构建多流流水线，让数据传输与计算重叠。但多流编程也增加了复杂度：必须确保同一流内操作有序，跨流依赖用事件显式声明，否则会出现数据竞争或结果不确定。对于初学者，建议先在单流模式下保证正确性，再逐步引入多流优化。