3. CUDA编程模型入门：Host与Device的概念，Kernel函数的定义与调用，理解<<>>执行配置的物理意义

好，咱们正式开始接触CUDA编程模型的核心了。

说实话，我第一次接触CUDA的时候，最困惑的就是这个Host和Device的概念。明明都是代码，怎么还分两套？后来踩过几次坑才明白——这其实是理解GPU编程的钥匙。

3.1 Host与Device：CPU和GPU的分工

简单来说：

• Host：就是CPU及其内存。负责控制流程、逻辑判断、I/O操作。
• Device：就是GPU及其显存。负责大规模并行计算。

你想想看，CPU擅长处理复杂、串行的任务。GPU呢？它擅长做简单、重复、并行的计算。所以CUDA程序天然就是"双核"的——CPU当总指挥，GPU当苦力。

我在项目中遇到过一个问题：有个同事把所有逻辑都写在GPU上，结果调试了三天。后来我帮他改成Host控制流程、Device只做矩阵运算，半天就搞定了。说白了，别让GPU做它不擅长的事。

3.2 Kernel函数：在GPU上运行的代码

Kernel函数，就是你在GPU上执行的函数。它有几个关键特征：

• 用 __global__ 修饰符声明
• 返回类型必须是 void
• 由Host调用，在Device上执行
• 所有线程执行同一份代码（SIMT模型）

来看一个最简单的例子：

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

int main() {
    // 调用Kernel函数
    hello_from_gpu<<<1, 5>>>();
    
    // 等待GPU执行完成
    cudaDeviceSynchronize();
    return 0;
}

这里有个细节：printf 在Kernel里能用，但只在计算能力2.0以上的GPU才支持。嗯，这个坑我踩过——在旧卡上跑，啥输出都没有，还以为程序挂了。

3.3 <<>>：执行配置的物理意义

这个语法，说实话，是CUDA里最让人迷惑的地方之一。我第一次看到 <<<1, 5>>> 的时候，心想：这什么鬼？三个尖括号？

其实它的物理意义非常直观：

• grid：线程格，包含多个线程块
• block：线程块，包含多个线程
• thread：线程，最小的执行单元

你可以这样理解：

• 一个Kernel启动，相当于一个grid
• grid由多个block组成
• 每个block包含多个thread

来看一个更具体的例子：

// 启动一个包含2个block、每个block有4个线程的Kernel
my_kernel<<<2, 4>>>();

// 等价于启动了 2 × 4 = 8 个线程
// 线程索引如下：
// Block 0: thread 0, 1, 2, 3
// Block 1: thread 4, 5, 6, 7

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

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

变量	含义	示例值
threadIdx.x	当前线程在block内的索引	0, 1, 2, ..., blockDim.x-1
blockIdx.x	当前block在grid内的索引	0, 1, 2, ..., gridDim.x-1
blockDim.x	每个block包含的线程数	由<<<grid, block>>>的第二个参数决定
gridDim.x	grid包含的block数	由<<<grid, block>>>的第一个参数决定

有了这些变量，你就可以计算全局线程ID：

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

这个公式我建议你背下来。我在项目中至少写过几百次。它把二维的block-grid结构，映射成了一维的线性索引。

3.5 一个完整的例子：向量加法

咱们把上面这些概念串起来，写一个完整的向量加法：

#include <stdio.h>

// Kernel函数：向量加法
__global__ void vec_add_kernel(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 = 1000;
    size_t size = n * sizeof(float);
    
    // 1. 在Host上分配内存
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    
    // 初始化数据...
    
    // 2. 在Device上分配显存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 3. 把数据从Host拷贝到Device
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 4. 配置执行参数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    // 5. 启动Kernel
    vec_add_kernel<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 6. 把结果拷回Host
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 7. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

这个模式你记住了：分配→拷贝→执行→拷回→清理。我刚开始写CUDA程序时，经常忘记 cudaMemcpy 的方向，结果GPU算了个寂寞。

3.6 执行配置的物理意义再思考

你可能会问：为什么要搞grid和block两层结构？直接启动N个线程不行吗？

原因有两点：

• 硬件限制：一个block内的线程共享shared memory，且数量有限（通常最多1024个）
• 调度效率：GPU以block为单位调度到SM上，block越多，越容易隐藏延迟

说白了，grid是逻辑上的组织方式，block是硬件上的调度单位。你设计执行配置时，要同时考虑这两个层面。

我记得有一次优化一个图像处理程序，把block大小从128改成256，性能提升了30%。原因就是每个SM能同时容纳更多线程，更好地隐藏了内存访问延迟。

3.7 本章小结

• Host是CPU，Device是GPU，各司其职
• Kernel函数用 __global__ 声明，在Device上执行
• <<<grid, block>>> 定义了线程的组织结构
• 用 threadIdx、blockIdx 等内置变量定位线程
• 向量加法是CUDA的"Hello World"，掌握它就掌握了基本流程

下一章，咱们会深入探讨线程同步和shared memory的使用。到时候你会发现，理解了grid和block，后面的内容就顺理成章了。