Rename everything

16 files changed, 6844 insertions, 0 deletions

diff --git a/zluda/Cargo.toml b/zluda/Cargo.toml
new file mode 100644
index 0000000..218fe8f
--- /dev/null
+++ b/zluda/Cargo.toml
@@ -0,0 +1,21 @@
+[package]
+name = "zluda"
+version = "0.0.0"
+authors = ["Andrzej Janik <[email protected]>"]
+edition = "2018"
+
+[lib]
+name = "zluda"
+crate-type = ["cdylib"]
+
+[dependencies]
+ptx = { path = "../ptx" }
+level_zero = { path = "../level_zero" }
+level_zero-sys = { path = "../level_zero-sys" }
+lazy_static = "1.4"
+num_enum = "0.4"
+lz4 = "1.23"
+
+[dev-dependencies]
+cuda-driver-sys = "0.3.0"
+paste = "1.0"
\ No newline at end of file
diff --git a/zluda/README b/zluda/README
new file mode 100644
index 0000000..089ddcd
--- /dev/null
+++ b/zluda/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/zluda/build.rs b/zluda/build.rs
new file mode 100644
index 0000000..3b8999f
--- /dev/null
+++ b/zluda/build.rs
@@ -0,0 +1,27 @@
+// HACK ALERT

+// This buidl script has been copy-pasted from cl-sys to avoid CUDA libraries

+// overriding path to OpenCL

+

+ fn main() {

+    if cfg!(windows) {

+        let known_sdk = [

+            // E.g. "c:\Program Files (x86)\Intel\OpenCL SDK\lib\x86\"

+            ("INTELOCLSDKROOT", "x64", "x86"),

+            // E.g. "c:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\Win32\"

+            ("CUDA_PATH", "x64", "Win32"),

+            // E.g. "C:\Program Files (x86)\AMD APP SDK\3.0\lib\x86\"

+            ("AMDAPPSDKROOT", "x86_64", "x86"),            

+        ];

+

+        for info in known_sdk.iter() {

+            if let Ok(sdk) = std::env::var(info.0) {

+                let mut path = std::path::PathBuf::from(sdk);

+                path.push("lib");

+                path.push(if cfg!(target_arch="x86_64") { info.1 } else { info.2 });

+                println!("cargo:rustc-link-search=native={}", path.display());

+            }

+        }

+

+        println!("cargo:rustc-link-search=native=C:\\Program Files (x86)\\OCL_SDK_Light\\lib\\x86_64");        

+    }

+}
\ No newline at end of file
diff --git a/zluda/src/cuda.rs b/zluda/src/cuda.rs
new file mode 100644
index 0000000..2dc26f5
--- /dev/null
+++ b/zluda/src/cuda.rs
@@ -0,0 +1,4496 @@
+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;
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGetErrorString(
+    error: CUresult,
+    pStr: *mut *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGetErrorName(
+    error: CUresult,
+    pStr: *mut *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuInit(Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::init().encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDriverGetVersion(driverVersion: *mut ::std::os::raw::c_int) -> CUresult {
+    unsafe { *driverVersion = r#impl::driver_get_version() };
+    CUresult::CUDA_SUCCESS
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGet(device: *mut CUdevice, ordinal: ::std::os::raw::c_int) -> CUresult {
+    r#impl::device::get(device.decuda(), ordinal).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetCount(count: *mut ::std::os::raw::c_int) -> CUresult {
+    r#impl::device::get_count(count).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: CUdevice) -> CUresult {
+    r#impl::device::get_uuid(uuid, dev.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceTotalMem_v2(bytes: *mut usize, dev: CUdevice) -> CUresult {
+    r#impl::device::total_mem_v2(bytes, dev.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetNvSciSyncAttributes(
+    nvSciSyncAttrList: *mut ::std::os::raw::c_void,
+    dev: CUdevice,
+    flags: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetProperties(prop: *mut CUdevprop, dev: CUdevice) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxRetain(pctx: *mut CUcontext, dev: CUdevice) -> CUresult {
+    r#impl::device::primary_ctx_retain(pctx.decuda(), dev.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxRelease(dev: CUdevice) -> CUresult {
+    cuDevicePrimaryCtxRelease_v2(dev)
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxRelease_v2(dev: CUdevice) -> CUresult {
+    r#impl::device::primary_ctx_release_v2(dev.decuda())
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxSetFlags(
+    dev: CUdevice,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    cuDevicePrimaryCtxSetFlags_v2(dev, flags)
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxSetFlags_v2(
+    dev: CUdevice,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxReset(dev: CUdevice) -> CUresult {
+    cuDevicePrimaryCtxReset_v2(dev)
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDevicePrimaryCtxReset_v2(dev: CUdevice) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxDestroy_v2(ctx: CUcontext) -> CUresult {
+    r#impl::context::destroy_v2(ctx.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxPushCurrent_v2(ctx: CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxPopCurrent_v2(pctx: *mut CUcontext) -> CUresult {
+    r#impl::context::pop_current_v2(pctx.decuda())
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxSetCurrent(ctx: CUcontext) -> CUresult {
+    r#impl::context::set_current(ctx.decuda())
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetCurrent(pctx: *mut CUcontext) -> CUresult {
+    r#impl::context::get_current(pctx.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetDevice(device: *mut CUdevice) -> CUresult {
+    r#impl::context::get_device(device.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetFlags(flags: *mut ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxSynchronize() -> CUresult {
+    r#impl::context::synchronize()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxSetLimit(limit: CUlimit, value: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetLimit(pvalue: *mut usize, limit: CUlimit) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetCacheConfig(pconfig: *mut CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxSetCacheConfig(config: CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetSharedMemConfig(pConfig: *mut CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxSetSharedMemConfig(config: CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetApiVersion(
+    ctx: CUcontext,
+    version: *mut ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::context::get_api_version(ctx.decuda(), version).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxGetStreamPriorityRange(
+    leastPriority: *mut ::std::os::raw::c_int,
+    greatestPriority: *mut ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxResetPersistingL2Cache() -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxAttach(pctx: *mut CUcontext, flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::context::attach(pctx.decuda(), flags).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxDetach(ctx: CUcontext) -> CUresult {
+    r#impl::context::detach(ctx.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleLoad(
+    module: *mut CUmodule,
+    fname: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleLoadData(
+    module: *mut CUmodule,
+    image: *const ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::module::load_data(module.decuda(), image).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleLoadFatBinary(
+    module: *mut CUmodule,
+    fatCubin: *const ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleUnload(hmod: CUmodule) -> CUresult {
+    r#impl::module::unload(hmod.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleGetFunction(
+    hfunc: *mut CUfunction,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::module::get_function(hfunc.decuda(), hmod.decuda(), name).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleGetTexRef(
+    pTexRef: *mut CUtexref,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuModuleGetSurfRef(
+    pSurfRef: *mut CUsurfref,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub 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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuLinkComplete(
+    state: CUlinkState,
+    cubinOut: *mut *mut ::std::os::raw::c_void,
+    sizeOut: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuLinkDestroy(state: CUlinkState) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult {
+    r#impl::memory::alloc_v2(dptr.decuda(), bytesize).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult {
+    r#impl::memory::free_v2(dptr.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemGetAddressRange_v2(
+    pbase: *mut CUdeviceptr,
+    psize: *mut usize,
+    dptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAllocHost_v2(
+    pp: *mut *mut ::std::os::raw::c_void,
+    bytesize: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemHostGetFlags(
+    pFlags: *mut ::std::os::raw::c_uint,
+    p: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAllocManaged(
+    dptr: *mut CUdeviceptr,
+    bytesize: usize,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetByPCIBusId(
+    dev: *mut CUdevice,
+    pciBusId: *const ::std::os::raw::c_char,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetPCIBusId(
+    pciBusId: *mut ::std::os::raw::c_char,
+    len: ::std::os::raw::c_int,
+    dev: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuIpcGetEventHandle(pHandle: *mut CUipcEventHandle, event: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuIpcOpenEventHandle(
+    phEvent: *mut CUevent,
+    handle: CUipcEventHandle,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuIpcGetMemHandle(pHandle: *mut CUipcMemHandle, dptr: CUdeviceptr) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuIpcOpenMemHandle(
+    pdptr: *mut CUdeviceptr,
+    handle: CUipcMemHandle,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuIpcCloseMemHandle(dptr: CUdeviceptr) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemHostUnregister(p: *mut ::std::os::raw::c_void) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy(dst: CUdeviceptr, src: CUdeviceptr, ByteCount: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyPeer(
+    dstDevice: CUdeviceptr,
+    dstContext: CUcontext,
+    srcDevice: CUdeviceptr,
+    srcContext: CUcontext,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyDtoD_v2(
+    dstDevice: CUdeviceptr,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyDtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyAtoD_v2(
+    dstDevice: CUdeviceptr,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyHtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyAtoH_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyAtoA_v2(
+    dstArray: CUarray,
+    dstOffset: usize,
+    srcArray: CUarray,
+    srcOffset: usize,
+    ByteCount: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy2D_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy2DUnaligned_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy3D_v2(pCopy: *const CUDA_MEMCPY3D) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy3DPeer(pCopy: *const CUDA_MEMCPY3D_PEER) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyAsync(
+    dst: CUdeviceptr,
+    src: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyPeerAsync(
+    dstDevice: CUdeviceptr,
+    dstContext: CUcontext,
+    srcDevice: CUdeviceptr,
+    srcContext: CUcontext,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyHtoDAsync_v2(
+    dstDevice: CUdeviceptr,
+    srcHost: *const ::std::os::raw::c_void,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyDtoHAsync_v2(
+    dstHost: *mut ::std::os::raw::c_void,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpyDtoDAsync_v2(
+    dstDevice: CUdeviceptr,
+    srcDevice: CUdeviceptr,
+    ByteCount: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy2DAsync_v2(pCopy: *const CUDA_MEMCPY2D, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy3DAsync_v2(pCopy: *const CUDA_MEMCPY3D, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemcpy3DPeerAsync(
+    pCopy: *const CUDA_MEMCPY3D_PEER,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD8_v2(
+    dstDevice: CUdeviceptr,
+    uc: ::std::os::raw::c_uchar,
+    N: usize,
+) -> CUresult {
+    r#impl::memory::set_d8_v2(dstDevice.decuda(), uc, N).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD16_v2(
+    dstDevice: CUdeviceptr,
+    us: ::std::os::raw::c_ushort,
+    N: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD32_v2(
+    dstDevice: CUdeviceptr,
+    ui: ::std::os::raw::c_uint,
+    N: usize,
+) -> CUresult {
+    r#impl::memory::set_d32_v2(dstDevice.decuda(), ui, N).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D8_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    uc: ::std::os::raw::c_uchar,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D16_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    us: ::std::os::raw::c_ushort,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D32_v2(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    ui: ::std::os::raw::c_uint,
+    Width: usize,
+    Height: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD8Async(
+    dstDevice: CUdeviceptr,
+    uc: ::std::os::raw::c_uchar,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD16Async(
+    dstDevice: CUdeviceptr,
+    us: ::std::os::raw::c_ushort,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD32Async(
+    dstDevice: CUdeviceptr,
+    ui: ::std::os::raw::c_uint,
+    N: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D8Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    uc: ::std::os::raw::c_uchar,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D16Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    us: ::std::os::raw::c_ushort,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemsetD2D32Async(
+    dstDevice: CUdeviceptr,
+    dstPitch: usize,
+    ui: ::std::os::raw::c_uint,
+    Width: usize,
+    Height: usize,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuArrayCreate_v2(
+    pHandle: *mut CUarray,
+    pAllocateArray: *const CUDA_ARRAY_DESCRIPTOR,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuArrayGetDescriptor_v2(
+    pArrayDescriptor: *mut CUDA_ARRAY_DESCRIPTOR,
+    hArray: CUarray,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuArrayDestroy(hArray: CUarray) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuArray3DCreate_v2(
+    pHandle: *mut CUarray,
+    pAllocateArray: *const CUDA_ARRAY3D_DESCRIPTOR,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuArray3DGetDescriptor_v2(
+    pArrayDescriptor: *mut CUDA_ARRAY3D_DESCRIPTOR,
+    hArray: CUarray,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMipmappedArrayCreate(
+    pHandle: *mut CUmipmappedArray,
+    pMipmappedArrayDesc: *const CUDA_ARRAY3D_DESCRIPTOR,
+    numMipmapLevels: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMipmappedArrayGetLevel(
+    pLevelArray: *mut CUarray,
+    hMipmappedArray: CUmipmappedArray,
+    level: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMipmappedArrayDestroy(hMipmappedArray: CUmipmappedArray) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAddressReserve(
+    ptr: *mut CUdeviceptr,
+    size: usize,
+    alignment: usize,
+    addr: CUdeviceptr,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAddressFree(ptr: CUdeviceptr, size: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemCreate(
+    handle: *mut CUmemGenericAllocationHandle,
+    size: usize,
+    prop: *const CUmemAllocationProp,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemRelease(handle: CUmemGenericAllocationHandle) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemMap(
+    ptr: CUdeviceptr,
+    size: usize,
+    offset: usize,
+    handle: CUmemGenericAllocationHandle,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemUnmap(ptr: CUdeviceptr, size: usize) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemSetAccess(
+    ptr: CUdeviceptr,
+    size: usize,
+    desc: *const CUmemAccessDesc,
+    count: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemGetAccess(
+    flags: *mut ::std::os::raw::c_ulonglong,
+    location: *const CUmemLocation,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemExportToShareableHandle(
+    shareableHandle: *mut ::std::os::raw::c_void,
+    handle: CUmemGenericAllocationHandle,
+    handleType: CUmemAllocationHandleType,
+    flags: ::std::os::raw::c_ulonglong,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemImportFromShareableHandle(
+    handle: *mut CUmemGenericAllocationHandle,
+    osHandle: *mut ::std::os::raw::c_void,
+    shHandleType: CUmemAllocationHandleType,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemGetAllocationGranularity(
+    granularity: *mut usize,
+    prop: *const CUmemAllocationProp,
+    option: CUmemAllocationGranularity_flags,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemGetAllocationPropertiesFromHandle(
+    prop: *mut CUmemAllocationProp,
+    handle: CUmemGenericAllocationHandle,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemRetainAllocationHandle(
+    handle: *mut CUmemGenericAllocationHandle,
+    addr: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuPointerGetAttribute(
+    data: *mut ::std::os::raw::c_void,
+    attribute: CUpointer_attribute,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemPrefetchAsync(
+    devPtr: CUdeviceptr,
+    count: usize,
+    dstDevice: CUdevice,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuMemAdvise(
+    devPtr: CUdeviceptr,
+    count: usize,
+    advice: CUmem_advise,
+    device: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuPointerSetAttribute(
+    value: *const ::std::os::raw::c_void,
+    attribute: CUpointer_attribute,
+    ptr: CUdeviceptr,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamCreate(
+    phStream: *mut CUstream,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::stream::create(phStream.decuda(), Flags).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamCreateWithPriority(
+    phStream: *mut CUstream,
+    flags: ::std::os::raw::c_uint,
+    priority: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamGetPriority(
+    hStream: CUstream,
+    priority: *mut ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamGetFlags(
+    hStream: CUstream,
+    flags: *mut ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUresult {
+    r#impl::stream::get_ctx(hStream.decuda(), pctx.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamWaitEvent(
+    hStream: CUstream,
+    hEvent: CUevent,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamBeginCapture_v2(
+    hStream: CUstream,
+    mode: CUstreamCaptureMode,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuThreadExchangeStreamCaptureMode(mode: *mut CUstreamCaptureMode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamEndCapture(hStream: CUstream, phGraph: *mut CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamIsCapturing(
+    hStream: CUstream,
+    captureStatus: *mut CUstreamCaptureStatus,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamGetCaptureInfo(
+    hStream: CUstream,
+    captureStatus: *mut CUstreamCaptureStatus,
+    id: *mut cuuint64_t,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamAttachMemAsync(
+    hStream: CUstream,
+    dptr: CUdeviceptr,
+    length: usize,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamQuery(hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamSynchronize(hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamDestroy_v2(hStream: CUstream) -> CUresult {
+    r#impl::stream::destroy_v2(hStream.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamCopyAttributes(dst: CUstream, src: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamGetAttribute(
+    hStream: CUstream,
+    attr: CUstreamAttrID,
+    value_out: *mut CUstreamAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamSetAttribute(
+    hStream: CUstream,
+    attr: CUstreamAttrID,
+    value: *const CUstreamAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventCreate(phEvent: *mut CUevent, Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventRecord(hEvent: CUevent, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventQuery(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventSynchronize(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventDestroy_v2(hEvent: CUevent) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuEventElapsedTime(
+    pMilliseconds: *mut f32,
+    hStart: CUevent,
+    hEnd: CUevent,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuImportExternalMemory(
+    extMem_out: *mut CUexternalMemory,
+    memHandleDesc: *const CUDA_EXTERNAL_MEMORY_HANDLE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuExternalMemoryGetMappedBuffer(
+    devPtr: *mut CUdeviceptr,
+    extMem: CUexternalMemory,
+    bufferDesc: *const CUDA_EXTERNAL_MEMORY_BUFFER_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuExternalMemoryGetMappedMipmappedArray(
+    mipmap: *mut CUmipmappedArray,
+    extMem: CUexternalMemory,
+    mipmapDesc: *const CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDestroyExternalMemory(extMem: CUexternalMemory) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuImportExternalSemaphore(
+    extSem_out: *mut CUexternalSemaphore,
+    semHandleDesc: *const CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDestroyExternalSemaphore(extSem: CUexternalSemaphore) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamWaitValue32(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint32_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamWaitValue64(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint64_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamWriteValue32(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint32_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamWriteValue64(
+    stream: CUstream,
+    addr: CUdeviceptr,
+    value: cuuint64_t,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuStreamBatchMemOp(
+    stream: CUstream,
+    count: ::std::os::raw::c_uint,
+    paramArray: *mut CUstreamBatchMemOpParams,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncGetAttribute(
+    pi: *mut ::std::os::raw::c_int,
+    attrib: CUfunction_attribute,
+    hfunc: CUfunction,
+) -> CUresult {
+    r#impl::function::get_attribute(pi, attrib, hfunc.decuda()).encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncSetAttribute(
+    hfunc: CUfunction,
+    attrib: CUfunction_attribute,
+    value: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncSetCacheConfig(hfunc: CUfunction, config: CUfunc_cache) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncSetSharedMemConfig(hfunc: CUfunction, config: CUsharedconfig) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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::function::launch_kernel(
+        f.decuda(),
+        gridDimX,
+        gridDimY,
+        gridDimZ,
+        blockDimX,
+        blockDimY,
+        blockDimZ,
+        sharedMemBytes,
+        hStream.decuda(),
+        kernelParams,
+        extra,
+    )
+    .encuda()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuLaunchHostFunc(
+    hStream: CUstream,
+    fn_: CUhostFn,
+    userData: *mut ::std::os::raw::c_void,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncSetSharedSize(
+    hfunc: CUfunction,
+    bytes: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuParamSetSize(hfunc: CUfunction, numbytes: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuParamSeti(
+    hfunc: CUfunction,
+    offset: ::std::os::raw::c_int,
+    value: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuParamSetf(
+    hfunc: CUfunction,
+    offset: ::std::os::raw::c_int,
+    value: f32,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuLaunch(f: CUfunction) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuParamSetTexRef(
+    hfunc: CUfunction,
+    texunit: ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphCreate(phGraph: *mut CUgraph, flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphKernelNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphKernelNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphMemcpyNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_MEMCPY3D,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphMemcpyNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_MEMCPY3D,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphMemsetNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_MEMSET_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphMemsetNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_MEMSET_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphHostNodeGetParams(
+    hNode: CUgraphNode,
+    nodeParams: *mut CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphHostNodeSetParams(
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphAddChildGraphNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+    childGraph: CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphChildGraphNodeGetGraph(
+    hNode: CUgraphNode,
+    phGraph: *mut CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphAddEmptyNode(
+    phGraphNode: *mut CUgraphNode,
+    hGraph: CUgraph,
+    dependencies: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphClone(phGraphClone: *mut CUgraph, originalGraph: CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphNodeFindInClone(
+    phNode: *mut CUgraphNode,
+    hOriginalNode: CUgraphNode,
+    hClonedGraph: CUgraph,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphNodeGetType(hNode: CUgraphNode, type_: *mut CUgraphNodeType) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphGetNodes(
+    hGraph: CUgraph,
+    nodes: *mut CUgraphNode,
+    numNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphGetRootNodes(
+    hGraph: CUgraph,
+    rootNodes: *mut CUgraphNode,
+    numRootNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphGetEdges(
+    hGraph: CUgraph,
+    from: *mut CUgraphNode,
+    to: *mut CUgraphNode,
+    numEdges: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphNodeGetDependencies(
+    hNode: CUgraphNode,
+    dependencies: *mut CUgraphNode,
+    numDependencies: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphNodeGetDependentNodes(
+    hNode: CUgraphNode,
+    dependentNodes: *mut CUgraphNode,
+    numDependentNodes: *mut usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphAddDependencies(
+    hGraph: CUgraph,
+    from: *const CUgraphNode,
+    to: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphRemoveDependencies(
+    hGraph: CUgraph,
+    from: *const CUgraphNode,
+    to: *const CUgraphNode,
+    numDependencies: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphDestroyNode(hNode: CUgraphNode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecKernelNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecMemcpyNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    copyParams: *const CUDA_MEMCPY3D,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecMemsetNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    memsetParams: *const CUDA_MEMSET_NODE_PARAMS,
+    ctx: CUcontext,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecHostNodeSetParams(
+    hGraphExec: CUgraphExec,
+    hNode: CUgraphNode,
+    nodeParams: *const CUDA_HOST_NODE_PARAMS,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphLaunch(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecDestroy(hGraphExec: CUgraphExec) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphDestroy(hGraph: CUgraph) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphExecUpdate(
+    hGraphExec: CUgraphExec,
+    hGraph: CUgraph,
+    hErrorNode_out: *mut CUgraphNode,
+    updateResult_out: *mut CUgraphExecUpdateResult,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphKernelNodeCopyAttributes(dst: CUgraphNode, src: CUgraphNode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphKernelNodeGetAttribute(
+    hNode: CUgraphNode,
+    attr: CUkernelNodeAttrID,
+    value_out: *mut CUkernelNodeAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphKernelNodeSetAttribute(
+    hNode: CUgraphNode,
+    attr: CUkernelNodeAttrID,
+    value: *const CUkernelNodeAttrValue,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetArray(
+    hTexRef: CUtexref,
+    hArray: CUarray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetMipmappedArray(
+    hTexRef: CUtexref,
+    hMipmappedArray: CUmipmappedArray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetAddress_v2(
+    ByteOffset: *mut usize,
+    hTexRef: CUtexref,
+    dptr: CUdeviceptr,
+    bytes: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetAddress2D_v3(
+    hTexRef: CUtexref,
+    desc: *const CUDA_ARRAY_DESCRIPTOR,
+    dptr: CUdeviceptr,
+    Pitch: usize,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetFormat(
+    hTexRef: CUtexref,
+    fmt: CUarray_format,
+    NumPackedComponents: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetAddressMode(
+    hTexRef: CUtexref,
+    dim: ::std::os::raw::c_int,
+    am: CUaddress_mode,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetMipmapFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetMipmapLevelBias(hTexRef: CUtexref, bias: f32) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetMipmapLevelClamp(
+    hTexRef: CUtexref,
+    minMipmapLevelClamp: f32,
+    maxMipmapLevelClamp: f32,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetMaxAnisotropy(
+    hTexRef: CUtexref,
+    maxAniso: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetBorderColor(hTexRef: CUtexref, pBorderColor: *mut f32) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefSetFlags(hTexRef: CUtexref, Flags: ::std::os::raw::c_uint) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetAddress_v2(pdptr: *mut CUdeviceptr, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetArray(phArray: *mut CUarray, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetMipmappedArray(
+    phMipmappedArray: *mut CUmipmappedArray,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetAddressMode(
+    pam: *mut CUaddress_mode,
+    hTexRef: CUtexref,
+    dim: ::std::os::raw::c_int,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetFilterMode(pfm: *mut CUfilter_mode, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetFormat(
+    pFormat: *mut CUarray_format,
+    pNumChannels: *mut ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetMipmapFilterMode(
+    pfm: *mut CUfilter_mode,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetMipmapLevelBias(pbias: *mut f32, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetMipmapLevelClamp(
+    pminMipmapLevelClamp: *mut f32,
+    pmaxMipmapLevelClamp: *mut f32,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetMaxAnisotropy(
+    pmaxAniso: *mut ::std::os::raw::c_int,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetBorderColor(pBorderColor: *mut f32, hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefGetFlags(
+    pFlags: *mut ::std::os::raw::c_uint,
+    hTexRef: CUtexref,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefCreate(pTexRef: *mut CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexRefDestroy(hTexRef: CUtexref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuSurfRefSetArray(
+    hSurfRef: CUsurfref,
+    hArray: CUarray,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuSurfRefGetArray(phArray: *mut CUarray, hSurfRef: CUsurfref) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexObjectDestroy(texObject: CUtexObject) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexObjectGetResourceDesc(
+    pResDesc: *mut CUDA_RESOURCE_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexObjectGetTextureDesc(
+    pTexDesc: *mut CUDA_TEXTURE_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuTexObjectGetResourceViewDesc(
+    pResViewDesc: *mut CUDA_RESOURCE_VIEW_DESC,
+    texObject: CUtexObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuSurfObjectCreate(
+    pSurfObject: *mut CUsurfObject,
+    pResDesc: *const CUDA_RESOURCE_DESC,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuSurfObjectDestroy(surfObject: CUsurfObject) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuSurfObjectGetResourceDesc(
+    pResDesc: *mut CUDA_RESOURCE_DESC,
+    surfObject: CUsurfObject,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceCanAccessPeer(
+    canAccessPeer: *mut ::std::os::raw::c_int,
+    dev: CUdevice,
+    peerDev: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxEnablePeerAccess(
+    peerContext: CUcontext,
+    Flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuCtxDisablePeerAccess(peerContext: CUcontext) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuDeviceGetP2PAttribute(
+    value: *mut ::std::os::raw::c_int,
+    attrib: CUdevice_P2PAttribute,
+    srcDevice: CUdevice,
+    dstDevice: CUdevice,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsUnregisterResource(resource: CUgraphicsResource) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+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()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsResourceGetMappedMipmappedArray(
+    pMipmappedArray: *mut CUmipmappedArray,
+    resource: CUgraphicsResource,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsResourceGetMappedPointer_v2(
+    pDevPtr: *mut CUdeviceptr,
+    pSize: *mut usize,
+    resource: CUgraphicsResource,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsResourceSetMapFlags_v2(
+    resource: CUgraphicsResource,
+    flags: ::std::os::raw::c_uint,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsMapResources(
+    count: ::std::os::raw::c_uint,
+    resources: *mut CUgraphicsResource,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGraphicsUnmapResources(
+    count: ::std::os::raw::c_uint,
+    resources: *mut CUgraphicsResource,
+    hStream: CUstream,
+) -> CUresult {
+    r#impl::unimplemented()
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuGetExportTable(
+    ppExportTable: *mut *const ::std::os::raw::c_void,
+    pExportTableId: *const CUuuid,
+) -> CUresult {
+    r#impl::export_table::get(ppExportTable, pExportTableId)
+}
+
+#[cfg_attr(not(test), no_mangle)]
+pub extern "C" fn cuFuncGetModule(hmod: *mut CUmodule, hfunc: CUfunction) -> CUresult {
+    r#impl::unimplemented()
+}
diff --git a/zluda/src/cuda_impl/mod.rs b/zluda/src/cuda_impl/mod.rs
new file mode 100644
index 0000000..63b9049
--- /dev/null
+++ b/zluda/src/cuda_impl/mod.rs
@@ -0,0 +1 @@
+pub mod rt;
\ No newline at end of file
diff --git a/zluda/src/cuda_impl/rt.rs b/zluda/src/cuda_impl/rt.rs
new file mode 100644
index 0000000..3931bc3
--- /dev/null
+++ b/zluda/src/cuda_impl/rt.rs
@@ -0,0 +1,2 @@
+pub enum ContextState {}
+pub enum ContextStateManager {}
diff --git a/zluda/src/impl/context.rs b/zluda/src/impl/context.rs
new file mode 100644
index 0000000..873fc47
--- /dev/null
+++ b/zluda/src/impl/context.rs
@@ -0,0 +1,359 @@
+use super::{device, stream::Stream, stream::StreamData, HasLivenessCookie, LiveCheck};
+use super::{CUresult, GlobalState};
+use crate::{cuda::CUcontext, cuda_impl};
+use l0::sys::ze_result_t;
+use std::{cell::RefCell, num::NonZeroU32, os::raw::c_uint, ptr, sync::atomic::AtomicU32};
+use std::{
+    collections::HashSet,
+    mem::{self},
+};
+
+thread_local! {
+    pub static CONTEXT_STACK: RefCell<Vec<*mut 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;
+
+    const LIVENESS_FAIL: CUresult = CUresult::CUDA_ERROR_INVALID_CONTEXT;
+
+    fn try_drop(&mut self) -> Result<(), CUresult> {
+        for stream in self.streams.iter() {
+            let stream = unsafe { &mut **stream };
+            stream.context = ptr::null_mut();
+            Stream::destroy_impl(unsafe { Stream::ptr_from_inner(stream) })?;
+        }
+        Ok(())
+    }
+}
+
+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
+            }
+        }
+    }
+}
+
+pub struct ContextData {
+    pub flags: AtomicU32,
+    // This pointer is null only for a moment when constructing primary context
+    pub device: *mut device::Device,
+    ref_count: ContextRefCount,
+    pub default_stream: StreamData,
+    pub streams: HashSet<*mut StreamData>,
+    // All the fields below are here to support internal CUDA driver API
+    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(
+        l0_ctx: &mut l0::Context,
+        l0_dev: &l0::Device,
+        flags: c_uint,
+        is_primary: bool,
+        dev: *mut device::Device,
+    ) -> Result<Self, CUresult> {
+        let default_stream = StreamData::new_unitialized(l0_ctx, l0_dev)?;
+        Ok(ContextData {
+            flags: AtomicU32::new(flags),
+            device: dev,
+            ref_count: ContextRefCount::new(is_primary),
+            default_stream,
+            streams: HashSet::new(),
+            cuda_manager: ptr::null_mut(),
+            cuda_state: ptr::null_mut(),
+            cuda_dtor_cb: None,
+        })
+    }
+}
+
+impl Context {
+    pub fn late_init(&mut self) {
+        let ctx_data = self.as_option_mut().unwrap();
+        ctx_data.default_stream.context = ctx_data as *mut _;
+    }
+}
+
+pub fn create_v2(
+    pctx: *mut *mut Context,
+    flags: u32,
+    dev_idx: device::Index,
+) -> Result<(), CUresult> {
+    if pctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let mut ctx_box = GlobalState::lock_device(dev_idx, |dev| {
+        let dev_ptr = dev as *mut _;
+        let mut ctx_box = Box::new(LiveCheck::new(ContextData::new(
+            &mut dev.l0_context,
+            &dev.base,
+            flags,
+            false,
+            dev_ptr as *mut _,
+        )?));
+        ctx_box.late_init();
+        Ok::<_, CUresult>(ctx_box)
+    })??;
+    let ctx_ref = ctx_box.as_mut() as *mut Context;
+    unsafe { *pctx = ctx_ref };
+    mem::forget(ctx_box);
+    CONTEXT_STACK.with(|stack| stack.borrow_mut().push(ctx_ref));
+    Ok(())
+}
+
+pub fn destroy_v2(ctx: *mut Context) -> Result<(), CUresult> {
+    if ctx == ptr::null_mut() {
+        return Err(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 *mut _),
+            None => false,
+        };
+        if should_pop {
+            stack.pop();
+        }
+    });
+    GlobalState::lock(|_| Context::destroy_impl(ctx))?
+}
+
+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 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) -> Result<(), CUresult> {
+    if ctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    GlobalState::lock(|_| {
+        unsafe { &*ctx }.as_result()?;
+        Ok::<_, CUresult>(())
+    })??;
+    //TODO: query device for properties roughly matching CUDA API version
+    unsafe { *version = 1100 };
+    Ok(())
+}
+
+pub fn get_device(dev: *mut device::Index) -> Result<(), CUresult> {
+    let dev_idx = GlobalState::lock_current_context(|ctx| unsafe { &*ctx.device }.index)?;
+    unsafe { *dev = dev_idx };
+    Ok(())
+}
+
+pub fn attach(pctx: *mut *mut Context, _flags: c_uint) -> Result<(), CUresult> {
+    if pctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let ctx = GlobalState::lock_current_context_unchecked(|unchecked_ctx| {
+        let ctx = unchecked_ctx.as_result_mut()?;
+        ctx.ref_count.incr()?;
+        Ok::<_, CUresult>(unchecked_ctx as *mut _)
+    })??;
+    unsafe { *pctx = ctx };
+    Ok(())
+}
+
+pub fn detach(pctx: *mut Context) -> Result<(), CUresult> {
+    if pctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    GlobalState::lock_current_context_unchecked(|unchecked_ctx| {
+        let ctx = unchecked_ctx.as_result_mut()?;
+        if ctx.ref_count.decr() {
+            Context::destroy_impl(unchecked_ctx)?;
+        }
+        Ok::<_, CUresult>(())
+    })?
+}
+
+pub(crate) fn synchronize() -> CUresult {
+    // TODO: change the implementation once we do async stream operations
+    CUresult::CUDA_SUCCESS
+}
+
+#[cfg(test)]
+mod test {
+    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/zluda/src/impl/device.rs b/zluda/src/impl/device.rs
new file mode 100644
index 0000000..23b75f0
--- /dev/null
+++ b/zluda/src/impl/device.rs
@@ -0,0 +1,397 @@
+use super::{context, CUresult, GlobalState};
+use crate::cuda;
+use cuda::{CUdevice_attribute, CUuuid_st};
+use std::{
+    cmp, mem,
+    os::raw::{c_char, c_int},
+    ptr,
+    sync::atomic::{AtomicU32, Ordering},
+};
+
+const PROJECT_URL_SUFFIX: &'static str = " [github.com/vosen/ZLUDA]";
+
+#[repr(transparent)]
+#[derive(Clone, Copy, Eq, PartialEq, Hash)]
+pub struct Index(pub c_int);
+
+pub struct Device {
+    pub index: Index,
+    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>>,
+    compute_properties: Option<Box<l0::sys::ze_device_compute_properties_t>>,
+}
+
+unsafe impl Send for Device {}
+
+impl Device {
+    // Unsafe because it does not fully initalize primary_context
+    unsafe fn new(drv: &l0::Driver, l0_dev: l0::Device, idx: usize) -> Result<Self, CUresult> {
+        let mut ctx = l0::Context::new(drv)?;
+        let queue = l0::CommandQueue::new(&mut ctx, &l0_dev)?;
+        let primary_context = context::Context::new(context::ContextData::new(
+            &mut ctx,
+            &l0_dev,
+            0,
+            true,
+            ptr::null_mut(),
+        )?);
+        Ok(Self {
+            index: Index(idx as c_int),
+            base: l0_dev,
+            default_queue: queue,
+            l0_context: ctx,
+            primary_context: primary_context,
+            properties: None,
+            image_properties: None,
+            memory_properties: None,
+            compute_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),
+        }
+    }
+
+    fn get_compute_properties(&mut self) -> l0::Result<&l0::sys::ze_device_compute_properties_t> {
+        if let Some(ref prop) = self.compute_properties {
+            return Ok(prop);
+        }
+        match self.base.get_compute_properties() {
+            Ok(prop) => Ok(self.compute_properties.get_or_insert(prop)),
+            Err(e) => Err(e),
+        }
+    }
+
+    pub fn late_init(&mut self) {
+        self.primary_context.as_option_mut().unwrap().device = self as *mut _;
+    }
+
+    fn get_max_simd(&mut self) -> l0::Result<u32> {
+        let props = self.get_compute_properties()?;
+        Ok(*props.subGroupSizes[0..props.numSubGroupSizes as usize]
+            .iter()
+            .max()
+            .unwrap())
+    }
+}
+
+pub fn init(driver: &l0::Driver) -> Result<Vec<Device>, CUresult> {
+    let ze_devices = driver.devices()?;
+    let mut devices = ze_devices
+        .into_iter()
+        .enumerate()
+        .map(|(idx, d)| unsafe { Device::new(driver, d, idx) })
+        .collect::<Result<Vec<_>, _>>()?;
+    for dev in devices.iter_mut() {
+        dev.late_init();
+        dev.primary_context.late_init();
+    }
+    Ok(devices)
+}
+
+pub fn get_count(count: *mut c_int) -> Result<(), CUresult> {
+    let len = GlobalState::lock(|state| state.devices.len())?;
+    unsafe { *count = len as c_int };
+    Ok(())
+}
+
+pub fn get(device: *mut Index, ordinal: c_int) -> Result<(), CUresult> {
+    if device == ptr::null_mut() || ordinal < 0 {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let len = GlobalState::lock(|state| state.devices.len())?;
+    if ordinal < (len as i32) {
+        unsafe { *device = Index(ordinal) };
+        Ok(())
+    } else {
+        Err(CUresult::CUDA_ERROR_INVALID_VALUE)
+    }
+}
+
+pub fn get_name(name: *mut c_char, len: i32, dev_idx: Index) -> Result<(), CUresult> {
+    if name == ptr::null_mut() || len < 0 {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let name_ptr = GlobalState::lock_device(dev_idx, |dev| {
+        let props = dev.get_properties()?;
+        Ok::<_, l0::sys::ze_result_t>(props.name.as_ptr())
+    })??;
+    let name_len = (0..256)
+        .position(|i| unsafe { *name_ptr.add(i) } == 0)
+        .unwrap_or(256);
+    let mut dst_null_pos = cmp::min((len - 1) as usize, name_len);
+    unsafe { std::ptr::copy_nonoverlapping(name_ptr, name, dst_null_pos) };
+    if name_len + PROJECT_URL_SUFFIX.len() < (len as usize) {
+        unsafe {
+            std::ptr::copy_nonoverlapping(
+                PROJECT_URL_SUFFIX.as_ptr(),
+                name.add(name_len) as *mut _,
+                PROJECT_URL_SUFFIX.len(),
+            )
+        };
+        dst_null_pos += PROJECT_URL_SUFFIX.len();
+    }
+    unsafe { *(name.add(dst_null_pos)) = 0 };
+    Ok(())
+}
+
+pub fn total_mem_v2(bytes: *mut usize, dev_idx: Index) -> Result<(), CUresult> {
+    if bytes == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let mem_props = GlobalState::lock_device(dev_idx, |dev| {
+        let mem_props = dev.get_memory_properties()?;
+        Ok::<_, l0::sys::ze_result_t>(mem_props)
+    })??;
+    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_idx: Index,
+) -> Result<(), CUresult> {
+    if pi == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    if let Some(value) = attrib.get_static_value() {
+        unsafe { *pi = value };
+        return Ok(());
+    }
+    let value = match attrib {
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(props.maxHardwareContexts as i32)
+            })??
+        }
+        // Streaming Multiprocessor corresponds roughly to a sub-slice (thread group can't cross either)
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_properties()?;
+                Ok::<_, l0::sys::ze_result_t>((props.numSlices * props.numSubslicesPerSlice) as i32)
+            })??
+        }
+        // I honestly don't know how to answer this query
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let max_simd = dev.get_max_simd()?;
+                let props = dev.get_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(
+                    (props.numEUsPerSubslice * props.numThreadsPerEU * max_simd) as i32,
+                )
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(cmp::min(
+                    i32::max_value() as u32,
+                    props.maxTotalGroupSize,
+                ) as i32)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_image_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(cmp::min(
+                    props.maxImageDims1D,
+                    c_int::max_value() as u32,
+                ) as c_int)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(cmp::min(
+                    i32::max_value() as u32,
+                    props.maxGroupCountX,
+                ) as i32)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(cmp::min(
+                    i32::max_value() as u32,
+                    props.maxGroupCountY,
+                ) as i32)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(cmp::min(
+                    i32::max_value() as u32,
+                    props.maxGroupCountZ,
+                ) as i32)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(
+                    cmp::min(i32::max_value() as u32, props.maxGroupSizeX) as i32,
+                )
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(
+                    cmp::min(i32::max_value() as u32, props.maxGroupSizeY) as i32,
+                )
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(
+                    cmp::min(i32::max_value() as u32, props.maxGroupSizeZ) as i32,
+                )
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK => {
+            GlobalState::lock_device(dev_idx, |dev| {
+                let props = dev.get_compute_properties()?;
+                Ok::<_, l0::sys::ze_result_t>(props.maxSharedLocalMemory as i32)
+            })??
+        }
+        CUdevice_attribute::CU_DEVICE_ATTRIBUTE_WARP_SIZE => {
+            GlobalState::lock_device(dev_idx, |dev| Ok::<_, CUresult>(dev.get_max_simd()? as i32))??
+        }
+        _ => {
+            // 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_idx: Index) -> Result<(), CUresult> {
+    let ze_uuid = GlobalState::lock_device(dev_idx, |dev| {
+        let props = dev.get_properties()?;
+        Ok::<_, l0::sys::ze_result_t>(props.uuid)
+    })??;
+    unsafe {
+        *uuid = CUuuid_st {
+            bytes: mem::transmute(ze_uuid.id),
+        }
+    };
+    Ok(())
+}
+
+pub fn primary_ctx_get_state(
+    dev_idx: Index,
+    flags: *mut u32,
+    active: *mut i32,
+) -> Result<(), CUresult> {
+    let (is_active, flags_value) = GlobalState::lock_device(dev_idx, |dev| {
+        // This is safe because primary context can't be dropped
+        let ctx_ptr = &mut dev.primary_context as *mut _;
+        let flags_ptr =
+            (&unsafe { dev.primary_context.as_ref_unchecked() }.flags) as *const AtomicU32;
+        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);
+        Ok::<_, l0::sys::ze_result_t>((is_active, flags_value))
+    })??;
+    unsafe { *active = if is_active { 1 } else { 0 } };
+    unsafe { *flags = flags_value };
+    Ok(())
+}
+
+pub fn primary_ctx_retain(
+    pctx: *mut *mut context::Context,
+    dev_idx: Index,
+) -> Result<(), CUresult> {
+    let ctx_ptr = GlobalState::lock_device(dev_idx, |dev| &mut dev.primary_context as *mut _)?;
+    unsafe { *pctx = ctx_ptr };
+    Ok(())
+}
+
+// TODO: allow for retain/reset/release of primary context
+pub(crate) fn primary_ctx_release_v2(_dev_idx: Index) -> CUresult {
+    CUresult::CUDA_SUCCESS
+}
+
+#[cfg(test)]
+mod test {
+    use super::super::test::CudaDriverFns;
+    use super::super::CUresult;
+
+    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/zluda/src/impl/export_table.rs b/zluda/src/impl/export_table.rs
new file mode 100644
index 0000000..87d7f40
--- /dev/null
+++ b/zluda/src/impl/export_table.rs
@@ -0,0 +1,372 @@
+use crate::cuda::CUresult;

+use crate::{

+    cuda::{CUcontext, CUdevice, CUmodule, CUuuid},

+    cuda_impl,

+};

+

+use super::{context, context::ContextData, device, module, Decuda, Encuda, GlobalState};

+use std::mem;

+use std::os::raw::{c_uint, c_ulong, c_ushort};

+use std::{

+    ffi::{c_void, CStr},

+    ptr, slice,

+};

+

+pub fn get(table: *mut *const std::os::raw::c_void, id: *const CUuuid) -> CUresult {

+    if table == ptr::null_mut() || id == ptr::null_mut() {

+        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: CUuuid = CUuuid {

+    bytes: [

+        0xa0, 0x94, 0x79, 0x8c, 0x2e, 0x74, 0x2e, 0x74, 0x93, 0xf2, 0x08, 0x00, 0x20, 0x0c, 0x0a,

+        0x66,

+    ],

+};

+#[repr(C)]

+union VTableEntry {

+    ptr: *const (),

+    length: usize,

+}

+unsafe impl Sync for VTableEntry {}

+const TOOLS_RUNTIME_CALLBACK_HOOKS_LENGTH: usize = 7;

+static TOOLS_RUNTIME_CALLBACK_HOOKS_VTABLE: [VTableEntry; TOOLS_RUNTIME_CALLBACK_HOOKS_LENGTH] = [

+    VTableEntry {

+        length: mem::size_of::<[VTableEntry; TOOLS_RUNTIME_CALLBACK_HOOKS_LENGTH]>(),

+    },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry {

+        ptr: runtime_callback_hooks_fn1 as *const (),

+    },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry {

+        ptr: runtime_callback_hooks_fn5 as *const (),

+    },

+];

+static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN1_SPACE: [usize; 512] = [0; 512];

+

+unsafe extern "C" fn runtime_callback_hooks_fn1(ptr: *mut *mut usize, size: *mut usize) {

+    *ptr = TOOLS_RUNTIME_CALLBACK_HOOKS_FN1_SPACE.as_mut_ptr();

+    *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN1_SPACE.len();

+}

+

+static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN5_SPACE: [u8; 2] = [0; 2];

+

+unsafe extern "C" fn runtime_callback_hooks_fn5(ptr: *mut *mut u8, size: *mut usize) -> *mut u8 {

+    *ptr = TOOLS_RUNTIME_CALLBACK_HOOKS_FN5_SPACE.as_mut_ptr();

+    *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN5_SPACE.len();

+    return TOOLS_RUNTIME_CALLBACK_HOOKS_FN5_SPACE.as_mut_ptr();

+}

+

+const CUDART_INTERFACE_GUID: CUuuid = CUuuid {

+    bytes: [

+        0x6b, 0xd5, 0xfb, 0x6c, 0x5b, 0xf4, 0xe7, 0x4a, 0x89, 0x87, 0xd9, 0x39, 0x12, 0xfd, 0x9d,

+        0xf9,

+    ],

+};

+

+const CUDART_INTERFACE_LENGTH: usize = 10;

+static CUDART_INTERFACE_VTABLE: [VTableEntry; CUDART_INTERFACE_LENGTH] = [

+    VTableEntry {

+        length: mem::size_of::<[VTableEntry; CUDART_INTERFACE_LENGTH]>(),

+    },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry {

+        ptr: cudart_interface_fn1 as *const (),

+    },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry {

+        ptr: get_module_from_cubin as *const (),

+    },

+    VTableEntry {

+        ptr: cudart_interface_fn6 as *const (),

+    },

+    VTableEntry { ptr: ptr::null() },

+    VTableEntry { ptr: ptr::null() },

+];

+

+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 = GlobalState::lock_device(dev, |d| &mut d.primary_context as *mut _)?;

+    unsafe { *pctx = ctx_ptr };

+    Ok(())

+}

+

+/*

+fat_cubin:

+typedef struct {

+  int magic;

+  int version;

+  const unsigned long long* data;

+  void *filename_or_fatbins;  /* version 1: offline filename,

+                               * version 2: array of prelinked fatbins */

+} __fatBinC_Wrapper_t;

+

+data start with this header:

+#define FATBIN_MAGIC 0xBA55ED50U

+#define OLD_STYLE_FATBIN_MAGIC 0x1EE55A01U

+#define FATBIN_VERSION 0x0001U

+

+struct fatbinary_ALIGN_(8) fatBinaryHeader

+{

+  unsigned int           magic;   // FATBIN_MAGIC

+  unsigned short         version; // FATBIN_VERSION

+  unsigned short         headerSize;

+  unsigned long long int fatSize; // size of the entire fat binary excluding this header

+};

+

+there's binary data after header

+

+*/

+

+const FATBINC_MAGIC: c_uint = 0x466243B1;

+const FATBINC_VERSION: c_uint = 0x1;

+

+#[repr(C)]

+struct FatbincWrapper {

+    magic: c_uint,

+    version: c_uint,

+    data: *const FatbinHeader,

+    filename_or_fatbins: *const c_void,

+}

+

+const FATBIN_MAGIC: c_uint = 0xBA55ED50;

+const FATBIN_VERSION: c_ushort = 0x01;

+

+#[repr(C, align(8))]

+struct FatbinHeader {

+    magic: c_uint,

+    version: c_ushort,

+    header_size: c_ushort,

+    files_size: c_ulong, // excluding frame header, size of all blocks framed by this frame

+}

+

+const FATBIN_FILE_HEADER_KIND_PTX: c_ushort = 0x01;

+const FATBIN_FILE_HEADER_VERSION_CURRENT: c_ushort = 0x101;

+

+// assembly file header is a bit different, but we don't care

+#[repr(C)]

+struct FatbinFileHeader {

+    kind: c_ushort,

+    version: c_ushort,

+    header_size: c_uint,

+    padded_payload_size: c_uint,

+    unknown0: c_uint, // check if it's written into separately

+    payload_size: c_uint,

+    unknown1: c_uint,

+    unknown2: c_uint,

+    sm_version: c_uint,

+    bit_width: c_uint,

+    unknown3: c_uint,

+    unknown4: c_ulong,

+    unknown5: c_ulong,

+    uncompressed_payload: c_ulong,

+}

+

+unsafe extern "C" fn get_module_from_cubin(

+    result: *mut CUmodule,

+    fatbinc_wrapper: *const FatbincWrapper,

+    ptr1: *mut c_void,

+    ptr2: *mut c_void,

+) -> CUresult {

+    // Not sure what those two parameters are actually used for,

+    // they are somehow involved in __cudaRegisterHostVar

+    if ptr1 != ptr::null_mut() || ptr2 != ptr::null_mut() {

+        return CUresult::CUDA_ERROR_NOT_SUPPORTED;

+    }

+    if result == ptr::null_mut()

+        || (*fatbinc_wrapper).magic != FATBINC_MAGIC

+        || (*fatbinc_wrapper).version != FATBINC_VERSION

+    {

+        return CUresult::CUDA_ERROR_INVALID_VALUE;

+    }

+    let result = result.decuda();

+    let fatbin_header = (*fatbinc_wrapper).data;

+    if (*fatbin_header).magic != FATBIN_MAGIC || (*fatbin_header).version != FATBIN_VERSION {

+        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);

+    let mut ptx_files = get_ptx_files(file, end);

+    ptx_files.sort_unstable_by_key(|f| c_uint::max_value() - (**f).sm_version);

+    for file in ptx_files {

+        let slice = slice::from_raw_parts(

+            (file as *const u8).add((*file).header_size as usize),

+            (*file).payload_size as usize,

+        );

+        let kernel_text =

+            lz4::block::decompress(slice, Some((*file).uncompressed_payload as i32)).unwrap();

+        let kernel_text_string = match CStr::from_bytes_with_nul(&kernel_text) {

+            Ok(c_str) => match c_str.to_str() {

+                Ok(s) => s,

+                Err(_) => continue,

+            },

+            Err(_) => continue,

+        };

+        let module = module::SpirvModule::new(kernel_text_string);

+        match module {

+            Ok(module) => {

+                match module::load_data_impl(result, module) {

+                    Ok(()) => {}

+                    Err(err) => return err,

+                }

+                return CUresult::CUDA_SUCCESS;

+            }

+            Err(_) => continue,

+        }

+    }

+    CUresult::CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE

+}

+

+unsafe fn get_ptx_files(file: *const u8, end: *const u8) -> Vec<*const FatbinFileHeader> {

+    let mut index = file;

+    let mut result = Vec::new();

+    while index < end {

+        let file = index as *const FatbinFileHeader;

+        if (*file).kind == FATBIN_FILE_HEADER_KIND_PTX

+            && (*file).version == FATBIN_FILE_HEADER_VERSION_CURRENT

+        {

+            result.push(file)

+        }

+        index = index.add((*file).header_size as usize + (*file).padded_payload_size as usize);

+    }

+    result

+}

+

+unsafe extern "C" fn cudart_interface_fn6(_: u64) {}

+

+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: CUuuid = CUuuid {

+    bytes: [

+        0xc6, 0x93, 0x33, 0x6e, 0x11, 0x21, 0xdf, 0x11, 0xa8, 0xc3, 0x68, 0xf3, 0x55, 0xd8, 0x95,

+        0x93,

+    ],

+};

+

+// the table is much bigger and starts earlier

+static CONTEXT_LOCAL_STORAGE_INTERFACE_V0301_VTABLE: [VTableEntry; 4] = [

+    VTableEntry {

+        ptr: context_local_storage_ctor as *const (),

+    },

+    VTableEntry {

+        ptr: context_local_storage_dtor as *const (),

+    },

+    VTableEntry {

+        ptr: context_local_storage_get_state as *const (),

+    },

+    VTableEntry { ptr: ptr::null() },

+];

+

+// some kind of ctor

+unsafe extern "C" fn context_local_storage_ctor(

+    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: 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> {

+    lock_context(cu_ctx, |ctx: &mut ContextData| {

+        ctx.cuda_manager = mgr;

+        ctx.cuda_state = ctx_state;

+        ctx.cuda_dtor_cb = dtor_cb;

+    })

+}

+

+// some kind of dtor

+unsafe extern "C" fn context_local_storage_dtor(_: *mut usize, _: *mut ()) -> u32 {

+    0

+}

+

+unsafe extern "C" fn context_local_storage_get_state(

+    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> {

+    let cuda_state = lock_context(cu_ctx, |ctx: &mut ContextData| ctx.cuda_state)?;

+    if cuda_state == ptr::null_mut() {

+        Err(CUresult::CUDA_ERROR_INVALID_VALUE)

+    } else {

+        unsafe { *ctx_state = cuda_state };

+        Ok(())

+    }

+}

+

+fn lock_context<T>(

+    cu_ctx: *mut context::Context,

+    fn_impl: impl FnOnce(&mut ContextData) -> T,

+) -> Result<T, CUresult> {

+    if cu_ctx == ptr::null_mut() {

+        GlobalState::lock_current_context(fn_impl)

+    } else {

+        GlobalState::lock(|_| {

+            let ctx = unsafe { &mut *cu_ctx }.as_result_mut()?;

+            Ok(fn_impl(ctx))

+        })?

+    }

+}

diff --git a/zluda/src/impl/function.rs b/zluda/src/impl/function.rs
new file mode 100644
index 0000000..27bf9b6
--- /dev/null
+++ b/zluda/src/impl/function.rs
@@ -0,0 +1,112 @@
+use ::std::os::raw::{c_uint, c_void};
+use std::{hint, ptr};
+
+use crate::cuda::CUfunction_attribute;
+
+use super::{stream::Stream, CUresult, GlobalState, HasLivenessCookie, LiveCheck};
+
+pub type Function = LiveCheck<FunctionData>;
+
+impl HasLivenessCookie for FunctionData {
+    #[cfg(target_pointer_width = "64")]
+    const COOKIE: usize = 0x5e2ab14d5840678e;
+
+    #[cfg(target_pointer_width = "32")]
+    const COOKIE: usize = 0x33e6a1e6;
+
+    const LIVENESS_FAIL: CUresult = CUresult::CUDA_ERROR_INVALID_HANDLE;
+
+    fn try_drop(&mut self) -> Result<(), CUresult> {
+        Ok(())
+    }
+}
+
+pub struct FunctionData {
+    pub base: l0::Kernel<'static>,
+    pub arg_size: Vec<usize>,
+    pub use_shared_mem: bool,
+    pub properties: Option<Box<l0::sys::ze_kernel_properties_t>>,
+}
+
+impl FunctionData {
+    fn get_properties(&mut self) -> Result<&l0::sys::ze_kernel_properties_t, l0::sys::ze_result_t> {
+        if let None = self.properties {
+            self.properties = Some(self.base.get_properties()?)
+        }
+        match self.properties {
+            Some(ref props) => Ok(props.as_ref()),
+            None => unsafe { hint::unreachable_unchecked() },
+        }
+    }
+}
+
+pub fn launch_kernel(
+    f: *mut Function,
+    grid_dim_x: c_uint,
+    grid_dim_y: c_uint,
+    grid_dim_z: c_uint,
+    block_dim_x: c_uint,
+    block_dim_y: c_uint,
+    block_dim_z: c_uint,
+    shared_mem_bytes: c_uint,
+    hstream: *mut Stream,
+    kernel_params: *mut *mut c_void,
+    extra: *mut *mut c_void,
+) -> Result<(), CUresult> {
+    if f == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    if extra != ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_NOT_SUPPORTED);
+    }
+    GlobalState::lock_stream(hstream, |stream| {
+        let func: &mut FunctionData = unsafe { &mut *f }.as_result_mut()?;
+        for (i, arg_size) in func.arg_size.iter().enumerate() {
+            unsafe {
+                func.base
+                    .set_arg_raw(i as u32, *arg_size, *kernel_params.add(i))?
+            };
+        }
+        if func.use_shared_mem {
+            unsafe {
+                func.base.set_arg_raw(
+                    func.arg_size.len() as u32,
+                    shared_mem_bytes as usize,
+                    ptr::null(),
+                )?
+            };
+        }
+        func.base
+            .set_group_size(block_dim_x, block_dim_y, block_dim_z)?;
+        let mut cmd_list = stream.command_list()?;
+        cmd_list.append_launch_kernel(
+            &mut func.base,
+            &[grid_dim_x, grid_dim_y, grid_dim_z],
+            None,
+            &mut [],
+        )?;
+        stream.queue.execute(cmd_list)?;
+        Ok(())
+    })?
+}
+
+pub(crate) fn get_attribute(
+    pi: *mut i32,
+    attrib: CUfunction_attribute,
+    func: *mut Function,
+) -> Result<(), CUresult> {
+    if pi == ptr::null_mut() || func == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    match attrib {
+        CUfunction_attribute::CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK => {
+            let max_threads = GlobalState::lock_function(func, |func| {
+                let props = func.get_properties()?;
+                Ok::<_, CUresult>(props.maxSubgroupSize * props.maxNumSubgroups)
+            })??;
+            unsafe { *pi = max_threads as i32 };
+            Ok(())
+        }
+        _ => Err(CUresult::CUDA_ERROR_NOT_SUPPORTED),
+    }
+}
diff --git a/zluda/src/impl/memory.rs b/zluda/src/impl/memory.rs
new file mode 100644
index 0000000..f33a08c
--- /dev/null
+++ b/zluda/src/impl/memory.rs
@@ -0,0 +1,100 @@
+use super::{stream, CUresult, GlobalState};

+use std::{ffi::c_void, mem};

+

+pub fn alloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> Result<(), CUresult> {

+    let ptr = GlobalState::lock_current_context(|ctx| {

+        let dev = unsafe { &mut *ctx.device };

+        Ok::<_, CUresult>(unsafe { dev.base.mem_alloc_device(&mut dev.l0_context, bytesize, 0) }?)

+    })??;

+    unsafe { *dptr = ptr };

+    Ok(())

+}

+

+pub fn copy_v2(dst: *mut c_void, src: *const c_void, bytesize: usize) -> Result<(), CUresult> {

+    GlobalState::lock_stream(stream::CU_STREAM_LEGACY, |stream| {

+        let mut cmd_list = stream.command_list()?;

+        unsafe { cmd_list.append_memory_copy_unsafe(dst, src, bytesize, None, &mut []) }?;

+        stream.queue.execute(cmd_list)?;

+        Ok::<_, CUresult>(())

+    })?

+}

+

+pub fn free_v2(ptr: *mut c_void) -> Result<(), CUresult> {

+    GlobalState::lock_current_context(|ctx| {

+        let dev = unsafe { &mut *ctx.device };

+        Ok::<_, CUresult>(unsafe { dev.l0_context.mem_free(ptr) }?)

+    })

+    .map_err(|_| CUresult::CUDA_ERROR_INVALID_VALUE)?

+}

+

+pub(crate) fn set_d32_v2(dst: *mut c_void, ui: u32, n: usize) -> Result<(), CUresult> {

+    GlobalState::lock_stream(stream::CU_STREAM_LEGACY, |stream| {

+        let mut cmd_list = stream.command_list()?;

+        unsafe {

+            cmd_list.append_memory_fill_unsafe(dst, &ui, mem::size_of::<u32>() * n, None, &mut [])

+        }?;

+        stream.queue.execute(cmd_list)?;

+        Ok::<_, CUresult>(())

+    })?

+}

+

+pub(crate) fn set_d8_v2(dst: *mut c_void, uc: u8, n: usize) -> Result<(), CUresult> {

+    GlobalState::lock_stream(stream::CU_STREAM_LEGACY, |stream| {

+        let mut cmd_list = stream.command_list()?;

+        unsafe {

+            cmd_list.append_memory_fill_unsafe(dst, &uc, mem::size_of::<u8>() * n, None, &mut [])

+        }?;

+        stream.queue.execute(cmd_list)?;

+        Ok::<_, CUresult>(())

+    })?

+}

+

+#[cfg(test)]

+mod test {

+    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);

+    }

+

+    cuda_driver_test!(free_without_ctx);

+

+    fn free_without_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);

+        assert_eq!(T::cuMemFree_v2(mem), CUresult::CUDA_ERROR_INVALID_VALUE);

+    }

+}

diff --git a/zluda/src/impl/mod.rs b/zluda/src/impl/mod.rs
new file mode 100644
index 0000000..086d260
--- /dev/null
+++ b/zluda/src/impl/mod.rs
@@ -0,0 +1,351 @@
+use crate::{
+    cuda::{CUctx_st, CUdevice, CUdeviceptr, CUfunc_st, CUmod_st, CUresult, CUstream_st},
+    r#impl::device::Device,
+};
+use std::{
+    ffi::c_void,
+    mem::{self, ManuallyDrop},
+    os::raw::c_int,
+    ptr,
+    sync::Mutex,
+    sync::TryLockError,
+};
+
+#[cfg(test)]
+#[macro_use]
+pub mod test;
+pub mod context;
+pub mod device;
+pub mod export_table;
+pub mod function;
+pub mod memory;
+pub mod module;
+pub mod stream;
+
+#[cfg(debug_assertions)]
+pub fn unimplemented() -> CUresult {
+    unimplemented!()
+}
+
+#[cfg(not(debug_assertions))]
+pub fn unimplemented() -> CUresult {
+    CUresult::CUDA_ERROR_NOT_SUPPORTED
+}
+
+pub trait HasLivenessCookie: Sized {
+    const COOKIE: usize;
+    const LIVENESS_FAIL: CUresult;
+
+    fn try_drop(&mut self) -> Result<(), CUresult>;
+}
+
+// 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),
+        }
+    }
+
+    fn destroy_impl(this: *mut Self) -> Result<(), CUresult> {
+        let mut ctx_box = ManuallyDrop::new(unsafe { Box::from_raw(this) });
+        ctx_box.try_drop()?;
+        unsafe { ManuallyDrop::drop(&mut ctx_box) };
+        Ok(())
+    }
+
+    unsafe fn ptr_from_inner(this: *mut T) -> *mut Self {
+        let outer_ptr = (this as *mut u8).sub(mem::size_of::<usize>());
+        outer_ptr as *mut Self
+    }
+
+    pub unsafe fn as_ref_unchecked(&self) -> &T {
+        &self.data
+    }
+
+    pub fn as_option_mut(&mut self) -> Option<&mut T> {
+        if self.cookie == T::COOKIE {
+            Some(&mut self.data)
+        } else {
+            None
+        }
+    }
+
+    pub fn as_result(&self) -> Result<&T, CUresult> {
+        if self.cookie == T::COOKIE {
+            Ok(&self.data)
+        } else {
+            Err(T::LIVENESS_FAIL)
+        }
+    }
+
+    pub fn as_result_mut(&mut self) -> Result<&mut T, CUresult> {
+        if self.cookie == T::COOKIE {
+            Ok(&mut self.data)
+        } else {
+            Err(T::LIVENESS_FAIL)
+        }
+    }
+
+    #[must_use]
+    pub fn try_drop(&mut self) -> Result<(), CUresult> {
+        if self.cookie == T::COOKIE {
+            self.cookie = 0;
+            self.data.try_drop()?;
+            unsafe { ManuallyDrop::drop(&mut self.data) };
+            return Ok(());
+        }
+        Err(T::LIVENESS_FAIL)
+    }
+}
+
+impl<T: HasLivenessCookie> Drop for LiveCheck<T> {
+    fn drop(&mut self) {
+        self.cookie = 0;
+    }
+}
+
+pub trait CudaRepr: Sized {
+    type Impl: Sized;
+}
+
+impl<T: CudaRepr> CudaRepr for *mut T {
+    type Impl = *mut T::Impl;
+}
+
+pub trait Decuda<To> {
+    fn decuda(self: Self) -> To;
+}
+
+impl<T: CudaRepr> Decuda<*mut T::Impl> for *mut T {
+    fn decuda(self: Self) -> *mut T::Impl {
+        self as *mut _
+    }
+}
+
+impl 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,
+        }
+    }
+}
+
+impl<T> From<TryLockError<T>> for CUresult {
+    fn from(_: TryLockError<T>) -> Self {
+        CUresult::CUDA_ERROR_ILLEGAL_STATE
+    }
+}
+
+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(),
+        }
+    }
+}
+
+lazy_static! {
+    static ref GLOBAL_STATE: Mutex<Option<GlobalState>> = Mutex::new(None);
+}
+
+struct GlobalState {
+    devices: Vec<Device>,
+}
+
+unsafe impl Send for GlobalState {}
+
+impl GlobalState {
+    fn lock<T>(f: impl FnOnce(&mut GlobalState) -> T) -> Result<T, CUresult> {
+        let mut mutex = GLOBAL_STATE
+            .lock()
+            .unwrap_or_else(|poison| poison.into_inner());
+        let global_state = mutex.as_mut().ok_or(CUresult::CUDA_ERROR_ILLEGAL_STATE)?;
+        Ok(f(global_state))
+    }
+
+    fn lock_device<T>(
+        device::Index(dev_idx): device::Index,
+        f: impl FnOnce(&'static mut device::Device) -> T,
+    ) -> Result<T, CUresult> {
+        if dev_idx < 0 {
+            return Err(CUresult::CUDA_ERROR_INVALID_DEVICE);
+        }
+        Self::lock(|global_state| {
+            if dev_idx >= global_state.devices.len() as c_int {
+                Err(CUresult::CUDA_ERROR_INVALID_DEVICE)
+            } else {
+                Ok(f(unsafe {
+                    transmute_lifetime_mut(&mut global_state.devices[dev_idx as usize])
+                }))
+            }
+        })?
+    }
+
+    fn lock_current_context<F: FnOnce(&mut context::ContextData) -> R, R>(
+        f: F,
+    ) -> Result<R, CUresult> {
+        Self::lock_current_context_unchecked(|ctx| Ok(f(ctx.as_result_mut()?)))?
+    }
+
+    fn lock_current_context_unchecked<F: FnOnce(&mut context::Context) -> R, R>(
+        f: F,
+    ) -> Result<R, CUresult> {
+        context::CONTEXT_STACK.with(|stack| {
+            stack
+                .borrow_mut()
+                .last_mut()
+                .ok_or(CUresult::CUDA_ERROR_INVALID_CONTEXT)
+                .map(|ctx| GlobalState::lock(|_| f(unsafe { &mut **ctx })))?
+        })
+    }
+
+    fn lock_stream<T>(
+        stream: *mut stream::Stream,
+        f: impl FnOnce(&mut stream::StreamData) -> T,
+    ) -> Result<T, CUresult> {
+        if stream == ptr::null_mut()
+            || stream == stream::CU_STREAM_LEGACY
+            || stream == stream::CU_STREAM_PER_THREAD
+        {
+            Self::lock_current_context(|ctx| Ok(f(&mut ctx.default_stream)))?
+        } else {
+            Self::lock(|_| {
+                let stream = unsafe { &mut *stream }.as_result_mut()?;
+                Ok(f(stream))
+            })?
+        }
+    }
+
+    fn lock_function<T>(
+        func: *mut function::Function,
+        f: impl FnOnce(&mut function::FunctionData) -> T,
+    ) -> Result<T, CUresult> {
+        if func == ptr::null_mut() {
+            return Err(CUresult::CUDA_ERROR_INVALID_HANDLE);
+        }
+        Self::lock(|_| {
+            let func = unsafe { &mut *func }.as_result_mut()?;
+            Ok(f(func))
+        })?
+    }
+}
+
+// TODO: implement
+fn is_intel_gpu_driver(_: &l0::Driver) -> bool {
+    true
+}
+
+pub fn init() -> Result<(), CUresult> {
+    let mut global_state = GLOBAL_STATE
+        .lock()
+        .map_err(|_| CUresult::CUDA_ERROR_UNKNOWN)?;
+    if global_state.is_some() {
+        return Ok(());
+    }
+    l0::init()?;
+    let drivers = l0::Driver::get()?;
+    let devices = match drivers.into_iter().find(is_intel_gpu_driver) {
+        None => return Err(CUresult::CUDA_ERROR_UNKNOWN),
+        Some(driver) => device::init(&driver)?,
+    };
+    *global_state = Some(GlobalState { devices });
+    drop(global_state);
+    Ok(())
+}
+
+unsafe fn transmute_lifetime_mut<'a, 'b, T: ?Sized>(t: &'a mut T) -> &'b mut T {
+    mem::transmute(t)
+}
+
+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 _
+    }
+}
+
+impl<'a> CudaRepr for CUmod_st {
+    type Impl = module::Module;
+}
+
+impl<'a> CudaRepr for CUfunc_st {
+    type Impl = function::Function;
+}
+
+impl<'a> CudaRepr for CUstream_st {
+    type Impl = stream::Stream;
+}
diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs
new file mode 100644
index 0000000..cba030e
--- /dev/null
+++ b/zluda/src/impl/module.rs
@@ -0,0 +1,188 @@
+use std::{
+    collections::hash_map, collections::HashMap, ffi::c_void, ffi::CStr, ffi::CString, mem,
+    os::raw::c_char, ptr, slice,
+};
+
+use super::{
+    device, function::Function, function::FunctionData, CUresult, GlobalState, HasLivenessCookie,
+    LiveCheck,
+};
+use ptx;
+
+pub type Module = LiveCheck<ModuleData>;
+
+impl HasLivenessCookie for ModuleData {
+    #[cfg(target_pointer_width = "64")]
+    const COOKIE: usize = 0xf1313bd46505f98a;
+
+    #[cfg(target_pointer_width = "32")]
+    const COOKIE: usize = 0xbdbe3f15;
+
+    const LIVENESS_FAIL: CUresult = CUresult::CUDA_ERROR_INVALID_HANDLE;
+
+    fn try_drop(&mut self) -> Result<(), CUresult> {
+        Ok(())
+    }
+}
+
+pub struct ModuleData {
+    pub spirv: SpirvModule,
+    // This should be a Vec<>, but I'm feeling lazy
+    pub device_binaries: HashMap<device::Index, CompiledModule>,
+}
+
+pub struct SpirvModule {
+    pub binaries: Vec<u32>,
+    pub kernel_info: HashMap<String, ptx::KernelInfo>,
+    pub should_link_ptx_impl: Option<&'static [u8]>,
+    pub build_options: CString,
+}
+
+pub struct CompiledModule {
+    pub base: l0::Module,
+    pub kernels: HashMap<CString, Box<Function>>,
+}
+
+impl<L, T, E> From<ptx::ParseError<L, T, E>> for CUresult {
+    fn from(_: ptx::ParseError<L, T, E>) -> Self {
+        CUresult::CUDA_ERROR_INVALID_PTX
+    }
+}
+
+impl From<ptx::TranslateError> for CUresult {
+    fn from(_: ptx::TranslateError) -> Self {
+        CUresult::CUDA_ERROR_INVALID_PTX
+    }
+}
+
+impl SpirvModule {
+    pub fn new_raw<'a>(text: *const c_char) -> Result<Self, CUresult> {
+        let u8_text = unsafe { CStr::from_ptr(text) };
+        let ptx_text = u8_text
+            .to_str()
+            .map_err(|_| CUresult::CUDA_ERROR_INVALID_PTX)?;
+        Self::new(ptx_text)
+    }
+
+    pub fn new<'a>(ptx_text: &str) -> Result<Self, CUresult> {
+        let mut errors = Vec::new();
+        let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_text)?;
+        let spirv_module = ptx::to_spirv_module(ast)?;
+        Ok(SpirvModule {
+            binaries: spirv_module.assemble(),
+            kernel_info: spirv_module.kernel_info,
+            should_link_ptx_impl: spirv_module.should_link_ptx_impl,
+            build_options: spirv_module.build_options,
+        })
+    }
+
+    pub fn compile(&self, ctx: &mut l0::Context, dev: &l0::Device) -> Result<l0::Module, CUresult> {
+        let byte_il = unsafe {
+            slice::from_raw_parts(
+                self.binaries.as_ptr() as *const u8,
+                self.binaries.len() * mem::size_of::<u32>(),
+            )
+        };
+        let l0_module = match self.should_link_ptx_impl {
+            None => {
+                l0::Module::build_spirv(ctx, dev, byte_il, Some(self.build_options.as_c_str())).0
+            }
+            Some(ptx_impl) => {
+                l0::Module::build_link_spirv(
+                    ctx,
+                    &dev,
+                    &[ptx_impl, byte_il],
+                    Some(self.build_options.as_c_str()),
+                )
+                .0
+            }
+        };
+        Ok(l0_module?)
+    }
+}
+
+pub fn get_function(
+    hfunc: *mut *mut Function,
+    hmod: *mut Module,
+    name: *const c_char,
+) -> Result<(), CUresult> {
+    if hfunc == ptr::null_mut() || hmod == ptr::null_mut() || name == ptr::null() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let name = unsafe { CStr::from_ptr(name) }.to_owned();
+    let function: *mut Function = GlobalState::lock_current_context(|ctx| {
+        let module = unsafe { &mut *hmod }.as_result_mut()?;
+        let device = unsafe { &mut *ctx.device };
+        let compiled_module = match module.device_binaries.entry(device.index) {
+            hash_map::Entry::Occupied(entry) => entry.into_mut(),
+            hash_map::Entry::Vacant(entry) => {
+                let new_module = CompiledModule {
+                    base: module.spirv.compile(&mut device.l0_context, &device.base)?,
+                    kernels: HashMap::new(),
+                };
+                entry.insert(new_module)
+            }
+        };
+        let kernel = match compiled_module.kernels.entry(name) {
+            hash_map::Entry::Occupied(entry) => entry.into_mut().as_mut(),
+            hash_map::Entry::Vacant(entry) => {
+                let kernel_info = module
+                    .spirv
+                    .kernel_info
+                    .get(unsafe {
+                        std::str::from_utf8_unchecked(entry.key().as_c_str().to_bytes())
+                    })
+                    .ok_or(CUresult::CUDA_ERROR_NOT_FOUND)?;
+                let mut kernel =
+                    l0::Kernel::new_resident(&compiled_module.base, entry.key().as_c_str())?;
+                kernel.set_indirect_access(
+                    l0::sys::ze_kernel_indirect_access_flags_t::ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE
+                    | l0::sys::ze_kernel_indirect_access_flags_t::ZE_KERNEL_INDIRECT_ACCESS_FLAG_HOST
+                    | l0::sys::ze_kernel_indirect_access_flags_t::ZE_KERNEL_INDIRECT_ACCESS_FLAG_SHARED
+                )?;
+                entry.insert(Box::new(Function::new(FunctionData {
+                    base: kernel,
+                    arg_size: kernel_info.arguments_sizes.clone(),
+                    use_shared_mem: kernel_info.uses_shared_mem,
+                    properties: None,
+                })))
+            }
+        };
+        Ok::<_, CUresult>(kernel as *mut _)
+    })??;
+    unsafe { *hfunc = function };
+    Ok(())
+}
+
+pub(crate) fn load_data(pmod: *mut *mut Module, image: *const c_void) -> Result<(), CUresult> {
+    let spirv_data = SpirvModule::new_raw(image as *const _)?;
+    load_data_impl(pmod, spirv_data)
+}
+
+pub fn load_data_impl(pmod: *mut *mut Module, spirv_data: SpirvModule) -> Result<(), CUresult> {
+    let module = GlobalState::lock_current_context(|ctx| {
+        let device = unsafe { &mut *ctx.device };
+        let l0_module = spirv_data.compile(&mut device.l0_context, &device.base)?;
+        let mut device_binaries = HashMap::new();
+        let compiled_module = CompiledModule {
+            base: l0_module,
+            kernels: HashMap::new(),
+        };
+        device_binaries.insert(device.index, compiled_module);
+        let module_data = ModuleData {
+            spirv: spirv_data,
+            device_binaries,
+        };
+        Ok::<_, CUresult>(module_data)
+    })??;
+    let module_ptr = Box::into_raw(Box::new(Module::new(module)));
+    unsafe { *pmod = module_ptr };
+    Ok(())
+}
+
+pub(crate) fn unload(module: *mut Module) -> Result<(), CUresult> {
+    if module == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    GlobalState::lock(|_| Module::destroy_impl(module))?
+}
diff --git a/zluda/src/impl/stream.rs b/zluda/src/impl/stream.rs
new file mode 100644
index 0000000..e212dfc
--- /dev/null
+++ b/zluda/src/impl/stream.rs
@@ -0,0 +1,242 @@
+use super::{
+    context::{Context, ContextData},
+    CUresult, GlobalState,
+};
+use std::{mem, ptr};
+
+use super::{HasLivenessCookie, LiveCheck};
+
+pub type Stream = LiveCheck<StreamData>;
+
+pub const CU_STREAM_LEGACY: *mut Stream = 1 as *mut _;
+pub const CU_STREAM_PER_THREAD: *mut Stream = 2 as *mut _;
+
+impl HasLivenessCookie for StreamData {
+    #[cfg(target_pointer_width = "64")]
+    const COOKIE: usize = 0x512097354de18d35;
+
+    #[cfg(target_pointer_width = "32")]
+    const COOKIE: usize = 0x77d5cc0b;
+
+    const LIVENESS_FAIL: CUresult = CUresult::CUDA_ERROR_INVALID_HANDLE;
+
+    fn try_drop(&mut self) -> Result<(), CUresult> {
+        if self.context != ptr::null_mut() {
+            let context = unsafe { &mut *self.context };
+            if !context.streams.remove(&(self as *mut _)) {
+                return Err(CUresult::CUDA_ERROR_UNKNOWN);
+            }
+        }
+        Ok(())
+    }
+}
+
+pub struct StreamData {
+    pub context: *mut ContextData,
+    pub queue: l0::CommandQueue,
+}
+
+impl StreamData {
+    pub fn new_unitialized(ctx: &mut l0::Context, dev: &l0::Device) -> Result<Self, CUresult> {
+        Ok(StreamData {
+            context: ptr::null_mut(),
+            queue: l0::CommandQueue::new(ctx, dev)?,
+        })
+    }
+    pub fn new(ctx: &mut ContextData) -> Result<Self, CUresult> {
+        let l0_ctx = &mut unsafe { &mut *ctx.device }.l0_context;
+        let l0_dev = &unsafe { &*ctx.device }.base;
+        Ok(StreamData {
+            context: ctx as *mut _,
+            queue: l0::CommandQueue::new(l0_ctx, l0_dev)?,
+        })
+    }
+
+    pub fn command_list(&self) -> Result<l0::CommandList, l0::sys::_ze_result_t> {
+        let ctx = unsafe { &mut *self.context };
+        let dev = unsafe { &mut *ctx.device };
+        l0::CommandList::new(&mut dev.l0_context, &dev.base)
+    }
+}
+
+impl Drop for StreamData {
+    fn drop(&mut self) {
+        if self.context == ptr::null_mut() {
+            return;
+        }
+        unsafe { (&mut *self.context).streams.remove(&(&mut *self as *mut _)) };
+    }
+}
+
+pub(crate) fn get_ctx(hstream: *mut Stream, pctx: *mut *mut Context) -> Result<(), CUresult> {
+    if pctx == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    let ctx_ptr = GlobalState::lock_stream(hstream, |stream| stream.context)?;
+    if ctx_ptr == ptr::null_mut() {
+        return Err(CUresult::CUDA_ERROR_CONTEXT_IS_DESTROYED);
+    }
+    unsafe { *pctx = Context::ptr_from_inner(ctx_ptr) };
+    Ok(())
+}
+
+pub(crate) fn create(phstream: *mut *mut Stream, _flags: u32) -> Result<(), CUresult> {
+    let stream_ptr = GlobalState::lock_current_context(|ctx| {
+        let mut stream_box = Box::new(Stream::new(StreamData::new(ctx)?));
+        let stream_ptr = stream_box.as_mut().as_option_mut().unwrap() as *mut _;
+        if !ctx.streams.insert(stream_ptr) {
+            return Err(CUresult::CUDA_ERROR_UNKNOWN);
+        }
+        mem::forget(stream_box);
+        Ok::<_, CUresult>(stream_ptr)
+    })??;
+    unsafe { *phstream = Stream::ptr_from_inner(stream_ptr) };
+    Ok(())
+}
+
+pub(crate) fn destroy_v2(pstream: *mut Stream) -> Result<(), CUresult> {
+    if pstream == ptr::null_mut() || pstream == CU_STREAM_LEGACY || pstream == CU_STREAM_PER_THREAD
+    {
+        return Err(CUresult::CUDA_ERROR_INVALID_VALUE);
+    }
+    GlobalState::lock(|_| Stream::destroy_impl(pstream))?
+}
+
+#[cfg(test)]
+mod test {
+    use crate::cuda::CUstream;
+
+    use super::super::test::CudaDriverFns;
+    use super::super::CUresult;
+    use std::{ptr, thread};
+
+    const CU_STREAM_LEGACY: CUstream = 1 as *mut _;
+    const CU_STREAM_PER_THREAD: CUstream = 2 as *mut _;
+
+    cuda_driver_test!(default_stream_uses_current_ctx_legacy);
+    cuda_driver_test!(default_stream_uses_current_ctx_ptsd);
+
+    fn default_stream_uses_current_ctx_legacy<T: CudaDriverFns>() {
+        default_stream_uses_current_ctx_impl::<T>(CU_STREAM_LEGACY);
+    }
+
+    fn default_stream_uses_current_ctx_ptsd<T: CudaDriverFns>() {
+        default_stream_uses_current_ctx_impl::<T>(CU_STREAM_PER_THREAD);
+    }
+
+    fn default_stream_uses_current_ctx_impl<T: CudaDriverFns>(stream: CUstream) {
+        assert_eq!(T::cuInit(0), CUresult::CUDA_SUCCESS);
+        let mut ctx1 = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx1, 0, 0), CUresult::CUDA_SUCCESS);
+        let mut stream_ctx1 = ptr::null_mut();
+        assert_eq!(
+            T::cuStreamGetCtx(stream, &mut stream_ctx1),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(ctx1, stream_ctx1);
+        let mut ctx2 = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx2, 0, 0), CUresult::CUDA_SUCCESS);
+        assert_ne!(ctx1, ctx2);
+        let mut stream_ctx2 = ptr::null_mut();
+        assert_eq!(
+            T::cuStreamGetCtx(stream, &mut stream_ctx2),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(ctx2, stream_ctx2);
+        //  Cleanup
+        assert_eq!(T::cuCtxDestroy_v2(ctx1), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxDestroy_v2(ctx2), CUresult::CUDA_SUCCESS);
+    }
+
+    cuda_driver_test!(stream_context_destroyed);
+
+    fn stream_context_destroyed<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 stream = ptr::null_mut();
+        assert_eq!(T::cuStreamCreate(&mut stream, 0), CUresult::CUDA_SUCCESS);
+        let mut stream_ctx1 = ptr::null_mut();
+        assert_eq!(
+            T::cuStreamGetCtx(stream, &mut stream_ctx1),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(stream_ctx1, ctx);
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+        let mut stream_ctx2 = ptr::null_mut();
+        // When a context gets destroyed, its streams are also destroyed
+        let cuda_result = T::cuStreamGetCtx(stream, &mut stream_ctx2);
+        assert!(
+            cuda_result == CUresult::CUDA_ERROR_INVALID_HANDLE
+                || cuda_result == CUresult::CUDA_ERROR_INVALID_CONTEXT
+                || cuda_result == CUresult::CUDA_ERROR_CONTEXT_IS_DESTROYED
+        );
+        assert_eq!(
+            T::cuStreamDestroy_v2(stream),
+            CUresult::CUDA_ERROR_INVALID_HANDLE
+        );
+        // Check if creating another context is possible
+        let mut ctx2 = ptr::null_mut();
+        assert_eq!(T::cuCtxCreate_v2(&mut ctx2, 0, 0), CUresult::CUDA_SUCCESS);
+        //  Cleanup
+        assert_eq!(T::cuCtxDestroy_v2(ctx2), CUresult::CUDA_SUCCESS);
+    }
+
+    cuda_driver_test!(stream_moves_context_to_another_thread);
+
+    fn stream_moves_context_to_another_thread<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 stream = ptr::null_mut();
+        assert_eq!(T::cuStreamCreate(&mut stream, 0), CUresult::CUDA_SUCCESS);
+        let mut stream_ctx1 = ptr::null_mut();
+        assert_eq!(
+            T::cuStreamGetCtx(stream, &mut stream_ctx1),
+            CUresult::CUDA_SUCCESS
+        );
+        assert_eq!(stream_ctx1, ctx);
+        let stream_ptr = stream as usize;
+        let stream_ctx_on_thread = thread::spawn(move || {
+            let mut stream_ctx2 = ptr::null_mut();
+            assert_eq!(
+                T::cuStreamGetCtx(stream_ptr as *mut _, &mut stream_ctx2),
+                CUresult::CUDA_SUCCESS
+            );
+            stream_ctx2 as usize
+        })
+        .join()
+        .unwrap();
+        assert_eq!(stream_ctx1, stream_ctx_on_thread as *mut _);
+        //  Cleanup
+        assert_eq!(T::cuStreamDestroy_v2(stream), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+    }
+
+    cuda_driver_test!(can_destroy_stream);
+
+    fn can_destroy_stream<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 stream = ptr::null_mut();
+        assert_eq!(T::cuStreamCreate(&mut stream, 0), CUresult::CUDA_SUCCESS);
+        assert_eq!(T::cuStreamDestroy_v2(stream), CUresult::CUDA_SUCCESS);
+        // Cleanup
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+    }
+
+    cuda_driver_test!(cant_destroy_default_stream);
+
+    fn cant_destroy_default_stream<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_ne!(
+            T::cuStreamDestroy_v2(super::CU_STREAM_LEGACY as *mut _),
+            CUresult::CUDA_SUCCESS
+        );
+        // Cleanup
+        assert_eq!(T::cuCtxDestroy_v2(ctx), CUresult::CUDA_SUCCESS);
+    }
+}
diff --git a/zluda/src/impl/test.rs b/zluda/src/impl/test.rs
new file mode 100644
index 0000000..b36ccd8
--- /dev/null
+++ b/zluda/src/impl/test.rs
@@ -0,0 +1,157 @@
+#![allow(non_snake_case)]
+
+use crate::cuda as zluda;
+use crate::cuda::CUstream;
+use crate::cuda::CUuuid;
+use crate::{
+    cuda::{CUdevice, CUdeviceptr},
+    r#impl::CUresult,
+};
+use ::std::{
+    ffi::c_void,
+    os::raw::{c_int, c_uint},
+};
+use cuda_driver_sys as cuda;
+
+#[macro_export]
+macro_rules! cuda_driver_test {
+    ($func:ident) => {
+        paste! {
+            #[test]
+            fn [<$func _zluda>]() {
+                $func::<crate::r#impl::test::Zluda>()
+            }
+
+            #[test]
+            fn [<$func _cuda>]() {
+                $func::<crate::r#impl::test::Cuda>()
+            }
+        }
+    };
+}
+
+pub trait CudaDriverFns {
+    fn cuInit(flags: c_uint) -> CUresult;
+    fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult;
+    fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult;
+    fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult;
+    fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult;
+    fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult;
+    fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult;
+    fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult;
+    fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult;
+    fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult;
+    fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult;
+    fn cuMemFree_v2(mem: *mut c_void) -> CUresult;
+    fn cuStreamDestroy_v2(stream: CUstream) -> CUresult;
+}
+
+pub struct Zluda();
+
+impl CudaDriverFns for Zluda {
+    fn cuInit(_flags: c_uint) -> CUresult {
+        zluda::cuInit(_flags as _)
+    }
+
+    fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult {
+        zluda::cuCtxCreate_v2(pctx as *mut _, flags, CUdevice(dev))
+    }
+
+    fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult {
+        zluda::cuCtxDestroy_v2(ctx as *mut _)
+    }
+
+    fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult {
+        zluda::cuCtxPopCurrent_v2(pctx as *mut _)
+    }
+
+    fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult {
+        zluda::cuCtxGetApiVersion(ctx as *mut _, version)
+    }
+
+    fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult {
+        zluda::cuCtxGetCurrent(pctx as *mut _)
+    }
+    fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult {
+        zluda::cuMemAlloc_v2(dptr as *mut _, bytesize)
+    }
+
+    fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult {
+        zluda::cuDeviceGetUuid(uuid, CUdevice(dev))
+    }
+
+    fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult {
+        zluda::cuDevicePrimaryCtxGetState(CUdevice(dev), flags, active)
+    }
+
+    fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult {
+        zluda::cuStreamGetCtx(hStream, pctx as _)
+    }
+
+    fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult {
+        zluda::cuStreamCreate(stream, flags)
+    }
+
+    fn cuMemFree_v2(dptr: *mut c_void) -> CUresult {
+        zluda::cuMemFree_v2(CUdeviceptr(dptr as _))
+    }
+
+    fn cuStreamDestroy_v2(stream: CUstream) -> CUresult {
+        zluda::cuStreamDestroy_v2(stream)
+    }
+}
+
+pub struct Cuda();
+
+impl CudaDriverFns for Cuda {
+    fn cuInit(flags: c_uint) -> CUresult {
+        unsafe { CUresult(cuda::cuInit(flags) as c_uint) }
+    }
+
+    fn cuCtxCreate_v2(pctx: *mut *mut c_void, flags: c_uint, dev: c_int) -> CUresult {
+        unsafe { CUresult(cuda::cuCtxCreate_v2(pctx as *mut _, flags, dev) as c_uint) }
+    }
+
+    fn cuCtxDestroy_v2(ctx: *mut c_void) -> CUresult {
+        unsafe { CUresult(cuda::cuCtxDestroy_v2(ctx as *mut _) as c_uint) }
+    }
+
+    fn cuCtxPopCurrent_v2(pctx: *mut *mut c_void) -> CUresult {
+        unsafe { CUresult(cuda::cuCtxPopCurrent_v2(pctx as *mut _) as c_uint) }
+    }
+
+    fn cuCtxGetApiVersion(ctx: *mut c_void, version: *mut c_uint) -> CUresult {
+        unsafe { CUresult(cuda::cuCtxGetApiVersion(ctx as *mut _, version) as c_uint) }
+    }
+
+    fn cuCtxGetCurrent(pctx: *mut *mut c_void) -> CUresult {
+        unsafe { CUresult(cuda::cuCtxGetCurrent(pctx as *mut _) as c_uint) }
+    }
+    fn cuMemAlloc_v2(dptr: *mut *mut c_void, bytesize: usize) -> CUresult {
+        unsafe { CUresult(cuda::cuMemAlloc_v2(dptr as *mut _, bytesize) as c_uint) }
+    }
+
+    fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: c_int) -> CUresult {
+        unsafe { CUresult(cuda::cuDeviceGetUuid(uuid as *mut _, dev) as c_uint) }
+    }
+
+    fn cuDevicePrimaryCtxGetState(dev: c_int, flags: *mut c_uint, active: *mut c_int) -> CUresult {
+        unsafe { CUresult(cuda::cuDevicePrimaryCtxGetState(dev, flags, active) as c_uint) }
+    }
+
+    fn cuStreamGetCtx(hStream: CUstream, pctx: *mut *mut c_void) -> CUresult {
+        unsafe { CUresult(cuda::cuStreamGetCtx(hStream as _, pctx as _) as c_uint) }
+    }
+
+    fn cuStreamCreate(stream: *mut CUstream, flags: c_uint) -> CUresult {
+        unsafe { CUresult(cuda::cuStreamCreate(stream as _, flags as _) as c_uint) }
+    }
+
+    fn cuMemFree_v2(mem: *mut c_void) -> CUresult {
+        unsafe { CUresult(cuda::cuMemFree_v2(mem as _) as c_uint) }
+    }
+
+    fn cuStreamDestroy_v2(stream: CUstream) -> CUresult {
+        unsafe { CUresult(cuda::cuStreamDestroy_v2(stream as _) as c_uint) }
+    }
+}
diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs
new file mode 100644
index 0000000..0f7d014
--- /dev/null
+++ b/zluda/src/lib.rs
@@ -0,0 +1,16 @@
+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;
+extern crate ptx;
+
+#[allow(warnings)]
+mod cuda;
+mod cuda_impl;
+pub(crate) mod r#impl;