Refactor main library, implement some more functionality

author: Andrzej Janik <[email protected]> 2020-09-01 01:43:09 +0200
committer: Andrzej Janik <[email protected]> 2020-09-01 01:43:09 +0200
commit: 2e4cadc2ab061c61bacd43fab9a375b5492a1897 (patch)
tree: c5ca5d149c8621437a9d861e1a9872052792fda9
parent: 34dc149be1cc83ecd931c0da742befe3eb67feaf (diff)
download: ZLUDA-2e4cadc2ab061c61bacd43fab9a375b5492a1897.tar.gz
ZLUDA-2e4cadc2ab061c61bacd43fab9a375b5492a1897.zip
18 files changed, 6019 insertions, 851 deletions
diff --git a/level_zero/src/ze.rs b/level_zero/src/ze.rs
index 3e7d3ce..cee736c 100644
--- a/level_zero/src/ze.rs
+++ b/level_zero/src/ze.rs
@@ -1,7 +1,7 @@
 use crate::sys;
 use std::{
     ffi::{c_void, CStr},
-    fmt::{Debug, Display},
+    fmt::Debug,
     marker::PhantomData,
     mem, ptr,
 };
@@ -12,7 +12,7 @@ macro_rules! check {
         {
             let err = unsafe { $expr };
             if err != crate::sys::ze_result_t::ZE_RESULT_SUCCESS {
-                return Result::Err(Error(err));
+                return Result::Err(err);
             }
         }
     };
@@ -27,39 +27,24 @@ macro_rules! check_panic {
     };
 }
 
-pub type Result<T> = std::result::Result<T, Error>;
+pub type Result<T> = std::result::Result<T, sys::ze_result_t>;
 
 #[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)]
 pub struct Error(pub sys::ze_result_t);
 
-impl Error {
-    fn new<T>(res: sys::ze_result_t, default: T) -> Result<T> {
-        if res == sys::ze_result_t::ZE_RESULT_SUCCESS {
-            Ok(default)
-        } else {
-            Err(Self(res))
-        }
-    }
-}
-
-impl Display for Error {
-    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
-        Debug::fmt(self, f)
-    }
-}
-
-impl std::error::Error for Error {}
-
 pub fn init() -> Result<()> {
-    Error::new(
-        unsafe { sys::zeInit(sys::ze_init_flags_t::ZE_INIT_FLAG_GPU_ONLY) },
-        (),
-    )
+    match unsafe { sys::zeInit(sys::ze_init_flags_t::ZE_INIT_FLAG_GPU_ONLY) } {
+        sys::ze_result_t::ZE_RESULT_SUCCESS => Ok(()),
+        e => Err(e),
+    }
 }
 
 #[repr(transparent)]
 pub struct Driver(sys::ze_driver_handle_t);
 
+unsafe impl Send for Driver {}
+unsafe impl Sync for Driver {}
+
 impl Driver {
     pub unsafe fn as_ffi(&self) -> sys::ze_driver_handle_t {
         self.0
@@ -184,6 +169,13 @@ impl Context {
     }
 }
 
+impl Drop for Context {
+    #[allow(unused_must_use)]
+    fn drop(&mut self) {
+        check_panic! { sys::zeContextDestroy(self.0) };
+    }
+}
+
 #[repr(transparent)]
 pub struct CommandQueue(sys::ze_command_queue_handle_t);
 
diff --git a/notcuda/Cargo.toml b/notcuda/Cargo.toml
index b941a90..56ec4d6 100644
--- a/notcuda/Cargo.toml
+++ b/notcuda/Cargo.toml
@@ -5,7 +5,7 @@ authors = ["Andrzej Janik <[email protected]>"]
 edition = "2018"
 
 [lib]
-name = "nvcuda"
+name = "notcuda"
 crate-type = ["cdylib"]
 
 [dependencies]
@@ -14,4 +14,8 @@ level_zero = { path = "../level_zero" }
 level_zero-sys = { path = "../level_zero-sys" }
 lazy_static = "1.4"
 num_enum = "0.4"
-lz4 = "1.23"
-\ No newline at end of file
+lz4 = "1.23"
+
+[dev-dependencies]
+cuda-driver-sys = "0.3.0"
+paste = "1.0"
+\ No newline at end of file
diff --git a/notcuda/README b/notcuda/README
new file mode 100644
index 0000000..089ddcd
--- /dev/null
+++ b/notcuda/README
@@ -0,0 +1,3 @@
+bindgen /usr/local/cuda/include/cuda.h -o cuda.rs --whitelist-function="^cu.*" --size_t-is-usize --default-enum-style=newtype --no-layout-tests --no-doc-comments --no-derive-debug --new-type-alias "^CUdevice$|^CUdeviceptr$"
+sed -i -e 's/extern "C" {//g' -e 's/-> CUresult;/-> CUresult { impl_::unsupported()/g' -e 's/pub fn /#[no_mangle] pub extern "C" fn /g' cuda.rs
+rustfmt cuda.rs
+\ No newline at end of file
diff --git a/notcuda/src/cu.rs b/notcuda/src/cu.rs
deleted file mode 100644
index 3bb5237..0000000
--- a/notcuda/src/cu.rs
+++ /dev/null
@@ -1,220 +0,0 @@
-use num_enum::TryFromPrimitive;
-use std::convert::TryFrom;
-use std::os::raw::c_int;
-use std::{mem, ptr};
-use crate::ze;
-
-#[repr(u32)]
-#[allow(non_camel_case_types)]
-pub enum Result {
-    SUCCESS = 0,
-    ERROR_INVALID_VALUE = 1,
-    ERROR_OUT_OF_MEMORY = 2,
-    ERROR_NOT_INITIALIZED = 3,
-    ERROR_DEINITIALIZED = 4,
-    ERROR_PROFILER_DISABLED = 5,
-    ERROR_PROFILER_NOT_INITIALIZED = 6,
-    ERROR_PROFILER_ALREADY_STARTED = 7,
-    ERROR_PROFILER_ALREADY_STOPPED = 8,
-    ERROR_NO_DEVICE = 100,
-    ERROR_INVALID_DEVICE = 101,
-    ERROR_INVALID_IMAGE = 200,
-    ERROR_INVALID_CONTEXT = 201,
-    ERROR_CONTEXT_ALREADY_CURRENT = 202,
-    ERROR_MAP_FAILED = 205,
-    ERROR_UNMAP_FAILED = 206,
-    ERROR_ARRAY_IS_MAPPED = 207,
-    ERROR_ALREADY_MAPPED = 208,
-    ERROR_NO_BINARY_FOR_GPU = 209,
-    ERROR_ALREADY_ACQUIRED = 210,
-    ERROR_NOT_MAPPED = 211,
-    ERROR_NOT_MAPPED_AS_ARRAY = 212,
-    ERROR_NOT_MAPPED_AS_POINTER = 213,
-    ERROR_ECC_UNCORRECTABLE = 214,
-    ERROR_UNSUPPORTED_LIMIT = 215,
-    ERROR_CONTEXT_ALREADY_IN_USE = 216,
-    ERROR_PEER_ACCESS_UNSUPPORTED = 217,
-    ERROR_INVALID_PTX = 218,
-    ERROR_INVALID_GRAPHICS_CONTEXT = 219,
-    ERROR_NVLINK_UNCORRECTABLE = 220,
-    ERROR_JIT_COMPILER_NOT_FOUND = 221,
-    ERROR_INVALID_SOURCE = 300,
-    ERROR_FILE_NOT_FOUND = 301,
-    ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302,
-    ERROR_SHARED_OBJECT_INIT_FAILED = 303,
-    ERROR_OPERATING_SYSTEM = 304,
-    ERROR_INVALID_HANDLE = 400,
-    ERROR_ILLEGAL_STATE = 401,
-    ERROR_NOT_FOUND = 500,
-    ERROR_NOT_READY = 600,
-    ERROR_ILLEGAL_ADDRESS = 700,
-    ERROR_LAUNCH_OUT_OF_RESOURCES = 701,
-    ERROR_LAUNCH_TIMEOUT = 702,
-    ERROR_LAUNCH_INCOMPATIBLE_TEXTURING = 703,
-    ERROR_PEER_ACCESS_ALREADY_ENABLED = 704,
-    ERROR_PEER_ACCESS_NOT_ENABLED = 705,
-    ERROR_PRIMARY_CONTEXT_ACTIVE = 708,
-    ERROR_CONTEXT_IS_DESTROYED = 709,
-    ERROR_ASSERT = 710,
-    ERROR_TOO_MANY_PEERS = 711,
-    ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712,
-    ERROR_HOST_MEMORY_NOT_REGISTERED = 713,
-    ERROR_HARDWARE_STACK_ERROR = 714,
-    ERROR_ILLEGAL_INSTRUCTION = 715,
-    ERROR_MISALIGNED_ADDRESS = 716,
-    ERROR_INVALID_ADDRESS_SPACE = 717,
-    ERROR_INVALID_PC = 718,
-    ERROR_LAUNCH_FAILED = 719,
-    ERROR_COOPERATIVE_LAUNCH_TOO_LARGE = 720,
-    ERROR_NOT_PERMITTED = 800,
-    ERROR_NOT_SUPPORTED = 801,
-    ERROR_SYSTEM_NOT_READY = 802,
-    ERROR_SYSTEM_DRIVER_MISMATCH = 803,
-    ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE = 804,
-    ERROR_STREAM_CAPTURE_UNSUPPORTED = 900,
-    ERROR_STREAM_CAPTURE_INVALIDATED = 901,
-    ERROR_STREAM_CAPTURE_MERGE = 902,
-    ERROR_STREAM_CAPTURE_UNMATCHED = 903,
-    ERROR_STREAM_CAPTURE_UNJOINED = 904,
-    ERROR_STREAM_CAPTURE_ISOLATION = 905,
-    ERROR_STREAM_CAPTURE_IMPLICIT = 906,
-    ERROR_CAPTURED_EVENT = 907,
-    ERROR_STREAM_CAPTURE_WRONG_THREAD = 908,
-    ERROR_TIMEOUT = 909,
-    ERROR_GRAPH_EXEC_UPDATE_FAILURE = 910,
-    ERROR_UNKNOWN = 999,
-}
-
-pub enum DeviceAttribute {
-    Static(DeviceStaticAttribute),
-    Dynamic(DeviceDynamicAttribute)
-}
-
-impl DeviceAttribute {
-    pub fn try_new(e: u8) -> Option<DeviceAttribute> {
-        DeviceStaticAttribute::try_from(e).map(DeviceAttribute::Static)
-            .or_else(|_| DeviceGeneralAttribute::try_from(e).map(DeviceDynamicAttribute::General).map(DeviceAttribute::Dynamic))
-            .or_else(|_| DeviceTextureAttribute::try_from(e).map(DeviceDynamicAttribute::Texture).map(DeviceAttribute::Dynamic))
-            .ok()
-    }
-}
-
-#[repr(u8)]
-#[derive(TryFromPrimitive)]
-#[allow(non_camel_case_types)]
-pub enum DeviceStaticAttribute {
-    GPU_OVERLAP = 15,
-    KERNEL_EXEC_TIMEOUT = 17,
-    INTEGRATED = 18,
-    COMPUTE_CAPABILITY_MAJOR = 75,
-    COMPUTE_CAPABILITY_MINOR = 76,
-}
-
-pub enum DeviceDynamicAttribute {
-    General(DeviceGeneralAttribute),
-    Texture(DeviceTextureAttribute)
-}
-
-
-#[repr(u8)]
-#[derive(TryFromPrimitive)]
-#[allow(non_camel_case_types)]
-pub enum DeviceGeneralAttribute {
-    MULTIPROCESSOR_COUNT = 16,
-    CAN_MAP_HOST_MEMORY = 19,
-    ASYNC_ENGINE_COUNT = 40,
-}
-
-
-#[repr(u8)]
-#[derive(TryFromPrimitive)]
-#[allow(non_camel_case_types)]
-pub enum DeviceTextureAttribute {
-    MAXIMUM_TEXTURE1D_WIDTH = 21
-}
-
-
-impl Result {
-    pub fn from_l0(result: l0_sys::ze_result_t) -> Result {
-        match result {
-            l0_sys::ze_result_t::ZE_RESULT_SUCCESS => Result::SUCCESS,
-            l0_sys::ze_result_t::ZE_RESULT_ERROR_UNINITIALIZED => Result::ERROR_NOT_INITIALIZED,
-            l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ENUMERATION => Result::ERROR_INVALID_VALUE,
-            l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ARGUMENT => Result::ERROR_INVALID_VALUE,
-            l0_sys::ze_result_t::ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY => Result::ERROR_OUT_OF_MEMORY,
-            l0_sys::ze_result_t::ZE_RESULT_ERROR_UNSUPPORTED_FEATURE => Result::ERROR_NOT_SUPPORTED,
-            _ => Result::ERROR_UNKNOWN
-        }
-    }
-}
-
-#[repr(C)]
-#[derive(PartialEq, Eq)]
-pub struct Uuid {
-    pub x: [std::os::raw::c_uchar; 16]
-}
-
-#[repr(transparent)]
-pub struct Device(pub c_int);
-
-#[repr(transparent)]
-pub struct DevicePtr(pub usize);
-
-#[repr(transparent)]
-#[derive(Clone, Copy, PartialEq)]
-pub struct Context(pub *mut ze::Context);
-impl Context {
-    pub fn null() -> Context {
-        Context(ptr::null_mut())
-    }
-}
-
-#[repr(transparent)]
-#[derive(Clone, Copy, PartialEq)]
-pub struct Module(*mut ze::Module);
-
-impl Module {
-    pub fn null() -> Module {
-        Module(ptr::null_mut())
-    }
-
-    pub fn new(inner: ze::Module) -> Module {
-        let mut boxed = Box::new(inner);
-        let result = Module(boxed.as_mut());
-        mem::forget(boxed);
-        result
-    }
-}
-
-#[repr(transparent)]
-pub struct Function(*mut ());
-
-#[repr(transparent)]
-pub struct Stream(*mut ());
-
-#[repr(i32)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash, PartialOrd, Ord)]
-#[allow(non_camel_case_types)]
-pub enum JitOption {
-    MAX_REGISTERS = 0,
-    THREADS_PER_BLOCK = 1,
-    WALL_TIME = 2,
-    INFO_LOG_BUFFER = 3,
-    INFO_LOG_BUFFER_SIZE_BYTES = 4,
-    ERROR_LOG_BUFFER = 5,
-    ERROR_LOG_BUFFER_SIZE_BYTES = 6,
-    OPTIMIZATION_LEVEL = 7,
-    TARGET_FROM_CUCONTEXT = 8,
-    TARGET = 9,
-    FALLBACK_STRATEGY = 10,
-    GENERATE_DEBUG_INFO = 11,
-    LOG_VERBOSE = 12,
-    GENERATE_LINE_INFO = 13,
-    CACHE_MODE = 14,
-    NEW_SM3X_OPT = 15,
-    FAST_COMPILE = 16,
-    GLOBAL_SYMBOL_NAMES = 17,
-    GLOBAL_SYMBOL_ADDRESSES = 18,
-    GLOBAL_SYMBOL_COUNT = 19,
-    NUM_OPTIONS = 20,
-}
-\ No newline at end of file
diff --git a/notcuda/src/cuda.rs b/notcuda/src/cuda.rs
new file mode 100644
index 0000000..3267042
--- /dev/null
+++ b/notcuda/src/cuda.rs
@@ -0,0 +1,4751 @@
+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 ::std::os::raw::c_ulonglong);
+#[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 "C" 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 "C" fn(
+        hStream: CUstream,
+        status: CUresult,
+        userData: *mut ::std::os::raw::c_void,
+    ),
+>;
+pub type CUoccupancyB2DSize =
+    ::std::option::Option<unsafe extern "C" 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;
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGetErrorString(
+    error: CUresult,
+    pStr: *mut *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGetErrorName(
+    error: CUresult,
+    pStr: *mut *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuInit(Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::init().encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDriverGetVersion(driverVersion: *mut ::std::os::raw::c_int) -> CUresult {
+    unsafe { *driverVersion = r#impl::driver_get_version() };
+    CUresult::CUDA_SUCCESS
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGet(device: *mut CUdevice, ordinal: ::std::os::raw::c_int) -> CUresult {
+    r#impl::device::get(device.decuda(), ordinal)
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetCount(count: *mut ::std::os::raw::c_int) -> CUresult {
+    r#impl::device::get_count(count)
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetName(
+    name: *mut ::std::os::raw::c_char,
+    len: ::std::os::raw::c_int,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::device::get_name(name, len, dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: CUdevice) -> CUresult {
+    r#impl::device::get_uuid(uuid, dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceTotalMem_v2(bytes: *mut usize, dev: CUdevice) -> CUresult {
+    r#impl::device::total_mem_v2(bytes, dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetAttribute(
+    pi: *mut ::std::os::raw::c_int,
+    attrib: CUdevice_attribute,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::device::get_attribute(pi, attrib, dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetNvSciSyncAttributes(
+    nvSciSyncAttrList: *mut ::std::os::raw::c_void,
+    dev: CUdevice,
+    flags: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetProperties(prop: *mut CUdevprop, dev: CUdevice) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceComputeCapability(
+    major: *mut ::std::os::raw::c_int,
+    minor: *mut ::std::os::raw::c_int,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDevicePrimaryCtxRetain(pctx: *mut CUcontext, dev: CUdevice) -> CUresult {
+    r#impl::device::primary_ctx_retain(pctx.decuda(), dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDevicePrimaryCtxRelease_v2(dev: CUdevice) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDevicePrimaryCtxSetFlags_v2(
+    dev: CUdevice,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDevicePrimaryCtxGetState(
+    dev: CUdevice,
+    flags: *mut ::std::os::raw::c_uint,
+    active: *mut ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::device::primary_ctx_get_state(dev.decuda(), flags, active).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDevicePrimaryCtxReset_v2(dev: CUdevice) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxCreate_v2(
+    pctx: *mut CUcontext,
+    flags: ::std::os::raw::c_uint,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::context::create_v2(pctx.decuda(), flags, dev.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxDestroy_v2(ctx: CUcontext) -> CUresult {
+    r#impl::context::destroy_v2(ctx.decuda())
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxPushCurrent_v2(ctx: CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxPopCurrent_v2(pctx: *mut CUcontext) -> CUresult {
+    r#impl::context::pop_current_v2(pctx.decuda())
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxSetCurrent(ctx: CUcontext) -> CUresult {
+    r#impl::context::set_current(ctx.decuda())
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetCurrent(pctx: *mut CUcontext) -> CUresult {
+    r#impl::context::get_current(pctx.decuda()).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetDevice(device: *mut CUdevice) -> CUresult {
+    r#impl::context::get_device(device.decuda())
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetFlags(flags: *mut ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxSynchronize() -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxSetLimit(limit: CUlimit, value: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetLimit(pvalue: *mut usize, limit: CUlimit) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetCacheConfig(pconfig: *mut CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxSetCacheConfig(config: CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetSharedMemConfig(pConfig: *mut CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxSetSharedMemConfig(config: CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetApiVersion(
+    ctx: CUcontext,
+    version: *mut ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::context::get_api_version(ctx.decuda(), version)
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxGetStreamPriorityRange(
+    leastPriority: *mut ::std::os::raw::c_int,
+    greatestPriority: *mut ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxResetPersistingL2Cache() -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxAttach(pctx: *mut CUcontext, flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxDetach(ctx: CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleLoad(
+    module: *mut CUmodule,
+    fname: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleLoadData(
+    module: *mut CUmodule,
+    image: *const ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleLoadFatBinary(
+    module: *mut CUmodule,
+    fatCubin: *const ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleUnload(hmod: CUmodule) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleGetFunction(
+    hfunc: *mut CUfunction,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleGetGlobal_v2(
+    dptr: *mut CUdeviceptr,
+    bytes: *mut usize,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleGetTexRef(
+    pTexRef: *mut CUtexref,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuModuleGetSurfRef(
+    pSurfRef: *mut CUsurfref,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLinkCreate_v2(
+    numOptions: ::std::os::raw::c_uint,
+    options: *mut CUjit_option,
+    optionValues: *mut *mut ::std::os::raw::c_void,
+    stateOut: *mut CUlinkState,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLinkComplete(
+    state: CUlinkState,
+    cubinOut: *mut *mut ::std::os::raw::c_void,
+    sizeOut: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLinkDestroy(state: CUlinkState) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult {
+    r#impl::memory::alloc_v2(dptr.decuda(), bytesize)
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAllocPitch_v2(
+    dptr: *mut CUdeviceptr,
+    pPitch: *mut usize,
+    WidthInBytes: usize,
+    Height: usize,
+    ElementSizeBytes: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemGetAddressRange_v2(
+    pbase: *mut CUdeviceptr,
+    psize: *mut usize,
+    dptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAllocHost_v2(
+    pp: *mut *mut ::std::os::raw::c_void,
+    bytesize: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemHostAlloc(
+    pp: *mut *mut ::std::os::raw::c_void,
+    bytesize: usize,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemHostGetDevicePointer_v2(
+    pdptr: *mut CUdeviceptr,
+    p: *mut ::std::os::raw::c_void,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemHostGetFlags(
+    pFlags: *mut ::std::os::raw::c_uint,
+    p: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAllocManaged(
+    dptr: *mut CUdeviceptr,
+    bytesize: usize,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetByPCIBusId(
+    dev: *mut CUdevice,
+    pciBusId: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetPCIBusId(
+    pciBusId: *mut ::std::os::raw::c_char,
+    len: ::std::os::raw::c_int,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuIpcGetEventHandle(pHandle: *mut CUipcEventHandle, event: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuIpcOpenEventHandle(
+    phEvent: *mut CUevent,
+    handle: CUipcEventHandle,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuIpcGetMemHandle(pHandle: *mut CUipcMemHandle, dptr: CUdeviceptr) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuIpcOpenMemHandle(
+    pdptr: *mut CUdeviceptr,
+    handle: CUipcMemHandle,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuIpcCloseMemHandle(dptr: CUdeviceptr) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemHostRegister_v2(
+    p: *mut ::std::os::raw::c_void,
+    bytesize: usize,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemHostUnregister(p: *mut ::std::os::raw::c_void) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy(dst: CUdeviceptr, src: CUdeviceptr, ByteCount: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyPeer(
+    dstDevice: CUdeviceptr,
+    dstContext: CUcontext,
+    srcDevice: CUdeviceptr,
+    srcContext: CUcontext,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyHtoD_v2(
+    dstDevice: CUdeviceptr,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::memory::copy_v2(dstDevice.decuda(), srcHost, ByteCount).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyDtoH_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::memory::copy_v2(dstHost, srcDevice.decuda(), ByteCount).encuda()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyDtoD_v2(
+    dstDevice: CUdeviceptr,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyDtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyAtoD_v2(
+    dstDevice: CUdeviceptr,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyHtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyAtoH_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyAtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy2D_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy2DUnaligned_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy3D_v2(pCopy: *const CUDA_MEMCPY3D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy3DPeer(pCopy: *const CUDA_MEMCPY3D_PEER) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyAsync(
+    dst: CUdeviceptr,
+    src: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyPeerAsync(
+    dstDevice: CUdeviceptr,
+    dstContext: CUcontext,
+    srcDevice: CUdeviceptr,
+    srcContext: CUcontext,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyHtoDAsync_v2(
+    dstDevice: CUdeviceptr,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyDtoHAsync_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyDtoDAsync_v2(
+    dstDevice: CUdeviceptr,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyHtoAAsync_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpyAtoHAsync_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy2DAsync_v2(pCopy: *const CUDA_MEMCPY2D, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy3DAsync_v2(pCopy: *const CUDA_MEMCPY3D, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemcpy3DPeerAsync(
+    pCopy: *const CUDA_MEMCPY3D_PEER,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD8_v2(
+    dstDevice: CUdeviceptr,
+    uc: ::std::os::raw::c_uchar,
+    N: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD16_v2(
+    dstDevice: CUdeviceptr,
+    us: ::std::os::raw::c_ushort,
+    N: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD32_v2(
+    dstDevice: CUdeviceptr,
+    ui: ::std::os::raw::c_uint,
+    N: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D8_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    uc: ::std::os::raw::c_uchar,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D16_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    us: ::std::os::raw::c_ushort,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D32_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    ui: ::std::os::raw::c_uint,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD8Async(
+    dstDevice: CUdeviceptr,
+    uc: ::std::os::raw::c_uchar,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD16Async(
+    dstDevice: CUdeviceptr,
+    us: ::std::os::raw::c_ushort,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD32Async(
+    dstDevice: CUdeviceptr,
+    ui: ::std::os::raw::c_uint,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D8Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    uc: ::std::os::raw::c_uchar,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D16Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    us: ::std::os::raw::c_ushort,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemsetD2D32Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    ui: ::std::os::raw::c_uint,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuArrayCreate_v2(
+    pHandle: *mut CUarray,
+    pAllocateArray: *const CUDA_ARRAY_DESCRIPTOR,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuArrayGetDescriptor_v2(
+    pArrayDescriptor: *mut CUDA_ARRAY_DESCRIPTOR,
+    hArray: CUarray,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuArrayDestroy(hArray: CUarray) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuArray3DCreate_v2(
+    pHandle: *mut CUarray,
+    pAllocateArray: *const CUDA_ARRAY3D_DESCRIPTOR,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuArray3DGetDescriptor_v2(
+    pArrayDescriptor: *mut CUDA_ARRAY3D_DESCRIPTOR,
+    hArray: CUarray,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMipmappedArrayCreate(
+    pHandle: *mut CUmipmappedArray,
+    pMipmappedArrayDesc: *const CUDA_ARRAY3D_DESCRIPTOR,
+    numMipmapLevels: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMipmappedArrayGetLevel(
+    pLevelArray: *mut CUarray,
+    hMipmappedArray: CUmipmappedArray,
+    level: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMipmappedArrayDestroy(hMipmappedArray: CUmipmappedArray) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAddressReserve(
+    ptr: *mut CUdeviceptr,
+    size: usize,
+    alignment: usize,
+    addr: CUdeviceptr,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAddressFree(ptr: CUdeviceptr, size: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemCreate(
+    handle: *mut CUmemGenericAllocationHandle,
+    size: usize,
+    prop: *const CUmemAllocationProp,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemRelease(handle: CUmemGenericAllocationHandle) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemMap(
+    ptr: CUdeviceptr,
+    size: usize,
+    offset: usize,
+    handle: CUmemGenericAllocationHandle,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemUnmap(ptr: CUdeviceptr, size: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemSetAccess(
+    ptr: CUdeviceptr,
+    size: usize,
+    desc: *const CUmemAccessDesc,
+    count: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemGetAccess(
+    flags: *mut ::std::os::raw::c_ulonglong,
+    location: *const CUmemLocation,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemExportToShareableHandle(
+    shareableHandle: *mut ::std::os::raw::c_void,
+    handle: CUmemGenericAllocationHandle,
+    handleType: CUmemAllocationHandleType,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemImportFromShareableHandle(
+    handle: *mut CUmemGenericAllocationHandle,
+    osHandle: *mut ::std::os::raw::c_void,
+    shHandleType: CUmemAllocationHandleType,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemGetAllocationGranularity(
+    granularity: *mut usize,
+    prop: *const CUmemAllocationProp,
+    option: CUmemAllocationGranularity_flags,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemGetAllocationPropertiesFromHandle(
+    prop: *mut CUmemAllocationProp,
+    handle: CUmemGenericAllocationHandle,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemRetainAllocationHandle(
+    handle: *mut CUmemGenericAllocationHandle,
+    addr: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuPointerGetAttribute(
+    data: *mut ::std::os::raw::c_void,
+    attribute: CUpointer_attribute,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemPrefetchAsync(
+    devPtr: CUdeviceptr,
+    count: usize,
+    dstDevice: CUdevice,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemAdvise(
+    devPtr: CUdeviceptr,
+    count: usize,
+    advice: CUmem_advise,
+    device: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuMemRangeGetAttribute(
+    data: *mut ::std::os::raw::c_void,
+    dataSize: usize,
+    attribute: CUmem_range_attribute,
+    devPtr: CUdeviceptr,
+    count: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuPointerSetAttribute(
+    value: *const ::std::os::raw::c_void,
+    attribute: CUpointer_attribute,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamCreate(
+    phStream: *mut CUstream,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamCreateWithPriority(
+    phStream: *mut CUstream,
+    flags: ::std::os::raw::c_uint,
+    priority: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamGetPriority(
+    hStream: CUstream,
+    priority: *mut ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamGetFlags(
+    hStream: CUstream,
+    flags: *mut ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamWaitEvent(
+    hStream: CUstream,
+    hEvent: CUevent,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamAddCallback(
+    hStream: CUstream,
+    callback: CUstreamCallback,
+    userData: *mut ::std::os::raw::c_void,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamBeginCapture_v2(
+    hStream: CUstream,
+    mode: CUstreamCaptureMode,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuThreadExchangeStreamCaptureMode(mode: *mut CUstreamCaptureMode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamEndCapture(hStream: CUstream, phGraph: *mut CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamIsCapturing(
+    hStream: CUstream,
+    captureStatus: *mut CUstreamCaptureStatus,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamGetCaptureInfo(
+    hStream: CUstream,
+    captureStatus: *mut CUstreamCaptureStatus,
+    id: *mut cuuint64_t,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamAttachMemAsync(
+    hStream: CUstream,
+    dptr: CUdeviceptr,
+    length: usize,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamQuery(hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamSynchronize(hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamDestroy_v2(hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamCopyAttributes(dst: CUstream, src: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamGetAttribute(
+    hStream: CUstream,
+    attr: CUstreamAttrID,
+    value_out: *mut CUstreamAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamSetAttribute(
+    hStream: CUstream,
+    attr: CUstreamAttrID,
+    value: *const CUstreamAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventCreate(phEvent: *mut CUevent, Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventRecord(hEvent: CUevent, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventQuery(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventSynchronize(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventDestroy_v2(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuEventElapsedTime(
+    pMilliseconds: *mut f32,
+    hStart: CUevent,
+    hEnd: CUevent,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuImportExternalMemory(
+    extMem_out: *mut CUexternalMemory,
+    memHandleDesc: *const CUDA_EXTERNAL_MEMORY_HANDLE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuExternalMemoryGetMappedBuffer(
+    devPtr: *mut CUdeviceptr,
+    extMem: CUexternalMemory,
+    bufferDesc: *const CUDA_EXTERNAL_MEMORY_BUFFER_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuExternalMemoryGetMappedMipmappedArray(
+    mipmap: *mut CUmipmappedArray,
+    extMem: CUexternalMemory,
+    mipmapDesc: *const CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDestroyExternalMemory(extMem: CUexternalMemory) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuImportExternalSemaphore(
+    extSem_out: *mut CUexternalSemaphore,
+    semHandleDesc: *const CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSignalExternalSemaphoresAsync(
+    extSemArray: *const CUexternalSemaphore,
+    paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS,
+    numExtSems: ::std::os::raw::c_uint,
+    stream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuWaitExternalSemaphoresAsync(
+    extSemArray: *const CUexternalSemaphore,
+    paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS,
+    numExtSems: ::std::os::raw::c_uint,
+    stream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDestroyExternalSemaphore(extSem: CUexternalSemaphore) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamWaitValue32(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint32_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamWaitValue64(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint64_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamWriteValue32(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint32_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamWriteValue64(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint64_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuStreamBatchMemOp(
+    stream: CUstream,
+    count: ::std::os::raw::c_uint,
+    paramArray: *mut CUstreamBatchMemOpParams,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncGetAttribute(
+    pi: *mut ::std::os::raw::c_int,
+    attrib: CUfunction_attribute,
+    hfunc: CUfunction,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncSetAttribute(
+    hfunc: CUfunction,
+    attrib: CUfunction_attribute,
+    value: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncSetCacheConfig(hfunc: CUfunction, config: CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncSetSharedMemConfig(hfunc: CUfunction, config: CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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 {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLaunchCooperativeKernelMultiDevice(
+    launchParamsList: *mut CUDA_LAUNCH_PARAMS,
+    numDevices: ::std::os::raw::c_uint,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLaunchHostFunc(
+    hStream: CUstream,
+    fn_: CUhostFn,
+    userData: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncSetSharedSize(
+    hfunc: CUfunction,
+    bytes: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuParamSetSize(hfunc: CUfunction, numbytes: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuParamSeti(
+    hfunc: CUfunction,
+    offset: ::std::os::raw::c_int,
+    value: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuParamSetf(
+    hfunc: CUfunction,
+    offset: ::std::os::raw::c_int,
+    value: f32,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLaunch(f: CUfunction) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLaunchGrid(
+    f: CUfunction,
+    grid_width: ::std::os::raw::c_int,
+    grid_height: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuLaunchGridAsync(
+    f: CUfunction,
+    grid_width: ::std::os::raw::c_int,
+    grid_height: ::std::os::raw::c_int,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuParamSetTexRef(
+    hfunc: CUfunction,
+    texunit: ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphCreate(phGraph: *mut CUgraph, flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddKernelNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphKernelNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphKernelNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddMemcpyNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    copyParams: *const CUDA_MEMCPY3D,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphMemcpyNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_MEMCPY3D,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphMemcpyNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_MEMCPY3D,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddMemsetNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    memsetParams: *const CUDA_MEMSET_NODE_PARAMS,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphMemsetNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_MEMSET_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphMemsetNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_MEMSET_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddHostNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    nodeParams: *const CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphHostNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphHostNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddChildGraphNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    childGraph: CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphChildGraphNodeGetGraph(
+    hNode: CUgraphNode,
+    phGraph: *mut CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddEmptyNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphClone(phGraphClone: *mut CUgraph, originalGraph: CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphNodeFindInClone(
+    phNode: *mut CUgraphNode,
+    hOriginalNode: CUgraphNode,
+    hClonedGraph: CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphNodeGetType(hNode: CUgraphNode, type_: *mut CUgraphNodeType) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphGetNodes(
+    hGraph: CUgraph,
+    nodes: *mut CUgraphNode,
+    numNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphGetRootNodes(
+    hGraph: CUgraph,
+    rootNodes: *mut CUgraphNode,
+    numRootNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphGetEdges(
+    hGraph: CUgraph,
+    from: *mut CUgraphNode,
+    to: *mut CUgraphNode,
+    numEdges: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphNodeGetDependencies(
+    hNode: CUgraphNode,
+    dependencies: *mut CUgraphNode,
+    numDependencies: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphNodeGetDependentNodes(
+    hNode: CUgraphNode,
+    dependentNodes: *mut CUgraphNode,
+    numDependentNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphAddDependencies(
+    hGraph: CUgraph,
+    from: *const CUgraphNode,
+    to: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphRemoveDependencies(
+    hGraph: CUgraph,
+    from: *const CUgraphNode,
+    to: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphDestroyNode(hNode: CUgraphNode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphInstantiate_v2(
+    phGraphExec: *mut CUgraphExec,
+    hGraph: CUgraph,
+    phErrorNode: *mut CUgraphNode,
+    logBuffer: *mut ::std::os::raw::c_char,
+    bufferSize: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecKernelNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecMemcpyNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    copyParams: *const CUDA_MEMCPY3D,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecMemsetNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    memsetParams: *const CUDA_MEMSET_NODE_PARAMS,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecHostNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphLaunch(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecDestroy(hGraphExec: CUgraphExec) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphDestroy(hGraph: CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphExecUpdate(
+    hGraphExec: CUgraphExec,
+    hGraph: CUgraph,
+    hErrorNode_out: *mut CUgraphNode,
+    updateResult_out: *mut CUgraphExecUpdateResult,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphKernelNodeCopyAttributes(dst: CUgraphNode, src: CUgraphNode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphKernelNodeGetAttribute(
+    hNode: CUgraphNode,
+    attr: CUkernelNodeAttrID,
+    value_out: *mut CUkernelNodeAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphKernelNodeSetAttribute(
+    hNode: CUgraphNode,
+    attr: CUkernelNodeAttrID,
+    value: *const CUkernelNodeAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuOccupancyMaxActiveBlocksPerMultiprocessor(
+    numBlocks: *mut ::std::os::raw::c_int,
+    func: CUfunction,
+    blockSize: ::std::os::raw::c_int,
+    dynamicSMemSize: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuOccupancyAvailableDynamicSMemPerBlock(
+    dynamicSmemSize: *mut usize,
+    func: CUfunction,
+    numBlocks: ::std::os::raw::c_int,
+    blockSize: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetArray(
+    hTexRef: CUtexref,
+    hArray: CUarray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetMipmappedArray(
+    hTexRef: CUtexref,
+    hMipmappedArray: CUmipmappedArray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetAddress_v2(
+    ByteOffset: *mut usize,
+    hTexRef: CUtexref,
+    dptr: CUdeviceptr,
+    bytes: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetAddress2D_v3(
+    hTexRef: CUtexref,
+    desc: *const CUDA_ARRAY_DESCRIPTOR,
+    dptr: CUdeviceptr,
+    Pitch: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetFormat(
+    hTexRef: CUtexref,
+    fmt: CUarray_format,
+    NumPackedComponents: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetAddressMode(
+    hTexRef: CUtexref,
+    dim: ::std::os::raw::c_int,
+    am: CUaddress_mode,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetMipmapFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetMipmapLevelBias(hTexRef: CUtexref, bias: f32) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetMipmapLevelClamp(
+    hTexRef: CUtexref,
+    minMipmapLevelClamp: f32,
+    maxMipmapLevelClamp: f32,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetMaxAnisotropy(
+    hTexRef: CUtexref,
+    maxAniso: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetBorderColor(hTexRef: CUtexref, pBorderColor: *mut f32) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefSetFlags(hTexRef: CUtexref, Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetAddress_v2(pdptr: *mut CUdeviceptr, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetArray(phArray: *mut CUarray, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetMipmappedArray(
+    phMipmappedArray: *mut CUmipmappedArray,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetAddressMode(
+    pam: *mut CUaddress_mode,
+    hTexRef: CUtexref,
+    dim: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetFilterMode(pfm: *mut CUfilter_mode, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetFormat(
+    pFormat: *mut CUarray_format,
+    pNumChannels: *mut ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetMipmapFilterMode(
+    pfm: *mut CUfilter_mode,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetMipmapLevelBias(pbias: *mut f32, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetMipmapLevelClamp(
+    pminMipmapLevelClamp: *mut f32,
+    pmaxMipmapLevelClamp: *mut f32,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetMaxAnisotropy(
+    pmaxAniso: *mut ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetBorderColor(pBorderColor: *mut f32, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefGetFlags(
+    pFlags: *mut ::std::os::raw::c_uint,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefCreate(pTexRef: *mut CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexRefDestroy(hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSurfRefSetArray(
+    hSurfRef: CUsurfref,
+    hArray: CUarray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSurfRefGetArray(phArray: *mut CUarray, hSurfRef: CUsurfref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" 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()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexObjectDestroy(texObject: CUtexObject) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexObjectGetResourceDesc(
+    pResDesc: *mut CUDA_RESOURCE_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexObjectGetTextureDesc(
+    pTexDesc: *mut CUDA_TEXTURE_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuTexObjectGetResourceViewDesc(
+    pResViewDesc: *mut CUDA_RESOURCE_VIEW_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSurfObjectCreate(
+    pSurfObject: *mut CUsurfObject,
+    pResDesc: *const CUDA_RESOURCE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSurfObjectDestroy(surfObject: CUsurfObject) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuSurfObjectGetResourceDesc(
+    pResDesc: *mut CUDA_RESOURCE_DESC,
+    surfObject: CUsurfObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceCanAccessPeer(
+    canAccessPeer: *mut ::std::os::raw::c_int,
+    dev: CUdevice,
+    peerDev: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxEnablePeerAccess(
+    peerContext: CUcontext,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuCtxDisablePeerAccess(peerContext: CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuDeviceGetP2PAttribute(
+    value: *mut ::std::os::raw::c_int,
+    attrib: CUdevice_P2PAttribute,
+    srcDevice: CUdevice,
+    dstDevice: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsUnregisterResource(resource: CUgraphicsResource) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsSubResourceGetMappedArray(
+    pArray: *mut CUarray,
+    resource: CUgraphicsResource,
+    arrayIndex: ::std::os::raw::c_uint,
+    mipLevel: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsResourceGetMappedMipmappedArray(
+    pMipmappedArray: *mut CUmipmappedArray,
+    resource: CUgraphicsResource,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsResourceGetMappedPointer_v2(
+    pDevPtr: *mut CUdeviceptr,
+    pSize: *mut usize,
+    resource: CUgraphicsResource,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsResourceSetMapFlags_v2(
+    resource: CUgraphicsResource,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsMapResources(
+    count: ::std::os::raw::c_uint,
+    resources: *mut CUgraphicsResource,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGraphicsUnmapResources(
+    count: ::std::os::raw::c_uint,
+    resources: *mut CUgraphicsResource,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuGetExportTable(
+    ppExportTable: *mut *const ::std::os::raw::c_void,
+    pExportTableId: *const CUuuid,
+) -> CUresult {
+    r#impl::export_table::get(ppExportTable, pExportTableId)
+}
+
+#[no_mangle]
+#[cfg(not(test))]
+pub extern "C" fn cuFuncGetModule(hmod: *mut CUmodule, hfunc: CUfunction) -> CUresult {
+    r#impl::unimplemented()
+}
diff --git a/notcuda/src/cuda/mod.rs b/notcuda/src/cuda_impl/mod.rs
index 63b9049..63b9049 100644
--- a/notcuda/src/cuda/mod.rs
+++ b/notcuda/src/cuda_impl/mod.rs
diff --git a/notcuda/src/cuda/rt.rs b/notcuda/src/cuda_impl/rt.rs
index 3931bc3..3931bc3 100644
--- a/notcuda/src/cuda/rt.rs
+++ b/notcuda/src/cuda_impl/rt.rs
diff --git a/notcuda/src/impl/context.rs b/notcuda/src/impl/context.rs
new file mode 100644
index 0000000..91d4460
--- /dev/null
+++ b/notcuda/src/impl/context.rs
@@ -0,0 +1,339 @@
+use super::CUresult;
+use super::{device, HasLivenessCookie, LiveCheck};
+use crate::{cuda::CUcontext, cuda_impl};
+use l0::sys::ze_result_t;
+use std::mem::{self, ManuallyDrop};
+use std::{
+    cell::RefCell,
+    num::NonZeroU32,
+    os::raw::c_uint,
+    ptr,
+    sync::{atomic::AtomicU32, Mutex},
+};
+
+thread_local! {
+    pub static CONTEXT_STACK: RefCell<Vec<*const Context>> = RefCell::new(Vec::new());
+}
+
+pub type Context = LiveCheck<ContextData>;
+
+impl HasLivenessCookie for ContextData {
+    #[cfg(target_pointer_width = "64")]
+    const COOKIE: usize = 0x5f0119560b643ffb;
+
+    #[cfg(target_pointer_width = "32")]
+    const COOKIE: usize = 0x0b643ffb;
+}
+
+enum ContextRefCount {
+    Primary,
+    NonPrimary(NonZeroU32),
+}
+
+impl ContextRefCount {
+    fn new(is_primary: bool) -> Self {
+        if is_primary {
+            ContextRefCount::Primary
+        } else {
+            ContextRefCount::NonPrimary(unsafe { NonZeroU32::new_unchecked(1) })
+        }
+    }
+
+    fn incr(&mut self) -> Result<(), CUresult> {
+        match self {
+            ContextRefCount::Primary => Ok(()),
+            ContextRefCount::NonPrimary(c) => {
+                let (new_count, overflow) = c.get().overflowing_add(1);
+                if overflow {
+                    Err(CUresult::CUDA_ERROR_INVALID_VALUE)
+                } else {
+                    *c = unsafe { NonZeroU32::new_unchecked(new_count) };
+                    Ok(())
+                }
+            }
+        }
+    }
+
+    #[must_use]
+    fn decr(&mut self) -> bool {
+        match self {
+            ContextRefCount::Primary => false,
+            ContextRefCount::NonPrimary(c) => {
+                if c.get() == 1 {
+                    return true;
+                }
+                *c = unsafe { NonZeroU32::new_unchecked(c.get() - 1) };
+                false
+            }
+        }
+    }
+
+    fn is_primary(&self) -> bool {
+        match self {
+            ContextRefCount::Primary => true,
+            ContextRefCount::NonPrimary(_) => false,
+        }
+    }
+}
+
+pub struct ContextData {
+    pub flags: AtomicU32,
+    pub device_index: device::Index,
+    // This pointer is null only for a moment when constructing primary context
+    pub device: *const Mutex<device::Device>,
+    // The split between mutable / non-mutable is mainly to avoid recursive locking in cuDevicePrimaryCtxGetState
+    pub mutable: Mutex<ContextDataMutable>,
+}
+
+pub struct ContextDataMutable {
+    ref_count: ContextRefCount,
+    pub cuda_manager: *mut cuda_impl::rt::ContextStateManager,
+    pub cuda_state: *mut cuda_impl::rt::ContextState,
+    pub cuda_dtor_cb: Option<
+        extern "C" fn(
+            CUcontext,
+            *mut cuda_impl::rt::ContextStateManager,
+            *mut cuda_impl::rt::ContextState,
+        ),
+    >,
+}
+
+impl ContextData {
+    pub fn new(
+        flags: c_uint,
+        is_primary: bool,
+        dev_index: device::Index,
+        dev: *const Mutex<device::Device>,
+    ) -> Self {
+        ContextData {
+            flags: AtomicU32::new(flags),
+            device_index: dev_index,
+            device: dev,
+            mutable: Mutex::new(ContextDataMutable {
+                ref_count: ContextRefCount::new(is_primary),
+                cuda_manager: ptr::null_mut(),
+                cuda_state: ptr::null_mut(),
+                cuda_dtor_cb: None,
+            }),
+        }
+    }
+}
+
+pub fn create_v2(pctx: *mut *mut Context, flags: u32, dev_idx: device::Index) -> CUresult {
+    if pctx == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let dev = device::get_device_ref(dev_idx);
+    let dev = match dev {
+        Ok(d) => d,
+        Err(e) => return e,
+    };
+    let mut ctx = Box::new(LiveCheck::new(ContextData::new(flags, false, dev_idx, dev)));
+    let ctx_ref = ctx.as_mut() as *mut Context;
+    unsafe { *pctx = ctx_ref };
+    mem::forget(ctx);
+    CONTEXT_STACK.with(|stack| stack.borrow_mut().push(ctx_ref));
+    CUresult::CUDA_SUCCESS
+}
+
+pub fn destroy_v2(ctx: *mut Context) -> CUresult {
+    if ctx == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    CONTEXT_STACK.with(|stack| {
+        let mut stack = stack.borrow_mut();
+        let should_pop = match stack.last() {
+            Some(active_ctx) => *active_ctx == (ctx as *const _),
+            None => false,
+        };
+        if should_pop {
+            stack.pop();
+        }
+    });
+    let mut ctx_box = ManuallyDrop::new(unsafe { Box::from_raw(ctx) });
+    if !ctx_box.try_drop() {
+        CUresult::CUDA_ERROR_INVALID_CONTEXT
+    } else {
+        unsafe { ManuallyDrop::drop(&mut ctx_box) };
+        CUresult::CUDA_SUCCESS
+    }
+}
+
+pub fn pop_current_v2(pctx: *mut *mut Context) -> CUresult {
+    if pctx == ptr::null_mut() {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let mut ctx = CONTEXT_STACK.with(|stack| stack.borrow_mut().pop());
+    let ctx_ptr = match &mut ctx {
+        Some(ctx) => *ctx as *mut _,
+        None => return CUresult::CUDA_ERROR_INVALID_CONTEXT,
+    };
+    unsafe { *pctx = ctx_ptr };
+    CUresult::CUDA_SUCCESS
+}
+
+pub fn with_current<F: FnOnce(&ContextData) -> R, R>(f: F) -> Result<R, CUresult> {
+    CONTEXT_STACK.with(|stack| {
+        stack
+            .borrow()
+            .last()
+            .and_then(|c| unsafe { &**c }.as_ref())
+            .ok_or(CUresult::CUDA_ERROR_INVALID_CONTEXT)
+            .map(f)
+    })
+}
+
+pub fn get_current(pctx: *mut *mut Context) -> l0::Result<()> {
+    if pctx == ptr::null_mut() {
+        return Err(ze_result_t::ZE_RESULT_ERROR_INVALID_ARGUMENT);
+    }
+    let ctx = CONTEXT_STACK.with(|stack| match stack.borrow().last() {
+        Some(ctx) => *ctx as *mut _,
+        None => ptr::null_mut(),
+    });
+    unsafe { *pctx = ctx };
+    Ok(())
+}
+
+pub fn set_current(ctx: *mut Context) -> CUresult {
+    if ctx == ptr::null_mut() {
+        CONTEXT_STACK.with(|stack| stack.borrow_mut().pop());
+        CUresult::CUDA_SUCCESS
+    } else {
+        CONTEXT_STACK.with(|stack| stack.borrow_mut().push(ctx));
+        CUresult::CUDA_SUCCESS
+    }
+}
+
+pub fn get_api_version(ctx: *mut Context, version: *mut u32) -> CUresult {
+    let _ctx = match unsafe { ctx.as_mut() } {
+        None => return CUresult::CUDA_ERROR_INVALID_VALUE,
+        Some(ctx) => match ctx.as_mut() {
+            None => return CUresult::CUDA_ERROR_INVALID_CONTEXT,
+            Some(ctx) => ctx,
+        },
+    };
+    //TODO: query device for properties roughly matching CUDA API version
+    unsafe { *version = 1100 };
+    CUresult::CUDA_SUCCESS
+}
+
+pub fn get_device(dev: *mut device::Index) -> CUresult {
+    let dev_idx = with_current(|ctx| ctx.device_index);
+    match dev_idx {
+        Ok(idx) => {
+            unsafe { *dev = idx }
+            CUresult::CUDA_SUCCESS
+        }
+        Err(err) => err,
+    }
+}
+
+#[cfg(test)]
+pub fn is_context_stack_empty() -> bool {
+    CONTEXT_STACK.with(|stack| stack.borrow().is_empty())
+}
+
+#[cfg(test)]
+mod tests {
+    use super::super::test::CudaDriverFns;
+    use super::super::CUresult;
+    use std::{ffi::c_void, ptr};
+
+    cuda_driver_test!(destroy_leaves_zombie_context);
+
+    fn destroy_leaves_zombie_context<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx1 = ptr::null_mut();
+        let mut ctx2 = ptr::null_mut();
+        let mut ctx3 = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx1, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx2, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx3, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxDestroy_v2(ctx2), CUresult::CUDA_SUCCESS);
+        let mut popped_ctx1 = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut popped_ctx1),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(popped_ctx1, ctx3);
+        let mut popped_ctx2 = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut popped_ctx2),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(popped_ctx2, ctx2);
+        let mut popped_ctx3 = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut popped_ctx3),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(popped_ctx3, ctx1);
+        let mut temp = 0;
+        assert_eq!(
+            T::cuCtxGetApiVersion(ctx2, &mut temp),
+            CUresult::CUDA_ERROR_INVALID_CONTEXT
+        );
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut ptr::null_mut()),
+            CUresult::CUDA_ERROR_INVALID_CONTEXT
+        );
+    }
+
+    cuda_driver_test!(empty_pop_fails);
+
+    fn empty_pop_fails<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut ctx),
+            CUresult::CUDA_ERROR_INVALID_CONTEXT
+        );
+    }
+
+    cuda_driver_test!(destroy_pops_top_of_stack);
+
+    fn destroy_pops_top_of_stack<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx1 = ptr::null_mut();
+        let mut ctx2 = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx1, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx2, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxDestroy_v2(ctx2), CUresult::CUDA_SUCCESS);
+        let mut popped_ctx1 = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut popped_ctx1),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(popped_ctx1, ctx1);
+        let mut popped_ctx2 = ptr::null_mut();
+        assert_eq!(
+            T::cuCtxPopCurrent_v2(&mut popped_ctx2),
+            CUresult::CUDA_ERROR_INVALID_CONTEXT
+        );
+    }
+
+    cuda_driver_test!(double_destroy_fails);
+
+    fn double_destroy_fails<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+        let destroy_result = T::cuCtxDestroy_v2(ctx);
+        // original CUDA impl returns randomly one or the other
+        assert!(
+            destroy_result == CUresult::CUDA_ERROR_INVALID_CONTEXT
+                || destroy_result == CUresult::CUDA_ERROR_CONTEXT_IS_DESTROYED
+        );
+    }
+
+    cuda_driver_test!(no_current_on_init);
+
+    fn no_current_on_init<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx = 1 as *mut c_void;
+        assert_eq!(T::cuCtxGetCurrent(&mut ctx), CUresult::CUDA_SUCCESS);
+        assert_eq!(ctx, ptr::null_mut());
+    }
+}
diff --git a/notcuda/src/impl/device.rs b/notcuda/src/impl/device.rs
new file mode 100644
index 0000000..8a8f2f8
--- /dev/null
+++ b/notcuda/src/impl/device.rs
@@ -0,0 +1,330 @@
+use super::{context, CUresult, Error};
+use crate::cuda;
+use cuda::{CUdevice_attribute, CUuuid_st};
+use std::{
+    cmp, mem,
+    os::raw::{c_char, c_int},
+    ptr,
+    sync::{
+        atomic::{AtomicU32, Ordering},
+        Mutex, MutexGuard,
+    },
+};
+
+static mut DEVICES: Option<Vec<Mutex<Device>>> = None;
+
+#[repr(transparent)]
+#[derive(Clone, Copy)]
+pub struct Index(pub c_int);
+
+pub struct Device {
+    pub base: l0::Device,
+    pub default_queue: l0::CommandQueue,
+    pub l0_context: l0::Context,
+    pub primary_context: context::Context,
+    properties: Option<Box<l0::sys::ze_device_properties_t>>,
+    image_properties: Option<Box<l0::sys::ze_device_image_properties_t>>,
+    memory_properties: Option<Vec<l0::sys::ze_device_memory_properties_t>>,
+}
+
+unsafe impl Send for Device {}
+
+impl Device {
+    // Unsafe because it does not fully initalize primary_context
+    unsafe fn new(drv: &l0::Driver, d: l0::Device, idx: usize) -> l0::Result<Self> {
+        let mut ctx = l0::Context::new(drv)?;
+        let queue = l0::CommandQueue::new(&mut ctx, &d)?;
+        let primary_context = context::Context::new(context::ContextData::new(
+            0,
+            true,
+            Index(idx as c_int),
+            ptr::null(),
+        ));
+        Ok(Self {
+            base: d,
+            default_queue: queue,
+            l0_context: ctx,
+            primary_context: primary_context,
+            properties: None,
+            image_properties: None,
+            memory_properties: None,
+        })
+    }
+
+    fn get_properties<'a>(&'a mut self) -> l0::Result<&'a l0::sys::ze_device_properties_t> {
+        if let Some(ref prop) = self.properties {
+            return Ok(prop);
+        }
+        match self.base.get_properties() {
+            Ok(prop) => Ok(self.properties.get_or_insert(prop)),
+            Err(e) => Err(e),
+        }
+    }
+
+    fn get_image_properties(&mut self) -> l0::Result<&l0::sys::ze_device_image_properties_t> {
+        if let Some(ref prop) = self.image_properties {
+            return Ok(prop);
+        }
+        match self.base.get_image_properties() {
+            Ok(prop) => Ok(self.image_properties.get_or_insert(prop)),
+            Err(e) => Err(e),
+        }
+    }
+
+    fn get_memory_properties(&mut self) -> l0::Result<&[l0::sys::ze_device_memory_properties_t]> {
+        if let Some(ref prop) = self.memory_properties {
+            return Ok(prop);
+        }
+        match self.base.get_memory_properties() {
+            Ok(prop) => Ok(self.memory_properties.get_or_insert(prop)),
+            Err(e) => Err(e),
+        }
+    }
+}
+
+pub fn init(driver: &l0::Driver) -> l0::Result<()> {
+    let ze_devices = driver.devices()?;
+    let mut devices = ze_devices
+        .into_iter()
+        .enumerate()
+        .map(|(idx, d)| unsafe { Device::new(driver, d, idx) }.map(Mutex::new))
+        .collect::<Result<Vec<_>, _>>()?;
+    for d in devices.iter_mut() {
+        d.get_mut()
+            .unwrap()
+            .primary_context
+            .as_mut()
+            .unwrap()
+            .device = d;
+    }
+    unsafe { DEVICES = Some(devices) };
+    Ok(())
+}
+
+fn devices() -> Result<&'static Vec<Mutex<Device>>, CUresult> {
+    match unsafe { &DEVICES } {
+        Some(devs) => Ok(devs),
+        None => Err(CUresult::CUDA_ERROR_NOT_INITIALIZED),
+    }
+}
+
+pub fn get_device_ref(Index(dev_idx): Index) -> Result<&'static Mutex<Device>, CUresult> {
+    let devs = devices()?;
+    if dev_idx < 0 || dev_idx >= devs.len() as c_int {
+        return Err(CUresult::CUDA_ERROR_INVALID_DEVICE);
+    }
+    Ok(&devs[dev_idx as usize])
+}
+
+pub fn get_device(dev_idx: Index) -> Result<MutexGuard<'static, Device>, CUresult> {
+    let dev = get_device_ref(dev_idx)?;
+    dev.lock().map_err(|_| CUresult::CUDA_ERROR_ILLEGAL_STATE)
+}
+
+pub fn get_count(count: *mut c_int) -> CUresult {
+    let len = devices().map(|d| d.len());
+    match len {
+        Ok(len) => {
+            unsafe { *count = len as c_int };
+            CUresult::CUDA_SUCCESS
+        }
+        Err(e) => e,
+    }
+}
+
+pub fn get(device: *mut Index, ordinal: c_int) -> CUresult {
+    if device == ptr::null_mut() || ordinal < 0 {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let len = devices().map(|d| d.len());
+    match len {
+        Ok(len) if ordinal < (len as i32) => {
+            unsafe { *device = Index(ordinal) };
+            CUresult::CUDA_SUCCESS
+        }
+        Ok(_) => CUresult::CUDA_ERROR_INVALID_VALUE,
+        Err(e) => e,
+    }
+}
+
+pub fn get_name(name: *mut c_char, len: i32, dev: Index) -> Result<(), CUresult> {
+    if name == ptr::null_mut() || len < 0 {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    // This is safe because devices are 'static
+    let name_ptr = {
+        let mut dev = get_device(dev)?;
+        let props = dev.get_properties().map_err(Into::<CUresult>::into)?;
+        props.name.as_ptr()
+    };
+    let name_len = (0..256)
+        .position(|i| unsafe { *name_ptr.add(i) } == 0)
+        .unwrap_or(256);
+    let dst_null_pos = cmp::min((len - 1) as usize, name_len);
+    unsafe { std::ptr::copy_nonoverlapping(name_ptr, name, dst_null_pos) };
+    unsafe { *(name.add(dst_null_pos)) = 0 };
+    Ok(())
+}
+
+unsafe fn transmute_lifetime<'a, 'b, T: ?Sized>(t: &'a T) -> &'b T {
+    mem::transmute(t)
+}
+
+pub fn total_mem_v2(bytes: *mut usize, dev: Index) -> Result<(), CUresult> {
+    if bytes == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    // This is safe because devices are 'static
+    let mem_props = {
+        let mut dev = get_device(dev)?;
+        unsafe {
+            transmute_lifetime(
+                dev.get_memory_properties()
+                    .map_err(Into::<CUresult>::into)?,
+            )
+        }
+    };
+    let max_mem = mem_props
+        .iter()
+        .map(|p| p.totalSize)
+        .max()
+        .ok_or(CUresult::CUDA_ERROR_ILLEGAL_STATE)?;
+    unsafe { *bytes = max_mem as usize };
+    Ok(())
+}
+
+impl CUdevice_attribute {
+    fn get_static_value(self) -> Option<i32> {
+        match self {
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_GPU_OVERLAP => Some(1),
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT => Some(1),
+            // TODO: fix this for DG1
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_INTEGRATED => Some(1),
+            // TODO: go back to this once we have more funcitonality implemented
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR => Some(8),
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR => Some(0),
+            CUdevice_attribute::CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY => Some(1),
+            _ => None,
+        }
+    }
+}
+
+pub fn get_attribute(pi: *mut i32, attrib: CUdevice_attribute, dev: Index) -> Result<(), Error> {
+    if pi == ptr::null_mut() {
+        return Err(Error::Cuda(CUresult::CUDA_ERROR_INVALID_VALUE));
+    }
+    if let Some(value) = attrib.get_static_value() {
+        unsafe { *pi = value };
+        return Ok(());
+    }
+    let mut dev = get_device(dev).map_err(Error::Cuda)?;
+    let value = match attrib {
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT => {
+            dev.get_properties().map_err(Error::L0)?.maxHardwareContexts as i32
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT => {
+            let props = dev.get_properties().map_err(Error::L0)?;
+            (props.numSlices * props.numSubslicesPerSlice * props.numEUsPerSubslice) as i32
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH => cmp::min(
+            dev.get_image_properties()
+                .map_err(Error::L0)?
+                .maxImageDims1D,
+            c_int::max_value() as u32,
+        ) as c_int,
+        _ => {
+            // TODO: support more attributes for CUDA runtime
+            /*
+            return Err(l0::Error(
+                l0::sys::ze_result_t::ZE_RESULT_ERROR_UNSUPPORTED_FEATURE,
+            ))
+            */
+            return Ok(());
+        }
+    };
+    unsafe { *pi = value };
+    Ok(())
+}
+
+pub fn get_uuid(uuid: *mut CUuuid_st, dev: Index) -> Result<(), Error> {
+    let ze_uuid = {
+        get_device(dev)
+            .map_err(Error::Cuda)?
+            .get_properties()
+            .map_err(Error::L0)?
+            .uuid
+    };
+    unsafe {
+        *uuid = CUuuid_st {
+            bytes: mem::transmute(ze_uuid.id),
+        }
+    };
+    Ok(())
+}
+
+pub fn with_current_exclusive<F: FnOnce(&mut Device) -> R, R>(f: F) -> Result<R, CUresult> {
+    let dev = super::context::with_current(|ctx| ctx.device);
+    dev.and_then(|dev| {
+        unsafe { &*dev }
+            .try_lock()
+            .map(|mut dev| f(&mut dev))
+            .map_err(|_| CUresult::CUDA_ERROR_ILLEGAL_STATE)
+    })
+}
+
+pub fn with_exclusive<F: FnOnce(&mut Device) -> R, R>(dev: Index, f: F) -> Result<R, CUresult> {
+    let dev = get_device_ref(dev)?;
+    dev.try_lock()
+        .map(|mut dev| f(&mut dev))
+        .map_err(|_| CUresult::CUDA_ERROR_ILLEGAL_STATE)
+}
+
+pub fn primary_ctx_get_state(
+    idx: Index,
+    flags: *mut u32,
+    active: *mut i32,
+) -> Result<(), CUresult> {
+    let (ctx_ptr, flags_ptr) = with_exclusive(idx, |dev| {
+        // This is safe because primary context can't be dropped
+        let ctx_ptr = &dev.primary_context as *const _;
+        let flags_ptr =
+            (&unsafe { dev.primary_context.as_ref_unchecked() }.flags) as *const AtomicU32;
+        (ctx_ptr, flags_ptr)
+    })?;
+    let is_active = context::CONTEXT_STACK
+        .with(|stack| stack.borrow().last().map(|x| *x))
+        .map(|current| current == ctx_ptr)
+        .unwrap_or(false);
+    let flags_value = unsafe { &*flags_ptr }.load(Ordering::Relaxed);
+    unsafe { *flags = flags_value };
+    unsafe { *active = if is_active { 1 } else { 0 } };
+    Ok(())
+}
+
+pub fn primary_ctx_retain(pctx: *mut *mut context::Context, dev: Index) -> Result<(), CUresult> {
+    let ctx_ptr = with_exclusive(dev, |dev| &mut dev.primary_context as *mut _)?;
+    unsafe { *pctx = ctx_ptr };
+    Ok(())
+}
+
+#[cfg(test)]
+mod tests {
+    use super::super::test::CudaDriverFns;
+    use super::super::CUresult;
+    use crate::cuda::CUuuid;
+    use std::{ffi::c_void, mem, ptr};
+
+    cuda_driver_test!(primary_ctx_default_inactive);
+
+    fn primary_ctx_default_inactive<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut flags = u32::max_value();
+        let mut active = i32::max_value();
+        assert_eq!(
+            T::cuDevicePrimaryCtxGetState(0, &mut flags, &mut active),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(flags, 0);
+        assert_eq!(active, 0);
+    }
+}
diff --git a/notcuda/src/export_table.rs b/notcuda/src/impl/export_table.rs
index b8b0cb7..3df07ed 100644
--- a/notcuda/src/export_table.rs
+++ b/notcuda/src/impl/export_table.rs
@@ -1,36 +1,42 @@
-use crate::cu;
-use crate::{cuda, ze};
+use crate::cuda::CUresult;
+use crate::{
+    cuda::{CUcontext, CUdevice, CUmodule, CUuuid},
+    cuda_impl,
+};
 
+use super::{context, device, Decuda, Encuda};
 use std::mem;
-use std::os::raw::{c_int, c_uint, c_ulong, c_ushort};
+use std::os::raw::{c_uint, c_ulong, c_ushort};
 use std::{ffi::c_void, ptr, slice};
 
-#[no_mangle]
-pub unsafe extern "C" fn cuGetExportTable(
-    table: *mut *const std::os::raw::c_void,
-    id: *const cu::Uuid,
-) -> cu::Result {
+pub fn get(table: *mut *const std::os::raw::c_void, id: *const CUuuid) -> CUresult {
     if table == ptr::null_mut() || id == ptr::null_mut() {
-        cu::Result::ERROR_INVALID_VALUE
-    } else if *id == TOOLS_RUNTIME_CALLBACK_HOOKS_GUID {
-        *table = TOOLS_RUNTIME_CALLBACK_HOOKS_VTABLE.as_ptr() as *const _;
-        cu::Result::SUCCESS
-    } else if *id == CUDART_INTERFACE_GUID {
-        *table = CUDART_INTERFACE_VTABLE.as_ptr() as *const _;
-        cu::Result::SUCCESS
-    } else if *id == TOOLS_TLS_GUID {
-        *table = 1 as _;
-        cu::Result::SUCCESS
-    } else if *id == CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_GUID {
-        *table = CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_VTABLE.as_ptr() as *const _;
-        cu::Result::SUCCESS
-    } else {
-        cu::Result::ERROR_NOT_SUPPORTED
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let id = unsafe { *id };
+    match id {
+        TOOLS_RUNTIME_CALLBACK_HOOKS_GUID => {
+            unsafe { *table = TOOLS_RUNTIME_CALLBACK_HOOKS_VTABLE.as_ptr() as *const _ };
+            CUresult::CUDA_SUCCESS
+        }
+        CUDART_INTERFACE_GUID => {
+            unsafe { *table = CUDART_INTERFACE_VTABLE.as_ptr() as *const _ };
+            CUresult::CUDA_SUCCESS
+        }
+        TOOLS_TLS_GUID => {
+            unsafe { *table = 1 as _ };
+            CUresult::CUDA_SUCCESS
+        }
+        CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_GUID => {
+            unsafe { *table = CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_VTABLE.as_ptr() as *const _ };
+            CUresult::CUDA_SUCCESS
+        }
+        _ => CUresult::CUDA_ERROR_NOT_SUPPORTED,
     }
 }
 
-const TOOLS_RUNTIME_CALLBACK_HOOKS_GUID: cu::Uuid = cu::Uuid {
-    x: [
+const TOOLS_RUNTIME_CALLBACK_HOOKS_GUID: CUuuid = CUuuid {
+    bytes: [
         0xa0, 0x94, 0x79, 0x8c, 0x2e, 0x74, 0x2e, 0x74, 0x93, 0xf2, 0x08, 0x00, 0x20, 0x0c, 0x0a,
         0x66,
     ],
@@ -73,8 +79,8 @@ unsafe extern "C" fn runtime_callback_hooks_fn5(ptr: *mut *mut u8, size: *mut us
     return TOOLS_RUNTIME_CALLBACK_HOOKS_FN5_SPACE.as_mut_ptr();
 }
 
-const CUDART_INTERFACE_GUID: cu::Uuid = cu::Uuid {
-    x: [
+const CUDART_INTERFACE_GUID: CUuuid = CUuuid {
+    bytes: [
         0x6b, 0xd5, 0xfb, 0x6c, 0x5b, 0xf4, 0xe7, 0x4a, 0x89, 0x87, 0xd9, 0x39, 0x12, 0xfd, 0x9d,
         0xf9,
     ],
@@ -102,8 +108,17 @@ static CUDART_INTERFACE_VTABLE: [VTableEntry; CUDART_INTERFACE_LENGTH] = [
     VTableEntry { ptr: ptr::null() },
 ];
 
-unsafe extern "C" fn cudart_interface_fn1(_: *mut c_ulong, _: c_int) -> c_int {
-    0
+unsafe extern "C" fn cudart_interface_fn1(pctx: *mut CUcontext, dev: CUdevice) -> CUresult {
+    cudart_interface_fn1_impl(pctx.decuda(), dev.decuda()).encuda()
+}
+
+fn cudart_interface_fn1_impl(
+    pctx: *mut *mut context::Context,
+    dev: device::Index,
+) -> Result<(), CUresult> {
+    let ctx_ptr = device::with_exclusive(dev, |d| &mut d.primary_context as *mut context::Context)?;
+    unsafe { *pctx = ctx_ptr };
+    Ok(())
 }
 
 /*
@@ -178,17 +193,20 @@ struct FatbinFileHeader {
 }
 
 unsafe extern "C" fn get_module_from_cubin(
-    result: *mut cu::Module,
+    result: *mut CUmodule,
     fatbinc_wrapper: *const FatbincWrapper,
     _: *mut c_void,
     _: *mut c_void,
-) -> cu::Result {
-    if result == ptr::null_mut() || (*fatbinc_wrapper).magic != FATBINC_MAGIC || (*fatbinc_wrapper).version != FATBINC_VERSION {
-        return cu::Result::ERROR_INVALID_VALUE;
+) -> CUresult {
+    if result == ptr::null_mut()
+        || (*fatbinc_wrapper).magic != FATBINC_MAGIC
+        || (*fatbinc_wrapper).version != FATBINC_VERSION
+    {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
     }
     let fatbin_header = (*fatbinc_wrapper).data;
     if (*fatbin_header).magic != FATBIN_MAGIC || (*fatbin_header).version != FATBIN_VERSION {
-        return cu::Result::ERROR_INVALID_VALUE;
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
     }
     let file = (fatbin_header as *const u8).add((*fatbin_header).header_size as usize);
     let end = file.add((*fatbin_header).files_size as usize);
@@ -201,13 +219,9 @@ unsafe extern "C" fn get_module_from_cubin(
         );
         let kernel_text =
             lz4::block::decompress(slice, Some((*file).uncompressed_payload as i32)).unwrap();
-        let module = ze::Module {
-            ptx_text: kernel_text,
-        };
-        *result = cu::Module::new(module);
-        return cu::Result::SUCCESS
+        return CUresult::CUDA_SUCCESS;
     }
-    cu::Result::ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE
+    CUresult::CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE
 }
 
 unsafe fn get_ptx_files(file: *const u8, end: *const u8) -> Vec<*const FatbinFileHeader> {
@@ -227,15 +241,15 @@ unsafe fn get_ptx_files(file: *const u8, end: *const u8) -> Vec<*const FatbinFil
 
 unsafe extern "C" fn cudart_interface_fn6(_: u64) {}
 
-const TOOLS_TLS_GUID: cu::Uuid = cu::Uuid {
-    x: [
+const TOOLS_TLS_GUID: CUuuid = CUuuid {
+    bytes: [
         0x42, 0xd8, 0x5a, 0x81, 0x23, 0xf6, 0xcb, 0x47, 0x82, 0x98, 0xf6, 0xe7, 0x8a, 0x3a, 0xec,
         0xdc,
     ],
 };
 
-const CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_GUID: cu::Uuid = cu::Uuid {
-    x: [
+const CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_GUID: CUuuid = CUuuid {
+    bytes: [
         0xc6, 0x93, 0x33, 0x6e, 0x11, 0x21, 0xdf, 0x11, 0xa8, 0xc3, 0x68, 0xf3, 0x55, 0xd8, 0x95,
         0x93,
     ],
@@ -257,23 +271,50 @@ static CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_VTABLE: [VTableEntry; 4] = [
 
 // some kind of ctor
 unsafe extern "C" fn context_local_storage_ctor(
-    cu_ctx: cu::Context, // always zero
-    mgr: *mut cuda::rt::ContextStateManager,
-    ctx_state: *mut cuda::rt::ContextState,
+    cu_ctx: CUcontext, // always zero
+    mgr: *mut cuda_impl::rt::ContextStateManager,
+    ctx_state: *mut cuda_impl::rt::ContextState,
     // clsContextDestroyCallback,  have to be called on cuDevicePrimaryCtxReset
-    dtor_cb: extern "C" fn(
-        cu::Context,
-        *mut cuda::rt::ContextStateManager,
-        *mut cuda::rt::ContextState,
-    ),
-) -> cu::Result {
-    if cu_ctx.0 == ptr::null_mut() {
-        return cu::Result::ERROR_NOT_SUPPORTED;
+    dtor_cb: Option<
+        extern "C" fn(
+            CUcontext,
+            *mut cuda_impl::rt::ContextStateManager,
+            *mut cuda_impl::rt::ContextState,
+        ),
+    >,
+) -> CUresult {
+    context_local_storage_ctor_impl(cu_ctx.decuda(), mgr, ctx_state, dtor_cb).encuda()
+}
+
+fn context_local_storage_ctor_impl(
+    cu_ctx: *mut context::Context,
+    mgr: *mut cuda_impl::rt::ContextStateManager,
+    ctx_state: *mut cuda_impl::rt::ContextState,
+    dtor_cb: Option<
+        extern "C" fn(
+            CUcontext,
+            *mut cuda_impl::rt::ContextStateManager,
+            *mut cuda_impl::rt::ContextState,
+        ),
+    >,
+) -> Result<(), CUresult> {
+    if cu_ctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_NOT_SUPPORTED);
     }
-    (*cu_ctx.0).cuda_manager = mgr;
-    (*cu_ctx.0).cuda_state = ctx_state;
-    (*cu_ctx.0).cuda_dtor_cb = dtor_cb;
-    cu::Result::SUCCESS
+    unsafe { &*cu_ctx }
+        .as_ref()
+        .ok_or(CUresult::CUDA_ERROR_INVALID_CONTEXT)
+        .and_then(|ctx| {
+            ctx.mutable
+                .try_lock()
+                .map_err(|_| CUresult::CUDA_ERROR_ILLEGAL_STATE)
+                .map(|mut mutable| {
+                    mutable.cuda_manager = mgr;
+                    mutable.cuda_state = ctx_state;
+                    mutable.cuda_dtor_cb = dtor_cb;
+                })
+        })?;
+    Ok(())
 }
 
 // some kind of dtor
@@ -282,13 +323,30 @@ unsafe extern "C" fn context_local_storage_dtor(_: *mut usize, _: *mut ()) -> u3
 }
 
 unsafe extern "C" fn context_local_storage_get_state(
-    ctx_state: *mut *mut cuda::rt::ContextState,
-    cu_ctx: cu::Context,
-    _: *mut cuda::rt::ContextStateManager,
-) -> cu::Result {
-    if cu_ctx == cu::Context::null() {
-        return cu::Result::ERROR_INVALID_CONTEXT;
+    ctx_state: *mut *mut cuda_impl::rt::ContextState,
+    cu_ctx: CUcontext,
+    state_mgr: *mut cuda_impl::rt::ContextStateManager,
+) -> CUresult {
+    context_local_storage_get_state_impl(ctx_state, cu_ctx.decuda(), state_mgr).encuda()
+}
+
+fn context_local_storage_get_state_impl(
+    ctx_state: *mut *mut cuda_impl::rt::ContextState,
+    cu_ctx: *mut context::Context,
+    _: *mut cuda_impl::rt::ContextStateManager,
+) -> Result<(), CUresult> {
+    if cu_ctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
     }
-    *ctx_state = (*cu_ctx.0).cuda_state;
-    cu::Result::SUCCESS
+    let cuda_state = unsafe { &*cu_ctx }
+        .as_ref()
+        .ok_or(CUresult::CUDA_ERROR_INVALID_CONTEXT)
+        .and_then(|ctx| {
+            ctx.mutable
+                .try_lock()
+                .map_err(|_| CUresult::CUDA_ERROR_ILLEGAL_STATE)
+                .map(|mutable| mutable.cuda_state)
+        })?;
+    unsafe { *ctx_state = cuda_state };
+    Ok(())
 }
diff --git a/notcuda/src/impl/memory.rs b/notcuda/src/impl/memory.rs
new file mode 100644
index 0000000..52e269d
--- /dev/null
+++ b/notcuda/src/impl/memory.rs
@@ -0,0 +1,81 @@
+use super::CUresult;
+use std::ffi::c_void;
+
+pub fn alloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult {
+    let alloc_result = super::device::with_current_exclusive(|dev| unsafe {
+        dev.base.mem_alloc_device(&mut dev.l0_context, bytesize, 0)
+    });
+    match alloc_result {
+        Ok(Ok(alloc)) => {
+            unsafe { *dptr = alloc };
+            CUresult::CUDA_SUCCESS
+        }
+        Ok(Err(e)) => e.into(),
+        Err(e) => e,
+    }
+}
+
+pub fn copy_v2(
+    dst: *mut c_void,
+    src: *const c_void,
+    bytesize: usize,
+) -> Result<Result<(), l0::sys::ze_result_t>, CUresult> {
+    super::device::with_current_exclusive(|dev| unsafe {
+        memcpy_impl(
+            &mut dev.l0_context,
+            dst,
+            src,
+            bytesize,
+            &dev.base,
+            &mut dev.default_queue,
+        )
+    })
+}
+
+unsafe fn memcpy_impl(
+    ctx: &mut l0::Context,
+    dst: *mut c_void,
+    src: *const c_void,
+    bytes_count: usize,
+    dev: &l0::Device,
+    queue: &mut l0::CommandQueue,
+) -> l0::Result<()> {
+    let mut cmd_list = l0::CommandList::new(ctx, &dev)?;
+    cmd_list.append_memory_copy_unsafe(dst, src, bytes_count, None, &mut [])?;
+    queue.execute(cmd_list)?;
+    Ok(())
+}
+
+#[cfg(test)]
+mod tests {
+    use super::super::test::CudaDriverFns;
+    use super::super::CUresult;
+    use std::ptr;
+
+    cuda_driver_test!(alloc_without_ctx);
+
+    fn alloc_without_ctx<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut mem = ptr::null_mut();
+        assert_eq!(
+            T::cuMemAlloc_v2(&mut mem, std::mem::size_of::<usize>()),
+            CUresult::CUDA_ERROR_INVALID_CONTEXT
+        );
+        assert_eq!(mem, ptr::null_mut());
+    }
+
+    cuda_driver_test!(alloc_with_ctx);
+
+    fn alloc_with_ctx<T: CudaDriverFns>() {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx, 0, 0), CUresult::CUDA_SUCCESS);
+        let mut mem = ptr::null_mut();
+        assert_eq!(
+            T::cuMemAlloc_v2(&mut mem, std::mem::size_of::<usize>()),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_ne!(mem, ptr::null_mut());
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+    }
+}
diff --git a/notcuda/src/impl/mod.rs b/notcuda/src/impl/mod.rs
new file mode 100644
index 0000000..ded2ff4
--- /dev/null
+++ b/notcuda/src/impl/mod.rs
@@ -0,0 +1,234 @@
+use crate::cuda::{CUctx_st, CUdevice, CUdeviceptr, CUresult};
+use std::{ffi::c_void, mem::ManuallyDrop, os::raw::c_int, sync::Mutex};
+
+#[cfg(test)]
+#[macro_use]
+pub mod test;
+pub mod context;
+pub mod device;
+pub mod export_table;
+pub mod memory;
+
+#[cfg(debug_assertions)]
+pub fn unimplemented() -> CUresult {
+    unimplemented!()
+}
+
+#[cfg(not(debug_assertions))]
+pub fn unimplemented() -> CUresult {
+    CUresult::CUDA_ERROR_NOT_SUPPORTED
+}
+
+pub trait HasLivenessCookie {
+    const COOKIE: usize;
+}
+
+// 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> {
+    cookie: usize,
+    data: ManuallyDrop<T>,
+}
+
+impl<T: HasLivenessCookie> LiveCheck<T> {
+    pub fn new(data: T) -> Self {
+        LiveCheck {
+            cookie: T::COOKIE,
+            data: ManuallyDrop::new(data),
+        }
+    }
+
+    pub unsafe fn as_ref_unchecked(&self) -> &T {
+        &self.data
+    }
+
+    pub fn as_ref(&self) -> Option<&T> {
+        if self.cookie == T::COOKIE {
+            Some(&self.data)
+        } else {
+            None
+        }
+    }
+
+    pub fn as_mut(&mut self) -> Option<&mut T> {
+        if self.cookie == T::COOKIE {
+            Some(&mut self.data)
+        } else {
+            None
+        }
+    }
+
+    #[must_use]
+    pub fn try_drop(&mut self) -> bool {
+        if self.cookie == T::COOKIE {
+            self.cookie = 0;
+            unsafe { ManuallyDrop::drop(&mut self.data) };
+            return true;
+        }
+        false
+    }
+}
+
+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 From<l0::sys::ze_result_t> for CUresult {
+    fn from(result: l0::sys::ze_result_t) -> Self {
+        match result {
+            l0::sys::ze_result_t::ZE_RESULT_SUCCESS => CUresult::CUDA_SUCCESS,
+            l0_sys::ze_result_t::ZE_RESULT_ERROR_UNINITIALIZED => {
+                CUresult::CUDA_ERROR_NOT_INITIALIZED
+            }
+            l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ENUMERATION => {
+                CUresult::CUDA_ERROR_INVALID_VALUE
+            }
+            l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ARGUMENT => {
+                CUresult::CUDA_ERROR_INVALID_VALUE
+            }
+            l0_sys::ze_result_t::ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY => {
+                CUresult::CUDA_ERROR_OUT_OF_MEMORY
+            }
+            l0_sys::ze_result_t::ZE_RESULT_ERROR_UNSUPPORTED_FEATURE => {
+                CUresult::CUDA_ERROR_NOT_SUPPORTED
+            }
+            _ => CUresult::CUDA_ERROR_UNKNOWN,
+        }
+    }
+}
+
+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 l0::sys::ze_result_t {
+    type To = CUresult;
+    fn encuda(self: Self) -> Self::To {
+        self.into()
+    }
+}
+
+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(),
+        }
+    }
+}
+
+pub enum Error {
+    L0(l0::sys::ze_result_t),
+    Cuda(CUresult),
+}
+
+impl Encuda for Error {
+    type To = CUresult;
+    fn encuda(self: Self) -> Self::To {
+        match self {
+            Error::L0(e) => e.into(),
+            Error::Cuda(e) => e,
+        }
+    }
+}
+
+lazy_static! {
+    static ref GLOBAL_STATE: Mutex<Option<GlobalState>> = Mutex::new(None);
+}
+
+struct GlobalState {
+    driver: l0::Driver,
+}
+
+unsafe impl Send for GlobalState {}
+
+// TODO: implement
+fn is_intel_gpu_driver(_: &l0::Driver) -> bool {
+    true
+}
+
+pub fn init() -> l0::Result<()> {
+    let mut global_state = GLOBAL_STATE
+        .lock()
+        .map_err(|_| l0::sys::ze_result_t::ZE_RESULT_ERROR_UNKNOWN)?;
+    if global_state.is_some() {
+        return Ok(());
+    }
+    l0::init()?;
+    let drivers = l0::Driver::get()?;
+    let driver = match drivers.into_iter().find(is_intel_gpu_driver) {
+        None => return Err(l0::sys::ze_result_t::ZE_RESULT_ERROR_UNKNOWN),
+        Some(driver) => {
+            device::init(&driver)?;
+            driver
+        }
+    };
+    *global_state = Some(GlobalState { driver });
+    drop(global_state);
+    Ok(())
+}
+
+pub fn driver_get_version() -> c_int {
+    i32::max_value()
+}
+
+impl<'a> CudaRepr for CUctx_st {
+    type Impl = context::Context;
+}
+
+impl<'a> CudaRepr for CUdevice {
+    type Impl = device::Index;
+}
+
+impl Decuda<device::Index> for CUdevice {
+    fn decuda(self) -> device::Index {
+        device::Index(self.0)
+    }
+}
+
+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 _
+    }
+}
diff --git a/notcuda/src/impl/test.rs b/notcuda/src/impl/test.rs
new file mode 100644
index 0000000..0ad625b
--- /dev/null
+++ b/notcuda/src/impl/test.rs
@@ -0,0 +1,118 @@
+#![allow(non_snake_case)]
+
+use crate::r#impl as notcuda;
+use crate::r#impl::CUresult;
+use crate::{cuda::CUuuid, r#impl::Encuda};
+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 _notcuda>]() {
+                $func::<crate::r#impl::test::NotCuda>()
+            }
+
+            #[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;
+}
+
+pub struct NotCuda();
+
+impl CudaDriverFns for NotCuda {
+    fn cuInit(_flags: c_uint) -> CUresult {
+        assert!(notcuda::context::is_context_stack_empty());
+        notcuda::init().encuda()
+    }
+
+    fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult {
+        notcuda::context::create_v2(pctx as *mut _, flags, notcuda::device::Index(dev)).encuda()
+    }
+
+    fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult {
+        notcuda::context::destroy_v2(ctx as *mut _)
+    }
+
+    fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult {
+        notcuda::context::pop_current_v2(pctx as *mut _)
+    }
+
+    fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult {
+        notcuda::context::get_api_version(ctx as *mut _, version)
+    }
+
+    fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult {
+        notcuda::context::get_current(pctx as *mut _).encuda()
+    }
+    fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult {
+        notcuda::memory::alloc_v2(dptr as *mut _, bytesize)
+    }
+
+    fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult {
+        notcuda::device::get_uuid(uuid, notcuda::device::Index(dev)).encuda()
+    }
+
+    fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult {
+        notcuda::device::primary_ctx_get_state(notcuda::device::Index(dev), flags, active).encuda()
+    }
+}
+
+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) }
+    }
+}
diff --git a/notcuda/src/lib.rs b/notcuda/src/lib.rs
index 7e3b895..48f5f85 100644
--- a/notcuda/src/lib.rs
+++ b/notcuda/src/lib.rs
@@ -2,396 +2,14 @@ extern crate level_zero as l0;
 extern crate level_zero_sys as l0_sys;
 #[macro_use]
 extern crate lazy_static;
+#[cfg(test)]
+extern crate cuda_driver_sys;
 extern crate lz4;
+#[cfg(test)]
+#[macro_use]
+extern crate paste;
 
-use std::cell::RefCell;
-use std::convert::TryFrom;
-use std::os::raw::{c_char, c_int, c_uint};
-use std::ptr;
-use std::{ffi::c_void, sync::Mutex};
-
-mod cu;
+#[allow(warnings)]
 mod cuda;
-mod export_table;
-mod ze;
-
-thread_local! {
-    static CONTEXT_STACK: RefCell<Vec<cu::Context>> = RefCell::new(Vec::new());
-}
-
-lazy_static! {
-    static ref GLOBAL_STATE: Mutex<Option<Context>> = Mutex::new(None);
-}
-
-struct Context {
-    base: l0::Context,
-    devices: Vec<ze::Device>,
-}
-unsafe impl Send for Context {}
-unsafe impl Sync for Context {}
-
-impl Context {
-    fn new() -> l0::Result<Context> {
-        let driver = l0::Driver::get()?;
-        let ze_devices = driver[0].devices()?;
-        let mut ctx = l0::Context::new(&driver[0])?;
-        let devices = ze_devices
-            .into_iter()
-            .map(|d| ze::Device::new(&mut ctx, d))
-            .collect::<Result<_, _>>()?;
-        Ok(Context {
-            base: ctx,
-            devices: devices,
-        })
-    }
-
-    fn call<F: FnOnce(&mut Context) -> l0_sys::ze_result_t>(f: F) -> cu::Result {
-        let mut lock = GLOBAL_STATE.try_lock();
-        if let Ok(ref mut mutex) = lock {
-            match **mutex {
-                None => return cu::Result::ERROR_NOT_INITIALIZED,
-                Some(ref mut driver) => {
-                    return cu::Result::from_l0(f(driver));
-                }
-            }
-        } else {
-            return cu::Result::ERROR_UNKNOWN;
-        }
-    }
-
-    fn call2<F: FnOnce(&mut Context) -> l0::Result<()>>(f: F) -> cu::Result {
-        Context::call(|ctx| match f(ctx) {
-            Ok(()) => l0::sys::ze_result_t::ZE_RESULT_SUCCESS,
-            Err(err) => err.0,
-        })
-    }
-
-    fn call_device<F: FnOnce(&mut ze::Device) -> l0::Result<()>>(
-        cu::Device(dev): cu::Device,
-        f: F,
-    ) -> cu::Result {
-        if dev < 0 {
-            return cu::Result::ERROR_INVALID_VALUE;
-        }
-        let dev = dev as usize;
-        Context::call(|driver| {
-            if dev >= driver.devices.len() {
-                return l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ARGUMENT;
-            }
-            match f(&mut driver.devices[dev]) {
-                Ok(()) => l0::sys::ze_result_t::ZE_RESULT_SUCCESS,
-                Err(err) => err.0,
-            }
-        })
-    }
-
-    fn device_get_count(&self, count: *mut i32) -> l0_sys::ze_result_t {
-        unsafe { *count = self.devices.len() as i32 };
-        l0_sys::ze_result_t::ZE_RESULT_SUCCESS
-    }
-
-    fn device_get(&self, device: *mut cu::Device, ordinal: c_int) -> l0_sys::ze_result_t {
-        if (ordinal as usize) >= self.devices.len() {
-            return l0_sys::ze_result_t::ZE_RESULT_ERROR_INVALID_ARGUMENT;
-        }
-        unsafe { *device = cu::Device(ordinal) };
-        l0_sys::ze_result_t::ZE_RESULT_SUCCESS
-    }
-}
-
-#[no_mangle]
-pub unsafe extern "C" fn cuDriverGetVersion(version: *mut c_int) -> cu::Result {
-    if version == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    *version = i32::max_value();
-    return cu::Result::SUCCESS;
-}
-
-#[no_mangle]
-pub unsafe extern "C" fn cuInit(_: c_uint) -> cu::Result {
-    let l0_sys_init = l0_sys::zeInit(l0_sys::ze_init_flags_t::ZE_INIT_FLAG_GPU_ONLY);
-    if l0_sys_init != l0_sys::ze_result_t::ZE_RESULT_SUCCESS {
-        return cu::Result::from_l0(l0_sys_init);
-    }
-    let mut lock = GLOBAL_STATE.try_lock();
-    if let Ok(ref mut mutex) = lock {
-        if let None = **mutex {
-            match Context::new() {
-                Ok(state) => **mutex = Some(state),
-                Err(err) => return cu::Result::from_l0(err.0),
-            }
-        }
-    } else {
-        return cu::Result::ERROR_UNKNOWN;
-    }
-    cu::Result::SUCCESS
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceGetCount(count: *mut c_int) -> cu::Result {
-    if count == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call(|driver| driver.device_get_count(count))
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceGet(device: *mut cu::Device, ordinal: c_int) -> cu::Result {
-    if ordinal < 0 || device == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call(|driver| driver.device_get(device, ordinal))
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceGetName(
-    name: *mut c_char,
-    len: c_int,
-    dev_idx: cu::Device,
-) -> cu::Result {
-    if name == ptr::null_mut() || len <= 0 {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call_device(dev_idx, |dev| dev.get_name(name, len))
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceTotalMem_v2(bytes: *mut usize, dev_idx: cu::Device) -> cu::Result {
-    if bytes == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call_device(dev_idx, |dev| dev.total_mem(bytes))
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceGetAttribute(
-    pi: *mut c_int,
-    attrib: c_int,
-    dev_idx: cu::Device,
-) -> cu::Result {
-    if pi == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    let attrib = match u8::try_from(attrib) {
-        Ok(a) => a,
-        Err(_) => return cu::Result::ERROR_INVALID_VALUE,
-    };
-    match cu::DeviceAttribute::try_new(attrib) {
-        Some(cu::DeviceAttribute::Static(a)) => {
-            unsafe { *pi = ze::Device::get_attribute_static(a) };
-            cu::Result::SUCCESS
-        }
-        Some(cu::DeviceAttribute::Dynamic(a)) => {
-            Context::call_device(dev_idx, |dev| dev.get_attribute(pi, a))
-        }
-        // TODO: add support for more properties
-        None => cu::Result::SUCCESS,
-    }
-}
-
-#[no_mangle]
-pub extern "C" fn cuDeviceGetUuid(uuid: *mut cu::Uuid, dev_idx: cu::Device) -> cu::Result {
-    if uuid == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call_device(dev_idx, |dev| dev.get_uuid(uuid))
-}
-
-#[no_mangle]
-pub extern "C" fn cuCtxGetCurrent(pctx: *mut cu::Context) -> cu::Result {
-    let ctx = CONTEXT_STACK.with(|stack| match stack.borrow().last() {
-        Some(ctx) => ctx.clone(),
-        None => cu::Context::null(),
-    });
-    unsafe { *pctx = ctx };
-    cu::Result::SUCCESS
-}
-
-#[no_mangle]
-pub extern "C" fn cuCtxSetCurrent(ctx: cu::Context) -> cu::Result {
-    CONTEXT_STACK.with(|stack| {
-        let mut stack = stack.borrow_mut();
-        stack.pop();
-        if ctx != cu::Context::null() {
-            stack.push(ctx);
-        }
-    });
-    cu::Result::SUCCESS
-}
-
-#[no_mangle]
-pub extern "C" fn cuMemAlloc_v2(dptr: *mut cu::DevicePtr, bytesize: usize) -> cu::Result {
-    if dptr == ptr::null_mut() || bytesize == 0 {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    Context::call2(|drv| {
-        unsafe {
-            *dptr = cu::DevicePtr(drv.devices[0].base.mem_alloc_device(
-                &mut drv.base,
-                bytesize,
-                0,
-            )? as usize)
-        };
-        Ok(())
-    })
-}
-
-#[no_mangle]
-pub extern "C" fn cuCtxDestroy_v2(ctx: cu::Context) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuMemcpyDtoH_v2(
-    dst_host: *mut c_void,
-    src_device: cu::DevicePtr,
-    byte_count: usize,
-) -> cu::Result {
-    if dst_host == ptr::null_mut() || src_device.0 == 0 || byte_count == 0 {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    // TODO: choose the right device from device ptr
-    Context::call(|drv| {
-        let cu_dev = &mut drv.devices[0];
-        // Safe, because there's no Drop impl for device
-        let dev = unsafe { l0::Device::from_ffi(cu_dev.base.as_ffi()) };
-        let queue = &mut cu_dev.default_queue;
-        let result = unsafe {
-            memcpy_impl(
-                &mut drv.base,
-                dst_host,
-                src_device.0 as *mut _,
-                byte_count,
-                &dev,
-                queue,
-            )
-        };
-        match result {
-            Ok(()) => l0_sys::ze_result_t::ZE_RESULT_SUCCESS,
-            Err(e) => e.0,
-        }
-    })
-}
-
-#[no_mangle]
-pub extern "C" fn cuMemFree_v2(srcDevice: cu::DevicePtr) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuModuleLoad(module: *mut cu::Module, fname: *const c_char) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuGetErrorString(error: cu::Result, pStr: *mut *const c_char) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuLaunchKernel(
-    f: cu::Function,
-    gridDimX: c_uint,
-    gridDimY: c_uint,
-    gridDimZ: c_uint,
-    blockDimX: c_uint,
-    blockDimY: c_uint,
-    blockDimZ: c_uint,
-    sharedMemBytes: c_uint,
-    hStream: cu::Stream,
-    kernelParams: *mut *mut (),
-    extra: *mut *mut (),
-) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-#[allow(non_snake_case)]
-pub extern "C" fn cuModuleLoadDataEx(
-    module: *mut cu::Module,
-    image: *const (),
-    numOptions: c_uint,
-    options: *mut cu::JitOption,
-    optionValues: *mut *mut (),
-) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuMemcpyHtoD_v2(
-    dst_device: cu::DevicePtr,
-    src_host: *const c_void,
-    byte_count: usize,
-) -> cu::Result {
-    if dst_device.0 == 0 || src_host == ptr::null_mut() || byte_count == 0 {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    // TODO: choose the right device from device ptr
-    Context::call(|drv| {
-        let cu_dev = &mut drv.devices[0];
-        // Safe, because there's no Drop impl for device
-        let dev = unsafe { l0::Device::from_ffi(cu_dev.base.as_ffi()) };
-        let queue = &mut cu_dev.default_queue;
-        let result = unsafe {
-            memcpy_impl(
-                &mut drv.base,
-                dst_device.0 as *mut _,
-                src_host,
-                byte_count,
-                &dev,
-                queue,
-            )
-        };
-        match result {
-            Ok(()) => l0_sys::ze_result_t::ZE_RESULT_SUCCESS,
-            Err(e) => e.0,
-        }
-    })
-}
-
-unsafe fn memcpy_impl(
-    ctx: &mut l0::Context,
-    dst: *mut c_void,
-    src: *const c_void,
-    bytes_count: usize,
-    dev: &l0::Device,
-    queue: &mut l0::CommandQueue,
-) -> l0::Result<()> {
-    let mut cmd_list = l0::CommandList::new(ctx, &dev)?;
-    cmd_list.append_memory_copy_unsafe(dst, src, bytes_count, None, &mut [])?;
-    queue.execute(cmd_list)?;
-    Ok(())
-}
-
-#[no_mangle]
-pub extern "C" fn cuCtxCreate_v2(
-    pctx: *mut cu::Context,
-    flags: c_uint,
-    dev: cu::Device,
-) -> cu::Result {
-    if pctx == ptr::null_mut() {
-        return cu::Result::ERROR_INVALID_VALUE;
-    }
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuModuleGetFunction(
-    hfunc: *mut cu::Function,
-    hmod: cu::Module,
-    name: *const c_char,
-) -> cu::Result {
-    cu::Result::ERROR_NOT_SUPPORTED
-}
-
-#[no_mangle]
-pub extern "C" fn cuDevicePrimaryCtxRetain(pctx: *mut cu::Context, dev: cu::Device) -> cu::Result {
-    cu::Result::SUCCESS
-}
-
-#[no_mangle]
-pub extern "C" fn cuCtxGetDevice(dev: *mut cu::Device) -> cu::Result {
-    unsafe { *dev = cu::Device(0) };
-    cu::Result::SUCCESS
-}
+mod cuda_impl;
+pub(crate) mod r#impl;
diff --git a/notcuda/src/ze.rs b/notcuda/src/ze.rs
deleted file mode 100644
index c41fd2b..0000000
--- a/notcuda/src/ze.rs
+++ /dev/null
@@ -1,140 +0,0 @@
-use super::cu;
-use crate::cuda;
-
-use std::cmp;
-use std::os::raw::{c_char, c_int};
-
-pub struct Device {
-    pub base: l0::Device,
-    pub default_queue: l0::CommandQueue,
-    properties: Option<Box<l0::sys::ze_device_properties_t>>,
-    image_properties: Option<Box<l0::sys::ze_device_image_properties_t>>,
-    memory_properties: Option<Vec<l0::sys::ze_device_memory_properties_t>>,
-}
-
-impl Device {
-    pub fn new(ctx: &mut l0::Context, d: l0::Device) -> l0::Result<Self> {
-        let queue = l0::CommandQueue::new(ctx, &d)?;
-        Ok(Self {
-            base: d,
-            default_queue: queue,
-            properties: None,
-            image_properties: None,
-            memory_properties: None,
-        })
-    }
-
-    fn get_properties<'a>(&'a mut self) -> l0::Result<&'a l0::sys::ze_device_properties_t> {
-        if let Some(ref prop) = self.properties {
-            return Ok(prop);
-        }
-        match self.base.get_properties() {
-            Ok(prop) => Ok(self.properties.get_or_insert(prop)),
-            Err(e) => Err(e),
-        }
-    }
-
-    fn get_image_properties(&mut self) -> l0::Result<&l0::sys::ze_device_image_properties_t> {
-        if let Some(ref prop) = self.image_properties {
-            return Ok(prop);
-        }
-        match self.base.get_image_properties() {
-            Ok(prop) => Ok(self.image_properties.get_or_insert(prop)),
-            Err(e) => Err(e),
-        }
-    }
-
-    fn get_memory_properties(&mut self) -> l0::Result<&[l0::sys::ze_device_memory_properties_t]> {
-        if let Some(ref prop) = self.memory_properties {
-            return Ok(prop);
-        }
-        match self.base.get_memory_properties() {
-            Ok(prop) => Ok(self.memory_properties.get_or_insert(prop)),
-            Err(e) => Err(e),
-        }
-    }
-
-    pub fn get_name(&mut self, name: *mut c_char, len: c_int) -> l0::Result<()> {
-        let props = self.get_properties()?;
-        let null_pos = props.name.iter().position(|&c| c == 0).unwrap_or(0);
-        let dst_null_pos = cmp::min((len - 1) as usize, null_pos);
-        unsafe { *(name.add(dst_null_pos)) = 0 };
-        unsafe { std::ptr::copy_nonoverlapping(props.name.as_ptr(), name, dst_null_pos) };
-        Ok(())
-    }
-
-    pub fn total_mem(&mut self, bytes: *mut usize) -> l0::Result<()> {
-        let props = self.get_memory_properties()?;
-        let max_mem = props.iter().map(|p| p.totalSize).max().unwrap();
-        unsafe { *bytes = max_mem as usize };
-        Ok(())
-    }
-
-    pub fn get_attribute_static(attr: cu::DeviceStaticAttribute) -> c_int {
-        match attr {
-            cu::DeviceStaticAttribute::GPU_OVERLAP => 1,
-            cu::DeviceStaticAttribute::KERNEL_EXEC_TIMEOUT => 0,
-            cu::DeviceStaticAttribute::INTEGRATED => 1,
-            cu::DeviceStaticAttribute::COMPUTE_CAPABILITY_MAJOR => c_int::max_value(),
-            cu::DeviceStaticAttribute::COMPUTE_CAPABILITY_MINOR => c_int::max_value(),
-        }
-    }
-
-    fn get_attribute_general(
-        attr: cu::DeviceGeneralAttribute,
-        props: &l0_sys::ze_device_properties_t,
-    ) -> c_int {
-        match attr {
-            cu::DeviceGeneralAttribute::CAN_MAP_HOST_MEMORY => 1,
-            cu::DeviceGeneralAttribute::ASYNC_ENGINE_COUNT => props.maxHardwareContexts as i32,
-            cu::DeviceGeneralAttribute::MULTIPROCESSOR_COUNT => {
-                (props.numSlices * props.numSubslicesPerSlice * props.numEUsPerSubslice) as i32
-            }
-        }
-    }
-
-    fn get_attribute_texture(
-        attr: cu::DeviceTextureAttribute,
-        props: &l0_sys::ze_device_image_properties_t,
-    ) -> c_int {
-        match attr {
-            cu::DeviceTextureAttribute::MAXIMUM_TEXTURE1D_WIDTH => {
-                cmp::min(props.maxImageDims1D, c_int::max_value() as u32) as c_int
-            }
-        }
-    }
-
-    pub fn get_attribute(
-        &mut self,
-        pi: *mut c_int,
-        attr: cu::DeviceDynamicAttribute,
-    ) -> l0::Result<()> {
-        let value = match attr {
-            cu::DeviceDynamicAttribute::General(a) => {
-                Device::get_attribute_general(a, self.get_properties()?)
-            }
-            cu::DeviceDynamicAttribute::Texture(a) => {
-                Device::get_attribute_texture(a, self.get_image_properties()?)
-            }
-        };
-        unsafe { *pi = value };
-        Ok(())
-    }
-
-    pub fn get_uuid(&mut self, uuid: *mut cu::Uuid) -> l0::Result<()> {
-        let props = self.get_properties()?;
-        unsafe { *uuid = cu::Uuid { x: props.uuid.id } };
-        Ok(())
-    }
-}
-
-pub struct Context {
-    pub cuda_manager: *mut cuda::rt::ContextStateManager,
-    pub cuda_state: *mut cuda::rt::ContextState,
-    pub cuda_dtor_cb:
-        extern "C" fn(cu::Context, *mut cuda::rt::ContextStateManager, *mut cuda::rt::ContextState),
-}
-
-pub struct Module {
-    pub ptx_text: Vec<u8>,
-}
diff --git a/ptx/Cargo.toml b/ptx/Cargo.toml
index e7be3b7..9842e27 100644
--- a/ptx/Cargo.toml
+++ b/ptx/Cargo.toml
@@ -13,7 +13,6 @@ rspirv = "0.6"
 spirv_headers = "1.4"
 quick-error = "1.2"
 bit-vec = "0.6"
-paste = "0.1"
 half ="1.6"
 
 [build-dependencies.lalrpop]
@@ -24,3 +23,4 @@ features = ["lexer"]
 level_zero-sys = { path = "../level_zero-sys" }
 level_zero = { path = "../level_zero" }
 spirv_tools-sys = { path = "../spirv_tools-sys" }
+paste = "1.0"
diff --git a/ptx/src/lib.rs b/ptx/src/lib.rs
index 15302ff..6912d92 100644
--- a/ptx/src/lib.rs
+++ b/ptx/src/lib.rs
@@ -1,4 +1,4 @@
-#[cfg_attr(macro_use, test)]
+#[cfg(test)]
 extern crate paste;
 #[macro_use]
 extern crate lalrpop_util;
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 7de2c54..257fad2 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -51,23 +51,23 @@ test_ptx!(shl, [11u64], [44u64]);
 test_ptx!(cvt_sat_s_u, [-1i32], [0i32]);
 test_ptx!(cvta, [3.0f32], [3.0f32]);
 
-struct DisplayError<T: Display + Debug> {
+struct DisplayError<T: Debug> {
     err: T,
 }
 
-impl<T: Display + Debug> Display for DisplayError<T> {
+impl<T: Debug> Display for DisplayError<T> {
     fn fmt(&self, f: &mut Formatter<'_>) -> fmt::Result {
-        Display::fmt(&self.err, f)
+        Debug::fmt(&self.err, f)
     }
 }
 
-impl<T: Display + Debug> Debug for DisplayError<T> {
+impl<T: Debug> Debug for DisplayError<T> {
     fn fmt(&self, f: &mut Formatter<'_>) -> fmt::Result {
         Debug::fmt(&self.err, f)
     }
 }
 
-impl<T: Display + Debug> error::Error for DisplayError<T> {}
+impl<T: Debug> error::Error for DisplayError<T> {}
 
 fn test_ptx_assert<'a, T: From<u8> + ze::SafeRepr + Debug + Copy + PartialEq>(
     name: &str,
author	Andrzej Janik <[email protected]>	2020-09-01 01:43:09 +0200
committer	Andrzej Janik <[email protected]>	2020-09-01 01:43:09 +0200
commit	2e4cadc2ab061c61bacd43fab9a375b5492a1897 (patch)
tree	c5ca5d149c8621437a9d861e1a9872052792fda9
parent	34dc149be1cc83ecd931c0da742befe3eb67feaf (diff)
download	ZLUDA-2e4cadc2ab061c61bacd43fab9a375b5492a1897.tar.gz ZLUDA-2e4cadc2ab061c61bacd43fab9a375b5492a1897.zip