ArrayFire has no HIP/ROCm backend; AMD GPUs are reachable only through the
OpenCL path. This adds a native HIP backend as a new sibling of the CUDA
backend (src/backend/hip, cloned from src/backend/cuda), keeping the
NVIDIA/CUDA path byte-for-byte unchanged. The HIP backend reports
AF_BACKEND_CUDA and builds as the library afcuda, so the unified dispatcher and
the gtest suite treat it as the CUDA-model backend on AMD; no public ABI/enum
change. The full afcuda shared library plus all 132 CUDA-tagged test binaries
build for gfx90a (CPU backend kept on as the in-process reference).

Review order: the defining work is the runtime-JIT engine in
src/backend/hip/compile_module.cpp, moved from NVRTC + the CUDA driver link
step (nvrtcGetPTX -> cuLinkCreate/cuLinkAddData/cuLinkComplete ->
cuModuleLoadData) to hipRTC's direct-code-object flow (hiprtcGetCode ->
hipModuleLoadData), arch from the device gcnArchName (--offload-arch). The
hipRTC compile needed: per-token splitting of the shared DefineValue macros'
" -D NAME=val" options (clang rejects the NVRTC-style joined form), the clang
resource-dir on -isystem (stddef.h), -D__CUDACC_RTC__ (so af/defines.h takes its
RTC path and skips host includes), and -D__CUDA_ARCH__ (so the embedded device
headers use their intrinsic path). The embedded JIT-source headers
(cuComplex.h / cuda_fp16.h / math_constants.h / vector_types.h) are HIP shims
under nvrtc_shims/; hip_compat.h is force-included on every TU and aliases the
cudaXxx / CUxxx / cuComplex surface.

hipRTC vs NVRTC device-code rules: NVRTC compiles the whole JIT TU as device
code, so unattributed helpers are implicitly device; clang/hipRTC treats them as
host. Helpers reachable from JIT kernels therefore need explicit device
attributes that NVRTC never required: common/half.hpp half2int / the member
half::infinity() (__DH__), hip/math.hpp division() and hip/minmax_op.hpp cabs /
MinMaxOp (__DH__), the kernel-local helpers diff_this / select getOffset /
convolve3 index (__device__), and the sparse-arith arith_op<T,op>::operator()
(__device__; the SSD/DSD csr/coo kernels call it). The hipRTC std shim in
half.hpp also gains numeric_limits<double> and std::isnan/isinf(float|double)
(hipRTC's bundled std is smaller than NVRTC's and injects only a hip_bfloat16
isnan). Half-precision transcendentals are emitted by the JIT as the bare math
name (sin/cos/...); __half converts to both float and double so the call is
ambiguous on HIP -- jit.cuh adds __half overloads (native h* intrinsic or
float-promoted). A shared .cuh that defines several kernel templates must pass
every -D either template's body references as a non-dependent identifier on ALL
launchers that compile it (clang does phase-1 lookup on the uninstantiated
template; NVRTC does not): fixed scan_first/scan_dim bcast and ireduce, lookup,
and sparse_arith (csrArith* use TX/TY, cooArith* use THREADS, so every launcher
passes all three). memCopyLoop13 had an upstream g1/id1 typo that only the HIP
dispatch reaches.

Library swaps: hipBLAS (function pointers reinterpret-cast because hipBLAS's
hipblasHalf / hipComplex element types differ from the backend's __half / POD
cfloat), hipSOLVER via its cuSOLVER-compatible hipsolverDn* API, hipFFT,
hipSPARSE (generic API + legacy sort/conversion + csrgeam2), rocThrust / hipCUB.
The void* handle aliasing (hipblas/hipsolver/hipsparse, AND the hipSPARSE
descriptors -- DnVec and DnMat are both typedef void*) is solved with a tag-keyed
RAII (hip_unique_handle.hpp). cfloat/cdouble are plain PODs on HIP for both the
host backend AND the JIT templated path (not HIP_vector_type, whose
componentwise friend operators would tie with arrayfire's complex operators).
Wave64: shfl_intrinsics 64-bit mask; reduce.hpp keeps 32-lane logical groups for
the row-packed reduce_first/all while reduce_by_key uses
kWarpSize and a kWarpSize-sized per-warp result buffer.

Sparse is implemented on hipSPARSE (it replaces the earlier
AF_ERR_NOT_SUPPORTED stubs). The CUDA backend's generic-API path
(cusparseCreateCsr/Csc/Coo, SpMV/SpMM, DenseToSparse/SparseToDense,
SpMatGetSize, Csr/CscSetPointers) plus the legacy
Xcsrsort/Xcoosort/Xcsr2coo/Xcoo2csr/CreateIdentityPermutation and the typed
csrgeam2 surface all map 1:1 to hipSPARSE 4.2. The .cu keep their cuSPARSE
spelling via a forwarding shim (nvrtc_shims/cusparse_v2.h -> hipsparse, on the
HIP include path only); unlike the NVIDIA build's runtime-dlopen cusparseModule
plugin, the HIP build links roc::hipsparse and calls the functions directly. Two
non-1:1 deltas: hipsparseSpMV/SpMM take the compute type as a hipDataType
(getType<T>(), not the hipblasComputeType_t getComputeType<T>() returns for the
dense gemm Ex path), and the typed complex csrgeam2 takes hipComplex* /
hipDoubleComplex* so the complex value/alpha/beta pointers are reinterpret_cast
at the call boundary (cfloat/cdouble are distinct layout-compatible PODs).

The int8 (schar) gemm is closed: rocBLAS rejects int8-in/float32-out, but
gfx9/CDNA supports int8 x int8 -> int32 accumulate, so the schar path runs
hipblasGemmEx with HIP_R_8I in + HIP_R_32I out + HIPBLAS_COMPUTE_32I and casts
the int32 result into the f32 output (if constexpr-guarded to the schar
instantiation). FreeImage is enabled (AF_WITH_IMAGEIO=ON) so confidence_connected
and imageio reach the GPU path.

Templated complex kernels needed three fixes so the complex-element JIT kernels
are correct: a bare `a * b` on a complex T must be a complex (not componentwise)
product -- the runtime-JIT cuComplex.h shim defines POD cuFloatComplex/
cuDoubleComplex (the host path keeps the hipFloatComplex aliases) and the
convolve kernels spell the product out via a local convMul -- and the JIT
complex == / != must live in the GLOBAL namespace beside the POD type (the
shim) so ADL finds them from any namespace; arrayfire::cuda's equality operators
in math.hpp are unreachable by ADL for a global-namespace POD, which made the
where-over-complex count-scan (common::Transform<cuFloatComplex,uint,
af_notzero_t>) fail overload resolution under hipRTC (math.hpp drops its complex
==/!= on the RTC path so the shim's are unambiguous). This surfaced as an
AF_ERR_INTERNAL in `where` for cfloat/cdouble; it is a host-set name-lookup bug,
not arch-specific, so the fix is arch-unified.

GPU-validated on gfx90a (CDNA2, wave64): the full CUDA.* gtest suite is
132/132 binaries passing (ctest -R '_cuda$'), no residual failures. The JIT
engine (jit 1781/1781), transpose, scan/scan_by_key, fft, reduce (incl. ragged
+ by-key), ireduce, cholesky/lu/qr/svd dense (hipSOLVER), complex, math (incl.
all half transcendentals), norm, binary, approx, convolve, medfilt, random, set,
dot, reorder, sort, the sparse suite (sparse 86/86, sparse_convert 41/41,
sparse_arith 123/123, threading 9/9), blas 127/127 (incl. the int8 schar case),
confidence_connected 36/36, topk 110/110 and nearest_neighbour 122/122. The
topk hipCUB-BlockRadixSort LDS-aliasing fault and the nearest_neighbour/hamming
faults are fixed.

On RDNA3 (gfx1100, wave32) the FP32-complex POTRF reconstruction of a large
(n=1024) matrix drifts ~0.073 vs the 0.05 cfloat cholesky test eps -- the
recovered factor matches a double reference to FP32 precision (relative factor
error ~3e-9), so it is genuine FP32 accumulation drift (RDNA vs CDNA FMA order),
not a defect. test/cholesky_dense.cpp widens only the cfloat large-matrix eps to
0.1 on the RDNA HIP backend (detected at runtime via the device compute major);
float/double/cdouble and CUDA/gfx90a keep the strict 0.05.

Authored with the assistance of Claude (Anthropic).

Test Plan:
- Full afcuda + test build for gfx90a and for gfx1100
  (-DCMAKE_HIP_ARCHITECTURES=<arch>, -DAF_WITH_IMAGEIO=ON): PASS.
- Full CUDA.* gtest suite on one isolated GPU (gfx90a), 132/132:
  HIP_VISIBLE_DEVICES=2 ctest -R '_cuda$' -j1 --output-on-failure
  => 100% tests passed, 0 tests failed out of 132
- Sparse + the two closed residuals specifically (gfx90a):
  HIP_VISIBLE_DEVICES=2 ./test/sparse_cuda             # 86/86
  HIP_VISIBLE_DEVICES=2 ./test/sparse_convert_cuda     # 41/41
  HIP_VISIBLE_DEVICES=2 ./test/sparse_arith_cuda       # 123/123
  HIP_VISIBLE_DEVICES=2 ./test/threading_cuda          # 9/9 (Threading.Sparse)
  HIP_VISIBLE_DEVICES=2 ./test/blas_cuda               # 127/127 (incl. schar int8)
  HIP_VISIBLE_DEVICES=2 ./test/confidence_connected_cuda  # 36/36 (FreeImage)
- No regression on the previously-faulting suites (gfx90a):
  HIP_VISIBLE_DEVICES=2 ./test/topk_cuda               # 110/110
  HIP_VISIBLE_DEVICES=2 ./test/nearest_neighbour_cuda  # 122/122

The previous default pinned CMAKE_HIP_ARCHITECTURES to gfx90a before
enable_language(HIP), which preempted CMake's own host-GPU detection. A
user on a non-gfx90a AMD GPU who omitted -DCMAKE_HIP_ARCHITECTURES would
silently build gfx90a code objects that then fail to load on their card
at runtime ("no kernel image is available for execution on the device").

Remove the pin and rely on enable_language(HIP): it honors an explicit
-DCMAKE_HIP_ARCHITECTURES when given, otherwise auto-detects the host
GPU(s), and errors out when no GPU is found rather than guessing. Builds
that pass an explicit architecture are unchanged.

Authored with assistance from Claude.

The HIP backend's hipBLAS GemmEx call uses hipblasGemmEx with hipDataType
and hipblasComputeType_t (blas.cu, cudaDataType.hpp). That signature was
introduced in hipBLAS 3, which ships with ROCm 7.0. ROCm 6.x ships
hipBLAS 2, whose hipblasGemmEx takes hipblasDatatype_t and the older
compute-type enum, so the backend does not compile there (reported on
ROCm 6.4.4; ROCm 7.0.3 builds). State the minimum next to the
AF_BUILD_HIP wiring where the backend is documented.

Documentation only; no source or device code changes.

Authored with the assistance of Claude (Anthropic).

Test Plan:
Documentation-only change (a CMake comment). No build or GPU run needed;
verified inert with the MOAT regression classifier (comment/doc-only,
carries prior validation forward on every platform).