[ROCm] Add native HIP backend for AMD GPUs

[jeffdaily]

Summary

This adds a native HIP backend (DEVICE=HIP) so AutoDock-GPU runs on AMD GPUs through ROCm, by porting the maintained CUDA kernels (cuda/) to HIP. A new Makefile.Hip builds cuda/kernels.cu with hipcc, and a new compat header cuda/cuda_to_hip.h aliases the CUDA spellings to HIP so the CUDA-spelled device sources compile unchanged. AMD users get the same kernels NVIDIA users get; the existing OpenCL backend is a separate code path. The CUDA and OpenCL/CPU build paths are byte-for-byte unchanged.

Build it with make DEVICE=HIP NUMWI=<N> HIP_ARCH=<arch>. The offload arch is configurable (HIP_ARCH ?= gfx90a), so other AMD targets build with HIP_ARCH=<arch> and no source edit. The build is documented in the README alongside DEVICE=CUDA.

What is implemented

• Makefile.Hip + cuda/cuda_to_hip.h: a HIP build of the CUDA kernels via a single compat header (the CUDA cuda/ kernels are reused as-is).
• Wave64 reductions: AMD's wavefront is 64 lanes on gfx90a, so the host derives warpmask/warpbits from the runtime warpSize (31/5 on NVIDIA, 63/6 on gfx90a) and the device reduction macros gain a stride-32 step under a 64-lane wavefront (a 6-step XOR-butterfly all-reduce; CUDA keeps the upstream 5-step). Per-warp partials recombine block-wide via atomicAdd. __shfl_sync/__any_sync use a 64-bit full mask, and two inline-PTX mov.b64 helpers become an equivalent memcpy bit-cast.
• TENSOR=ON (rocWMMA): the optional tensor-core build (USE_NVTENSOR, NUMWI >= 64) accelerates the energy/gradient sum-reduction with the GPU matrix cores through rocWMMA, mirroring the CUDA tensor-core path. CDNA (gfx90a) uses the native fp32 WMMA; RDNA, which has no fp32 WMMA, uses a bf16-split emulation. Results match the standard reduction.
• Windows: the HIP device pass defines _WIN32 and __HIP_DEVICE_COMPILE__ but not a host-architecture macro, so the Windows SDK processthreadsapi.h errors with "No Target Architecture". It is guarded with !__HIP_DEVICE_COMPILE__ (host compilation only), with a device-side stub processid() for header-inlined references.

Validation

The bundled streptavidin-biotin case (input/1stp) was docked with -nrun 10 and fixed seeds on a single pinned GPU, for both the default and TENSOR=ON builds, recovering the native biotin pose (all 10 runs in one cluster):

GPU	Arch	OS / ROCm	1stp result (64wi, seed 42)
Instinct MI250X	gfx90a (CDNA2, wave64)	Linux, ROCm 7.2.1	-8.28 kcal/mol, reference RMSD 0.48 A
Radeon Pro W7800	gfx1100 (RDNA3, wave32)	Linux, ROCm 7.2.1	-8.28 kcal/mol, reference RMSD 0.40 A
Radeon RX 9070 XT	gfx1201 (RDNA4, wave32)	Windows, ROCm 7.14	-8.28 kcal/mol, reference RMSD 0.59 A

Results reproduce across seeds (seed 7: -8.37 kcal/mol on each) and block sizes (NUMWI=128, which also exercises the cross-warp final reduction); best energies agree within 0.05 kcal/mol across architectures and reference RMSD stays sub-0.6 A. The TENSOR=ON (rocWMMA) build produces the same poses and energies as the default reduction. This confirms the GA min-reduction (kernel4) and the energy/gradient sum-reductions are numerically correct on both wave64 and wave32.

The CUDA/OpenCL/CPU paths are unchanged and not affected by this backend.

This work was authored with the assistance of Claude, an AI assistant.

[Young-TW]

TL;DR — Validated on gfx1201 / RDNA4 (RX 9070 XT, Linux, ROCm 7.13.0): builds clean, results
correct, NUMWI=64 is the sweet spot. Profiling shows the run is entirely dominated by
gpu_gradient_minAD_kernel, which is VGPR-bound (152 VGPR → 9/16 waves); a compiler-flag
attempt to relieve that didn't help. Details below.

Testing report: HIP backend on RDNA4 (gfx1201 / RX 9070 XT)

Gave this branch a spin on a consumer AMD card — nice to confirm the arch-configurable design works
well beyond the gfx90a default.

Setup: AMD Radeon RX 9070 XT (gfx1201, RDNA4, wave32), ROCm 7.13.0 (clang/LLVM 23), Linux.
Built with make DEVICE=HIP NUMWI=<N> HIP_ARCH=gfx1201. Builds cleanly — only benign
-Wunused-result warnings from the cudaDeviceReset() mapping through the nodiscard HIP headers.

Correctness

All work-group sizes converge to the same energy on the bundled inputs — e.g. 3ce3 → −14.58,
7cpa → ~−21 kcal/mol. WARP_SIZE correctly resolves to 32 on gfx1201 (non-__GFX9__), so the
warp-shuffle reductions are right on RDNA.

NUMWI sweep — best-of-N end-to-end job-set time (relative to NUMWI=64; higher % = slower)

NUMWI	waves/block	3ce3 (5 tors.)	7cpa (15 tors.)
32	1	+16%	+18%
64	2	fastest	fastest
128	4	+17%	+7%
256	8	+64%	+55%

Absolute anchor for cross-GPU comparison: 7cpa at NUMWI=64 ≈ 8.6 s end-to-end job-set
(-nrun 100, best of 5). NUMWI=64 is the sweet spot on RDNA4. Larger ligands narrow the 128 gap
(more work per task, as expected), but 64 still wins across the board; 256 is consistently poor.

Kernel profiling (rocprofv3 --kernel-trace, NUMWI=64, 7cpa)

gpu_gradient_minAD_kernel (ADADELTA local search) completely dominates:

Kernel	Calls	GPU time	%
gpu_gradient_minAD_kernel	144	7.54 s	99.6%
gpu_gen_and_eval_newpops_kernel	144	27.5 ms	0.36%
gpu_sum_evals_kernel	145	0.49 ms	<0.01%
gpu_calc_initpop_kernel	1	0.23 ms	<0.01%

Resource usage for the hot kernel: 152 VGPR, 0 scratch spill, 7680 B LDS/block. That caps
occupancy at 9 of 16 waves/SIMD — it's VGPR-bound, not LDS-bound. This also explains the
sweep: at NUMWI=256 (8 waves/block) occupancy collapses, hence the ~55–64% regression.

VGPR-reduction attempt (no luck, reporting for completeness)

Tried to trade registers for occupancy via __attribute__((amdgpu_waves_per_eu(10/12/16))) on the
hot kernel (this LLVM has no -mllvm -amdgpu-waves-per-eu global flag). VGPR stayed pinned at 152
with zero spill and no runtime change — the existing
__launch_bounds__(NUM_OF_THREADS_PER_BLOCK, 1024 / NUM_OF_THREADS_PER_BLOCK) already requests max
occupancy, and the allocator correctly declines to go lower (spilling would be net-negative). So
152 VGPR looks intrinsic to the kernel's live-set; further gains would need source-level changes
(recompute vs. store, narrower types), not compiler flags.

Minor note (pre-existing, unrelated to this PR)

input/7cpa doesn't run out of the box: its .maps.fld references rec.A.map (etc.) but the map
files are named 7cpa_rec.A.map. I symlinked around it to test. Might be worth fixing separately.

Bottom line: HIP backend works great on RDNA4 / gfx1201 — correct results, clean build,
sensible perf. Suggest documenting NUMWI=64 as the default guidance for consumer AMD cards. 👍

[jeffdaily]

Thanks for testing this on RDNA4 and profiling the hot kernel. The NUMWI=64 guidance is now in the build docs (2966dcd); it was also the fastest setting in our measurements on the datacenter default, so the recommendation holds across both. The VGPR-bound profile for gpu_gradient_minAD_kernel matches what we saw here: the 152-VGPR live set is intrinsic to the kernel, so occupancy gains would need source-level changes rather than compiler flags. On 7cpa, the .maps.fld referencing rec.A.map while the shipped maps are named 7cpa_rec.A.map is a pre-existing issue in the input data (this PR does not touch input/), so it is unrelated here.