Synchronization, streams, and profiling
Barriers make shared memory correct, divergence makes lanes idle, streams overlap independent work — and a profiler tells you which of them actually costs you.
A block of threads shares a small scratchpad of memory but has no shared
program counter. So the moment one thread writes to shared memory and another
reads it, you have a race — the same lost-update problem you will meet on the
CPU, just on silicon. CUDA's fix is a barrier: __syncthreads(). Every thread
in the block must arrive before any proceeds. It is the synchronisation point
that makes the load-tile, compute, load-tile pattern correct.
Synchronisation is where performance leaks, though. A barrier stalls every warp in the block until the slowest arrives, and warp divergence — lanes of a warp choosing different branch paths — serialises the paths while idle lanes wait. Both are the tax you pay for structure, and both are invisible until you measure them. This lesson is about keeping the structure correct while finding where the time actually goes.
Three independent costs
Barriers cost you by stalling. Divergence costs you by idling lanes. Serial kernel launches cost you by leaving the machine empty between them. They are three different problems with three different fixes — smaller barriers, branch-free code, and streams — and a profiler is how you find out which one you have.
A checkpoint in a relay
__syncthreads() is a checkpoint where every runner in a team must arrive
before anyone continues. It keeps the team together; it also means the
slowest runner sets the pace for the whole team at every checkpoint. A
barrier you did not need is pure waiting.
That checkpoint image is worth holding onto. A barrier is correct by construction — no read can overtake the write it depends on — and expensive by construction, because the block advances at the speed of its slowest warp.
Split the warp and then overlap kernels below.
Split a warp on a branch, then overlap independent kernels with streams
warp divergence · 2 passes · divergent
50% lane efficiency · 32 wasted lane-slots
execution passes
stream overlap · 6 kernels on 2 streams
6 ms vs 12 ms serial · 2.0×
Divergence is why you restructure a kernel so neighbouring threads take the same branch, or replace the branch with predication. Streams are the other lever: two independent kernels on one stream run one after another, but on separate streams the scheduler can fill idle units. The analogy and timings are illustrative — real overlap depends on resource use.
Barriers are block-scoped
__syncthreads() synchronises one block, nothing more. Blocks are independent
and may run in any order, so there is no global barrier inside a kernel — if
you need one, you finish the kernel and launch another. This is deliberate: it
is what lets the scheduler run thousands of blocks without a global
coordination cost. The rules that follow:
- Never call
__syncthreads()inside divergent control flow. If only some lanes reach the barrier, the others never will and the block hangs. - Barrier count must be identical for every thread in the block.
- Shared memory overwritten before the barrier is a race even if it looks ordered.
Divergence
A warp executes in lockstep. If lanes need different branches, the hardware runs the taken path for the active lanes while the rest wait, then runs the other path — effectively serialising the branches. The measured cost is the serialisation of distinct paths, not the number of branches:
A warp where every lane takes a different path degenerates to 32 serial passes, 1/32 efficiency. The fixes are structural: sort or partition data so neighbouring lanes follow the same path, replace short branches with predication, or choose a different algorithm. Irregular problems (sparse matrices, graph traversal, variable-length sequences) are where this bites hardest, and it is the main reason a GPU can be slower than a CPU on graph-like data.
Streams and overlap
By default everything you launch goes on one stream, in order: kernel A must finish before kernel B starts. If A and B are independent, that serialisation wastes the machine. Multiple streams let the scheduler overlap them — B's blocks fill idle units while A finishes — and let you overlap host-to-device copies with compute. The limit is resources: two kernels that both saturate the arithmetic units will not overlap, they will queue.
Profiling before optimising
The discipline is the same as CPU optimisation, and the same as Amdahl's law: find the bottleneck before touching code. Nsight Compute reports, per kernel, utilisation of compute versus memory, achieved occupancy, warp stall reasons, and divergence. The workflow:
- Measure the end-to-end time and find the dominant kernel.
- Classify it: compute bound, memory bound, or latency bound (occupancy).
- Apply the matching fix — more arithmetic per byte, better coalescing, or more resident warps.
- Re-measure. Optimisation without a before/after number is folklore.
Careful
A barrier is not free and not always safe. __syncthreads() inside a branch
that not all lanes take is a hang, and putting barriers in a loop that
executes too often can cost more than the race it prevents. Equally, more
streams are not always faster: oversubscribing the device makes kernels
contend for the same units and adds launch overhead without overlap. Measure.
Illustrative vs real
The lane grid assumes a 32-wide warp with a single two-way branch, and the stream timeline assumes each kernel is independent and evenly sized. Real divergence can be nested and partially predicated, and real overlap depends on registers, shared memory and scheduler policy. The formulas are exact within the toy model; the hardware is messier.
Check yourself
Eduspheria wiki · Systems for AI, GPU computing
0 / 5 answered
From the mid-term paper
Modeled on NITJ AI-619, Mid-Term March 2025
0 / 5 answered
Where next: the same correctness problem on the CPU — threads, races, and the synchronization primitives that fix them.