jeffdaily opened a new pull request, #1399: URL: https://github.com/apache/mahout/pull/1399
## Summary Adds an AMD GPU build of the QDP (Quantum Data Plane) native engine under `qdp/`, behind a Cargo `hip` feature (and `QDP_USE_HIP=1`). The default `cuda` feature is unchanged, so the NVIDIA build is byte-for-byte identical and nothing on the AMD path is reachable without opting in. The separate Triton AMD backend is orthogonal and untouched. This gives AMD parity on the native engine the project is built around (pinned-buffer pool, dual-stream overlap, in-Rust DLPack ownership). This change was authored with the assistance of Claude (Anthropic) and validated on real AMD GPU hardware (see Test Plan). Review it in two layers. ## Kernels (`qdp-kernels`) `build.rs` gains a HIP branch that compiles the same six `.cu` with hipcc, taking `--offload-arch` from `QDP_HIP_ARCH_LIST` (default `gfx90a` only when unset, never a literal that overrides the env, so other AMD targets build the same source by setting `QDP_HIP_ARCH_LIST` alone) and linking the AMD HIP runtime; the CUDA branch (nvcc) is untouched. hipcc ships no `<cuda_runtime.h>` / `<cuComplex.h>` / `<vector_types.h>`, so `qdp-kernels/hip_compat/` holds forwarding shim headers of those exact names, added to the include path ONLY on the HIP build, that map the small `cuda*` runtime + cuComplex surface the kernels use onto HIP. A CUDA build never sees that directory and pulls the real toolkit headers, so the `.cu` keep their CUDA spellings unchanged. `amplitude.cu` is the only kernel needing source fixes, both arch-unified: the `__shfl_down_sync` mask becomes 64-bit on HIP (ROCm static_asserts `sizeof(mask)==8`; the 32-bit literal fails to compile) while staying `0xffffffffu` on CUDA, and the warp-id `threadIdx.x >> 5` becomes `threadIdx.x / warpSize`. The latter is a genuine wave64 correctness fix: `>> 5` assumes 32-lane warps, so on a 64-lane (CDNA) warp the per-warp L2-norm partial landed in the wrong shared slot; `/ warpSize` is identical to `>> 5` on 32-lane hardware and correct on 64-lane hardware. ## Host (`qdp-core`) The cudarc crate is CUDA-only with no ROCm backend, so on the HIP build it is displaced by a thin HIP-runtime shim with the SAME type names and method signatures (`qdp-kernels/src/device.rs`: `CudaDevice`, `CudaSlice`, `CudaStream`, `DevicePtr`/`DevicePtrMut`/`DeviceSlice`, `DeviceRepr`/`ValidAsZeroBits`, backed by the AMD HIP runtime). `qdp-core/src/gpu_rt.rs` re-exports it as the single import point; every `use cudarc::driver` became `use crate::gpu_rt`, so the call sites compile unchanged on either vendor. The runtime FFI (`gpu/cuda_ffi.rs`) keeps its `cuda*` names and binds the matching `hip*` entry points under the `hip` feature. The async H2D path maps `cudaMemcpyAsync` to `hipMemcpyAsync` (its exact enqueue-and-return 1:1), NOT `hipMemcpyWithStream`, which synchronizes the stream and would block the host and silently serialize the dual-stream overlap pipeline; the copy-done event plus stream-wait still order copy before compute, so correctness is unchanged. DLPack tags exported tensors `kDLROCM` on the HIP build (a ROCm PyTorch's `from_dlpack` rejects a CUDA tag); the two device-type tests are made arch-aware. The GPU stack is selected by a build-script `qdp_gpu_platform` cfg (Linux always; Windows when the `hip` feature is on) rather than a `target_os == "linux"` proxy, which enables the same code on Windows ROCm; on Linux this cfg is always set, so the Linux output is identical. ## Docs and attribution `qdp/DEVELOPMENT.md` documents the ROCm/HIP build next to the existing CUDA flow, and `qdp/qdp-python/README.md` notes the native engine's AMD build. New and substantially-extended files for the AMD build carry an AMD copyright line below the Apache header and name the author. ## Test Plan Built and tested on gfx90a (MI250X, ROCm 7.2.1), gfx1100 (Radeon Pro W7800), and on Windows ROCm gfx1201 (RX 9070 XT) and gfx1151 (Radeon 8060S): ``` export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx90a ROCM_PATH=/opt/rocm cd qdp cargo build -p qdp-core -p qdp-kernels --no-default-features --features hip cargo test -p qdp-core -p qdp-kernels --no-default-features --features hip -- --test-threads=1 ``` All Rust tests RUN (they previously skipped on AMD, since cudarc found no device) and PASS, 0 failures: qdp-kernels 31 (amplitude 21, angle 10); qdp-core lib 77; GPU suites gpu_angle 12, gpu_api_workflow 8, gpu_basis 7, gpu_dlpack 9, gpu_fidelity 17, gpu_iqp 22, gpu_memory_safety 4, gpu_norm_f32 2, gpu_ptr_encoding 64, gpu_validation 8; plus the non-GPU regression suites (arrow/null/numpy/parquet/preprocessing/tensorflow/torch/types). The dual-stream async-pipeline tests pass with `QDP_ENABLE_OVERLAP_TRACKING=1`, confirming the non-blocking `hipMemcpyAsync` H2D path. The default `cuda` feature build is unchanged and remains the default; the AMD path is reachable only via `--features hip` / `QDP_USE_HIP=1`. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
