第4章 CUDA编程基础

第1章：CUDA编程基础：Kernel函数定义、执行配置与线程索引计算

各位同学，欢迎来到《CUDA内核优化：量化策略回测效率提升实战》的第一章。

说实话，很多做量化的人一听到CUDA就觉得头大。觉得那是搞图形学、搞AI训练的人才需要碰的东西。但我要告诉你，在量化回测这个场景里，GPU并行计算的价值被严重低估了。我见过太多团队，用Python单线程跑回测，一个策略优化要跑三天三夜。后来我帮他们改成CUDA内核，同样的逻辑，十分钟搞定。

嗯，咱们不画大饼。今天就从最基础的东西讲起。你不需要有GPU编程经验，但最好懂一点C语言基础。咱们一步步来。

1.1 什么是Kernel函数？

Kernel函数，说白了就是在GPU上执行的函数。它和CPU上的普通函数最大的区别是什么？是它会被成千上万个线程同时执行。

我刚开始学CUDA时，总把Kernel当成普通函数来写。结果发现，一个Kernel里如果写了printf，控制台会刷出几万行输出——因为每个线程都执行了一次。这就是Kernel的“并行”本质。

核心定义：Kernel函数使用 __global__ 修饰符声明，调用时使用 <<<grid, block>>> 语法指定执行配置。它没有返回值，必须返回void。

// 一个最简单的Kernel函数
__global__ void hello_kernel() {
    // 每个线程都会执行这里的代码
    printf("Hello from thread!\n");
}

int main() {
    // 启动Kernel：1个grid，1个block
    hello_kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

你看，这个Kernel只启动了一个线程。但实际项目中，我们通常会启动成千上万个线程。比如处理100万条行情数据，我们就启动100万个线程，每个线程处理一条。

个人经验：我习惯在Kernel函数名前面加个 kernel_ 前缀，比如 kernel_calc_ma。这样在代码里一眼就能看出哪些是GPU函数，哪些是CPU函数。团队协作时特别有用。

3.2 执行配置：<<<grid, block>>> 到底怎么配？

这是新手最容易懵的地方。为什么要有grid和block两层结构？直接启动N个线程不行吗？

答案是：不行。GPU的硬件架构决定了线程必须分层组织。

我打个比方你就明白了。想象你是一家量化公司的老板，手下有1000个员工（线程）。你不能让这1000个人都直接向你汇报，那你会疯掉。你得把他们分成小组（block），每个小组有个组长。然后这些小组再组成大组（grid），由大组长统一调度。

在CUDA里：

grid：线程块（block）的集合。一个grid包含多个block。
block：线程的集合。一个block包含多个线程。
线程：实际执行计算的最小单元。

// 执行配置示例
dim3 gridDim(2, 2);    // 2x2的grid，共4个block
dim3 blockDim(256);    // 每个block有256个线程
// 总线程数 = 4 * 256 = 1024个线程

my_kernel<<<gridDim, blockDim>>>();

避坑指南：我曾经犯过一个低级错误——把blockDim设成了1024，结果程序直接崩溃。为什么？因为大多数GPU的每个block最多支持1024个线程。超过这个数，CUDA会报错。所以，blockDim一般取32的倍数（warp大小），常用值有128、256、512。

你可能会问：那grid和block的大小怎么选？

没有标准答案。但有个经验法则：

如果数据量小（比如几千条），用1个grid、1个block就够了。
如果数据量大（比如几百万条），让block大小固定为256，然后根据数据量计算grid大小。

// 实际项目中的常见写法
int N = 1000000;  // 数据量
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;  // 向上取整

kernel_process_data<<<gridSize, blockSize>>>(d_data, N);

3.3 线程索引计算：找到你的“工位”

好了，现在我们有成千上万个线程在GPU上跑。每个线程都需要知道自己处理哪条数据。这就需要线程索引计算。

CUDA提供了几个内置变量：

变量名	含义	类型
`threadIdx.x`	线程在block内的索引	uint3
`blockIdx.x`	block在grid内的索引	uint3
`blockDim.x`	每个block的线程数	dim3
`gridDim.x`	grid内的block数	dim3

计算全局线程索引的公式很简单：

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

这个公式我建议你背下来。在量化回测里，我们经常用它来索引数据数组。

// 完整的Kernel示例：计算移动平均线
__global__ void kernel_sma(float* prices, float* sma, int N, int window) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查：防止越界
    if (tid >= N) return;
    
    // 计算SMA
    float sum = 0.0f;
    for (int i = 0; i < window; i++) {
        int idx = tid - i;
        if (idx >= 0) {
            sum += prices[idx];
        }
    }
    sma[tid] = sum / window;
}

注意：边界检查（if (tid >= N) return;）是必须的。因为我们的gridSize是向上取整的，所以会有一些多余的线程。如果不做检查，这些线程会访问到数组外的内存，轻则数据错误，重则程序崩溃。

我记得有一次帮一个团队优化回测代码，他们没做边界检查，结果回测结果时对时错。排查了整整两天才发现是这个问题。从那以后，我写Kernel的第一件事就是加边界检查。

3.4 本章知识体系总览

下面这张图总结了本章的核心内容。你可以把它当作一个速查表。

我的建议：刚开始学的时候，别急着写复杂的Kernel。先写一个最简单的——比如把两个数组相加。跑通它，理解每个线程是怎么工作的。然后再逐步增加复杂度。我在带新人时，都是让他们先写10个不同的简单Kernel，把基础打牢了再上项目。

好了，第一章的内容就到这里。记住：Kernel函数定义、执行配置、线程索引计算，这三者是CUDA编程的基石。后面的所有优化技巧，都是建立在这个基础之上的。

📚 课程目录

第1章：CUDA编程基础：Kernel函数定义、执行配置与线程索引计算

1.1 什么是Kernel函数？

3.2 执行配置：<<<grid, block>>> 到底怎么配？

3.3 线程索引计算：找到你的“工位”

3.4 本章知识体系总览