You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mahout.apache.org by "jeffdaily (via GitHub)" <gi...@apache.org> on 2026/06/11 06:53:49 UTC
[PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397737008
##########
qdp/qdp-kernels/Cargo.toml:
##########
@@ -4,11 +4,18 @@ version.workspace = true
edition.workspace = true
[dependencies]
-cudarc = { workspace = true }
+cudarc = { workspace = true, optional = true }
[build-dependencies]
cc = { workspace = true }
+[features]
+# Default build: NVIDIA CUDA via cudarc + nvcc-compiled kernels.
+default = ["cuda"]
+cuda = ["dep:cudarc"]
+# AMD build: hipcc-compiled kernels + the in-crate HIP device traits; no cudarc.
+hip = []
Review Comment:
Added in dbee1e1 -- a note in both qdp-kernels and qdp-core `[features]` that `cuda` and `hip` are mutually exclusive and `hip` takes precedence (kernels build for HIP; cudarc is compiled but unused).
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "rich7420 (via GitHub)" <gi...@apache.org>.
rich7420 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3394439508
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
Review Comment:
`hipFree` hits the calling thread's *current* device, but Drop never re-binds (unlike `alloc`, which does `self.bind()?`). Multi-GPU: dropping while another device is current frees the wrong device — silent leak or wrong-pointer free. cudarc binds in Drop. A best-effort `self._device.bind()` before `hipFree` fixes it.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
+
+ impl Drop for CudaStream {
+ fn drop(&mut self) {
+ if !self.stream.is_null() {
+ unsafe {
+ let _ = hipStreamDestroy(self.stream);
+ }
+ }
+ }
+ }
+
+ /// HIP device handle. Stand-in for `cudarc::CudaDevice`; created via
+ /// `CudaDevice::new(ordinal)` and shared as `Arc<CudaDevice>` exactly like
+ /// cudarc (whose `new` already returns the `Arc`).
+ pub struct CudaDevice {
+ ordinal: usize,
+ }
+
+ impl CudaDevice {
+ /// Select device `ordinal` and return a shared handle, or an error if no
+ /// such device exists. Matches `cudarc::CudaDevice::new`.
+ pub fn new(ordinal: usize) -> Result<Arc<Self>, DriverError> {
+ unsafe {
+ let mut count: i32 = 0;
+ check(hipGetDeviceCount(&mut count))?;
+ if ordinal as i32 >= count {
+ return Err(DriverError(101)); // hipErrorInvalidDevice
+ }
+ check(hipSetDevice(ordinal as i32))?;
+ }
+ Ok(Arc::new(Self { ordinal }))
+ }
+
+ pub fn ordinal(&self) -> usize {
+ self.ordinal
+ }
+
+ fn bind(&self) -> Result<(), DriverError> {
+ unsafe { check(hipSetDevice(self.ordinal as i32)) }
+ }
+
+ /// Allocate `len` uninitialized elements of `T` on the device.
+ ///
+ /// # Safety
+ /// Contents are uninitialized until written, mirroring `cudarc`'s
+ /// `unsafe fn alloc`.
+ pub unsafe fn alloc<T>(self: &Arc<Self>, len: usize) -> Result<CudaSlice<T>, DriverError> {
+ self.bind()?;
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ let mut ptr: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipMalloc(&mut ptr, bytes.max(1)))?;
+ }
+ Ok(CudaSlice {
+ ptr: ptr as u64,
+ len,
+ _device: Arc::clone(self),
+ _marker: PhantomData,
+ })
+ }
+
+ /// Allocate `len` zero-initialized elements of `T` on the device.
+ pub fn alloc_zeros<T: ValidAsZeroBits>(
+ self: &Arc<Self>,
+ len: usize,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let slice = unsafe { self.alloc::<T>(len)? };
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemset(slice.raw_ptr(), 0, bytes))?;
+ }
+ }
+ Ok(slice)
+ }
+
+ /// Copy a host slice to a freshly allocated device buffer (blocking).
+ pub fn htod_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &[T],
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let mut slice = unsafe { self.alloc::<T>(src.len())? };
+ self.htod_sync_copy_into(src, &mut slice)?;
+ Ok(slice)
+ }
+
+ /// Copy an owned host Vec to a freshly allocated device buffer. cudarc's
+ /// `htod_copy` keeps the Vec alive until an async copy completes; our
+ /// copy is synchronous (blocking hipMemcpy), so the Vec can be dropped
+ /// on return with identical observable behavior.
+ pub fn htod_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: Vec<T>,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ self.htod_sync_copy(&src)
+ }
+
+ /// Copy a host slice into an existing device buffer or sub-view
+ /// (blocking). Accepts any `DevicePtrMut` target so both `CudaSlice` and
+ /// the `slice_mut` view work, matching cudarc's generic destination.
+ pub fn htod_sync_copy_into<T: DeviceRepr, D: DevicePtrMut<T> + DeviceSlice<T>>(
+ self: &Arc<Self>,
+ src: &[T],
+ dst: &mut D,
+ ) -> Result<(), DriverError> {
+ assert_eq!(
+ dst.len(),
+ src.len(),
+ "htod_sync_copy_into: dst.len() != src.len()"
+ );
+ self.bind()?;
+ let bytes = std::mem::size_of_val(src);
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ (*dst.device_ptr_mut()) as *mut c_void,
+ src.as_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_HOST_TO_DEVICE,
+ ))?;
+ }
+ }
+ Ok(())
+ }
+
+ /// Copy a device buffer back to a freshly allocated host Vec (blocking).
+ ///
+ /// Matches cudarc's bound of just `DeviceRepr` (no `Default`): the Vec is
+ /// allocated uninitialized and every byte is written by the copy before
+ /// its length is set, which is sound because `DeviceRepr` is plain data.
+ pub fn dtoh_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &CudaSlice<T>,
+ ) -> Result<Vec<T>, DriverError> {
+ self.bind()?;
+ let len = src.len;
+ let mut out: Vec<T> = Vec::with_capacity(len);
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ out.as_mut_ptr() as *mut c_void,
+ src.raw_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_DEVICE_TO_HOST,
+ ))?;
+ }
+ }
+ unsafe {
+ out.set_len(len);
+ }
+ Ok(out)
+ }
+
+ /// Block until all work on the device's default stream completes.
+ pub fn synchronize(&self) -> Result<(), DriverError> {
+ self.bind()?;
+ unsafe { check(hipDeviceSynchronize()) }
+ }
+
+ /// Create a new stream tied to this device. Mirrors
+ /// `cudarc::CudaDevice::fork_default_stream`.
+ pub fn fork_default_stream(self: &Arc<Self>) -> Result<CudaStream, DriverError> {
+ self.bind()?;
+ let mut stream: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipStreamCreate(&mut stream))?;
Review Comment:
`hipStreamCreate` is the *blocking* stream ctor; cudarc forks a non-blocking one (`hipStreamCreateWithFlags(.., hipStreamNonBlocking)`). Could quietly serialize the H2D/compute overlap. Your overlap test passes, so just worth a confirm.
##########
qdp/qdp-kernels/build.rs:
##########
@@ -154,9 +171,76 @@ fn apply_default_arch_targets(build: &mut cc::Build) {
}
}
+fn hip_requested() -> bool {
+ if cfg!(feature = "hip") {
+ return true;
+ }
+ env::var("QDP_USE_HIP")
Review Comment:
`QDP_USE_HIP` flips the *kernel* build to hipcc even with the cargo `hip` feature off → a default `cargo build` with this env set = AMD kernels + cudarc host = mismatch. Maybe error when `QDP_USE_HIP` and `CARGO_FEATURE_HIP` disagree, so it fails loud.
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +50,201 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ // hipPointerAttribute_t has the same first fields we read (memoryType,
+ // device, devicePointer, hostPointer); a #[repr(C)] alias suffices since we
+ // only read memory_type/device.
+ unsafe extern "C" {
+ fn hipHostMalloc(ptr: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ fn hipHostFree(ptr: *mut c_void) -> i32;
+ fn hipPointerGetAttributes(attributes: *mut c_void, ptr: *const c_void) -> i32;
+ fn hipMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+ fn hipMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, count: usize, kind: u32) -> i32;
+ fn hipEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ fn hipEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ fn hipEventDestroy(event: *mut c_void) -> i32;
+ fn hipStreamWaitEvent(stream: *mut c_void, event: *mut c_void, flags: u32) -> i32;
+ fn hipStreamSynchronize(stream: *mut c_void) -> i32;
+ fn hipMemsetAsync(dst: *mut c_void, value: i32, count: usize, stream: *mut c_void) -> i32;
+ fn hipEventQuery(event: *mut c_void) -> i32;
+ fn hipEventSynchronize(event: *mut c_void) -> i32;
+ fn hipEventElapsedTime(ms: *mut f32, start: *mut c_void, end: *mut c_void) -> i32;
+ }
+
+ // hipHostMallocDefault == 0, matching cudaHostAllocDefault used by callers.
+ pub(crate) unsafe fn cudaHostAlloc(p: *mut *mut c_void, size: usize, flags: u32) -> i32 {
+ unsafe { hipHostMalloc(p, size, flags) }
+ }
+ pub(crate) unsafe fn cudaFreeHost(ptr: *mut c_void) -> i32 {
+ unsafe { hipHostFree(ptr) }
+ }
#[allow(dead_code)]
- pub(crate) fn cudaPointerGetAttributes(
+ pub(crate) unsafe fn cudaPointerGetAttributes(
Review Comment:
Heads up — this is live (lib.rs:82), not dead code. The cast assumes `hipMemoryType` enum values == CUDA's (`DEVICE=2`), which has differed across ROCm versions. Fine on 7.2.1, but a version landmine — safer to read the real `hipPointerAttribute_t` and map the enum explicitly.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4684240992
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. However, I finally root-caused the blocking vs non-blocking stream issues. This commit https://github.com/apache/mahout/pull/1399/commits/d97db73adeac70b377c22e17fde6af7ff4ff3057 should take care of it. See the commit message there for details.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397540559
##########
qdp/qdp-core/src/gpu/encodings/amplitude.rs:
##########
@@ -877,6 +877,13 @@ impl AmplitudeEncoder {
{
crate::profile_scope!("GPU::NormValidationF32");
+ // The norm kernel ran on the caller's stream, but dtoh_sync_copy
+ // reads back on the default stream. Synchronize the caller's stream
+ // first so the result is visible: with a non-blocking stream (which
+ // does not implicitly order against the default stream) the readback
+ // would otherwise race and observe the zero-initialized buffer. This
+ // mirrors the single-sample path (calculate_inv_norm_gpu_with_stream).
+ sync_cuda_stream(stream, "Norm stream synchronize failed (batch f32)")?;
Review Comment:
Good fix — but worth being precise that this executes on the **CUDA path too**, so it is a (correct) behavior change there, not just HIP enablement. The norm kernel runs on the caller's non-blocking stream while `dtoh_sync_copy` reads back on the legacy default stream, which non-blocking streams don't order against — so the readback could race and observe the zeroed buffer on NVIDIA as well. The second commit message states this accurately ("latent cross-stream ordering hazard shared with the CUDA path"), but the PR description's "the NVIDIA build is behavior-preserving (no functional change)" slightly oversells it. Suggest flagging this as a CUDA-path bug fix in release notes.
Verified on NVIDIA at this head (RTX 3090 Ti): full suite 316 passed / 0 failed, including the dual-stream tests with `QDP_ENABLE_OVERLAP_TRACKING=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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397247837
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
Review Comment:
Fixed in 0b5042e -- Drop now does a best-effort `self._device.bind()` before `hipFree` (result swallowed, since Drop can't fail), matching the alloc path and cudarc.
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +50,201 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ // hipPointerAttribute_t has the same first fields we read (memoryType,
+ // device, devicePointer, hostPointer); a #[repr(C)] alias suffices since we
+ // only read memory_type/device.
+ unsafe extern "C" {
+ fn hipHostMalloc(ptr: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ fn hipHostFree(ptr: *mut c_void) -> i32;
+ fn hipPointerGetAttributes(attributes: *mut c_void, ptr: *const c_void) -> i32;
+ fn hipMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+ fn hipMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, count: usize, kind: u32) -> i32;
+ fn hipEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ fn hipEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ fn hipEventDestroy(event: *mut c_void) -> i32;
+ fn hipStreamWaitEvent(stream: *mut c_void, event: *mut c_void, flags: u32) -> i32;
+ fn hipStreamSynchronize(stream: *mut c_void) -> i32;
+ fn hipMemsetAsync(dst: *mut c_void, value: i32, count: usize, stream: *mut c_void) -> i32;
+ fn hipEventQuery(event: *mut c_void) -> i32;
+ fn hipEventSynchronize(event: *mut c_void) -> i32;
+ fn hipEventElapsedTime(ms: *mut f32, start: *mut c_void, end: *mut c_void) -> i32;
+ }
+
+ // hipHostMallocDefault == 0, matching cudaHostAllocDefault used by callers.
+ pub(crate) unsafe fn cudaHostAlloc(p: *mut *mut c_void, size: usize, flags: u32) -> i32 {
+ unsafe { hipHostMalloc(p, size, flags) }
+ }
+ pub(crate) unsafe fn cudaFreeHost(ptr: *mut c_void) -> i32 {
+ unsafe { hipHostFree(ptr) }
+ }
#[allow(dead_code)]
- pub(crate) fn cudaPointerGetAttributes(
+ pub(crate) unsafe fn cudaPointerGetAttributes(
Review Comment:
Fixed in 0b5042e -- reads `hipPointerAttribute_t` and compares the `type` field against the named `hipMemoryTypeDevice`/`hipMemoryTypeManaged` constants, instead of assuming the enum value matches CUDA's `2`.
##########
qdp/qdp-kernels/build.rs:
##########
@@ -154,9 +171,76 @@ fn apply_default_arch_targets(build: &mut cc::Build) {
}
}
+fn hip_requested() -> bool {
+ if cfg!(feature = "hip") {
+ return true;
+ }
+ env::var("QDP_USE_HIP")
Review Comment:
Fixed in 0b5042e -- added a check that panics with a clear message when `QDP_USE_HIP` and `CARGO_FEATURE_HIP` disagree in either direction, so the kernel and host builds can't silently mismatch.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
+
+ impl Drop for CudaStream {
+ fn drop(&mut self) {
+ if !self.stream.is_null() {
+ unsafe {
+ let _ = hipStreamDestroy(self.stream);
+ }
+ }
+ }
+ }
+
+ /// HIP device handle. Stand-in for `cudarc::CudaDevice`; created via
+ /// `CudaDevice::new(ordinal)` and shared as `Arc<CudaDevice>` exactly like
+ /// cudarc (whose `new` already returns the `Arc`).
+ pub struct CudaDevice {
+ ordinal: usize,
+ }
+
+ impl CudaDevice {
+ /// Select device `ordinal` and return a shared handle, or an error if no
+ /// such device exists. Matches `cudarc::CudaDevice::new`.
+ pub fn new(ordinal: usize) -> Result<Arc<Self>, DriverError> {
+ unsafe {
+ let mut count: i32 = 0;
+ check(hipGetDeviceCount(&mut count))?;
+ if ordinal as i32 >= count {
+ return Err(DriverError(101)); // hipErrorInvalidDevice
+ }
+ check(hipSetDevice(ordinal as i32))?;
+ }
+ Ok(Arc::new(Self { ordinal }))
+ }
+
+ pub fn ordinal(&self) -> usize {
+ self.ordinal
+ }
+
+ fn bind(&self) -> Result<(), DriverError> {
+ unsafe { check(hipSetDevice(self.ordinal as i32)) }
+ }
+
+ /// Allocate `len` uninitialized elements of `T` on the device.
+ ///
+ /// # Safety
+ /// Contents are uninitialized until written, mirroring `cudarc`'s
+ /// `unsafe fn alloc`.
+ pub unsafe fn alloc<T>(self: &Arc<Self>, len: usize) -> Result<CudaSlice<T>, DriverError> {
+ self.bind()?;
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ let mut ptr: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipMalloc(&mut ptr, bytes.max(1)))?;
+ }
+ Ok(CudaSlice {
+ ptr: ptr as u64,
+ len,
+ _device: Arc::clone(self),
+ _marker: PhantomData,
+ })
+ }
+
+ /// Allocate `len` zero-initialized elements of `T` on the device.
+ pub fn alloc_zeros<T: ValidAsZeroBits>(
+ self: &Arc<Self>,
+ len: usize,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let slice = unsafe { self.alloc::<T>(len)? };
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemset(slice.raw_ptr(), 0, bytes))?;
+ }
+ }
+ Ok(slice)
+ }
+
+ /// Copy a host slice to a freshly allocated device buffer (blocking).
+ pub fn htod_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &[T],
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let mut slice = unsafe { self.alloc::<T>(src.len())? };
+ self.htod_sync_copy_into(src, &mut slice)?;
+ Ok(slice)
+ }
+
+ /// Copy an owned host Vec to a freshly allocated device buffer. cudarc's
+ /// `htod_copy` keeps the Vec alive until an async copy completes; our
+ /// copy is synchronous (blocking hipMemcpy), so the Vec can be dropped
+ /// on return with identical observable behavior.
+ pub fn htod_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: Vec<T>,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ self.htod_sync_copy(&src)
+ }
+
+ /// Copy a host slice into an existing device buffer or sub-view
+ /// (blocking). Accepts any `DevicePtrMut` target so both `CudaSlice` and
+ /// the `slice_mut` view work, matching cudarc's generic destination.
+ pub fn htod_sync_copy_into<T: DeviceRepr, D: DevicePtrMut<T> + DeviceSlice<T>>(
+ self: &Arc<Self>,
+ src: &[T],
+ dst: &mut D,
+ ) -> Result<(), DriverError> {
+ assert_eq!(
+ dst.len(),
+ src.len(),
+ "htod_sync_copy_into: dst.len() != src.len()"
+ );
+ self.bind()?;
+ let bytes = std::mem::size_of_val(src);
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ (*dst.device_ptr_mut()) as *mut c_void,
+ src.as_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_HOST_TO_DEVICE,
+ ))?;
+ }
+ }
+ Ok(())
+ }
+
+ /// Copy a device buffer back to a freshly allocated host Vec (blocking).
+ ///
+ /// Matches cudarc's bound of just `DeviceRepr` (no `Default`): the Vec is
+ /// allocated uninitialized and every byte is written by the copy before
+ /// its length is set, which is sound because `DeviceRepr` is plain data.
+ pub fn dtoh_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &CudaSlice<T>,
+ ) -> Result<Vec<T>, DriverError> {
+ self.bind()?;
+ let len = src.len;
+ let mut out: Vec<T> = Vec::with_capacity(len);
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ out.as_mut_ptr() as *mut c_void,
+ src.raw_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_DEVICE_TO_HOST,
+ ))?;
+ }
+ }
+ unsafe {
+ out.set_len(len);
+ }
+ Ok(out)
+ }
+
+ /// Block until all work on the device's default stream completes.
+ pub fn synchronize(&self) -> Result<(), DriverError> {
+ self.bind()?;
+ unsafe { check(hipDeviceSynchronize()) }
+ }
+
+ /// Create a new stream tied to this device. Mirrors
+ /// `cudarc::CudaDevice::fork_default_stream`.
+ pub fn fork_default_stream(self: &Arc<Self>) -> Result<CudaStream, DriverError> {
+ self.bind()?;
+ let mut stream: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipStreamCreate(&mut stream))?;
Review Comment:
Fixed in 0b5042e -- now `hipStreamCreateWithFlags(.., hipStreamNonBlocking)` to match cudarc. Making it non-blocking exposed a latent race in the batch-f32 amplitude readback (default-stream copy without syncing the caller's stream, previously masked by the blocking stream), fixed in the same commit. Overlap tests still pass.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5320133772
> @jeffdaily any updates?
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:
> This pull request was prepared with the help of an AI assistant acting as a coding agent and was read and approved by a person before it was opened. It comes from an ongoing effort to add AMD GPU support to widely used CUDA projects, one repository at a time: https://github.com/AMD-Ecosystem/moat -- that repository describes how the work is done and what a person checks before anything is submitted.
>
> If you would rather not receive pull requests from this effort, say so here or open an issue at https://github.com/AMD-Ecosystem/moat/issues/new/choose and we will close this and stop. That can cover this repository alone or everything you own, whichever you prefer.
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.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "viiccwen (via GitHub)" <gi...@apache.org>.
viiccwen commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5307553266
@jeffdaily any updates?
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5319996318
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.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397540559
##########
qdp/qdp-core/src/gpu/encodings/amplitude.rs:
##########
@@ -877,6 +877,13 @@ impl AmplitudeEncoder {
{
crate::profile_scope!("GPU::NormValidationF32");
+ // The norm kernel ran on the caller's stream, but dtoh_sync_copy
+ // reads back on the default stream. Synchronize the caller's stream
+ // first so the result is visible: with a non-blocking stream (which
+ // does not implicitly order against the default stream) the readback
+ // would otherwise race and observe the zero-initialized buffer. This
+ // mirrors the single-sample path (calculate_inv_norm_gpu_with_stream).
+ sync_cuda_stream(stream, "Norm stream synchronize failed (batch f32)")?;
Review Comment:
Good fix — but worth being precise that this executes on the **CUDA path too**, so it is a (correct) behavior change there, not just HIP enablement. The norm kernel runs on the caller's non-blocking stream while `dtoh_sync_copy` reads back on the legacy default stream, which non-blocking streams don't order against — so the readback could race and observe the zeroed buffer on NVIDIA as well. The second commit message states this accurately ("latent cross-stream ordering hazard shared with the CUDA path"), but the PR description's "the NVIDIA build is behavior-preserving (no functional change)" slightly oversells it. Suggest flagging this as a CUDA-path bug fix in release notes.
Verified on NVIDIA at this head (RTX 3090 Ti): full suite 316 passed / 0 failed, including the dual-stream tests with `QDP_ENABLE_OVERLAP_TRACKING=1`.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,418 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
Review Comment:
**ASF source-header policy (applies to all 10 files carrying this header):** the PR adds `Portions Copyright (c) 2026 Advanced Micro Devices, Inc.` and `Author:` lines beneath the ASF header here and in `gpu_rt.rs`, `cuda_ffi.rs`, `qdp-kernels/build.rs`, `qdp-python/build.rs`, `amplitude.cu`, `kernel_compat.h`, and the three `hip_compat/` headers.
Per https://www.apache.org/legal/src-headers.html, contributions submitted with copyright notices should have those notices removed or relocated to the project `NOTICE` file, and author tags in source files are discouraged across ASF projects. There is no precedent for either in this repo, and `NOTICE` currently has no AMD entry.
This needs a deliberate call before merge — likely either move the copyright line to `NOTICE` or drop it, and drop the `Author:` lines (git history preserves attribution).
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +50,247 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::{CUDA_MEMORY_TYPE_DEVICE, CUDA_MEMORY_TYPE_MANAGED, CudaPointerAttributes};
+ use std::ffi::c_void;
+
+ // hipMemoryType enum values are NOT guaranteed equal to CUDA's across ROCm
+ // releases (older HIP used Host=0/Device=1; the hip_runtime_api.h note flags
+ // this explicitly). So we read the real hipPointerAttribute_t and compare its
+ // `type` against the named hipMemoryType* constants rather than a magic
+ // number, then translate to the CUDA convention the caller expects.
+ const HIP_MEMORY_TYPE_DEVICE: i32 = 2; // hipMemoryTypeDevice
Review Comment:
These constants pin the **ROCm 6+** `hipMemoryType` convention — as the comment above notes, ROCm 5.x used `Device=1` — and this is live code: `validate_cuda_input_ptr` calls `cudaPointerGetAttributes` on every `encode*_from_gpu_ptr` entry point (the 64 `gpu_ptr_encoding` tests exercise it). On ROCm 5.x, valid device pointers would be rejected as "not device memory".
So the effective floor is ROCm >= 6.0. Suggest stating that in the DEVELOPMENT.md prerequisites (testing was on 7.2.1).
##########
qdp/qdp-kernels/src/lib.rs:
##########
@@ -20,6 +20,9 @@
use std::ffi::c_void;
+pub mod device;
+use device::{DeviceRepr, ValidAsZeroBits};
Review Comment:
With neither `cuda` nor `hip` enabled, this import is unresolved: `cargo check -p qdp-kernels --no-default-features` fails with a raw `E0432` (reproduced locally). Before this PR that configuration built, since cudarc was unconditional.
Consider a guard in `device.rs` for a clean diagnostic:
```rust
#[cfg(not(any(feature = "cuda", feature = "hip")))]
compile_error!("qdp-kernels requires exactly one of the `cuda` or `hip` features");
```
Non-blocking — that config was only incidentally buildable before.
##########
qdp/qdp-kernels/Cargo.toml:
##########
@@ -4,11 +4,18 @@ version.workspace = true
edition.workspace = true
[dependencies]
-cudarc = { workspace = true }
+cudarc = { workspace = true, optional = true }
[build-dependencies]
cc = { workspace = true }
+[features]
+# Default build: NVIDIA CUDA via cudarc + nvcc-compiled kernels.
+default = ["cuda"]
+cuda = ["dep:cudarc"]
+# AMD build: hipcc-compiled kernels + the in-crate HIP device traits; no cudarc.
+hip = []
Review Comment:
`cuda` and `hip` aren't additive: with both enabled (e.g. `--features hip` without `--no-default-features`, or via workspace feature unification), `hip` silently wins everywhere (`device.rs`, `cuda_ffi.rs`, the kernel build) while cudarc still compiles, unused. The `QDP_USE_HIP` consistency panic in build.rs keeps the outcome coherent, which is good — but a one-line note here (and in qdp-core's `[features]`) that the two are mutually exclusive with `hip` taking precedence would prevent surprises.
##########
qdp/DEVELOPMENT.md:
##########
@@ -114,6 +114,43 @@ cd ..
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)
Review 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.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,418 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+ // hipStreamNonBlocking: the new stream does not implicitly synchronize with
+ // the NULL/default stream, matching cudarc's fork_default_stream.
+ const HIP_STREAM_NON_BLOCKING: u32 = 1;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32) -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreateWithFlags(stream: *mut *mut c_void, flags: u32) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(
+ range.start <= range.end && range.end <= self.len,
+ "slice_mut out of bounds"
+ );
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ // hipFree releases on the calling thread's current device, so
+ // re-bind the owning device first (cudarc does the same in Drop):
+ // on multi-GPU a different device may be current, which would
+ // otherwise free against the wrong device. Best-effort -- Drop
+ // cannot report an error, so a failed bind is swallowed.
+ let _ = self._device.bind();
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
+
+ impl Drop for CudaStream {
+ fn drop(&mut self) {
+ if !self.stream.is_null() {
+ unsafe {
+ let _ = hipStreamDestroy(self.stream);
+ }
+ }
+ }
+ }
+
+ /// HIP device handle. Stand-in for `cudarc::CudaDevice`; created via
+ /// `CudaDevice::new(ordinal)` and shared as `Arc<CudaDevice>` exactly like
+ /// cudarc (whose `new` already returns the `Arc`).
+ pub struct CudaDevice {
+ ordinal: usize,
+ }
+
+ impl CudaDevice {
+ /// Select device `ordinal` and return a shared handle, or an error if no
+ /// such device exists. Matches `cudarc::CudaDevice::new`.
+ pub fn new(ordinal: usize) -> Result<Arc<Self>, DriverError> {
+ unsafe {
+ let mut count: i32 = 0;
+ check(hipGetDeviceCount(&mut count))?;
+ if ordinal as i32 >= count {
+ return Err(DriverError(101)); // hipErrorInvalidDevice
+ }
+ check(hipSetDevice(ordinal as i32))?;
+ }
+ Ok(Arc::new(Self { ordinal }))
+ }
+
+ pub fn ordinal(&self) -> usize {
+ self.ordinal
+ }
+
+ fn bind(&self) -> Result<(), DriverError> {
Review Comment:
Not a regression — just flagging unchanged territory: the shim re-binds the owning device per call here, but the kernel launches in the `.cu` host wrappers still use the calling thread's current device, which is the same implicit single-device assumption the existing CUDA path makes (tests run single-device). Fine for now; multi-GPU on either vendor would need a broader audit.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397540596
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,418 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+ // hipStreamNonBlocking: the new stream does not implicitly synchronize with
+ // the NULL/default stream, matching cudarc's fork_default_stream.
+ const HIP_STREAM_NON_BLOCKING: u32 = 1;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32) -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreateWithFlags(stream: *mut *mut c_void, flags: u32) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(
+ range.start <= range.end && range.end <= self.len,
+ "slice_mut out of bounds"
+ );
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ // hipFree releases on the calling thread's current device, so
+ // re-bind the owning device first (cudarc does the same in Drop):
+ // on multi-GPU a different device may be current, which would
+ // otherwise free against the wrong device. Best-effort -- Drop
+ // cannot report an error, so a failed bind is swallowed.
+ let _ = self._device.bind();
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
+
+ impl Drop for CudaStream {
+ fn drop(&mut self) {
+ if !self.stream.is_null() {
+ unsafe {
+ let _ = hipStreamDestroy(self.stream);
+ }
+ }
+ }
+ }
+
+ /// HIP device handle. Stand-in for `cudarc::CudaDevice`; created via
+ /// `CudaDevice::new(ordinal)` and shared as `Arc<CudaDevice>` exactly like
+ /// cudarc (whose `new` already returns the `Arc`).
+ pub struct CudaDevice {
+ ordinal: usize,
+ }
+
+ impl CudaDevice {
+ /// Select device `ordinal` and return a shared handle, or an error if no
+ /// such device exists. Matches `cudarc::CudaDevice::new`.
+ pub fn new(ordinal: usize) -> Result<Arc<Self>, DriverError> {
+ unsafe {
+ let mut count: i32 = 0;
+ check(hipGetDeviceCount(&mut count))?;
+ if ordinal as i32 >= count {
+ return Err(DriverError(101)); // hipErrorInvalidDevice
+ }
+ check(hipSetDevice(ordinal as i32))?;
+ }
+ Ok(Arc::new(Self { ordinal }))
+ }
+
+ pub fn ordinal(&self) -> usize {
+ self.ordinal
+ }
+
+ fn bind(&self) -> Result<(), DriverError> {
Review Comment:
Not a regression — just flagging unchanged territory: the shim re-binds the owning device per call here, but the kernel launches in the `.cu` host wrappers still use the calling thread's current device, which is the same implicit single-device assumption the existing CUDA path makes (tests run single-device). Fine for now; multi-GPU on either vendor would need a broader audit.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397736468
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,418 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
Review Comment:
Done in 088041c -- per ASF policy, dropped the per-file AMD copyright and `Author:` lines from all 10 files (the ASF Apache-2.0 header is untouched) and added the AMD attribution to `NOTICE`. Git history preserves authorship.
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +50,247 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::{CUDA_MEMORY_TYPE_DEVICE, CUDA_MEMORY_TYPE_MANAGED, CudaPointerAttributes};
+ use std::ffi::c_void;
+
+ // hipMemoryType enum values are NOT guaranteed equal to CUDA's across ROCm
+ // releases (older HIP used Host=0/Device=1; the hip_runtime_api.h note flags
+ // this explicitly). So we read the real hipPointerAttribute_t and compare its
+ // `type` against the named hipMemoryType* constants rather than a magic
+ // number, then translate to the CUDA convention the caller expects.
+ const HIP_MEMORY_TYPE_DEVICE: i32 = 2; // hipMemoryTypeDevice
Review Comment:
Documented in dbee1e1 -- added a ROCm >= 6.0 prerequisite to DEVELOPMENT.md noting the `hipMemoryType` device-pointer check uses the ROCm 6+ convention (ROCm 5.x would reject valid device pointers). Tested on 7.2.1.
##########
qdp/qdp-kernels/src/lib.rs:
##########
@@ -20,6 +20,9 @@
use std::ffi::c_void;
+pub mod device;
+use device::{DeviceRepr, ValidAsZeroBits};
Review Comment:
Added in dbee1e1 -- a `#[cfg(not(any(feature = "cuda", feature = "hip")))] compile_error!(...)` in device.rs, so that config now gives a clean diagnostic instead of the raw E0432.
##########
qdp/DEVELOPMENT.md:
##########
@@ -114,6 +114,43 @@ cd ..
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)
Review 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.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682218305
I'm also curious about the more complex optimizations can or cannot be converted this easily. I think recently someone open some PRs that had some real profound ones.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "github-actions[bot] (via GitHub)" <gi...@apache.org>.
github-actions[bot] closed pull request #1399: QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine
URL: https://github.com/apache/mahout/pull/1399
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3448788190
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +47,247 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::{CUDA_MEMORY_TYPE_DEVICE, CUDA_MEMORY_TYPE_MANAGED, CudaPointerAttributes};
+ use std::ffi::c_void;
+
+ // hipMemoryType enum values are NOT guaranteed equal to CUDA's across ROCm
+ // releases (older HIP used Host=0/Device=1; the hip_runtime_api.h note flags
+ // this explicitly). So we read the real hipPointerAttribute_t and compare its
+ // `type` against the named hipMemoryType* constants rather than a magic
+ // number, then translate to the CUDA convention the caller expects.
+ const HIP_MEMORY_TYPE_DEVICE: i32 = 2; // hipMemoryTypeDevice
+ const HIP_MEMORY_TYPE_MANAGED: i32 = 3; // hipMemoryTypeManaged
+
+ // Mirror of hipPointerAttribute_t (ROCm hip_runtime_api.h): the leading
+ // `type` field is the hipMemoryType enum read by cudaPointerGetAttributes.
+ #[repr(C)]
+ struct HipPointerAttributes {
Review Comment:
+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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4683015613
@ryankert01 Thanks for running the full suite on the 3090 Ti -- good to have the default CUDA path vendor-confirmed.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5154768507
/unstale
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "400Ping (via GitHub)" <gi...@apache.org>.
400Ping commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5707151619
Will take a look at this PR when I have time. I am a little busy these days with limited bandwidth, will try my best to find time to review this.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "rich7420 (via GitHub)" <gi...@apache.org>.
rich7420 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3448782155
##########
qdp/qdp-core/src/gpu/validation.rs:
##########
@@ -23,7 +23,7 @@
#![allow(unused_unsafe)]
use crate::error::{MahoutError, Result, cuda_error_to_string};
-use cudarc::driver::{CudaDevice, CudaSlice, DevicePtrMut};
+use crate::gpu_rt::{CudaDevice, CudaSlice, DevicePtrMut};
Review Comment:
heads up — this `gpu_rt` swap is what pulls the HIP backend into this file, and I think it quietly exposes a stream race in the four validators below (`check_finite_batch_f32`/`_f64`, `validate_and_cast_basis_indices_f32`, `assert_basis_indices_in_range_usize`).
Each one launches its kernel on the caller's `stream` and then reads the flag straight back with `device.dtoh_sync_copy(&flag)` (around lines 74 / 125 / 210 / 262) without syncing `stream` first. On CUDA the legacy default stream hides it, but on HIP a non-blocking stream (the PyTorch/DLPack path, or a forked stream) isn't ordered against the null-stream copy, so the host can read the still-zero flag and validation silently passes. The basis ones are the scary case — an out-of-range index sails through and the basis kernel then does an OOB device write.
Funny enough the doc comment on line 165 already says it "Synchronizes on `stream`..." but the code doesn't. You already added exactly this fix over in `amplitude.rs` (the `sync_cuda_stream(stream, ...)` right before `dtoh_sync_copy`) — looks like these four just got missed. Mind dropping the same sync in before each readback here? 🙏
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +47,247 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::{CUDA_MEMORY_TYPE_DEVICE, CUDA_MEMORY_TYPE_MANAGED, CudaPointerAttributes};
+ use std::ffi::c_void;
+
+ // hipMemoryType enum values are NOT guaranteed equal to CUDA's across ROCm
+ // releases (older HIP used Host=0/Device=1; the hip_runtime_api.h note flags
+ // this explicitly). So we read the real hipPointerAttribute_t and compare its
+ // `type` against the named hipMemoryType* constants rather than a magic
+ // number, then translate to the CUDA convention the caller expects.
+ const HIP_MEMORY_TYPE_DEVICE: i32 = 2; // hipMemoryTypeDevice
+ const HIP_MEMORY_TYPE_MANAGED: i32 = 3; // hipMemoryTypeManaged
+
+ // Mirror of hipPointerAttribute_t (ROCm hip_runtime_api.h): the leading
+ // `type` field is the hipMemoryType enum read by cudaPointerGetAttributes.
+ #[repr(C)]
+ struct HipPointerAttributes {
Review Comment:
one to double-check on real hardware: this `HipPointerAttributes` is a hand-rolled mirror of ROCm's `hipPointerAttribute_t`, and that struct's field order/names actually shifted between ROCm 5.x and 6.x. If the layout doesn't line up with the `hip_runtime_api.h` on the target ROCm, `hipPointerGetAttributes` fills it at the wrong offsets and `validate_cuda_input_ptr` ends up reading garbage for `memory_type`/`device` — which would either reject every valid device ptr or (worse) wave a host ptr through to a kernel.
Couldn't verify it here without the ROCm headers, so might be worth pinning the supported ROCm floor and sanity-checking this layout against it. 👀
##########
qdp/qdp-kernels/hip_compat/cuComplex.h:
##########
@@ -0,0 +1,46 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+// HIP forwarding shim for <cuComplex.h> (HIP build path only; see
+// cuda_runtime.h in this directory for how it is selected). hipcc does not
+// ship a <cuComplex.h>; HIP's <hip/hip_complex.h> provides the same complex
+// layout and helpers under hip* names. The aliases below let the .cu sources
+// keep their cuComplex / cuDoubleComplex / make_cu* / cuC* spellings unchanged.
+
+#pragma once
+#include <hip/hip_complex.h>
+
+typedef hipDoubleComplex cuDoubleComplex;
+typedef hipFloatComplex cuComplex;
+
+#define make_cuDoubleComplex make_hipDoubleComplex
+#define make_cuComplex make_hipFloatComplex
Review Comment:
small thing — these object-like macros (`make_cuComplex`, `cuCreal`, etc.) get defined right after `#include <hip/hip_complex.h>`, and newer ROCm's `hip_complex.h` already ships CUDA-compat aliases with these exact names (as inline functions). When that happens the macro rewrites the function's own declaration token and hipcc errors out on a redefinition.
It'd be a hard build break rather than anything silent, and only on the ROCm versions that ship those aliases — but might want to `#undef` first or wrap each in `#ifndef` to be safe.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,443 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(not(any(feature = "cuda", feature = "hip")))]
+compile_error!("qdp-kernels requires exactly one of the `cuda` or `hip` features");
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+ // hipStreamNonBlocking: the new stream does not implicitly synchronize with
+ // the NULL/default stream, matching cudarc's fork_default_stream.
+ const HIP_STREAM_NON_BLOCKING: u32 = 1;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32) -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreateWithFlags(stream: *mut *mut c_void, flags: u32) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Synchronize the NULL/default stream so its prior work is ordered before any
+ /// other stream observes the affected buffers.
+ ///
+ /// The blocking shim copies (htod/alloc_zeros) issue hipMemcpy/hipMemset on the
+ /// default (NULL) stream. CUDA's legacy default stream is synchronizing, so on
+ /// NVIDIA that work is implicitly ordered before a kernel launched on a forked
+ /// non-blocking stream that reads the same buffer. HIP's default stream is NOT
+ /// synchronizing relative to a hipStreamNonBlocking stream, so without this an
+ /// encoder that sets up input/output on the default stream and then launches
+ /// the norm/encode kernels on the caller's forked stream would race the setup
+ /// (the kernel reads stale/zero data). A default-stream synchronize after the
+ /// blocking copy restores the CUDA-equivalent ordering while preserving the
+ /// dual-stream copy/compute overlap (which uses async copies on explicit
+ /// streams, not these blocking paths).
+ fn sync_default_stream() -> Result<(), DriverError> {
+ unsafe { check(hipStreamSynchronize(std::ptr::null_mut())) }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(
+ range.start <= range.end && range.end <= self.len,
+ "slice_mut out of bounds"
+ );
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
Review Comment:
nit / latent: `range.start * size_of::<T>()` is done in `usize` before the cast to `u64`, so a big enough `range.start` could wrap before the offset is even applied, and the `range.end <= self.len` assert above wouldn't catch it. Not reachable at today's buffer sizes, but the encoding code uses `checked_mul` for this kind of thing elsewhere — might as well match it here.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,443 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(not(any(feature = "cuda", feature = "hip")))]
+compile_error!("qdp-kernels requires exactly one of the `cuda` or `hip` features");
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+ // hipStreamNonBlocking: the new stream does not implicitly synchronize with
+ // the NULL/default stream, matching cudarc's fork_default_stream.
+ const HIP_STREAM_NON_BLOCKING: u32 = 1;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32) -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreateWithFlags(stream: *mut *mut c_void, flags: u32) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Synchronize the NULL/default stream so its prior work is ordered before any
+ /// other stream observes the affected buffers.
+ ///
+ /// The blocking shim copies (htod/alloc_zeros) issue hipMemcpy/hipMemset on the
+ /// default (NULL) stream. CUDA's legacy default stream is synchronizing, so on
+ /// NVIDIA that work is implicitly ordered before a kernel launched on a forked
+ /// non-blocking stream that reads the same buffer. HIP's default stream is NOT
+ /// synchronizing relative to a hipStreamNonBlocking stream, so without this an
+ /// encoder that sets up input/output on the default stream and then launches
+ /// the norm/encode kernels on the caller's forked stream would race the setup
+ /// (the kernel reads stale/zero data). A default-stream synchronize after the
+ /// blocking copy restores the CUDA-equivalent ordering while preserving the
+ /// dual-stream copy/compute overlap (which uses async copies on explicit
+ /// streams, not these blocking paths).
+ fn sync_default_stream() -> Result<(), DriverError> {
+ unsafe { check(hipStreamSynchronize(std::ptr::null_mut())) }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(
+ range.start <= range.end && range.end <= self.len,
+ "slice_mut out of bounds"
+ );
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ // hipFree releases on the calling thread's current device, so
+ // re-bind the owning device first (cudarc does the same in Drop):
+ // on multi-GPU a different device may be current, which would
+ // otherwise free against the wrong device. Best-effort -- Drop
+ // cannot report an error, so a failed bind is swallowed.
+ let _ = self._device.bind();
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
Review Comment:
fwiw cudarc's own `CudaStream` is `Send` but not `Sync`, and this shim adds `Sync` on top. Nothing shares a stream by `&` across threads today so it's latent, but it does widen the contract the CUDA side leans on — if some future code ever hands one `CudaStream` to two threads, `hipStream_t` isn't safe to enqueue onto concurrently and the compiler won't stop it anymore. Could just drop the `Sync` impl to keep it matching cudarc.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "github-actions[bot] (via GitHub)" <gi...@apache.org>.
github-actions[bot] commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-6244351004
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](https://lists.apache.org/list.html?dev@mahout.apache.org) or in [GitHub Discussions](https://github.com/apache/mahout/discussions).
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682322151
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.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682482362
@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."
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "rich7420 (via GitHub)" <gi...@apache.org>.
rich7420 commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4678743699
btw, please check the precommit errors
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "github-actions[bot] (via GitHub)" <gi...@apache.org>.
github-actions[bot] commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5707129888
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](https://lists.apache.org/list.html?dev@mahout.apache.org) or in [GitHub Discussions](https://github.com/apache/mahout/discussions).
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "jeffdaily (via GitHub)" <gi...@apache.org>.
jeffdaily commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682482171
@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.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "rich7420 (via GitHub)" <gi...@apache.org>.
rich7420 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3394439511
##########
qdp/qdp-core/src/gpu/cuda_ffi.rs:
##########
@@ -37,65 +50,201 @@ pub(crate) struct CudaPointerAttributes {
pub allocation_flags: u32,
}
-// CUDA error codes
+// CUDA/HIP error codes (numerically identical for the codes used).
pub(crate) const CUDA_SUCCESS: i32 = 0;
-// Note: CUDA_ERROR_NOT_READY may be used in future optimizations for non-blocking event checks
-// Reference: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g3f51e3575c2178246db0a94a430e0028
#[allow(dead_code)]
pub(crate) const CUDA_ERROR_NOT_READY: i32 = 34;
-unsafe extern "C" {
- pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
- pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+// ---- CUDA backend: bind libcudart directly ----
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub(crate) use cuda_rt::*;
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+mod cuda_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ unsafe extern "C" {
+ pub(crate) fn cudaHostAlloc(pHost: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ pub(crate) fn cudaFreeHost(ptr: *mut c_void) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaPointerGetAttributes(
+ attributes: *mut CudaPointerAttributes,
+ ptr: *const c_void,
+ ) -> i32;
+
+ pub(crate) fn cudaMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+
+ pub(crate) fn cudaMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaMemcpy(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ ) -> i32;
+
+ pub(crate) fn cudaEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ pub(crate) fn cudaEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ pub(crate) fn cudaEventDestroy(event: *mut c_void) -> i32;
+ pub(crate) fn cudaStreamWaitEvent(
+ stream: *mut c_void,
+ event: *mut c_void,
+ flags: u32,
+ ) -> i32;
+ pub(crate) fn cudaStreamSynchronize(stream: *mut c_void) -> i32;
+
+ pub(crate) fn cudaMemsetAsync(
+ devPtr: *mut c_void,
+ value: i32,
+ count: usize,
+ stream: *mut c_void,
+ ) -> i32;
+
+ #[allow(dead_code)]
+ pub(crate) fn cudaEventQuery(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventSynchronize(event: *mut c_void) -> i32;
+ pub(crate) fn cudaEventElapsedTime(
+ ms: *mut f32,
+ start: *mut c_void,
+ end: *mut c_void,
+ ) -> i32;
+ }
+}
+
+// ---- HIP backend: bind libamdhip64, expose the same cuda* names ----
+#[cfg(feature = "hip")]
+pub(crate) use hip_rt::*;
+
+// The wrapper functions deliberately keep the cuda* spelling so call sites are
+// vendor-agnostic; suppress the snake_case lint for that intentional naming.
+#[cfg(feature = "hip")]
+#[allow(non_snake_case)]
+mod hip_rt {
+ use super::CudaPointerAttributes;
+ use std::ffi::c_void;
+
+ // hipPointerAttribute_t has the same first fields we read (memoryType,
+ // device, devicePointer, hostPointer); a #[repr(C)] alias suffices since we
+ // only read memory_type/device.
+ unsafe extern "C" {
+ fn hipHostMalloc(ptr: *mut *mut c_void, size: usize, flags: u32) -> i32;
+ fn hipHostFree(ptr: *mut c_void) -> i32;
+ fn hipPointerGetAttributes(attributes: *mut c_void, ptr: *const c_void) -> i32;
+ fn hipMemGetInfo(free: *mut usize, total: *mut usize) -> i32;
+ fn hipMemcpyAsync(
+ dst: *mut c_void,
+ src: *const c_void,
+ count: usize,
+ kind: u32,
+ stream: *mut c_void,
+ ) -> i32;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, count: usize, kind: u32) -> i32;
+ fn hipEventCreateWithFlags(event: *mut *mut c_void, flags: u32) -> i32;
+ fn hipEventRecord(event: *mut c_void, stream: *mut c_void) -> i32;
+ fn hipEventDestroy(event: *mut c_void) -> i32;
+ fn hipStreamWaitEvent(stream: *mut c_void, event: *mut c_void, flags: u32) -> i32;
+ fn hipStreamSynchronize(stream: *mut c_void) -> i32;
+ fn hipMemsetAsync(dst: *mut c_void, value: i32, count: usize, stream: *mut c_void) -> i32;
+ fn hipEventQuery(event: *mut c_void) -> i32;
+ fn hipEventSynchronize(event: *mut c_void) -> i32;
+ fn hipEventElapsedTime(ms: *mut f32, start: *mut c_void, end: *mut c_void) -> i32;
+ }
+
+ // hipHostMallocDefault == 0, matching cudaHostAllocDefault used by callers.
+ pub(crate) unsafe fn cudaHostAlloc(p: *mut *mut c_void, size: usize, flags: u32) -> i32 {
+ unsafe { hipHostMalloc(p, size, flags) }
+ }
+ pub(crate) unsafe fn cudaFreeHost(ptr: *mut c_void) -> i32 {
+ unsafe { hipHostFree(ptr) }
+ }
#[allow(dead_code)]
- pub(crate) fn cudaPointerGetAttributes(
+ pub(crate) unsafe fn cudaPointerGetAttributes(
Review Comment:
This reinterpret-cast onto `hipPointerAttribute_t` is live (lib.rs:82 reads `memory_type`/`device`), and it assumes `hipMemoryType` enum values equal CUDA's (`CUDA_MEMORY_TYPE_DEVICE = 2`). `hipMemoryType*` has differed from `cudaMemoryType*` across ROCm versions, so a valid device pointer could be misclassified on some ROCm. Tests pass on 7.2.1, but consider reading `hipPointerGetAttributes` into the real `hipPointerAttribute_t` and mapping the enum explicitly, rather than relying on layout+value coincidence.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
+ }
+ }
+ }
+ }
+
+ /// A HIP stream. The public `stream` field mirrors cudarc's
+ /// `CudaStream { stream: sys::CUstream, .. }` so existing call sites that do
+ /// `ctx.stream_compute.stream as *mut c_void` keep working.
+ pub struct CudaStream {
+ pub stream: *mut c_void,
+ _device: Arc<CudaDevice>,
+ }
+
+ unsafe impl Send for CudaStream {}
+ unsafe impl Sync for CudaStream {}
+
+ impl Drop for CudaStream {
+ fn drop(&mut self) {
+ if !self.stream.is_null() {
+ unsafe {
+ let _ = hipStreamDestroy(self.stream);
+ }
+ }
+ }
+ }
+
+ /// HIP device handle. Stand-in for `cudarc::CudaDevice`; created via
+ /// `CudaDevice::new(ordinal)` and shared as `Arc<CudaDevice>` exactly like
+ /// cudarc (whose `new` already returns the `Arc`).
+ pub struct CudaDevice {
+ ordinal: usize,
+ }
+
+ impl CudaDevice {
+ /// Select device `ordinal` and return a shared handle, or an error if no
+ /// such device exists. Matches `cudarc::CudaDevice::new`.
+ pub fn new(ordinal: usize) -> Result<Arc<Self>, DriverError> {
+ unsafe {
+ let mut count: i32 = 0;
+ check(hipGetDeviceCount(&mut count))?;
+ if ordinal as i32 >= count {
+ return Err(DriverError(101)); // hipErrorInvalidDevice
+ }
+ check(hipSetDevice(ordinal as i32))?;
+ }
+ Ok(Arc::new(Self { ordinal }))
+ }
+
+ pub fn ordinal(&self) -> usize {
+ self.ordinal
+ }
+
+ fn bind(&self) -> Result<(), DriverError> {
+ unsafe { check(hipSetDevice(self.ordinal as i32)) }
+ }
+
+ /// Allocate `len` uninitialized elements of `T` on the device.
+ ///
+ /// # Safety
+ /// Contents are uninitialized until written, mirroring `cudarc`'s
+ /// `unsafe fn alloc`.
+ pub unsafe fn alloc<T>(self: &Arc<Self>, len: usize) -> Result<CudaSlice<T>, DriverError> {
+ self.bind()?;
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ let mut ptr: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipMalloc(&mut ptr, bytes.max(1)))?;
+ }
+ Ok(CudaSlice {
+ ptr: ptr as u64,
+ len,
+ _device: Arc::clone(self),
+ _marker: PhantomData,
+ })
+ }
+
+ /// Allocate `len` zero-initialized elements of `T` on the device.
+ pub fn alloc_zeros<T: ValidAsZeroBits>(
+ self: &Arc<Self>,
+ len: usize,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let slice = unsafe { self.alloc::<T>(len)? };
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemset(slice.raw_ptr(), 0, bytes))?;
+ }
+ }
+ Ok(slice)
+ }
+
+ /// Copy a host slice to a freshly allocated device buffer (blocking).
+ pub fn htod_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &[T],
+ ) -> Result<CudaSlice<T>, DriverError> {
+ let mut slice = unsafe { self.alloc::<T>(src.len())? };
+ self.htod_sync_copy_into(src, &mut slice)?;
+ Ok(slice)
+ }
+
+ /// Copy an owned host Vec to a freshly allocated device buffer. cudarc's
+ /// `htod_copy` keeps the Vec alive until an async copy completes; our
+ /// copy is synchronous (blocking hipMemcpy), so the Vec can be dropped
+ /// on return with identical observable behavior.
+ pub fn htod_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: Vec<T>,
+ ) -> Result<CudaSlice<T>, DriverError> {
+ self.htod_sync_copy(&src)
+ }
+
+ /// Copy a host slice into an existing device buffer or sub-view
+ /// (blocking). Accepts any `DevicePtrMut` target so both `CudaSlice` and
+ /// the `slice_mut` view work, matching cudarc's generic destination.
+ pub fn htod_sync_copy_into<T: DeviceRepr, D: DevicePtrMut<T> + DeviceSlice<T>>(
+ self: &Arc<Self>,
+ src: &[T],
+ dst: &mut D,
+ ) -> Result<(), DriverError> {
+ assert_eq!(
+ dst.len(),
+ src.len(),
+ "htod_sync_copy_into: dst.len() != src.len()"
+ );
+ self.bind()?;
+ let bytes = std::mem::size_of_val(src);
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ (*dst.device_ptr_mut()) as *mut c_void,
+ src.as_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_HOST_TO_DEVICE,
+ ))?;
+ }
+ }
+ Ok(())
+ }
+
+ /// Copy a device buffer back to a freshly allocated host Vec (blocking).
+ ///
+ /// Matches cudarc's bound of just `DeviceRepr` (no `Default`): the Vec is
+ /// allocated uninitialized and every byte is written by the copy before
+ /// its length is set, which is sound because `DeviceRepr` is plain data.
+ pub fn dtoh_sync_copy<T: DeviceRepr>(
+ self: &Arc<Self>,
+ src: &CudaSlice<T>,
+ ) -> Result<Vec<T>, DriverError> {
+ self.bind()?;
+ let len = src.len;
+ let mut out: Vec<T> = Vec::with_capacity(len);
+ let bytes = len.saturating_mul(std::mem::size_of::<T>());
+ if bytes > 0 {
+ unsafe {
+ check(hipMemcpy(
+ out.as_mut_ptr() as *mut c_void,
+ src.raw_ptr() as *const c_void,
+ bytes,
+ HIP_MEMCPY_DEVICE_TO_HOST,
+ ))?;
+ }
+ }
+ unsafe {
+ out.set_len(len);
+ }
+ Ok(out)
+ }
+
+ /// Block until all work on the device's default stream completes.
+ pub fn synchronize(&self) -> Result<(), DriverError> {
+ self.bind()?;
+ unsafe { check(hipDeviceSynchronize()) }
+ }
+
+ /// Create a new stream tied to this device. Mirrors
+ /// `cudarc::CudaDevice::fork_default_stream`.
+ pub fn fork_default_stream(self: &Arc<Self>) -> Result<CudaStream, DriverError> {
+ self.bind()?;
+ let mut stream: *mut c_void = std::ptr::null_mut();
+ unsafe {
+ check(hipStreamCreate(&mut stream))?;
Review Comment:
`hipStreamCreate` is the blocking-stream constructor (implicitly synchronizes with the default/NULL stream); cudarc's `fork_default_stream` creates a *non-blocking* stream. If the dual-stream H2D/compute overlap matters here, this may quietly serialize it. `hipStreamCreateWithFlags(&mut stream, hipStreamNonBlocking)` matches cudarc. (Your overlap-tracking tests reportedly pass, so flagging for confirmation rather than asserting a regression.)
##########
qdp/qdp-kernels/build.rs:
##########
@@ -154,9 +171,76 @@ fn apply_default_arch_targets(build: &mut cc::Build) {
}
}
+fn hip_requested() -> bool {
+ if cfg!(feature = "hip") {
+ return true;
+ }
+ env::var("QDP_USE_HIP")
Review Comment:
`QDP_USE_HIP` flips the *kernel* build to hipcc even when the cargo `hip` feature is off. A plain `cargo build` (default `cuda`) with this env set would compile AMD kernels against a cudarc Rust host → mismatch. Since the host side keys off the cargo feature only, consider erroring when `QDP_USE_HIP` is truthy but `CARGO_FEATURE_HIP` is unset (or vice versa) so it fails loud instead of producing a desynced build.
##########
qdp/qdp-kernels/src/device.rs:
##########
@@ -0,0 +1,401 @@
+//
+// Licensed to the Apache Software Foundation (ASF) under one or more
+// contributor license agreements. See the NOTICE file distributed with
+// this work for additional information regarding copyright ownership.
+// The ASF licenses this file to You under the Apache License, Version 2.0
+// (the "License"); you may not use this file except in compliance with
+// the License. You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// Author: Jeff Daily <je...@amd.com>
+
+//! Device runtime surface, vendor-selected at compile time.
+//!
+//! `cudarc` is CUDA-only with no ROCm backend, so the AMD build cannot depend
+//! on it. This module is the seam: on the default (`cuda`) feature it simply
+//! re-exports the slice of `cudarc::driver` the crates use; on the `hip`
+//! feature it provides a thin HIP-runtime shim with the SAME type names and
+//! method signatures, so every call site (`device.alloc`, `htod_sync_copy`,
+//! `slice.device_ptr()`, ...) compiles unchanged on both vendors.
+//!
+//! The marker traits `DeviceRepr` / `ValidAsZeroBits` live here (not in
+//! qdp-core) because qdp-kernels implements them on its complex structs and is
+//! the lowest crate in the workspace.
+
+#[cfg(all(feature = "cuda", not(feature = "hip")))]
+pub use cudarc::driver::{
+ CudaDevice, CudaSlice, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice, ValidAsZeroBits,
+ safe::CudaStream,
+};
+
+#[cfg(feature = "hip")]
+pub use hip::{
+ CudaDevice, CudaSlice, CudaStream, DevicePtr, DevicePtrMut, DeviceRepr, DeviceSlice,
+ DriverError, ValidAsZeroBits,
+};
+
+#[cfg(feature = "hip")]
+mod hip {
+ use std::ffi::c_void;
+ use std::marker::PhantomData;
+ use std::sync::Arc;
+
+ // ---- HIP runtime FFI (subset used by the device abstraction) ----
+ // hip* names map 1:1 to the cuda* runtime entry points cudarc wraps; HIP
+ // error codes match CUDA's numerically for the codes we surface.
+ #[allow(non_camel_case_types)]
+ type hipError_t = i32;
+
+ const HIP_SUCCESS: hipError_t = 0;
+ const HIP_MEMCPY_HOST_TO_DEVICE: u32 = 1;
+ const HIP_MEMCPY_DEVICE_TO_HOST: u32 = 2;
+
+ unsafe extern "C" {
+ fn hipSetDevice(device: i32) -> hipError_t;
+ fn hipGetDeviceCount(count: *mut i32) -> hipError_t;
+ fn hipMalloc(ptr: *mut *mut c_void, size: usize) -> hipError_t;
+ fn hipFree(ptr: *mut c_void) -> hipError_t;
+ fn hipMemset(ptr: *mut c_void, value: i32, size: usize) -> hipError_t;
+ fn hipMemcpy(dst: *mut c_void, src: *const c_void, size: usize, kind: u32)
+ -> hipError_t;
+ fn hipDeviceSynchronize() -> hipError_t;
+ fn hipStreamCreate(stream: *mut *mut c_void) -> hipError_t;
+ fn hipStreamDestroy(stream: *mut c_void) -> hipError_t;
+ fn hipStreamSynchronize(stream: *mut c_void) -> hipError_t;
+ }
+
+ /// Mirrors the role of `cudarc::driver::DriverError`: an opaque, `Debug`able
+ /// wrapper over a runtime status code. Call sites only ever `{:?}`-format it.
+ #[derive(Debug, Clone, Copy, PartialEq, Eq)]
+ pub struct DriverError(pub hipError_t);
+
+ fn check(code: hipError_t) -> Result<(), DriverError> {
+ if code == HIP_SUCCESS {
+ Ok(())
+ } else {
+ Err(DriverError(code))
+ }
+ }
+
+ /// Marker: type is safe to byte-copy to/from the device. Mirrors
+ /// `cudarc::driver::DeviceRepr`.
+ ///
+ /// # Safety
+ /// Implementor must be `#[repr(C)]`/`#[repr(transparent)]` plain-old-data
+ /// with no padding that would expose uninitialized bytes.
+ pub unsafe trait DeviceRepr: Copy {}
+ unsafe impl DeviceRepr for f32 {}
+ unsafe impl DeviceRepr for f64 {}
+ unsafe impl DeviceRepr for i32 {}
+ unsafe impl DeviceRepr for u32 {}
+ unsafe impl DeviceRepr for usize {}
+
+ /// Marker: an all-zero bit pattern is a valid value (enables alloc_zeros).
+ /// Mirrors `cudarc::driver::ValidAsZeroBits`.
+ ///
+ /// # Safety
+ /// All-zero bytes must be a valid inhabitant of the type.
+ pub unsafe trait ValidAsZeroBits {}
+ unsafe impl ValidAsZeroBits for f32 {}
+ unsafe impl ValidAsZeroBits for f64 {}
+ unsafe impl ValidAsZeroBits for i32 {}
+ unsafe impl ValidAsZeroBits for u32 {}
+ unsafe impl ValidAsZeroBits for usize {}
+
+ /// Raw device-pointer accessors, matching cudarc's traits. The returned
+ /// reference is to the device address stored as `u64`, so the existing
+ /// `*slice.device_ptr() as *mut T` call sites work verbatim.
+ pub trait DevicePtr<T> {
+ fn device_ptr(&self) -> &u64;
+ }
+ pub trait DevicePtrMut<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64;
+ }
+ /// Length accessor, matching cudarc's `DeviceSlice`.
+ pub trait DeviceSlice<T> {
+ fn len(&self) -> usize;
+ fn is_empty(&self) -> bool {
+ self.len() == 0
+ }
+ }
+
+ /// Owned device allocation; frees on drop. Stand-in for `cudarc::CudaSlice`.
+ pub struct CudaSlice<T> {
+ ptr: u64,
+ len: usize,
+ _device: Arc<CudaDevice>,
+ _marker: PhantomData<T>,
+ }
+
+ // The device address is just an integer; ownership/lifetime is enforced by
+ // the held Arc<CudaDevice>. Safe to move across threads like cudarc's slice.
+ unsafe impl<T: Send> Send for CudaSlice<T> {}
+ unsafe impl<T: Sync> Sync for CudaSlice<T> {}
+
+ impl<T> CudaSlice<T> {
+ fn raw_ptr(&self) -> *mut c_void {
+ self.ptr as *mut c_void
+ }
+
+ /// Mutable sub-view `[range.start, range.end)`. Mirrors
+ /// `cudarc::CudaSlice::slice_mut`; the returned view borrows this slice
+ /// and is itself a `DevicePtrMut`/`DeviceSlice` copy target.
+ pub fn slice_mut(&mut self, range: std::ops::Range<usize>) -> CudaViewMut<'_, T> {
+ assert!(range.start <= range.end && range.end <= self.len, "slice_mut out of bounds");
+ let offset_ptr = self.ptr + (range.start * std::mem::size_of::<T>()) as u64;
+ CudaViewMut {
+ ptr: offset_ptr,
+ len: range.end - range.start,
+ _parent: PhantomData,
+ }
+ }
+ }
+
+ /// Borrowed mutable view into a `CudaSlice`, returned by `slice_mut`.
+ pub struct CudaViewMut<'a, T> {
+ ptr: u64,
+ len: usize,
+ _parent: PhantomData<&'a mut T>,
+ }
+
+ impl<T> DevicePtr<T> for CudaViewMut<'_, T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaViewMut<'_, T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaViewMut<'_, T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> DevicePtr<T> for CudaSlice<T> {
+ fn device_ptr(&self) -> &u64 {
+ &self.ptr
+ }
+ }
+ impl<T> DevicePtrMut<T> for CudaSlice<T> {
+ fn device_ptr_mut(&mut self) -> &mut u64 {
+ &mut self.ptr
+ }
+ }
+ impl<T> DeviceSlice<T> for CudaSlice<T> {
+ fn len(&self) -> usize {
+ self.len
+ }
+ }
+
+ impl<T> Drop for CudaSlice<T> {
+ fn drop(&mut self) {
+ if self.ptr != 0 {
+ unsafe {
+ let _ = hipFree(self.raw_ptr());
Review Comment:
`hipFree` runs against the calling thread's *current* device, but Drop never re-binds `self._device` (compare `alloc` above, which does `self.bind()?`). On multi-GPU, if another device is current when this drops, you free against the wrong device — `hipErrorInvalidDevicePointer` (swallowed by `let _ =` → leak) or a wrong-pointer free. cudarc binds the device in its slice Drop. A `self._device.bind()` (best-effort, ignore error) before `hipFree` restores parity.
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "ryankert01 (via GitHub)" <gi...@apache.org>.
ryankert01 commented on code in PR #1399:
URL: https://github.com/apache/mahout/pull/1399#discussion_r3397678116
##########
qdp/DEVELOPMENT.md:
##########
@@ -114,6 +114,43 @@ cd ..
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)
Review Comment:
I think we should add a build github action workflow for hip!
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "github-actions[bot] (via GitHub)" <gi...@apache.org>.
github-actions[bot] commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-5154619605
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](https://lists.apache.org/list.html?dev@mahout.apache.org) or in [GitHub Discussions](https://github.com/apache/mahout/discussions).
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] QDP: add an AMD GPU (ROCm/HIP) build for the native encoder engine [mahout]
Posted by "github-actions[bot] (via GitHub)" <gi...@apache.org>.
github-actions[bot] commented on PR #1399:
URL: https://github.com/apache/mahout/pull/1399#issuecomment-6392068878
This pull request has been automatically closed because there has been no more activity in the 7 days since being marked stale.
Please feel free to reopen or open a new pull request if you'd still like this to be addressed. You can always ask for help on the [Mahout dev mailing list](https://lists.apache.org/list.html?dev@mahout.apache.org) or in [GitHub Discussions](https://github.com/apache/mahout/discussions).
Thanks again for your contribution!
--
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: commits-unsubscribe@mahout.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org