3、CUDA编程基础：核函数(Kernel)概念、线程层次结构(grid/block/thread)、<<<>>>语法详解

好，咱们今天聊点硬核的。CUDA编程里最核心、最绕不开的东西——核函数（Kernel）。

说白了，核函数就是跑在GPU上的函数。你平时写的C/C++函数，是在CPU上串行执行的。而核函数，是让成千上万个GPU线程同时去执行的那段代码。我刚开始接触时，总觉得这玩意儿跟普通函数没啥区别，不就是加个__global__修饰符嘛。直到第一次调试，发现核函数里不能随便调用printf（老版本不行），也不能用malloc，才意识到事情没那么简单。

3.1 核函数长什么样？

先看个最简单的例子：

// 定义一个核函数
__global__ void hello_from_gpu() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    // 调用核函数，启动10个线程
    hello_from_gpu<<<1, 10>>>();
    
    // 同步，确保GPU执行完毕
    cudaDeviceSynchronize();
    return 0;
}

看到那个<<<1, 10>>>了吗？这就是CUDA里最标志性的语法。它告诉运行时系统：我要启动1个线程块（block），每个块里有10个线程（thread）。

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

你想想看，GPU里动辄几千个核心，怎么组织这些线程？CUDA用了三层结构：

• Thread（线程）：最基础的执行单元
• Block（线程块）：一组线程，可以协作（共享内存、同步）
• Grid（网格）：一组Block，构成整个任务

我习惯用军队来类比：Thread是士兵，Block是班，Grid是连。士兵之间可以互相配合（同一个Block内），但班与班之间相对独立。

层级	最大维度	最大数量	可访问资源
Thread	3 (x,y,z)	1024/Block	寄存器、局部内存
Block	3 (x,y,z)	65535/Grid (各维度)	共享内存、同步
Grid	3 (x,y,z)	取决于GPU	全局内存

3.3 <<<>>> 语法详解

这个三括号语法，其实是CUDA的扩展。它的完整形式是：

kernel_name<<<grid_dim, block_dim, shared_mem_size, stream>>>(args...);

四个参数：

1. grid_dim：网格维度，可以是整数或dim3类型。比如<<<2, 256>>>表示2个Block。
2. block_dim：块维度，每个Block里的线程数。比如<<<1, 512>>>表示512个线程。
3. shared_mem_size：动态分配的共享内存大小（字节），默认0。
4. stream：CUDA流，用于异步执行，默认0。

实际项目中，我经常用dim3来处理二维或三维数据：

// 处理1024x1024的图像
dim3 grid(32, 32);      // 32x32 = 1024个Block
dim3 block(32, 32);     // 32x32 = 1024个线程/Block

image_kernel<<<grid, block>>>(d_input, d_output, width, height);

这里每个线程处理一个像素。为什么这么分？因为总线程数 = 1024 * 1024 = 1,048,576，刚好覆盖所有像素。你想想看，如果图像是1920x1080，该怎么分？

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

在核函数内部，CUDA提供了几个内置变量，帮你定位当前线程：

• threadIdx：当前线程在Block内的索引（dim3类型）
• blockIdx：当前Block在Grid内的索引
• blockDim：Block的维度（线程数）
• gridDim：Grid的维度（Block数）

计算全局线程ID的公式：

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

// 二维情况
int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
int tid_y = blockIdx.y * blockDim.y + threadIdx.y;

3.5 实战：向量加法

来，咱们写个完整的例子：

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

int main() {
    int n = 1000000;
    size_t size = n * sizeof(float);
    
    // 分配主机内存
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);
    
    // 初始化数据...
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 配置核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 拷贝结果回CPU
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 清理...
    return 0;
}

注意那个if (idx < n)的边界检查。为什么需要？因为blocks_per_grid * threads_per_block可能大于n，多出来的线程什么都不做。我曾经忘了加这个判断，结果数组越界，查了半天bug。

3.6 避坑指南

嗯，今天就到这儿。核函数是CUDA的基石，理解透了后面学共享内存、流、TensorRT都会轻松很多。下一章咱们聊聊内存管理——为什么cudaMalloc和malloc不能混用。