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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
