can alltoall use zero copy (only send without recv) by ncclCommRegister?

[ieeechuhaoyun]

// nWarpPerWork = nWarps/nWorks
int nWarpPerWork = __popc(__ballot_sync(~0u, nWorks*(lane+1) <= nWarps));
int nRecvWarpPerWork = nWarpPerWork<=4 ? nWarpPerWork/2 : (nWarpPerWork-1)/2;
int nSendWarpPerWork = nWarpPerWork<=4 ? nRecvWarpPerWork : nRecvWarpPerWork+1;

i found these codes in nccl/src/device/sendrecv.h, which means recv alway use half warps , maybe nccl-tests -R param can not use in ALL2ALL?

[AddyLaddy]

NCCL Send/Recv cannot make use of the Symmetric memory registration.
However, NCCL 2.28 now includes a new Collective API: ncclAlltoAll
And there is also a CE based implementation of the AlltoAll collective call. You need to use both Symmetric registration at CTA Policy NCCL_CTA_POLICY_ZERO to utilize it on certain platforms.

[alokprasad]

@AddyLaddy

1. "You need to use both Symmetric registration at CTA Policy NCCL_CTA_POLICY_ZERO to utilize it on certain platforms."
   What do you mean by platform ?
   if i use NCCL_CTA_POLICY_ZERO then with NCCL 2.28/Cuda 13 will nccl-tests alltoall will use CE ?
   how can i know amount of SM and CE used ?

[zhenhaohe]

We have the requirements to use CE specified here:
https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/bufferreg.html#zero-cta-optimization
AlltoAll API uses CE in NCCL perf tests if the above requirements are met.

[alokprasad]

@zhenhaohe thanks and how to confirm from nccl logs,

[zhenhaohe]

You could add

• NCCL_DEBUG=INFO
• NCCL_DEBUG_SUBSYS=INIT

And you will see information on CE initialization on the first run of your workload:
https://github.com/NVIDIA/nccl/blob/master/src/ce_coll.cc#L48