[Question]: Any reason the `smem` requested by NCCL's kernels is based on MAX_THREADS instead of actual threads for the kernel?
Open
Nobody has claimed this yet.
question
- Dominant language
- C++
- Stars
- 5.1k
- Forks
- 1.4k
- Avg merge
- 2h 3m
- Merged PRs (30d)
- 2
Description
Question
Specifically on this line here https://github.com/NVIDIA/nccl/blob/master/src/enqueue.cc#L1679C51-L1680C8
curious why does NCCL use
int smem = ncclShmemDynamicSize(comm->cudaArch);
instead of something like
int smem = (plan->threadPerBlock / WARP_SIZE) * smem_warp;
Contributor guide
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Read the allocation logic around src/enqueue.cc lines 1679-1680, including ncclShmemDynamicSize(comm->cudaArch) and the plan thread count. Trace how the requested shared memory is used by NCCL's kernels and compare it with the proposed per-warp calculation. Done means documenting the reason for the MAX_THREADS-based request, or identifying a confirmed change needed.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- distributed-systems
- Issue type
- Documentation
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100