atomicAdd 及相关操作
硬件级读取—修改—写入操作
atomicAdd 及相关操作 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 2 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
一个不可分割的步骤
原子操作会将读取、修改和写入值作为一个步骤完成，其他线程无法在中途打断它。这样就能消除竞争。⚛️
认识 atomicAdd
atomicAdd 接收一个指针和一个值，安全地执行加法，并返回旧值。多个线程可以访问同一个地址，而不会丢失更新。
atomicAdd(total, 1);修复存在竞争的计数器
将普通的递增操作替换为 atomicAdd 后，无论运行多少线程，计数结果每次都会正确。
__global__ void count(int* total) {
atomicAdd(total, 1);
}返回值很有用
atomicAdd 返回您执行加法之前所保存的值。这个旧值是一个唯一位置，您可以将它用作索引或票号。
int slot = atomicAdd(counter, 1);完整的操作家族
atomicAdd 有许多伙伴：atomicSub、atomicMax、atomicMin、atomicExch、atomicAnd、atomicOr 和 atomicXor 都以相同的原子方式工作。
交换一个值
atomicExch 会在一个步骤中存储新值并返回旧值。它非常适合安全地获取标志的控制权或读取之前的状态。
int prev = atomicExch(flag, 1);原子最大值和最小值
只有当您的值更大时，atomicMax 才会更新目标值。它非常适合在没有竞争的情况下查找所有线程中的全局最大值。
atomicMax(best, myValue);支持的类型
原子操作支持 int、unsigned，在较新的硬件上还支持 float 和 double。请始终检查您的类型和架构是否支持所需的调用。
原子操作也有代价
原子操作会将访问同一地址的线程串行化。一个位置上的严重争用会成为瓶颈，因此请只在需要时使用原子操作。
分散负载
让线程更新不同的地址，然后再合并结果，可以减少争用。冲突越少，原子操作就越快。
原子操作与归约
对于超大型数组求和，树形归约通常优于大量原子操作。当更新稀疏或不规则时，再选择原子操作。
快速检查
检查一下您对 atomicAdd 工作方式的掌握程度。
总结：atomicAdd 及其伙伴
您认识了原子操作家族，修复了存在竞争的计数器，并了解到原子操作会返回旧值，但在争用严重时需要付出时间代价。✅
常见问题解答
「atomicAdd 及相关操作」课时是免费的吗？
是的 — 「atomicAdd 及相关操作」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「atomicAdd 及相关操作」这节课中我会学到什么？
硬件级读取—修改—写入操作 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 2 节课，共 4 节。
「atomicAdd 及相关操作」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- GPU 上的竞态条件
- atomicAdd 及相关操作
- 构建直方图
- 使用 atomicCAS 自定义原子操作