Add missing V-RAY host functions

8 files changed, 306 insertions, 51 deletions

diff --git a/zluda/src/cuda.rs b/zluda/src/cuda.rs
index e66ee75..bceb8bc 100644
--- a/zluda/src/cuda.rs
+++ b/zluda/src/cuda.rs
@@ -2184,11 +2184,11 @@ pub struct CUgraphExecUpdateResult_enum(pub ::std::os::raw::c_uint);
 pub use self::CUgraphExecUpdateResult_enum as CUgraphExecUpdateResult;
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuGetErrorString(
+pub unsafe extern "system" fn cuGetErrorString(
     CUresult(e): CUresult,
     pStr: *mut *const ::std::os::raw::c_char,
 ) -> CUresult {
-    unsafe { *pStr = hipGetErrorString(hipError_t(e)) };
+    *pStr = hipGetErrorString(hipError_t(e));
     CUresult::CUDA_SUCCESS
 }
 
@@ -2273,17 +2273,20 @@ pub extern "system" fn cuDeviceGetNvSciSyncAttributes(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuDeviceGetProperties(prop: *mut CUdevprop, dev: CUdevice) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuDeviceGetProperties(
+    prop: *mut CUdevprop,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::device::get_properties(prop, dev).encuda()
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuDeviceComputeCapability(
+pub unsafe extern "system" fn cuDeviceComputeCapability(
     major: *mut ::std::os::raw::c_int,
     minor: *mut ::std::os::raw::c_int,
     dev: CUdevice,
 ) -> CUresult {
-    r#impl::unimplemented()
+    hipDeviceComputeCapability(major, minor, dev.0).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2309,7 +2312,7 @@ pub extern "system" fn cuDevicePrimaryCtxSetFlags(
     dev: CUdevice,
     flags: ::std::os::raw::c_uint,
 ) -> CUresult {
-    cuDevicePrimaryCtxSetFlags_v2(dev, flags)
+    CUresult::CUDA_SUCCESS
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2317,7 +2320,7 @@ pub extern "system" fn cuDevicePrimaryCtxSetFlags_v2(
     dev: CUdevice,
     flags: ::std::os::raw::c_uint,
 ) -> CUresult {
-    r#impl::unimplemented()
+    cuDevicePrimaryCtxSetFlags(dev, flags)
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2391,12 +2394,12 @@ pub extern "system" fn cuCtxSynchronize() -> CUresult {
 
 #[cfg_attr(not(test), no_mangle)]
 pub extern "system" fn cuCtxSetLimit(limit: CUlimit, value: usize) -> CUresult {
-    r#impl::unimplemented()
+    r#impl::context::set_limit(limit, value)
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuCtxGetLimit(pvalue: *mut usize, limit: CUlimit) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuCtxGetLimit(pvalue: *mut usize, limit: CUlimit) -> CUresult {
+    r#impl::context::get_limit(pvalue, limit)
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2406,7 +2409,7 @@ pub extern "system" fn cuCtxGetCacheConfig(pconfig: *mut CUfunc_cache) -> CUresu
 
 #[cfg_attr(not(test), no_mangle)]
 pub extern "system" fn cuCtxSetCacheConfig(config: CUfunc_cache) -> CUresult {
-    r#impl::unimplemented()
+    CUresult::CUDA_SUCCESS
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2504,13 +2507,13 @@ pub extern "system" fn cuModuleGetFunction(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuModuleGetGlobal_v2(
+pub unsafe extern "system" fn cuModuleGetGlobal_v2(
     dptr: *mut CUdeviceptr,
     bytes: *mut usize,
     hmod: CUmodule,
     name: *const ::std::os::raw::c_char,
 ) -> CUresult {
-    r#impl::unimplemented()
+    hipModuleGetGlobal(dptr as _, bytes, hmod as _, name).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2532,17 +2535,17 @@ pub extern "system" fn cuModuleGetSurfRef(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuLinkCreate_v2(
+pub unsafe extern "system" fn cuLinkCreate_v2(
     numOptions: ::std::os::raw::c_uint,
     options: *mut CUjit_option,
     optionValues: *mut *mut ::std::os::raw::c_void,
     stateOut: *mut CUlinkState,
 ) -> CUresult {
-    r#impl::unimplemented()
+    r#impl::link::create(numOptions, options, optionValues, stateOut)
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuLinkAddData_v2(
+pub unsafe extern "system" fn cuLinkAddData_v2(
     state: CUlinkState,
     type_: CUjitInputType,
     data: *mut ::std::os::raw::c_void,
@@ -2552,7 +2555,16 @@ pub extern "system" fn cuLinkAddData_v2(
     options: *mut CUjit_option,
     optionValues: *mut *mut ::std::os::raw::c_void,
 ) -> CUresult {
-    r#impl::unimplemented()
+    r#impl::link::add_data(
+        state,
+        type_,
+        data,
+        size,
+        name,
+        numOptions,
+        options,
+        optionValues,
+    )
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2573,17 +2585,17 @@ pub extern "system" fn cuLinkComplete(
     cubinOut: *mut *mut ::std::os::raw::c_void,
     sizeOut: *mut usize,
 ) -> CUresult {
-    r#impl::unimplemented()
+    r#impl::link::complete(state, cubinOut, sizeOut)
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuLinkDestroy(state: CUlinkState) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuLinkDestroy(state: CUlinkState) -> CUresult {
+    r#impl::link::destroy(state)
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult {
+    hipMemGetInfo(free, total).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2641,17 +2653,17 @@ pub extern "system" fn cuMemAllocHost_v2(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult {
+    hipFreeHost(p).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuMemHostAlloc(
+pub unsafe extern "system" fn cuMemHostAlloc(
     pp: *mut *mut ::std::os::raw::c_void,
     bytesize: usize,
     Flags: ::std::os::raw::c_uint,
 ) -> CUresult {
-    r#impl::unimplemented()
+    hipMemAllocHost(pp, bytesize).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2694,7 +2706,7 @@ pub extern "system" fn cuDeviceGetPCIBusId(
     len: ::std::os::raw::c_int,
     dev: CUdevice,
 ) -> CUresult {
-    r#impl::unimplemented()
+    unsafe { hipDeviceGetPCIBusId(pciBusId, len, dev.0) }.into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2883,8 +2895,8 @@ pub extern "system" fn cuMemcpy2DUnaligned_v2(pCopy: *const CUDA_MEMCPY2D) -> CU
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuMemcpy3D_v2(pCopy: *const CUDA_MEMCPY3D) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuMemcpy3D_v2(pCopy: *const CUDA_MEMCPY3D) -> CUresult {
+    r#impl::memory::copy_3d(pCopy).encuda()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -2915,13 +2927,13 @@ pub extern "system" fn cuMemcpyPeerAsync(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuMemcpyHtoDAsync_v2(
+pub unsafe extern "system" fn cuMemcpyHtoDAsync_v2(
     dstDevice: CUdeviceptr,
     srcHost: *const ::std::os::raw::c_void,
     ByteCount: usize,
     hStream: CUstream,
 ) -> CUresult {
-    r#impl::unimplemented()
+    hipMemcpyHtoDAsync(dstDevice.0 as _, srcHost as _, ByteCount, hStream as _).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -3153,16 +3165,16 @@ pub extern "system" fn cuArrayGetDescriptor_v2(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuArrayDestroy(hArray: CUarray) -> CUresult {
-    r#impl::unimplemented()
+pub unsafe extern "system" fn cuArrayDestroy(hArray: CUarray) -> CUresult {
+    hipArrayDestroy(hArray as _).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuArray3DCreate_v2(
+pub unsafe extern "system" fn cuArray3DCreate_v2(
     pHandle: *mut CUarray,
     pAllocateArray: *const CUDA_ARRAY3D_DESCRIPTOR,
 ) -> CUresult {
-    r#impl::unimplemented()
+    hipArray3DCreate(pHandle as _, pAllocateArray as _).into()
 }
 
 #[cfg_attr(not(test), no_mangle)]
@@ -3307,12 +3319,12 @@ pub extern "system" fn cuMemRetainAllocationHandle(
 }
 
 #[cfg_attr(not(test), no_mangle)]
-pub extern "system" fn cuPointerGetAttribute(
+pub unsafe extern "system" fn cuPointerGetAttribute(
     data: *mut ::std::os::raw::c_void,
     attribute: CUpointer_attribute,
     ptr: CUdeviceptr,
 ) -> CUresult {
-    r#impl::unimplemented()
+    r#impl::pointer::get_attribute(data, attribute, ptr).encuda()
 }
 
 #[cfg_attr(not(test), no_mangle)]
diff --git a/zluda/src/impl/context.rs b/zluda/src/impl/context.rs
new file mode 100644
index 0000000..fffceb8
--- /dev/null
+++ b/zluda/src/impl/context.rs
@@ -0,0 +1,24 @@
+use std::ptr;
+
+use crate::cuda::CUlimit;
+use crate::cuda::CUresult;
+
+pub(crate) unsafe fn get_limit(pvalue: *mut usize, limit: CUlimit) -> CUresult {
+    if pvalue == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    if limit == CUlimit::CU_LIMIT_STACK_SIZE {
+        *pvalue = 512; // GTX 1060 reports 1024
+        CUresult::CUDA_SUCCESS
+    } else {
+        CUresult::CUDA_ERROR_NOT_SUPPORTED
+    }
+}
+
+pub(crate) fn set_limit(limit: CUlimit, value: usize) -> CUresult {
+    if limit == CUlimit::CU_LIMIT_STACK_SIZE {
+        CUresult::CUDA_SUCCESS
+    } else {
+        CUresult::CUDA_ERROR_NOT_SUPPORTED
+    }
+}
diff --git a/zluda/src/impl/device.rs b/zluda/src/impl/device.rs
index f234f0b..0c63494 100644
--- a/zluda/src/impl/device.rs
+++ b/zluda/src/impl/device.rs
@@ -1,5 +1,8 @@
 use super::{transmute_lifetime, transmute_lifetime_mut, CUresult};
-use crate::cuda;
+use crate::{
+    cuda::{self, CUdevice, CUdevprop},
+    hip_call,
+};
 use cuda::{CUdevice_attribute, CUuuid_st};
 use hip_runtime_sys::{
     hipDeviceAttribute_t, hipDeviceGetAttribute, hipError_t, hipGetDeviceProperties,
@@ -325,3 +328,21 @@ pub fn get_luid(
     unsafe { *dev_node_mask = 0 };
     Ok(())
 }
+
+pub(crate) unsafe fn get_properties(prop: *mut CUdevprop, dev: CUdevice) -> Result<(), hipError_t> {
+    if prop == ptr::null_mut() {
+        return Err(hipError_t::hipErrorInvalidValue);
+    }
+    let mut hip_props = mem::zeroed();
+    hip_call! { hipGetDeviceProperties(&mut hip_props, dev.0) };
+    (*prop).maxThreadsPerBlock = hip_props.maxThreadsPerBlock;
+    (*prop).maxThreadsDim = hip_props.maxThreadsDim;
+    (*prop).maxGridSize = hip_props.maxGridSize;
+    (*prop).totalConstantMemory = usize::min(hip_props.totalConstMem, i32::MAX as usize) as i32;
+    (*prop).SIMDWidth = hip_props.warpSize;
+    (*prop).memPitch = usize::min(hip_props.memPitch, i32::MAX as usize) as i32;
+    (*prop).regsPerBlock = hip_props.regsPerBlock;
+    (*prop).clockRate = hip_props.clockRate;
+    (*prop).textureAlign = usize::min(hip_props.textureAlignment, i32::MAX as usize) as i32;
+    Ok(())
+}
diff --git a/zluda/src/impl/link.rs b/zluda/src/impl/link.rs
new file mode 100644
index 0000000..928180d
--- /dev/null
+++ b/zluda/src/impl/link.rs
@@ -0,0 +1,67 @@
+use std::{
+    ffi::{c_void, CStr},
+    mem, ptr, slice,
+};
+
+use crate::cuda::{CUjitInputType, CUjit_option, CUlinkState, CUresult};
+
+struct LinkState {
+    modules: Vec<String>,
+}
+
+pub(crate) unsafe fn create(
+    num_options: u32,
+    options: *mut CUjit_option,
+    option_values: *mut *mut c_void,
+    state_out: *mut CUlinkState,
+) -> CUresult {
+    if state_out == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let state = Box::new(LinkState {
+        modules: Vec::new(),
+    });
+    *state_out = mem::transmute(state);
+    CUresult::CUDA_SUCCESS
+}
+
+pub(crate) unsafe fn add_data(
+    state: CUlinkState,
+    type_: CUjitInputType,
+    data: *mut c_void,
+    size: usize,
+    name: *const i8,
+    num_options: u32,
+    options: *mut CUjit_option,
+    option_values: *mut *mut c_void,
+) -> CUresult {
+    if state == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let state: *mut LinkState = mem::transmute(state);
+    let state = &mut *state;
+    // V-RAY specific hack
+    if state.modules.len() == 2 {
+        return CUresult::CUDA_SUCCESS;
+    }
+    let ptx = slice::from_raw_parts(data as *mut u8, size);
+    state.modules.push(
+        CStr::from_bytes_with_nul_unchecked(ptx)
+            .to_string_lossy()
+            .to_string(),
+    );
+    CUresult::CUDA_SUCCESS
+}
+
+pub(crate) fn complete(
+    state: CUlinkState,
+    cubin_out: *mut *mut c_void,
+    size_out: *mut usize,
+) -> CUresult {
+    CUresult::CUDA_SUCCESS
+}
+
+pub(crate) unsafe fn destroy(state: CUlinkState) -> CUresult {
+    let state: Box<LinkState> = mem::transmute(state);
+    CUresult::CUDA_SUCCESS
+}
diff --git a/zluda/src/impl/memory.rs b/zluda/src/impl/memory.rs
new file mode 100644
index 0000000..6041623
--- /dev/null
+++ b/zluda/src/impl/memory.rs
@@ -0,0 +1,55 @@
+use hip_runtime_sys::{
+    hipDrvMemcpy3D, hipError_t, hipMemcpy3D, hipMemcpy3DParms, hipMemoryType, hipPitchedPtr,
+    hipPos, HIP_MEMCPY3D,
+};
+use std::ptr;
+
+use crate::{
+    cuda::{CUDA_MEMCPY3D_st, CUdeviceptr, CUmemorytype, CUresult},
+    hip_call,
+};
+
+// TODO change HIP impl to 64 bits
+pub(crate) unsafe fn copy_3d(cu_copy: *const CUDA_MEMCPY3D_st) -> Result<(), hipError_t> {
+    if cu_copy == ptr::null() {
+        return Err(hipError_t::hipErrorInvalidValue);
+    }
+    let cu_copy = *cu_copy;
+    let hip_copy = HIP_MEMCPY3D {
+        srcXInBytes: cu_copy.srcXInBytes as u32,
+        srcY: cu_copy.srcY as u32,
+        srcZ: cu_copy.srcZ as u32,
+        srcLOD: cu_copy.srcLOD as u32,
+        srcMemoryType: memory_type(cu_copy.srcMemoryType)?,
+        srcHost: cu_copy.srcHost,
+        srcDevice: cu_copy.srcDevice.0 as _,
+        srcArray: cu_copy.srcArray as _,
+        srcPitch: cu_copy.srcPitch as u32,
+        srcHeight: cu_copy.srcHeight as u32,
+        dstXInBytes: cu_copy.dstXInBytes as u32,
+        dstY: cu_copy.dstY as u32,
+        dstZ: cu_copy.dstZ as u32,
+        dstLOD: cu_copy.dstLOD as u32,
+        dstMemoryType: memory_type(cu_copy.dstMemoryType)?,
+        dstHost: cu_copy.dstHost,
+        dstDevice: cu_copy.dstDevice.0 as _,
+        dstArray: cu_copy.dstArray as _,
+        dstPitch: cu_copy.dstPitch as u32,
+        dstHeight: cu_copy.dstHeight as u32,
+        WidthInBytes: cu_copy.WidthInBytes as u32,
+        Height: cu_copy.Height as u32,
+        Depth: cu_copy.Depth as u32,
+    };
+    hip_call! { hipDrvMemcpy3D(&hip_copy) };
+    Ok(())
+}
+
+pub(crate) fn memory_type(cu: CUmemorytype) -> Result<hipMemoryType, hipError_t> {
+    match cu {
+        CUmemorytype::CU_MEMORYTYPE_HOST => Ok(hipMemoryType::hipMemoryTypeHost),
+        CUmemorytype::CU_MEMORYTYPE_DEVICE => Ok(hipMemoryType::hipMemoryTypeDevice),
+        CUmemorytype::CU_MEMORYTYPE_ARRAY => Ok(hipMemoryType::hipMemoryTypeArray),
+        CUmemorytype::CU_MEMORYTYPE_UNIFIED => Ok(hipMemoryType::hipMemoryTypeUnified),
+        _ => Err(hipError_t::hipErrorInvalidValue),
+    }
+}
diff --git a/zluda/src/impl/mod.rs b/zluda/src/impl/mod.rs
index e0d19ae..1335ef6 100644
--- a/zluda/src/impl/mod.rs
+++ b/zluda/src/impl/mod.rs
@@ -20,6 +20,10 @@ pub mod function;
 #[cfg_attr(not(windows), path = "os_unix.rs")]
 pub(crate) mod os;
 pub(crate) mod module;
+pub(crate) mod context;
+pub(crate) mod memory;
+pub(crate) mod link;
+pub(crate) mod pointer;
 
 #[cfg(debug_assertions)]
 pub fn unimplemented() -> CUresult {
@@ -31,6 +35,19 @@ pub fn unimplemented() -> CUresult {
     CUresult::CUDA_ERROR_NOT_SUPPORTED
 }
 
+#[macro_export]
+macro_rules! hip_call {
+    ($expr:expr) => {
+        #[allow(unused_unsafe)]
+        {
+            let err = unsafe { $expr };
+            if err != hip_runtime_sys::hipError_t::hipSuccess {
+                return Result::Err(err);
+            }
+        }
+    };
+}
+
 pub trait HasLivenessCookie: Sized {
     const COOKIE: usize;
     const LIVENESS_FAIL: CUresult;
diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs
index 6575d96..ba09869 100644
--- a/zluda/src/impl/module.rs
+++ b/zluda/src/impl/module.rs
@@ -87,7 +87,7 @@ pub fn load_data_impl(pmod: *mut CUmodule, spirv_data: SpirvModule) -> Result<()
     let err = unsafe { hipGetDeviceProperties(&mut props, dev) };
     let arch_binary = compile_amd(
         &props,
-        &spirv_data.binaries[..],
+        &[&spirv_data.binaries[..]],
         spirv_data.should_link_ptx_impl,
     )
     .map_err(|_| hipError_t::hipErrorUnknown)?;
@@ -115,7 +115,7 @@ const AMDGPU_BITCODE_DEVICE_PREFIX: &'static str = "oclc_isa_version_";
 
 fn compile_amd(
     device_pros: &hipDeviceProp_t,
-    spirv_il: &[u32],
+    spirv_il: &[&[u32]],
     ptx_lib: Option<(&'static [u8], &'static [u8])>,
 ) -> io::Result<Vec<u8>> {
     let null_terminator = device_pros
@@ -133,24 +133,30 @@ fn compile_amd(
             return Err(io::Error::new(io::ErrorKind::Other, ""));
         };
     let dir = tempfile::tempdir()?;
-    let mut spirv = NamedTempFile::new_in(&dir)?;
-    let llvm = NamedTempFile::new_in(&dir)?;
-    let spirv_il_u8 = unsafe {
-        slice::from_raw_parts(
-            spirv_il.as_ptr() as *const u8,
-            spirv_il.len() * mem::size_of::<u32>(),
-        )
-    };
-    spirv.write_all(spirv_il_u8)?;
+    let spirv_files = spirv_il
+        .iter()
+        .map(|spirv| {
+            let mut spirv = NamedTempFile::new_in(&dir)?;
+            let spirv_il_u8 = unsafe {
+                slice::from_raw_parts(
+                    spirv_il.as_ptr() as *const u8,
+                    spirv_il.len() * mem::size_of::<u32>(),
+                )
+            };
+            spirv.write_all(spirv_il_u8)?;
+            Ok::<_, io::Error>(spirv)
+        })
+        .collect::<Result<Vec<_>, _>>()?;
     let llvm_spirv_path = match env::var("LLVM_SPIRV") {
         Ok(path) => Cow::Owned(path),
         Err(_) => Cow::Borrowed(LLVM_SPIRV),
     };
+    let llvm = NamedTempFile::new_in(&dir)?;
     let to_llvm_cmd = Command::new(&*llvm_spirv_path)
         .arg("-r")
         .arg("-o")
         .arg(llvm.path())
-        .arg(spirv.path())
+        .args(spirv_files.iter().map(|f| f.path()))
         .status()?;
     assert!(to_llvm_cmd.success());
     if cfg!(debug_assertions) {
diff --git a/zluda/src/impl/pointer.rs b/zluda/src/impl/pointer.rs
new file mode 100644
index 0000000..2b925cd
--- /dev/null
+++ b/zluda/src/impl/pointer.rs
@@ -0,0 +1,53 @@
+use std::{ffi::c_void, mem, ptr};
+
+use hip_runtime_sys::{hipError_t, hipMemoryType, hipPointerGetAttributes};
+
+use crate::{
+    cuda::{CUdeviceptr, CUmemorytype, CUpointer_attribute},
+    hip_call,
+};
+
+pub(crate) unsafe fn get_attribute(
+    data: *mut c_void,
+    attribute: CUpointer_attribute,
+    ptr: CUdeviceptr,
+) -> Result<(), hipError_t> {
+    if data == ptr::null_mut() {
+        return Err(hipError_t::hipErrorInvalidValue);
+    }
+    let mut attribs = mem::zeroed();
+    hip_call! { hipPointerGetAttributes(&mut attribs, ptr.0 as _) };
+    match attribute {
+        CUpointer_attribute::CU_POINTER_ATTRIBUTE_CONTEXT => {
+            *(data as *mut _) = attribs.device;
+            Ok(())
+        }
+        CUpointer_attribute::CU_POINTER_ATTRIBUTE_MEMORY_TYPE => {
+            *(data as *mut _) = memory_type(attribs.memoryType)?;
+            Ok(())
+        }
+        CUpointer_attribute::CU_POINTER_ATTRIBUTE_DEVICE_POINTER => {
+            *(data as *mut _) = attribs.devicePointer;
+            Ok(())
+        }
+        CUpointer_attribute::CU_POINTER_ATTRIBUTE_HOST_POINTER => {
+            *(data as *mut _) = attribs.hostPointer;
+            Ok(())
+        }
+        CUpointer_attribute::CU_POINTER_ATTRIBUTE_IS_MANAGED => {
+            *(data as *mut _) = attribs.isManaged;
+            Ok(())
+        }
+        _ => Err(hipError_t::hipErrorNotSupported),
+    }
+}
+
+pub(crate) fn memory_type(cu: hipMemoryType) -> Result<CUmemorytype, hipError_t> {
+    match cu {
+        hipMemoryType::hipMemoryTypeHost => Ok(CUmemorytype::CU_MEMORYTYPE_HOST),
+        hipMemoryType::hipMemoryTypeDevice => Ok(CUmemorytype::CU_MEMORYTYPE_DEVICE),
+        hipMemoryType::hipMemoryTypeArray => Ok(CUmemorytype::CU_MEMORYTYPE_ARRAY),
+        hipMemoryType::hipMemoryTypeUnified => Ok(CUmemorytype::CU_MEMORYTYPE_UNIFIED),
+        _ => Err(hipError_t::hipErrorInvalidValue),
+    }
+}