第一章：核函数基础——GPU编程的“灵魂”

各位同学，欢迎来到《核函数优化实战案例集》。我是你们的老朋友，一个在GPU优化领域摸爬滚打多年的工程师。今天咱们聊聊最基础、也最核心的东西——核函数。

说实话，我见过不少新手，一上来就急着写复杂的并行算法，结果连核函数的基本概念都没搞明白。最后代码跑出来，要么比CPU还慢，要么直接报错。嗯，咱们今天就把这块地基打牢。

1.1 什么是核函数？

核函数（Kernel Function），说白了，就是一段在GPU上执行的函数。它跟普通的C/C++函数最大的区别在于：它会被成千上万个线程同时执行。

你想想看，CPU上的函数，一次调用只跑一次。但GPU上的核函数，一次调用，成千上万个线程一起跑。这就是并行计算的核心。

核函数的定义：在CUDA中，用 __global__ 修饰符声明的函数，就是核函数。它运行在GPU设备端，由CPU端调用。

举个例子，我当年刚接触CUDA时，写了一个向量加法的核函数：

// 核函数定义
__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];
    }
}

// CPU端调用
int main() {
    // ... 数据准备 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // ... 结果回收 ...
}

看到没？__global__ 就是核函数的“身份证”。<<<gridSize, blockSize>>> 是调用时的“调度参数”，告诉GPU要开多少个线程块、每个块里有多少线程。

个人经验：我习惯把核函数想象成一个“工厂流水线”。<<<grid, block>>> 就是告诉工厂：我要开几条产线（grid），每条产线有多少工人（block）。每个工人（线程）干同样的活，但处理不同的数据。

1.2 核函数在GPU编程中的角色

核函数在GPU编程里，到底扮演什么角色？我总结了三点：

并行执行的载体：没有核函数，GPU就是一块废铁。所有并行计算任务，最终都要落到核函数上。
CPU与GPU的桥梁：CPU负责逻辑控制、数据准备，GPU通过核函数执行计算密集型任务。两者配合，才能发挥最大效能。
性能优化的核心：你写的核函数好不好，直接决定了程序能跑多快。我见过同一个算法，不同人写的核函数，性能能差10倍以上。

为什么会这样？因为核函数的设计，直接影响GPU的线程利用率、内存访问模式、计算吞吐量。这些咱们后面会详细讲。

我记得有一次，帮一个团队优化图像处理算法。他们原来的核函数，每个线程只处理一个像素，结果GPU利用率不到30%。我改成每个线程处理4个像素，并优化了内存访问，性能直接翻了3倍。嗯，这就是核函数优化的魅力。

1.3 CUDA编程模型回顾

要理解核函数，必须回顾一下CUDA的编程模型。说白了，CUDA把GPU看作一个协处理器，CPU是主机（Host），GPU是设备（Device）。

整个流程是这样的：

数据准备：在CPU上分配内存，初始化数据。
数据拷贝：把数据从CPU内存拷贝到GPU显存。
执行核函数：CPU调用核函数，GPU上成千上万的线程并行执行。
结果回传：把计算结果从GPU显存拷贝回CPU内存。
清理资源：释放GPU显存和CPU内存。

这个模型，我称之为“搬运-计算-搬运”模式。很多新手容易忽略第一步和第二步，觉得数据拷贝不重要。其实，数据拷贝往往是性能瓶颈。我曾经遇到一个项目，计算只花了10毫秒，数据拷贝却花了100毫秒。你说气不气人？

避坑指南：我曾经犯过一个低级错误——在核函数里用 printf 打印调试信息。结果核函数执行时间从1毫秒变成了100毫秒。记住：核函数里尽量别用 printf，它会把GPU的并行执行串行化。

下面这张图，是我自己画的CUDA编程模型流程图，帮你理清思路：

这张图里，箭头代表数据流动的方向。注意看，数据从CPU到GPU，再从GPU回CPU，中间就是核函数在干活。这个流程，我建议你背下来，因为后面所有的优化，都是围绕这个流程展开的。

1.4 核函数的基本语法

咱们再深入一点，看看核函数的语法细节。CUDA里，核函数用 __global__ 修饰，返回类型必须是 void。为什么？因为核函数是异步执行的，CPU不会等它返回结果。

核函数内部，有几个内置变量，用来区分不同的线程：

变量名	含义	取值范围
`threadIdx.x`	线程在块内的索引	0 到 blockDim.x - 1
`blockIdx.x`	块在网格内的索引	0 到 gridDim.x - 1
`blockDim.x`	每个块的线程数	由调用时的 blockSize 决定
`gridDim.x`	网格中的块数	由调用时的 gridSize 决定

这些变量，说白了就是每个线程的“身份证号”。通过它们，每个线程知道自己该处理哪部分数据。

我的习惯：我一般用 int tid = blockIdx.x * blockDim.x + threadIdx.x; 来计算全局线程ID。这样每个线程都有一个唯一的ID，方便索引数据。这个写法，几乎每个核函数里都会出现，建议你形成肌肉记忆。

1.5 核函数的限制

核函数不是万能的，它有一些限制，我列出来给你看：

不能有返回值：核函数必须是 void 类型。想返回结果？通过参数传指针。
不能使用静态变量：核函数里不能声明 static 变量。
不能使用函数指针：核函数不支持函数指针调用。
不能使用变长参数：像 printf 那样的变长参数，核函数里不能用（除非是计算能力2.0以上的设备）。
递归调用有限制：虽然CUDA支持递归，但深度有限，而且性能很差。我建议你尽量别用。

这些限制，说白了就是GPU的“脾气”。你顺着它，它给你高性能；你逆着它，它就给你报错。我曾经在核函数里试图用 malloc 动态分配内存，结果程序直接崩溃。嗯，从那以后我再也不在核函数里搞动态分配了。

1.6 小结

好了，咱们今天把核函数的基础讲完了。总结一下：

核函数是GPU上并行执行的函数，用 __global__ 修饰。
它是CPU和GPU协作的桥梁，负责计算密集型任务。
CUDA编程模型是“搬运-计算-搬运”模式，数据拷贝不可忽视。
核函数有语法限制，要顺着GPU的“脾气”来。

这些内容，看起来简单，但都是后面优化的基础。你想想看，如果连核函数的基本概念都没搞懂，怎么去优化它的性能？

下一章，咱们会深入讨论线程层次结构，以及如何合理配置线程块大小。到时候，我会分享一些我在实际项目中踩过的坑，保证让你少走弯路。

想要更多资料，联系可可

📚 课程目录

第一章：核函数基础——GPU编程的“灵魂”

1.1 什么是核函数？

1.2 核函数在GPU编程中的角色

1.3 CUDA编程模型回顾

1.4 核函数的基本语法

1.5 核函数的限制

1.6 小结