[ROCm] Add native HIP backend for AMD GPUs

Makefile

@@ -26,6 +26,16 @@
 
 OVERLAP = ON
 
+# HIP/ROCm backend: native AMD GPU path ported from the CUDA kernels. Selected
+# explicitly with DEVICE=HIP; bypasses the nvcc CUDA probe entirely.
+ifeq ($(DEVICE), HIP)
+override DEVICE:=GPU
+export
+$(info Using HIP)
+$(info )
+include Makefile.Hip
+else
+
 ifeq ($(DEVICE), $(filter $(DEVICE),GPU CUDA))
 MIN_COMPUTE:=50
 ifeq ($(TENSOR), ON)
@@ -64,3 +74,5 @@ $(info GPU_INCLUDE_PATH and GPU_LIBRARY_PATH)
 $(info )
 include Makefile.OpenCL
 endif
+
+endif # DEVICE=HIP

Makefile.Hip

@@ -0,0 +1,201 @@
+# AutoDock-GPU HIP Makefile (ROCm/AMD backend, DEVICE=HIP)
+#
+# Mirrors Makefile.Cuda but compiles the shared cuda/ kernel translation unit
+# with hipcc and links the HIP runtime. The host .cpp are compiled by g++ with
+# the HIP runtime headers (host-side HIP API only). Selected via DEVICE=HIP in
+# the top-level Makefile.
+#
+# Arch is configurable and never hardcoded, so any GPU target builds
+# with only HIP_ARCH=<arch> and no source/Makefile edit:
+#   make DEVICE=HIP NUMWI=64 HIP_ARCH=gfx90a
+
+ROCM_PATH ?= /opt/rocm
+HIPCC = $(ROCM_PATH)/bin/hipcc
+CPP = g++
+UNAME := $(shell uname)
+
+# Default the target GPU arch ONLY when unset; a passed HIP_ARCH must win.
+HIP_ARCH ?= gfx90a
+
+# TENSOR=ON accelerates the energy/gradient sum-reduction with matrix-core units.
+# On AMD this is rocWMMA over native fp32 MFMA (header-only, no extra link); it
+# activates the same USE_NVTENSOR call sites the CUDA backend uses. The reduction
+# fills a wavefront-wide collective, so it requires NUMWI > 32 (as on CUDA).
+ifeq ($(TENSOR), ON)
+NVTENSOR = -DUSE_NVTENSOR
+ifeq ($(filter 64 128 256,$(NUMWI)),)
+$(error TENSOR=ON requires NUMWI greater than 32 (use NUMWI=64, 128, or 256))
+endif
+endif
+
+ifeq ($(DEVICE), CPU)
+	DEV =-DCPU_DEVICE
+else ifeq ($(DEVICE), GPU)
+	DEV =-DGPU_DEVICE
+endif
+
+# ------------------------------------------------------
+# Project directories
+COMMON_DIR=./common
+HOST_INC_DIR=./host/inc
+HOST_SRC_DIR=./host/src
+KRNL_DIR=./cuda
+KCMN_DIR=$(COMMON_DIR)
+BIN_DIR=./bin
+
+TARGET := autodock
+TOOL_TARGET := adgpu_analysis
+
+# HIP runtime headers/libs for both the hipcc kernel TU and the g++ host TUs.
+HIP_INCLUDES = -I$(ROCM_PATH)/include
+HIP_HOST_FLAGS = -DUSE_HIP -D__HIP_PLATFORM_AMD__ $(HIP_INCLUDES)
+LIB_HIP = kernels.o -L$(ROCM_PATH)/lib -lamdhip64 -DUSE_HIP
+
+IFLAGS=-I$(COMMON_DIR) -I$(HOST_INC_DIR) -I$(KRNL_DIR) $(HIP_INCLUDES)
+LFLAGS=-Wl,-rpath=$(ROCM_PATH)/lib
+CFLAGS=-std=c++17 $(IFLAGS) $(LFLAGS)
+TOOL_CFLAGS=-std=c++17 -I$(COMMON_DIR) -I$(HOST_INC_DIR)
+
+ifeq ($(DEVICE), CPU)
+	TARGET:=$(TARGET)_cpu
+else ifeq ($(DEVICE), GPU)
+	NWI=-DN64WI
+	TARGET:=$(TARGET)_gpu
+endif
+
+ifeq ($(OVERLAP), ON)
+	PIPELINE=-DUSE_PIPELINE -fopenmp
+endif
+
+BIN := $(wildcard $(TARGET)*)
+
+# ------------------------------------------------------
+# Number of work-items (wi)
+# Valid values: 32, 64, 128, 256
+NUMWI=
+
+ifeq ($(NUMWI), 32)
+	NWI=-DN32WI
+	TARGET:=$(TARGET)_32wi
+else ifeq ($(NUMWI), 64)
+	NWI=-DN64WI
+	TARGET:=$(TARGET)_64wi
+else ifeq ($(NUMWI), 128)
+	NWI=-DN128WI
+	TARGET:=$(TARGET)_128wi
+else ifeq ($(NUMWI), 256)
+		NWI=-DN256WI
+		TARGET:=$(TARGET)_256wi
+else
+	ifeq ($(DEVICE), CPU)
+		NWI=-DN16WI
+		TARGET:=$(TARGET)_16wi
+	else ifeq ($(DEVICE), GPU)
+		NWI=-DN128WI
+		TARGET:=$(TARGET)_128wi
+	endif
+endif
+
+# ------------------------------------------------------
+# Configuration
+# FDEBUG (full) : enables debugging on both host + device
+# LDEBUG (light): enables debugging on host
+# RELEASE
+CONFIG=RELEASE
+HIP_FLAGS = --offload-arch=$(HIP_ARCH) -std=c++17
+
+ifeq ($(CONFIG),FDEBUG)
+	OPT =-O0 -g3 -Wall -DDOCK_DEBUG
+	HIP_FLAGS += -g -ffast-math
+else ifeq ($(CONFIG),LDEBUG)
+	OPT =-O0 -g3 -Wall
+	HIP_FLAGS += -ffast-math
+else ifeq ($(CONFIG),RELEASE)
+	OPT =-O3
+	HIP_FLAGS += -O3 -ffast-math
+else
+	OPT =
+	HIP_FLAGS += -ffast-math
+endif
+
+# ------------------------------------------------------
+# Reproduce results (remove randomness)
+REPRO=NO
+
+ifeq ($(REPRO),YES)
+	REP =-DREPRO
+else
+	REP =
+endif
+# ------------------------------------------------------
+
+all: otool odock
+
+check-env-dev:
+	@if test -z "$$DEVICE"; then \
+		echo "Please set DEVICE to either CPU, GPU, CUDA, OCLGPU, or HIP to build docking software."; \
+		exit 1; \
+	else \
+		if [ "$$DEVICE" = "GPU" ]; then \
+			echo "DEVICE is set to $$DEVICE"; \
+		else \
+			echo "DEVICE value is invalid. Please set DEVICE to either CPU, GPU, CUDA, OCLGPU, or HIP"; \
+			exit 1; \
+		fi; \
+	fi; \
+	echo " "
+
+check-env-all: check-env-dev
+
+# ------------------------------------------------------
+# Printing out its git version hash
+
+GIT_VERSION := $(shell ./version_string.sh)
+
+CFLAGS+=-DVERSION=\"$(GIT_VERSION)\"
+TOOL_CFLAGS+=-DVERSION=\"$(GIT_VERSION)\"
+
+# ------------------------------------------------------
+
+kernels: $(KERNEL_SRC)
+	$(HIPCC) $(NWI) $(NVTENSOR) $(REP) $(HIP_FLAGS) -fPIE -DUSE_HIP $(IFLAGS) -c $(KRNL_DIR)/kernels.cu -o kernels.o
+
+otool:
+	@echo "Building" $(TOOL_TARGET) "..."
+	$(CPP) \
+	$(shell ls $(HOST_SRC_DIR)/*.cpp) \
+	$(TOOL_CFLAGS) \
+	-o$(BIN_DIR)/$(TOOL_TARGET) \
+	$(PIPELINE) $(OPT) -DTOOLMODE $(REP)
+
+odock: check-env-all kernels
+	@echo "Building" $(TARGET) "..."
+	$(CPP) \
+	$(shell ls $(HOST_SRC_DIR)/*.cpp) \
+	$(CFLAGS) \
+	$(HIP_HOST_FLAGS) \
+	$(LIB_HIP) \
+	-o$(BIN_DIR)/$(TARGET) \
+	$(DEV) $(NWI) $(NVTENSOR) $(PIPELINE) $(OPT) $(DD) $(REP) $(KFLAGS)
+
+PDB      := 3ce3
+NRUN     := 100
+NGEN     := 27000
+POPSIZE  := 150
+TESTNAME := test
+TESTLS   := ad
+
+test: odock
+	$(BIN_DIR)/$(TARGET) \
+	-ffile ./input/$(PDB)/derived/$(PDB)_protein.maps.fld \
+	-lfile ./input/$(PDB)/derived/$(PDB)_ligand.pdbqt \
+	-nrun $(NRUN) \
+	-ngen $(NGEN) \
+	-psize $(POPSIZE) \
+	-resnam $(TESTNAME) \
+	-gfpop 0 \
+	-lsmet $(TESTLS)
+
+.PHONY: clean
+clean:
+	rm -f kernels.o $(BIN_DIR)/$(TARGET)* $(BIN_DIR)/$(TOOL_TARGET)*

README.md

@@ -45,11 +45,11 @@ make DEVICE=<TYPE> NUMWI=<NWI>
 
 | Parameters | Description                  | Values                                             |
 |:----------:|:----------------------------:|:--------------------------------------------------:|
-| `<TYPE>`   | Accelerator chosen           | `CPU`, `GPU`, `CUDA`, `OCLGPU`, `OPENCL`           |
+| `<TYPE>`   | Accelerator chosen           | `CPU`, `GPU`, `CUDA`, `HIP`, `OCLGPU`, `OPENCL`    |
 | `<NWI>`    | work-group/thread block size | `1`, `2`, `4`, `8`, `16`, `32`, `64`, `128`, `256` |
 
-When `DEVICE=GPU` is chosen, the Makefile will automatically tests if it can compile Cuda succesfully. To override, use `DEVICE=CUDA` or `DEVICE=OCLGPU`. The cpu target is only supported using OpenCL. Furthermore, an OpenMP-enabled overlapped pipeline (for setup and processing) can be compiled with `OVERLAP=ON`.
-Hints: The best work-group size depends on the GPU and workload. Try `NUMWI=128` or `NUMWI=64` for modern cards with the example workloads. On macOS, use `NUMWI=1` for CPUs.
+When `DEVICE=GPU` is chosen, the Makefile will automatically tests if it can compile Cuda succesfully. To override, use `DEVICE=CUDA` or `DEVICE=OCLGPU`. For AMD GPUs, use `DEVICE=HIP` with `HIP_ARCH=<gfx-target>` (for example `HIP_ARCH=gfx90a`), which builds the same CUDA kernels through ROCm/HIP. On AMD GPUs with matrix cores, adding `TENSOR=ON` (requires `NUMWI` of 64 or more) accelerates the energy/gradient sum-reduction through rocWMMA, mirroring the CUDA `TENSOR=ON` path. The cpu target is only supported using OpenCL. Furthermore, an OpenMP-enabled overlapped pipeline (for setup and processing) can be compiled with `OVERLAP=ON`.
+Hints: The best work-group size depends on the GPU and workload. Try `NUMWI=128` or `NUMWI=64` for modern cards with the example workloads. On AMD GPUs, `NUMWI=64` was fastest in our tests on both consumer (RDNA) and datacenter cards. On macOS, use `NUMWI=1` for CPUs.
 
 After successful compilation, the host binary **autodock_&lt;type&gt;_&lt;N&gt;wi** is placed under [bin](./bin).
 

cuda/GpuData.h

@@ -26,6 +26,7 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA
 #ifndef GPUDATADOTH
 #define GPUDATADOTH
 #include <float.h>
+#include "cuda_to_hip.h"
 
 
 static const int   TERMBITS         = 10;

cuda/cuda_to_hip.h

@@ -0,0 +1,78 @@
+/*
+
+AutoDock-GPU, an OpenCL implementation of AutoDock 4.2 running a Lamarckian Genetic Algorithm
+Copyright (C) 2017 TU Darmstadt, Embedded Systems and Applications Group, Germany. All rights reserved.
+For some of the code, Copyright (C) 2019 Computational Structural Biology Center, the Scripps Research Institute.
+For the ROCm/HIP port, Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+AutoDock is a Trade Mark of the Scripps Research Institute.
+
+This library is free software; you can redistribute it and/or
+modify it under the terms of the GNU Lesser General Public
+License as published by the Free Software Foundation; either
+version 2.1 of the License, or (at your option) any later version.
+
+This library is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+Lesser General Public License for more details.
+
+You should have received a copy of the GNU Lesser General Public
+License along with this library; if not, write to the Free Software
+Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA
+
+*/
+
+// Single CUDA-to-HIP compatibility shim for the HIP backend (DEVICE=HIP).
+// This is the only file that knows HIP symbol names: on ROCm it aliases the
+// CUDA runtime spellings the project uses to their HIP equivalents; everywhere
+// else it is a plain CUDA-runtime include, so the CUDA build is unchanged.
+// Authoritative cuda->hip name source: pytorch
+// torch/utils/hipify/cuda_to_hip_mappings.py.
+
+#ifndef CUDA_TO_HIP_H
+#define CUDA_TO_HIP_H
+
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+
+// libc host decls must win over HIP's __device__ memcpy/memset overloads once
+// <hip/hip_runtime.h> is in scope inside a .cu compiled as HIP (the host helper
+// code in this TU calls memcpy); include them first.
+#include <cstring>
+#include <cstdlib>
+#include <hip/hip_runtime.h>
+
+// Runtime API
+#define cudaError_t                  hipError_t
+#define cudaSuccess                  hipSuccess
+#define cudaGetErrorString           hipGetErrorString
+#define cudaGetLastError             hipGetLastError
+#define cudaDeviceSynchronize        hipDeviceSynchronize
+#define cudaDeviceReset              hipDeviceReset
+#define cudaSetDevice                hipSetDevice
+#define cudaGetDevice                hipGetDevice
+#define cudaGetDeviceCount           hipGetDeviceCount
+#define cudaGetDeviceProperties      hipGetDeviceProperties
+#define cudaDeviceProp               hipDeviceProp_t
+#define cudaDeviceSetLimit           hipDeviceSetLimit
+#define cudaLimitPrintfFifoSize      hipLimitPrintfFifoSize
+#define cudaMemGetInfo               hipMemGetInfo
+
+// Memory
+#define cudaMalloc                   hipMalloc
+#define cudaMallocManaged            hipMallocManaged
+#define cudaMemAttachGlobal          hipMemAttachGlobal
+#define cudaFree                     hipFree
+#define cudaMemcpy                   hipMemcpy
+#define cudaMemcpyToSymbol           hipMemcpyToSymbol
+#define cudaMemcpyFromSymbol         hipMemcpyFromSymbol
+#define cudaMemcpyHostToDevice       hipMemcpyHostToDevice
+#define cudaMemcpyDeviceToHost       hipMemcpyDeviceToHost
+
+#else // CUDA
+
+#include <cuda_runtime.h>
+
+#endif
+
+#endif // CUDA_TO_HIP_H

cuda/kernel4.cu

@@ -90,7 +90,7 @@ gpu_gen_and_eval_newpops_kernel(
 		// Perform final reduction in warp 0
 		if (warpID == 0)
 		{
-			int blocks = blockDim.x / 32;
+			int blocks = blockDim.x / WARP_SIZE;
 			if (tgx < blocks)
 			{
 				bestID = sBestID[tgx];