aboutsummaryrefslogtreecommitdiffhomepage
path: root/zluda
diff options
context:
space:
mode:
Diffstat (limited to 'zluda')
-rw-r--r--zluda/src/impl/context.rs2
-rw-r--r--zluda/src/impl/device.rs123
-rw-r--r--zluda/src/impl/driver.rs4
-rw-r--r--zluda/src/impl/memory.rs9
-rw-r--r--zluda/src/impl/mod.rs7
-rw-r--r--zluda/src/impl/module.rs2
-rw-r--r--zluda/src/impl/pointer.rs2
-rw-r--r--zluda/src/lib.rs8
8 files changed, 143 insertions, 14 deletions
diff --git a/zluda/src/impl/context.rs b/zluda/src/impl/context.rs
index 973febc..e812e85 100644
--- a/zluda/src/impl/context.rs
+++ b/zluda/src/impl/context.rs
@@ -1,5 +1,5 @@
use super::{driver, FromCuda, ZludaObject};
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use rustc_hash::FxHashSet;
use std::{cell::RefCell, ptr, sync::Mutex};
diff --git a/zluda/src/impl/device.rs b/zluda/src/impl/device.rs
index 8836c1e..13bce63 100644
--- a/zluda/src/impl/device.rs
+++ b/zluda/src/impl/device.rs
@@ -1,4 +1,4 @@
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use std::{mem, ptr};
@@ -70,6 +70,16 @@ pub(crate) fn get_attribute(
attrib: CUdevice_attribute,
dev_idx: hipDevice_t,
) -> hipError_t {
+ fn get_device_prop(
+ pi: &mut i32,
+ dev_idx: hipDevice_t,
+ f: impl FnOnce(&hipDeviceProp_tR0600) -> i32,
+ ) -> hipError_t {
+ let mut props = unsafe { mem::zeroed() };
+ unsafe { hipGetDevicePropertiesR0600(&mut props, dev_idx)? };
+ *pi = f(&props);
+ Ok(())
+ }
match attrib {
CUdevice_attribute::CU_DEVICE_ATTRIBUTE_WARP_SIZE => {
*pi = 32;
@@ -79,6 +89,110 @@ pub(crate) fn get_attribute(
*pi = 0;
return Ok(());
}
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLayered[2])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture1DLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture1DLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_CAN_TEX2D_GATHER => {
+ return get_device_prop(pi, dev_idx, |props| {
+ (props.maxTexture2DGather[0] > 0 && props.maxTexture2DGather[1] > 0) as i32
+ })
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DGather[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DGather[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture3DAlt[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture3DAlt[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture3DAlt[2])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTextureCubemap)
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTextureCubemapLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTextureCubemapLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface1D)
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface2D[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface2D[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface3D[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface3D[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface3D[2])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface1DLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface1DLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface2DLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface2DLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurface2DLayered[2])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurfaceCubemap)
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurfaceCubemapLayered[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS => {
+ return get_device_prop(pi, dev_idx, |props| props.maxSurfaceCubemapLayered[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture1DLinear)
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLinear[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLinear[1])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DLinear[2])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DMipmap[0])
+ }
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture2DMipmap[1])
+ }
CUdevice_attribute::CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR => {
*pi = COMPUTE_CAPABILITY_MAJOR;
return Ok(());
@@ -87,6 +201,9 @@ pub(crate) fn get_attribute(
*pi = COMPUTE_CAPABILITY_MINOR;
return Ok(());
}
+ CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH => {
+ return get_device_prop(pi, dev_idx, |props| props.maxTexture1DMipmap)
+ }
_ => {}
}
let attrib = remap_attribute! {
@@ -260,7 +377,7 @@ pub(crate) fn get_name(
name: *mut ::core::ffi::c_char,
len: ::core::ffi::c_int,
dev: hipDevice_t,
-) -> cuda_types::CUresult {
+) -> CUresult {
unsafe { hipDeviceGetName(name, len, dev) }?;
let len = len as usize;
let buffer = unsafe { std::slice::from_raw_parts(name, len) };
@@ -287,7 +404,7 @@ pub(crate) fn total_mem_v2(bytes: *mut usize, dev: hipDevice_t) -> hipError_t {
unsafe { hipDeviceTotalMem(bytes, dev) }
}
-pub(crate) fn get_properties(prop: &mut cuda_types::CUdevprop, dev: hipDevice_t) -> hipError_t {
+pub(crate) fn get_properties(prop: &mut CUdevprop, dev: hipDevice_t) -> hipError_t {
let mut hip_props = unsafe { mem::zeroed() };
unsafe { hipGetDevicePropertiesR0600(&mut hip_props, dev) }?;
prop.maxThreadsPerBlock = hip_props.maxThreadsPerBlock;
diff --git a/zluda/src/impl/driver.rs b/zluda/src/impl/driver.rs
index 7ff2f54..5b15afb 100644
--- a/zluda/src/impl/driver.rs
+++ b/zluda/src/impl/driver.rs
@@ -1,4 +1,4 @@
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use std::{
ffi::{CStr, CString},
@@ -74,6 +74,6 @@ pub(crate) fn init(flags: ::core::ffi::c_uint) -> CUresult {
}
pub(crate) fn get_version(version: &mut ::core::ffi::c_int) -> CUresult {
- *version = cuda_types::CUDA_VERSION as i32;
+ *version = cuda_types::cuda::CUDA_VERSION as i32;
Ok(())
}
diff --git a/zluda/src/impl/memory.rs b/zluda/src/impl/memory.rs
index 3843776..18e58e7 100644
--- a/zluda/src/impl/memory.rs
+++ b/zluda/src/impl/memory.rs
@@ -1,4 +1,5 @@
use hip_runtime_sys::*;
+use std::mem;
pub(crate) fn alloc_v2(dptr: *mut hipDeviceptr_t, bytesize: usize) -> hipError_t {
unsafe { hipMalloc(dptr.cast(), bytesize) }?;
@@ -33,3 +34,11 @@ pub(crate) fn get_address_range_v2(
) -> hipError_t {
unsafe { hipMemGetAddressRange(pbase, psize, dptr) }
}
+
+pub(crate) fn set_d32_v2(dst: hipDeviceptr_t, ui: ::core::ffi::c_uint, n: usize) -> hipError_t {
+ unsafe { hipMemsetD32(dst, mem::transmute(ui), n) }
+}
+
+pub(crate) fn set_d8_v2(dst: hipDeviceptr_t, value: ::core::ffi::c_uchar, n: usize) -> hipError_t {
+ unsafe { hipMemsetD8(dst, value, n) }
+}
diff --git a/zluda/src/impl/mod.rs b/zluda/src/impl/mod.rs
index 766b4a5..4d8bc83 100644
--- a/zluda/src/impl/mod.rs
+++ b/zluda/src/impl/mod.rs
@@ -1,4 +1,4 @@
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use std::mem::{self, ManuallyDrop, MaybeUninit};
@@ -107,10 +107,11 @@ from_cuda_nop!(
*const ::core::ffi::c_char,
*mut ::core::ffi::c_void,
*mut *mut ::core::ffi::c_void,
+ u8,
i32,
u32,
usize,
- cuda_types::CUdevprop,
+ cuda_types::cuda::CUdevprop,
CUdevice_attribute
);
from_cuda_transmute!(
@@ -136,7 +137,7 @@ impl<'a> FromCuda<'a, CUlimit> for hipLimit_t {
pub(crate) trait ZludaObject: Sized + Send + Sync {
const COOKIE: usize;
- const LIVENESS_FAIL: CUerror = cuda_types::CUerror::INVALID_VALUE;
+ const LIVENESS_FAIL: CUerror = cuda_types::cuda::CUerror::INVALID_VALUE;
type CudaHandle: Sized;
diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs
index 8b19c1b..b469a89 100644
--- a/zluda/src/impl/module.rs
+++ b/zluda/src/impl/module.rs
@@ -1,5 +1,5 @@
use super::ZludaObject;
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use std::{ffi::CStr, mem};
diff --git a/zluda/src/impl/pointer.rs b/zluda/src/impl/pointer.rs
index 6b458a0..e620bab 100644
--- a/zluda/src/impl/pointer.rs
+++ b/zluda/src/impl/pointer.rs
@@ -1,4 +1,4 @@
-use cuda_types::*;
+use cuda_types::cuda::*;
use hip_runtime_sys::*;
use std::{ffi::c_void, ptr};
diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs
index 1568f47..e058bd7 100644
--- a/zluda/src/lib.rs
+++ b/zluda/src/lib.rs
@@ -1,7 +1,7 @@
pub(crate) mod r#impl;
macro_rules! unimplemented {
- ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => {
+ ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:ty;)*) => {
$(
#[cfg_attr(not(test), no_mangle)]
#[allow(improper_ctypes)]
@@ -14,7 +14,7 @@ macro_rules! unimplemented {
}
macro_rules! implemented {
- ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => {
+ ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:ty;)*) => {
$(
#[cfg_attr(not(test), no_mangle)]
#[allow(improper_ctypes)]
@@ -28,7 +28,7 @@ macro_rules! implemented {
}
macro_rules! implemented_in_function {
- ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => {
+ ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:ty;)*) => {
$(
#[cfg_attr(not(test), no_mangle)]
#[allow(improper_ctypes)]
@@ -72,6 +72,8 @@ cuda_base::cuda_function_declarations!(
cuModuleUnload,
cuPointerGetAttribute,
cuMemGetAddressRange_v2,
+ cuMemsetD32_v2,
+ cuMemsetD8_v2
],
implemented_in_function <= [
cuLaunchKernel,
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-25 01:20:41 +0000