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));内核启动错误检查globalvoid 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;globalvoid 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.hglobalvoid 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 的旧值globalvoid 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); }};globalvoid kernelWithLock(int *data, Lock lock) {lock.lock();// 临界区代码(*data);lock.unlock();}这段代码是 CUDAGPU 编程中非常经典的一种锁机制实现叫做“自旋锁”Spinlock。要理解这段代码需要弄懂两个核心概念atomicCAS 是什么以及 while 循环在干什么。核心概念atomicCASatomicCAS 全称是 Atomic Compare-And-Swap原子比较并交换。在这个函数中atomicCAS(mutex, 0, 1) 接收三个参数参数 1 (mutex)你要操作的那个变量锁的状态。参数 2 (0)你期望此时锁的值是多少0 表示锁当前是空闲的。参数 3 (1)如果锁真的像你期望的一样是空闲的为 0你就把它改成新值1 表示你占用了这个锁。⚠️ 最容易产生误解的地方必须记住atomicCAS 的返回值永远是 mutex 改变之前的“旧值”。它并不是返回一个 True 或 False“原子操作”意味着这个动作是瞬间完成的绝对不可被打断。就算有 1000 个 GPU 线程同时执行这行代码硬件也会保证它们一个一个排队执行这个判断和交换的过程。场景推演它是怎么锁住的我们假设有线程 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 循环里只能再次执行 atomicCAS 进行判断。只要线程 A 不放开锁线程 B 就会一直在 while 里面疯狂打转这就是为什么叫“自旋锁”它在原地自旋等待。3. 如何解锁虽然你没贴出解锁的代码但配合起来看更容易理解。解锁的代码通常非常简单devicevoid unlock() {atomicExch(mutex, 0); // 释放锁直接把 mutex 变回 0}一旦线程 A 执行了 unlock()把 mutex 改回了 0。还在 while 里苦苦循环的线程 B 在下一次执行 atomicCAS(mutex, 0, 1) 时就会发现旧值变成 0 了于是它成功把 mutex 变成 1返回值变为 0跳出循环成功接管这把锁。 通俗的比喻把 mutex 想象成公共厕所门上的那个指示牌0 代表绿色的“无人”1 代表红色的“有人”。atomicCAS 就是你跑过去看一眼并锁门的连贯动作。你走过去发现是绿色0你立刻进去并把牌子翻到红色1。因为你进来前是绿色返回 0所以你不用排队跳出循环。另一个人走过来发现是红色1他想翻牌子失败了。因为他看到的是红色返回 1所以他只能在门外一直转圈踱步while 循环死死盯着牌子直到你出来把牌子翻回绿色0为止。原子操作性能考虑原子操作比普通操作慢多个线程对同一地址原子操作会串行化尽量减少原子操作的使用考虑使用共享内存减少全局内存原子操作练习题 10为什么多线程环境下普通递增操作 (*count) 会产生错误结果atomicAdd(addr, val) 的作用是什么返回值是什么如何使用原子操作实现一个简单的互斥锁第十一课CUDA 流与异步执行知识点什么是 CUDA 流CUDA 流是一系列按顺序执行的命令队列。不同流中的命令可以并发执行。默认流Stream 0┌─────────────────────────────────────┐│ 内核A → 拷贝1 → 内核B → 拷贝2 │ 串行执行└─────────────────────────────────────┘多流并发Stream 1: ┌─────────────────────────────┐│ 内核A → 拷贝1 │└─────────────────────────────┘Stream 2: ┌─────────────────────────────┐│ 内核B → 拷贝2 │ 并发执行└─────────────────────────────┘创建和使用流cudaStream_t stream1, stream2;// 创建流cudaStreamCreate(stream1);cudaStreamCreate(stream2);// 在指定流中执行操作cudaMemcpyAsync(d_a1, h_a1, size, cudaMemcpyHostToDevice, stream1);kernel1grid, block, 0, stream1(d_a1, d_c1);cudaMemcpyAsync(h_c1, d_c1, size, cudaMemcpyDeviceToHost, stream1);cudaMemcpyAsync(d_a2, h_a2, size, cudaMemcpyHostToDevice, stream2);kernel2grid, block, 0, stream2(d_a2, d_c2);cudaMemcpyAsync(h_c2, d_c2, size, cudaMemcpyDeviceToHost, stream2);// 同步流cudaStreamSynchronize(stream1);cudaStreamSynchronize(stream2);// 销毁流cudaStreamDestroy(stream1);cudaStreamDestroy(stream2);异步内存拷贝// 同步拷贝阻塞cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);// 异步拷贝非阻塞cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);流同步// 同步单个流cudaStreamSynchronize(stream);// 同步所有流cudaDeviceSynchronize();// 等待多个流cudaStreamWaitEvent(stream, event);流优先级// 创建高优先级流int priority_high, priority_low;cudaDeviceGetStreamPriorityRange(priority_low, priority_high);cudaStream_t stream_high;cudaStreamCreateWithPriority(stream_high, cudaStreamNonBlocking, priority_high);完整示例多流并发#include stdio.h#define N_STREAMS 4#define N 1000000globalvoid vectorAdd(float *a, 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 N;size_t size n * sizeof(float);// 分配主机内存页锁定内存用于异步传输 float *h_a[N_STREAMS], *h_b[N_STREAMS], *h_c[N_STREAMS]; for (int i 0; i N_STREAMS; i) { cudaMallocHost(h_a[i], size); cudaMallocHost(h_b[i], size); cudaMallocHost(h_c[i], size); } // 分配设备内存 float *d_a[N_STREAMS], *d_b[N_STREAMS], *d_c[N_STREAMS]; for (int i 0; i N_STREAMS; i) { cudaMalloc(d_a[i], size); cudaMalloc(d_b[i], size); cudaMalloc(d_c[i], size); } // 创建流 cudaStream_t streams[N_STREAMS]; for (int i 0; i N_STREAMS; i) { cudaStreamCreate(streams[i]); } // 并发执行 int blockSize 256; int gridSize (n blockSize - 1) / blockSize; for (int i 0; i N_STREAMS; i) { cudaMemcpyAsync(d_a[i], h_a[i], size, cudaMemcpyHostToDevice, streams[i]); cudaMemcpyAsync(d_b[i], h_b[i], size, cudaMemcpyHostToDevice, streams[i]); vectorAddgridSize, blockSize, 0, streams[i](d_a[i], d_b[i], d_c[i], n); cudaMemcpyAsync(h_c[i], d_c[i], size, cudaMemcpyDeviceToHost, streams[i]); } // 同步所有流 cudaDeviceSynchronize(); // 清理 for (int i 0; i N_STREAMS; i) { cudaStreamDestroy(streams[i]); cudaFree(d_a[i]); cudaFree(d_b[i]); cudaFree(d_c[i]); cudaFreeHost(h_a[i]); cudaFreeHost(h_b[i]); cudaFreeHost(h_c[i]); } return 0;}页锁定内存Pinned Memory// 普通主机内存可分页floath_data (float)malloc(size);// 页锁定主机内存不可分页用于异步传输float *h_data_pinned;cudaMallocHost(h_data_pinned, size);// 释放cudaFreeHost(h_data_pinned);优点支持异步传输传输速度更快DMA 直接访问缺点占用物理内存分配速度较慢练习题 11CUDA 流的作用是什么cudaMemcpy 和 cudaMemcpyAsync 的区别是什么为什么异步传输需要使用页锁定内存第十二课CUDA 事件与性能计时知识点什么是 CUDA 事件CUDA 事件是 GPU 上的时间标记用于测量内核执行时间流同步性能分析创建和使用事件cudaEvent_t start, stop;// 创建事件cudaEventCreate(start);cudaEventCreate(stop);// 记录事件cudaEventRecord(start);myKernelgrid, block(…);cudaEventRecord(stop);// 等待事件完成cudaEventSynchronize(stop);// 计算时间float milliseconds 0;cudaEventElapsedTime(milliseconds, start, stop);printf(“执行时间: %.3f ms\n”, milliseconds);// 销毁事件cudaEventDestroy(start);cudaEventDestroy(stop);完整的性能测试示例#include stdio.hglobalvoid vectorAdd(float *a, 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 10000000;size_t size n * sizeof(float);// 分配内存 float *h_a (float*)malloc(size); float *h_b (float*)malloc(size); float *h_c (float*)malloc(size); float *d_a, *d_b, *d_c; cudaMalloc(d_a, size); cudaMalloc(d_b, size); cudaMalloc(d_c, size); // 初始化数据 for (int i 0; i n; i) { h_a[i] (float)i; h_b[i] (float)(i * 2); } // 拷贝数据 cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice); // 创建事件 cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); // 测试不同 Block 大小 int blockSizes[] {32, 64, 128, 256, 512, 1024}; int numTests sizeof(blockSizes) / sizeof(int); for (int i 0; i numTests; i) { int blockSize blockSizes[i]; int gridSize (n blockSize - 1) / blockSize; // 预热 vectorAddgridSize, blockSize(d_a, d_b, d_c, n); cudaDeviceSynchronize(); // 计时 cudaEventRecord(start); vectorAddgridSize, blockSize(d_a, d_b, d_c, n); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms; cudaEventElapsedTime(ms, start, stop); printf(Block%4d, Grid%6d, 时间%.3f ms\n, blockSize, gridSize, ms); } // 清理 cudaEventDestroy(start); cudaEventDestroy(stop); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0;}事件同步// 等待单个事件cudaEventSynchronize(event);// 流等待事件cudaStreamWaitEvent(stream, event, 0);// 事件完成检查cudaError_t err cudaEventQuery(event);if (err cudaSuccess) {printf(“事件已完成\n”);}流间同步cudaStream_t stream1, stream2;cudaEvent_t event;cudaStreamCreate(stream1);cudaStreamCreate(stream2);cudaEventCreate(event);// Stream1 记录事件kernel1grid, block, 0, stream1(…);cudaEventRecord(event, stream1);// Stream2 等待事件cudaStreamWaitEvent(stream2, event, 0);kernel2grid, block, 0, stream2(…); // 等待 kernel1 完成练习题 12CUDA 事件的主要用途是什么cudaEventRecord() 和 cudaEventSynchronize() 的区别如何使用事件实现两个流之间的同步第十三课统一内存知识点什么是统一内存统一内存Unified Memory创建一个在 CPU 和 GPU 之间共享的内存池自动管理数据传输。传统内存模型┌─────────┐ ┌─────────┐│ CPU 内存 │ ←─────→ │ GPU 内存 │└─────────┘ 手动传输 └─────────┘统一内存模型┌─────────────────────────────┐│ 统一内存池 ││ CPU 和 GPU 共享访问 ││ 自动管理数据迁移 │└─────────────────────────────┘创建统一内存// 分配统一内存float *data;cudaMallocManaged(data, size);// CPU 和 GPU 都可以直接访问data[0] 1.0f; // CPU 写入myKernelgrid, block(data); // GPU 读取和修改cudaDeviceSynchronize();printf(“%f\n”, data[0]); // CPU 读取 GPU 修改后的值// 释放cudaFree(data);完整示例#include stdio.hglobalvoid add(float *a, 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 1000000;size_t size n * sizeof(float);// 分配统一内存 float *a, *b, *c; cudaMallocManaged(a, size); cudaMallocManaged(b, size); cudaMallocManaged(c, size); // CPU 初始化 for (int i 0; i n; i) { a[i] (float)i; b[i] (float)(i * 2); } // GPU 计算 int blockSize 256; int gridSize (n blockSize - 1) / blockSize; addgridSize, blockSize(a, b, c, n); // 等待 GPU 完成 cudaDeviceSynchronize(); // CPU 验证结果 bool success true; for (int i 0; i n; i) { if (c[i] ! a[i] b[i]) { success false; break; } } printf(验证: %s\n, success ? 成功 : 失败); // 释放 cudaFree(a); cudaFree(b); cudaFree(c); return 0;}内存迁移提示// 提示内存将在设备上访问cudaMemPrefetchAsync(data, size, deviceId, stream);// 提示内存访问模式cudaMemAdvise(data, size, cudaMemAdviseSetPreferredLocation, deviceId);// 预取示例int deviceId;cudaGetDevice(deviceId);// 初始化数据CPUfor (int i 0; i n; i) {data[i] i;}// 预取到 GPUcudaMemPrefetchAsync(data, size, deviceId);// GPU 计算kernelgrid, block(data, n);统一内存的优势优势 说明简化编程 无需手动管理内存拷贝减少代码 不需要 cudaMalloc/cudaMemcpy自动迁移 数据按需在 CPU/GPU 间迁移超额订阅 可使用超过 GPU 内存的数据统一内存的限制性能可能不如手动优化需要调用 cudaDeviceSynchronize() 同步频繁迁移会有开销旧 GPU 不支持某些特性练习题 13统一内存与传统内存模型的主要区别是什么cudaMallocManaged() 分配的内存可以被谁访问为什么使用统一内存后还需要调用 cudaDeviceSynchronize()第十四课常量内存知识点什么是常量内存常量内存是只读内存具有缓存优化适合存储不会改变的数据。常量内存特点大小限制64KB只读有缓存适合广播读取所有线程读取相同地址声明和使用常量内存// 声明常量内存全局作用域constantfloat constData[256];// 从主机拷贝到常量内存cudaMemcpyToSymbol(constData, h_data, size);// 在内核中使用globalvoid kernel(float *output) {int idx blockIdx.x * blockDim.x threadIdx.x;output[idx] constData[idx % 256]; // 读取常量内存}完整示例使用常量内存存储滤波器#include stdio.h#define FILTER_SIZE 5// 声明常量内存constantfloat filter[FILTER_SIZE];globalvoid convolution(float *input, float *output, int n) {int idx blockIdx.x * blockDim.x threadIdx.x;if (idx FILTER_SIZE / 2 idx n - FILTER_SIZE / 2) { float sum 0.0f; for (int i 0; i FILTER_SIZE; i) { sum input[idx - FILTER_SIZE / 2 i] * filter[i]; } output[idx] sum; }}int main() {int n 1000;size_t size n * sizeof(float);// 准备滤波器数据 float h_filter[FILTER_SIZE] {0.1f, 0.2f, 0.4f, 0.2f, 0.1f}; // 拷贝到常量内存 cudaMemcpyToSymbol(filter, h_filter, FILTER_SIZE * sizeof(float)); // 分配内存 float *h_input (float*)malloc(size); float *h_output (float*)malloc(size); float *d_input, *d_output; cudaMalloc(d_input, size); cudaMalloc(d_output, size); // 初始化输入 for (int i 0; i n; i) { h_input[i] (float)i; } // 拷贝数据 cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice); // 执行卷积 int blockSize 256; int gridSize (n blockSize - 1) / blockSize; convolutiongridSize, blockSize(d_input, d_output, n); // 拷贝结果 cudaMemcpy(h_output, d_output, size, cudaMemcpyDeviceToHost); // 清理 cudaFree(d_input); cudaFree(d_output); free(h_input); free(h_output); return 0;}常量内存 vs 全局内存特性 常量内存 全局内存访问权限 只读 读写大小限制 64KB 大缓存 有常量缓存 有L1/L2广播优化 是 否适用场景 常量数据、滤波器、查找表 通用数据何时使用常量内存数据不会改变所有线程读取相同数据广播数据量小于 64KB需要缓存优化练习题 14常量内存的大小限制是多少使用什么函数将数据拷贝到常量内存常量内存适合什么场景第十五课纹理内存知识点什么是纹理内存纹理内存是专门为图像处理优化的只读内存支持硬件插值边界处理缓存优化纹理内存特点纹理内存特性只读支持插值线性、最近邻支持边界模式截断、环绕、镜像2D 空间局部性缓存优化适合图像处理使用纹理内存// 声明纹理引用旧方式CUDA 10 之前texturefloat, 2, cudaReadModeElementType texRef;// 绑定纹理cudaBindTextureToArray(texRef, texArray);// 在内核中读取globalvoid kernel(float *output, int width, int height) {int x blockIdx.x * blockDim.x threadIdx.x;int y blockIdx.y * blockDim.y threadIdx.y;if (x width y height) { // 使用纹理读取支持插值 float val tex2D(texRef, x, y); output[y * width x] val; }}// 解绑纹理cudaUnbindTexture(texRef);现代方式纹理对象CUDA 11// 创建纹理对象cudaResourceDesc resDesc;memset(resDesc, 0, sizeof(resDesc));resDesc.resType cudaResourceTypeLinear;resDesc.res.linear.devPtr devPtr;resDesc.res.linear.desc cudaCreateChannelDesc();resDesc.res.linear.sizeInBytes size;cudaTextureDesc texDesc;memset(texDesc, 0, sizeof(texDesc));texDesc.addressMode[0] cudaAddressModeClamp;texDesc.addressMode[1] cudaAddressModeClamp;texDesc.filterMode cudaFilterModeLinear;texDesc.readMode cudaReadModeElementType;cudaTextureObject_t texObj;cudaCreateTextureObject(texObj, resDesc, texDesc, NULL);// 在内核中使用globalvoid kernel(cudaTextureObject_t texObj, float *output, int n) {int idx blockIdx.x * blockDim.x threadIdx.x;if (idx n) {output[idx] tex1Dfetch(texObj, idx);}}// 销毁纹理对象cudaDestroyTextureObject(texObj);