support rocm7.2

.azure-pipelines/codecov.yml

@@ -80,6 +80,8 @@ jobs:
     matrix:
       rocm6_2:
         containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm6.2
+      rocm7_2:
+        containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm7.2
 
   container:
     image: $(containerImage)

.azure-pipelines/rccl-api-test.yml

@@ -35,6 +35,8 @@ jobs:
     matrix:
       rocm6_2:
         containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm6.2
+      rocm7_2:
+        containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm7.2
 
   container:
     image: $(containerImage)

.azure-pipelines/templates/rccl-test.yml

@@ -46,25 +46,25 @@ steps:
     name: RunRcclAllGatherTest
     displayName: Run RCCL AllGather Test with or without MSCCLPP Lib
     remoteScript: |
-      mpirun -np 8 --bind-to numa --allow-run-as-root -x LD_PRELOAD=/root/mscclpp/build/lib/libmscclpp_nccl.so -x MSCCLPP_NCCL_SYMMETRIC_MEMORY=1 -x NCCL_DEBUG=WARN /root/rocm-systems/projects/rccl-tests/build/all_gather_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
-      mpirun -np 8 --bind-to numa --allow-run-as-root /root/rocm-systems/projects/rccl-tests/build/all_gather_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
+      mpirun -np 8 --bind-to numa --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 -x LD_PRELOAD=/root/mscclpp/build/lib/libmscclpp_nccl.so -x MSCCLPP_NCCL_SYMMETRIC_MEMORY=1 -x NCCL_DEBUG=WARN /root/rocm-systems/projects/rccl-tests/build/all_gather_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
+      mpirun -np 8 --bind-to numa --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 /root/rocm-systems/projects/rccl-tests/build/all_gather_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
 
 - template: run-remote-task.yml
   parameters:
     name: RunRcclAllReduceTest
     displayName: Run RCCL AllReduce Test with or without MSCCLPP Lib
     remoteScript: |
-      mpirun -np 8 --bind-to numa --allow-run-as-root -x LD_PRELOAD=/root/mscclpp/build/lib/libmscclpp_nccl.so -x MSCCLPP_NCCL_SYMMETRIC_MEMORY=1 -x NCCL_DEBUG=WARN /root/rocm-systems/projects/rccl-tests/build/all_reduce_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
-      mpirun -np 8 --bind-to numa --allow-run-as-root /root/rocm-systems/projects/rccl-tests/build/all_reduce_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
+      mpirun -np 8 --bind-to numa --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 -x LD_PRELOAD=/root/mscclpp/build/lib/libmscclpp_nccl.so -x MSCCLPP_NCCL_SYMMETRIC_MEMORY=1 -x NCCL_DEBUG=WARN /root/rocm-systems/projects/rccl-tests/build/all_reduce_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
+      mpirun -np 8 --bind-to numa --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 /root/rocm-systems/projects/rccl-tests/build/all_reduce_perf -b 1K -e 1G -f 2 -d half -G 20 -w 10 -n 20
 
 - template: run-remote-task.yml
   parameters:
     name: PyBench
     displayName: Run Collective Benchmarks
     remoteScript: |
-      mpirun --allow-run-as-root -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allreduce --dtype float8_e4m3b15 --accum-type float32 --autotune
-      mpirun --allow-run-as-root -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allreduce --dtype float8_e4m3fnuz --accum-type float32 --autotune
-      mpirun --allow-run-as-root -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allgather --dtype float8_e4m3b15 --autotune --buffer-mode out-of-place
+      mpirun --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allreduce --dtype float8_e4m3b15 --accum-type float32 --autotune
+      mpirun --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allreduce --dtype float8_e4m3fnuz --accum-type float32 --autotune
+      mpirun --allow-run-as-root -x HSA_NO_SCRATCH_RECLAIM=1 -x GPU_MAX_HW_QUEUES=8 -np 8 python3 -m mscclpp_benchmark.bench_collective --collective allgather --dtype float8_e4m3b15 --autotune --buffer-mode out-of-place
 
 - template: stop.yml
   parameters:

.azure-pipelines/ut.yml

@@ -137,6 +137,8 @@ jobs:
     matrix:
       rocm6_2:
         containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm6.2
+      rocm7_2:
+        containerImage: ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm7.2
 
   container:
     image: $(containerImage)

.github/workflows/codeql-analysis.yml

@@ -85,7 +85,7 @@ jobs:
       fail-fast: false
       matrix:
         language: [ 'cpp', 'python' ]
-        version: [ 'rocm6.2' ]
+        version: [ 'rocm6.2', 'rocm7.2' ]
 
     steps:
     - name: Checkout repository

CMakeLists.txt

@@ -59,6 +59,7 @@ option(MSCCLPP_BYPASS_GPU_CHECK "Bypass GPU check." OFF)
 option(MSCCLPP_NPKIT_FLAGS "Set NPKIT flags" OFF)
 option(MSCCLPP_ENABLE_COVERAGE "Enable code coverage" OFF)
 option(MSCCLPP_DISABLE_NB_LEAK_WARNINGS "Disable Nanobind leak warnings" ON)
+option(MSCCLPP_ROCM_USE_FNUZ_FP8 "Use ROCm FNUZ native FP8 types." ON)
 set(MSCCLPP_GPU_ARCHS "" CACHE STRING "Specify GPU architectures with delimiters (comma, space, or semicolon).")
 
 if(MSCCLPP_BYPASS_GPU_CHECK)
@@ -183,11 +184,20 @@ elseif(MSCCLPP_USE_CUDA)
 	    endif()
     endif()
 elseif(MSCCLPP_USE_ROCM)
-    set(MSCCLPP_GPU_ARCHS gfx90a gfx941 gfx942)
+    set(MSCCLPP_GPU_ARCHS gfx90a gfx942)
 endif()
 
 message(STATUS "GPU architectures: ${MSCCLPP_GPU_ARCHS}")
 
+if(MSCCLPP_USE_ROCM)
+    if(MSCCLPP_ROCM_USE_FNUZ_FP8)
+        add_compile_definitions(MSCCLPP_ROCM_FP8_FNUZ)
+        message(STATUS "ROCm native FP8 aliases: FNUZ (MSCCLPP_ROCM_USE_FNUZ_FP8=ON)")
+    else()
+        message(STATUS "ROCm native FP8 aliases: OCP (MSCCLPP_ROCM_USE_FNUZ_FP8=OFF)")
+    endif()
+endif()
+
 # Declare project
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wextra")

docker/base-dev-x.dockerfile

@@ -70,7 +70,7 @@ RUN if echo "$TARGET" | grep -q "^cuda"; then \
 # Install ROCm-specific packages if building for ROCm
 RUN if echo "$TARGET" | grep -q "^rocm"; then \
         apt-get update -y && \
-        apt-get install -y hipblas hipsparse rocsparse rocrand hiprand rocthrust rocsolver rocfft hipfft hipcub rocprim rccl roctracer-dev && \
+        apt-get install -y hipblas hipsparse rocsparse rocrand hiprand rocthrust-dev rocsolver rocfft hipfft hipcub-dev rocprim-dev rccl roctracer-dev && \
         apt-get autoremove -y && \
         apt-get clean && \
         rm -rf /var/lib/apt/lists/* /tmp/*; \
@@ -84,7 +84,7 @@ ENV PATH="/root/venv/bin:${PATH}"
 # Install Python dependencies
 ADD . /tmp/mscclpp
 WORKDIR /tmp/mscclpp
-RUN target_type=$(echo $TARGET | sed 's/\.[0-9]*$//') && \
+RUN target_type=$(echo "$TARGET" | sed -E 's/^([[:alpha:]]+[0-9]+).*/\1/') && \
     if echo "$TARGET" | grep -q "^rocm"; then \
         export CUPY_INSTALL_USE_HIP=1 && export ROCM_HOME=/opt/rocm; \
     fi && \

docker/build.sh

@@ -10,11 +10,13 @@ baseImageTable=(
     ["cuda12.9"]="nvidia/cuda:12.9.1-devel-ubuntu24.04"
     ["cuda13.0"]="nvidia/cuda:13.0.2-devel-ubuntu24.04"
     ["rocm6.2"]="rocm/dev-ubuntu-22.04:6.2.2"
+    ["rocm7.2"]="rocm/dev-ubuntu-24.04:7.2.4"
 )
 
 declare -A extraLdPathTable
 extraLdPathTable=(
     ["rocm6.2"]="/opt/rocm/lib"
+    ["rocm7.2"]="/opt/rocm/lib"
 )
 
 declare -A ofedVersionTable
@@ -25,13 +27,14 @@ ofedVersionTable=(
     ["cuda12.9"]="24.10-1.1.4.0"
     ["cuda13.0"]="24.10-3.2.5.0"
     ["rocm6.2"]="24.10-1.1.4.0"
+    ["rocm7.2"]="24.10-3.2.5.0"
 )
 
 TARGET=${1}
 OS_ARCH=$(uname -m)
 
 print_usage() {
-    echo "Usage: $0 [cuda11.8|cuda12.4|cuda12.8|cuda12.9|cuda13.0|rocm6.2]"
+    echo "Usage: $0 [cuda11.8|cuda12.4|cuda12.8|cuda12.9|cuda13.0|rocm6.2|rocm7.2]"
 }
 
 if [[ ! -v "baseImageTable[${TARGET}]" ]]; then

docs/quickstart.md

@@ -47,7 +47,7 @@ We provide docker images which package all prerequisites for MSCCL++. You can se
 # For NVIDIA platforms
 $ docker run -it --privileged --net=host --ipc=host --gpus all --name mscclpp-dev ghcr.io/microsoft/mscclpp/mscclpp:base-dev-cuda12.9 bash
 # For AMD platforms
-$ docker run -it --privileged --net=host --ipc=host --security-opt=seccomp=unconfined --group-add=video --name mscclpp-dev ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm6.2 bash
+$ docker run -it --privileged --net=host --ipc=host --security-opt=seccomp=unconfined --group-add=video --name mscclpp-dev ghcr.io/microsoft/mscclpp/mscclpp:base-dev-rocm7.2 bash
 ```
 
 See all available images [here](https://github.com/microsoft/mscclpp/pkgs/container/mscclpp%2Fmscclpp).
@@ -106,16 +106,16 @@ Python 3.10 or later is required.
 # For NVIDIA platforms (specify your CUDA version)
 $ python -m pip install ".[cuda12]"
 # For AMD platforms
-$ CXX=/opt/rocm/bin/hipcc python -m pip install ".[rocm6]"
+$ CXX=/opt/rocm/bin/hipcc python -m pip install ".[rocm7]"
 ```
 
-> **Note:** A platform extra (`cuda11`, `cuda12`, `cuda13`, or `rocm6`) is required to install CuPy.
-> The CUDA extras install pre-built CuPy wheels and CUDA Python bindings. The `rocm6` extra installs CuPy from source
-> and HIP Python 6.x, which require ROCm and may take longer. Running `pip install .` without an extra will not install CuPy.
+> **Note:** A platform extra (`cuda11`, `cuda12`, `cuda13`, `rocm6`, or `rocm7`) is required to install CuPy.
+> The CUDA extras install pre-built CuPy wheels and CUDA Python bindings. The ROCm extras install CuPy from source
+> and HIP Python for the matching ROCm major version, which require ROCm and may take longer. Running `pip install .` without an extra will not install CuPy.
 
 Optional extras can be installed by specifying them in brackets. Available extras:
 - **`cuda11`**, **`cuda12`**, **`cuda13`**: Install a pre-built CuPy package and CUDA Python bindings for your CUDA version.
-- **`rocm6`**: Install CuPy from source and HIP Python 6.x for AMD ROCm platforms.
+- **`rocm6`**, **`rocm7`**: Install CuPy from source and HIP Python for AMD ROCm platforms.
 - **`benchmark`**: Install benchmark dependencies (mpi4py, prettytable, netifaces, matplotlib).
 - **`test`**: Install test dependencies (pytest, mpi4py, netifaces).
 
@@ -215,7 +215,7 @@ $ mpirun -np 16 -npernode 8 -hostfile hostfile ./bin/mp_unit_tests -ip_port 10.0
 
 ```bash
 # Install with benchmark dependencies and the appropriate CUDA/ROCm extras.
-# Replace `cuda12` with your platform: cuda11, cuda12, cuda13, or rocm6.
+# Replace `cuda12` with your platform: cuda11, cuda12, cuda13, rocm6, or rocm7.
 $ python3 -m pip install ".[cuda12,benchmark,test]"
 
 ```

include/mscclpp/gpu_data_types.hpp

@@ -16,16 +16,15 @@ using __bfloat16 = __hip_bfloat16;
 using __bfloat162 = __hip_bfloat162;
 #define __CUDA_BF16_TYPES_EXIST__
 
-// AMD FP8 support - Use fnuz types for HIP 6.0 or when HIP_FP8_TYPE_FNUZ is enabled and HIP_FP8_TYPE_OCP is not
-// enabled. Otherwise, use the standard FP8 types.
+// AMD FP8 support. The CMake MSCCLPP_ROCM_USE_FNUZ_FP8 option controls whether native FP8 aliases use FNUZ.
 #if defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR >= 6)
 #include <hip/hip_fp8.h>
 
 // Create aliases matching CUDA naming convention for cross-platform compatibility.
 // Define __FP8_E4M3_IS_FNUZ__ / __FP8_E5M2_IS_FNUZ__ when the platform-native FP8 is the
 // "fnuz" variant (no infinities, NaN-only at 0x80, bias differs from OCP). Dispatch layers
 // use these macros to throw on unsupported variants requested via DataType.
-#if (HIP_VERSION_MAJOR == 6) || (HIP_VERSION_MAJOR > 6 && HIP_FP8_TYPE_FNUZ && !HIP_FP8_TYPE_OCP)
+#if defined(MSCCLPP_ROCM_FP8_FNUZ)
 using __fp8_e4m3 = __hip_fp8_e4m3_fnuz;
 using __fp8_e5m2 = __hip_fp8_e5m2_fnuz;
 using __fp8x2_e4m3 = __hip_fp8x2_e4m3_fnuz;

pyproject.toml

@@ -37,6 +37,10 @@ rocm6 = [
     "cupy",
     "hip-python>=6,<7",
 ]
+rocm7 = [
+    "cupy",
+    "hip-python>=7,<8",
+]
 benchmark = [
     "mpi4py",
     "prettytable",

python/mscclpp_benchmark/bench_collective.py

@@ -18,7 +18,7 @@
     check_correctness as _check_correctness,
     fill_case_for_benchmark as _fill_case_for_benchmark,
 )
-from mscclpp_benchmark.gpu import capture_graph, init_runtime
+from mscclpp_benchmark.gpu import capture_graph, init_runtime, runtime_name, version
 from mscclpp_benchmark.tuner import OfflineTuner
 from mscclpp_benchmark.tuning_config import HardwareProfile, TunedConfig, TunedConfigStore, normalize_sku
 
@@ -376,7 +376,6 @@ def _measure_case(
     n_ops_per_graph: int,
 ) -> float:
     _fill_case_for_benchmark(case, comm.rank)
-    comm.comm_group.barrier()
     if comm.run(case, config) != 0:
         raise RuntimeError("algorithm returned non-zero status")
     cp.cuda.runtime.deviceSynchronize()
@@ -607,6 +606,10 @@ def main(argv: list[str] | None = None) -> None:
             )
             if comm.rank == 0:
                 print(".", end="", flush=True)
+            if runtime_name() == "hip" and version()[:2] == (7, 2):
+                # TODO: remove this after ROCm 7.2 HIP IPC export issue is fixed.
+                del case
+                comm.comm_group.barrier()
 
         if args.write_config and comm.rank == 0:
             config_store.write_path(args.write_config)

python/mscclpp_benchmark/correctness.py

@@ -62,7 +62,6 @@ def check_correctness(
     local_total = 0
     for iteration in range(niter):
         _fill_case_for_correctness(case, comm.rank, iteration)
-        comm.comm_group.barrier()
         ret = comm.run(case, config)
         cp.cuda.runtime.deviceSynchronize()
         comm.comm_group.barrier()

python/mscclpp_benchmark/gpu.py

@@ -18,6 +18,7 @@
     "graph_destroy": ("hipGraphDestroy", "cudaGraphDestroy"),
     "graph_exec_destroy": ("hipGraphExecDestroy", "cudaGraphExecDestroy"),
     "get_error_string": ("hipGetErrorString", "cudaGetErrorString"),
+    "runtime_get_version": ("hipRuntimeGetVersion", "cudaRuntimeGetVersion"),
 }
 
 
@@ -118,6 +119,21 @@ def init_runtime() -> None:
     return None
 
 
+def runtime_name() -> str:
+    return _RUNTIME.name
+
+
+def _runtime_version_raw() -> int:
+    return int(_api("runtime_get_version")()[0])
+
+
+def version() -> tuple[int, int, int]:
+    version_value = _runtime_version_raw()
+    if _RUNTIME.name == "hip":
+        return version_value // 10_000_000, (version_value // 100_000) % 100, version_value % 100_000
+    return version_value // 1000, (version_value % 1000) // 10, version_value % 10
+
+
 def capture_graph(stream: Any, capture_fn: Callable[[], None]) -> Graph:
     _api("set_device")(current_device())
     stream_ptr = _stream_ptr(stream)

python/mscclpp_benchmark/tuner.py

@@ -59,7 +59,6 @@ def tune(self, case: Any) -> TunedConfig | None:
                     if not self._check_correctness(self.comm, case, config):
                         self.comm.reset(config)
                         continue
-                    self.comm.reset(config)
                     time_us = self._measure(
                         self.comm,
                         case,

python/requirements_rocm7.txt

@@ -0,0 +1,11 @@
+mpi4py
+cupy
+prettytable
+netifaces
+pytest
+numpy
+matplotlib
+sortedcontainers
+blake3
+pybind11
+hip-python>=7,<8