atomicAdd and Friends
Hardware read-modify-write operations.
atomicAdd and Friends is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
One Indivisible Step
An atomic operation reads, modifies, and writes a value as a single step that no other thread can interrupt. That kills the race. ⚛️
Meet atomicAdd
atomicAdd takes a pointer and a value, adds safely, and returns the old value. Many threads can hit the same address with no lost updates.
atomicAdd(total, 1);Fixing the Racy Counter
Swap the plain increment for atomicAdd and the count becomes correct every single time, no matter how many threads run.
__global__ void count(int* total) {
atomicAdd(total, 1);
}The Return Value Is Useful
atomicAdd returns the value held before your add. That old value is a unique slot you can use as an index or a ticket number.
int slot = atomicAdd(counter, 1);A Whole Family
atomicAdd has many friends: atomicSub, atomicMax, atomicMin, atomicExch, atomicAnd, atomicOr, and atomicXor all work the same atomic way.
Swapping a Value
atomicExch stores a new value and returns the old one in one step. It is handy for claiming flags or grabbing a previous state safely.
int prev = atomicExch(flag, 1);Atomic Max and Min
atomicMax updates the target only if your value is larger. It is perfect for finding a global maximum across all threads without races.
atomicMax(best, myValue);Supported Types
Atomics cover int, unsigned, and on newer hardware float and double too. Always check that your type and arch support the call you want.
Atomics Have a Cost
Atomics serialize threads that target the same address. Heavy contention on one spot becomes a bottleneck, so use them only where needed.
Spread the Load
Reduce contention by having threads update different addresses, then combine results. Fewer collisions means atomics stay fast.
Atomics vs Reductions
For summing huge arrays, a tree reduction often beats a flood of atomics. Reach for atomics when updates are sparse or irregular.
Quick Check
Test your grip on atomicAdd's behavior.
Recap: atomicAdd and Friends
You met the atomic family, fixed the racy counter, and learned atomics return the old value but cost time under contention. ✅
Frequently asked questions
Is the “atomicAdd and Friends” lesson free?
Yes — the full text of “atomicAdd and Friends” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “atomicAdd and Friends”?
Hardware read-modify-write operations. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “atomicAdd and Friends” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Race Conditions on the GPU
- atomicAdd and Friends
- Building a Histogram
- Custom Atomics with atomicCAS