Skip to content
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

Jump to bottom

ncclAllReduce may require additional barriers #1587

Open
suranap opened this issue Jan 27, 2025 · 0 comments
Open

ncclAllReduce may require additional barriers #1587

suranap opened this issue Jan 27, 2025 · 0 comments
Assignees
@jiazhihao
Labels
bug Something isn't working

Comments

@suranap
Copy link
Collaborator

suranap commented Jan 27, 2025

flexflow-train/lib/kernels/src/cuda/optimizer_kernel.cu

Line 83 in de7fa32

checkNCCL(ncclAllReduce(w_grad_ptr,

A thread on Zulip mentioned some additional care required for NCCL within a Legion task. Rohan spotted a problem in FlexFlow's use of ncclAllReduce. You may need to add concurrent_task_barrier before and after the call, and call set_concurrent_barrier on the task. More info is in the comment for that barrier.
@suranap suranap added the bug Something isn't working label Jan 27, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@jiazhihao jiazhihao

Labels
bug Something isn't working
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
2 participants
@suranap @jiazhihao