一、CUDA编程模型与GPU架构基础

大家好，我是你们这门课的主讲。今天咱们聊聊CUDA编程最核心的东西——编程模型和GPU架构。

说实话，我刚开始接触GPU计算时，也踩过不少坑。那时候总觉得CPU编程那一套直接搬过来就行，结果性能惨不忍睹。后来才明白，GPU和CPU完全是两种思维方式。你想想看，CPU是「一个聪明人干复杂活」，GPU是「一万个普通人干简单活」。这个思维转变，是学好CUDA的第一步。

1.1 GPU并行计算的核心思想

GPU为什么能并行？说白了，它把大量晶体管用在了计算单元上，而不是缓存和控制逻辑。我习惯用一个比喻：CPU像法拉利，跑得快但只能坐俩人；GPU像大巴车，跑得慢点但能拉几十号人。

GPU适合处理什么？数据并行任务。就是同样的操作，作用在不同数据上。比如图像处理里每个像素做同样的滤波，矩阵运算里每个元素做同样的乘法。这种「单指令多数据」的模式，正是GPU的强项。

我在项目中遇到过一位同事，非要用GPU做链表操作。结果呢？性能还不如CPU。为什么？因为链表是串行访问，GPU的并行优势根本发挥不出来。所以记住：GPU不是万能药，它只适合数据密集且可并行的任务。

1.2 CUDA编程的基本流程

CUDA编程其实就三步：搬数据、算数据、搬回来。嗯，听起来很简单，但细节里全是坑。

1.2.1 主机端与设备端的内存管理

CPU这边叫主机（Host），GPU那边叫设备（Device）。两者内存是物理隔离的，不能直接互相访问。所以你得显式地做数据搬运。

我建议你记住这几个API，它们是基本功：

API函数	作用
cudaMalloc	在设备端分配显存
cudaMemcpy	在主机和设备之间拷贝数据
cudaFree	释放设备端显存
cudaMallocHost	分配主机端页锁定内存（提升传输速度）

这里有个坑：cudaMalloc分配的是显存，不是内存。我曾经见过有人忘了释放显存，结果跑几次程序后GPU直接OOM。所以养成好习惯：谁分配，谁释放。

// 典型的内存管理流程
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));  // 设备端分配

float *h_data = (float*)malloc(N * sizeof(float));  // 主机端分配
// ... 填充h_data ...

cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);  // 搬到设备

// ... 在GPU上计算 ...

cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);  // 搬回主机

cudaFree(d_data);   // 释放显存
free(h_data);       // 释放内存

1.2.2 核函数的编写与启动

核函数（Kernel）就是在GPU上执行的函数。写法上有个关键点：前面要加__global__修饰符，返回值必须是void。

我个人习惯把核函数命名成xxx_kernel，这样一看就知道是GPU代码。当然这只是风格问题，但团队协作时统一命名规范能省不少事。

// 一个简单的向量加法核函数
__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 threads_per_block = 256;
int blocks_per_grid = (N + threads_per_block - 1) / threads_per_block;
vec_add_kernel<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, N);

注意那个<<<...>>>语法，这是CUDA特有的核函数启动方式。第一个参数是网格中的block数量，第二个参数是每个block中的thread数量。

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

这是CUDA最核心的概念，也是初学者最容易懵的地方。我画个图帮你理解：

这张图展示了CUDA的线程层次结构。最外层是Grid，里面包含多个Block，每个Block又包含多个Thread。你想想看，这就像军队编制：Grid是军，Block是师，Thread是士兵。

每个线程都有唯一的ID，通过threadIdx、blockIdx、blockDim这三个内置变量来定位。我习惯用这个公式：

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

这个公式把一维的线程索引映射到全局。如果是二维或三维，道理一样，只是多几个维度而已。

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

光说不练假把式。咱们写个完整的向量加法，把上面讲的知识串起来。

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

// 核函数：向量加法
__global__ void vector_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];
    }
}

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;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
    
    // 5. 等待GPU完成并检查错误
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(err));
        return -1;
    }
    
    // 6. 结果从设备搬回主机
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 7. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("h_c[%d] = %f\n", i, h_c[i]);
    }
    
    // 8. 清理资源
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

这个例子虽然简单，但包含了CUDA编程的所有关键步骤。我建议你亲手敲一遍，感受一下整个流程。

嗯，第一章的内容就到这儿。这些基础概念虽然简单，但它们是后面所有高级话题的基石。我在实际项目中见过太多人因为基础不牢，写出来的CUDA代码要么跑不对，要么跑不快。所以，别急着跳，把这一章吃透了再说。