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

好，咱们正式开始动手写CUDA代码了。这一章是地基，地基不稳，后面全白搭。我见过不少同学，一上来就调库、跑模型，结果出了bug连在哪查都不知道——说白了，就是没搞懂CPU和GPU怎么配合的。

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

在CUDA的世界里，CPU叫主机（Host），GPU叫设备（Device）。主机负责逻辑控制、数据准备、串行任务；设备负责大规模并行计算。

你想想看，CPU就像项目经理，GPU就像一群搬砖工。项目经理不搬砖，但他得告诉工人们：砖在哪、搬到哪、怎么搬。搬砖工只管闷头干，效率极高。

核心原则：

主机管理内存分配、数据传输、核函数启动
设备执行核函数，处理海量数据
主机和设备有各自独立的内存空间

我在项目中遇到过最典型的坑：新手把主机指针直接传给核函数，结果GPU访问不到，程序直接崩。记住，主机内存和设备内存是两套系统，必须显式拷贝。

2.2 核函数（Kernel）编写

核函数就是在GPU上执行的函数。写法上有个标志性的东西——__global__修饰符。

// 一个最简单的核函数
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

注意几个要点：

__global__ 表示该函数在设备上执行，从主机调用
核函数返回类型必须是 void
核函数不能有可变参数（...）
核函数内部可以调用设备函数（__device__），但不能调用主机函数

我的习惯：核函数命名加个 _kernel 后缀，比如 matmul_kernel。这样一看就知道是GPU代码，方便维护。

2.3 线程层次结构：grid、block、thread

这是CUDA最核心的概念，没有之一。我当年学的时候，花了整整一周才真正理解这三个东西的关系。

简单来说：

thread：最小的执行单元，每个线程执行核函数的一份拷贝
block：一组线程的集合，可以协作（共享内存、同步）
grid：一组block的集合，构成整个计算任务

用个比喻：

thread = 一个工人
block = 一个班组（共享工具、可以互相帮忙）
grid = 整个工地（所有班组一起干活）

每个线程都有唯一的ID，通过内置变量获取：

threadIdx.x：线程在block内的索引
blockIdx.x：block在grid内的索引
blockDim.x：每个block有多少线程
gridDim.x：每个grid有多少block

全局线程ID的计算公式：

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

嗯，这里要注意：block内的线程数量有限制，最大1024（取决于GPU架构）。我刚开始写的时候，一上来就设了2048，结果程序跑不起来，查了半天才发现是block太大。

2.4 <<< >>> 语法：启动核函数

核函数写好了，怎么让它跑起来？用三括号语法：

// 定义grid和block的维度
int threads_per_block = 256;
int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;

// 启动核函数
vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);

三括号里的参数：

第一个参数：grid的维度（block数量）
第二个参数：block的维度（thread数量）
可选第三个参数：共享内存大小（字节）
可选第四个参数：流（stream）

我曾经踩过的坑：核函数启动是异步的！主机不会等GPU算完就继续执行下一行。如果你紧接着就拷贝结果回主机，拿到的可能是垃圾数据。必须加 cudaDeviceSynchronize() 或者用 cudaMemcpy 隐式同步。

2.5 完整示例：向量加法

把上面所有知识点串起来，写一个完整的向量加法：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vec_add_kernel(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // 1048576个元素
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }

    // 2. 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. 拷贝数据到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vec_add_kernel<<<blocks, threads>>>(d_a, d_b, d_c, n);

    // 5. 拷贝结果回主机
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }

    // 7. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

这个流程你记牢了：分配主机内存 → 分配设备内存 → 拷贝到设备 → 启动核函数 → 拷贝回主机 → 清理。几乎所有CUDA程序都是这个套路。

2.6 知识体系结构图

下面这张图把本章的核心逻辑串起来了，建议你多看几遍：

我的建议：刚开始学，别急着调优。先把主机-设备的数据流跑通，把核函数写对，把线程ID算明白。性能优化是后面的事。

好了，这一章的内容就到这。记住：主机管控制，设备管计算，核函数是桥梁，线程ID是灵魂。把这四个点吃透，后面学起来就顺了。

想要更多资料，联系可可

📚 课程目录

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

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

2.2 核函数（Kernel）编写

2.3 线程层次结构：grid、block、thread

2.4 <<< >>> 语法：启动核函数

2.5 完整示例：向量加法

2.6 知识体系结构图