CUDA Academy · 课时

使用 racecheck 排查竞态

检测共享内存数据竞态

第 3 / 4 课13 个步骤

使用 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 反馈 — 无需本地设置。

此课程中的所有课时

  1. 在 cuda-gdb 中单步执行内核
  2. 使用 memcheck 查找内存泄漏
  3. 使用 racecheck 排查竞态
  4. 使用 synccheck 捕获同步错误
← 返回 CUDA Academy