3、CUDA编程入门：CUDA环境搭建、核函数编写、线程层次结构
好，咱们直接进入正题。这一章是CUDA编程的起手式，说白了就是让你能写第一行GPU代码。我当年刚接触CUDA时，光装环境就折腾了两天，后来发现很多坑其实是可以避免的。今天我把这些经验都揉碎了讲给你听。
3.1 CUDA环境搭建——别让工具卡住你
环境搭建这事儿，说简单也简单，说复杂也复杂。我个人习惯是：先确认硬件，再选驱动版本，最后装CUDA Toolkit。顺序错了，后面全是坑。
3.1.1 硬件要求
你得有一块NVIDIA显卡。别笑，我真见过有人拿着AMD显卡来问我为什么CUDA跑不起来。嗯，这确实不行。
| 显卡系列 | 计算能力(Compute Capability) | 推荐程度 |
|---|---|---|
| RTX 30/40系列 | 8.0+ | 强烈推荐 |
| GTX 10/16系列 | 6.0-7.5 | 可用，但显存偏小 |
| Tesla/Quadro系列 | 7.0+ | 专业级，适合服务器 |
3.1.2 驱动与CUDA Toolkit安装
这里有个经典误区：很多人以为装了CUDA Toolkit就自带驱动了。其实不是。你得先装NVIDIA驱动，再装Toolkit。
# 检查驱动是否装好
nvidia-smi
# 如果看到类似下面的输出，说明驱动OK
# +-----------------------------------------------------------------------------+
# | NVIDIA-SMI 535.86.05 Driver Version: 535.86.05 CUDA Version: 12.2 |
# +-----------------------------------------------------------------------------+3.1.3 验证安装
装完之后，跑一下官方示例是最稳妥的验证方式。
# 编译并运行deviceQuery
cd /usr/local/cuda/samples/1_Utilities/deviceQuery
make
./deviceQuery
# 如果看到 "Detected 1 CUDA Capable device(s)" 就说明环境OK了3.2 核函数编写——GPU的"入口函数"
核函数（Kernel Function）是运行在GPU上的函数。你想想看，CPU和GPU是两套不同的处理器，CPU上跑的是普通函数，GPU上跑的就是核函数。写法上有个关键标识符：__global__。
3.2.1 第一个核函数
咱们写一个最简单的向量加法。别小看这个例子，我当年面试量化公司时，手撕的就是这个。
// 核函数定义
__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];
}
}
// 主函数中调用
int main() {
// ... 分配内存、拷贝数据 ...
// 启动核函数
vecAdd<<<1, 256>>>(d_A, d_B, d_C, N);
// ... 拷贝结果回CPU ...
}<<<gridDim, blockDim>>> 语法启动。gridDim是网格维度，blockDim是线程块维度。这两个参数决定了有多少线程在干活。
3.2.2 核函数的限制
核函数和普通函数有几个重要区别，我在项目中吃过亏，给你列出来：
- 不能有返回值——必须用
void类型 - 不能调用C++标准库——比如
printf在核函数里能用，但std::vector不行 - 不能使用静态变量——每个线程都是独立的
- 递归调用有限制——GPU的栈空间很小
3.3 线程层次结构——理解GPU的"组织架构"
GPU的线程组织方式，说白了就是三层结构：网格(Grid) → 线程块(Block) → 线程(Thread)。你想想看，一个公司有部门，部门里有小组，小组里有员工。GPU也是这个道理。
3.3.1 三层结构详解
| 层级 | 说明 | 最大数量 |
|---|---|---|
| Grid (网格) | 一个核函数启动时创建 | 2^31-1 (一维) |
| Block (线程块) | 线程的集合，共享共享内存 | 1024个线程/块 |
| Thread (线程) | 最小的执行单元 | 取决于硬件 |
每个线程都有一个唯一的ID，通过 threadIdx、blockIdx、blockDim 这三个内置变量来计算。公式很简单：
// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 二维情况（处理图像时常用）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;3.3.2 线程块大小怎么选？
这是个经典问题。我刚开始时也纠结过，后来总结出几条经验：
- 线程块大小最好是32的倍数——因为warp（线程束）是32个线程一组
- 常用值：128、256、512——太小了隐藏不了延迟，太大了寄存器不够用
- 不要超过1024——这是硬件上限
举个例子，处理100万条数据：
int N = 1000000;
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize; // 向上取整
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);3.3.3 线程层次与硬件映射
理解了这个映射关系，你才算真正懂了GPU。我画了一张图帮你理解：
这张图展示了从Grid到Block再到Thread的层级关系。每个Block内的线程可以通过共享内存通信，但不同Block之间的线程是隔离的。这个特性在量化系统的并行回测中特别有用——每个Block处理一只股票，互不干扰。
3.4 实战：第一个CUDA程序
光说不练假把式。咱们写一个完整的向量加法程序，把今天学的串起来。
#include <iostream>
#include <cuda_runtime.h>
// 核函数：向量加法
__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
C[idx] = A[idx] + B[idx];
}
}
int main() {
int N = 1 << 20; // 1048576个元素
size_t bytes = N * sizeof(float);
// 1. 分配主机内存
float *h_A = new float[N];
float *h_B = new float[N];
float *h_C = new float[N];
// 初始化数据
for (int i = 0; i < N; i++) {
h_A[i] = i * 1.0f;
h_B[i] = (N - i) * 1.0f;
}
// 2. 分配设备内存
float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, bytes);
cudaMalloc(&d_B, bytes);
cudaMalloc(&d_C, bytes);
// 3. 拷贝数据到GPU
cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
// 4. 启动核函数
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
// 5. 拷贝结果回CPU
cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
// 6. 验证结果
bool correct = true;
for (int i = 0; i < N; i++) {
if (fabs(h_C[i] - (h_A[i] + h_B[i])) > 1e-5) {
correct = false;
break;
}
}
std::cout << "结果" << (correct ? "正确" : "错误") << std::endl;
// 7. 清理
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
delete[] h_A;
delete[] h_B;
delete[] h_C;
return 0;
}cudaFree 释放设备内存。我刚开始时经常忘记，结果显存泄漏，跑几次回测就卡死了。养成好习惯，申请和释放成对出现。
3.5 常见错误与调试技巧
写CUDA程序，调试是个麻烦事。我分享几个实用技巧：
- 检查错误码： 每个CUDA API调用后，用
cudaGetLastError()检查 - 使用cuda-memcheck： 检测越界访问和内存泄漏
- printf调试： 核函数里可以用
printf，但别打太多，会影响性能
// 错误检查宏，我每个项目都这么用
#define CUDA_CHECK(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
// 使用示例
CUDA_CHECK(cudaMalloc(&d_A, bytes));嗯，这一章的内容就到这里。环境搭好了，核函数会写了，线程层次也理解了，接下来就可以开始真正的GPU编程了。记住，多写多练，这些概念自然就熟了。
公众号:蓝海资料掘金营，微信deep3321