Skip to content

Allreduce performance optimization and correctness fix #360

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

Open
wants to merge 14 commits into
base: main
Choose a base branch
from
Open

Allreduce performance optimization and correctness fix #360

Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
a8c30c0
clipping fp16/bf16 addition
chhwang Sep 16, 2024
780f0f8
rccl-tests pass
chhwang Sep 16, 2024
c6cf400
align with msccl clipping
chhwang Sep 25, 2024
d929d25
revert
chhwang Sep 25, 2024
dee8fe2
final barrier for allreduce8
chhwang Sep 25, 2024
abe69b8
this is weird
chhwang Sep 25, 2024
ca6741c
apps/nccl: 16B LLPacket for allreduce7
nusislam Sep 26, 2024
6484dce
apps/nccl: performance optimization for allreduce7
nusislam Sep 26, 2024
01e105b
apps/nccl: add unroll in allred8
nusislam Sep 30, 2024
ea4f77d
Update allgather.hpp
chhwang Oct 2, 2024
99b5997
apps/nccl: Add barrier in allgather
nusislam Oct 2, 2024
f9def85
apps/nccl: fix allreduce7 correctness issue for non power of 2 sizes
nusislam Oct 5, 2024
cdbb2de
apps/nccl: allgather tuning
nusislam Oct 14, 2024
45d40a5
Merge branch 'main' into clip-llpacket
chhwang Oct 22, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 1 addition & 2 deletions apps/nccl/src/allgather.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,12 +103,11 @@ __global__ void __launch_bounds__(1024, 1)
}
}

deviceSyncer.sync(gridDim.x);

if (threadIdx.x < nPeer) {
smChans[threadIdx.x].relaxedSignal();
smChans[threadIdx.x].wait();
}
__syncthreads();
}

template <bool IsOutOfPlace, typename T>
Expand Down
Loading