当前位置： 首页 > news >正文

CUDA 补充教程 - 进阶与深入

news 2026/6/26 9:07:26

第九课：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

  1. CUDA 错误码cudaSuccess的值是什么？
  2. cudaGetLastError()和cudaDeviceSynchronize()分别检查什么错误？
  3. 为什么内核启动后需要调用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 先到达
  1. 线程 A 执行atomicCAS(mutex, 0, 1)。
  2. 硬件一看，当前的mutex确实是0（没人占用）。
  3. 于是硬件把mutex改成了1（表示被线程 A 锁上了）。
  4. 返回值：返回mutex被修改前的旧值，也就是0。
  5. 来看while判断条件：while( 0 != 0 )。
  6. 这个条件是假 (False)！所以线程 A跳出while循环，成功拿到锁，去执行后面的代码了。
场景二：线程 B 紧接着到达（此时线程 A 还没释放锁）
  1. 此时mutex已经被线程 A 变成了1。
  2. 线程 B 执行atomicCAS(mutex, 0, 1)。
  3. 硬件一看，当前的mutex是1，跟你期望的0不相等！
  4. 所以硬件什么都不做（不会把值改成 1）。
  5. 返回值：依然返回mutex此时的旧值，也就是1。
  6. 来看while判断条件：while( 1 != 0 )。
  7. 这个条件是真 (True)！所以线程 B 被困在了while循环里，只能再次执行
http://www.jsqmd.com/news/1080423/

相关文章：

  • 小白程序员快收藏！低成本AI挖网络安全漏洞实战干货
  • RAG实战指南：构建可落地的检索增强生成系统
  • 【VMware+K8s双栈架构终极手册】：打通vCenter API自动化纳管、Tanzu Kubernetes Grid深度集成与GitOps交付流水线
  • VMware vSphere测试环境部署全流程：从零到上线仅需90分钟，附自动化脚本下载链接
  • 百度网盘解析工具完整教程：免费获取高速下载链接的终极指南
  • dbx-数据库管理神器
  • YOLO26瓶子罐子识别检测系统：7967张标注图像+PyQt5界面+模型权重+远程环境部署（项目源码+数据集+模型权重+UI界面+python+深度学习+远程环境部署）
  • 8 Ball Pool 精准瞄准开源工具：从理论到实战的完整指南
  • DLSS Swapper深度解析：专业级游戏DLSS版本管理实战指南
  • EtherNet/IP 转 Modbus 网关你用过吗？
  • 进程放后台运行，异常退出，如何排查
  • YOLO26扑克牌识别检测系统（项目源码+数据集+模型权重+UI界面+python+深度学习+远程环境部署）
  • VMware中Kubernetes集群搭建失败的7大隐性原因，第4个连资深工程师都曾忽略（附诊断脚本+日志解析速查表）
  • GetQzonehistory：3分钟掌握QQ空间数据备份，永久保存你的青春记忆
  • 重新定义Windows桌面美学：TranslucentTB深度解析与创新实践
  • SchoolCMS开源教务系统：5分钟搭建专业级学校管理平台
  • 2026年南宁市AI获客公司，哪家更受青睐？
  • 易语言调用Java实现3DES加解密：跨语言整合实战指南
  • VMware测试环境搭建实战手册（含ESXi 8.0+Workstation 17双路径详解）
  • HACS集成部署与故障排除技术指南：架构解析与性能优化方案
  • mac安装homebrew
  • Windows 11终极清理指南：3分钟告别系统臃肿，找回纯净体验
  • 【VMware Hadoop集群搭建终极指南】：20年架构师亲授5大避坑要点与3节点高可用部署实录
  • 飞凌嵌入式ElfBoard-线程之线程ID
  • RAG系统抗令牌擦除：基于语义感知冗余的检索增强生成优化
  • 【VMware Python开发环境搭建黄金法则】：20年运维专家亲授5步极速部署法，避开99%新手踩坑雷区
  • 16位海明码硬件实现：从原理到Verilog电路设计全解析
  • 01. 速通Linux内核喂饭版教程
  • 低成本ECC安全芯片—LKT2412
  • Transformer 全面介绍：从原理到应用