QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine

[jeffdaily]

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 behavior-preserving (no functional change) 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.

[rich7420]

@jeffdaily thanks for the patch!!! Welcome to Mahout
some comments left:

Nice port — the gpu_rt indirection + scoped hip_compat shims + the wave64 >>5→/warpSize fix are clean.

I ran the cuda path on a 2080 Ti: build + lib/fidelity/amplitude/angle all green, so default path's healthy. No AMD HW here so the HIP notes are read-only — flagging for your call. 4 inline comments; (Drop frees without binding the device) is the only one I'd want fixed before merge. Rest are minor / sanity-checks.

(Tiny nit: "CUDA byte-for-byte unchanged" isn't quite true — metrics.rs swapped driver cuMemcpyDtoH_v2→runtime cudaMemcpy, and the kernel >>5→/warpSize changes SASS. Both verified behavior-identical, so no action.)

[rich7420]

btw, please check the precommit errors

[ryankert01]

I'm also curious about complex optimizations can or cannot be converted this easy. I think recently someone open PRs that have some real profound ones. #1390

[jeffdaily]

Thanks for the review, and for running the CUDA path on the 2080 Ti. All four addressed in 0b5042e:

• Drop now binds the owning device before hipFree (the merge-blocker).
• The hipMemoryType check reads hipPointerAttribute_t and compares against hipMemoryTypeDevice instead of the hardcoded 2.
• build.rs fails loudly when QDP_USE_HIP and the hip feature disagree.
• The stream is now non-blocking (hipStreamCreateWithFlags(.., hipStreamNonBlocking)) to match cudarc.

On #4: making the stream non-blocking surfaced a pre-existing latent race in the batch-f32 amplitude path -- it read the norm back on the default stream without syncing the caller's stream, masked until now by the blocking stream. Fixed with the same stream-sync the other batch paths use; it's shared (non-HIP) code, so it applies to the CUDA path too.

On byte-for-byte: that's overstated, now corrected in the PR description. metrics.rs swapped the driver cuMemcpyDtoH_v2 for the runtime cudaMemcpy, the kernel >>5 became /warpSize, and the #4 sync adds one more -- all SASS-changing but behavior-identical. The CUDA path is behavior-preserving, not literally byte/SASS-identical.

[jeffdaily]

@rich7420 Fixed in 6d2de29 -- ran cargo fmt over the HIP-path files (import order + line wrapping). Clippy is clean on the Rust side; the one remaining compiler warning (iqp.cu unused parameter) is pre-existing upstream, untouched by this PR.

[jeffdaily]

@ryankert01 This PR was the easy end: hand-written kernels plus the CUDA runtime/driver API, which hipify mostly mechanically -- the one real subtlety was wave size (the >>5 -> /warpSize fix, since CDNA warps are 64-wide).

#1390's implicit-Hadamard Ozaki engine is the hard end. It uses nvcuda::wmma and raw inline PTX (mma.sync.aligned.m16n8k32...s8.s8.s32) -- int8 tensor-core MMA -- to do an Ozaki accurate-FWT. None of that hipifies: inline PTX doesn't compile under hipcc, and it has to be rewritten against AMD matrix cores (MFMA on CDNA / rocWMMA), whose fragment shapes differ (gfx90a MFMA isn't m16n8k32), so the Ozaki tiling and the pre/post layout shuffles have to be re-derived for the AMD fragment geometry. The Ozaki scheme as an algorithm is portable; its whole payoff is fast low-precision matrix-core matmul, so it's a real reimplementation of that path, not a translation.

So: runtime/memory/plain-kernel code ports about as easily as this PR did; anything built on tensor cores, mma/PTX, or CUTLASS/CuTe needs a from-scratch AMD matrix-core path. Convertible, but not "this easy."

[ryankert01]

I checked out this head and ran the complete Rust suite on NVIDIA hardware (RTX 3090 Ti, CUDA toolkit): 316 passed / 0 failed, including all GPU suites and the dual-stream pipeline tests with QDP_ENABLE_OVERLAP_TRACKING=1; cargo clippy --all-targets is clean. So the refactored default CUDA path is verified on the vendor side the author could not test on.

Inline comments below cover the issues and risks found in review. The only one I'd hold merge for is the ASF header/NOTICE question; the rest are small follow-ups that could land here or later.

Please also fix the pre-commit hook!

[ryankert01]

Maintenance risk worth acknowledging: CI has no AMD runners, so after merge the HIP path is never exercised and can silently rot as the CUDA path evolves. The gpu_rt seam minimizes divergence pressure, but regressions will only surface when someone rebuilds on AMD hardware. A compile-only hipcc job (ROCm apt packages on ubuntu runners) could be a cheap follow-up.

[ryankert01]

I think we should add a build github action workflow for hip!

[jeffdaily]

Agreed -- the gpu_rt seam limits divergence, but without AMD runners the HIP path can still bit-rot. I'll put up a compile-only hipcc job (ROCm apt packages on an ubuntu runner) as a separate follow-up PR so the build is at least guarded.

[jeffdaily]

Acknowledged. With no AMD runners a CPU-only hipcc job would compile-check but never exercise the GPU path where the real regressions show up, so I would rather not add a workflow that looks like coverage it cannot provide. The gpu_rt seam keeps the CUDA and HIP paths close, and the change is validated on gfx90a hardware. Not filing a tracking issue.

[jeffdaily]

@ryankert01 Thanks for running the full suite on the 3090 Ti -- good to have the default CUDA path vendor-confirmed.

[jeffdaily]

Wanted to apologize for the recent churn here. Normally I would have done this BEFORE opening the PR, but I didn't fully understand the problem. I would make a fix on Linux-gfx90a, then validated it on Win-gfx1201, then Win-gfx1201 would make a change and send it back to Linux-gfx90a, and so on. However, I finally root-caused the blocking vs non-blocking stream issues. This commit d97db73 should take care of it. See the commit message there for details.

[ryankert01]

btw, I think you can have a issue tracking this topic, because there's follow-ups worth taking track of.

[ryankert01]

Need to resolve conflicts