* Update unit tests for alt_rsmi impl
- Create distinct test executable for alt_rsmi testing
- Updated alt_rsmi tests to use public methods
- Compiles alt_rsmi.cc with ARSMI_TEST_BUILD
- Enables external linkage of internal variables
- Only for AltRsmiTests.cpp that manipulates internals
- Clean separation for test behavior
* Address review comments
* restore hidden symbol visibility
* remove node-count and threshold restrictions from p2p-batching
* remove batching threshold usage, fix typo for using batching-enablement flag
---------
Co-authored-by: Mustafa Abduljabbar <mustafa.abduljabbar@amd.com>
* Add ncclCommDump API
* remove trailing whitespace changes
* Add more proxy trace timestamps
* Add facebook_rccl namespace before proxyTrace timestamp call
* Clean up ProxyTrae construction
* Move updateProxyOpCounter to member function
* Move setProxyOpTimestamp to member function
* Move addNewProxyOp to member function
* Make internal methods private
* Make ProxyTrace thread safe
* Fix unit tests
* Fix overwritten ProxyTrace DONE setting in net.cc
* Use one side stream per process
* Handle multiple GPUs per process
* Reset stream when not found
* Address review comments
* Fix missing mutex initializer
* for multinode gfx950, extend AR LL128 up to 256MB, extend RS LL128 up to 8MB per rank, extend AG LL up to 64KB per rank
* dont override direct allgather threshold if set to -1
* restore 2-node AR simple at earlier message sizes than higher multi-node AR
* extend range of LL for single-node RS on gfx950
* update algo/proto for multi-node allreduce on gfx942
* set single-node AR on gfx950 to Tree LL for KB message sizes
* decrease threshold for single node Tree for gfx950 AR
* Fail the job if compiler flag HIP_HOST_UNCACHED_MEMORY is not turned on on mi350x
Place the check after initTransportsRank as the GPU arch info in comm->topo->nodes info is populated after that.
* Update src/init.cc to use ERROR instead of WARN
Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com>
* Enhance logging in NCCL initialization
It's convenient to log comms obj and default channels together for debugging
* Add opCount to collDevWork and update increment logic
Added opCount to collDevWork and incremented it when proxyOpQueue is empty (e.g., for intra-node comms)
* Clarify opCount increment logic in enqueue.cc
Updated comment to clarify incrementing opCount for intranode communications.
* Refactor NCCL_INIT logging format
Updated logging format for NCCL_INIT to improve clarity.
* Remove duplicate INFO logging in init.cc
* Added ERROR message class to handle fatal error messages.
New ERROR message class will print the message in all debug level,
including none.
Change some of the fatal error message to be in ERROR instead of WARN.
Added new error handler function to print out more meaningful error
message in the future.
* Added CHANGELOG entry.
* Update CHANGELOG.md
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
* Change to no longer reuse NONE as ERROR. ERROR is now a separated class.
* Update CHANGELOG.md
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
---------
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
* prevent batching when send/recv bytes dont match, restore bit reversal for channel to part mapping, prevent batching beyond 32-nodes
* correct computation for channel to part mapping
* update changelog
* disabling p2p-batching by default
* Add initial commit to increase tb size to 512
* Fix LL perf issue when subset of NCCL_MAX_NTHREADS is used
Adding a constant to barrier_generic logic from using fallback logic when nthreads < NCCL_MAX_NTHREADS and nthreads == blockDim.X
* Adjust nthreads for LL
* Opt threads for reduce_scatter upper small range
* Add macro for single node
* Restrict MSCCL to 256 threads to prevent mem access fault
* Support pre-MI350 compatibility
* Partially refactor threadblock size override
* Use const macros instead of numerals
* opt out of unused function
* Implements casting key loads and stores to address_space(1) so that vector global load and store instructions are emitted by the compiler instead of more costly flat loads and stores
* Removes nontemporality from some key stores for gfx950.
* Adding more path to the kernel load and an environment variable to force enable DMABUF
---------
Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>
- Earlier fix PR1788 is no longer necessary after ROCr fix and pre-ROCr fix workaround
- Inserts an s_waitcnt vmcnt(0), which fixes a data corruption issue in MSCCL
* Gate based on ROCM version, safe for ROCm 7.0.2 and beyond.
* Updates naming to gfx9CheapFenceOff since we use this for gfx942 and gfx950. Thanks Nilesh.
* Add info logging statement to NCCL_INIT to print whether enabled when INFO logging is enabled.
Atul Kulkarni
Mustafa Abduljabbar
isaki001
Ahmed Khan
Wenkai Du
corey-derochie-amd
AbandiGa
Matt Williams
Pedram Alizadeh
Bertan Dogancay
gilbertlee-amd
Dingming Wu
Gheorghe-Teodor Bercea
Kapil S. Pawar
Ghadeer Ahmed H Alabandi
alex-breslow-amd
Arm Patinyasakdikul
nawrinsu
Nilesh M Negi
David DeBonis
mberenjk
Rahul Vaidya