GPU 上的竞态条件
了解并发写入为何会破坏数据
GPU 上的竞态条件 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 1 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
本课时的部分内容尚未翻译,以英文显示。
Many Hands, One Counter
On the GPU, thousands of threads can touch the same memory at once. When several write to one spot together, you get a race condition. 🏁
What a Race Really Is
A race condition happens when the final result depends on the unpredictable order in which threads run. Same code, different answer each launch.
Read-Modify-Write
Incrementing a counter is really three steps: read the value, add one, write it back. This read-modify-write sequence is where things break.
counter = counter + 1;How the Steps Interleave
Two threads both read 5, both compute 6, both write 6. You expected 7 but got 6. One increment silently vanished.
The Lost Update
That vanished increment is called a lost update. With thousands of threads racing, dozens or hundreds of updates can disappear at once.
A Kernel That Looks Fine
This kernel looks correct, but every thread races on the same address. The final count will be wrong and will vary between runs.
__global__ void count(int* total) {
*total = *total + 1;
}Why It Is Nondeterministic
The hardware never promises a thread order. So a racy kernel is nondeterministic: it may even pass on small inputs and fail on big ones.
Reads Alone Are Safe
Many threads reading the same value is perfectly fine. Trouble starts only when at least one thread writes while others read or write.
Disjoint Writes Are Safe Too
If each thread writes its own unique slot, like out[i], there is no conflict. A race needs threads aiming at the same location.
out[i] = a[i] + b[i];The Fix Preview
The cure is to make read-modify-write happen as one indivisible step. That is an atomic operation, coming up in the next lesson. ⚛️
Spotting Races in Review
When reviewing a kernel, ask: do two threads write the same address without protection? If yes, you almost certainly have a data race.
Quick Check
Let's make sure the race idea clicked.
Recap: Races on the GPU
You learned that unguarded read-modify-write on shared data causes race conditions and lost updates. Atomics, up next, make those steps indivisible. ✅
常见问题解答
「GPU 上的竞态条件」课时是免费的吗?
是的 — 「GPU 上的竞态条件」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「GPU 上的竞态条件」这节课中我会学到什么?
了解并发写入为何会破坏数据 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 1 节课,共 4 节。
「GPU 上的竞态条件」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- GPU 上的竞态条件
- atomicAdd 及相关操作
- 构建直方图
- 使用 atomicCAS 自定义原子操作