Merge ~bullwinkle-team/ubuntu/+source/rccl:bullwinkle/llvm-21/ubuntu/devel into ubuntu/+source/rccl:ubuntu/devel

[Bojan Aleksovski]

Update to new upstream version 7.1.0

[Bojan Aleksovski (b0b0a)]

Uploaded package to this PPA: https://launchpad.net/~b0b0a/+archive/ubuntu/rccl-2140330

(-proposed and archs enabled)

Will trigger autopkgtest when it is built and published

[Andreas Hasenack (ahasenack)]

Please show the reverse dependency analysis on what else will have to be dropped on ppc64el:
--- a/debian/control
+++ b/debian/control
@@ -31,7 +38,7 @@ Rules-Requires-Root: no

 Package: librccl1
 Section: libs
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: ${misc:Depends}, ${shlibs:Depends},
 Description: ROCm Communication Collectives Library - library
  RCCL (pronounced "Rickle") is a library of collective communication routines
@@ -45,7 +52,7 @@ Description: ROCm Communication Collectives Library - library

 Package: librccl-dev
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: librccl1 (= ${binary:Version}),${misc:Depends}, ${shlibs:Depends},
          libamdhip64-dev,
 Suggests: librccl-doc
@@ -61,7 +68,7 @@ Description: ROCm Communication Collectives Library - headers

 Package: librccl1-tests
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: librccl1 (= ${binary:Version}), ${misc:Depends}, ${shlibs:Depends},
 Build-Profiles: <!nocheck>
 Description: ROCm Communication Collectives Library - tests

review: Needs Fixing

[Bojan Aleksovski (b0b0a)]

`reverse-depends --arch ppc64el {librccl1,librccl-dev,librccl1-tests}`:
- librccl1: reverse-depends librccl-dev and librccl1-tests
- librccl-dev and librccl1-tests: No reverse dependencies found
thus nothing additionally needs to be dropped on ppc64el

On a further note, this MP is blocked by rocm-hipamd.

[Bojan Aleksovski (b0b0a)]

We are evaluating a change that might bump rccl to 7.1.0-0ubuntu2 regarding rocm-hipamd dependency.
Please postpone further review until further notice.
Thank you!

[Talha Can Havadar (tchavadar)]

As long as there is no direct dependency to rocm-hipamd please refrain to add it to build-depends. As of today there is a transitive dependency to rocm-hipamd through hipcc. But if we add this to rccl wven though it doesnt depend on it directly then in the future hipcc may drop hipamd dependency and then this would make us update rccl again (ideally)

Please investigate the source code of rccl and check if we have dependency to hipamd before trying to add it to build-depends

[Bojan Aleksovski (b0b0a)]

Andreas please continue with the review.
libamdhip64-dev (=> 7.1~) has been added as a build-dep (reasoning: https://code.launchpad.net/~bullwinkle-team/ubuntu/+source/rccl/+git/rccl/+merge/499891).
Thank you!

[Andreas Hasenack (ahasenack)]

--- a/debian/librccl1.install
+++ b/debian/librccl1.install
@@ -1 +1,2 @@
-usr/lib/*/librccl.so.*
+usr/lib/librccl.so.*
+usr/bin/rcclras

Why are the libraries now installed directly in /usr/lib, instead of /usr/lib/<arch>? Very few libraries are installed directly in /usr/lib:

$ l /usr/lib/*.so.*|wc -l
8

$ l /usr/lib/x86_64-linux-gnu/*.so.*|wc -l
2666

review: Needs Information

[Andreas Hasenack (ahasenack)]

I was checking if we needed to change the soname, but looks like it's tied to NCCL_MAJOR, defined in makefiles/version.mk, and that has remained at 2. Which just begs the question, why wasn't our soname already 2?

./makefiles/version.mk:NCCL_MAJOR := 2
./src/Makefile:LIBSONAME := $(LIBNAME:%=%.$(NCCL_MAJOR))

Perhaps check that in parallel, for another future PR if needed.

[Andreas Hasenack (ahasenack)]

So, symbols. I see the symbols in the previous package were not demangled, and the new symbols for this version are also not demangled. This makes it hard to see if symbols were dropped, in which case we would have to either bump the soname.

Perhaps just use d/shlibs like rocalution was doing (
https://code.launchpad.net/~bullwinkle-team/ubuntu/+source/rocalution/+git/rocalution/+merge/499090).

Or else we need proper demangling of symbols in the old (current) package, and then the new one here, and let dpkg-gensymbols tell us which symbols were dropped/added.

Right now I understand the only rdeps are within the package, but this may change.

79a6749... by Bojan Aleksovski on 2026-02-09

d/patches: add new fix-install-libdir.patch

[Bojan Aleksovski (b0b0a)]

Hello Andreas,

Please find the effort to have the comments fixed:
- Regarding the installation of /usr/lib instead of /usr/lib/<arch>, I've added fix-install-libdir.patch that has been forwarded upstream (it seems they did this in upstream unintentionally as part of a non-related commit)
- Regarding the soname, in CMakeLists.txt L1316 it is set to 1.0:
`## Setup librccl.so version
rocm_set_soversion(rccl "1.0")`
the one that showed up for you is for unrelated nccl.
- Regarding the symbols, I've added d/shlibs instead as per your recommendation until we get them demangled in the future.

The easiest place to view the mentioned fixes is https://launchpadlibrarian.net/846853345/rccl_7.1.0-0ubuntu2~ppa1_7.1.0-0ubuntu2~ppa2.diff.gz

Thank you!

[Andreas Hasenack (ahasenack)]

+1

review: Approve

[Andreas Hasenack (ahasenack)]

Uploading rccl_7.1.0-0ubuntu2.dsc
Uploading rccl_7.1.0.orig.tar.gz
Uploading rccl_7.1.0-0ubuntu2.debian.tar.xz
Uploading rccl_7.1.0-0ubuntu2_source.buildinfo
Uploading rccl_7.1.0-0ubuntu2_source.changes

Unmerged commits

148928e... by Bojan Aleksovski on 2026-02-05

d/changelog: add 7.1.0-0ubuntu2 entry

a7f7989... by Bojan Aleksovski on 2026-02-05

d/control: add libamdhip64-dev (=> 7.1~) as a build dep

It is a direct hard dependency of rccl

e8b5890... by Bojan Aleksovski on 2026-01-30

d/changelog: add resolute entry for 7.1.0

2d112ca... by Bojan Aleksovski on 2026-02-03

d/rules: [Temporary] remove gfx908 and gfx90a ROCm ISAs

Remove gfx908 and gfx90a ISAs due to LP amd64 and amd64v3 builders
resetting themselves and failing to build without any buildlog.
*This shall be reverted when LP infra problem is solved*

851a82e... by Bojan Aleksovski on 2026-01-30

d/control: remove ppc64el build

79a6749... by Bojan Aleksovski on 2026-02-09

d/patches: add new fix-install-libdir.patch

5a5a124... by Bojan Aleksovski on 2026-01-30

d/{librccl1.symbols,shlibs}: remove symbols file and instead add shlibs

Otherwise proper demangling of symbols will be needed

b9e72fe... by Bojan Aleksovski on 2026-01-30

d/librccl-doc.install: add doc install for msccl-algorithms

309a4a8... by Bojan Aleksovski on 2026-01-30

d/librccl1-tests.install: fix path and add msccl-unit-test-algorithms

Tests reference the algorithms

23c6bee... by Bojan Aleksovski on 2026-01-30

d/librccl1.install: add rcclras install

diff --git a/.azuredevops/multinode-ci-nightly.yml b/.azuredevops/multinode-ci-nightly.yml
new file mode 100644
index 0000000..4b9b143
--- /dev/null
+++ b/.azuredevops/multinode-ci-nightly.yml
@@ -0,0 +1,71 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+- name: pytestFolder
+  value: '.azuredevops/tests/pytest'
+
+parameters:
+- name: pytestList
+  type: object
+  default:
+    - HelloWorld
+
+trigger: none
+pr: none
+schedules:
+  - cron: "0 5 * 11-3 *"  # 11 PM CST (November - March)
+    displayName: "Nightly Build (CST)"
+    branches:
+      include:
+        - develop
+    always: false
+
+  - cron: "0 4 * 4-10 *"  # 11 PM CDT (April - October)
+    displayName: "Nightly Build (CDT)"
+    branches:
+      include:
+        - develop
+    always: false
+
+jobs:
+- job: rccl
+  timeoutInMinutes: 180
+  pool: rocm-ci_rccl_pool
+  workspace:
+    clean: all
+  steps:
+  - task: DeleteFiles@1
+    inputs:
+      Contents: '**/*'
+  - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/checkout.yml@pipelines_repo
+    parameters:
+      submoduleBehaviour: recursive
+  - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/build-cmake.yml@pipelines_repo
+    parameters:
+      installEnabled: false
+      printDiskSpace: false
+      extraBuildFlags: >-
+        -DCMAKE_BUILD_TYPE=Release
+        -DBUILD_TESTS=ON
+        -GNinja
+  - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/test.yml@pipelines_repo
+    parameters:
+      componentName: rccl
+      testDir: $(Build.SourcesDirectory)/build/test
+      testExecutable: 'LD_LIBRARY_PATH=$(Build.SourcesDirectory)/build:${LD_LIBRARY_PATH} NCCL_DEBUG=INFO RCCL_ENABLE_SIGNALHANDLER=1 ./rccl-UnitTests'
+      testParameters: '--gtest_output=xml:./test_output.xml --gtest_color=yes'
+  - ${{ each pytestScript in parameters.pytestList }}:
+    - task: Bash@3
+      displayName: Test ${{ pytestScript }}
+      continueOnError: true
+      inputs:
+        targetType: inline
+        workingDirectory: $(Build.SourcesDirectory)/$(pytestFolder)
+        script: pytest ${{ pytestScript }}.py
diff --git a/.azuredevops/multinode-ci-pr.yml b/.azuredevops/multinode-ci-pr.yml
new file mode 100644
index 0000000..040377e
--- /dev/null
+++ b/.azuredevops/multinode-ci-pr.yml
@@ -0,0 +1,77 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+- name: pytestFolder
+  value: '.azuredevops/tests/pytest'
+
+parameters:
+- name: pytestList
+  type: object
+  default:
+    - HelloWorld
+
+trigger: none
+pr:
+  autoCancel: true
+  branches:
+    include:
+    - develop
+  paths:
+    exclude:
+    - .github
+    - .jenkins
+    - docs
+    - '*.md'
+    - LICENSE.txt
+    - NOTICES.txt
+  drafts: false
+
+stages:
+- stage: rcclStage
+  displayName: 'RCCL develop PR'
+  jobs:
+  - deployment: rccl_pr_approval
+    displayName: "CI Run Requires Approval"
+    environment: rccl
+  - job: rccl
+    timeoutInMinutes: 180
+    pool: rocm-ci_rccl_pool
+    workspace:
+      clean: all
+    steps:
+    - task: DeleteFiles@1
+      inputs:
+        Contents: '**/*'
+    - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/checkout.yml@pipelines_repo
+      parameters:
+        submoduleBehaviour: recursive
+    - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/build-cmake.yml@pipelines_repo
+      parameters:
+        installEnabled: false
+        printDiskSpace: false
+        extraBuildFlags: >-
+          -DCMAKE_BUILD_TYPE=Release
+          -DBUILD_TESTS=ON
+          -DGPU_TARGETS=gfx942
+          -GNinja
+    - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/test.yml@pipelines_repo
+      parameters:
+        componentName: rccl
+        testDir: $(Build.SourcesDirectory)/build/test
+        testExecutable: 'LD_LIBRARY_PATH=$(Build.SourcesDirectory)/build:${LD_LIBRARY_PATH} NCCL_DEBUG=INFO RCCL_ENABLE_SIGNALHANDLER=1 ./rccl-UnitTests'
+        testParameters: '--gtest_output=xml:./test_output.xml --gtest_color=yes'
+    - ${{ each pytestScript in parameters.pytestList }}:
+      - task: Bash@3
+        displayName: Test ${{ pytestScript }}
+        continueOnError: true
+        inputs:
+          targetType: inline
+          workingDirectory: $(Build.SourcesDirectory)/$(pytestFolder)
+          script: pytest ${{ pytestScript }}.py
diff --git a/.azuredevops/multinode-ci-slurm-nightly.yml b/.azuredevops/multinode-ci-slurm-nightly.yml
new file mode 100644
index 0000000..b7438f3
--- /dev/null
+++ b/.azuredevops/multinode-ci-slurm-nightly.yml
@@ -0,0 +1,44 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+
+trigger: none
+pr: none
+schedules:
+  - cron: "0 5 * 11-3 *"  # 11 PM CST (November - March)
+    displayName: "Nightly Build (CST)"
+    branches:
+      include:
+        - develop
+    always: false
+
+  - cron: "0 4 * 4-10 *"  # 11 PM CDT (April - October)
+    displayName: "Nightly Build (CDT)"
+    branches:
+      include:
+        - develop
+    always: false
+
+jobs:
+- job: rccl
+  timeoutInMinutes: 180
+  pool: rocm-ci_rccl_slurm_pool
+  workspace:
+    clean: all
+  steps:
+  - task: DeleteFiles@1
+    inputs:
+      Contents: '**/*'
+  - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/checkout.yml@pipelines_repo
+    parameters:
+      submoduleBehaviour: recursive
+  - template: templates/build.yml
+  - template: templates/test_rccl-UnitTests.yml
+  - template: templates/test_rccl-tests.yml
diff --git a/.azuredevops/multinode-ci-slurm-pr.yml b/.azuredevops/multinode-ci-slurm-pr.yml
new file mode 100644
index 0000000..c3d89ff
--- /dev/null
+++ b/.azuredevops/multinode-ci-slurm-pr.yml
@@ -0,0 +1,49 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+
+trigger: none
+pr:
+  autoCancel: true
+  branches:
+    include:
+    - develop
+  paths:
+    exclude:
+    - .github
+    - .jenkins
+    - docs
+    - '*.md'
+    - LICENSE.txt
+    - NOTICES.txt
+  drafts: false
+
+stages:
+- stage: rcclStage
+  displayName: 'RCCL develop PR'
+  jobs:
+  - deployment: rccl_pr_approval
+    displayName: "CI Run Requires Approval"
+    environment: rccl
+  - job: rccl
+    timeoutInMinutes: 180
+    pool: rocm-ci_rccl_slurm_pool
+    workspace:
+      clean: all
+    steps:
+    - task: DeleteFiles@1
+      inputs:
+        Contents: '**/*'
+    - template: ${{ variables.CI_TEMPLATE_PATH }}/steps/checkout.yml@pipelines_repo
+      parameters:
+        submoduleBehaviour: recursive
+    - template: templates/build.yml
+    - template: templates/test_rccl-UnitTests.yml
+    - template: templates/test_rccl-tests.yml
diff --git a/.azuredevops/rocm-ci.yml b/.azuredevops/rocm-ci.yml
new file mode 100644
index 0000000..e87ae14
--- /dev/null
+++ b/.azuredevops/rocm-ci.yml
@@ -0,0 +1,46 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+
+trigger:
+  batch: true
+  branches:
+    include:
+    - develop
+    - mainline
+  paths:
+    exclude:
+    - .github
+    - .jenkins
+    - docs
+    - '.*.y*ml'
+    - '*.md'
+    - LICENSE.txt
+    - NOTICES.txt
+
+pr:
+  autoCancel: true
+  branches:
+    include:
+    - develop
+    - mainline
+  paths:
+    exclude:
+    - .github
+    - .jenkins
+    - docs
+    - '.*.y*ml'
+    - '*.md'
+    - LICENSE.txt
+    - NOTICES.txt
+  drafts: false
+
+jobs:
+  - template: ${{ variables.CI_COMPONENT_PATH }}/rccl.yml@pipelines_repo
diff --git a/.azuredevops/slurm/build.sh b/.azuredevops/slurm/build.sh
new file mode 100644
index 0000000..9ebcefd
--- /dev/null
+++ b/.azuredevops/slurm/build.sh
@@ -0,0 +1,51 @@
+#!/bin/bash
+#SBATCH --job-name=rccl-build
+#SBATCH --output=rccl-build-%j.out
+#SBATCH --error=rccl-build-%j.out
+#SBATCH --time=60
+#SBATCH --nodes=1
+#SBATCH --exclusive
+#SBATCH --partition=gt
+
+short_id=$(hostname | cut -d'.' -f1 | cut -d'-' -f3-)
+echo "Node identifier: $short_id"
+
+source /etc/profile.d/lmod.sh
+module load rocm/6.4.1
+
+# Setup local binary path
+export PATH="$HOME/.local/bin:$PATH"
+mkdir -p "$HOME/.local/bin"
+
+# Install Ninja if not already available
+if ! command -v ninja &>/dev/null; then
+  echo "Ninja not found. Installing locally..."
+  wget -q https://github.com/ninja-build/ninja/releases/download/v1.11.1/ninja-linux.zip -O /tmp/ninja.zip
+  unzip -q /tmp/ninja.zip -d "$HOME/.local/bin"
+  chmod +x "$HOME/.local/bin/ninja"
+fi
+
+echo "Using Ninja at: $(which ninja)"
+ninja --version
+
+# Define GPU target
+export GPU_TARGETS="gfx942"
+
+cd "${SLURM_SUBMIT_DIR:-$PWD}"
+## Building RCCL
+mkdir -p build
+cd build
+cmake -G Ninja -DCMAKE_INSTALL_PREFIX="$BINARIES_DIR" -DCMAKE_BUILD_TYPE=Release -DGPU_TARGETS=${GPU_TARGETS} -DBUILD_TESTS=ON -DROCM_PATH="$ROCM_PATH" ..
+cmake --build .
+cmake --build . --target install
+
+
+cd "${SLURM_SUBMIT_DIR:-$PWD}"
+## Building RCCL-Tests
+git clone https://github.com/ROCm/rccl-tests
+cd rccl-tests
+mkdir -p build
+cd build
+cmake -DCMAKE_PREFIX_PATH="$BINARIES_DIR;$MPI_HOME" -DUSE_MPI=ON -DCMAKE_INSTALL_PREFIX="$BINARIES_DIR" -DCMAKE_BUILD_TYPE=Release -DGPU_TARGETS=${GPU_TARGETS} -DROCM_PATH="$ROCM_PATH" ..
+cmake --build .
+cmake --build . --target install
diff --git a/.azuredevops/slurm/test_rccl-UnitTests.sh b/.azuredevops/slurm/test_rccl-UnitTests.sh
new file mode 100644
index 0000000..f397fab
--- /dev/null
+++ b/.azuredevops/slurm/test_rccl-UnitTests.sh
@@ -0,0 +1,16 @@
+#!/bin/bash
+#SBATCH --job-name=rccl-UnitTests
+#SBATCH --output=%x-%j.out
+#SBATCH --error=%x-%j.out
+#SBATCH --time=120
+#SBATCH --nodes=1
+#SBATCH --exclusive
+#SBATCH --partition=gt
+
+short_id=$(hostname | cut -d'.' -f1 | cut -d'-' -f3-)
+echo "Node identifier: $short_id"
+
+source /etc/profile.d/lmod.sh
+module load rocm/6.4.1
+cd "$BINARIES_DIR/bin"
+LD_LIBRARY_PATH="$BINARIES_DIR/lib:$LD_LIBRARY_PATH" NCCL_DEBUG=INFO RCCL_ENABLE_SIGNALHANDLER=1 HSA_NO_SCRATCH_RECLAIM=1 ./rccl-UnitTests --gtest_output=xml:$PIPELINE_WORKSPACE/rccl-UnitTests_output.xml --gtest_color=yes
diff --git a/.azuredevops/slurm/test_rccl-tests.sh b/.azuredevops/slurm/test_rccl-tests.sh
new file mode 100644
index 0000000..a4509ce
--- /dev/null
+++ b/.azuredevops/slurm/test_rccl-tests.sh
@@ -0,0 +1,62 @@
+#!/bin/bash
+#SBATCH --job-name=rccl-tests
+#SBATCH --output=%x-%j.out
+#SBATCH --error=%x-%j.out
+#SBATCH --time=60
+#SBATCH --nodes=1
+#SBATCH --exclusive
+#SBATCH --partition=gt
+
+short_id=$(hostname | cut -d'.' -f1 | cut -d'-' -f3-)
+echo "Node identifier: $short_id"
+
+source /etc/profile.d/lmod.sh
+module load rocm/6.4.1
+
+cd ${PIPELINE_WORKSPACE}/TestResults
+mkdir -p ${PIPELINE_WORKSPACE}/TestResults/rccl-tests_logs
+export WORKDIR=${PIPELINE_WORKSPACE}/TestResults/rccl-tests_logs
+
+export PATH="$BINARIES_DIR/bin:$PATH"
+export LD_LIBRARY_PATH="$BINARIES_DIR/lib:$LD_LIBRARY_PATH"
+
+### create hostlist
+#nodelist=($(scontrol show hostnames))
+#echo "SLURM nodes:"
+#echo ${nodelist[@]}
+#echo ""
+#
+#hosts_8ppn=()
+#for node in "${nodelist[@]}"
+#do
+#    hosts_8ppn+=("${node}:8")
+#done
+#echo ${hosts_8ppn[@]}
+
+### Run multi- and single-node RCCL-Tests
+## Run single-node RCCL-Tests
+for n in 1
+do
+    total=$((n*8))
+    #h_8ppn=`echo ${hosts_8ppn[@]:0:${n}} | tr ' ' ','`
+
+    for coll in all_reduce all_gather reduce_scatter alltoall alltoallv broadcast gather reduce scatter sendrecv
+    do
+        for dtype in float bfloat16 half fp8_e5m2
+        do
+            out_filename="${WORKDIR}/rccl-tests_${coll}_1KB-16GB_nodes${n}_gpus${total}_${dtype}.log"
+            #cmd="${MPI_HOME}/bin/mpirun -np ${total} --host ${h_8ppn} -mca pml ob1 -mca btl ^openib -mca oob_tcp_if_exclude docker,lo -mca btl_tcp_if_exclude docker,lo -x PATH -x LD_LIBRARY_PATH -x NCCL_DEBUG=VERSION -x NCCL_IB_HCA=bnxt_re0,bnxt_re1,bnxt_re2,bnxt_re3,bnxt_re4,bnxt_re5,bnxt_re6,bnxt_re7 -x NCCL_IGNORE_CPU_AFFINITY=1 -x HSA_NO_SCRATCH_RECLAIM=1 -x NCCL_IB_GID_INDEX=3 ${BINARIES_DIR}/bin/${coll}_perf -b 1K -e 16G -f 2 -g 1 -n 100 -w 50 -d ${dtype} -Z json -x ${WORKDIR}/rccl-tests_${coll}_nodes${n}_gpus${total}_${dtype}.json"
+            cmd="${MPI_HOME}/bin/mpirun -np ${total} -mca pml ^ucx -mca osc ^ucx -mca btl ^openib -mca oob_tcp_if_exclude docker,lo -mca btl_tcp_if_exclude docker,lo -x PATH -x LD_LIBRARY_PATH -x NCCL_DEBUG=VERSION -x NCCL_IGNORE_CPU_AFFINITY=1 -x HSA_NO_SCRATCH_RECLAIM=1 ${BINARIES_DIR}/bin/${coll}_perf -b 1K -e 16G -f 2 -g 1 -n 100 -w 50 -d ${dtype} -Z json -x ${WORKDIR}/rccl-tests_${coll}_nodes${n}_gpus${total}_${dtype}.json"
+
+            echo "Running ${coll}" 2>&1 | tee ${out_filename}
+            echo "Run cmd: ${cmd}" 2>&1 | tee -a ${out_filename}
+            eval ${cmd} 2>&1 | tee -a ${out_filename}
+
+            sleep 2
+        done
+    done
+done
+
+## To add
+### Summarize results
+### Convert to junit
diff --git a/.azuredevops/templates/build.yml b/.azuredevops/templates/build.yml
new file mode 100644
index 0000000..fc671a7
--- /dev/null
+++ b/.azuredevops/templates/build.yml
@@ -0,0 +1,86 @@
+# small subset of files to check for install to determine pass/fail
+parameters:
+- name: expectedInstallFiles
+  type: object
+  default:
+    - bin/rccl-UnitTests
+    - include/rccl/rccl.h
+    - lib/cmake/rccl/rccl-config.cmake
+    - lib/librccl.so
+    - share/doc/rccl/LICENSE.txt
+    - share/rccl/msccl-algorithms
+    - share/rccl/msccl-unit-test-algorithms
+
+steps:
+  - task: Bash@3
+    displayName: Build Job
+    env:
+      BINARIES_DIR: $(Build.BinariesDirectory)
+    inputs:
+      targetType: inline
+      script: |
+        echo "##[section]Starting build job..."
+
+        rm -rf $(Build.BinariesDirectory)/*
+
+        echo "Submitting build job..."
+        mkdir -p $(Build.BinariesDirectory)
+        BUILD_JOB_ID=$(sbatch --export=ALL --parsable $(Build.SourcesDirectory)/.azuredevops/slurm/build.sh)
+        echo "Submitted build job: $BUILD_JOB_ID"
+        echo "##vso[task.setvariable variable=BUILD_JOB_ID]$BUILD_JOB_ID"
+
+        echo "Waiting for build job to start..."
+        while squeue -j $BUILD_JOB_ID 2>/dev/null | grep -q $BUILD_JOB_ID; do
+          echo "##[section]Build job $BUILD_JOB_ID is still running..."
+          sleep 60
+        done
+
+        echo "Waiting for final status via sacct..."
+        LOOP_COUNT=0
+        MAX_LOOPS=30  # Maximum of 30 loops (30 minutes)
+        while true; do
+          STATE=$(sacct -j "$BUILD_JOB_ID" --format=JobID,State --noheader | awk '$1 ~ /\.batch$/ { print $2; exit }' | xargs)
+          echo "##[section]Build job state: $STATE"
+          if [[ "$STATE" == "COMPLETED" ]]; then
+            break
+          elif [[ "$STATE" =~ ^(FAILED|CANCELLED|TIMEOUT)$ ]]; then
+            echo "Build failed with state $STATE"
+            break
+          fi
+          sleep 60
+          LOOP_COUNT=$((LOOP_COUNT + 1))
+          if [ $LOOP_COUNT -ge $MAX_LOOPS ]; then
+            echo "Time limit reached while waiting for final status."
+            exit 1  # Exit with an error code if time limit is reached
+          fi
+        done
+
+        echo "Checking for expected installed files..."
+        MISSING_FILES=0
+
+        expectedFiles="${{ join(' ', parameters.expectedInstallFiles) }}"
+        i=1
+        total=$(echo "$expectedFiles" | wc -w)
+        while [ $i -le $total ]; do
+          relpath=$(echo "$expectedFiles" | cut -d ' ' -f"$i")
+          fullpath="$BINARIES_DIR/$relpath"
+          if [ ! -e "$fullpath" ]; then
+            echo "##vso[task.logissue type=error]Missing expected file: $fullpath"
+            MISSING_FILES=1
+          fi
+          i=$((i + 1))
+        done
+
+        if [ "$MISSING_FILES" -eq 1 ]; then
+          echo "One or more expected files are missing from the install directory."
+          exit 1
+        else
+          echo "All expected files are present in the install directory."
+        fi
+  - task: Bash@3
+    displayName: Build Logs
+    condition: always()
+    inputs:
+      targetType: inline
+      script: |
+        cat rccl-build-${BUILD_JOB_ID}.out || echo "No log found"
diff --git a/.azuredevops/templates/test_rccl-UnitTests.yml b/.azuredevops/templates/test_rccl-UnitTests.yml
new file mode 100644
index 0000000..8d195b1
--- /dev/null
+++ b/.azuredevops/templates/test_rccl-UnitTests.yml
@@ -0,0 +1,69 @@
+steps:
+  - task: Bash@3
+    displayName: RCCL UnitTests
+    env:
+      BINARIES_DIR: $(Build.BinariesDirectory)
+      PIPELINE_WORKSPACE: $(Pipeline.Workspace)
+    inputs:
+      targetType: inline
+      script: |
+        echo "Submitting test job..."
+        TEST_JOB_ID=$(sbatch --export=ALL --parsable $(Build.SourcesDirectory)/.azuredevops/slurm/test_rccl-UnitTests.sh)
+        echo "Submitted test job: $TEST_JOB_ID"
+        echo "##vso[task.setvariable variable=TEST_JOB_ID]$TEST_JOB_ID"
+
+        echo "Waiting for test job to start..."
+        while squeue -j $TEST_JOB_ID 2>/dev/null | grep -q $TEST_JOB_ID; do
+          echo "##[section]Test job $TEST_JOB_ID is still running..."
+          sleep 60
+        done
+
+        echo "Waiting for final status via sacct..."
+        LOOP_COUNT=0
+        MAX_LOOPS=120  # Maximum of 120 loops (120 minutes)
+        while true; do
+          STATE=$(sacct -j "$TEST_JOB_ID" --format=JobID,State --noheader | awk '$1 ~ /\.batch$/ { print $2; exit }' | xargs)
+          echo "##[section]Test job state: $STATE"
+          if [[ "$STATE" == "COMPLETED" ]]; then
+            break
+          elif [[ "$STATE" =~ ^(FAILED|CANCELLED|TIMEOUT)$ ]]; then
+            echo "Test failed with state $STATE"
+            break
+          fi
+          sleep 60
+          LOOP_COUNT=$((LOOP_COUNT + 1))
+          if [ $LOOP_COUNT -ge $MAX_LOOPS ]; then
+            echo "Time limit reached while waiting for final status."
+            exit 1  # Exit with an error code if time limit is reached
+          fi
+        done
+
+        echo "Checking test result XML for failures..."
+        TEST_XML=$(find "$(Pipeline.Workspace)" -name 'rccl-UnitTests_output.xml' | head -n1)
+        if [ -z "$TEST_XML" ]; then
+          echo "##vso[task.logissue type=error]No $TEST_XML file found"
+          echo "##vso[task.complete result=Failed;]DONE"
+          exit 1
+        fi
+
+        if grep -q 'failures="[^0]' "$TEST_XML"; then
+          echo "##vso[task.logissue type=error]Test failures detected in $TEST_XML"
+          echo "##vso[task.complete result=Failed;]DONE"
+          exit 1
+        else
+          echo "No test failures detected."
+        fi
+  - task: Bash@3
+    displayName: Test Logs
+    condition: always()
+    inputs:
+      targetType: inline
+      script: |
+        cat rccl-UnitTests-${TEST_JOB_ID}.out || echo "No log found"
+  - task: PublishTestResults@2
+    displayName: 'Publish Results'
+    condition: succeededOrFailed()
+    inputs:
+      searchFolder: $(Pipeline.Workspace)
+      testResultsFormat: JUnit
+      testResultsFiles: '**/rccl-UnitTests_output.xml'
diff --git a/.azuredevops/templates/test_rccl-tests.yml b/.azuredevops/templates/test_rccl-tests.yml
new file mode 100644
index 0000000..9b047c4
--- /dev/null
+++ b/.azuredevops/templates/test_rccl-tests.yml
@@ -0,0 +1,77 @@
+steps:
+  - task: Bash@3
+    displayName: RCCL-Tests
+    env:
+      BINARIES_DIR: $(Build.BinariesDirectory)
+      PIPELINE_WORKSPACE: $(Pipeline.Workspace)
+    inputs:
+      targetType: inline
+      script: |
+        echo "Submitting test job..."
+        TEST_JOB_ID=$(sbatch --export=ALL --parsable $(Build.SourcesDirectory)/.azuredevops/slurm/test_rccl-tests.sh)
+        echo "Submitted test job: $TEST_JOB_ID"
+        echo "##vso[task.setvariable variable=TEST_JOB_ID]$TEST_JOB_ID"
+
+        echo "Waiting for test job to start..."
+        while squeue -j $TEST_JOB_ID 2>/dev/null | grep -q $TEST_JOB_ID; do
+          echo "##[section]Test job $TEST_JOB_ID is still running..."
+          sleep 60
+        done
+
+        echo "Waiting for final status via sacct..."
+        LOOP_COUNT=0
+        MAX_LOOPS=120  # Maximum of 120 loops (120 minutes)
+        while true; do
+          STATE=$(sacct -j "$TEST_JOB_ID" --format=JobID,State --noheader | awk '$1 ~ /\.batch$/ { print $2; exit }' | xargs)
+          echo "##[section]Test job state: $STATE"
+          if [[ "$STATE" == "COMPLETED" ]]; then
+            break
+          elif [[ "$STATE" =~ ^(FAILED|CANCELLED|TIMEOUT)$ ]]; then
+            echo "Test failed with state $STATE"
+            break
+          fi
+          sleep 60
+          LOOP_COUNT=$((LOOP_COUNT + 1))
+          if [ $LOOP_COUNT -ge $MAX_LOOPS ]; then
+            echo "Time limit reached while waiting for final status."
+            exit 1  # Exit with an error code if time limit is reached
+          fi
+        done
+
+        echo "Checking test result json for failures..."
+        TEST_JSON=$(find "$(Pipeline.Workspace)" -name 'rccl-tests*.json')
+        if [ -z "$TEST_JSON" ]; then
+          echo "##vso[task.logissue type=error]No $TEST_JSON file(s) found"
+          echo "##vso[task.complete result=Failed;]DONE"
+          exit 1
+        fi
+
+        #echo "Checking test result XML for failures..."
+        #TEST_XML=$(find "$(Pipeline.Workspace)" -name 'rccl-tests_output.xml' | head -n1)
+        #if [ -z "$TEST_XML" ]; then
+        #  echo "##vso[task.logissue type=error]No $TES_XML file found"
+        #  echo "##vso[task.complete result=Failed;]DONE"
+        #  exit 1
+        #fi
+
+        #if grep -q 'failures="[^0]' "$TEST_XML"; then
+        #  echo "##vso[task.logissue type=error]Test failures detected in $TEST_XML"
+        #  echo "##vso[task.complete result=Failed;]DONE"
+        #  exit 1
+        #else
+        #  echo "No test failures detected."
+        #fi
+  - task: Bash@3
+    displayName: Test Logs
+    condition: always()
+    inputs:
+      targetType: inline
+      script: |
+        cat rccl-tests-${TEST_JOB_ID}.out || echo "No log found"
+#  - task: PublishTestResults@2
+#    displayName: 'Publish Results'
+#    condition: succeededOrFailed()
+#    inputs:
+#      searchFolder: $(Pipeline.Workspace)
+#      testResultsFormat: JUnit
+#      testResultsFiles: '**/rccl-tests_output.xml'
diff --git a/.azuredevops/tests/pytest/HelloWorld.py b/.azuredevops/tests/pytest/HelloWorld.py
new file mode 100644
index 0000000..52f05a1
--- /dev/null
+++ b/.azuredevops/tests/pytest/HelloWorld.py
@@ -0,0 +1,5 @@
+import pytest
+
+def test_HelloWorld():
+    greeting = "Hello, World!"
+    assert greeting == "Hello, World!"
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
new file mode 100755
index 0000000..35c9980
--- /dev/null
+++ b/.github/CODEOWNERS
@@ -0,0 +1,5 @@
+* @wenkaidu @gilbertlee-amd @PedramAlizadeh @nusislam @nileshnegi @KawtharShafie @AtlantaPepsi @mberenjk @corey-derochie-amd @mustafabar @thananon @JhaShweta1 @BertanDogancay @rahulvaidya20 @isaki001 @PJAvinash @AbandiGa @Nikhil-Nunna @haripriya-amd @atulkulk @ddebonis-amd @amd-mengshwu @Kapil-Shyam-Pawar @weilewei @nawrinsu @speriaswamy-amd# Documentation files
+docs/ @ROCm/rocm-documentation
+*.md @ROCm/rocm-documentation
+*.rst @ROCm/rocm-documentation
+.readthedocs.yaml @ROCm/rocm-documentation
diff --git a/.github/PULL_REQUEST_TEMPLATE.md b/.github/PULL_REQUEST_TEMPLATE.md
new file mode 100644
index 0000000..636de79
--- /dev/null
+++ b/.github/PULL_REQUEST_TEMPLATE.md
@@ -0,0 +1,23 @@
+## Details
+___Do not mention proprietary info or link to internal work items in this PR.___
+
+**Work item:** _"Internal", or link to GitHub issue (if applicable)._
+
+**What were the changes?**  
+_One sentence describing the work done._
+
+**Why were the changes made?**  
+_Explain the motivation behind the work. Provide any publicly-available historical context._
+
+**How was the outcome achieved?**  
+_Technical details behind the work. Explain any publicly-available hardware peculiarities._
+
+**Additional Documentation:**  
+_What else should the reviewer know?_
+
+## Approval Checklist
+___Do not approve until these items are satisfied.___
+- [ ] Verify the CHANGELOG has been updated, if
+  - there are any NCCL API version changes,
+  - any changes impact library users, and/or
+  - any changes impact any other ROCm library.
diff --git a/.github/dependabot.yml b/.github/dependabot.yml
new file mode 100644
index 0000000..848c7f5
--- /dev/null
+++ b/.github/dependabot.yml
@@ -0,0 +1,17 @@
+# To get started with Dependabot version updates, you'll need to specify which
+# package ecosystems to update and where the package manifests are located.
+# Please see the documentation for all configuration options:
+# https://docs.github.com/github/administering-a-repository/configuration-options-for-dependency-updates
+
+version: 2
+updates:
+  - package-ecosystem: "pip" # See documentation for possible values
+    directory: "/docs/sphinx" # Location of package manifests
+    open-pull-requests-limit: 10
+    schedule:
+      interval: "daily"
+    labels:
+      - "dependencies"
+      - "ci:docs-only"
+    reviewers:
+      - "samjwu"
diff --git a/.github/scripts/therock_configure_ci.py b/.github/scripts/therock_configure_ci.py
new file mode 100644
index 0000000..2afff17
--- /dev/null
+++ b/.github/scripts/therock_configure_ci.py
@@ -0,0 +1,131 @@
+import fnmatch
+import json
+import os
+from pathlib import Path
+import subprocess
+import sys
+from typing import Iterable, Optional, Mapping
+
+def gha_set_output(vars: Mapping[str, str | Path]):
+    """Sets values in a step's output parameters.
+
+    This appends to the file located at the $GITHUB_OUTPUT environment variable.
+
+    See
+      * https://docs.github.com/en/actions/reference/workflow-commands-for-github-actions#setting-an-output-parameter
+      * https://docs.github.com/en/actions/writing-workflows/choosing-what-your-workflow-does/passing-information-between-jobs
+    """
+    print(f"Setting github output:\n{vars}")
+
+    step_output_file = os.getenv("GITHUB_OUTPUT")
+    if not step_output_file:
+        print("  Warning: GITHUB_OUTPUT env var not set, can't set github outputs")
+        return
+
+    with open(step_output_file, "a") as f:
+        f.writelines(f"{k}={str(v)}" + "\n" for k, v in vars.items())
+
+def get_modified_paths(base_ref: str) -> Optional[Iterable[str]]:
+    """Returns the paths of modified files relative to the base reference."""
+    try:
+        return subprocess.run(
+            ["git", "diff", "--name-only", base_ref],
+            stdout=subprocess.PIPE,
+            check=True,
+            text=True,
+            timeout=60,
+        ).stdout.splitlines()
+    except TimeoutError:
+        print(
+            "Computing modified files timed out. Not using PR diff to determine"
+            " jobs to run.",
+            file=sys.stderr,
+        )
+        return None
+    
+GITHUB_WORKFLOWS_CI_PATTERNS = [
+    "therock*.yml",
+]
+
+
+def is_path_workflow_file_related_to_ci(path: str) -> bool:
+    return any(
+        fnmatch.fnmatch(path, ".github/workflows/" + pattern)
+        for pattern in GITHUB_WORKFLOWS_CI_PATTERNS
+    )
+    
+def check_for_workflow_file_related_to_ci(paths: Optional[Iterable[str]]) -> bool:
+    if paths is None:
+        return False
+    return any(is_path_workflow_file_related_to_ci(p) for p in paths)
+
+# Paths matching any of these patterns are considered to have no influence over
+# build or test workflows so any related jobs can be skipped if all paths
+# modified by a commit/PR match a pattern in this list.
+SKIPPABLE_PATH_PATTERNS = [
+    "docs/*",
+    "*.gitignore",
+    "*.md",
+    "*LICENSE*",
+    "*NOTICES*",
+    '.github/CODEOWNERS',
+    '.github/*.md',
+    '.github/dependabot.yml',
+    '.azuredevops*',
+]
+
+def is_path_skippable(path: str) -> bool:
+    """Determines if a given relative path to a file matches any skippable patterns."""
+    return any(fnmatch.fnmatch(path, pattern) for pattern in SKIPPABLE_PATH_PATTERNS)
+
+def check_for_non_skippable_path(paths: Optional[Iterable[str]]) -> bool:
+    """Returns true if at least one path is not in the skippable set."""
+    if paths is None:
+        return False
+    return any(not is_path_skippable(p) for p in paths)
+
+def should_ci_run_given_modified_paths(paths: Optional[Iterable[str]]) -> bool:
+    """Returns true if CI workflows should run given a list of modified paths."""
+
+    if paths is None:
+        print("No files were modified, skipping TheRock CI jobs")
+        return False
+
+    paths_set = set(paths)
+    github_workflows_paths = set(
+        [p for p in paths if p.startswith(".github/workflows")]
+    )
+    other_paths = paths_set - github_workflows_paths
+    
+    related_to_ci = check_for_workflow_file_related_to_ci(github_workflows_paths)
+    contains_other_non_skippable_files = check_for_non_skippable_path(other_paths)
+
+    print("should_ci_run_given_modified_paths findings:")
+    print(f"  contains_other_non_skippable_files: {contains_other_non_skippable_files}")
+
+    if related_to_ci:
+        print("Enabling build jobs since a related workflow file was modified")
+        return True
+    elif contains_other_non_skippable_files:
+        print("Enabling TheRock CI jobs since a non-skippable path was modified")
+        return True
+    else:
+        print(
+            "Only unrelated and/or skippable paths were modified, skipping TheRock CI jobs"
+        )
+        return False
+
+def main(args):
+    base_ref = args.get("base_ref")
+    modified_paths = get_modified_paths(base_ref)
+    print("modified_paths (max 200):", modified_paths[:200])
+    enable_jobs = should_ci_run_given_modified_paths(modified_paths)
+    output = {
+        'enable_therock_ci': json.dumps(enable_jobs)
+    }
+    gha_set_output(output)
+
+if __name__ == "__main__":
+    args = {}
+    args["base_ref"] = os.environ.get("BASE_REF", "HEAD^1")
+    main(args)
diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml
new file mode 100644
index 0000000..d390446
--- /dev/null
+++ b/.github/workflows/therock-ci-linux.yml
@@ -0,0 +1,126 @@
+name: TheRock CI Linux
+
+on:
+  workflow_call:
+    inputs:
+      amdgpu_families:
+        type: string
+      extra_cmake_options:
+        type: string
+
+permissions:
+  contents: read
+
+jobs:
+  therock-build-linux:
+    name: Build Linux Packages
+    runs-on: azure-linux-scale-rocm
+    permissions:
+      id-token: write
+    container:
+      image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:283673fe3e1bf498d079e3f386b794af1b4f71845a9a0107c6cf7aa304dce050
+      options: -v /runner/config:/home/awsconfig/
+    env:
+      AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }}
+      TEATIME_FORCE_INTERACTIVE: 0
+      AWS_SHARED_CREDENTIALS_FILE: /home/awsconfig/credentials.ini
+    steps:
+      - name: Checkout TheRock repository
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          repository: "ROCm/TheRock"
+          ref: 16ee54fb580a4dde62dc4133f978e73370a545af
+
+      - name: Checkout rccl repository
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          repository: "ROCm/rccl"
+          path: rccl
+
+      - name: Checkout rccl-tests repository
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          repository: "ROCm/rccl-tests"
+          path: rccl-tests
+
+      - name: Runner Health Settings
+        run: |
+          df -h
+          echo cmake --version
+          echo "Installed Python versions:"
+          ls -d /opt/python
+          echo "python: $(which python), python3: $(which python3)"
+          echo "Git version: $(git --version)"
+          git config --global --add safe.directory $PWD
+          git config fetch.parallel 10
+
+      - name: Fetch sources
+        run: |
+          ./build_tools/fetch_sources.py --jobs 12
+
+      - name: Install python deps
+        run: |
+          pip install -r requirements.txt
+          pip freeze
+
+      - name: Configure Projects
+        env:
+          amdgpu_families: ${{ env.AMDGPU_FAMILIES }}
+          package_version: ADHOCBUILD
+          extra_cmake_options: ${{ inputs.extra_cmake_options }}
+          BUILD_DIR: build
+        run: |
+          python3 build_tools/github_actions/build_configure.py
+
+      - name: Build therock-dist
+        run: cmake --build build
+
+      - name: Build therock-archives
+        run: cmake --build build --target therock-archives
+
+      - name: Report
+        #if: ${{ !cancelled() }}
+        run: |
+          echo "Full SDK du:"
+          echo "------------"
+          du -h -d 1 build/dist/rocm
+          echo "Artifact Archives:"
+          echo "------------------"
+          ls -lh build/artifacts/*.tar.xz
+          echo "Artifacts:"
+          echo "----------"
+          du -h -d 1 build/artifacts
+
+      - name: Configure AWS Credentials for non-forked repos
+        if: ${{ always() && !github.event.pull_request.head.repo.fork }}
+        uses: aws-actions/configure-aws-credentials@7474bc4690e29a8392af63c5b98e7449536d5c3a # v4.3.1
+        with:
+          aws-region: us-east-2
+          role-to-assume: arn:aws:iam::692859939525:role/therock-artifacts-external
+
+      - name: Post Build Upload
+        if: always()
+        run: |
+          python3 build_tools/github_actions/post_build_upload.py \
+            --run-id ${{ github.run_id }} \
+            --amdgpu-family ${{ env.AMDGPU_FAMILIES }} \
+            --build-dir build \
+            --upload
+
+  therock-test-linux-multi-node:
+    name: "Test multi-node"
+    needs: [therock-build-linux]
+    uses: ./.github/workflows/therock-test-packages-multi-node.yml
+    with:
+      amdgpu_families: ${{ inputs.amdgpu_families }}
+      test_runs_on: vultr-linux-rocm
+      artifact_run_id: ${{ github.run_id }}
+
+  therock-test-linux-single-node:
+    name: "Test single-node"
+    needs: [therock-build-linux]
+    uses: ./.github/workflows/therock-test-packages-single-node.yml
+    with:
+      amdgpu_families: ${{ inputs.amdgpu_families }}
+      test_runs_on: linux-mi325-1gpu-ossci-rocm
+      artifact_run_id: ${{ github.run_id }}
diff --git a/.github/workflows/therock-ci.yml b/.github/workflows/therock-ci.yml
new file mode 100644
index 0000000..1d866b1
--- /dev/null
+++ b/.github/workflows/therock-ci.yml
@@ -0,0 +1,81 @@
+name: TheRock CI for rccl
+
+on:
+  push:
+    branches:
+      - develop
+  workflow_dispatch:
+
+permissions:
+  contents: read
+
+concurrency:
+  # A PR number if a pull request and otherwise the commit hash. This cancels
+  # queued and in-progress runs for the same PR (presubmit) or commit
+  # (postsubmit). The workflow name is prepended to avoid conflicts between
+  # different workflows.
+  group: ${{ github.workflow }}-${{ github.event.number || github.sha }}
+  cancel-in-progress: true
+
+jobs:
+  setup:
+    runs-on: ubuntu-24.04
+    env:
+      # The commit being checked out is the merge commit for a PR. Its first
+      # parent will be the tip of the base branch.
+      BASE_REF: HEAD^
+    outputs:
+      enable_therock_ci: ${{ steps.configure.outputs.enable_therock_ci }}
+    steps:
+      - name: "Checking out repository"
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
+        with:
+          # We need the parent commit to do a diff
+          fetch-depth: 2
+
+      - name: "Configuring CI options"
+        id: configure
+        run: python .github/scripts/therock_configure_ci.py
+
+  therock-ci-linux:
+    name: TheRock CI Linux
+    needs: setup
+    if: ${{ needs.setup.outputs.enable_therock_ci == 'true' }}
+    permissions:
+      contents: read
+      id-token: write
+    uses: ./.github/workflows/therock-ci-linux.yml
+    secrets: inherit
+    with:
+      amdgpu_families: "gfx94X-dcgpu"
+      extra_cmake_options: >
+        -DTHEROCK_ENABLE_ALL=OFF 
+        -DTHEROCK_BUILD_TESTING=ON 
+        -DTHEROCK_BUNDLE_SYSDEPS=ON 
+        -DTHEROCK_ENABLE_COMM_LIBS=ON 
+        -DTHEROCK_ENABLE_ROCPROFV3=ON 
+        -DTHEROCK_USE_EXTERNAL_RCCL=ON 
+        -DTHEROCK_USE_EXTERNAL_RCCL_TESTS=ON 
+        -DTHEROCK_RCCL_SOURCE_DIR=./rccl 
+        -DTHEROCK_RCCL_TESTS_SOURCE_DIR=./rccl-tests
+        -DTHEROCK_ENABLE_MPI=ON
+
+  therock_ci_summary:
+    name: TheRock CI Summary
+    if: always()
+    needs:
+      - setup
+      - therock-ci-linux
+    runs-on: ubuntu-24.04
+    steps:
+      - name: Output failed jobs
+        run: |
+          echo '${{ toJson(needs) }}'
+          FAILED_JOBS="$(echo '${{ toJson(needs) }}' \
+            | jq --raw-output \
+            'map_values(select(.result!="success" and .result!="skipped")) | keys | join(",")' \
+          )"
+          if [[ "${FAILED_JOBS}" != "" ]]; then
+            echo "The following jobs failed: ${FAILED_JOBS}"
+            exit 1
+          fi
diff --git a/.github/workflows/therock-test-packages-multi-node.yml b/.github/workflows/therock-test-packages-multi-node.yml
new file mode 100644
index 0000000..0aa2def
--- /dev/null
+++ b/.github/workflows/therock-test-packages-multi-node.yml
@@ -0,0 +1,57 @@
+name: TheRock Test Packages multi-node
+
+on:
+  workflow_call:
+    inputs:
+      amdgpu_families:
+        type: string
+      test_runs_on:
+        type: string
+      artifact_run_id:
+        type: string
+  workflow_dispatch:
+    inputs:
+      amdgpu_families:
+        type: string
+      test_runs_on:
+        type: string
+      artifact_run_id:
+        type: string
+
+permissions:
+  contents: read
+
+jobs:
+  test_rccl_multi_node:
+    name: 'Test multi-node'
+    runs-on: ${{ inputs.test_runs_on }}
+    defaults:
+      run:
+        shell: bash
+    env:
+      VENV_DIR: ${{ github.workspace }}/.venv
+      ARTIFACT_RUN_ID: "${{ inputs.artifact_run_id }}"
+      OUTPUT_ARTIFACTS_DIR: /home/arravikum/dist_new/dist/rocm
+      THEROCK_BIN_DIR: "./build/bin"
+    steps:
+      - name: Checkout Repository
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          repository: "ROCm/TheRock"
+          ref: 890c856134d955441790c8ed2d60ad4fb027f4e5
+
+      - name: Run setup test environment workflow
+        uses: './.github/actions/setup_test_environment'
+        with:
+          ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }}
+          AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }}
+          OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }}
+          VENV_DIR: ${{ env.VENV_DIR }}
+          FETCH_ARTIFACT_ARGS: "--rccl"
+          IS_PR_FROM_FORK: ${{ github.event.pull_request.head.repo.fork }}
+
+      - name: Test
+        run: |
+          source /home/arravikum/TheRock/.venv/bin/activate
+          cd /home/arravikum/cvs
+          pytest -vvv --log-file=/tmp/rccl_log.log -s ./tests/rccl/rccl_multinode_cvs.py --cluster_file ./input/cluster.json --config_file ./input/mi300_config.json --html=/var/www/html/cvs/ci_test_report.html --capture=tee-sys --self-contained-html
diff --git a/.github/workflows/therock-test-packages-single-node.yml b/.github/workflows/therock-test-packages-single-node.yml
new file mode 100644
index 0000000..809ae81
--- /dev/null
+++ b/.github/workflows/therock-test-packages-single-node.yml
@@ -0,0 +1,69 @@
+name: TheRock Test Packages single-node
+
+on:
+  workflow_call:
+    inputs:
+      amdgpu_families:
+        type: string
+      test_runs_on:
+        type: string
+      artifact_run_id:
+        type: string
+  workflow_dispatch:
+    inputs:
+      amdgpu_families:
+        type: string
+      test_runs_on:
+        type: string
+      artifact_run_id:
+        type: string
+
+permissions:
+  contents: read
+
+jobs:
+  test_rccl_single_node:
+    name: 'Test single-node'
+    runs-on: ${{ inputs.test_runs_on }}
+    container:
+      image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:405945a40deaff9db90b9839c0f41d4cba4a383c1a7459b28627047bf6302a26
+      options: --ipc host
+        --group-add video
+        --device /dev/kfd
+        --device /dev/dri
+        --group-add 992
+        --env-file /etc/podinfo/gha-gpu-isolation-settings
+    defaults:
+      run:
+        shell: bash
+    env:
+      VENV_DIR: ${{ github.workspace }}/.venv
+      ARTIFACT_RUN_ID: "${{ inputs.artifact_run_id }}"
+      OUTPUT_ARTIFACTS_DIR: "./build"
+      THEROCK_BIN_DIR: "./build/bin"
+    steps:
+      - name: Checkout Repository
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+        with:
+          repository: "ROCm/TheRock"
+          ref: f89dcd5c5625baecb467b9287e952c5c819073fd
+
+      - name: Run setup test environment workflow
+        uses: './.github/actions/setup_test_environment'
+        with:
+          ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }}
+          AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }}
+          OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }}
+          VENV_DIR: ${{ env.VENV_DIR }}
+          FETCH_ARTIFACT_ARGS: "--rccl --tests"
+          IS_PR_FROM_FORK: ${{ github.event.pull_request.head.repo.fork }}
+
+      - name: Test
+        timeout-minutes: 15
+        # Currently, TheRock CI in RCCL always builds with MPI-supported enabled which causes the
+        # RCCL correctness tests to fail on the mi325 runners which don't have MPI pre-installed.
+        # TODO (geomin12): Rebuild rccl-tests without MPI to enable RCCL correctness tests.
+        run: |
+          pytest ./build_tools/github_actions/test_executable_scripts/test_rccl.py -v -s \
+            --log-cli-level=info \
+            -k "not test_rccl_correctness_tests"
diff --git a/.gitignore b/.gitignore
index 87ad949..d3dedfb 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,4 +1,8 @@
 # Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
-/build
 *.gcov
 /coverage/
+build/
+ext/
+
+# Visual Studio Code
+.vscode
\ No newline at end of file
diff --git a/.gitmodules b/.gitmodules
new file mode 100644
index 0000000..eae2c51
--- /dev/null
+++ b/.gitmodules
@@ -0,0 +1,10 @@
+[submodule "ext-src/mscclpp"]
+	path = ext-src/mscclpp
+	url = https://github.com/microsoft/mscclpp.git
+	ignore = dirty
+	shallow = true
+[submodule "ext-src/json"]
+	path = ext-src/json
+	url = https://github.com/nlohmann/json.git
+	ignore = dirty
+	shallow = true
diff --git a/.jenkins/common.groovy b/.jenkins/common.groovy
deleted file mode 100644
index d3ad3a2..0000000
--- a/.jenkins/common.groovy
+++ /dev/null
@@ -1,41 +0,0 @@
-// This file is for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-def runCompileCommand(platform, project, jobName)
-{
-    project.paths.construct_build_prefix()
-
-    def command = """#!/usr/bin/env bash
-                set -x
-                cd ${project.paths.project_build_prefix}
-                LD_LIBRARY_PATH=/opt/rocm/hcc/lib ${project.paths.build_command}
-            """
-
-    platform.runCommand(this,command)
-}
-
-def runTestCommand (platform, project, gfilter)
-{
-    String sudo = auxiliary.sudo(platform.jenkinsLabel)
-
-    def command = """#!/usr/bin/env bash
-                set -x
-                cd ${project.paths.project_build_prefix}/build/release/test
-                ${sudo} ulimit -l unlimited
-                ulimit -a
-                ${sudo} RCCL_ENABLE_SIGNALHANDLER=1 NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./UnitTests --gtest_filter=${gfilter} --gtest_output=xml --gtest_color=yes
-            """
-
-   platform.runCommand(this, command)
-   junit "${project.paths.project_build_prefix}/build/release/test/*.xml"
-}
-
-def runPackageCommand(platform, project, jobName)
-{
-    def packageHelper = platform.makePackage(platform.jenkinsLabel,"${project.paths.project_build_prefix}/build/release")
-
-    platform.runCommand(this, packageHelper[0])
-    platform.archiveArtifacts(this, packageHelper[1])
-}
-
-return this
diff --git a/.jenkins/extended.groovy b/.jenkins/extended.groovy
deleted file mode 100644
index 647cd2d..0000000
--- a/.jenkins/extended.groovy
+++ /dev/null
@@ -1,88 +0,0 @@
-#!/usr/bin/env groovy
-// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
-@Library('rocJenkins@pong') _
-
-// This is file for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI = 
-{
-    nodeDetails, jobName->
-
-    def prj  = new rocProject('rccl', 'Extended')
-
-    prj.timeout.test = 1440
-    prj.paths.build_command = './install.sh -t '
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = false
-
-    def commonGroovy
-
-    def compileCommand =
-    {
-        platform, project->
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-    
-    def testCommand =
-    {
-        platform, project->
-
-        commonGroovy.runTestCommand(platform, project, "*")
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-        
-        commonGroovy.runPackageCommand(platform, project, jobName)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi-hipclang":[pipelineTriggers([cron('0 1 * * 0')])]]
-
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([centos8:['8gfx906']])]
-    
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-    
-    propertyList.each 
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    jobNameList.each 
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(jobName) {
-                runCI(nodeDetails, jobName)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(urlJobName) {
-            runCI([ubuntu18:['4gfx906']], urlJobName)
-        }
-    }
-}
diff --git a/.jenkins/precheckin.groovy b/.jenkins/precheckin.groovy
deleted file mode 100644
index 61dc43f..0000000
--- a/.jenkins/precheckin.groovy
+++ /dev/null
@@ -1,88 +0,0 @@
-#!/usr/bin/env groovy
-// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
-@Library('rocJenkins@pong') _
-
-// This is file for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCI =
-{
-    nodeDetails, jobName->
-
-    def prj  = new rocProject('rccl', 'PreCheckin')
-
-    prj.timeout.test = 1440
-    prj.paths.build_command = './install.sh -t '
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = false
-
-    def commonGroovy
-
-    def compileCommand =
-    {
-        platform, project->
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-
-    def testCommand =
-    {
-        platform, project->
-
-        commonGroovy.runTestCommand(platform, project, "-*ManagedMem")
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project, jobName)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-ci: {
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi-hipclang":[pipelineTriggers([cron('0 1 * * 0')])]]
-
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([sles15sp1:['4gfx906'],centos8:['8gfx908'],centos7:['8gfx906'],ubuntu18:['4gfx906', '4gfx908']])]
-
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    propertyList.each
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(jobName) {
-                runCI(nodeDetails, jobName)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(urlJobName) {
-            runCI([ubuntu18:['4gfx906']], urlJobName)
-        }
-    }
-}
diff --git a/.jenkins/staticanalysis.groovy b/.jenkins/staticanalysis.groovy
deleted file mode 100644
index aa5fec3..0000000
--- a/.jenkins/staticanalysis.groovy
+++ /dev/null
@@ -1,68 +0,0 @@
-#!/usr/bin/env groovy
-// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
-@Library('rocJenkins@pong') _
-
-// This is file for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCompileCommand(platform, project, jobName, boolean debug=false)
-{
-    project.paths.construct_build_prefix()
-    
-    def command = """#!/usr/bin/env bash
-            set -x
-            ${project.paths.project_build_prefix}/docs/run_doc.sh
-            """
-
-    try
-    {
-        platform.runCommand(this, command)
-    }
-    catch(e)
-    {
-        throw e
-    }
-
-    publishHTML([allowMissing: false,
-                alwaysLinkToLastBuild: false,
-                keepAll: false,
-                reportDir: "${project.paths.project_build_prefix}/docs/source/_build/html",
-                reportFiles: "index.html",
-                reportName: "Documentation",
-                reportTitles: "Documentation"])
-}
-
-def runCI = 
-{
-    nodeDetails, jobName->
-
-    def prj  = new rocProject('rccl-internal', 'StaticAnalysis')
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = false
-    boolean staticAnalysis = true
-
-    def compileCommand =
-    {
-        platform, project->
-
-        runCompileCommand(platform, project, jobName, false)
-    }
-
-    buildProject(prj , formatCheck, nodes.dockerArray, compileCommand, null, null, staticAnalysis)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * 6')])]))
-    stage(urlJobName) {
-        runCI([ubuntu20:['cpu']], urlJobName)
-    }
-}
diff --git a/.jenkins/staticlibrary.groovy b/.jenkins/staticlibrary.groovy
deleted file mode 100644
index 2ab6696..0000000
--- a/.jenkins/staticlibrary.groovy
+++ /dev/null
@@ -1,85 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path;
-
-def runCI =
-{
-    nodeDetails, jobName->
-    
-    def prj = new rocProject('rccl', 'Static Library PreCheckin')
-
-    prj.timeout.test = 1440
-    prj.paths.build_command = './install.sh -t --static'
-    
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-     
-    def compileCommand =
-    {
-        platform, project->
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-
-    
-    def testCommand =
-    {
-        platform, project->
-
-        commonGroovy.runTestCommand(platform, project, "*sum_float32*")
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project, jobName)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi":[pipelineTriggers([cron('0 1 * * 0')])], 
-                        "compute-rocm-dkms-no-npi-hipclang":[pipelineTriggers([cron('0 1 * * 0')])],
-                        "rocm-docker":[]]
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["compute-rocm-dkms-no-npi":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]), 
-                       "compute-rocm-dkms-no-npi-hipclang":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]), 
-                       "rocm-docker":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    propertyList.each 
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(jobName) {
-                runCI(nodeDetails, jobName)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(urlJobName) {
-            runCI([ubuntu16:['4gfx906']], urlJobName)
-        }
-    }
-}
diff --git a/.readthedocs.yaml b/.readthedocs.yaml
new file mode 100644
index 0000000..ced8a98
--- /dev/null
+++ b/.readthedocs.yaml
@@ -0,0 +1,18 @@
+# Read the Docs configuration file
+# See https://docs.readthedocs.io/en/stable/config-file/v2.html for details
+
+version: 2
+
+build:
+   os: ubuntu-22.04
+   tools:
+      python: "3.10"
+
+sphinx:
+   configuration: docs/conf.py
+
+formats: [htmlzip, pdf, epub]
+
+python:
+   install:
+   - requirements: docs/sphinx/requirements.txt
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 448c973..6da80ac 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,7 +1,209 @@
-# Change Log for RCCL
+# Changelog for RCCL
 
 Full documentation for RCCL is available at [https://rccl.readthedocs.io](https://rccl.readthedocs.io)
 
+## Unreleased - RCCL 2.27.7 for ROCm 7.1.0
+
+### Added
+* Added `RCCL_P2P_BATCH_THRESHOLD` to set the message size limit for batching P2P operations. This mainly affects small message performance for alltoall at a large scale but also applies to alltoallv.
+* Added `RCCL_P2P_BATCH_ENABLE` to enable batching P2P operations to receive performance gains for smaller messages up to 4MB for alltoall when the workload requires it. This is to avoid performance dips for larger messages.
+
+### Changed
+
+* The MSCCL++ feature is now disabled by default. The `--disable-mscclpp` build flag is replaced with `--enable-mscclpp` in the `rccl/install.sh` script.
+* Compatibility with NCCL 2.27.7
+
+### Resolved issues
+* Improve small message performance for alltoall by enabling and optimizing batched P2P operations. 
+
+### Known issues
+* Symmetric memory kernels are currently disabled due to ongoing CUMEM enablement work.
+
+## RCCL 2.26.6 for ROCm 7.0.0
+
+### Resolved issues
+
+* Resolved an issue when using more than 64 channels when multiple collectives are used in the same `ncclGroup()` call.
+* Fixed unit test failures in tests ending with `ManagedMem` and `ManagedMemGraph` suffixes.
+* Suboptimal algorithmic switching point for AllReduce on MI300x.
+* Fixed the known issue "When splitting a communicator using `ncclCommSplit` in some GPU configurations, MSCCL initialization can cause a segmentation fault." with a design change to use `comm` instead of `rank` for `mscclStatus`. The Global map for `comm` to `mscclStatus` is still not thread safe but should be explicitly handled by mutexes for read writes. This is tested for correctness, but there is a plan to use a thread-safe map data structure in upcoming changes.
+
+### Added
+
+* Added new GPU target `gfx950`.
+* Added support for `unroll=1` in device-code generation to improve performance,
+* Set a default of 112 channels for a single node with `8 * gfx950`,
+* Enabled LL128 protocol on `gfx950`.
+* Added MSCCL support for AllGather multinode gfx942/gfx950 (i.e., 16 and 32 GPUs). To enable, set the environment variable `RCCL_MSCCL_FORCE_ENABLE=1`. Max message size for MSCCL AllGather usage is `12292 * sizeof(datatype) * nGPUs`.
+* Thread thresholds for LL/LL128 are selected in Tuning Models for the MI300X. This impacts the number of channels used for AG and RS. Channel tuning model is bypassed if `NCCL_THREAD_THRESHOLDS`, `NCCL_MIN_NCHANNELS', or 'NCCL_MAX_NCHANNELS` are set.
+* Multi-node tuning for AllGather, AllReduce, and ReduceScatter that leverages LL/LL64/LL128 protocol to use nontemporal vector load/store for tunable message size ranges.
+* LL/LL128 usage ranges for AR, AG, and RS are part of the tuning models, which enable architecture-specific tuning in conjunction with the existing Rome Models scheme in RCCL.
+* Two new APIs are exposed as part of an initiative to separate RCCL code. These APIs are `rcclGetAlgoInfo` and `rcclFuncMaxSendRecvCount`. However, user-level invocation requires that RCCL be built with `RCCL_EXPOSE_STATIC` enabled.
+* Enabled double-buffering in `reduceCopyPacks` to trigger pipelining, especially to overlap `bf16` arithmetic and bridge the gap between `fp32` performance and `bf16` for both `gfx942` and `gfx950`. Pipelining has been made tunable via `rcclSetPipelining`, similar to algorithms/protocols so that regression is avoided in certain message sizes.
+* Added a direct allgather algorithm. This is enabled by default for multi-node if there are 16 nodes or fewer. The message size threshold is 4MB.
+* Added `RCCL_OVERRIDE_PROTO` and `RCCL_OVERRIDE_ALGO` to allow direct replacement of protocol and algorithm choices. Unlike `NCCL_PROTO` and `NCCL_ALGO`, which re-run the model across enabled combinations and may not guarantee the intended override, these new options enforce the specified selections explicitly.
+
+### Changed
+
+* Compatibility with NCCL 2.23.4
+* Compatibility with NCCL 2.24.3
+* Compatibility with NCCL 2.25.1
+* Compatibility with NCCL 2.26.6
+
+## RCCL 2.22.3 for ROCm 6.4.2
+
+### Added
+
+* Added support for the LL128 protocol on gfx942.
+
+## RCCL 2.22.3 for ROCm 6.4.1
+
+### Resolved issues
+
+* Fixed the accuracy issue for MSCCLPP `allreduce7` kernel in graph mode.
+* Fixed IntraNet performance.
+* Fixed an issue where, in rare circumstances, the application could stop responding due to a proxy thread synchronization issue.
+
+### Known issues
+
+* When splitting a communicator using `ncclCommSplit` in some GPU configurations, MSCCL initialization can cause a segmentation fault.
+  The recommended workaround is to disable MSCCL with `export RCCL_MSCCL_ENABLE=0`.
+* Within the RCCL-UnitTests test suite, failures occur in tests ending with the `ManagedMem` and `ManagedMemGraph` suffixes. These failures only affect the test results and do not affect the RCCL component itself. This issue will be resolved in the next major release.
+
+## RCCL 2.22.3 for ROCm 6.4.0
+
+### Added
+
+* `RCCL_SOCKET_REUSEADDR` and `RCCL_SOCKET_LINGER` environment parameters.
+* Setting `NCCL_DEBUG=TRACE NCCL_DEBUG_SUBSYS=VERBS` will generate traces for fifo and data `ibv_post_sends`.
+* Added `--log-trace` flag to enable traces through the install.sh script (e.g. `./install.sh --log-trace`).
+
+### Changed
+
+* Compatibility with NCCL 2.22.3
+* Added support for the rail-optimized tree algorithm for the MI300 series. This feature requires the use of all eight GPUs within
+  each node. It limits NIC traffic to use only GPUs of the same index across nodes and should not impact performance
+  on non-rail-optimized network topologies. The original method of building trees can be enabled by setting the
+  environment variable `RCCL_DISABLE_RAIL_TREES=1`.
+* Additional debug information about how the trees are built can be logged to the GRAPH logging subsys by setting
+  `RCCL_OUTPUT_TREES=1`.
+* Added documentation about the NPS4 and CPX partition modes performance benefits on the MI300X.
+
+## RCCL 2.21.5 for ROCm 6.3.1
+
+### Added
+
+### Changed
+
+* Enhanced user documentation
+
+### Resolved issues
+
+* Corrected user help strings in `install.sh`
+
+## RCCL 2.21.5 for ROCm 6.3.0
+
+### Added
+
+* MSCCL++ integration for AllReduce and AllGather on gfx942
+* Performance collection to rccl_replayer
+* Tuner Plugin example for MI300
+* Tuning table for large number of nodes
+* Support for amdclang++
+* Allow NIC ID remapping using `NCCL_RINGS_REMAP` environment variable
+
+### Changed
+
+* Compatibility with NCCL 2.21.5
+* Increased channel count for MI300X multi-node
+* Enabled MSCCL for single-process multi-threaded contexts
+* Enabled gfx12
+* Enabled CPX mode for MI300X
+* Enabled tracing with rocprof
+* Improved version reporting
+* Enabled GDRDMA for Linux kernel 6.4.0+
+
+### Resolved issues
+
+* Fixed model matching with PXN enable
+
+## RCCL 2.20.5 for ROCm 6.2.1
+### Fixed
+- GDR support flag now set with DMABUF
+### Known issues
+- On systems running Linux kernel 6.8.0, such as Ubuntu 24.04, Direct Memory Access (DMA) transfers between the GPU and NIC are disabled and impacts multi-node RCCL performance.
+  - This issue was reproduced with RCCL 2.20.5 (ROCm 6.2.0 and 6.2.1) on systems with Broadcom Thor-2 NICs and affects other systems with RoCE networks using Linux 6.8.0 or newer.
+  - Older RCCL versions are also impacted.
+  - This issue will be addressed in a future ROCm release.
+
+## RCCL 2.20.5 for ROCm 6.2.0
+### Changed
+- Compatibility with NCCL 2.20.5
+- Compatibility with NCCL 2.19.4
+- Performance tuning for some collective operations on MI300
+- Enabled NVTX code in RCCL
+- Replaced rccl_bfloat16 with hip_bfloat16
+- NPKit updates:
+  - Removed warm-up iteration removal by default, need to opt in now
+  - Doubled the size of buffers to accommodate for more channels
+- Modified rings to be rail-optimized topology friendly
+- Replaced ROCmSoftwarePlatform links with ROCm links
+### Added
+- Support for fp8 and rccl_bfloat8
+- Support for using HIP contiguous memory
+- Implemented ROC-TX for host-side profiling
+- Enabled static build
+- Added new rome model
+- Added fp16 and fp8 cases to unit tests
+- New unit test for main kernel stack size
+- New -n option for topo_expl to override # of nodes
+- Improved debug messages of memory allocations
+### Fixed
+- Bug when configuring RCCL for only LL128 protocol
+- Scratch memory allocation after API change for MSCCL
+
+## RCCL 2.18.6 for ROCm 6.1.0
+### Changed
+- Compatibility with NCCL 2.18.6
+
+## RCCL 2.18.3 for ROCm 6.0.0
+### Changed
+- Compatibility with NCCL 2.18.3
+
+## RCCL 2.17.1-1 for ROCm 5.7.0
+### Changed
+- Compatibility with NCCL 2.17.1-1
+- Performance tuning for some collective operations
+### Added
+- Minor improvements to MSCCL codepath
+- NCCL_NCHANNELS_PER_PEER support
+- Improved compilation performance
+- Support for gfx94x
+### Fixed
+- Potential race-condition during ncclSocketClose()
+
+## RCCL 2.16.2 for ROCm 5.6.0
+### Changed
+- Compatibility with NCCL 2.16.2
+### Fixed
+- Remove workaround and use indirect function call
+
+## RCCL 2.15.5 for ROCm 5.5.0
+### Changed
+- Compatibility with NCCL 2.15.5
+- Unit test executable renamed to rccl-UnitTests
+### Added
+- HW-topology aware binary tree implementation
+- Experimental support for MSCCL
+- New unit tests for hipGraph support
+- NPKit integration
+### Fixed
+- rocm-smi ID conversion
+- Support for HIP_VISIBLE_DEVICES for unit tests
+- Support for p2p transfers to non (HIP) visible devices
+### Removed
+- Removed TransferBench from tools.  Exists in standalone repo: https://github.com/ROCm/TransferBench
+
 ## RCCL-2.13.4 for ROCm 5.4.0
 ### Changed
 - Compatibility with NCCL 2.13.4
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 67a2069..c8281f2 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,61 +1,388 @@
-# Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2019-2023 Advanced Micro Devices, Inc. All rights reserved.
+# Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
 
-cmake_minimum_required(VERSION 3.5)
-INCLUDE(CheckIncludeFiles)
-INCLUDE(CheckSymbolExists)
+# CMake version minimum requirements
+#==================================================================================================
+cmake_minimum_required(VERSION 3.16)
 
-# We use C++14 features, this will add compile option: -std=c++14
-set( CMAKE_CXX_STANDARD 14 )
-# Without this line, it will add -std=gnu++14 instead, which has some issues.
-set( CMAKE_CXX_EXTENSIONS OFF )
+# CMake Toolchain file to define compilers and path to ROCm
+#==================================================================================================
+if (NOT CMAKE_TOOLCHAIN_FILE)
+  set(CMAKE_TOOLCHAIN_FILE "${CMAKE_CURRENT_SOURCE_DIR}/toolchain-linux.cmake")
+  message(STATUS "CMAKE_TOOLCHAIN_FILE: ${CMAKE_TOOLCHAIN_FILE}")
+endif()
 
+# RCCL project
+#==================================================================================================
 project(rccl CXX)
 
-include(cmake/Dependencies.cmake)
+# Build options
+#==================================================================================================
+option(BUILD_ADDRESS_SANITIZER                 "Enable address sanitizer"                      OFF)
+option(BUILD_BFD                               "Enable custom backtrace (if bfd.h exists)"     OFF)
+option(BUILD_LOCAL_GPU_TARGET_ONLY             "Build only for GPUs detected on this machine"  OFF)
+option(BUILD_SHARED_LIBS                       "Build as shared library"                       ON)
+option(BUILD_TESTS                             "Build unit test programs"                      OFF)
+option(COLLTRACE                               "Collective Trace Option"                       ON)
+option(DUMP_ASM                                "Disassemble and dump"                          OFF)
+option(ENABLE_CODE_COVERAGE                    "Enable code coverage"                          OFF)
+option(ENABLE_MSCCL_KERNEL                     "Enable MSCCL while compiling"                  ON)
+option(ENABLE_MSCCLPP                          "Enable MSCCL++"                                OFF)
+option(ENABLE_MSCCLPP_CLIP                     "Enable MSCCL++ CLIP"                           OFF)
+option(ENABLE_MSCCLPP_EXECUTOR                 "Enable MSCCL++ Executor"                       OFF)
+option(ENABLE_MSCCLPP_FORMAT_CHECKS            "Enable formatting checks in MSCCL++"           OFF)
+option(ENABLE_NPKIT                            "Enable NPKit"                                  OFF)
+option(ENABLE_IFC                              "Enable indirect function call"                 OFF)
+option(GENERATE_SYM_KERNELS                    "Generate symmetric memory kernels"             OFF)
+option(INSTALL_DEPENDENCIES                    "Force install dependencies"                    OFF)
+option(ROCTX                                   "Enable ROCTX"                                  ON)
+option(PROFILE                                 "Enable profiling"                              OFF)
+option(TIMETRACE                               "Enable time-trace during compilation"          OFF)
+option(TRACE                                   "Enable additional tracing"                     OFF)
+option(FAULT_INJECTION                         "Enable fault injection"                        ON)
+
+# Default GPU architectures to build
+#==================================================================================================
+set(DEFAULT_GPUS
+      gfx906
+      gfx908
+      gfx90a
+      gfx942
+      gfx950
+      gfx1030
+      gfx1100
+      gfx1101
+      gfx1102
+      gfx1200
+      gfx1201)
+
+# Load CMake modules
+#==================================================================================================
+include(CheckIncludeFiles)
+include(CheckSymbolExists)
+include(cmake/Dependencies.cmake) # GTest, rocm-cmake, rocm_local_targets
+include(cmake/CheckSymbolExistsNoWarn.cmake)
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
-# Detect compiler support for target ID
-# This section is deprecated. Please use rocm_check_target_ids for future use.
-if( CMAKE_CXX_COMPILER MATCHES ".*/hipcc$" )
-    execute_process(COMMAND ${CMAKE_CXX_COMPILER} "--help"
-        OUTPUT_VARIABLE CXX_OUTPUT
-        OUTPUT_STRIP_TRAILING_WHITESPACE
-        ERROR_STRIP_TRAILING_WHITESPACE)
-    string(REGEX MATCH ".mcode\-object\-version" TARGET_ID_SUPPORT ${CXX_OUTPUT})
+# Build only for local GPU architecture
+if (BUILD_LOCAL_GPU_TARGET_ONLY)
+  message(STATUS "Building only for local GPU target")
+  if (COMMAND rocm_local_targets)
+    rocm_local_targets(DEFAULT_GPUS)
+  else()
+    message(WARNING "Unable to determine local GPU targets. Falling back to default GPUs.")
+  endif()
 endif()
 
+# Determine which GPU architectures to build for
+set(GPU_TARGETS "${DEFAULT_GPUS}" CACHE STRING "Target default GPUs if GPU_TARGETS is not defined.")
+
+# Modify GPU architectures for Address Sanitizer builds by appending "xnack+"
+if (BUILD_ADDRESS_SANITIZER)
+  SET(amdgpu_targets "")
+  foreach(amdgpu_target IN LISTS GPU_TARGETS)
+    if(NOT amdgpu_target STREQUAL "")
+      string(FIND "${amdgpu_target}" ":xnack+" HAS_XNACK_SUFFIX)
+      if(HAS_XNACK_SUFFIX EQUAL -1)
+        list(APPEND amdgpu_targets "${amdgpu_target}:xnack+")
+      else()
+        list(APPEND amdgpu_targets "${amdgpu_target}")
+      endif()
+    endif()
+  endforeach()
+  SET(GPU_TARGETS "${amdgpu_targets}")
+endif()
+
+# Check if clang compiler can offload to GPU_TARGETS
+if (COMMAND rocm_check_target_ids)
+  message(STATUS "Checking for ROCm support for GPU targets: " "${GPU_TARGETS}")
+  rocm_check_target_ids(SUPPORTED_GPUS TARGETS ${GPU_TARGETS})
+else()
+  message(WARNING "Unable to check for supported GPU targets. Falling back to default GPUs.")
+  set(SUPPORTED_GPUS ${DEFAULT_GPUS})
+endif()
+
+set(GPU_TARGETS "${SUPPORTED_GPUS}")
+message(STATUS "Compiling for ${GPU_TARGETS}")
+
+## NOTE: Reload rocm-cmake in order to update GPU_TARGETS
+include(cmake/Dependencies.cmake) # Reloading to use desired GPU_TARGETS instead of defaults
+
+# Try to establish ROCM_PATH (for find_package)
+#==================================================================================================
 if(NOT DEFINED ROCM_PATH)
-  get_filename_component(_real_path ${CMAKE_CXX_COMPILER} REALPATH)
-  get_filename_component(_new_path "${_real_path}" DIRECTORY)
-  get_filename_component(ROCM_PATH "${_new_path}/../.." REALPATH)
+  # Guess default location
+  set(ROCM_PATH "/opt/rocm")
+  message(WARNING "Unable to find ROCM_PATH: Falling back to ${ROCM_PATH}")
+else()
+  message(STATUS "ROCM_PATH found: ${ROCM_PATH}")
 endif()
+set(ENV{ROCM_PATH} ${ROCM_PATH})
 
+if("${CMAKE_CXX_COMPILER}" MATCHES ".*amdclang\\+\\+")
+  message(STATUS "Compiling with amdclang++")
+  set(COMPILER_EXE_NAME amdclang++)
+  set(COMPILER_GREP_STRING "AMD clang version")
+  set(COMPILER_AWK_CMD "awk -F\" \" '{ printf $4}'")
+elseif("${CMAKE_CXX_COMPILER}" MATCHES ".*clang\\+\\+")
+  message(STATUS "Compiling with clang++")
+  set(COMPILER_EXE_NAME clang++)
+  set(COMPILER_GREP_STRING "AMD clang version")
+  set(COMPILER_AWK_CMD "awk -F\" \" '{ printf $4}'")
+elseif("${CMAKE_CXX_COMPILER}" MATCHES ".*hipcc$")
+  message(STATUS "Compiling with hipcc")
+  set(COMPILER_EXE_NAME hipcc)
+  set(COMPILER_GREP_STRING "HIP version")
+  set(COMPILER_AWK_CMD "awk -F\" \" '{ printf $3}' | awk -F\"-\" '{ printf $1}'")
+else()
+  message(FATAL_ERROR "RCCL can be built only with hipcc or amdclang++")
+endif()
+
+# Set CMAKE flags
+#==================================================================================================
 set(CMAKE_INSTALL_PREFIX "${ROCM_PATH}" CACHE PATH "")
+set(CMAKE_CXX_STANDARD   17)   # We use C++17 features, this will add compile option: -std=c++17
+set(CMAKE_CXX_EXTENSIONS OFF)  # Without this line, it will add -std=gnu++17 instead, which has some issues.
+if(ROCM_PATH)
+  list(APPEND CMAKE_PREFIX_PATH  # Add ROCM_PATH to CMake search paths (for finding HIP / HSA
+              ${ROCM_PATH}
+              ${ROCM_PATH}/hip
+              ${ROCM_PATH}/llvm)
+endif()
+
+# Check for required dependencies
+#==================================================================================================
+## Check for Threads
+set(THREADS_PREFER_PTHREAD_FLAG ON)
+find_package(Threads REQUIRED)
 
-#Set the AMDGPU_TARGETS with backward compatiblity
-if(COMMAND rocm_check_target_ids)
-   rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-       TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx1030;gfx1100;gfx1101;gfx1102"
-   )
+## Check for HIP
+find_package(hip REQUIRED)
+message(STATUS "HIP compiler:     ${HIP_COMPILER}")
+message(STATUS "HIP runtime:      ${HIP_RUNTIME}")
+if(NOT "${HIP_COMPILER}" MATCHES "clang")
+  message(FATAL_ERROR "RCCL requires clang-based compiler (amdclang++ or hipcc)")
+endif()
+
+## Check for compiler version
+find_program(compiler_executable ${COMPILER_EXE_NAME})
+message(STATUS "${COMPILER_EXE_NAME} executable: ${compiler_executable}")
+execute_process(
+  COMMAND         bash "-c" "${compiler_executable} --version | grep \"${COMPILER_GREP_STRING}\" | ${COMPILER_AWK_CMD}"
+  OUTPUT_VARIABLE compiler_version_string)
+message(STATUS "${COMPILER_EXE_NAME} version:    ${compiler_version_string}")
+
+## Check for HIP version
+find_program(hipconfig_executable hipconfig)
+message(STATUS "hipconfig executable: ${hipconfig_executable}")
+execute_process(
+  COMMAND         bash "-c" "${hipconfig_executable} -v | awk -F\"-\" '{ printf $1 }'"
+  OUTPUT_VARIABLE hip_version_string)
+message(STATUS "${COMPILER_EXE_NAME} HIP version:    ${hip_version_string}")
+
+## Check for ROCm version
+set(EXPLICIT_ROCM_VERSION "" CACHE STRING "Explicit ROCM version to compile to (auto detect if empty)")
+if(EXPLICIT_ROCM_VERSION)
+  set(rocm_version_string "${EXPLICIT_ROCM_VERSION}")
+elseif(ROCM_PATH)
+  message(STATUS "Reading ROCM version from ${ROCM_PATH}/.info/version")
+  file(READ "${ROCM_PATH}/.info/version" rocm_version_string)
 else()
-    # Use target ID syntax if supported for AMDGPU_TARGETS
-    if(TARGET_ID_SUPPORT)
-        set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx1030;gfx1100;gfx1101;gfx1102")
+  message(FATAL_ERROR "Could not determine ROCM version (set EXPLICIT_ROCM_VERSION or set ROCM_PATH to a valid installation)")
+endif()
+string(REGEX MATCH "([0-9]+)\\.([0-9]+)\\.([0-9]+)" rocm_version_matches ${rocm_version_string})
+if (rocm_version_matches)
+    set(ROCM_MAJOR_VERSION ${CMAKE_MATCH_1})
+    set(ROCM_MINOR_VERSION ${CMAKE_MATCH_2})
+    set(ROCM_PATCH_VERSION ${CMAKE_MATCH_3})
+
+    message(STATUS "ROCm version: ${ROCM_MAJOR_VERSION}.${ROCM_MINOR_VERSION}.${ROCM_PATCH_VERSION}")
+
+    # Convert the version components to int for comparison
+    math(EXPR ROCM_VERSION "(10000 * ${ROCM_MAJOR_VERSION}) + (100 * ${ROCM_MINOR_VERSION}) + ${ROCM_PATCH_VERSION}")
+    add_definitions("-DROCM_VERSION=${ROCM_VERSION}")
+else()
+    message(WARNING "Failed to extract ROCm version.")
+endif()
+
+### Required for checking HIP device symbols when building with amdclang++
+set(CMAKE_REQUIRED_LIBRARIES hip::device)
+
+### Check for hipDeviceMallocUncached support
+check_symbol_exists("hipDeviceMallocUncached" "hip/hip_runtime_api.h" HIP_UNCACHED_MEMORY)
+
+### Check for hipHostMallocUncached support
+check_symbol_exists("hipHostMallocUncached" "hip/hip_runtime_api.h" HIP_HOST_UNCACHED_MEMORY)
+
+### Check for hipDeviceMallocContiguous support
+check_symbol_exists("hipDeviceMallocContiguous" "hip/hip_runtime_api.h" HIP_CONTIGUOUS_MEMORY)
+
+unset(CMAKE_REQUIRED_LIBRARIES)
+
+### Check for indirect function call support
+if(ENABLE_IFC)
+  if("${hip_version_string}" VERSION_GREATER_EQUAL "5.5.30201")
+    set(IFC_ENABLED ON)
+    message(STATUS "Indirect function call enabled")
+  else()
+    set(IFC_ENABLED OFF)
+    message(WARNING "Indirect function call disabled - requires HIP version >= 5.5.30201")
+  endif()
+else()
+  set(IFC_ENABLED OFF)
+endif()
+
+## Check for LL128 support
+if("${hip_version_string}" VERSION_GREATER_EQUAL "6.1.33591")
+  set(LL128_ENABLED ON)
+  message(STATUS "RCCL LL128 protocol enabled")
+else()
+  message(STATUS "RCCL LL128 protocol disabled - requires HIP version >= 6.1.33591")
+endif()
+
+## Check for hsa-runtime64
+find_package(hsa-runtime64 REQUIRED)
+get_target_property(HSA_INCLUDE_PATH hsa-runtime64::hsa-runtime64 INTERFACE_INCLUDE_DIRECTORIES)
+message(STATUS "HSA runtime: ${HSA_INCLUDE_PATH}")
+
+## Check for ROCM-smi
+find_package(rocm_smi PATHS ${ROCM_PATH}/lib/cmake/rocm_smi)
+if (rocm_smi_FOUND)
+  message(STATUS "Found rocm_smi at ${ROCM_SMI_INCLUDE_DIR}")
+else()
+  message(STATUS "Checking old include directory structure for rocm_smi")
+  set(ROCM_SMI_INCLUDE_DIR "${ROCM_PATH}/rocm_smi/include")
+  set(ROCM_SMI_LIB_DIR     "${ROCM_PATH}/rocm_smi/lib")
+  set(ROCM_SMI_LIBRARIES   rocm_smi64)
+endif()
+check_include_file_cxx("${ROCM_SMI_INCLUDE_DIR}/rocm_smi/rocm_smi64Config.h" HAVE_ROCM_SMI64CONFIG)
+### Check for RSMI_INIT_FLAG_THRAD_ONLY_MUTEX support
+file(READ "${ROCM_SMI_INCLUDE_DIR}/rocm_smi/rocm_smi.h" rocm_smi_incl)
+string(FIND "${rocm_smi_incl}" "RSMI_INIT_FLAG_THRAD_ONLY_MUTEX" matchres)
+if(${matchres} EQUAL -1)
+  message(STATUS "RSMI_INIT_FLAG_THRAD_ONLY_MUTEX not supported")
+else()
+  message(STATUS "RSMI_INIT_FLAG_THRAD_ONLY_MUTEX supported")
+  set(HAVE_ROCM_SMI_THREAD_ONLY_MUTEX True)
+endif ()
+
+## Check for BFD library if custom backtrace is requested
+if(BUILD_BFD)
+  enable_language(C)
+  check_include_files(bfd.h HAVE_BFD)
+  if (HAVE_BFD)
+    message(STATUS "-- Found BFD support")
+
+    ### Required for checking HIP device symbols when building with amdclang++
+    set(CMAKE_REQUIRED_LIBRARIES hip::device)
+
+    # Check for specific BFD feature support
+    CHECK_SYMBOL_EXISTS(bfd_get_section_flags "bfd.h" HAVE_DECL_BFD_GET_SECTION_FLAGS)
+    CHECK_SYMBOL_EXISTS(bfd_get_section_vma   "bfd.h" HAVE_DECL_BFD_GET_SECTION_VMA)
+    CHECK_CXX_SOURCE_COMPILES(
+      "#include <bfd.h>
+
+       int main (int argc, char **argv){
+           bfd_size_type size;
+           bfd abfd;
+           asection sec;
+           size = bfd_section_size(&abfd, &sec);
+           return (int)(size);
+       }"
+      HAVE_TWO_ARG_BFD_SECTION_SIZE)
+
+    unset(CMAKE_REQUIRED_LIBRARIES)
+
+    # Check for iberty support
+    find_library(HAVE_IBERTY iberty PATHS /usr/lib64 /usr/lib/ PATH_SUFFIXES x86_64-linux-gnu)
+    if(HAVE_IBERTY)
+      message(STATUS "iberty found @ ${HAVE_IBERTY}")
+    endif()
+
+    # Check for demangle support
+    find_path(DEMANGLE_DIR demangle.h PATHS /usr/include PATH_SUFFIXES libiberty)
+    if(NOT DEMANGLE_DIR)
+      message(WARNING "Could not find demangle.h ${DEMANGLE_DIR}")
     else()
-        set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900;gfx906;gfx908")
+      message(STATUS "Found demangle.h in ${DEMANGLE_DIR}")
     endif()
+  else()
+    message(WARNING "bfd.h header not found - Disabling custom backtrace")
+  endif()
+endif()
+
+# Check for --amdgpu-kernarg-preload-count
+check_cxx_compiler_flag("-mllvm --amdgpu-kernarg-preload-count=16" HAVE_KERNARG_PRELOAD)
+if (HAVE_KERNARG_PRELOAD)
+  message(STATUS "Kernarg preloading to SGPR enabled")
+endif()
+
+check_cxx_compiler_flag("-parallel-jobs=12" HAVE_PARALLEL_JOBS)
+if (HAVE_PARALLEL_JOBS)
+  message(STATUS "Parallel jobs enabled")
+endif()
+
+## Disable building MSCCL++ if the build environment is invalid
+## Currently MSCCL++ is supported only on gfx942 and gfx950, and only on Ubuntu and CentOS
+set(MSCCLPP_SUPPORTED_ARCHS "gfx942" "gfx942:xnack-" "gfx942:xnack+" "gfx950" "gfx950:xnack-" "gfx950:xnack+")
+
+# Check if any of the supported architectures are in GPU_TARGETS
+set(ARCH_MATCH_FOUND OFF)
+set(MSCCLPP_GPU_TARGETS "")
+foreach(ARCH IN LISTS GPU_TARGETS)
+  if(ARCH IN_LIST MSCCLPP_SUPPORTED_ARCHS)
+    set(ARCH_MATCH_FOUND ON)
+    list(APPEND MSCCLPP_GPU_TARGETS "${ARCH}")
+  endif()
+endforeach()
+set(MSCCLPP_GPU_TARGETS "${MSCCLPP_GPU_TARGETS}" CACHE STRING "GPU Targets supported by MSCCL++" FORCE)
+
+if (ENABLE_MSCCLPP AND NOT ARCH_MATCH_FOUND)
+  set(ENABLE_MSCCLPP OFF)
+  message(WARNING "Can only build MSCCL++ for supported GPU_TARGETS: ${MSCCLPP_SUPPORTED_ARCHS}; current GPU_TARGETS: ${GPU_TARGETS}; so disabling MSCCL++ build")
+endif()
+
+# MSCCL++ is only supported on ROCm 6.2.0 or newer
+if (ENABLE_MSCCLPP AND ROCM_VERSION VERSION_LESS "60200")
+  set(ENABLE_MSCCLPP OFF)
+  message(WARNING "MSCCL++ integration only supported on ROCm 6.2.0 or greater; disabling MSCCL++ build")
+endif()
+
+# cmake_host_system_information(RESULT HOST_OS_ID QUERY DISTRIB_ID) ## Requires cmake 3.22
+execute_process(
+  COMMAND bash -c "grep '^ID=' /etc/os-release | cut -d'=' -f2 | cut -d'\"' -f2"
+  OUTPUT_VARIABLE HOST_OS_ID
+  OUTPUT_STRIP_TRAILING_WHITESPACE
+)
+
+execute_process(
+  COMMAND bash -c "grep '^ID_LIKE=' /etc/os-release | cut -d'=' -f2 | cut -d'\"' -f2"
+  OUTPUT_VARIABLE HOST_OS_FAMILY
+  OUTPUT_STRIP_TRAILING_WHITESPACE
+)
+
+if (ENABLE_MSCCLPP AND NOT(${HOST_OS_ID} STREQUAL "ubuntu" OR ${HOST_OS_ID} STREQUAL "centos"))
+  set(ENABLE_MSCCLPP OFF)
+  message(WARNING "MSCCL++ integration not supported on this OS (${HOST_OS_ID}); disabling MSCCL++ build")
 endif()
-set(AMDGPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
 
-option(BUILD_TESTS "Build test programs" OFF)
-option(INSTALL_DEPENDENCIES "Force install dependencies" OFF)
-option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
-option(BUILD_ALLREDUCE_ONLY "Build AllReduce + sum + float kernel only" OFF)
-#Set the header wrapper ON by default.
-option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" ON)
+# Check for ROCTX
+if(ROCTX)
+  find_library(ROCTX_LIB NAMES roctx64)
+  if(ROCTX_LIB)
+    set(ROCTX_ENABLE ON)
+    message(STATUS "ROCTX library found: ${ROCTX_LIB}")
+  else()
+    message(WARNING "ROCTX library not found. Skipping ROCTX linking.")
+  endif()
+endif()
 
-# parse version from Makefile NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH must exist
-# NCCL_SUFFIX is optional NCCL_VERSION formatting is ((X) * 1000 + (Y) * 100 +
-# (Z)) so we must first detect one or two digits first
+# Determine version from makefiles/version.mk and fill in templates
+#==================================================================================================
+## parse version from Makefile NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH must exist
+## NCCL_SUFFIX is optional
+## NCCL_VERSION formatting is ((X) * 1000 + (Y) * 100 + (Z)) so we must first detect one or two digits first
 file(READ makefiles/version.mk version_mk_text)
 if("${version_mk_text}" MATCHES "NCCL_MAJOR *:= *([0-9]*)")
   set(NCCL_MAJOR ${CMAKE_MATCH_1})
@@ -88,295 +415,966 @@ else()
   set(NCCL_VERSION "${NCCL_MAJOR}${NCCL_MINOR}0${NCCL_PATCH}")
 endif()
 
-# Setup VERSION
+## Setup VERSION
 set(VERSION_STRING "${NCCL_MAJOR}.${NCCL_MINOR}.${NCCL_PATCH}")
 rocm_setup_version(VERSION ${VERSION_STRING})
 
-list(APPEND CMAKE_PREFIX_PATH
-            ${ROCM_PATH}
-            ${ROCM_PATH}/hip
-            ${ROCM_PATH}/llvm
-            ${ROCM_PATH}/hcc)
+## Fill in version information for main header file
+configure_file(src/nccl.h.in ${PROJECT_BINARY_DIR}/include/rccl/rccl.h) # For external linking
+configure_file(src/nccl.h.in ${PROJECT_BINARY_DIR}/include/nccl.h)      # Used by some internal files
 
-find_package(hip REQUIRED)
-message(STATUS "HIP compiler: ${HIP_COMPILER}")
-message(STATUS "HIP runtime: ${HIP_RUNTIME}")
+# Collect list of all source files
+#==================================================================================================
+# E.g: find src -type f \( -name "*.cc" -o -name "*.h" -o -name "*.hpp" \) | sort
+set(SRC_FILES
+  src/allocator.cc
+  src/bootstrap.cc
+  src/channel.cc
+  src/collectives.cc
+  src/debug.cc
+  src/enqueue.cc
+  src/group.cc
+  src/init.cc
+  src/init_nvtx.cc
+  src/mnnvl.cc
+  src/msccl.cc
+  src/proxy.cc
+  src/rccl_wrap.cc
+  src/symmetric.cc
+  src/transport.cc
+  src/device/all_gather.h
+  src/device/all_reduce.h
+  src/device/alltoall_pivot.h
+  src/device/broadcast.h
+  src/device/common.h
+  src/device/common_kernel.h
+  src/device/op128.h
+  src/device/primitives.h
+  src/device/prims_ll128.h
+  src/device/prims_ll.h
+  src/device/prims_simple.h
+  src/device/reduce.h
+  src/device/reduce_kernel.h
+  src/device/reduce_scatter.h
+  src/device/rccl_metadata.h
+  src/device/sendrecv.h
+  src/device/common.cu
+  src/device/onerank.cu
+  src/device/network/unpack/unpack_defs.h
+  src/device/network/unpack/unpack.h
+  src/device/symmetric/all_gather.cuh
+  src/device/symmetric/all_reduce.cuh
+  src/device/symmetric/kernel.cuh
+  src/device/symmetric/primitives.cuh
+  src/device/symmetric/reduce_scatter.cuh
+  src/graph/connect.cc
+  src/graph/paths.cc
+  src/graph/rings.cc
+  src/graph/rings.h
+  src/graph/rome_models.cc
+  src/graph/rome_models.h
+  src/graph/search.cc
+  src/graph/topo.cc
+  src/graph/topo.h
+  src/graph/trees.cc
+  src/graph/tuning.cc
+  src/graph/xml.cc
+  src/graph/xml.h
+  src/include/alloc.h
+  src/include/allocator.h
+  src/include/alt_rsmi.h
+  src/include/archinfo.h
+  src/include/api_trace.h
+  src/include/argcheck.h
+  src/include/BfdBacktrace.hpp
+  src/include/bitops.h
+  src/include/bootstrap.h
+  src/include/channel.h
+  src/include/checks.h
+  src/include/collectives.h
+  src/include/coll_net.h
+  src/include/comm.h
+  src/include/core.h
+  src/include/cpuset.h
+# src/include/cudawrap.h
+  src/include/debug.h
+  src/include/device.h
+  src/include/enqueue.h
+  src/include/gdrwrap.h
+  src/include/git_version.h
+  src/include/graph.h
+  src/include/group.h
+  src/include/hip_rocm_version_info.h
+  src/include/ibvcore.h
+  src/include/ibvsymbols.h
+  src/include/ibvwrap.h
+  src/include/info.h
+  src/include/ipcsocket.h
+  src/include/mnnvl.h
+  src/include/nccl_common.h
+  src/include/net_device.h
+  src/include/net.h
+  src/include/nvmlwrap.h
+  src/include/nvtx.h
+  src/include/nvtx_payload_schemas.h
+  src/include/nvtx_stub.h
+  src/include/p2p.h
+  src/include/param.h
+  src/include/profiler.h
+  src/include/proxy.h
+  src/include/ras.h
+  src/include/rccl_common.h
+  src/include/rccl_vars.h
+  src/include/register.h
+  src/include/register_inline.h
+  src/include/rccl_float8.h
+  src/include/rocm_smi_wrap.h
+  src/include/rocmwrap.h
+  src/include/roctx.h
+  src/include/recorder.h
+  src/include/shm.h
+  src/include/shmutils.h
+  src/include/signals.h
+  src/include/socket.h
+  src/include/strongstream.h
+  src/include/symmetric.h
+  src/include/timer.h
+  src/include/transport.h
+  src/include/trees.h
+  src/include/tuner.h
+  src/include/utils.h
+  src/include/mlx5/mlx5dvcore.h
+  src/include/mlx5/mlx5dvsymbols.h
+  src/include/mlx5/mlx5dvwrap.h
+  src/include/msccl/msccl_lifecycle.h
+  src/include/msccl/msccl_parser.h
+  src/include/msccl/msccl_scheduler.h
+  src/include/msccl/msccl_setup.h
+  src/include/msccl/msccl_status.h
+  src/include/msccl/msccl_struct.h
+  src/include/npkit/npkit.h
+  src/include/npkit/npkit_event.h
+  src/include/npkit/npkit_struct.h
+  src/include/nvtx3/nvToolsExt.h
+  src/include/nvtx3/nvToolsExtCounters.h
+  src/include/nvtx3/nvToolsExtCuda.h
+  src/include/nvtx3/nvToolsExtCudaRt.h
+  src/include/nvtx3/nvToolsExtMem.h
+  src/include/nvtx3/nvToolsExtMemCudaRt.h
+  src/include/nvtx3/nvToolsExtOpenCL.h
+  src/include/nvtx3/nvToolsExtPayload.h
+  src/include/nvtx3/nvToolsExtPayloadHelper.h
+  src/include/nvtx3/nvToolsExtSemanticsCounters.h
+  src/include/nvtx3/nvToolsExtSemanticsScope.h
+  src/include/nvtx3/nvToolsExtSync.h
+  src/include/nvtx3/nvtx3.hpp
+  src/include/nvtx3/nvtxDetail/nvtxExtHelperMacros.h
+  src/include/nvtx3/nvtxDetail/nvtxExtImpl.h
+  src/include/nvtx3/nvtxDetail/nvtxExtImplCounters_v1.h
+  src/include/nvtx3/nvtxDetail/nvtxExtImplMem_v1.h
+  src/include/nvtx3/nvtxDetail/nvtxExtImplMemCudaRt_v1.h
+  src/include/nvtx3/nvtxDetail/nvtxExtImplPayload_v1.h
+  src/include/nvtx3/nvtxDetail/nvtxExtInit.h
+  src/include/nvtx3/nvtxDetail/nvtxExtPayloadHelperInternal.h
+  src/include/nvtx3/nvtxDetail/nvtxExtPayloadTypeInfo.h
+  src/include/nvtx3/nvtxDetail/nvtxExtTypes.h
+  src/include/nvtx3/nvtxDetail/nvtxImpl.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCore.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCuda_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCudaRt_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxImplOpenCL_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxImplSync_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxInit.h
+  src/include/nvtx3/nvtxDetail/nvtxInitDecls.h
+  src/include/nvtx3/nvtxDetail/nvtxInitDefs.h
+  src/include/nvtx3/nvtxDetail/nvtxLinkOnce.h
+  src/include/nvtx3/nvtxDetail/nvtxTypes.h
+  src/include/proxy_trace/proxy_trace.h
+  src/include/plugin/nccl_net.h
+  src/include/plugin/nccl_profiler.h
+  src/include/plugin/nccl_tuner.h
+  src/include/plugin/plugin.h
+  src/include/plugin/net/net_v6.h
+  src/include/plugin/net/net_v7.h
+  src/include/plugin/net/net_v8.h
+  src/include/plugin/net/net_v9.h
+  src/include/plugin/net/net_v10.h
+  src/include/plugin/profiler/net_ib_v1.h
+  src/include/plugin/profiler/net_ib.h
+  src/include/plugin/profiler/net_socket_v1.h
+  src/include/plugin/profiler/net_socket.h
+  src/include/plugin/profiler/profiler_v1.h
+  src/include/plugin/profiler/profiler_v2.h
+  src/include/plugin/profiler/profiler_v3.h
+  src/include/plugin/profiler/profiler_v4.h
+  src/include/plugin/tuner/tuner_v2.h
+  src/include/plugin/tuner/tuner_v3.h
+  src/include/plugin/tuner/tuner_v4.h
+  src/misc/alt_rsmi.cc
+  src/misc/archinfo.cc
+  src/misc/argcheck.cc
+  src/misc/api_trace.c
+  src/misc/api_trace.cc
+# src/misc/cudawrap.cc
+# src/misc/gdrwrap.cc
+  src/misc/ibvsymbols.cc
+  src/misc/ibvwrap.cc
+  src/misc/ipcsocket.cc
+  src/misc/mlx5dvsymbols.cc
+  src/misc/mlx5dvwrap.cc
+  src/misc/npkit.cc
+# src/misc/nvmlwrap.cc
+  src/misc/nvmlwrap_stub.cc
+  src/misc/param.cc
+  src/misc/rocm_smi_wrap.cc
+  src/misc/rocmwrap.cc
+  src/misc/roctx.cc
+  src/misc/recorder.cc
+  src/misc/shmutils.cc
+  src/misc/signals.cc
+  src/misc/socket.cc
+  src/misc/strongstream.cc
+  src/misc/utils.cc
+  src/misc/msccl/msccl_lifecycle.cc
+  src/misc/msccl/msccl_parser.cc
+  src/misc/msccl/msccl_setup.cc
+  src/misc/msccl/msccl_status.cc
+  src/misc/proxy_trace/proxy_trace.cc
+  src/plugin/net.cc
+  src/plugin/plugin_open.cc
+  src/plugin/profiler.cc
+  src/plugin/tuner.cc
+  src/plugin/net/net_v6.cc
+  src/plugin/net/net_v7.cc
+  src/plugin/net/net_v8.cc
+  src/plugin/net/net_v9.cc
+  src/plugin/net/net_v10.cc
+  src/plugin/profiler/profiler_v1.cc
+  src/plugin/profiler/profiler_v2.cc
+  src/plugin/profiler/profiler_v3.cc
+  src/plugin/profiler/profiler_v4.cc
+  src/plugin/tuner/tuner_v2.cc
+  src/plugin/tuner/tuner_v3.cc
+  src/plugin/tuner/tuner_v4.cc
+  src/ras/client.cc
+  src/ras/client_support.cc
+  src/ras/collectives.cc
+  src/ras/peers.cc
+  src/ras/ras.cc
+  src/ras/ras_internal.h
+  src/ras/rasnet.cc
+  src/register/coll_reg.cc
+  src/register/register.cc
+  src/register/sendrecv_reg.cc
+  src/transport/coll_net.cc
+  src/transport/generic.cc
+  src/transport/net.cc
+  src/transport/net_ib.cc
+  src/transport/net_socket.cc
+  src/transport/nvls.cc
+  src/transport/p2p.cc
+  src/transport/profiler.cc
+  src/transport/shm.cc
+  src/include/latency_profiler/CollTrace.h
+  src/include/latency_profiler/CollTraceEvent.h
+  src/include/latency_profiler/CollTraceFunc.h
+  src/include/latency_profiler/CollTraceUtils.h
+  src/include/latency_profiler/EventQueue.h
+  src/misc/latency_profiler/CollTrace.cc
+  src/misc/latency_profiler/CollTraceEvent.cc
+  src/misc/latency_profiler/CollTraceFunc.cc
+  src/misc/latency_profiler/CollTraceUtils.cc
+)
 
-if(BUILD_STATIC)
-  option(BUILD_SHARED_LIBS "Build as a shared library" OFF)
-else()
-  option(BUILD_SHARED_LIBS "Build as a shared library" ON)
+if (ENABLE_MSCCL_KERNEL)
+  set(MSCCL_KERNEL_SOURCES
+    src/device/msccl_kernel_impl.h
+    src/include/msccl/msccl_kernel.h
+  )
+  list(APPEND SRC_FILES ${MSCCL_KERNEL_SOURCES})
 endif()
 
-if(BUILD_ADDRESS_SANITIZER)
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsanitize=address -shared-libasan")
-  set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address -shared-libasan")
-  add_link_options(-fuse-ld=lld)
-endif()
-
-configure_file(src/nccl.h.in ${PROJECT_BINARY_DIR}/include/rccl/rccl.h)
-configure_file(src/nccl.h.in ${PROJECT_BINARY_DIR}/include/rccl/nccl.h)
-
-include_directories(${PROJECT_BINARY_DIR}/include) # for generated rccl.h header
-include_directories(${PROJECT_BINARY_DIR}/include/rccl) # for generated rccl.h header
-include_directories(src)
-include_directories(src/include)
-include_directories(src/collectives)
-include_directories(src/collectives/device)
-
-if (BUILD_ALLREDUCE_ONLY)
-  add_definitions(-DBUILD_ALLREDUCE_ONLY)
-  set(CU_SOURCES
-      src/collectives/device/all_reduce.cu
-      src/collectives/device/sendrecv.cu
-      src/collectives/device/functions.cu)
-else()
-  set(CU_SOURCES
-      src/collectives/device/all_reduce.cu
-      src/collectives/device/all_gather.cu
-      src/collectives/device/alltoall_pivot.cu
-      src/collectives/device/reduce.cu
-      src/collectives/device/broadcast.cu
-      src/collectives/device/reduce_scatter.cu
-      src/collectives/device/sendrecv.cu
-      src/collectives/device/onerank_reduce.cu
-      src/collectives/device/functions.cu)
-endif()
-
-set(CPP_SOURCES)
-foreach(filename ${CU_SOURCES})
-  string(REPLACE ".cu"
-                 ".cpp"
-                 cpp_filename
-                 ${filename})
-  configure_file(${filename} ${cpp_filename} COPYONLY)
-  list(APPEND CPP_SOURCES ${cpp_filename})
-endforeach(filename)
-
-set(CC_SOURCES
-    src/init.cc
-    src/graph/trees.cc
-    src/graph/rings.cc
-    src/graph/paths.cc
-    src/graph/search.cc
-    src/graph/connect.cc
-    src/graph/tuning.cc
-    src/graph/topo.cc
-    src/graph/xml.cc
-    src/graph/rome_models.cc
-    src/collectives/all_reduce_api.cc
-    src/collectives/all_gather_api.cc
-    src/collectives/reduce_api.cc
-    src/collectives/broadcast_api.cc
-    src/collectives/reduce_scatter_api.cc
-    src/collectives/sendrecv_api.cc
-    src/collectives/gather_api.cc
-    src/collectives/scatter_api.cc
-    src/collectives/all_to_all_api.cc
-    src/collectives/all_to_allv_api.cc
-    src/channel.cc
-    src/misc/argcheck.cc
-    src/misc/nvmlwrap_stub.cc
-    src/misc/utils.cc
-    src/misc/ibvwrap.cc
-    src/misc/nvmlwrap_stub.cc
-    src/misc/rocm_smi_wrap.cc
-    src/misc/profiler.cc
-    src/misc/npkit.cc
-    src/misc/shmutils.cc
-    src/misc/signals.cc              # RCCL
-    src/misc/socket.cc
-    src/misc/param.cc
-    src/misc/rocmwrap.cc
-    src/misc/strongstream.cc
-    src/transport/coll_net.cc
-    src/transport/net.cc
-    src/transport/net_ib.cc
-    src/transport/net_socket.cc
-    src/transport/p2p.cc
-    src/transport/shm.cc
-    src/transport.cc
-    src/debug.cc
-    src/group.cc
-    src/bootstrap.cc
-    src/proxy.cc
-    src/net.cc
-    src/enqueue.cc
-    ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp)
-
-foreach(filename ${CC_SOURCES})
-  list(APPEND CPP_SOURCES ${filename})
-endforeach(filename)
-
-add_library(rccl ${CPP_SOURCES})
-
-# Create a custom target that creates/updates git_version.cpp
-# that executes whenever rccl is built
+if (ENABLE_MSCCLPP)
+  set(MSCCLPP_SOURCES
+    src/include/mscclpp/mscclpp_nccl.h
+    src/misc/mscclpp/mscclpp_nccl.cc
+  )
+  list(APPEND SRC_FILES ${MSCCLPP_SOURCES})
+endif()
+
+# Hipify source files (copy of source generated into hipify directory)
+#==================================================================================================
+find_program(hipify-perl_executable hipify-perl)
+if(NOT hipify-perl_executable)
+  message(FATAL_ERROR "hipify-perl not found")
+endif()
+set(HIPIFY_DIR "${CMAKE_CURRENT_BINARY_DIR}/hipify")
+
+## Loop over each source file to hipify
+foreach(SRC_FILE ${SRC_FILES})
+  # Check that file exists
+  if (NOT EXISTS ${CMAKE_SOURCE_DIR}/${SRC_FILE})
+    message(FATAL_ERROR "Unable to find file listed in CMakeLists.txt: ${CMAKE_SOURCE_DIR}/${SRC_FILE}")
+  endif()
+
+  # Establish hipified copy of the source file
+  set(HIP_FILE "${HIPIFY_DIR}/${SRC_FILE}")
+  get_filename_component(HIP_FILE_DIR ${HIP_FILE} DIRECTORY)
+
+  # Make sure the file name is unique and there is no duplicate
+  add_file_unique(HIP_SOURCES ${HIP_FILE})
+
+  # Convert .cu files to .cpp so that they get processed properly
+  string(REPLACE "\.cuh" "\.h" HIP_FILE ${HIP_FILE})
+  string(REPLACE "\.cu" "\.cu.cpp" HIP_FILE ${HIP_FILE})
+  list(APPEND HIP_SOURCES ${HIP_FILE})
+
+  # Create a custom command to create hipified source code
+  if (FAULT_INJECTION)
+    add_custom_command(
+      OUTPUT ${HIP_FILE}
+      COMMAND mkdir -p ${HIP_FILE_DIR}
+              && ${hipify-perl_executable} -quiet-warnings ${CMAKE_SOURCE_DIR}/${SRC_FILE} -o ${HIP_FILE}
+              && ${CMAKE_COMMAND} -E env bash ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/add_unroll.sh ${HIP_FILE}
+              && ${CMAKE_COMMAND} -E env bash ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/add_faults.sh ${HIP_FILE}
+      MAIN_DEPENDENCY ${SRC_FILE}
+      COMMENT "Hipifying ${SRC_FILE} -> ${HIP_FILE}"
+    )
+  else()
+    add_custom_command(
+      OUTPUT ${HIP_FILE}
+      COMMAND mkdir -p ${HIP_FILE_DIR}
+              && ${hipify-perl_executable} -quiet-warnings ${CMAKE_SOURCE_DIR}/${SRC_FILE} -o ${HIP_FILE}
+              && ${CMAKE_COMMAND} -E env bash ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/add_unroll.sh ${HIP_FILE}
+      MAIN_DEPENDENCY ${SRC_FILE}
+      COMMENT "Hipifying ${SRC_FILE} -> ${HIP_FILE}"
+    )
+  endif()
+endforeach()
+
+# Adding custom target to hipify all the source files
+# This is required to make sure that all the hipified source files are
+# available before compiling the unit tests executable(s)
+add_custom_target(hipify_all DEPENDS ${HIP_SOURCES})
+
+if (BUILD_TESTS)
+  if (ROCM_VERSION VERSION_GREATER_EQUAL "60400" AND CMAKE_BUILD_TYPE MATCHES "Debug")
+    ## Set definition for exposing rccl static function
+    add_definitions(-DRCCL_EXPOSE_STATIC)
+
+    set(HIPIFY_SRC_DIR "${PROJECT_BINARY_DIR}/hipify/src")
+    set(REPLACE_SCRIPT "${CMAKE_SOURCE_DIR}/tools/scripts/replace_static.sh")
+    message ("Replacing static functions in ${HIPIFY_SRC_DIR} with ${REPLACE_SCRIPT} for unit tests")
+    # Create a list of files which needs to be modified to remove static
+    set(TEST_NONSTATIC_SOURCE_FILES
+      ${HIPIFY_SRC_DIR}/misc/alt_rsmi.cc
+      ${HIPIFY_SRC_DIR}/register/coll_reg.cc
+      ${HIPIFY_SRC_DIR}/transport/shm.cc
+      ${HIPIFY_SRC_DIR}/transport/p2p.cc
+    )
+
+    set(EXCLUDE_STATIC_FILE "${CMAKE_SOURCE_DIR}/tools/scripts/exclude_static_list.txt")
+    # Read the exclude list file into a CMake variable
+    file(READ "${EXCLUDE_STATIC_FILE}" EXCLUDE_STATIC_CONTENTS)
+    string(REPLACE "\n" ";" EXCLUDE_STATIC_LINES "${EXCLUDE_STATIC_CONTENTS}")
+    # Create a mapping from full/relative filename to exclude list
+    unset(EXCLUDE_MAP)
+    foreach(line ${EXCLUDE_STATIC_LINES})
+      if(line MATCHES "^([a-zA-Z0-9_./-]+):([a-zA-Z0-9_,]*)")
+        set(fname "${CMAKE_MATCH_1}")
+        set(exlist "${CMAKE_MATCH_2}")
+        # Map both the basename and the full/relative path for flexibility
+        get_filename_component(basename "${fname}" NAME)
+        set(EXCLUDE_MAP_${fname} "${exlist}")
+        set(EXCLUDE_MAP_${basename} "${exlist}")
+      endif()
+    endforeach()
+
+    # Now, for each file, get the exclude list and pass to the script
+    # Create a custom command to backup the original files and remove static
+    # Always run replace script on hipified files, but preserve original backups
+    foreach(srcfile ${TEST_NONSTATIC_SOURCE_FILES})
+      # Try to match using the full/relative path first, then fallback to basename
+      set(exclude_list "")
+      if(DEFINED EXCLUDE_MAP_${srcfile})
+        set(exclude_list "${EXCLUDE_MAP_${srcfile}}")
+      else()
+        get_filename_component(basename "${srcfile}" NAME)
+        if(DEFINED EXCLUDE_MAP_${basename})
+          set(exclude_list "${EXCLUDE_MAP_${basename}}")
+        endif()
+      endif()
+      add_custom_command(
+        OUTPUT "${srcfile}.staticbak"
+        COMMAND bash -c "\
+            ${CMAKE_COMMAND} -E echo 'Processing ${srcfile} for static replacement' && \
+            if [ ! -f '${srcfile}.staticbak' ]; then \
+              ${CMAKE_COMMAND} -E copy '${srcfile}' '${srcfile}.staticbak' && \
+              ${CMAKE_COMMAND} -E echo 'Created backup: ${srcfile}.staticbak'; \
+            fi && \
+            ${CMAKE_COMMAND} -E echo 'Exposing internal functions/variables in ${srcfile}' && \
+            ${REPLACE_SCRIPT} ${srcfile} --replace-vars --exclude-list=${exclude_list} 2>&1 \
+          "
+        DEPENDS "${srcfile}" "${REPLACE_SCRIPT}" hipify_all
+        COMMENT "Removing static from ${srcfile} (backup preserved at ${srcfile}.staticbak)"
+        VERBATIM
+      )
+      list(APPEND STATIC_BAK_FILES "${srcfile}.staticbak")
+    endforeach()
+    add_custom_target(replace_static_in_hipify ALL DEPENDS ${STATIC_BAK_FILES})
+    add_dependencies(replace_static_in_hipify hipify_all)
+  else()
+    # Restore originals in the build directory if backup exists
+    foreach(srcfile ${HIPIFY_SRC_FILES})
+      if(EXISTS "${srcfile}.staticbak")
+        execute_process(
+          COMMAND ${CMAKE_COMMAND} -E copy "${srcfile}.staticbak" "${srcfile}"
+        )
+        execute_process(
+          COMMAND ${CMAKE_COMMAND} -E remove "${srcfile}.staticbak"
+        )
+      endif()
+    endforeach()
+  endif()
+endif()
+
+# Generate device/host tables and all the collective functions that are going to be in librccl.so
+#==================================================================================================
+find_package(Python3 COMPONENTS Interpreter REQUIRED)
+if (NOT Python3_FOUND)
+  message(FATAL_ERROR "RCCL requires Python3 for generating host/device tables")
+endif()
+
+set(GEN_DIR "${HIPIFY_DIR}/gensrc")
+set(GEN_SYM_DIR "${GEN_DIR}/symmetric")
+
+if(ONLY_FUNCS)
+  message(WARNING "Using ONLY_FUNCS = ${ONLY_FUNCS}. Not meant for release builds.")
+endif()
+
+# Execute the python script to generate required collective functions
+execute_process(
+    COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/src/device/generate.py ${GEN_DIR} ${IFC_ENABLED} ${COLLTRACE} ${ENABLE_MSCCL_KERNEL} ${BUILD_LOCAL_GPU_TARGET_ONLY} ${ONLY_FUNCS}
+    WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}
+    RESULT_VARIABLE gen_py_result
+    ERROR_VARIABLE gen_py_error
+)
+if (gen_py_result)
+  message(SEND_ERROR "Error: ${gen_py_error}")
+  message(FATAL_ERROR "${CMAKE_SOURCE_DIR}/src/device/generate.py failed")
+endif()
+
+if (GENERATE_SYM_KERNELS)
+  # Execute the python script to generate required symmetric memory kernels
+  execute_process(
+    COMMAND ${Python3_EXECUTABLE} ${CMAKE_SOURCE_DIR}/src/device/symmetric/generate.py ${GEN_SYM_DIR}
+    WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}
+    RESULT_VARIABLE gen_sym_py_result
+    ERROR_VARIABLE gen_sym_py_error
+  )
+  if (gen_sym_py_result)
+    message(SEND_ERROR "Error: ${gen_sym_py_error}")
+    message(FATAL_ERROR "${CMAKE_SOURCE_DIR}/src/device/symmetric/generate.py failed")
+  endif()
+endif()
+
+# Find the generated files in the output directory
+file(GLOB_RECURSE GENERATED_FILES "${GEN_DIR}/*")
+
+# Append all found generated files to the list
+foreach(file ${GENERATED_FILES})
+  list(APPEND HIP_SOURCES ${file})
+endforeach()
+
+# Create an initial git_version.cpp file (that will be updated with latest git version)
+#==================================================================================================
+file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp "")
+list(APPEND HIP_SOURCES ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp)
+
+# Create a custom target that updates git_version.cpp and executes whenever rccl is built
 add_custom_target(git_version_check
   COMMENT "Updating git_version.cpp if necessary"
-  COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/git_version.cmake
+  COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/git_version.cmake
   VERBATIM
 )
 
-# Create a dummy git_version.cpp file in case it doesn't exist
-configure_file(src/nccl.h.in ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp)
 
-# Execute git_version_check whenever rccl library is built
-add_dependencies(rccl git_version_check)
+# Set up RCCL library
+#==================================================================================================
+## Set RCCL source files
+add_library(rccl ${HIP_SOURCES})
 
-if(TRACE)
-  add_definitions(-DENABLE_TRACE)
-endif()
+## Set RCCL dependencies
+add_dependencies(rccl git_version_check)                                      # Execute git_version_check during build
 
-if(PROFILE)
-  add_definitions(-DENABLE_PROFILING)
+if (BUILD_TESTS AND ROCM_VERSION VERSION_GREATER_EQUAL "60400" AND CMAKE_BUILD_TYPE MATCHES "Debug")
+  ## Set static replacement dependency for fixture unit tests
+  add_dependencies(rccl replace_static_in_hipify)
 endif()
 
-if(NPKIT_FLAGS)
-  add_definitions(${NPKIT_FLAGS})
+## Set RCCL include directories
+target_include_directories(rccl PRIVATE ${PROJECT_BINARY_DIR}/include)        # for generated rccl.h header
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src)                    # for hipfied headers
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src/device)
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src/device/network/unpack)
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src/include)
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src/include/mlx5)
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/src/include/plugin)
+target_include_directories(rccl PRIVATE ${HIPIFY_DIR}/gensrc)
+target_include_directories(rccl PRIVATE ${HSA_INCLUDE_PATH})
+target_include_directories(rccl PRIVATE ${ROCM_SMI_INCLUDE_DIR})
+if(DEMANGLE_DIR)
+  target_include_directories(rccl PRIVATE ${DEMANGLE_DIR})
 endif()
 
-set(COLLTRACE 1 CACHE BOOL "Collective Trace Option")
+## Set RCCL compile definitions
 if(COLLTRACE)
-  add_definitions(-DENABLE_COLLTRACE)
+  target_compile_definitions(rccl PRIVATE ENABLE_COLLTRACE)
+endif()
+if(ENABLE_MSCCL_KERNEL)
+  target_compile_definitions(rccl PRIVATE COMPILE_MSCCL_KERNEL)
+endif()
+if(ENABLE_MSCCLPP)
+  target_compile_definitions(rccl PRIVATE ENABLE_MSCCLPP)
+endif()
+if(HAVE_ROCM_SMI64CONFIG)
+  target_compile_definitions(rccl PRIVATE USE_ROCM_SMI64CONFIG)
+endif()
+if(HAVE_ROCM_SMI_THREAD_ONLY_MUTEX)
+  target_compile_definitions(rccl PRIVATE USE_ROCM_SMI_THREAD_ONLY_MUTEX)
 endif()
 
-enable_language(C)
-CHECK_INCLUDE_FILES(bfd.h HAVE_BFD)
-if (HAVE_BFD)
-  add_definitions(-DHAVE_BFD)
-  message ("-- Found BFD")
-  CHECK_SYMBOL_EXISTS(bfd_get_section_flags "bfd.h" HAVE_DECL_BFD_GET_SECTION_FLAGS)
+# NPKit flags
+## May be better to move these to a separate file
+if(ENABLE_NPKIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_TIME_SYNC_GPU)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_TIME_SYNC_CPU)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_REDUCE_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_RECV_REDUCE_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_SEND_FROM_OUTPUT_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_DIRECT_SEND_FROM_OUTPUT_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_COPY_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_COPY_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_RECV_REDUCE_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_FROM_OUTPUT_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_FROM_OUTPUT_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_SIMPLE_WAIT_PEER_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_SIMPLE_WAIT_PEER_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_SIMPLE_REDUCE_OR_COPY_MULTI_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_SIMPLE_REDUCE_OR_COPY_MULTI_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL_WAIT_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL_WAIT_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL_DATA_PROCESS_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL_DATA_PROCESS_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL128_WAIT_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL128_WAIT_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL128_DATA_PROCESS_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_PRIM_LL128_DATA_PROCESS_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_TEST_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_TEST_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_NET_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_LOCAL_COPY_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_LOCAL_COPY_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_SEND_RECV_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_RECV_COPY_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_RECV_COPY_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_DIRECT_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_ALL_GATHER_RING_DIRECT_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_GENERIC_OP_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_GENERIC_OP_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_REDUCE_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_REDUCE_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RECV_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RECV_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RUN_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RUN_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RECV_REDUCE_COPY_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_RECV_REDUCE_COPY_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_INIT_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_MSCCL_INIT_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_BROADCAST_RING_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_BROADCAST_RING_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_RECV_REDUCE_SEND_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_RECV_REDUCE_SEND_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_RECV_REDUCE_COPY_ENTRY)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_EVENT_REDUCE_SCATTER_RING_RECV_REDUCE_COPY_EXIT)
+  target_compile_definitions(rccl PRIVATE ENABLE_NPKIT_PRIM_COLLECT_DATA_PROCESS_TIME)
+endif()
+
+if(PROFILE)
+  target_compile_definitions(rccl PRIVATE ENABLE_PROFILING)
+endif()
+if(ROCTX_ENABLE)
+  target_compile_definitions(rccl PRIVATE ROCTX_ENABLE)
+else()
+  target_compile_definitions(rccl PRIVATE NVTX_NO_IMPL)
+  target_compile_definitions(rccl PRIVATE NVTX_DISABLE)
+endif()
+if(TRACE)
+  target_compile_definitions(rccl PRIVATE ENABLE_TRACE)
+endif()
+if(${HIP_CONTIGUOUS_MEMORY})
+  target_compile_definitions(rccl PRIVATE HIP_CONTIGUOUS_MEMORY)
+  message(STATUS "HIP_CONTIGUOUS_MEMORY enabled")
+else()
+  message(STATUS "HIP_CONTIGUOUS_MEMORY disabled")
+endif()
+if("${hip_version_string}" VERSION_GREATER_EQUAL "5.7.31920")
+  target_compile_definitions(rccl PRIVATE HIP_UNCACHED_MEMORY)
+  message(STATUS "HIP_UNCACHED_MEMORY enabled")
+else()
+  message(STATUS "HIP_UNCACHED_MEMORY disabled - requires HIP version >= 5.7.31920")
+  # keep --hipcc-func-supp on older HIP and compiler
+  if(NOT IFC_ENABLED)
+    target_compile_options(rccl PRIVATE --hipcc-func-supp)
+    message(STATUS "--hipcc-func-supp enabled")
+  else()
+    message(STATUS "--hipcc-func-supp disabled")
+  endif()
+endif()
+if (HIP_HOST_UNCACHED_MEMORY)
+  target_compile_definitions(rccl PRIVATE HIP_HOST_UNCACHED_MEMORY)
+  message(STATUS "HIP_HOST_UNCACHED_MEMORY enabled")
+else()
+  message(STATUS "HIP_HOST_UNCACHED_MEMORY disabled")
+endif()
+if (BUILD_BFD)
+  if (HAVE_BFD)
+    target_compile_definitions(rccl PRIVATE HAVE_BFD)
+  endif()
   if (HAVE_DECL_BFD_GET_SECTION_FLAGS)
-    add_definitions(-DHAVE_DECL_BFD_GET_SECTION_FLAGS)
+    target_compile_definitions(rccl PRIVATE HAVE_DECL_BFD_GET_SECTION_FLAGS)
   endif()
-  CHECK_SYMBOL_EXISTS(bfd_get_section_vma "bfd.h" HAVE_DECL_BFD_GET_SECTION_VMA)
   if (HAVE_DECL_BFD_GET_SECTION_VMA)
-    add_definitions(-DHAVE_DECL_BFD_GET_SECTION_VMA)
+    target_compile_definitions(rccl PRIVATE HAVE_DECL_BFD_GET_SECTION_VMA)
   endif()
-  CHECK_CXX_SOURCE_COMPILES(
-    "#include <bfd.h>
-
-       int main (int argc, char **argv) {
-         bfd_size_type size;
-         bfd abfd;
-         asection sec;
-         size = bfd_section_size(&abfd, &sec);
-         return (int)(size);
-       }"
-    HAVE_TWO_ARG_BFD_SECTION_SIZE)
   if (HAVE_TWO_ARG_BFD_SECTION_SIZE)
-    add_definitions(-DHAVE_TWO_ARG_BFD_SECTION_SIZE)
-  endif()
-  find_path(DEMANGLE_HEADER demangle.h PATHS /usr/include PATH_SUFFIXES libiberty)
-  if(NOT DEMANGLE_HEADER)
-    message("Could not find demangle.h ${DEMANGLE_HEADER}")
-  else()
-    add_definitions(-DHAVE_CPLUS_DEMANGLE)
-    message("Found demangle.h in ${DEMANGLE_HEADER}")
-    set (HAVE_CPLUS_DEMANGLE  1)
-    set (HAVE_DECL_BASENAME  "1")
-    INCLUDE_DIRECTORIES(${DEMANGLE_HEADER})
+    target_compile_definitions(rccl PRIVATE HAVE_TWO_ARG_BFD_SECTION_SIZE)
   endif()
 endif()
+if (IFC_ENABLED)
+  target_compile_definitions(rccl PRIVATE USE_INDIRECT_FUNCTION_CALL)
+endif()
+if(DEMANGLE_DIR)
+  target_compile_definitions(rccl PRIVATE "HAVE_CPLUS_DEMANGLE=1")
+  target_compile_definitions(rccl PRIVATE "HAVE_DECL_BASENAME=1")
+endif()
+if(LL128_ENABLED)
+  target_compile_definitions(rccl PRIVATE ENABLE_LL128)
+endif()
 
-find_package(rocm_smi PATHS ${ROCM_PATH}/lib/cmake/rocm_smi)
-if (rocm_smi_FOUND)
-  message ("-- Found rocm_smi at ${ROCM_SMI_INCLUDE_DIR}")
-  CHECK_INCLUDE_FILE_CXX("${ROCM_SMI_INCLUDE_DIR}/rocm_smi/rocm_smi64Config.h" HAVE_ROCM_SMI64CONFIG)
+## Set RCCL compile options
+if (HAVE_PARALLEL_JOBS)
+  target_compile_options(rccl PRIVATE -parallel-jobs=12)
+endif()
+
+if (ROCM_VERSION VERSION_GREATER_EQUAL "60200")
+  target_compile_options(rccl PRIVATE --offload-compress)    # Compress GPU code at compile time.
+  target_link_libraries(rccl PRIVATE --offload-compress)     # Compress GPU code at link time.
+  message(STATUS "--offload-compress enabled - ROCm version >= 6.2.0")
 else()
-  message ("-- Checking old include directory structure for rocm_smi")
-  set(ROCM_SMI_INCLUDE_DIR "${ROCM_PATH}/rocm_smi/include")
-  set(ROCM_SMI_LIB_DIR "${ROCM_PATH}/rocm_smi/lib")
-  set(ROCM_SMI_LIBRARIES rocm_smi64)
-  CHECK_INCLUDE_FILE_CXX("${ROCM_SMI_INCLUDE_DIR}/rocm_smi/rocm_smi64Config.h" HAVE_ROCM_SMI64CONFIG)
+  message(STATUS "--offload-compress disabled - ROCm version < 6.2.0")
 endif()
-IF(HAVE_ROCM_SMI64CONFIG)
-  add_definitions(-DUSE_ROCM_SMI64CONFIG)
-ENDIF()
 
-foreach(target ${AMDGPU_TARGETS})
-  target_link_libraries(rccl PRIVATE --amdgpu-target=${target})
-endforeach()
+target_compile_options(rccl PRIVATE -Werror=uninitialized)
+target_compile_options(rccl PRIVATE -Werror=sometimes-uninitialized)
+target_compile_options(rccl PRIVATE -Wall)
+target_compile_options(rccl PRIVATE -Werror=deprecated-copy-with-user-provided-copy)
+target_compile_options(rccl PRIVATE -Wno-format-nonliteral)
+target_compile_options(rccl PRIVATE -Wno-unused-function)
+target_compile_options(rccl PRIVATE -fgpu-rdc)
 
-if("${HIP_COMPILER}" MATCHES "clang")
-  target_compile_options(rccl PRIVATE -fvisibility=hidden --hipcc-func-supp)
-  foreach(target ${AMDGPU_TARGETS})
-    target_compile_options(rccl PRIVATE -fgpu-rdc)
-  endforeach()
-  target_link_libraries(rccl PRIVATE -fgpu-rdc)
-  target_include_directories(rccl PRIVATE ${ROCM_PATH}/hsa/include)
-  find_program( hipcc_executable hipcc )
-  execute_process(COMMAND bash "-c" "${hipcc_executable} -help | grep 'parallel-jobs'" OUTPUT_VARIABLE hipcc_parallel_jobs)
-  if("${hipcc_parallel_jobs}" MATCHES "parallel-jobs")
-    target_compile_options(rccl PRIVATE -parallel-jobs=8 PRIVATE -Wno-format-nonliteral)
-    target_link_libraries(rccl PRIVATE -parallel-jobs=8)
+## Set RCCL compile and linker options for unit tests and code coverage
+if(ENABLE_CODE_COVERAGE)
+  if(NOT CMAKE_BUILD_TYPE MATCHES "Debug")
+    message(FATAL_ERROR "Code coverage is enabled, but the build type is '${CMAKE_BUILD_TYPE}'. "
+                        "Code coverage requires 'Debug' build types to expose internal symbols. "
+                        "Please set CMAKE_BUILD_TYPE to 'Debug' and reconfigure.")
   endif()
 
-  # RCCL static lib uses -fgpu-rdc which requires hipcc as the linker and archiver
-  if(BUILD_STATIC)
-    target_link_libraries(rccl PRIVATE --emit-static-lib)
-    set(CMAKE_AR "${hipcc_executable}")
-    get_property(link_libraries TARGET rccl PROPERTY LINK_LIBRARIES)
-    string (REPLACE ";" " " LINK_PROPS "${link_libraries}")
-    set(CMAKE_CXX_ARCHIVE_CREATE "<CMAKE_AR> -o <TARGET> ${LINK_PROPS} <LINK_FLAGS> <OBJECTS>")
+  message(STATUS "Code coverage is enabled with build type '${CMAKE_BUILD_TYPE}'.")
+
+  target_compile_options(rccl PRIVATE
+    -fvisibility=default -Xarch_host -fprofile-instr-generate
+    -Xarch_host -fcoverage-mapping)
+
+  set(COVERAGE_SHARED_LINKER_FLAGS
+    -fprofile-generate
+    -Wl,--enable-new-dtags,--build-id=sha1,--rpath,$ORIGIN
+  )
+
+  set(COVERAGE_EXE_LINKER_FLAGS
+    -fprofile-generate
+    -Wl,--enable-new-dtags,--build-id=sha1,--rpath,$ORIGIN/../lib
+  )
+
+  target_link_options(rccl PRIVATE ${COVERAGE_SHARED_LINKER_FLAGS})
+  target_link_options(rccl PRIVATE ${COVERAGE_EXE_LINKER_FLAGS})
+elseif(BUILD_TESTS) # Enable default/hidden visibility based on build type and ROCM_VERSION
+  if (ROCM_VERSION VERSION_GREATER_EQUAL "60400" AND CMAKE_BUILD_TYPE MATCHES "Debug")
+    target_compile_options(rccl PRIVATE -fvisibility=default)
+  else()
+    target_compile_options(rccl PRIVATE -fvisibility=hidden)
   endif()
+else()  # Enable hidden visibility for library without tests/code coverage enabled
+  target_compile_options(rccl PRIVATE -fvisibility=hidden)
 endif()
 
-if("${HIP_COMPILER}" MATCHES "hcc")
-  find_program( hcc_executable hcc )
-  execute_process(COMMAND bash "-c" "${hcc_executable} --version | sed -e '1!d' -e 's/.*based on HCC\\s*//'" OUTPUT_VARIABLE hcc_version_string)
-  execute_process(COMMAND bash "-c" "echo \"${hcc_version_string}\" | awk -F\".\" '{ printf $1}'" OUTPUT_VARIABLE hcc_major_version)
-  execute_process(COMMAND bash "-c" "echo \"${hcc_version_string}\" | awk -F\".\" '{ printf $2}'" OUTPUT_VARIABLE hcc_minor_version)
-  if ("${hcc_major_version}.${hcc_minor_version}" VERSION_LESS "4.0")
-    target_link_libraries(rccl PRIVATE -hc-function-calls)
-  endif()
+if (HAVE_KERNARG_PRELOAD)
+  target_compile_options(rccl PRIVATE -mllvm --amdgpu-kernarg-preload-count=16)
 endif()
 
-target_include_directories(rccl PRIVATE ${ROCM_SMI_INCLUDE_DIR})
-target_link_libraries(rccl PRIVATE hip::device dl -l${ROCM_SMI_LIBRARIES} -L${ROCM_SMI_LIB_DIR})
-target_link_libraries(rccl INTERFACE hip::host)
+if (DUMP_ASM) # Save temporary files from kernel compilation
+  message(STATUS "Disassembling librccl.so to asm")
+  # Maintain symbols but without changing code.  Keep additional data in dwarf section of binary.
+  target_compile_options(rccl PRIVATE -gline-tables-only)
+  set(OBJ_DUMP ${ROCM_PATH}/llvm/bin/llvm-objdump)
+
+  add_custom_command(TARGET rccl POST_BUILD
+    COMMENT "Disassembling RCCL library"
+    COMMAND /bin/bash -c "${OBJ_DUMP} --offload-fatbin librccl.so"
+    VERBATIM
+  )
+  foreach(GPUARCH ${GPU_TARGETS})
+    add_custom_command(TARGET rccl POST_BUILD
+        COMMENT "Disassembling RCCL library to dump assembly for ${GPUARCH}"
+        COMMAND /bin/bash -c "${OBJ_DUMP} -d -l --source --symbolize-operands librccl.so.0.hipv4-amdgcn-amd-amdhsa--${GPUARCH} > librccl.${GPUARCH}.s"
+        VERBATIM
+    )
+  endforeach()
+endif()
+
+## NOTE: This is currently being handled by rocm-cmake, however may need to be re-enabled in the future
+#foreach(target ${GPU_TARGETS})
+#  target_compile_options(rccl PRIVATE --offload-arch=${target})
+#endforeach()
+
+if(BUILD_ADDRESS_SANITIZER)
+  target_compile_options(rccl PRIVATE -fsanitize=address -shared-libasan)
+endif()
+if(TIMETRACE)
+  target_compile_options(rccl PRIVATE -ftime-trace)
+endif()
+if (FAULT_INJECTION)
+  target_compile_definitions(rccl PRIVATE ENABLE_FAULT_INJECTION)
+  message(STATUS "Fault injection enabled")
+endif()
+
+## Set RCCL linked library directories
+target_link_directories(rccl PRIVATE ${ROCM_SMI_LIB_DIR})
 
-if(HAVE_BFD)
-  target_link_libraries(rccl PRIVATE bfd dl z)
-  find_library(HAVE_IBERTY iberty PATHS /usr/lib64 /usr/lib/
-    PATH_SUFFIXES x86_64-linux-gnu)
+if (ROCM_VERSION VERSION_GREATER_EQUAL "60100")
+    option(RCCL_ROCPROFILER_REGISTER "Enable rocprofiler-register support" ON)
+else()
+    if(RCCL_ROCPROFILER_REGISTER)
+        message(AUTHOR_WARNING "RCCL_ROCPROFILER_REGISTER is not valid option for ROCm < 6.2. Current ROCm version: ${ROCM_VERSION}")
+    endif()
+    set(RCCL_ROCPROFILER_REGISTER OFF CACHE BOOL "" FORCE)
+endif()
+if(RCCL_ROCPROFILER_REGISTER)
+  find_package(rocprofiler-register REQUIRED)
+  target_compile_definitions(rccl PRIVATE RCCL_ROCPROFILER_REGISTER=1)
+  target_link_libraries(
+      rccl PRIVATE rocprofiler-register::rocprofiler-register)
+endif()
+
+## Set RCCL linked libraries
+if (HAVE_BFD)
+  target_link_libraries(rccl PRIVATE bfd)
   if(HAVE_IBERTY)
-    message("iberty found @  ${HAVE_IBERTY} ")
-    target_link_libraries(rccl PRIVATE iberty dl z)
+    target_link_libraries(rccl PRIVATE iberty z)
+  endif()
+endif()
+if (ROCTX_ENABLE)
+  target_link_libraries(rccl PRIVATE -lroctx64)
+endif()
+target_link_libraries(rccl PRIVATE   -fgpu-rdc)             # Required when linking relocatable device code
+target_link_libraries(rccl PRIVATE   Threads::Threads)
+target_link_libraries(rccl INTERFACE hip::host)
+target_link_libraries(rccl PRIVATE   hip::device)
+target_link_libraries(rccl PRIVATE   dl)
+target_link_libraries(rccl PRIVATE   ${ROCM_SMI_LIBRARIES})
+target_link_libraries(rccl PRIVATE fmt::fmt-header-only)
+if(ENABLE_MSCCLPP)
+  target_link_libraries(rccl PRIVATE mscclpp_nccl)
+endif()
+
+## Set RCCL link options
+## Find out available memory
+execute_process(
+  COMMAND bash "-c" "cat /sys/fs/cgroup/memory.max"
+  OUTPUT_VARIABLE memory_max_string)
+if (${memory_max_string} MATCHES "^[0-9]+")
+  math(EXPR memory_in_gb "${memory_max_string} / (1024 * 1024 * 1024)")
+else()
+  execute_process(
+    COMMAND bash "-c" "free | grep -o '[[:digit:]]*' | head -1"
+    OUTPUT_VARIABLE memory_max_string)
+  ## memory_max_string holds the free memory in KB
+  if (${memory_max_string} MATCHES "^[0-9]+")
+    math(EXPR memory_in_gb "${memory_max_string} / (1024 * 1024)") ## KB to GB conversion
+  else()
+    cmake_host_system_information(RESULT memory_max_string QUERY AVAILABLE_PHYSICAL_MEMORY )
+    math(EXPR memory_in_gb "${memory_max_string} / 1024")
   endif()
 endif()
+## Reserve 16GB for each linker job. Limit max number of linker jobs to 16
+if (HAVE_PARALLEL_JOBS)
+  math(EXPR num_linker_jobs "(${memory_in_gb} + 15) / 16")
+  if (${num_linker_jobs} GREATER_EQUAL "16")
+    set(num_linker_jobs "16")
+  endif()
+  message(STATUS "Use ${num_linker_jobs} jobs for linking")
+  target_link_options(rccl PRIVATE -parallel-jobs=${num_linker_jobs})       # Use multiple threads to link
+endif()
+if(BUILD_ADDRESS_SANITIZER)
+  target_link_options(rccl PRIVATE -fuse-ld=lld)
+endif()
+if(TIMETRACE)
+  target_link_options(rccl PRIVATE -ftime-trace)
+endif()
+
+if(NOT BUILD_SHARED_LIBS)
+  message(STATUS "Building static RCCL library")
+else()
+  message(STATUS "Building shared RCCL library")
+endif()
+if (HAVE_KERNARG_PRELOAD)
+  target_link_options(rccl PRIVATE "SHELL:-Xoffload-linker -mllvm=-amdgpu-kernarg-preload-count=16")
+endif()
+
+if(ENABLE_MSCCLPP)
+  include(cmake/MSCCLPP.cmake)
+endif()
 
-#Setup librccl.so version
+## Track linking time
+set_property(TARGET rccl PROPERTY RULE_LAUNCH_LINK "${CMAKE_COMMAND} -E time")
+
+## Setup librccl.so version
 rocm_set_soversion(rccl "1.0")
 
-rocm_install_targets(TARGETS
-                     rccl
-		     )
-rocm_install(FILES ${PROJECT_BINARY_DIR}/include/rccl/rccl.h src/include/nccl_net.h
-        DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rccl)
-
-rocm_export_targets(NAMESPACE
-                    roc::
-                    TARGETS
-                    rccl
-                    DEPENDS
-                    hip)
-if(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY)
-  #Create wrapper files
-  rocm_wrap_header_dir( "${PROJECT_BINARY_DIR}/include/rccl"
-            PATTERNS "rccl.h"
-            GUARDS SYMLINK WRAPPER
-            WRAPPER_LOCATIONS ${CMAKE_INSTALL_INCLUDEDIR} rccl/${CMAKE_INSTALL_INCLUDEDIR})
-  #install the wrapper header file to package
-  rocm_install( FILES ${PROJECT_BINARY_DIR}/rccl/include/rccl.h src/include/nccl_net.h
-	        DESTINATION "./rccl/${CMAKE_INSTALL_INCLUDEDIR}/" )
-  rocm_install( FILES ${PROJECT_BINARY_DIR}/include/rccl.h src/include/nccl_net.h
-                DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}/" )
-endif()
-
-rocm_package_add_dependencies(DEPENDS "hip-rocclr >= 3.5.0" "rocm-smi-lib >= 4.0.0")
+if(NOT BUILD_SHARED_LIBS)
+  # To create a static lib with `-fgpu-rdc`, you need `--emit-static-lib` and `--hip-link`.
+  # You also need to invoke amdclang++ again to trigger GPU code generation.
+  set(static_link_flags
+    ${CXXFLAGS}
+    --hip-link
+    -fgpu-rdc
+    --emit-static-lib
+  )
+
+  # Find all the libraries we need to link at link time to include them in the clang link
+  # command line.
+  get_target_property(rccl_libs rccl LINK_LIBRARIES)
+  foreach(target ${rccl_libs})
+    if(TARGET ${target})
+      get_target_property(location ${target} LOCATION)
+      if(location)
+        LIST(APPEND static_link_flags -l${location})
+      endif()
+    endif()
+  endforeach()
+
+  foreach(target ${GPU_TARGETS})
+    list(APPEND static_link_flags --offload-arch=${target})
+  endforeach()
+  list(JOIN static_link_flags " " flags_str)
+
+  # Invoking amdclang++ this way will produce a static archive, so just override ARCHIVE_CREATE.
+  set(CMAKE_CXX_ARCHIVE_CREATE "<CMAKE_CXX_COMPILER> ${flags_str} -o <TARGET> <OBJECTS>")
+endif()
+
+# Install settings
+#==================================================================================================
+## Specify install targets
+rocm_install_targets(TARGETS rccl)
+rocm_install(FILES       ${PROJECT_BINARY_DIR}/include/rccl/rccl.h src/include/plugin/nccl_net.h
+             DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rccl)
+rocm_install(FILES       src/include/api_trace.h
+             DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rccl/amd_detail)
+file(COPY tools/msccl-algorithms DESTINATION ${PROJECT_BINARY_DIR})
+file(COPY tools/msccl-unit-test-algorithms DESTINATION ${PROJECT_BINARY_DIR})
+## Install Algorithm files under share folder
+rocm_install(DIRECTORY ${PROJECT_BINARY_DIR}/msccl-algorithms DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
+rocm_install(DIRECTORY ${PROJECT_BINARY_DIR}/msccl-unit-test-algorithms DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
+
+rocm_export_targets(
+  NAMESPACE roc::
+  TARGETS   rccl
+  DEPENDS   hip)
+
+## Set package dependencies
+if(BUILD_ADDRESS_SANITIZER)
+  set(DEPENDS_HIP_RUNTIME "hip-runtime-amd-asan" )
+else()
+  set(DEPENDS_HIP_RUNTIME "hip-runtime-amd" )
+endif()
+rocm_package_add_dependencies(DEPENDS "${DEPENDS_HIP_RUNTIME} >= 4.5.0" "rocm-smi-lib >= 4.0.0")
+set(CPACK_DEB_COMPONENT_INSTALL ON)
 set(CPACK_DEBIAN_PACKAGE_SHLIBDEPS ON)
+set(CPACK_RPM_COMPONENT_INSTALL ON)
 set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt" "${ROCM_PATH}")
 
 find_file (DEBIAN debian_version debconf.conf PATHS /etc)
@@ -385,13 +1383,14 @@ if(DEBIAN)
   file(WRITE "${CMAKE_BINARY_DIR}/copyright"
   "Format: https://www.debian.org/doc/packaging-manuals/copyright-format/1.0/
 Upstream-Name: rccl
-Source: https://github.com/ROCmSoftwarePlatform/rccl
+Source: https://github.com/ROCm/rccl
 
 Files: *
 Copyright: (c) 2016-2020, NVIDIA CORPORATION. All rights reserved.
-Modifications Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+Modifications Copyright (c) 2020-2023 Advanced Micro Devices, Inc. All rights reserved.
+Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
 License: See LICENSE.txt for license information\n")
-  install(FILES "${CMAKE_BINARY_DIR}/copyright" DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
+  rocm_install(FILES "${CMAKE_BINARY_DIR}/copyright" DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
   # Write changelog file
   find_program( date_executable date )
   execute_process(COMMAND ${date_executable} -R OUTPUT_VARIABLE TIMESTAMP)
@@ -402,24 +1401,32 @@ License: See LICENSE.txt for license information\n")
 
  -- RCCL Maintainer <[e-mail]>  ${TIMESTAMP}\n")
   find_program( gzip_executable gzip )
-  execute_process(COMMAND bash "-c" "${gzip_executable} -9 -c ${CMAKE_BINARY_DIR}/changelog"
+  execute_process(COMMAND bash "-c" "${gzip_executable} -9 -c -n ${CMAKE_BINARY_DIR}/changelog"
     WORKING_DIRECTORY ${CMAKE_BINARY_DIR} OUTPUT_FILE "${CMAKE_BINARY_DIR}/changelog.Debian.gz")
-  install(FILES "${CMAKE_BINARY_DIR}/changelog.Debian.gz" DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
+  rocm_install(FILES "${CMAKE_BINARY_DIR}/changelog.Debian.gz" DESTINATION ${CMAKE_INSTALL_DATADIR}/rccl)
   set(CPACK_DEBIAN_PACKAGE_DESCRIPTION "ROCm Communication Collectives Library
   Optimized primitives for collective multi-GPU communication")
 endif()
 
+## Building RCCL RAS
+include(cmake/rcclRAS.cmake)
+
 if(BUILD_TESTS)
   rocm_package_setup_component(clients)
-  rocm_package_setup_client_component(tests)
+  rocm_package_setup_client_component(tests PACKAGE_NAME unittests)
   add_subdirectory(test)
+
+  if(BUILD_SHARED_LIBS)
+    add_custom_command(TARGET rccl POST_BUILD
+      COMMENT "Extracting metadata from librccl.so"
+      COMMAND COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/extract_metadata.cmake
+      VERBATIM
+    )
+  endif()
 endif()
 
 rocm_create_package(
-  NAME
-  rccl
-  DESCRIPTION
-  "ROCm Communication Collectives Library"
-  MAINTAINER
-  "RCCL Maintainer <[e-mail]>"
+  NAME        rccl
+  DESCRIPTION "ROCm Communication Collectives Library"
+  MAINTAINER  "RCCL Maintainer <[e-mail]>"
   LDCONFIG)
diff --git a/LICENSE.txt b/LICENSE.txt
index 11b343f..da5cc8f 100644
--- a/LICENSE.txt
+++ b/LICENSE.txt
@@ -4,41 +4,43 @@ Attributions
 Contains contributions from NVIDIA.
 
 Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
-Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
-
- Redistribution and use in source and binary forms, with or without
- modification, are permitted provided that the following conditions
- are met:
-  * Redistributions of source code must retain the above copyright
-    notice, this list of conditions and the following disclaimer.
-  * Redistributions in binary form must reproduce the above copyright
-    notice, this list of conditions and the following disclaimer in the
-    documentation and/or other materials provided with the distribution.
-  * Neither the name of NVIDIA CORPORATION, Lawrence Berkeley National
-    Laboratory, the U.S. Department of Energy, nor the names of their
-    contributors may be used to endorse or promote products derived
-    from this software without specific prior written permission.
-
- THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
- EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
- CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
- The U.S. Department of Energy funded the development of this software
- under subcontract 7078610 with Lawrence Berkeley National Laboratory.
+Modifications Copyright (c) 2019-2025 Advanced Micro Devices, Inc. All rights reserved.
+Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+*  Redistributions of source code must retain the above copyright
+   notice, this list of conditions and the following disclaimer.
+*  Redistributions in binary form must reproduce the above copyright
+   notice, this list of conditions and the following disclaimer in the
+   documentation and/or other materials provided with the distribution.
+*  Neither the name of NVIDIA CORPORATION, Lawrence Berkeley National
+   Laboratory, the U.S. Department of Energy, nor the names of their
+   contributors may be used to endorse or promote products derived
+   from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ''AS IS'' AND ANY
+EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+The U.S. Department of Energy funded the development of this software
+under subcontract 7078610 with Lawrence Berkeley National Laboratory.
 
 
 This code also includes files from the NVIDIA Tools Extension SDK project.
 
 See:
 
-   https://github.com/NVIDIA/NVTX
+https://github.com/NVIDIA/NVTX
 
 for more information and license details.
diff --git a/NOTICES.txt b/NOTICES.txt
index 75794b0..372651b 100644
--- a/NOTICES.txt
+++ b/NOTICES.txt
@@ -1,66 +1,128 @@
 Notices and Licenses file
 _______________________________________________________________
 
-Dependencies on nvidia-nccl v2.3.7-1 (BSD3)
-Copyright (c) 2015-2018, NVIDIA CORPORATION.
-Modifications Copyright (c) 2019-2020 Advanced Micro Devices, Inc.
-
-Redistribution and use in source and binary forms, with or without
-modification, are permitted provided that the following conditions
-are met:
- * Redistributions of source code must retain the above copyright
-   notice, this list of conditions and the following disclaimer.
- * Redistributions in binary form must reproduce the above copyright
-   notice, this list of conditions and the following disclaimer in the
-   documentation and/or other materials provided with the distribution.
- * Neither the name of NVIDIA CORPORATION, Lawrence Berkeley National
-   Laboratory, the U.S. Department of Energy, nor the names of their
-   contributors may be used to endorse or promote products derived
-   from this software without specific prior written permission.
-
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
-EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
-PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
-CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
-EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
-PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
-PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
-OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-The U.S. Department of Energy funded the development of this software
-under subcontract 7078610 with Lawrence Berkeley National Laboratory.
-
-
-nvidia-nccl v2.3.7-1 (BSD2)
-Copyright (c) 2015-2018, NVIDIA CORPORATION. All rights reserved.
-
-Redistribution and use in source and binary forms, with or without
-modification, are permitted provided that the following conditions
-are met:
- * Redistributions of source code must retain the above copyright
-   notice, this list of conditions and the following disclaimer.
- * Redistributions in binary form must reproduce the above copyright
-   notice, this list of conditions and the following disclaimer in the
-   documentation and/or other materials provided with the distribution.
- * Neither the name of NVIDIA CORPORATION, Lawrence Berkeley National
-   Laboratory, the U.S. Department of Energy, nor the names of their
-   contributors may be used to endorse or promote products derived
-   from this software without specific prior written permission.
-
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
-EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
-PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
-CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
-EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
-PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
-PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
-OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-The U.S. Department of Energy funded the development of this software
-under subcontract 7078610 with Lawrence Berkeley National Laboratory.
+Dependencies on nvidia-nccl v2.27.3-1 (BSD3)
+
+Copyright (c) 2015-2025, NVIDIA CORPORATION. All rights reserved.
+Modifications Copyright (c) 2019-2024 Advanced Micro Devices, Inc. All rights reserved.
+Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
+
+ Redistribution and use in source and binary forms, with or without
+ modification, are permitted provided that the following conditions
+ are met:
+  * Redistributions of source code must retain the above copyright
+    notice, this list of conditions and the following disclaimer.
+  * Redistributions in binary form must reproduce the above copyright
+    notice, this list of conditions and the following disclaimer in the
+    documentation and/or other materials provided with the distribution.
+  * Neither the name of NVIDIA CORPORATION, Lawrence Berkeley National
+    Laboratory, the U.S. Department of Energy, nor the names of their
+    contributors may be used to endorse or promote products derived
+    from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ The U.S. Department of Energy funded the development of this software
+ under subcontract 7078610 with Lawrence Berkeley National Laboratory.
+
+
+This code also includes files from the NVIDIA Tools Extension SDK project.
+
+See:
+
+   https://github.com/NVIDIA/NVTX
+
+for more information and license details.
+
+_______________________________________________________________
+
+Dependencies on NPKit (MIT License)
+
+    Copyright (c) Microsoft Corporation.
+
+    Permission is hereby granted, free of charge, to any person obtaining a copy
+    of this software and associated documentation files (the "Software"), to deal
+    in the Software without restriction, including without limitation the rights
+    to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+    copies of the Software, and to permit persons to whom the Software is
+    furnished to do so, subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be included in all
+    copies or substantial portions of the Software.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+    IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+    FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+    AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+    LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+    OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+    SOFTWARE
+
+_______________________________________________________________
+
+Dependencies on MSCCL++ (MIT License)
+
+    Copyright (c) Microsoft Corporation.
+
+    Permission is hereby granted, free of charge, to any person obtaining a copy
+    of this software and associated documentation files (the "Software"), to deal
+    in the Software without restriction, including without limitation the rights
+    to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+    copies of the Software, and to permit persons to whom the Software is
+    furnished to do so, subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be included in all
+    copies or substantial portions of the Software.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+    IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+    FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+    AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+    LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+    OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+    SOFTWARE
+
+See:
+
+    https://github.com/microsoft/mscclpp
+
+for more information and license details.
+
+_______________________________________________________________
+
+Dependencies on Latency Profiler (MIT License)
+
+    Copyright (c) Meta Platforms, Inc. and affiliates.
+
+    Permission is hereby granted, free of charge, to any person obtaining a copy
+    of this software and associated documentation files (the "Software"), to deal
+    in the Software without restriction, including without limitation the rights
+    to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+    copies of the Software, and to permit persons to whom the Software is
+    furnished to do so, subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be included in all
+    copies or substantial portions of the Software.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+    IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+    FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+    AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+    LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+    OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+    SOFTWARE.
+
+See:
+
+    src/include/latency_profiler
+    src/misc/latency_profiler
diff --git a/README.md b/README.md
index 8bdce66..d25f661 100644
--- a/README.md
+++ b/README.md
@@ -2,6 +2,11 @@
 
 ROCm Communication Collectives Library
 
+[![RCCL](https://dev.azure.com/ROCm-CI/ROCm-CI/_apis/build/status%2Frccl?repoName=ROCm%2Frccl&branchName=develop)](https://dev.azure.com/ROCm-CI/ROCm-CI/_build/latest?definitionId=107&repoName=ROCm%2Frccl&branchName=develop)
+[![TheRock CI](https://github.com/ROCm/rccl/actions/workflows/therock-ci.yml/badge.svg?branch=develop&event=push)](https://github.com/ROCm/rccl/actions/workflows/therock-ci.yml)
+
+> **Note:** The published documentation is available at [RCCL](https://rocm.docs.amd.com/projects/rccl/en/latest/index.html) in an organized easy-to-read format that includes a table of contents and search functionality. The documentation source files reside in the [rccl/docs](https://github.com/ROCm/rccl/tree/develop/docs) folder in this repository. As with all ROCm projects, the documentation is open source. For more information, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
+
 ## Introduction
 
 RCCL (pronounced "Rickle") is a stand-alone library of standard collective communication routines for GPUs, implementing all-reduce, all-gather, reduce, broadcast, reduce-scatter, gather, scatter, and all-to-all. There is also initial support for direct GPU-to-GPU send and receive operations.  It has been optimized to achieve high bandwidth on platforms using PCIe, xGMI as well as networking using InfiniBand Verbs or TCP/IP sockets. RCCL supports an arbitrary number of GPUs installed in a single node or multiple nodes, and can be used in either single- or multi-process (e.g., MPI) applications.
@@ -11,44 +16,81 @@ The collective operations are implemented using ring and tree algorithms and hav
 ## Requirements
 
 1. ROCm supported GPUs
-2. ROCm stack installed on the system (HIP runtime & HCC or HIP-Clang)
+2. ROCm stack installed on the system (HIP runtime & HIP-Clang)
 
 ## Quickstart RCCL Build
 
-RCCL directly depends on HIP runtime, plus the HCC C++ compiler or the HIP-Clang compiler which are part of the ROCm software stack.
-For ROCm installation instructions, see https://github.com/RadeonOpenCompute/ROCm.
+RCCL directly depends on HIP runtime plus the HIP-Clang compiler, which are part of the ROCm software stack.
+For ROCm installation instructions, see https://github.com/ROCm/ROCm.
+
+The root of this repository has a helper script `install.sh` to build and install RCCL with a single command. It hard-codes configurations that can be specified through invoking cmake directly, but it's a great way to get started quickly and can serve as an example of how to build/install RCCL.
+
+### To build the library using the install script:
+
+```shell
+./install.sh
+```
 
-The root of this repository has a helper script 'install.sh' to build and install RCCL on Ubuntu with a single command.  It does not take a lot of options and hard-codes configuration that can be specified through invoking cmake directly, but it's a great way to get started quickly and can serve as an example of how to build/install.
+For more info on build options/flags when using the install script, use `./install.sh --help`
+```shell
+./install.sh --help
+RCCL build & installation helper script
+ Options:
+       --address-sanitizer     Build with address sanitizer enabled
+    -c|--enable-code-coverage  Enable code coverage
+    -d|--dependencies          Install RCCL dependencies
+       --debug                 Build debug library
+       --enable_backtrace      Build with custom backtrace support
+       --disable-colltrace     Build without collective trace
+       --disable-msccl-kernel  Build without MSCCL kernels
+       --enable-mscclpp        Build with MSCCL++ support
+       --enable-mscclpp-clip   Build MSCCL++ with clip wrapper on bfloat16 and half addition routines
+       --disable-roctx         Build without ROCTX logging
+    -f|--fast                  Quick-build RCCL (local gpu arch only, no backtrace, and collective trace support)
+    -h|--help                  Prints this help message
+    -i|--install               Install RCCL library (see --prefix argument below)
+    -j|--jobs                  Specify how many parallel compilation jobs to run ($nproc by default)
+    -l|--local_gpu_only        Only compile for local GPU architecture
+       --amdgpu_targets        Only compile for specified GPU architecture(s). For multiple targets, separate by ';' (builds for all supported GPU architectures by default)
+       --no_clean              Don't delete files if they already exist
+       --npkit-enable          Compile with npkit enabled
+       --log-trace             Build with log trace enabled (i.e. NCCL_DEBUG=TRACE)
+       --openmp-test-enable    Enable OpenMP in rccl unit tests
+    -p|--package_build         Build RCCL package
+       --prefix                Specify custom directory to install RCCL to (default: `/opt/rocm`)
+       --run_tests_all         Run all rccl unit tests (must be built already)
+    -r|--run_tests_quick       Run small subset of rccl unit tests (must be built already)
+       --static                Build RCCL as a static library instead of shared library
+    -t|--tests_build           Build rccl unit tests, but do not run
+       --time-trace            Plot the build time of RCCL (requires `ninja-build` package installed on the system)
+       --verbose               Show compile commands
+```
 
-*  `./install.sh` -- builds library including unit tests
-*  `./install.sh -i` -- builds and installs the library to /opt/rocm/rccl; installation path can be changed with --prefix argument (see below.)
-*  `./install.sh -d` -- installs all necessary dependencies for RCCL.  Should be re-invoked if the build folder is removed.
-*  `./install.sh -h` -- shows help
-*  `./install.sh -t` -- builds library including unit tests
-*  `./install.sh -r` -- runs unit tests (must be already built)
-*  `./install.sh -p` -- builds RCCL package
-*  `./install.sh -s` -- builds RCCL as a static library (default: shared)
-*  `./install.sh -hcc` -- builds RCCL with hcc compiler; note that hcc is now deprecated. (default:hip-clang)
-*  `./install.sh --prefix` -- specify custom path to install RCCL to (default:/opt/rocm)
+By default, RCCL builds for all GPU targets defined in `DEFAULT_GPUS` in `CMakeLists.txt`. To target specific GPU(s), and potentially reduce build time, use `--amdgpu_targets` as a `;` separated string listing GPU(s) to target.
 
 ## Manual build
-#### To build the library :
+
+### To build the library using CMake:
 
 ```shell
-$ git clone https://github.com/ROCmSoftwarePlatform/rccl.git
+$ git clone --recursive https://github.com/ROCm/rccl.git
 $ cd rccl
 $ mkdir build
 $ cd build
-$ CXX=/opt/rocm/bin/hipcc cmake -DCMAKE_PREFIX_PATH=/opt/rocm/ ..
-$ make -j
+$ cmake ..
+$ make -j 16      # Or some other suitable number of parallel jobs
 ```
-You may substitute an installation path of your own choosing by passing CMAKE_INSTALL_PREFIX. For example:
+If you have already cloned, you can checkout the external submodules manually.
 ```shell
-$ CXX=/opt/rocm/bin/hipcc cmake -DCMAKE_PREFIX_PATH=/opt/rocm/ -DCMAKE_INSTALL_PREFIX=$PWD/rccl-install ..
+$ git submodule update --init --recursive --depth=1
+```
+You may substitute an installation path of your own choosing by passing `CMAKE_INSTALL_PREFIX`. For example:
+```shell
+$ cmake -DCMAKE_INSTALL_PREFIX=$PWD/rccl-install -DCMAKE_BUILD_TYPE=Release ..
 ```
 Note: ensure rocm-cmake is installed, `apt install rocm-cmake`.
 
-#### To build the RCCL package and install package :
+### To build the RCCL package and install package :
 
 Assuming you have already cloned this repository and built the library as shown in the previous section:
 
@@ -58,49 +100,48 @@ $ make package
 $ sudo dpkg -i *.deb
 ```
 
-RCCL package install requires sudo/root access because it creates a directory called "rccl" under /opt/rocm/. This is an optional step and RCCL can be used directly by including the path containing librccl.so.
+RCCL package install requires sudo/root access because it installs under `/opt/rocm/`. This is an optional step as RCCL can instead be used directly by including the path containing `librccl.so`.
+
+## Docker build
 
-## Enabling peer-to-peer transport
-In order to enable peer-to-peer access on machines with PCIe-connected GPUs, the HSA environment variable HSA_FORCE_FINE_GRAIN_PCIE=1 is required to be set, on top of requiring GPUs that support peer-to-peer access and proper large BAR addressing support.
+Refer to [docker/README.md](docker/README.md "docker/README.md")
 
 ## Tests
 
-There are unit tests implemented with the Googletest framework in RCCL.  The unit tests require Googletest 1.10 or higher to build and execute properly (installed with the -d option to install.sh).
-To invoke the unit tests, go to the build folder, then the test subfolder, and execute the appropriate unit test executable(s).
+There are rccl unit tests implemented with the Googletest framework in RCCL.  The rccl unit tests require Googletest 1.10 or higher to build and execute properly (installed with the -d option to install.sh).
+To invoke the rccl unit tests, go to the build folder, then the test subfolder, and execute the appropriate rccl unit test executable(s).
 
-Unit test names are now of the format:
+rccl unit test names are now of the format:
 
     CollectiveCall.[Type of test]
 
-Filtering of unit tests should be done with environment variable and by passing the --gtest_filter command line flag, for example:
+Filtering of rccl unit tests should be done with environment variable and by passing the `--gtest_filter` command line flag, for example:
 
 ```shell
-UT_DATATYPES=ncclBfloat16 UT_REDOPS=prod ./UnitTests --gtest_filter="AllReduce.C*"
+UT_DATATYPES=ncclBfloat16 UT_REDOPS=prod ./rccl-UnitTests --gtest_filter="AllReduce.C*"
 ```
-will run only AllReduce correctness tests with float16 datatype. A list of available filtering environment variables appears at the top of every run. See "Running a Subset of the Tests" at https://chromium.googlesource.com/external/github.com/google/googletest/+/HEAD/googletest/docs/advanced.md for more information on how to form more advanced filters.
 
+will run only AllReduce correctness tests with float16 datatype. A list of available filtering environment variables appears at the top of every run. See "Running a Subset of the Tests" at https://google.github.io/googletest/advanced.html#running-a-subset-of-the-tests for more information on how to form more advanced filters.
 
-There are also other performance and error-checking tests for RCCL.  These are maintained separately at https://github.com/ROCmSoftwarePlatform/rccl-tests.
+There are also other performance and error-checking tests for RCCL.  These are maintained separately at https://github.com/ROCm/rccl-tests.
 See the rccl-tests README for more information on how to build and run those tests.
 
-## NPKit
-
-RCCL integrates [NPKit](https://github.com/microsoft/npkit), a profiler framework that enables collecting fine-grained trace events in RCCL components, especially in giant collective GPU kernels.
-
-Please check [NPKit sample workflow for RCCL](https://github.com/microsoft/NPKit/tree/main/rccl_samples) as a fully automated usage example. It also provides good templates for the following manual instructions.
-
-To manually build RCCL with NPKit enabled, pass `-DNPKIT_FLAGS="-DENABLE_NPKIT -DENABLE_NPKIT_...(other NPKit compile-time switches)"` with cmake command. All NPKit compile-time switches are declared in the RCCL code base as macros with prefix `ENABLE_NPKIT_`, and they control which information will be collected. Also note that currently NPKit only supports collecting non-overlapped events on GPU, and `-DNPKIT_FLAGS` should follow this rule.
+## Library and API Documentation
 
-To manually run RCCL with NPKit enabled, environment variable `NPKIT_DUMP_DIR` needs to be set as the NPKit event dump directory. Also note that currently NPKit only supports 1 GPU per process.
+Please refer to the [RCCL Documentation Site](https://rocm.docs.amd.com/projects/rccl/en/latest/) for current documentation.
 
-To manually analyze NPKit dump results, please leverage [npkit_trace_generator.py](https://github.com/microsoft/NPKit/blob/main/rccl_samples/npkit_trace_generator.py).
+### How to build documentation
 
-## Library and API Documentation
+Run the steps below to build documentation locally.
 
-Please refer to the [Library documentation](https://rccl.readthedocs.io/) for current documentation.
+```shell
+cd docs
+pip3 install -r sphinx/requirements.txt
+python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
+```
 
 ## Copyright
 
-All source code and accompanying documentation is copyright (c) 2015-2022, NVIDIA CORPORATION. All rights reserved.
+All source code and accompanying documentation is copyright (c) 2015-2025, NVIDIA CORPORATION. All rights reserved.
 
-All modifications are copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
+All modifications are copyright (c) 2019-2025 Advanced Micro Devices, Inc. All rights reserved.
diff --git a/cmake/CheckSymbolExistsNoWarn.cmake b/cmake/CheckSymbolExistsNoWarn.cmake
new file mode 100644
index 0000000..b478f59
--- /dev/null
+++ b/cmake/CheckSymbolExistsNoWarn.cmake
@@ -0,0 +1,40 @@
+# MIT License
+#
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+# These overrides are due to CMake CHECK_SYMBOL_EXISTS modifying CMAKE_CXX_FLAGS to do a test compile,
+# while ROCMChecks gives a warning if this variable is modified manually without a target.
+
+# We now choose to disable ROCMChecks for this one case.
+
+set(DISABLE_ROCM_CHECK OFF)
+
+function(rocm_check_toolchain_var var access value list_file)
+  if(NOT DISABLE_ROCM_CHECK)
+    _rocm_check_toolchain_var("${var}" "${access}" "${value}" "${list_file}")
+  endif()
+endfunction()
+
+macro(CHECK_SYMBOL_EXISTS)
+  set(DISABLE_ROCM_CHECK ON)
+  _check_symbol_exists(${ARGN})
+  set(DISABLE_ROCM_CHECK OFF)
+endmacro()
diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
index 124c268..fc084ec 100644
--- a/cmake/Dependencies.cmake
+++ b/cmake/Dependencies.cmake
@@ -32,6 +32,8 @@
 # For downloading, building, and installing required dependencies
 include(cmake/DownloadProject.cmake)
 
+include(FetchContent)
+
 if(NOT INSTALL_DEPENDENCIES)
     find_package(GTest 1.11)
 endif()
@@ -50,7 +52,7 @@ if(NOT GTest_FOUND AND BUILD_TESTS OR INSTALL_DEPENDENCIES)
 
     download_project(PROJ                googletest
                      GIT_REPOSITORY      https://github.com/google/googletest.git
-                     GIT_TAG             release-1.11.0
+                     GIT_TAG             release-1.12.0
                      INSTALL_DIR         ${GTEST_ROOT}
                      CMAKE_ARGS          -DBUILD_GTEST=ON -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> ${COMPILER_OVERRIDE} -DBUILD_SHARED_LIBS=OFF
                      LOG_DOWNLOAD        TRUE
@@ -60,31 +62,37 @@ if(NOT GTest_FOUND AND BUILD_TESTS OR INSTALL_DEPENDENCIES)
                      UPDATE_DISCONNECTED TRUE
     )
     set(GTEST_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/gtest/include CACHE PATH "")
+    set(GMOCK_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/gmock/include CACHE PATH "")
     if(EXISTS ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib)
         set(GTEST_BOTH_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib/libgtest.a;${CMAKE_CURRENT_BINARY_DIR}/gtest/lib/libgtest_main.a CACHE PATH "")
+        set(GMOCK_BOTH_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib/libgmock.a;${CMAKE_CURRENT_BINARY_DIR}/gtest/lib/libgmock_main.a CACHE PATH "")
     elseif(EXISTS ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib64)
         set(GTEST_BOTH_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib64/libgtest.a;${CMAKE_CURRENT_BINARY_DIR}/gtest/lib64/libgtest_main.a CACHE PATH "")
+        set(GMOCK_BOTH_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/gtest/lib64/libgmock.a;${CMAKE_CURRENT_BINARY_DIR}/gtest/lib64/libgmock_main.a CACHE PATH "")
     else()
         message(FATAL_ERROR "Cannot find gtest library installation path.")
     find_package(GTest REQUIRED CONFIG PATHS ${GTEST_ROOT})
+    find_package(GMock REQUIRED CONFIG PATHS ${GTEST_ROOT})
     endif()
+elseif(GTest_FOUND AND BUILD_TESTS)
+    set(GTEST_BOTH_LIBRARIES "GTest::gtest;GTest::gtest_main")
+    set(GMOCK_BOTH_LIBRARIES "GTest::gmock;GTest::gmock_main")
 endif()
 
-
 # Find or download/install rocm-cmake project
 set( PROJECT_EXTERN_DIR ${CMAKE_CURRENT_BINARY_DIR}/extern )
 find_package(ROCM 0.7.3 QUIET CONFIG PATHS /opt/rocm)
 if(NOT ROCM_FOUND)
     set(rocm_cmake_tag "master" CACHE STRING "rocm-cmake tag to download")
     file(
-        DOWNLOAD https://github.com/RadeonOpenCompute/rocm-cmake/archive/${rocm_cmake_tag}.zip
+        DOWNLOAD https://github.com/ROCm/rocm-cmake/archive/${rocm_cmake_tag}.zip
         ${PROJECT_EXTERN_DIR}/rocm-cmake-${rocm_cmake_tag}.zip
         STATUS rocm_cmake_download_status LOG rocm_cmake_download_log
     )
     list(GET rocm_cmake_download_status 0 rocm_cmake_download_error_code)
     if(rocm_cmake_download_error_code)
         message(FATAL_ERROR "Error: downloading "
-            "https://github.com/RadeonOpenCompute/rocm-cmake/archive/${rocm_cmake_tag}.zip failed "
+            "https://github.com/ROCm/rocm-cmake/archive/${rocm_cmake_tag}.zip failed "
             "error_code: ${rocm_cmake_download_error_code} "
             "log: ${rocm_cmake_download_log} "
         )
@@ -106,6 +114,74 @@ if(NOT ROCM_FOUND)
     find_package( ROCM 0.7.3 REQUIRED CONFIG PATHS ${PROJECT_EXTERN_DIR}/rocm-cmake )
 endif()
 
+set(CMAKE_INSTALL_LIBDIR lib CACHE STRING "Define install directory for libraries" FORCE)
+
+# Find or download/install fmt
+find_package(fmt QUIET)
+if(NOT fmt_FOUND)
+    set(FMT_INSTALL OFF)
+    message(STATUS "fmt not found, fetching from source...")
+    FetchContent_Declare(
+        fmt
+        GIT_REPOSITORY https://github.com/fmtlib/fmt
+        GIT_TAG        e69e5f977d458f2650bb346dadf2ad30c5320281 # 10.2.1
+    )
+    FetchContent_MakeAvailable(fmt)
+else()
+    message(STATUS "Using system fmt")
+    get_target_property(FMT_INCLUDE_DIRS fmt::fmt-header-only INTERFACE_INCLUDE_DIRECTORIES)
+    message(STATUS "fmt include directories: ${FMT_INCLUDE_DIRS}")
+endif()
+
+# Find available local ROCM targets
+# NOTE: This will eventually be part of ROCm-CMake and should be removed at that time
+function(rocm_local_targets VARIABLE)
+  set(${VARIABLE} "NOTFOUND" PARENT_SCOPE)
+  find_program(_rocm_agent_enumerator rocm_agent_enumerator HINTS /opt/rocm/bin ENV ROCM_PATH)
+  if(NOT _rocm_agent_enumerator STREQUAL "_rocm_agent_enumerator-NOTFOUND")
+    execute_process(
+      COMMAND "${_rocm_agent_enumerator}"
+      RESULT_VARIABLE _found_agents
+      OUTPUT_VARIABLE _rocm_agents
+      ERROR_QUIET
+      )
+    if (_found_agents EQUAL 0)
+      string(REPLACE "\n" ";" _rocm_agents "${_rocm_agents}")
+      unset(result)
+      foreach (agent IN LISTS _rocm_agents)
+        if (NOT agent STREQUAL "gfx000")
+          list(APPEND result "${agent}")
+        endif()
+      endforeach()
+      if(result)
+        list(REMOVE_DUPLICATES result)
+        set(${VARIABLE} "${result}" PARENT_SCOPE)
+      endif()
+    endif()
+  endif()
+endfunction()
+
+# Iterate over the "source" list and check if there is a duplicate file name
+# NOTE: This is due to compiler bug '--save-temps' and can be removed when fix availabe
+function(add_file_unique FILE_LIST FILE)
+  get_filename_component(FILE_NAME "${FILE}" NAME)
+
+  # Iterate over whatever is in the list so far
+  foreach(curr_file IN LISTS ${FILE_LIST})
+    get_filename_component(curr_file_name ${curr_file} NAME)
+
+    # Check if duplicate
+    if(${FILE_NAME} STREQUAL ${curr_file_name})
+      get_filename_component(DIR_PATH "${FILE}" DIRECTORY)
+      get_filename_component(FILE_NAME_WE "${FILE}" NAME_WE)
+      get_filename_component(FILE_EXT "${FILE}" EXT)
+
+      # Construct a new file name by adding _tmp
+      set(HIP_FILE "${DIR_PATH}/${FILE_NAME_WE}_tmp${FILE_EXT}" PARENT_SCOPE)
+    endif()
+  endforeach()
+endfunction()
+
 include(ROCMSetupVersion)
 include(ROCMCreatePackage)
 include(ROCMInstallTargets)
@@ -113,4 +189,4 @@ include(ROCMPackageConfigHelpers)
 include(ROCMInstallSymlinks)
 include(ROCMCheckTargetIds)
 include(ROCMClients)
-include( ROCMHeaderWrapper )
+include(ROCMHeaderWrapper)
diff --git a/cmake/DownloadProject.cmake b/cmake/DownloadProject.cmake
index 110bbd5..765b1f6 100644
--- a/cmake/DownloadProject.cmake
+++ b/cmake/DownloadProject.cmake
@@ -158,7 +158,7 @@ function(download_project)
     if(result)
         message(FATAL_ERROR "CMake step for ${DL_ARGS_PROJ} failed: ${result}")
     endif()
-    execute_process(COMMAND ${CMAKE_COMMAND} --build .
+    execute_process(COMMAND ${CMAKE_COMMAND} --build . -j16
                     RESULT_VARIABLE result
                     ${OUTPUT_QUIET}
                     WORKING_DIRECTORY "${DL_ARGS_DOWNLOAD_DIR}"
@@ -167,4 +167,4 @@ function(download_project)
         message(FATAL_ERROR "Build step for ${DL_ARGS_PROJ} failed: ${result}")
     endif()
 
-endfunction()
\ No newline at end of file
+endfunction()
diff --git a/cmake/FindIBVerbs.cmake b/cmake/FindIBVerbs.cmake
new file mode 100644
index 0000000..d02c0e1
--- /dev/null
+++ b/cmake/FindIBVerbs.cmake
@@ -0,0 +1,39 @@
+# MIT License
+#
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+find_path(IBVERBS_INCLUDE_DIRS
+  NAMES infiniband/verbs.h
+  HINTS
+  ${IBVERBS_INCLUDE_DIR}
+  ${IBVERBS_ROOT_DIR}
+  ${IBVERBS_ROOT_DIR}/include)
+
+find_library(IBVERBS_LIBRARIES
+  NAMES ibverbs
+  HINTS
+  ${IBVERBS_LIB_DIR}
+  ${IBVERBS_ROOT_DIR}
+  ${IBVERBS_ROOT_DIR}/lib)
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(IBVerbs DEFAULT_MSG IBVERBS_INCLUDE_DIRS IBVERBS_LIBRARIES)
+mark_as_advanced(IBVERBS_INCLUDE_DIR IBVERBS_LIBRARIES)
diff --git a/cmake/Findmscclpp_nccl.cmake b/cmake/Findmscclpp_nccl.cmake
new file mode 100644
index 0000000..313b0c6
--- /dev/null
+++ b/cmake/Findmscclpp_nccl.cmake
@@ -0,0 +1,36 @@
+# MIT License
+#
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+find_path(MSCCLPP_INCLUDE_DIRS
+    NAMES mscclpp/gpu.hpp
+    HINTS
+    ${MSCCLPP_ROOT}/include)
+
+find_library(MSCCLPP_LIBRARIES
+    NAMES mscclpp_nccl
+    HINTS
+    ${MSCCLPP_ROOT}/lib)
+
+include (FindPackageHandleStandardArgs)
+find_package_handle_standard_args(mscclpp_nccl DEFAULT_MSG MSCCLPP_INCLUDE_DIRS MSCCLPP_LIBRARIES)
+mark_as_advanced(MSCCLPP_INCLUDE_DIRS MSCCLPP_LIBRARIES)
+    
\ No newline at end of file
diff --git a/cmake/MSCCLPP.cmake b/cmake/MSCCLPP.cmake
new file mode 100644
index 0000000..21fbc64
--- /dev/null
+++ b/cmake/MSCCLPP.cmake
@@ -0,0 +1,229 @@
+# MIT License
+#
+# Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+# Dependencies
+
+# HIP dependency is handled earlier in the project cmake file
+# when VerifyCompiler.cmake is included.
+
+# GIT
+
+# Test dependencies
+
+# For downloading, building, and installing required dependencies
+include(cmake/DownloadProject.cmake)
+
+if(ENABLE_MSCCLPP)
+    # Try to find the mscclpp install
+    set(MSCCLPP_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/ext/mscclpp CACHE PATH "")
+    execute_process(
+        COMMAND mkdir -p ${MSCCLPP_ROOT}
+    )
+    list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+    find_package(mscclpp_nccl)
+
+    #if(NOT mscclpp_nccl_FOUND)
+        # Ensure the source code is checked out
+        set(MSCCLPP_SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/mscclpp CACHE PATH "")
+        set(JSON_SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/json CACHE PATH "")
+        if((NOT EXISTS ${MSCCLPP_SOURCE}/CMakeLists.txt) OR (NOT EXISTS ${JSON_SOURCE}/CMakeLists.txt))
+            message(STATUS "Checking out external code")
+            execute_process(
+                COMMAND git submodule update --init --recursive
+                WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
+            )
+        endif()
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/cpx.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/read-allred.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/mscclpp_ibv_access_relaxed_ordering.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/mem-reg.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/non-multiple-128-fix.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/bf16-tuning.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/reg-fix.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/no-cache.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/device-flag.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/remove-clip.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/disable-executor.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/disable-format-checks.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        set(CMAKE_INHERITED_ARGS "")
+        set(CMAKE_ARGS_LIST "CMAKE_PREFIX_PATH;CMAKE_INSTALL_RPATH_USE_LINK_PATH;HIP_COMPILER")
+        foreach(arg IN LISTS CMAKE_ARGS_LIST)
+            if(DEFINED ${arg})
+                string(REPLACE ";" "%" ARG_VALUE "${${arg}}") # Replace ; with new list separator symbol % to avoid CMake errors
+                string(STRIP "${ARG_VALUE}" ARG_VALUE) # Eliminate whitespace, reducing to empty string if necessary
+
+                # Only add a cmake argument if it has a value
+                if("${ARG_VALUE}" STREQUAL "")
+                    continue()
+                endif()
+                string(APPEND CMAKE_INHERITED_ARGS "-D${arg}=\"${ARG_VALUE}\" ")
+            endif()
+        endforeach()
+
+        if(NOT DEFINED CACHE{MSCCLPP_GPU_TARGETS})
+            message(STATUS "Building MSCCL++ only for supported variants: gfx942;gfx950")
+            set(MSCCLPP_GPU_TARGETS "gfx942;gfx950")
+            if(BUILD_ADDRESS_SANITIZER)
+                set(MSCCLPP_GPU_TARGETS "gfx942:xnack+;gfx950:xnack+")
+            endif()
+        else()
+            message(STATUS "Building MSCCL++ for ${MSCCLPP_GPU_TARGETS}")
+        endif()
+
+        string(REPLACE ";" "%" MSCCLPP_GPU_TARGETS "${MSCCLPP_GPU_TARGETS}")
+
+        download_project(PROJ                mscclpp_nccl
+                         #GIT_REPOSITORY      https://github.com/microsoft/mscclpp.git
+                         #GIT_TAG             4ee15b7ad085daaf74349d4c49c9b8480d28f0dc
+                         INSTALL_DIR         ${MSCCLPP_ROOT}
+                         LIST_SEPARATOR      %
+			 CMAKE_ARGS          "-DGPU_TARGETS=${MSCCLPP_GPU_TARGETS}" -DMSCCLPP_BYPASS_GPU_CHECK=ON -DMSCCLPP_USE_ROCM=ON -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} -DMSCCLPP_BUILD_APPS_NCCL=ON -DMSCCLPP_BUILD_PYTHON_BINDINGS=OFF -DMSCCLPP_BUILD_TESTS=OFF -DMSCCLPP_CLIP_ENABLED=${ENABLE_MSCCLPP_CLIP} -DMSCCLPP_ENABLE_EXECUTOR=${ENABLE_MSCCLPP_EXECUTOR} -DMSCCLPP_ENABLE_FORMAT_CHECKS=${ENABLE_MSCCLPP_FORMAT_CHECKS} -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_VERBOSE_MAKEFILE=1 "${CMAKE_INHERITED_ARGS}" -DFETCHCONTENT_SOURCE_DIR_JSON=${JSON_SOURCE}
+                         LOG_DOWNLOAD        FALSE
+                         LOG_CONFIGURE       FALSE
+                         LOG_BUILD           FALSE
+                         LOG_INSTALL         FALSE
+                         UPDATE_DISCONNECTED TRUE
+                         SOURCE_DIR          ${MSCCLPP_SOURCE}
+        )
+
+        find_package(mscclpp_nccl REQUIRED)
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/disable-format-checks.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/disable-executor.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/remove-clip.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/device-flag.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/no-cache.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/reg-fix.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/bf16-tuning.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/non-multiple-128-fix.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/mem-reg.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/mscclpp_ibv_access_relaxed_ordering.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/read-allred.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+        execute_process(
+            COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/cpx.patch
+            WORKING_DIRECTORY ${MSCCLPP_SOURCE}
+        )
+
+    #endif()
+
+    execute_process(COMMAND objcopy
+                    --redefine-syms=${CMAKE_CURRENT_SOURCE_DIR}/src/misc/mscclpp/mscclpp_nccl_syms.txt
+                    "${MSCCLPP_ROOT}/lib/libmscclpp_nccl_static.a"
+                    "${PROJECT_BINARY_DIR}/libmscclpp_nccl.a"
+    )
+    add_library(mscclpp_nccl STATIC IMPORTED)
+    set_target_properties(mscclpp_nccl PROPERTIES IMPORTED_LOCATION ${PROJECT_BINARY_DIR}/libmscclpp_nccl.a)
+
+endif()
diff --git a/cmake/rcclRAS.cmake b/cmake/rcclRAS.cmake
new file mode 100644
index 0000000..255623e
--- /dev/null
+++ b/cmake/rcclRAS.cmake
@@ -0,0 +1,24 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+
+cmake_minimum_required(VERSION 3.16)
+
+message("Building rccl RAS client executable")
+
+add_executable(rcclras "${PROJECT_BINARY_DIR}/hipify/src/ras/client.cc")
+
+target_include_directories(rcclras PRIVATE ${PROJECT_BINARY_DIR}/include)
+target_include_directories(rcclras PRIVATE ${HIPIFY_DIR}/src)
+target_include_directories(rcclras PRIVATE ${HIPIFY_DIR}/src/include)
+
+target_link_libraries(rcclras PRIVATE hip::host)
+target_link_libraries(rcclras PRIVATE dl)
+
+if(BUILD_SHARED_LIBS)
+  target_link_libraries(rcclras PRIVATE rccl hip::device)
+else()
+  add_dependencies(rcclras rccl)
+  target_link_libraries(rcclras PRIVATE dl rt -lrccl -L${CMAKE_BINARY_DIR} -lamdhip64 -L${ROCM_PATH}/lib)
+endif()
+
+
+rocm_install(TARGETS rcclras)
diff --git a/cmake/scripts/add_faults.sh b/cmake/scripts/add_faults.sh
new file mode 100755
index 0000000..5d6c59f
--- /dev/null
+++ b/cmake/scripts/add_faults.sh
@@ -0,0 +1,27 @@
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+HIP_FILE=$1
+
+if [[ "$HIP_FILE" =~ .*/src/device/.*\.h ]]; then
+  sed -i "s/__syncthreads()/__syncthreads(); insert_random_delay_per_warp()/" "$HIP_FILE"
+
+  echo "Added fault injection to $HIP_FILE"
+fi
\ No newline at end of file
diff --git a/cmake/scripts/add_unroll.sh b/cmake/scripts/add_unroll.sh
new file mode 100755
index 0000000..3b324b3
--- /dev/null
+++ b/cmake/scripts/add_unroll.sh
@@ -0,0 +1,42 @@
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+HIP_FILE=$1
+
+if [[ "$HIP_FILE" =~ .*/src/device/.*\.h ]]; then
+  perl -pi -e 's/(template<typename T, typename RedOp(?:, typename Proto)?)(, bool isNetOffload.*?)?>/\1, int USE_ACC, int COLL_UNROLL, int Pipeline\2>/g' "$HIP_FILE"
+  perl -pi -e 's/(template<typename T, typename RedOp(?:, typename Proto)?(?:, int RCCLMetadata)?)(, bool isNetOffload.*?)?>/\1, int USE_ACC, int COLL_UNROLL, int Pipeline\2>/g' "$HIP_FILE"
+  perl -pi -e 's/(ProtoSimple<[^,]*?,[^,]+?)>/\1, USE_ACC, COLL_UNROLL>/g' "$HIP_FILE"
+  perl -pi -e 's/(runRing<T.*?)((, (true|false))?>\()/\1, USE_ACC, COLL_UNROLL\2/g' "$HIP_FILE"
+  perl -pi -e 's/(runTreeUpDown<T.*?)>\(/\1, USE_ACC, COLL_UNROLL>(/' "$HIP_FILE"
+  perl -pi -e 's/(runTreeSplit<T.*?)>\(/\1, USE_ACC, COLL_UNROLL>(/' "$HIP_FILE"
+
+  perl -pi -e 's/(runTreeSplit<T, RedOp, (ProtoLL|ProtoLL128), USE_ACC, COLL_UNROLL.*?)>/\1, 0>/' "$HIP_FILE"
+  perl -pi -e 's/(runTreeUpDown<T, RedOp, (ProtoLL|ProtoLL128), USE_ACC, COLL_UNROLL.*?)>/\1, 0>/' "$HIP_FILE"
+  perl -pi -e 's/(runRing<T, RedOp, (ProtoLL|ProtoLL128), USE_ACC, COLL_UNROLL.*?)>/\1, 0>/' "$HIP_FILE"
+  perl -pi -e 's/(runRing<T, RedOp, (ProtoLL|ProtoLL128), (RCCL_ONE_NODE_RING_SIMPLE|RCCL_METADATA_EMPTY), USE_ACC, COLL_UNROLL.*?)>/\1, 0>/' "$HIP_FILE"
+
+  perl -pi -e 's/(runRing<T, RedOp, Proto, (RCCL_ONE_NODE_RING_SIMPLE|RCCL_METADATA_EMPTY), USE_ACC, COLL_UNROLL.*?)>/\1, Pipeline>/' "$HIP_FILE"
+  perl -pi -e 's/(runRing<T, RedOp, Proto, USE_ACC, COLL_UNROLL.*?)>/\1, Pipeline>/' "$HIP_FILE"
+  perl -pi -e 's/(runTreeSplit<T, RedOp, Proto, USE_ACC, COLL_UNROLL.*?)>/\1, Pipeline>/' "$HIP_FILE"
+  perl -pi -e 's/(runTreeUpDown<T, RedOp, Proto, USE_ACC, COLL_UNROLL.*?)>/\1, Pipeline>/' "$HIP_FILE"
+  sed -i "s/\\(struct RunWorkBatch<ncclFunc[^>]*\\)>*/\\1, USE_ACC, COLL_UNROLL, Pipeline>/" "$HIP_FILE"
+  sed -i "s/\\(RunWorkColl<[^,]*,[^,]*,[^,]*,[^,]*,[^>]*\\)>/\\1, USE_ACC, COLL_UNROLL, Pipeline>/" "$HIP_FILE"
+fi
\ No newline at end of file
diff --git a/cmake/scripts/extract_metadata.cmake b/cmake/scripts/extract_metadata.cmake
new file mode 100644
index 0000000..7e70124
--- /dev/null
+++ b/cmake/scripts/extract_metadata.cmake
@@ -0,0 +1,81 @@
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+set(EXTRACT_TIMEOUT 5 CACHE STRING "Timeout in seconds for roc-obj-* calls")
+
+## List the objects for each gfx architecture
+execute_process( COMMAND roc-obj-ls librccl.so
+    RESULT_VARIABLE list_result
+    OUTPUT_VARIABLE cmd_output
+    ERROR_VARIABLE cmd_error
+    OUTPUT_STRIP_TRAILING_WHITESPACE
+    ERROR_STRIP_TRAILING_WHITESPACE
+    TIMEOUT ${EXTRACT_TIMEOUT}
+)
+
+if(list_result EQUAL 0)
+    ## Convert cmd output to list of lines
+    string(REGEX REPLACE "\n$" "" cmd_output "${cmd_output}")
+    string(REPLACE "\n" ";" cmd_output "${cmd_output}")
+
+    ## Extract file paths for the selected gfx archs
+    foreach(line ${cmd_output})
+        if(line MATCHES "(gfx90a|gfx942|gfx950)")
+            string(REGEX MATCH "\\file://(.*)" file_match ${line})
+            if(file_match)
+                list(APPEND file_paths ${file_match})
+            endif()
+        endif()
+    endforeach()
+
+    ## Extract objects from files
+    foreach(file ${file_paths})
+        execute_process(
+          COMMAND roc-obj-extract ${file}
+          RESULT_VARIABLE extraction_result
+          ERROR_VARIABLE extraction_error
+          OUTPUT_STRIP_TRAILING_WHITESPACE
+          ERROR_STRIP_TRAILING_WHITESPACE
+          TIMEOUT ${EXTRACT_TIMEOUT}
+        )
+        if(extraction_result STREQUAL "TIMEOUT")
+          message(
+            WARNING
+              "[Timeout] Extraction of '${file}' did not finish within ${EXTRACT_TIMEOUT}s. stderr: ${extraction_error}.
+                    Timeouts have been known to happen as a result of mismatched ROCm versions/executables/etc."
+          )
+        elseif(NOT extraction_result EQUAL 0)
+          message(
+            WARNING
+              "[Error ${extraction_result}] Could not extract objects from '${file}'. stderr: ${extraction_error}"
+          )
+        endif()
+    endforeach()
+
+elseif(list_result STREQUAL "TIMEOUT")
+  message(
+    WARNING
+      "[Timeout] roc-obj-ls did not finish within ${EXTRACT_TIMEOUT}s. stderr: ${cmd_error}.
+                     Timeouts have been known to happen as a result of mismatched ROCm versions/executables/etc"
+  )
+else()
+    ## We don't want to stop building unit-tests if this command fails.
+    message(WARNING "[Error ${list_result}] roc-obj-ls failed. stderr: ${cmd_error}")
+endif()
diff --git a/cmake/git_version.cmake b/cmake/scripts/git_version.cmake
similarity index 99%
rename from cmake/git_version.cmake
rename to cmake/scripts/git_version.cmake
index e070739..fde9485 100755
--- a/cmake/git_version.cmake
+++ b/cmake/scripts/git_version.cmake
@@ -19,6 +19,7 @@
 # SOFTWARE.
 
 # Attempt to collect the latest git hash
+set(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
 execute_process(COMMAND git log --pretty=format:'%h' -n 1
                 OUTPUT_VARIABLE GIT_REV
                 ERROR_QUIET)
diff --git a/custom.properties b/custom.properties
deleted file mode 100644
index 45e5972..0000000
--- a/custom.properties
+++ /dev/null
@@ -1,3 +0,0 @@
-booktitle=RCCL API Guide
-spreadsheet.xml=docs/classification-map.xml
-document.locale=enus
\ No newline at end of file
diff --git a/debian/changelog b/debian/changelog
index fcb999b..c1a6a4f 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -1,3 +1,32 @@
+rccl (7.1.0-0ubuntu2) resolute; urgency=medium
+
+  * d/control: add libamdhip64-dev (=> 7.1~) as a build dep
+
+ -- Bojan Aleksovski <[e-mail]>  Thu, 05 Feb 2026 14:15:33 +0100
+
+rccl (7.1.0-0ubuntu1) resolute; urgency=medium
+
+  * d/p/0004: [Temporary] Patch rccl to remove rocm-core as build-dep
+  * d/control: Update build deps for ROCm v7.1.0
+  * d/control: update maintainer to ubuntu devel
+  * d/rules: align ROCM_ISAS with expected list syntax
+  * d/librccl1.install: add rcclras install
+  * d/librccl1-tests.install: fix path and add msccl-unit-test-algorithms
+  * d/librccl-doc.install: add doc install for msccl-algorithms
+  * d/{librccl1.symbols,shlibs}: remove symbols file and instead add shlibs
+  * d/patches: add new fix-install-libdir.patch
+  * d/control: remove ppc64el build
+  * d/rules: [Temporary] remove gfx908 and gfx90a ROCm ISAs
+
+ -- Bojan Aleksovski <[e-mail]>  Fri, 30 Jan 2026 10:07:24 +0100
+
+rccl (7.1.0-0ubuntu1~exp1) resolute; urgency=medium
+
+  * new upstream 7.1.0 release
+  * drop 0002-expand-architecture-support.patch as it was accepted upstream
+
+ -- Tim Flink <[e-mail]>  Mon, 19 Jan 2026 17:42:17 -0600
+
 rccl (5.4.3-3ubuntu2) resolute; urgency=medium
 
   * No change rebuild for rocm-hipamd.
diff --git a/debian/control b/debian/control
index 12c6691..c0c8d8e 100644
--- a/debian/control
+++ b/debian/control
@@ -5,24 +5,32 @@ Priority: optional
 Standards-Version: 4.6.2
 Vcs-Git: https://salsa.debian.org/rocm-team/rccl.git
 Vcs-Browser: https://salsa.debian.org/rocm-team/rccl
-Maintainer: Debian ROCm Team <[e-mail]>
+Maintainer: Ubuntu Developers <[e-mail]>
+XSBC-Original-Maintainer: Debian ROCm Team <[e-mail]>
 Uploaders: Cordell Bloor <[e-mail]>,
            Christian Kastner <[e-mail]>,
+           Tim Flink <[e-mail]>
 Build-Depends: debhelper-compat (= 13),
                cmake,
-               hipcc (>= 5.6.1~),
-               libamd-comgr-dev (>= 6.0~),
-               libhsa-runtime-dev (>= 5.7.1~),
+               hipcc (>= 7.1~),
+               libamd-comgr-dev (>= 7.1~),
+               libhsa-runtime-dev (>= 7.1~),
+               libamdhip64-dev (>= 7.1~),
                rocm-cmake,
                librocm-smi-dev,
                liboam-dev,
                chrpath <!nocheck>,
+               pkg-rocm-tools (>= 0.9.4),
+               libfmt-dev,
+               hipify-perl (>= 7.0.0),
+               libdrm-dev,
                libgtest-dev <!nocheck>
 Build-Depends-Indep: dh-sequence-sphinxdoc <!nodoc>,
                doxygen <!nodoc>,
                python3-breathe <!nodoc>,
                python3-sphinx <!nodoc>,
                python3-sphinx-rtd-theme <!nodoc>,
+               python3-rocm-docs (>= 1.20.0-1~) <!nodoc>,
                libjs-jquery <!nodoc>,
                libjs-mathjax <!nodoc>,
                libjs-sphinxdoc <!nodoc>,
@@ -31,7 +39,7 @@ Rules-Requires-Root: no
 
 Package: librccl1
 Section: libs
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: ${misc:Depends}, ${shlibs:Depends},
 Description: ROCm Communication Collectives Library - library
  RCCL (pronounced "Rickle") is a library of collective communication routines
@@ -45,7 +53,7 @@ Description: ROCm Communication Collectives Library - library
 
 Package: librccl-dev
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: librccl1 (= ${binary:Version}),${misc:Depends}, ${shlibs:Depends},
          libamdhip64-dev,
 Suggests: librccl-doc
@@ -61,7 +69,7 @@ Description: ROCm Communication Collectives Library - headers
 
 Package: librccl1-tests
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 Depends: librccl1 (= ${binary:Version}), ${misc:Depends}, ${shlibs:Depends},
 Build-Profiles: <!nocheck>
 Description: ROCm Communication Collectives Library - tests
diff --git a/debian/librccl-doc.install b/debian/librccl-doc.install
new file mode 100644
index 0000000..865d42f
--- /dev/null
+++ b/debian/librccl-doc.install
@@ -0,0 +1 @@
+usr/share/rccl/msccl-algorithms/*
diff --git a/debian/librccl1-tests.install b/debian/librccl1-tests.install
index 0e38109..623ee95 100644
--- a/debian/librccl1-tests.install
+++ b/debian/librccl1-tests.install
@@ -1,2 +1,3 @@
-usr/bin/UnitTests      usr/libexec/rocm/librccl1-tests
+usr/bin/rccl-UnitTests      usr/libexec/rocm/librccl1-tests
+usr/share/rccl/msccl-unit-test-algorithms/*
 
diff --git a/debian/librccl1.install b/debian/librccl1.install
index b97befa..8080107 100644
--- a/debian/librccl1.install
+++ b/debian/librccl1.install
@@ -1 +1,2 @@