Problems installing gromacs 2025.2 on AMD 9070XT GPU with HIP

User discussions
21

Hello, thank you for the report. Did you build your version of GROMACS from the 2025.3 release archive, or from the experimental branch?
It sounds quite slow, and if this is from the feature branch I’ll need to investigate

22

Thanks for your reply, I have tried two versions. First, the official version of Gromacs 2025.3 compiled with ROCM and AdaptiveCPP, which is exactly the situation I described above. Additionally, I also tried the 4947-hip-feature-enablement Gromacs based on HIP. Not only is its speed similar to the one mentioned above, but it also frequently causes my computer to freeze.

23

Hi Paul (@pbauer),

Following up on the RX 9070 XT slowness reported earlier in this thread (~1/3 of an RTX 3090): I spent some time characterising it on the same PC with an RX 7800 XT as a control. I believe there are two independent gfx1201 issues, one in float atomics and one in HIP queue scheduling. Working around both makes the 9070 XT go from 3.6x slower than the 7800 XT to faster than it.

TL;DR

  1. Global float atomicAdd (global_atomic_add_f32) is capped at ~3.5 Gops/s on gfx1201 — roughly one op per clock for the whole device, independent of contention. Integer atomics run at 33–44 Gops/s, and an equivalent CAS loop at 18–25 Gops/s. This makes PME spread ~10x slower and the non-bonded and bonded kernels 3–5x slower than on gfx1101. OpenMM is not affected because it accumulates forces in 64-bit fixed point.
  2. Per-dispatch penalty when ≥2 other streams hold pending hipStreamWaitEvents: a tiny kernel goes from 3.3 µs to 42 µs. This matches ROCm issue ROCm/legacy-rocm-build#6685, which was reported on an R9700. GPU_MAX_HW_QUEUES=2 removes it.
  3. A small patch that uses a CAS loop for global float atomics only on gfx120x, combined with GPU_MAX_HW_QUEUES=2, gives:
ns/day, all on GPU RX 7800 XT RX 9070 XT stock 9070 XT + CAS patch 9070 XT + CAS patch + HWQ=2
small, 23,873 atoms 718 196 463 729
medium, 71,580 atoms 347 95 240 370
large, 144,910 atoms 190 53 173 212

make check passes 106/106 (including the 43 GPU tests) with the patch.

System

  • Ryzen 7 5700X, B550M-K, 32 GB; Ubuntu 24.04.5, kernel 7.0.0-38-generic
  • ROCm 10.1 (HIP 7.16.26385), amdgpu-install 31.60
  • GPU A: RX 7800 XT (gfx1101), MES 0x52, SMU 80.74.0
  • GPU B: RX 9070 XT (gfx1201, Gigabyte Gaming OC)
    • stock Ubuntu firmware: SMU 104.76.0, MES 0x89
    • amdgpu-dkms 7.1.9 + amdgpu-dkms-firmware 31.60: SMU 104.80.0, MES 0x93
    • results are identical with both firmware sets
  • GROMACS: branch 4947-hip-feature-enablement (2026.4-2026_HIP_ENABLEMENT-dev), one binary for both cards:
    -DGMX_GPU=HIP -DGMX_HIP_TARGET_ARCH="gfx1101,gfx1201" -DGMX_ENABLE_AMD_RDNA_SUPPORT=ON -DGMX_GPU_FFT_LIBRARY=VkFFT -DGMX_BUILD_OWN_FFTW=ON, no hwloc, no MPI, internal BLAS/LAPACK
  • OpenMM 8.6 (master) HIP platform, mixed precision, as a reference

Benchmarks: lysozyme 1AKI, AMBER99SB-ILDN + TIP3P, 0.15 M NaCl, PME 1.0 nm cut-offs, dt 2 fs, V-rescale + C-rescale.

  • small: dodecahedron, 1.0 nm; medium: cubic, 2.0 nm; large: cubic, 3.2 nm. The same .tpr files are used on both cards.
  • gmx mdrun -ntmpi 1 -ntomp 8 -pin on -nb gpu -pme gpu -bonded gpu -update gpu -notunepme -resethway.
  • small 30k steps, medium 20k, large 10k.

1. Float atomics on gfx1201

The compiler emits the native instruction as expected: global_atomic_add_f32 on both gfx1101 and gfx1201. Note that on gfx1201 it is emitted even with -mno-unsafe-fp-atomics, whereas gfx1101 falls back to CAS without the flag. What differs is the throughput. Measured with atomics_bench.hip (attached): 524k threads × 8 adds into S distinct addresses, values in Gops/s.

S (addresses) store f32 atomicAdd f32 CAS loop u32 atomicAdd u64 atomicAdd
256 11.6 0.34 0.62 21.7 24.0
8,192 43.9 2.7 16.1 49.0 45.7
262,144 44.1 3.49 24.8 44.2 43.2
4,194,304 20.0 3.49 5.2 11.9 9.3

The f32 kernel takes exactly 1.200 ms for both 262k and 4M addresses. That is a fixed rate of about 1 op/clock at ~3.3 GHz, not a contention effect, and it is unchanged with the newer firmware. (RDNA3 control numbers for this benchmark: to be added.)

Kernel times from rocprofv3, small system. With GPU_MAX_HW_QUEUES=1 there is no kernel overlap, so these are pure execution times; bonded is from the default-queue run.

kernel 7800 XT 9070 XT stock 9070 XT + CAS patch
nbnxmKernel 94 µs 305 µs 106 µs
pmeSplineAndSpreadKernel 26 µs 258 µs 42 µs
bonded_kernel_gpu 11 µs 61 µs 15 µs
VkFFT_main 7 µs 15 µs 8 µs
OpenMM computeNonbonded 106 µs 74 µs —

2. Queue / pending-wait penalty

Measured with queue_repro.hip (attached): 200 tiny kernels on one stream while K other streams each hold a pending hipStreamWaitEvent. Values are µs per kernel.

pending waits 7800 XT 9070 XT 9070 XT, HWQ=1 9070 XT, HWQ=2
0 2.76 3.38 3.33 2.95
1 2.74 3.31 3.33 3.27
2 2.74 42.1 3.33 3.27
4 2.72 42.3 3.45 3.26

The same happens in eager mode. In mdrun the default configuration on gfx1201 is also noisy: two identical patched runs gave 463 and 388 ns/day, consistent with spawn-time stream-to-queue mapping.

Effect of GPU_MAX_HW_QUEUES (small, all on GPU, ns/day):

default HWQ=1 HWQ=2
7800 XT 718 579 668
9070 XT stock 196 233 247
9070 XT + CAS patch 463 599 729

So HWQ=2 is a clear win on gfx1201, but it costs 1–7 % on gfx1101 (and on OpenMM HWQ=1 costs 10–20 % on both), so it should not become a global default.

3. The patch

The patch adds a helper gmxAtomicAddF32() (header attached):

  • on __gfx1200__ / __gfx1201__, a CAS loop for non-LDS addresses (__builtin_amdgcn_is_shared);
  • everywhere else, plain atomicAdd, so gfx1101 code is unchanged.

It replaces the global float atomics in:

  • pme_spread_hip.cpp (grid flush)
  • nbnxm_hip_kernel_utils.h (f, fshift, energies)
  • nbnxm_hip_kernel_sum.cpp
  • the float3 overloads in vectype_ops_cuda_hip_shared_declarations.h (behind a macro, so CUDA is unchanged)
  • vectype_ops_hip.h, amdFastAtomicAddForce
  • listed-forces and PME-solve energy/virial sums

The diff is attached. I’m not suggesting this exact form for upstream — it’s a proof of concept. But it shows that the gap is entirely float atomics plus queue scheduling. The remaining difference in PME spread (42 vs 26 µs) may come from the LDS float atomics I left native.

Questions

  1. Is the ~1 op/clock rate of global_atomic_add_f32 on gfx12 known or expected? If it is, would a gfx12-specific CAS path (or another accumulation scheme) be acceptable in the HIP backend?
  2. For the queue penalty: would it make sense for GROMACS to limit the number of streams with pending cross-stream waits on gfx12, or to document GPU_MAX_HW_QUEUES=2 for RDNA4 until ROCm addresses it?

I’m reporting the atomic throughput to ROCm separately and adding the 9070 XT data to #6685. I’m happy to test patches or collect more data — both cards are in the same machine.

Thanks for all the work on the HIP backend!

Bryan

Attached: both reproducers (queue_repro.hip, atomics_bench.hip), the helper header and the diff, concatenated into one .txt file because of the new-user upload/link limits.

Note: the include is placed before gmxpre.h in the .cpp files; this is just a quick PoC and would need the usual include ordering for an actual MR.

gfx1201_reproducers_and_patch.txt (24.8 KB)