3、CUDA编程入门：核函数编写、线程层次结构、内存模型、第一个向量加法程序

好，咱们正式开始动手写CUDA代码了。

说实话，很多搞金融量化的人一听到「GPU编程」就觉得头大，觉得那是图形学或者AI训练才用的东西。其实不然。我当年第一次在回测系统里用CUDA做并行计算时，一个原本跑20分钟的策略，优化后只用了40秒。那一刻我就知道——这条路走对了。

这一章，咱们不讲虚的。直接上手写核函数，搞懂线程怎么排布，内存怎么分配。最后，我会带你写出你的第一个CUDA程序：向量加法。

3.1 核函数：GPU上跑的函数长什么样？

核函数，说白了就是在GPU上执行的函数。CPU上跑的叫「普通函数」，GPU上跑的叫「核函数」。写法上，就多了一个关键字：__global__。

核函数的基本规则：

  • 用 __global__ 修饰
  • 返回类型必须是 void
  • 调用时用尖括号语法：kernel_name<<<grid, block>>>(args)

举个例子，一个最简单的核函数：

__global__ void hello_from_gpu() {
    printf("Hello from GPU!\n");
}

调用它也很简单：

hello_from_gpu<<<1, 1>>>();
cudaDeviceSynchronize();

嗯，这里要注意：cudaDeviceSynchronize() 是必须的。为什么？因为CPU和GPU是异步执行的。你不加这一行，CPU可能已经跑完了，GPU还没开始干活。我曾经在项目里踩过这个坑，调试了半天发现结果不对，最后发现是没同步。

3.2 线程层次结构：别让GPU闲着

GPU最擅长的就是「人多力量大」。一个核函数启动后，会创建成千上万个线程。这些线程怎么组织？CUDA用了两层结构：网格（Grid） 和 线程块（Block）。

你想想看，一个Grid就像一个大工厂，里面有很多个Block（车间），每个Block里又有多个Thread（工人）。

我习惯用这张图来理解：

Grid（网格） Block (0,0) Thread Thread Thread Thread Block (1,0) Thread Thread Thread Thread Block (2,0) Thread Thread Thread Thread 每个Block包含多个Thread 每个Grid包含多个Block 内置变量： threadIdx.x, blockIdx.x, blockDim.x, gridDim.x

每个线程都有一个唯一的ID。怎么算？公式很简单：

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

这个公式我建议你背下来。在金融计算中，我们经常用这个ID来索引数据。比如处理100万个股票价格数据，每个线程处理一个数据点。

我的经验：Block的大小一般设为256或512。太小了浪费SM资源，太大了可能超出寄存器限制。我一般从256开始调，效果不好再试512。

3.3 内存模型：数据在GPU里怎么放？

GPU的内存不像CPU那么简单。它有好多层，速度差异巨大。我画个图你就明白了：

全局内存（Global Memory） 容量大（GB级），速度慢（~400 GB/s），所有线程可访问 共享内存（Shared Memory） 容量小（几十KB），速度快，同一Block内线程共享 寄存器（Register） 最快，每个线程私有 速度：寄存器 > 共享内存 > 全局内存

在金融计算中，最常用的就是全局内存和共享内存。全局内存用来存原始数据，比如行情数据、因子矩阵。共享内存用来做中间计算，比如归约求和。

避坑指南：我曾经在计算协方差矩阵时，把所有数据都放在全局内存里，结果带宽成了瓶颈。后来把部分中间结果放到共享内存，速度提升了3倍。记住：能放共享内存的，就别放全局内存。

3.4 第一个CUDA程序：向量加法

好了，理论讲完了，咱们动手写代码。向量加法是CUDA的「Hello World」。两个数组对应元素相加。

先看CPU版本：

void vector_add_cpu(float *a, float *b, float *c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

再看GPU版本：

__global__ void vector_add_gpu(float *a, float *b, float *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

注意到那个 if (tid < n) 了吗？这是边界检查。因为数据长度不一定正好是线程数的整数倍，多出来的线程不能干活。

完整的调用代码：

int main() {
    int n = 1000000;
    size_t bytes = n * sizeof(float);
    
    // 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);
    
    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    // 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    
    // 启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    vector_add_gpu<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 拷贝结果回CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("h_c[%d] = %f\n", i, h_c[i]);
    }
    
    // 释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

关键点总结：

  • cudaMalloc 在GPU上分配内存
  • cudaMemcpy 负责CPU和GPU之间的数据传输
  • 核函数用 <<<grid, block>>> 启动
  • 别忘了 cudaDeviceSynchronize() 或 cudaMemcpy 隐式同步

说实话，第一次跑通这个程序的时候，我心里还是挺激动的。虽然只是简单的加法，但背后是成千上万个线程在同时工作。你想想看，如果数据量是1亿，CPU要循环1亿次，GPU只需要分块并行处理，速度差距是数量级的。

我的习惯：每次写CUDA程序，我都会先写一个CPU版本做基准测试。然后对比GPU结果，确保数值一致。别问我为什么——我曾经因为忘记初始化GPU内存，查了整整两天的bug。

好了，这一章的内容就到这里。核函数怎么写、线程怎么排、内存怎么管、向量加法怎么实现，你应该都清楚了。下一章，我们会深入内存优化，看看怎么用共享内存把性能再提一个台阶。

专注资料整理

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