Skip to content

Commit

Permalink
feat: detail usage of compute-sanitizer in 2.apsp.md
Browse files Browse the repository at this point in the history
  • Loading branch information
oliverYoung2001 committed May 10, 2024
1 parent 9c39819 commit 44fbf6f
Showing 1 changed file with 1 addition and 1 deletion.
2 changes: 1 addition & 1 deletion docs/exp/2.apsp.md
Original file line number Diff line number Diff line change
Expand Up @@ -106,7 +106,7 @@ for (int k = p * b; k < (p + 1) * b; k++)
- 记得在适当位置添加 `__syncthreads()`(潜在优化点:最短路算法的某一步中,就算访问旧的 shared memory 值,也可证明答案是正确的,此时可以不做同步;如实现此项优化,鼓励在报告中提供证明);
- 注意检查 kernel 和 CUDA API(例如 `cudaMemcpy`)调用中的错误(kernel 参数不合法、访存越界等错误时常体现为答案错误而非运行错误,此时需要注意排查):
- 注意检查错误返回值(可使用 `cuda_utils.h` 中的 `CHK_CUDA_ERR` 宏);
- 可使用 `compute-sanitizer`(旧称 `cuda-memcheck`)调试工具自动检查错误;
- 可使用 `compute-sanitizer`(旧称 `cuda-memcheck`)调试工具自动检查错误,编译时添加`-lineinfo`可定位到具体行号
- 对于 kernel 调用时产生的错误,非法参数(kernel 根本没跑起来)可以在 kernel 调用语句后用 `cudaGetLastError` 检查;而 kernel 运行时错误(kernel 跑起来了但出错了)会体现在 kernel 调用后的第一个同步语句(包括 `cudaDeviceSynchronize``cudaMemcpy`)的返回值中;
- 不需要 global memory 上的同步时,尽量将循环置于 kernel 内,以减少不必要的 kernel 启动次数;
- 调整 `threadIdx.x``threadIdx.y` 的顺序,或者调整数组的行优先/列优先方向,更好地适应全局内存的缓存行(cache line),以及避免(对性能的影响极其严重的)共享内存的 bank conflict 现象。
Expand Down

0 comments on commit 44fbf6f

Please sign in to comment.