第九课CUDA 错误处理知识点为什么需要错误处理CUDA API 调用可能失败常见原因内存不足设备不存在内核启动失败驱动程序错误不检查错误会导致程序崩溃结果错误难以调试CUDA 错误类型typedef enum cudaError {cudaSuccess 0, // 成功cudaErrorInvalidValue 1, // 无效参数cudaErrorMemoryAllocation 2, // 内存分配失败cudaErrorInvalidDevice 10, // 无效设备cudaErrorInvalidMemcpyDirection 21, // 无效拷贝方向// ... 更多错误码} cudaError;错误检查函数// 基本错误检查cudaError_t err cudaMalloc(d_data, size);if (err ! cudaSuccess) {printf(CUDA 错误: %s\n, cudaGetErrorString(err));exit(1);}封装错误检查宏// 定义错误检查宏#define CUDA_CHECK(call) \do { \cudaError_t err call; \if (err ! cudaSuccess) { \fprintf(stderr, CUDA 错误 at %s:%d: %s\n, \__FILE__, __LINE__, cudaGetErrorString(err)); \exit(1); \} \} while(0)// 使用宏CUDA_CHECK(cudaMalloc(d_data, size));CUDA_CHECK(cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice));内核启动错误检查__global__ void myKernel(int *data, int n) {int idx blockIdx.x * blockDim.x threadIdx.x;if (idx n) {data[idx] idx * 2;}}int main() {// 启动内核myKernelgrid, block(d_data, n);// 检查内核启动错误cudaError_t err cudaGetLastError();if (err ! cudaSuccess) {printf(内核启动失败: %s\n, cudaGetErrorString(err));return -1;}// 等待内核完成并检查执行错误err cudaDeviceSynchronize();if (err ! cudaSuccess) {printf(内核执行失败: %s\n, cudaGetErrorString(err));return -1;}return 0;}完整的错误处理模板#include stdio.h#include stdlib.h#define CUDA_CHECK(call) \do { \cudaError_t err call; \if (err ! cudaSuccess) { \fprintf(stderr, CUDA 错误 at %s:%d: %s\n, \__FILE__, __LINE__, cudaGetErrorString(err)); \exit(1); \} \} while(0)#define CUDA_KERNEL_CHECK() \do { \cudaError_t err cudaGetLastError(); \if (err ! cudaSuccess) { \fprintf(stderr, 内核启动错误 at %s:%d: %s\n, \__FILE__, __LINE__, cudaGetErrorString(err)); \exit(1); \} \err cudaDeviceSynchronize(); \if (err ! cudaSuccess) { \fprintf(stderr, 内核执行错误 at %s:%d: %s\n, \__FILE__, __LINE__, cudaGetErrorString(err)); \exit(1); \} \} while(0)int main() {int n 1000;size_t size n * sizeof(float);float *d_data;CUDA_CHECK(cudaMalloc(d_data, size));myKernelgrid, block(d_data, n);CUDA_KERNEL_CHECK();CUDA_CHECK(cudaFree(d_data));return 0;}练习题 9CUDA 错误码cudaSuccess的值是什么cudaGetLastError()和cudaDeviceSynchronize()分别检查什么错误为什么内核启动后需要调用cudaDeviceSynchronize()才能检测到执行错误第十课原子操作知识点什么是原子操作原子操作是不可分割的操作在多线程环境下保证数据一致性。问题场景// 非原子操作危险int count 0;__global__ void increment(int *count) {(*count); // 多个线程同时执行结果不确定}解决方案使用原子操作CUDA 原子函数函数操作说明atomicAdd()加法*addr valatomicSub()减法*addr - valatomicExch()交换*addr valatomicMin()最小值*addr min(*addr, val)atomicMax()最大值*addr max(*addr, val)atomicInc()递增*addr (*addr val) ? 0 : *addr 1atomicDec()递减addr (addr 0)atomicCAS()比较并交换条件交换atomicAnd()与运算*addr valatomicOr()或运算*addr | valatomicXor()异或运算*addr ^ valatomicAdd 示例#include stdio.h__global__ void atomicAddKernel(int *count, int n) {int idx blockIdx.x * blockDim.x threadIdx.x;if (idx n) {atomicAdd(count, 1); // 原子递增}}int main() {int n 10000;int h_count 0;int *d_count;cudaMalloc(d_count, sizeof(int));cudaMemcpy(d_count, h_count, sizeof(int), cudaMemcpyHostToDevice);int blockSize 256;int gridSize (n blockSize - 1) / blockSize;atomicAddKernelgridSize, blockSize(d_count, n);cudaMemcpy(h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost);printf(计数结果: %d (预期: %d)\n, h_count, n);cudaFree(d_count);return 0;}atomicCAS比较并交换// atomicCAS(int *addr, int compare, int val)// 如果 *addr compare则 *addr val// 返回 *addr 的旧值__global__ void casExample(int *data, int old_val, int new_val) {int idx blockIdx.x * blockDim.x threadIdx.x;if (idx 0) {int old atomicCAS(data, old_val, new_val);printf(旧值: %d, 新值: %d\n, old, new_val);}}原子操作实现锁struct Lock {int *mutex;Lock() {cudaMalloc(mutex, sizeof(int));cudaMemset(mutex, 0, sizeof(int));}~Lock() {cudaFree(mutex);}__device__ void lock() {while (atomicCAS(mutex, 0, 1) ! 0) {// 等待锁释放}}__device__ void unlock() {atomicExch(mutex, 0);}};__global__ void kernelWithLock(int *data, Lock lock) {lock.lock();// 临界区代码(*data);lock.unlock();}这段代码是 CUDAGPU 编程中非常经典的一种锁机制实现叫做“自旋锁”Spinlock。要理解这段代码需要弄懂两个核心概念atomicCAS是什么以及while循环在干什么。1. 核心概念atomicCASatomicCAS全称是Atomic Compare-And-Swap原子比较并交换。在这个函数中atomicCAS(mutex, 0, 1)接收三个参数参数 1 (mutex)你要操作的那个变量锁的状态。参数 2 (0)你期望此时锁的值是多少0 表示锁当前是空闲的。参数 3 (1)如果锁真的像你期望的一样是空闲的为 0你就把它改成新值1 表示你占用了这个锁。⚠️最容易产生误解的地方必须记住atomicCAS的返回值永远是mutex改变之前的“旧值”。它并不是返回一个 True 或 False“原子操作”意味着这个动作是瞬间完成的绝对不可被打断。就算有 1000 个 GPU 线程同时执行这行代码硬件也会保证它们一个一个排队执行这个判断和交换的过程。2. 场景推演它是怎么锁住的我们假设有线程 A和线程 B同时想要获取这个锁。初始状态下锁是解开的也就是mutex 0。场景一线程 A 先到达线程 A 执行atomicCAS(mutex, 0, 1)。硬件一看当前的mutex确实是0没人占用。于是硬件把mutex改成了1表示被线程 A 锁上了。返回值返回mutex被修改前的旧值也就是0。来看while判断条件while( 0 ! 0 )。这个条件是假 (False)所以线程 A跳出while循环成功拿到锁去执行后面的代码了。场景二线程 B 紧接着到达此时线程 A 还没释放锁此时mutex已经被线程 A 变成了1。线程 B 执行atomicCAS(mutex, 0, 1)。硬件一看当前的mutex是1跟你期望的0不相等所以硬件什么都不做不会把值改成 1。返回值依然返回mutex此时的旧值也就是1。来看while判断条件while( 1 ! 0 )。这个条件是真 (True)所以线程 B 被困在了while循环里只能再次执行