第一章：CUDA编程入门——从零开始理解GPU编程

大家好，我是你们这门课的主讲。今天咱们聊聊CUDA编程入门。

说实话，我第一次接触CUDA的时候，也是一头雾水。CPU编程玩得好好的，突然要面对成千上万个线程，感觉就像从开小轿车直接跳到开火车。但别担心，只要理解了核心模型，你会发现GPU编程其实很直观。

1.1 CUDA编程模型：CPU+GPU的协作模式

CUDA的编程模型，说白了就是「CPU负责调度，GPU负责干活」。CPU我们叫它Host（主机），GPU叫Device（设备）。

我习惯把GPU想象成一个超级打工仔——你（CPU）把任务分好，告诉它「去，把这一万个数字加起来」，它就会派成千上万个线程同时干活。干完了再把结果还给你。

这个模型的核心就三个步骤：

  • 数据准备：在CPU上准备好输入数据
  • 数据传输：把数据从CPU内存拷贝到GPU显存
  • 执行计算：在GPU上启动Kernel函数，让成千上万个线程并行计算
  • 结果回传：把计算结果从GPU显存拷回CPU内存

重要概念：CPU和GPU有各自独立的内存空间。你不能直接在GPU上访问CPU的变量，反之亦然。所有数据都得通过显式的拷贝操作来传输。

1.2 第一个Kernel函数：Hello World的GPU版

写一个CUDA程序，其实就是写一个能在GPU上运行的函数——我们叫它Kernel函数。来看个最简单的例子：

// 这是一个Kernel函数，前面加上了 __global__ 修饰符
__global__ void hello_from_gpu() {
    // 每个线程都会执行这里的代码
    printf("Hello from GPU! Thread ID: %d\n", threadIdx.x);
}

int main() {
    // 启动Kernel：<<<1, 5>>> 表示用1个block，每个block里5个线程
    hello_from_gpu<<<1, 5>>>();
    
    // 等待GPU执行完毕
    cudaDeviceSynchronize();
    return 0;
}

运行这段代码，你会看到5行输出，每行来自不同的线程。嗯，这里要注意：<<<grid_size, block_size>>> 这个语法是CUDA特有的，它告诉GPU「我要启动多少个线程」。

我的经验：刚开始写Kernel函数时，我总忘记加 cudaDeviceSynchronize()。结果CPU跑完了，GPU还没算完，程序就退出了。记住：Kernel是异步的，CPU不会等GPU。

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

为什么要有Grid和Block？直接搞一堆线程不行吗？

你想想看，GPU里有几千个计算核心，如果所有线程都乱跑，那不就乱套了？CUDA用了一个三层结构来组织线程：

  • Thread（线程）：最基础的计算单元，每个线程执行相同的代码
  • Block（线程块）：一组线程的集合，同一个Block里的线程可以共享数据
  • Grid（网格）：一组Block的集合，一个Kernel对应一个Grid

我画了一张图，帮你理解这个层次结构：

Grid（网格） Block (0,0) Block (1,0) T(0,0) T(1,0) T(0,1) T(1,1) 共享内存 T(0,0) T(1,0) T(0,1) T(1,1) 共享内存 全局内存（Global Memory）—— 所有线程都能访问

每个线程都有自己唯一的ID，通过 threadIdx、blockIdx、blockDim 这些内置变量来获取。比如：

// 获取全局唯一的线程ID
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式我建议你背下来。我在项目中几乎每个Kernel都会用到它。

1.4 内存管理：cudaMalloc 和 cudaMemcpy

GPU不能直接用CPU的内存，所以我们需要在显存上分配空间。CUDA提供了两个核心函数：

函数 作用 类比CPU
cudaMalloc() 在GPU显存上分配内存 malloc()
cudaMemcpy() 在CPU和GPU之间拷贝数据 memcpy()
cudaFree() 释放GPU显存 free()

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

// Kernel：两个向量相加
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1000;
    size_t size = n * sizeof(float);
    
    // 1. 在CPU上分配内存并初始化
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    // 初始化 h_a 和 h_b...
    
    // 2. 在GPU上分配显存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 3. 把数据从CPU拷贝到GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 4. 启动Kernel
    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);
    
    // 5. 把结果拷回CPU
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 6. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

我曾经踩过的坑：有一次我忘了检查 cudaMalloc 的返回值，结果显存不够用，程序直接崩了。后来我养成了习惯，每次分配完都加一句错误检查：

cudaError_t err = cudaMalloc(&d_a, size);
if (err != cudaSuccess) {
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
    return -1;
}

1.5 线程块大小的选择：经验之谈

你可能会问：一个Block里放多少个线程合适？

这个问题没有标准答案，但我可以分享一些经验：

  • 线程数最好是32的倍数：因为GPU的warp（线程束）大小就是32。不是32的倍数会浪费计算资源
  • 常见选择是128、256、512：太小了隐藏不了内存延迟，太大了寄存器不够用
  • 一个Block的线程数不要超过1024：这是硬件限制，超过了Kernel启动会失败

我的习惯：一般先用256试试，然后根据性能调大调小。如果计算密集型的任务，我会用128；如果是内存密集型的，我会用256或512。具体还得看你的GPU型号。

1.6 本章小结

这一章我们走通了CUDA编程的完整流程：从编程模型到Kernel编写，从线程层次到内存管理。说白了，就是三件事：

  1. 写一个带 __global__ 的函数
  2. 用 <<<grid, block>>> 启动它
  3. 用 cudaMalloc 和 cudaMemcpy 管理数据

下一章我们会深入Kernel函数的优化技巧，包括如何利用共享内存、如何避免bank conflict。但今天先把基础打牢——回去写个向量加法的程序，跑通了再说别的。

可可精品资料库

← 上一章 📖 返回目录 下一章 →