第九课：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() { |
// 启动内核 |
myKernel<<<grid, 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)); |
myKernel<<<grid, block>>>(d_data, n); |
CUDA_KERNEL_CHECK(); |
CUDA_CHECK(cudaFree(d_data)); |
return 0; |
} |
练习题 9
- CUDA 错误码
cudaSuccess的值是什么？ cudaGetLastError()和cudaDeviceSynchronize()分别检查什么错误？- 为什么内核启动后需要调用
cudaDeviceSynchronize()才能检测到执行错误？
第十课：原子操作
知识点
什么是原子操作？
原子操作是不可分割的操作，在多线程环境下保证数据一致性。
问题场景：
// 非原子操作（危险！） |
int count = 0; |
__global__ void increment(int *count) { |
(*count)++; // 多个线程同时执行，结果不确定 |
} |
解决方案：使用原子操作
CUDA 原子函数
| 函数 | 操作 | 说明 |
|---|---|---|
atomicAdd() | 加法 | *addr += val |
atomicSub() | 减法 | *addr -= val |
atomicExch() | 交换 | *addr = val |
atomicMin() | 最小值 | *addr = min(*addr, val) |
atomicMax() | 最大值 | *addr = max(*addr, val) |
atomicInc() | 递增 | *addr = (*addr >= val) ? 0 : *addr + 1 |
atomicDec() | 递减 | `addr = (addr == 0) |
atomicCAS() | 比较并交换 | 条件交换 |
atomicAnd() | 与运算 | *addr &= val |
atomicOr() | 或运算 | *addr |= val |
atomicXor() | 异或运算 | *addr ^= val |
atomicAdd 示例
#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; |
atomicAddKernel<<<gridSize, 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(); |
} |
这段代码是 CUDA（GPU 编程）中非常经典的一种锁机制实现，叫做“自旋锁”（Spinlock）。
要理解这段代码，需要弄懂两个核心概念：atomicCAS是什么，以及while循环在干什么。
1. 核心概念：atomicCAS
atomicCAS全称是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循环里，只能再次执行
（支持资料、图片参考_相关定制）_可以扫码)
