CUDA Academy · 课时

使用 racecheck 排查竞态

检测共享内存数据竞态

使用 racecheck 排查竞态 是 CoddyKit 上的免费 CUDA Academy 课时。这是第 3 节课，共 4 节。你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

本课时的部分内容尚未翻译，以英文显示。

What a Data Race Is

A data race happens when two threads touch the same shared memory and at least one writes, without ordering between them. 😬

Why Races Are Sneaky

Races often give the right answer in testing and the wrong one in production, because the result depends on unpredictable thread timing.

Meet racecheck

The racecheck tool of compute-sanitizer analyzes shared memory accesses and reports hazards that a normal run would hide.

compute-sanitizer --tool racecheck ./reduce

A Missing Barrier

Writing to a shared array and then reading a neighbor's slot without a barrier creates a race, since the neighbor may not be written yet.

tile[tid] = in[i];
out[tid] = tile[tid + 1];

The Fix: __syncthreads

Insert __syncthreads between the write and the read so every thread in the block finishes writing before anyone reads.

tile[tid] = in[i];
__syncthreads();
out[tid] = tile[tid + 1];

Write-After-Write Hazards

racecheck also catches a write-after-write hazard, where two threads store to the same shared slot with no order between them.

Read-After-Write Hazards

The most common hazard is read-after-write: a thread reads a value another thread is still writing, so it reads stale data.

Reading the Hazard Report

racecheck names the hazard type plus the two threads and shared address involved, so you know exactly which barrier is missing.

It Watches Shared Memory

racecheck focuses on __shared__ memory within a block; global races usually need careful logic or atomics instead.

Add Line Info

Just like memcheck, compiling with -lineinfo lets racecheck point at the precise source line of each hazard.

nvcc -lineinfo reduce.cu -o reduce

Verify the Fix

Rerun racecheck after adding the barrier. A report of zero hazards means your shared-memory access is now properly ordered.

Quick Check

What is the usual fix when racecheck reports a shared-memory hazard?

Recap

You used racecheck to find shared-memory hazards, read its reports, and fixed them with __syncthreads barriers. Timing bugs, exposed. 🎯

常见问题解答

「使用 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 Academy