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