3、CUDA运行时API详解：设备管理与查询、内存管理、流与事件机制、错误处理与调试技巧
好，咱们直接进入正题。CUDA运行时API，说白了就是你跟GPU打交道的那层“翻译官”。你写的核函数、你分配显存、你同步数据，最后都得通过它来落地。我这些年调过不少GPU程序，说实话，很多性能瓶颈和隐蔽的bug，根源都在API使用不当上。这一章，咱们就把这些核心API掰开揉碎，讲清楚。
3.1 设备管理与查询：先摸清你的GPU家底
写CUDA程序的第一步，不是写核函数，而是先问问GPU：“你行不行？” 我习惯在程序启动时，先做一次设备探测。
核心函数：
cudaGetDeviceCount()：看看机器上有几块GPU。cudaSetDevice()：指定用哪块卡干活。多卡环境里，这个调用顺序很重要。cudaGetDeviceProperties()：获取GPU的详细参数，比如计算能力、显存大小、SM数量、最大线程数等。
避坑指南： 我曾经在服务器上跑一个多卡任务，忘了调用 cudaSetDevice，结果所有线程都默认抢了0号卡，显存直接爆了。嗯，从那以后，我每个线程启动第一件事就是绑定设备。
举个例子，查询设备属性的代码片段：
int deviceCount;
cudaGetDeviceCount(&deviceCount);
for (int i = 0; i < deviceCount; ++i) {
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, i);
printf("设备 %d: %s, 计算能力 %d.%d, 显存 %zu MB\n",
i, prop.name, prop.major, prop.minor, prop.totalGlobalMem / 1024 / 1024);
}你想想看，如果不做这一步，你写的核函数可能根本跑不起来——比如你用了计算能力8.0的特性，结果用户手里是块6.1的卡，那直接就崩了。
3.2 内存管理：显存操作的“三板斧”
CUDA内存管理，核心就三个函数：cudaMalloc、cudaMemcpy、cudaFree。听起来简单，但坑不少。
3.2.1 cudaMalloc：分配显存
这个函数负责在GPU上划一块内存出来。注意，它返回的是设备指针，你不能在CPU端直接解引用。
float *d_data;
cudaMalloc((void**)&d_data, size * sizeof(float));我个人习惯，每次分配后立刻检查返回值，或者用 cudaError_t 捕获错误。别偷懒，显存分配失败是常有的事。
3.2.2 cudaMemcpy：CPU与GPU的数据搬运工
这个函数负责在主机和设备之间拷贝数据。方向由最后一个参数指定：
cudaMemcpyHostToDevice：CPU → GPUcudaMemcpyDeviceToHost：GPU → CPUcudaMemcpyDeviceToDevice：GPU内部拷贝
cudaMemcpy(d_data, h_data, size * sizeof(float), cudaMemcpyHostToDevice);注意： 这个函数是同步的。也就是说，它要等数据全部搬完才返回。如果你在性能敏感路径上频繁调用它，程序会卡死。我见过有人把 cudaMemcpy 放在循环里，结果带宽利用率极低，性能惨不忍睹。
3.2.3 cudaFree：用完记得还
显存是稀缺资源。分配了不释放，就是内存泄漏。在GPU上，泄漏的后果更严重——整个程序可能因为显存不足而崩溃。
cudaFree(d_data);嗯，这里要注意：释放空指针是安全的，但释放已经释放过的指针，会报错。所以，我一般会在释放后把指针置为 nullptr。
3.3 流（Stream）与事件（Event）机制：让GPU忙起来
默认情况下，CUDA操作是串行的。但GPU是个并行怪兽，你得用流和事件来喂饱它。
3.3.1 流：任务并行的基本单位
流，你可以理解成GPU上的一条“流水线”。不同流里的操作可以并发执行。比如，你可以在流A里跑核函数，同时在流B里拷贝数据。
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在流1中执行核函数
kernelA<<<grid, block, 0, stream1>>>(d_data);
// 在流2中执行核函数
kernelB<<<grid, block, 0, stream2>>>(d_data);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);个人经验： 我在做视频处理时，把解码、预处理、推理分别放到三个流里，吞吐量直接翻倍。但要注意，流不是越多越好。流的数量超过硬件能并行处理的极限，反而会增加调度开销。
3.3.2 事件：精确的时间戳与同步点
事件用来标记流中的某个时间点。你可以用它来测量核函数执行时间，或者让一个流等待另一个流的某个事件完成。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
kernel<<<grid, block>>>(d_data);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("核函数耗时: %f ms\n", milliseconds);
cudaEventDestroy(start);
cudaEventDestroy(stop);为什么会用事件而不是CPU计时？因为事件是在GPU端记录的，不受CPU调度干扰，精度更高。
3.4 错误处理与调试技巧：别让bug过夜
CUDA程序调试起来比CPU程序麻烦，因为GPU端崩溃了，CPU端往往毫无察觉。所以，错误处理必须做到“逢错必查”。
3.4.1 检查每个API的返回值
几乎所有CUDA运行时API都返回 cudaError_t。我习惯写一个宏来封装检查逻辑：
#define CUDA_CHECK(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA错误: %s, 位置: %s:%d\n", \
cudaGetErrorString(err), __FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
} while(0)
// 使用示例
CUDA_CHECK(cudaMalloc(&d_data, size));避坑指南： 我曾经在调试一个复杂的流水线时，某个 cudaMemcpy 悄悄失败了，但程序没报错，只是结果不对。我花了整整两天，一行一行加检查，才发现是之前一个核函数越界写坏了显存。从那以后，我所有API调用都套上了 CUDA_CHECK。
3.4.2 使用cuda-gdb和Nsight
如果 printf 在核函数里不好使（其实从计算能力2.0开始就支持了，但输出顺序不可控），那就上专业工具：
- cuda-gdb：GDB的CUDA版本，可以断点、单步调试核函数，查看设备变量。
- Nsight Systems：看时间线，分析流和事件的执行情况，定位性能瓶颈。
- Nsight Compute：分析核函数内部的指令和带宽利用率。
说白了，调试CUDA程序，三分靠代码检查，七分靠工具。别硬扛，该上工具就上工具。
3.5 本章知识体系总览
下面这张图，是我梳理的CUDA运行时API核心脉络。你可以把它当成一张“作战地图”。
这张图把四个核心模块串起来了。你从设备管理开始，搞清楚硬件能力；然后用内存管理分配好数据；接着用流和事件把任务并行化；最后用错误处理兜底。每一步都做扎实了，GPU程序才能既快又稳。