Shared memory and barriers

Phases

Consecutive local loops in one global loop are phases. Between two phases, every thread of the block has finished the first before any thread starts the second. Run as plain loops, that's automatic: the first loop ends before the second starts. On a GPU, every thread runs both phases one after the other, so the compiler puts a barrier, __syncthreads(), between them.

Here is a block of 8 reversing its part of x through shared memory:

%sh = alloc 8 dtype=f32                # the block's shared memory
%b = range 4 axis=global
# phase 1: thread t copies x[8b + t] into sh[t]
%t = range 8 axis=local
...
%e1 = end %s1 %t
%sh1 = after %sh %e1
# phase 2: thread u reads sh[7 - u] and writes o[8b + u]
%u = range 8 axis=local
...

In phase 2, thread u reads a slot another thread wrote. Inside one phase that would be a race. Across the barrier it's safe: every slot was written before any thread got past it. So the rule from lesson 2 is checked one phase at a time, and a later phase may read anything an earlier phase wrote to shared memory.

Buffers are different. Other blocks run at the same time, and no barrier waits for them, so the kernel check covers all of a block's phases at once: every access to a buffer in the kernel uses the same offset expression. That's why the reverse goes through sh instead of reading back o.

Shared memory, in .uir

An alloc whose every load and store is inside one global loop is that block's shared memory. Each block starts with its own copy, all zeros. run re-zeroes it for every value of g, just as week 7 re-zeroed a thread loop's private alloc. The test block_shared_starts_at_zero checks that, and plain_loop_alloc_carries_over runs the same program with plain loops, where the alloc keeps its contents.

On the GPU, shared memory isn't zero when a block starts. It may hold an earlier block's leftovers. So the rendered kernel zeroes it first, with every thread clearing part of it, and then a barrier. Skip that and a test like c_shared_dirty_then_fresh reads another block's leftovers.

The kernel check, which proves blocks don't collide, skips shared memory: each block has its own, so block 0's sh[0] and block 1's sh[0] are different memory.

Phases in a loop

A tiled GEMM puts its phases in a plain loop over tiles of k: load a tile of A and B into shared memory, multiply it in, then load the next. That needs a barrier at the end of the last phase too. Otherwise a fast thread starts loading tile k + 1 while a slow one is still multiplying tile k. That's why the CUDA renderer puts a barrier after every phase, including the last.

A missing barrier

A missing barrier gives wrong answers only sometimes. On a real GPU the threads of a warp usually run together, though nothing promises it, so a small block of one or two warps often gets lucky, and the test passes on your machine. That's why the suite includes a block of 1024 threads, c_reverse_shared_1024, whose 32 warps really do run at different times. The emulator runs every thread of a block as a real operating-system thread, with a real barrier, so a missing barrier shows up there too.