第4章 CUDA编程入门

第4章：CUDA编程入门：CUDA C基础、核函数编写、<<<>>>语法、线程索引计算

好，咱们终于要开始写真正的CUDA代码了。

前面几章我们把GPU的硬件架构、内存模型、并行原理都聊了一遍。说实话，那些都是理论基础。但做工程的人都知道，光懂理论没用，得能上手写代码才行。这一章，我就带你亲手写第一个CUDA程序。

4.1 CUDA C，它到底是个啥？

说白了，CUDA C就是标准C语言加上一小撮扩展。你别把它想得太神秘。

我当年第一次接触CUDA时，心里也在打鼓：是不是又要学一门新语言？结果翻了几页文档就放心了——if、for、while这些老朋友全都在，只是多了几个修饰符和语法糖。

具体来说，CUDA C在标准C的基础上增加了这么几样东西：

函数类型限定符：告诉编译器这个函数跑在CPU上还是GPU上
内置变量：比如threadIdx、blockIdx，帮你定位当前线程
<<<>>>语法：用来启动GPU上的函数（也就是核函数）
一些API函数：比如cudaMalloc、cudaMemcpy，负责管理GPU内存

嗯，就这些。你如果会写C，那CUDA C你已经会了八成。

4.2 核函数：跑在GPU上的函数

核函数（Kernel Function）是CUDA编程的核心。你可以把它理解成——一个被GPU成千上万个线程同时执行的函数。

写核函数有几个硬性规定，我列出来：

前面必须加__global__修饰符
返回值必须是void
不能使用可变参数（就是不能像printf那样传不定个数参数）
可以调用设备函数（用__device__修饰的函数），但不能调用主机函数

来看个最简单的例子：

// 定义一个核函数：让每个线程打印自己的ID
__global__ void hello_from_gpu()
{
    printf("Hello from thread %d in block %d\n", 
           threadIdx.x, blockIdx.x);
}

你看，除了前面那个__global__，其他跟普通C函数没啥两样。

小提示：我刚开始写核函数时，总忘记加__global__。编译报错后找半天才发现——嗯，这种低级错误犯过一次就不会再犯了。

4.3 <<<>>>语法：启动核函数的钥匙

核函数写好了，怎么让它跑起来？答案就是那个看起来有点奇怪的尖括号语法。

调用格式是这样的：

kernel_name<<<grid_dim, block_dim>>>(参数列表);

这里的grid_dim和block_dim，分别指定了网格的维度和线程块的维度。说白了，就是告诉GPU：你要开多少个线程块，每个块里放多少个线程。

举个例子：

// 启动核函数：2个线程块，每个块4个线程
hello_from_gpu<<<2, 4>>>();

这行代码一执行，GPU上就会启动2×4=8个线程。每个线程都跑一遍hello_from_gpu函数里的代码。

  重要：核函数调用是异步的。也就是说，CPU发出启动命令后不会等GPU执行完，而是直接继续往下走。如果你需要等GPU干完活，得手动调用cudaDeviceSynchronize()。

为什么会这样？因为GPU是协处理器，它和CPU是并行工作的。CPU把任务派给GPU后，自己还能干别的事。这个设计很巧妙，但也容易踩坑——我曾经在项目里忘了同步，结果CPU去读GPU的输出，读到的全是垃圾数据。

4.4 线程索引计算：找到你的位置

现在问题来了：8个线程同时跑同一段代码，它们怎么知道自己该处理哪部分数据？

答案就是内置变量。CUDA提供了几个关键变量，帮每个线程找到自己的“身份”：

threadIdx：线程在块内的索引（从0开始）
blockIdx：线程块在网格内的索引（从0开始）
blockDim：线程块的维度（即每个块有多少线程）
gridDim：网格的维度（即有多少个块）

要计算一个线程的全局唯一ID，公式是这样的：

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

这个公式我建议你记牢。做GPU编程，几乎每个核函数里都会用到它。

来看一个完整的例子：用GPU给数组每个元素加1。

// 核函数：数组每个元素加1
__global__ void add_one(int *data, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 防止越界
    if (tid < n)
    {
        data[tid] += 1;
    }
}

// 主机端调用
int main()
{
    int n = 1000;
    int *d_data;
    
    // 在GPU上分配内存
    cudaMalloc(&d_data, n * sizeof(int));
    
    // 启动核函数：使用256个线程的块，计算需要多少个块
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    
    add_one<<<grid_size, block_size>>>(d_data, n);
    
    // 等待GPU完成
    cudaDeviceSynchronize();
    
    // 清理
    cudaFree(d_data);
    return 0;
}

避坑指南：我曾经在项目里犯过一个低级错误——忘记做越界检查。当数据量不是线程数的整数倍时，最后几个线程会访问到数组外面。嗯，那种bug特别难查，因为不是每次都崩，只在特定数据量下才出问题。从那以后，我每个核函数里都会加一句if (tid < n)。

4.5 线程块的配置：选多少合适？

你可能会问：block_size设成256，这个数字是怎么来的？

其实没有标准答案，但有经验法则：

线程块大小最好是32的倍数（因为warp大小是32）
常见选择：128、256、512
不要超过1024（硬件限制）
网格大小 = (数据量 + 块大小 - 1) / 块大小

我个人习惯用256。这个数字不大不小，大多数场景下表现都不错。当然，追求极致性能时，还是得做profiling来调。

4.6 二维和三维索引

上面讲的都是一维情况。但实际应用中，图像、矩阵这些数据天然就是二维的。CUDA也支持二维和三维的线程索引。

举个例子，处理一张1024×768的图片：

// 二维核函数
__global__ void process_image(int *image, int width, int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x < width && y < height)
    {
        int idx = y * width + x;
        // 处理像素...
        image[idx] = image[idx] * 2;  // 亮度加倍
    }
}

// 调用时
dim3 block_size(16, 16);
dim3 grid_size((1024 + 15) / 16, (768 + 15) / 16);
process_image<<<grid_size, block_size>>>(d_image, 1024, 768);

你看，用dim3类型可以方便地指定多维尺寸。每个线程通过threadIdx.x和threadIdx.y拿到自己的坐标，然后计算出在数组中的线性索引。

小技巧：处理二维数据时，我建议用16×16或32×8的线程块。这两个尺寸在大多数GPU上都能充分利用SM（流式多处理器）的资源。

4.7 本章小结

这一章我们走完了CUDA编程的“Hello World”流程。你想想看，其实核心就三件事：

用__global__写一个核函数
用<<<grid, block>>>启动它
在核函数里用threadIdx和blockIdx计算线程ID

嗯，就这么简单。下一章我们会深入内存管理，聊聊怎么把数据从CPU搬到GPU，再搬回来。那才是真正干活的地方。

📚 课程目录