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

Why does NCCL not support one GPU multiple NICs in all2all collective communication? #804

Closed
zhangmenghao opened this issue Mar 17, 2023 · 5 comments
Closed

Why does NCCL not support one GPU multiple NICs in all2all collective communication? #804

zhangmenghao opened this issue Mar 17, 2023 · 5 comments

Comments

@zhangmenghao
Copy link

I was testing NCCL between two node, and each node has one GPU and two NICs in the same NUMA, connecting directing with one switch. Each NIC is assigned with a different IP address.
I found that in all_reduce, NCCL already support one GPU two NICs, and the bandwidth of these two NICs can be fully utilized.
However, in all2all, only one NIC is used, and the other NIC is left used.

Therefore, my question is, why does NCCL not support one GPU two NICs in alltoall? Is there some difficulty that prevent you from supporting such a feature?
@sjeaugey
Copy link
Member

What version are you using? We had that issue on old versions but it was supposed to be fixed now.

@zhangmenghao
Copy link
Author

zhangmenghao commented Mar 17, 2023 •
edited
Loading

We used NCCL 2.14.3. Which version are you supposed to fix that problem? @sjeaugey

@zhangmenghao
Copy link
Author

Hi, Sylvain, could you please answer this question for us? @sjeaugey

@sjeaugey
Copy link
Member

Sorry for the delay. Turns out I was wrong and it wasn't fixed. We're working on a fix right now, we'll post updates here when ready.

@zhangmenghao
Copy link
Author

zhangmenghao commented Mar 22, 2023 •
edited
Loading

Really great news! We are looking forward to your new NCCL version.

In NCCL 2.14.3, we have hacked the function ncclTopoGetLocalNet() in topo.cc, and make alltoall support one GPU two NICs finally. Especially, we assigned netDev based on not only rr, but also on channelId. However, such a change only uses half of these channels, i.e., half of channels are assigned to these two NICs equally. And we are confused by such a phenomenon. Could you please explain a few words about how the setting-up channels are used in the subsequent all2all communication?

sjeaugey added a commit that referenced this issue Apr 19, 2023
@sjeaugey
2.18.1-1
d97a32f 
Add support for IB SHARP to NVLS (NVLink SHARP algorithm).
Add NVLS+Tree algorithm.
Add support for memory management using cuMem* functions.
Use all NICs for Send/Receive operations on systems with more than
one NIC per GPU (#804).
Add ncclCommSplit primitive, with resource sharing option in config.
Fix alltoallv hang (#788)
Increase number of channels on H100 when we're not limited by NVLink.
Improve error reporting in case of IB failure, printing local and
remote ID (#779).
Add build option to allow compilation against RDMA includes instead
of dynamically loading IB verbs symbols (#802).
Fix context creation for progress thread (#803).
NET/IB: add option to use multiple QPs in round-robin mode.
Fix tree performance issue when NVB is disabled on HCM topologies.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
2 participants
@zhangmenghao @sjeaugey