From 99813a3288dc4e3a3d7b940fbf42755a01181677 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Wed, 11 Jun 2025 09:09:57 -0500 Subject: [PATCH] msccl: adjust msccl threshold for bf16 (#1736) * msccl: adjust msccl threshold for bf16 * Update src/collectives.cc Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com> --------- Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com> [ROCm/rccl commit: 75c3c8215c2cfe3733194bb26c7a5170cb76c5f8] --- projects/rccl/src/collectives.cc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/rccl/src/collectives.cc b/projects/rccl/src/collectives.cc index d51394ad49..37860cc25d 100644 --- a/projects/rccl/src/collectives.cc +++ b/projects/rccl/src/collectives.cc @@ -125,9 +125,11 @@ ncclResult_t ncclAllReduce_impl(const void* sendbuff, void* recvbuff, size_t cou } if (mscclAvailable(comm) && !mscclIsCaller()) { - return mscclEnqueueCheck( - sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr, - count, datatype, 0, 0, op, mscclFuncAllReduce, comm, stream); + if (datatype != ncclBfloat16 || (count * ncclTypeSize(datatype) <= 8388608)) { + return mscclEnqueueCheck( + sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr, + count, datatype, 0, 0, op, mscclFuncAllReduce, comm, stream); + } } return ncclEnqueueCheck(&info);