QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine - #1399
QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine#1399jeffdaily wants to merge 4 commits into
Conversation
There was a problem hiding this comment.
@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.)
|
btw, please check the precommit errors |
|
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 |
|
Thanks for the review, and for running the CUDA path on the 2080 Ti. All four addressed in 0b5042e:
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. |
|
@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 #1390's implicit-Hadamard Ozaki engine is the hard end. It uses So: runtime/memory/plain-kernel code ports about as easily as this PR did; anything built on tensor cores, |
There was a problem hiding this comment.
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!
| The first command is what `maturin develop --release` runs on CI; the | ||
| second verifies tests type-check in the CUDA build. | ||
|
|
||
| ### AMD GPU build (ROCm / HIP) |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
I think we should add a build github action workflow for hip!
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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.
|
@ryankert01 Thanks for running the full suite on the 3090 Ti -- good to have the default CUDA path vendor-confirmed. |
|
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
left a comment
ryankert01
left a comment
There was a problem hiding this comment.
btw, I think you can have a issue tracking this topic, because there's follow-ups worth taking track of.
Add an AMD ROCm/HIP build of the QDP native engine (qdp-kernels + qdp-core, with the qdp-python extension) behind a Cargo `hip` feature and the QDP_USE_HIP=1 build env. The default NVIDIA CUDA build is unchanged. To review: start with qdp-kernels/build.rs (the hipcc compile branch, selected by the `hip` feature or QDP_USE_HIP, targeting --offload-arch from QDP_HIP_ARCH_LIST) and qdp-kernels/hip_compat/ (forwarding shim headers named cuda_runtime.h / cuComplex.h / vector_types.h so the .cu sources keep their CUDA spelling; the shim dir is only on the HIP include path, so the CUDA build is untouched). Then qdp-kernels/src/device.rs and qdp-core/src/gpu_rt.rs: cudarc is CUDA-only with no ROCm backend, so on the `hip` feature a small self-contained shim provides the same CudaDevice/CudaSlice/CudaStream type names and method signatures backed by libamdhip64, and every call site swaps `cudarc::driver::` for `crate::gpu_rt::` with byte-identical bodies. Finally qdp-core/src/gpu/cuda_ffi.rs selects libcudart, the no-CUDA stubs, or the libamdhip64 wrappers at compile time. Two kernel correctness/compile fixes for AMD live in amplitude.cu, keyed on __HIP_PLATFORM_AMD__: the warp-reduction shuffle mask is 64-bit on HIP (ROCm static_asserts a 64-bit mask; the 32-bit literal fails to compile) and the per-warp shared-memory slot index uses threadIdx.x / warpSize instead of a hardcoded >> 5 (a real wave64 bug: the >> 5 assumes 32-lane warps, so the per-warp partial lands in the wrong slot on gfx90a). Both are arch-unified: identical to the original on 32-lane hardware, correct on 64-lane hardware. DLPack tags exported tensors kDLROCM on the HIP build. Reviewer follow-ups folded in: Drop rebinds the owning device before hipFree (multi-GPU safety); the pinned-pointer query reads the real hipMemoryType and hipPointerAttribute_t (not a magic enum value); build.rs rejects a QDP_USE_HIP vs `hip`-feature mismatch; forked streams are non-blocking (hipStreamNonBlocking) to match cudarc, with a default-stream synchronize on the blocking-copy paths to preserve CUDA-equivalent ordering. Authored with assistance from Claude (Anthropic). Test Plan: Built and tested on Linux gfx90a (MI250X, ROCm 7.2.1); the CUDA path is compile-checked without an NVIDIA toolkit via the qdp_no_cuda stubs. ``` export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx90a ROCM_PATH=/opt/rocm cargo build -p qdp-core -p qdp-kernels --no-default-features --features hip -j 16 cargo test -p qdp-core -p qdp-kernels --no-default-features --features hip -- --test-threads=1 cargo check -p qdp-kernels --no-default-features # neither backend: clean compile_error cargo check # default cuda feature, qdp_no_cuda stubs ```
Follow-up to the AMD/HIP port addressing reviewer findings. The GPU input validators (check_finite_batch_f32/_f64, validate_and_cast_basis_indices_f32, assert_basis_indices_in_range_usize) launch their flag-writing kernel on the caller's stream, then read the flag back with a default-stream copy. With the forked streams now non-blocking (hipStreamNonBlocking, to match cudarc), the null-stream readback is no longer ordered against the kernel, so the host could read a still-zero flag and let an out-of-range index through to a kernel that then writes out of bounds. Each validator now synchronizes the caller's stream before the readback, the same guard the amplitude encoder already uses; the CUDA path gains the same explicit ordering. The pinned-pointer query keeps a hand-rolled Rust mirror of hipPointerAttribute_t and the hipMemoryType enum values, both of which have shifted across ROCm major versions. A new hipcc-compiled translation unit (verify_pointer_attrs.cpp) static_asserts the enum values and the struct field offsets/size against the installed hip_runtime_api.h, so a header change fails the build loudly instead of silently misreading pointer attributes. Smaller fixes: cuComplex.h guards its make_cuComplex/cuCreal/... aliases with #ifndef so newer ROCm hip_complex.h (which ships those names) does not hit a redefinition; the device slice_mut offset widens to u64 before multiplying by the element size to avoid a usize wrap; CudaStream drops its Sync impl to match cudarc's Send-only contract; and the per-file AMD copyright/author lines and the NOTICE entry are removed per ASF source-header policy (git history preserves authorship). Authored with assistance from Claude (Anthropic). Test Plan: Linux gfx90a (MI250X, ROCm 7.2.1). ``` export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx90a ROCM_PATH=/opt/rocm cargo build -p qdp-core -p qdp-kernels --no-default-features --features hip -j 16 HIP_VISIBLE_DEVICES=2 cargo test -p qdp-core -p qdp-kernels \ --no-default-features --features hip -- --test-threads=1 # 358 passed, 0 failed cargo check -p qdp-kernels --no-default-features # clean "exactly one of cuda/hip" error QDP_NO_CUDA=1 cargo check -p qdp-core -p qdp-kernels # CUDA path (qdp_no_cuda stubs) cargo check -p qdp-kernels --features hip # cuda+hip both: hip precedence, no error ```
|
This pull request has been automatically marked as stale because it has not had any activity for 30 days. It will be closed in another 7 days if no further activity occurs. Thank you for your contribution. If you'd like to keep this open, leave any comment and the stale label will be removed. You can always ask for help on the Mahout dev mailing list or in GitHub Discussions. |
|
/unstale |
Upstream moved on since this branch was last synchronized, leaving the pull request conflicting. Merge current main and resolve the four conflicts so it applies cleanly again. Nothing is rebased: the previously reviewed commits stay exactly where they were. qdp-core/src/gpu/cuda_ffi.rs: upstream added a cudaGetDeviceCount binding and the cuda_runtime_available() helper that reports whether a usable device exists. Both are kept, and the HIP backend gains the matching hipGetDeviceCount wrapper, so the helper answers for AMD devices on the ROCm build instead of leaving an unresolved name. The three-way backend split (real CUDA runtime, no-toolkit stubs, HIP) is otherwise unchanged, and the stub backend already carried upstream's new entry point. qdp-core/src/gpu/encodings/amplitude.rs: kept upstream's shared validate_qubit_count import next to the qdp_gpu_platform cfg this branch uses in place of a bare target_os check, so the new qubit-limit validation applies on both vendors. qdp-core/tests/gpu_ptr_encoding.rs: kept upstream's new MAX_QUBITS imports for the excessive-qubit test. The device-pointer traits keep coming from qdp_core::gpu_rt, which resolves to cudarc on CUDA and to the HIP shim on ROCm. qdp-python/src/lib.rs: registered upstream's new cuda_available() function and kept the qdp_gpu_platform cfg on the loader and pipeline registrations. Authored with assistance from Claude (Anthropic). Test Plan: Linux gfx1100 (Radeon Pro W7800, ROCm 7.2.1) for the HIP path, and a CUDA 12.8 compile check on the same host (no NVIDIA GPU present, so CUDA is build-only). ``` export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx1100 ROCM_PATH=/opt/rocm cargo build -p qdp-core -p qdp-kernels --no-default-features --features hip -j 16 HIP_VISIBLE_DEVICES=0 cargo test -p qdp-core -p qdp-kernels \ --no-default-features --features hip -- --test-threads=1 # 368 passed, 0 failed ``` ``` export CUDA_PATH=/opt/conda/envs/cuda-12.8 PATH=$CUDA_PATH/bin:$PATH cargo build -p qdp-core -p qdp-kernels -j 16 # nvcc 12.8, exit 0 cargo test -p qdp-core -p qdp-kernels --no-run -j 16 # 23 test binaries ``` The AMD run includes upstream's newly merged suites: estimate (9) and parquet_f32_fidelity (4, including the three GPU fidelity cases), plus the new excessive-qubit case in gpu_ptr_encoding (69, was 68) and the estimate_memory doc example. That is 368 passing tests on this GPU, up from 353 before the merge.
The Parquet f32 fidelity suite that just merged gates its GPU module on target_os = "linux". Every other GPU test in the crate uses the qdp_gpu_platform cfg the build script emits, which is set on Linux and also on Windows when the HIP feature is on, so this one file would silently drop its three GPU cases on a Windows ROCm build while the rest of the suite runs. Use the same cfg here, and drop the "Linux + CUDA" wording from the header comments since the module now covers either vendor's GPU backend. Test only; no library or kernel code changes. Authored with assistance from Claude (Anthropic). Test Plan: Linux gfx1100 (Radeon Pro W7800, ROCm 7.2.1). ``` export QDP_USE_HIP=1 QDP_HIP_ARCH_LIST=gfx1100 ROCM_PATH=/opt/rocm HIP_VISIBLE_DEVICES=0 cargo test -p qdp-core --no-default-features --features hip \ --test parquet_f32_fidelity -- --test-threads=1 ``` 4 passed, 0 failed: the CPU reader-consistency case plus the three GPU fidelity cases at 8, 12 and 16 qubits.
|
@jeffdaily any updates? |
|
Note: this reply was drafted by an AI assistant. Conflicts resolved by merging current main into the branch (merge, not rebase, so the review history stands). Four conflicts, all at the seams between the new cuda_available()/max-qubit-validation work and the HIP path; each resolution keeps both sides. One addition was needed to compile: the new cuda_runtime_available() helper calls cudaGetDeviceCount, so the HIP build now maps that to hipGetDeviceCount alongside the other runtime entry points -- it correctly reports AMD devices too. One small follow-up commit: the new Parquet f32 fidelity tests gated their GPU module on target_os = "linux", while the rest of the GPU tests use the build-script cfg that also covers Windows ROCm builds; switched it to match so those three cases run there as well. Re-ran everything at the merged tip: Rust suite 368 passed / 0 failed on Linux (MI250X gfx90a, Radeon Pro W7800 gfx1100) and Windows (Radeon 8060S gfx1151); Python bindings rebuilt on Windows, 286 passed / 0 failed, with the new cuda_available() returning true on the HIP build; the default CUDA feature still compiles clean with nvcc 12.8. All of the new upstream tests (estimate, Parquet fidelity, the excessive-qubit case) run and pass on AMD hardware. |
Sorry for the delay, sincerely. This PR was prepared by our new MOAT project when it was still a prototype under my jeffdaily github org in June. It graduated to a supported AMD product last week after a brief outage. I've been trying to follow up with open PRs like yours as soon as I can, using MOAT's agent flow to address review comments, rerun all builds and tests across sufficient AMD hardware hosted on various servers, babysitting (agentsitting?) results, etc. Also, this PR was opened before it would have gotten the (now usual) disclosure:
Some of the comments I made above were done with AI assistance and weren't disclosed, sorry. Like I said, you caught me between prototype and release of MOAT, and we've refined our human-in-the-loop policy since then. This is actually me responding to you, to all of you reviewers, and I greatly appreciate your time on this PR review. |
5 participants
Summary
Adds an AMD GPU build of the QDP (Quantum Data Plane) native engine under
qdp/, behind a Cargohipfeature (andQDP_USE_HIP=1). The defaultcudafeature 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.rsgains a HIP branch that compiles the same six.cuwith hipcc, taking--offload-archfromQDP_HIP_ARCH_LIST(defaultgfx90aonly when unset, never a literal that overrides the env, so other AMD targets build the same source by settingQDP_HIP_ARCH_LISTalone) and linking the AMD HIP runtime; the CUDA branch (nvcc) is untouched.hipcc ships no
<cuda_runtime.h>/<cuComplex.h>/<vector_types.h>, soqdp-kernels/hip_compat/holds forwarding shim headers of those exact names, added to the include path ONLY on the HIP build, that map the smallcuda*runtime + cuComplex surface the kernels use onto HIP. A CUDA build never sees that directory and pulls the real toolkit headers, so the.cukeep their CUDA spellings unchanged.amplitude.cuis the only kernel needing source fixes, both arch-unified: the__shfl_down_syncmask becomes 64-bit on HIP (ROCm static_assertssizeof(mask)==8; the 32-bit literal fails to compile) while staying0xffffffffuon CUDA, and the warp-idthreadIdx.x >> 5becomesthreadIdx.x / warpSize. The latter is a genuine wave64 correctness fix:>> 5assumes 32-lane warps, so on a 64-lane (CDNA) warp the per-warp L2-norm partial landed in the wrong shared slot;/ warpSizeis identical to>> 5on 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.rsre-exports it as the single import point; everyuse cudarc::driverbecameuse crate::gpu_rt, so the call sites compile unchanged on either vendor.The runtime FFI (
gpu/cuda_ffi.rs) keeps itscuda*names and binds the matchinghip*entry points under thehipfeature. The async H2D path mapscudaMemcpyAsynctohipMemcpyAsync(its exact enqueue-and-return 1:1), NOThipMemcpyWithStream, 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
kDLROCMon the HIP build (a ROCm PyTorch'sfrom_dlpackrejects a CUDA tag); the two device-type tests are made arch-aware. The GPU stack is selected by a build-scriptqdp_gpu_platformcfg (Linux always; Windows when thehipfeature is on) rather than atarget_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.mddocuments the ROCm/HIP build next to the existing CUDA flow, andqdp/qdp-python/README.mdnotes 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):
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-blockinghipMemcpyAsyncH2D path.The default
cudafeature build is unchanged and remains the default; the AMD path is reachable only via--features hip/QDP_USE_HIP=1.