CUDA核函数入门：从零开始理解并行计算

说实话，很多初学者第一次接触CUDA时，最困惑的就是核函数到底是个什么东西。我当年也是这样，看了半天文档，感觉就是一堆奇怪的语法。今天咱们就把这个事儿彻底讲明白。

什么是核函数？

核函数（Kernel Function），说白了就是在GPU上执行的函数。它和普通的C++函数最大的区别在于——它会被成千上万个线程同时执行。

// 一个最简单的核函数
__global__ void hello_from_gpu() {
    printf("Hello from thread %d\n", threadIdx.x);
}

int main() {
    // 启动1个网格，每个网格1个块，每个块10个线程
    hello_from_gpu<<<1, 10>>>();
    cudaDeviceSynchronize();
    return 0;
}

嗯，这里要注意：cudaDeviceSynchronize() 不能少。我刚开始写CUDA时就吃过这个亏——核函数是异步执行的，不加同步，主线程跑完了GPU还没开始干活呢。

线程层次结构：网格、块与线程

CUDA的线程组织方式，其实就是一个三层结构。我个人习惯把它想象成一个军队：

• 网格（Grid）：整个任务，相当于一个军团
• 块（Block）：任务的一个子集，相当于一个连队
• 线程（Thread）：最小的执行单元，相当于一个士兵

为什么要有这种层次？我在项目中遇到过一个问题：如果所有线程都平铺开，那管理起来太乱了。而且GPU硬件本身就有层级——线程束（warp）是32个线程一组，块内的线程可以共享内存，块之间则不行。

线程索引计算：找到你的位置

每个线程都需要知道自己该处理哪部分数据。CUDA提供了几个内置变量：

变量名	含义	取值范围
threadIdx	线程在块内的索引	0 到 blockDim - 1
blockIdx	块在网格内的索引	0 到 gridDim - 1
blockDim	块的大小（线程数）	由调用时指定
gridDim	网格的大小（块数）	由调用时指定

计算全局线程ID的公式很简单：

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

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

你想想看，这个公式其实就是一个线性化过程。把二维的网格和块，映射到一维的连续地址空间上。

实战：向量加法

咱们写个完整的例子，把上面这些概念串起来：

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查，防止越界
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

int main() {
    int n = 1000000;
    float *d_a, *d_b, *d_c;
    
    // 分配显存
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));
    
    // 配置执行参数
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    
    // 启动核函数
    vec_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);
    
    cudaDeviceSynchronize();
    // ... 拷贝结果回主机 ...
}

这里有个细节：grid_size = (n + block_size - 1) / block_size 这个写法，是为了保证即使数据量不是块大小的整数倍，也能覆盖所有元素。多出来的线程会在核函数里被 if (tid < n) 过滤掉。

显存释放：别让GPU内存泄漏

写CUDA程序，最怕的就是显存泄漏。GPU的内存是有限的，你分配了不释放，跑几次程序就炸了。

// 正确的释放方式
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

我个人习惯在写 cudaMalloc 的同时就把 cudaFree 写好，就像写括号一样，成对出现。这样就不会漏掉。

知识体系总览

下面这张图，把咱们今天讲的核心内容串起来了：

这张图展示了网格、块和线程的关系。每个块里有多个线程，线程通过 blockIdx 和 threadIdx 的组合来定位自己。

总结一下今天的关键点

• 核函数用 __global__ 声明，返回 void
• 线程组织：Grid → Block → Thread，三层结构
• 全局索引 = blockIdx * blockDim + threadIdx
• 块大小建议32的倍数，常用128、256、512
• 显存分配和释放要成对出现，避免泄漏

说实话，这些概念刚开始可能觉得有点绕。但只要你动手写几个例子，跑一跑，很快就上手了。我在带新人时发现，最容易出问题的地方就是索引计算和显存管理。这两个坑踩过去，后面就顺了。

好了，今天就到这儿。记住：多写代码，多调试，CUDA没那么神秘。