企业官网建设流程全解析

学习目标
学完本节你将能够：

• 识别 CUDA 入门阶段最高频的编码错误和运行时错误
• 理解每类错误产生的根本原因
• 掌握对应的排查方法和修复方案
• 形成良好的 CUDA 编程习惯，避免在后续算子开发中重复踩坑

1. 错误分类总览

入门阶段的常见错误可以分为以下几大类：

2. 高频错误详解与修复

2.1 索引越界

错误场景：

__global__ void badIndex(int *arr, int N) { int idx = blockIdx.x * blockDim.x + threadIdx.x; arr[idx] = idx * 2; // 没有边界判断！ } // 如果 N 不是 blockSize 的整数倍，部分线程会访问 arr[N] 及之后的内存

后果：

• 写入非法地址，触发illegal memory access
• 可能污染其他数据，导致结果错误但不报错

修复：

__global__ void goodIndex(int *arr, int N) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < N) { // 必须加边界判断 arr[idx] = idx * 2; } }

2.2 线程数分配不合理

错误场景 1：Block 大小不是 32 的倍数

kernel<<<gridSize, 100>>>(...); // 最后一个 Warp 只有 4 个活跃线程

浪费硬件资源，降低 SM 占用率。

错误场景 2：Grid 过大或过小

kernel<<<1000000, 32>>>(...); // 启动海量 Block，调度开销大 kernel<<<1, 1024>>>(...); // 只用一个 Block，无法充分利用多 SM

修复原则：

• Block 大小设为 32 的倍数，推荐 128、256、512
• Grid 大小根据数据规模和 Block 大小计算：gridSize = ceil(N / blockSize)
• 超大数组使用 Grid‑Stride Loop 控制 Block 总量

2.3 同步问题

错误场景 1：条件分支中的__syncthreads()

if (threadIdx.x < 16) { __syncthreads(); // 死锁！ }

错误场景 2：忘记__syncthreads()

__shared__ float tile[256]; tile[tid] = in[tid]; // 没有 __syncthreads() float result = tile[255]; // 可能读到未写入的数据

错误场景 3：循环中提前退出

for (int i = 0; i < N; i++) { __syncthreads(); if (condition) return; // 部分线程提前退出，后续 __syncthreads 死锁 }

修复原则：

• __syncthreads()必须在所有线程都能执行到的路径上
• 共享内存写入后、读取前必须同步
• 循环内禁止条件性退出，边界判断放在循环体内部

2.4 异步错误未捕获

错误场景：

kernel<<<grid, block>>>(...); // 没有 cudaDeviceSynchronize()，直接检查错误 cudaError_t err = cudaGetLastError(); // 可能捕获不到运行时错误

后果：

• Kernel 执行期错误（如越界访问）不会被及时发现
• 错误状态可能被后续 CUDA API 调用覆盖

修复：

kernel<<<grid, block>>>(...); CUDA_CHECK(cudaGetLastError()); // 启动错误 CUDA_CHECK(cudaDeviceSynchronize()); // 执行期错误

2.5 内存泄漏

错误场景：

for (int iter = 0; iter < 1000; iter++) { float *d_data; cudaMalloc(&d_data, 1 << 20); // 使用 d_data... // 忘记 cudaFree(d_data) ！ }

后果：

• 显存逐渐耗尽，最终out of memory
• 长期运行的服务程序尤其严重

修复原则：

• 每次cudaMalloc必须有对应的cudaFree
• 使用 RAII 封装（如thrust::device_vector）自动管理显存
• 定期用cudaMemGetInfo检查显存使用量

2.6 指针混用

错误场景 1：设备指针传给主机函数

int *d_arr; cudaMalloc(&d_arr, N * sizeof(int)); printf("%d", d_arr[0]); // 错误！主机不能直接解引用设备指针

错误场景 2：主机指针传给 Kernel

int h_arr[100]; kernel<<<1, 1>>>(h_arr); // 错误！Kernel 不能访问主机内存

错误场景 3：cudaMemcpy 方向传反

cudaMemcpy(h_arr, d_arr, bytes, cudaMemcpyHostToDevice); // 错误！源是设备指针，目标主机指针，方向应为 DeviceToHost

修复原则：

• 设备指针只能通过 CUDA API 操作
• 主机指针和设备指针不能直接互换
• cudaMemcpy方向必须与源/目标指针类型匹配

2.7 寄存器溢出

错误场景：

__global__ void manyVariables(float *out) { float a0 = 0, a1 = 1, a2 = 2, ...; // 上百个变量 // 编译器可能将部分变量放到局部内存 out[threadIdx.x] = a0 + a1 + ...; }

检测：使用--ptxas-options=-v编译，查看 local memory 是否大于 0。

后果：溢出变量存储在设备内存中，访问延迟远高于寄存器，性能大幅下降。

修复：

• 减少不必要的局部变量
• 复用变量
• 将数据放入共享内存
• 使用-maxrregcount限制寄存器（但需谨慎，可能导致更多溢出）

2.8 共享内存 Bank Conflict

错误场景：

__shared__ float sdata[32][32]; int idx = threadIdx.x; float val = sdata[idx][idx * 2]; // 多个线程访问同一 Bank 的不同地址

后果：共享内存访问串行化，性能下降数倍。

检测：使用 ncu（Nsight Compute）分析 Bank Conflict。

修复：

• 使用连续访问模式（sdata[idx]）
• 添加 padding 避免 Bank 冲突（如sdata[32][33]）
• 使用洗牌函数替代共享内存（Warp 内数据交换）

2.9 浮点精度问题

错误场景：

// Kernel 中直接累加 float sum = 0.0f; for (int i = 0; i < 1000000; i++) { sum += data[i]; // 大数吃小数，精度损失 }

后果：GPU 结果与 CPU 串行累加结果不一致，误差随累加次数增大。

修复：

• 使用double提高精度
• 使用分块归约减少累加误差
• 使用__fadd_rn等显式舍入函数（高级用法）

3. 排查方法总结

遇到 CUDA 错误时，按以下顺序排查：

1. 编译错误：查看 NVCC 输出，定位语法和类型错误
2. 启动错误：cudaGetLastError()检查配置是否合法（Block 大小、Grid 维度）
3. 执行期错误：cudaDeviceSynchronize()后检查错误
4. 结果错误：
   • 检查边界判断
   • 检查同步是否缺失
   • 检查数据竞争
5. 性能问题：
   • 使用--ptxas-options=-v查看寄存器溢出
   • 使用cudaEvent或 ncu 定位瓶颈
   • 检查 Block 大小是否合理
   • 检查内存访问模式是否为合并访问

4. 课后练习

练习1：越界复现

编写一个 Kernel，故意去掉边界判断，处理 N = 1000、blockSize = 256 的数据，运行并观察错误信息。

练习2：死锁复现

将__syncthreads()放在if (threadIdx.x < 16)中，运行程序，观察是否卡死（死锁）。使用 timeout 命令限制运行时间，避免程序一直挂起。

练习3：寄存器溢出检测

编写一个 Kernel，声明 200 个 float 局部变量并累加。使用--ptxas-options=-v编译，查看寄存器数量和局部内存溢出量。然后优化变量使用，再次查看。

练习4：内存泄漏模拟

在循环中分配设备内存但不释放，运行程序，使用nvidia‑smi观察显存占用逐渐增加，最终触发 out of memory 错误。

练习5：Bank Conflict 实验

编写一个 Kernel，使用共享内存，分别测试连续访问模式和跨步访问模式，用 cudaEvent 计时对比性能差异。

5. 下一步

恭喜你完成 CUDA 基础执行模型入门板块的全部 8 节内容！

下一板块将进入 GPU 存储层级与内存访问规则，深入讲解：

• 全局内存合并访问的硬件机制
• 共享内存 Bank Conflict 的详细规则与优化
• 常量内存和纹理内存的高级用法
• 内存访问模式对性能的定量影响

掌握这些内容后，你将能够编写出高效利用 GPU 内存带宽的 Kernel。