2. CUDA编程模型：从Kernel到内存，一个都不能少

好，咱们进入正题。CUDA编程模型，说白了就是一套规则——告诉GPU怎么干活、数据放哪儿、怎么搬数据。我刚开始接触CUDA时，觉得这东西不就是写个函数嘛，结果一跑就崩。后来才明白，不理解这套模型，写出来的代码就是碰运气。

2.1 Kernel函数定义：GPU的入口函数

Kernel函数，就是跑在GPU上的函数。定义它有个固定套路：

__global__ void myKernel(float *d_data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        d_data[idx] = d_data[idx] * 2.0f;
    }
}

注意几个关键点：

• __global__ 修饰符——告诉编译器这个函数是给GPU用的
• 返回类型必须是 void——你不能从GPU直接返回一个值
• 调用时用 <<<grid, block>>> 语法——指定线程组织方式

你可能会问：为什么不能直接传一个普通数组进去？嗯，这里有个坑——Kernel函数只能访问设备内存（GPU显存）。你传个CPU指针进去，它根本找不到地址。

2.2 线程层次结构：Grid/Block/Thread

CUDA的线程组织方式，我管它叫"三层嵌套"：

• Thread（线程）：最基础的计算单元，每个线程执行相同的Kernel代码
• Block（线程块）：一组线程的集合，共享一块共享内存（Shared Memory）
• Grid（网格）：所有Block的集合，对应一次Kernel调用

用一张图来理解：

每个线程都有一个唯一的ID，通过 threadIdx、blockIdx、blockDim 计算出来。我常用的公式：

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

这个公式，你写100个Kernel，99个都会用到它。

2.3 内存层次结构：速度与容量的博弈

CUDA的内存模型，说白了就是"越靠近计算单元，越快但越小"。我整理了一张表：

你想想看，全局内存访问一次要400个周期，而寄存器只要1个周期。这差距有多大？所以优化的核心思路就是：尽量把数据往快的存里放。

2.4 主机与设备交互：数据搬运的艺术

主机（CPU）和设备（GPU）之间的数据交互，说白了就是"搬数据"和"下指令"。常用的API：

• cudaMalloc：在设备上分配显存
• cudaMemcpy：在主机和设备之间拷贝数据
• cudaFree：释放显存

一个典型的流程：

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

// 2. 分配设备内存
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 3. 把数据从主机拷贝到设备
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 4. 启动Kernel
myKernel<<<gridSize, blockSize>>>(d_data, N);

// 5. 把结果从设备拷贝回主机
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// 6. 清理
cudaFree(d_data);
free(h_data);

还有一个容易被忽略的点：cudaDeviceSynchronize()。这个函数会等待所有Kernel执行完毕。我一般在调试时会在每个Kernel后面加一句，确保不会出现"数据还没算完就开始拷贝"的情况。

嗯，CUDA编程模型的核心就这些。理解了这个模型，你写出来的代码至少不会"跑不起来"。至于怎么跑得快，那是后面章节要聊的话题。