Repository navigation
Fix amdsmiFabricMaxDevices limit (ROCM-26920) - #7766
Merged
cmcknigh merged 1 commit intoJun 24, 2026
Merged
Conversation
The issue is that a wrong constant is hardcoded in the RCCL source tree. More details can be found in this JIRA ticket (https://amd-hub.atlassian.net/browse/ROCM-26920). modified: projects/rccl/src/include/amdsmi_wrap.h modified: projects/rccl/src/misc/amdsmi_wrap.cc
isaki001
approved these changes
Jun 24, 2026
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adjusts RCCL’s AMD SMI UALoE fabric device-info caching to scale with the library’s maximum supported local GPU count and to error out explicitly when the detected device count exceeds that internal bound.
Changes:
- Increase
amdsmiFabricMaxDevicesfrom a fixed 32 toNCCL_MAX_LOCAL_RANKS(viadevice.h) so all local GPUs can be cached. - Change fabric initialization behavior to return
ncclInternalError(instead of truncating) when the detected device count exceedsamdsmiFabricMaxDevices, and update the warning message accordingly. - Make the assignment to
amdsmiFabricDeviceCountan explicit cast toint.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated no comments.
| File | Description |
|---|---|
| projects/rccl/src/misc/amdsmi_wrap.cc | Switch overflow handling from truncation to a hard failure and adjust logging/type assignment. |
| projects/rccl/src/include/amdsmi_wrap.h | Resize the fabric device cache limit to NCCL_MAX_LOCAL_RANKS and include device.h to obtain the macro. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
1 task
Kapil-Shyam-Pawar
added a commit
that referenced
this pull request
Jul 2, 2026
…peed) (#8077) ## Motivation On gfx950 CPX+NPS2 (64 ranks), all_gather_perf aborts during comm init even after #7766. Issues identified: - A bad-pointer `Bad address` in the topology bootstrap AllGather ``` NCCL WARN ncclOsSocketProgressOpt: Call to send to 10.235.24.26<60509> failed : Bad address ``` - A segfault from WarpSpeed over-allocating channels. ``` NCCL INFO WarpSpeed enabled: warpSpeedChannelMultiplier 4, maxNchannels 1760, nc 1760 ``` ## Technical Details - bootstrap.cc: ring slice offset slice * size overflowed 32-bit int with the ~102 MB per-rank topology buffer at 64 ranks - connect.cc: single-node WarpSpeed channel count was unclamped (nc=1760), overrunning comm->channels[MAXCHANNELS] ## JIRA ID JIRA ID : ROCM-27078 ## Test Plan - Run all_gather_perf on gfx950 CPX+NPS2 (64 GPUs) - Add host unit tests for the 64-bit offset and the channel clamp ## Test Result - AllGather on CPX+NPS2 passes end-to-end with no bad address/segfault - New unit tests (`BootstrapBidir.SliceOffsetIsComputedIn64Bit`, `Rcclwrap.WarpSpeedChannelClampCapsAtHalfMaxChannels`) pass ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
1 task
marbre
added a commit
that referenced
this pull request
Jul 7, 2026
…peed) (#8077) (#8194) Cherry-picks commit 831fad2 into the 7.14 release branch. --- ## Motivation On gfx950 CPX+NPS2 (64 ranks), all_gather_perf aborts during comm init even after #7766. Issues identified: - A bad-pointer `Bad address` in the topology bootstrap AllGather ``` NCCL WARN ncclOsSocketProgressOpt: Call to send to 10.235.24.26<60509> failed : Bad address ``` - A segfault from WarpSpeed over-allocating channels. ``` NCCL INFO WarpSpeed enabled: warpSpeedChannelMultiplier 4, maxNchannels 1760, nc 1760 ``` ## Technical Details - bootstrap.cc: ring slice offset slice * size overflowed 32-bit int with the ~102 MB per-rank topology buffer at 64 ranks - connect.cc: single-node WarpSpeed channel count was unclamped (nc=1760), overrunning comm->channels[MAXCHANNELS] ## JIRA ID JIRA ID : ROCM-27078 ## Test Plan - Run all_gather_perf on gfx950 CPX+NPS2 (64 GPUs) - Add host unit tests for the 64-bit offset and the channel clamp ## Test Result - AllGather on CPX+NPS2 passes end-to-end with no bad address/segfault - New unit tests (`BootstrapBidir.SliceOffsetIsComputedIn64Bit`, `Rcclwrap.WarpSpeedChannelClampCapsAtHalfMaxChannels`) pass ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Kapil S. Pawar <KapilShyam.Pawar@amd.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Raise fabric device cache to NCCL_MAX_LOCAL_RANKS and fail on overflow instead of truncating.
Technical Details
Errors:
Observing ncclOsSocketProgressOpt reported bad address when running pyt_huggingface bert, pyt_huggingface_xbloom, vllm.
Fix:
JIRA ID
ROCM-26920
Test Plan
Follow the detailed reproducing steps documented in the original ticket (JIRA ROCM-26920).
Test Result
Each of the applications listed in the original JIRA ticket completes its run successfully.