When is shared memory ready to read?

Imagine a block of 64 threads reversing a tile of 64 numbers. Thread 0 needs the value that thread 63 brought in; thread 63 needs the value that thread 0 brought in. They agree on the addresses. What tells one thread that another thread’s value is ready?

A place to share a value

A register belongs to a thread. If thread 63 loads a number into one of its registers, thread 0 cannot read that register by naming it. For this example the block uses a shared array, tile[64]. Thread t writes its input into tile[t]. Now every thread in this block can name the slot where its partner’s value belongs. The next block has its own tile.

Visibility has a timing component. Individual threads in a block can reach the load and store instructions at different times. There are two 32-thread warps here; the fact that both belong to one block does not tell us which one gets to its store first.

A time to read it

// One block of 64 threads; one complete tile.
__shared__ int tile[64];
int t = threadIdx.x;
tile[t] = input[t];
__syncthreads();
output[t] = tile[63 - t];

Suppose input[0] is 12 and input[63] is 90. The first phase leaves 12 in slot 0 and 90 in slot 63. The block barrier separates this phase from the reads: when the threads proceed, the shared writes before the barrier are visible to their peers. Thread 0 therefore reads 90. Without the barrier, a warp might reach the read of slot 63 before the other warp has initialized that slot. The code has a race, even if one particular run happens to return 90.

Why edge tiles need care

A short final tile changes who owns data. It need not change who reaches the barrier. For a block with 60 valid inputs, one design has all 64 threads participate, initializes the four unused slots to a defined value, and lets only the 60 valid threads write output. If the goal is to reverse just those 60 entries, the partner index is 59 - t. Keeping 63 - t would feed padding to valid threads. A correct barrier cannot fix the wrong partner index.

The scope matters too. This barrier coordinates this block; it does not mean every block in the grid has finished its tile. If a later phase needs results from every block, a simple design is a second kernel launched after the first in the same stream. See NVIDIA’s thread hierarchy and synchronization guide.