4、GPU编程模型与编译前端：CUDA C/C++语法基础、kernel函数定义、<<>>执行配置、host与device代码分离

好，我们进入第四章。这一章，咱们要真正动手写CUDA代码了。

说实话，我第一次接触CUDA的时候，最困惑的不是语法本身，而是「这玩意儿到底怎么跟CPU代码配合？」。你想想看，一个程序里既有CPU跑的代码，又有GPU跑的代码，编译器怎么知道哪段该给谁？

嗯，这就是本章要解决的核心问题。

4.1 CUDA C/C++：不是新语言，是扩展

先澄清一个概念：CUDA C/C++ 并不是一门全新的语言。它就是在标准的C/C++基础上，加了一小撮「私货」——几个关键字、几个语法糖、一个运行时库。

我个人习惯把CUDA C/C++ 看作「带GPU加速功能的C++」。你已有的C++知识，90%以上可以直接用。剩下的10%，就是我们要讲的这些。

核心要点：CUDA编译器（nvcc）会解析这些扩展语法，把CPU代码和GPU代码分开编译。CPU代码走主机编译器（gcc/clang），GPU代码走设备编译器（ptxas）。最后再链接到一起。

4.2 kernel函数定义：GPU的入口函数

什么叫kernel函数？说白了，就是在GPU上执行的函数。CPU调用它，GPU执行它。

定义kernel函数，你需要用到一个关键字：__global__。

// 一个简单的kernel函数
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

这里有几个要点：

  • __global__ 告诉编译器：这个函数是给GPU用的，从CPU端调用。
  • kernel函数必须返回 void。你不能让它返回一个值。
  • 函数体里，你可以用 threadIdx、blockIdx、blockDim 这些内置变量。它们告诉你「我是哪个线程」。我在项目中遇到过新手直接写死索引，结果数据全乱套了——嗯，一定要动态计算。

小技巧：我习惯在kernel函数名前加 kernel_ 前缀，比如 kernel_vecAdd。这样一眼就能看出哪些函数是GPU代码。纯个人习惯，但确实能减少混淆。

4.3 <<>> 执行配置：告诉GPU怎么干活

定义好kernel函数，怎么调用它？不是普通函数调用，而是用三括号语法：

// 调用kernel函数
int N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这个 <<<grid, block>>> 就是执行配置。它告诉GPU两件事：

  • grid：有多少个线程块（block）。
  • block：每个线程块里有多少个线程（thread）。

你想想看，GPU里有成千上万个核心。你通过这个配置，告诉它「我要启动这么多线程，你看着办」。实际上，GPU的调度单元会把线程块分配到不同的SM（流多处理器）上执行。

参数 含义 典型值
grid 线程块的数量 几十到几千
block 每个块内的线程数 128、256、512

注意：block内的线程数不能超过GPU硬件的限制。老一点的GPU最多512，新一点的1024。我曾经在GTX 680上设了1024，结果kernel直接启动失败——嗯，查了半天文档才发现是硬件限制。所以，写代码前先查一下你GPU的 maxThreadsPerBlock。

4.4 host与device代码分离：CPU和GPU各司其职

这是CUDA编程模型的核心思想。一个CUDA程序，天然分成两部分：

  • host代码：在CPU上执行。负责数据准备、kernel调用、结果回收。
  • device代码：在GPU上执行。就是kernel函数本身。

编译器怎么区分？靠关键字：

  • __global__：在device上执行，从host调用。
  • __device__：在device上执行，只能从device调用（比如被其他kernel函数调用）。
  • __host__：在host上执行，就是普通CPU函数。其实不加也默认是host。

举个例子：

// host函数：CPU上执行
void initData(float* data, int N) {
    for (int i = 0; i < N; i++) {
        data[i] = (float)i;
    }
}

// device函数：GPU上执行，被kernel调用
__device__ float square(float x) {
    return x * x;
}

// kernel函数：GPU上执行，从host调用
__global__ void kernel_square(float* in, float* out, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        out[i] = square(in[i]);
    }
}

int main() {
    // host代码：数据准备
    float* h_in = new float[1024];
    initData(h_in, 1024);

    // ... 分配device内存，拷贝数据 ...

    // host代码：调用kernel
    kernel_square<<<4, 256>>>(d_in, d_out, 1024);

    // ... 拷贝结果回host，释放内存 ...
    return 0;
}

你看，代码结构非常清晰。CPU负责「管理」，GPU负责「计算」。这种分离，说白了就是分工合作——CPU擅长控制流，GPU擅长数据并行。

4.5 知识体系总览

为了让你更直观地理解这一章的知识结构，我画了一张图：

CUDA编程模型与编译前端知识体系 CUDA C/C++ 编程模型 语法基础 C/C++扩展语法 __global__ / __device__ / __host__ kernel函数定义 void返回类型 内置变量：threadIdx/blockIdx 执行配置 <<<grid, block>>>语法 grid/block/thread层级 host与device代码分离 host：CPU管理 device：GPU计算 nvcc编译器分离

这张图把本章的核心知识点串起来了。从语法基础到kernel定义，再到执行配置和代码分离，每一步都是环环相扣的。

4.6 避坑指南

最后，分享几个我踩过的坑：

  • 忘记加 __global__：编译器不会报错，但函数会在CPU上执行。结果就是——慢得离谱。我曾经debug了一下午才发现。
  • kernel里用 printf：CUDA支持device端的printf，但输出是异步的。你调用完kernel立刻打印，可能什么都看不到。加个 cudaDeviceSynchronize() 就好。
  • 误解 threadIdx 和 blockIdx：它们是三维的！threadIdx.x、threadIdx.y、threadIdx.z。如果你只处理一维数据，记得只用 .x，别把其他维度搞混了。

嗯，这一章的内容就到这里。代码写起来其实不难，关键是理解背后的模型——CPU和GPU怎么分工，数据怎么流动。把这些想清楚了，后面的优化章节你会轻松很多。

公众号:蓝海资料掘金营，微信deep3321

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