Skip to content

Fix producer-consumer locking mechanism in Example 17 using atomic_add and memory fences #267

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Copilot wants to merge 5 commits into
base: main
Choose a base branch
from
Open

Fix producer-consumer locking mechanism in Example 17 using atomic_add and memory fences #267

Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
beafa1c
Initial plan
Copilot Oct 19, 2025
3fa7d2b
Fix producer-consumer locking with atomic_add and proper synchronization
Copilot Oct 19, 2025
fd45d66
Fix synchronization to wait for world_size instead of world_size-1
Copilot Oct 19, 2025
cf9e546
Simplify consumer wait loop to avoid resetting tile_ready
Copilot Oct 19, 2025
8d7c7e1
Add memory fences to ensure proper visibility of local_C writes
Copilot Oct 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 21 additions & 13 deletions examples/17_gemm_one_shot_all_reduce_pc/gemm_one_shot_all_reduce_pc.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,14 +127,16 @@ def persistent_gemm(
# Write to local buffer
tl.store(local_C + local_offset, acc, mask=mask, cache_modifier=".wt")

# Ensure local_C write is visible before signaling
tl.atomic_fence(sem="release", scope="sys")

# Signal that this tile is ready
tl.debug_barrier()
tl.store(locks + tile_id, 1, cache_modifier=".wt")

# Signal to all remote ranks that this tile is ready
# Signal to all remote ranks that this tile is ready by incrementing their counter
for remote_rank in range(world_size):
if remote_rank != cur_rank:
iris.atomic_xchg(tile_ready + tile_id, 1, cur_rank, remote_rank, heap_bases, sem="release", scope="sys")
iris.atomic_add(tile_ready + tile_id, 1, cur_rank, remote_rank, heap_bases, sem="release", scope="sys")

if COLLECT_TIMESTAMPS:
timestamp = read_realtime()
Expand Down Expand Up @@ -213,16 +215,22 @@ def persistent_all_reduce(
while tl.load(locks + tile_id, cache_modifier=".cv", volatile=True) != 1:
pass

# Wait for remote ranks
for remote_rank in range(world_size):
if remote_rank != cur_rank:
while (
iris.atomic_cas(
tile_ready + tile_id, 0, 0, cur_rank, remote_rank, heap_bases, sem="acquire", scope="sys"
)
!= 1
):
pass
# Ensure local producer's writes are visible
tl.atomic_fence(sem="acquire", scope="gpu")

# Wait for remote ranks - each remote rank increments tile_ready when done
# We expect (world_size - 1) increments from all other ranks
while iris.atomic_cas(
tile_ready + tile_id,
0, # Never matches when ready, so acts as atomic read
0,
cur_rank,
cur_rank,
heap_bases,
sem="acquire",
scope="sys",
) < (world_size - 1):
pass

# Map tile_id to (pid_m, pid_n)
num_pid_in_group = GROUP_SIZE_M * num_pid_n
Expand Down