4.8 KiB

Raw Permalink Blame History

常见错误与排错指南

编译错误

nvcc 版本不匹配

症状：CMake Error: nvcc not found 或 nvcc fatal: Unsupported gpu architecture

原因：CUDA Toolkit 版本与 CMake 预设的架构参数不兼容。

解决：

检查 nvcc 版本：nvcc --version
在 CMakePresets.json 中将 CMAKE_CUDA_ARCHITECTURES 改为 "native"，或显式指定你的 GPU 架构编号
确保 nvcc 在 PATH 中

CUDA 架构参数错误

症状：cudaErrorNoKernelImageForDevice 或运行时 kernel launch 失败

原因：编译时指定的 GPU 架构（如 89 对应 L40）与运行时 GPU 不匹配。

解决：

将 CMAKE_CUDA_ARCHITECTURES 设为 "native" 让 CMake 自动检测
或根据 GPU 型号查表设置正确的架构编号

CUTLASS 拉取失败

症状：FetchContent failed to download CUTLASS

原因：GitHub 网络不可达或代理问题。

解决：

检查网络：git ls-remote https://github.com/NVIDIA/cutlass.git
配置代理：git config --global http.proxy http://your-proxy:port
或手动下载 CUTLASS v3.7.0 放到 third_party/cutlass/

运行时错误

越界访问（grid-stride loop 边界条件）

症状：随机数值错误、CUDA illegal memory access 错误

原因：idx < N 条件缺失或写错，导致线程访问超出 tensor 范围的内存。

解决：

// 正确写法：每次循环都检查 idx < N
for (int idx = blockIdx.x * blockDim.x + threadIdx.x;
     idx < N;
     idx += gridDim.x * blockDim.x) {
    out[idx] = in[idx];  // 安全：idx 始终 < N
}

`__syncthreads()` 在条件分支内

症状：kernel 在 reduce_sum 或 softmax 中挂起（hang），或结果错误

原因：__syncthreads() 放在 if 分支内。CUDA 要求 block 内所有线程都到达同一个 __syncthreads()，如果有线程走 else 分支跳过同步点，整个 block 就会死锁。

解决：

// 错误
if (threadIdx.x < N) {
    smem[tid] = val;
    __syncthreads();  // 部分线程不执行，死锁！
}

// 正确
smem[tid] = (threadIdx.x < N) ? val : 0;
__syncthreads();  // 全部线程都到达

shared memory 大小不足

症状：编译错误 uses too much shared data

原因：申请的 shared memory 超过 GPU 的物理限制（通常 48KB-164KB/block）。

解决：

检查 extern __shared__ 声明的数组大小
减小 block size 或 shared memory 使用量
分多轮处理

数值错误

softmax 未减 max（大数值溢出）

症状：softmax 输出全为 NaN 或 inf

原因：exp(x) 在 x > 88 时溢出为 inf，直接除 inf 得到 NaN。

解决：先减行最大值再做 exp：

float max_val = row[0];
for (int j = 1; j < N; j++) {
    max_val = fmaxf(max_val, row[j * stride]);
}
float sum = 0;
for (int j = 0; j < N; j++) {
    sum += expf(row[j * stride] - max_val);
}

浮点精度差异（多后端对比）

症状：NVIDIA 和 TileLang 的输出在小数点后几位不一致，测试报 not close

原因：不同后端使用不同数学库（exp2/log2 vs exp/log），或规约顺序不同导致浮点累加误差。

解决：

检查测试容差设置是否合理（FP16 容差应比 FP32 宽松）
确认两边的算法逻辑一致（如 online vs 三趟 softmax）
规约顺序差异导致的误差在合理范围内（rtol=1e-3 for FP16）则正常

规约顺序影响结果

症状：reduce_sum 结果每次运行都略有不同

原因：浮点加法不满足结合律，tree reduction 的配对顺序影响结果。

解决：这是正常的浮点行为，只要误差在容差范围内即可接受。

性能问题

bank conflict

症状：reduce_sum kernel 带宽利用率远低于预期（如 < 50%）

原因：shared memory 访问时多个线程命中同一 bank。

解决：

添加 padding 偏移访问地址
使用 __shared__ float smem[BLOCK_SIZE + PADDING] 并错开 bank
将 stride 从 1 改为 2 开始归约

线程利用率低

症状：benchmark 带宽利用率低，但代码逻辑正确

原因：grid/block 尺寸选择不当，GPU SM 上的线程不够填满所有 core。

解决：

block size 建议 128/256/512（取 32 的倍数）
grid size 建议 (N + block_size - 1) / block_size 或更大
用 cudaOccupancyMaxPotentialBlockSize 自动计算最佳配置

全局内存未合并访问

症状：copy 或 vector_add 带宽利用率远低于峰值

原因：线程访问的内存地址不连续，导致多次内存事务而非一次合并访问。

解决：

确保相邻线程访问相邻内存地址（lane 0→地址 0，lane 1→地址 1…）
检查 stride 是否为 1
检查数据类型是否与访问模式对齐

4.8 KiB Raw Permalink Blame History Unescape Escape