CUDA核函数编写入门

可视化演示

观察粒子在马尔可夫链状态间的随机转移过程，理解转移概率矩阵与稳态分布的形成。

核函数概述：让GPU替你打工

装修公司接了一个大工程：给一栋10层大楼装灯泡，每层100个房间，一共1000个灯泡。老板一个人装，得装三天三夜。但他手下有100个工人，只要写一份施工手册，让每个工人照着手册去装自己负责的房间，一上午就能干完。GPU就是这支施工队，而那份手册，就是CUDA核函数。

在CUDA里，核函数（kernel）是一个看上去很普通的C++函数，唯一的不同是它前面多了一个印章__global__。这个印章告诉编译器：“把这份函数编译成GPU工人们能执行的代码，而不是CPU的普通指令。”正因为如此，核函数运行在GPU（设备端）上，但由CPU（主机端）负责“下发”。

GPU的厉害之处在于“人多”：一个核函数被启动后，GPU会创建成百上千个线程（工人），每个线程执行同样的代码，但处理不同的数据。比如装灯泡的活，1000个灯泡对应1000个线程，每个线程装一个。为了保证每个工人知道该装哪个灯泡，我们需要给线程编号——这就是后面要讲的线程索引。

所以，核函数就是施工手册，GPU就是施工队。接下来我们要学的，就是怎么把手册写好——包括封面盖什么章、工人怎么编号、队伍怎么分组、口令怎么喊、有什么禁忌，以及出错了怎么算账。

声明限定符：手册上的三种印章

写施工手册前，先要搞清楚这份手册给谁看。有的手册是给施工队看的，有的手册是施工现场内部传阅的，还有的是总部办公室存档的。CUDA用三种“印章”来区分代码的运行位置：__global__、__device__、__host__。

__global__是核函数的专属印章。盖了这枚章的函数在GPU上执行，但由CPU调用。它就像老板站在楼下向全队广播的施工总则——发号施令的是老板，卖力干活的是工人。

__device__盖着的函数，是工人们内部使用的工具函数，只能在核函数里被调用。__host__则是普通CPU函数，我们平时写的main函数就是一个__host__函数。还有一种组合写法：__host__ __device__，表示这个函数CPU和GPU都能跑——就像“干湿两用”的多功能工具。

盖好__global__这枚印章，你的施工手册才算是被GPU施工队认可的正式文件。漏了印章，编译器会直接报错——施工队压根不认这份手册。

线程层次结构：工人的工号体系

工地上有1000个工人，怎么让每个人准确找到自己该装的那个灯泡？答案是工号。CUDA把工人分成了两级组织：整个施工队叫grid（网格），施工队里又分成若干个小组叫block（线程块）。每个工人有两个编号：在小组内的编号叫threadIdx.x，所在小组的编号叫blockIdx.x。

为什么需要两级编号？因为GPU的硬件架构就是这样设计的：线程块（block）可以被调度到不同的流处理器上执行，而同一个块内的线程可以通过共享内存快速协作。就像施工队分成小队，小队内部可以快速互相借工具，不同小队之间的交流则要慢得多。

在编程时，最常用的操作是把两个编号换算成一个唯一的全局编号（gid）。假设每组4人，你是第2组的第3号工人，那么你的全局编号就是 2×4+3=11。每一位工人的全局编号都不一样，这就是他认领数据房间的“门牌号”。

别小看这套工号体系。GPU能让上千个工人并行干活，靠的就是这套“对号入座”的机制——每个工人拿着工号，找到自己的灯泡，绝不跑错房间。

维度配置：施工队怎么分组

老板在广播里喊“开工”之前，得先想清楚：施工队分成多少组？每组多少人？这就是grid和block的维度配置。

CUDA提供了dim3类型来定义三维尺寸。处理一维数组时，只用到x方向；处理图片（二维）时，可以同时指定x和y；处理三维体素数据时，x、y、z三个方向一起用。不过入门阶段，我们几乎都只用一维x方向。

配置的核心思想是“人多力量大，但要够用”。假设有N个数据要处理，每个小组安排blockDim.x个工人，那么小组数就是 gridDim.x = ceil(N / blockDim.x)。向上取整是为了保证最后剩几个数据也有工人处理。例如N=1000，blockDim.x=128，gridDim.x=ceil(7.8125)=8，总线程数=8×128=1024，比1000稍多，多出的线程做“空转”或边界检查跳过。

分组分得好，工人不白跑。维度配置就像装修前规划楼层分工：一层楼派一组人，组里再分房间，分配得当，工程就顺。

启动与限制：开工口令与手册铁律

手册写好、队伍分好，老板怎么开工？CUDA的答案是：一行特殊的函数调用。老板在CPU端写下这行“口令”，GPU施工队听到口令就立刻行动。

这行口令的语法是：kernel_name<<<grid, block>>>(参数列表)。其中<<<grid, block>>>就是启动配置——grid是小组数，block是每组人数。这行代码是异步的：CPU喊完口令后不会站在旁边等工人干完，而是继续干自己手头的其他活。

施工手册上还有几条“铁律”：核函数不能有返回值（必须void）；参数按值传递，传指针时传的是显存地址；不能递归调用；不能是类的普通成员函数（可以是static成员函数）；由于每条线程的栈空间和寄存器有限，核函数里也不能开太大的局部数组。

开工口令一下，工人们就冲出去了。但手册上写得很清楚：工人干完活不许跑回办公室口头汇报——这就是核函数没有返回值的底层逻辑。工人把结果直接写进现场（显存），老板要数据时再去现场取（cudaMemcpy）。

核函数调试：监工的上岗指南

工人多了，出错也难免。可能某个工人把灯泡装在走廊，可能有人数错房间号。作为老板，你需要一套工具来定位问题。

入门阶段最直接的监工是printf。没错，核函数里可以写printf！它会打印当前线程的信息。比如printf("block %d, thread %d\n", blockIdx.x, threadIdx.x);就能看到每个线程的执行情况。但注意，GPU的printf是异步的，很可能程序跑完了输出还没刷新，这时候在核函数启动后调用cudaDeviceSynchronize()可以强制同步。

更专业的工具是cuda-gdb。先用nvcc -g -G编译，再用cuda-gdb启动程序，就能像调试普通程序一样打断点、单步执行、查看每个线程的变量值。此外，几乎每个CUDA API调用都应该检查返回值：cudaGetLastError()可以告诉你最近一次调用是否出错。

监工的职责是让问题在萌芽期暴露。核函数里的printf、CPU端对错误码的检查、cuda-gdb里的断点，就是工地上的三班监工——24小时盯着施工队，bug一冒头就揪出来。

---

AI 应用场景

Python 演示

class KernelSimulator:    """Simulate CUDA kernel thread organization and vector add."""     def __init__(self, block_dim):        self.block_dim = block_dim  # threads per block     def calc_grid_dim(self, n):        # ceil(n / block_dim)        return (n + self.block_dim - 1) // self.block_dim     def global_id(self, block_idx, thread_idx):        # gid = blockIdx.x * blockDim.x + threadIdx.x        return block_idx * self.block_dim + thread_idx     def vector_add(self, a, b):        n = len(a)        grid_dim = self.calc_grid_dim(n)        c = [0] * n        print("gridDim=%d, blockDim=%d, total threads=%d" % (grid_dim, self.block_dim, grid_dim * self.block_dim))        for bi in range(grid_dim):            for ti in range(self.block_dim):                gid = self.global_id(bi, ti)                if gid < n:                    c[gid] = a[gid] + b[gid]                    print("block %d thread %d -> gid %d: %d + %d = %d" % (bi, ti, gid, a[gid], b[gid], c[gid]))                else:                    print("block %d thread %d -> gid %d out of range, skip" % (bi, ti, gid))        return c  if __name__ == "__main__":    sim = KernelSimulator(block_dim=4)    a = [1, 2, 3, 4, 5, 6, 7, 8, 9]    b = [10, 20, 30, 40, 50, 60, 70, 80, 90]    result = sim.vector_add(a, b)    print("result:", result)

1. grid_dim = (n + self.block_dim - 1) // self.block_dim等价于 ceil(N / blockDim)，也就是向上取整算出需要开多少个block。
2. gid = block_idx * self.block_dim + thread_idx这就是CUDA的核心索引公式：全局线程编号 gid = blockIdx.x × blockDim.x + threadIdx.x。
3. if gid < n:边界检查。总线程数可能比数据多，超出的线程什么都不做，保证不越界访问。
4. c[gid] = a[gid] + b[gid]每个线程处理自己责任范围内的数据，这正是GPU并行加速的微观缩影。

常见误区

面试题