Rebindgen to emit send,sync,hash

6 files changed, 301 insertions, 230 deletions

diff --git a/cuda_base/src/lib.rs b/cuda_base/src/lib.rs
index 765af71..0cc1f53 100644
--- a/cuda_base/src/lib.rs
+++ b/cuda_base/src/lib.rs
@@ -161,17 +161,8 @@ pub fn cuda_normalize_fn(tokens: TokenStream) -> TokenStream {
         .0
         .ident
         .to_string();
-    let known_modules = [
-        ("ctx", "context"),
-        ("device", "device"),
-        ("function", "function"),
-        ("link", "link"),
-        ("memory", "memory"),
-        ("module", "module"),
-        ("pointer", "pointer"),
-    ];
-    let segments: Vec<String> = split(&fn_[2..]);
-    let fn_path = join(segments, &known_modules);
+    let segments: Vec<String> = split(&fn_[2..]); // skip "cu"
+    let fn_path = join(segments);
     quote! {
         #path #fn_path
     }
@@ -190,17 +181,42 @@ fn split(fn_: &str) -> Vec<String> {
     result
 }
 
-fn join(fn_: Vec<String>, known_modules: &[(&str, &str)]) -> Punctuated<Ident, Token![::]> {
-    let (prefix, suffix) = fn_.split_at(1);
-    if let Some((_, mod_name)) = known_modules
-        .iter()
-        .find(|(mod_prefix, _)| mod_prefix == &prefix[0])
-    {
-        [*mod_name, &suffix.join("_")]
-            .into_iter()
-            .map(|seg| Ident::new(seg, Span::call_site()))
-            .collect()
-    } else {
-        iter::once(Ident::new(&fn_.join("_"), Span::call_site())).collect()
+fn join(fn_: Vec<String>) -> Punctuated<Ident, Token![::]> {
+    fn full_form(segment: &str) -> Option<&[&str]> {
+        Some(match segment {
+            "ctx" => &["context"],
+            "memcpy" => &["memory", "copy"],
+            _ => return None,
+        })
     }
+    const MODULES: &[&str] = &[
+        "context",
+        "device",
+        "function",
+        "link",
+        "memory",
+        "module",
+        "pointer"
+    ];
+    let mut normalized: Vec<&str> = Vec::new();
+    for segment in fn_.iter() {
+        match full_form(segment) {
+            Some(segments) => normalized.extend(segments.into_iter()),
+            None => normalized.push(&*segment),
+        }
+    }
+    if !MODULES.contains(&normalized[0]) {
+        let mut globalized = vec!["global"];
+        globalized.extend(normalized);
+        normalized = globalized;
+    }
+    let (module, path) = normalized.split_first().unwrap();
+    let path = path.join("_");
+    let mut result = Punctuated::new();
+    result.extend(
+        [module, &&*path]
+            .into_iter()
+            .map(|s| Ident::new(s, Span::call_site())),
+    );
+    result
 }
diff --git a/cuda_types/src/lib.rs b/cuda_types/src/lib.rs
index 945c0a7..2c2716a 100644
--- a/cuda_types/src/lib.rs
+++ b/cuda_types/src/lib.rs
@@ -48,37 +48,43 @@ pub const CUDA_EGL_INFINITE_TIMEOUT: u32 = 4294967295;
 pub type cuuint32_t = u32;
 pub type cuuint64_t = u64;
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUdeviceptr_v2(pub *mut ::core::ffi::c_void);
 pub type CUdeviceptr = CUdeviceptr_v2;
-#[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
-pub struct CUdevice_v1(pub ::core::ffi::c_int);
+pub type CUdevice_v1 = ::core::ffi::c_int;
 pub type CUdevice = CUdevice_v1;
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUctx_st {
     _unused: [u8; 0],
 }
-pub type CUcontext = *mut CUctx_st;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUcontext(pub *mut CUctx_st);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUmod_st {
     _unused: [u8; 0],
 }
-pub type CUmodule = *mut CUmod_st;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUmodule(pub *mut CUmod_st);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUfunc_st {
     _unused: [u8; 0],
 }
-pub type CUfunction = *mut CUfunc_st;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUfunction(pub *mut CUfunc_st);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUlib_st {
     _unused: [u8; 0],
 }
-pub type CUlibrary = *mut CUlib_st;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUlibrary(pub *mut CUlib_st);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUkern_st {
@@ -120,7 +126,9 @@ pub type CUevent = *mut CUevent_st;
 pub struct CUstream_st {
     _unused: [u8; 0],
 }
-pub type CUstream = *mut CUstream_st;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct CUstream(pub *mut CUstream_st);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct CUgraphicsResource_st {
@@ -187,7 +195,7 @@ pub struct CUasyncCallbackEntry_st {
 }
 pub type CUasyncCallbackHandle = *mut CUasyncCallbackEntry_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUuuid_st {
     pub bytes: [::core::ffi::c_uchar; 16usize],
 }
@@ -197,7 +205,7 @@ pub type CUuuid = CUuuid_st;
  between processes on different nodes they must be connected via the
  NVSwitch fabric.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemFabricHandle_st {
     pub data: [::core::ffi::c_uchar; 64usize],
 }
@@ -213,7 +221,7 @@ pub type CUmemFabricHandle_v1 = CUmemFabricHandle_st;
 pub type CUmemFabricHandle = CUmemFabricHandle_v1;
 /// CUDA IPC event handle
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUipcEventHandle_st {
     pub reserved: [::core::ffi::c_char; 64usize],
 }
@@ -223,7 +231,7 @@ pub type CUipcEventHandle_v1 = CUipcEventHandle_st;
 pub type CUipcEventHandle = CUipcEventHandle_v1;
 /// CUDA IPC mem handle
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUipcMemHandle_st {
     pub reserved: [::core::ffi::c_char; 64usize],
 }
@@ -600,13 +608,13 @@ pub union CUstreamBatchMemOpParams_union_CUstreamMemOpWriteValueParams_st__bindg
     pub value64: cuuint64_t,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUstreamBatchMemOpParams_union_CUstreamMemOpFlushRemoteWritesParams_st {
     pub operation: CUstreamBatchMemOpType,
     pub flags: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUstreamBatchMemOpParams_union_CUstreamMemOpMemoryBarrierParams_st {
     pub operation: CUstreamBatchMemOpType,
     pub flags: ::core::ffi::c_uint,
@@ -616,7 +624,7 @@ pub type CUstreamBatchMemOpParams_v1 = CUstreamBatchMemOpParams_union;
 /// Per-operation parameters for ::cuStreamBatchMemOp
 pub type CUstreamBatchMemOpParams = CUstreamBatchMemOpParams_v1;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st {
     pub ctx: CUcontext,
     pub count: ::core::ffi::c_uint,
@@ -627,7 +635,7 @@ pub type CUDA_BATCH_MEM_OP_NODE_PARAMS_v1 = CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st;
 pub type CUDA_BATCH_MEM_OP_NODE_PARAMS = CUDA_BATCH_MEM_OP_NODE_PARAMS_v1;
 /// Batch memory operation node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st {
     ///< Context to use for the operations.
     pub ctx: CUcontext,
@@ -698,7 +706,7 @@ pub union CUasyncNotificationInfo_st__bindgen_ty_1 {
     pub overBudget: CUasyncNotificationInfo_st__bindgen_ty_1__bindgen_ty_1,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUasyncNotificationInfo_st__bindgen_ty_1__bindgen_ty_1 {
     pub bytesOverBudget: ::core::ffi::c_ulonglong,
 }
@@ -1756,7 +1764,7 @@ pub struct CUdevice_attribute_enum(pub ::core::ffi::c_uint);
 pub use self::CUdevice_attribute_enum as CUdevice_attribute;
 /// Legacy device properties
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUdevprop_st {
     ///< Maximum number of threads per block
     pub maxThreadsPerBlock: ::core::ffi::c_int,
@@ -3015,7 +3023,7 @@ pub type CUaccessPolicyWindow_v1 = CUaccessPolicyWindow_st;
 pub type CUaccessPolicyWindow = CUaccessPolicyWindow_v1;
 /// GPU kernel node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_KERNEL_NODE_PARAMS_st {
     ///< Kernel to launch
     pub func: CUfunction,
@@ -3042,7 +3050,7 @@ pub struct CUDA_KERNEL_NODE_PARAMS_st {
 pub type CUDA_KERNEL_NODE_PARAMS_v1 = CUDA_KERNEL_NODE_PARAMS_st;
 /// GPU kernel node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_KERNEL_NODE_PARAMS_v2_st {
     ///< Kernel to launch
     pub func: CUfunction,
@@ -3075,7 +3083,7 @@ pub type CUDA_KERNEL_NODE_PARAMS_v2 = CUDA_KERNEL_NODE_PARAMS_v2_st;
 pub type CUDA_KERNEL_NODE_PARAMS = CUDA_KERNEL_NODE_PARAMS_v2;
 /// GPU kernel node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_KERNEL_NODE_PARAMS_v3_st {
     ///< Kernel to launch
     pub func: CUfunction,
@@ -3106,7 +3114,7 @@ pub struct CUDA_KERNEL_NODE_PARAMS_v3_st {
 pub type CUDA_KERNEL_NODE_PARAMS_v3 = CUDA_KERNEL_NODE_PARAMS_v3_st;
 /// Memset node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMSET_NODE_PARAMS_st {
     ///< Destination device pointer
     pub dst: CUdeviceptr,
@@ -3127,7 +3135,7 @@ pub type CUDA_MEMSET_NODE_PARAMS_v1 = CUDA_MEMSET_NODE_PARAMS_st;
 pub type CUDA_MEMSET_NODE_PARAMS = CUDA_MEMSET_NODE_PARAMS_v1;
 /// Memset node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMSET_NODE_PARAMS_v2_st {
     ///< Destination device pointer
     pub dst: CUdeviceptr,
@@ -3148,7 +3156,7 @@ pub struct CUDA_MEMSET_NODE_PARAMS_v2_st {
 pub type CUDA_MEMSET_NODE_PARAMS_v2 = CUDA_MEMSET_NODE_PARAMS_v2_st;
 /// Host node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone)]
+#[derive(Debug, Copy, Clone, Hash)]
 pub struct CUDA_HOST_NODE_PARAMS_st {
     ///< The function to call when the node executes
     pub fn_: CUhostFn,
@@ -3161,7 +3169,7 @@ pub type CUDA_HOST_NODE_PARAMS_v1 = CUDA_HOST_NODE_PARAMS_st;
 pub type CUDA_HOST_NODE_PARAMS = CUDA_HOST_NODE_PARAMS_v1;
 /// Host node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_HOST_NODE_PARAMS_v2_st {
     ///< The function to call when the node executes
     pub fn_: CUhostFn,
@@ -3190,7 +3198,7 @@ pub struct CUgraphConditionalNodeType_enum(pub ::core::ffi::c_uint);
 pub use self::CUgraphConditionalNodeType_enum as CUgraphConditionalNodeType;
 /// Conditional node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_CONDITIONAL_NODE_PARAMS {
     /**< Conditional node handle.
 Handles must be created in advance of creating the node
@@ -3334,7 +3342,7 @@ pub use self::CUgraphDependencyType_enum as CUgraphDependencyType;
  default to a zero-initialized value if not specified. A zero-initialized struct indicates a
  standard full serialization of two nodes with memory visibility.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUgraphEdgeData_st {
     /**< This indicates when the dependency is triggered from the upstream
 node on the edge. The meaning is specfic to the node type. A value
@@ -3405,7 +3413,7 @@ pub struct CUgraphInstantiateResult_enum(pub ::core::ffi::c_uint);
 pub use self::CUgraphInstantiateResult_enum as CUgraphInstantiateResult;
 /// Graph instantiation parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_GRAPH_INSTANTIATE_PARAMS_st {
     ///< Instantiation flags
     pub flags: cuuint64_t,
@@ -3520,7 +3528,7 @@ pub use self::CUlaunchMemSyncDomain_enum as CUlaunchMemSyncDomain;
 
  Domain ID range is available through ::CU_DEVICE_ATTRIBUTE_MEM_SYNC_DOMAIN_COUNT.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchMemSyncDomainMap_st {
     ///< The default domain ID to use for designated kernels
     pub default_: ::core::ffi::c_uchar,
@@ -3765,14 +3773,14 @@ scheduling policy preference for the kernel.*/
       - \p z - The Z dimension of the cluster, in blocks. Must be a divisor
                of the grid Z dimension.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchAttributeValue_union__bindgen_ty_1 {
     pub x: ::core::ffi::c_uint,
     pub y: ::core::ffi::c_uint,
     pub z: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchAttributeValue_union__bindgen_ty_2 {
     ///< Event to fire when all blocks trigger it
     pub event: CUevent,
@@ -3783,7 +3791,7 @@ pub struct CUlaunchAttributeValue_union__bindgen_ty_2 {
     pub triggerAtBlockStart: ::core::ffi::c_int,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchAttributeValue_union__bindgen_ty_3 {
     ///< Event to fire when the last block launches
     pub event: CUevent,
@@ -3791,7 +3799,7 @@ pub struct CUlaunchAttributeValue_union__bindgen_ty_3 {
     pub flags: ::core::ffi::c_int,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchAttributeValue_union__bindgen_ty_4 {
     ///< Whether or not the resulting kernel node should be device-updatable.
     pub deviceUpdatable: ::core::ffi::c_int,
@@ -3814,7 +3822,7 @@ pub struct CUlaunchAttribute_st {
 pub type CUlaunchAttribute = CUlaunchAttribute_st;
 /// CUDA extensible launch configuration
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlaunchConfig_st {
     ///< Width of grid in blocks
     pub gridDimX: ::core::ffi::c_uint,
@@ -3966,7 +3974,7 @@ pub struct CUexecAffinityType_enum(pub ::core::ffi::c_uint);
 pub use self::CUexecAffinityType_enum as CUexecAffinityType;
 /// Value for ::CU_EXEC_AFFINITY_TYPE_SM_COUNT
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUexecAffinitySmCount_st {
     ///< The number of SMs the context is limited to use.
     pub val: ::core::ffi::c_uint,
@@ -4025,7 +4033,7 @@ pub struct CUlibraryOption_enum(pub ::core::ffi::c_uint);
 /// Library options to be specified with ::cuLibraryLoadData() or ::cuLibraryLoadFromFile()
 pub use self::CUlibraryOption_enum as CUlibraryOption;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUlibraryHostUniversalFunctionAndDataTable_st {
     pub functionTable: *mut ::core::ffi::c_void,
     pub functionWindowSize: usize,
@@ -4092,7 +4100,7 @@ pub type CUoccupancyB2DSize = ::core::option::Option<
 >;
 /// 2D memory copy parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY2D_st {
     ///< Source X in bytes
     pub srcXInBytes: usize,
@@ -4133,7 +4141,7 @@ pub type CUDA_MEMCPY2D_v2 = CUDA_MEMCPY2D_st;
 pub type CUDA_MEMCPY2D = CUDA_MEMCPY2D_v2;
 /// 3D memory copy parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY3D_st {
     ///< Source X in bytes
     pub srcXInBytes: usize,
@@ -4192,7 +4200,7 @@ pub type CUDA_MEMCPY3D_v2 = CUDA_MEMCPY3D_st;
 pub type CUDA_MEMCPY3D = CUDA_MEMCPY3D_v2;
 /// 3D memory cross-context copy parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY3D_PEER_st {
     ///< Source X in bytes
     pub srcXInBytes: usize,
@@ -4251,7 +4259,7 @@ pub type CUDA_MEMCPY3D_PEER_v1 = CUDA_MEMCPY3D_PEER_st;
 pub type CUDA_MEMCPY3D_PEER = CUDA_MEMCPY3D_PEER_v1;
 /// Memcpy node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY_NODE_PARAMS_st {
     ///< Must be zero
     pub flags: ::core::ffi::c_int,
@@ -4266,7 +4274,7 @@ pub struct CUDA_MEMCPY_NODE_PARAMS_st {
 pub type CUDA_MEMCPY_NODE_PARAMS = CUDA_MEMCPY_NODE_PARAMS_st;
 /// Array descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY_DESCRIPTOR_st {
     ///< Width of array
     pub Width: usize,
@@ -4283,7 +4291,7 @@ pub type CUDA_ARRAY_DESCRIPTOR_v2 = CUDA_ARRAY_DESCRIPTOR_st;
 pub type CUDA_ARRAY_DESCRIPTOR = CUDA_ARRAY_DESCRIPTOR_v2;
 /// 3D array descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY3D_DESCRIPTOR_st {
     ///< Width of 3D array
     pub Width: usize,
@@ -4304,7 +4312,7 @@ pub type CUDA_ARRAY3D_DESCRIPTOR_v2 = CUDA_ARRAY3D_DESCRIPTOR_st;
 pub type CUDA_ARRAY3D_DESCRIPTOR = CUDA_ARRAY3D_DESCRIPTOR_v2;
 /// CUDA array sparse properties
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY_SPARSE_PROPERTIES_st {
     pub tileExtent: CUDA_ARRAY_SPARSE_PROPERTIES_st__bindgen_ty_1,
     /// First mip level at which the mip tail begins.
@@ -4316,7 +4324,7 @@ pub struct CUDA_ARRAY_SPARSE_PROPERTIES_st {
     pub reserved: [::core::ffi::c_uint; 4usize],
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY_SPARSE_PROPERTIES_st__bindgen_ty_1 {
     ///< Width of sparse tile in elements
     pub width: ::core::ffi::c_uint,
@@ -4331,7 +4339,7 @@ pub type CUDA_ARRAY_SPARSE_PROPERTIES_v1 = CUDA_ARRAY_SPARSE_PROPERTIES_st;
 pub type CUDA_ARRAY_SPARSE_PROPERTIES = CUDA_ARRAY_SPARSE_PROPERTIES_v1;
 /// CUDA array memory requirements
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY_MEMORY_REQUIREMENTS_st {
     ///< Total required memory size
     pub size: usize,
@@ -4363,19 +4371,19 @@ pub union CUDA_RESOURCE_DESC_st__bindgen_ty_1 {
     pub reserved: CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_1 {
     ///< CUDA array
     pub hArray: CUarray,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_2 {
     ///< CUDA mipmapped array
     pub hMipmappedArray: CUmipmappedArray,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3 {
     ///< Device pointer
     pub devPtr: CUdeviceptr,
@@ -4387,7 +4395,7 @@ pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3 {
     pub sizeInBytes: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4 {
     ///< Device pointer
     pub devPtr: CUdeviceptr,
@@ -4403,7 +4411,7 @@ pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4 {
     pub pitchInBytes: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5 {
     pub reserved: [::core::ffi::c_int; 32usize],
 }
@@ -4657,7 +4665,7 @@ pub struct CUresourceViewFormat_enum(pub ::core::ffi::c_uint);
 pub use self::CUresourceViewFormat_enum as CUresourceViewFormat;
 /// Resource view descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_RESOURCE_VIEW_DESC_st {
     ///< Resource view format
     pub format: CUresourceViewFormat,
@@ -4684,7 +4692,7 @@ pub type CUDA_RESOURCE_VIEW_DESC = CUDA_RESOURCE_VIEW_DESC_v1;
 /// Tensor map descriptor. Requires compiler support for aligning to 64 bytes.
 #[repr(C)]
 #[repr(align(64))]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUtensorMap_st {
     pub opaque: [cuuint64_t; 16usize],
 }
@@ -4852,7 +4860,7 @@ pub struct CUtensorMapFloatOOBfill_enum(pub ::core::ffi::c_uint);
 pub use self::CUtensorMapFloatOOBfill_enum as CUtensorMapFloatOOBfill;
 /// GPU Direct v3 tokens
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_st {
     pub p2pToken: ::core::ffi::c_ulonglong,
     pub vaSpaceToken: ::core::ffi::c_uint,
@@ -4889,7 +4897,7 @@ pub struct CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS_enum(pub ::core::ffi::c_uint);
 pub use self::CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS_enum as CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS;
 /// Kernel launch parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_LAUNCH_PARAMS_st {
     ///< Kernel to launch
     pub function: CUfunction,
@@ -5009,7 +5017,7 @@ pub union CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1 {
  ::CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT
  then 'name' must be NULL.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1 {
     /// Valid NT handle. Must be NULL if 'name' is non-NULL
     pub handle: *mut ::core::ffi::c_void,
@@ -5023,7 +5031,7 @@ pub type CUDA_EXTERNAL_MEMORY_HANDLE_DESC_v1 = CUDA_EXTERNAL_MEMORY_HANDLE_DESC_
 pub type CUDA_EXTERNAL_MEMORY_HANDLE_DESC = CUDA_EXTERNAL_MEMORY_HANDLE_DESC_v1;
 /// External memory buffer descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st {
     /// Offset into the memory object where the buffer's base is
     pub offset: ::core::ffi::c_ulonglong,
@@ -5039,7 +5047,7 @@ pub type CUDA_EXTERNAL_MEMORY_BUFFER_DESC_v1 = CUDA_EXTERNAL_MEMORY_BUFFER_DESC_
 pub type CUDA_EXTERNAL_MEMORY_BUFFER_DESC = CUDA_EXTERNAL_MEMORY_BUFFER_DESC_v1;
 /// External memory mipmap descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st {
     /** Offset into the memory object where the base level of the
  mipmap chain is.*/
@@ -5157,7 +5165,7 @@ pub union CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1 {
  - ::CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT
  then 'name' must be NULL.*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st__bindgen_ty_1__bindgen_ty_1 {
     /// Valid NT handle. Must be NULL if 'name' is non-NULL
     pub handle: *mut ::core::ffi::c_void,
@@ -5195,7 +5203,7 @@ pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1 {
 }
 /// Parameters for fence objects
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_1 {
     /// Value of fence to be signaled
     pub value: ::core::ffi::c_ulonglong,
@@ -5210,7 +5218,7 @@ pub union CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_2 {
 }
 /// Parameters for keyed mutex objects
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st__bindgen_ty_1__bindgen_ty_3 {
     /// Value of key to release the mutex with
     pub key: ::core::ffi::c_ulonglong,
@@ -5244,7 +5252,7 @@ pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1 {
 }
 /// Parameters for fence objects
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_1 {
     /// Value of fence to be waited on
     pub value: ::core::ffi::c_ulonglong,
@@ -5259,7 +5267,7 @@ pub union CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_2 {
 }
 /// Parameters for keyed mutex objects
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st__bindgen_ty_1__bindgen_ty_3 {
     /// Value of key to acquire the mutex with
     pub key: ::core::ffi::c_ulonglong,
@@ -5272,7 +5280,7 @@ pub type CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_v1 = CUDA_EXTERNAL_SEMAPHORE_WAIT_P
 pub type CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS = CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_v1;
 /// Semaphore signal node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st {
     ///< Array of external semaphore handles.
     pub extSemArray: *mut CUexternalSemaphore,
@@ -5287,7 +5295,7 @@ pub type CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v1 = CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st
 pub type CUDA_EXT_SEM_SIGNAL_NODE_PARAMS = CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v1;
 /// Semaphore signal node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st {
     ///< Array of external semaphore handles.
     pub extSemArray: *mut CUexternalSemaphore,
@@ -5300,7 +5308,7 @@ pub struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st {
 pub type CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2 = CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st;
 /// Semaphore wait node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_st {
     ///< Array of external semaphore handles.
     pub extSemArray: *mut CUexternalSemaphore,
@@ -5315,7 +5323,7 @@ pub type CUDA_EXT_SEM_WAIT_NODE_PARAMS_v1 = CUDA_EXT_SEM_WAIT_NODE_PARAMS_st;
 pub type CUDA_EXT_SEM_WAIT_NODE_PARAMS = CUDA_EXT_SEM_WAIT_NODE_PARAMS_v1;
 /// Semaphore wait node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2_st {
     ///< Array of external semaphore handles.
     pub extSemArray: *mut CUexternalSemaphore,
@@ -5575,7 +5583,7 @@ pub union CUarrayMapInfo_st__bindgen_ty_2 {
     pub miptail: CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_2,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_1 {
     ///< For CUDA mipmapped arrays must a valid mipmap level. For CUDA arrays must be zero
     pub level: ::core::ffi::c_uint,
@@ -5595,7 +5603,7 @@ pub struct CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_1 {
     pub extentDepth: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUarrayMapInfo_st__bindgen_ty_2__bindgen_ty_2 {
     ///< For CUDA layered arrays must be a valid layer index. Otherwise, must be zero
     pub layer: ::core::ffi::c_uint,
@@ -5615,7 +5623,7 @@ pub type CUarrayMapInfo_v1 = CUarrayMapInfo_st;
 pub type CUarrayMapInfo = CUarrayMapInfo_v1;
 /// Specifies a memory location.
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemLocation_st {
     ///< Specifies the location type, which modifies the meaning of id.
     pub type_: CUmemLocationType,
@@ -5646,7 +5654,7 @@ pub struct CUmemAllocationCompType_enum(pub ::core::ffi::c_uint);
 pub use self::CUmemAllocationCompType_enum as CUmemAllocationCompType;
 /// Specifies the allocation properties for a allocation.
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemAllocationProp_st {
     /// Allocation type
     pub type_: CUmemAllocationType,
@@ -5663,7 +5671,7 @@ pub struct CUmemAllocationProp_st {
     pub allocFlags: CUmemAllocationProp_st__bindgen_ty_1,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemAllocationProp_st__bindgen_ty_1 {
     /** Allocation hint for requesting compressible memory.
  On devices that support Compute Data Compression, compressible
@@ -5703,7 +5711,7 @@ pub struct CUmulticastGranularity_flags_enum(pub ::core::ffi::c_uint);
 pub use self::CUmulticastGranularity_flags_enum as CUmulticastGranularity_flags;
 /// Specifies the properties for a multicast object.
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmulticastObjectProp_st {
     /** The number of devices in the multicast team that will bind memory to this
  object*/
@@ -5723,7 +5731,7 @@ pub type CUmulticastObjectProp_v1 = CUmulticastObjectProp_st;
 pub type CUmulticastObjectProp = CUmulticastObjectProp_v1;
 /// Memory access descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemAccessDesc_st {
     ///< Location on which the request is to change it's accessibility
     pub location: CUmemLocation,
@@ -5796,7 +5804,7 @@ pub struct CUgraphExecUpdateResult_enum(pub ::core::ffi::c_uint);
 pub use self::CUgraphExecUpdateResult_enum as CUgraphExecUpdateResult;
 /// Result information returned by cuGraphExecUpdate
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUgraphExecUpdateResultInfo_st {
     /// Gives more specific detail when a cuda graph update fails.
     pub result: CUgraphExecUpdateResult,
@@ -5888,7 +5896,7 @@ pub struct CUmemPool_attribute_enum(pub ::core::ffi::c_uint);
 pub use self::CUmemPool_attribute_enum as CUmemPool_attribute;
 /// Specifies the properties of allocations made from the pool.
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemPoolProps_st {
     ///< Allocation type. Currently must be specified as CU_MEM_ALLOCATION_TYPE_PINNED
     pub allocType: CUmemAllocationType,
@@ -5912,7 +5920,7 @@ pub type CUmemPoolProps_v1 = CUmemPoolProps_st;
 pub type CUmemPoolProps = CUmemPoolProps_v1;
 /// Opaque data for exporting a pool allocation
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUmemPoolPtrExportData_st {
     pub reserved: [::core::ffi::c_uchar; 64usize],
 }
@@ -5922,7 +5930,7 @@ pub type CUmemPoolPtrExportData_v1 = CUmemPoolPtrExportData_st;
 pub type CUmemPoolPtrExportData = CUmemPoolPtrExportData_v1;
 /// Memory allocation node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEM_ALLOC_NODE_PARAMS_v1_st {
     /** in: location where the allocation should reside (specified in ::location).
  ::handleTypes must be ::CU_MEM_HANDLE_TYPE_NONE. IPC is not supported.*/
@@ -5942,7 +5950,7 @@ pub type CUDA_MEM_ALLOC_NODE_PARAMS_v1 = CUDA_MEM_ALLOC_NODE_PARAMS_v1_st;
 pub type CUDA_MEM_ALLOC_NODE_PARAMS = CUDA_MEM_ALLOC_NODE_PARAMS_v1;
 /// Memory allocation node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEM_ALLOC_NODE_PARAMS_v2_st {
     /** in: location where the allocation should reside (specified in ::location).
  ::handleTypes must be ::CU_MEM_HANDLE_TYPE_NONE. IPC is not supported.*/
@@ -5960,7 +5968,7 @@ pub struct CUDA_MEM_ALLOC_NODE_PARAMS_v2_st {
 pub type CUDA_MEM_ALLOC_NODE_PARAMS_v2 = CUDA_MEM_ALLOC_NODE_PARAMS_v2_st;
 /// Memory free node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEM_FREE_NODE_PARAMS_st {
     ///< in: the pointer to free
     pub dptr: CUdeviceptr,
@@ -6004,7 +6012,7 @@ pub struct CUgraphMem_attribute_enum(pub ::core::ffi::c_uint);
 pub use self::CUgraphMem_attribute_enum as CUgraphMem_attribute;
 /// Child graph node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_CHILD_GRAPH_NODE_PARAMS_st {
     /**< The child graph to clone into the node for node creation, or
 a handle to the graph owned by the node for node query*/
@@ -6014,7 +6022,7 @@ a handle to the graph owned by the node for node query*/
 pub type CUDA_CHILD_GRAPH_NODE_PARAMS = CUDA_CHILD_GRAPH_NODE_PARAMS_st;
 /// Event record node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EVENT_RECORD_NODE_PARAMS_st {
     ///< The event to record when the node executes
     pub event: CUevent,
@@ -6023,7 +6031,7 @@ pub struct CUDA_EVENT_RECORD_NODE_PARAMS_st {
 pub type CUDA_EVENT_RECORD_NODE_PARAMS = CUDA_EVENT_RECORD_NODE_PARAMS_st;
 /// Event wait node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_EVENT_WAIT_NODE_PARAMS_st {
     ///< The event to wait on from the node
     pub event: CUevent,
@@ -6442,7 +6450,7 @@ pub struct CUdevResourceType(pub ::core::ffi::c_uint);
 /** \struct CUdevSmResource
  Data for SM-related resources*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUdevSmResource_st {
     ///< The amount of streaming multiprocessors available in this resource. This is an output parameter only, do not write to this field.
     pub smCount: ::core::ffi::c_uint,
@@ -6507,10 +6515,10 @@ pub type CUdevResource_v1 = CUdevResource_st;
  \p sm.smCount will reflect the amount of streaming multiprocessors available in this resource.*/
 pub type CUdevResource = CUdevResource_v1;
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUdeviceptr_v1(pub ::core::ffi::c_uint);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY2D_v1_st {
     ///< Source X in bytes
     pub srcXInBytes: ::core::ffi::c_uint,
@@ -6547,7 +6555,7 @@ pub struct CUDA_MEMCPY2D_v1_st {
 }
 pub type CUDA_MEMCPY2D_v1 = CUDA_MEMCPY2D_v1_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_MEMCPY3D_v1_st {
     ///< Source X in bytes
     pub srcXInBytes: ::core::ffi::c_uint,
@@ -6602,7 +6610,7 @@ pub struct CUDA_MEMCPY3D_v1_st {
 }
 pub type CUDA_MEMCPY3D_v1 = CUDA_MEMCPY3D_v1_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY_DESCRIPTOR_v1_st {
     ///< Width of array
     pub Width: ::core::ffi::c_uint,
@@ -6615,7 +6623,7 @@ pub struct CUDA_ARRAY_DESCRIPTOR_v1_st {
 }
 pub type CUDA_ARRAY_DESCRIPTOR_v1 = CUDA_ARRAY_DESCRIPTOR_v1_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct CUDA_ARRAY3D_DESCRIPTOR_v1_st {
     ///< Width of 3D array
     pub Width: ::core::ffi::c_uint,
@@ -7870,7 +7878,7 @@ impl CUerror {
     });
 }
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Hash, Copy, Clone, PartialEq, Eq)]
 pub struct CUerror(pub ::core::num::NonZeroU32);
 pub trait CUresultConsts {
     const SUCCESS: CUresult = CUresult::Ok(());
@@ -8088,3 +8096,15 @@ impl From<hip_runtime_sys::hipErrorCode_t> for CUerror {
         Self(error.0)
     }
 }
+unsafe impl Send for CUdeviceptr {}
+unsafe impl Sync for CUdeviceptr {}
+unsafe impl Send for CUcontext {}
+unsafe impl Sync for CUcontext {}
+unsafe impl Send for CUstream {}
+unsafe impl Sync for CUstream {}
+unsafe impl Send for CUmodule {}
+unsafe impl Sync for CUmodule {}
+unsafe impl Send for CUfunction {}
+unsafe impl Sync for CUfunction {}
+unsafe impl Send for CUlibrary {}
+unsafe impl Sync for CUlibrary {}
diff --git a/ext/hip_runtime-sys/src/lib.rs b/ext/hip_runtime-sys/src/lib.rs
index 81e2db5..4aad7e6 100644
--- a/ext/hip_runtime-sys/src/lib.rs
+++ b/ext/hip_runtime-sys/src/lib.rs
@@ -151,7 +151,7 @@ pub const hipGraphKernelNodePortProgrammatic: u32 = 1;
 #[doc = " @defgroup GlobalDefs Global enum and defines\n @{\n\n/\n/**\n hipDeviceArch_t\n"]
 #[repr(C)]
 #[repr(align(4))]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipDeviceArch_t {
     pub _bitfield_align_1: [u8; 0],
     pub _bitfield_1: __BindgenBitfieldUnit<[u8; 3usize]>,
@@ -553,7 +553,7 @@ impl hipDeviceArch_t {
     }
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipUUID_t {
     pub bytes: [::core::ffi::c_char; 16usize],
 }
@@ -561,7 +561,7 @@ pub type hipUUID = hipUUID_t;
 /** hipDeviceProp
 */
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipDeviceProp_tR0600 {
     ///< Device name.
     pub name: [::core::ffi::c_char; 256usize],
@@ -836,7 +836,7 @@ impl hipMemoryType {
 pub struct hipMemoryType(pub ::core::ffi::c_uint);
 /// Pointer attributes
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipPointerAttribute_t {
     pub type_: hipMemoryType,
     pub device: ::core::ffi::c_int,
@@ -1606,7 +1606,7 @@ impl hipGPUDirectRDMAWritesOrdering {
 #[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipGPUDirectRDMAWritesOrdering(pub ::core::ffi::c_uint);
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipDeviceptr_t(pub *mut ::core::ffi::c_void);
 impl hipChannelFormatKind {
     pub const hipChannelFormatKindSigned: hipChannelFormatKind = hipChannelFormatKind(0);
@@ -1626,7 +1626,7 @@ impl hipChannelFormatKind {
 #[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipChannelFormatKind(pub ::core::ffi::c_uint);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipChannelFormatDesc {
     pub x: ::core::ffi::c_int,
     pub y: ::core::ffi::c_int,
@@ -1669,7 +1669,7 @@ impl hipArray_Format {
 #[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipArray_Format(pub ::core::ffi::c_uint);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_ARRAY_DESCRIPTOR {
     pub Width: usize,
     pub Height: usize,
@@ -1677,7 +1677,7 @@ pub struct HIP_ARRAY_DESCRIPTOR {
     pub NumChannels: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_ARRAY3D_DESCRIPTOR {
     pub Width: usize,
     pub Height: usize,
@@ -1687,7 +1687,7 @@ pub struct HIP_ARRAY3D_DESCRIPTOR {
     pub Flags: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hip_Memcpy2D {
     pub srcXInBytes: usize,
     pub srcY: usize,
@@ -1707,7 +1707,7 @@ pub struct hip_Memcpy2D {
     pub Height: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMipmappedArray {
     pub data: *mut ::core::ffi::c_void,
     pub desc: hipChannelFormatDesc,
@@ -2214,24 +2214,24 @@ pub union hipResourceDesc__bindgen_ty_1 {
     pub pitch2D: hipResourceDesc__bindgen_ty_1__bindgen_ty_4,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipResourceDesc__bindgen_ty_1__bindgen_ty_1 {
     pub array: hipArray_t,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipResourceDesc__bindgen_ty_1__bindgen_ty_2 {
     pub mipmap: hipMipmappedArray_t,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipResourceDesc__bindgen_ty_1__bindgen_ty_3 {
     pub devPtr: *mut ::core::ffi::c_void,
     pub desc: hipChannelFormatDesc,
     pub sizeInBytes: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipResourceDesc__bindgen_ty_1__bindgen_ty_4 {
     pub devPtr: *mut ::core::ffi::c_void,
     pub desc: hipChannelFormatDesc,
@@ -2258,19 +2258,19 @@ pub union HIP_RESOURCE_DESC_st__bindgen_ty_1 {
     pub reserved: HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_1 {
     ///< HIP array
     pub hArray: hipArray_t,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_2 {
     ///< HIP mipmapped array
     pub hMipmappedArray: hipMipmappedArray_t,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3 {
     ///< Device pointer
     pub devPtr: hipDeviceptr_t,
@@ -2282,7 +2282,7 @@ pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_3 {
     pub sizeInBytes: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4 {
     ///< Device pointer
     pub devPtr: hipDeviceptr_t,
@@ -2298,14 +2298,14 @@ pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_4 {
     pub pitchInBytes: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_DESC_st__bindgen_ty_1__bindgen_ty_5 {
     pub reserved: [::core::ffi::c_int; 32usize],
 }
 pub type HIP_RESOURCE_DESC = HIP_RESOURCE_DESC_st;
 /// hip resource view descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipResourceViewDesc {
     pub format: hipResourceViewFormat,
     pub width: usize,
@@ -2318,7 +2318,7 @@ pub struct hipResourceViewDesc {
 }
 /// Resource view descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_RESOURCE_VIEW_DESC_st {
     ///< Resource view format
     pub format: HIPresourceViewFormat,
@@ -2369,7 +2369,7 @@ impl hipMemcpyKind {
 #[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemcpyKind(pub ::core::ffi::c_uint);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipPitchedPtr {
     pub ptr: *mut ::core::ffi::c_void,
     pub pitch: usize,
@@ -2377,21 +2377,21 @@ pub struct hipPitchedPtr {
     pub ysize: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExtent {
     pub width: usize,
     pub height: usize,
     pub depth: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipPos {
     pub x: usize,
     pub y: usize,
     pub z: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemcpy3DParms {
     pub srcArray: hipArray_t,
     pub srcPos: hipPos,
@@ -2403,7 +2403,7 @@ pub struct hipMemcpy3DParms {
     pub kind: hipMemcpyKind,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_MEMCPY3D {
     pub srcXInBytes: usize,
     pub srcY: usize,
@@ -2739,15 +2739,17 @@ pub struct hipDeviceP2PAttr(pub ::core::ffi::c_uint);
 pub struct ihipStream_t {
     _unused: [u8; 0],
 }
-pub type hipStream_t = *mut ihipStream_t;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct hipStream_t(pub *mut ihipStream_t);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipIpcMemHandle_st {
     pub reserved: [::core::ffi::c_char; 64usize],
 }
 pub type hipIpcMemHandle_t = hipIpcMemHandle_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipIpcEventHandle_st {
     pub reserved: [::core::ffi::c_char; 64usize],
 }
@@ -2758,14 +2760,16 @@ pub struct ihipModule_t {
     _unused: [u8; 0],
 }
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipModule_t(pub *mut ihipModule_t);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct ihipModuleSymbol_t {
     _unused: [u8; 0],
 }
-pub type hipFunction_t = *mut ihipModuleSymbol_t;
+#[repr(transparent)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
+pub struct hipFunction_t(pub *mut ihipModuleSymbol_t);
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct ihipMemPoolHandle_t {
@@ -2774,7 +2778,7 @@ pub struct ihipMemPoolHandle_t {
 /// HIP memory pool
 pub type hipMemPool_t = *mut ihipMemPoolHandle_t;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipFuncAttributes {
     pub binaryVersion: ::core::ffi::c_int,
     pub cacheModeCA: ::core::ffi::c_int,
@@ -3004,7 +3008,7 @@ pub struct hipMemLocationType(pub ::core::ffi::c_uint);
 
  To specify a gpu, set type = @p hipMemLocationTypeDevice and set id = the gpu's device ID*/
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemLocation {
     ///< Specifies the location type, which describes the meaning of id
     pub type_: hipMemLocationType,
@@ -3030,7 +3034,7 @@ impl hipMemAccessFlags {
 pub struct hipMemAccessFlags(pub ::core::ffi::c_uint);
 /// Memory access descriptor
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemAccessDesc {
     ///< Location on which the accessibility has to change
     pub location: hipMemLocation,
@@ -3089,7 +3093,7 @@ impl hipMemAllocationHandleType {
 pub struct hipMemAllocationHandleType(pub ::core::ffi::c_uint);
 /// Specifies the properties of allocations made from the pool.
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemPoolProps {
     ///< Allocation type. Currently must be specified as @p hipMemAllocationTypePinned
     pub allocType: hipMemAllocationType,
@@ -3106,7 +3110,7 @@ pub struct hipMemPoolProps {
 }
 /// Opaque data structure for exporting a pool allocation
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemPoolPtrExportData {
     pub reserved: [::core::ffi::c_uchar; 64usize],
 }
@@ -3225,7 +3229,7 @@ impl hipSharedMemConfig {
 pub struct hipSharedMemConfig(pub ::core::ffi::c_uint);
 /// Struct for data in 3D
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct dim3 {
     ///< x
     pub x: u32,
@@ -3236,7 +3240,7 @@ pub struct dim3 {
 }
 /// struct hipLaunchParams_t
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipLaunchParams_t {
     ///< Device function symbol
     pub func: *mut ::core::ffi::c_void,
@@ -3255,7 +3259,7 @@ pub struct hipLaunchParams_t {
 pub type hipLaunchParams = hipLaunchParams_t;
 /// struct hipFunctionLaunchParams_t
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipFunctionLaunchParams_t {
     ///< Kernel to launch
     pub function: hipFunction_t,
@@ -3341,14 +3345,14 @@ pub union hipExternalMemoryHandleDesc_st__bindgen_ty_1 {
     pub nvSciBufObject: *const ::core::ffi::c_void,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalMemoryHandleDesc_st__bindgen_ty_1__bindgen_ty_1 {
     pub handle: *mut ::core::ffi::c_void,
     pub name: *const ::core::ffi::c_void,
 }
 pub type hipExternalMemoryHandleDesc = hipExternalMemoryHandleDesc_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalMemoryBufferDesc_st {
     pub offset: ::core::ffi::c_ulonglong,
     pub size: ::core::ffi::c_ulonglong,
@@ -3357,7 +3361,7 @@ pub struct hipExternalMemoryBufferDesc_st {
 }
 pub type hipExternalMemoryBufferDesc = hipExternalMemoryBufferDesc_st;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalMemoryMipmappedArrayDesc_st {
     pub offset: ::core::ffi::c_ulonglong,
     pub formatDesc: hipChannelFormatDesc,
@@ -3437,7 +3441,7 @@ pub union hipExternalSemaphoreHandleDesc_st__bindgen_ty_1 {
     pub NvSciSyncObj: *const ::core::ffi::c_void,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreHandleDesc_st__bindgen_ty_1__bindgen_ty_1 {
     pub handle: *mut ::core::ffi::c_void,
     pub name: *const ::core::ffi::c_void,
@@ -3460,7 +3464,7 @@ pub struct hipExternalSemaphoreSignalParams_st__bindgen_ty_1 {
     pub reserved: [::core::ffi::c_uint; 12usize],
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreSignalParams_st__bindgen_ty_1__bindgen_ty_1 {
     pub value: ::core::ffi::c_ulonglong,
 }
@@ -3471,7 +3475,7 @@ pub union hipExternalSemaphoreSignalParams_st__bindgen_ty_1__bindgen_ty_2 {
     pub reserved: ::core::ffi::c_ulonglong,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreSignalParams_st__bindgen_ty_1__bindgen_ty_3 {
     pub key: ::core::ffi::c_ulonglong,
 }
@@ -3493,7 +3497,7 @@ pub struct hipExternalSemaphoreWaitParams_st__bindgen_ty_1 {
     pub reserved: [::core::ffi::c_uint; 10usize],
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreWaitParams_st__bindgen_ty_1__bindgen_ty_1 {
     pub value: ::core::ffi::c_ulonglong,
 }
@@ -3504,7 +3508,7 @@ pub union hipExternalSemaphoreWaitParams_st__bindgen_ty_1__bindgen_ty_2 {
     pub reserved: ::core::ffi::c_ulonglong,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreWaitParams_st__bindgen_ty_1__bindgen_ty_3 {
     pub key: ::core::ffi::c_ulonglong,
     pub timeoutMs: ::core::ffi::c_uint,
@@ -3644,13 +3648,13 @@ pub type hipHostFn_t = ::core::option::Option<
     unsafe extern "C" fn(userData: *mut ::core::ffi::c_void),
 >;
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipHostNodeParams {
     pub fn_: hipHostFn_t,
     pub userData: *mut ::core::ffi::c_void,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipKernelNodeParams {
     pub blockDim: dim3,
     pub extra: *mut *mut ::core::ffi::c_void,
@@ -3660,7 +3664,7 @@ pub struct hipKernelNodeParams {
     pub sharedMemBytes: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemsetParams {
     pub dst: *mut ::core::ffi::c_void,
     pub elementSize: ::core::ffi::c_uint,
@@ -3670,7 +3674,7 @@ pub struct hipMemsetParams {
     pub width: usize,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemAllocNodeParams {
     /**< Pool properties, which contain where
 < the location should reside*/
@@ -3741,7 +3745,7 @@ priority of kernel.*/
 }
 /// Memset node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct HIP_MEMSET_NODE_PARAMS {
     ///< Destination pointer on device
     pub dst: hipDeviceptr_t,
@@ -4021,7 +4025,7 @@ due to the nodes belonging to different contexts*/
 pub struct hipGraphInstantiateResult(pub ::core::ffi::c_uint);
 /// Graph Instantiation parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipGraphInstantiateParams {
     ///< The node which caused instantiation to fail, if any
     pub errNode_out: hipGraphNode_t,
@@ -4035,7 +4039,7 @@ If it failed, the reason why*/
 }
 /// Memory allocation properties
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemAllocationProp {
     ///< Memory allocation type
     pub type_: hipMemAllocationType,
@@ -4048,7 +4052,7 @@ pub struct hipMemAllocationProp {
     pub allocFlags: hipMemAllocationProp__bindgen_ty_1,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemAllocationProp__bindgen_ty_1 {
     ///< Compression type
     pub compressionType: ::core::ffi::c_uchar,
@@ -4059,7 +4063,7 @@ pub struct hipMemAllocationProp__bindgen_ty_1 {
 }
 /// External semaphore signal node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreSignalNodeParams {
     pub extSemArray: *mut hipExternalSemaphore_t,
     pub paramsArray: *const hipExternalSemaphoreSignalParams,
@@ -4067,7 +4071,7 @@ pub struct hipExternalSemaphoreSignalNodeParams {
 }
 /// External semaphore wait node parameters
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipExternalSemaphoreWaitNodeParams {
     pub extSemArray: *mut hipExternalSemaphore_t,
     pub paramsArray: *const hipExternalSemaphoreWaitParams,
@@ -4169,7 +4173,7 @@ pub union hipArrayMapInfo__bindgen_ty_2 {
     pub miptail: hipArrayMapInfo__bindgen_ty_2__bindgen_ty_2,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipArrayMapInfo__bindgen_ty_2__bindgen_ty_1 {
     ///< For mipmapped arrays must be a valid mipmap level. For arrays must be zero
     pub level: ::core::ffi::c_uint,
@@ -4189,7 +4193,7 @@ pub struct hipArrayMapInfo__bindgen_ty_2__bindgen_ty_1 {
     pub extentDepth: ::core::ffi::c_uint,
 }
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipArrayMapInfo__bindgen_ty_2__bindgen_ty_2 {
     ///< For layered arrays must be a valid layer index. Otherwise, must be zero
     pub layer: ::core::ffi::c_uint,
@@ -4205,7 +4209,7 @@ pub union hipArrayMapInfo__bindgen_ty_3 {
 }
 /// Memcpy node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemcpyNodeParams {
     ///< Must be zero.
     pub flags: ::core::ffi::c_int,
@@ -4216,7 +4220,7 @@ pub struct hipMemcpyNodeParams {
 }
 /// Child graph node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipChildGraphNodeParams {
     /**< Either the child graph to clone into the node, or
 < a handle to the graph possesed by the node used during query*/
@@ -4224,21 +4228,21 @@ pub struct hipChildGraphNodeParams {
 }
 /// Event record node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipEventWaitNodeParams {
     ///< Event to wait on
     pub event: hipEvent_t,
 }
 /// Event record node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipEventRecordNodeParams {
     ///< The event to be recorded when node executes
     pub event: hipEvent_t,
 }
 /// Memory free node params
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipMemFreeNodeParams {
     ///< the pointer to be freed
     pub dptr: *mut ::core::ffi::c_void,
@@ -4282,7 +4286,7 @@ impl hipGraphDependencyType {
 #[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipGraphDependencyType(pub ::core::ffi::c_uint);
 #[repr(C)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Copy, Clone, Hash, PartialEq, Eq)]
 pub struct hipGraphEdgeData {
     /**< This indicates when the dependency is triggered from the upstream node on the
 < edge. The meaning is specfic to the node type. A value of 0 in all cases
@@ -12444,7 +12448,7 @@ impl hipErrorCode_t {
     });
 }
 #[repr(transparent)]
-#[derive(Debug, Copy, Clone, PartialEq, Eq)]
+#[derive(Debug, Hash, Copy, Clone, PartialEq, Eq)]
 pub struct hipErrorCode_t(pub ::core::num::NonZeroU32);
 pub trait hipError_tConsts {
     const Success: hipError_t = hipError_t::Ok(());
@@ -12654,5 +12658,11 @@ pub type hipError_t = ::core::result::Result<(), hipErrorCode_t>;
 const _: fn() = || {
     let _ = std::mem::transmute::<hipError_t, u32>;
 };
+unsafe impl Send for hipDeviceptr_t {}
+unsafe impl Sync for hipDeviceptr_t {}
+unsafe impl Send for hipStream_t {}
+unsafe impl Sync for hipStream_t {}
 unsafe impl Send for hipModule_t {}
 unsafe impl Sync for hipModule_t {}
+unsafe impl Send for hipFunction_t {}
+unsafe impl Sync for hipFunction_t {}
diff --git a/zluda_bindgen/src/main.rs b/zluda_bindgen/src/main.rs
index 3d7ea2e..7332254 100644
--- a/zluda_bindgen/src/main.rs
+++ b/zluda_bindgen/src/main.rs
@@ -3,9 +3,9 @@ use quote::{format_ident, quote, ToTokens};
 use rustc_hash::{FxHashMap, FxHashSet};
 use std::{collections::hash_map, fs::File, io::Write, iter, path::PathBuf, str::FromStr};
 use syn::{
-    parse_quote, punctuated::Punctuated, visit_mut::VisitMut, Abi, Fields, FnArg, ForeignItem,
-    ForeignItemFn, Ident, Item, ItemConst, ItemForeignMod, ItemUse, LitStr, Path, PathArguments,
-    Signature, Type, TypePath, UseTree,
+    parse_quote, punctuated::Punctuated, visit_mut::VisitMut, Abi, Fields, FieldsUnnamed, FnArg,
+    ForeignItem, ForeignItemFn, Ident, Item, ItemConst, ItemForeignMod, ItemUse, LitStr, Path,
+    PathArguments, Signature, Type, TypePath, UseTree,
 };
 
 fn main() {
@@ -22,6 +22,7 @@ fn main() {
             is_bitfield: false,
             is_global: false,
         })
+        .derive_hash(true)
         .derive_eq(true)
         .header_contents("cuda_wrapper.h", include_str!("../build/cuda_wrapper.h"))
         .allowlist_type("^CU.*")
@@ -30,8 +31,12 @@ fn main() {
         .must_use_type("cudaError_enum")
         .constified_enum("cudaError_enum")
         .no_partialeq("CUDA_HOST_NODE_PARAMS_st")
-        .new_type_alias(r"^CUdevice_v\d+$")
         .new_type_alias(r"^CUdeviceptr_v\d+$")
+        .new_type_alias(r"^CUcontext$")
+        .new_type_alias(r"^CUstream$")
+        .new_type_alias(r"^CUmodule$")
+        .new_type_alias(r"^CUfunction$")
+        .new_type_alias(r"^CUlibrary$")
         .clang_args(["-I/usr/local/cuda/include"])
         .generate()
         .unwrap()
@@ -56,6 +61,7 @@ fn generate_hip_runtime(output: &PathBuf, path: &[&str]) {
             is_bitfield: false,
             is_global: false,
         })
+        .derive_hash(true)
         .derive_eq(true)
         .header("/opt/rocm/include/hip/hip_runtime_api.h")
         .allowlist_type("^hip.*")
@@ -64,7 +70,9 @@ fn generate_hip_runtime(output: &PathBuf, path: &[&str]) {
         .must_use_type("hipError_t")
         .constified_enum("hipError_t")
         .new_type_alias("^hipDeviceptr_t$")
+        .new_type_alias("^hipStream_t$")
         .new_type_alias("^hipModule_t$")
+        .new_type_alias("^hipFunction_t$")
         .clang_args(["-I/opt/rocm/include", "-D__HIP_PLATFORM_AMD__"])
         .generate()
         .unwrap()
@@ -89,7 +97,15 @@ fn generate_hip_runtime(output: &PathBuf, path: &[&str]) {
         })
         .collect::<Vec<_>>();
     converter.flush(&mut module.items);
-    add_send_sync(&mut module.items, &["hipModule_t"]);
+    add_send_sync(
+        &mut module.items,
+        &[
+            "hipDeviceptr_t",
+            "hipStream_t",
+            "hipModule_t",
+            "hipFunction_t",
+        ],
+    );
     let mut output = output.clone();
     output.extend(path);
     write_rust_to_file(output, &prettyplease::unparse(&module))
@@ -176,6 +192,17 @@ fn generate_types(output: &PathBuf, path: &[&str], module: &syn::File) {
             }
         }
     });
+    add_send_sync(
+        &mut module.items,
+        &[
+            "CUdeviceptr",
+            "CUcontext",
+            "CUstream",
+            "CUmodule",
+            "CUfunction",
+            "CUlibrary",
+        ],
+    );
     syn::visit_mut::visit_file_mut(&mut FixAbi, &mut module);
     let mut output = output.clone();
     output.extend(path);
@@ -252,7 +279,7 @@ impl ConvertIntoRustResult {
                 #(#error_variants)*
             }
             #[repr(transparent)]
-            #[derive(Debug, Copy, Clone, PartialEq, Eq)]
+            #[derive(Debug, Hash, Copy, Clone, PartialEq, Eq)]
             pub struct #new_error_type(pub ::core::num::NonZeroU32);
 
             pub trait #type_trait {
@@ -327,6 +354,8 @@ fn generate_display(
     module: &syn::File,
 ) {
     let ignore_types = [
+        "CUdevice",
+        "CUdeviceptr_v1",
         "CUarrayMapInfo_st",
         "CUDA_RESOURCE_DESC_st",
         "CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st",
@@ -545,9 +574,9 @@ fn cuda_derive_display_trait_for_item<'a>(
                 })
             } else {
                 let struct_ = &item_struct.ident;
-                let (first_field, rest_of_fields) = match item_struct.fields {
+                match item_struct.fields {
                     Fields::Named(ref fields) => {
-                        let mut all_idents = fields.named.iter().filter_map(|f| {
+                        let mut rest_of_fields = fields.named.iter().filter_map(|f| {
                             let f_ident = f.ident.as_ref().unwrap();
                             let name = f_ident.to_string();
                             if name.starts_with("reserved") || name == "_unused" {
@@ -556,27 +585,35 @@ fn cuda_derive_display_trait_for_item<'a>(
                                 Some(f_ident)
                             }
                         });
-                        let first = match all_idents.next() {
+                        let first_field = match rest_of_fields.next() {
                             Some(f) => f,
                             None => return None,
                         };
-                        (first, all_idents)
+                        Some(parse_quote! {
+                            impl crate::format::CudaDisplay for #path_prefix :: #struct_ {
+                                fn write(&self, _fn_name: &'static str, _index: usize, writer: &mut (impl std::io::Write + ?Sized)) -> std::io::Result<()> {
+                                    writer.write_all(concat!("{ ", stringify!(#first_field), ": ").as_bytes())?;
+                                    crate::format::CudaDisplay::write(&self.#first_field, "", 0, writer)?;
+                                    #(
+                                        writer.write_all(concat!(", ", stringify!(#rest_of_fields), ": ").as_bytes())?;
+                                        crate::format::CudaDisplay::write(&self.#rest_of_fields, "", 0, writer)?;
+                                    )*
+                                    writer.write_all(b" }")
+                                }
+                            }
+                        })
                     }
-                    _ => return None,
-                };
-                Some(parse_quote! {
-                    impl crate::format::CudaDisplay for #path_prefix :: #struct_ {
-                        fn write(&self, _fn_name: &'static str, _index: usize, writer: &mut (impl std::io::Write + ?Sized)) -> std::io::Result<()> {
-                            writer.write_all(concat!("{ ", stringify!(#first_field), ": ").as_bytes())?;
-                            crate::format::CudaDisplay::write(&self.#first_field, "", 0, writer)?;
-                            #(
-                                writer.write_all(concat!(", ", stringify!(#rest_of_fields), ": ").as_bytes())?;
-                                crate::format::CudaDisplay::write(&self.#rest_of_fields, "", 0, writer)?;
-                            )*
-                            writer.write_all(b" }")
-                        }
+                    Fields::Unnamed(FieldsUnnamed { ref unnamed, .. }) if unnamed.len() == 1 => {
+                        Some(parse_quote! {
+                            impl crate::format::CudaDisplay for #path_prefix :: #struct_ {
+                                fn write(&self, _fn_name: &'static str, _index: usize, writer: &mut (impl std::io::Write + ?Sized)) -> std::io::Result<()> {
+                                    write!(writer, "{:p}", self.0)
+                                }
+                            }
+                        })
                     }
-                })
+                    _ => return None,
+                }
             }
         }
         Item::Type(item_type) => {
diff --git a/zluda_dump/src/format.rs b/zluda_dump/src/format.rs
index a6fc885..c1aac61 100644
--- a/zluda_dump/src/format.rs
+++ b/zluda_dump/src/format.rs
@@ -26,28 +26,6 @@ impl CudaDisplay for cuda_types::CUuuid {
     }

 }

 

-impl CudaDisplay for cuda_types::CUdevice {

-    fn write(

-        &self,

-        _fn_name: &'static str,

-        _index: usize,

-        writer: &mut (impl std::io::Write + ?Sized),

-    ) -> std::io::Result<()> {

-        write!(writer, "{}", self.0)

-    }

-}

-

-impl CudaDisplay for cuda_types::CUdeviceptr {

-    fn write(

-        &self,

-        _fn_name: &'static str,

-        _index: usize,

-        writer: &mut (impl std::io::Write + ?Sized),

-    ) -> std::io::Result<()> {

-        write!(writer, "{:p}", self.0)

-    }

-}

-

 impl CudaDisplay for cuda_types::CUdeviceptr_v1 {

     fn write(

         &self,

diff --git a/zluda_dump/src/format_generated.rs b/zluda_dump/src/format_generated.rs
index 5d28997..742fb57 100644
--- a/zluda_dump/src/format_generated.rs
+++ b/zluda_dump/src/format_generated.rs
@@ -1,6 +1,16 @@
 // Generated automatically by zluda_bindgen
 // DO NOT EDIT MANUALLY
 #![allow(warnings)]
+impl crate::format::CudaDisplay for cuda_types::CUdeviceptr_v2 {
+    fn write(
+        &self,
+        _fn_name: &'static str,
+        _index: usize,
+        writer: &mut (impl std::io::Write + ?Sized),
+    ) -> std::io::Result<()> {
+        write!(writer, "{:p}", self.0)
+    }
+}
 impl crate::format::CudaDisplay for cuda_types::CUcontext {
     fn write(
         &self,
@@ -8,7 +18,7 @@ impl crate::format::CudaDisplay for cuda_types::CUcontext {
         _index: usize,
         writer: &mut (impl std::io::Write + ?Sized),
     ) -> std::io::Result<()> {
-        write!(writer, "{:p}", *self)
+        write!(writer, "{:p}", self.0)
     }
 }
 impl crate::format::CudaDisplay for cuda_types::CUmodule {
@@ -18,7 +28,7 @@ impl crate::format::CudaDisplay for cuda_types::CUmodule {
         _index: usize,
         writer: &mut (impl std::io::Write + ?Sized),
     ) -> std::io::Result<()> {
-        write!(writer, "{:p}", *self)
+        write!(writer, "{:p}", self.0)
     }
 }
 impl crate::format::CudaDisplay for cuda_types::CUfunction {
@@ -28,7 +38,7 @@ impl crate::format::CudaDisplay for cuda_types::CUfunction {
         _index: usize,
         writer: &mut (impl std::io::Write + ?Sized),
     ) -> std::io::Result<()> {
-        write!(writer, "{:p}", *self)
+        write!(writer, "{:p}", self.0)
     }
 }
 impl crate::format::CudaDisplay for cuda_types::CUlibrary {
@@ -38,7 +48,7 @@ impl crate::format::CudaDisplay for cuda_types::CUlibrary {
         _index: usize,
         writer: &mut (impl std::io::Write + ?Sized),
     ) -> std::io::Result<()> {
-        write!(writer, "{:p}", *self)
+        write!(writer, "{:p}", self.0)
     }
 }
 impl crate::format::CudaDisplay for cuda_types::CUkernel {
@@ -108,7 +118,7 @@ impl crate::format::CudaDisplay for cuda_types::CUstream {
         _index: usize,
         writer: &mut (impl std::io::Write + ?Sized),
     ) -> std::io::Result<()> {
-        write!(writer, "{:p}", *self)
+        write!(writer, "{:p}", self.0)
     }
 }
 impl crate::format::CudaDisplay for cuda_types::CUgraphicsResource {