CUDA编程基础：主机与设备、核函数编写、线程层次结构

大家好，我是你们的老朋友。今天咱们聊聊CUDA编程最核心的几个概念。说实话，很多新手一上来就被各种术语搞懵了——主机、设备、核函数、Grid、Block、Thread……别急，我当年也是这么过来的。咱们一个一个拆开讲。

主机与设备：CPU和GPU的分工

先搞清楚一个基本事实：CUDA程序里，CPU是“主机”（Host），GPU是“设备”（Device）。

主机负责逻辑控制、数据准备、任务调度。设备负责大规模并行计算。说白了，CPU是项目经理，GPU是搬砖工头，带着成千上万个线程一起干活。

我在项目中遇到过不少新手，上来就把所有计算都扔给GPU，结果发现性能还不如CPU。为什么？因为数据在主机和设备之间传输是有开销的。你想想看，如果数据量很小，传输时间比计算时间还长，那还不如用CPU算。

举个简单的例子：

// 主机端分配内存
float *h_data = (float*)malloc(N * sizeof(float));

// 设备端分配显存
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 主机到设备的数据传输
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 在设备上执行核函数
myKernel<<<grid, block>>>(d_data);

// 设备到主机的数据传输
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

嗯，这里要注意：cudaMalloc和cudaMemcpy是CUDA编程中最常用的API，一定要记得检查返回值。我曾经因为忘了检查cudaMalloc是否成功，结果在显存不足时程序直接崩溃，排查了半天才发现是显存分配失败。

核函数：在GPU上运行的函数

核函数（Kernel）就是在GPU上执行的函数。它的写法跟普通C函数很像，但有几个关键区别：

• 前面要加__global__修饰符
• 返回类型必须是void
• 不能使用可变参数（比如printf那种）
• 不能是类的成员函数（静态成员函数可以）

我个人的习惯是，核函数命名时加个_kernel后缀，比如vectorAdd_kernel，这样一看就知道是GPU代码。当然这只是个人风格，不强求。

来看一个最简单的核函数：

__global__ void hello_kernel() {
    printf("Hello from thread %d in block %d\n", 
           threadIdx.x, blockIdx.x);
}

调用它的时候，要用<<<grid, block>>>语法指定线程配置：

hello_kernel<<<2, 4>>>();

这行代码的意思是：启动2个Block，每个Block里有4个Thread。总共8个线程并行执行。

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

这是CUDA最核心的概念，也是很多人一开始最晕的地方。我尽量用大白话讲清楚。

想象一下：你有一家公司（GPU），公司里有多个部门（Grid），每个部门里有多个小组（Block），每个小组里有多个员工（Thread）。

具体来说：

• Thread（线程）：最小的执行单元，每个线程执行核函数中的同一段代码，但处理不同的数据
• Block（线程块）：一组线程的集合，Block内的线程可以共享显存（Shared Memory），也可以通过__syncthreads()同步
• Grid（线程网格）：一组Block的集合，一个核函数对应一个Grid

每个线程都有一个唯一的ID，由threadIdx和blockIdx决定。比如：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式我建议你背下来。它把一维的Block和Thread索引映射成全局唯一的线程ID。我刚开始学的时候，每次写核函数都要先写这一行，都快成肌肉记忆了。

为什么要有这种层次结构？说白了是为了硬件效率。GPU的SM（流多处理器）以Block为单位分配计算资源，Block内的线程可以高效协作。你想想看，如果所有线程都平铺开，没有层次结构，那共享内存和同步机制就没法实现了。

<<< >>>语法：配置执行参数

这个语法是CUDA特有的，用来指定核函数启动时的线程配置。格式是：

kernel_name<<<grid_dim, block_dim, shared_mem, stream>>>(args);

四个参数：

1. grid_dim：Grid的维度，可以是dim3类型，比如dim3(2, 3, 1)
2. block_dim：Block的维度，也是dim3类型
3. shared_mem：每个Block动态分配的共享内存大小（字节），默认为0
4. stream：CUDA流，用于异步执行，默认为0（默认流）

举个例子：

// 一维配置
int N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

// 二维配置
dim3 gridDim(32, 32);
dim3 blockDim(16, 16);
matrixMul_kernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, width);

嗯，这里有个小技巧：计算Block数量时，用(N + threadsPerBlock - 1) / threadsPerBlock这个公式，可以保证即使N不能被threadsPerBlock整除，也能覆盖所有数据。我刚开始写的时候经常忘了加这个，结果数组末尾的数据没被处理，查了半天才发现。

知识体系总览

下面这张图总结了本章的核心内容，我画成了SVG，方便你理解整个知识结构：

这张图把主机与设备、核函数、线程层次结构、<<<>>>语法串在了一起。你可以看到，核函数是连接主机和设备的桥梁，而线程层次结构是核函数执行时的组织方式。

最后总结一下：

• 主机负责控制，设备负责计算
• 核函数用__global__修饰，在设备上执行
• 线程按Grid→Block→Thread三层组织
• <<<grid, block>>>配置执行参数

这些概念是CUDA编程的基石。我建议你动手写几个小例子，比如向量加法、矩阵乘法，把今天讲的内容跑一遍。光看是学不会的，得亲手敲代码才行。

好了，今天就到这里。记住，CUDA编程的核心就是理解线程怎么组织、数据怎么流动。把这个搞明白了，后面的优化技巧才能派上用场。