Add AMD GPU support (ROCm/HIP)

[jeffdaily]

This adds AMD GPU support to Marian by porting the CUDA GPU backend to HIP, targeting ROCm. There is no AMD path in the tree today, so this is new hardware support rather than a change to existing behavior. The NVIDIA build is untouched except for three shared fixes called out in their own section below: everything else is additive or guarded, and with USE_HIP off the CUDA sources compile as before.

An earlier revision of this branch left sparse matrix multiply, multi-GPU collectives and the cuDNN convolution path out of scope. All three are now ported and exercised on real GPUs, which is what most of the new work below is.

Build integration

A new USE_HIP CMake option selects the HIP path. Marian compiles its GPU sources through the legacy cuda_add_library from FindCUDA, so under USE_HIP that call is satisfied by a shim macro that creates an ordinary library and marks the GPU translation units as HIP, driving CMAKE_HIP_ARCHITECTURES instead of a -gencode list. No architecture is hardcoded; when the option is not set the architecture is auto-detected, and the build stops with a clear message rather than guessing if detection fails.

Source changes

A single compat header aliases the CUDA runtime symbols Marian uses to their HIP equivalents. A small directory of forwarding headers (cuda_fp16.h, cublas_v2.h, cublasLt.h, cusparse.h, curand.h, nccl.h, cooperative_groups.h and friends) sits on the include path for HIP builds only, so the sources keep their CUDA spelling and the diff stays small; on NVIDIA that directory is absent and the real toolkit headers are used.

Library mappings are one-to-one: cuBLAS to hipBLAS, cuBLASLt to hipBLASLt, cuSPARSE to hipSPARSE, cuRAND to hipRAND, Thrust to rocThrust, cuDNN to MIOpen, NCCL to RCCL. The option defaults are upstream's, unchanged: USE_NCCL is on and now works under HIP rather than having to be turned off, and USE_CUDNN stays off by default but builds and runs when enabled.

The one change worth close review

src/tensors/gpu/topk.cu and src/translator/nth_element.cu reduce shared memory with a __syncthreads() tree down to 32 elements and then finish with an unrolled warp-synchronous tail that has no barriers and reads non-volatile shared memory. That is correct only where the hardware warp is 32 lanes wide and executes in lockstep. AMD CDNA parts have a 64-lane wavefront, so the tail's assumption does not hold there: the low 32 lanes are not implicitly synchronized across those steps and non-volatile shared memory may be register-cached, which can make the beam-search argmax silently wrong and non-deterministic. On the HIP path the tree now runs all the way down to a single element under __syncthreads(), preserving the original comparison order.

Because this sits on the beam-search top-k, where a wrong result corrupts translations rather than crashing, it was checked with fixed-seed determinism runs (two runs bit-identical, and identical to a CPU decode of the same model) and not by inspection alone.

Sparse matrix multiply

CSRProd in src/tensors/gpu/prod_sparse_cu11.h now produces correct results under hipSPARSE for both fp32 and fp16. The existing CUSPARSE_ORDER_ROW and CSR_ALG2 selection is correct on hipSPARSE too and is unchanged; the two fixes below are independent of each other.

For fp32, the multiply was never issued. The code called cusparseSpMM only inside if(bufferSize > 0); cuSPARSE always asks for scratch for CSR_ALG2 so on NVIDIA the guard never fires, but hipSPARSE reports zero for these shapes and C kept whatever it already held. The allocation is now conditional and the multiply unconditional, which is the one behavioral difference this change makes on the NVIDIA path and is inert there for the reason just given.

For fp16, hipSPARSE has no uniform half SpMM: half sparse and dense operands are accepted only against a float result matrix with float compute. The half instantiation therefore multiplies into a float scratch matrix and casts the result back, seeding that scratch from C when beta is non-zero so accumulation still means what it did. Scalars on this path are float to match the compute type. This costs a cast round-trip per call; it is correct but untuned, and worth revisiting if fp16 sparse attention throughput ever matters.

Multi-GPU collectives

find_package(rccl) and a nccl.h forwarding header are enough for src/training/communicator_nccl.h to compile unchanged, since RCCL keeps the NCCL API and symbol names. The vendored NCCL ExternalProject (which does not build under hipcc) is skipped under USE_HIP, and RCCL is linked from ROCm instead.

This is exercised on four GPUs rather than compile-checked: synchronous SGD across --devices 0 1 2 3 converges on both a 4x MI250X and a 4x Radeon Pro W7800 host, the communicators construct, and decoding the resulting model reproduces the task exactly.

Convolution and pooling

src/tensors/gpu/cudnn_wrappers.{h,cu} keep one set of class declarations, with the handle and descriptor types aliased to MIOpen's and the MIOpen implementation as a self-contained block ahead of the untouched cuDNN one. MIOpen uses an ordinary tensor descriptor for convolution weights rather than a distinct filter type, which the aliasing accounts for.

Three MIOpen differences shaped the implementation. Workspaces are explicit and the algorithm is chosen by benchmarking, and the miopenFindConvolution*Algorithm calls write to whatever tensor is named as their output, so each search is given a destination that is about to be overwritten anyway rather than the live input or weight buffer; results are cached per input shape. Only alpha=1 and beta=0 are honored in 2D, so the accumulating backward form is a convolution into a staging buffer followed by miopenOpTensor, and forward bias is likewise an add rather than the overwriting miopenConvolutionForwardBias. Pooling backward reads the indices its forward recorded, so that workspace outlives the call.

Three fixes the NVIDIA build needs too

These are the only changes here that alter CUDA behavior, and they are the reason USE_CUDNN=ON works at all.

• src/common/options.cpp never instantiated Get<std::pair<int, int>>, which is how src/layers/convolution.cpp reads kernel-dims, paddings and strides, so nothing linked with USE_CUDNN=ON on any platform before this. The instantiation is added, next to the existing scalar ones; fastopt.cpp already had the matching As<std::pair<int, int>> and needs no change.
• PoolingOp in src/graph/node_operators_unary.h kept its input's shape instead of asking the wrapper for the pooled one, so any window that actually shrinks the input left most of the output unwritten. It now calls getOutputShape the way ConvolutionOp does.
• The sparse zero-scratch guard described above, which is inert on cuSPARSE.

One caveat on the CUDA side: cudnn_wrappers.cu targets the convolution algorithm-selection API that cuDNN 8 removed, so the CUDA cuDNN path does not compile against a current cuDNN with or without this change. The two fixes above are in shared files that are not cuDNN-specific, and they compile and link in the CUDA build.

Validation

Linux, ROCm 7.2, USE_HIP=ON with USE_CUDNN=ON and USE_NCCL=ON, on AMD Instinct MI250X (gfx90a, 64-lane wavefront) and Radeon Pro W7800 (gfx1100, 32-lane wavefront), four GPUs of each. 294 of 294 targets build on both, and the numbers below are identical on the two architectures.

• operator_tests: 603 of 603 assertions, all four test cases pass, fp32 GPU, fp16 GPU and CPU. The sparse and fp16 affine sections are included; previously the sparse exception aborted the case and everything after it never ran on any platform.
• graph 10/10, attention 6/6, transformer 3/3, binary 9/9, fastopt 23/23, utils 8/8.
• rnn_tests 21 of 24, the three unchanged reference-value mismatches described under limitations.
• Convolution and pooling numerics: a standalone program linked against the built library checks convolution forward and its gradients with respect to input, kernel and bias, plus max and average pooling forward and backward, against hand-computed references, and a multi-output-channel convolution of the character encoder's shape against a host reference. All agree. The in-tree pooling app test only constructs a graph, so it is not evidence on its own.
• The character-CNN encoder (--type char-s2s), the only consumer of the convolution path, trains on GPU with cost falling steadily and no NaN or Inf. It needs --char-highway 0 because the highway network calls sigmoid(vector<Expr>), which aborts with "Not implemented" on every platform; that is unrelated to this change and not addressed here.
• End-to-end determinism: a transformer trained to convergence on a toy task, decoded with beam 6 twice on GPU and once on CPU, gives byte-identical output all three times, on both architectures.
• The CUDA build is unaffected: 161 of 161 targets compile and link with nvcc 12.8 and USE_HIP off.

Windows, on Radeon 8060S (gfx1151) with a current ROCm toolchain: 277 of 277 targets build, the fp32 GPU unit suites pass (graph 10/10, attention 2/2, transformer 1/1, operator 202/202 fp32 GPU and 198/198 CPU, binary 9/9, fastopt 23/23, utils 8/8, test_prod), training runs on GPU, and beam-6 decode succeeds both on a toy model and on a production-size one (512-dimensional embeddings, six layers, 22k vocabulary, mini-batch 32). Radeon RX 9070 XT (gfx1201) passed the same checks on the earlier revision of this branch. MIOpen and RCCL are validated on Linux only.

On Windows the fused GEMM-plus-bias path is computed as a plain GEMM followed by the existing broadcast bias add, under a _WIN32 guard, because the current hipBLASLt release faults inside algorithm selection for every such call there. Results are identical; this is the same fallback the pre-CUDA-11 path uses.

Known limitations

• fp16 is unavailable in Windows builds regardless of GPU vendor, because src/common/types.h stubs float16 out under _MSC_VER; the fp16 test sections abort there as they do on a Windows CUDA build. Not addressed here.
• rnn_tests has three failures in its "Simple RNN" case on every AMD architecture. The test compares against reference values hard-coded from cuRAND's generator; hipRAND produces a different, equally valid random stream, so the initial weights differ. It is a test-fixture assumption, not a computation error.
• The CUB segmented-sort path in topk.cu is not ported; on HIP Sort() uses the rocThrust sort-by-key path instead. Only the tests reach it.
• One Windows configuration does not work yet for a reason outside Marian: on gfx1101 a batched GEMM falls back to a generic kernel image that the Windows loader rejects, because that architecture's math library ships no compiled kernel for the layout Marian's batched product selects. The gfx1201 and gfx1151 configurations above are unaffected.

Build recipe

cmake -S . -B build-hip -G Ninja \
  -DUSE_HIP=ON -DCOMPILE_CUDA=ON \
  -DUSE_CUDNN=ON -DUSE_NCCL=ON -DUSE_FBGEMM=OFF \
  -DCOMPILE_CPU=ON -DCOMPILE_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build-hip -j$(nproc)

CMAKE_HIP_ARCHITECTURES can be set to target a specific GPU; left unset, it is auto-detected from the host.

---

This pull request was prepared with the help of an AI assistant (Claude) and was read and approved by a person before it was opened. It comes from an ongoing effort to add AMD GPU support to widely used CUDA projects, one repository at a time: https://github.com/AMD-Ecosystem/moat -- that repository describes how the work is done and what a person checks before anything is submitted.

If you would rather not receive pull requests from this effort, say so here or open an issue at https://github.com/AMD-Ecosystem/moat/issues/new/choose and we will close this and stop. That can cover this repository alone or everything you own, whichever you prefer.

[jeffdaily]

Converting to draft while we port additional features (sparse matrix, NCCL, and cuDNN paths). Will mark ready for review when complete.