Add a library for dumping kernels arguments before and after launch (#18)

16 files changed, 4951 insertions, 66 deletions

diff --git a/.github/workflows/rust.yml b/.github/workflows/rust.yml
index 734052f..435d25f 100644
--- a/.github/workflows/rust.yml
+++ b/.github/workflows/rust.yml
@@ -48,7 +48,7 @@ jobs:
         sudo apt update
         sudo apt install ocl-icd-opencl-dev
     - name: Build
-      run: cargo build --verbose
+      run: cargo build --workspace --verbose
     # TODO(take-cheeze): Support testing
     # - name: Run tests
     #   run: cargo test --verbose
diff --git a/Cargo.toml b/Cargo.toml
index 9562731..038af15 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -6,6 +6,7 @@ members = [
     "level_zero",

     "spirv_tools-sys",

     "zluda",

+    "zluda_dump",

     "zluda_lib",

     "zluda_inject",

     "zluda_redirect",

diff --git a/detours-sys/build.rs b/detours-sys/build.rs
index f3da0e6..20bb342 100644
--- a/detours-sys/build.rs
+++ b/detours-sys/build.rs
@@ -1,10 +1,8 @@
-use std::error::Error;
-
 #[cfg(not(target_os = "windows"))]
 fn main() {}
 
 #[cfg(target_os = "windows")]
-fn main() -> Result<(), Box<dyn Error>> {
+fn main() -> Result<(), Box<dyn std::error::Error>> {
     windows::main_impl()
 }
 
@@ -37,18 +35,15 @@ mod windows {
         .try_compile("detours")?;
         Ok(())
     }
-
-    #[cfg(target_env = "msvc")]
-    fn add_target_options(build: &mut cc::Build) -> &mut cc::Build {
-        build
-    }
-
-    #[cfg(not(target_env = "msvc"))]
     fn add_target_options(build: &mut cc::Build) -> &mut cc::Build {
-        build
-            .compiler("clang")
-            .cpp(true)
-            .flag("-fms-extensions")
-            .flag("-Wno-everything")
+        if std::env::var("CARGO_CFG_TARGET_ENV").unwrap() != "msvc" {
+            build
+                .compiler("clang")
+                .cpp(true)
+                .flag("-fms-extensions")
+                .flag("-Wno-everything")
+        } else {
+            build
+        }
     }
 }
diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs
index aba6bda..d81cd66 100644
--- a/ptx/src/ast.rs
+++ b/ptx/src/ast.rs
@@ -294,6 +294,17 @@ pub enum KernelArgumentType {
     Shared,
 }
 
+impl From<KernelArgumentType> for Type {
+    fn from(this: KernelArgumentType) -> Self {
+        match this {
+            KernelArgumentType::Normal(typ) => typ.into(),
+            KernelArgumentType::Shared => {
+                Type::Pointer(PointerType::Scalar(ScalarType::B8), LdStateSpace::Shared)
+            }
+        }
+    }
+}
+
 impl FnArgumentType {
     pub fn to_type(&self, is_kernel: bool) -> Type {
         if is_kernel {
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 700bdcd..471a2d7 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -6039,7 +6039,7 @@ impl ast::Type {
         }

     }

 

-    fn size_of(&self) -> usize {

+    pub fn size_of(&self) -> usize {

         match self {

             ast::Type::Scalar(typ) => typ.size_of() as usize,

             ast::Type::Vector(typ, len) => (typ.size_of() as usize) * (*len as usize),

@@ -6253,18 +6253,6 @@ impl<'a> ast::Instruction<ast::ParsedArgParams<'a>> {
     }

 }

 

-impl From<ast::KernelArgumentType> for ast::Type {

-    fn from(this: ast::KernelArgumentType) -> Self {

-        match this {

-            ast::KernelArgumentType::Normal(typ) => typ.into(),

-            ast::KernelArgumentType::Shared => ast::Type::Pointer(

-                ast::PointerType::Scalar(ast::ScalarType::B8),

-                ast::LdStateSpace::Shared,

-            ),

-        }

-    }

-}

-

 impl<T: ArgParamsEx> ast::Arg1<T> {

     fn map<U: ArgParamsEx, V: ArgumentMapVisitor<T, U>>(

         self,

diff --git a/zluda/Cargo.toml b/zluda/Cargo.toml
index 52a3445..6e0d077 100644
--- a/zluda/Cargo.toml
+++ b/zluda/Cargo.toml
@@ -13,7 +13,7 @@ level_zero = { path = "../level_zero" }
 level_zero-sys = { path = "../level_zero-sys" }
 lazy_static = "1.4"
 num_enum = "0.4"
-lz4 = "1.23"
+lz4-sys = "1.9"
 
 [dev-dependencies]
 cuda-driver-sys = "0.3.0"
diff --git a/zluda/src/impl/export_table.rs b/zluda/src/impl/export_table.rs
index 87d7f40..d3ae82d 100644
--- a/zluda/src/impl/export_table.rs
+++ b/zluda/src/impl/export_table.rs
@@ -5,12 +5,12 @@ use crate::{
 };

 

 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,

+    ptr,

 };

+use std::{mem, os::raw::c_int};

 

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

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

@@ -177,6 +177,7 @@ const FATBIN_FILE_HEADER_VERSION_CURRENT: c_ushort = 0x101;
 

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

 #[repr(C)]

+#[derive(Debug)]

 struct FatbinFileHeader {

     kind: c_ushort,

     version: c_ushort,

@@ -221,12 +222,10 @@ unsafe extern "C" fn get_module_from_cubin(
     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 = match decompress_kernel_module(file) {

+            None => continue,

+            Some(vec) => vec,

+        };

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

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

                 Ok(s) => s,

@@ -264,6 +263,33 @@ unsafe fn get_ptx_files(file: *const u8, end: *const u8) -> Vec<*const FatbinFil
     result

 }

 

+const MAX_PTX_MODULE_DECOMPRESSION_BOUND: usize = 16 * 1024 * 1024;

+

+unsafe fn decompress_kernel_module(file: *const FatbinFileHeader) -> Option<Vec<u8>> {

+    let decompressed_size = usize::max(1024, (*file).uncompressed_payload as usize);

+    let mut decompressed_vec = vec![0u8; decompressed_size];

+    loop {

+        match lz4_sys::LZ4_decompress_safe(

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

+            decompressed_vec.as_mut_ptr() as *mut _,

+            (*file).payload_size as c_int,

+            decompressed_vec.len() as c_int,

+        ) {

+            error if error < 0 => {

+                let new_size = decompressed_vec.len() * 2;

+                if new_size > MAX_PTX_MODULE_DECOMPRESSION_BOUND {

+                    return None;

+                }

+                decompressed_vec.resize(decompressed_vec.len() * 2, 0);

+            }

+            real_decompressed_size => {

+                decompressed_vec.truncate(real_decompressed_size as usize);

+                return Some(decompressed_vec);

+            }

+        }

+    }

+}

+

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

 

 const TOOLS_TLS_GUID: CUuuid = CUuuid {

diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs
index ed0bf68..c0ddd5b 100644
--- a/zluda/src/lib.rs
+++ b/zluda/src/lib.rs
@@ -4,7 +4,6 @@ extern crate level_zero_sys as l0_sys;
 extern crate lazy_static;
 #[cfg(test)]
 extern crate cuda_driver_sys;
-extern crate lz4;
 #[cfg(test)]
 #[macro_use]
 extern crate paste;
diff --git a/zluda_dump/Cargo.toml b/zluda_dump/Cargo.toml
new file mode 100644
index 0000000..b81ef13
--- /dev/null
+++ b/zluda_dump/Cargo.toml
@@ -0,0 +1,22 @@
+[package]

+name = "zluda_dump"

+version = "0.0.0"

+authors = ["Andrzej Janik <[email protected]>"]

+edition = "2018"

+

+[lib]

+name = "zluda_dump"

+crate-type = ["cdylib"]

+

+[dependencies]

+ptx = { path = "../ptx" }

+lz4-sys = "1.9"

+regex = "1.4"

+

+[target.'cfg(windows)'.dependencies]

+winapi = { version = "0.3", features = ["libloaderapi", "debugapi"] }

+wchar = "0.6"

+detours-sys = { path = "../detours-sys" }

+

+[target.'cfg(not(windows))'.dependencies]

+libc = "0.2"

diff --git a/zluda_dump/src/cuda.rs b/zluda_dump/src/cuda.rs
new file mode 100644
index 0000000..208b59b
--- /dev/null
+++ b/zluda_dump/src/cuda.rs
@@ -0,0 +1,4072 @@
+/* automatically generated by rust-bindgen 0.56.0 */
+
+pub type __uint32_t = ::std::os::raw::c_uint;
+pub type __uint64_t = ::std::os::raw::c_ulong;
+pub type cuuint32_t = u32;
+pub type cuuint64_t = u64;
+#[repr(transparent)]
+#[derive(Copy, Clone)]
+pub struct CUdeviceptr(pub usize);
+#[repr(transparent)]
+#[derive(Copy, Clone)]
+pub struct CUdevice(pub ::std::os::raw::c_int);
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUctx_st {
+    _unused: [u8; 0],
+}
+pub type CUcontext = *mut CUctx_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUmod_st {
+    _unused: [u8; 0],
+}
+pub type CUmodule = *mut CUmod_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUfunc_st {
+    _unused: [u8; 0],
+}
+pub type CUfunction = *mut CUfunc_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUarray_st {
+    _unused: [u8; 0],
+}
+pub type CUarray = *mut CUarray_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUmipmappedArray_st {
+    _unused: [u8; 0],
+}
+pub type CUmipmappedArray = *mut CUmipmappedArray_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUtexref_st {
+    _unused: [u8; 0],
+}
+pub type CUtexref = *mut CUtexref_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUsurfref_st {
+    _unused: [u8; 0],
+}
+pub type CUsurfref = *mut CUsurfref_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUevent_st {
+    _unused: [u8; 0],
+}
+pub type CUevent = *mut CUevent_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUstream_st {
+    _unused: [u8; 0],
+}
+pub type CUstream = *mut CUstream_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUgraphicsResource_st {
+    _unused: [u8; 0],
+}
+pub type CUgraphicsResource = *mut CUgraphicsResource_st;
+pub type CUtexObject = ::std::os::raw::c_ulonglong;
+pub type CUsurfObject = ::std::os::raw::c_ulonglong;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUextMemory_st {
+    _unused: [u8; 0],
+}
+pub type CUexternalMemory = *mut CUextMemory_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUextSemaphore_st {
+    _unused: [u8; 0],
+}
+pub type CUexternalSemaphore = *mut CUextSemaphore_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUgraph_st {
+    _unused: [u8; 0],
+}
+pub type CUgraph = *mut CUgraph_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUgraphNode_st {
+    _unused: [u8; 0],
+}
+pub type CUgraphNode = *mut CUgraphNode_st;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUgraphExec_st {
+    _unused: [u8; 0],
+}
+pub type CUgraphExec = *mut CUgraphExec_st;
+#[repr(C)]
+#[derive(Copy, Clone, PartialEq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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_SPARSE_CUDA_ARRAY_SUPPORTED: CUdevice_attribute_enum =
+        CUdevice_attribute_enum(112);
+}
+impl CUdevice_attribute_enum {
+    pub const CU_DEVICE_ATTRIBUTE_READ_ONLY_HOST_REGISTER_SUPPORTED: CUdevice_attribute_enum =
+        CUdevice_attribute_enum(113);
+}
+impl CUdevice_attribute_enum {
+    pub const CU_DEVICE_ATTRIBUTE_MAX: CUdevice_attribute_enum = CUdevice_attribute_enum(114);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+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);
+}
+impl CUpointer_attribute_enum {
+    pub const CU_POINTER_ATTRIBUTE_ACCESS_FLAGS: CUpointer_attribute_enum =
+        CUpointer_attribute_enum(16);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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);
+}
+impl CUgraphNodeType_enum {
+    pub const CU_GRAPH_NODE_TYPE_WAIT_EVENT: CUgraphNodeType_enum = CUgraphNodeType_enum(6);
+}
+impl CUgraphNodeType_enum {
+    pub const CU_GRAPH_NODE_TYPE_EVENT_RECORD: CUgraphNodeType_enum = CUgraphNodeType_enum(7);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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_STUB_LIBRARY: cudaError_enum = cudaError_enum(34);
+}
+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_DEVICE_NOT_LICENSED: cudaError_enum = cudaError_enum(102);
+}
+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_UNSUPPORTED_PTX_VERSION: cudaError_enum = cudaError_enum(222);
+}
+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(Copy, Clone, Hash, PartialEq, Eq, Debug)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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_ARRAY_SPARSE_PROPERTIES_st {
+    pub tileExtent: CUDA_ARRAY_SPARSE_PROPERTIES_st__bindgen_ty_1,
+    pub miptailFirstLevel: ::std::os::raw::c_uint,
+    pub miptailSize: ::std::os::raw::c_ulonglong,
+    pub flags: ::std::os::raw::c_uint,
+    pub reserved: [::std::os::raw::c_uint; 4usize],
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUDA_ARRAY_SPARSE_PROPERTIES_st__bindgen_ty_1 {
+    pub width: ::std::os::raw::c_uint,
+    pub height: ::std::os::raw::c_uint,
+    pub depth: ::std::os::raw::c_uint,
+}
+pub type CUDA_ARRAY_SPARSE_PROPERTIES = CUDA_ARRAY_SPARSE_PROPERTIES_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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+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(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUmemAllocationGranularity_flags_enum(pub ::std::os::raw::c_uint);
+pub use self::CUmemAllocationGranularity_flags_enum as CUmemAllocationGranularity_flags;
+impl CUarraySparseSubresourceType_enum {
+    pub const CU_ARRAY_SPARSE_SUBRESOURCE_TYPE_SPARSE_LEVEL: CUarraySparseSubresourceType_enum =
+        CUarraySparseSubresourceType_enum(0);
+}
+impl CUarraySparseSubresourceType_enum {
+    pub const CU_ARRAY_SPARSE_SUBRESOURCE_TYPE_MIPTAIL: CUarraySparseSubresourceType_enum =
+        CUarraySparseSubresourceType_enum(1);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUarraySparseSubresourceType_enum(pub ::std::os::raw::c_uint);
+pub use self::CUarraySparseSubresourceType_enum as CUarraySparseSubresourceType;
+impl CUmemOperationType_enum {
+    pub const CU_MEM_OPERATION_TYPE_MAP: CUmemOperationType_enum = CUmemOperationType_enum(1);
+}
+impl CUmemOperationType_enum {
+    pub const CU_MEM_OPERATION_TYPE_UNMAP: CUmemOperationType_enum = CUmemOperationType_enum(2);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUmemOperationType_enum(pub ::std::os::raw::c_uint);
+pub use self::CUmemOperationType_enum as CUmemOperationType;
+impl CUmemHandleType_enum {
+    pub const CU_MEM_HANDLE_TYPE_GENERIC: CUmemHandleType_enum = CUmemHandleType_enum(0);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUmemHandleType_enum(pub ::std::os::raw::c_uint);
+pub use self::CUmemHandleType_enum as CUmemHandleType;
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUarrayMapInfo_st {
+    pub resourceType: CUresourcetype,
+    pub resource: CUarrayMapInfo_st__bindgen_ty_1,
+    pub subresourceType: CUarraySparseSubresourceType,
+    pub subresource: CUarrayMapInfo_st__bindgen_ty_2,
+    pub memOperationType: CUmemOperationType,
+    pub memHandleType: CUmemHandleType,
+    pub memHandle: CUarrayMapInfo_st__bindgen_ty_3,
+    pub offset: ::std::os::raw::c_ulonglong,
+    pub deviceBitMask: ::std::os::raw::c_uint,
+    pub flags: ::std::os::raw::c_uint,
+    pub reserved: [::std::os::raw::c_uint; 2usize],
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub union CUarrayMapInfo_st__bindgen_ty_1 {
+    pub mipmap: CUmipmappedArray,
+    pub array: CUarray,
+    _bindgen_union_align: u64,
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub union CUarrayMapInfo_st__bindgen_ty_2 {
+    pub sparseLevel: CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_1,
+    pub miptail: CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_2,
+    _bindgen_union_align: [u64; 4usize],
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_1 {
+    pub level: ::std::os::raw::c_uint,
+    pub layer: ::std::os::raw::c_uint,
+    pub offsetX: ::std::os::raw::c_uint,
+    pub offsetY: ::std::os::raw::c_uint,
+    pub offsetZ: ::std::os::raw::c_uint,
+    pub extentWidth: ::std::os::raw::c_uint,
+    pub extentHeight: ::std::os::raw::c_uint,
+    pub extentDepth: ::std::os::raw::c_uint,
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_2 {
+    pub layer: ::std::os::raw::c_uint,
+    pub offset: ::std::os::raw::c_ulonglong,
+    pub size: ::std::os::raw::c_ulonglong,
+}
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub union CUarrayMapInfo_st__bindgen_ty_3 {
+    pub memHandle: CUmemGenericAllocationHandle,
+    _bindgen_union_align: u64,
+}
+pub type CUarrayMapInfo = CUarrayMapInfo_st;
+#[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 usage: ::std::os::raw::c_ushort,
+    pub reserved: [::std::os::raw::c_uchar; 4usize],
+}
+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(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUgraphExecUpdateResult_enum(pub ::std::os::raw::c_uint);
+pub use self::CUgraphExecUpdateResult_enum as CUgraphExecUpdateResult;
+extern_redirect! {
+    pub fn cuGetErrorString(error: CUresult, pStr: *mut *const ::std::os::raw::c_char) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGetErrorName(error: CUresult, pStr: *mut *const ::std::os::raw::c_char) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuInit(Flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDriverGetVersion(driverVersion: *mut ::std::os::raw::c_int) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGet(device: *mut CUdevice, ordinal: ::std::os::raw::c_int) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetCount(count: *mut ::std::os::raw::c_int) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetName(
+        name: *mut ::std::os::raw::c_char,
+        len: ::std::os::raw::c_int,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetUuid(uuid: *mut CUuuid, dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetLuid(
+        luid: *mut ::std::os::raw::c_char,
+        deviceNodeMask: *mut ::std::os::raw::c_uint,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceTotalMem_v2(bytes: *mut usize, dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetTexture1DLinearMaxWidth(
+        maxWidthInElements: *mut usize,
+        format: CUarray_format,
+        numChannels: ::std::os::raw::c_uint,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetAttribute(
+        pi: *mut ::std::os::raw::c_int,
+        attrib: CUdevice_attribute,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetNvSciSyncAttributes(
+        nvSciSyncAttrList: *mut ::std::os::raw::c_void,
+        dev: CUdevice,
+        flags: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetProperties(prop: *mut CUdevprop, dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceComputeCapability(
+        major: *mut ::std::os::raw::c_int,
+        minor: *mut ::std::os::raw::c_int,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxRetain(pctx: *mut CUcontext, dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxRelease(dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxRelease_v2(dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxSetFlags(dev: CUdevice, flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxSetFlags_v2(dev: CUdevice, flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxGetState(
+        dev: CUdevice,
+        flags: *mut ::std::os::raw::c_uint,
+        active: *mut ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxReset(dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDevicePrimaryCtxReset_v2(dev: CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxCreate_v2(
+        pctx: *mut CUcontext,
+        flags: ::std::os::raw::c_uint,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxDestroy_v2(ctx: CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxPushCurrent_v2(ctx: CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxPopCurrent_v2(pctx: *mut CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxSetCurrent(ctx: CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetCurrent(pctx: *mut CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetDevice(device: *mut CUdevice) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetFlags(flags: *mut ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxSynchronize() -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxSetLimit(limit: CUlimit, value: usize) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetLimit(pvalue: *mut usize, limit: CUlimit) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetCacheConfig(pconfig: *mut CUfunc_cache) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxSetCacheConfig(config: CUfunc_cache) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetSharedMemConfig(pConfig: *mut CUsharedconfig) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxSetSharedMemConfig(config: CUsharedconfig) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetApiVersion(ctx: CUcontext, version: *mut ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxGetStreamPriorityRange(
+        leastPriority: *mut ::std::os::raw::c_int,
+        greatestPriority: *mut ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxResetPersistingL2Cache() -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxAttach(pctx: *mut CUcontext, flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxDetach(ctx: CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuModuleLoad(module: *mut CUmodule, fname: *const ::std::os::raw::c_char) -> CUresult;
+}
+extern_redirect_with! {
+    pub fn cuModuleLoadData(
+        module: *mut CUmodule,
+        image: *const ::std::os::raw::c_void,
+    ) -> CUresult;
+    super::cuModuleLoadData;
+}
+extern_redirect_with! {
+    pub 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;
+    super::cuModuleLoadDataEx;
+}
+extern_redirect! {
+    pub fn cuModuleLoadFatBinary(
+        module: *mut CUmodule,
+        fatCubin: *const ::std::os::raw::c_void,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuModuleUnload(hmod: CUmodule) -> CUresult;
+}
+extern_redirect_with! {
+    pub fn cuModuleGetFunction(
+        hfunc: *mut CUfunction,
+        hmod: CUmodule,
+        name: *const ::std::os::raw::c_char,
+    ) -> CUresult;
+    super::cuModuleGetFunction;
+}
+extern_redirect! {
+    pub fn cuModuleGetGlobal_v2(
+        dptr: *mut CUdeviceptr,
+        bytes: *mut usize,
+        hmod: CUmodule,
+        name: *const ::std::os::raw::c_char,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuModuleGetTexRef(
+        pTexRef: *mut CUtexref,
+        hmod: CUmodule,
+        name: *const ::std::os::raw::c_char,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuModuleGetSurfRef(
+        pSurfRef: *mut CUsurfref,
+        hmod: CUmodule,
+        name: *const ::std::os::raw::c_char,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub fn cuLinkComplete(
+        state: CUlinkState,
+        cubinOut: *mut *mut ::std::os::raw::c_void,
+        sizeOut: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuLinkDestroy(state: CUlinkState) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> CUresult;
+}
+extern_redirect_with! {
+    pub fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult;
+    super::cuMemAlloc_v2;
+}
+extern_redirect! {
+    pub fn cuMemAllocPitch_v2(
+        dptr: *mut CUdeviceptr,
+        pPitch: *mut usize,
+        WidthInBytes: usize,
+        Height: usize,
+        ElementSizeBytes: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemGetAddressRange_v2(
+        pbase: *mut CUdeviceptr,
+        psize: *mut usize,
+        dptr: CUdeviceptr,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemAllocHost_v2(pp: *mut *mut ::std::os::raw::c_void, bytesize: usize) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemFreeHost(p: *mut ::std::os::raw::c_void) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemHostAlloc(
+        pp: *mut *mut ::std::os::raw::c_void,
+        bytesize: usize,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemHostGetDevicePointer_v2(
+        pdptr: *mut CUdeviceptr,
+        p: *mut ::std::os::raw::c_void,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemHostGetFlags(
+        pFlags: *mut ::std::os::raw::c_uint,
+        p: *mut ::std::os::raw::c_void,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemAllocManaged(
+        dptr: *mut CUdeviceptr,
+        bytesize: usize,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetByPCIBusId(
+        dev: *mut CUdevice,
+        pciBusId: *const ::std::os::raw::c_char,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetPCIBusId(
+        pciBusId: *mut ::std::os::raw::c_char,
+        len: ::std::os::raw::c_int,
+        dev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcGetEventHandle(pHandle: *mut CUipcEventHandle, event: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcOpenEventHandle(phEvent: *mut CUevent, handle: CUipcEventHandle) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcGetMemHandle(pHandle: *mut CUipcMemHandle, dptr: CUdeviceptr) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcOpenMemHandle(
+        pdptr: *mut CUdeviceptr,
+        handle: CUipcMemHandle,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcOpenMemHandle_v2(
+        pdptr: *mut CUdeviceptr,
+        handle: CUipcMemHandle,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuIpcCloseMemHandle(dptr: CUdeviceptr) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemHostRegister_v2(
+        p: *mut ::std::os::raw::c_void,
+        bytesize: usize,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemHostUnregister(p: *mut ::std::os::raw::c_void) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy(dst: CUdeviceptr, src: CUdeviceptr, ByteCount: usize) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyPeer(
+        dstDevice: CUdeviceptr,
+        dstContext: CUcontext,
+        srcDevice: CUdeviceptr,
+        srcContext: CUcontext,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyHtoD_v2(
+        dstDevice: CUdeviceptr,
+        srcHost: *const ::std::os::raw::c_void,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyDtoH_v2(
+        dstHost: *mut ::std::os::raw::c_void,
+        srcDevice: CUdeviceptr,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyDtoD_v2(
+        dstDevice: CUdeviceptr,
+        srcDevice: CUdeviceptr,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyDtoA_v2(
+        dstArray: CUarray,
+        dstOffset: usize,
+        srcDevice: CUdeviceptr,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyAtoD_v2(
+        dstDevice: CUdeviceptr,
+        srcArray: CUarray,
+        srcOffset: usize,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyHtoA_v2(
+        dstArray: CUarray,
+        dstOffset: usize,
+        srcHost: *const ::std::os::raw::c_void,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyAtoH_v2(
+        dstHost: *mut ::std::os::raw::c_void,
+        srcArray: CUarray,
+        srcOffset: usize,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyAtoA_v2(
+        dstArray: CUarray,
+        dstOffset: usize,
+        srcArray: CUarray,
+        srcOffset: usize,
+        ByteCount: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy2D_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy2DUnaligned_v2(pCopy: *const CUDA_MEMCPY2D) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy3D_v2(pCopy: *const CUDA_MEMCPY3D) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy3DPeer(pCopy: *const CUDA_MEMCPY3D_PEER) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyAsync(
+        dst: CUdeviceptr,
+        src: CUdeviceptr,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyPeerAsync(
+        dstDevice: CUdeviceptr,
+        dstContext: CUcontext,
+        srcDevice: CUdeviceptr,
+        srcContext: CUcontext,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyHtoDAsync_v2(
+        dstDevice: CUdeviceptr,
+        srcHost: *const ::std::os::raw::c_void,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyDtoHAsync_v2(
+        dstHost: *mut ::std::os::raw::c_void,
+        srcDevice: CUdeviceptr,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyDtoDAsync_v2(
+        dstDevice: CUdeviceptr,
+        srcDevice: CUdeviceptr,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyHtoAAsync_v2(
+        dstArray: CUarray,
+        dstOffset: usize,
+        srcHost: *const ::std::os::raw::c_void,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpyAtoHAsync_v2(
+        dstHost: *mut ::std::os::raw::c_void,
+        srcArray: CUarray,
+        srcOffset: usize,
+        ByteCount: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy2DAsync_v2(pCopy: *const CUDA_MEMCPY2D, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy3DAsync_v2(pCopy: *const CUDA_MEMCPY3D, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemcpy3DPeerAsync(pCopy: *const CUDA_MEMCPY3D_PEER, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD8_v2(dstDevice: CUdeviceptr, uc: ::std::os::raw::c_uchar, N: usize)
+        -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD16_v2(
+        dstDevice: CUdeviceptr,
+        us: ::std::os::raw::c_ushort,
+        N: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD32_v2(dstDevice: CUdeviceptr, ui: ::std::os::raw::c_uint, N: usize)
+        -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D8_v2(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        uc: ::std::os::raw::c_uchar,
+        Width: usize,
+        Height: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D16_v2(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        us: ::std::os::raw::c_ushort,
+        Width: usize,
+        Height: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D32_v2(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        ui: ::std::os::raw::c_uint,
+        Width: usize,
+        Height: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD8Async(
+        dstDevice: CUdeviceptr,
+        uc: ::std::os::raw::c_uchar,
+        N: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD16Async(
+        dstDevice: CUdeviceptr,
+        us: ::std::os::raw::c_ushort,
+        N: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD32Async(
+        dstDevice: CUdeviceptr,
+        ui: ::std::os::raw::c_uint,
+        N: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D8Async(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        uc: ::std::os::raw::c_uchar,
+        Width: usize,
+        Height: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D16Async(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        us: ::std::os::raw::c_ushort,
+        Width: usize,
+        Height: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemsetD2D32Async(
+        dstDevice: CUdeviceptr,
+        dstPitch: usize,
+        ui: ::std::os::raw::c_uint,
+        Width: usize,
+        Height: usize,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArrayCreate_v2(
+        pHandle: *mut CUarray,
+        pAllocateArray: *const CUDA_ARRAY_DESCRIPTOR,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArrayGetDescriptor_v2(
+        pArrayDescriptor: *mut CUDA_ARRAY_DESCRIPTOR,
+        hArray: CUarray,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArrayGetSparseProperties(
+        sparseProperties: *mut CUDA_ARRAY_SPARSE_PROPERTIES,
+        array: CUarray,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMipmappedArrayGetSparseProperties(
+        sparseProperties: *mut CUDA_ARRAY_SPARSE_PROPERTIES,
+        mipmap: CUmipmappedArray,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArrayDestroy(hArray: CUarray) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArray3DCreate_v2(
+        pHandle: *mut CUarray,
+        pAllocateArray: *const CUDA_ARRAY3D_DESCRIPTOR,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuArray3DGetDescriptor_v2(
+        pArrayDescriptor: *mut CUDA_ARRAY3D_DESCRIPTOR,
+        hArray: CUarray,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMipmappedArrayCreate(
+        pHandle: *mut CUmipmappedArray,
+        pMipmappedArrayDesc: *const CUDA_ARRAY3D_DESCRIPTOR,
+        numMipmapLevels: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMipmappedArrayGetLevel(
+        pLevelArray: *mut CUarray,
+        hMipmappedArray: CUmipmappedArray,
+        level: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMipmappedArrayDestroy(hMipmappedArray: CUmipmappedArray) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemAddressReserve(
+        ptr: *mut CUdeviceptr,
+        size: usize,
+        alignment: usize,
+        addr: CUdeviceptr,
+        flags: ::std::os::raw::c_ulonglong,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemAddressFree(ptr: CUdeviceptr, size: usize) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemCreate(
+        handle: *mut CUmemGenericAllocationHandle,
+        size: usize,
+        prop: *const CUmemAllocationProp,
+        flags: ::std::os::raw::c_ulonglong,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemRelease(handle: CUmemGenericAllocationHandle) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemMap(
+        ptr: CUdeviceptr,
+        size: usize,
+        offset: usize,
+        handle: CUmemGenericAllocationHandle,
+        flags: ::std::os::raw::c_ulonglong,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemMapArrayAsync(
+        mapInfoList: *mut CUarrayMapInfo,
+        count: ::std::os::raw::c_uint,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemUnmap(ptr: CUdeviceptr, size: usize) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemSetAccess(
+        ptr: CUdeviceptr,
+        size: usize,
+        desc: *const CUmemAccessDesc,
+        count: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemGetAccess(
+        flags: *mut ::std::os::raw::c_ulonglong,
+        location: *const CUmemLocation,
+        ptr: CUdeviceptr,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemExportToShareableHandle(
+        shareableHandle: *mut ::std::os::raw::c_void,
+        handle: CUmemGenericAllocationHandle,
+        handleType: CUmemAllocationHandleType,
+        flags: ::std::os::raw::c_ulonglong,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemImportFromShareableHandle(
+        handle: *mut CUmemGenericAllocationHandle,
+        osHandle: *mut ::std::os::raw::c_void,
+        shHandleType: CUmemAllocationHandleType,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemGetAllocationGranularity(
+        granularity: *mut usize,
+        prop: *const CUmemAllocationProp,
+        option: CUmemAllocationGranularity_flags,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemGetAllocationPropertiesFromHandle(
+        prop: *mut CUmemAllocationProp,
+        handle: CUmemGenericAllocationHandle,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemRetainAllocationHandle(
+        handle: *mut CUmemGenericAllocationHandle,
+        addr: *mut ::std::os::raw::c_void,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuPointerGetAttribute(
+        data: *mut ::std::os::raw::c_void,
+        attribute: CUpointer_attribute,
+        ptr: CUdeviceptr,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemPrefetchAsync(
+        devPtr: CUdeviceptr,
+        count: usize,
+        dstDevice: CUdevice,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemAdvise(
+        devPtr: CUdeviceptr,
+        count: usize,
+        advice: CUmem_advise,
+        device: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuMemRangeGetAttribute(
+        data: *mut ::std::os::raw::c_void,
+        dataSize: usize,
+        attribute: CUmem_range_attribute,
+        devPtr: CUdeviceptr,
+        count: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub fn cuPointerSetAttribute(
+        value: *const ::std::os::raw::c_void,
+        attribute: CUpointer_attribute,
+        ptr: CUdeviceptr,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuPointerGetAttributes(
+        numAttributes: ::std::os::raw::c_uint,
+        attributes: *mut CUpointer_attribute,
+        data: *mut *mut ::std::os::raw::c_void,
+        ptr: CUdeviceptr,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamCreate(phStream: *mut CUstream, Flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamCreateWithPriority(
+        phStream: *mut CUstream,
+        flags: ::std::os::raw::c_uint,
+        priority: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamGetPriority(hStream: CUstream, priority: *mut ::std::os::raw::c_int)
+        -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamGetFlags(hStream: CUstream, flags: *mut ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamWaitEvent(
+        hStream: CUstream,
+        hEvent: CUevent,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamAddCallback(
+        hStream: CUstream,
+        callback: CUstreamCallback,
+        userData: *mut ::std::os::raw::c_void,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamBeginCapture_v2(hStream: CUstream, mode: CUstreamCaptureMode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuThreadExchangeStreamCaptureMode(mode: *mut CUstreamCaptureMode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamEndCapture(hStream: CUstream, phGraph: *mut CUgraph) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamIsCapturing(
+        hStream: CUstream,
+        captureStatus: *mut CUstreamCaptureStatus,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamGetCaptureInfo(
+        hStream: CUstream,
+        captureStatus: *mut CUstreamCaptureStatus,
+        id: *mut cuuint64_t,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamAttachMemAsync(
+        hStream: CUstream,
+        dptr: CUdeviceptr,
+        length: usize,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamQuery(hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamSynchronize(hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamDestroy_v2(hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamCopyAttributes(dst: CUstream, src: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamGetAttribute(
+        hStream: CUstream,
+        attr: CUstreamAttrID,
+        value_out: *mut CUstreamAttrValue,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamSetAttribute(
+        hStream: CUstream,
+        attr: CUstreamAttrID,
+        value: *const CUstreamAttrValue,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventCreate(phEvent: *mut CUevent, Flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventRecord(hEvent: CUevent, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventRecordWithFlags(
+        hEvent: CUevent,
+        hStream: CUstream,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventQuery(hEvent: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventSynchronize(hEvent: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventDestroy_v2(hEvent: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuEventElapsedTime(pMilliseconds: *mut f32, hStart: CUevent, hEnd: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuImportExternalMemory(
+        extMem_out: *mut CUexternalMemory,
+        memHandleDesc: *const CUDA_EXTERNAL_MEMORY_HANDLE_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuExternalMemoryGetMappedBuffer(
+        devPtr: *mut CUdeviceptr,
+        extMem: CUexternalMemory,
+        bufferDesc: *const CUDA_EXTERNAL_MEMORY_BUFFER_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuExternalMemoryGetMappedMipmappedArray(
+        mipmap: *mut CUmipmappedArray,
+        extMem: CUexternalMemory,
+        mipmapDesc: *const CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDestroyExternalMemory(extMem: CUexternalMemory) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuImportExternalSemaphore(
+        extSem_out: *mut CUexternalSemaphore,
+        semHandleDesc: *const CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSignalExternalSemaphoresAsync(
+        extSemArray: *const CUexternalSemaphore,
+        paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS,
+        numExtSems: ::std::os::raw::c_uint,
+        stream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuWaitExternalSemaphoresAsync(
+        extSemArray: *const CUexternalSemaphore,
+        paramsArray: *const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS,
+        numExtSems: ::std::os::raw::c_uint,
+        stream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDestroyExternalSemaphore(extSem: CUexternalSemaphore) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamWaitValue32(
+        stream: CUstream,
+        addr: CUdeviceptr,
+        value: cuuint32_t,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamWaitValue64(
+        stream: CUstream,
+        addr: CUdeviceptr,
+        value: cuuint64_t,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamWriteValue32(
+        stream: CUstream,
+        addr: CUdeviceptr,
+        value: cuuint32_t,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamWriteValue64(
+        stream: CUstream,
+        addr: CUdeviceptr,
+        value: cuuint64_t,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuStreamBatchMemOp(
+        stream: CUstream,
+        count: ::std::os::raw::c_uint,
+        paramArray: *mut CUstreamBatchMemOpParams,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncGetAttribute(
+        pi: *mut ::std::os::raw::c_int,
+        attrib: CUfunction_attribute,
+        hfunc: CUfunction,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncSetAttribute(
+        hfunc: CUfunction,
+        attrib: CUfunction_attribute,
+        value: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncSetCacheConfig(hfunc: CUfunction, config: CUfunc_cache) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncSetSharedMemConfig(hfunc: CUfunction, config: CUsharedconfig) -> CUresult;
+}
+extern_redirect_with! {
+    pub 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;
+    super::cuLaunchKernel;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub fn cuLaunchCooperativeKernelMultiDevice(
+        launchParamsList: *mut CUDA_LAUNCH_PARAMS,
+        numDevices: ::std::os::raw::c_uint,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuLaunchHostFunc(
+        hStream: CUstream,
+        fn_: CUhostFn,
+        userData: *mut ::std::os::raw::c_void,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncSetBlockShape(
+        hfunc: CUfunction,
+        x: ::std::os::raw::c_int,
+        y: ::std::os::raw::c_int,
+        z: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuFuncSetSharedSize(hfunc: CUfunction, bytes: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuParamSetSize(hfunc: CUfunction, numbytes: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuParamSeti(
+        hfunc: CUfunction,
+        offset: ::std::os::raw::c_int,
+        value: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuParamSetf(hfunc: CUfunction, offset: ::std::os::raw::c_int, value: f32) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuParamSetv(
+        hfunc: CUfunction,
+        offset: ::std::os::raw::c_int,
+        ptr: *mut ::std::os::raw::c_void,
+        numbytes: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuLaunch(f: CUfunction) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuLaunchGrid(
+        f: CUfunction,
+        grid_width: ::std::os::raw::c_int,
+        grid_height: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuLaunchGridAsync(
+        f: CUfunction,
+        grid_width: ::std::os::raw::c_int,
+        grid_height: ::std::os::raw::c_int,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuParamSetTexRef(
+        hfunc: CUfunction,
+        texunit: ::std::os::raw::c_int,
+        hTexRef: CUtexref,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphCreate(phGraph: *mut CUgraph, flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddKernelNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphKernelNodeGetParams(
+        hNode: CUgraphNode,
+        nodeParams: *mut CUDA_KERNEL_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphKernelNodeSetParams(
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddMemcpyNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        copyParams: *const CUDA_MEMCPY3D,
+        ctx: CUcontext,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphMemcpyNodeGetParams(
+        hNode: CUgraphNode,
+        nodeParams: *mut CUDA_MEMCPY3D,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphMemcpyNodeSetParams(
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_MEMCPY3D,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddMemsetNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        memsetParams: *const CUDA_MEMSET_NODE_PARAMS,
+        ctx: CUcontext,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphMemsetNodeGetParams(
+        hNode: CUgraphNode,
+        nodeParams: *mut CUDA_MEMSET_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphMemsetNodeSetParams(
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_MEMSET_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddHostNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        nodeParams: *const CUDA_HOST_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphHostNodeGetParams(
+        hNode: CUgraphNode,
+        nodeParams: *mut CUDA_HOST_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphHostNodeSetParams(
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_HOST_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddChildGraphNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        childGraph: CUgraph,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphChildGraphNodeGetGraph(hNode: CUgraphNode, phGraph: *mut CUgraph) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddEmptyNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddEventRecordNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        event: CUevent,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphEventRecordNodeGetEvent(hNode: CUgraphNode, event_out: *mut CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphEventRecordNodeSetEvent(hNode: CUgraphNode, event: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddEventWaitNode(
+        phGraphNode: *mut CUgraphNode,
+        hGraph: CUgraph,
+        dependencies: *const CUgraphNode,
+        numDependencies: usize,
+        event: CUevent,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphEventWaitNodeGetEvent(hNode: CUgraphNode, event_out: *mut CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphEventWaitNodeSetEvent(hNode: CUgraphNode, event: CUevent) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphClone(phGraphClone: *mut CUgraph, originalGraph: CUgraph) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphNodeFindInClone(
+        phNode: *mut CUgraphNode,
+        hOriginalNode: CUgraphNode,
+        hClonedGraph: CUgraph,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphNodeGetType(hNode: CUgraphNode, type_: *mut CUgraphNodeType) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphGetNodes(
+        hGraph: CUgraph,
+        nodes: *mut CUgraphNode,
+        numNodes: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphGetRootNodes(
+        hGraph: CUgraph,
+        rootNodes: *mut CUgraphNode,
+        numRootNodes: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphGetEdges(
+        hGraph: CUgraph,
+        from: *mut CUgraphNode,
+        to: *mut CUgraphNode,
+        numEdges: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphNodeGetDependencies(
+        hNode: CUgraphNode,
+        dependencies: *mut CUgraphNode,
+        numDependencies: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphNodeGetDependentNodes(
+        hNode: CUgraphNode,
+        dependentNodes: *mut CUgraphNode,
+        numDependentNodes: *mut usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphAddDependencies(
+        hGraph: CUgraph,
+        from: *const CUgraphNode,
+        to: *const CUgraphNode,
+        numDependencies: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphRemoveDependencies(
+        hGraph: CUgraph,
+        from: *const CUgraphNode,
+        to: *const CUgraphNode,
+        numDependencies: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphDestroyNode(hNode: CUgraphNode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphInstantiate_v2(
+        phGraphExec: *mut CUgraphExec,
+        hGraph: CUgraph,
+        phErrorNode: *mut CUgraphNode,
+        logBuffer: *mut ::std::os::raw::c_char,
+        bufferSize: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecKernelNodeSetParams(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_KERNEL_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecMemcpyNodeSetParams(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        copyParams: *const CUDA_MEMCPY3D,
+        ctx: CUcontext,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecMemsetNodeSetParams(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        memsetParams: *const CUDA_MEMSET_NODE_PARAMS,
+        ctx: CUcontext,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecHostNodeSetParams(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        nodeParams: *const CUDA_HOST_NODE_PARAMS,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecChildGraphNodeSetParams(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        childGraph: CUgraph,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecEventRecordNodeSetEvent(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        event: CUevent,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecEventWaitNodeSetEvent(
+        hGraphExec: CUgraphExec,
+        hNode: CUgraphNode,
+        event: CUevent,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphUpload(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphLaunch(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecDestroy(hGraphExec: CUgraphExec) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphDestroy(hGraph: CUgraph) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphExecUpdate(
+        hGraphExec: CUgraphExec,
+        hGraph: CUgraph,
+        hErrorNode_out: *mut CUgraphNode,
+        updateResult_out: *mut CUgraphExecUpdateResult,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphKernelNodeCopyAttributes(dst: CUgraphNode, src: CUgraphNode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphKernelNodeGetAttribute(
+        hNode: CUgraphNode,
+        attr: CUkernelNodeAttrID,
+        value_out: *mut CUkernelNodeAttrValue,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphKernelNodeSetAttribute(
+        hNode: CUgraphNode,
+        attr: CUkernelNodeAttrID,
+        value: *const CUkernelNodeAttrValue,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuOccupancyMaxActiveBlocksPerMultiprocessor(
+        numBlocks: *mut ::std::os::raw::c_int,
+        func: CUfunction,
+        blockSize: ::std::os::raw::c_int,
+        dynamicSMemSize: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub 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;
+}
+extern_redirect! {
+    pub fn cuOccupancyAvailableDynamicSMemPerBlock(
+        dynamicSmemSize: *mut usize,
+        func: CUfunction,
+        numBlocks: ::std::os::raw::c_int,
+        blockSize: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetArray(
+        hTexRef: CUtexref,
+        hArray: CUarray,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetMipmappedArray(
+        hTexRef: CUtexref,
+        hMipmappedArray: CUmipmappedArray,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetAddress_v2(
+        ByteOffset: *mut usize,
+        hTexRef: CUtexref,
+        dptr: CUdeviceptr,
+        bytes: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetAddress2D_v3(
+        hTexRef: CUtexref,
+        desc: *const CUDA_ARRAY_DESCRIPTOR,
+        dptr: CUdeviceptr,
+        Pitch: usize,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetFormat(
+        hTexRef: CUtexref,
+        fmt: CUarray_format,
+        NumPackedComponents: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetAddressMode(
+        hTexRef: CUtexref,
+        dim: ::std::os::raw::c_int,
+        am: CUaddress_mode,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetMipmapFilterMode(hTexRef: CUtexref, fm: CUfilter_mode) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetMipmapLevelBias(hTexRef: CUtexref, bias: f32) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetMipmapLevelClamp(
+        hTexRef: CUtexref,
+        minMipmapLevelClamp: f32,
+        maxMipmapLevelClamp: f32,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetMaxAnisotropy(
+        hTexRef: CUtexref,
+        maxAniso: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetBorderColor(hTexRef: CUtexref, pBorderColor: *mut f32) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefSetFlags(hTexRef: CUtexref, Flags: ::std::os::raw::c_uint) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetAddress_v2(pdptr: *mut CUdeviceptr, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetArray(phArray: *mut CUarray, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetMipmappedArray(
+        phMipmappedArray: *mut CUmipmappedArray,
+        hTexRef: CUtexref,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetAddressMode(
+        pam: *mut CUaddress_mode,
+        hTexRef: CUtexref,
+        dim: ::std::os::raw::c_int,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetFilterMode(pfm: *mut CUfilter_mode, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetFormat(
+        pFormat: *mut CUarray_format,
+        pNumChannels: *mut ::std::os::raw::c_int,
+        hTexRef: CUtexref,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetMipmapFilterMode(pfm: *mut CUfilter_mode, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetMipmapLevelBias(pbias: *mut f32, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetMipmapLevelClamp(
+        pminMipmapLevelClamp: *mut f32,
+        pmaxMipmapLevelClamp: *mut f32,
+        hTexRef: CUtexref,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetMaxAnisotropy(
+        pmaxAniso: *mut ::std::os::raw::c_int,
+        hTexRef: CUtexref,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetBorderColor(pBorderColor: *mut f32, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefGetFlags(pFlags: *mut ::std::os::raw::c_uint, hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefCreate(pTexRef: *mut CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexRefDestroy(hTexRef: CUtexref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSurfRefSetArray(
+        hSurfRef: CUsurfref,
+        hArray: CUarray,
+        Flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSurfRefGetArray(phArray: *mut CUarray, hSurfRef: CUsurfref) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexObjectCreate(
+        pTexObject: *mut CUtexObject,
+        pResDesc: *const CUDA_RESOURCE_DESC,
+        pTexDesc: *const CUDA_TEXTURE_DESC,
+        pResViewDesc: *const CUDA_RESOURCE_VIEW_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexObjectDestroy(texObject: CUtexObject) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexObjectGetResourceDesc(
+        pResDesc: *mut CUDA_RESOURCE_DESC,
+        texObject: CUtexObject,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexObjectGetTextureDesc(
+        pTexDesc: *mut CUDA_TEXTURE_DESC,
+        texObject: CUtexObject,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuTexObjectGetResourceViewDesc(
+        pResViewDesc: *mut CUDA_RESOURCE_VIEW_DESC,
+        texObject: CUtexObject,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSurfObjectCreate(
+        pSurfObject: *mut CUsurfObject,
+        pResDesc: *const CUDA_RESOURCE_DESC,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSurfObjectDestroy(surfObject: CUsurfObject) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuSurfObjectGetResourceDesc(
+        pResDesc: *mut CUDA_RESOURCE_DESC,
+        surfObject: CUsurfObject,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceCanAccessPeer(
+        canAccessPeer: *mut ::std::os::raw::c_int,
+        dev: CUdevice,
+        peerDev: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxEnablePeerAccess(peerContext: CUcontext, Flags: ::std::os::raw::c_uint)
+        -> CUresult;
+}
+extern_redirect! {
+    pub fn cuCtxDisablePeerAccess(peerContext: CUcontext) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuDeviceGetP2PAttribute(
+        value: *mut ::std::os::raw::c_int,
+        attrib: CUdevice_P2PAttribute,
+        srcDevice: CUdevice,
+        dstDevice: CUdevice,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsUnregisterResource(resource: CUgraphicsResource) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsSubResourceGetMappedArray(
+        pArray: *mut CUarray,
+        resource: CUgraphicsResource,
+        arrayIndex: ::std::os::raw::c_uint,
+        mipLevel: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsResourceGetMappedMipmappedArray(
+        pMipmappedArray: *mut CUmipmappedArray,
+        resource: CUgraphicsResource,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsResourceGetMappedPointer_v2(
+        pDevPtr: *mut CUdeviceptr,
+        pSize: *mut usize,
+        resource: CUgraphicsResource,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsResourceSetMapFlags_v2(
+        resource: CUgraphicsResource,
+        flags: ::std::os::raw::c_uint,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsMapResources(
+        count: ::std::os::raw::c_uint,
+        resources: *mut CUgraphicsResource,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect! {
+    pub fn cuGraphicsUnmapResources(
+        count: ::std::os::raw::c_uint,
+        resources: *mut CUgraphicsResource,
+        hStream: CUstream,
+    ) -> CUresult;
+}
+extern_redirect_with! {
+    pub fn cuGetExportTable(
+        ppExportTable: *mut *const ::std::os::raw::c_void,
+        pExportTableId: *const CUuuid,
+    ) -> CUresult;
+    super::cuGetExportTable;
+}
+extern_redirect! {
+    pub fn cuFuncGetModule(hmod: *mut CUmodule, hfunc: CUfunction) -> CUresult;
+}
diff --git a/zluda_dump/src/lib.rs b/zluda_dump/src/lib.rs
new file mode 100644
index 0000000..409606e
--- /dev/null
+++ b/zluda_dump/src/lib.rs
@@ -0,0 +1,670 @@
+use std::{
+    collections::HashMap,
+    env,
+    error::Error,
+    ffi::{c_void, CStr},
+    fs,
+    io::prelude::*,
+    mem,
+    os::raw::{c_int, c_uint, c_ulong, c_ushort},
+    path::PathBuf,
+    rc::Rc,
+    slice,
+};
+use std::{fs::File, ptr};
+
+use cuda::{CUdeviceptr, CUfunction, CUjit_option, CUmodule, CUresult, CUstream, CUuuid};
+use ptx::ast;
+use regex::Regex;
+
+#[cfg_attr(windows, path = "os_win.rs")]
+#[cfg_attr(not(windows), path = "os_unix.rs")]
+mod os;
+
+macro_rules! extern_redirect {
+    (pub fn $fn_name:ident ( $($arg_id:ident: $arg_type:ty),* $(,)? ) -> $ret_type:ty ;) => {
+        #[no_mangle]
+        pub fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type {
+            unsafe { $crate::init_libcuda_handle() };
+            let name = std::ffi::CString::new(stringify!($fn_name)).unwrap();
+            let fn_ptr = unsafe { crate::os::get_proc_address($crate::LIBCUDA_HANDLE, &name) };
+            if fn_ptr == std::ptr::null_mut() {
+                return CUresult::CUDA_ERROR_UNKNOWN;
+            }
+            let typed_fn = unsafe { std::mem::transmute::<_, fn( $( $arg_id : $arg_type),* ) -> $ret_type>(fn_ptr) };
+            typed_fn($( $arg_id ),*)
+        }
+    };
+}
+
+macro_rules! extern_redirect_with {
+    (
+        pub fn $fn_name:ident ( $($arg_id:ident: $arg_type:ty),* $(,)? ) -> $ret_type:ty ;
+        $receiver:path ;
+    ) => {
+        #[no_mangle]
+        pub fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type {
+            unsafe { $crate::init_libcuda_handle() };
+            let continuation = |$( $arg_id : $arg_type),* | {
+                let name = std::ffi::CString::new(stringify!($fn_name)).unwrap();
+                let fn_ptr = unsafe { crate::os::get_proc_address($crate::LIBCUDA_HANDLE, &name) };
+                if fn_ptr == std::ptr::null_mut() {
+                    return CUresult::CUDA_ERROR_UNKNOWN;
+                }
+                let typed_fn = unsafe { std::mem::transmute::<_, fn( $( $arg_id : $arg_type),* ) -> $ret_type>(fn_ptr) };
+                typed_fn($( $arg_id ),*)
+            };
+            unsafe { $receiver($( $arg_id ),* , continuation) }
+        }
+    };
+}
+
+#[allow(warnings)]
+mod cuda;
+
+pub static mut LIBCUDA_HANDLE: *mut c_void = ptr::null_mut();
+pub static mut MODULES: Option<HashMap<CUmodule, ModuleDump>> = None;
+pub static mut KERNELS: Option<HashMap<CUfunction, KernelDump>> = None;
+pub static mut BUFFERS: Vec<(usize, usize)> = Vec::new();
+pub static mut LAUNCH_COUNTER: usize = 0;
+pub static mut KERNEL_PATTERN: Option<Regex> = None;
+
+pub struct ModuleDump {
+    content: Rc<String>,
+    kernels_args: HashMap<String, Vec<usize>>,
+}
+
+pub struct KernelDump {
+    module_content: Rc<String>,
+    name: String,
+    arguments: Vec<usize>,
+}
+
+// We are doing dlopen here instead of just using LD_PRELOAD,
+// it's because CUDA Runtime API does dlopen to open libcuda.so, which ignores LD_PRELOAD
+pub unsafe fn init_libcuda_handle() {
+    if LIBCUDA_HANDLE == ptr::null_mut() {
+        let libcuda_handle = os::load_cuda_library();
+        assert_ne!(libcuda_handle, ptr::null_mut());
+        LIBCUDA_HANDLE = libcuda_handle;
+        match env::var("ZLUDA_DUMP_KERNEL") {
+            Ok(kernel_filter) => match Regex::new(&kernel_filter) {
+                Ok(r) => KERNEL_PATTERN = Some(r),
+                Err(err) => {
+                    eprintln!(
+                        "[ZLUDA_DUMP] Env variable ZLUDA_DUMP_KERNEL is not a regex: {}",
+                        err
+                    );
+                }
+            },
+            Err(_) => (),
+        }
+        eprintln!("[ZLUDA_DUMP] Initialized");
+    }
+}
+
+#[allow(non_snake_case)]
+pub unsafe fn cuModuleLoadData(
+    module: *mut CUmodule,
+    raw_image: *const ::std::os::raw::c_void,
+    cont: impl FnOnce(*mut CUmodule, *const c_void) -> CUresult,
+) -> CUresult {
+    let result = cont(module, raw_image);
+    if result == CUresult::CUDA_SUCCESS {
+        record_module_image_raw(*module, raw_image);
+    }
+    result
+}
+
+unsafe fn record_module_image_raw(module: CUmodule, raw_image: *const ::std::os::raw::c_void) {
+    let image = to_str(raw_image);
+    match image {
+        None => eprintln!("[ZLUDA_DUMP] Malformed module image: {:?}", raw_image),
+        Some(image) => record_module_image(module, image),
+    };
+}
+
+unsafe fn record_module_image(module: CUmodule, image: &str) {
+    if !image.contains(&".address_size") {
+        eprintln!("[ZLUDA_DUMP] Malformed module image: {:?}", module)
+    } else {
+        let mut errors = Vec::new();
+        let ast = ptx::ModuleParser::new().parse(&mut errors, image);
+        match (&*errors, ast) {
+            (&[], Ok(ast)) => {
+                let kernels_args = ast
+                    .directives
+                    .iter()
+                    .filter_map(directive_to_kernel)
+                    .collect::<HashMap<_, _>>();
+                let modules = MODULES.get_or_insert_with(|| HashMap::new());
+                modules.insert(
+                    module,
+                    ModuleDump {
+                        content: Rc::new(image.to_string()),
+                        kernels_args,
+                    },
+                );
+            }
+            (errs, ast) => {
+                let err_string = errs
+                    .iter()
+                    .map(|e| format!("{:?}", e))
+                    .chain(ast.err().iter().map(|e| format!("{:?}", e)))
+                    .collect::<Vec<_>>()
+                    .join("\n");
+                eprintln!(
+                    "[ZLUDA_DUMP] Errors when parsing module:\n---ERRORS---\n{}\n---MODULE---\n{}",
+                    err_string, image
+                );
+            }
+        }
+    }
+}
+
+unsafe fn to_str<T>(image: *const T) -> Option<&'static str> {
+    let ptr = image as *const u8;
+    let mut offset = 0;
+    loop {
+        let c = *ptr.add(offset);
+        if !c.is_ascii() {
+            return None;
+        }
+        if c == 0 {
+            return Some(std::str::from_utf8_unchecked(slice::from_raw_parts(
+                ptr, offset,
+            )));
+        }
+        offset += 1;
+    }
+}
+
+fn directive_to_kernel(dir: &ast::Directive<ast::ParsedArgParams>) -> Option<(String, Vec<usize>)> {
+    match dir {
+        ast::Directive::Method(ast::Function {
+            func_directive: ast::MethodDecl::Kernel { name, in_args },
+            ..
+        }) => {
+            let arg_sizes = in_args
+                .iter()
+                .map(|arg| ast::Type::from(arg.v_type.clone()).size_of())
+                .collect();
+            Some((name.to_string(), arg_sizes))
+        }
+        _ => None,
+    }
+}
+
+#[allow(non_snake_case)]
+pub unsafe fn cuModuleLoadDataEx(
+    module: *mut CUmodule,
+    image: *const c_void,
+    numOptions: c_uint,
+    options: *mut CUjit_option,
+    optionValues: *mut *mut c_void,
+    cont: impl FnOnce(
+        *mut CUmodule,
+        *const c_void,
+        c_uint,
+        *mut CUjit_option,
+        *mut *mut c_void,
+    ) -> CUresult,
+) -> CUresult {
+    let result = cont(module, image, numOptions, options, optionValues);
+    if result == CUresult::CUDA_SUCCESS {
+        record_module_image_raw(*module, image);
+    }
+    result
+}
+
+#[allow(non_snake_case)]
+unsafe fn cuModuleGetFunction(
+    hfunc: *mut CUfunction,
+    hmod: CUmodule,
+    name: *const ::std::os::raw::c_char,
+    cont: impl FnOnce(*mut CUfunction, CUmodule, *const ::std::os::raw::c_char) -> CUresult,
+) -> CUresult {
+    let result = cont(hfunc, hmod, name);
+    if result != CUresult::CUDA_SUCCESS {
+        return result;
+    }
+    if let Some(modules) = &MODULES {
+        if let Some(module_dump) = modules.get(&hmod) {
+            if let Some(kernel) = to_str(name) {
+                if let Some(args) = module_dump.kernels_args.get(kernel) {
+                    let kernel_args = KERNELS.get_or_insert_with(|| HashMap::new());
+                    kernel_args.insert(
+                        *hfunc,
+                        KernelDump {
+                            module_content: module_dump.content.clone(),
+                            name: kernel.to_string(),
+                            arguments: args.clone(),
+                        },
+                    );
+                } else {
+                    eprintln!("[ZLUDA_DUMP] Unknown kernel: {}", kernel);
+                }
+            } else {
+                eprintln!("[ZLUDA_DUMP] Unknown kernel name at: {:?}", hfunc);
+            }
+        } else {
+            eprintln!("[ZLUDA_DUMP] Unknown module: {:?}", hmod);
+        }
+    } else {
+        eprintln!("[ZLUDA_DUMP] Unknown module: {:?}", hmod);
+    }
+    CUresult::CUDA_SUCCESS
+}
+
+#[allow(non_snake_case)]
+pub unsafe fn cuMemAlloc_v2(
+    dptr: *mut CUdeviceptr,
+    bytesize: usize,
+    cont: impl FnOnce(*mut CUdeviceptr, usize) -> CUresult,
+) -> CUresult {
+    let result = cont(dptr, bytesize);
+    assert_eq!(result, CUresult::CUDA_SUCCESS);
+    let start = (*dptr).0 as usize;
+    BUFFERS.push((start, bytesize));
+    CUresult::CUDA_SUCCESS
+}
+
+#[allow(non_snake_case)]
+pub unsafe 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,
+    cont: impl FnOnce(
+        CUfunction,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        ::std::os::raw::c_uint,
+        CUstream,
+        *mut *mut ::std::os::raw::c_void,
+        *mut *mut ::std::os::raw::c_void,
+    ) -> CUresult,
+) -> CUresult {
+    let mut error;
+    let dump_env = match create_dump_dir(f, LAUNCH_COUNTER) {
+        Ok(dump_env) => dump_env,
+        Err(err) => {
+            eprintln!("[ZLUDA_DUMP] {:#?}", err);
+            None
+        }
+    };
+    if let Some(dump_env) = &dump_env {
+        dump_pre_data(
+            gridDimX,
+            gridDimY,
+            gridDimZ,
+            blockDimX,
+            blockDimY,
+            blockDimZ,
+            sharedMemBytes,
+            kernelParams,
+            dump_env,
+        )
+        .unwrap_or_else(|err| eprintln!("[ZLUDA_DUMP] {:#?}", err));
+    };
+    error = cont(
+        f,
+        gridDimX,
+        gridDimY,
+        gridDimZ,
+        blockDimX,
+        blockDimY,
+        blockDimZ,
+        sharedMemBytes,
+        hStream,
+        kernelParams,
+        extra,
+    );
+    assert_eq!(error, CUresult::CUDA_SUCCESS);
+    error = cuda::cuStreamSynchronize(hStream);
+    assert_eq!(error, CUresult::CUDA_SUCCESS);
+    if let Some((_, kernel_dump)) = &dump_env {
+        dump_arguments(
+            kernelParams,
+            "post",
+            &kernel_dump.name,
+            LAUNCH_COUNTER,
+            &kernel_dump.arguments,
+        )
+        .unwrap_or_else(|err| eprintln!("[ZLUDA_DUMP] {:#?}", err));
+    }
+    LAUNCH_COUNTER += 1;
+    CUresult::CUDA_SUCCESS
+}
+
+#[allow(non_snake_case)]
+fn dump_launch_arguments(
+    gridDimX: u32,
+    gridDimY: u32,
+    gridDimZ: u32,
+    blockDimX: u32,
+    blockDimY: u32,
+    blockDimZ: u32,
+    sharedMemBytes: u32,
+    dump_dir: &PathBuf,
+) -> Result<(), Box<dyn Error>> {
+    let mut module_file_path = dump_dir.clone();
+    module_file_path.push("launch.txt");
+    let mut module_file = File::create(module_file_path)?;
+    write!(&mut module_file, "{}\n", gridDimX)?;
+    write!(&mut module_file, "{}\n", gridDimY)?;
+    write!(&mut module_file, "{}\n", gridDimZ)?;
+    write!(&mut module_file, "{}\n", blockDimX)?;
+    write!(&mut module_file, "{}\n", blockDimY)?;
+    write!(&mut module_file, "{}\n", blockDimZ)?;
+    write!(&mut module_file, "{}\n", sharedMemBytes)?;
+    Ok(())
+}
+
+unsafe fn should_dump_kernel(name: &str) -> bool {
+    match &KERNEL_PATTERN {
+        Some(pattern) => pattern.is_match(name),
+        None => true,
+    }
+}
+
+unsafe fn create_dump_dir(
+    f: CUfunction,
+    counter: usize,
+) -> Result<Option<(PathBuf, &'static KernelDump)>, Box<dyn Error>> {
+    match KERNELS.as_ref().and_then(|kernels| kernels.get(&f)) {
+        Some(kernel_dump) => {
+            if !should_dump_kernel(&kernel_dump.name) {
+                return Ok(None);
+            }
+            let mut dump_dir = get_dump_dir()?;
+            dump_dir.push(format!("{:04}_{}", counter, kernel_dump.name));
+            fs::create_dir_all(&dump_dir)?;
+            Ok(Some((dump_dir, kernel_dump)))
+        }
+        None => Err("Unknown kernel: {:?}")?,
+    }
+}
+
+#[allow(non_snake_case)]
+unsafe fn dump_pre_data(
+    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,
+    kernelParams: *mut *mut ::std::os::raw::c_void,
+    (dump_dir, kernel_dump): &(PathBuf, &'static KernelDump),
+) -> Result<(), Box<dyn Error>> {
+    dump_launch_arguments(
+        gridDimX,
+        gridDimY,
+        gridDimZ,
+        blockDimX,
+        blockDimY,
+        blockDimZ,
+        sharedMemBytes,
+        dump_dir,
+    )?;
+    let mut module_file_path = dump_dir.clone();
+    module_file_path.push("module.ptx");
+    let mut module_file = File::create(module_file_path)?;
+    module_file.write_all(kernel_dump.module_content.as_bytes())?;
+    dump_arguments(
+        kernelParams,
+        "pre",
+        &kernel_dump.name,
+        LAUNCH_COUNTER,
+        &kernel_dump.arguments,
+    )?;
+    Ok(())
+}
+
+unsafe fn dump_arguments(
+    kernel_params: *mut *mut ::std::os::raw::c_void,
+    prefix: &str,
+    kernel_name: &str,
+    counter: usize,
+    args: &[usize],
+) -> Result<(), Box<dyn Error>> {
+    let mut dump_dir = get_dump_dir()?;
+    dump_dir.push(format!("{:04}_{}", counter, kernel_name));
+    dump_dir.push(prefix);
+    if dump_dir.exists() {
+        fs::remove_dir_all(&dump_dir)?;
+    }
+    fs::create_dir_all(&dump_dir)?;
+    for (i, arg_len) in args.iter().enumerate() {
+        let dev_ptr = *(*kernel_params.add(i) as *mut usize);
+        match BUFFERS.iter().find(|(start, _)| *start == dev_ptr as usize) {
+            Some((start, len)) => {
+                let mut output = vec![0u8; *len];
+                let error =
+                    cuda::cuMemcpyDtoH_v2(output.as_mut_ptr() as *mut _, CUdeviceptr(*start), *len);
+                assert_eq!(error, CUresult::CUDA_SUCCESS);
+                let mut path = dump_dir.clone();
+                path.push(format!("arg_{:03}.buffer", i));
+                let mut file = File::create(path)?;
+                file.write_all(&mut output)?;
+            }
+            None => {
+                let mut path = dump_dir.clone();
+                path.push(format!("arg_{:03}", i));
+                let mut file = File::create(path)?;
+                file.write_all(slice::from_raw_parts(
+                    *kernel_params.add(i) as *mut u8,
+                    *arg_len,
+                ))?;
+            }
+        }
+    }
+    Ok(())
+}
+
+fn get_dump_dir() -> Result<PathBuf, Box<dyn Error>> {
+    let dir = env::var("ZLUDA_DUMP_DIR")?;
+    let mut main_dir = PathBuf::from(dir);
+    let current_exe = env::current_exe()?;
+    main_dir.push(current_exe.file_name().unwrap());
+    fs::create_dir_all(&main_dir)?;
+    Ok(main_dir)
+}
+
+// TODO make this more common with ZLUDA implementation
+const CUDART_INTERFACE_GUID: CUuuid = CUuuid {
+    bytes: [
+        0x6b, 0xd5, 0xfb, 0x6c, 0x5b, 0xf4, 0xe7, 0x4a, 0x89, 0x87, 0xd9, 0x39, 0x12, 0xfd, 0x9d,
+        0xf9,
+    ],
+};
+
+const GET_MODULE_OFFSET: usize = 6;
+static mut CUDART_INTERFACE_VTABLE: Vec<*const c_void> = Vec::new();
+static mut ORIGINAL_GET_MODULE_FROM_CUBIN: Option<
+    unsafe extern "C" fn(
+        result: *mut CUmodule,
+        fatbinc_wrapper: *const FatbincWrapper,
+        ptr1: *mut c_void,
+        ptr2: *mut c_void,
+    ) -> CUresult,
+> = None;
+
+#[allow(non_snake_case)]
+pub unsafe fn cuGetExportTable(
+    ppExportTable: *mut *const ::std::os::raw::c_void,
+    pExportTableId: *const CUuuid,
+    cont: impl FnOnce(*mut *const ::std::os::raw::c_void, *const CUuuid) -> CUresult,
+) -> CUresult {
+    if *pExportTableId == CUDART_INTERFACE_GUID {
+        if CUDART_INTERFACE_VTABLE.len() == 0 {
+            let mut base_table = ptr::null();
+            let base_result = cont(&mut base_table, pExportTableId);
+            if base_result != CUresult::CUDA_SUCCESS {
+                return base_result;
+            }
+            let len = *(base_table as *const usize);
+            CUDART_INTERFACE_VTABLE = vec![ptr::null(); len];
+            ptr::copy_nonoverlapping(
+                base_table as *const _,
+                CUDART_INTERFACE_VTABLE.as_mut_ptr(),
+                len,
+            );
+            if GET_MODULE_OFFSET >= len {
+                return CUresult::CUDA_ERROR_UNKNOWN;
+            }
+            ORIGINAL_GET_MODULE_FROM_CUBIN =
+                mem::transmute(CUDART_INTERFACE_VTABLE[GET_MODULE_OFFSET]);
+            CUDART_INTERFACE_VTABLE[GET_MODULE_OFFSET] = get_module_from_cubin as *const _;
+        }
+        *ppExportTable = CUDART_INTERFACE_VTABLE.as_ptr() as *const _;
+        return CUresult::CUDA_SUCCESS;
+    } else {
+        cont(ppExportTable, pExportTableId)
+    }
+}
+
+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)]
+#[derive(Debug)]
+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(
+    module: *mut CUmodule,
+    fatbinc_wrapper: *const FatbincWrapper,
+    ptr1: *mut c_void,
+    ptr2: *mut c_void,
+) -> CUresult {
+    if module == ptr::null_mut()
+        || (*fatbinc_wrapper).magic != FATBINC_MAGIC
+        || (*fatbinc_wrapper).version != FATBINC_VERSION
+    {
+        return CUresult::CUDA_ERROR_INVALID_VALUE;
+    }
+    let fatbin_header = (*fatbinc_wrapper).data;
+    if (*fatbin_header).magic != FATBIN_MAGIC || (*fatbin_header).version != FATBIN_VERSION {
+        return 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);
+    let mut maybe_kernel_text = None;
+    for file in ptx_files {
+        match decompress_kernel_module(file) {
+            None => continue,
+            Some(vec) => {
+                maybe_kernel_text = Some(vec);
+                break;
+            }
+        };
+    }
+    let result = ORIGINAL_GET_MODULE_FROM_CUBIN.unwrap()(module, fatbinc_wrapper, ptr1, ptr2);
+    if result != CUresult::CUDA_SUCCESS {
+        return result;
+    }
+    if let Some(text) = maybe_kernel_text {
+        match CStr::from_bytes_with_nul(&text) {
+            Ok(cstr) => match cstr.to_str() {
+                Ok(utf8_str) => record_module_image(*module, utf8_str),
+                Err(_) => {}
+            },
+            Err(_) => {}
+        }
+    }
+    result
+}
+
+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
+}
+
+const MAX_PTX_MODULE_DECOMPRESSION_BOUND: usize = 16 * 1024 * 1024;
+
+unsafe fn decompress_kernel_module(file: *const FatbinFileHeader) -> Option<Vec<u8>> {
+    let decompressed_size = usize::max(1024, (*file).uncompressed_payload as usize);
+    let mut decompressed_vec = vec![0u8; decompressed_size];
+    loop {
+        match lz4_sys::LZ4_decompress_safe(
+            (file as *const u8).add((*file).header_size as usize) as *const _,
+            decompressed_vec.as_mut_ptr() as *mut _,
+            (*file).payload_size as c_int,
+            decompressed_vec.len() as c_int,
+        ) {
+            error if error < 0 => {
+                let new_size = decompressed_vec.len() * 2;
+                if new_size > MAX_PTX_MODULE_DECOMPRESSION_BOUND {
+                    return None;
+                }
+                decompressed_vec.resize(decompressed_vec.len() * 2, 0);
+            }
+            real_decompressed_size => {
+                decompressed_vec.truncate(real_decompressed_size as usize);
+                return Some(decompressed_vec);
+            }
+        }
+    }
+}
diff --git a/zluda_dump/src/os_unix.rs b/zluda_dump/src/os_unix.rs
new file mode 100644
index 0000000..b3d9343
--- /dev/null
+++ b/zluda_dump/src/os_unix.rs
@@ -0,0 +1,14 @@
+use std::ffi::{c_void, CStr};

+

+const NVCUDA_DEFAULT_PATH: &'static [u8] = b"/usr/lib/x86_64-linux-gnu/libcuda.so.1\0";

+

+pub unsafe fn load_cuda_library() -> *mut c_void {

+    libc::dlopen(

+        NVCUDA_DEFAULT_PATH.as_ptr() as *const _,

+        libc::RTLD_LOCAL | libc::RTLD_NOW,

+    )

+}

+

+pub unsafe fn get_proc_address(handle: *mut c_void, func: &CStr) -> *mut c_void {

+    libc::dlsym(handle, func.as_ptr() as *const _)

+}

diff --git a/zluda_dump/src/os_win.rs b/zluda_dump/src/os_win.rs
new file mode 100644
index 0000000..7f985c5
--- /dev/null
+++ b/zluda_dump/src/os_win.rs
@@ -0,0 +1,68 @@
+use std::{

+    ffi::{c_void, CStr},

+    mem,

+    os::raw::c_ushort,

+    ptr,

+};

+

+use wchar::wch_c;

+use winapi::{

+    shared::minwindef::HMODULE,

+    um::libloaderapi::{GetProcAddress, LoadLibraryW},

+};

+

+const NVCUDA_DEFAULT_PATH: &[u16] = wch_c!(r"C:\Windows\System32\nvcuda.dll");

+const LOAD_LIBRARY_NO_REDIRECT: &'static [u8] = b"ZludaLoadLibraryW_NoRedirect\0";

+

+include!("../../zluda_redirect/src/payload_guid.rs");

+

+pub unsafe fn load_cuda_library() -> *mut c_void {

+    let load_lib = if is_detoured() {

+        match get_non_detoured_load_library() {

+            Some(load_lib) => load_lib,

+            None => return ptr::null_mut(),

+        }

+    } else {

+        LoadLibraryW

+    };

+    load_lib(NVCUDA_DEFAULT_PATH.as_ptr()) as *mut _

+}

+

+unsafe fn is_detoured() -> bool {

+    let mut module = ptr::null_mut();

+    loop {

+        module = detours_sys::DetourEnumerateModules(module);

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

+            break;

+        }

+        let mut size = 0;

+        let payload = detours_sys::DetourFindPayload(module, &PAYLOAD_GUID, &mut size);

+        if payload != ptr::null_mut() {

+            return true;

+        }

+    }

+    false

+}

+

+unsafe fn get_non_detoured_load_library(

+) -> Option<unsafe extern "system" fn(*const c_ushort) -> HMODULE> {

+    let mut module = ptr::null_mut();

+    loop {

+        module = detours_sys::DetourEnumerateModules(module);

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

+            break;

+        }

+        let result = GetProcAddress(

+            module as *mut _,

+            LOAD_LIBRARY_NO_REDIRECT.as_ptr() as *mut _,

+        );

+        if result != ptr::null_mut() {

+            return Some(mem::transmute(result));

+        }

+    }

+    None

+}

+

+pub unsafe fn get_proc_address(handle: *mut c_void, func: &CStr) -> *mut c_void {

+    GetProcAddress(handle as *mut _, func.as_ptr()) as *mut _

+}

diff --git a/zluda_inject/Cargo.toml b/zluda_inject/Cargo.toml
index 193c36e..7576e08 100644
--- a/zluda_inject/Cargo.toml
+++ b/zluda_inject/Cargo.toml
@@ -9,5 +9,5 @@ name = "zluda_with"
 path = "src/main.rs"
 
 [target.'cfg(windows)'.dependencies]
-winapi = { version = "0.3", features = ["jobapi2", "processthreadsapi", "std", "synchapi"] }
+winapi = { version = "0.3", features = ["jobapi2", "processthreadsapi", "std", "synchapi", "winbase"] }
 detours-sys = { path = "../detours-sys" }
diff --git a/zluda_inject/src/bin.rs b/zluda_inject/src/bin.rs
index af44d74..ce83fe9 100644
--- a/zluda_inject/src/bin.rs
+++ b/zluda_inject/src/bin.rs
@@ -1,8 +1,7 @@
-use std::env;

-use std::env::Args;

 use std::mem;

 use std::path::Path;

 use std::ptr;

+use std::{env, ops::Deref};

 use std::{error::Error, process};

 

 use mem::size_of_val;

@@ -25,15 +24,15 @@ static ZLUDA_DLL: &'static str = "nvcuda.dll";
 include!("../../zluda_redirect/src/payload_guid.rs");

 

 pub fn main_impl() -> Result<(), Box<dyn Error>> {

-    let args = env::args();

-    if args.len() == 0 {

-        print_help();

-        process::exit(1);

+    let args = env::args().collect::<Vec<_>>();

+    if args.len() <= 1 {

+        print_help_and_exit();

     }

-    let mut cmd_line = construct_command_line(args);

     let injector_path = env::current_exe()?;

     let injector_dir = injector_path.parent().unwrap();

     let redirect_path = create_redirect_path(injector_dir);

+    let (mut inject_path, cmd) = create_inject_path(&args[1..], injector_dir);

+    let mut cmd_line = construct_command_line(cmd);

     let mut startup_info = unsafe { mem::zeroed::<detours_sys::_STARTUPINFOW>() };

     let mut proc_info = unsafe { mem::zeroed::<detours_sys::_PROCESS_INFORMATION>() };

     os_call!(

@@ -54,13 +53,12 @@ pub fn main_impl() -> Result<(), Box<dyn Error>> {
         |x| x != 0

     );

     kill_child_on_process_exit(proc_info.hProcess)?;

-    let mut zluda_path = create_zluda_path(injector_dir);

     os_call!(

         detours_sys::DetourCopyPayloadToProcess(

             proc_info.hProcess,

             &PAYLOAD_GUID,

-            zluda_path.as_mut_ptr() as *mut _,

-            (zluda_path.len() * mem::size_of::<u16>()) as u32

+            inject_path.as_mut_ptr() as *mut _,

+            (inject_path.len() * mem::size_of::<u16>()) as u32

         ),

         |x| x != 0

     );

@@ -93,22 +91,29 @@ fn kill_child_on_process_exit(child: HANDLE) -> Result<(), Box<dyn Error>> {
     Ok(())

 }

 

-fn print_help() {

+fn print_help_and_exit() -> ! {

+    let current_exe = env::current_exe().unwrap();

+    let exe_name = current_exe.file_name().unwrap().to_string_lossy();

     println!(

         "USAGE:

-    zluda <EXE> [ARGS]...

+    {0} -- <EXE> [ARGS]...

+    {0} <DLL> -- <EXE> [ARGS]...

 ARGS:

-    <EXE>        Path to the executable to be injected with ZLUDA

+    <DLL>        DLL to ne injected instead of system nvcuda.dll, if not provided

+                 will use nvcuda.dll from the directory where {0} is located

+    <EXE>        Path to the executable to be injected with <DLL>

     <ARGS>...    Arguments that will be passed to <EXE>

-"

+",

+        exe_name

     );

+    process::exit(1)

 }

 

 // Adapted from https://docs.microsoft.com/en-us/archive/blogs/twistylittlepassagesallalike/everyone-quotes-command-line-arguments-the-wrong-way

-fn construct_command_line(args: Args) -> Vec<u16> {

+fn construct_command_line(args: &[String]) -> Vec<u16> {

     let mut cmd_line = Vec::new();

     let args_len = args.len();

-    for (idx, arg) in args.enumerate().skip(1) {

+    for (idx, arg) in args.iter().enumerate() {

         if !arg.contains(&[' ', '\t', '\n', '\u{2B7F}', '\"'][..]) {

             cmd_line.extend(arg.encode_utf16());

         } else {

@@ -168,14 +173,22 @@ fn create_redirect_path(injector_dir: &Path) -> Vec<u8> {
     result

 }

 

-fn create_zluda_path(injector_dir: &Path) -> Vec<u16> {

-    let mut injector_dir = injector_dir.to_path_buf();

-    injector_dir.push(ZLUDA_DLL);

-    let mut result = injector_dir

-        .to_string_lossy()

-        .as_ref()

-        .encode_utf16()

-        .collect::<Vec<_>>();

-    result.push(0);

-    result

+fn create_inject_path<'a>(args: &'a [String], injector_dir: &Path) -> (Vec<u16>, &'a [String]) {

+    if args.get(0).map(Deref::deref) == Some("--") {

+        let mut injector_dir = injector_dir.to_path_buf();

+        injector_dir.push(ZLUDA_DLL);

+        let mut result = injector_dir

+            .to_string_lossy()

+            .as_ref()

+            .encode_utf16()

+            .collect::<Vec<_>>();

+        result.push(0);

+        (result, &args[1..])

+    } else if args.get(1).map(Deref::deref) == Some("--") {

+        let mut dll_path = args[0].encode_utf16().collect::<Vec<_>>();

+        dll_path.push(0);

+        (dll_path, &args[2..])

+    } else {

+        print_help_and_exit()

+    }

 }

diff --git a/zluda_redirect/src/lib.rs b/zluda_redirect/src/lib.rs
index d0497a3..04b2413 100644
--- a/zluda_redirect/src/lib.rs
+++ b/zluda_redirect/src/lib.rs
@@ -46,6 +46,12 @@ static mut LOAD_LIBRARY_EX_W: unsafe extern "system" fn(
     dwFlags: DWORD,
 ) -> HMODULE = LoadLibraryExW;
 
+#[no_mangle]
+#[allow(non_snake_case)]
+unsafe extern "system" fn ZludaLoadLibraryW_NoRedirect(lpLibFileName: LPCWSTR) -> HMODULE {
+    (LOAD_LIBRARY_W)(lpLibFileName)
+}
+
 #[allow(non_snake_case)]
 unsafe extern "system" fn ZludaLoadLibraryA(lpLibFileName: LPCSTR) -> HMODULE {
     let nvcuda_file_name = if is_nvcuda_dll_utf8(lpLibFileName as *const _) {