[ROCm] Add AMD GPU support to the simulation engine

[jeffdaily]

alien's simulation engine was CUDA-only and its device-selection gate rejected any GPU below compute capability 6.0, so it did not run on AMD hardware at all. This adds a HIP build (USE_HIP, default OFF) alongside the existing CUDA build: only the .cu translation units see the HIP toolchain, host C++ is untouched, and the NVIDIA build is byte-for-byte unchanged.

Review order: start with the top-level CMakeLists.txt (the USE_HIP option, language gating, the .cu->HIP retag, and the relocatable-device-code wiring), then source/EngineKernels/cuda_to_hip.h (the single compat header, force-included on every HIP and host TU) and source/hip_compat/ (forwarding shims for the toolkit angle-bracket includes, on the HIP include path only). The remaining files are the handful of real semantic fixes and the Windows build support.

Relocatable device code is load-bearing: __constant__ cudaSimulationParameters (ConstantMemory.cu) is read by device code in sibling EngineKernels translation units, so the CUDA path already uses CUDA_SEPARABLE_COMPILATION. On HIP, CMake 3.31 neither injects -fgpu-rdc from the property nor emits a separate device-link step, so -fgpu-rdc is added to the .cu compiles and -fgpu-rdc --hip-link to the executables that link the engine libraries; without both the device link fails with undefined cudaSimulationParameters / _hip_fatbin* symbols.

Semantic fixes (each correct on wave32 and wave64, NVIDIA path unchanged):

• USE_HIP-guard the compute-capability gate so AMD GPUs are accepted.
• Replace cg::reduce in the SPH fluid kernels with a tile shfl_xor butterfly in the cooperative_groups/reduce.h shim (HIP cooperative groups has no cg::reduce/cg::plus). tiled_partition<32> slices the wavefront into fixed 32-lane tiles and every tile shuffle is tile-relative, so the code is identical on wave32 and wave64. The two fluid kernels launch at blockDim 25 and 81, so their last tile is partial; the butterfly substitutes the reduction identity for any partner read from a non-resident lane (rank ^ offset >= the tile's active-lane count), so lane 0 -- the only lane the call sites consume -- yields exactly the resident-lane sum without relying on the value a shuffle returns for an inactive lane.
• Guard out the float2/float3/int2 operators HIP's vector types already define component-wise (same semantics) to avoid an ambiguous overload; keep the mixed float2-int2 helper HIP lacks.
• Cast max(int, uint32_t) operands to a common type (HIP's mixed overload returns double, breaking a uint64_t modulo).
• Give a muscle-activation lambda an explicit float return type.

Windows support (gated on the compiler frontend variant, the Linux/GCC build unaffected):

• clang-cl (MSVC frontend): /FI force-include, /clang:-ffast-math, NOMINMAX on the HIP TUs.
• Work around a clang-offload-bundler input==output failure (override CMAKE_HIP_COMPILE_OBJECT to emit -o instead of /Fo) and drive the -fgpu-rdc device link through a clang-driver wrapper (cmake/hip_link_win.py) for both the MSVC-frontend and GCC-frontend clang++ link paths, since lld-link cannot perform the HIP device link.

Building on AMD GPUs (also documented in the README's build instructions):

mkdir build && cd build
cmake .. -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a
cmake --build . --config Release -j8

Test Plan:

Built and validated on AMD Instinct MI250X (gfx90a, CDNA2 wave64), Radeon Pro W7800 (gfx1100, RDNA3 wave32), Radeon PRO V710 (gfx1101, RDNA3 wave32, Windows), and Radeon RX 9070 XT (gfx1201, RDNA4 wave32, Windows).

• EngineTests on gfx90a: 2978 tests, 2973 passed, 3 skipped, 2 failed -- the same tally on both the debug (direct-launch) and CUDA-graph paths, including the three 10000-20000-timestep long-runners. CLI smoke (graph path, headless): 1000 timesteps, fault-free, total energy conserved.
• The two failures are not port defects. CommunicatorTests.sender_signalPriority_lowerNumTimesSentWins exercises an upstream last-writer-wins race in CommunicatorProcessor::tryTransmitSignal (no priority compare; the winner depends on block execution order, which differs across vendors). DataTransferTests.multipleCells_genome_multipleGenes_multipleNodes is a pure set/get round-trip compared with exact float ==, where one cell-position path rounds ~1 ULP (2^-17 at coordinate ~100) off CUDA -- physically identical to 8 significant figures.
• GeometryTests is the GL-interop render path and needs an on-screen GL context, separate from the headless compute path these tests cover.

Authored with the assistance of Claude (Anthropic).

[chrxh]

First of all, thank you for your PR and sorry for the late response!

I'm still undecided about integrating the changes and will leave the PR open for now. It's already possible to compile ALIEN for AMD GPUs using SCALE (see #366).
However, I haven't yet gained any personal experience or made any comparisons regarding how this works on AMD GPUs...

[jeffdaily]

Thanks, and no rush.

SCALE is one option, though it and this PR take different routes. SCALE is a separate proprietary toolchain that recompiles the CUDA sources, whereas this PR makes AMD a native in-tree build target through the standard open-source ROCm/HIP toolchain, behind USE_HIP (default OFF). Only the .cu translation units see the HIP compiler, so the default CUDA build is unaffected; I rebuilt it with nvcc to confirm it still compiles and links unchanged.

Going native rather than through a translation layer avoids depending on an external proprietary compiler and tends to be faster; #366 has a community data point where mauri870's separate port reported roughly 130 TPS against about 53 for SCALE on a 7900 XTX. I have not compared the two myself, but I can run a direct measurement on the same GPU if that would help you decide. It also runs on Windows ROCm (validated on gfx1201 and gfx1101), which SCALE does not target.

I have also rebased the branch onto the current develop, so it is up to date and conflict-free if you do want to try it.

[chrxh]

Thanks again for the PR!
Being able to compile and run the engine with HIP on AMD GPUs is a great addition :)

My main concern before this can be merged is the code organization: I don't want the HIP-specific code to live under source/. The engine source tree should stay CUDA-only. Right now the port adds a whole set of HIP-only files under source/ and source/EngineKernels/cuda_to_hip.h and the entire source/hip_compat/ directory.

What I'd like to see all HIP-only files (the compat shim and the whole hip_compat/ tree) out of source/ e.g. into external/ or a dedicated top-level hip/ directory. Keep source/EngineKernels/ free of HIP.

Once the HIP-specific parts are relocated out of source/, I'm happy to take another look.

[jeffdaily]

The HIP-only files now live under a top-level hip/ directory (hip/cuda_to_hip.h and hip/hip_compat/); source/ is CUDA-only again. Done in 444d12b, re-validated on gfx90a and confirmed the default CUDA build is unchanged.

[chrxh]

Merged. Thanks again for the contribution!