From 6fa27aac42ae59b985ed96093fed4a5604d3dee5 Mon Sep 17 00:00:00 2001 From: Ryan Houdek Date: Tue, 10 Mar 2026 12:42:08 -0700 Subject: [PATCH] Thunks: Implement support for CUDA This is enough to get less complex cuda applications running, and is a good starting spot to slowly finish off the remaining implementation. Some information: - 429 functions in total - 153 only compiled for 64-bit (35.6%) - 7 functions disabled entirely (1.6%) The main thing /not/ working with this initial implementation is .cu files compiled in to an ELF using the static cuda runtime. This is due to the `cuGetExportTable` function being stubbed out and the static cuda RT requires at least two interfaces from that function before it continues. This function isn't publicly documented by NVIDIA but has been publicly reverse engineered to be fairly trivial. It's just a jump table with the first element being the size of the table in bytes. That will be the next step of the implementation. --- Data/ThunksDB.json | 7 + ThunkLibs/GuestLibs/CMakeLists.txt | 3 + ThunkLibs/HostLibs/CMakeLists.txt | 3 + ThunkLibs/libcuda/cuda_defines.h | 1017 +++++++++++++++ ThunkLibs/libcuda/fixup_defines.h | 15 + ThunkLibs/libcuda/libcuda_Guest.cpp | 84 ++ ThunkLibs/libcuda/libcuda_Host.cpp | 81 ++ ThunkLibs/libcuda/libcuda_interface.cpp | 1571 +++++++++++++++++++++++ 8 files changed, 2781 insertions(+) create mode 100644 ThunkLibs/libcuda/cuda_defines.h create mode 100644 ThunkLibs/libcuda/fixup_defines.h create mode 100644 ThunkLibs/libcuda/libcuda_Guest.cpp create mode 100644 ThunkLibs/libcuda/libcuda_Host.cpp create mode 100644 ThunkLibs/libcuda/libcuda_interface.cpp diff --git a/Data/ThunksDB.json b/Data/ThunksDB.json index 67f7ef4d2..d52ff0481 100644 --- a/Data/ThunksDB.json +++ b/Data/ThunksDB.json @@ -46,6 +46,13 @@ "@PREFIX_LIB@/libwayland-client.so.0", "@PREFIX_LIB@/libwayland-client.so.0.20.0" ] + }, + "cuda": { + "Library" : "libcuda-guest.so", + "Overlay": [ + "@PREFIX_LIB@/libcuda.so", + "@PREFIX_LIB@/libcuda.so.1" + ] } } } diff --git a/ThunkLibs/GuestLibs/CMakeLists.txt b/ThunkLibs/GuestLibs/CMakeLists.txt index 661971215..0b3069204 100644 --- a/ThunkLibs/GuestLibs/CMakeLists.txt +++ b/ThunkLibs/GuestLibs/CMakeLists.txt @@ -266,3 +266,6 @@ add_library(PlaceholderX11 SHARED ../libX11/libX11_NativeGuest.cpp) target_link_options(PlaceholderX11 PRIVATE "LINKER:-soname,libX11.so.6") set_target_properties(PlaceholderX11 PROPERTIES NO_SONAME ON) target_link_libraries(GL-guest PRIVATE PlaceholderX11) + +generate(libcuda ${CMAKE_CURRENT_SOURCE_DIR}/../libcuda/libcuda_interface.cpp) +add_guest_lib(cuda "libcuda.so.1") diff --git a/ThunkLibs/HostLibs/CMakeLists.txt b/ThunkLibs/HostLibs/CMakeLists.txt index 28b423420..a526c72ff 100644 --- a/ThunkLibs/HostLibs/CMakeLists.txt +++ b/ThunkLibs/HostLibs/CMakeLists.txt @@ -167,6 +167,9 @@ foreach(GUEST_BITNESS IN LISTS BITNESS_LIST) find_package(OpenGL REQUIRED) target_link_libraries(GL-host-${GUEST_BITNESS} PRIVATE OpenGL::GL) + + generate(libcuda ${CMAKE_CURRENT_SOURCE_DIR}/../libcuda/libcuda_interface.cpp ${GUEST_BITNESS}) + add_host_lib(cuda ${GUEST_BITNESS}) endforeach() if (BUILD_FEX_LINUX_TESTS) diff --git a/ThunkLibs/libcuda/cuda_defines.h b/ThunkLibs/libcuda/cuda_defines.h new file mode 100644 index 000000000..7a0522156 --- /dev/null +++ b/ThunkLibs/libcuda/cuda_defines.h @@ -0,0 +1,1017 @@ +// SPDX-License-Identifier: MIT +// Generated file. Be careful if modifying +#pragma once +#include +#include "fixup_defines.h" +#ifdef __cplusplus +extern "C" { +#endif + +// Type forward declarations +typedef uint32_t cuuint32_t; +typedef uint64_t cuuint64_t; +typedef unsigned long long CUdeviceptr_v2; +typedef CUdeviceptr_v2 CUdeviceptr; +typedef int CUdevice_v1; +typedef CUdevice_v1 CUdevice; +struct CUctx_st; +typedef struct CUctx_st* CUcontext; +struct CUmod_st; +typedef struct CUmod_st* CUmodule; +struct CUfunc_st; +typedef struct CUfunc_st* CUfunction; +struct CUlib_st; +typedef struct CUlib_st* CUlibrary; +struct CUkern_st; +typedef struct CUkern_st* CUkernel; +struct CUarray_st; +typedef struct CUarray_st* CUarray; +struct CUmipmappedArray_st; +typedef struct CUmipmappedArray_st* CUmipmappedArray; +struct CUtexref_st; +typedef struct CUtexref_st* CUtexref; +struct CUsurfref_st; +typedef struct CUsurfref_st* CUsurfref; +struct CUevent_st; +typedef struct CUevent_st* CUevent; +struct CUstream_st; +typedef struct CUstream_st* CUstream; +struct CUgraphicsResource_st; +typedef struct CUgraphicsResource_st* CUgraphicsResource; +typedef unsigned long long CUtexObject_v1; +typedef CUtexObject_v1 CUtexObject; +typedef unsigned long long CUsurfObject_v1; +typedef CUsurfObject_v1 CUsurfObject; +struct CUextMemory_st; +typedef struct CUextMemory_st* CUexternalMemory; +struct CUextSemaphore_st; +typedef struct CUextSemaphore_st* CUexternalSemaphore; +struct CUgraph_st; +typedef struct CUgraph_st* CUgraph; +struct CUgraphNode_st; +typedef struct CUgraphNode_st* CUgraphNode; +struct CUgraphExec_st; +typedef struct CUgraphExec_st* CUgraphExec; +struct CUmemPoolHandle_st; +typedef struct CUmemPoolHandle_st* CUmemoryPool; +struct CUuserObject_st; +typedef struct CUuserObject_st* CUuserObject; +typedef cuuint64_t CUgraphConditionalHandle; +struct CUgraphDeviceUpdatableNode_st; +typedef struct CUgraphDeviceUpdatableNode_st* CUgraphDeviceNode; +struct CUasyncCallbackEntry_st; +typedef struct CUasyncCallbackEntry_st* CUasyncCallbackHandle; +struct CUuuid_st { + char _0[16]; +}; +typedef struct CUuuid_st CUuuid; +struct CUmemFabricHandle_st { + unsigned char _0[64]; +}; +typedef struct CUmemFabricHandle_st CUmemFabricHandle_v1; +typedef CUmemFabricHandle_v1 CUmemFabricHandle; +struct CUipcEventHandle_st { + char _0[64]; +}; +typedef struct CUipcEventHandle_st CUipcEventHandle_v1; +typedef CUipcEventHandle_v1 CUipcEventHandle; +struct CUipcMemHandle_st { + char _0[64]; +}; +typedef struct CUipcMemHandle_st CUipcMemHandle_v1; +typedef CUipcMemHandle_v1 CUipcMemHandle; +enum CUipcMem_flags_enum : unsigned int {}; +typedef enum CUipcMem_flags_enum CUipcMem_flags; +enum CUmemAttach_flags_enum : unsigned int {}; +typedef enum CUmemAttach_flags_enum CUmemAttach_flags; +enum CUctx_flags_enum : unsigned int {}; +typedef enum CUctx_flags_enum CUctx_flags; +enum CUevent_sched_flags_enum : unsigned int {}; +typedef enum CUevent_sched_flags_enum CUevent_sched_flags; +enum CUstream_flags_enum : unsigned int {}; +typedef enum CUstream_flags_enum CUstream_flags; +enum CUevent_flags_enum : unsigned int {}; +typedef enum CUevent_flags_enum CUevent_flags; +enum CUevent_record_flags_enum : unsigned int {}; +typedef enum CUevent_record_flags_enum CUevent_record_flags; +enum CUevent_wait_flags_enum : unsigned int {}; +typedef enum CUevent_wait_flags_enum CUevent_wait_flags; +enum CUstreamWaitValue_flags_enum : unsigned int {}; +typedef enum CUstreamWaitValue_flags_enum CUstreamWaitValue_flags; +enum CUstreamWriteValue_flags_enum : unsigned int {}; +typedef enum CUstreamWriteValue_flags_enum CUstreamWriteValue_flags; +enum CUstreamBatchMemOpType_enum : unsigned int {}; +typedef enum CUstreamBatchMemOpType_enum CUstreamBatchMemOpType; +enum CUstreamMemoryBarrier_flags_enum : unsigned int {}; +typedef enum CUstreamMemoryBarrier_flags_enum CUstreamMemoryBarrier_flags; +union CUstreamBatchMemOpParams_union; +typedef union CUstreamBatchMemOpParams_union CUstreamBatchMemOpParams_v1; +typedef CUstreamBatchMemOpParams_v1 CUstreamBatchMemOpParams; +struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st; +typedef struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st CUDA_BATCH_MEM_OP_NODE_PARAMS_v1; +typedef CUDA_BATCH_MEM_OP_NODE_PARAMS_v1 CUDA_BATCH_MEM_OP_NODE_PARAMS; +struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st; +typedef struct CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st CUDA_BATCH_MEM_OP_NODE_PARAMS_v2; +enum CUoccupancy_flags_enum : unsigned int {}; +typedef enum CUoccupancy_flags_enum CUoccupancy_flags; +enum CUstreamUpdateCaptureDependencies_flags_enum : unsigned int {}; +typedef enum CUstreamUpdateCaptureDependencies_flags_enum CUstreamUpdateCaptureDependencies_flags; +enum CUasyncNotificationType_enum : unsigned int {}; +typedef enum CUasyncNotificationType_enum CUasyncNotificationType; +struct CUasyncNotificationInfo_st; +typedef struct CUasyncNotificationInfo_st CUasyncNotificationInfo; +enum CUarray_format_enum : unsigned int {}; +typedef enum CUarray_format_enum CUarray_format; +enum CUaddress_mode_enum : unsigned int {}; +typedef enum CUaddress_mode_enum CUaddress_mode; +enum CUfilter_mode_enum : unsigned int {}; +typedef enum CUfilter_mode_enum CUfilter_mode; +enum CUdevice_attribute_enum : unsigned int {}; +typedef enum CUdevice_attribute_enum CUdevice_attribute; +struct CUdevprop_st { + int _0; + int _1[3]; + int _2[3]; + int _3; + int _4; + int _5; + int _6; + int _7; + int _8; + int _9; +}; +typedef struct CUdevprop_st CUdevprop_v1; +typedef CUdevprop_v1 CUdevprop; +enum CUpointer_attribute_enum : unsigned int {}; +typedef enum CUpointer_attribute_enum CUpointer_attribute; +enum CUfunction_attribute_enum : unsigned int {}; +typedef enum CUfunction_attribute_enum CUfunction_attribute; +enum CUfunc_cache_enum : unsigned int {}; +typedef enum CUfunc_cache_enum CUfunc_cache; +enum CUsharedconfig_enum : unsigned int {}; +typedef enum CUsharedconfig_enum CUsharedconfig; +enum CUshared_carveout_enum : int {}; +typedef enum CUshared_carveout_enum CUshared_carveout; +enum CUmemorytype_enum : unsigned int {}; +typedef enum CUmemorytype_enum CUmemorytype; +enum CUcomputemode_enum : unsigned int {}; +typedef enum CUcomputemode_enum CUcomputemode; +enum CUmem_advise_enum : unsigned int {}; +typedef enum CUmem_advise_enum CUmem_advise; +enum CUmem_range_attribute_enum : unsigned int {}; +typedef enum CUmem_range_attribute_enum CUmem_range_attribute; +enum CUjit_option_enum : unsigned int {}; +typedef enum CUjit_option_enum CUjit_option; +enum CUjit_target_enum : unsigned int {}; +typedef enum CUjit_target_enum CUjit_target; +enum CUjit_fallback_enum : unsigned int {}; +typedef enum CUjit_fallback_enum CUjit_fallback; +enum CUjit_cacheMode_enum : unsigned int {}; +typedef enum CUjit_cacheMode_enum CUjit_cacheMode; +enum CUjitInputType_enum : unsigned int {}; +typedef enum CUjitInputType_enum CUjitInputType; +struct CUlinkState_st; +typedef struct CUlinkState_st* CUlinkState; +enum CUgraphicsRegisterFlags_enum : unsigned int {}; +typedef enum CUgraphicsRegisterFlags_enum CUgraphicsRegisterFlags; +enum CUgraphicsMapResourceFlags_enum : unsigned int {}; +typedef enum CUgraphicsMapResourceFlags_enum CUgraphicsMapResourceFlags; +enum CUarray_cubemap_face_enum : unsigned int {}; +typedef enum CUarray_cubemap_face_enum CUarray_cubemap_face; +enum CUlimit_enum : unsigned int {}; +typedef enum CUlimit_enum CUlimit; +enum CUresourcetype_enum : unsigned int {}; +typedef enum CUresourcetype_enum CUresourcetype; +enum CUaccessProperty_enum : unsigned int {}; +typedef enum CUaccessProperty_enum CUaccessProperty; +struct CUaccessPolicyWindow_st; +typedef struct CUaccessPolicyWindow_st CUaccessPolicyWindow_v1; +typedef CUaccessPolicyWindow_v1 CUaccessPolicyWindow; +struct CUDA_KERNEL_NODE_PARAMS_st; +typedef struct CUDA_KERNEL_NODE_PARAMS_st CUDA_KERNEL_NODE_PARAMS_v1; +struct CUDA_KERNEL_NODE_PARAMS_v2_st; +typedef struct CUDA_KERNEL_NODE_PARAMS_v2_st CUDA_KERNEL_NODE_PARAMS_v2; +typedef CUDA_KERNEL_NODE_PARAMS_v2 CUDA_KERNEL_NODE_PARAMS; +struct CUDA_KERNEL_NODE_PARAMS_v3_st; +typedef struct CUDA_KERNEL_NODE_PARAMS_v3_st CUDA_KERNEL_NODE_PARAMS_v3; +struct CUDA_MEMSET_NODE_PARAMS_st; +typedef struct CUDA_MEMSET_NODE_PARAMS_st CUDA_MEMSET_NODE_PARAMS_v1; +typedef CUDA_MEMSET_NODE_PARAMS_v1 CUDA_MEMSET_NODE_PARAMS; +struct CUDA_MEMSET_NODE_PARAMS_v2_st; +typedef struct CUDA_MEMSET_NODE_PARAMS_v2_st CUDA_MEMSET_NODE_PARAMS_v2; +struct CUDA_HOST_NODE_PARAMS_st { + CUhostFn _0; + void* _1; +}; +typedef struct CUDA_HOST_NODE_PARAMS_st CUDA_HOST_NODE_PARAMS_v1; +typedef CUDA_HOST_NODE_PARAMS_v1 CUDA_HOST_NODE_PARAMS; +struct CUDA_HOST_NODE_PARAMS_v2_st { + CUhostFn _0; + void* _1; +}; +typedef struct CUDA_HOST_NODE_PARAMS_v2_st CUDA_HOST_NODE_PARAMS_v2; +enum CUgraphConditionalNodeType_enum : unsigned int {}; +typedef enum CUgraphConditionalNodeType_enum CUgraphConditionalNodeType; +enum CUgraphNodeType_enum : unsigned int {}; +typedef enum CUgraphNodeType_enum CUgraphNodeType; +enum CUgraphDependencyType_enum : unsigned int {}; +typedef enum CUgraphDependencyType_enum CUgraphDependencyType; +struct CUgraphEdgeData_st { + unsigned char _0; + unsigned char _1; + unsigned char _2; + unsigned char _3[5]; +}; +typedef struct CUgraphEdgeData_st CUgraphEdgeData; +enum CUgraphInstantiateResult_enum : unsigned int {}; +typedef enum CUgraphInstantiateResult_enum CUgraphInstantiateResult; +struct CUDA_GRAPH_INSTANTIATE_PARAMS_st; +typedef struct CUDA_GRAPH_INSTANTIATE_PARAMS_st CUDA_GRAPH_INSTANTIATE_PARAMS; +enum CUsynchronizationPolicy_enum : unsigned int {}; +typedef enum CUsynchronizationPolicy_enum CUsynchronizationPolicy; +enum CUclusterSchedulingPolicy_enum : unsigned int {}; +typedef enum CUclusterSchedulingPolicy_enum CUclusterSchedulingPolicy; +enum CUlaunchMemSyncDomain_enum : unsigned int {}; +typedef enum CUlaunchMemSyncDomain_enum CUlaunchMemSyncDomain; +struct CUlaunchMemSyncDomainMap_st { + unsigned char _0; + unsigned char _1; +}; +typedef struct CUlaunchMemSyncDomainMap_st CUlaunchMemSyncDomainMap; +enum CUlaunchAttributeID_enum : unsigned int {}; +typedef enum CUlaunchAttributeID_enum CUlaunchAttributeID; +union CUlaunchAttributeValue_union; +typedef union CUlaunchAttributeValue_union CUlaunchAttributeValue; +struct CUlaunchAttribute_st; +typedef struct CUlaunchAttribute_st CUlaunchAttribute; +struct CUlaunchConfig_st; +typedef struct CUlaunchConfig_st CUlaunchConfig; +typedef CUlaunchAttributeID CUkernelNodeAttrID; +typedef CUlaunchAttributeValue CUkernelNodeAttrValue_v1; +typedef CUkernelNodeAttrValue_v1 CUkernelNodeAttrValue; +enum CUstreamCaptureStatus_enum : unsigned int {}; +typedef enum CUstreamCaptureStatus_enum CUstreamCaptureStatus; +enum CUstreamCaptureMode_enum : unsigned int {}; +typedef enum CUstreamCaptureMode_enum CUstreamCaptureMode; +typedef CUlaunchAttributeID CUstreamAttrID; +typedef CUlaunchAttributeValue CUstreamAttrValue_v1; +typedef CUstreamAttrValue_v1 CUstreamAttrValue; +enum CUdriverProcAddress_flags_enum : unsigned int {}; +typedef enum CUdriverProcAddress_flags_enum CUdriverProcAddress_flags; +enum CUdriverProcAddressQueryResult_enum : unsigned int {}; +typedef enum CUdriverProcAddressQueryResult_enum CUdriverProcAddressQueryResult; +enum CUexecAffinityType_enum : unsigned int {}; +typedef enum CUexecAffinityType_enum CUexecAffinityType; +struct CUexecAffinitySmCount_st { + unsigned int _0; +}; +typedef struct CUexecAffinitySmCount_st CUexecAffinitySmCount_v1; +typedef CUexecAffinitySmCount_v1 CUexecAffinitySmCount; +struct CUexecAffinityParam_st; +typedef struct CUexecAffinityParam_st CUexecAffinityParam_v1; +typedef CUexecAffinityParam_v1 CUexecAffinityParam; +enum CUlibraryOption_enum : unsigned int {}; +typedef enum CUlibraryOption_enum CUlibraryOption; +struct CUlibraryHostUniversalFunctionAndDataTable_st { + void* _0; + size_t _1; + void* _2; + size_t _3; +}; +typedef struct CUlibraryHostUniversalFunctionAndDataTable_st CUlibraryHostUniversalFunctionAndDataTable; +typedef enum cudaError_enum CUresult; +enum CUdevice_P2PAttribute_enum : unsigned int {}; +typedef enum CUdevice_P2PAttribute_enum CUdevice_P2PAttribute; +struct CUDA_MEMCPY2D_st; +typedef struct CUDA_MEMCPY2D_st CUDA_MEMCPY2D_v2; +typedef CUDA_MEMCPY2D_v2 CUDA_MEMCPY2D; +struct CUDA_MEMCPY3D_st { + size_t _0; + size_t _1; + size_t _2; + size_t _3; + CUmemorytype _4; + const void* _5; + CUdeviceptr _6; + CUarray _7; + void* _8; + size_t _9; + size_t _10; + size_t _11; + size_t _12; + size_t _13; + size_t _14; + CUmemorytype _15; + void* _16; + CUdeviceptr _17; + CUarray _18; + void* _19; + size_t _20; + size_t _21; + size_t _22; + size_t _23; + size_t _24; +}; +typedef struct CUDA_MEMCPY3D_st CUDA_MEMCPY3D_v2; +typedef CUDA_MEMCPY3D_v2 CUDA_MEMCPY3D; +struct CUDA_MEMCPY3D_PEER_st { + size_t _0; + size_t _1; + size_t _2; + size_t _3; + CUmemorytype _4; + const void* _5; + CUdeviceptr _6; + CUarray _7; + CUcontext _8; + size_t _9; + size_t _10; + size_t _11; + size_t _12; + size_t _13; + size_t _14; + CUmemorytype _15; + void* _16; + CUdeviceptr _17; + CUarray _18; + CUcontext _19; + size_t _20; + size_t _21; + size_t _22; + size_t _23; + size_t _24; +}; +typedef struct CUDA_MEMCPY3D_PEER_st CUDA_MEMCPY3D_PEER_v1; +typedef CUDA_MEMCPY3D_PEER_v1 CUDA_MEMCPY3D_PEER; +struct CUDA_MEMCPY_NODE_PARAMS_st; +typedef struct CUDA_MEMCPY_NODE_PARAMS_st CUDA_MEMCPY_NODE_PARAMS; +struct CUDA_ARRAY_DESCRIPTOR_st; +typedef struct CUDA_ARRAY_DESCRIPTOR_st CUDA_ARRAY_DESCRIPTOR_v2; +typedef CUDA_ARRAY_DESCRIPTOR_v2 CUDA_ARRAY_DESCRIPTOR; +struct CUDA_ARRAY3D_DESCRIPTOR_st; +typedef struct CUDA_ARRAY3D_DESCRIPTOR_st CUDA_ARRAY3D_DESCRIPTOR_v2; +typedef CUDA_ARRAY3D_DESCRIPTOR_v2 CUDA_ARRAY3D_DESCRIPTOR; +struct CUDA_ARRAY_SPARSE_PROPERTIES_st; +typedef struct CUDA_ARRAY_SPARSE_PROPERTIES_st CUDA_ARRAY_SPARSE_PROPERTIES_v1; +typedef CUDA_ARRAY_SPARSE_PROPERTIES_v1 CUDA_ARRAY_SPARSE_PROPERTIES; +struct CUDA_ARRAY_MEMORY_REQUIREMENTS_st { + size_t _0; + size_t _1; + unsigned int _2[4]; +}; +typedef struct CUDA_ARRAY_MEMORY_REQUIREMENTS_st CUDA_ARRAY_MEMORY_REQUIREMENTS_v1; +typedef CUDA_ARRAY_MEMORY_REQUIREMENTS_v1 CUDA_ARRAY_MEMORY_REQUIREMENTS; +struct CUDA_RESOURCE_DESC_st; +typedef struct CUDA_RESOURCE_DESC_st CUDA_RESOURCE_DESC_v1; +typedef CUDA_RESOURCE_DESC_v1 CUDA_RESOURCE_DESC; +struct CUDA_TEXTURE_DESC_st; +typedef struct CUDA_TEXTURE_DESC_st CUDA_TEXTURE_DESC_v1; +typedef CUDA_TEXTURE_DESC_v1 CUDA_TEXTURE_DESC; +enum CUresourceViewFormat_enum : unsigned int {}; +typedef enum CUresourceViewFormat_enum CUresourceViewFormat; +struct CUDA_RESOURCE_VIEW_DESC_st; +typedef struct CUDA_RESOURCE_VIEW_DESC_st CUDA_RESOURCE_VIEW_DESC_v1; +typedef CUDA_RESOURCE_VIEW_DESC_v1 CUDA_RESOURCE_VIEW_DESC; +struct CUtensorMap_st { + cuuint64_t _0[16]; +}; +typedef struct CUtensorMap_st CUtensorMap; +enum CUtensorMapDataType_enum : unsigned int {}; +typedef enum CUtensorMapDataType_enum CUtensorMapDataType; +enum CUtensorMapInterleave_enum : unsigned int {}; +typedef enum CUtensorMapInterleave_enum CUtensorMapInterleave; +enum CUtensorMapSwizzle_enum : unsigned int {}; +typedef enum CUtensorMapSwizzle_enum CUtensorMapSwizzle; +enum CUtensorMapL2promotion_enum : unsigned int {}; +typedef enum CUtensorMapL2promotion_enum CUtensorMapL2promotion; +enum CUtensorMapFloatOOBfill_enum : unsigned int {}; +typedef enum CUtensorMapFloatOOBfill_enum CUtensorMapFloatOOBfill; +struct CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_st { + unsigned long long _0; + unsigned int _1; +}; +typedef struct CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_st CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_v1; +typedef CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_v1 CUDA_POINTER_ATTRIBUTE_P2P_TOKENS; +enum CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS_enum : unsigned int {}; +typedef enum CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS_enum CUDA_POINTER_ATTRIBUTE_ACCESS_FLAGS; +struct CUDA_LAUNCH_PARAMS_st { + CUfunction _0; + unsigned int _1; + unsigned int _2; + unsigned int _3; + unsigned int _4; + unsigned int _5; + unsigned int _6; + unsigned int _7; + CUstream _8; + void** _9; +}; +typedef struct CUDA_LAUNCH_PARAMS_st CUDA_LAUNCH_PARAMS_v1; +typedef CUDA_LAUNCH_PARAMS_v1 CUDA_LAUNCH_PARAMS; +enum CUexternalMemoryHandleType_enum : unsigned int {}; +typedef enum CUexternalMemoryHandleType_enum CUexternalMemoryHandleType; +struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st; +typedef struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st CUDA_EXTERNAL_MEMORY_HANDLE_DESC_v1; +typedef CUDA_EXTERNAL_MEMORY_HANDLE_DESC_v1 CUDA_EXTERNAL_MEMORY_HANDLE_DESC; +struct CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st { + unsigned long long _0; + unsigned long long _1; + unsigned int _2; + unsigned int _3[16]; +}; +typedef struct CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st CUDA_EXTERNAL_MEMORY_BUFFER_DESC_v1; +typedef CUDA_EXTERNAL_MEMORY_BUFFER_DESC_v1 CUDA_EXTERNAL_MEMORY_BUFFER_DESC; +struct CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st; +typedef struct CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_v1; +typedef CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_v1 CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC; +enum CUexternalSemaphoreHandleType_enum : unsigned int {}; +typedef enum CUexternalSemaphoreHandleType_enum CUexternalSemaphoreHandleType; +struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st; +typedef struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_v1; +typedef CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_v1 CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC; +struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st; +typedef struct CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_st CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_v1; +typedef CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS_v1 CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS; +struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st; +typedef struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_v1; +typedef CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_v1 CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS; +struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st { + CUexternalSemaphore* _0; + const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS* _1; + unsigned int _2; +}; +typedef struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v1; +typedef CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v1 CUDA_EXT_SEM_SIGNAL_NODE_PARAMS; +struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st { + CUexternalSemaphore* _0; + const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS* _1; + unsigned int _2; +}; +typedef struct CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2; +struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_st { + CUexternalSemaphore* _0; + const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS* _1; + unsigned int _2; +}; +typedef struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_st CUDA_EXT_SEM_WAIT_NODE_PARAMS_v1; +typedef CUDA_EXT_SEM_WAIT_NODE_PARAMS_v1 CUDA_EXT_SEM_WAIT_NODE_PARAMS; +struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2_st { + CUexternalSemaphore* _0; + const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS* _1; + unsigned int _2; +}; +typedef struct CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2_st CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2; +typedef unsigned long long CUmemGenericAllocationHandle_v1; +typedef CUmemGenericAllocationHandle_v1 CUmemGenericAllocationHandle; +enum CUmemAllocationHandleType_enum : unsigned int {}; +typedef enum CUmemAllocationHandleType_enum CUmemAllocationHandleType; +enum CUmemAccess_flags_enum : unsigned int {}; +typedef enum CUmemAccess_flags_enum CUmemAccess_flags; +enum CUmemLocationType_enum : unsigned int {}; +typedef enum CUmemLocationType_enum CUmemLocationType; +enum CUmemAllocationType_enum : unsigned int {}; +typedef enum CUmemAllocationType_enum CUmemAllocationType; +enum CUmemAllocationGranularity_flags_enum : unsigned int {}; +typedef enum CUmemAllocationGranularity_flags_enum CUmemAllocationGranularity_flags; +enum CUmemRangeHandleType_enum : unsigned int {}; +typedef enum CUmemRangeHandleType_enum CUmemRangeHandleType; +enum CUarraySparseSubresourceType_enum : unsigned int {}; +typedef enum CUarraySparseSubresourceType_enum CUarraySparseSubresourceType; +enum CUmemOperationType_enum : unsigned int {}; +typedef enum CUmemOperationType_enum CUmemOperationType; +enum CUmemHandleType_enum : unsigned int {}; +typedef enum CUmemHandleType_enum CUmemHandleType; +struct CUarrayMapInfo_st; +typedef struct CUarrayMapInfo_st CUarrayMapInfo_v1; +typedef CUarrayMapInfo_v1 CUarrayMapInfo; +struct CUmemLocation_st { + CUmemLocationType _0; + int _1; +}; +typedef struct CUmemLocation_st CUmemLocation_v1; +typedef CUmemLocation_v1 CUmemLocation; +enum CUmemAllocationCompType_enum : unsigned int {}; +typedef enum CUmemAllocationCompType_enum CUmemAllocationCompType; +struct CUmemAllocationProp_st; +typedef struct CUmemAllocationProp_st CUmemAllocationProp_v1; +typedef CUmemAllocationProp_v1 CUmemAllocationProp; +enum CUmulticastGranularity_flags_enum : unsigned int {}; +typedef enum CUmulticastGranularity_flags_enum CUmulticastGranularity_flags; +struct CUmulticastObjectProp_st { + unsigned int _0; + size_t _1; + unsigned long long _2; + unsigned long long _3; +}; +typedef struct CUmulticastObjectProp_st CUmulticastObjectProp_v1; +typedef CUmulticastObjectProp_v1 CUmulticastObjectProp; +struct CUmemAccessDesc_st; +typedef struct CUmemAccessDesc_st CUmemAccessDesc_v1; +typedef CUmemAccessDesc_v1 CUmemAccessDesc; +enum CUgraphExecUpdateResult_enum : unsigned int {}; +typedef enum CUgraphExecUpdateResult_enum CUgraphExecUpdateResult; +struct CUgraphExecUpdateResultInfo_st; +typedef struct CUgraphExecUpdateResultInfo_st CUgraphExecUpdateResultInfo_v1; +typedef CUgraphExecUpdateResultInfo_v1 CUgraphExecUpdateResultInfo; +enum CUmemPool_attribute_enum : unsigned int {}; +typedef enum CUmemPool_attribute_enum CUmemPool_attribute; +struct CUmemPoolProps_st; +typedef struct CUmemPoolProps_st CUmemPoolProps_v1; +typedef CUmemPoolProps_v1 CUmemPoolProps; +struct CUmemPoolPtrExportData_st { + unsigned char _0[64]; +}; +typedef struct CUmemPoolPtrExportData_st CUmemPoolPtrExportData_v1; +typedef CUmemPoolPtrExportData_v1 CUmemPoolPtrExportData; +struct CUDA_MEM_ALLOC_NODE_PARAMS_v1_st; +typedef struct CUDA_MEM_ALLOC_NODE_PARAMS_v1_st CUDA_MEM_ALLOC_NODE_PARAMS_v1; +typedef CUDA_MEM_ALLOC_NODE_PARAMS_v1 CUDA_MEM_ALLOC_NODE_PARAMS; +struct CUDA_MEM_ALLOC_NODE_PARAMS_v2_st; +typedef struct CUDA_MEM_ALLOC_NODE_PARAMS_v2_st CUDA_MEM_ALLOC_NODE_PARAMS_v2; +struct CUDA_MEM_FREE_NODE_PARAMS_st; +typedef struct CUDA_MEM_FREE_NODE_PARAMS_st CUDA_MEM_FREE_NODE_PARAMS; +enum CUgraphMem_attribute_enum : unsigned int {}; +typedef enum CUgraphMem_attribute_enum CUgraphMem_attribute; +struct CUDA_CHILD_GRAPH_NODE_PARAMS_st; +typedef struct CUDA_CHILD_GRAPH_NODE_PARAMS_st CUDA_CHILD_GRAPH_NODE_PARAMS; +struct CUDA_EVENT_RECORD_NODE_PARAMS_st; +typedef struct CUDA_EVENT_RECORD_NODE_PARAMS_st CUDA_EVENT_RECORD_NODE_PARAMS; +struct CUDA_EVENT_WAIT_NODE_PARAMS_st; +typedef struct CUDA_EVENT_WAIT_NODE_PARAMS_st CUDA_EVENT_WAIT_NODE_PARAMS; +struct CUgraphNodeParams_st; +typedef struct CUgraphNodeParams_st CUgraphNodeParams; +enum CUflushGPUDirectRDMAWritesOptions_enum : unsigned int {}; +typedef enum CUflushGPUDirectRDMAWritesOptions_enum CUflushGPUDirectRDMAWritesOptions; +enum CUGPUDirectRDMAWritesOrdering_enum : unsigned int {}; +typedef enum CUGPUDirectRDMAWritesOrdering_enum CUGPUDirectRDMAWritesOrdering; +enum CUflushGPUDirectRDMAWritesScope_enum : unsigned int {}; +typedef enum CUflushGPUDirectRDMAWritesScope_enum CUflushGPUDirectRDMAWritesScope; +enum CUflushGPUDirectRDMAWritesTarget_enum : unsigned int {}; +typedef enum CUflushGPUDirectRDMAWritesTarget_enum CUflushGPUDirectRDMAWritesTarget; +enum CUgraphDebugDot_flags_enum : unsigned int {}; +typedef enum CUgraphDebugDot_flags_enum CUgraphDebugDot_flags; +enum CUuserObject_flags_enum : unsigned int {}; +typedef enum CUuserObject_flags_enum CUuserObject_flags; +enum CUuserObjectRetain_flags_enum : unsigned int {}; +typedef enum CUuserObjectRetain_flags_enum CUuserObjectRetain_flags; +enum CUgraphInstantiate_flags_enum : unsigned int {}; +typedef enum CUgraphInstantiate_flags_enum CUgraphInstantiate_flags; +enum CUdeviceNumaConfig_enum : unsigned int {}; +typedef enum CUdeviceNumaConfig_enum CUdeviceNumaConfig; +enum CUmoduleLoadingMode_enum : unsigned int {}; +typedef enum CUmoduleLoadingMode_enum CUmoduleLoadingMode; +enum CUfunctionLoadingState_enum : unsigned int {}; +typedef enum CUfunctionLoadingState_enum CUfunctionLoadingState; +enum CUcoredumpSettings_enum : unsigned int {}; +typedef enum CUcoredumpSettings_enum CUcoredumpSettings; +struct CUgreenCtx_st; +typedef struct CUgreenCtx_st* CUgreenCtx; +struct CUdevResourceDesc_st; +typedef struct CUdevResourceDesc_st* CUdevResourceDesc; +typedef enum CUgreenCtxCreate_flags CUgreenCtxCreate_flags; +typedef enum CUdevResourceType CUdevResourceType; +struct CUdevSmResource_st { + unsigned int _0; +}; +typedef struct CUdevSmResource_st CUdevSmResource; +struct CUdevResource_st; +typedef struct CUdevResource_st CUdevResource_v1; +typedef CUdevResource_v1 CUdevResource; + +// Function forward declarations +CUresult cuGetErrorString(CUresult, const char**); +CUresult cuGetErrorName(CUresult, const char**); +CUresult cuInit(unsigned int); +CUresult cuDriverGetVersion(int*); +CUresult cuDeviceGet(CUdevice*, int); +CUresult cuDeviceGetCount(int*); +CUresult cuDeviceGetName(char*, int, CUdevice); +CUresult cuDeviceGetUuid(CUuuid*, CUdevice); +CUresult cuDeviceGetUuid_v2(CUuuid*, CUdevice); +CUresult cuDeviceGetLuid(char*, unsigned int*, CUdevice); +CUresult cuDeviceTotalMem_v2(size_t*, CUdevice); +CUresult cuDeviceGetTexture1DLinearMaxWidth(size_t*, CUarray_format, unsigned int, CUdevice); +CUresult cuDeviceGetAttribute(int*, CUdevice_attribute, CUdevice); +CUresult cuDeviceGetNvSciSyncAttributes(void*, CUdevice, int); +CUresult cuDeviceSetMemPool(CUdevice, CUmemoryPool); +CUresult cuDeviceGetMemPool(CUmemoryPool*, CUdevice); +CUresult cuDeviceGetDefaultMemPool(CUmemoryPool*, CUdevice); +CUresult cuDeviceGetExecAffinitySupport(int*, CUexecAffinityType, CUdevice); +CUresult cuFlushGPUDirectRDMAWrites(CUflushGPUDirectRDMAWritesTarget, CUflushGPUDirectRDMAWritesScope); +CUresult cuDeviceGetProperties(CUdevprop*, CUdevice); +CUresult cuDeviceComputeCapability(int*, int*, CUdevice); +CUresult cuDevicePrimaryCtxRetain(CUcontext*, CUdevice); +CUresult cuDevicePrimaryCtxRelease_v2(CUdevice); +CUresult cuDevicePrimaryCtxSetFlags_v2(CUdevice, unsigned int); +CUresult cuDevicePrimaryCtxGetState(CUdevice, unsigned int*, int*); +CUresult cuDevicePrimaryCtxReset_v2(CUdevice); +CUresult cuCtxCreate_v2(CUcontext*, unsigned int, CUdevice); +CUresult cuCtxCreate_v3(CUcontext*, CUexecAffinityParam*, int, unsigned int, CUdevice); +CUresult cuCtxDestroy_v2(CUcontext); +CUresult cuCtxPushCurrent_v2(CUcontext); +CUresult cuCtxPopCurrent_v2(CUcontext*); +CUresult cuCtxSetCurrent(CUcontext); +CUresult cuCtxGetCurrent(CUcontext*); +CUresult cuCtxGetDevice(CUdevice*); +CUresult cuCtxGetFlags(unsigned int*); +CUresult cuCtxSetFlags(unsigned int); +CUresult cuCtxGetId(CUcontext, unsigned long long*); +CUresult cuCtxSynchronize(); +CUresult cuCtxSetLimit(CUlimit, size_t); +CUresult cuCtxGetLimit(size_t*, CUlimit); +CUresult cuCtxGetCacheConfig(CUfunc_cache*); +CUresult cuCtxSetCacheConfig(CUfunc_cache); +CUresult cuCtxGetApiVersion(CUcontext, unsigned int*); +CUresult cuCtxGetStreamPriorityRange(int*, int*); +CUresult cuCtxResetPersistingL2Cache(); +CUresult cuCtxGetExecAffinity(CUexecAffinityParam*, CUexecAffinityType); +CUresult cuCtxAttach(CUcontext*, unsigned int); +CUresult cuCtxDetach(CUcontext); +CUresult cuCtxGetSharedMemConfig(CUsharedconfig*); +CUresult cuCtxSetSharedMemConfig(CUsharedconfig); +CUresult cuModuleLoad(CUmodule*, const char*); +CUresult cuModuleLoadData(CUmodule*, const void*); +CUresult cuModuleLoadDataEx(CUmodule*, const void*, unsigned int, CUjit_option*, void**); +CUresult cuModuleLoadFatBinary(CUmodule*, const void*); +CUresult cuModuleUnload(CUmodule); +CUresult cuModuleGetLoadingMode(CUmoduleLoadingMode*); +CUresult cuModuleGetFunction(CUfunction*, CUmodule, const char*); +CUresult cuModuleGetFunctionCount(unsigned int*, CUmodule); +CUresult cuModuleEnumerateFunctions(CUfunction*, unsigned int, CUmodule); +CUresult cuModuleGetGlobal_v2(CUdeviceptr*, size_t*, CUmodule, const char*); +CUresult cuLinkCreate_v2(unsigned int, CUjit_option*, void**, CUlinkState*); +CUresult cuLinkAddData_v2(CUlinkState, CUjitInputType, void*, size_t, const char*, unsigned int, CUjit_option*, void**); +CUresult cuLinkAddFile_v2(CUlinkState, CUjitInputType, const char*, unsigned int, CUjit_option*, void**); +CUresult cuLinkComplete(CUlinkState, void**, size_t*); +CUresult cuLinkDestroy(CUlinkState); +CUresult cuModuleGetTexRef(CUtexref*, CUmodule, const char*); +CUresult cuModuleGetSurfRef(CUsurfref*, CUmodule, const char*); +CUresult cuLibraryLoadData(CUlibrary*, const void*, CUjit_option*, void**, unsigned int, CUlibraryOption*, void**, unsigned int); +CUresult cuLibraryLoadFromFile(CUlibrary*, const char*, CUjit_option*, void**, unsigned int, CUlibraryOption*, void**, unsigned int); +CUresult cuLibraryUnload(CUlibrary); +CUresult cuLibraryGetKernel(CUkernel*, CUlibrary, const char*); +CUresult cuLibraryGetKernelCount(unsigned int*, CUlibrary); +CUresult cuLibraryEnumerateKernels(CUkernel*, unsigned int, CUlibrary); +CUresult cuLibraryGetModule(CUmodule*, CUlibrary); +CUresult cuKernelGetFunction(CUfunction*, CUkernel); +CUresult cuLibraryGetGlobal(CUdeviceptr*, size_t*, CUlibrary, const char*); +CUresult cuLibraryGetManaged(CUdeviceptr*, size_t*, CUlibrary, const char*); +CUresult cuLibraryGetUnifiedFunction(void**, CUlibrary, const char*); +CUresult cuKernelGetAttribute(int*, CUfunction_attribute, CUkernel, CUdevice); +CUresult cuKernelSetAttribute(CUfunction_attribute, int, CUkernel, CUdevice); +CUresult cuKernelSetCacheConfig(CUkernel, CUfunc_cache, CUdevice); +CUresult cuKernelGetName(const char**, CUkernel); +CUresult cuKernelGetParamInfo(CUkernel, size_t, size_t*, size_t*); +CUresult cuMemGetInfo_v2(size_t*, size_t*); +CUresult cuMemAlloc_v2(CUdeviceptr*, size_t); +CUresult cuMemAllocPitch_v2(CUdeviceptr*, size_t*, size_t, size_t, unsigned int); +CUresult cuMemFree_v2(CUdeviceptr); +CUresult cuMemGetAddressRange_v2(CUdeviceptr*, size_t*, CUdeviceptr); +CUresult cuMemAllocHost_v2(void**, size_t); +CUresult cuMemFreeHost(void*); +CUresult cuMemHostAlloc(void**, size_t, unsigned int); +CUresult cuMemHostGetDevicePointer_v2(CUdeviceptr*, void*, unsigned int); +CUresult cuMemHostGetFlags(unsigned int*, void*); +CUresult cuMemAllocManaged(CUdeviceptr*, size_t, unsigned int); +// CUresult cuDeviceRegisterAsyncNotification(CUdevice, CUasyncCallback, void *, CUasyncCallbackHandle *); +CUresult cuDeviceUnregisterAsyncNotification(CUdevice, CUasyncCallbackHandle); +CUresult cuDeviceGetByPCIBusId(CUdevice*, const char*); +CUresult cuDeviceGetPCIBusId(char*, int, CUdevice); +CUresult cuIpcGetEventHandle(CUipcEventHandle*, CUevent); +CUresult cuIpcOpenEventHandle(CUevent*, CUipcEventHandle); +CUresult cuIpcGetMemHandle(CUipcMemHandle*, CUdeviceptr); +CUresult cuIpcOpenMemHandle_v2(CUdeviceptr*, CUipcMemHandle, unsigned int); +CUresult cuIpcCloseMemHandle(CUdeviceptr); +CUresult cuMemHostRegister_v2(void*, size_t, unsigned int); +CUresult cuMemHostUnregister(void*); +CUresult cuMemcpy(CUdeviceptr, CUdeviceptr, size_t); +CUresult cuMemcpyPeer(CUdeviceptr, CUcontext, CUdeviceptr, CUcontext, size_t); +CUresult cuMemcpyHtoD_v2(CUdeviceptr, const void*, size_t); +CUresult cuMemcpyDtoH_v2(void*, CUdeviceptr, size_t); +CUresult cuMemcpyDtoD_v2(CUdeviceptr, CUdeviceptr, size_t); +CUresult cuMemcpyDtoA_v2(CUarray, size_t, CUdeviceptr, size_t); +CUresult cuMemcpyAtoD_v2(CUdeviceptr, CUarray, size_t, size_t); +CUresult cuMemcpyHtoA_v2(CUarray, size_t, const void*, size_t); +CUresult cuMemcpyAtoH_v2(void*, CUarray, size_t, size_t); +CUresult cuMemcpyAtoA_v2(CUarray, size_t, CUarray, size_t, size_t); +CUresult cuMemcpy2D_v2(const CUDA_MEMCPY2D*); +CUresult cuMemcpy2DUnaligned_v2(const CUDA_MEMCPY2D*); +CUresult cuMemcpy3D_v2(const CUDA_MEMCPY3D*); +CUresult cuMemcpy3DPeer(const CUDA_MEMCPY3D_PEER*); +CUresult cuMemcpyAsync(CUdeviceptr, CUdeviceptr, size_t, CUstream); +CUresult cuMemcpyPeerAsync(CUdeviceptr, CUcontext, CUdeviceptr, CUcontext, size_t, CUstream); +CUresult cuMemcpyHtoDAsync_v2(CUdeviceptr, const void*, size_t, CUstream); +CUresult cuMemcpyDtoHAsync_v2(void*, CUdeviceptr, size_t, CUstream); +CUresult cuMemcpyDtoDAsync_v2(CUdeviceptr, CUdeviceptr, size_t, CUstream); +CUresult cuMemcpyHtoAAsync_v2(CUarray, size_t, const void*, size_t, CUstream); +CUresult cuMemcpyAtoHAsync_v2(void*, CUarray, size_t, size_t, CUstream); +CUresult cuMemcpy2DAsync_v2(const CUDA_MEMCPY2D*, CUstream); +CUresult cuMemcpy3DAsync_v2(const CUDA_MEMCPY3D*, CUstream); +CUresult cuMemcpy3DPeerAsync(const CUDA_MEMCPY3D_PEER*, CUstream); +CUresult cuMemsetD8_v2(CUdeviceptr, unsigned char, size_t); +CUresult cuMemsetD16_v2(CUdeviceptr, unsigned short, size_t); +CUresult cuMemsetD32_v2(CUdeviceptr, unsigned int, size_t); +CUresult cuMemsetD2D8_v2(CUdeviceptr, size_t, unsigned char, size_t, size_t); +CUresult cuMemsetD2D16_v2(CUdeviceptr, size_t, unsigned short, size_t, size_t); +CUresult cuMemsetD2D32_v2(CUdeviceptr, size_t, unsigned int, size_t, size_t); +CUresult cuMemsetD8Async(CUdeviceptr, unsigned char, size_t, CUstream); +CUresult cuMemsetD16Async(CUdeviceptr, unsigned short, size_t, CUstream); +CUresult cuMemsetD32Async(CUdeviceptr, unsigned int, size_t, CUstream); +CUresult cuMemsetD2D8Async(CUdeviceptr, size_t, unsigned char, size_t, size_t, CUstream); +CUresult cuMemsetD2D16Async(CUdeviceptr, size_t, unsigned short, size_t, size_t, CUstream); +CUresult cuMemsetD2D32Async(CUdeviceptr, size_t, unsigned int, size_t, size_t, CUstream); +CUresult cuArrayCreate_v2(CUarray*, const CUDA_ARRAY_DESCRIPTOR*); +CUresult cuArrayGetDescriptor_v2(CUDA_ARRAY_DESCRIPTOR*, CUarray); +CUresult cuArrayGetSparseProperties(CUDA_ARRAY_SPARSE_PROPERTIES*, CUarray); +CUresult cuMipmappedArrayGetSparseProperties(CUDA_ARRAY_SPARSE_PROPERTIES*, CUmipmappedArray); +CUresult cuArrayGetMemoryRequirements(CUDA_ARRAY_MEMORY_REQUIREMENTS*, CUarray, CUdevice); +CUresult cuMipmappedArrayGetMemoryRequirements(CUDA_ARRAY_MEMORY_REQUIREMENTS*, CUmipmappedArray, CUdevice); +CUresult cuArrayGetPlane(CUarray*, CUarray, unsigned int); +CUresult cuArrayDestroy(CUarray); +CUresult cuArray3DCreate_v2(CUarray*, const CUDA_ARRAY3D_DESCRIPTOR*); +CUresult cuArray3DGetDescriptor_v2(CUDA_ARRAY3D_DESCRIPTOR*, CUarray); +CUresult cuMipmappedArrayCreate(CUmipmappedArray*, const CUDA_ARRAY3D_DESCRIPTOR*, unsigned int); +CUresult cuMipmappedArrayGetLevel(CUarray*, CUmipmappedArray, unsigned int); +CUresult cuMipmappedArrayDestroy(CUmipmappedArray); +CUresult cuMemGetHandleForAddressRange(void*, CUdeviceptr, size_t, CUmemRangeHandleType, unsigned long long); +CUresult cuMemAddressReserve(CUdeviceptr*, size_t, size_t, CUdeviceptr, unsigned long long); +CUresult cuMemAddressFree(CUdeviceptr, size_t); +CUresult cuMemCreate(CUmemGenericAllocationHandle*, size_t, const CUmemAllocationProp*, unsigned long long); +CUresult cuMemRelease(CUmemGenericAllocationHandle); +CUresult cuMemMap(CUdeviceptr, size_t, size_t, CUmemGenericAllocationHandle, unsigned long long); +CUresult cuMemMapArrayAsync(CUarrayMapInfo*, unsigned int, CUstream); +CUresult cuMemUnmap(CUdeviceptr, size_t); +CUresult cuMemSetAccess(CUdeviceptr, size_t, const CUmemAccessDesc*, size_t); +CUresult cuMemGetAccess(unsigned long long*, const CUmemLocation*, CUdeviceptr); +CUresult cuMemExportToShareableHandle(void*, CUmemGenericAllocationHandle, CUmemAllocationHandleType, unsigned long long); +CUresult cuMemImportFromShareableHandle(CUmemGenericAllocationHandle*, void*, CUmemAllocationHandleType); +CUresult cuMemGetAllocationGranularity(size_t*, const CUmemAllocationProp*, CUmemAllocationGranularity_flags); +CUresult cuMemGetAllocationPropertiesFromHandle(CUmemAllocationProp*, CUmemGenericAllocationHandle); +CUresult cuMemRetainAllocationHandle(CUmemGenericAllocationHandle*, void*); +CUresult cuMemFreeAsync(CUdeviceptr, CUstream); +CUresult cuMemAllocAsync(CUdeviceptr*, size_t, CUstream); +CUresult cuMemPoolTrimTo(CUmemoryPool, size_t); +CUresult cuMemPoolSetAttribute(CUmemoryPool, CUmemPool_attribute, void*); +CUresult cuMemPoolGetAttribute(CUmemoryPool, CUmemPool_attribute, void*); +CUresult cuMemPoolSetAccess(CUmemoryPool, const CUmemAccessDesc*, size_t); +CUresult cuMemPoolGetAccess(CUmemAccess_flags*, CUmemoryPool, CUmemLocation*); +CUresult cuMemPoolCreate(CUmemoryPool*, const CUmemPoolProps*); +CUresult cuMemPoolDestroy(CUmemoryPool); +CUresult cuMemAllocFromPoolAsync(CUdeviceptr*, size_t, CUmemoryPool, CUstream); +CUresult cuMemPoolExportToShareableHandle(void*, CUmemoryPool, CUmemAllocationHandleType, unsigned long long); +CUresult cuMemPoolImportFromShareableHandle(CUmemoryPool*, void*, CUmemAllocationHandleType, unsigned long long); +CUresult cuMemPoolExportPointer(CUmemPoolPtrExportData*, CUdeviceptr); +CUresult cuMemPoolImportPointer(CUdeviceptr*, CUmemoryPool, CUmemPoolPtrExportData*); +CUresult cuMulticastCreate(CUmemGenericAllocationHandle*, const CUmulticastObjectProp*); +CUresult cuMulticastAddDevice(CUmemGenericAllocationHandle, CUdevice); +CUresult cuMulticastBindMem(CUmemGenericAllocationHandle, size_t, CUmemGenericAllocationHandle, size_t, size_t, unsigned long long); +CUresult cuMulticastBindAddr(CUmemGenericAllocationHandle, size_t, CUdeviceptr, size_t, unsigned long long); +CUresult cuMulticastUnbind(CUmemGenericAllocationHandle, CUdevice, size_t, size_t); +CUresult cuMulticastGetGranularity(size_t*, const CUmulticastObjectProp*, CUmulticastGranularity_flags); +CUresult cuPointerGetAttribute(void*, CUpointer_attribute, CUdeviceptr); +CUresult cuMemPrefetchAsync(CUdeviceptr, size_t, CUdevice, CUstream); +CUresult cuMemPrefetchAsync_v2(CUdeviceptr, size_t, CUmemLocation, unsigned int, CUstream); +CUresult cuMemAdvise(CUdeviceptr, size_t, CUmem_advise, CUdevice); +CUresult cuMemAdvise_v2(CUdeviceptr, size_t, CUmem_advise, CUmemLocation); +CUresult cuMemRangeGetAttribute(void*, size_t, CUmem_range_attribute, CUdeviceptr, size_t); +CUresult cuMemRangeGetAttributes(void**, size_t*, CUmem_range_attribute*, size_t, CUdeviceptr, size_t); +CUresult cuPointerSetAttribute(const void*, CUpointer_attribute, CUdeviceptr); +CUresult cuPointerGetAttributes(unsigned int, CUpointer_attribute*, void**, CUdeviceptr); +CUresult cuStreamCreate(CUstream*, unsigned int); +CUresult cuStreamCreateWithPriority(CUstream*, unsigned int, int); +CUresult cuStreamGetPriority(CUstream, int*); +CUresult cuStreamGetFlags(CUstream, unsigned int*); +CUresult cuStreamGetId(CUstream, unsigned long long*); +CUresult cuStreamGetCtx(CUstream, CUcontext*); +CUresult cuStreamWaitEvent(CUstream, CUevent, unsigned int); +// CUresult cuStreamAddCallback(CUstream, CUstreamCallback, void *, unsigned int); +CUresult cuStreamBeginCapture_v2(CUstream, CUstreamCaptureMode); +CUresult cuStreamBeginCaptureToGraph(CUstream, CUgraph, const CUgraphNode*, const CUgraphEdgeData*, size_t, CUstreamCaptureMode); +CUresult cuThreadExchangeStreamCaptureMode(CUstreamCaptureMode*); +CUresult cuStreamEndCapture(CUstream, CUgraph*); +CUresult cuStreamIsCapturing(CUstream, CUstreamCaptureStatus*); +// CUresult cuStreamGetCaptureInfo_v2(CUstream, CUstreamCaptureStatus *, cuuint64_t *, CUgraph *, const CUgraphNode **, size_t *); +// CUresult cuStreamGetCaptureInfo_v3(CUstream, CUstreamCaptureStatus *, cuuint64_t *, CUgraph *, const CUgraphNode **, const CUgraphEdgeData **, size_t *); +CUresult cuStreamUpdateCaptureDependencies(CUstream, CUgraphNode*, size_t, unsigned int); +CUresult cuStreamUpdateCaptureDependencies_v2(CUstream, CUgraphNode*, const CUgraphEdgeData*, size_t, unsigned int); +CUresult cuStreamAttachMemAsync(CUstream, CUdeviceptr, size_t, unsigned int); +CUresult cuStreamQuery(CUstream); +CUresult cuStreamSynchronize(CUstream); +CUresult cuStreamDestroy_v2(CUstream); +CUresult cuStreamCopyAttributes(CUstream, CUstream); +CUresult cuStreamGetAttribute(CUstream, CUstreamAttrID, CUstreamAttrValue*); +CUresult cuStreamSetAttribute(CUstream, CUstreamAttrID, const CUstreamAttrValue*); +CUresult cuEventCreate(CUevent*, unsigned int); +CUresult cuEventRecord(CUevent, CUstream); +CUresult cuEventRecordWithFlags(CUevent, CUstream, unsigned int); +CUresult cuEventQuery(CUevent); +CUresult cuEventSynchronize(CUevent); +CUresult cuEventDestroy_v2(CUevent); +CUresult cuEventElapsedTime(float*, CUevent, CUevent); +CUresult cuImportExternalMemory(CUexternalMemory*, const CUDA_EXTERNAL_MEMORY_HANDLE_DESC*); +CUresult cuExternalMemoryGetMappedBuffer(CUdeviceptr*, CUexternalMemory, const CUDA_EXTERNAL_MEMORY_BUFFER_DESC*); +CUresult cuExternalMemoryGetMappedMipmappedArray(CUmipmappedArray*, CUexternalMemory, const CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC*); +CUresult cuDestroyExternalMemory(CUexternalMemory); +CUresult cuImportExternalSemaphore(CUexternalSemaphore*, const CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC*); +CUresult cuSignalExternalSemaphoresAsync(const CUexternalSemaphore*, const CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS*, unsigned int, CUstream); +CUresult cuWaitExternalSemaphoresAsync(const CUexternalSemaphore*, const CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS*, unsigned int, CUstream); +CUresult cuDestroyExternalSemaphore(CUexternalSemaphore); +CUresult cuStreamWaitValue32_v2(CUstream, CUdeviceptr, cuuint32_t, unsigned int); +CUresult cuStreamWaitValue64_v2(CUstream, CUdeviceptr, cuuint64_t, unsigned int); +CUresult cuStreamWriteValue32_v2(CUstream, CUdeviceptr, cuuint32_t, unsigned int); +CUresult cuStreamWriteValue64_v2(CUstream, CUdeviceptr, cuuint64_t, unsigned int); +CUresult cuStreamBatchMemOp_v2(CUstream, unsigned int, CUstreamBatchMemOpParams*, unsigned int); +CUresult cuFuncGetAttribute(int*, CUfunction_attribute, CUfunction); +CUresult cuFuncSetAttribute(CUfunction, CUfunction_attribute, int); +CUresult cuFuncSetCacheConfig(CUfunction, CUfunc_cache); +CUresult cuFuncGetModule(CUmodule*, CUfunction); +CUresult cuFuncGetName(const char**, CUfunction); +CUresult cuFuncGetParamInfo(CUfunction, size_t, size_t*, size_t*); +CUresult cuFuncIsLoaded(CUfunctionLoadingState*, CUfunction); +CUresult cuFuncLoad(CUfunction); +CUresult cuLaunchKernel(CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, + CUstream, void**, void**); +CUresult cuLaunchKernelEx(const CUlaunchConfig*, CUfunction, void**, void**); +CUresult cuLaunchCooperativeKernel(CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, + unsigned int, CUstream, void**); +CUresult cuLaunchCooperativeKernelMultiDevice(CUDA_LAUNCH_PARAMS*, unsigned int, unsigned int); +CUresult cuLaunchHostFunc(CUstream, CUhostFn, void*); +CUresult cuFuncSetBlockShape(CUfunction, int, int, int); +CUresult cuFuncSetSharedSize(CUfunction, unsigned int); +CUresult cuParamSetSize(CUfunction, unsigned int); +CUresult cuParamSeti(CUfunction, int, unsigned int); +CUresult cuParamSetf(CUfunction, int, float); +CUresult cuParamSetv(CUfunction, int, void*, unsigned int); +CUresult cuLaunch(CUfunction); +CUresult cuLaunchGrid(CUfunction, int, int); +CUresult cuLaunchGridAsync(CUfunction, int, int, CUstream); +CUresult cuParamSetTexRef(CUfunction, int, CUtexref); +CUresult cuFuncSetSharedMemConfig(CUfunction, CUsharedconfig); +CUresult cuGraphCreate(CUgraph*, unsigned int); +CUresult cuGraphAddKernelNode_v2(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_KERNEL_NODE_PARAMS*); +CUresult cuGraphKernelNodeGetParams_v2(CUgraphNode, CUDA_KERNEL_NODE_PARAMS*); +CUresult cuGraphKernelNodeSetParams_v2(CUgraphNode, const CUDA_KERNEL_NODE_PARAMS*); +CUresult cuGraphAddMemcpyNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_MEMCPY3D*, CUcontext); +CUresult cuGraphMemcpyNodeGetParams(CUgraphNode, CUDA_MEMCPY3D*); +CUresult cuGraphMemcpyNodeSetParams(CUgraphNode, const CUDA_MEMCPY3D*); +CUresult cuGraphAddMemsetNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_MEMSET_NODE_PARAMS*, CUcontext); +CUresult cuGraphMemsetNodeGetParams(CUgraphNode, CUDA_MEMSET_NODE_PARAMS*); +CUresult cuGraphMemsetNodeSetParams(CUgraphNode, const CUDA_MEMSET_NODE_PARAMS*); +CUresult cuGraphAddHostNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_HOST_NODE_PARAMS*); +CUresult cuGraphHostNodeGetParams(CUgraphNode, CUDA_HOST_NODE_PARAMS*); +CUresult cuGraphHostNodeSetParams(CUgraphNode, const CUDA_HOST_NODE_PARAMS*); +CUresult cuGraphAddChildGraphNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUgraph); +CUresult cuGraphChildGraphNodeGetGraph(CUgraphNode, CUgraph*); +CUresult cuGraphAddEmptyNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t); +CUresult cuGraphAddEventRecordNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUevent); +CUresult cuGraphEventRecordNodeGetEvent(CUgraphNode, CUevent*); +CUresult cuGraphEventRecordNodeSetEvent(CUgraphNode, CUevent); +CUresult cuGraphAddEventWaitNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUevent); +CUresult cuGraphEventWaitNodeGetEvent(CUgraphNode, CUevent*); +CUresult cuGraphEventWaitNodeSetEvent(CUgraphNode, CUevent); +CUresult cuGraphAddExternalSemaphoresSignalNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_EXT_SEM_SIGNAL_NODE_PARAMS*); +CUresult cuGraphExternalSemaphoresSignalNodeGetParams(CUgraphNode, CUDA_EXT_SEM_SIGNAL_NODE_PARAMS*); +CUresult cuGraphExternalSemaphoresSignalNodeSetParams(CUgraphNode, const CUDA_EXT_SEM_SIGNAL_NODE_PARAMS*); +CUresult cuGraphAddExternalSemaphoresWaitNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_EXT_SEM_WAIT_NODE_PARAMS*); +CUresult cuGraphExternalSemaphoresWaitNodeGetParams(CUgraphNode, CUDA_EXT_SEM_WAIT_NODE_PARAMS*); +CUresult cuGraphExternalSemaphoresWaitNodeSetParams(CUgraphNode, const CUDA_EXT_SEM_WAIT_NODE_PARAMS*); +CUresult cuGraphAddBatchMemOpNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, const CUDA_BATCH_MEM_OP_NODE_PARAMS*); +CUresult cuGraphBatchMemOpNodeGetParams(CUgraphNode, CUDA_BATCH_MEM_OP_NODE_PARAMS*); +CUresult cuGraphBatchMemOpNodeSetParams(CUgraphNode, const CUDA_BATCH_MEM_OP_NODE_PARAMS*); +CUresult cuGraphExecBatchMemOpNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_BATCH_MEM_OP_NODE_PARAMS*); +CUresult cuGraphAddMemAllocNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUDA_MEM_ALLOC_NODE_PARAMS*); +CUresult cuGraphMemAllocNodeGetParams(CUgraphNode, CUDA_MEM_ALLOC_NODE_PARAMS*); +CUresult cuGraphAddMemFreeNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUdeviceptr); +CUresult cuGraphMemFreeNodeGetParams(CUgraphNode, CUdeviceptr*); +CUresult cuDeviceGraphMemTrim(CUdevice); +CUresult cuDeviceGetGraphMemAttribute(CUdevice, CUgraphMem_attribute, void*); +CUresult cuDeviceSetGraphMemAttribute(CUdevice, CUgraphMem_attribute, void*); +CUresult cuGraphClone(CUgraph*, CUgraph); +CUresult cuGraphNodeFindInClone(CUgraphNode*, CUgraphNode, CUgraph); +CUresult cuGraphNodeGetType(CUgraphNode, CUgraphNodeType*); +CUresult cuGraphGetNodes(CUgraph, CUgraphNode*, size_t*); +CUresult cuGraphGetRootNodes(CUgraph, CUgraphNode*, size_t*); +CUresult cuGraphGetEdges(CUgraph, CUgraphNode*, CUgraphNode*, size_t*); +CUresult cuGraphGetEdges_v2(CUgraph, CUgraphNode*, CUgraphNode*, CUgraphEdgeData*, size_t*); +CUresult cuGraphNodeGetDependencies(CUgraphNode, CUgraphNode*, size_t*); +CUresult cuGraphNodeGetDependencies_v2(CUgraphNode, CUgraphNode*, CUgraphEdgeData*, size_t*); +CUresult cuGraphNodeGetDependentNodes(CUgraphNode, CUgraphNode*, size_t*); +CUresult cuGraphNodeGetDependentNodes_v2(CUgraphNode, CUgraphNode*, CUgraphEdgeData*, size_t*); +CUresult cuGraphAddDependencies(CUgraph, const CUgraphNode*, const CUgraphNode*, size_t); +CUresult cuGraphAddDependencies_v2(CUgraph, const CUgraphNode*, const CUgraphNode*, const CUgraphEdgeData*, size_t); +CUresult cuGraphRemoveDependencies(CUgraph, const CUgraphNode*, const CUgraphNode*, size_t); +CUresult cuGraphRemoveDependencies_v2(CUgraph, const CUgraphNode*, const CUgraphNode*, const CUgraphEdgeData*, size_t); +CUresult cuGraphDestroyNode(CUgraphNode); +CUresult cuGraphInstantiateWithFlags(CUgraphExec*, CUgraph, unsigned long long); +CUresult cuGraphInstantiateWithParams(CUgraphExec*, CUgraph, CUDA_GRAPH_INSTANTIATE_PARAMS*); +CUresult cuGraphExecGetFlags(CUgraphExec, cuuint64_t*); +CUresult cuGraphExecKernelNodeSetParams_v2(CUgraphExec, CUgraphNode, const CUDA_KERNEL_NODE_PARAMS*); +CUresult cuGraphExecMemcpyNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_MEMCPY3D*, CUcontext); +CUresult cuGraphExecMemsetNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_MEMSET_NODE_PARAMS*, CUcontext); +CUresult cuGraphExecHostNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_HOST_NODE_PARAMS*); +CUresult cuGraphExecChildGraphNodeSetParams(CUgraphExec, CUgraphNode, CUgraph); +CUresult cuGraphExecEventRecordNodeSetEvent(CUgraphExec, CUgraphNode, CUevent); +CUresult cuGraphExecEventWaitNodeSetEvent(CUgraphExec, CUgraphNode, CUevent); +CUresult cuGraphExecExternalSemaphoresSignalNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_EXT_SEM_SIGNAL_NODE_PARAMS*); +CUresult cuGraphExecExternalSemaphoresWaitNodeSetParams(CUgraphExec, CUgraphNode, const CUDA_EXT_SEM_WAIT_NODE_PARAMS*); +CUresult cuGraphNodeSetEnabled(CUgraphExec, CUgraphNode, unsigned int); +CUresult cuGraphNodeGetEnabled(CUgraphExec, CUgraphNode, unsigned int*); +CUresult cuGraphUpload(CUgraphExec, CUstream); +CUresult cuGraphLaunch(CUgraphExec, CUstream); +CUresult cuGraphExecDestroy(CUgraphExec); +CUresult cuGraphDestroy(CUgraph); +CUresult cuGraphExecUpdate_v2(CUgraphExec, CUgraph, CUgraphExecUpdateResultInfo*); +CUresult cuGraphKernelNodeCopyAttributes(CUgraphNode, CUgraphNode); +CUresult cuGraphKernelNodeGetAttribute(CUgraphNode, CUkernelNodeAttrID, CUkernelNodeAttrValue*); +CUresult cuGraphKernelNodeSetAttribute(CUgraphNode, CUkernelNodeAttrID, const CUkernelNodeAttrValue*); +CUresult cuGraphDebugDotPrint(CUgraph, const char*, unsigned int); +CUresult cuUserObjectCreate(CUuserObject*, void*, CUhostFn, unsigned int, unsigned int); +CUresult cuUserObjectRetain(CUuserObject, unsigned int); +CUresult cuUserObjectRelease(CUuserObject, unsigned int); +CUresult cuGraphRetainUserObject(CUgraph, CUuserObject, unsigned int, unsigned int); +CUresult cuGraphReleaseUserObject(CUgraph, CUuserObject, unsigned int); +CUresult cuGraphAddNode(CUgraphNode*, CUgraph, const CUgraphNode*, size_t, CUgraphNodeParams*); +CUresult cuGraphAddNode_v2(CUgraphNode*, CUgraph, const CUgraphNode*, const CUgraphEdgeData*, size_t, CUgraphNodeParams*); +CUresult cuGraphNodeSetParams(CUgraphNode, CUgraphNodeParams*); +CUresult cuGraphExecNodeSetParams(CUgraphExec, CUgraphNode, CUgraphNodeParams*); +CUresult cuGraphConditionalHandleCreate(CUgraphConditionalHandle*, CUgraph, CUcontext, unsigned int, unsigned int); +CUresult cuOccupancyMaxActiveBlocksPerMultiprocessor(int*, CUfunction, int, size_t); +CUresult cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int*, CUfunction, int, size_t, unsigned int); +// CUresult cuOccupancyMaxPotentialBlockSize(int *, int *, CUfunction, CUoccupancyB2DSize, size_t, int); +// CUresult cuOccupancyMaxPotentialBlockSizeWithFlags(int *, int *, CUfunction, CUoccupancyB2DSize, size_t, int, unsigned int); +CUresult cuOccupancyAvailableDynamicSMemPerBlock(size_t*, CUfunction, int, int); +CUresult cuOccupancyMaxPotentialClusterSize(int*, CUfunction, const CUlaunchConfig*); +CUresult cuOccupancyMaxActiveClusters(int*, CUfunction, const CUlaunchConfig*); +CUresult cuTexRefSetArray(CUtexref, CUarray, unsigned int); +CUresult cuTexRefSetMipmappedArray(CUtexref, CUmipmappedArray, unsigned int); +CUresult cuTexRefSetAddress_v2(size_t*, CUtexref, CUdeviceptr, size_t); +CUresult cuTexRefSetAddress2D_v3(CUtexref, const CUDA_ARRAY_DESCRIPTOR*, CUdeviceptr, size_t); +CUresult cuTexRefSetFormat(CUtexref, CUarray_format, int); +CUresult cuTexRefSetAddressMode(CUtexref, int, CUaddress_mode); +CUresult cuTexRefSetFilterMode(CUtexref, CUfilter_mode); +CUresult cuTexRefSetMipmapFilterMode(CUtexref, CUfilter_mode); +CUresult cuTexRefSetMipmapLevelBias(CUtexref, float); +CUresult cuTexRefSetMipmapLevelClamp(CUtexref, float, float); +CUresult cuTexRefSetMaxAnisotropy(CUtexref, unsigned int); +CUresult cuTexRefSetBorderColor(CUtexref, float*); +CUresult cuTexRefSetFlags(CUtexref, unsigned int); +CUresult cuTexRefGetAddress_v2(CUdeviceptr*, CUtexref); +CUresult cuTexRefGetArray(CUarray*, CUtexref); +CUresult cuTexRefGetMipmappedArray(CUmipmappedArray*, CUtexref); +CUresult cuTexRefGetAddressMode(CUaddress_mode*, CUtexref, int); +CUresult cuTexRefGetFilterMode(CUfilter_mode*, CUtexref); +CUresult cuTexRefGetFormat(CUarray_format*, int*, CUtexref); +CUresult cuTexRefGetMipmapFilterMode(CUfilter_mode*, CUtexref); +CUresult cuTexRefGetMipmapLevelBias(float*, CUtexref); +CUresult cuTexRefGetMipmapLevelClamp(float*, float*, CUtexref); +CUresult cuTexRefGetMaxAnisotropy(int*, CUtexref); +CUresult cuTexRefGetBorderColor(float*, CUtexref); +CUresult cuTexRefGetFlags(unsigned int*, CUtexref); +CUresult cuTexRefCreate(CUtexref*); +CUresult cuTexRefDestroy(CUtexref); +CUresult cuSurfRefSetArray(CUsurfref, CUarray, unsigned int); +CUresult cuSurfRefGetArray(CUarray*, CUsurfref); +CUresult cuTexObjectCreate(CUtexObject*, const CUDA_RESOURCE_DESC*, const CUDA_TEXTURE_DESC*, const CUDA_RESOURCE_VIEW_DESC*); +CUresult cuTexObjectDestroy(CUtexObject); +CUresult cuTexObjectGetResourceDesc(CUDA_RESOURCE_DESC*, CUtexObject); +CUresult cuTexObjectGetTextureDesc(CUDA_TEXTURE_DESC*, CUtexObject); +CUresult cuTexObjectGetResourceViewDesc(CUDA_RESOURCE_VIEW_DESC*, CUtexObject); +CUresult cuSurfObjectCreate(CUsurfObject*, const CUDA_RESOURCE_DESC*); +CUresult cuSurfObjectDestroy(CUsurfObject); +CUresult cuSurfObjectGetResourceDesc(CUDA_RESOURCE_DESC*, CUsurfObject); +CUresult cuTensorMapEncodeTiled(CUtensorMap*, CUtensorMapDataType, cuuint32_t, void*, const cuuint64_t*, const cuuint64_t*, const cuuint32_t*, + const cuuint32_t*, CUtensorMapInterleave, CUtensorMapSwizzle, CUtensorMapL2promotion, CUtensorMapFloatOOBfill); +// CUresult cuTensorMapEncodeIm2col(CUtensorMap *, CUtensorMapDataType, cuuint32_t, void *, const cuuint64_t *, const cuuint64_t *, const int +// *, const int *, cuuint32_t, cuuint32_t, const cuuint32_t *, CUtensorMapInterleave, CUtensorMapSwizzle, CUtensorMapL2promotion, CUtensorMapFloatOOBfill); +CUresult cuTensorMapReplaceAddress(CUtensorMap*, void*); +CUresult cuDeviceCanAccessPeer(int*, CUdevice, CUdevice); +CUresult cuCtxEnablePeerAccess(CUcontext, unsigned int); +CUresult cuCtxDisablePeerAccess(CUcontext); +CUresult cuDeviceGetP2PAttribute(int*, CUdevice_P2PAttribute, CUdevice, CUdevice); +CUresult cuGraphicsUnregisterResource(CUgraphicsResource); +CUresult cuGraphicsSubResourceGetMappedArray(CUarray*, CUgraphicsResource, unsigned int, unsigned int); +CUresult cuGraphicsResourceGetMappedMipmappedArray(CUmipmappedArray*, CUgraphicsResource); +CUresult cuGraphicsResourceGetMappedPointer_v2(CUdeviceptr*, size_t*, CUgraphicsResource); +CUresult cuGraphicsResourceSetMapFlags_v2(CUgraphicsResource, unsigned int); +CUresult cuGraphicsMapResources(unsigned int, CUgraphicsResource*, CUstream); +CUresult cuGraphicsUnmapResources(unsigned int, CUgraphicsResource*, CUstream); +CUresult cuGetProcAddress_v2(const char*, void**, int, cuuint64_t, CUdriverProcAddressQueryResult*); +CUresult cuCoredumpGetAttribute(CUcoredumpSettings, void*, size_t*); +CUresult cuCoredumpGetAttributeGlobal(CUcoredumpSettings, void*, size_t*); +CUresult cuCoredumpSetAttribute(CUcoredumpSettings, void*, size_t*); +CUresult cuCoredumpSetAttributeGlobal(CUcoredumpSettings, void*, size_t*); +CUresult cuGetExportTable(const void**, const CUuuid*); +CUresult cuGreenCtxCreate(CUgreenCtx*, CUdevResourceDesc, CUdevice, unsigned int); +CUresult cuGreenCtxDestroy(CUgreenCtx); +CUresult cuCtxFromGreenCtx(CUcontext*, CUgreenCtx); +CUresult cuDeviceGetDevResource(CUdevice, CUdevResource*, CUdevResourceType); +CUresult cuCtxGetDevResource(CUcontext, CUdevResource*, CUdevResourceType); +CUresult cuGreenCtxGetDevResource(CUgreenCtx, CUdevResource*, CUdevResourceType); +CUresult cuDevSmResourceSplitByCount(CUdevResource*, unsigned int*, const CUdevResource*, CUdevResource*, unsigned int, unsigned int); +CUresult cuDevResourceGenerateDesc(CUdevResourceDesc*, CUdevResource*, unsigned int); +CUresult cuGreenCtxRecordEvent(CUgreenCtx, CUevent); +CUresult cuGreenCtxWaitEvent(CUgreenCtx, CUevent); +CUresult cuStreamGetGreenCtx(CUstream, CUgreenCtx*); +#ifdef __cplusplus +} +#endif diff --git a/ThunkLibs/libcuda/fixup_defines.h b/ThunkLibs/libcuda/fixup_defines.h new file mode 100644 index 000000000..cc8a94990 --- /dev/null +++ b/ThunkLibs/libcuda/fixup_defines.h @@ -0,0 +1,15 @@ +// SPDX-License-Identifier: MIT +#pragma once +#include +#ifdef __cplusplus +extern "C" { +#endif + +typedef void (*CUhostFn)(void*); +enum cudaError_enum {}; +enum CUgreenCtxCreate_flags {}; +enum CUdevResourceType {}; + +#ifdef __cplusplus +} +#endif diff --git a/ThunkLibs/libcuda/libcuda_Guest.cpp b/ThunkLibs/libcuda/libcuda_Guest.cpp new file mode 100644 index 000000000..beac01ed9 --- /dev/null +++ b/ThunkLibs/libcuda/libcuda_Guest.cpp @@ -0,0 +1,84 @@ +// SPDX-License-Identifier: MIT +#include "common/Guest.h" +#include "cuda_defines.h" + +#include "thunkgen_guest_libcuda.inl" +#include +#include +#include +#include +#include + +#include + +// Maps cuda API function names to the address of a guest function which is +// linked to the corresponding host function pointer +const std::unordered_map HostPtrInvokers = std::invoke([]() { +#define PAIR(name, unused) Ret[#name] = reinterpret_cast(GetCallerForHostFunction(name)); + std::unordered_map Ret; + FOREACH_internal_SYMBOL(PAIR); + return Ret; +#undef PAIR +}); + +extern "C" { + +// This variable controls the behavior of cuGetProcAddress for functions we don't know the signature of: +// - if false (default), we return a nullptr (since the application might have a fallback code path) +// - if true, we return a stub function that fatally errors upon being called +constexpr bool stub_unknown_functions = false; + +// Fatally erroring function with a thunk-like interface. This is used as a placeholder for unknown CUDA functions +[[noreturn]] +static void FatalError(void* raw_args) { + auto called_function = reinterpret_cast*>(raw_args)->a0; + fprintf(stderr, "FATAL: Called unknown CUDA function at address %p\n", reinterpret_cast(called_function)); + __builtin_trap(); +} + +static void* MakeGuestCallable(const char* origin, void* func, const char* name) { + auto It = HostPtrInvokers.find(name); + if (It == HostPtrInvokers.end()) { + fprintf(stderr, "%s: Unknown cuda function at address %p: %s\n", origin, func, name); + if (stub_unknown_functions) { + const auto StubHostPtrInvoker = CallHostFunction; + LinkAddressToFunction((uintptr_t)func, reinterpret_cast(StubHostPtrInvoker)); + return func; + } + return nullptr; + } + LinkAddressToFunction((uintptr_t)func, It->second); + return func; +} + +struct override_entry { + std::string_view name; + void* ptr; +}; + +constexpr static std::array proc_override = { + {{"cuGetProcAddress", (void*)cuGetProcAddress_v2}, {"cuGetProcAddress_v2", (void*)cuGetProcAddress_v2}}}; + +CUresult cuGetProcAddress_v2(const char* symbol, void** pfn, int cudaVersion, cuuint64_t flags, CUdriverProcAddressQueryResult* symbolStatus) { + + for (auto& over : proc_override) { + if (symbol == over.name) { + *pfn = over.ptr; + if (symbolStatus) { + *symbolStatus = (CUdriverProcAddressQueryResult)0; + } + return (CUresult)0; // CUDA_SUCCESS + } + } + + void* ptr {}; + auto Ret = fexfn_pack_cuGetProcAddress_v2(symbol, &ptr, cudaVersion, flags, symbolStatus); + + if (!Ret) { + *pfn = MakeGuestCallable(__FUNCTION__, ptr, symbol); + } + return Ret; +} +} + +LOAD_LIB(libcuda) diff --git a/ThunkLibs/libcuda/libcuda_Host.cpp b/ThunkLibs/libcuda/libcuda_Host.cpp new file mode 100644 index 000000000..2ba10f675 --- /dev/null +++ b/ThunkLibs/libcuda/libcuda_Host.cpp @@ -0,0 +1,81 @@ +// SPDX-License-Identifier: MIT +#include "common/Host.h" +#include "cuda_defines.h" + +#include +#include +#include +#include +#include +#include "thunkgen_host_libcuda.inl" + +#define FEXFN_IMPL(fn) fexfn_impl_libcuda_##fn +#define LDR_PTR(fn) fexldr_ptr_libcuda_##fn + +struct override_entry { + std::string_view name; + void* ptr; +}; + +constexpr static std::array proc_override = { + {{"cuCtxCreate_v2", (void*)FEXFN_IMPL(cuCtxCreate_v2)}, {"cuGetExportTable", (void*)FEXFN_IMPL(cuGetExportTable)}}}; + +static CUresult FEXFN_IMPL(cuGetProcAddress_v2)(const char* symbol, guest_layout pfn, int cudaVersion, cuuint64_t flags, + CUdriverProcAddressQueryResult* symbolStatus) { + host_layout host_data {}; + void* ptr {}; + CUresult ret {}; + for (auto& over : proc_override) { + if (symbol == over.name) { + ptr = over.ptr; + if (symbolStatus) { + *symbolStatus = (CUdriverProcAddressQueryResult)0; // CU_GET_PROC_ADDRESS_SUCCESS + } + ret = (CUresult)0; // CUDA_SUCCESS + break; + } + } + + if (!ptr) { + ret = LDR_PTR(cuGetProcAddress_v2)(symbol, &ptr, cudaVersion, flags, symbolStatus); + } + + host_data.data = ptr; + *pfn.get_pointer() = to_guest(host_data); + return ret; +} + +CUresult FEXFN_IMPL(cuCtxCreate_v2)(guest_layout pctx, unsigned int flags, CUdevice dev) { + host_layout host_data {}; + CUcontext ctx; + auto ret = LDR_PTR(cuCtxCreate_v2)(&ctx, flags, dev); + host_data.data = ctx; + *pctx.get_pointer() = to_guest(host_data); + return ret; +} + +CUresult FEXFN_IMPL(cuGetExportTable)(guest_layout ppExportTable, const CUuuid* pExportTableId) { + // This function returns a pointer to an driver internal export table that is undocumented publicly. + // + // Some documentation about the UUIDs that has been reversed. + // https://github.com/vosen/ZLUDA/blob/1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf/zluda_dark_api/src/lib.rs#L197 + // + // These two UUIDs are required by the static cuda runtime at startup. Might require more but it currently halts without these + // implemented. UUID: [0x6b, 0xd5, 0xfb, 0x6c, 0x5b, 0xf4, 0xe7, 0x4a, 0x89, 0x87, 0xd9, 0x39, 0x12, 0xfd, 0x9d, 0xf9] + // - cudart interface (0x68 bytes?) + // UUID: [0xa0, 0x94, 0x79, 0x8c, 0x2e, 0x74, 0x2e, 0x74, 0x93, 0xf2, 0x08, 0x00, 0x20, 0x0c, 0x0a, 0x66] + // - Tools runtime callback hooks (0x38 bytes?) + + std::string uuid {}; + for (size_t i = 0; i < sizeof(CUuuid); ++i) { + const bool last = (i + 1) == sizeof(CUuuid); + char tmp[8]; + auto size = snprintf(tmp, 8, "0x%02x%s", pExportTableId->_0[i], last ? "" : ", "); + uuid += std::string_view(tmp, size); + } + fprintf(stderr, "cuGetExportTable not implemented\n"); + fprintf(stderr, "UUID: %s\n", uuid.c_str()); + return (CUresult)4; // CUDA_ERROR_DEINITIALIZED +} + +EXPORTS(libcuda) diff --git a/ThunkLibs/libcuda/libcuda_interface.cpp b/ThunkLibs/libcuda/libcuda_interface.cpp new file mode 100644 index 000000000..57d4c67c4 --- /dev/null +++ b/ThunkLibs/libcuda/libcuda_interface.cpp @@ -0,0 +1,1571 @@ +// SPDX-License-Identifier: MIT +// Generated file. Be careful if modifying +#include +#include "cuda_defines.h" +template +struct fex_gen_config { + unsigned version = 1; +}; + +template +struct fex_gen_type {}; + +// Type definitions +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; +template<> +struct fex_gen_type {}; +template<> +struct fex_gen_type : fexgen::opaque_type {}; + +// Function definitions +namespace internal { +template +struct fex_gen_config : fexgen::generate_guest_symtable, fexgen::indirect_guest_calls {}; + +// Function, parameter index, parameter type [optional] +template +struct fex_gen_param {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config : fexgen::custom_host_impl {}; +template<> +struct fex_gen_param : fexgen::ptr_passthrough {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +// Disabled due to callbacks +// template<> +// struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_param : fexgen::assume_compatible_data_layout {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +// Disabled due to callbacks +// template<> +// struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +// template<> +// struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +// template<> +// struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +// template<> +// struct fex_gen_config {}; +// template<> +// struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +// template<> +// struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config : fexgen::custom_host_impl, fexgen::custom_guest_entrypoint, fexgen::returns_guest_pointer {}; +template<> +struct fex_gen_param : fexgen::ptr_passthrough {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config : fexgen::custom_host_impl {}; +template<> +struct fex_gen_param : fexgen::ptr_passthrough {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +template<> +struct fex_gen_config {}; +template<> +struct fex_gen_config {}; +#ifndef IS_32BIT_THUNK +template<> +struct fex_gen_config {}; +#endif +} // namespace internal