diff options
Diffstat (limited to 'zluda/src')
| -rw-r--r-- | zluda/src/cuda.rs | 4720 | ||||
| -rw-r--r-- | zluda/src/impl/context.rs | 99 | ||||
| -rw-r--r-- | zluda/src/impl/device.rs | 579 | ||||
| -rw-r--r-- | zluda/src/impl/driver.rs | 79 | ||||
| -rw-r--r-- | zluda/src/impl/function.rs | 62 | ||||
| -rw-r--r-- | zluda/src/impl/link.rs | 86 | ||||
| -rw-r--r-- | zluda/src/impl/memory.rs | 80 | ||||
| -rw-r--r-- | zluda/src/impl/mod.rs | 349 | ||||
| -rw-r--r-- | zluda/src/impl/module.rs | 280 | ||||
| -rw-r--r-- | zluda/src/impl/pointer.rs | 57 | ||||
| -rw-r--r-- | zluda/src/impl/test.rs | 157 | ||||
| -rw-r--r-- | zluda/src/lib.rs | 90 |
12 files changed, 819 insertions, 5819 deletions
diff --git a/zluda/src/cuda.rs b/zluda/src/cuda.rs deleted file mode 100644 index e7f5e42..0000000 --- a/zluda/src/cuda.rs +++ /dev/null @@ -1,4720 +0,0 @@ -use hip_runtime_sys::*; - -use super::r#impl; -use super::r#impl::{Decuda, Encuda}; - -/* automatically generated by rust-bindgen 0.55.1 */ - -pub type __uint32_t = ::std::os::raw::c_uint; -pub type __uint64_t = ::std::os::raw::c_ulong; -pub type cuuint32_t = u32; -pub type cuuint64_t = u64; -#[repr(transparent)] -#[derive(Copy, Clone)] -pub struct CUdeviceptr(pub usize); -#[repr(transparent)] -#[derive(Copy, Clone)] -pub struct CUdevice(pub ::std::os::raw::c_int); -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUctx_st { - _unused: [u8; 0], -} -pub type CUcontext = *mut CUctx_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmod_st { - _unused: [u8; 0], -} -pub type CUmodule = *mut CUmod_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUfunc_st { - _unused: [u8; 0], -} -pub type CUfunction = *mut CUfunc_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUarray_st { - _unused: [u8; 0], -} -pub type CUarray = *mut CUarray_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmipmappedArray_st { - _unused: [u8; 0], -} -pub type CUmipmappedArray = *mut CUmipmappedArray_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUtexref_st { - _unused: [u8; 0], -} -pub type CUtexref = *mut CUtexref_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUsurfref_st { - _unused: [u8; 0], -} -pub type CUsurfref = *mut CUsurfref_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUevent_st { - _unused: [u8; 0], -} -pub type CUevent = *mut CUevent_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUstream_st { - _unused: [u8; 0], -} -pub type CUstream = *mut CUstream_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUgraphicsResource_st { - _unused: [u8; 0], -} -pub type CUgraphicsResource = *mut CUgraphicsResource_st; -pub type CUtexObject = ::std::os::raw::c_ulonglong; -pub type CUsurfObject = ::std::os::raw::c_ulonglong; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUextMemory_st { - _unused: [u8; 0], -} -pub type CUexternalMemory = *mut CUextMemory_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUextSemaphore_st { - _unused: [u8; 0], -} -pub type CUexternalSemaphore = *mut CUextSemaphore_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUgraph_st { - _unused: [u8; 0], -} -pub type CUgraph = *mut CUgraph_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUgraphNode_st { - _unused: [u8; 0], -} -pub type CUgraphNode = *mut CUgraphNode_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUgraphExec_st { - _unused: [u8; 0], -} -pub type CUgraphExec = *mut CUgraphExec_st; -#[repr(C)] -#[derive(Copy, Clone, PartialEq, Eq)] -pub struct CUuuid_st { - pub bytes: [::std::os::raw::c_uchar; 16usize], -} -pub type CUuuid = CUuuid_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUipcEventHandle_st { - pub reserved: [::std::os::raw::c_char; 64usize], -} -pub type CUipcEventHandle = CUipcEventHandle_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUipcMemHandle_st { - pub reserved: [::std::os::raw::c_char; 64usize], -} -pub type CUipcMemHandle = CUipcMemHandle_st; -impl CUstreamBatchMemOpType_enum { - pub const CU_STREAM_MEM_OP_WAIT_VALUE_32: CUstreamBatchMemOpType_enum = - CUstreamBatchMemOpType_enum(1); -} -impl CUstreamBatchMemOpType_enum { - pub const CU_STREAM_MEM_OP_WRITE_VALUE_32: CUstreamBatchMemOpType_enum = - CUstreamBatchMemOpType_enum(2); -} -impl CUstreamBatchMemOpType_enum { - pub const CU_STREAM_MEM_OP_WAIT_VALUE_64: CUstreamBatchMemOpType_enum = - CUstreamBatchMemOpType_enum(4); -} -impl CUstreamBatchMemOpType_enum { - pub const CU_STREAM_MEM_OP_WRITE_VALUE_64: CUstreamBatchMemOpType_enum = - CUstreamBatchMemOpType_enum(5); -} -impl CUstreamBatchMemOpType_enum { - pub const CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES: CUstreamBatchMemOpType_enum = - CUstreamBatchMemOpType_enum(3); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUstreamBatchMemOpType_enum(pub ::std::os::raw::c_uint); -pub use self::CUstreamBatchMemOpType_enum as CUstreamBatchMemOpType; -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUstreamBatchMemOpParams_union { - pub operation: CUstreamBatchMemOpType, - pub waitValue: CUstreamBatchMemOpParams_union_CUstreamMemOpWaitValueParams_st, - pub writeValue: CUstreamBatchMemOpParams_union_CUstreamMemOpWriteValueParams_st, - pub flushRemoteWrites: CUstreamBatchMemOpParams_union_CUstreamMemOpFlushRemoteWritesParams_st, - pub pad: [cuuint64_t; 6usize], - _bindgen_union_align: [u64; 6usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUstreamBatchMemOpParams_union_CUstreamMemOpWaitValueParams_st { - pub operation: CUstreamBatchMemOpType, - pub address: CUdeviceptr, - pub __bindgen_anon_1: - CUstreamBatchMemOpParams_union_CUstreamMemOpWaitValueParams_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, - pub alias: CUdeviceptr, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUstreamBatchMemOpParams_union_CUstreamMemOpWaitValueParams_st__bindgen_ty_1 { - pub value: cuuint32_t, - pub value64: cuuint64_t, - _bindgen_union_align: u64, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUstreamBatchMemOpParams_union_CUstreamMemOpWriteValueParams_st { - pub operation: CUstreamBatchMemOpType, - pub address: CUdeviceptr, - pub __bindgen_anon_1: - CUstreamBatchMemOpParams_union_CUstreamMemOpWriteValueParams_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, - pub alias: CUdeviceptr, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUstreamBatchMemOpParams_union_CUstreamMemOpWriteValueParams_st__bindgen_ty_1 { - pub value: cuuint32_t, - pub value64: cuuint64_t, - _bindgen_union_align: u64, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUstreamBatchMemOpParams_union_CUstreamMemOpFlushRemoteWritesParams_st { - pub operation: CUstreamBatchMemOpType, - pub flags: ::std::os::raw::c_uint, -} -pub type CUstreamBatchMemOpParams = CUstreamBatchMemOpParams_union; -impl CUarray_format_enum { - pub const CU_AD_FORMAT_UNSIGNED_INT8: CUarray_format_enum = CUarray_format_enum(1); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_UNSIGNED_INT16: CUarray_format_enum = CUarray_format_enum(2); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_UNSIGNED_INT32: CUarray_format_enum = CUarray_format_enum(3); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_SIGNED_INT8: CUarray_format_enum = CUarray_format_enum(8); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_SIGNED_INT16: CUarray_format_enum = CUarray_format_enum(9); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_SIGNED_INT32: CUarray_format_enum = CUarray_format_enum(10); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_HALF: CUarray_format_enum = CUarray_format_enum(16); -} -impl CUarray_format_enum { - pub const CU_AD_FORMAT_FLOAT: CUarray_format_enum = CUarray_format_enum(32); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUarray_format_enum(pub ::std::os::raw::c_uint); -pub use self::CUarray_format_enum as CUarray_format; -impl CUaddress_mode_enum { - pub const CU_TR_ADDRESS_MODE_WRAP: CUaddress_mode_enum = CUaddress_mode_enum(0); -} -impl CUaddress_mode_enum { - pub const CU_TR_ADDRESS_MODE_CLAMP: CUaddress_mode_enum = CUaddress_mode_enum(1); -} -impl CUaddress_mode_enum { - pub const CU_TR_ADDRESS_MODE_MIRROR: CUaddress_mode_enum = CUaddress_mode_enum(2); -} -impl CUaddress_mode_enum { - pub const CU_TR_ADDRESS_MODE_BORDER: CUaddress_mode_enum = CUaddress_mode_enum(3); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUaddress_mode_enum(pub ::std::os::raw::c_uint); -pub use self::CUaddress_mode_enum as CUaddress_mode; -impl CUfilter_mode_enum { - pub const CU_TR_FILTER_MODE_POINT: CUfilter_mode_enum = CUfilter_mode_enum(0); -} -impl CUfilter_mode_enum { - pub const CU_TR_FILTER_MODE_LINEAR: CUfilter_mode_enum = CUfilter_mode_enum(1); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUfilter_mode_enum(pub ::std::os::raw::c_uint); -pub use self::CUfilter_mode_enum as CUfilter_mode; -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(1); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X: CUdevice_attribute_enum = - CUdevice_attribute_enum(2); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y: CUdevice_attribute_enum = - CUdevice_attribute_enum(3); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z: CUdevice_attribute_enum = - CUdevice_attribute_enum(4); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X: CUdevice_attribute_enum = - CUdevice_attribute_enum(5); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y: CUdevice_attribute_enum = - CUdevice_attribute_enum(6); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z: CUdevice_attribute_enum = - CUdevice_attribute_enum(7); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(8); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(8); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY: CUdevice_attribute_enum = - CUdevice_attribute_enum(9); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_WARP_SIZE: CUdevice_attribute_enum = CUdevice_attribute_enum(10); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_PITCH: CUdevice_attribute_enum = CUdevice_attribute_enum(11); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(12); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_REGISTERS_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(12); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CLOCK_RATE: CUdevice_attribute_enum = CUdevice_attribute_enum(13); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT: CUdevice_attribute_enum = - CUdevice_attribute_enum(14); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_GPU_OVERLAP: CUdevice_attribute_enum = - CUdevice_attribute_enum(15); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT: CUdevice_attribute_enum = - CUdevice_attribute_enum(16); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT: CUdevice_attribute_enum = - CUdevice_attribute_enum(17); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_INTEGRATED: CUdevice_attribute_enum = CUdevice_attribute_enum(18); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY: CUdevice_attribute_enum = - CUdevice_attribute_enum(19); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COMPUTE_MODE: CUdevice_attribute_enum = - CUdevice_attribute_enum(20); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(21); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(22); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(23); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(24); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(25); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(26); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(27); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(28); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(29); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(27); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(28); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES: CUdevice_attribute_enum = - CUdevice_attribute_enum(29); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT: CUdevice_attribute_enum = - CUdevice_attribute_enum(30); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS: CUdevice_attribute_enum = - CUdevice_attribute_enum(31); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_ECC_ENABLED: CUdevice_attribute_enum = - CUdevice_attribute_enum(32); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_PCI_BUS_ID: CUdevice_attribute_enum = CUdevice_attribute_enum(33); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID: CUdevice_attribute_enum = - CUdevice_attribute_enum(34); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_TCC_DRIVER: CUdevice_attribute_enum = CUdevice_attribute_enum(35); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE: CUdevice_attribute_enum = - CUdevice_attribute_enum(36); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(37); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE: CUdevice_attribute_enum = - CUdevice_attribute_enum(38); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(39); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT: CUdevice_attribute_enum = - CUdevice_attribute_enum(40); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING: CUdevice_attribute_enum = - CUdevice_attribute_enum(41); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(42); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(43); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_TEX2D_GATHER: CUdevice_attribute_enum = - CUdevice_attribute_enum(44); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(45); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(46); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE: CUdevice_attribute_enum = - CUdevice_attribute_enum(47); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE: CUdevice_attribute_enum = - CUdevice_attribute_enum(48); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE: CUdevice_attribute_enum = - CUdevice_attribute_enum(49); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID: CUdevice_attribute_enum = - CUdevice_attribute_enum(50); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT: CUdevice_attribute_enum = - CUdevice_attribute_enum(51); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(52); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(53); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(54); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(55); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(56); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(57); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(58); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(59); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(60); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(61); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(62); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(63); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(64); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(65); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(66); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(67); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS: CUdevice_attribute_enum = - CUdevice_attribute_enum(68); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(69); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(70); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(71); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH: CUdevice_attribute_enum = - CUdevice_attribute_enum(72); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(73); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT: CUdevice_attribute_enum = - CUdevice_attribute_enum(74); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(75); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(76); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH: CUdevice_attribute_enum = - CUdevice_attribute_enum(77); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_STREAM_PRIORITIES_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(78); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_GLOBAL_L1_CACHE_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(79); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_LOCAL_L1_CACHE_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(80); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(81); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(82); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MANAGED_MEMORY: CUdevice_attribute_enum = - CUdevice_attribute_enum(83); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD: CUdevice_attribute_enum = - CUdevice_attribute_enum(84); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD_GROUP_ID: CUdevice_attribute_enum = - CUdevice_attribute_enum(85); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_HOST_NATIVE_ATOMIC_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(86); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_SINGLE_TO_DOUBLE_PRECISION_PERF_RATIO: CUdevice_attribute_enum = - CUdevice_attribute_enum(87); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS: CUdevice_attribute_enum = - CUdevice_attribute_enum(88); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS: CUdevice_attribute_enum = - CUdevice_attribute_enum(89); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COMPUTE_PREEMPTION_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(90); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_USE_HOST_POINTER_FOR_REGISTERED_MEM: CUdevice_attribute_enum = - CUdevice_attribute_enum(91); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_MEM_OPS: CUdevice_attribute_enum = - CUdevice_attribute_enum(92); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS: CUdevice_attribute_enum = - CUdevice_attribute_enum(93); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_WAIT_VALUE_NOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(94); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH: CUdevice_attribute_enum = - CUdevice_attribute_enum(95); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH: CUdevice_attribute_enum = - CUdevice_attribute_enum(96); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN: CUdevice_attribute_enum = - CUdevice_attribute_enum(97); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_CAN_FLUSH_REMOTE_WRITES: CUdevice_attribute_enum = - CUdevice_attribute_enum(98); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_HOST_REGISTER_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(99); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS_USES_HOST_PAGE_TABLES: - CUdevice_attribute_enum = CUdevice_attribute_enum(100); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_DIRECT_MANAGED_MEM_ACCESS_FROM_HOST: CUdevice_attribute_enum = - CUdevice_attribute_enum(101); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(102); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED: - CUdevice_attribute_enum = CUdevice_attribute_enum(103); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(104); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_KMT_HANDLE_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(105); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_BLOCKS_PER_MULTIPROCESSOR: CUdevice_attribute_enum = - CUdevice_attribute_enum(106); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(107); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_PERSISTING_L2_CACHE_SIZE: CUdevice_attribute_enum = - CUdevice_attribute_enum(108); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX_ACCESS_POLICY_WINDOW_SIZE: CUdevice_attribute_enum = - CUdevice_attribute_enum(109); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_GPU_DIRECT_RDMA_WITH_CUDA_VMM_SUPPORTED: CUdevice_attribute_enum = - CUdevice_attribute_enum(110); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_RESERVED_SHARED_MEMORY_PER_BLOCK: CUdevice_attribute_enum = - CUdevice_attribute_enum(111); -} -impl CUdevice_attribute_enum { - pub const CU_DEVICE_ATTRIBUTE_MAX: CUdevice_attribute_enum = CUdevice_attribute_enum(112); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUdevice_attribute_enum(pub ::std::os::raw::c_uint); -pub use self::CUdevice_attribute_enum as CUdevice_attribute; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUdevprop_st { - pub maxThreadsPerBlock: ::std::os::raw::c_int, - pub maxThreadsDim: [::std::os::raw::c_int; 3usize], - pub maxGridSize: [::std::os::raw::c_int; 3usize], - pub sharedMemPerBlock: ::std::os::raw::c_int, - pub totalConstantMemory: ::std::os::raw::c_int, - pub SIMDWidth: ::std::os::raw::c_int, - pub memPitch: ::std::os::raw::c_int, - pub regsPerBlock: ::std::os::raw::c_int, - pub clockRate: ::std::os::raw::c_int, - pub textureAlign: ::std::os::raw::c_int, -} -pub type CUdevprop = CUdevprop_st; -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_CONTEXT: CUpointer_attribute_enum = CUpointer_attribute_enum(1); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_MEMORY_TYPE: CUpointer_attribute_enum = - CUpointer_attribute_enum(2); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_DEVICE_POINTER: CUpointer_attribute_enum = - CUpointer_attribute_enum(3); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_HOST_POINTER: CUpointer_attribute_enum = - CUpointer_attribute_enum(4); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_P2P_TOKENS: CUpointer_attribute_enum = - CUpointer_attribute_enum(5); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_SYNC_MEMOPS: CUpointer_attribute_enum = - CUpointer_attribute_enum(6); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_BUFFER_ID: CUpointer_attribute_enum = - CUpointer_attribute_enum(7); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_IS_MANAGED: CUpointer_attribute_enum = - CUpointer_attribute_enum(8); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL: CUpointer_attribute_enum = - CUpointer_attribute_enum(9); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE: CUpointer_attribute_enum = - CUpointer_attribute_enum(10); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_RANGE_START_ADDR: CUpointer_attribute_enum = - CUpointer_attribute_enum(11); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_RANGE_SIZE: CUpointer_attribute_enum = - CUpointer_attribute_enum(12); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_MAPPED: CUpointer_attribute_enum = CUpointer_attribute_enum(13); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES: CUpointer_attribute_enum = - CUpointer_attribute_enum(14); -} -impl CUpointer_attribute_enum { - pub const CU_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE: CUpointer_attribute_enum = - CUpointer_attribute_enum(15); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUpointer_attribute_enum(pub ::std::os::raw::c_uint); -pub use self::CUpointer_attribute_enum as CUpointer_attribute; -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK: CUfunction_attribute_enum = - CUfunction_attribute_enum(0); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES: CUfunction_attribute_enum = - CUfunction_attribute_enum(1); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES: CUfunction_attribute_enum = - CUfunction_attribute_enum(2); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES: CUfunction_attribute_enum = - CUfunction_attribute_enum(3); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_NUM_REGS: CUfunction_attribute_enum = CUfunction_attribute_enum(4); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_PTX_VERSION: CUfunction_attribute_enum = - CUfunction_attribute_enum(5); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_BINARY_VERSION: CUfunction_attribute_enum = - CUfunction_attribute_enum(6); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_CACHE_MODE_CA: CUfunction_attribute_enum = - CUfunction_attribute_enum(7); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: CUfunction_attribute_enum = - CUfunction_attribute_enum(8); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: CUfunction_attribute_enum = - CUfunction_attribute_enum(9); -} -impl CUfunction_attribute_enum { - pub const CU_FUNC_ATTRIBUTE_MAX: CUfunction_attribute_enum = CUfunction_attribute_enum(10); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUfunction_attribute_enum(pub ::std::os::raw::c_uint); -pub use self::CUfunction_attribute_enum as CUfunction_attribute; -impl CUfunc_cache_enum { - pub const CU_FUNC_CACHE_PREFER_NONE: CUfunc_cache_enum = CUfunc_cache_enum(0); -} -impl CUfunc_cache_enum { - pub const CU_FUNC_CACHE_PREFER_SHARED: CUfunc_cache_enum = CUfunc_cache_enum(1); -} -impl CUfunc_cache_enum { - pub const CU_FUNC_CACHE_PREFER_L1: CUfunc_cache_enum = CUfunc_cache_enum(2); -} -impl CUfunc_cache_enum { - pub const CU_FUNC_CACHE_PREFER_EQUAL: CUfunc_cache_enum = CUfunc_cache_enum(3); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUfunc_cache_enum(pub ::std::os::raw::c_uint); -pub use self::CUfunc_cache_enum as CUfunc_cache; -impl CUsharedconfig_enum { - pub const CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE: CUsharedconfig_enum = CUsharedconfig_enum(0); -} -impl CUsharedconfig_enum { - pub const CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE: CUsharedconfig_enum = - CUsharedconfig_enum(1); -} -impl CUsharedconfig_enum { - pub const CU_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE: CUsharedconfig_enum = - CUsharedconfig_enum(2); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUsharedconfig_enum(pub ::std::os::raw::c_uint); -pub use self::CUsharedconfig_enum as CUsharedconfig; -impl CUmemorytype_enum { - pub const CU_MEMORYTYPE_HOST: CUmemorytype_enum = CUmemorytype_enum(1); -} -impl CUmemorytype_enum { - pub const CU_MEMORYTYPE_DEVICE: CUmemorytype_enum = CUmemorytype_enum(2); -} -impl CUmemorytype_enum { - pub const CU_MEMORYTYPE_ARRAY: CUmemorytype_enum = CUmemorytype_enum(3); -} -impl CUmemorytype_enum { - pub const CU_MEMORYTYPE_UNIFIED: CUmemorytype_enum = CUmemorytype_enum(4); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemorytype_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemorytype_enum as CUmemorytype; -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_SET_READ_MOSTLY: CUmem_advise_enum = CUmem_advise_enum(1); -} -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_UNSET_READ_MOSTLY: CUmem_advise_enum = CUmem_advise_enum(2); -} -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_SET_PREFERRED_LOCATION: CUmem_advise_enum = CUmem_advise_enum(3); -} -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_UNSET_PREFERRED_LOCATION: CUmem_advise_enum = CUmem_advise_enum(4); -} -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_SET_ACCESSED_BY: CUmem_advise_enum = CUmem_advise_enum(5); -} -impl CUmem_advise_enum { - pub const CU_MEM_ADVISE_UNSET_ACCESSED_BY: CUmem_advise_enum = CUmem_advise_enum(6); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmem_advise_enum(pub ::std::os::raw::c_uint); -pub use self::CUmem_advise_enum as CUmem_advise; -impl CUmem_range_attribute_enum { - pub const CU_MEM_RANGE_ATTRIBUTE_READ_MOSTLY: CUmem_range_attribute_enum = - CUmem_range_attribute_enum(1); -} -impl CUmem_range_attribute_enum { - pub const CU_MEM_RANGE_ATTRIBUTE_PREFERRED_LOCATION: CUmem_range_attribute_enum = - CUmem_range_attribute_enum(2); -} -impl CUmem_range_attribute_enum { - pub const CU_MEM_RANGE_ATTRIBUTE_ACCESSED_BY: CUmem_range_attribute_enum = - CUmem_range_attribute_enum(3); -} -impl CUmem_range_attribute_enum { - pub const CU_MEM_RANGE_ATTRIBUTE_LAST_PREFETCH_LOCATION: CUmem_range_attribute_enum = - CUmem_range_attribute_enum(4); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmem_range_attribute_enum(pub ::std::os::raw::c_uint); -pub use self::CUmem_range_attribute_enum as CUmem_range_attribute; -impl CUjit_option_enum { - pub const CU_JIT_MAX_REGISTERS: CUjit_option_enum = CUjit_option_enum(0); -} -impl CUjit_option_enum { - pub const CU_JIT_THREADS_PER_BLOCK: CUjit_option_enum = CUjit_option_enum(1); -} -impl CUjit_option_enum { - pub const CU_JIT_WALL_TIME: CUjit_option_enum = CUjit_option_enum(2); -} -impl CUjit_option_enum { - pub const CU_JIT_INFO_LOG_BUFFER: CUjit_option_enum = CUjit_option_enum(3); -} -impl CUjit_option_enum { - pub const CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES: CUjit_option_enum = CUjit_option_enum(4); -} -impl CUjit_option_enum { - pub const CU_JIT_ERROR_LOG_BUFFER: CUjit_option_enum = CUjit_option_enum(5); -} -impl CUjit_option_enum { - pub const CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES: CUjit_option_enum = CUjit_option_enum(6); -} -impl CUjit_option_enum { - pub const CU_JIT_OPTIMIZATION_LEVEL: CUjit_option_enum = CUjit_option_enum(7); -} -impl CUjit_option_enum { - pub const CU_JIT_TARGET_FROM_CUCONTEXT: CUjit_option_enum = CUjit_option_enum(8); -} -impl CUjit_option_enum { - pub const CU_JIT_TARGET: CUjit_option_enum = CUjit_option_enum(9); -} -impl CUjit_option_enum { - pub const CU_JIT_FALLBACK_STRATEGY: CUjit_option_enum = CUjit_option_enum(10); -} -impl CUjit_option_enum { - pub const CU_JIT_GENERATE_DEBUG_INFO: CUjit_option_enum = CUjit_option_enum(11); -} -impl CUjit_option_enum { - pub const CU_JIT_LOG_VERBOSE: CUjit_option_enum = CUjit_option_enum(12); -} -impl CUjit_option_enum { - pub const CU_JIT_GENERATE_LINE_INFO: CUjit_option_enum = CUjit_option_enum(13); -} -impl CUjit_option_enum { - pub const CU_JIT_CACHE_MODE: CUjit_option_enum = CUjit_option_enum(14); -} -impl CUjit_option_enum { - pub const CU_JIT_NEW_SM3X_OPT: CUjit_option_enum = CUjit_option_enum(15); -} -impl CUjit_option_enum { - pub const CU_JIT_FAST_COMPILE: CUjit_option_enum = CUjit_option_enum(16); -} -impl CUjit_option_enum { - pub const CU_JIT_GLOBAL_SYMBOL_NAMES: CUjit_option_enum = CUjit_option_enum(17); -} -impl CUjit_option_enum { - pub const CU_JIT_GLOBAL_SYMBOL_ADDRESSES: CUjit_option_enum = CUjit_option_enum(18); -} -impl CUjit_option_enum { - pub const CU_JIT_GLOBAL_SYMBOL_COUNT: CUjit_option_enum = CUjit_option_enum(19); -} -impl CUjit_option_enum { - pub const CU_JIT_NUM_OPTIONS: CUjit_option_enum = CUjit_option_enum(20); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUjit_option_enum(pub ::std::os::raw::c_uint); -pub use self::CUjit_option_enum as CUjit_option; -impl CUjitInputType_enum { - pub const CU_JIT_INPUT_CUBIN: CUjitInputType_enum = CUjitInputType_enum(0); -} -impl CUjitInputType_enum { - pub const CU_JIT_INPUT_PTX: CUjitInputType_enum = CUjitInputType_enum(1); -} -impl CUjitInputType_enum { - pub const CU_JIT_INPUT_FATBINARY: CUjitInputType_enum = CUjitInputType_enum(2); -} -impl CUjitInputType_enum { - pub const CU_JIT_INPUT_OBJECT: CUjitInputType_enum = CUjitInputType_enum(3); -} -impl CUjitInputType_enum { - pub const CU_JIT_INPUT_LIBRARY: CUjitInputType_enum = CUjitInputType_enum(4); -} -impl CUjitInputType_enum { - pub const CU_JIT_NUM_INPUT_TYPES: CUjitInputType_enum = CUjitInputType_enum(5); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUjitInputType_enum(pub ::std::os::raw::c_uint); -pub use self::CUjitInputType_enum as CUjitInputType; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUlinkState_st { - _unused: [u8; 0], -} -pub type CUlinkState = *mut CUlinkState_st; -impl CUlimit_enum { - pub const CU_LIMIT_STACK_SIZE: CUlimit_enum = CUlimit_enum(0); -} -impl CUlimit_enum { - pub const CU_LIMIT_PRINTF_FIFO_SIZE: CUlimit_enum = CUlimit_enum(1); -} -impl CUlimit_enum { - pub const CU_LIMIT_MALLOC_HEAP_SIZE: CUlimit_enum = CUlimit_enum(2); -} -impl CUlimit_enum { - pub const CU_LIMIT_DEV_RUNTIME_SYNC_DEPTH: CUlimit_enum = CUlimit_enum(3); -} -impl CUlimit_enum { - pub const CU_LIMIT_DEV_RUNTIME_PENDING_LAUNCH_COUNT: CUlimit_enum = CUlimit_enum(4); -} -impl CUlimit_enum { - pub const CU_LIMIT_MAX_L2_FETCH_GRANULARITY: CUlimit_enum = CUlimit_enum(5); -} -impl CUlimit_enum { - pub const CU_LIMIT_PERSISTING_L2_CACHE_SIZE: CUlimit_enum = CUlimit_enum(6); -} -impl CUlimit_enum { - pub const CU_LIMIT_MAX: CUlimit_enum = CUlimit_enum(7); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUlimit_enum(pub ::std::os::raw::c_uint); -pub use self::CUlimit_enum as CUlimit; -impl CUresourcetype_enum { - pub const CU_RESOURCE_TYPE_ARRAY: CUresourcetype_enum = CUresourcetype_enum(0); -} -impl CUresourcetype_enum { - pub const CU_RESOURCE_TYPE_MIPMAPPED_ARRAY: CUresourcetype_enum = CUresourcetype_enum(1); -} -impl CUresourcetype_enum { - pub const CU_RESOURCE_TYPE_LINEAR: CUresourcetype_enum = CUresourcetype_enum(2); -} -impl CUresourcetype_enum { - pub const CU_RESOURCE_TYPE_PITCH2D: CUresourcetype_enum = CUresourcetype_enum(3); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUresourcetype_enum(pub ::std::os::raw::c_uint); -pub use self::CUresourcetype_enum as CUresourcetype; -pub type CUhostFn = - ::std::option::Option<unsafe extern "system" fn(userData: *mut ::std::os::raw::c_void)>; -impl CUaccessProperty_enum { - pub const CU_ACCESS_PROPERTY_NORMAL: CUaccessProperty_enum = CUaccessProperty_enum(0); -} -impl CUaccessProperty_enum { - pub const CU_ACCESS_PROPERTY_STREAMING: CUaccessProperty_enum = CUaccessProperty_enum(1); -} -impl CUaccessProperty_enum { - pub const CU_ACCESS_PROPERTY_PERSISTING: CUaccessProperty_enum = CUaccessProperty_enum(2); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUaccessProperty_enum(pub ::std::os::raw::c_uint); -pub use self::CUaccessProperty_enum as CUaccessProperty; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUaccessPolicyWindow_st { - pub base_ptr: *mut ::std::os::raw::c_void, - pub num_bytes: usize, - pub hitRatio: f32, - pub hitProp: CUaccessProperty, - pub missProp: CUaccessProperty, -} -pub type CUaccessPolicyWindow = CUaccessPolicyWindow_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_KERNEL_NODE_PARAMS_st { - pub func: CUfunction, - pub gridDimX: ::std::os::raw::c_uint, - pub gridDimY: ::std::os::raw::c_uint, - pub gridDimZ: ::std::os::raw::c_uint, - pub blockDimX: ::std::os::raw::c_uint, - pub blockDimY: ::std::os::raw::c_uint, - pub blockDimZ: ::std::os::raw::c_uint, - pub sharedMemBytes: ::std::os::raw::c_uint, - pub kernelParams: *mut *mut ::std::os::raw::c_void, - pub extra: *mut *mut ::std::os::raw::c_void, -} -pub type CUDA_KERNEL_NODE_PARAMS = CUDA_KERNEL_NODE_PARAMS_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_MEMSET_NODE_PARAMS_st { - pub dst: CUdeviceptr, - pub pitch: usize, - pub value: ::std::os::raw::c_uint, - pub elementSize: ::std::os::raw::c_uint, - pub width: usize, - pub height: usize, -} -pub type CUDA_MEMSET_NODE_PARAMS = CUDA_MEMSET_NODE_PARAMS_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_HOST_NODE_PARAMS_st { - pub fn_: CUhostFn, - pub userData: *mut ::std::os::raw::c_void, -} -pub type CUDA_HOST_NODE_PARAMS = CUDA_HOST_NODE_PARAMS_st; -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_KERNEL: CUgraphNodeType_enum = CUgraphNodeType_enum(0); -} -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_MEMCPY: CUgraphNodeType_enum = CUgraphNodeType_enum(1); -} -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_MEMSET: CUgraphNodeType_enum = CUgraphNodeType_enum(2); -} -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_HOST: CUgraphNodeType_enum = CUgraphNodeType_enum(3); -} -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_GRAPH: CUgraphNodeType_enum = CUgraphNodeType_enum(4); -} -impl CUgraphNodeType_enum { - pub const CU_GRAPH_NODE_TYPE_EMPTY: CUgraphNodeType_enum = CUgraphNodeType_enum(5); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUgraphNodeType_enum(pub ::std::os::raw::c_uint); -pub use self::CUgraphNodeType_enum as CUgraphNodeType; -impl CUsynchronizationPolicy_enum { - pub const CU_SYNC_POLICY_AUTO: CUsynchronizationPolicy_enum = CUsynchronizationPolicy_enum(1); -} -impl CUsynchronizationPolicy_enum { - pub const CU_SYNC_POLICY_SPIN: CUsynchronizationPolicy_enum = CUsynchronizationPolicy_enum(2); -} -impl CUsynchronizationPolicy_enum { - pub const CU_SYNC_POLICY_YIELD: CUsynchronizationPolicy_enum = CUsynchronizationPolicy_enum(3); -} -impl CUsynchronizationPolicy_enum { - pub const CU_SYNC_POLICY_BLOCKING_SYNC: CUsynchronizationPolicy_enum = - CUsynchronizationPolicy_enum(4); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUsynchronizationPolicy_enum(pub ::std::os::raw::c_uint); -pub use self::CUsynchronizationPolicy_enum as CUsynchronizationPolicy; -impl CUkernelNodeAttrID_enum { - pub const CU_KERNEL_NODE_ATTRIBUTE_ACCESS_POLICY_WINDOW: CUkernelNodeAttrID_enum = - CUkernelNodeAttrID_enum(1); -} -impl CUkernelNodeAttrID_enum { - pub const CU_KERNEL_NODE_ATTRIBUTE_COOPERATIVE: CUkernelNodeAttrID_enum = - CUkernelNodeAttrID_enum(2); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUkernelNodeAttrID_enum(pub ::std::os::raw::c_uint); -pub use self::CUkernelNodeAttrID_enum as CUkernelNodeAttrID; -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUkernelNodeAttrValue_union { - pub accessPolicyWindow: CUaccessPolicyWindow, - pub cooperative: ::std::os::raw::c_int, - _bindgen_union_align: [u64; 4usize], -} -pub type CUkernelNodeAttrValue = CUkernelNodeAttrValue_union; -impl CUstreamCaptureStatus_enum { - pub const CU_STREAM_CAPTURE_STATUS_NONE: CUstreamCaptureStatus_enum = - CUstreamCaptureStatus_enum(0); -} -impl CUstreamCaptureStatus_enum { - pub const CU_STREAM_CAPTURE_STATUS_ACTIVE: CUstreamCaptureStatus_enum = - CUstreamCaptureStatus_enum(1); -} -impl CUstreamCaptureStatus_enum { - pub const CU_STREAM_CAPTURE_STATUS_INVALIDATED: CUstreamCaptureStatus_enum = - CUstreamCaptureStatus_enum(2); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUstreamCaptureStatus_enum(pub ::std::os::raw::c_uint); -pub use self::CUstreamCaptureStatus_enum as CUstreamCaptureStatus; -impl CUstreamCaptureMode_enum { - pub const CU_STREAM_CAPTURE_MODE_GLOBAL: CUstreamCaptureMode_enum = CUstreamCaptureMode_enum(0); -} -impl CUstreamCaptureMode_enum { - pub const CU_STREAM_CAPTURE_MODE_THREAD_LOCAL: CUstreamCaptureMode_enum = - CUstreamCaptureMode_enum(1); -} -impl CUstreamCaptureMode_enum { - pub const CU_STREAM_CAPTURE_MODE_RELAXED: CUstreamCaptureMode_enum = - CUstreamCaptureMode_enum(2); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUstreamCaptureMode_enum(pub ::std::os::raw::c_uint); -pub use self::CUstreamCaptureMode_enum as CUstreamCaptureMode; -impl CUstreamAttrID_enum { - pub const CU_STREAM_ATTRIBUTE_ACCESS_POLICY_WINDOW: CUstreamAttrID_enum = - CUstreamAttrID_enum(1); -} -impl CUstreamAttrID_enum { - pub const CU_STREAM_ATTRIBUTE_SYNCHRONIZATION_POLICY: CUstreamAttrID_enum = - CUstreamAttrID_enum(3); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUstreamAttrID_enum(pub ::std::os::raw::c_uint); -pub use self::CUstreamAttrID_enum as CUstreamAttrID; -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUstreamAttrValue_union { - pub accessPolicyWindow: CUaccessPolicyWindow, - pub syncPolicy: CUsynchronizationPolicy, - _bindgen_union_align: [u64; 4usize], -} -pub type CUstreamAttrValue = CUstreamAttrValue_union; -impl cudaError_enum { - pub const CUDA_SUCCESS: cudaError_enum = cudaError_enum(0); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_VALUE: cudaError_enum = cudaError_enum(1); -} -impl cudaError_enum { - pub const CUDA_ERROR_OUT_OF_MEMORY: cudaError_enum = cudaError_enum(2); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_INITIALIZED: cudaError_enum = cudaError_enum(3); -} -impl cudaError_enum { - pub const CUDA_ERROR_DEINITIALIZED: cudaError_enum = cudaError_enum(4); -} -impl cudaError_enum { - pub const CUDA_ERROR_PROFILER_DISABLED: cudaError_enum = cudaError_enum(5); -} -impl cudaError_enum { - pub const CUDA_ERROR_PROFILER_NOT_INITIALIZED: cudaError_enum = cudaError_enum(6); -} -impl cudaError_enum { - pub const CUDA_ERROR_PROFILER_ALREADY_STARTED: cudaError_enum = cudaError_enum(7); -} -impl cudaError_enum { - pub const CUDA_ERROR_PROFILER_ALREADY_STOPPED: cudaError_enum = cudaError_enum(8); -} -impl cudaError_enum { - pub const CUDA_ERROR_NO_DEVICE: cudaError_enum = cudaError_enum(100); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_DEVICE: cudaError_enum = cudaError_enum(101); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_IMAGE: cudaError_enum = cudaError_enum(200); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_CONTEXT: cudaError_enum = cudaError_enum(201); -} -impl cudaError_enum { - pub const CUDA_ERROR_CONTEXT_ALREADY_CURRENT: cudaError_enum = cudaError_enum(202); -} -impl cudaError_enum { - pub const CUDA_ERROR_MAP_FAILED: cudaError_enum = cudaError_enum(205); -} -impl cudaError_enum { - pub const CUDA_ERROR_UNMAP_FAILED: cudaError_enum = cudaError_enum(206); -} -impl cudaError_enum { - pub const CUDA_ERROR_ARRAY_IS_MAPPED: cudaError_enum = cudaError_enum(207); -} -impl cudaError_enum { - pub const CUDA_ERROR_ALREADY_MAPPED: cudaError_enum = cudaError_enum(208); -} -impl cudaError_enum { - pub const CUDA_ERROR_NO_BINARY_FOR_GPU: cudaError_enum = cudaError_enum(209); -} -impl cudaError_enum { - pub const CUDA_ERROR_ALREADY_ACQUIRED: cudaError_enum = cudaError_enum(210); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_MAPPED: cudaError_enum = cudaError_enum(211); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_MAPPED_AS_ARRAY: cudaError_enum = cudaError_enum(212); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_MAPPED_AS_POINTER: cudaError_enum = cudaError_enum(213); -} -impl cudaError_enum { - pub const CUDA_ERROR_ECC_UNCORRECTABLE: cudaError_enum = cudaError_enum(214); -} -impl cudaError_enum { - pub const CUDA_ERROR_UNSUPPORTED_LIMIT: cudaError_enum = cudaError_enum(215); -} -impl cudaError_enum { - pub const CUDA_ERROR_CONTEXT_ALREADY_IN_USE: cudaError_enum = cudaError_enum(216); -} -impl cudaError_enum { - pub const CUDA_ERROR_PEER_ACCESS_UNSUPPORTED: cudaError_enum = cudaError_enum(217); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_PTX: cudaError_enum = cudaError_enum(218); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_GRAPHICS_CONTEXT: cudaError_enum = cudaError_enum(219); -} -impl cudaError_enum { - pub const CUDA_ERROR_NVLINK_UNCORRECTABLE: cudaError_enum = cudaError_enum(220); -} -impl cudaError_enum { - pub const CUDA_ERROR_JIT_COMPILER_NOT_FOUND: cudaError_enum = cudaError_enum(221); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_SOURCE: cudaError_enum = cudaError_enum(300); -} -impl cudaError_enum { - pub const CUDA_ERROR_FILE_NOT_FOUND: cudaError_enum = cudaError_enum(301); -} -impl cudaError_enum { - pub const CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND: cudaError_enum = cudaError_enum(302); -} -impl cudaError_enum { - pub const CUDA_ERROR_SHARED_OBJECT_INIT_FAILED: cudaError_enum = cudaError_enum(303); -} -impl cudaError_enum { - pub const CUDA_ERROR_OPERATING_SYSTEM: cudaError_enum = cudaError_enum(304); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_HANDLE: cudaError_enum = cudaError_enum(400); -} -impl cudaError_enum { - pub const CUDA_ERROR_ILLEGAL_STATE: cudaError_enum = cudaError_enum(401); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_FOUND: cudaError_enum = cudaError_enum(500); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_READY: cudaError_enum = cudaError_enum(600); -} -impl cudaError_enum { - pub const CUDA_ERROR_ILLEGAL_ADDRESS: cudaError_enum = cudaError_enum(700); -} -impl cudaError_enum { - pub const CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES: cudaError_enum = cudaError_enum(701); -} -impl cudaError_enum { - pub const CUDA_ERROR_LAUNCH_TIMEOUT: cudaError_enum = cudaError_enum(702); -} -impl cudaError_enum { - pub const CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING: cudaError_enum = cudaError_enum(703); -} -impl cudaError_enum { - pub const CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED: cudaError_enum = cudaError_enum(704); -} -impl cudaError_enum { - pub const CUDA_ERROR_PEER_ACCESS_NOT_ENABLED: cudaError_enum = cudaError_enum(705); -} -impl cudaError_enum { - pub const CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE: cudaError_enum = cudaError_enum(708); -} -impl cudaError_enum { - pub const CUDA_ERROR_CONTEXT_IS_DESTROYED: cudaError_enum = cudaError_enum(709); -} -impl cudaError_enum { - pub const CUDA_ERROR_ASSERT: cudaError_enum = cudaError_enum(710); -} -impl cudaError_enum { - pub const CUDA_ERROR_TOO_MANY_PEERS: cudaError_enum = cudaError_enum(711); -} -impl cudaError_enum { - pub const CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED: cudaError_enum = cudaError_enum(712); -} -impl cudaError_enum { - pub const CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED: cudaError_enum = cudaError_enum(713); -} -impl cudaError_enum { - pub const CUDA_ERROR_HARDWARE_STACK_ERROR: cudaError_enum = cudaError_enum(714); -} -impl cudaError_enum { - pub const CUDA_ERROR_ILLEGAL_INSTRUCTION: cudaError_enum = cudaError_enum(715); -} -impl cudaError_enum { - pub const CUDA_ERROR_MISALIGNED_ADDRESS: cudaError_enum = cudaError_enum(716); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_ADDRESS_SPACE: cudaError_enum = cudaError_enum(717); -} -impl cudaError_enum { - pub const CUDA_ERROR_INVALID_PC: cudaError_enum = cudaError_enum(718); -} -impl cudaError_enum { - pub const CUDA_ERROR_LAUNCH_FAILED: cudaError_enum = cudaError_enum(719); -} -impl cudaError_enum { - pub const CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE: cudaError_enum = cudaError_enum(720); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_PERMITTED: cudaError_enum = cudaError_enum(800); -} -impl cudaError_enum { - pub const CUDA_ERROR_NOT_SUPPORTED: cudaError_enum = cudaError_enum(801); -} -impl cudaError_enum { - pub const CUDA_ERROR_SYSTEM_NOT_READY: cudaError_enum = cudaError_enum(802); -} -impl cudaError_enum { - pub const CUDA_ERROR_SYSTEM_DRIVER_MISMATCH: cudaError_enum = cudaError_enum(803); -} -impl cudaError_enum { - pub const CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE: cudaError_enum = cudaError_enum(804); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED: cudaError_enum = cudaError_enum(900); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_INVALIDATED: cudaError_enum = cudaError_enum(901); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_MERGE: cudaError_enum = cudaError_enum(902); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_UNMATCHED: cudaError_enum = cudaError_enum(903); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_UNJOINED: cudaError_enum = cudaError_enum(904); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_ISOLATION: cudaError_enum = cudaError_enum(905); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_IMPLICIT: cudaError_enum = cudaError_enum(906); -} -impl cudaError_enum { - pub const CUDA_ERROR_CAPTURED_EVENT: cudaError_enum = cudaError_enum(907); -} -impl cudaError_enum { - pub const CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD: cudaError_enum = cudaError_enum(908); -} -impl cudaError_enum { - pub const CUDA_ERROR_TIMEOUT: cudaError_enum = cudaError_enum(909); -} -impl cudaError_enum { - pub const CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE: cudaError_enum = cudaError_enum(910); -} -impl cudaError_enum { - pub const CUDA_ERROR_UNKNOWN: cudaError_enum = cudaError_enum(999); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct cudaError_enum(pub ::std::os::raw::c_uint); -pub use self::cudaError_enum as CUresult; -impl CUdevice_P2PAttribute_enum { - pub const CU_DEVICE_P2P_ATTRIBUTE_PERFORMANCE_RANK: CUdevice_P2PAttribute_enum = - CUdevice_P2PAttribute_enum(1); -} -impl CUdevice_P2PAttribute_enum { - pub const CU_DEVICE_P2P_ATTRIBUTE_ACCESS_SUPPORTED: CUdevice_P2PAttribute_enum = - CUdevice_P2PAttribute_enum(2); -} -impl CUdevice_P2PAttribute_enum { - pub const CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED: CUdevice_P2PAttribute_enum = - CUdevice_P2PAttribute_enum(3); -} -impl CUdevice_P2PAttribute_enum { - pub const CU_DEVICE_P2P_ATTRIBUTE_ACCESS_ACCESS_SUPPORTED: CUdevice_P2PAttribute_enum = - CUdevice_P2PAttribute_enum(4); -} -impl CUdevice_P2PAttribute_enum { - pub const CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED: CUdevice_P2PAttribute_enum = - CUdevice_P2PAttribute_enum(4); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUdevice_P2PAttribute_enum(pub ::std::os::raw::c_uint); -pub use self::CUdevice_P2PAttribute_enum as CUdevice_P2PAttribute; -pub type CUstreamCallback = ::std::option::Option< - unsafe extern "system" fn( - hStream: CUstream, - status: CUresult, - userData: *mut ::std::os::raw::c_void, - ), ->; -pub type CUoccupancyB2DSize = - ::std::option::Option<unsafe extern "system" fn(blockSize: ::std::os::raw::c_int) -> usize>; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_MEMCPY2D_st { - pub srcXInBytes: usize, - pub srcY: usize, - pub srcMemoryType: CUmemorytype, - pub srcHost: *const ::std::os::raw::c_void, - pub srcDevice: CUdeviceptr, - pub srcArray: CUarray, - pub srcPitch: usize, - pub dstXInBytes: usize, - pub dstY: usize, - pub dstMemoryType: CUmemorytype, - pub dstHost: *mut ::std::os::raw::c_void, - pub dstDevice: CUdeviceptr, - pub dstArray: CUarray, - pub dstPitch: usize, - pub WidthInBytes: usize, - pub Height: usize, -} -pub type CUDA_MEMCPY2D = CUDA_MEMCPY2D_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_MEMCPY3D_st { - pub srcXInBytes: usize, - pub srcY: usize, - pub srcZ: usize, - pub srcLOD: usize, - pub srcMemoryType: CUmemorytype, - pub srcHost: *const ::std::os::raw::c_void, - pub srcDevice: CUdeviceptr, - pub srcArray: CUarray, - pub reserved0: *mut ::std::os::raw::c_void, - pub srcPitch: usize, - pub srcHeight: usize, - pub dstXInBytes: usize, - pub dstY: usize, - pub dstZ: usize, - pub dstLOD: usize, - pub dstMemoryType: CUmemorytype, - pub dstHost: *mut ::std::os::raw::c_void, - pub dstDevice: CUdeviceptr, - pub dstArray: CUarray, - pub reserved1: *mut ::std::os::raw::c_void, - pub dstPitch: usize, - pub dstHeight: usize, - pub WidthInBytes: usize, - pub Height: usize, - pub Depth: usize, -} -pub type CUDA_MEMCPY3D = CUDA_MEMCPY3D_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_MEMCPY3D_PEER_st { - pub srcXInBytes: usize, - pub srcY: usize, - pub srcZ: usize, - pub srcLOD: usize, - pub srcMemoryType: CUmemorytype, - pub srcHost: *const ::std::os::raw::c_void, - pub srcDevice: CUdeviceptr, - pub srcArray: CUarray, - pub srcContext: CUcontext, - pub srcPitch: usize, - pub srcHeight: usize, - pub dstXInBytes: usize, - pub dstY: usize, - pub dstZ: usize, - pub dstLOD: usize, - pub dstMemoryType: CUmemorytype, - pub dstHost: *mut ::std::os::raw::c_void, - pub dstDevice: CUdeviceptr, - pub dstArray: CUarray, - pub dstContext: CUcontext, - pub dstPitch: usize, - pub dstHeight: usize, - pub WidthInBytes: usize, - pub Height: usize, - pub Depth: usize, -} -pub type CUDA_MEMCPY3D_PEER = CUDA_MEMCPY3D_PEER_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_ARRAY_DESCRIPTOR_st { - pub Width: usize, - pub Height: usize, - pub Format: CUarray_format, - pub NumChannels: ::std::os::raw::c_uint, -} -pub type CUDA_ARRAY_DESCRIPTOR = CUDA_ARRAY_DESCRIPTOR_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_ARRAY3D_DESCRIPTOR_st { - pub Width: usize, - pub Height: usize, - pub Depth: usize, - pub Format: CUarray_format, - pub NumChannels: ::std::os::raw::c_uint, - pub Flags: ::std::os::raw::c_uint, -} -pub type CUDA_ARRAY3D_DESCRIPTOR = CUDA_ARRAY3D_DESCRIPTOR_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st { - pub resType: CUresourcetype, - pub res: CUDA_RESOURCE_DESC_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUDA_RESOURCE_DESC_st__bindgen_ty_1 { - pub array: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_1, - pub mipmap: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_2, - pub linear: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3, - pub pitch2D: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4, - pub reserved: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5, - _bindgen_union_align: [u64; 16usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_1 { - pub hArray: CUarray, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_2 { - pub hMipmappedArray: CUmipmappedArray, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3 { - pub devPtr: CUdeviceptr, - pub format: CUarray_format, - pub numChannels: ::std::os::raw::c_uint, - pub sizeInBytes: usize, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4 { - pub devPtr: CUdeviceptr, - pub format: CUarray_format, - pub numChannels: ::std::os::raw::c_uint, - pub width: usize, - pub height: usize, - pub pitchInBytes: usize, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5 { - pub reserved: [::std::os::raw::c_int; 32usize], -} -pub type CUDA_RESOURCE_DESC = CUDA_RESOURCE_DESC_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_TEXTURE_DESC_st { - pub addressMode: [CUaddress_mode; 3usize], - pub filterMode: CUfilter_mode, - pub flags: ::std::os::raw::c_uint, - pub maxAnisotropy: ::std::os::raw::c_uint, - pub mipmapFilterMode: CUfilter_mode, - pub mipmapLevelBias: f32, - pub minMipmapLevelClamp: f32, - pub maxMipmapLevelClamp: f32, - pub borderColor: [f32; 4usize], - pub reserved: [::std::os::raw::c_int; 12usize], -} -pub type CUDA_TEXTURE_DESC = CUDA_TEXTURE_DESC_st; -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_NONE: CUresourceViewFormat_enum = CUresourceViewFormat_enum(0); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_1X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(1); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_2X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(2); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_4X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(3); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_1X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(4); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_2X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(5); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_4X8: CUresourceViewFormat_enum = CUresourceViewFormat_enum(6); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_1X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(7); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_2X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(8); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_4X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(9); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_1X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(10); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_2X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(11); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_4X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(12); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_1X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(13); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_2X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(14); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UINT_4X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(15); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_1X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(16); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_2X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(17); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SINT_4X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(18); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_1X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(19); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_2X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(20); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_4X16: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(21); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_1X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(22); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_2X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(23); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_FLOAT_4X32: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(24); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC1: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(25); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC2: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(26); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC3: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(27); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC4: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(28); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SIGNED_BC4: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(29); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC5: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(30); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SIGNED_BC5: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(31); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC6H: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(32); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_SIGNED_BC6H: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(33); -} -impl CUresourceViewFormat_enum { - pub const CU_RES_VIEW_FORMAT_UNSIGNED_BC7: CUresourceViewFormat_enum = - CUresourceViewFormat_enum(34); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUresourceViewFormat_enum(pub ::std::os::raw::c_uint); -pub use self::CUresourceViewFormat_enum as CUresourceViewFormat; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_RESOURCE_VIEW_DESC_st { - pub format: CUresourceViewFormat, - pub width: usize, - pub height: usize, - pub depth: usize, - pub firstMipmapLevel: ::std::os::raw::c_uint, - pub lastMipmapLevel: ::std::os::raw::c_uint, - pub firstLayer: ::std::os::raw::c_uint, - pub lastLayer: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -pub type CUDA_RESOURCE_VIEW_DESC = CUDA_RESOURCE_VIEW_DESC_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_LAUNCH_PARAMS_st { - pub function: CUfunction, - pub gridDimX: ::std::os::raw::c_uint, - pub gridDimY: ::std::os::raw::c_uint, - pub gridDimZ: ::std::os::raw::c_uint, - pub blockDimX: ::std::os::raw::c_uint, - pub blockDimY: ::std::os::raw::c_uint, - pub blockDimZ: ::std::os::raw::c_uint, - pub sharedMemBytes: ::std::os::raw::c_uint, - pub hStream: CUstream, - pub kernelParams: *mut *mut ::std::os::raw::c_void, -} -pub type CUDA_LAUNCH_PARAMS = CUDA_LAUNCH_PARAMS_st; -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(1); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(2); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(3); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_HEAP: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(4); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(5); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(6); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(7); -} -impl CUexternalMemoryHandleType_enum { - pub const CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF: CUexternalMemoryHandleType_enum = - CUexternalMemoryHandleType_enum(8); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUexternalMemoryHandleType_enum(pub ::std::os::raw::c_uint); -pub use self::CUexternalMemoryHandleType_enum as CUexternalMemoryHandleType; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st { - pub type_: CUexternalMemoryHandleType, - pub handle: CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1, - pub size: ::std::os::raw::c_ulonglong, - pub flags: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1 { - pub fd: ::std::os::raw::c_int, - pub win32: CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1, - pub nvSciBufObject: *const ::std::os::raw::c_void, - _bindgen_union_align: [u64; 2usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1 { - pub handle: *mut ::std::os::raw::c_void, - pub name: *const ::std::os::raw::c_void, -} -pub type CUDA_EXTERNAL_MEMORY_HANDLE_DESC = CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st { - pub offset: ::std::os::raw::c_ulonglong, - pub size: ::std::os::raw::c_ulonglong, - pub flags: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -pub type CUDA_EXTERNAL_MEMORY_BUFFER_DESC = CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st { - pub offset: ::std::os::raw::c_ulonglong, - pub arrayDesc: CUDA_ARRAY3D_DESCRIPTOR, - pub numLevels: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -pub type CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC = CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st; -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD: CUexternalSemaphoreHandleType_enum = - CUexternalSemaphoreHandleType_enum(1); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32: CUexternalSemaphoreHandleType_enum = - CUexternalSemaphoreHandleType_enum(2); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: - CUexternalSemaphoreHandleType_enum = CUexternalSemaphoreHandleType_enum(3); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE: CUexternalSemaphoreHandleType_enum = - CUexternalSemaphoreHandleType_enum(4); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE: CUexternalSemaphoreHandleType_enum = - CUexternalSemaphoreHandleType_enum(5); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC: CUexternalSemaphoreHandleType_enum = - CUexternalSemaphoreHandleType_enum(6); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX: - CUexternalSemaphoreHandleType_enum = CUexternalSemaphoreHandleType_enum(7); -} -impl CUexternalSemaphoreHandleType_enum { - pub const CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT: - CUexternalSemaphoreHandleType_enum = CUexternalSemaphoreHandleType_enum(8); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUexternalSemaphoreHandleType_enum(pub ::std::os::raw::c_uint); -pub use self::CUexternalSemaphoreHandleType_enum as CUexternalSemaphoreHandleType; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st { - pub type_: CUexternalSemaphoreHandleType, - pub handle: CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1 { - pub fd: ::std::os::raw::c_int, - pub win32: CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1, - pub nvSciSyncObj: *const ::std::os::raw::c_void, - _bindgen_union_align: [u64; 2usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1 { - pub handle: *mut ::std::os::raw::c_void, - pub name: *const ::std::os::raw::c_void, -} -pub type CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC = CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st { - pub params: CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1 { - pub fence: CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_1, - pub nvSciSync: CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_2, - pub keyedMutex: CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_3, - pub reserved: [::std::os::raw::c_uint; 12usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_1 { - pub value: ::std::os::raw::c_ulonglong, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_2 { - pub fence: *mut ::std::os::raw::c_void, - pub reserved: ::std::os::raw::c_ulonglong, - _bindgen_union_align: u64, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_3 { - pub key: ::std::os::raw::c_ulonglong, -} -pub type CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS = CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st { - pub params: CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1, - pub flags: ::std::os::raw::c_uint, - pub reserved: [::std::os::raw::c_uint; 16usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1 { - pub fence: CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_1, - pub nvSciSync: CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_2, - pub keyedMutex: CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_3, - pub reserved: [::std::os::raw::c_uint; 10usize], -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_1 { - pub value: ::std::os::raw::c_ulonglong, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub union CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_2 { - pub fence: *mut ::std::os::raw::c_void, - pub reserved: ::std::os::raw::c_ulonglong, - _bindgen_union_align: u64, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_3 { - pub key: ::std::os::raw::c_ulonglong, - pub timeoutMs: ::std::os::raw::c_uint, -} -pub type CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS = CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st; -pub type CUmemGenericAllocationHandle = ::std::os::raw::c_ulonglong; -impl CUmemAllocationHandleType_enum { - pub const CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR: CUmemAllocationHandleType_enum = - CUmemAllocationHandleType_enum(1); -} -impl CUmemAllocationHandleType_enum { - pub const CU_MEM_HANDLE_TYPE_WIN32: CUmemAllocationHandleType_enum = - CUmemAllocationHandleType_enum(2); -} -impl CUmemAllocationHandleType_enum { - pub const CU_MEM_HANDLE_TYPE_WIN32_KMT: CUmemAllocationHandleType_enum = - CUmemAllocationHandleType_enum(4); -} -impl CUmemAllocationHandleType_enum { - pub const CU_MEM_HANDLE_TYPE_MAX: CUmemAllocationHandleType_enum = - CUmemAllocationHandleType_enum(4294967295); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemAllocationHandleType_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemAllocationHandleType_enum as CUmemAllocationHandleType; -impl CUmemAccess_flags_enum { - pub const CU_MEM_ACCESS_FLAGS_PROT_NONE: CUmemAccess_flags_enum = CUmemAccess_flags_enum(0); -} -impl CUmemAccess_flags_enum { - pub const CU_MEM_ACCESS_FLAGS_PROT_READ: CUmemAccess_flags_enum = CUmemAccess_flags_enum(1); -} -impl CUmemAccess_flags_enum { - pub const CU_MEM_ACCESS_FLAGS_PROT_READWRITE: CUmemAccess_flags_enum = - CUmemAccess_flags_enum(3); -} -impl CUmemAccess_flags_enum { - pub const CU_MEM_ACCESS_FLAGS_PROT_MAX: CUmemAccess_flags_enum = - CUmemAccess_flags_enum(4294967295); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemAccess_flags_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemAccess_flags_enum as CUmemAccess_flags; -impl CUmemLocationType_enum { - pub const CU_MEM_LOCATION_TYPE_INVALID: CUmemLocationType_enum = CUmemLocationType_enum(0); -} -impl CUmemLocationType_enum { - pub const CU_MEM_LOCATION_TYPE_DEVICE: CUmemLocationType_enum = CUmemLocationType_enum(1); -} -impl CUmemLocationType_enum { - pub const CU_MEM_LOCATION_TYPE_MAX: CUmemLocationType_enum = CUmemLocationType_enum(4294967295); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemLocationType_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemLocationType_enum as CUmemLocationType; -impl CUmemAllocationType_enum { - pub const CU_MEM_ALLOCATION_TYPE_INVALID: CUmemAllocationType_enum = - CUmemAllocationType_enum(0); -} -impl CUmemAllocationType_enum { - pub const CU_MEM_ALLOCATION_TYPE_PINNED: CUmemAllocationType_enum = CUmemAllocationType_enum(1); -} -impl CUmemAllocationType_enum { - pub const CU_MEM_ALLOCATION_TYPE_MAX: CUmemAllocationType_enum = - CUmemAllocationType_enum(4294967295); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemAllocationType_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemAllocationType_enum as CUmemAllocationType; -impl CUmemAllocationGranularity_flags_enum { - pub const CU_MEM_ALLOC_GRANULARITY_MINIMUM: CUmemAllocationGranularity_flags_enum = - CUmemAllocationGranularity_flags_enum(0); -} -impl CUmemAllocationGranularity_flags_enum { - pub const CU_MEM_ALLOC_GRANULARITY_RECOMMENDED: CUmemAllocationGranularity_flags_enum = - CUmemAllocationGranularity_flags_enum(1); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUmemAllocationGranularity_flags_enum(pub ::std::os::raw::c_uint); -pub use self::CUmemAllocationGranularity_flags_enum as CUmemAllocationGranularity_flags; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmemLocation_st { - pub type_: CUmemLocationType, - pub id: ::std::os::raw::c_int, -} -pub type CUmemLocation = CUmemLocation_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmemAllocationProp_st { - pub type_: CUmemAllocationType, - pub requestedHandleTypes: CUmemAllocationHandleType, - pub location: CUmemLocation, - pub win32HandleMetaData: *mut ::std::os::raw::c_void, - pub allocFlags: CUmemAllocationProp_st__bindgen_ty_1, -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmemAllocationProp_st__bindgen_ty_1 { - pub compressionType: ::std::os::raw::c_uchar, - pub gpuDirectRDMACapable: ::std::os::raw::c_uchar, - pub reserved: [::std::os::raw::c_uchar; 6usize], -} -pub type CUmemAllocationProp = CUmemAllocationProp_st; -#[repr(C)] -#[derive(Copy, Clone)] -pub struct CUmemAccessDesc_st { - pub location: CUmemLocation, - pub flags: CUmemAccess_flags, -} -pub type CUmemAccessDesc = CUmemAccessDesc_st; -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_SUCCESS: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(0); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(1); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR_TOPOLOGY_CHANGED: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(2); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR_NODE_TYPE_CHANGED: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(3); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR_FUNCTION_CHANGED: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(4); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR_PARAMETERS_CHANGED: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(5); -} -impl CUgraphExecUpdateResult_enum { - pub const CU_GRAPH_EXEC_UPDATE_ERROR_NOT_SUPPORTED: CUgraphExecUpdateResult_enum = - CUgraphExecUpdateResult_enum(6); -} -#[repr(transparent)] -#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] -pub struct CUgraphExecUpdateResult_enum(pub ::std::os::raw::c_uint); -pub use self::CUgraphExecUpdateResult_enum as CUgraphExecUpdateResult; - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuGetErrorString( - CUresult(e): CUresult, - pStr: *mut *const ::std::os::raw::c_char, -) -> CUresult { - *pStr = hipGetErrorString(hipError_t(e)); - CUresult::CUDA_SUCCESS -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGetErrorName( - error: CUresult, - pStr: *mut *const ::std::os::raw::c_char, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuInit(Flags: ::std::os::raw::c_uint) -> CUresult { - unsafe { hipInit(Flags).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDriverGetVersion(driverVersion: *mut ::std::os::raw::c_int) -> CUresult { - // GeekBench checks this value - // TODO: encode something more sensible - unsafe { *driverVersion = r#impl::driver_get_version() }; - CUresult::CUDA_SUCCESS -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGet( - device: *mut CUdevice, - ordinal: ::std::os::raw::c_int, -) -> CUresult { - unsafe { hipDeviceGet(device as _, ordinal).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetCount(count: *mut ::std::os::raw::c_int) -> CUresult { - unsafe { hipGetDeviceCount(count).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetName( - name: *mut ::std::os::raw::c_char, - len: ::std::os::raw::c_int, - CUdevice(dev): CUdevice, -) -> CUresult { - unsafe { hipDeviceGetName(name, len, dev).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: CUdevice) -> CUresult { - r#impl::device::get_uuid(uuid, dev.0).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetLuid( - luid: *mut ::std::os::raw::c_char, - deviceNodeMask: *mut ::std::os::raw::c_uint, - dev: CUdevice, -) -> CUresult { - r#impl::device::get_luid(luid, deviceNodeMask, dev.0).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceTotalMem_v2(bytes: *mut usize, CUdevice(dev): CUdevice) -> CUresult { - unsafe { hipDeviceTotalMem(bytes, dev).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetAttribute( - pi: *mut ::std::os::raw::c_int, - attrib: CUdevice_attribute, - CUdevice(dev): CUdevice, -) -> CUresult { - r#impl::device::get_attribute(pi, attrib, dev).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetNvSciSyncAttributes( - nvSciSyncAttrList: *mut ::std::os::raw::c_void, - dev: CUdevice, - flags: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -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 unsafe extern "system" fn cuDeviceComputeCapability( - major: *mut ::std::os::raw::c_int, - minor: *mut ::std::os::raw::c_int, - dev: CUdevice, -) -> CUresult { - hipDeviceComputeCapability(major, minor, dev.0).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxRetain( - pctx: *mut CUcontext, - CUdevice(dev): CUdevice, -) -> CUresult { - unsafe { hipDevicePrimaryCtxRetain(pctx as _, dev).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxRelease(dev: CUdevice) -> CUresult { - cuDevicePrimaryCtxRelease_v2(dev) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxRelease_v2(CUdevice(dev): CUdevice) -> CUresult { - unsafe { hipDevicePrimaryCtxRelease(dev).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxSetFlags( - dev: CUdevice, - flags: ::std::os::raw::c_uint, -) -> CUresult { - CUresult::CUDA_SUCCESS -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxSetFlags_v2( - dev: CUdevice, - flags: ::std::os::raw::c_uint, -) -> CUresult { - cuDevicePrimaryCtxSetFlags(dev, flags) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxGetState( - CUdevice(dev): CUdevice, - flags: *mut ::std::os::raw::c_uint, - active: *mut ::std::os::raw::c_int, -) -> CUresult { - unsafe { hipDevicePrimaryCtxGetState(dev, flags, active).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxReset(dev: CUdevice) -> CUresult { - cuDevicePrimaryCtxReset_v2(dev) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDevicePrimaryCtxReset_v2(dev: CUdevice) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxCreate_v2( - pctx: *mut CUcontext, - flags: ::std::os::raw::c_uint, - CUdevice(dev): CUdevice, -) -> CUresult { - unsafe { hipCtxCreate(pctx as _, flags, dev).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxDestroy_v2(ctx: CUcontext) -> CUresult { - unsafe { hipCtxDestroy(ctx as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxPushCurrent_v2(ctx: CUcontext) -> CUresult { - unsafe { hipCtxPushCurrent(ctx as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxPopCurrent_v2(pctx: *mut CUcontext) -> CUresult { - unsafe { hipCtxPopCurrent(pctx as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxSetCurrent(ctx: CUcontext) -> CUresult { - unsafe { hipCtxSetCurrent(ctx as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetCurrent(pctx: *mut CUcontext) -> CUresult { - unsafe { hipCtxGetCurrent(pctx as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetDevice(device: *mut CUdevice) -> CUresult { - unsafe { hipCtxGetDevice(device as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetFlags(flags: *mut ::std::os::raw::c_uint) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxSynchronize() -> CUresult { - // hipCtxSynchronize is not implemented - unsafe { hipDeviceSynchronize().into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxSetLimit(limit: CUlimit, value: usize) -> CUresult { - r#impl::context::set_limit(limit, value) -} - -#[cfg_attr(not(test), no_mangle)] -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)] -pub extern "system" fn cuCtxGetCacheConfig(pconfig: *mut CUfunc_cache) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxSetCacheConfig(config: CUfunc_cache) -> CUresult { - CUresult::CUDA_SUCCESS -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetSharedMemConfig(pConfig: *mut CUsharedconfig) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxSetSharedMemConfig(config: CUsharedconfig) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetApiVersion( - ctx: CUcontext, - version: *mut ::std::os::raw::c_uint, -) -> CUresult { - unsafe { hipCtxGetApiVersion(ctx as _, version as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxGetStreamPriorityRange( - leastPriority: *mut ::std::os::raw::c_int, - greatestPriority: *mut ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxResetPersistingL2Cache() -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxAttach( - pctx: *mut CUcontext, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxDetach(ctx: CUcontext) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleLoad( - module: *mut CUmodule, - fname: *const ::std::os::raw::c_char, -) -> CUresult { - r#impl::module::load(module, fname).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleLoadData( - module: *mut CUmodule, - image: *const ::std::os::raw::c_void, -) -> CUresult { - r#impl::module::load_data(module, image).encuda() -} - -// TODO: parse jit options -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleLoadDataEx( - module: *mut CUmodule, - image: *const ::std::os::raw::c_void, - numOptions: ::std::os::raw::c_uint, - options: *mut CUjit_option, - optionValues: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::module::load_data(module, image).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleLoadFatBinary( - module: *mut CUmodule, - fatCubin: *const ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleUnload(hmod: CUmodule) -> CUresult { - unsafe { hipModuleUnload(hmod as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleGetFunction( - hfunc: *mut CUfunction, - hmod: CUmodule, - name: *const ::std::os::raw::c_char, -) -> CUresult { - unsafe { hipModuleGetFunction(hfunc as _, hmod as _, name).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuModuleGetGlobal_v2( - dptr: *mut CUdeviceptr, - bytes: *mut usize, - hmod: CUmodule, - name: *const ::std::os::raw::c_char, -) -> CUresult { - hipModuleGetGlobal(dptr as _, bytes, hmod as _, name).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleGetTexRef( - pTexRef: *mut CUtexref, - hmod: CUmodule, - name: *const ::std::os::raw::c_char, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuModuleGetSurfRef( - pSurfRef: *mut CUsurfref, - hmod: CUmodule, - name: *const ::std::os::raw::c_char, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -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::link::create(numOptions, options, optionValues, stateOut) -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuLinkAddData_v2( - state: CUlinkState, - type_: CUjitInputType, - data: *mut ::std::os::raw::c_void, - size: usize, - name: *const ::std::os::raw::c_char, - numOptions: ::std::os::raw::c_uint, - options: *mut CUjit_option, - optionValues: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::link::add_data( - state, - type_, - data, - size, - name, - numOptions, - options, - optionValues, - ) - .encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLinkAddFile_v2( - state: CUlinkState, - type_: CUjitInputType, - path: *const ::std::os::raw::c_char, - numOptions: ::std::os::raw::c_uint, - options: *mut CUjit_option, - optionValues: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuLinkComplete( - state: CUlinkState, - cubinOut: *mut *mut ::std::os::raw::c_void, - sizeOut: *mut usize, -) -> CUresult { - r#impl::link::complete(state, cubinOut, sizeOut).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuLinkDestroy(state: CUlinkState) -> CUresult { - r#impl::link::destroy(state) -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult { - hipMemGetInfo(free, total).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult { - let mut dev_ptr = std::ptr::null_mut(); - let err = unsafe { hipMalloc(&mut dev_ptr, bytesize) }.into(); - if err != CUresult::CUDA_SUCCESS { - return err; - } - // HACK ALERT: GeekBench is buggy and sometimes assumes that buffers are zeroed-out on creation - let err = unsafe { hipMemsetD8(dev_ptr, 0, bytesize) }.into(); - /* - let bytesize_rounded_down = bytesize & !3usize; - let bytes = usize::min(bytesize_rounded_down, 4096); - let err = unsafe { hipMemsetD32(dev_ptr, 0, bytes / 1024).into() }; - */ - if err != CUresult::CUDA_SUCCESS { - return err; - } - unsafe { *dptr = CUdeviceptr(dev_ptr as usize) }; - CUresult::CUDA_SUCCESS -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAllocPitch_v2( - dptr: *mut CUdeviceptr, - pPitch: *mut usize, - WidthInBytes: usize, - Height: usize, - ElementSizeBytes: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult { - unsafe { hipFree(dptr.0 as _).into() } -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemGetAddressRange_v2( - pbase: *mut CUdeviceptr, - psize: *mut usize, - dptr: CUdeviceptr, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAllocHost_v2( - pp: *mut *mut ::std::os::raw::c_void, - bytesize: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult { - hipFreeHost(p).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuMemHostAlloc( - pp: *mut *mut ::std::os::raw::c_void, - bytesize: usize, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - hipMemAllocHost(pp, bytesize).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemHostGetDevicePointer_v2( - pdptr: *mut CUdeviceptr, - p: *mut ::std::os::raw::c_void, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemHostGetFlags( - pFlags: *mut ::std::os::raw::c_uint, - p: *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAllocManaged( - dptr: *mut CUdeviceptr, - bytesize: usize, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetByPCIBusId( - dev: *mut CUdevice, - pciBusId: *const ::std::os::raw::c_char, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetPCIBusId( - pciBusId: *mut ::std::os::raw::c_char, - len: ::std::os::raw::c_int, - dev: CUdevice, -) -> CUresult { - unsafe { hipDeviceGetPCIBusId(pciBusId, len, dev.0) }.into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcGetEventHandle( - pHandle: *mut CUipcEventHandle, - event: CUevent, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcOpenEventHandle( - phEvent: *mut CUevent, - handle: CUipcEventHandle, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcGetMemHandle( - pHandle: *mut CUipcMemHandle, - dptr: CUdeviceptr, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcOpenMemHandle( - pdptr: *mut CUdeviceptr, - handle: CUipcMemHandle, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcOpenMemHandle_v2( - pdptr: *mut CUdeviceptr, - handle: CUipcMemHandle, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuIpcCloseMemHandle(dptr: CUdeviceptr) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemHostRegister_v2( - p: *mut ::std::os::raw::c_void, - bytesize: usize, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemHostUnregister(p: *mut ::std::os::raw::c_void) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy(dst: CUdeviceptr, src: CUdeviceptr, ByteCount: usize) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyPeer( - dstDevice: CUdeviceptr, - dstContext: CUcontext, - srcDevice: CUdeviceptr, - srcContext: CUcontext, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyHtoD_v2( - dstDevice: CUdeviceptr, - srcHost: *const ::std::os::raw::c_void, - ByteCount: usize, -) -> CUresult { - unsafe { hipMemcpyHtoD(dstDevice.0 as _, srcHost as _, ByteCount).into() } -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyHtoD_v2_ptds( - dstDevice: CUdeviceptr, - srcHost: *const ::std::os::raw::c_void, - ByteCount: usize, -) -> CUresult { - cuMemcpyHtoD_v2(dstDevice, srcHost, ByteCount) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoH_v2( - dstHost: *mut ::std::os::raw::c_void, - srcDevice: CUdeviceptr, - ByteCount: usize, -) -> CUresult { - unsafe { hipMemcpyDtoH(dstHost as _, srcDevice.0 as _, ByteCount).into() } -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoH_v2_ptds( - dstHost: *mut ::std::os::raw::c_void, - srcDevice: CUdeviceptr, - ByteCount: usize, -) -> CUresult { - cuMemcpyDtoH_v2(dstHost, srcDevice, ByteCount) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoD_v2( - dstDevice: CUdeviceptr, - srcDevice: CUdeviceptr, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoA_v2( - dstArray: CUarray, - dstOffset: usize, - srcDevice: CUdeviceptr, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyAtoD_v2( - dstDevice: CUdeviceptr, - srcArray: CUarray, - srcOffset: usize, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyHtoA_v2( - dstArray: CUarray, - dstOffset: usize, - srcHost: *const ::std::os::raw::c_void, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyAtoH_v2( - dstHost: *mut ::std::os::raw::c_void, - srcArray: CUarray, - srcOffset: usize, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyAtoA_v2( - dstArray: CUarray, - dstOffset: usize, - srcArray: CUarray, - srcOffset: usize, - ByteCount: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy2D_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy2DUnaligned_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -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)] -pub extern "system" fn cuMemcpy3DPeer(pCopy: *const CUDA_MEMCPY3D_PEER) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyAsync( - dst: CUdeviceptr, - src: CUdeviceptr, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyPeerAsync( - dstDevice: CUdeviceptr, - dstContext: CUcontext, - srcDevice: CUdeviceptr, - srcContext: CUcontext, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuMemcpyHtoDAsync_v2( - dstDevice: CUdeviceptr, - srcHost: *const ::std::os::raw::c_void, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - hipMemcpyHtoDAsync(dstDevice.0 as _, srcHost as _, ByteCount, hStream as _).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoHAsync_v2( - dstHost: *mut ::std::os::raw::c_void, - srcDevice: CUdeviceptr, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyDtoDAsync_v2( - dstDevice: CUdeviceptr, - srcDevice: CUdeviceptr, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyHtoAAsync_v2( - dstArray: CUarray, - dstOffset: usize, - srcHost: *const ::std::os::raw::c_void, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpyAtoHAsync_v2( - dstHost: *mut ::std::os::raw::c_void, - srcArray: CUarray, - srcOffset: usize, - ByteCount: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy2DAsync_v2( - pCopy: *const CUDA_MEMCPY2D, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy3DAsync_v2( - pCopy: *const CUDA_MEMCPY3D, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemcpy3DPeerAsync( - pCopy: *const CUDA_MEMCPY3D_PEER, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD8_v2( - dstDevice: CUdeviceptr, - uc: ::std::os::raw::c_uchar, - N: usize, -) -> CUresult { - unsafe { hipMemsetD8(dstDevice.0 as _, uc, N).into() } -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD8_v2_ptds( - dstDevice: CUdeviceptr, - uc: ::std::os::raw::c_uchar, - N: usize, -) -> CUresult { - cuMemsetD8_v2(dstDevice, uc, N) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD16_v2( - dstDevice: CUdeviceptr, - us: ::std::os::raw::c_ushort, - N: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD32_v2( - dstDevice: CUdeviceptr, - ui: ::std::os::raw::c_uint, - N: usize, -) -> CUresult { - unsafe { hipMemsetD32(dstDevice.0 as _, ui as _, N).into() } -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD32_v2_ptds( - dstDevice: CUdeviceptr, - ui: ::std::os::raw::c_uint, - N: usize, -) -> CUresult { - cuMemsetD32_v2(dstDevice, ui, N) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D8_v2( - dstDevice: CUdeviceptr, - dstPitch: usize, - uc: ::std::os::raw::c_uchar, - Width: usize, - Height: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D16_v2( - dstDevice: CUdeviceptr, - dstPitch: usize, - us: ::std::os::raw::c_ushort, - Width: usize, - Height: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D32_v2( - dstDevice: CUdeviceptr, - dstPitch: usize, - ui: ::std::os::raw::c_uint, - Width: usize, - Height: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD8Async( - dstDevice: CUdeviceptr, - uc: ::std::os::raw::c_uchar, - N: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD16Async( - dstDevice: CUdeviceptr, - us: ::std::os::raw::c_ushort, - N: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD32Async( - dstDevice: CUdeviceptr, - ui: ::std::os::raw::c_uint, - N: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D8Async( - dstDevice: CUdeviceptr, - dstPitch: usize, - uc: ::std::os::raw::c_uchar, - Width: usize, - Height: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D16Async( - dstDevice: CUdeviceptr, - dstPitch: usize, - us: ::std::os::raw::c_ushort, - Width: usize, - Height: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemsetD2D32Async( - dstDevice: CUdeviceptr, - dstPitch: usize, - ui: ::std::os::raw::c_uint, - Width: usize, - Height: usize, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuArrayCreate_v2( - pHandle: *mut CUarray, - pAllocateArray: *const CUDA_ARRAY_DESCRIPTOR, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuArrayGetDescriptor_v2( - pArrayDescriptor: *mut CUDA_ARRAY_DESCRIPTOR, - hArray: CUarray, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuArrayDestroy(hArray: CUarray) -> CUresult { - hipArrayDestroy(hArray as _).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuArray3DCreate_v2( - pHandle: *mut CUarray, - pAllocateArray: *const CUDA_ARRAY3D_DESCRIPTOR, -) -> CUresult { - hipArray3DCreate(pHandle as _, pAllocateArray as _).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuArray3DGetDescriptor_v2( - pArrayDescriptor: *mut CUDA_ARRAY3D_DESCRIPTOR, - hArray: CUarray, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMipmappedArrayCreate( - pHandle: *mut CUmipmappedArray, - pMipmappedArrayDesc: *const CUDA_ARRAY3D_DESCRIPTOR, - numMipmapLevels: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMipmappedArrayGetLevel( - pLevelArray: *mut CUarray, - hMipmappedArray: CUmipmappedArray, - level: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMipmappedArrayDestroy(hMipmappedArray: CUmipmappedArray) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAddressReserve( - ptr: *mut CUdeviceptr, - size: usize, - alignment: usize, - addr: CUdeviceptr, - flags: ::std::os::raw::c_ulonglong, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAddressFree(ptr: CUdeviceptr, size: usize) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemCreate( - handle: *mut CUmemGenericAllocationHandle, - size: usize, - prop: *const CUmemAllocationProp, - flags: ::std::os::raw::c_ulonglong, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemRelease(handle: CUmemGenericAllocationHandle) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemMap( - ptr: CUdeviceptr, - size: usize, - offset: usize, - handle: CUmemGenericAllocationHandle, - flags: ::std::os::raw::c_ulonglong, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemUnmap(ptr: CUdeviceptr, size: usize) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemSetAccess( - ptr: CUdeviceptr, - size: usize, - desc: *const CUmemAccessDesc, - count: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemGetAccess( - flags: *mut ::std::os::raw::c_ulonglong, - location: *const CUmemLocation, - ptr: CUdeviceptr, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemExportToShareableHandle( - shareableHandle: *mut ::std::os::raw::c_void, - handle: CUmemGenericAllocationHandle, - handleType: CUmemAllocationHandleType, - flags: ::std::os::raw::c_ulonglong, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemImportFromShareableHandle( - handle: *mut CUmemGenericAllocationHandle, - osHandle: *mut ::std::os::raw::c_void, - shHandleType: CUmemAllocationHandleType, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemGetAllocationGranularity( - granularity: *mut usize, - prop: *const CUmemAllocationProp, - option: CUmemAllocationGranularity_flags, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemGetAllocationPropertiesFromHandle( - prop: *mut CUmemAllocationProp, - handle: CUmemGenericAllocationHandle, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemRetainAllocationHandle( - handle: *mut CUmemGenericAllocationHandle, - addr: *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub unsafe extern "system" fn cuPointerGetAttribute( - data: *mut ::std::os::raw::c_void, - attribute: CUpointer_attribute, - ptr: CUdeviceptr, -) -> CUresult { - r#impl::pointer::get_attribute(data, attribute, ptr).encuda() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemPrefetchAsync( - devPtr: CUdeviceptr, - count: usize, - dstDevice: CUdevice, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemAdvise( - devPtr: CUdeviceptr, - count: usize, - advice: CUmem_advise, - device: CUdevice, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemRangeGetAttribute( - data: *mut ::std::os::raw::c_void, - dataSize: usize, - attribute: CUmem_range_attribute, - devPtr: CUdeviceptr, - count: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuMemRangeGetAttributes( - data: *mut *mut ::std::os::raw::c_void, - dataSizes: *mut usize, - attributes: *mut CUmem_range_attribute, - numAttributes: usize, - devPtr: CUdeviceptr, - count: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuPointerSetAttribute( - value: *const ::std::os::raw::c_void, - attribute: CUpointer_attribute, - ptr: CUdeviceptr, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuPointerGetAttributes( - numAttributes: ::std::os::raw::c_uint, - attributes: *mut CUpointer_attribute, - data: *mut *mut ::std::os::raw::c_void, - ptr: CUdeviceptr, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamCreate( - phStream: *mut CUstream, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - unsafe { hipStreamCreateWithFlags(phStream as _, Flags) }.into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamCreateWithPriority( - phStream: *mut CUstream, - flags: ::std::os::raw::c_uint, - priority: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetPriority( - hStream: CUstream, - priority: *mut ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetFlags( - hStream: CUstream, - flags: *mut ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUresult { - unsafe { hipStreamGetCtx(hStream as _, pctx as _) }.into() -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetCtx_ptsz(hStream: CUstream, pctx: *mut CUcontext) -> CUresult { - cuStreamGetCtx(hStream, pctx) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamWaitEvent( - hStream: CUstream, - hEvent: CUevent, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamAddCallback( - hStream: CUstream, - callback: CUstreamCallback, - userData: *mut ::std::os::raw::c_void, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamBeginCapture_v2( - hStream: CUstream, - mode: CUstreamCaptureMode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuThreadExchangeStreamCaptureMode( - mode: *mut CUstreamCaptureMode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamEndCapture(hStream: CUstream, phGraph: *mut CUgraph) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamIsCapturing( - hStream: CUstream, - captureStatus: *mut CUstreamCaptureStatus, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetCaptureInfo( - hStream: CUstream, - captureStatus: *mut CUstreamCaptureStatus, - id: *mut cuuint64_t, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamAttachMemAsync( - hStream: CUstream, - dptr: CUdeviceptr, - length: usize, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamQuery(hStream: CUstream) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamSynchronize(hStream: CUstream) -> CUresult { - unsafe { hipStreamSynchronize(hStream as _) }.into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamDestroy_v2(hStream: CUstream) -> CUresult { - unsafe { hipStreamDestroy(hStream as _) }.into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamCopyAttributes(dst: CUstream, src: CUstream) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamGetAttribute( - hStream: CUstream, - attr: CUstreamAttrID, - value_out: *mut CUstreamAttrValue, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamSetAttribute( - hStream: CUstream, - attr: CUstreamAttrID, - value: *const CUstreamAttrValue, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventCreate( - phEvent: *mut CUevent, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventRecord(hEvent: CUevent, hStream: CUstream) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventQuery(hEvent: CUevent) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventSynchronize(hEvent: CUevent) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventDestroy_v2(hEvent: CUevent) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuEventElapsedTime( - pMilliseconds: *mut f32, - hStart: CUevent, - hEnd: CUevent, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuImportExternalMemory( - extMem_out: *mut CUexternalMemory, - memHandleDesc: *const CUDA_EXTERNAL_MEMORY_HANDLE_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuExternalMemoryGetMappedBuffer( - devPtr: *mut CUdeviceptr, - extMem: CUexternalMemory, - bufferDesc: *const CUDA_EXTERNAL_MEMORY_BUFFER_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuExternalMemoryGetMappedMipmappedArray( - mipmap: *mut CUmipmappedArray, - extMem: CUexternalMemory, - mipmapDesc: *const CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDestroyExternalMemory(extMem: CUexternalMemory) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuImportExternalSemaphore( - extSem_out: *mut CUexternalSemaphore, - semHandleDesc: *const CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSignalExternalSemaphoresAsync( - extSemArray: *const CUexternalSemaphore, - paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS, - numExtSems: ::std::os::raw::c_uint, - stream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuWaitExternalSemaphoresAsync( - extSemArray: *const CUexternalSemaphore, - paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS, - numExtSems: ::std::os::raw::c_uint, - stream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDestroyExternalSemaphore(extSem: CUexternalSemaphore) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamWaitValue32( - stream: CUstream, - addr: CUdeviceptr, - value: cuuint32_t, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamWaitValue64( - stream: CUstream, - addr: CUdeviceptr, - value: cuuint64_t, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamWriteValue32( - stream: CUstream, - addr: CUdeviceptr, - value: cuuint32_t, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamWriteValue64( - stream: CUstream, - addr: CUdeviceptr, - value: cuuint64_t, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuStreamBatchMemOp( - stream: CUstream, - count: ::std::os::raw::c_uint, - paramArray: *mut CUstreamBatchMemOpParams, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncGetAttribute( - pi: *mut ::std::os::raw::c_int, - attrib: CUfunction_attribute, - hfunc: CUfunction, -) -> CUresult { - r#impl::function::get_attribute(pi, attrib, hfunc).into() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncSetAttribute( - hfunc: CUfunction, - attrib: CUfunction_attribute, - value: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncSetCacheConfig(hfunc: CUfunction, config: CUfunc_cache) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncSetSharedMemConfig( - hfunc: CUfunction, - config: CUsharedconfig, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchKernel( - f: CUfunction, - gridDimX: ::std::os::raw::c_uint, - gridDimY: ::std::os::raw::c_uint, - gridDimZ: ::std::os::raw::c_uint, - blockDimX: ::std::os::raw::c_uint, - blockDimY: ::std::os::raw::c_uint, - blockDimZ: ::std::os::raw::c_uint, - sharedMemBytes: ::std::os::raw::c_uint, - hStream: CUstream, - kernelParams: *mut *mut ::std::os::raw::c_void, - extra: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - unsafe { - hipModuleLaunchKernel( - f as _, - gridDimX, - gridDimY, - gridDimZ, - blockDimX, - blockDimY, - blockDimZ, - sharedMemBytes, - hStream as _, - kernelParams, - extra, - ) - } - .into() -} - -// TODO: implement default stream semantics -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchKernel_ptsz( - f: CUfunction, - gridDimX: ::std::os::raw::c_uint, - gridDimY: ::std::os::raw::c_uint, - gridDimZ: ::std::os::raw::c_uint, - blockDimX: ::std::os::raw::c_uint, - blockDimY: ::std::os::raw::c_uint, - blockDimZ: ::std::os::raw::c_uint, - sharedMemBytes: ::std::os::raw::c_uint, - hStream: CUstream, - kernelParams: *mut *mut ::std::os::raw::c_void, - extra: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - cuLaunchKernel( - f, - gridDimX, - gridDimY, - gridDimZ, - blockDimX, - blockDimY, - blockDimZ, - sharedMemBytes, - hStream, - kernelParams, - extra, - ) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchCooperativeKernel( - f: CUfunction, - gridDimX: ::std::os::raw::c_uint, - gridDimY: ::std::os::raw::c_uint, - gridDimZ: ::std::os::raw::c_uint, - blockDimX: ::std::os::raw::c_uint, - blockDimY: ::std::os::raw::c_uint, - blockDimZ: ::std::os::raw::c_uint, - sharedMemBytes: ::std::os::raw::c_uint, - hStream: CUstream, - kernelParams: *mut *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchCooperativeKernelMultiDevice( - launchParamsList: *mut CUDA_LAUNCH_PARAMS, - numDevices: ::std::os::raw::c_uint, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchHostFunc( - hStream: CUstream, - fn_: CUhostFn, - userData: *mut ::std::os::raw::c_void, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncSetBlockShape( - hfunc: CUfunction, - x: ::std::os::raw::c_int, - y: ::std::os::raw::c_int, - z: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncSetSharedSize( - hfunc: CUfunction, - bytes: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuParamSetSize( - hfunc: CUfunction, - numbytes: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuParamSeti( - hfunc: CUfunction, - offset: ::std::os::raw::c_int, - value: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuParamSetf( - hfunc: CUfunction, - offset: ::std::os::raw::c_int, - value: f32, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuParamSetv( - hfunc: CUfunction, - offset: ::std::os::raw::c_int, - ptr: *mut ::std::os::raw::c_void, - numbytes: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunch(f: CUfunction) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchGrid( - f: CUfunction, - grid_width: ::std::os::raw::c_int, - grid_height: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuLaunchGridAsync( - f: CUfunction, - grid_width: ::std::os::raw::c_int, - grid_height: ::std::os::raw::c_int, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuParamSetTexRef( - hfunc: CUfunction, - texunit: ::std::os::raw::c_int, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphCreate( - phGraph: *mut CUgraph, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddKernelNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, - nodeParams: *const CUDA_KERNEL_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphKernelNodeGetParams( - hNode: CUgraphNode, - nodeParams: *mut CUDA_KERNEL_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphKernelNodeSetParams( - hNode: CUgraphNode, - nodeParams: *const CUDA_KERNEL_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddMemcpyNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, - copyParams: *const CUDA_MEMCPY3D, - ctx: CUcontext, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphMemcpyNodeGetParams( - hNode: CUgraphNode, - nodeParams: *mut CUDA_MEMCPY3D, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphMemcpyNodeSetParams( - hNode: CUgraphNode, - nodeParams: *const CUDA_MEMCPY3D, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddMemsetNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, - memsetParams: *const CUDA_MEMSET_NODE_PARAMS, - ctx: CUcontext, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphMemsetNodeGetParams( - hNode: CUgraphNode, - nodeParams: *mut CUDA_MEMSET_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphMemsetNodeSetParams( - hNode: CUgraphNode, - nodeParams: *const CUDA_MEMSET_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddHostNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, - nodeParams: *const CUDA_HOST_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphHostNodeGetParams( - hNode: CUgraphNode, - nodeParams: *mut CUDA_HOST_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphHostNodeSetParams( - hNode: CUgraphNode, - nodeParams: *const CUDA_HOST_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddChildGraphNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, - childGraph: CUgraph, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphChildGraphNodeGetGraph( - hNode: CUgraphNode, - phGraph: *mut CUgraph, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddEmptyNode( - phGraphNode: *mut CUgraphNode, - hGraph: CUgraph, - dependencies: *const CUgraphNode, - numDependencies: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphClone( - phGraphClone: *mut CUgraph, - originalGraph: CUgraph, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphNodeFindInClone( - phNode: *mut CUgraphNode, - hOriginalNode: CUgraphNode, - hClonedGraph: CUgraph, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphNodeGetType( - hNode: CUgraphNode, - type_: *mut CUgraphNodeType, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphGetNodes( - hGraph: CUgraph, - nodes: *mut CUgraphNode, - numNodes: *mut usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphGetRootNodes( - hGraph: CUgraph, - rootNodes: *mut CUgraphNode, - numRootNodes: *mut usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphGetEdges( - hGraph: CUgraph, - from: *mut CUgraphNode, - to: *mut CUgraphNode, - numEdges: *mut usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphNodeGetDependencies( - hNode: CUgraphNode, - dependencies: *mut CUgraphNode, - numDependencies: *mut usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphNodeGetDependentNodes( - hNode: CUgraphNode, - dependentNodes: *mut CUgraphNode, - numDependentNodes: *mut usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphAddDependencies( - hGraph: CUgraph, - from: *const CUgraphNode, - to: *const CUgraphNode, - numDependencies: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphRemoveDependencies( - hGraph: CUgraph, - from: *const CUgraphNode, - to: *const CUgraphNode, - numDependencies: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphDestroyNode(hNode: CUgraphNode) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphInstantiate_v2( - phGraphExec: *mut CUgraphExec, - hGraph: CUgraph, - phErrorNode: *mut CUgraphNode, - logBuffer: *mut ::std::os::raw::c_char, - bufferSize: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecKernelNodeSetParams( - hGraphExec: CUgraphExec, - hNode: CUgraphNode, - nodeParams: *const CUDA_KERNEL_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecMemcpyNodeSetParams( - hGraphExec: CUgraphExec, - hNode: CUgraphNode, - copyParams: *const CUDA_MEMCPY3D, - ctx: CUcontext, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecMemsetNodeSetParams( - hGraphExec: CUgraphExec, - hNode: CUgraphNode, - memsetParams: *const CUDA_MEMSET_NODE_PARAMS, - ctx: CUcontext, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecHostNodeSetParams( - hGraphExec: CUgraphExec, - hNode: CUgraphNode, - nodeParams: *const CUDA_HOST_NODE_PARAMS, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphLaunch(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecDestroy(hGraphExec: CUgraphExec) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphDestroy(hGraph: CUgraph) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphExecUpdate( - hGraphExec: CUgraphExec, - hGraph: CUgraph, - hErrorNode_out: *mut CUgraphNode, - updateResult_out: *mut CUgraphExecUpdateResult, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphKernelNodeCopyAttributes( - dst: CUgraphNode, - src: CUgraphNode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphKernelNodeGetAttribute( - hNode: CUgraphNode, - attr: CUkernelNodeAttrID, - value_out: *mut CUkernelNodeAttrValue, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphKernelNodeSetAttribute( - hNode: CUgraphNode, - attr: CUkernelNodeAttrID, - value: *const CUkernelNodeAttrValue, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuOccupancyMaxActiveBlocksPerMultiprocessor( - numBlocks: *mut ::std::os::raw::c_int, - func: CUfunction, - blockSize: ::std::os::raw::c_int, - dynamicSMemSize: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - numBlocks: *mut ::std::os::raw::c_int, - func: CUfunction, - blockSize: ::std::os::raw::c_int, - dynamicSMemSize: usize, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuOccupancyMaxPotentialBlockSize( - minGridSize: *mut ::std::os::raw::c_int, - blockSize: *mut ::std::os::raw::c_int, - func: CUfunction, - blockSizeToDynamicSMemSize: CUoccupancyB2DSize, - dynamicSMemSize: usize, - blockSizeLimit: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuOccupancyMaxPotentialBlockSizeWithFlags( - minGridSize: *mut ::std::os::raw::c_int, - blockSize: *mut ::std::os::raw::c_int, - func: CUfunction, - blockSizeToDynamicSMemSize: CUoccupancyB2DSize, - dynamicSMemSize: usize, - blockSizeLimit: ::std::os::raw::c_int, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuOccupancyAvailableDynamicSMemPerBlock( - dynamicSmemSize: *mut usize, - func: CUfunction, - numBlocks: ::std::os::raw::c_int, - blockSize: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetArray( - hTexRef: CUtexref, - hArray: CUarray, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetMipmappedArray( - hTexRef: CUtexref, - hMipmappedArray: CUmipmappedArray, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetAddress_v2( - ByteOffset: *mut usize, - hTexRef: CUtexref, - dptr: CUdeviceptr, - bytes: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetAddress2D_v3( - hTexRef: CUtexref, - desc: *const CUDA_ARRAY_DESCRIPTOR, - dptr: CUdeviceptr, - Pitch: usize, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetFormat( - hTexRef: CUtexref, - fmt: CUarray_format, - NumPackedComponents: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetAddressMode( - hTexRef: CUtexref, - dim: ::std::os::raw::c_int, - am: CUaddress_mode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetMipmapFilterMode( - hTexRef: CUtexref, - fm: CUfilter_mode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetMipmapLevelBias(hTexRef: CUtexref, bias: f32) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetMipmapLevelClamp( - hTexRef: CUtexref, - minMipmapLevelClamp: f32, - maxMipmapLevelClamp: f32, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetMaxAnisotropy( - hTexRef: CUtexref, - maxAniso: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetBorderColor( - hTexRef: CUtexref, - pBorderColor: *mut f32, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefSetFlags( - hTexRef: CUtexref, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetAddress_v2( - pdptr: *mut CUdeviceptr, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetArray(phArray: *mut CUarray, hTexRef: CUtexref) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetMipmappedArray( - phMipmappedArray: *mut CUmipmappedArray, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetAddressMode( - pam: *mut CUaddress_mode, - hTexRef: CUtexref, - dim: ::std::os::raw::c_int, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetFilterMode( - pfm: *mut CUfilter_mode, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetFormat( - pFormat: *mut CUarray_format, - pNumChannels: *mut ::std::os::raw::c_int, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetMipmapFilterMode( - pfm: *mut CUfilter_mode, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetMipmapLevelBias(pbias: *mut f32, hTexRef: CUtexref) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetMipmapLevelClamp( - pminMipmapLevelClamp: *mut f32, - pmaxMipmapLevelClamp: *mut f32, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetMaxAnisotropy( - pmaxAniso: *mut ::std::os::raw::c_int, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetBorderColor( - pBorderColor: *mut f32, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefGetFlags( - pFlags: *mut ::std::os::raw::c_uint, - hTexRef: CUtexref, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefCreate(pTexRef: *mut CUtexref) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexRefDestroy(hTexRef: CUtexref) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSurfRefSetArray( - hSurfRef: CUsurfref, - hArray: CUarray, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSurfRefGetArray(phArray: *mut CUarray, hSurfRef: CUsurfref) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexObjectCreate( - pTexObject: *mut CUtexObject, - pResDesc: *const CUDA_RESOURCE_DESC, - pTexDesc: *const CUDA_TEXTURE_DESC, - pResViewDesc: *const CUDA_RESOURCE_VIEW_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexObjectDestroy(texObject: CUtexObject) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexObjectGetResourceDesc( - pResDesc: *mut CUDA_RESOURCE_DESC, - texObject: CUtexObject, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexObjectGetTextureDesc( - pTexDesc: *mut CUDA_TEXTURE_DESC, - texObject: CUtexObject, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuTexObjectGetResourceViewDesc( - pResViewDesc: *mut CUDA_RESOURCE_VIEW_DESC, - texObject: CUtexObject, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSurfObjectCreate( - pSurfObject: *mut CUsurfObject, - pResDesc: *const CUDA_RESOURCE_DESC, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSurfObjectDestroy(surfObject: CUsurfObject) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuSurfObjectGetResourceDesc( - pResDesc: *mut CUDA_RESOURCE_DESC, - surfObject: CUsurfObject, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceCanAccessPeer( - canAccessPeer: *mut ::std::os::raw::c_int, - dev: CUdevice, - peerDev: CUdevice, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxEnablePeerAccess( - peerContext: CUcontext, - Flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuCtxDisablePeerAccess(peerContext: CUcontext) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuDeviceGetP2PAttribute( - value: *mut ::std::os::raw::c_int, - attrib: CUdevice_P2PAttribute, - srcDevice: CUdevice, - dstDevice: CUdevice, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsUnregisterResource(resource: CUgraphicsResource) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsSubResourceGetMappedArray( - pArray: *mut CUarray, - resource: CUgraphicsResource, - arrayIndex: ::std::os::raw::c_uint, - mipLevel: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsResourceGetMappedMipmappedArray( - pMipmappedArray: *mut CUmipmappedArray, - resource: CUgraphicsResource, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsResourceGetMappedPointer_v2( - pDevPtr: *mut CUdeviceptr, - pSize: *mut usize, - resource: CUgraphicsResource, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsResourceSetMapFlags_v2( - resource: CUgraphicsResource, - flags: ::std::os::raw::c_uint, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsMapResources( - count: ::std::os::raw::c_uint, - resources: *mut CUgraphicsResource, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGraphicsUnmapResources( - count: ::std::os::raw::c_uint, - resources: *mut CUgraphicsResource, - hStream: CUstream, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuGetExportTable( - ppExportTable: *mut *const ::std::os::raw::c_void, - pExportTableId: *const CUuuid, -) -> CUresult { - r#impl::export_table::get(ppExportTable, pExportTableId) -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuFuncGetModule(hmod: *mut CUmodule, hfunc: CUfunction) -> CUresult { - r#impl::unimplemented() -} - -impl CUoutput_mode_enum { - pub const CU_OUT_KEY_VALUE_PAIR: CUoutput_mode_enum = CUoutput_mode_enum(0); -} -impl CUoutput_mode_enum { - pub const CU_OUT_CSV: CUoutput_mode_enum = CUoutput_mode_enum(1); -} -#[repr(transparent)] -#[derive(Copy, Clone, Hash, PartialEq, Eq)] -pub struct CUoutput_mode_enum(pub ::std::os::raw::c_uint); -pub use self::CUoutput_mode_enum as CUoutput_mode; - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuProfilerInitialize( - configFile: *const ::std::os::raw::c_char, - outputFile: *const ::std::os::raw::c_char, - outputMode: CUoutput_mode, -) -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuProfilerStart() -> CUresult { - r#impl::unimplemented() -} - -#[cfg_attr(not(test), no_mangle)] -pub extern "system" fn cuProfilerStop() -> CUresult { - r#impl::unimplemented() -} diff --git a/zluda/src/impl/context.rs b/zluda/src/impl/context.rs index fffceb8..973febc 100644 --- a/zluda/src/impl/context.rs +++ b/zluda/src/impl/context.rs @@ -1,24 +1,93 @@ -use std::ptr; +use super::{driver, FromCuda, ZludaObject}; +use cuda_types::*; +use hip_runtime_sys::*; +use rustc_hash::FxHashSet; +use std::{cell::RefCell, ptr, sync::Mutex}; -use crate::cuda::CUlimit; -use crate::cuda::CUresult; +thread_local! { + pub(crate) static CONTEXT_STACK: RefCell<Vec<(CUcontext, hipDevice_t)>> = RefCell::new(Vec::new()); +} + +pub(crate) struct Context { + pub(crate) device: hipDevice_t, + pub(crate) mutable: Mutex<OwnedByContext>, +} + +pub(crate) struct OwnedByContext { + pub(crate) ref_count: usize, // only used by primary context + pub(crate) _memory: FxHashSet<hipDeviceptr_t>, + pub(crate) _streams: FxHashSet<hipStream_t>, + pub(crate) _modules: FxHashSet<CUmodule>, +} -pub(crate) unsafe fn get_limit(pvalue: *mut usize, limit: CUlimit) -> CUresult { - if pvalue == ptr::null_mut() { - return CUresult::CUDA_ERROR_INVALID_VALUE; +impl ZludaObject for Context { + const COOKIE: usize = 0x5f867c6d9cb73315; + + type CudaHandle = CUcontext; + + fn drop_checked(&mut self) -> CUresult { + Ok(()) } - 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 new(device: hipDevice_t) -> Context { + Context { + device, + mutable: Mutex::new(OwnedByContext { + ref_count: 0, + _memory: FxHashSet::default(), + _streams: FxHashSet::default(), + _modules: FxHashSet::default(), + }), } } -pub(crate) fn set_limit(limit: CUlimit, value: usize) -> CUresult { - if limit == CUlimit::CU_LIMIT_STACK_SIZE { - CUresult::CUDA_SUCCESS +pub(crate) unsafe fn get_limit(pvalue: *mut usize, limit: hipLimit_t) -> hipError_t { + unsafe { hipDeviceGetLimit(pvalue, limit) } +} + +pub(crate) fn set_limit(limit: hipLimit_t, value: usize) -> hipError_t { + unsafe { hipDeviceSetLimit(limit, value) } +} + +pub(crate) fn synchronize() -> hipError_t { + unsafe { hipDeviceSynchronize() } +} + +pub(crate) fn get_primary(hip_dev: hipDevice_t) -> Result<(&'static Context, CUcontext), CUerror> { + let dev = driver::device(hip_dev)?; + Ok(dev.primary_context()) +} + +pub(crate) fn set_current(raw_ctx: CUcontext) -> CUresult { + let new_device = if raw_ctx.0 == ptr::null_mut() { + CONTEXT_STACK.with(|stack| { + let mut stack = stack.borrow_mut(); + if let Some((_, old_device)) = stack.pop() { + if let Some((_, new_device)) = stack.last() { + if old_device != *new_device { + return Some(*new_device); + } + } + } + None + }) } else { - CUresult::CUDA_ERROR_NOT_SUPPORTED + let ctx: &Context = FromCuda::from_cuda(&raw_ctx)?; + let device = ctx.device; + CONTEXT_STACK.with(move |stack| { + let mut stack = stack.borrow_mut(); + let last_device = stack.last().map(|(_, dev)| *dev); + stack.push((raw_ctx, device)); + match last_device { + None => Some(device), + Some(last_device) if last_device != device => Some(device), + _ => None, + } + }) + }; + if let Some(dev) = new_device { + unsafe { hipSetDevice(dev)? }; } + Ok(()) } diff --git a/zluda/src/impl/device.rs b/zluda/src/impl/device.rs index 0c63494..8836c1e 100644 --- a/zluda/src/impl/device.rs +++ b/zluda/src/impl/device.rs @@ -1,29 +1,27 @@ -use super::{transmute_lifetime, transmute_lifetime_mut, CUresult}; -use crate::{ - cuda::{self, CUdevice, CUdevprop}, - hip_call, -}; -use cuda::{CUdevice_attribute, CUuuid_st}; -use hip_runtime_sys::{ - hipDeviceAttribute_t, hipDeviceGetAttribute, hipError_t, hipGetDeviceProperties, -}; -use ocl_core::{ClDeviceIdPtr, ContextProperties, DeviceType}; -use paste::paste; -use std::{ - cmp, - collections::HashSet, - ffi::c_void, - mem, - os::raw::{c_char, c_int, c_uint}, - ptr, - sync::atomic::{AtomicU32, Ordering}, -}; +use cuda_types::*; +use hip_runtime_sys::*; +use std::{mem, ptr}; -const PROJECT_URL_SUFFIX_SHORT: &'static str = " [ZLUDA]"; -const PROJECT_URL_SUFFIX_LONG: &'static str = " [github.com/vosen/ZLUDA]"; +use super::context; + +const PROJECT_SUFFIX: &[u8] = b" [ZLUDA]\0"; +pub const COMPUTE_CAPABILITY_MAJOR: i32 = 8; +pub const COMPUTE_CAPABILITY_MINOR: i32 = 8; + +pub(crate) fn compute_capability(major: &mut i32, minor: &mut i32, _dev: hipDevice_t) -> CUresult { + *major = COMPUTE_CAPABILITY_MAJOR; + *minor = COMPUTE_CAPABILITY_MINOR; + Ok(()) +} + +pub(crate) fn get(device: *mut hipDevice_t, ordinal: i32) -> hipError_t { + unsafe { hipDeviceGet(device, ordinal) } +} #[allow(warnings)] -trait hipDeviceAttribute_t_ext { +trait DeviceAttributeNames { + const hipDeviceAttributeGpuOverlap: hipDeviceAttribute_t = + hipDeviceAttribute_t::hipDeviceAttributeDeviceOverlap; const hipDeviceAttributeMaximumTexture1DWidth: hipDeviceAttribute_t = hipDeviceAttribute_t::hipDeviceAttributeMaxTexture1DWidth; const hipDeviceAttributeMaximumTexture2DWidth: hipDeviceAttribute_t = @@ -42,307 +40,300 @@ trait hipDeviceAttribute_t_ext { hipDeviceAttribute_t::hipDeviceAttributeMaxThreadsPerMultiProcessor; const hipDeviceAttributeAsyncEngineCount: hipDeviceAttribute_t = hipDeviceAttribute_t::hipDeviceAttributeConcurrentKernels; + const hipDeviceAttributePciDomainId: hipDeviceAttribute_t = + hipDeviceAttribute_t::hipDeviceAttributePciDomainID; + const hipDeviceAttributeMultiGpuBoard: hipDeviceAttribute_t = + hipDeviceAttribute_t::hipDeviceAttributeIsMultiGpuBoard; + const hipDeviceAttributeMultiGpuBoardGroupId: hipDeviceAttribute_t = + hipDeviceAttribute_t::hipDeviceAttributeMultiGpuBoardGroupID; + const hipDeviceAttributeMaxSharedMemoryPerBlockOptin: hipDeviceAttribute_t = + hipDeviceAttribute_t::hipDeviceAttributeSharedMemPerBlockOptin; } -impl hipDeviceAttribute_t_ext for hipDeviceAttribute_t {} +impl DeviceAttributeNames for hipDeviceAttribute_t {} macro_rules! remap_attribute { ($attrib:expr => $([ $($word:expr)* ]),*,) => { match $attrib { $( - paste! { CUdevice_attribute:: [< CU_DEVICE_ATTRIBUTE $(_ $word:upper)* >] } => { - paste! { hipDeviceAttribute_t:: [< hipDeviceAttribute $($word:camel)* >] } + paste::paste! { CUdevice_attribute:: [< CU_DEVICE_ATTRIBUTE $(_ $word:upper)* >] } => { + paste::paste! { hipDeviceAttribute_t:: [< hipDeviceAttribute $($word:camel)* >] } } )* - _ => return hipError_t::hipErrorInvalidValue + _ => return Err(hipErrorCode_t::NotSupported) } } } -pub fn get_attribute(pi: *mut i32, attrib: CUdevice_attribute, dev_idx: c_int) -> hipError_t { - if pi == ptr::null_mut() { - return hipError_t::hipErrorInvalidValue; - } - //let mut props = unsafe { mem::zeroed() }; - let hip_attrib = match attrib { - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT => { - unsafe { *pi = 1 }; - return hipError_t::hipSuccess; +pub(crate) fn get_attribute( + pi: &mut i32, + attrib: CUdevice_attribute, + dev_idx: hipDevice_t, +) -> hipError_t { + match attrib { + CUdevice_attribute::CU_DEVICE_ATTRIBUTE_WARP_SIZE => { + *pi = 32; + return Ok(()); } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_GPU_OVERLAP - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_STREAM_PRIORITIES_SUPPORTED - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_GLOBAL_L1_CACHE_SUPPORTED - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_LOCAL_L1_CACHE_SUPPORTED => { - unsafe { *pi = 1 }; - return hipError_t::hipSuccess; - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_TCC_DRIVER - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_HEIGHT - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD_GROUP_ID => { - unsafe { *pi = 0 }; - return hipError_t::hipSuccess; + CUdevice_attribute::CU_DEVICE_ATTRIBUTE_TCC_DRIVER => { + *pi = 0; + return Ok(()); } CUdevice_attribute::CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR => { - unsafe { *pi = 8 }; - return hipError_t::hipSuccess; + *pi = COMPUTE_CAPABILITY_MAJOR; + return Ok(()); } CUdevice_attribute::CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR => { - unsafe { *pi = 0 }; - return hipError_t::hipSuccess; - } - // we assume that arrayed texts have the same limits - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_HEIGHT => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DHeight - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture1DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture1DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_HEIGHT => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DHeight - } - // we treat surface the same as texture - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT => { - hipDeviceAttribute_t::hipDeviceAttributeTextureAlignment - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture1DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DHeight - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture3DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture3DHeight - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture3DDepth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DWidth - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture2DHeight - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH => { - hipDeviceAttribute_t::hipDeviceAttributeMaxTexture1DWidth - } - // Totally made up - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_LAYERS - | CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_LAYERS => { - unsafe { *pi = u16::MAX as i32 }; - return hipError_t::hipSuccess; + *pi = COMPUTE_CAPABILITY_MINOR; + return Ok(()); } - // linear sizes - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH => { - let mut prop = unsafe { mem::zeroed() }; - let err = unsafe { hipGetDeviceProperties(&mut prop, dev_idx) }; - if err != hipError_t::hipSuccess { - return err; - } - unsafe { *pi = prop.maxTexture1DLinear }; - return hipError_t::hipSuccess; - } - CUdevice_attribute::CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID => { - let mut prop = unsafe { mem::zeroed() }; - let err = unsafe { hipGetDeviceProperties(&mut prop, dev_idx) }; - if err != hipError_t::hipSuccess { - return err; - } - unsafe { *pi = prop.pciDomainID }; - return hipError_t::hipSuccess; - } - attrib => remap_attribute! { - attrib => - [MAX THREADS PER BLOCK], - [MAX BLOCK DIM X], - [MAX BLOCK DIM Y], - [MAX BLOCK DIM Z], - [MAX GRID DIM X], - [MAX GRID DIM Y], - [MAX GRID DIM Z], - [MAX SHARED MEMORY PER BLOCK], - [TOTAL CONSTANT MEMORY], - [WARP SIZE], - [MAX PITCH], - [MAX REGISTERS PER BLOCK], - [CLOCK RATE], - [TEXTURE ALIGNMENT], - //[GPU OVERLAP], - [MULTIPROCESSOR COUNT], - [KERNEL EXEC TIMEOUT], - [INTEGRATED], - [CAN MAP HOST MEMORY], - [COMPUTE MODE], - [MAXIMUM TEXTURE1D WIDTH], - [MAXIMUM TEXTURE2D WIDTH], - [MAXIMUM TEXTURE2D HEIGHT], - [MAXIMUM TEXTURE3D WIDTH], - [MAXIMUM TEXTURE3D HEIGHT], - [MAXIMUM TEXTURE3D DEPTH], - //[MAXIMUM TEXTURE2D LAYERED WIDTH], - //[MAXIMUM TEXTURE2D LAYERED HEIGHT], - //[MAXIMUM TEXTURE2D LAYERED LAYERS], - //[MAXIMUM TEXTURE2D ARRAY WIDTH], - //[MAXIMUM TEXTURE2D ARRAY HEIGHT], - //[MAXIMUM TEXTURE2D ARRAY NUMSLICES], - //[SURFACE ALIGNMENT], - [CONCURRENT KERNELS], - [ECC ENABLED], - [PCI BUS ID], - [PCI DEVICE ID], - //[TCC DRIVER], - [MEMORY CLOCK RATE], - [GLOBAL MEMORY BUS WIDTH], - [L2 CACHE SIZE], - [MAX THREADS PER MULTIPROCESSOR], - [ASYNC ENGINE COUNT], - //[UNIFIED ADDRESSING], - //[MAXIMUM TEXTURE1D LAYERED WIDTH], - //[MAXIMUM TEXTURE1D LAYERED LAYERS], - //[CAN TEX2D GATHER], - //[MAXIMUM TEXTURE2D GATHER WIDTH], - //[MAXIMUM TEXTURE2D GATHER HEIGHT], - //[MAXIMUM TEXTURE3D WIDTH ALTERNATE], - //[MAXIMUM TEXTURE3D HEIGHT ALTERNATE], - //[MAXIMUM TEXTURE3D DEPTH ALTERNATE], - //[PCI DOMAIN ID], - [TEXTURE PITCH ALIGNMENT], - //[MAXIMUM TEXTURECUBEMAP WIDTH], - //[MAXIMUM TEXTURECUBEMAP LAYERED WIDTH], - //[MAXIMUM TEXTURECUBEMAP LAYERED LAYERS], - //[MAXIMUM SURFACE1D WIDTH], - //[MAXIMUM SURFACE2D WIDTH], - //[MAXIMUM SURFACE2D HEIGHT], - //[MAXIMUM SURFACE3D WIDTH], - //[MAXIMUM SURFACE3D HEIGHT], - //[MAXIMUM SURFACE3D DEPTH], - //[MAXIMUM SURFACE1D LAYERED WIDTH], - //[MAXIMUM SURFACE1D LAYERED LAYERS], - //[MAXIMUM SURFACE2D LAYERED WIDTH], - //[MAXIMUM SURFACE2D LAYERED HEIGHT], - //[MAXIMUM SURFACE2D LAYERED LAYERS], - //[MAXIMUM SURFACECUBEMAP WIDTH], - //[MAXIMUM SURFACECUBEMAP LAYERED WIDTH], - //[MAXIMUM SURFACECUBEMAP LAYERED LAYERS], - //[MAXIMUM TEXTURE1D LINEAR WIDTH], - //[MAXIMUM TEXTURE2D LINEAR WIDTH], - //[MAXIMUM TEXTURE2D LINEAR HEIGHT], - //[MAXIMUM TEXTURE2D LINEAR PITCH], - //[MAXIMUM TEXTURE2D MIPMAPPED WIDTH], - //[MAXIMUM TEXTURE2D MIPMAPPED HEIGHT], - //[COMPUTE CAPABILITY MAJOR], - //[COMPUTE CAPABILITY MINOR], - //[MAXIMUM TEXTURE1D MIPMAPPED WIDTH], - //[STREAM PRIORITIES SUPPORTED], - //[GLOBAL L1 CACHE SUPPORTED], - //[LOCAL L1 CACHE SUPPORTED], - [MAX SHARED MEMORY PER MULTIPROCESSOR], - //[MAX REGISTERS PER MULTIPROCESSOR], - [MANAGED MEMORY], - //[MULTI GPU BOARD], - //[MULTI GPU BOARD GROUP ID], - //[HOST NATIVE ATOMIC SUPPORTED], - //[SINGLE TO DOUBLE PRECISION PERF RATIO], - [PAGEABLE MEMORY ACCESS], - [CONCURRENT MANAGED ACCESS], - //[COMPUTE PREEMPTION SUPPORTED], - //[CAN USE HOST POINTER FOR REGISTERED MEM], - //[CAN USE STREAM MEM OPS], - //[CAN USE 64 BIT STREAM MEM OPS], - //[CAN USE STREAM WAIT VALUE NOR], - [COOPERATIVE LAUNCH], - [COOPERATIVE MULTI DEVICE LAUNCH], - //[MAX SHARED MEMORY PER BLOCK OPTIN], - //[CAN FLUSH REMOTE WRITES], - //[HOST REGISTER SUPPORTED], - [PAGEABLE MEMORY ACCESS USES HOST PAGE TABLES], - [DIRECT MANAGED MEM ACCESS FROM HOST], - //[VIRTUAL ADDRESS MANAGEMENT SUPPORTED], - //[VIRTUAL MEMORY MANAGEMENT SUPPORTED], - //[HANDLE TYPE POSIX FILE DESCRIPTOR SUPPORTED], - //[HANDLE TYPE WIN32 HANDLE SUPPORTED], - //[HANDLE TYPE WIN32 KMT HANDLE SUPPORTED], - //[MAX BLOCKS PER MULTIPROCESSOR], - //[GENERIC COMPRESSION SUPPORTED], - //[MAX PERSISTING L2 CACHE SIZE], - //[MAX ACCESS POLICY WINDOW SIZE], - //[GPU DIRECT RDMA WITH CUDA VMM SUPPORTED], - //[RESERVED SHARED MEMORY PER BLOCK], - //[SPARSE CUDA ARRAY SUPPORTED], - //[READ ONLY HOST REGISTER SUPPORTED], - //[TIMELINE SEMAPHORE INTEROP SUPPORTED], - //[MEMORY POOLS SUPPORTED], - }, + _ => {} + } + let attrib = remap_attribute! { + attrib => + [MAX THREADS PER BLOCK], + [MAX BLOCK DIM X], + [MAX BLOCK DIM Y], + [MAX BLOCK DIM Z], + [MAX GRID DIM X], + [MAX GRID DIM Y], + [MAX GRID DIM Z], + [MAX SHARED MEMORY PER BLOCK], + [TOTAL CONSTANT MEMORY], + //[WARP SIZE], + [MAX PITCH], + [MAX REGISTERS PER BLOCK], + [CLOCK RATE], + [TEXTURE ALIGNMENT], + [GPU OVERLAP], + [MULTIPROCESSOR COUNT], + [KERNEL EXEC TIMEOUT], + [INTEGRATED], + [CAN MAP HOST MEMORY], + [COMPUTE MODE], + [MAXIMUM TEXTURE1D WIDTH], + [MAXIMUM TEXTURE2D WIDTH], + [MAXIMUM TEXTURE2D HEIGHT], + [MAXIMUM TEXTURE3D WIDTH], + [MAXIMUM TEXTURE3D HEIGHT], + [MAXIMUM TEXTURE3D DEPTH], + //[MAXIMUM TEXTURE2D LAYERED WIDTH], + //[MAXIMUM TEXTURE2D LAYERED HEIGHT], + //[MAXIMUM TEXTURE2D LAYERED LAYERS], + //[MAXIMUM TEXTURE2D ARRAY WIDTH], + //[MAXIMUM TEXTURE2D ARRAY HEIGHT], + //[MAXIMUM TEXTURE2D ARRAY NUMSLICES], + [SURFACE ALIGNMENT], + [CONCURRENT KERNELS], + [ECC ENABLED], + [PCI BUS ID], + [PCI DEVICE ID], + //[TCC DRIVER], + [MEMORY CLOCK RATE], + [GLOBAL MEMORY BUS WIDTH], + [L2 CACHE SIZE], + [MAX THREADS PER MULTIPROCESSOR], + [ASYNC ENGINE COUNT], + [UNIFIED ADDRESSING], + //[MAXIMUM TEXTURE1D LAYERED WIDTH], + //[MAXIMUM TEXTURE1D LAYERED LAYERS], + //[CAN TEX2D GATHER], + //[MAXIMUM TEXTURE2D GATHER WIDTH], + //[MAXIMUM TEXTURE2D GATHER HEIGHT], + //[MAXIMUM TEXTURE3D WIDTH ALTERNATE], + //[MAXIMUM TEXTURE3D HEIGHT ALTERNATE], + //[MAXIMUM TEXTURE3D DEPTH ALTERNATE], + [PCI DOMAIN ID], + [TEXTURE PITCH ALIGNMENT], + //[MAXIMUM TEXTURECUBEMAP WIDTH], + //[MAXIMUM TEXTURECUBEMAP LAYERED WIDTH], + //[MAXIMUM TEXTURECUBEMAP LAYERED LAYERS], + //[MAXIMUM SURFACE1D WIDTH], + //[MAXIMUM SURFACE2D WIDTH], + //[MAXIMUM SURFACE2D HEIGHT], + //[MAXIMUM SURFACE3D WIDTH], + //[MAXIMUM SURFACE3D HEIGHT], + //[MAXIMUM SURFACE3D DEPTH], + //[MAXIMUM SURFACE1D LAYERED WIDTH], + //[MAXIMUM SURFACE1D LAYERED LAYERS], + //[MAXIMUM SURFACE2D LAYERED WIDTH], + //[MAXIMUM SURFACE2D LAYERED HEIGHT], + //[MAXIMUM SURFACE2D LAYERED LAYERS], + //[MAXIMUM SURFACECUBEMAP WIDTH], + //[MAXIMUM SURFACECUBEMAP LAYERED WIDTH], + //[MAXIMUM SURFACECUBEMAP LAYERED LAYERS], + //[MAXIMUM TEXTURE1D LINEAR WIDTH], + //[MAXIMUM TEXTURE2D LINEAR WIDTH], + //[MAXIMUM TEXTURE2D LINEAR HEIGHT], + //[MAXIMUM TEXTURE2D LINEAR PITCH], + //[MAXIMUM TEXTURE2D MIPMAPPED WIDTH], + //[MAXIMUM TEXTURE2D MIPMAPPED HEIGHT], + //[COMPUTE CAPABILITY MAJOR], + //[COMPUTE CAPABILITY MINOR], + //[MAXIMUM TEXTURE1D MIPMAPPED WIDTH], + [STREAM PRIORITIES SUPPORTED], + [GLOBAL L1 CACHE SUPPORTED], + [LOCAL L1 CACHE SUPPORTED], + [MAX SHARED MEMORY PER MULTIPROCESSOR], + [MAX REGISTERS PER MULTIPROCESSOR], + [MANAGED MEMORY], + [MULTI GPU BOARD], + [MULTI GPU BOARD GROUP ID], + [HOST NATIVE ATOMIC SUPPORTED], + [SINGLE TO DOUBLE PRECISION PERF RATIO], + [PAGEABLE MEMORY ACCESS], + [CONCURRENT MANAGED ACCESS], + [COMPUTE PREEMPTION SUPPORTED], + [CAN USE HOST POINTER FOR REGISTERED MEM], + //[CAN USE STREAM MEM OPS], + [COOPERATIVE LAUNCH], + [COOPERATIVE MULTI DEVICE LAUNCH], + [MAX SHARED MEMORY PER BLOCK OPTIN], + //[CAN FLUSH REMOTE WRITES], + [HOST REGISTER SUPPORTED], + [PAGEABLE MEMORY ACCESS USES HOST PAGE TABLES], + [DIRECT MANAGED MEM ACCESS FROM HOST], + //[VIRTUAL ADDRESS MANAGEMENT SUPPORTED], + [VIRTUAL MEMORY MANAGEMENT SUPPORTED], + //[HANDLE TYPE POSIX FILE DESCRIPTOR SUPPORTED], + //[HANDLE TYPE WIN32 HANDLE SUPPORTED], + //[HANDLE TYPE WIN32 KMT HANDLE SUPPORTED], + //[MAX BLOCKS PER MULTIPROCESSOR], + //[GENERIC COMPRESSION SUPPORTED], + //[MAX PERSISTING L2 CACHE SIZE], + //[MAX ACCESS POLICY WINDOW SIZE], + //[GPU DIRECT RDMA WITH CUDA VMM SUPPORTED], + //[RESERVED SHARED MEMORY PER BLOCK], + //[SPARSE CUDA ARRAY SUPPORTED], + //[READ ONLY HOST REGISTER SUPPORTED], + //[TIMELINE SEMAPHORE INTEROP SUPPORTED], + [MEMORY POOLS SUPPORTED], + //[GPU DIRECT RDMA SUPPORTED], + //[GPU DIRECT RDMA FLUSH WRITES OPTIONS], + //[GPU DIRECT RDMA WRITES ORDERING], + //[MEMPOOL SUPPORTED HANDLE TYPES], + //[CLUSTER LAUNCH], + //[DEFERRED MAPPING CUDA ARRAY SUPPORTED], + //[CAN USE 64 BIT STREAM MEM OPS], + //[CAN USE STREAM WAIT VALUE NOR], + //[DMA BUF SUPPORTED], + //[IPC EVENT SUPPORTED], + //[MEM SYNC DOMAIN COUNT], + //[TENSOR MAP ACCESS SUPPORTED], + //[HANDLE TYPE FABRIC SUPPORTED], + //[UNIFIED FUNCTION POINTERS], + //[NUMA CONFIG], + //[NUMA ID], + //[MULTICAST SUPPORTED], + //[MPS ENABLED], + //[HOST NUMA ID], }; - unsafe { hipDeviceGetAttribute(pi, hip_attrib, dev_idx) } + unsafe { hipDeviceGetAttribute(pi, attrib, dev_idx) } +} + +pub(crate) fn get_uuid(uuid: *mut hipUUID, device: hipDevice_t) -> hipError_t { + unsafe { hipDeviceGetUuid(uuid, device) } +} + +pub(crate) fn get_uuid_v2(uuid: *mut hipUUID, device: hipDevice_t) -> hipError_t { + get_uuid(uuid, device) } -pub fn get_uuid(uuid: *mut CUuuid_st, _dev_idx: c_int) -> Result<(), CUresult> { +pub(crate) fn get_luid( + luid: *mut ::core::ffi::c_char, + device_node_mask: &mut ::core::ffi::c_uint, + dev: hipDevice_t, +) -> hipError_t { + let luid = unsafe { + luid.cast::<[i8; 8]>() + .as_mut() + .ok_or(hipErrorCode_t::InvalidValue) + }?; + let mut properties = unsafe { mem::zeroed() }; + unsafe { hipGetDevicePropertiesR0600(&mut properties, dev) }?; + *luid = properties.luid; + *device_node_mask = properties.luidDeviceNodeMask; + Ok(()) +} + +pub(crate) fn get_name( + name: *mut ::core::ffi::c_char, + len: ::core::ffi::c_int, + dev: hipDevice_t, +) -> cuda_types::CUresult { + unsafe { hipDeviceGetName(name, len, dev) }?; + let len = len as usize; + let buffer = unsafe { std::slice::from_raw_parts(name, len) }; + let first_zero = buffer.iter().position(|c| *c == 0); + let first_zero = if let Some(x) = first_zero { + x + } else { + return Ok(()); + }; + if (first_zero + PROJECT_SUFFIX.len()) > len { + return Ok(()); + } unsafe { - *uuid = CUuuid_st { - bytes: mem::zeroed(), - } + ptr::copy_nonoverlapping( + PROJECT_SUFFIX.as_ptr() as _, + name.add(first_zero), + PROJECT_SUFFIX.len(), + ) }; Ok(()) } -// TODO: add support if Level 0 exposes it -pub fn get_luid( - luid: *mut c_char, - dev_node_mask: *mut c_uint, - _dev_idx: c_int, -) -> Result<(), CUresult> { - unsafe { ptr::write_bytes(luid, 0u8, 8) }; - unsafe { *dev_node_mask = 0 }; +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 { + let mut hip_props = unsafe { mem::zeroed() }; + unsafe { hipGetDevicePropertiesR0600(&mut hip_props, dev) }?; + prop.maxThreadsPerBlock = hip_props.maxThreadsPerBlock; + prop.maxThreadsDim = hip_props.maxThreadsDim; + prop.maxGridSize = hip_props.maxGridSize; + prop.totalConstantMemory = clamp_usize(hip_props.totalConstMem); + prop.SIMDWidth = 32; + prop.memPitch = clamp_usize(hip_props.memPitch); + prop.regsPerBlock = hip_props.regsPerBlock; + prop.clockRate = hip_props.clockRate; + prop.textureAlign = clamp_usize(hip_props.textureAlignment); + Ok(()) +} + +pub(crate) fn get_count(count: &mut ::core::ffi::c_int) -> hipError_t { + unsafe { hipGetDeviceCount(count) } +} + +fn clamp_usize(x: usize) -> i32 { + usize::min(x, i32::MAX as usize) as i32 +} + +pub(crate) fn primary_context_retain( + pctx: &mut CUcontext, + hip_dev: hipDevice_t, +) -> Result<(), CUerror> { + let (ctx, raw_ctx) = context::get_primary(hip_dev)?; + { + let mut mutable_ctx = ctx.mutable.lock().map_err(|_| CUerror::UNKNOWN)?; + mutable_ctx.ref_count += 1; + } + *pctx = raw_ctx; 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); +pub(crate) fn primary_context_release(hip_dev: hipDevice_t) -> Result<(), CUerror> { + let (ctx, _) = context::get_primary(hip_dev)?; + { + let mut mutable_ctx = ctx.mutable.lock().map_err(|_| CUerror::UNKNOWN)?; + if mutable_ctx.ref_count == 0 { + return Err(CUerror::INVALID_CONTEXT); + } + mutable_ctx.ref_count -= 1; + if mutable_ctx.ref_count == 0 { + // TODO: drop all children + } } - 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/driver.rs b/zluda/src/impl/driver.rs new file mode 100644 index 0000000..7ff2f54 --- /dev/null +++ b/zluda/src/impl/driver.rs @@ -0,0 +1,79 @@ +use cuda_types::*;
+use hip_runtime_sys::*;
+use std::{
+ ffi::{CStr, CString},
+ mem, slice,
+ sync::OnceLock,
+};
+
+use crate::r#impl::context;
+
+use super::LiveCheck;
+
+pub(crate) struct GlobalState {
+ pub devices: Vec<Device>,
+}
+
+pub(crate) struct Device {
+ pub(crate) _comgr_isa: CString,
+ primary_context: LiveCheck<context::Context>,
+}
+
+impl Device {
+ pub(crate) fn primary_context<'a>(&'a self) -> (&'a context::Context, CUcontext) {
+ unsafe {
+ (
+ self.primary_context.data.assume_init_ref(),
+ self.primary_context.as_handle(),
+ )
+ }
+ }
+}
+
+pub(crate) fn device(dev: i32) -> Result<&'static Device, CUerror> {
+ global_state()?
+ .devices
+ .get(dev as usize)
+ .ok_or(CUerror::INVALID_DEVICE)
+}
+
+pub(crate) fn global_state() -> Result<&'static GlobalState, CUerror> {
+ static GLOBAL_STATE: OnceLock<Result<GlobalState, CUerror>> = OnceLock::new();
+ fn cast_slice<'a>(bytes: &'a [i8]) -> &'a [u8] {
+ unsafe { slice::from_raw_parts(bytes.as_ptr().cast(), bytes.len()) }
+ }
+ GLOBAL_STATE
+ .get_or_init(|| {
+ let mut device_count = 0;
+ unsafe { hipGetDeviceCount(&mut device_count) }?;
+ Ok(GlobalState {
+ devices: (0..device_count)
+ .map(|i| {
+ let mut props = unsafe { mem::zeroed() };
+ unsafe { hipGetDevicePropertiesR0600(&mut props, i) }?;
+ Ok::<_, CUerror>(Device {
+ _comgr_isa: CStr::from_bytes_until_nul(cast_slice(
+ &props.gcnArchName[..],
+ ))
+ .map_err(|_| CUerror::UNKNOWN)?
+ .to_owned(),
+ primary_context: LiveCheck::new(context::new(i)),
+ })
+ })
+ .collect::<Result<Vec<_>, _>>()?,
+ })
+ })
+ .as_ref()
+ .map_err(|e| *e)
+}
+
+pub(crate) fn init(flags: ::core::ffi::c_uint) -> CUresult {
+ unsafe { hipInit(flags) }?;
+ global_state()?;
+ Ok(())
+}
+
+pub(crate) fn get_version(version: &mut ::core::ffi::c_int) -> CUresult {
+ *version = cuda_types::CUDA_VERSION as i32;
+ Ok(())
+}
diff --git a/zluda/src/impl/function.rs b/zluda/src/impl/function.rs index 7f35bb4..8d006ec 100644 --- a/zluda/src/impl/function.rs +++ b/zluda/src/impl/function.rs @@ -1,26 +1,46 @@ -use hip_runtime_sys::{hipError_t, hipFuncAttribute, hipFuncGetAttribute, hipFuncGetAttributes, hipFunction_attribute, hipLaunchKernel, hipModuleLaunchKernel}; - -use super::{CUresult, HasLivenessCookie, LiveCheck}; -use crate::cuda::{CUfunction, CUfunction_attribute, CUstream}; -use ::std::os::raw::{c_uint, c_void}; -use std::{mem, ptr}; +use hip_runtime_sys::*; pub(crate) fn get_attribute( - pi: *mut i32, - cu_attrib: CUfunction_attribute, - func: CUfunction, + pi: &mut i32, + cu_attrib: hipFunction_attribute, + func: hipFunction_t, +) -> hipError_t { + // TODO: implement HIP_FUNC_ATTRIBUTE_PTX_VERSION + // TODO: implement HIP_FUNC_ATTRIBUTE_BINARY_VERSION + unsafe { hipFuncGetAttribute(pi, cu_attrib, func) }?; + if cu_attrib == hipFunction_attribute::HIP_FUNC_ATTRIBUTE_NUM_REGS { + *pi = (*pi).max(1); + } + Ok(()) +} + +pub(crate) fn launch_kernel( + f: hipFunction_t, + grid_dim_x: ::core::ffi::c_uint, + grid_dim_y: ::core::ffi::c_uint, + grid_dim_z: ::core::ffi::c_uint, + block_dim_x: ::core::ffi::c_uint, + block_dim_y: ::core::ffi::c_uint, + block_dim_z: ::core::ffi::c_uint, + shared_mem_bytes: ::core::ffi::c_uint, + stream: hipStream_t, + kernel_params: *mut *mut ::core::ffi::c_void, + extra: *mut *mut ::core::ffi::c_void, ) -> hipError_t { - if pi == ptr::null_mut() || func == ptr::null_mut() { - return hipError_t::hipErrorInvalidValue; + // TODO: fix constants in extra + unsafe { + hipModuleLaunchKernel( + f, + grid_dim_x, + grid_dim_y, + grid_dim_z, + block_dim_x, + block_dim_y, + block_dim_z, + shared_mem_bytes, + stream, + kernel_params, + extra, + ) } - let attrib = match cu_attrib { - CUfunction_attribute::CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK => { - hipFunction_attribute::HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK - } - CUfunction_attribute::CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES => { - hipFunction_attribute::HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES - } - _ => return hipError_t::hipErrorInvalidValue, - }; - unsafe { hipFuncGetAttribute(pi, attrib, func as _) } } diff --git a/zluda/src/impl/link.rs b/zluda/src/impl/link.rs deleted file mode 100644 index d66608f..0000000 --- a/zluda/src/impl/link.rs +++ /dev/null @@ -1,86 +0,0 @@ -use std::{ - ffi::{c_void, CStr}, - mem, ptr, slice, -}; - -use hip_runtime_sys::{hipCtxGetDevice, hipError_t, hipGetDeviceProperties}; - -use crate::{ - cuda::{CUjitInputType, CUjit_option, CUlinkState, CUresult}, - hip_call, -}; - -use super::module::{self, SpirvModule}; - -struct LinkState { - modules: Vec<SpirvModule>, - result: Option<Vec<u8>>, -} - -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(), - result: None, - }); - *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, -) -> Result<(), hipError_t> { - if state == ptr::null_mut() { - return Err(hipError_t::hipErrorInvalidValue); - } - let state: *mut LinkState = mem::transmute(state); - let state = &mut *state; - // V-RAY specific hack - if state.modules.len() == 2 { - return Err(hipError_t::hipSuccess); - } - let spirv_data = SpirvModule::new_raw(data as *const _)?; - state.modules.push(spirv_data); - Ok(()) -} - -pub(crate) unsafe fn complete( - state: CUlinkState, - cubin_out: *mut *mut c_void, - size_out: *mut usize, -) -> Result<(), hipError_t> { - let mut dev = 0; - hip_call! { hipCtxGetDevice(&mut dev) }; - let mut props = unsafe { mem::zeroed() }; - hip_call! { hipGetDeviceProperties(&mut props, dev) }; - let state: &mut LinkState = mem::transmute(state); - let spirv_bins = state.modules.iter().map(|m| &m.binaries[..]); - let should_link_ptx_impl = state.modules.iter().find_map(|m| m.should_link_ptx_impl); - let mut arch_binary = module::compile_amd(&props, spirv_bins, should_link_ptx_impl) - .map_err(|_| hipError_t::hipErrorUnknown)?; - let ptr = arch_binary.as_mut_ptr(); - let size = arch_binary.len(); - state.result = Some(arch_binary); - *cubin_out = ptr as _; - *size_out = size; - Ok(()) -} - -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 index 6041623..3843776 100644 --- a/zluda/src/impl/memory.rs +++ b/zluda/src/impl/memory.rs @@ -1,55 +1,35 @@ -use hip_runtime_sys::{ - hipDrvMemcpy3D, hipError_t, hipMemcpy3D, hipMemcpy3DParms, hipMemoryType, hipPitchedPtr, - hipPos, HIP_MEMCPY3D, -}; -use std::ptr; +use hip_runtime_sys::*; -use crate::{ - cuda::{CUDA_MEMCPY3D_st, CUdeviceptr, CUmemorytype, CUresult}, - hip_call, -}; +pub(crate) fn alloc_v2(dptr: *mut hipDeviceptr_t, bytesize: usize) -> hipError_t { + unsafe { hipMalloc(dptr.cast(), bytesize) }?; + // TODO: parametrize for non-Geekbench + unsafe { hipMemsetD8(*dptr, 0, bytesize) } +} + +pub(crate) fn free_v2(dptr: hipDeviceptr_t) -> hipError_t { + unsafe { hipFree(dptr.0) } +} + +pub(crate) fn copy_dto_h_v2( + dst_host: *mut ::core::ffi::c_void, + src_device: hipDeviceptr_t, + byte_count: usize, +) -> hipError_t { + unsafe { hipMemcpyDtoH(dst_host, src_device, byte_count) } +} -// 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 copy_hto_d_v2( + dst_device: hipDeviceptr_t, + src_host: *const ::core::ffi::c_void, + byte_count: usize, +) -> hipError_t { + unsafe { hipMemcpyHtoD(dst_device, src_host.cast_mut(), byte_count) } } -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), - } +pub(crate) fn get_address_range_v2( + pbase: *mut hipDeviceptr_t, + psize: *mut usize, + dptr: hipDeviceptr_t, +) -> hipError_t { + unsafe { hipMemGetAddressRange(pbase, psize, dptr) } } diff --git a/zluda/src/impl/mod.rs b/zluda/src/impl/mod.rs index 1335ef6..766b4a5 100644 --- a/zluda/src/impl/mod.rs +++ b/zluda/src/impl/mod.rs @@ -1,230 +1,209 @@ -use hip_runtime_sys::hipError_t; - -use crate::cuda::{CUctx_st, CUdevice, CUdeviceptr, CUfunc_st, CUmod_st, CUresult, CUstream_st}; -use std::{ - ffi::c_void, - mem::{self, ManuallyDrop}, - os::raw::c_int, - ptr, - sync::Mutex, - sync::TryLockError, -}; - -#[cfg(test)] -#[macro_use] -pub mod test; -pub mod device; -pub mod export_table; -pub mod function; -#[cfg_attr(windows, path = "os_win.rs")] -#[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; +use cuda_types::*; +use hip_runtime_sys::*; +use std::mem::{self, ManuallyDrop, MaybeUninit}; + +pub(super) mod context; +pub(super) mod device; +pub(super) mod driver; +pub(super) mod function; +pub(super) mod memory; +pub(super) mod module; +pub(super) mod pointer; #[cfg(debug_assertions)] -pub fn unimplemented() -> CUresult { +pub(crate) fn unimplemented() -> CUresult { unimplemented!() } #[cfg(not(debug_assertions))] -pub fn unimplemented() -> CUresult { - CUresult::CUDA_ERROR_NOT_SUPPORTED +pub(crate) fn unimplemented() -> CUresult { + CUresult::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(crate) trait FromCuda<'a, T>: Sized { + fn from_cuda(t: &'a T) -> Result<Self, CUerror>; +} + +macro_rules! from_cuda_nop { + ($($type_:ty),*) => { + $( + impl<'a> FromCuda<'a, $type_> for $type_ { + fn from_cuda(x: &'a $type_) -> Result<Self, CUerror> { + Ok(*x) + } } - } + + impl<'a> FromCuda<'a, *mut $type_> for &'a mut $type_ { + fn from_cuda(x: &'a *mut $type_) -> Result<Self, CUerror> { + match unsafe { x.as_mut() } { + Some(x) => Ok(x), + None => Err(CUerror::INVALID_VALUE), + } + } + } + )* + }; +} + +macro_rules! from_cuda_transmute { + ($($from:ty => $to:ty),*) => { + $( + impl<'a> FromCuda<'a, $from> for $to { + fn from_cuda(x: &'a $from) -> Result<Self, CUerror> { + Ok(unsafe { std::mem::transmute(*x) }) + } + } + + impl<'a> FromCuda<'a, *mut $from> for &'a mut $to { + fn from_cuda(x: &'a *mut $from) -> Result<Self, CUerror> { + match unsafe { x.cast::<$to>().as_mut() } { + Some(x) => Ok(x), + None => Err(CUerror::INVALID_VALUE), + } + } + } + + impl<'a> FromCuda<'a, *mut $from> for * mut $to { + fn from_cuda(x: &'a *mut $from) -> Result<Self, CUerror> { + Ok(x.cast::<$to>()) + } + } + )* + }; +} + +macro_rules! from_cuda_object { + ($($type_:ty),*) => { + $( + impl<'a> FromCuda<'a, <$type_ as ZludaObject>::CudaHandle> for <$type_ as ZludaObject>::CudaHandle { + fn from_cuda(handle: &'a <$type_ as ZludaObject>::CudaHandle) -> Result<<$type_ as ZludaObject>::CudaHandle, CUerror> { + Ok(*handle) + } + } + + impl<'a> FromCuda<'a, *mut <$type_ as ZludaObject>::CudaHandle> for &'a mut <$type_ as ZludaObject>::CudaHandle { + fn from_cuda(handle: &'a *mut <$type_ as ZludaObject>::CudaHandle) -> Result<&'a mut <$type_ as ZludaObject>::CudaHandle, CUerror> { + match unsafe { handle.as_mut() } { + Some(x) => Ok(x), + None => Err(CUerror::INVALID_VALUE), + } + } + } + + impl<'a> FromCuda<'a, <$type_ as ZludaObject>::CudaHandle> for &'a $type_ { + fn from_cuda(handle: &'a <$type_ as ZludaObject>::CudaHandle) -> Result<&'a $type_, CUerror> { + Ok(as_ref(handle).as_result()?) + } + } + )* }; } -pub trait HasLivenessCookie: Sized { +from_cuda_nop!( + *mut i8, + *mut i32, + *mut usize, + *const ::core::ffi::c_void, + *const ::core::ffi::c_char, + *mut ::core::ffi::c_void, + *mut *mut ::core::ffi::c_void, + i32, + u32, + usize, + cuda_types::CUdevprop, + CUdevice_attribute +); +from_cuda_transmute!( + CUuuid => hipUUID, + CUfunction => hipFunction_t, + CUfunction_attribute => hipFunction_attribute, + CUstream => hipStream_t, + CUpointer_attribute => hipPointer_attribute, + CUdeviceptr_v2 => hipDeviceptr_t +); +from_cuda_object!(module::Module, context::Context); + +impl<'a> FromCuda<'a, CUlimit> for hipLimit_t { + fn from_cuda(limit: &'a CUlimit) -> Result<Self, CUerror> { + Ok(match *limit { + CUlimit::CU_LIMIT_STACK_SIZE => hipLimit_t::hipLimitStackSize, + CUlimit::CU_LIMIT_PRINTF_FIFO_SIZE => hipLimit_t::hipLimitPrintfFifoSize, + CUlimit::CU_LIMIT_MALLOC_HEAP_SIZE => hipLimit_t::hipLimitMallocHeapSize, + _ => return Err(CUerror::NOT_SUPPORTED), + }) + } +} + +pub(crate) trait ZludaObject: Sized + Send + Sync { const COOKIE: usize; - const LIVENESS_FAIL: CUresult; + const LIVENESS_FAIL: CUerror = cuda_types::CUerror::INVALID_VALUE; - fn try_drop(&mut self) -> Result<(), CUresult>; + type CudaHandle: Sized; + + fn drop_checked(&mut self) -> CUresult; + + fn wrap(self) -> Self::CudaHandle { + unsafe { mem::transmute_copy(&LiveCheck::wrap(self)) } + } } -// This struct is a best-effort check if wrapped value has been dropped, -// while it's inherently safe, its use coming from FFI is very unsafe #[repr(C)] -pub struct LiveCheck<T: HasLivenessCookie> { +pub(crate) struct LiveCheck<T: ZludaObject> { cookie: usize, - data: ManuallyDrop<T>, + data: MaybeUninit<T>, } -impl<T: HasLivenessCookie> LiveCheck<T> { - pub fn new(data: T) -> Self { +impl<T: ZludaObject> LiveCheck<T> { + fn new(data: T) -> Self { LiveCheck { cookie: T::COOKIE, - data: ManuallyDrop::new(data), + data: MaybeUninit::new(data), } } - fn destroy_impl(this: *mut Self) -> Result<(), CUresult> { - let mut ctx_box = ManuallyDrop::new(unsafe { Box::from_raw(this) }); - ctx_box.try_drop()?; - unsafe { ManuallyDrop::drop(&mut ctx_box) }; - Ok(()) + fn as_handle(&self) -> T::CudaHandle { + unsafe { mem::transmute_copy(&self) } } - unsafe fn ptr_from_inner(this: *mut T) -> *mut Self { - let outer_ptr = (this as *mut u8).sub(mem::size_of::<usize>()); - outer_ptr as *mut Self + fn wrap(data: T) -> *mut Self { + Box::into_raw(Box::new(Self::new(data))) } - pub unsafe fn as_ref_unchecked(&self) -> &T { - &self.data - } - - pub fn as_option_mut(&mut self) -> Option<&mut T> { + fn as_result(&self) -> Result<&T, CUerror> { if self.cookie == T::COOKIE { - Some(&mut self.data) - } else { - None - } - } - - pub fn as_result(&self) -> Result<&T, CUresult> { - if self.cookie == T::COOKIE { - Ok(&self.data) - } else { - Err(T::LIVENESS_FAIL) - } - } - - pub fn as_result_mut(&mut self) -> Result<&mut T, CUresult> { - if self.cookie == T::COOKIE { - Ok(&mut self.data) + Ok(unsafe { self.data.assume_init_ref() }) } else { Err(T::LIVENESS_FAIL) } } + // This looks like nonsense, but it's not. There are two cases: + // Err(CUerror) -> meaning that the object is invalid, this pointer does not point into valid memory + // Ok(maybe_error) -> meaning that the object is valid, we dropped everything, but there *might* + // an error in the underlying runtime that we want to propagate #[must_use] - pub fn try_drop(&mut self) -> Result<(), CUresult> { + fn drop_checked(&mut self) -> Result<Result<(), CUerror>, CUerror> { if self.cookie == T::COOKIE { self.cookie = 0; - self.data.try_drop()?; - unsafe { ManuallyDrop::drop(&mut self.data) }; - return Ok(()); - } - Err(T::LIVENESS_FAIL) - } -} - -impl<T: HasLivenessCookie> Drop for LiveCheck<T> { - fn drop(&mut self) { - self.cookie = 0; - } -} - -pub trait CudaRepr: Sized { - type Impl: Sized; -} - -impl<T: CudaRepr> CudaRepr for *mut T { - type Impl = *mut T::Impl; -} - -pub trait Decuda<To> { - fn decuda(self: Self) -> To; -} - -impl<T: CudaRepr> Decuda<*mut T::Impl> for *mut T { - fn decuda(self: Self) -> *mut T::Impl { - self as *mut _ - } -} - -impl<T> From<TryLockError<T>> for CUresult { - fn from(_: TryLockError<T>) -> Self { - CUresult::CUDA_ERROR_ILLEGAL_STATE - } -} - -impl From<ocl_core::Error> for CUresult { - fn from(result: ocl_core::Error) -> Self { - match result { - _ => CUresult::CUDA_ERROR_UNKNOWN, - } - } -} - -impl From<hip_runtime_sys::hipError_t> for CUresult { - fn from(result: hip_runtime_sys::hipError_t) -> Self { - match result { - hip_runtime_sys::hipError_t::hipErrorRuntimeMemory - | hip_runtime_sys::hipError_t::hipErrorRuntimeOther => CUresult::CUDA_ERROR_UNKNOWN, - hip_runtime_sys::hipError_t(e) => CUresult(e), - } - } -} - -pub trait Encuda { - type To: Sized; - fn encuda(self: Self) -> Self::To; -} - -impl Encuda for CUresult { - type To = CUresult; - fn encuda(self: Self) -> Self::To { - self - } -} - -impl Encuda for () { - type To = CUresult; - fn encuda(self: Self) -> Self::To { - CUresult::CUDA_SUCCESS - } -} - -impl<T1: Encuda<To = CUresult>, T2: Encuda<To = CUresult>> Encuda for Result<T1, T2> { - type To = CUresult; - fn encuda(self: Self) -> Self::To { - match self { - Ok(e) => e.encuda(), - Err(e) => e.encuda(), + let result = unsafe { self.data.assume_init_mut().drop_checked() }; + unsafe { MaybeUninit::assume_init_drop(&mut self.data) }; + Ok(result) + } else { + Err(T::LIVENESS_FAIL) } } } -impl Encuda for hipError_t { - type To = CUresult; - fn encuda(self: Self) -> Self::To { - self.into() - } -} - -unsafe fn transmute_lifetime<'a, 'b, T: ?Sized>(t: &'a T) -> &'b T { - mem::transmute(t) -} - -unsafe fn transmute_lifetime_mut<'a, 'b, T: ?Sized>(t: &'a mut T) -> &'b mut T { - mem::transmute(t) +pub fn as_ref<'a, T: ZludaObject>( + handle: &'a T::CudaHandle, +) -> &'a ManuallyDrop<Box<LiveCheck<T>>> { + unsafe { mem::transmute(handle) } } -pub fn driver_get_version() -> c_int { - i32::max_value() -} - -impl<'a> CudaRepr for CUdeviceptr { - type Impl = *mut c_void; -} - -impl Decuda<*mut c_void> for CUdeviceptr { - fn decuda(self) -> *mut c_void { - self.0 as *mut _ - } +pub fn drop_checked<T: ZludaObject>(handle: T::CudaHandle) -> Result<(), CUerror> { + let mut wrapped_object: ManuallyDrop<Box<LiveCheck<T>>> = + unsafe { mem::transmute_copy(&handle) }; + let underlying_error = LiveCheck::drop_checked(&mut wrapped_object)?; + unsafe { ManuallyDrop::drop(&mut wrapped_object) }; + underlying_error } diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs index 24fa88a..8b19c1b 100644 --- a/zluda/src/impl/module.rs +++ b/zluda/src/impl/module.rs @@ -1,261 +1,53 @@ -use std::borrow::Cow; -use std::collections::HashMap; -use std::ffi::{CStr, CString}; -use std::fs::File; -use std::io::{self, Read, Write}; -use std::ops::Add; -use std::os::raw::c_char; -use std::path::{Path, PathBuf}; -use std::process::Command; -use std::{env, fs, iter, mem, ptr, slice}; +use super::ZludaObject; +use cuda_types::*; +use hip_runtime_sys::*; +use std::{ffi::CStr, mem}; -use hip_runtime_sys::{ - hipCtxGetCurrent, hipCtxGetDevice, hipDeviceGetAttribute, hipDeviceGetName, hipDeviceProp_t, - hipError_t, hipGetDeviceProperties, hipGetStreamDeviceId, hipModuleLoadData, -}; -use tempfile::NamedTempFile; - -use crate::cuda::CUmodule; -use crate::hip_call; - -pub struct SpirvModule { - pub binaries: Vec<u32>, - pub kernel_info: HashMap<String, ptx::KernelInfo>, - pub should_link_ptx_impl: Option<(&'static [u8], &'static [u8])>, - pub build_options: CString, +pub(crate) struct Module { + base: hipModule_t, } -impl SpirvModule { - pub fn new_raw<'a>(text: *const c_char) -> Result<Self, hipError_t> { - let u8_text = unsafe { CStr::from_ptr(text) }; - let ptx_text = u8_text - .to_str() - .map_err(|_| hipError_t::hipErrorInvalidImage)?; - Self::new(ptx_text) - } +impl ZludaObject for Module { + const COOKIE: usize = 0xe9138bd040487d4a; - pub fn new<'a>(ptx_text: &str) -> Result<Self, hipError_t> { - let mut errors = Vec::new(); - let ast = ptx::ModuleParser::new() - .parse(&mut errors, ptx_text) - .map_err(|_| hipError_t::hipErrorInvalidImage)?; - if errors.len() > 0 { - return Err(hipError_t::hipErrorInvalidImage); - } - let spirv_module = - ptx::to_spirv_module(ast).map_err(|_| hipError_t::hipErrorInvalidImage)?; - Ok(SpirvModule { - binaries: spirv_module.assemble(), - kernel_info: spirv_module.kernel_info, - should_link_ptx_impl: spirv_module.should_link_ptx_impl, - build_options: spirv_module.build_options, - }) - } -} + type CudaHandle = CUmodule; -pub(crate) fn load(module: *mut CUmodule, fname: *const i8) -> Result<(), hipError_t> { - let file_name = unsafe { CStr::from_ptr(fname) } - .to_str() - .map_err(|_| hipError_t::hipErrorInvalidValue)?; - let mut file = File::open(file_name).map_err(|_| hipError_t::hipErrorFileNotFound)?; - let mut file_buffer = Vec::new(); - file.read_to_end(&mut file_buffer) - .map_err(|_| hipError_t::hipErrorUnknown)?; - let result = load_data(module, file_buffer.as_ptr() as _); - drop(file_buffer); - result -} - -pub(crate) fn load_data( - module: *mut CUmodule, - image: *const std::ffi::c_void, -) -> Result<(), hipError_t> { - if image == ptr::null() { - return Err(hipError_t::hipErrorInvalidValue); - } - if unsafe { *(image as *const u32) } == 0x464c457f { - return match unsafe { hipModuleLoadData(module as _, image) } { - hipError_t::hipSuccess => Ok(()), - e => Err(e), - }; + fn drop_checked(&mut self) -> CUresult { + unsafe { hipModuleUnload(self.base) }?; + Ok(()) } - let spirv_data = SpirvModule::new_raw(image as *const _)?; - load_data_impl(module, spirv_data) } -pub fn load_data_impl(pmod: *mut CUmodule, spirv_data: SpirvModule) -> Result<(), hipError_t> { +pub(crate) fn load_data(module: &mut CUmodule, image: *const std::ffi::c_void) -> CUresult { + let text = unsafe { CStr::from_ptr(image.cast()) } + .to_str() + .map_err(|_| CUerror::INVALID_VALUE)?; + let ast = ptx_parser::parse_module_checked(text).map_err(|_| CUerror::NO_BINARY_FOR_GPU)?; + let llvm_module = ptx::to_llvm_module(ast).map_err(|_| CUerror::UNKNOWN)?; let mut dev = 0; - hip_call! { hipCtxGetDevice(&mut dev) }; + unsafe { hipCtxGetDevice(&mut dev) }?; let mut props = unsafe { mem::zeroed() }; - hip_call! { hipGetDeviceProperties(&mut props, dev) }; - let arch_binary = compile_amd( - &props, - iter::once(&spirv_data.binaries[..]), - spirv_data.should_link_ptx_impl, + unsafe { hipGetDevicePropertiesR0600(&mut props, dev) }?; + let elf_module = comgr::compile_bitcode( + unsafe { CStr::from_ptr(props.gcnArchName.as_ptr()) }, + &*llvm_module.llvm_ir, + llvm_module.linked_bitcode(), ) - .map_err(|_| hipError_t::hipErrorUnknown)?; - hip_call! { hipModuleLoadData(pmod as _, arch_binary.as_ptr() as _) }; + .map_err(|_| CUerror::UNKNOWN)?; + let mut hip_module = unsafe { mem::zeroed() }; + unsafe { hipModuleLoadData(&mut hip_module, elf_module.as_ptr().cast()) }?; + *module = Module { base: hip_module }.wrap(); Ok(()) } -const LLVM_SPIRV: &'static str = "/home/vosen/amd/llvm-project/build/bin/llvm-spirv"; -const AMDGPU: &'static str = "/opt/rocm/"; -const AMDGPU_TARGET: &'static str = "amdgcn-amd-amdhsa"; -const AMDGPU_BITCODE: [&'static str; 8] = [ - "opencl.bc", - "ocml.bc", - "ockl.bc", - "oclc_correctly_rounded_sqrt_off.bc", - "oclc_daz_opt_on.bc", - "oclc_finite_only_off.bc", - "oclc_unsafe_math_off.bc", - "oclc_wavefrontsize64_off.bc", -]; -const AMDGPU_BITCODE_DEVICE_PREFIX: &'static str = "oclc_isa_version_"; - -pub(crate) fn compile_amd<'a>( - device_pros: &hipDeviceProp_t, - spirv_il: impl Iterator<Item = &'a [u32]>, - ptx_lib: Option<(&'static [u8], &'static [u8])>, -) -> io::Result<Vec<u8>> { - let null_terminator = device_pros - .gcnArchName - .iter() - .position(|&x| x == 0) - .unwrap(); - let gcn_arch_slice = unsafe { - slice::from_raw_parts(device_pros.gcnArchName.as_ptr() as _, null_terminator + 1) - }; - let device_name = - if let Ok(Ok(name)) = CStr::from_bytes_with_nul(gcn_arch_slice).map(|x| x.to_str()) { - name - } else { - return Err(io::Error::new(io::ErrorKind::Other, "")); - }; - let dir = tempfile::tempdir()?; - let llvm_spirv_path = match env::var("LLVM_SPIRV") { - Ok(path) => Cow::Owned(path), - Err(_) => Cow::Borrowed(LLVM_SPIRV), - }; - let llvm_files = spirv_il - .map(|spirv| { - let mut spirv_file = NamedTempFile::new_in(&dir)?; - let spirv_u8 = unsafe { - slice::from_raw_parts( - spirv.as_ptr() as *const u8, - spirv.len() * mem::size_of::<u32>(), - ) - }; - spirv_file.write_all(spirv_u8)?; - if cfg!(debug_assertions) { - persist_file(spirv_file.path())?; - } - let llvm = NamedTempFile::new_in(&dir)?; - let to_llvm_cmd = Command::new(&*llvm_spirv_path) - //.arg("--spirv-debug") - .arg("-r") - .arg("-o") - .arg(llvm.path()) - .arg(spirv_file.path()) - .status()?; - assert!(to_llvm_cmd.success()); - if cfg!(debug_assertions) { - persist_file(llvm.path())?; - } - Ok::<_, io::Error>(llvm) - }) - .collect::<Result<Vec<_>, _>>()?; - let linked_binary = NamedTempFile::new_in(&dir)?; - let mut llvm_link = PathBuf::from(AMDGPU); - llvm_link.push("llvm"); - llvm_link.push("bin"); - llvm_link.push("llvm-link"); - let mut linker_cmd = Command::new(&llvm_link); - linker_cmd - .arg("-o") - .arg(linked_binary.path()) - .args(llvm_files.iter().map(|f| f.path())) - .args(get_bitcode_paths(device_name)); - if cfg!(debug_assertions) { - linker_cmd.arg("-v"); - } - let status = linker_cmd.status()?; - assert!(status.success()); - if cfg!(debug_assertions) { - persist_file(linked_binary.path())?; - } - let mut ptx_lib_bitcode = NamedTempFile::new_in(&dir)?; - let compiled_binary = NamedTempFile::new_in(&dir)?; - let mut clang_exe = PathBuf::from(AMDGPU); - clang_exe.push("llvm"); - clang_exe.push("bin"); - clang_exe.push("clang"); - let mut compiler_cmd = Command::new(&clang_exe); - compiler_cmd - .arg(format!("-mcpu={}", device_name)) - .arg("-ffp-contract=off") - .arg("-nogpulib") - .arg("-mno-wavefrontsize64") - .arg("-O3") - .arg("-Xclang") - .arg("-O3") - .arg("-Xlinker") - .arg("--no-undefined") - .arg("-target") - .arg(AMDGPU_TARGET) - .arg("-o") - .arg(compiled_binary.path()) - .arg("-x") - .arg("ir") - .arg(linked_binary.path()); - if let Some((_, bitcode)) = ptx_lib { - ptx_lib_bitcode.write_all(bitcode)?; - compiler_cmd.arg(ptx_lib_bitcode.path()); - }; - if cfg!(debug_assertions) { - compiler_cmd.arg("-v"); - } - let status = compiler_cmd.status()?; - assert!(status.success()); - let mut result = Vec::new(); - let compiled_bin_path = compiled_binary.path(); - let mut compiled_binary = File::open(compiled_bin_path)?; - compiled_binary.read_to_end(&mut result)?; - if cfg!(debug_assertions) { - persist_file(compiled_bin_path)?; - } - Ok(result) -} - -fn persist_file(path: &Path) -> io::Result<()> { - let mut persistent = PathBuf::from("/tmp/zluda"); - std::fs::create_dir_all(&persistent)?; - persistent.push(path.file_name().unwrap()); - std::fs::copy(path, persistent)?; - Ok(()) +pub(crate) fn unload(hmod: CUmodule) -> CUresult { + super::drop_checked::<Module>(hmod) } -fn get_bitcode_paths(device_name: &str) -> impl Iterator<Item = PathBuf> { - let generic_paths = AMDGPU_BITCODE.iter().map(|x| { - let mut path = PathBuf::from(AMDGPU); - path.push("amdgcn"); - path.push("bitcode"); - path.push(x); - path - }); - let suffix = if let Some(suffix_idx) = device_name.find(':') { - suffix_idx - } else { - device_name.len() - }; - let mut additional_path = PathBuf::from(AMDGPU); - additional_path.push("amdgcn"); - additional_path.push("bitcode"); - additional_path.push(format!( - "{}{}{}", - AMDGPU_BITCODE_DEVICE_PREFIX, - &device_name[3..suffix], - ".bc" - )); - generic_paths.chain(std::iter::once(additional_path)) +pub(crate) fn get_function( + hfunc: &mut hipFunction_t, + hmod: &Module, + name: *const ::core::ffi::c_char, +) -> hipError_t { + unsafe { hipModuleGetFunction(hfunc, hmod.base, name) } } diff --git a/zluda/src/impl/pointer.rs b/zluda/src/impl/pointer.rs index 2b925cd..6b458a0 100644 --- a/zluda/src/impl/pointer.rs +++ b/zluda/src/impl/pointer.rs @@ -1,53 +1,40 @@ -use std::{ffi::c_void, mem, ptr}; - -use hip_runtime_sys::{hipError_t, hipMemoryType, hipPointerGetAttributes}; - -use crate::{ - cuda::{CUdeviceptr, CUmemorytype, CUpointer_attribute}, - hip_call, -}; +use cuda_types::*; +use hip_runtime_sys::*; +use std::{ffi::c_void, ptr}; pub(crate) unsafe fn get_attribute( data: *mut c_void, - attribute: CUpointer_attribute, - ptr: CUdeviceptr, -) -> Result<(), hipError_t> { + attribute: hipPointer_attribute, + ptr: hipDeviceptr_t, +) -> hipError_t { if data == ptr::null_mut() { - return Err(hipError_t::hipErrorInvalidValue); + return hipError_t::ErrorInvalidValue; } - 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; + // TODO: implement by getting device ordinal & allocation start, + // then go through every context for that device + hipPointer_attribute::HIP_POINTER_ATTRIBUTE_CONTEXT => hipError_t::ErrorNotSupported, + hipPointer_attribute::HIP_POINTER_ATTRIBUTE_MEMORY_TYPE => { + let mut hip_result = hipMemoryType(0); + hipPointerGetAttribute( + (&mut hip_result as *mut hipMemoryType).cast::<c_void>(), + attribute, + ptr, + )?; + let cuda_result = memory_type(hip_result)?; + unsafe { *(data.cast()) = cuda_result }; Ok(()) } - _ => Err(hipError_t::hipErrorNotSupported), + _ => unsafe { hipPointerGetAttribute(data, attribute, ptr) }, } } -pub(crate) fn memory_type(cu: hipMemoryType) -> Result<CUmemorytype, hipError_t> { +fn memory_type(cu: hipMemoryType) -> Result<CUmemorytype, hipErrorCode_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), + _ => Err(hipErrorCode_t::InvalidValue), } } diff --git a/zluda/src/impl/test.rs b/zluda/src/impl/test.rs deleted file mode 100644 index b36ccd8..0000000 --- a/zluda/src/impl/test.rs +++ /dev/null @@ -1,157 +0,0 @@ -#![allow(non_snake_case)] - -use crate::cuda as zluda; -use crate::cuda::CUstream; -use crate::cuda::CUuuid; -use crate::{ - cuda::{CUdevice, CUdeviceptr}, - r#impl::CUresult, -}; -use ::std::{ - ffi::c_void, - os::raw::{c_int, c_uint}, -}; -use cuda_driver_sys as cuda; - -#[macro_export] -macro_rules! cuda_driver_test { - ($func:ident) => { - paste! { - #[test] - fn [<$func _zluda>]() { - $func::<crate::r#impl::test::Zluda>() - } - - #[test] - fn [<$func _cuda>]() { - $func::<crate::r#impl::test::Cuda>() - } - } - }; -} - -pub trait CudaDriverFns { - fn cuInit(flags: c_uint) -> CUresult; - fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult; - fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult; - fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult; - fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult; - fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult; - fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult; - fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult; - fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult; - fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult; - fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult; - fn cuMemFree_v2(mem: *mut c_void) -> CUresult; - fn cuStreamDestroy_v2(stream: CUstream) -> CUresult; -} - -pub struct Zluda(); - -impl CudaDriverFns for Zluda { - fn cuInit(_flags: c_uint) -> CUresult { - zluda::cuInit(_flags as _) - } - - fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult { - zluda::cuCtxCreate_v2(pctx as *mut _, flags, CUdevice(dev)) - } - - fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult { - zluda::cuCtxDestroy_v2(ctx as *mut _) - } - - fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult { - zluda::cuCtxPopCurrent_v2(pctx as *mut _) - } - - fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult { - zluda::cuCtxGetApiVersion(ctx as *mut _, version) - } - - fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult { - zluda::cuCtxGetCurrent(pctx as *mut _) - } - fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult { - zluda::cuMemAlloc_v2(dptr as *mut _, bytesize) - } - - fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult { - zluda::cuDeviceGetUuid(uuid, CUdevice(dev)) - } - - fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult { - zluda::cuDevicePrimaryCtxGetState(CUdevice(dev), flags, active) - } - - fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult { - zluda::cuStreamGetCtx(hStream, pctx as _) - } - - fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult { - zluda::cuStreamCreate(stream, flags) - } - - fn cuMemFree_v2(dptr: *mut c_void) -> CUresult { - zluda::cuMemFree_v2(CUdeviceptr(dptr as _)) - } - - fn cuStreamDestroy_v2(stream: CUstream) -> CUresult { - zluda::cuStreamDestroy_v2(stream) - } -} - -pub struct Cuda(); - -impl CudaDriverFns for Cuda { - fn cuInit(flags: c_uint) -> CUresult { - unsafe { CUresult(cuda::cuInit(flags) as c_uint) } - } - - fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult { - unsafe { CUresult(cuda::cuCtxCreate_v2(pctx as *mut _, flags, dev) as c_uint) } - } - - fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult { - unsafe { CUresult(cuda::cuCtxDestroy_v2(ctx as *mut _) as c_uint) } - } - - fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult { - unsafe { CUresult(cuda::cuCtxPopCurrent_v2(pctx as *mut _) as c_uint) } - } - - fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult { - unsafe { CUresult(cuda::cuCtxGetApiVersion(ctx as *mut _, version) as c_uint) } - } - - fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult { - unsafe { CUresult(cuda::cuCtxGetCurrent(pctx as *mut _) as c_uint) } - } - fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult { - unsafe { CUresult(cuda::cuMemAlloc_v2(dptr as *mut _, bytesize) as c_uint) } - } - - fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult { - unsafe { CUresult(cuda::cuDeviceGetUuid(uuid as *mut _, dev) as c_uint) } - } - - fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult { - unsafe { CUresult(cuda::cuDevicePrimaryCtxGetState(dev, flags, active) as c_uint) } - } - - fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult { - unsafe { CUresult(cuda::cuStreamGetCtx(hStream as _, pctx as _) as c_uint) } - } - - fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult { - unsafe { CUresult(cuda::cuStreamCreate(stream as _, flags as _) as c_uint) } - } - - fn cuMemFree_v2(mem: *mut c_void) -> CUresult { - unsafe { CUresult(cuda::cuMemFree_v2(mem as _) as c_uint) } - } - - fn cuStreamDestroy_v2(stream: CUstream) -> CUresult { - unsafe { CUresult(cuda::cuStreamDestroy_v2(stream as _) as c_uint) } - } -} diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs index 72ca51c..1568f47 100644 --- a/zluda/src/lib.rs +++ b/zluda/src/lib.rs @@ -1,13 +1,79 @@ -#[macro_use] -extern crate lazy_static; -#[cfg(test)] -extern crate cuda_driver_sys; -#[cfg(test)] -#[macro_use] -extern crate paste; -extern crate ptx; - -#[allow(warnings)] -pub mod cuda; -mod cuda_impl; pub(crate) mod r#impl; + +macro_rules! unimplemented { + ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => { + $( + #[cfg_attr(not(test), no_mangle)] + #[allow(improper_ctypes)] + #[allow(improper_ctypes_definitions)] + pub unsafe extern $abi fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type { + crate::r#impl::unimplemented() + } + )* + }; +} + +macro_rules! implemented { + ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => { + $( + #[cfg_attr(not(test), no_mangle)] + #[allow(improper_ctypes)] + #[allow(improper_ctypes_definitions)] + pub unsafe extern $abi fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type { + cuda_base::cuda_normalize_fn!( crate::r#impl::$fn_name ) ($(crate::r#impl::FromCuda::from_cuda(&$arg_id)?),*)?; + Ok(()) + } + )* + }; +} + +macro_rules! implemented_in_function { + ($($abi:literal fn $fn_name:ident( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:path;)*) => { + $( + #[cfg_attr(not(test), no_mangle)] + #[allow(improper_ctypes)] + #[allow(improper_ctypes_definitions)] + pub unsafe extern $abi fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type { + cuda_base::cuda_normalize_fn!( crate::r#impl::function::$fn_name ) ($(crate::r#impl::FromCuda::from_cuda(&$arg_id)?),*)?; + Ok(()) + } + )* + }; +} + +cuda_base::cuda_function_declarations!( + unimplemented, + implemented <= [ + cuCtxGetLimit, + cuCtxSetCurrent, + cuCtxSetLimit, + cuCtxSynchronize, + cuDeviceComputeCapability, + cuDeviceGet, + cuDeviceGetAttribute, + cuDeviceGetCount, + cuDeviceGetLuid, + cuDeviceGetName, + cuDevicePrimaryCtxRelease, + cuDevicePrimaryCtxRetain, + cuDeviceGetProperties, + cuDeviceGetUuid, + cuDeviceGetUuid_v2, + cuDeviceTotalMem_v2, + cuDriverGetVersion, + cuFuncGetAttribute, + cuInit, + cuMemAlloc_v2, + cuMemFree_v2, + cuMemcpyDtoH_v2, + cuMemcpyHtoD_v2, + cuModuleGetFunction, + cuModuleLoadData, + cuModuleUnload, + cuPointerGetAttribute, + cuMemGetAddressRange_v2, + ], + implemented_in_function <= [ + cuLaunchKernel, + ] +);
\ No newline at end of file |