CUDA编程模型：核函数Kernel

可视化演示

展示Grid-Block-Thread层次结构、Warp内存合并访问模式及核函数启动流程，支持三模式切换交互。

核函数声明

老板把“红烧肉怎么做、几分熟出锅”写成一张标准操作卡，盖上主厨专用印章，贴在厨房每个灶台前。这张卡在CUDA里就叫核函数（Kernel），印章上的字就是“``__global__``”。凡是盖了这枚印章的函数，就表示它的执行者不是CPU，而是GPU上每一位并行工作的“厨师线程”。

从菜谱过渡到代码：你要写的核函数与普通C函数长得很像，但有三个显著区别。第一，声明时必须以“``__global__`”开头，告诉编译器“这段代码编译后要放到GPU上去执行”；第二，返回类型必须写成“`void``”，因为成千上万个线程同时执行时，谁也没有资格向CPU返回一个单独的值，所有计算结果必须通过参数指针写回内存；第三，核函数只能由主机端代码调用，不能由另一个核函数直接调用。

打个比方：核函数声明就是“操作卡抬头+印章”。至于卡上具体写什么步骤、派多少人去执行，则由下面几节决定。

执行配置

菜谱写好了，可老板还不能撒手。他需要再写一张排班表：菜谱要复制给几组人？每组多少人？在CUDA里，这张排班表就是调用核函数时的“执行配置”，用一对尖括号“``<<< >>>``”表达，紧跟在函数名之后。

执行配置的完整语法是“``kernel_name<<<grid_size, block_size>>>(args)`”。其中第一个数字“`grid_size`”表示启动多少个线程块（block），相当于厨房分成几个班组；第二个数字“`block_size`”表示每个班组里有多少个线程（thread），相当于每个班组几个厨师。例如“`vectorAdd<<<2, 32>>>(a, b, c, N)``”表示启动2个块、每块32个线程，总共64个线程并行工作。

回到厨房场景：排班表决定了一道菜的复制份数和人手规模。老板只需写一行“``kernel<<<2, 32>>>(...)``”，GPU就会自动安排2个班组、每班32人同时开工，完全不需要老板逐个厨师去叮嘱。

线程索引

排班表贴出去了，可每个厨师怎么知道自己是“第几组的第几号”呢？CUDA给每个线程发了两个身份牌：一个是``blockIdx.x`，表示这个线程所在的块是网格里的第几个块；另一个是`threadIdx.x``，表示线程在块内的第几个位置。两者都从0开始编号。

光有身份牌还不够，程序员需要把“组号+组内编号”换算成一个全局编号，才能对应到数组的下标。换算逻辑很直观：先数一数当前线程前面有几个完整的块，每个块里有``blockDim.x`个线程，加上自己在块内的`threadIdx.x``，就是全场第几个线程。

用一个具体例子：网格有2个块，每块8个线程。第1块（``blockIdx.x=1`）的第5号厨师（`threadIdx.x=5``）的全局索引就是“\(1\times8+5=13\)”。每个厨师只需记住自己两个号码，就能唯一确定全局工位。

网格与块

中央厨房的架构不止一层：整个厨房叫“网格”（grid），里面划分成若干“块”（block），每个块里坐着若干“厨师”（thread）。更妙的是，网格和块都可以是一维、二维或三维的，分别适用于处理数组、矩阵图像或立体数据。

在CUDA中，``gridDim.x`、`gridDim.y`、`gridDim.z`表示网格在三个方向上的块数量；`blockDim.x`、`blockDim.y`、`blockDim.z``表示块在三个方向上的线程数量。把三个方向的人数乘起来，就是网格中的总线程数。

以前面的例子来说，如果启动“``<<<(2,2), (4,4)>>>``”，表示网格有2×2个块，每块4×4个线程，总共64位厨师同时开工。理解了这个嵌套结构，就能明白为什么GPU吞吐量如此惊人——它相当于把一个大厨房分成无数小班组，每个班组独立又并行地运转。

并行映射

排班表和工位号都到位了，接下来最关键的问题：每个厨师到底做哪道菜？在CUDA里，最常见的做法是“一人一道菜”：把大数组均匀切分，让每个线程通过自己的全局索引去访问对应的数组元素。

比如做向量加法“``C[i] = A[i] + B[i]`”，传统CPU写法是写一个循环依次处理每个元素；CUDA则是让第`i`号线程直接去处理第`i``个元素。由于线程们同时开工，整个循环的时间被压缩成一次“并行爆发”。

回到厨房：老板把食材切成N份摆上案台，每个厨师按工位号拿起自己那份。若食材不够分（多出来的厨师），就站在一旁待命，绝不抢占别人的案板。这就是“数据并行”的精髓——人不动，数据动。

编译运行

菜谱和排班表都写好了，要怎么把整套方案落地？CUDA提供了专门的编译器“``nvcc`”，它能把一份混合了CPU代码和GPU代码的`.cu``源文件，自动分离并编译成主机端可执行文件与设备端二进制代码。

运行流程分两条线：主机端（CPU）像餐厅老板一样按顺序执行自己的逻辑，遇到“``kernel<<<...>>>``”时，把核函数交给GPU异步启动，然后立即继续做自己的事；设备端（GPU）的厨师们并行执行核函数体。由于核函数启动是异步的，程序员需要主动调用错误检查函数和同步函数来确认执行结果。

整个流程就像餐厅开业的完整链条：老板（CPU）写菜单、发通知、查验收；厨房（GPU）领菜谱、按排班表并行做菜、把菜摆到出餐口。如果排班表写错或某道菜没熟，老板必须主动检查，否则顾客只能干等。

AI 应用场景

Python 演示

def simulate_cuda_kernel(grid_x, grid_y, block_x, block_y):    """Simulate the thread index calculation of a CUDA kernel."""    print("=" * 60)    print("CUDA Kernel Thread Index Simulation")    print("grid: (%d, %d), block: (%d, %d)" % (grid_x, grid_y, block_x, block_y))    print("=" * 60)     block_size = block_x * block_y    grid_size = grid_x * grid_y    total = grid_size * block_size    print("Total threads = %d * %d = %d" % (grid_size, block_size, total))    print("-" * 60)     count = 0    for block_idx in range(grid_size):        block_x_id = block_idx % grid_x        block_y_id = block_idx // grid_x        for thread_idx in range(block_size):            thread_x_id = thread_idx % block_x            thread_y_id = thread_idx // block_x            global_x = block_x_id * block_x + thread_x_id            global_y = block_y_id * block_y + thread_y_id            if count < 16:                print("thread#%2d: block=(%d,%d) thread=(%d,%d) -> global=(%d,%d)"                      % (block_idx * block_size + thread_idx,                         block_x_id, block_y_id,                         thread_x_id, thread_y_id,                         global_x, global_y))            count += 1    print("... total %d threads" % total)  def vector_add_simulation(n, block_size):    """Simulate vector add with boundary check in CUDA style."""    a = [i * 2 for i in range(n)]    b = [i * 3 for i in range(n)]    c = [-1] * n     grid_size = (n + block_size - 1) // block_size    print("\nVector Add Simulation: N=%d, block_size=%d, grid_size=%d"          % (n, block_size, grid_size))     for i in range(n):        bid = i // block_size        tid = i % block_size        gid = bid * block_size + tid        if gid < n:            c[gid] = a[gid] + b[gid]     print("c[0] = %d + %d = %d" % (a[0], b[0], c[0]))    print("c[%d] = %d + %d = %d" % (n - 1, a[n - 1], b[n - 1], c[n - 1]))  if __name__ == "__main__":    simulate_cuda_kernel(2, 2, 4, 4)    vector_add_simulation(10, 4)

1. simulate_cuda_kernel(2, 2, 4, 4)模拟一个2×2网格、每块4×4线程的核函数启动，打印前16个线程的全局坐标。
2. for block_idx in range(grid_size)遍历所有线程块，相当于GPU硬件按块调度线程。
3. block_x_id = block_idx % grid_x从线性块编号反推出块在网格x方向的坐标，对应``blockIdx.x``。
4. global_x = block_x_id * block_x + thread_x_id核心换算公式：组编号×组大小+组内编号，得到全局x坐标。
5. if gid < n:边界检查，模拟CUDA中“数组长度不是线程数整数倍”时的越界保护。
6. grid_size = (n + block_size - 1) // block_size向上取整计算需要的块数，保证所有元素都有线程覆盖。

常见误区

面试题