diff --git a/.azure-pipelines/codecov.yml b/.azure-pipelines/codecov.yml
index c4abeaa78..5f3e531e0 100644
--- a/.azure-pipelines/codecov.yml
+++ b/.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)
diff --git a/.azure-pipelines/rccl-api-test.yml b/.azure-pipelines/rccl-api-test.yml
index 43841079e..fc793c88d 100644
--- a/.azure-pipelines/rccl-api-test.yml
+++ b/.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)
diff --git a/.azure-pipelines/templates/rccl-test.yml b/.azure-pipelines/templates/rccl-test.yml
index 63788ac27..65f1c9848 100644
--- a/.azure-pipelines/templates/rccl-test.yml
+++ b/.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:
diff --git a/.azure-pipelines/ut.yml b/.azure-pipelines/ut.yml
index 6b8c9eda1..0cabdaeb1 100644
--- a/.azure-pipelines/ut.yml
+++ b/.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)
diff --git a/.github/workflows/codeql-analysis.yml b/.github/workflows/codeql-analysis.yml
index fb0651415..982d8633a 100644
--- a/.github/workflows/codeql-analysis.yml
+++ b/.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
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 49154e0b0..928220123 100644
--- a/CMakeLists.txt
+++ b/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")
diff --git a/docker/base-dev-x.dockerfile b/docker/base-dev-x.dockerfile
index 474362020..89ea1ec7e 100644
--- a/docker/base-dev-x.dockerfile
+++ b/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 && \
diff --git a/docker/build.sh b/docker/build.sh
index b84eac9ad..73434cb92 100755
--- a/docker/build.sh
+++ b/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
diff --git a/docs/quickstart.md b/docs/quickstart.md
index 320a2db78..6bda2877b 100644
--- a/docs/quickstart.md
+++ b/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]"
 
 ```
diff --git a/include/mscclpp/gpu_data_types.hpp b/include/mscclpp/gpu_data_types.hpp
index 4a16628e3..0c1e6479e 100644
--- a/include/mscclpp/gpu_data_types.hpp
+++ b/include/mscclpp/gpu_data_types.hpp
@@ -16,8 +16,7 @@ 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>
 
@@ -25,7 +24,7 @@ using __bfloat162 = __hip_bfloat162;
 // 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;
diff --git a/pyproject.toml b/pyproject.toml
index b35b1b3a9..89ee4a731 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -37,6 +37,10 @@ rocm6 = [
     "cupy",
     "hip-python>=6,<7",
 ]
+rocm7 = [
+    "cupy",
+    "hip-python>=7,<8",
+]
 benchmark = [
     "mpi4py",
     "prettytable",
diff --git a/python/mscclpp_benchmark/bench_collective.py b/python/mscclpp_benchmark/bench_collective.py
index c526438da..0cb11739a 100644
--- a/python/mscclpp_benchmark/bench_collective.py
+++ b/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)
diff --git a/python/mscclpp_benchmark/correctness.py b/python/mscclpp_benchmark/correctness.py
index 0d9ab5c13..15de2f5ae 100644
--- a/python/mscclpp_benchmark/correctness.py
+++ b/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()
diff --git a/python/mscclpp_benchmark/gpu.py b/python/mscclpp_benchmark/gpu.py
index 1309a504e..248065efb 100644
--- a/python/mscclpp_benchmark/gpu.py
+++ b/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)
diff --git a/python/mscclpp_benchmark/tuner.py b/python/mscclpp_benchmark/tuner.py
index 8df3259b9..42c15ab57 100644
--- a/python/mscclpp_benchmark/tuner.py
+++ b/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,
diff --git a/python/requirements_rocm7.txt b/python/requirements_rocm7.txt
new file mode 100644
index 000000000..982f9ae6b
--- /dev/null
+++ b/python/requirements_rocm7.txt
@@ -0,0 +1,11 @@
+mpi4py
+cupy
+prettytable
+netifaces
+pytest
+numpy
+matplotlib
+sortedcontainers
+blake3
+pybind11
+hip-python>=7,<8
diff --git a/src/core/gpu_ipc_mem.cc b/src/core/gpu_ipc_mem.cc
index 0f58ed202..49dfcf822 100644
--- a/src/core/gpu_ipc_mem.cc
+++ b/src/core/gpu_ipc_mem.cc
@@ -81,8 +81,6 @@ bool isFabricMemHandleAvailable() {
 #endif  // !(CUDA_NVLS_API_AVAILABLE)
 }
 
-#if defined(MSCCLPP_DEVICE_HIP)
-
 // Custom hash and equality for cudaIpcMemHandle_t
 struct CudaIpcMemHandleHash {
   size_t operator()(const cudaIpcMemHandle_t& handle) const {
@@ -97,43 +95,59 @@ struct CudaIpcMemHandleEqual {
   }
 };
 
-static std::unordered_map<cudaIpcMemHandle_t, void*, CudaIpcMemHandleHash, CudaIpcMemHandleEqual>
+static std::unordered_map<cudaIpcMemHandle_t, std::weak_ptr<void>, CudaIpcMemHandleHash, CudaIpcMemHandleEqual>
     openCudaIpcMemHandleMap;
 
 static std::mutex openCudaIpcMemHandleMapMutex;
 
-// Cache open ipc handles to avoid opening multiple times (ROCm may exceed system limit on vm.max_map_count).
-static inline cudaError_t cudaIpcOpenMemHandleWrapper(void** addr, cudaIpcMemHandle_t ipcHandle) {
+static inline cudaError_t cudaIpcOpenMemHandleWrapper(std::shared_ptr<void>& basePtr, cudaIpcMemHandle_t ipcHandle) {
+#if defined(MSCCLPP_DEVICE_HIP)
+  // Cache open ipc handles to avoid opening multiple times (ROCm may exceed system limit on vm.max_map_count).
   std::lock_guard<std::mutex> lock(openCudaIpcMemHandleMapMutex);
   auto it = openCudaIpcMemHandleMap.find(ipcHandle);
   if (it != openCudaIpcMemHandleMap.end()) {
-    *addr = it->second;
-    return cudaSuccess;
-  }
-  cudaError_t err = cudaIpcOpenMemHandle(addr, ipcHandle, cudaIpcMemLazyEnablePeerAccess);
-  if (err == cudaSuccess) {
-    openCudaIpcMemHandleMap[ipcHandle] = *addr;
+    basePtr = it->second.lock();
+    if (basePtr) {
+      return cudaSuccess;
+    }
+    openCudaIpcMemHandleMap.erase(it);
   }
-  return err;
-}
-
-static inline cudaError_t cudaIpcCloseMemHandleWrapper(void* addr, cudaIpcMemHandle_t ipcHandle) {
-  std::lock_guard<std::mutex> lock(openCudaIpcMemHandleMapMutex);
-  openCudaIpcMemHandleMap.erase(ipcHandle);
-  return cudaIpcCloseMemHandle(addr);
-}
-
-#else  // !defined(MSCCLPP_DEVICE_HIP)
 
-static inline cudaError_t cudaIpcOpenMemHandleWrapper(void** addr, cudaIpcMemHandle_t ipcHandle) {
-  return cudaIpcOpenMemHandle(addr, ipcHandle, cudaIpcMemLazyEnablePeerAccess);
-}
-
-static inline cudaError_t cudaIpcCloseMemHandleWrapper(void* addr, [[maybe_unused]] cudaIpcMemHandle_t ipcHandle) {
-  return cudaIpcCloseMemHandle(addr);
-}
+  void* rawBasePtr = nullptr;
+  cudaError_t err = cudaIpcOpenMemHandle(&rawBasePtr, ipcHandle, cudaIpcMemLazyEnablePeerAccess);
+  if (err != cudaSuccess) {
+    return err;
+  }
 
+  basePtr = std::shared_ptr<void>(rawBasePtr, [ipcHandle](void* ptr) {
+    {
+      std::lock_guard<std::mutex> lock(openCudaIpcMemHandleMapMutex);
+      openCudaIpcMemHandleMap.erase(ipcHandle);
+    }
+    cudaError_t err = cudaIpcCloseMemHandle(ptr);
+    if (err != cudaSuccess) {
+      WARN(GPU, "Failed to close CUDA IPC handle at pointer ", ptr, ": ", cudaGetErrorString(err));
+    }
+  });
+  openCudaIpcMemHandleMap[ipcHandle] = basePtr;
+  return cudaSuccess;
+#else   // !defined(MSCCLPP_DEVICE_HIP)
+  void* rawBasePtr = nullptr;
+  cudaError_t err = cudaIpcOpenMemHandle(&rawBasePtr, ipcHandle, cudaIpcMemLazyEnablePeerAccess);
+  if (err != cudaSuccess) {
+    return err;
+  }
+
+  basePtr = std::shared_ptr<void>(rawBasePtr, [](void* ptr) {
+    cudaError_t err = cudaIpcCloseMemHandle(ptr);
+    if (err != cudaSuccess) {
+      WARN(GPU, "Failed to close CUDA IPC handle at pointer ", ptr, ": ", cudaGetErrorString(err));
+      (void)cudaGetLastError();
+    }
+  });
+  return cudaSuccess;
 #endif  // !defined(MSCCLPP_DEVICE_HIP)
+}
 
 void GpuIpcMemHandle::deleter(GpuIpcMemHandle* handle) {
   if (handle) {
@@ -172,7 +186,8 @@ UniqueGpuIpcMemHandle GpuIpcMemHandle::create(const CUdeviceptr ptr) {
   if (err == cudaSuccess) {
     handle->typeFlags |= GpuIpcMemHandle::Type::RuntimeIpc;
   } else {
-    (void)cudaGetLastError();
+    WARN(GPU, "Failed to create runtime CUDA IPC handle for pointer ", (void*)basePtr,
+         ": error=", static_cast<int>(err), " (", std::string(cudaGetErrorString(err)), ")");
   }
 
 #if !defined(MSCCLPP_DEVICE_HIP)  // Remove when HIP fully supports virtual memory management APIs
@@ -345,17 +360,10 @@ std::shared_ptr<void> GpuIpcMem::map() {
 
   if (type_ == GpuIpcMemHandle::Type::RuntimeIpc) {
     // RuntimeIpc: Open handle and return shared_ptr with cleanup in deleter
-    void* basePtr = nullptr;
-    MSCCLPP_CUDATHROW(cudaIpcOpenMemHandleWrapper(&basePtr, handle_.runtimeIpc.handle));
-    void* dataPtr = static_cast<void*>(static_cast<char*>(basePtr) + handle_.offsetFromBase);
-    cudaIpcMemHandle_t ipcHandle = handle_.runtimeIpc.handle;
-    return std::shared_ptr<void>(dataPtr, [self = shared_from_this(), basePtr, ipcHandle](void*) {
-      cudaError_t err = cudaIpcCloseMemHandleWrapper(basePtr, ipcHandle);
-      if (err != cudaSuccess) {
-        WARN(GPU, "Failed to close CUDA IPC handle at pointer ", basePtr, ": ", cudaGetErrorString(err));
-        (void)cudaGetLastError();
-      }
-    });
+    std::shared_ptr<void> basePtr;
+    MSCCLPP_CUDATHROW(cudaIpcOpenMemHandleWrapper(basePtr, handle_.runtimeIpc.handle));
+    void* rawDataPtr = static_cast<void*>(static_cast<char*>(basePtr.get()) + handle_.offsetFromBase);
+    return std::shared_ptr<void>(basePtr, rawDataPtr);
   }
 
   size_t pageSize = getpagesize();
diff --git a/src/core/registered_memory.cc b/src/core/registered_memory.cc
index f464de2a6..c82233cef 100644
--- a/src/core/registered_memory.cc
+++ b/src/core/registered_memory.cc
@@ -97,11 +97,14 @@ MSCCLPP_API_CPP std::vector<char> RegisteredMemory::serialize() const {
   for (auto& entry : pimpl_->transportInfos) {
     detail::serialize(result, entry.transport);
     if (entry.transport == Transport::CudaIpc) {
+      if (entry.gpuIpcMemHandle.typeFlags == GpuIpcMemHandle::Type::None) {
+        THROW(GPU, Error, ErrorCode::InternalError, "GpuIpcMemHandle type is None");
+      }
       detail::serialize(result, entry.gpuIpcMemHandle);
     } else if (AllIBTransports.has(entry.transport)) {
       detail::serialize(result, entry.ibMrInfo);
     } else {
-      throw Error("Unknown transport", ErrorCode::InternalError);
+      THROW(GPU, Error, ErrorCode::InternalError, "Unknown transport");
     }
   }
   return result;
diff --git a/src/ext/collectives/allreduce/allreduce_fullmesh.cu b/src/ext/collectives/allreduce/allreduce_fullmesh.cu
index 24d2a31c2..725bdb0df 100644
--- a/src/ext/collectives/allreduce/allreduce_fullmesh.cu
+++ b/src/ext/collectives/allreduce/allreduce_fullmesh.cu
@@ -10,7 +10,7 @@ namespace mscclpp {
 namespace collective {
 
 template <ReduceOp OpType, typename T, typename AccumT = T>
-__global__ void __launch_bounds__(512, 1)
+__global__ void __launch_bounds__(1024, 1)
     allreduceFullmesh(T* buff, T* scratch, T* resultBuff, DeviceHandle<MemoryChannel>* memoryChannels,
                       DeviceHandle<MemoryChannel>* memoryOutChannels, size_t channelOutDataOffset, int rank,
                       int nRanksPerNode, int worldSize, size_t nelems) {
@@ -194,17 +194,6 @@ CommResult AllreduceFullmesh::allreduceKernelFunc(
     MSCCLPP_CUTHROW(cuMemGetAddressRange(&recvBasePtr, &recvBytes, (CUdeviceptr)output));
     channelOutOffset = (char*)output - (char*)recvBasePtr;
   }
-  std::shared_ptr<DeviceHandle<MemoryChannel>> inputChannelHandles;
-  if (this->memoryChannelsMap_.find(input) != this->memoryChannelsMap_.end()) {
-    inputChannelHandles = this->memoryChannelsMap_[input].second;
-  } else {
-    RegisteredMemory localMemory = comm_->registerMemory(const_cast<void*>(input), inputSize, Transport::CudaIpc);
-    std::vector<MemoryChannel> channels =
-        setupMemoryChannels(this->conns_, this->inputScratchSemaphores_, this->remoteScratchMemories_, localMemory,
-                            nChannelsPerConnection_);
-    this->memoryChannelsMap_[input] = std::make_pair(channels, setupMemoryChannelDeviceHandles(channels));
-  }
-  inputChannelHandles = this->memoryChannelsMap_[input].second;
 
   AllreduceFunc allreduce = dispatch<AllreduceAllconnectAdapter>(op, dtype, accumDtype);
   if (!allreduce) {
@@ -220,10 +209,12 @@ CommResult AllreduceFullmesh::allreduceKernelFunc(
   if (numBlocksAndThreads.first == 0 || numBlocksAndThreads.second == 0) {
     numBlocksAndThreads = {35, 512};
   }
+  auto inputChannelDeviceHandles = ctx->memoryChannelDeviceHandles.get();
+  auto outputChannelDeviceHandles = inputChannelDeviceHandles + ctx->memoryChannels.size() / 2;
   cudaError_t error =
-      allreduce(input, this->scratchBuffer_, output, inputChannelHandles.get(), ctx->memoryChannelDeviceHandles.get(),
-                nullptr, nullptr, 0, channelOutOffset, 0, ctx->rank, ctx->nRanksPerNode, ctx->workSize, inputSize,
-                stream, nullptr, 0, 0, numBlocksAndThreads.first, numBlocksAndThreads.second);
+      allreduce(input, this->scratchBuffer_, output, inputChannelDeviceHandles, outputChannelDeviceHandles, nullptr,
+                nullptr, 0, channelOutOffset, 0, ctx->rank, ctx->nRanksPerNode, ctx->workSize, inputSize, stream,
+                nullptr, 0, 0, numBlocksAndThreads.first, numBlocksAndThreads.second);
   if (error != cudaSuccess) {
     WARN("AllreduceAllconnect failed with error: %s", cudaGetErrorString(error));
     return CommResult::CommUnhandledCudaError;
@@ -231,20 +222,21 @@ CommResult AllreduceFullmesh::allreduceKernelFunc(
   return CommResult::CommSuccess;
 }
 
-AlgorithmCtxKey AllreduceFullmesh::generateAllreduceContextKey(const void*, void* output, size_t, DataType,
+AlgorithmCtxKey AllreduceFullmesh::generateAllreduceContextKey(const void* input, void* output, size_t size, DataType,
                                                                bool symmetricMemory) {
-  static int tag = 0;
-  size_t recvBytes;
-  CUdeviceptr recvBasePtr;
-  MSCCLPP_CUTHROW(cuMemGetAddressRange(&recvBasePtr, &recvBytes, (CUdeviceptr)output));
   symmetricMemory_ = symmetricMemory;
   if (!symmetricMemory_) {
-    return AlgorithmCtxKey{nullptr, (void*)recvBasePtr, 0, recvBytes, tag++};
+    return AlgorithmCtxKey{const_cast<void*>(input), output, size, size, 0};
   }
-  return AlgorithmCtxKey{nullptr, (void*)recvBasePtr, 0, recvBytes, 0};
+
+  size_t sendBytes, recvBytes;
+  CUdeviceptr sendBasePtr, recvBasePtr;
+  MSCCLPP_CUTHROW(cuMemGetAddressRange(&sendBasePtr, &sendBytes, (CUdeviceptr)input));
+  MSCCLPP_CUTHROW(cuMemGetAddressRange(&recvBasePtr, &recvBytes, (CUdeviceptr)output));
+  return AlgorithmCtxKey{(void*)sendBasePtr, (void*)recvBasePtr, sendBytes, recvBytes, 0};
 }
 
-std::shared_ptr<void> AllreduceFullmesh::initAllreduceContext(std::shared_ptr<Communicator> comm, const void*,
+std::shared_ptr<void> AllreduceFullmesh::initAllreduceContext(std::shared_ptr<Communicator> comm, const void* input,
                                                               void* output, size_t size, DataType) {
   auto ctx = std::make_shared<AlgorithmCtx>();
   ctx->rank = comm->bootstrap()->getRank();
@@ -263,8 +255,19 @@ std::shared_ptr<void> AllreduceFullmesh::initAllreduceContext(std::shared_ptr<Co
   }
   RegisteredMemory localMemory = comm->registerMemory((void*)recvBasePtr, recvBytes, Transport::CudaIpc);
   ctx->registeredMemories = setupRemoteMemories(comm, ctx->rank, localMemory);
-  ctx->memoryChannels = setupMemoryChannels(this->conns_, ctx->memorySemaphores, ctx->registeredMemories, localMemory,
-                                            nChannelsPerConnection_);
+
+  RegisteredMemory inputMemory = comm->registerMemory(const_cast<void*>(input), size, TransportFlags());
+  std::vector<MemoryChannel> inputMemoryChannels = setupMemoryChannels(
+      this->conns_, this->inputScratchSemaphores_, this->remoteScratchMemories_, inputMemory, nChannelsPerConnection_);
+  std::vector<MemoryChannel> outputMemoryChannels = setupMemoryChannels(
+      this->conns_, this->outputSemaphores_, ctx->registeredMemories, localMemory, nChannelsPerConnection_);
+  ctx->memoryChannels.reserve(inputMemoryChannels.size() + outputMemoryChannels.size());
+  for (auto& channel : inputMemoryChannels) {
+    ctx->memoryChannels.emplace_back(std::move(channel));
+  }
+  for (auto& channel : outputMemoryChannels) {
+    ctx->memoryChannels.emplace_back(std::move(channel));
+  }
   ctx->memoryChannelDeviceHandles = setupMemoryChannelDeviceHandles(ctx->memoryChannels);
   return ctx;
 }
diff --git a/src/ext/collectives/include/allreduce/allreduce_fullmesh.hpp b/src/ext/collectives/include/allreduce/allreduce_fullmesh.hpp
index a54352b3f..e0c63a3d3 100644
--- a/src/ext/collectives/include/allreduce/allreduce_fullmesh.hpp
+++ b/src/ext/collectives/include/allreduce/allreduce_fullmesh.hpp
@@ -30,8 +30,6 @@ class AllreduceFullmesh : public mscclpp::AlgorithmBuilder {
   std::vector<std::shared_ptr<MemoryDevice2DeviceSemaphore>> inputScratchSemaphores_;
   std::vector<RegisteredMemory> remoteScratchMemories_;
   RegisteredMemory localScratchMemory_;
-  std::unordered_map<const void*, std::pair<std::vector<MemoryChannel>, std::shared_ptr<DeviceHandle<MemoryChannel>>>>
-      memoryChannelsMap_;
   bool symmetricMemory_ = false;
 };
 }  // namespace collective
diff --git a/test/deploy/setup.sh b/test/deploy/setup.sh
index 2a88a3100..fbcdc119d 100644
--- a/test/deploy/setup.sh
+++ b/test/deploy/setup.sh
@@ -68,7 +68,9 @@ elif [[ "${CUDA_VERSION}" == *"12."* ]]; then
 elif [[ "${CUDA_VERSION}" == *"13."* ]]; then
     pip3 install ".[cuda13,benchmark,test]"
 elif [ "${PLATFORM}" == "rocm" ]; then
-    pip3 install ".[rocm6,benchmark,test]"
+    ROCM_VERSION=$(cat /opt/rocm/.info/version)
+    ROCM_MAJOR="${ROCM_VERSION%%.*}"
+    pip3 install ".[rocm${ROCM_MAJOR},benchmark,test]"
 else
     pip3 install ".[benchmark,test]"
 fi