第一章：CUDA编程基础——主机与设备、核函数定义与调用、执行配置与内置变量

各位同学，欢迎来到《CUDA核函数算子优化实战精讲》。我是你们这门课的主讲。

今天咱们聊点最基础，但也最核心的东西。说白了，就是CUDA编程的“起手式”。

很多新手一上来就盯着复杂的算子优化，结果连threadIdx和blockIdx都搞混。嗯，我当年也犯过这个错。所以第一章，咱们把地基打牢。

1.1 主机与设备：CPU和GPU的分工

先问一个问题：你的程序跑在哪儿？

传统C/C++程序，代码跑在CPU上，我们叫它主机（Host）。而GPU，我们叫它设备（Device）。

我个人习惯把CPU想象成“项目经理”，GPU是“施工队”。项目经理负责逻辑判断、流程控制、数据准备；施工队负责大规模、重复性的体力活——也就是并行计算。

在CUDA里，主机端代码用普通的C/C++写，设备端代码用核函数（Kernel）写。两者通过CUDA API进行数据交换。

我在项目中遇到过不少新手，直接在核函数里用printf调试，结果发现控制台啥也没输出。为什么？因为核函数在GPU上跑，printf的输出需要特殊处理。这个坑，咱们后面会细讲。

1.2 核函数定义与调用：怎么写一个能在GPU上跑的函数

核函数，就是用__global__修饰符定义的函数。它告诉编译器：这段代码要在GPU上执行，并且可以被主机调用。

看个最简单的例子：

// 定义一个核函数
__global__ void helloFromGPU() {
    printf("Hello from GPU! blockIdx.x = %d, threadIdx.x = %d\n", 
           blockIdx.x, threadIdx.x);
}

int main() {
    // 调用核函数
    helloFromGPU<<<2, 4>>>();
    
    // 同步，确保GPU执行完毕
    cudaDeviceSynchronize();
    
    return 0;
}

注意几个关键点：

• __global__：返回类型必须是void
• 调用时用三括号<<<grid, block>>>指定执行配置
• 核函数是异步的，主机不会等它执行完

1.3 执行配置：<<>> 到底在配置什么？

三括号里的两个参数，决定了你的GPU有多少线程在干活。

咱们拆开来看：

• grid：线程网格，可以理解为一个“任务大组”
• block：线程块，可以理解为一个“任务小组”

每个block里包含多个线程。所有block组成一个grid。

举个例子：

// 2个block，每个block有4个线程
// 总共 2 * 4 = 8 个线程
helloFromGPU<<<2, 4>>>();

你想想看，如果我要处理一个1024个元素的数组，每个线程处理一个元素，那需要1024个线程。怎么配置？

// 8个block，每个block 128个线程
// 8 * 128 = 1024
myKernel<<<8, 128>>>(d_array);

这里有个经验值：block里的线程数最好是32的倍数。为什么？因为GPU的warp大小就是32。不是32的倍数，会有线程束浪费。我刚开始没注意这个，性能直接掉了30%。

1.4 内置变量：threadIdx, blockIdx, blockDim, gridDim

这四个变量，是CUDA编程的“身份证”。每个线程都能通过它们知道自己的位置。

变量名	含义	取值范围
threadIdx	线程在block内的索引	0 到 blockDim - 1
blockIdx	block在grid内的索引	0 到 gridDim - 1
blockDim	block的维度（包含的线程数）	由执行配置决定
gridDim	grid的维度（包含的block数）	由执行配置决定

它们都是dim3类型，有.x、.y、.z三个分量。不过大多数时候我们只用一维。

怎么用？看一个向量加法的例子：

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    // 计算全局线程ID
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

这里的关键公式：全局线程ID = blockIdx.x * blockDim.x + threadIdx.x。这个公式你最好刻在脑子里，后面所有的一维并行计算都离不开它。

为什么要有边界检查？因为线程总数可能超过数组长度。比如数组有1000个元素，你配了1024个线程，那多出来的24个线程就得跳过。我见过有人忘了加这个判断，结果数组越界，数据全乱了。

知识体系总览

下面这张图，把本章的核心逻辑串起来了。你可以把它当成一个“地图”，后面每学一个新概念，都往这张图上放一放。

这张图从左到右，从上到下，就是CUDA程序从主机启动到GPU线程执行的全流程。你写代码的时候，脑子里过一遍这个流程，基本不会跑偏。

本章小结

好，咱们捋一下今天讲了啥：

• 主机（CPU）和设备（GPU）的分工协作
• 用__global__定义核函数，用三括号调用
• 执行配置<<<grid, block>>>决定线程数量
• 四个内置变量帮你定位每个线程

这些是CUDA编程的“ABC”。后面所有的优化技巧——共享内存、线程束合并、流水线——都建立在这个基础上。

下一章，咱们会深入内存模型，聊聊全局内存、共享内存、寄存器这些概念。到时候你会发现，理解了内存，才算真正入了CUDA的门。

---