Fix host code and update to CUDA 12.4 (#299)

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