使用 racecheck 排查竞态
检测共享内存数据竞态
使用 racecheck 排查竞态 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
什么是数据竞争
当两个线程访问同一块共享内存且至少有一个线程执行写入,同时两者之间没有顺序约束时,就会发生数据竞争。😬
为什么竞争难以察觉
竞争在测试中经常给出正确答案,却在生产环境中给出错误答案,因为结果取决于不可预测的线程时序。
认识 racecheck
compute-sanitizer 的racecheck工具会分析共享内存访问,并报告普通运行中可能隐藏的危险。
compute-sanitizer --tool racecheck ./reduce缺少屏障
向共享数组写入后,在没有屏障的情况下读取相邻线程的槽位,会产生竞争,因为相邻位置可能尚未写入。
tile[tid] = in[i];
out[tid] = tile[tid + 1];修复方法:__syncthreads
在写入和读取之间插入__syncthreads,这样线程块中的所有线程都会先完成写入,然后才开始读取。
tile[tid] = in[i];
__syncthreads();
out[tid] = tile[tid + 1];写后写危险
racecheck 还会捕获写后写危险,即两个线程在没有顺序约束的情况下写入同一个共享槽位。
读后写危险
最常见的危险是读后写:一个线程读取另一个线程仍在写入的值,因此读到的是过时数据。
阅读危险报告
racecheck 会列出危险类型、涉及的两个线程以及共享地址,让您准确知道缺少哪个屏障。
它监视共享内存
racecheck 关注线程块内部的__shared__内存;全局内存竞争通常需要仔细设计逻辑或改用原子操作。
添加行信息
与 memcheck 一样,使用-lineinfo进行编译后,racecheck 就能指出每个危险所在的确切源代码行。
nvcc -lineinfo reduce.cu -o reduce验证修复结果
添加屏障后重新运行 racecheck。报告零个危险意味着共享内存访问现在已经具有正确的顺序。
快速检查
当 racecheck 报告共享内存危险时,通常应如何修复?
回顾
您使用racecheck查找共享内存危险,阅读了相关报告,并使用 __syncthreads 屏障修复了这些问题。时序错误已被揭示。🎯
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「使用 racecheck 排查竞态」课时是免费的吗?
是的 — 「使用 racecheck 排查竞态」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「使用 racecheck 排查竞态」这节课中我会学到什么?
检测共享内存数据竞态 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 3 节课,共 4 节。
「使用 racecheck 排查竞态」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- 在 cuda-gdb 中单步执行内核
- 使用 memcheck 查找内存泄漏
- 使用 racecheck 排查竞态
- 使用 synccheck 捕获同步错误