Skip to content

Gemm-allreduce loop optimization#170

Open
erieaton-amd wants to merge 3 commits into
ByteDance-Seed:mainfrom
erieaton-amd:arg-4
Open

Gemm-allreduce loop optimization#170
erieaton-amd wants to merge 3 commits into
ByteDance-Seed:mainfrom
erieaton-amd:arg-4

Conversation

@erieaton-amd
Copy link
Copy Markdown
Collaborator

@erieaton-amd erieaton-amd commented May 13, 2026

Removes some unnecessary loop iterations. This actually cuts the execution time roughly in half (7ms down to 3ms) on MI350. I also added the allreduce test to the CI.

remote_c_ptr = dl.consume_token(remote_c_ptr, token)
remote_c_ptrs = remote_c_ptr + offs_cm[:, None] * stride_cm + offs_cn[None, :] * stride_cn
final_acc = tl.load(remote_c_ptrs, mask=c_mask, other=0.0).to(tl.float32)
for i in range(1, world_size):
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this is still wait-all, reduce-all, if we decouple wait and reduce in separate loops, that can help as well.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants