Skip to content

Fix for allgather_fullmesh algo #813

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

Merged
Binyang2014 merged 1 commit into from
May 26, 2026
+3 −3
Merged

Fix for allgather_fullmesh algo #813

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
6 changes: 3 additions & 3 deletions src/ext/collectives/allgather/allgather_fullmesh_2.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,14 +17,14 @@ __global__ void __launch_bounds__(1024, 1)
const size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
const size_t lid = tid % WARP_SIZE;
const size_t wid = tid / WARP_SIZE;
const size_t nPeer = nRanksPerNode - 1;

// Round down to multiple of warp size
const size_t nThread = (blockDim.x * gridDim.x) / WARP_SIZE * WARP_SIZE;
// Round down to multiple of peer count.
const size_t nThread = (blockDim.x * gridDim.x) / WARP_SIZE / nPeer * nPeer * WARP_SIZE;
if (tid >= nThread) {
return;
}
const size_t nWarp = nThread / WARP_SIZE;
const size_t nPeer = nRanksPerNode - 1;
const size_t chanOffset = nPeer * blockIdx.x;
auto memChans = memoryChannels + chanOffset;

Expand Down
Loading