diff options
Diffstat (limited to 'userspace/cuda')
18 files changed, 1807 insertions, 32 deletions
diff --git a/userspace/cuda/570.144/annotations/README.md b/userspace/cuda/570.144/annotations/README.md index 868c503..85c14af 100644 --- a/userspace/cuda/570.144/annotations/README.md +++ b/userspace/cuda/570.144/annotations/README.md @@ -8,8 +8,9 @@ accumulates across driver versions. | File | Format | Purpose | |---|---|---| -| `names.tsv` | `file_vaddr<TAB>name<TAB>source` | symbol map to apply to a Ghidra project (or any address-keyed analysis). `source` is `export`, `wrapper`, `verified`, or `todo`. | +| `names.tsv` | `file_vaddr<TAB>name<TAB>source` | symbol map to apply to a Ghidra project (or any address-keyed analysis). `source` is `export`, `wrapper`, `wrapper-callee`, `verified`, or `todo`. | | `subsystems.tsv` | `file_vaddr<TAB>name<TAB>subsystem<TAB>role<TAB>evidence` | human-readable subsystem index: which function belongs to which layer and what it does, with the evidence for the claim. | +| `name-candidates.tsv` | `file_vaddr<TAB>proposed_name<TAB>confidence<TAB>evidence<TAB>status` | machine-proposed API-implementation names for wrapper callees (`tools/name_candidates.py`), with the confidence and the candidate sets behind each proposal. `status` is `unverified` until the body is read; only `verified` entries may be promoted to `names.tsv`. | Addresses are **file virtual addresses** (`libcuda+0x...`). Apply the map with `tools/ghidra/ApplyNameMap.java`; see `docs/06-tooling.md`. diff --git a/userspace/cuda/570.144/annotations/name-candidates.tsv b/userspace/cuda/570.144/annotations/name-candidates.tsv new file mode 100644 index 0000000..559633f --- /dev/null +++ b/userspace/cuda/570.144/annotations/name-candidates.tsv @@ -0,0 +1,334 @@ +# Open-IOV CUDA API-implementation candidates (our own work product; tracked). +# Generated by tools/name_candidates.py from the elf-index exports, +# callgraph and the version's names.tsv; see that script for the method. +# +# file_vaddr<TAB>proposed_name<TAB>confidence<TAB>evidence<TAB>status +# confidence: high = exactly one filtered candidate, shared by >=2 variants +# or a unique callee; medium = one survivor also called by another +# wrapper, or several survivors (best-ranked proposed). +# status: unverified (tool output). Verification = read the body in the +# decompiled corpus and mark verified/corrected here; only verified +# names may be copied into names.tsv. +# +0x001a36d0 impl_cuCtxGetApiVersion high group=cuCtxGetApiVersion variants=cuCtxGetApiVersion; before=0x1a36d0,0x2bb140,0x2bb420,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1a36d0; outside_wrappers=0 verified +0x001a3770 impl_cuCtxGetId medium group=cuCtxGetId variants=cuCtxGetId; before=0x1a3770,0x2bb070,0x2bb140,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1a3770,0x2bb070; outside_wrappers=0 unverified +0x001a4480 impl_cuArrayGetSparseProperties medium group=cuArrayGetSparseProperties variants=cuArrayGetSparseProperties; before=0x1a4480,0x257720,0x258750,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a4480,0x258750,0x257720; outside_wrappers=0 unverified +0x001a46d0 impl_cuMipmappedArrayGetSparseProperties medium group=cuMipmappedArrayGetSparseProperties variants=cuMipmappedArrayGetSparseProperties; before=0x1a46d0,0x257720,0x258750,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a46d0,0x258750,0x257720; outside_wrappers=0 unverified +0x001a4900 impl_cuArrayGetMemoryRequirements high group=cuArrayGetMemoryRequirements variants=cuArrayGetMemoryRequirements; before=0x1a4900,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a4900; outside_wrappers=0 verified +0x001a4ac0 impl_cuMipmappedArrayGetMemoryRequirements high group=cuMipmappedArrayGetMemoryRequirements variants=cuMipmappedArrayGetMemoryRequirements; before=0x1a4ac0,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a4ac0; outside_wrappers=0 unverified +0x001a5410 impl_cuArrayGetPlane high group=cuArrayGetPlane variants=cuArrayGetPlane; before=0x1a5410,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a5410; outside_wrappers=0 verified +0x001a5630 impl_cuMipmappedArrayGetLevel high group=cuMipmappedArrayGetLevel variants=cuMipmappedArrayGetLevel; before=0x1a5630,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a5630; outside_wrappers=0 unverified +0x001a5820 impl_cuDevicePrimaryCtxGetState high group=cuDevicePrimaryCtxGetState variants=cuDevicePrimaryCtxGetState; before=0x1792f0,0x179350,0x1a5820,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1a5820; outside_wrappers=0 verified +0x001a6150 impl_cuDeviceRegisterAsyncNotification high group=cuDeviceRegisterAsyncNotification variants=cuDeviceRegisterAsyncNotification; before=0x1792f0,0x179350,0x179a80,0x1a6150,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a6150; outside_wrappers=0 verified +0x001a62c0 impl_cuDeviceUnregisterAsyncNotification high group=cuDeviceUnregisterAsyncNotification variants=cuDeviceUnregisterAsyncNotification; before=0x179200,0x1792f0,0x179350,0x1a62c0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a62c0; outside_wrappers=0 unverified +0x001a6430 impl_cuCtxAttach high group=cuCtxAttach variants=cuCtxAttach; before=0x1a6430,0x2bb140,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a6430; outside_wrappers=0 verified +0x001a66f0 impl_cuCtxGetCurrent high group=cuCtxGetCurrent variants=cuCtxGetCurrent; before=0x1798f0,0x1a66f0,0x2fde50,0x2fe440,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a66f0; outside_wrappers=0 verified +0x001a6850 impl_cuCtxGetDevice high group=cuCtxGetDevice variants=cuCtxGetDevice; before=0x1a6850,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a6850; outside_wrappers=0 verified +0x001a6960 impl_cuCtxGetFlags high group=cuCtxGetFlags variants=cuCtxGetFlags; before=0x1a6960,0x2dc100,0x2fde50,0x2fe540,0x4a7e90,0xd4bd20,0xd4dc50; after=0x1a6960; outside_wrappers=0 verified +0x001a6af0 impl_cuFlushGPUDirectRDMAWrites high group=cuFlushGPUDirectRDMAWrites variants=cuFlushGPUDirectRDMAWrites; before=0x1a6af0,0x2baed0,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a6af0; outside_wrappers=0 verified +0x001a6c50 impl_cuCtxGetExecAffinity high group=cuCtxGetExecAffinity variants=cuCtxGetExecAffinity; before=0x1a6c50,0x2baed0,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a6c50; outside_wrappers=0 verified +0x001a7320 impl_cuDeviceGetLuid high group=cuDeviceGetLuid variants=cuDeviceGetLuid; before=0x1a7320,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1a7320; outside_wrappers=0 verified +0x001a78c0 impl_cuDeviceGetDevResource high group=cuDeviceGetDevResource variants=cuDeviceGetDevResource; before=0x1a78c0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1a78c0; outside_wrappers=0 unverified +0x001a7aa0 impl_cuModuleGetFunctionCount high group=cuModuleGetFunctionCount variants=cuModuleGetFunctionCount; before=0x1a7aa0,0x2baed0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a7aa0; outside_wrappers=0 unverified +0x001a7c40 impl_cuModuleEnumerateFunctions high group=cuModuleEnumerateFunctions variants=cuModuleEnumerateFunctions; before=0x1a7c40,0x2baed0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a7c40; outside_wrappers=0 unverified +0x001a8120 impl_cuArrayDestroy medium group=cuArrayDestroy variants=cuArrayDestroy; before=0x1792f0,0x179350,0x1a8120,0x2597f0,0x2bb140,0x2d8f40,0x2e5900,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50,0xd4f3e0,0xd51980; after=0x1a8120,0x2597f0,0x2e5900; outside_wrappers=0 unverified +0x001a82f0 impl_cuMipmappedArrayCreate high group=cuMipmappedArrayCreate variants=cuMipmappedArrayCreate; before=0x1a82f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1a82f0; outside_wrappers=0 unverified +0x001a8520 impl_cuMipmappedArrayDestroy medium group=cuMipmappedArrayDestroy variants=cuMipmappedArrayDestroy; before=0x1792f0,0x179350,0x1a8520,0x25a070,0x2bb140,0x2d8f40,0x2e5a50,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50,0xd4f3e0,0xd51980; after=0x1a8520,0x2e5a50,0x25a070; outside_wrappers=0 unverified +0x001aa6c0 impl_cuCtxGetCacheConfig high group=cuCtxGetCacheConfig variants=cuCtxGetCacheConfig; before=0x179cf0,0x1aa6c0,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1aa6c0; outside_wrappers=0 unverified +0x001aa880 impl_cuCtxGetSharedMemConfig high group=cuCtxGetSharedMemConfig variants=cuCtxGetSharedMemConfig; before=0x179cf0,0x1aa880,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1aa880; outside_wrappers=0 unverified +0x001aad70 impl_cuCheckpointProcessLock medium group=cuCheckpointProcessLock variants=cuCheckpointProcessLock; before=0x179200,0x1aad70,0x2623e0,0x2fe540,0x30a870,0x4a7e90,0xd19d70,0xd4dc50; after=0x1aad70,0xd19d70,0x2623e0,0x30a870; outside_wrappers=0 unverified +0x001aae30 impl_cuCheckpointProcessCheckpoint high group=cuCheckpointProcessCheckpoint variants=cuCheckpointProcessCheckpoint; before=0x1aae30,0x2fe540,0x4a7e90,0xd4dc50; after=0x1aae30; outside_wrappers=0 unverified +0x001aaee0 impl_cuCheckpointProcessRestore high group=cuCheckpointProcessRestore variants=cuCheckpointProcessRestore; before=0x1aaee0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1aaee0; outside_wrappers=0 unverified +0x001aaf90 impl_cuCheckpointProcessUnlock high group=cuCheckpointProcessUnlock variants=cuCheckpointProcessUnlock; before=0x1aaf90,0x2fe540,0x4a7e90,0xd4dc50; after=0x1aaf90; outside_wrappers=0 unverified +0x001ab040 impl_cuCoredumpGetAttribute medium group=cuCoredumpGetAttribute variants=cuCoredumpGetAttribute; before=0x179190,0x179470,0x1ab040,0x2bb420,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ab040,0x179190,0x179470; outside_wrappers=0 unverified +0x001ab1e0 impl_cuCoredumpGetAttributeGlobal medium group=cuCoredumpGetAttributeGlobal variants=cuCoredumpGetAttributeGlobal; before=0x179190,0x179470,0x1ab1e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ab1e0,0x179190,0x179470; outside_wrappers=0 unverified +0x001ac3f0 impl_cuCtxDetach high group=cuCtxDetach variants=cuCtxDetach; before=0x1ac3f0,0x26d8e0,0x2bb300,0x2fde50,0x2fe0f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ac3f0; outside_wrappers=0 unverified +0x001acc50 impl_cuCtxSetCurrent medium group=cuCtxSetCurrent variants=cuCtxSetCurrent; before=0x1acc50,0x24a4d0,0x24b930,0x24b9e0,0x2fde50,0x2fde90,0x2fe540,0x4a7e90,0xd4dc50; after=0x1acc50,0x2fde90; outside_wrappers=0 unverified +0x001acec0 impl_cuCtxSetFlags medium group=cuCtxSetFlags variants=cuCtxSetFlags; before=0x1acec0,0x26ab80,0x2d8f40,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1acec0,0x26ab80; outside_wrappers=0 unverified +0x001ad4d0 impl_cuCtxResetPersistingL2Cache medium group=cuCtxResetPersistingL2Cache variants=cuCtxResetPersistingL2Cache; before=0x179cf0,0x1ad4d0,0x26b920,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1ad4d0,0x26b920; outside_wrappers=0 unverified +0x001ad670 impl_cuCtxSetCacheConfig medium group=cuCtxSetCacheConfig variants=cuCtxSetCacheConfig; before=0x1ad670,0x26ae40,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1ad670,0x26ae40; outside_wrappers=0 unverified +0x001ad770 impl_cuCtxSetSharedMemConfig medium group=cuCtxSetSharedMemConfig variants=cuCtxSetSharedMemConfig; before=0x1ad770,0x26ae40,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1ad770,0x26ae40; outside_wrappers=0 unverified +0x001ad870 impl_cuCtxEnablePeerAccess medium group=cuCtxEnablePeerAccess variants=cuCtxEnablePeerAccess; before=0x1792f0,0x179350,0x1ad870,0x2baed0,0x2c1500,0x2e5bf0,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ad870,0x2e5bf0,0x2c1500; outside_wrappers=0 unverified +0x001ada20 impl_cuCtxDisablePeerAccess medium group=cuCtxDisablePeerAccess variants=cuCtxDisablePeerAccess; before=0x1792f0,0x179350,0x1ada20,0x259530,0x2baed0,0x2d8f40,0x2e4ff0,0x2fe230,0x2fe540,0x443950,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1ada20,0x2e4ff0,0x259530,0x443950; outside_wrappers=0 unverified +0x001adc60 impl_cuCtxRecordEvent high group=cuCtxRecordEvent variants=cuCtxRecordEvent; before=0x1792f0,0x179350,0x1adc60,0x2886b0,0x2895f0,0x2bb140,0x2f8e10,0x2f8fd0,0x2fe540,0x4a7e90,0xd1e3f0,0xd4d9a0,0xd4dc50; after=0x1adc60; outside_wrappers=0 unverified +0x001ade30 impl_cuCtxWaitEvent medium group=cuCtxWaitEvent variants=cuCtxWaitEvent; before=0x1792f0,0x179350,0x1ade30,0x2bb140,0x2ed7d0,0x2fe540,0x4a7e90,0xd1e090,0xd1e230,0xd4d9a0,0xd4dc50; after=0x1ade30,0xd1e230; outside_wrappers=0 unverified +0x001adfb0 impl_cuDeviceGetByPCIBusId high group=cuDeviceGetByPCIBusId variants=cuDeviceGetByPCIBusId; before=0x1adfb0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1adfb0; outside_wrappers=0 unverified +0x001ae150 impl_cuDeviceGetPCIBusId medium group=cuDeviceGetPCIBusId variants=cuDeviceGetPCIBusId; before=0x179560,0x1ae150,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ae150,0x179560; outside_wrappers=0 unverified +0x001ae290 impl_cuDeviceGetTexture1DLinearMaxWidth medium group=cuDeviceGetTexture1DLinearMaxWidth variants=cuDeviceGetTexture1DLinearMaxWidth; before=0x1ae290,0x257890,0x2583d0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ae290,0x257890; outside_wrappers=0 unverified +0x001ae3d0 impl_cuDeviceComputeCapability high group=cuDeviceComputeCapability variants=cuDeviceComputeCapability; before=0x1ae3d0,0x26bc10,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ae3d0; outside_wrappers=0 verified +0x001ae550 impl_cuDeviceGetProperties high group=cuDeviceGetProperties variants=cuDeviceGetProperties; before=0x1ae550,0x26bc10,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ae550; outside_wrappers=0 verified +0x001ae780 impl_cuDeviceGetAttribute high group=cuDeviceGetAttribute variants=cuDeviceGetAttribute; before=0x1ae780,0x26bc10,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ae780; outside_wrappers=0 verified +0x001aec20 impl_cuDeviceCanAccessPeer medium group=cuDeviceCanAccessPeer variants=cuDeviceCanAccessPeer; before=0x1aec20,0x2c1610,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1aec20,0x2c1610; outside_wrappers=0 unverified +0x001aed70 impl_cuDevicePrimaryCtxRetain medium group=cuDevicePrimaryCtxRetain variants=cuDevicePrimaryCtxRetain; before=0x1792f0,0x179350,0x1aed70,0x277570,0x2fe540,0x4a7e90,0xd4dc50; after=0x1aed70,0x277570; outside_wrappers=0 unverified +0x001af400 impl_cuDeviceSetMemPool medium group=cuDeviceSetMemPool variants=cuDeviceSetMemPool; before=0x1af400,0x2c3b00,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1af400,0x2c3b00; outside_wrappers=0 unverified +0x001af5a0 impl_cuDeviceGetMemPool medium group=cuDeviceGetMemPool variants=cuDeviceGetMemPool; before=0x1af5a0,0x26bab0,0x2c61f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1af5a0,0x2c61f0,0x26bab0; outside_wrappers=0 unverified +0x001af960 impl_cuDeviceGetGraphMemAttribute high group=cuDeviceGetGraphMemAttribute variants=cuDeviceGetGraphMemAttribute; before=0x1af960,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1af960; outside_wrappers=0 unverified +0x001afae0 impl_cuDeviceSetGraphMemAttribute high group=cuDeviceSetGraphMemAttribute variants=cuDeviceSetGraphMemAttribute; before=0x1afae0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1afae0; outside_wrappers=0 unverified +0x001afc50 impl_cuDevSmResourceSplitByCount high group=cuDevSmResourceSplitByCount variants=cuDevSmResourceSplitByCount; before=0x1afc50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1afc50; outside_wrappers=0 unverified +0x001afe80 impl_cuGraphicsEGLRegisterImage high group=cuGraphicsEGLRegisterImage variants=cuGraphicsEGLRegisterImage; before=0x1afe80,0x2fe540,0x4a7e90,0xd4dc50; after=0x1afe80; outside_wrappers=0 unverified +0x001b0130 impl_cuEGLStreamConsumerDisconnect medium group=cuEGLStreamConsumerDisconnect variants=cuEGLStreamConsumerDisconnect; before=0x1b0130,0x2fde50,0x2fe230,0x2fe540,0x350790,0x3c52d0,0x3c5370,0x400550,0x4a7e90,0xd4dc50; after=0x1b0130,0x350790; outside_wrappers=0 unverified +0x001b02b0 impl_cuEGLStreamProducerDisconnect medium group=cuEGLStreamProducerDisconnect variants=cuEGLStreamProducerDisconnect; before=0x1b02b0,0x2fde50,0x2fe230,0x2fe540,0x350790,0x3c52d0,0x3c5370,0x400550,0x4a7e90,0xd4dc50; after=0x1b02b0,0x350790; outside_wrappers=0 unverified +0x001b0440 impl_cuEGLStreamConsumerAcquireFrame high group=cuEGLStreamConsumerAcquireFrame variants=cuEGLStreamConsumerAcquireFrame; before=0x1b0440,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b0440; outside_wrappers=0 unverified +0x001b0960 impl_cuEGLStreamProducerConnect medium group=cuEGLStreamProducerConnect variants=cuEGLStreamProducerConnect; before=0x1b0960,0x26f650,0x270710,0x2fe230,0x2fe540,0x3c52d0,0x3c5370,0x3c54c0,0x4a7e90,0xd4da90,0xd4dc50; after=0x1b0960,0x270710,0x26f650,0x3c54c0; outside_wrappers=0 unverified +0x001b0af0 impl_cuEGLStreamProducerPresentFrame high group=cuEGLStreamProducerPresentFrame variants=cuEGLStreamProducerPresentFrame; before=0x1b0af0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b0af0; outside_wrappers=0 unverified +0x001b1370 impl_cuGraphDestroyNode high group=cuGraphDestroyNode variants=cuGraphDestroyNode; before=0x1b1370,0x2a2b90,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1b1370; outside_wrappers=0 unverified +0x001b1f20 impl_cuGraphAddEmptyNode high group=cuGraphAddEmptyNode variants=cuGraphAddEmptyNode; before=0x1b1f20,0x2a50f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b1f20; outside_wrappers=0 unverified +0x001b2070 impl_cuGraphAddEventRecordNode high group=cuGraphAddEventRecordNode variants=cuGraphAddEventRecordNode; before=0x1b2070,0x2a50f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b2070; outside_wrappers=0 unverified +0x001b2190 impl_cuGraphAddEventWaitNode high group=cuGraphAddEventWaitNode variants=cuGraphAddEventWaitNode; before=0x1b2190,0x2a50f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b2190; outside_wrappers=0 unverified +0x001b23c0 impl_cuEventQuery medium group=cuEventQuery variants=cuEventQuery; before=0x1b23c0,0x288ba0,0x2bb140,0x2fe230,0x2fe540,0x3034c0,0x4a7e90,0xd1e090,0xd4dc50; after=0x1b23c0,0x3034c0; outside_wrappers=0 unverified +0x001b3370 impl_cuIpcGetEventHandle medium group=cuIpcGetEventHandle variants=cuIpcGetEventHandle; before=0x1b3370,0x2bb420,0x2bbaa0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b3370,0x2bbaa0; outside_wrappers=0 unverified +0x001b37e0 impl_cuImportExternalMemory high group=cuImportExternalMemory variants=cuImportExternalMemory; before=0x1b37e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b37e0; outside_wrappers=0 unverified +0x001b3ab0 impl_cuExternalMemoryGetMappedBuffer medium group=cuExternalMemoryGetMappedBuffer variants=cuExternalMemoryGetMappedBuffer; before=0x1b3ab0,0x28a960,0x2d8f40,0x2fe540,0x4a7e90,0xd4d9a0,0xd4db70,0xd4dc50,0xd51980; after=0x1b3ab0,0x28a960; outside_wrappers=0 unverified +0x001b3cd0 impl_cuExternalMemoryGetMappedMipmappedArray high group=cuExternalMemoryGetMappedMipmappedArray variants=cuExternalMemoryGetMappedMipmappedArray; before=0x1b3cd0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b3cd0; outside_wrappers=0 unverified +0x001b3fb0 impl_cuDestroyExternalMemory medium group=cuDestroyExternalMemory variants=cuDestroyExternalMemory; before=0x1b3fb0,0x28b480,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b3fb0,0x28b480; outside_wrappers=0 unverified +0x001b4110 impl_cuImportExternalSemaphore high group=cuImportExternalSemaphore variants=cuImportExternalSemaphore; before=0x1b4110,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b4110; outside_wrappers=0 unverified +0x001b50a0 impl_cuFuncSetBlockShape medium group=cuFuncSetBlockShape variants=cuFuncSetBlockShape; before=0x1b50a0,0x2938b0,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b50a0,0x2938b0; outside_wrappers=0 unverified +0x001b5260 impl_cuFuncSetSharedSize medium group=cuFuncSetSharedSize variants=cuFuncSetSharedSize; before=0x1b5260,0x293ae0,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b5260,0x293ae0; outside_wrappers=0 unverified +0x001b53f0 impl_cuFuncSetCacheConfig medium group=cuFuncSetCacheConfig variants=cuFuncSetCacheConfig; before=0x1b53f0,0x262ec0,0x263b80,0x263e10,0x264110,0x2baed0,0x2d8f40,0x2e0770,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b53f0,0x262ec0,0x264110,0x263b80,0x263e10; outside_wrappers=0 unverified +0x001b5670 impl_cuFuncSetSharedMemConfig medium group=cuFuncSetSharedMemConfig variants=cuFuncSetSharedMemConfig; before=0x1b5670,0x262ec0,0x263b80,0x263e10,0x264110,0x2baed0,0x2d8f40,0x2e0770,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b5670,0x262ec0,0x264110,0x263b80,0x263e10; outside_wrappers=0 unverified +0x001b58f0 impl_cuFuncGetModule high group=cuFuncGetModule variants=cuFuncGetModule; before=0x1b58f0,0x2baed0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b58f0; outside_wrappers=0 unverified +0x001b5a20 impl_cuFuncIsLoaded high group=cuFuncIsLoaded variants=cuFuncIsLoaded; before=0x1b5a20,0x2baed0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b5a20; outside_wrappers=0 unverified +0x001b5b70 impl_cuFuncLoad high group=cuFuncLoad variants=cuFuncLoad; before=0x1b5b70,0x2baed0,0x2e0770,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b5b70; outside_wrappers=0 unverified +0x001b5cb0 impl_cuFuncGetAttribute medium group=cuFuncGetAttribute variants=cuFuncGetAttribute; before=0x1b5cb0,0x293ec0,0x2baed0,0x2dc100,0x2fe540,0x4a7e90,0xd4bd20,0xd4dc50; after=0x1b5cb0,0x293ec0; outside_wrappers=0 unverified +0x001b5e80 impl_cuFuncSetAttribute medium group=cuFuncSetAttribute variants=cuFuncSetAttribute; before=0x1b5e80,0x293b70,0x2baed0,0x2d8f40,0x2e0770,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1b5e80,0x293b70; outside_wrappers=0 unverified +0x001b6030 impl_cuModuleGetFunction high group=cuModuleGetFunction variants=cuModuleGetFunction; before=0x1b6030,0x2baed0,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980,0xd6c4a0; after=0x1b6030; outside_wrappers=0 unverified +0x001b6210 impl_cuGraphicsGLRegisterBuffer medium group=cuGraphicsGLRegisterBuffer variants=cuGraphicsGLRegisterBuffer; before=0x179200,0x1799e0,0x1b6210,0x294fd0,0x2fe230,0x2fe540,0x3c52d0,0x3c5370,0x3c54c0,0x4a7e90,0xd4da90,0xd4dc50; after=0x1b6210,0x294fd0,0x3c54c0; outside_wrappers=0 unverified +0x001b63e0 impl_cuGraphicsGLRegisterImage high group=cuGraphicsGLRegisterImage variants=cuGraphicsGLRegisterImage; before=0x1b63e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b63e0; outside_wrappers=0 unverified +0x001b66e0 impl_cuGLRegisterBufferObject high group=cuGLRegisterBufferObject variants=cuGLRegisterBufferObject; before=0x1b66e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b66e0; outside_wrappers=0 unverified +0x001b75f0 impl_cuGraphNodeFindInClone high group=cuGraphNodeFindInClone variants=cuGraphNodeFindInClone; before=0x1b75f0,0x2fe540,0x3aa550,0x4a7e90,0xd4d9a0,0xd4dc50,0xd663e0; after=0x1b75f0; outside_wrappers=0 unverified +0x001b7e60 impl_cuGLUnregisterBufferObject high group=cuGLUnregisterBufferObject variants=cuGLUnregisterBufferObject; before=0x1b7e60,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b7e60; outside_wrappers=0 unverified +0x001b8090 impl_cuGLSetBufferObjectMapFlags high group=cuGLSetBufferObjectMapFlags variants=cuGLSetBufferObjectMapFlags; before=0x1b8090,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b8090; outside_wrappers=0 unverified +0x001b8580 impl_cuUserObjectRetain high group=cuUserObjectRetain variants=cuUserObjectRetain; before=0x179200,0x1b8580,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b8580; outside_wrappers=0 unverified +0x001b8670 impl_cuUserObjectRelease high group=cuUserObjectRelease variants=cuUserObjectRelease; before=0x179200,0x1b8670,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b8670; outside_wrappers=0 unverified +0x001b8760 impl_cuGraphRetainUserObject medium group=cuGraphRetainUserObject variants=cuGraphRetainUserObject; before=0x1792f0,0x179350,0x1b8760,0x2fe540,0x4a7e90,0xd4dc50,0xd6a780; after=0x1b8760,0xd6a780; outside_wrappers=0 unverified +0x001b8900 impl_cuGraphReleaseUserObject medium group=cuGraphReleaseUserObject variants=cuGraphReleaseUserObject; before=0x1792f0,0x179350,0x1b8900,0x2fe540,0x4a7e90,0xd4dc50,0xd6a780; after=0x1b8900,0xd6a780; outside_wrappers=0 unverified +0x001b8ba0 impl_cuGraphGetEdges medium group=cuGraphGetEdges variants=cuGraphGetEdges,cuGraphGetEdges_v2; before=0x1b8ba0,0x2a69b0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1b8ba0,0x2a69b0; outside_wrappers=0 unverified +0x001b8fd0 impl_cuGraphExecDestroy medium group=cuGraphExecDestroy variants=cuGraphExecDestroy; before=0x1b8fd0,0x1c1c60,0x2a3530,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b8fd0,0x2a3530; outside_wrappers=0 unverified +0x001b96b0 impl_cuGraphExecMemcpyNodeSetParams medium group=cuGraphExecMemcpyNodeSetParams variants=cuGraphExecMemcpyNodeSetParams; before=0x1b96b0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b96b0; outside_wrappers=1 unverified +0x001b9830 impl_cuGraphExecMemsetNodeSetParams medium group=cuGraphExecMemsetNodeSetParams variants=cuGraphExecMemsetNodeSetParams; before=0x1b9830,0x2fe540,0x4a7e90,0xd4dc50; after=0x1b9830; outside_wrappers=1 unverified +0x001ba6f0 impl_cuGraphAddHostNode high group=cuGraphAddHostNode variants=cuGraphAddHostNode; before=0x1ba6f0,0x2a63a0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ba6f0; outside_wrappers=0 unverified +0x001baec0 impl_cuGraphAddDependencies high group=cuGraphAddDependencies variants=cuGraphAddDependencies,cuGraphAddDependencies_v2; before=0x1baec0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1baec0; outside_wrappers=0 unverified +0x001bb0c0 impl_cuGraphRemoveDependencies high group=cuGraphRemoveDependencies variants=cuGraphRemoveDependencies,cuGraphRemoveDependencies_v2; before=0x1bb0c0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1bb0c0; outside_wrappers=0 unverified +0x001bb960 impl_cuMemcpyBatchAsync medium group=cuMemcpyBatchAsync variants=cuMemcpyBatchAsync_ptsz,cuMemcpyBatchAsync; before=0x1bb8a0,0x1bb960,0x1bddb0,0x2baed0,0x2f16c0,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50; after=0x1bb960,0x1bddb0,0x1bb8a0; outside_wrappers=0 unverified +0x001be0e0 impl_cuMemcpy3DBatchAsync high group=cuMemcpy3DBatchAsync variants=cuMemcpy3DBatchAsync_ptsz,cuMemcpy3DBatchAsync; before=0x1be0e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1be0e0; outside_wrappers=0 unverified +0x001be900 impl_cuCtxDestroy high group=cuCtxDestroy variants=cuCtxDestroy_v2,cuCtxDestroy; before=0x1be900,0x26d8e0,0x2baed0,0x2fde50,0x2fe0f0,0x2fe540,0x4a7e90,0xd4da50,0xd4dc50; after=0x1be900; outside_wrappers=0 verified +0x001bea40 impl_cuCtxPushCurrent medium group=cuCtxPushCurrent variants=cuCtxPushCurrent_v2,cuCtxPushCurrent; before=0x1bea40,0x24a4d0,0x24b930,0x2fde50,0x2fe020,0x2fe540,0x4a7e90,0xd4dc50,0xd4f3e0; after=0x1bea40,0x24b930; outside_wrappers=0 unverified +0x001bec50 impl_cuCtxPopCurrent medium group=cuCtxPopCurrent variants=cuCtxPopCurrent_v2,cuCtxPopCurrent; before=0x1bec50,0x24a4d0,0x24b9e0,0x2fde50,0x2fe0f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1bec50,0x24b9e0; outside_wrappers=0 unverified +0x001bee20 impl_cuDeviceGetNvSciSyncAttributes medium group=cuDeviceGetNvSciSyncAttributes variants=cuDeviceGetNvSciSyncAttributes; before=0x1bee20,0x1bf270,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1bee20,0x1bf270; outside_wrappers=0 unverified +0x001bf3e0 impl_cuDeviceGetDefaultMemPool medium group=cuDeviceGetDefaultMemPool variants=cuDeviceGetDefaultMemPool; before=0x1798f0,0x1bf3e0,0x1bf560,0x2fe440,0x2fe540,0x4a7e90,0xd4dc50; after=0x1bf3e0,0x1bf560; outside_wrappers=0 unverified +0x001bf660 impl_cuCtxGetDevResource high group=cuCtxGetDevResource variants=cuCtxGetDevResource; before=0x1bf660,0x2baed0,0x2bb140,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1bf660; outside_wrappers=0 unverified +0x001bf820 impl_cuGreenCtxGetDevResource high group=cuGreenCtxGetDevResource variants=cuGreenCtxGetDevResource; before=0x1bf820,0x2baed0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1bf820; outside_wrappers=0 unverified +0x001bf970 impl_cuDevResourceGenerateDesc high group=cuDevResourceGenerateDesc variants=cuDevResourceGenerateDesc; before=0x1bf970,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1bf970; outside_wrappers=0 unverified +0x001c0120 impl_cuEGLStreamConsumerConnect medium group=cuEGLStreamConsumerConnect variants=cuEGLStreamConsumerConnect; before=0x1c0120,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c0120; outside_wrappers=1 unverified +0x001c0120 impl_cuEGLStreamConsumerConnectWithFlags medium group=cuEGLStreamConsumerConnectWithFlags variants=cuEGLStreamConsumerConnectWithFlags; before=0x1c0120,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c0120; outside_wrappers=1 unverified +0x001c0930 impl_cuStreamBatchMemOp high group=cuStreamBatchMemOp variants=cuStreamBatchMemOp_v2_ptsz,cuStreamBatchMemOp,cuStreamBatchMemOp_ptsz,cuStreamBatchMemOp_v2; before=0x1c0930,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c0930; outside_wrappers=0 unverified +0x001c0d00 impl_cuGraphBatchMemOpNodeSetParams medium group=cuGraphBatchMemOpNodeSetParams variants=cuGraphBatchMemOpNodeSetParams; before=0x179200,0x179a80,0x179c00,0x1c0740,0x1c0d00,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c0d00; outside_wrappers=1 unverified +0x001c0f70 impl_cuSignalExternalSemaphoresAsync high group=cuSignalExternalSemaphoresAsync variants=cuSignalExternalSemaphoresAsync_ptsz,cuSignalExternalSemaphoresAsync; before=0x1c0f70,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c0f70; outside_wrappers=0 unverified +0x001c12c0 impl_cuWaitExternalSemaphoresAsync high group=cuWaitExternalSemaphoresAsync variants=cuWaitExternalSemaphoresAsync_ptsz,cuWaitExternalSemaphoresAsync; before=0x1c12c0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c12c0; outside_wrappers=0 unverified +0x001c15f0 impl_cuDestroyExternalSemaphore medium group=cuDestroyExternalSemaphore variants=cuDestroyExternalSemaphore; before=0x1798f0,0x1c15f0,0x1c1730,0x2fe230,0x2fe440,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c15f0,0x1c1730; outside_wrappers=0 unverified +0x001c1820 impl_cuGLCtxCreate high group=cuGLCtxCreate variants=cuGLCtxCreate_v2,cuGLCtxCreate; before=0x1c1820,0x2689e0,0x26d8e0,0x277680,0x2fe020,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c1820; outside_wrappers=0 unverified +0x001c2060 impl_cuGraphAddKernelNode high group=cuGraphAddKernelNode variants=cuGraphAddKernelNode_v2,cuGraphAddKernelNode; before=0x1c2060,0x1c2b10,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c2060; outside_wrappers=0 unverified +0x001c2570 impl_cuGraphAddMemsetNode medium group=cuGraphAddMemsetNode variants=cuGraphAddMemsetNode; before=0x1c2570,0x2a64e0,0x2baed0,0x2bb140,0x2ded70,0x2fde50,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c2570,0x2ded70,0x2a64e0; outside_wrappers=0 unverified +0x001c2830 impl_cuGraphMemsetNodeSetParams medium group=cuGraphMemsetNodeSetParams variants=cuGraphMemsetNodeSetParams; before=0x1c2830,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c2830; outside_wrappers=1 unverified +0x001c3c80 impl_cuGraphKernelNodeCopyAttributes high group=cuGraphKernelNodeCopyAttributes variants=cuGraphKernelNodeCopyAttributes; before=0x1c2b10,0x1c3c80,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c3c80; outside_wrappers=0 unverified +0x001c3e90 impl_cuGraphKernelNodeGetAttribute high group=cuGraphKernelNodeGetAttribute variants=cuGraphKernelNodeGetAttribute; before=0x1c3e90,0x2d8f40,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50,0xd51980; after=0x1c3e90; outside_wrappers=0 unverified +0x001c4540 impl_cuGraphKernelNodeSetAttribute medium group=cuGraphKernelNodeSetAttribute variants=cuGraphKernelNodeSetAttribute; before=0x1c40f0,0x1c4540,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1c4540,0x1c40f0; outside_wrappers=0 unverified +0x001c4a90 impl_cuGraphExecNodeSetParams medium group=cuGraphExecNodeSetParams variants=cuGraphExecNodeSetParams; before=0x1a39d0,0x1b96b0,0x1b9830,0x1b9a90,0x1b9ea0,0x1c0e20,0x1c3870,0x1c4920,0x1c4a90,0x2ac1b0,0x2b4f80,0x2fe540,0x4a7e90,0xd183b0,0xd4dc50; after=0x1c4a90,0x1b9a90,0x1b9ea0,0x1a39d0,0x2b4f80,0x1c3870,0x1b96b0,0x1c0e20,0x1b9830,0x2ac1b0,0x1c4920; outside_wrappers=0 unverified +0x001c4d50 impl_cuGraphMemcpyNodeSetParams medium group=cuGraphMemcpyNodeSetParams variants=cuGraphMemcpyNodeSetParams; before=0x1c4d50,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c4d50; outside_wrappers=1 unverified +0x001c4e80 impl_cuGraphNodeSetParams medium group=cuGraphNodeSetParams variants=cuGraphNodeSetParams; before=0x1c0d00,0x1c2830,0x1c2a00,0x1c3750,0x1c4920,0x1c4d50,0x1c4e80,0x2a9f60,0x2aa190,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c4e80,0x1c2a00,0x1c2830,0x1c0d00,0x1c4d50,0x1c3750,0x1c4920; outside_wrappers=0 unverified +0x001c51c0 impl_cuGraphAddMemcpyNode medium group=cuGraphAddMemcpyNode variants=cuGraphAddMemcpyNode; before=0x1c51c0,0x2a60a0,0x2baed0,0x2bb140,0x2d7780,0x2fde50,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c51c0,0x2a60a0,0x2d7780; outside_wrappers=0 unverified +0x001c53e0 impl_cuGraphAddNode high group=cuGraphAddNode variants=cuGraphAddNode,cuGraphAddNode_v2; before=0x1c53e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c53e0; outside_wrappers=0 unverified +0x001c6360 impl_cuMemExportToShareableHandle high group=cuMemExportToShareableHandle variants=cuMemExportToShareableHandle; before=0x1c6360,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c6360; outside_wrappers=0 unverified +0x001c64f0 impl_cuLibraryGetKernel high group=cuLibraryGetKernel variants=cuLibraryGetKernel; before=0x1c64f0,0x2fe540,0x3aa550,0x4a7e90,0xd4d9a0,0xd4dc50,0xd663e0; after=0x1c64f0; outside_wrappers=0 unverified +0x001c6610 impl_cuLibraryGetManaged high group=cuLibraryGetManaged variants=cuLibraryGetManaged; before=0x1c6610,0x2fe540,0x3aa550,0x4a7e90,0xd4d9a0,0xd4dc50,0xd663e0; after=0x1c6610; outside_wrappers=0 unverified +0x001c69f0 impl_cuMemGetAllocationGranularity high group=cuMemGetAllocationGranularity variants=cuMemGetAllocationGranularity; before=0x1c69f0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c69f0; outside_wrappers=0 unverified +0x001c6ec0 impl_cuGreenCtxRecordEvent high group=cuGreenCtxRecordEvent variants=cuGreenCtxRecordEvent; before=0x1792f0,0x179350,0x1c6ec0,0x2886b0,0x2895f0,0x2baed0,0x2f8e10,0x2f8fd0,0x2fe540,0x4a7e90,0xd1e3f0,0xd4d9a0,0xd4dc50; after=0x1c6ec0; outside_wrappers=0 unverified +0x001c7000 impl_cuGreenCtxWaitEvent medium group=cuGreenCtxWaitEvent variants=cuGreenCtxWaitEvent; before=0x1c7000,0x2baed0,0x2ed7d0,0x2fe540,0x4a7e90,0xd1e230,0xd1e3f0,0xd4d9a0,0xd4dc50; after=0x1c7000,0xd1e230; outside_wrappers=0 unverified +0x001c7120 impl_cuGreenCtxStreamCreate medium group=cuGreenCtxStreamCreate variants=cuGreenCtxStreamCreate; before=0x1c7120,0x2baed0,0x2ef570,0x2f0fb0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50,0xd613b0; after=0x1c7120,0x2ef570,0x2f0fb0,0xd613b0; outside_wrappers=0 unverified +0x001c73e0 impl_cuGraphicsResourceGetMappedMipmappedArray high group=cuGraphicsResourceGetMappedMipmappedArray variants=cuGraphicsResourceGetMappedMipmappedArray; before=0x1798f0,0x1c73e0,0x2baed0,0x2fe440,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1c73e0; outside_wrappers=0 unverified +0x001c7530 impl_cuGraphicsSubResourceGetMappedArray high group=cuGraphicsSubResourceGetMappedArray variants=cuGraphicsSubResourceGetMappedArray; before=0x1c7530,0x2baed0,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1c7530; outside_wrappers=0 unverified +0x001c7cf0 impl_cuMemPoolGetAccess high group=cuMemPoolGetAccess variants=cuMemPoolGetAccess; before=0x1c7cf0,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c7cf0; outside_wrappers=0 unverified +0x001c7ed0 impl_cuGraphicsUnregisterResource high group=cuGraphicsUnregisterResource variants=cuGraphicsUnregisterResource; before=0x1c7ed0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1c7ed0; outside_wrappers=0 unverified +0x001c80b0 impl_cuGraphicsResourceGetMappedEglFrame high group=cuGraphicsResourceGetMappedEglFrame variants=cuGraphicsResourceGetMappedEglFrame; before=0x1c80b0,0x2baed0,0x2fde50,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1c80b0; outside_wrappers=0 unverified +0x001c94b0 impl_cuMemPoolExportPointer high group=cuMemPoolExportPointer variants=cuMemPoolExportPointer; before=0x1792f0,0x179350,0x1c94b0,0x2fe230,0x2fe540,0x3fe100,0x4a7e90,0xd4dc50; after=0x1c94b0; outside_wrappers=0 unverified +0x001c9700 impl_cuMemGetAllocationPropertiesFromHandle high group=cuMemGetAllocationPropertiesFromHandle variants=cuMemGetAllocationPropertiesFromHandle; before=0x1c9700,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c9700; outside_wrappers=0 unverified +0x001c9b90 impl_cuKernelGetParamInfo high group=cuKernelGetParamInfo variants=cuKernelGetParamInfo; before=0x1c9b90,0x2bb420,0x2bfad0,0x2fe540,0x4a7e90,0xd4dc50,0xd6c4a0; after=0x1c9b90; outside_wrappers=0 unverified +0x001c9d80 impl_cuLibraryLoadFromFile medium group=cuLibraryLoadFromFile variants=cuLibraryLoadFromFile; before=0x179200,0x1c9d80,0x2bbd30,0x2d3f60,0x2df390,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1c9d80,0x2d3f60; outside_wrappers=0 unverified +0x001c9f90 impl_cuLibraryGetGlobal high group=cuLibraryGetGlobal variants=cuLibraryGetGlobal; before=0x1c9f90,0x2bb420,0x2bfad0,0x2f9f00,0x2fe540,0x402500,0x4a7e90,0xd4dc50; after=0x1c9f90; outside_wrappers=0 unverified +0x001ca790 impl_cuCtxGetLimit medium group=cuCtxGetLimit variants=cuCtxGetLimit; before=0x1ca790,0x2bb420,0x2c0400,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ca790,0x2c0400; outside_wrappers=0 unverified +0x001ca940 impl_cuMemHostUnregister high group=cuMemHostUnregister variants=cuMemHostUnregister; before=0x1ca940,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ca940; outside_wrappers=0 unverified +0x001cab60 impl_cuMemAllocManaged high group=cuMemAllocManaged variants=cuMemAllocManaged; before=0x1cab60,0x2fe540,0x4a7e90,0xd4dc50; after=0x1cab60; outside_wrappers=0 unverified +0x001cb030 impl_cuMemRangeGetAttributes medium group=cuMemRangeGetAttributes variants=cuMemRangeGetAttributes; before=0x1cb030,0x2fe540,0x4a7e90,0xd4dc50; after=0x1cb030; outside_wrappers=1 unverified +0x001cbda0 impl_cuMemHostGetFlags high group=cuMemHostGetFlags variants=cuMemHostGetFlags; before=0x1cbda0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1cbda0; outside_wrappers=0 unverified +0x001cc6f0 impl_cuIpcGetMemHandle medium group=cuIpcGetMemHandle variants=cuIpcGetMemHandle; before=0x1cc6f0,0x2bb420,0x2bbaa0,0x2fe540,0x3fe700,0x402500,0x4a7e90,0xd4bd90,0xd4dc50; after=0x1cc6f0,0x2bbaa0; outside_wrappers=0 unverified +0x001ccb30 impl_cuIpcCloseMemHandle high group=cuIpcCloseMemHandle variants=cuIpcCloseMemHandle; before=0x1ccb30,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ccb30; outside_wrappers=0 unverified +0x001ccd30 impl_cuMemGetHandleForAddressRange high group=cuMemGetHandleForAddressRange variants=cuMemGetHandleForAddressRange; before=0x1795e0,0x1ccd30,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ccd30; outside_wrappers=0 unverified +0x001d2a60 impl_cuMemAddressReserve medium group=cuMemAddressReserve variants=cuMemAddressReserve; before=0x1792f0,0x179350,0x1795e0,0x1d2a60,0x2d8ae0,0x2fe540,0x3acc80,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1d2a60,0x3acc80,0x2d8ae0; outside_wrappers=0 unverified +0x001d2c30 impl_cuMemAddressFree medium group=cuMemAddressFree variants=cuMemAddressFree; before=0x1792f0,0x179350,0x1795e0,0x1d2c30,0x2d8c30,0x2fe540,0x306160,0x3ac0f0,0x3ac3f0,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1d2c30,0x2d8c30,0x3ac0f0; outside_wrappers=0 unverified +0x001d2f80 impl_cuMemPoolImportPointer high group=cuMemPoolImportPointer variants=cuMemPoolImportPointer; before=0x1d2f80,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d2f80; outside_wrappers=0 unverified +0x001d3bf0 impl_cuMemUnmap high group=cuMemUnmap variants=cuMemUnmap; before=0x1d3bf0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d3bf0; outside_wrappers=0 unverified +0x001d4010 impl_cuMemImportFromShareableHandle high group=cuMemImportFromShareableHandle variants=cuMemImportFromShareableHandle; before=0x1d4010,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d4010; outside_wrappers=0 unverified +0x001d4800 impl_cuMemRetainAllocationHandle medium group=cuMemRetainAllocationHandle variants=cuMemRetainAllocationHandle; before=0x1792f0,0x179350,0x1d4800,0x257620,0x2fe540,0x3ac3f0,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1d4800,0x257620; outside_wrappers=0 unverified +0x001d5020 impl_cuMemPoolTrimTo medium group=cuMemPoolTrimTo variants=cuMemPoolTrimTo; before=0x1792f0,0x179350,0x1d5020,0x2c51a0,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d5020,0x2c51a0; outside_wrappers=0 unverified +0x001d5180 impl_cuMemPoolSetAttribute medium group=cuMemPoolSetAttribute variants=cuMemPoolSetAttribute; before=0x1792f0,0x179350,0x1d5180,0x2db020,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d5180,0x2db020; outside_wrappers=0 unverified +0x001d52d0 impl_cuMemPoolGetAttribute medium group=cuMemPoolGetAttribute variants=cuMemPoolGetAttribute; before=0x1792f0,0x179350,0x1d52d0,0x2db110,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d52d0,0x2db110; outside_wrappers=0 unverified +0x001d5430 impl_cuMemPoolSetAccess high group=cuMemPoolSetAccess variants=cuMemPoolSetAccess; before=0x1d5430,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d5430; outside_wrappers=0 unverified +0x001d5a00 impl_cuMemPoolExportToShareableHandle medium group=cuMemPoolExportToShareableHandle variants=cuMemPoolExportToShareableHandle; before=0x1792f0,0x179350,0x1799e0,0x1d5a00,0x2c6e60,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d5a00,0x2c6e60; outside_wrappers=0 unverified +0x001d5c60 impl_cuMemPoolImportFromShareableHandle medium group=cuMemPoolImportFromShareableHandle variants=cuMemPoolImportFromShareableHandle; before=0x1792f0,0x179350,0x179450,0x1799e0,0x1d5c60,0x28d350,0x2c3b00,0x2fe230,0x2fe540,0x30a9d0,0x3aa550,0x3ab420,0x4a7e90,0xd4dc50,0xd663e0; after=0x1d5c60,0x28d350,0x3ab420,0x30a9d0,0x2c3b00; outside_wrappers=0 unverified +0x001d5ed0 impl_cuGraphExecUpdate high group=cuGraphExecUpdate variants=cuGraphExecUpdate_v2,cuGraphExecUpdate; before=0x1d5ed0,0x2fe540,0x4a7e90,0xd4da50,0xd4dc50; after=0x1d5ed0; outside_wrappers=0 unverified +0x001d61c0 impl_cuGraphConditionalHandleCreate high group=cuGraphConditionalHandleCreate variants=cuGraphConditionalHandleCreate; before=0x1d61c0,0x2baed0,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d61c0; outside_wrappers=0 unverified +0x001d6410 impl_cuGraphicsResourceSetMapFlags high group=cuGraphicsResourceSetMapFlags variants=cuGraphicsResourceSetMapFlags_v2,cuGraphicsResourceSetMapFlags; before=0x1d6410,0x2baed0,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1d6410; outside_wrappers=0 unverified +0x001d65a0 impl_cuGraphicsMapResources high group=cuGraphicsMapResources variants=cuGraphicsMapResources_ptsz,cuGraphicsMapResources; before=0x1d65a0,0x2baed0,0x2f16c0,0x2fe230,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1d65a0; outside_wrappers=0 unverified +0x001d67a0 impl_cuGraphicsUnmapResources high group=cuGraphicsUnmapResources variants=cuGraphicsUnmapResources_ptsz,cuGraphicsUnmapResources; before=0x1d67a0,0x2baed0,0x2f16c0,0x2fe230,0x2fe540,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50; after=0x1d67a0; outside_wrappers=0 unverified +0x001d69a0 impl_cuLinkAddFile medium group=cuLinkAddFile variants=cuLinkAddFile_v2,cuLinkAddFile; before=0x179200,0x1d69a0,0x2bbd30,0x2d2990,0x2df390,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d69a0,0x2df390; outside_wrappers=0 unverified +0x001d7d70 impl_cuLaunchCooperativeKernelMultiDevice high group=cuLaunchCooperativeKernelMultiDevice variants=cuLaunchCooperativeKernelMultiDevice; before=0x1d7d70,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d7d70; outside_wrappers=0 unverified +0x001d9540 impl_cuKernelSetCacheConfig high group=cuKernelSetCacheConfig variants=cuKernelSetCacheConfig; before=0x1d9540,0x2bb420,0x2fe540,0x4a7e90,0xd4dc50; after=0x1d9540; outside_wrappers=0 unverified +0x001da450 impl_cuMemFree high group=cuMemFree variants=cuMemFree_v2,cuMemFree; before=0x1da450,0x2fe540,0x4a7e90,0xd4dc50; after=0x1da450; outside_wrappers=0 verified +0x001da880 impl_cuMemFreeHost high group=cuMemFreeHost variants=cuMemFreeHost; before=0x1da880,0x2fe540,0x4a7e90,0xd4dc50; after=0x1da880; outside_wrappers=0 verified +0x001daad0 impl_cuMemAllocHost high group=cuMemAllocHost variants=cuMemAllocHost_v2,cuMemAllocHost; before=0x1daad0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1daad0; outside_wrappers=1 unverified +0x001dae90 impl_cuMemPrefetchAsync high group=cuMemPrefetchAsync variants=cuMemPrefetchAsync_ptsz,cuMemPrefetchAsync_v2_ptsz,cuMemPrefetchAsync,cuMemPrefetchAsync_v2; before=0x1dae90,0x2fe540,0x4a7e90,0xd4dc50; after=0x1dae90; outside_wrappers=0 unverified +0x001db0f0 impl_cuMemAdvise high group=cuMemAdvise variants=cuMemAdvise,cuMemAdvise_v2; before=0x1db0f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1db0f0; outside_wrappers=0 unverified +0x001dbdd0 impl_cuPointerGetAttributes medium group=cuPointerGetAttributes variants=cuPointerGetAttributes; before=0x1db310,0x1dbdd0,0x2bb300,0x2fe540,0x3fe100,0x402500,0x4a7e90,0xd4dc50; after=0x1dbdd0,0x1db310; outside_wrappers=0 unverified +0x001dc060 impl_cuMemHostGetDevicePointer medium group=cuMemHostGetDevicePointer variants=cuMemHostGetDevicePointer_v2,cuMemHostGetDevicePointer; before=0x1792f0,0x179350,0x1dc060,0x2bb300,0x2fe540,0x3ac3f0,0x4026d0,0x403470,0x4a7e90,0xd4dc50; after=0x1dc060,0x4026d0; outside_wrappers=0 unverified +0x001dc310 impl_cuMemHostRegister high group=cuMemHostRegister variants=cuMemHostRegister_v2,cuMemHostRegister; before=0x1dc310,0x2fe540,0x4a7e90,0xd4dc50; after=0x1dc310; outside_wrappers=0 unverified +0x001dc640 impl_cuIpcOpenMemHandle medium group=cuIpcOpenMemHandle variants=cuIpcOpenMemHandle_v2,cuIpcOpenMemHandle; before=0x1dc640,0x2bb420,0x2c1d40,0x2d8f40,0x2e7bf0,0x2fe540,0x400550,0x403330,0x4a7e90,0xd4dc50,0xd51980; after=0x1dc640,0x403330; outside_wrappers=0 unverified +0x001dc7c0 impl_cuPointerSetAttribute medium group=cuPointerSetAttribute variants=cuPointerSetAttribute; before=0x1dc7c0,0x1dcc50,0x2fde50,0x2fe230,0x2fe540,0x3fe100,0x402500,0x4a7e90,0xd4dc50; after=0x1dc7c0,0x1dcc50; outside_wrappers=0 unverified +0x001dc9b0 impl_cuPointerGetAttribute high group=cuPointerGetAttribute variants=cuPointerGetAttribute; before=0x1dc9b0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1dc9b0; outside_wrappers=0 unverified +0x001dfbe0 impl_cuMemSetAccess medium group=cuMemSetAccess variants=cuMemSetAccess; before=0x1dfbe0,0x1e0420,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1dfbe0,0x1e0420; outside_wrappers=0 unverified +0x001e0640 impl_cuMemGetAccess medium group=cuMemGetAccess variants=cuMemGetAccess; before=0x1e0640,0x1e0950,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1e0640,0x1e0950; outside_wrappers=0 unverified +0x001e0b40 impl_cuMemAllocFromPoolAsync high group=cuMemAllocFromPoolAsync variants=cuMemAllocFromPoolAsync_ptsz,cuMemAllocFromPoolAsync; before=0x1e0b40,0x2f16c0,0x2fe230,0x2fe540,0x3c7f40,0x4a7e90,0xd4dc50; after=0x1e0b40; outside_wrappers=0 unverified +0x001e0cb0 impl_cuMemCreate high group=cuMemCreate variants=cuMemCreate; before=0x1e0cb0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1e0cb0; outside_wrappers=0 unverified +0x001e18f0 impl_cuMemMap high group=cuMemMap variants=cuMemMap; before=0x1e18f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1e18f0; outside_wrappers=0 unverified +0x001e9ab0 impl_cuModuleLoad high group=cuModuleLoad variants=cuModuleLoad; before=0x1e9ab0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1e9ab0; outside_wrappers=0 verified +0x001e9d80 impl_cuModuleLoadData high group=cuModuleLoadData variants=cuModuleLoadData; before=0x1e9d80,0x2fe540,0x4a7e90,0xd4dc50; after=0x1e9d80; outside_wrappers=0 verified +0x001ea020 impl_cuModuleLoadFatBinary high group=cuModuleLoadFatBinary variants=cuModuleLoadFatBinary; before=0x1ea020,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ea020; outside_wrappers=0 unverified +0x001ea370 impl_cuTensorMapEncodeTiled high group=cuTensorMapEncodeTiled variants=cuTensorMapEncodeTiled; before=0x1ea370,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ea370; outside_wrappers=0 unverified +0x001ea530 impl_cuTensorMapEncodeIm2col high group=cuTensorMapEncodeIm2col variants=cuTensorMapEncodeIm2col; before=0x1ea530,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ea530; outside_wrappers=0 unverified +0x001ea6f0 impl_cuTensorMapEncodeIm2colWide high group=cuTensorMapEncodeIm2colWide variants=cuTensorMapEncodeIm2colWide; before=0x1ea6f0,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ea6f0; outside_wrappers=0 unverified +0x001ea8f0 impl_cuTensorMapReplaceAddress high group=cuTensorMapReplaceAddress variants=cuTensorMapReplaceAddress; before=0x1798f0,0x1ea8f0,0x2fe440,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ea8f0; outside_wrappers=0 unverified +0x001ea9d0 impl_cuModuleLoadDataEx high group=cuModuleLoadDataEx variants=cuModuleLoadDataEx; before=0x1ea9d0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ea9d0; outside_wrappers=0 unverified +0x001eb4d0 impl_cuSurfRefGetArray high group=cuSurfRefGetArray variants=cuSurfRefGetArray; before=0x1eb4d0,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x1eb4d0; outside_wrappers=0 unverified +0x001eb960 impl_cuModuleGetTexRef medium group=cuModuleGetTexRef variants=cuModuleGetTexRef; before=0x179980,0x1eb960,0x2baed0,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1eb960,0x179980; outside_wrappers=0 unverified +0x001ebb60 impl_cuModuleGetSurfRef medium group=cuModuleGetSurfRef variants=cuModuleGetSurfRef; before=0x179980,0x1ebb60,0x2baed0,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ebb60,0x179980; outside_wrappers=0 unverified +0x001ec6e0 impl_cuParamSetSize high group=cuParamSetSize variants=cuParamSetSize; before=0x1ec6e0,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ec6e0; outside_wrappers=0 unverified +0x001ec8b0 impl_cuParamSetTexRef high group=cuParamSetTexRef variants=cuParamSetTexRef; before=0x1ec8b0,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x1ec8b0; outside_wrappers=0 unverified +0x001eca70 impl_cuOccupancyAvailableDynamicSMemPerBlock high group=cuOccupancyAvailableDynamicSMemPerBlock variants=cuOccupancyAvailableDynamicSMemPerBlock; before=0x1eca70,0x2fe540,0x4a7e90,0xd4dc50; after=0x1eca70; outside_wrappers=0 unverified +0x001ecf80 impl_cuParamSetv high group=cuParamSetv variants=cuParamSetv; before=0x179c00,0x1ecf80,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ecf80; outside_wrappers=0 unverified +0x001ed180 impl_cuParamSeti high group=cuParamSeti variants=cuParamSeti; before=0x1ed180,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ed180; outside_wrappers=0 unverified +0x001ed340 impl_cuParamSetf high group=cuParamSetf variants=cuParamSetf; before=0x1ed340,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1ed340; outside_wrappers=0 unverified +0x001ed5c0 impl_cuProfilerStart medium group=cuProfilerStart variants=cuProfilerStart; before=0x1798f0,0x1ed5c0,0x2fe230,0x2fe440,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50; after=0x1ed5c0; outside_wrappers=1 unverified +0x001ed5c0 impl_cuProfilerStop medium group=cuProfilerStop variants=cuProfilerStop; before=0x1798f0,0x1ed5c0,0x2fe230,0x2fe440,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50; after=0x1ed5c0; outside_wrappers=1 unverified +0x001ed720 impl_cuSurfObjectGetResourceDesc high group=cuSurfObjectGetResourceDesc variants=cuSurfObjectGetResourceDesc; before=0x1ed720,0x2d8f40,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1ed720; outside_wrappers=0 unverified +0x001ed850 impl_cuTexObjectGetResourceDesc high group=cuTexObjectGetResourceDesc variants=cuTexObjectGetResourceDesc; before=0x1ed850,0x2d8f40,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1ed850; outside_wrappers=0 unverified +0x001eda80 impl_cuTexObjectGetTextureDesc high group=cuTexObjectGetTextureDesc variants=cuTexObjectGetTextureDesc; before=0x1eda80,0x2d8f40,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1eda80; outside_wrappers=0 unverified +0x001edc80 impl_cuTexObjectGetResourceViewDesc high group=cuTexObjectGetResourceViewDesc variants=cuTexObjectGetResourceViewDesc; before=0x1edc80,0x2d8f40,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1edc80; outside_wrappers=0 unverified +0x001f0fd0 impl_cuSurfObjectDestroy medium group=cuSurfObjectDestroy variants=cuSurfObjectDestroy; before=0x1f0fd0,0x2d8f40,0x2f9890,0x2fe230,0x2fe540,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1f0fd0,0x2f9890; outside_wrappers=0 unverified +0x001f11d0 impl_cuSurfRefSetArray high group=cuSurfRefSetArray variants=cuSurfRefSetArray; before=0x1f11d0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f11d0; outside_wrappers=0 unverified +0x001f1540 impl_cuSurfObjectCreate high group=cuSurfObjectCreate variants=cuSurfObjectCreate; before=0x1f1540,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f1540; outside_wrappers=0 unverified +0x001f2970 impl_cuTexObjectCreate high group=cuTexObjectCreate variants=cuTexObjectCreate; before=0x1f2970,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f2970; outside_wrappers=0 unverified +0x001f35d0 impl_cuTexObjectDestroy medium group=cuTexObjectDestroy variants=cuTexObjectDestroy; before=0x1f35d0,0x2d8f40,0x2fbad0,0x2fe230,0x2fe540,0x4a6e10,0x4a7e90,0xd4db70,0xd4dc50,0xd51980; after=0x1f35d0,0x4a6e10,0x2fbad0; outside_wrappers=0 unverified +0x001f3830 impl_cuTexRefDestroy medium group=cuTexRefDestroy variants=cuTexRefDestroy; before=0x1798f0,0x1f3830,0x2d8f40,0x2fbad0,0x2fe440,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f3830,0x2fbad0; outside_wrappers=0 unverified +0x001f3990 impl_cuTexRefSetArray high group=cuTexRefSetArray variants=cuTexRefSetArray; before=0x1f3990,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f3990; outside_wrappers=0 unverified +0x001f3ca0 impl_cuTexRefSetMipmappedArray high group=cuTexRefSetMipmappedArray variants=cuTexRefSetMipmappedArray; before=0x1f3ca0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f3ca0; outside_wrappers=0 unverified +0x001f4000 impl_cuTexRefSetFormat high group=cuTexRefSetFormat variants=cuTexRefSetFormat; before=0x1f4000,0x2583d0,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4000; outside_wrappers=0 unverified +0x001f4200 impl_cuTexRefSetAddressMode high group=cuTexRefSetAddressMode variants=cuTexRefSetAddressMode; before=0x1f4200,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4200; outside_wrappers=0 unverified +0x001f4370 impl_cuTexRefSetFilterMode high group=cuTexRefSetFilterMode variants=cuTexRefSetFilterMode; before=0x1f4370,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4370; outside_wrappers=0 unverified +0x001f44b0 impl_cuTexRefSetBorderColor high group=cuTexRefSetBorderColor variants=cuTexRefSetBorderColor; before=0x1f44b0,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f44b0; outside_wrappers=0 unverified +0x001f4620 impl_cuTexRefSetFlags high group=cuTexRefSetFlags variants=cuTexRefSetFlags; before=0x1f4620,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4620; outside_wrappers=0 unverified +0x001f4770 impl_cuTexRefSetMipmapFilterMode high group=cuTexRefSetMipmapFilterMode variants=cuTexRefSetMipmapFilterMode; before=0x1f4770,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4770; outside_wrappers=0 unverified +0x001f48b0 impl_cuTexRefSetMipmapLevelBias high group=cuTexRefSetMipmapLevelBias variants=cuTexRefSetMipmapLevelBias; before=0x1f48b0,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f48b0; outside_wrappers=0 unverified +0x001f49f0 impl_cuTexRefSetMipmapLevelClamp high group=cuTexRefSetMipmapLevelClamp variants=cuTexRefSetMipmapLevelClamp; before=0x1f49f0,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f49f0; outside_wrappers=0 unverified +0x001f4b60 impl_cuTexRefSetMaxAnisotropy high group=cuTexRefSetMaxAnisotropy variants=cuTexRefSetMaxAnisotropy; before=0x1f4b60,0x2d8f40,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50,0xd51980; after=0x1f4b60; outside_wrappers=0 unverified +0x001f4f60 impl_cuTexRefGetAddressMode high group=cuTexRefGetAddressMode variants=cuTexRefGetAddressMode; before=0x1798f0,0x1f4f60,0x2fe440,0x2fe540,0x4a7e90,0xd11b90,0xd4dc50; after=0x1f4f60; outside_wrappers=0 unverified +0x001f5c20 impl_cuGraphicsVDPAURegisterVideoSurface high group=cuGraphicsVDPAURegisterVideoSurface variants=cuGraphicsVDPAURegisterVideoSurface; before=0x179200,0x1799e0,0x1f5c20,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1f5c20; outside_wrappers=0 unverified +0x001f5e10 impl_cuGraphicsVDPAURegisterOutputSurface high group=cuGraphicsVDPAURegisterOutputSurface variants=cuGraphicsVDPAURegisterOutputSurface; before=0x179200,0x1799e0,0x1f5e10,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x1f5e10; outside_wrappers=0 unverified +0x001f8c90 impl_cuMemPoolDestroy medium group=cuMemPoolDestroy variants=cuMemPoolDestroy; before=0x1798f0,0x1d5910,0x1f8c90,0x2fe230,0x2fe440,0x2fe540,0x4a7e90,0xd4dc50; after=0x1f8c90,0x1d5910; outside_wrappers=0 unverified +0x001f9780 impl_cuMulticastUnbind high group=cuMulticastUnbind variants=cuMulticastUnbind; before=0x1f94c0,0x1f9780,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1f9780; outside_wrappers=0 unverified +0x001f9950 impl_cuMulticastAddDevice high group=cuMulticastAddDevice variants=cuMulticastAddDevice; before=0x1f94c0,0x1f9950,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1f9950; outside_wrappers=0 unverified +0x001f9a90 impl_cuMulticastGetGranularity high group=cuMulticastGetGranularity variants=cuMulticastGetGranularity; before=0x1f94c0,0x1f9a90,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1f9a90; outside_wrappers=0 unverified +0x001f9f70 impl_cuOccupancyMaxActiveBlocksPerMultiprocessor medium group=cuOccupancyMaxActiveBlocksPerMultiprocessor variants=cuOccupancyMaxActiveBlocksPerMultiprocessor; before=0x1f9e20,0x1f9f70,0x2bb420,0x2dc100,0x2e4340,0x2fe540,0x4a7e90,0xd4bd20,0xd4dc50; after=0x1f9f70,0x2e4340; outside_wrappers=1 unverified +0x001f9f70 impl_cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags medium group=cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags variants=cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags; before=0x1f9e20,0x1f9f70,0x2bb420,0x2dc100,0x2e4340,0x2fe540,0x4a7e90,0xd4bd20,0xd4dc50; after=0x1f9f70,0x2e4340; outside_wrappers=1 unverified +0x001fa330 impl_cuOccupancyMaxPotentialBlockSize medium group=cuOccupancyMaxPotentialBlockSize variants=cuOccupancyMaxPotentialBlockSize; before=0x1f9e20,0x1fa330,0x2baed0,0x2d8f40,0x2e4570,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1fa330,0x2e4570; outside_wrappers=1 unverified +0x001fa330 impl_cuOccupancyMaxPotentialBlockSizeWithFlags medium group=cuOccupancyMaxPotentialBlockSizeWithFlags variants=cuOccupancyMaxPotentialBlockSizeWithFlags; before=0x1f9e20,0x1fa330,0x2baed0,0x2d8f40,0x2e4570,0x2fde50,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x1fa330,0x2e4570; outside_wrappers=1 unverified +0x001fa6d0 impl_cuOccupancyMaxPotentialClusterSize high group=cuOccupancyMaxPotentialClusterSize variants=cuOccupancyMaxPotentialClusterSize; before=0x1fa6d0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1fa6d0; outside_wrappers=0 unverified +0x001faa00 impl_cuOccupancyMaxActiveClusters high group=cuOccupancyMaxActiveClusters variants=cuOccupancyMaxActiveClusters; before=0x1faa00,0x2fe540,0x4a7e90,0xd4dc50; after=0x1faa00; outside_wrappers=0 unverified +0x001facd0 impl_cuStreamCreate medium group=cuStreamCreate variants=cuStreamCreate; before=0x1facd0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1facd0; outside_wrappers=1 unverified +0x001facd0 impl_cuStreamCreateWithPriority medium group=cuStreamCreateWithPriority variants=cuStreamCreateWithPriority; before=0x1facd0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1facd0; outside_wrappers=1 unverified +0x001fb3d0 impl_cuStreamWaitEvent high group=cuStreamWaitEvent variants=cuStreamWaitEvent_ptsz,cuStreamWaitEvent; before=0x1792f0,0x179350,0x1fb3d0,0x2a2b90,0x2a50f0,0x2baed0,0x2ed7d0,0x2fe540,0x3abad0,0x4a7e90,0xd1d500,0xd1e090,0xd4dc50; after=0x1fb3d0; outside_wrappers=0 unverified +0x001fb7f0 impl_cuStreamGetCaptureInfo high group=cuStreamGetCaptureInfo variants=cuStreamGetCaptureInfo_v2_ptsz,cuStreamGetCaptureInfo_v3_ptsz,cuStreamGetCaptureInfo,cuStreamGetCaptureInfo_ptsz,cuStreamGetCaptureInfo_v2,cuStreamGetCaptureInfo_v3; before=0x1fb7f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1fb7f0; outside_wrappers=0 unverified +0x001fbe30 impl_cuStreamBeginCaptureToGraph high group=cuStreamBeginCaptureToGraph variants=cuStreamBeginCaptureToGraph_ptsz,cuStreamBeginCaptureToGraph; before=0x1fba70,0x1fbe30,0x2d8f40,0x2eff60,0x2f16c0,0x2fe540,0x4a7e90,0xd1e090,0xd4dc50,0xd51980; after=0x1fbe30; outside_wrappers=0 unverified +0x001fd190 impl_cuStreamCopyAttributes high group=cuStreamCopyAttributes variants=cuStreamCopyAttributes_ptsz,cuStreamCopyAttributes; before=0x1fd190,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50; after=0x1fd190; outside_wrappers=0 unverified +0x001fd2e0 impl_cuStreamGetAttribute high group=cuStreamGetAttribute variants=cuStreamGetAttribute_ptsz,cuStreamGetAttribute; before=0x1fd2e0,0x2d8f40,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50,0xd51980; after=0x1fd2e0; outside_wrappers=0 unverified +0x001fd460 impl_cuVDPAUCtxCreate high group=cuVDPAUCtxCreate variants=cuVDPAUCtxCreate_v2,cuVDPAUCtxCreate; before=0x1fd460,0x2689e0,0x26d8e0,0x277680,0x2fe020,0x2fe540,0x4a7e90,0xd4dc50; after=0x1fd460; outside_wrappers=0 unverified +0x001ff270 impl_cuMulticastBindMem medium group=cuMulticastBindMem variants=cuMulticastBindMem; before=0x1f94c0,0x1ff270,0x2fde40,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x1ff270,0x2fde40; outside_wrappers=0 unverified +0x001ff3e0 impl_cuMulticastBindAddr high group=cuMulticastBindAddr variants=cuMulticastBindAddr; before=0x1ff3e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x1ff3e0; outside_wrappers=0 unverified +0x0024a6d0 impl_cuCoredumpSetAttributeGlobal high group=cuCoredumpSetAttributeGlobal variants=cuCoredumpSetAttributeGlobal; before=0x24a4d0,0x24a6d0,0x2fe540,0x4a7e90,0xd4dc50; after=0x24a6d0; outside_wrappers=0 unverified +0x00259480 impl_cuTexRefSetAddress2D high group=cuTexRefSetAddress2D variants=cuTexRefSetAddress2D_v3,cuTexRefSetAddress2D_v2,cuTexRefSetAddress2D; before=0x259480,0x2baed0,0x2d8f40,0x2fbe80,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x259480; outside_wrappers=0 unverified +0x00261f70 impl_cuCheckpointProcessGetRestoreThreadId high group=cuCheckpointProcessGetRestoreThreadId variants=cuCheckpointProcessGetRestoreThreadId; before=0x261f70,0x2fe540,0x4a7e90,0xd4dc50; after=0x261f70; outside_wrappers=0 unverified +0x00262690 impl_cuCheckpointProcessGetState medium group=cuCheckpointProcessGetState variants=cuCheckpointProcessGetState; before=0x179200,0x2623e0,0x262690,0x2fe540,0x30a870,0x4a7e90,0xd4dc50; after=0x262690,0x2623e0,0x30a870; outside_wrappers=0 unverified +0x00266fc0 impl_cuCoredumpSetAttribute high group=cuCoredumpSetAttribute variants=cuCoredumpSetAttribute; before=0x266fc0,0x2fe540,0x4a7e90,0xd4dc50; after=0x266fc0; outside_wrappers=0 unverified +0x0026a9b0 impl_cuCtxSynchronize medium group=cuCtxSynchronize variants=cuCtxSynchronize; before=0x1ad290,0x26a9b0,0x2baed0,0x2fde50,0x2fe540,0x4a7e90,0xd1e3f0,0xd4dc50; after=0x26a9b0,0x1ad290; outside_wrappers=0 unverified +0x0026ec30 impl_cuEGLStreamProducerReturnFrame medium group=cuEGLStreamProducerReturnFrame variants=cuEGLStreamProducerReturnFrame; before=0x1b0db0,0x26ec30,0x2fde50,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50,0xd6e3d0; after=0x26ec30,0x1b0db0,0xd6e3d0; outside_wrappers=0 unverified +0x0026fa30 impl_cuEGLStreamConsumerReleaseFrame medium group=cuEGLStreamConsumerReleaseFrame variants=cuEGLStreamConsumerReleaseFrame; before=0x1b07f0,0x26e7c0,0x26fa30,0x2fde50,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x26fa30,0x26e7c0,0x1b07f0; outside_wrappers=0 unverified +0x00287bb0 impl_cuGetErrorName medium group=cuGetErrorName variants=cuGetErrorName; before=0x287bb0,0x2fe540,0x4a7e90,0xd4dc50; after=0x287bb0; outside_wrappers=1 unverified +0x00287bb0 impl_cuGetErrorString medium group=cuGetErrorString variants=cuGetErrorString; before=0x287bb0,0x2fe540,0x4a7e90,0xd4dc50; after=0x287bb0; outside_wrappers=1 unverified +0x00287cb0 impl_cuEventElapsedTime medium group=cuEventElapsedTime variants=cuEventElapsedTime,cuEventElapsedTime_v2; before=0x287cb0,0x288ba0,0x2bb140,0x2fe540,0x487640,0x4a7e90,0xd4dc50; after=0x287cb0,0x487640,0x288ba0; outside_wrappers=0 unverified +0x00287ed0 impl_cuEventCreate medium group=cuEventCreate variants=cuEventCreate; before=0x1b0fb0,0x287ed0,0x2bb420,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x287ed0,0x1b0fb0; outside_wrappers=0 unverified +0x00288cd0 impl_cuEventSynchronize medium group=cuEventSynchronize variants=cuEventSynchronize; before=0x1b2600,0x288cd0,0x2bb140,0x2fe230,0x2fe540,0x3034c0,0x4a7e90,0xd1e090,0xd4dc50; after=0x288cd0,0x1b2600,0x3034c0; outside_wrappers=0 unverified +0x002891d0 impl_cuEventDestroy high group=cuEventDestroy variants=cuEventDestroy_v2,cuEventDestroy; before=0x2891d0,0x2baed0,0x2bb140,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x2891d0; outside_wrappers=0 verified +0x00294db0 impl_cuGLGetDevices medium group=cuGLGetDevices variants=cuGLGetDevices_v2,cuGLGetDevices; before=0x294db0,0x29d320,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x294db0,0x29d320; outside_wrappers=0 unverified +0x00294fd0 impl_cuGLInit medium group=cuGLInit variants=cuGLInit; before=0x294fd0,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50; after=0x294fd0; outside_wrappers=1 unverified +0x0029d480 impl_cuGraphCreate high group=cuGraphCreate variants=cuGraphCreate; before=0x29d480,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x29d480; outside_wrappers=0 unverified +0x002a1b10 impl_cuGraphAddMemAllocNode medium group=cuGraphAddMemAllocNode variants=cuGraphAddMemAllocNode; before=0x1baa70,0x2a1930,0x2a1b10,0x2a5430,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2a1b10,0x2a5430,0x1baa70,0x2a1930; outside_wrappers=0 unverified +0x002a3530 impl_cuGraphDestroy medium group=cuGraphDestroy variants=cuGraphDestroy; before=0x2a3530,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2a3530; outside_wrappers=1 unverified +0x002a5a20 impl_cuGraphAddMemFreeNode medium group=cuGraphAddMemFreeNode variants=cuGraphAddMemFreeNode; before=0x1c1db0,0x2a5a20,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2a5a20,0x1c1db0; outside_wrappers=0 unverified +0x002a66e0 impl_cuGraphAddBatchMemOpNode medium group=cuGraphAddBatchMemOpNode variants=cuGraphAddBatchMemOpNode; before=0x1c0740,0x1c0bc0,0x2a66e0,0x2bb140,0x2fe540,0x4a7e90,0xd4dc50; after=0x2a66e0,0x1c0bc0; outside_wrappers=0 unverified +0x002a6bf0 impl_cuGraphDebugDotPrint high group=cuGraphDebugDotPrint variants=cuGraphDebugDotPrint; before=0x2a6bf0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2a6bf0; outside_wrappers=0 verified +0x002aa0a0 impl_cuGraphAddExternalSemaphoresSignalNode medium group=cuGraphAddExternalSemaphoresSignalNode variants=cuGraphAddExternalSemaphoresSignalNode; before=0x1b9cc0,0x2aa0a0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2aa0a0,0x1b9cc0; outside_wrappers=0 unverified +0x002aa2d0 impl_cuGraphAddExternalSemaphoresWaitNode medium group=cuGraphAddExternalSemaphoresWaitNode variants=cuGraphAddExternalSemaphoresWaitNode; before=0x1ba0d0,0x2aa2d0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2aa2d0,0x1ba0d0; outside_wrappers=0 unverified +0x002aabc0 impl_cuGraphClone medium group=cuGraphClone variants=cuGraphClone; before=0x1ba8d0,0x2aabc0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2aabc0,0x1ba8d0; outside_wrappers=0 unverified +0x002ac120 impl_cuGraphExecBatchMemOpNodeSetParams medium group=cuGraphExecBatchMemOpNodeSetParams variants=cuGraphExecBatchMemOpNodeSetParams; before=0x179200,0x179a80,0x179c00,0x1c0740,0x1c0e20,0x2ac120,0x2baed0,0x2fe540,0x4a7e90,0xd183b0,0xd4dc50; after=0x2ac120,0x1c0e20; outside_wrappers=0 unverified +0x002aca10 impl_cuGraphKernelNodeSetParams high group=cuGraphKernelNodeSetParams variants=cuGraphKernelNodeSetParams_v2,cuGraphKernelNodeSetParams; before=0x1c2b10,0x2aca10,0x2fe540,0x4a7e90,0xd4dc50; after=0x2aca10; outside_wrappers=0 unverified +0x002afcf0 impl_cuGraphNodeSetEnabled high group=cuGraphNodeSetEnabled variants=cuGraphNodeSetEnabled; before=0x2afcf0,0x2fe540,0x4a7e90,0xd183b0,0xd4dc50; after=0x2afcf0; outside_wrappers=0 unverified +0x002b01c0 impl_cuGraphExecKernelNodeSetParams medium group=cuGraphExecKernelNodeSetParams variants=cuGraphExecKernelNodeSetParams_v2,cuGraphExecKernelNodeSetParams; before=0x1c2b10,0x2b01c0,0x2fe540,0x4a7e90,0xd183b0,0xd4dc50,0xd6a9f0; after=0x2b01c0,0xd6a9f0; outside_wrappers=0 unverified +0x002b4f80 impl_cuGraphExecChildGraphNodeSetParams medium group=cuGraphExecChildGraphNodeSetParams variants=cuGraphExecChildGraphNodeSetParams; before=0x2ac1b0,0x2b4f80,0x2fe540,0x4a7e90,0xd183b0,0xd4dc50; after=0x2b4f80,0x2ac1b0; outside_wrappers=1 unverified +0x002b9ff0 impl_cuGreenCtxDestroy medium group=cuGreenCtxDestroy variants=cuGreenCtxDestroy; before=0x1792f0,0x179350,0x1c6be0,0x26d680,0x2b9ff0,0x2bb140,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2b9ff0,0x1c6be0; outside_wrappers=0 unverified +0x002ba3e0 impl_cuGreenCtxCreate medium group=cuGreenCtxCreate variants=cuGreenCtxCreate; before=0x1792f0,0x179350,0x1c6830,0x26d8a0,0x277570,0x2ba3e0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2ba3e0,0x1c6830,0x277570,0x26d8a0; outside_wrappers=0 unverified +0x002bdf80 impl_cuLinkDestroy high group=cuLinkDestroy variants=cuLinkDestroy; before=0x2bdf80,0x2fe540,0x4a7e90,0xd4dc50; after=0x2bdf80; outside_wrappers=0 unverified +0x002bf2f0 impl_cuKernelGetAttribute high group=cuKernelGetAttribute variants=cuKernelGetAttribute; before=0x2bb420,0x2bf2f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2bf2f0; outside_wrappers=0 unverified +0x002bfd50 impl_cuKernelSetAttribute medium group=cuKernelSetAttribute variants=cuKernelSetAttribute; before=0x179730,0x1ca190,0x2bb420,0x2bfd50,0x2d8f40,0x2fe540,0x4a7e90,0xd4dc50,0xd51f50; after=0x2bfd50,0x1ca190,0xd51f50,0x179730; outside_wrappers=0 unverified +0x002c05e0 impl_cuCtxSetLimit medium group=cuCtxSetLimit variants=cuCtxSetLimit; before=0x1ca690,0x2c05e0,0x2d8f40,0x2fe230,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x2c05e0,0x1ca690; outside_wrappers=0 unverified +0x002c1040 impl_cuDeviceGetP2PAttribute medium group=cuDeviceGetP2PAttribute variants=cuDeviceGetP2PAttribute; before=0x1af7e0,0x2c1040,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2c1040,0x1af7e0; outside_wrappers=0 unverified +0x002c2890 impl_cuLibraryUnload medium group=cuLibraryUnload variants=cuLibraryUnload; before=0x1c9e80,0x2c2890,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x2c2890,0x1c9e80; outside_wrappers=0 unverified +0x002c4050 impl_cuMemPoolCreate medium group=cuMemPoolCreate variants=cuMemPoolCreate; before=0x1d5780,0x2c4050,0x2c6030,0x2cee80,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50; after=0x2c4050,0x1d5780,0x2cee80,0x2c6030; outside_wrappers=0 unverified +0x002c71c0 impl_cuDeviceGraphMemTrim high group=cuDeviceGraphMemTrim variants=cuDeviceGraphMemTrim; before=0x2c71c0,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2c71c0; outside_wrappers=0 unverified +0x002cb940 impl_cuInit medium group=cuInit variants=cuInit; before=0x2ba8f0,0x2cb940,0x2fe540,0x4a7e90,0xd4dc50; after=0x2cb940,0x2ba8f0; outside_wrappers=0 unverified +0x002ce4e0 impl_cuLinkCreate high group=cuLinkCreate variants=cuLinkCreate_v2,cuLinkCreate; before=0x2bb420,0x2bbd30,0x2ce4e0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2ce4e0; outside_wrappers=0 verified +0x002d09f0 impl_cuLinkComplete high group=cuLinkComplete variants=cuLinkComplete; before=0x2d09f0,0x2fe540,0x4a7e90,0xd4dc50; after=0x2d09f0; outside_wrappers=0 unverified +0x002d3f60 impl_cuLibraryLoadData medium group=cuLibraryLoadData variants=cuLibraryLoadData; before=0x2bbd30,0x2d3f60,0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x2d3f60; outside_wrappers=1 unverified +0x002d81b0 impl_cuMemBatchDecompressAsync medium group=cuMemBatchDecompressAsync variants=cuMemBatchDecompressAsync_ptsz,cuMemBatchDecompressAsync; before=0x26bc10,0x2d8060,0x2d81b0,0x2f8e10,0x2f8fd0,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50,0xd73560; after=0x2d81b0,0xd73560,0x2d8060; outside_wrappers=0 unverified +0x002e84d0 impl_cuModuleUnload medium group=cuModuleUnload variants=cuModuleUnload; before=0x1eb2e0,0x2baed0,0x2d8f40,0x2e84d0,0x2fe230,0x2fe540,0x301460,0x4a7e90,0xd4dc50,0xd4f3e0,0xd51980; after=0x2e84d0,0x1eb2e0,0x301460; outside_wrappers=0 unverified +0x002eb720 impl_cuMemMapArrayAsync medium group=cuMemMapArrayAsync variants=cuMemMapArrayAsync_ptsz,cuMemMapArrayAsync; before=0x1dfae0,0x2d8f40,0x2eb720,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50,0xd51980; after=0x2eb720,0x1dfae0; outside_wrappers=0 unverified +0x002ec4a0 impl_cuStreamAttachMemAsync medium group=cuStreamAttachMemAsync variants=cuStreamAttachMemAsync_ptsz,cuStreamAttachMemAsync; before=0x1fce40,0x2d8f40,0x2ec4a0,0x2f16c0,0x2fe230,0x2fe540,0x4a7e90,0xd4dc50,0xd51980; after=0x2ec4a0,0x1fce40; outside_wrappers=0 unverified +0x002f0110 impl_cuStreamEndCapture medium group=cuStreamEndCapture variants=cuStreamEndCapture_ptsz,cuStreamEndCapture; before=0x1fb6b0,0x2f0110,0x2f16c0,0x2fe540,0x4a7e90,0xd1e090,0xd4dc50; after=0x2f0110,0x1fb6b0; outside_wrappers=0 unverified +0x002f0630 impl_cuStreamDestroy medium group=cuStreamDestroy variants=cuStreamDestroy_v2,cuStreamDestroy; before=0x1792f0,0x179350,0x1fb130,0x2d8f40,0x2f0630,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50,0xd51980; after=0x2f0630,0x1fb130; outside_wrappers=0 unverified +0x002f0fe0 impl_cuGLUnmapBufferObject medium group=cuGLUnmapBufferObject variants=cuGLUnmapBufferObject; before=0x1c1a20,0x2baed0,0x2f0fe0,0x2fde50,0x2fe230,0x2fe540,0x3aa550,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50,0xd663e0; after=0x2f0fe0,0x1c1a20; outside_wrappers=1 unverified +0x002f0fe0 impl_cuGLUnmapBufferObjectAsync medium group=cuGLUnmapBufferObjectAsync variants=cuGLUnmapBufferObjectAsync; before=0x1c1a20,0x2baed0,0x2f0fe0,0x2fde50,0x2fe230,0x2fe540,0x3aa550,0x3c52d0,0x3c5370,0x4a7e90,0xd4dc50,0xd663e0; after=0x2f0fe0,0x1c1a20; outside_wrappers=1 unverified +0x002f1880 impl_cuStreamSetAttribute high group=cuStreamSetAttribute variants=cuStreamSetAttribute_ptsz,cuStreamSetAttribute; before=0x2d8f40,0x2f1880,0x2fe540,0x4a7e90,0xd1d500,0xd4dc50,0xd51980; after=0x2f1880; outside_wrappers=0 unverified +0x002f9f00 impl_cuModuleGetGlobal high group=cuModuleGetGlobal variants=cuModuleGetGlobal_v2,cuModuleGetGlobal; before=0x2baed0,0x2d8f40,0x2f9f00,0x2fe540,0x402500,0x4a7e90,0xd4dc50,0xd51980; after=0x2f9f00; outside_wrappers=1 verified +0x002fb8e0 impl_cuTexRefCreate medium group=cuTexRefCreate variants=cuTexRefCreate; before=0x1f2870,0x2d8f40,0x2fb8e0,0x2fe540,0x4a7e90,0xd4da90,0xd4dc50,0xd51980; after=0x2fb8e0,0x1f2870; outside_wrappers=0 unverified +0x002fffa0 impl_cuMemRangeGetAttribute medium group=cuMemRangeGetAttribute variants=cuMemRangeGetAttribute; before=0x1cb030,0x2fe540,0x2fffa0,0x306160,0x3fe700,0x4a7e90,0xd4bd90,0xd4db70,0xd4dc50; after=0x2fffa0,0x1cb030; outside_wrappers=0 unverified +0x00304420 impl_cuStreamQuery medium group=cuStreamQuery variants=cuStreamQuery_ptsz,cuStreamQuery; before=0x1fcfe0,0x2ea920,0x2f16c0,0x2fe230,0x2fe540,0x304420,0x4a7e90,0xd4dc50; after=0x304420,0x1fcfe0,0x2ea920; outside_wrappers=0 unverified +0x003aa910 impl_cuLibraryEnumerateKernels medium group=cuLibraryEnumerateKernels variants=cuLibraryEnumerateKernels; before=0x2fe540,0x3aa910,0x3aad50,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x3aa910,0x3aad50; outside_wrappers=1 unverified +0x003aa910 impl_cuLibraryGetKernelCount medium group=cuLibraryGetKernelCount variants=cuLibraryGetKernelCount; before=0x2fe540,0x3aa910,0x3aad50,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x3aa910,0x3aad50; outside_wrappers=1 unverified +0x003c9250 impl_cuMemFreeAsync high group=cuMemFreeAsync variants=cuMemFreeAsync_ptsz,cuMemFreeAsync; before=0x1798f0,0x2fe440,0x2fe540,0x3c9250,0x4a7e90,0xd4dc50; after=0x3c9250; outside_wrappers=0 verified +0x003fb7b0 impl_cuMemRelease medium group=cuMemRelease variants=cuMemRelease; before=0x1792f0,0x179350,0x1d2dc0,0x2fe540,0x3fb7b0,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x3fb7b0,0x1d2dc0; outside_wrappers=0 unverified +0x003fcfa0 impl_cuMulticastCreate medium group=cuMulticastCreate variants=cuMulticastCreate; before=0x1792f0,0x179350,0x1f94c0,0x1f9580,0x26bc10,0x2fe540,0x3fcfa0,0x4a7e90,0xd4d9a0,0xd4dc50; after=0x3fcfa0,0x1f9580; outside_wrappers=0 unverified +0x00400ec0 impl_cuMemAlloc high group=cuMemAlloc variants=cuMemAlloc_v2,cuMemAlloc; before=0x1d9cb0,0x2bb300,0x2c1d40,0x2d8f40,0x2e7bf0,0x2fe230,0x2fe540,0x400550,0x400ec0,0x4a7e90,0xd4dc50,0xd51980; after=0x400ec0; outside_wrappers=0 verified +0x00402a30 impl_cuMemGetAddressRange high group=cuMemGetAddressRange variants=cuMemGetAddressRange_v2,cuMemGetAddressRange; before=0x1792f0,0x179350,0x2bb300,0x2fe540,0x3ac3f0,0x402a30,0x403470,0x4a7e90,0xd4bd90,0xd4dc50; after=0x402a30; outside_wrappers=0 verified +0x00d11d90 impl_cuStreamBeginCapture high group=cuStreamBeginCapture variants=cuStreamBeginCapture_v2_ptsz,cuStreamBeginCapture,cuStreamBeginCapture_ptsz,cuStreamBeginCapture_v2; before=0x2d8f40,0x2eff60,0x2f16c0,0x2fe540,0x4a7e90,0xd11d90,0xd4dc50,0xd51980; after=0xd11d90; outside_wrappers=0 verified +0x00d11e90 impl_cuStreamSynchronize medium group=cuStreamSynchronize variants=cuStreamSynchronize_ptsz,cuStreamSynchronize; before=0x1fb280,0x2f16c0,0x2fe230,0x2fe540,0x4a7e90,0xd11e90,0xd4dc50; after=0xd11e90,0x1fb280; outside_wrappers=0 unverified +0x00d121d0 impl_cuGetProcAddress high group=cuGetProcAddress variants=cuGetProcAddress_v2,cuGetProcAddress; before=0x2fe540,0x4a7e90,0xd121d0,0xd4dc50; after=0xd121d0; outside_wrappers=0 unverified +0x00d180a0 impl_cuGraphAddChildGraphNode medium group=cuGraphAddChildGraphNode variants=cuGraphAddChildGraphNode; before=0x1ba7c0,0x2fe540,0x4a7e90,0xd180a0,0xd4d9a0,0xd4dc50; after=0xd180a0,0x1ba7c0; outside_wrappers=0 unverified +0x00d20520 impl_cuIpcOpenEventHandle medium group=cuIpcOpenEventHandle variants=cuIpcOpenEventHandle; before=0x179200,0x1799e0,0x1b3530,0x2bb420,0x2d8f40,0x2da1f0,0x2fe540,0x483550,0x4a7e90,0xd20520,0xd4dc50,0xd4e540,0xd51980; after=0xd20520,0x1b3530,0xd4e540,0x483550,0x2da1f0; outside_wrappers=0 unverified +0x00d668f0 impl_cuLibraryGetUnifiedFunction high group=cuLibraryGetUnifiedFunction variants=cuLibraryGetUnifiedFunction; before=0x2fe540,0x4a7e90,0xd4d9a0,0xd4dc50,0xd668f0; after=0xd668f0; outside_wrappers=0 verified diff --git a/userspace/cuda/570.144/annotations/names.tsv b/userspace/cuda/570.144/annotations/names.tsv index 403bef3..e668055 100644 --- a/userspace/cuda/570.144/annotations/names.tsv +++ b/userspace/cuda/570.144/annotations/names.tsv @@ -1302,3 +1302,42 @@ 0x4676f0 cudaEngineContextCreateHook verified 0x481360 rmEngineContextAlloc verified 0x453cc0 cudaEngineContextCleanup verified +# --- engine-record path verified in the 2026-09 type-recovery pass (see docs/05 section 3.7) --- +0x450ff0 cudaDeviceBuildEngineRecords verified +0x462770 cudaRmGetEngines verified +0x44a170 cudaDeviceSupportsClass verified +0x458040 cudaDeviceAbiTableInit verified +0x44ac20 cudaDeviceInitCaps verified +# +# --- wrapper callees verified in the 2026-09 name-candidate pass --- +# Method: tools/name_candidates.py; per-entry evidence and confidence in +# name-candidates.tsv; roles/evidence also recorded in subsystems.tsv. +0x1a36d0 impl_cuCtxGetApiVersion wrapper-callee +0x1a4900 impl_cuArrayGetMemoryRequirements wrapper-callee +0x1a5410 impl_cuArrayGetPlane wrapper-callee +0x1a5820 impl_cuDevicePrimaryCtxGetState wrapper-callee +0x1a6150 impl_cuDeviceRegisterAsyncNotification wrapper-callee +0x1a6430 impl_cuCtxAttach wrapper-callee +0x1a66f0 impl_cuCtxGetCurrent wrapper-callee +0x1a6850 impl_cuCtxGetDevice wrapper-callee +0x1a6960 impl_cuCtxGetFlags wrapper-callee +0x1a6af0 impl_cuFlushGPUDirectRDMAWrites wrapper-callee +0x1a6c50 impl_cuCtxGetExecAffinity wrapper-callee +0x1a7320 impl_cuDeviceGetLuid wrapper-callee +0x1ae3d0 impl_cuDeviceComputeCapability wrapper-callee +0x1ae550 impl_cuDeviceGetProperties wrapper-callee +0x1ae780 impl_cuDeviceGetAttribute wrapper-callee +0x1be900 impl_cuCtxDestroy wrapper-callee +0x1da450 impl_cuMemFree wrapper-callee +0x1da880 impl_cuMemFreeHost wrapper-callee +0x1e9ab0 impl_cuModuleLoad wrapper-callee +0x1e9d80 impl_cuModuleLoadData wrapper-callee +0x2891d0 impl_cuEventDestroy wrapper-callee +0x2a6bf0 impl_cuGraphDebugDotPrint wrapper-callee +0x2ce4e0 impl_cuLinkCreate wrapper-callee +0x2f9f00 impl_cuModuleGetGlobal wrapper-callee +0x3c9250 impl_cuMemFreeAsync wrapper-callee +0x400ec0 impl_cuMemAlloc wrapper-callee +0x402a30 impl_cuMemGetAddressRange wrapper-callee +0xd11d90 impl_cuStreamBeginCapture wrapper-callee +0xd668f0 impl_cuLibraryGetUnifiedFunction wrapper-callee diff --git a/userspace/cuda/570.144/annotations/subsystems.tsv b/userspace/cuda/570.144/annotations/subsystems.tsv index 5a512bd..4d0c423 100644 --- a/userspace/cuda/570.144/annotations/subsystems.tsv +++ b/userspace/cuda/570.144/annotations/subsystems.tsv @@ -25,13 +25,52 @@ # --- engine/context bring-up (the cuCtxCreate path: crash solved, engine topology blocking; see docs/05) --- 0x3a9c20 cudaHandlePoolAlloc engine bitmap handle allocator: lock, scan 64-bit words for first zero bit, return base + index*64 + bit decompiled 0x3a2f50 cudaHandlePoolFree engine clears the handle bit under the pool mutex decompiled -0x453ef0 cudaEngineContextRegisterHelper engine-engines per-engine registration: alloc handle, call create hook, call bind hook, publish engine id + descriptor into context arrays; called twice for engine index 0 (GR + secondary async CE) decompiled+trace -0x45a670 cudaDeviceInitEngineContexts engine-engines walks 0x14-byte engine records (base+0x2fec for GR; +0x300c+(idx-1)*0x14; +0x350c; +0x351c+(idx-0x16)*0x14), checks engine existence via [table+0x4e00], calls the helper; sees only GR + CE0..4 because the GET_ENGINES_V2/classDB filter drops the async CE (see docs/05 section 3.6) decompiled+trace -0x4676f0 cudaEngineContextCreateHook engine-engines hook at [table+0x40]+0x530; selects instance/class from the engine descriptor and calls rmEngineContextAlloc; its saved rbp was zeroed by the alloc reply's 16-byte class-params copy-back (crash solved, see docs/05 section 3.5) decompiled+trace -0x481360 rmEngineContextAlloc rm builds NV_ESC_RM_ALLOC requests per class (0x80 device, 0x5/0x79 per-class table, 0x2080 subdevice array of 32, default generic) and returns small status codes; with a zeroed secondary engine record the class is rewritten to NV01_ROOT_CLIENT (0x41), status 0x19, and the context rolls back (rc=999) decompiled+trace +0x453ef0 cudaEngineContextRegisterHelper engine-engines per-engine registration: idx = record+0 kind (1 normalizes to 0), alloc handle, call create hook with the record, call bind hook, publish engine id at ctx+0x50+idx*4 and the record pointer at ctx[0x9488]+0x88+idx*8; called twice for engine index 0 (GR + secondary async CE) decompiled+trace +0x45a670 cudaDeviceInitEngineContexts engine-engines per-context engine-init hook installed at ctx+0x94b0; selects the 0x14-byte engine record for ctx+0x38 (GR at +0x2fec; copy at +0x300c+(idx-1)*0x14, +0x300c+(idx-2)*0x14 for idx 0xc..0x15; SEC2 at +0x350c; decomp copy at +0x351c+(idx-0x16)*0x14), checks existence via [table+0x4e00](table, record+0xc), calls the helper; for index 0 also registers the secondary record at *(table+0x4af0)+0x75c decompiled+trace +0x4676f0 cudaEngineContextCreateHook engine-engines create hook at [table+0x40]+0x530; reads record kind +0, instance +8, class +0xc; copies instance into RM alloc-param dword 1 only for kinds 2/3/5; passes record+0xc as the alloc class to rmEngineContextAlloc; its saved rbp was zeroed by the alloc reply's 16-byte class-params copy-back (crash solved, see docs/05 section 3.5); the typed pass recovered its 4 parameters and the field logic (run/analysis/typed/fn_0x4676f0.c) decompiled+trace +0x481360 rmEngineContextAlloc rm builds NV_ESC_RM_ALLOC requests per class (0x80 device, 0x5/0x79 per-class table, 0x2080 subdevice array of 32, default generic with the record as params) and returns small status codes; class comes from the engine record, so a zeroed secondary record means class 0 is rewritten to NV01_ROOT_CLIENT (0x41), status 0x19, and the context rolls back (rc=999) decompiled+trace 0x453cc0 cudaEngineContextCleanup engine-engines failure-path cleanup of an engine context; runs on the rolled-back context (rc=999) decompiled +# --- engine record construction (the record identity and the COPY14 blocker; see docs/05 section 3.7) --- +0x450ff0 cudaDeviceBuildEngineRecords engine-engines builds the 0x14-byte per-engine records from the GET_ENGINES_V2 list at device+0x2aa8 (count +0x2fe8): kind 1 GR -> +0x2fec (class device+0xd38), kind 3 copy -> +0x300c+n*0x14 (class +0xd44, caps bit0 = CE_GRCE also seeds the secondary record at *(device+0x4af0)+0x75c), kind 4 SEC2 -> +0x350c (class +0xd4c), kind 5 decomp copy -> +0x351c+n*0x14 (class +0xd50); fields +0 kind, +4 NV2080 engine type, +8 instance, +0xc class, +0x10 caps decompiled+trace +0x462770 cudaRmGetEngines engine-engines RM control AL wrapper for NV2080_CTRL_CMD_GPU_GET_ENGINES_V2: reads the engine list and classifies each raw NV2080 engine type into {kind, type, instance} records at 0x10-byte stride (GR 1..8 -> kind 1; copy 9..0x12 and 0x34..0x3d -> kind 3; decomp copy 0x40..0x53 -> kind 5; SEC2 0x26 -> kind 4; others skipped); called from 0x44ac20 via AL slot 0x1d8 decompiled +0x44a170 cudaDeviceSupportsClass engine-engines device class-list membership callback installed at device+0x4e00 by the ABI table setup (0x458040 and 0x4494d0): returns true when class is in the class list at device+0xd90 (count +0xd98); the walker calls it with (device, record+0xc) before registering an engine decompiled+trace +0x458040 cudaDeviceAbiTableInit engine-engines builds the RM device state and installs the device ABI function table at device+0x4d18.. (engine builder +0x4d88, class callback +0x4e00 = 0x44a170, ctx table setup +0x4ef0, init caps +0x4d38); class list copied from the config block into device+0xd90/+0xd98 decompiled +0x44ac20 cudaDeviceInitCaps engine-engines device init-caps path: calls the AL getEngines slot 0x1d8 (0x462770) to fill device+0x2aa8/+0x2fe8, then the ABI static-info hooks; the engine list it produces is what the record builder consumes decompiled # --- CURRENT BLOCKER (2026-09-23): engine list / classDB consistency (async CE COPY14) --- -0x45a670 cudaDeviceInitEngineContexts engine-engines CURRENT BLOCKER: classDB-filtered engine list has 6 engines in the guest (GR + CE0..4) vs 7 on the host oracle (adds async CE COPY14: engDesc 0x793ceb0e, instance 14, client type 0x38, class 0xc9b5); without a classDB entry the walker reads a zeroed secondary record and the RM alloc is rejected (class 0x41, st 0x19). Record layout above; see docs/05 section 3.6 decompiled+trace +0x45a670 cudaDeviceInitEngineContexts engine-engines CURRENT BLOCKER: the GET_ENGINES_V2 list has 6 engines in the guest (GR + CE0..4) vs 7 on the host oracle (adds async CE COPY14: engDesc 0x793ceb0e, NV2080 type 0x38, instance 14, paired class 0xc9b5); the classDB filter drops COPY14, so no copy record carries CE_GRCE (caps bit0, live guest 0xbc/0xfc) and the secondary record at *(table+0x4af0)+0x75c stays zero (host t=2 cls=0xc5b5); the zero class is rewritten to 0x41 (st 0x19) and the context rolls back. Record field map and builder chain in docs/05 section 3.7 decompiled+trace # --- open questions --- 0x4a7e90 FUN_004a7e90 unknown called by every generated wrapper; likely tool/trace dispatch or thread-data helper todo 0x2cb940 FUN_002cb940 unknown large function (10 KB) reachable from the cuInit wrapper; candidate init implementation todo +# --- wrapper callees verified in the 2026-09 name-candidate pass --- +# Names proposed by tools/name_candidates.py (see name-candidates.tsv), each +# verified by reading its body in the 570.144 corpus; source `wrapper-callee` +# in names.tsv. Evidence: decompiled. +0x1a36d0 impl_cuCtxGetApiVersion api current-context API version getter: resolves the thread context and returns its API-version field decompiled (wrapper-callee) +0x1a4900 impl_cuArrayGetMemoryRequirements mem array memory-requirements getter: refreshes the backing-store state under the context mutex and returns the allocation size decompiled (wrapper-callee) +0x1a5410 impl_cuArrayGetPlane mem array plane getter: validates the plane index and returns the plane handle decompiled (wrapper-callee) +0x1a5820 impl_cuDevicePrimaryCtxGetState api primary-context state getter: reads flags/active state under the device mutex decompiled (wrapper-callee) +0x1a6150 impl_cuDeviceRegisterAsyncNotification api device async-notification registration: allocates and links a notification record under the device mutex decompiled (wrapper-callee) +0x1a6430 impl_cuCtxAttach api context attach: increments the context attach refcount under the context lock decompiled (wrapper-callee) +0x1a66f0 impl_cuCtxGetCurrent api current-context getter: returns the calling thread's context from TLS decompiled (wrapper-callee) +0x1a6850 impl_cuCtxGetDevice api context device getter: resolves the current context and returns its device ordinal decompiled (wrapper-callee) +0x1a6960 impl_cuCtxGetFlags api context flags getter: composes flags from context fields under the context lock decompiled (wrapper-callee) +0x1a6af0 impl_cuFlushGPUDirectRDMAWrites api GPUDirect RDMA write flush: validates args and dispatches through the device op table decompiled (wrapper-callee) +0x1a6c50 impl_cuCtxGetExecAffinity api context exec-affinity getter: returns the affinity value under the context lock decompiled (wrapper-callee) +0x1a7320 impl_cuDeviceGetLuid api device LUID getter: bounds-checks the ordinal and calls the device op at table+0x5180 decompiled (wrapper-callee) +0x1ae3d0 impl_cuDeviceComputeCapability api compute-capability getter: reads device attributes 0x4b/0x4c decompiled (wrapper-callee) +0x1ae550 impl_cuDeviceGetProperties api device properties getter: fills CUDA device properties from the device record decompiled (wrapper-callee) +0x1ae780 impl_cuDeviceGetAttribute api device attribute getter: validates the ordinal and reads one attribute (fast path for attribute 0x20080403) decompiled (wrapper-callee) +0x1be900 impl_cuCtxDestroy api context destroy: validates the handle, takes the context lock and tears the context down decompiled (wrapper-callee) +0x1da450 impl_cuMemFree mem device memory free: resolves the allocation under the context lock and frees it decompiled (wrapper-callee) +0x1da880 impl_cuMemFreeHost mem host memory free: frees pinned host memory (or routes to the device free path) decompiled (wrapper-callee) +0x1e9ab0 impl_cuModuleLoad module module load from file: reads the file and parses it into the context decompiled (wrapper-callee) +0x1e9d80 impl_cuModuleLoadData module module load from memory image: same path as file load with an in-memory image decompiled (wrapper-callee) +0x2891d0 impl_cuEventDestroy sync event destroy: unlinks and frees an event under the context mutex decompiled (wrapper-callee) +0x2a6bf0 impl_cuGraphDebugDotPrint launch graph DOT dump: writes the graph as a DOT debug file decompiled (wrapper-callee) +0x2ce4e0 impl_cuLinkCreate jit linker create: allocates and initializes the linker state decompiled (wrapper-callee) +0x2f9f00 impl_cuModuleGetGlobal module module global getter: resolves a named global through the module symbol table decompiled (wrapper-callee) +0x3c9250 impl_cuMemFreeAsync mem async device memory free: resolves the allocation and frees it on a stream decompiled (wrapper-callee) +0x400ec0 impl_cuMemAlloc mem device memory alloc: validates/derives the allocation descriptor and allocates decompiled (wrapper-callee) +0x402a30 impl_cuMemGetAddressRange mem allocation lookup: resolves the allocation containing an address under the context mutex decompiled (wrapper-callee) +0xd11d90 impl_cuStreamBeginCapture stream stream begin capture: validates the mode and starts capture under the stream lock decompiled (wrapper-callee) +0xd668f0 impl_cuLibraryGetUnifiedFunction module library function getter: resolves a named function through the library symbol table decompiled (wrapper-callee) diff --git a/userspace/cuda/570.144/ghidra/DecompileAt.java b/userspace/cuda/570.144/ghidra/DecompileAt.java index 5092ee9..f1572c7 100644 --- a/userspace/cuda/570.144/ghidra/DecompileAt.java +++ b/userspace/cuda/570.144/ghidra/DecompileAt.java @@ -1,5 +1,18 @@ // SPDX-License-Identifier: BSD-3-Clause // +// ########################################################################### +// # WARNING: DEPRECATED - DO NOT USE FOR NEW ANALYSIS (kept for history). # +// # # +// # It treats its argument as a Ghidra program address, but the analysis # +// # convention here is FILE virtual addresses (libcuda+0x...) and Ghidra # +// # loads these binaries at image base 0x100000. Called with a file vaddr # +// # it decompiles the function 0x100000 bytes too low (the bug behind the # +// # old `helper_453ef0.c` / `fn_0x481360.c` artifacts). # +// # # +// # Replacement: tools/ghidra/DecompileAt.java (adds the image base) used # +// # through tools/ghidra_decompile_at.sh. # +// ########################################################################### +// // DecompileAt.java - Ghidra post-script: decompile the function containing a // given address (no auto-analysis required). // diff --git a/userspace/cuda/570.144/scripts/ghidra_analyze.sh b/userspace/cuda/570.144/scripts/ghidra_analyze.sh index b844300..eb8422b 100755 --- a/userspace/cuda/570.144/scripts/ghidra_analyze.sh +++ b/userspace/cuda/570.144/scripts/ghidra_analyze.sh @@ -1,6 +1,21 @@ #!/bin/bash # SPDX-License-Identifier: BSD-3-Clause # +# ########################################################################### +# # WARNING: DEPRECATED - DO NOT USE FOR NEW ANALYSIS (kept for history). # +# # # +# # Two bugs make its output unusable: # +# # 1. It passes FILE virtual addresses to Ghidra, which loads these # +# # binaries at image base 0x100000. Every address is 0x100000 too low, # +# # so the decompiled "functions" are the wrong ones. # +# # 2. `-analysisTimeoutPerFile 0` does not mean "no timeout" in this # +# # Ghidra build: it aborts auto-analysis immediately, producing an # +# # analysed-looking project with no analysis at all. # +# # # +# # Replacement: tools/ghidra_full_analysis.sh (plus # +# # tools/ghidra/ConfigureAnalysis.java). # +# ########################################################################### +# # ghidra_analyze.sh - import a library into a Ghidra project and run the full # auto-analysis once. The project is kept, so later decompilations are cheap: # diff --git a/userspace/cuda/570.144/scripts/ghidra_decompile_project.sh b/userspace/cuda/570.144/scripts/ghidra_decompile_project.sh index 66d7c90..b44b2b2 100755 --- a/userspace/cuda/570.144/scripts/ghidra_decompile_project.sh +++ b/userspace/cuda/570.144/scripts/ghidra_decompile_project.sh @@ -1,6 +1,19 @@ #!/bin/bash # SPDX-License-Identifier: BSD-3-Clause # +# ########################################################################### +# # WARNING: DEPRECATED - DO NOT USE FOR NEW ANALYSIS (kept for history). # +# # # +# # It passes a FILE virtual address straight to 570.144/ghidra/ # +# # DecompileAt.java, but Ghidra loads these binaries at image base # +# # 0x100000: the script resolves (and decompiles) the function 0x100000 # +# # bytes too low. It also depends on a project produced by the broken # +# # ghidra_analyze.sh above. # +# # # +# # Replacement: tools/ghidra_decompile_at.sh, which translates the file # +# # vaddr with tools/ghidra/DecompileAt.java (file-vaddr aware). # +# ########################################################################### +# # ghidra_decompile_project.sh - decompile a function from an already-analyzed # Ghidra project (run scripts/ghidra_analyze.sh once first). No re-import and # no analysis, so this takes seconds. diff --git a/userspace/cuda/READ-ME-FIRST.md b/userspace/cuda/READ-ME-FIRST.md index 180907f..2abb797 100644 --- a/userspace/cuda/READ-ME-FIRST.md +++ b/userspace/cuda/READ-ME-FIRST.md @@ -112,17 +112,22 @@ userspace/cuda/ ├── README.md index + honest status ├── docs/01..07 public analysis (tracked) ├── tools/ reusable tooling (tracked) -├── LOCAL-SOURCES.md local source map (gitignored) +│ ├── cudatool.py ELF exports/FDEs/wrappers/delta/fdematch +│ ├── callwalk.py callee/caller/path walker +│ ├── name_candidates.py API-implementation name candidates +│ ├── subsystem_digest.py per-subsystem call-tree digest +│ ├── diff_corpus.py corpus diff across driver versions +│ └── ghidra_*.sh, ghidra/ Ghidra analysis/decompile/name-map scripts └── 570.144/ ├── README.md, docs/, scripts/ version-specific helpers + old write-up ├── annotations/ name map + subsystem index (tracked) - ├── LOCAL-ANALYSIS.md local version notes (gitignored) └── run/ ├── libcuda.so.570.144 the guest binary (md5 ed5b2fc0…) ├── ghidra-full/ analyzed project (29,494 functions) ├── analysis/elf-index/ functions/exports/callgraph/strings TSVs ├── analysis/decompiled/ raw corpus (29,254 functions, 0 failures) - └── analysis/decompiled-annotated/ corpus + names + call context + ├── analysis/decompiled-annotated/ corpus + names + call context + └── analysis/subsystem-digest.txt, fdematch-*.txt generated reports ``` ## 5. Rules while you work in this area diff --git a/userspace/cuda/README.md b/userspace/cuda/README.md index 38ff6f5..9c3163d 100644 --- a/userspace/cuda/README.md +++ b/userspace/cuda/README.md @@ -70,8 +70,11 @@ userspace/cuda/ ├── READ-ME-FIRST.md reading guide for agents/debuggers ├── docs/ public analysis documents ├── tools/ reusable, version-agnostic tooling -│ ├── cudatool.py ELF/export/FDE/wrapper/version-delta toolkit +│ ├── cudatool.py ELF/export/FDE/wrapper/delta/fdematch toolkit │ ├── callwalk.py callgraph callee/caller/path walker +│ ├── name_candidates.py propose impl names for wrapper callees +│ ├── subsystem_digest.py per-subsystem call-tree digest +│ ├── diff_corpus.py diff two decompiled corpus indexes │ ├── annotate_corpus.py annotate + rename the decompiled corpus │ ├── ghidra_full_analysis.sh │ ├── ghidra_export_index.sh diff --git a/userspace/cuda/docs/05-570.144-decompilation.md b/userspace/cuda/docs/05-570.144-decompilation.md index 84b377e..4e23866 100644 --- a/userspace/cuda/docs/05-570.144-decompilation.md +++ b/userspace/cuda/docs/05-570.144-decompilation.md @@ -162,9 +162,18 @@ engine index >= 0x16 -> table + 0x351c + (index-0x16)*0x14 before the helper, the walker calls the callback at `table+0x4e00` with `(table, desc->[0xc])` — a -"does this engine/class exist" check — and only proceeds when it returns -true. The helper then uses `desc->[0]` (engine kind), `desc->[8]` (instance) -and `desc->[0xc]` (class) for the registration. +"does this engine/class exist" check (a class-list membership test; the +callback pointer is installed by the device ABI table setup at +`libcuda+0x458040`/`libcuda+0x4494d0`, target `libcuda+0x44a170`) — and only +proceeds when it returns true. The helper then uses `desc->[0]` (engine kind), +`desc->[8]` (instance) and `desc->[0xc]` (class) for the registration. + +The records are built by `libcuda+0x450ff0` from the engine list that the +`GET_ENGINES_V2` wrapper `libcuda+0x462770` returns (`{kind, engine type, +instance}` triples at 0x10-byte stride). The record layout and the full field +map are in §3.7. Note the walker's index ranges cover copy slots `1..0xa` and +`0xc..0x15`, with `0xb` used for the SEC2 slot and `>=0x16` for decompression +copy slots; the engine list is a different, dense array. ### 3.4 The engine create/register function @@ -232,8 +241,8 @@ path on the host's real driver): * libcuda asks the driver for the GPU's engines (`GET_ENGINES_V2`). The host oracle's filtered list has **7** engines: GR, CE0..CE4 and `COPY14`, the async copy engine. The guest list has only **6** — the same minus `COPY14`. -* `COPY14`'s engine descriptor is `0x793ceb0e`, instance 14, client engine - type `0x38`, paired with engine class `0xc9b5`. The GSP-RM firmware's +* `COPY14`'s engine descriptor is `0x793ceb0e`, instance 14, engine type + `0x38`, paired with engine class `0xc9b5`. The GSP-RM firmware's `{class, engDesc}` pair table pairs OBJCE instances `0..7` with class `0xcab5` and instances `0x0b..0x13` (including `0x0e` = 14) with class `0xc9b5`. @@ -256,24 +265,56 @@ static info and add a 13th device-info entry with descriptor `0x793ceb0e`, instance 14, RM engine type `0x17`, FIFO tag 6, runlist 5) did not change the guest list, so the missing input is elsewhere: what builds the driver's classDB entry for descriptor `0x793ceb0e` (class list/static-info class data -versus the device-info-table fields) is the open question. The record layout -the walker reads is in §3.3; the per-engine fields it consumes are -`+0` (kind), `+8` (instance) and `+0xc` (class). - -### 3.7 What the emulation should check first - -* **Engine-list/classDB consistency — the active blocker (§3.6).** The engine - records at the offsets in §3.3 must be present and self-consistent for every - engine index the walker visits, the "does this engine exist" callback - (`[table+0x4e00]`) must agree with them, and the whole set must agree with - the class-filtered engine list the driver builds from the GSP static info - (`GET_ENGINES_V2` + classDB). A missing entry is silent: the walker then - reads a zeroed record. +versus the device-info-table fields) is the open question. + +### 3.7 The engine descriptor record, and how it explains the blocker + +The device state holds a list of 0x10-byte engine entries plus a parallel set +of 0x14-byte **per-engine descriptor records**; the walker (§3.3) reads the +records, the helper and create hook consume them. The record was recovered +from the record builder (`libcuda+0x450ff0`), the `GET_ENGINES_V2` wrapper +(`libcuda+0x462770`), the create hook disassembly (`libcuda+0x4676f0`) and the +live guest dump (`run/gdb32b.out`): + +| Offset | Field | Meaning | Evidence | +|---|---|---|---| +| `+0x00` | engine kind | `1` = graphics (GR), `3` = copy (CE), `4` = SEC2, `5` = decompression copy. The helper normalizes `1` to `0` for the per-context slot index. | builder branches on it; `GET_ENGINES_V2` wrapper classifies each engine type into it; live: GR `t=1`, CE `t=3` | +| `+0x04` | engine type | raw `NV2080_ENGINE_TYPE` from the engine list (`1` = GR, `9..13` = COPY0..4, `0x38` = COPY14) | builder copies the first 8 bytes of the list entry into `+0/+4`; not read by the walker/helper | +| `+0x08` | instance | engine instance (`COPY0..4` → `0..4`, `COPY14` → `14`) | helper + create hook read it; create hook copies it into the RM alloc parameters only for kinds 2/3/5; live dump | +| `+0x0c` | class | RM allocation class for this engine: GR → the device's compute class (live `0xc5c0`), copy → the DMA copy class (live `0xc5b5`), SEC2/decomp copy → their classes | builder stores the per-kind class; the create hook passes `+0xc` straight to `libcuda+0x481360` as the `NV_ESC_RM_ALLOC` class; host capture shows `class=0xc5b5` | +| `+0x10` | engine caps | for copy records: the CE capability bitmask; **bit 0 (`CE_GRCE`) marks the async/GRCE copy engine**. A copy record with bit 0 set is also copied to the *secondary record* slot and does not take a normal CE slot. | builder tests `+0x10 & 1` and copies to the secondary; live caps `0xbc`/`0xfc` (bit 0 clear) | + +Secondary record: the GR context registers a second engine from a dedicated +0x14-byte slot (same layout) that is filled from the GRCE-capable copy record +and then normalized to kind `2` and the async copy class. GR's own record is +larger (a 0x20-byte slot with an unrecovered tail at `+0x10..+0x1f`). + +**How this explains the blocker.** The builder derives the records from the +engine list the driver returns for `GET_ENGINES_V2`; the guest list has 6 +entries (GR + CE0..4) because the class-database filter drops `COPY14`. With +no `COPY14` entry there is no copy record with the `CE_GRCE` capability bit +set, so the secondary record slot is never filled and the walker sees class +`0`. libcuda then allocates the engine object with class `0`, the escape path +rewrites that to `NV01_ROOT_CLIENT` (`0x41`), the allocation fails with status +`0x19`, the context rolls back and `cuCtxCreate` returns 999. In other words: +the guest is not merely missing one entry in a list — it is missing the +*async copy engine record that seeds the secondary record*. + +**What the emulation should check first (revised)**: + +* **`COPY14` in `GET_ENGINES_V2`** — the engine's classDB entry must exist so + that `gpuUpdateEngineTable`'s class-list filter keeps OBJCE instance 14. + The class/advertised-engine-table consistency is the root cause; the + secondary-record zeroing is its symptom. +* **CE capabilities.** The copy records' caps word must carry the `CE_GRCE` + bit for the async CE. If `COPY14` is restored but its emulated CE caps still + return bit 0 clear, the secondary record stays empty and the class-0 + allocation returns. * **Consumed-control reply fidelity.** The controls answered generically (the host log lists `0x20803039`, `0x20802a12`, `0x20808162`, `0x2080182b`, `0x20800a70`, `0x20800a93`, …) should be checked for consumed fields, because the walker's engine table is built from data the driver collected earlier - (device caps, engine list, class list). + (device caps, engine list, class list, CE caps). * The `NV_ESC_RM_ALLOC` parameter block built by `libcuda+0x481360` for class `0xc5c0` (`TURING_COMPUTE_A`) must be answered with a coherent object: echo the request's params back (§3.5) and return a coherent handle; the driver @@ -307,7 +348,15 @@ citations in docs and gdb sessions use the same numbers. holds our name map (exports are exact, internals grow over time). * The decompiler does not recover struct types; field offsets are named `[obj+0xNN]` and must be interpreted with the descriptions in `docs/01` - and `docs/03`. + and `docs/03`. A targeted Decompiler Parameter ID pass + (`tools/ghidra/AnalyzeFunctions.java`, see `docs/06` §2d) re-ran over 80 + key functions (crash/engine path, resolved wrapper implementations, largest + `.text`) and committed their signatures into a project copy + (`run/ghidra-typed/`); the re-decompiled bodies are in + `run/analysis/typed/` and the before/after report in + `run/analysis/typed/paramid_report.tsv`. 16 of 80 bodies changed materially + (parameter/type recovery), the rest now carry committed signatures that + improve downstream decompilation. * Indirect calls dominate the API dispatch; the callgraph contains direct edges only (77,742). Cross-referencing indirect targets requires reading the wrapper/stub relocations (as in §2) or live traces. diff --git a/userspace/cuda/docs/06-tooling.md b/userspace/cuda/docs/06-tooling.md index 85d1c8c..72fc8ac 100644 --- a/userspace/cuda/docs/06-tooling.md +++ b/userspace/cuda/docs/06-tooling.md @@ -8,6 +8,9 @@ tools are stdlib-only or Ghidra scripts; nothing needs network access. userspace/cuda/tools/ ├── cudatool.py ELF/interface toolkit (python3, no deps) ├── callwalk.py callee/caller trees and shortest call path +├── name_candidates.py propose impl names for wrapper callees +├── subsystem_digest.py per-subsystem call-tree digest +├── diff_corpus.py diff two decompiled corpus indexes ├── annotate_corpus.py annotated + name-substituted corpus ├── ghidra_full_analysis.sh import + full auto-analysis into a project ├── ghidra_export_index.sh functions/exports/callgraph/strings index @@ -19,6 +22,7 @@ userspace/cuda/tools/ ├── ExportFunctionIndex.java index exporter ├── BatchDecompile.java chunk decompiler ├── DecompileAt.java single function (file-vaddr aware) + ├── AnalyzeFunctions.java targeted Decompiler Parameter ID pass ├── ApplyNameMap.java apply/transfer a name map └── PrintProgramInfo.java address convention + stats ``` @@ -34,12 +38,33 @@ tools/cudatool.py crash <lib> <map_base> <ip> # dmesg -> file vaddr tools/cudatool.py strings <lib> [--grep RE] tools/cudatool.py note <lib> # .note.cuda.fwd_compatibility JSON tools/cudatool.py delta <libA> <libB> # see §4 +tools/cudatool.py fdematch <libA> <libB> [--names names.tsv] [--window 0x100] ``` Useful for: checking what changed between two driver package versions before spending time on decompilation; converting a guest dmesg crash to a file vaddr without gdb; getting a function list in a second. +### 1b. Cross-version function matching — `fdematch` + +``` +tools/cudatool.py fdematch <libA> <libB> [--names <names.tsv>] [--window 0x100] +``` + +`fdematch` matches function ranges from the two `.eh_frame_hdr` tables. A +single global address shift is not enough across patch versions, so it uses +the common exports as **local anchors**: for each A function it looks for a B +function with the same FDE size within `--window` bytes of +`old_addr + delta(nearest common export)`. The report gives exact / window / +unique match counts and the overall match rate, and with `--names` it breaks +the name map down by its `source` column (exports, wrappers, verified, +wrapper-callee), which is exactly what you need to know before transferring +names to a new version. Example result for 570.144 → 570.211.01: 85.5% of +functions match (58.3% uniquely), **all 659 exports transfer exactly, but +only ~1% of the generated wrapper addresses do** — the wrapper region is +regenerated and must be re-resolved with `cudatool.py wrappers`, never by +address arithmetic. + ### Address conventions * `cudatool.py` prints **file virtual addresses** — the `libcuda+0x...` used @@ -102,6 +127,67 @@ tools/callwalk.py <indexdir>/callgraph.tsv chain 0x31b7f0 0x2cb940 Addresses are file vaddrs; `names.tsv` next to the callgraph is used for labels. +## 2d. Targeted type recovery + +The full analysis disables the Decompiler Parameter ID analyzer (§2), so +function signatures stay at the default `FUN_<addr>(...)`. `AnalyzeFunctions.java` +re-runs that analyzer over a caller-supplied function list instead of the whole +program: it takes an addresses file (one `0x<file_vaddr>` per line, comments +allowed), translates through the image base like `DecompileAt.java`, resolves +each address to its function and runs the analyzer on the resulting address +set (batched, with a per-function decompiler timeout). Run it against a +project copy, not the canonical one, because it commits parameters into the +project: + +``` +cp -a <proj_dir> <proj_dir>-typed +analyzeHeadless <proj_dir>-typed <proj_name> -process <program> -noanalysis \ + -scriptPath tools/ghidra \ + -postScript AnalyzeFunctions.java <addresses.txt> \ + --timeout 120 --batch 20 --out <report.tsv> +``` + +Then re-decompile the affected functions with `ghidra_decompile_at.sh` from the +typed project; the report's `params_before/after` and `GAIN` columns show +which functions gained structure. `--list` resolves the list without running +the analyzer (dry run). + +## 2d. Subsystem call-tree digest — `subsystem_digest.py` + +``` +tools/subsystem_digest.py <subsystems.tsv> <callgraph.tsv> [--out FILE] + [--depth 2] [--max-per-node 8] [--subsystem NAME] +``` + +For every subsystem in `annotations/subsystems.tsv` it prints the annotated +members, the subsystem **entry functions** (members not called by another +member of the same subsystem) and a depth-limited call tree from each entry +whose direct callees are grouped by subsystem. It closes with a +cross-subsystem edge count matrix. This is the quickest way to see a +subsystem's shape and where its boundary with the rest of the driver is, +without opening Ghidra. Example output for 570.144 is saved under +`570.144/run/analysis/subsystem-digest.txt` (gitignored). + +## 2e. API-implementation candidates — `name_candidates.py` + +``` +tools/name_candidates.py <exports.tsv> <callgraph.tsv> <names.tsv> <out.tsv> + [--max-outside 2] +``` + +Exports are generated stubs → generated wrappers → the hand-written +implementation of that API. `name_candidates.py` groups the wrappers by base +name (stripping `_v<N>`, `_ptsz`, `_ptds`), takes the intersection of their +direct callees and filters out helpers that are also called by wrappers +outside the group (the trace/init/lock helpers). A single surviving candidate +is `high` confidence when the group has several variants or the callee is +unique; otherwise the best-ranked candidate is `medium`. Output rows carry +the full candidate set and the outside-wrapper count as evidence, plus a +`status` column that starts as `unverified`. Verification means reading the +body in the corpus and marking `verified`/`corrected`; only verified entries +may be copied into `names.tsv` (source `wrapper-callee`). The 570.144 run is +tracked in `570.144/annotations/name-candidates.tsv`. + ## 3. Looking things up ``` @@ -131,6 +217,23 @@ first thing to run on a new driver release; it tells you immediately whether the update is a patch (nothing moved), a feature release (new exports) or a reshuffle (many moves). +### 4b. Corpus deltas — `diff_corpus.py` + +``` +tools/diff_corpus.py <indexA.tsv> <indexB.tsv> [--label-a A] [--label-b B] + [--out FILE] [--max-list 40] [--json] +``` + +Diffs two `decompiled_index.tsv` corpora (the merged output of +`ghidra_batch_decompile.sh` + `ghidra_index_merge.sh`) and reports +added/removed/moved/resized functions and a summary. Named functions match by +name; the address shift is estimated from them and anonymous `FUN_*` +functions match at `old_addr + shift` (a size difference there is a resize). +This makes a driver upgrade measurable at the function level: run it after +decompiling the new version to see how much of the previous corpus still +corresponds. It needs a decompiled index for the other version; if only the +binary exists, use `cudatool.py fdematch` (no decompilation required). + ## 5. Name transfer — `ApplyNameMap.java` ``` diff --git a/userspace/cuda/docs/07-version-upgrade-playbook.md b/userspace/cuda/docs/07-version-upgrade-playbook.md index 337f5eb..23569bd 100644 --- a/userspace/cuda/docs/07-version-upgrade-playbook.md +++ b/userspace/cuda/docs/07-version-upgrade-playbook.md @@ -49,6 +49,33 @@ with `PrintProgramInfo.java` before citing addresses. ## 3. Transfer names and knowledge +Measure the transfer first, before copying anything: + +``` +tools/cudatool.py fdematch <previous>/run/libcuda.so.<old> <libcuda.so.<new>> \ + --names <previous>/annotations/names.tsv +``` + +`fdematch` matches function ranges from both `.eh_frame_hdr` tables using the +common exports as local address anchors (same FDE size within a small window +of `old_addr + nearest-export delta`), and breaks the previous name map down +by its `source` column. Measured for **570.144 → 570.211.01** +(`570.144/run/analysis/fdematch-570.144-570.211.txt`, gitignored): + +| Map part | Transferable | How | +|---|---|---| +| 659 exports | **100%**, exact | same name, same size, address = old + local delta | +| verified internals | 5/7 (71%) | anchored same-size window | +| wrapper-callee internals | 21/29 (72%) | anchored same-size window | +| generated `wrapper_*` | **1.1%** | **do not transfer by address** | +| all A functions | 85.5% matched (58.3% unique) | anchored same-size window | + +The wrapper result is the important one: the generated wrapper region is +re-laid-out between patch versions, so wrapper addresses must be re-resolved +on the new binary with `tools/cudatool.py wrappers` (and then +`annotate_corpus.py` re-run), never by adding a delta to the old map. Export +and hand-written internal names do carry over. + ``` cp <previous>/annotations/names.tsv <new>/annotations/names.tsv # translate addresses: the map is file-vaddr keyed, so for unchanged @@ -61,7 +88,9 @@ analyzeHeadless <proj> cuda-full -process <program> -noanalysis \ Also transfer the *subsystem index* (`annotations/subsystems.tsv`: address, our name, subsystem, evidence) and re-validate the entries that the delta -marked as moved. +marked as moved. For corpus-level measurement (added/removed/resized +functions, not just name transfer) use `tools/diff_corpus.py` on the old and +new `decompiled_index.tsv` once the new version has been decompiled. ## 4. Validate against the guest diff --git a/userspace/cuda/docs/08-virtualization-approaches.md b/userspace/cuda/docs/08-virtualization-approaches.md index 3523ad1..00ca13f 100644 --- a/userspace/cuda/docs/08-virtualization-approaches.md +++ b/userspace/cuda/docs/08-virtualization-approaches.md @@ -2,7 +2,7 @@ One-page decision summary for how a guest's CUDA work gets executed in Open-IOV. Companion to `docs/04-virtualization-and-emulation.md` (design -principles) and to the project implementation docs (`docs/refactored-architecture.md`, +principles) and to the project implementation docs (`docs/03-architecture/06-host-module-and-hybrid-model.md`, `docs/04-implementation/02-practical-challenges-and-adversarial-interop.md`). ## 1. The options diff --git a/userspace/cuda/tools/cudatool.py b/userspace/cuda/tools/cudatool.py index 2e4c1c6..3417d2b 100755 --- a/userspace/cuda/tools/cudatool.py +++ b/userspace/cuda/tools/cudatool.py @@ -21,6 +21,7 @@ # strings <lib> [--grep RE] defined strings with file offsets # delta <libA> <libB> interface-level diff between two versions # wrappers <lib> export stub -> generated wrapper target +# fdematch <libA> <libB> function-range match report across versions # note <lib> CUDA forward-compatibility note (JSON) # # Address convention: every address printed by this tool is a FILE VIRTUAL @@ -476,6 +477,151 @@ def cmd_wrappers(args): print(f"0x{w['stub']:08x} -> 0x{w['wrapper']:08x} {w['name']}") +def cmd_fdematch(args): + """Match function ranges between two libraries by unwind-table (FDE) + start/size, using the common exports as local address anchors. + + A single global address shift is not enough across patch versions: the + library is partly re-laid-out (function insertions/removals drift the + addresses locally, and the generated wrapper region is regenerated). The + matcher therefore uses the nearest common export by address as the local + anchor and looks for a B function of the same FDE size within a small + window around old_start + anchor_delta. This measures how transferable a + name map is between two driver versions, and the --names breakdown shows + which parts of the map (exports, generated wrappers, internals) survive.""" + import bisect + from collections import Counter, defaultdict + + a, b = ELF(args.libA), ELF(args.libB) + fa = {s: (e - s if e is not None else None) + for s, e in a.function_ranges()} + fb = {s: (e - s if e is not None else None) + for s, e in b.function_ranges()} + + ea = {s["name"]: s["value"] for s in a.exports()} + eb = {s["name"]: s["value"] for s in b.exports()} + common = sorted(set(ea) & set(eb)) + anchors = sorted((ea[n], eb[n] - ea[n]) for n in common) + anchor_addrs = [x for x, _ in anchors] + deltas = [d for _, d in anchors] + d_counter = Counter(d // 0x1000 * 0x1000 for d in deltas) + + def anchor_delta(v): + i = bisect.bisect_left(anchor_addrs, v) + cand = [] + if i < len(anchors): + cand.append(anchors[i]) + if i > 0: + cand.append(anchors[i - 1]) + if not cand: + return 0 + return min(cand, key=lambda t: abs(t[0] - v))[1] + + size_b = defaultdict(list) + for s, sz in fb.items(): + size_b[sz].append(s) + for k in size_b: + size_b[k].sort() + + def window_matches(v, size, w): + if size is None or size not in size_b: + return [] + target = v + anchor_delta(v) + starts = size_b[size] + lo = bisect.bisect_left(starts, target - w) + hi = bisect.bisect_right(starts, target + w) + return starts[lo:hi] + + exact = window = unique = unmatched = 0 + for v, size in fa.items(): + m = window_matches(v, size, args.window) + if not m: + unmatched += 1 + elif v + anchor_delta(v) in m: + exact += 1 + if len(m) == 1: + unique += 1 + else: + window += 1 + if len(m) == 1: + unique += 1 + matched = exact + window + total = len(fa) + rate = matched / total if total else 0.0 + + transfer = None + if args.names: + transfer = defaultdict(lambda: dict(entries=0, exact=0, window=0, + missing=0, not_in_a=0)) + with open(args.names) as f: + for line in f: + line = line.strip() + if not line or line.startswith("#"): + continue + p = line.split("\t") + if len(p) < 2: + continue + try: + v = int(p[0], 16) + except ValueError: + continue + src = p[2] if len(p) > 2 else "?" + t = transfer[src] + t["entries"] += 1 + if v not in fa: + t["not_in_a"] += 1 + continue + m = window_matches(v, fa[v], args.window) + if not m: + t["missing"] += 1 + elif v + anchor_delta(v) in m: + t["exact"] += 1 + else: + t["window"] += 1 + + if args.json: + print(json.dumps(dict( + lib_a=args.libA, lib_b=args.libB, + functions_a=total, functions_b=len(fb), + anchors=len(anchors), + delta_median=hex(deltas[len(deltas) // 2]) if deltas else "0x0", + delta_clusters=[[hex(k), v] for k, v in d_counter.most_common(6)], + window=hex(args.window), + exact=exact, window_matches=window, unique=unique, + unmatched=unmatched, match_rate=round(rate, 4), + name_map=transfer), indent=1)) + return + + print(f"fdematch A: {args.libA}") + print(f" B: {args.libB}") + print(f"functions : A={total} B={len(fb)} " + f"(ranges from .eh_frame_hdr)") + print(f"anchors : {len(anchors)} common exports") + if deltas: + print(f"delta clusters : " + ", ".join( + f"{'+' if k >= 0 else '-'}0x{abs(k):x} ({v})" + for k, v in d_counter.most_common(5))) + print(f"matching window : +/-0x{args.window:x} around " + f"nearest-anchor-shifted address, same FDE size") + print(f"exact match : {exact} ({exact/total*100:.1f}%)" + if total else "exact match : 0") + print(f"window match : {window} ({window/total*100:.1f}%)" + if total else "window match : 0") + print(f"unique match : {unique} ({unique/total*100:.1f}%)" + if total else "unique match : 0") + print(f"unmatched A : {unmatched}") + print(f"match rate : {rate*100:.1f}% of A's {total} functions") + if transfer: + print(f"name map : {args.names} (per source column)") + for src in sorted(transfer): + t = transfer[src] + got = t["exact"] + t["window"] + pct = got / t["entries"] * 100 if t["entries"] else 0.0 + print(f" {src:<14} : {got}/{t['entries']} ({pct:.1f}%) " + f"exact={t['exact']} window={t['window']} " + f"missing={t['missing']} not-in-A={t['not_in_a']}") + + def cmd_note(args): elf = ELF(args.lib) print(json.dumps(elf.cuda_note(), indent=2)) @@ -502,6 +648,14 @@ def main(): p = sub.add_parser("note"); p.add_argument("lib"); p.set_defaults(fn=cmd_note) p = sub.add_parser("wrappers"); p.add_argument("lib") p.add_argument("--json", action="store_true"); p.set_defaults(fn=cmd_wrappers) + p = sub.add_parser("fdematch"); p.add_argument("libA"); p.add_argument("libB") + p.add_argument("--names", default=None, + help="names.tsv to report transferability for") + p.add_argument("--window", type=lambda x: int(x, 0), default=0x100, + help="same-size match window around the anchored address " + "(default 0x100)") + p.add_argument("--json", action="store_true") + p.set_defaults(fn=cmd_fdematch) args = ap.parse_args() try: diff --git a/userspace/cuda/tools/diff_corpus.py b/userspace/cuda/tools/diff_corpus.py new file mode 100755 index 0000000..7224ac7 --- /dev/null +++ b/userspace/cuda/tools/diff_corpus.py @@ -0,0 +1,221 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: BSD-3-Clause +# +# diff_corpus.py - diff two decompiled corpus indexes across driver versions. +# +# Inputs are `decompiled_index.tsv` files (the merged output of the batch +# decompiler): a header line and one row per function: +# addr<TAB>file_addr<TAB>size<TAB>name<TAB>file<TAB>line<TAB>status +# `addr` is the Ghidra program address, `file_addr` the file virtual address +# (libcuda+0x...) used everywhere in the analysis. +# +# Matching (function identity across versions): +# 1. Named functions (anything whose name does not start with FUN_) match by +# name; the same name in both indexes is the same function. +# 2. The address shift between the two builds is estimated as the median of +# (file_addrB - file_addrA) over name-matched functions. +# 3. Anonymous FUN_* functions match when B has a function at +# file_addrA + shift; a size difference there is reported as a resize. +# Functions that do not land on a B function are `removed`; B functions +# that no A function maps to are `added`. +# +# This is deliberately address/size based (the corpus has no structural hash), +# so it measures exactly what matters for name-map transfer: did the function +# move, grow, disappear or appear. +# +# Usage: +# diff_corpus.py <indexA.tsv> <indexB.tsv> [--label-a A] [--label-b B] +# [--out FILE] [--max-list N] [--json] +# +# Exit status is 0 for a normal report (differences do not fail the tool). +# +# SPDX-License-Identifier: BSD-3-Clause + +import argparse +import json +import statistics +import sys +from collections import Counter, defaultdict + + +def load_index(path): + funcs = {} + with open(path) as f: + f.readline() # header + for line in f: + p = line.rstrip("\n").split("\t") + if len(p) < 4: + continue + try: + funcs[int(p[1], 16)] = dict(addr=int(p[0], 16), size=int(p[2]), + name=p[3]) + except ValueError: + continue + return funcs + + +def is_named(name): + return bool(name) and not name.startswith("FUN_") + + +def label(f, names): + if f is None: + return "<none>" + return f"{names.get(f['name'], f['name'])} (libcuda+0x{f['file']:x}, {f['size']}B)" + + +def main(): + ap = argparse.ArgumentParser( + description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter) + ap.add_argument("index_a") + ap.add_argument("index_b") + ap.add_argument("--label-a", default=None) + ap.add_argument("--label-b", default=None) + ap.add_argument("--out", default=None, help="write the report here " + "(default: stdout)") + ap.add_argument("--max-list", type=int, default=40, + help="max entries listed per category (default 40)") + ap.add_argument("--json", action="store_true", help="emit JSON") + args = ap.parse_args() + + a = load_index(args.index_a) + b = load_index(args.index_b) + la = args.label_a or args.index_a + lb = args.label_b or args.index_b + + named_a = defaultdict(list) + named_b = defaultdict(list) + for fa, f in a.items(): + if is_named(f["name"]): + named_a[f["name"]].append(fa) + for fb, f in b.items(): + if is_named(f["name"]): + named_b[f["name"]].append(fb) + + matched_named = [] # (name, fa, fb) + for name in sorted(set(named_a) & set(named_b)): + # duplicate symbol names (multiple thunks) match in address order + for fa, fb in zip(sorted(named_a[name]), sorted(named_b[name])): + matched_named.append((name, fa, fb)) + + deltas = [fb - fa for _, fa, fb in matched_named] + shift = int(statistics.median(deltas)) if deltas else 0 + dominant = (sum(1 for d in deltas if abs(d - shift) <= 0x40) / + len(deltas)) if deltas else 0.0 + + used_a = set(fa for _, fa, _ in matched_named) + used_b = set(fb for _, _, fb in matched_named) + moved = [(n, fa, fb) for n, fa, fb in matched_named if fa != fb] + resized = [(n, fa, fb) for n, fa, fb in matched_named + if a[fa]["size"] != b[fb]["size"]] + + shifted = [] + for fa, f in sorted(a.items()): + if fa in used_a or is_named(f["name"]): + continue + fb = fa + shift + if fb in b and fb not in used_b: + shifted.append((fa, fb)) + used_a.add(fa) + used_b.add(fb) + shifted_resized = [(fa, fb) for fa, fb in shifted + if a[fa]["size"] != b[fb]["size"]] + shifted_ok = [(fa, fb) for fa, fb in shifted + if a[fa]["size"] == b[fb]["size"]] + + removed = sorted(fa for fa in a if fa not in used_a) + added = sorted(fb for fb in b if fb not in used_b) + + anon_a = sum(1 for f in a.values() if not is_named(f["name"])) + anon_b = sum(1 for f in b.values() if not is_named(f["name"])) + named_count_a = sum(len(v) for v in named_a.values()) + named_count_b = sum(len(v) for v in named_b.values()) + matched_total = len(matched_named) + len(shifted) + + result = dict( + label_a=la, label_b=lb, + total_a=len(a), total_b=len(b), + named_a=named_count_a, named_b=named_count_b, + anon_a=anon_a, anon_b=anon_b, + shift=shift, shift_dominant=round(dominant, 4), + matched_named=len(matched_named), matched_shifted=len(shifted_ok), + matched_shifted_resized=len(shifted_resized), + moved=len(moved), resized_named=len(resized), + added=len(added), removed=len(removed), + match_rate_a=round(matched_total / len(a), 4) if a else 0.0, + match_rate_b=round(matched_total / len(b), 4) if b else 0.0, + named_transfer_rate=round(len(matched_named) / named_count_a, 4) + if named_count_a else 0.0, + ) + + def hexlist(addrs): + return [f"libcuda+0x{x:x} {a[x]['name']} ({a[x]['size']}B)" + if x in a else f"libcuda+0x{x:x} {b[x]['name']} ({b[x]['size']}B)" + for x in addrs[:args.max_list]] + + if args.json: + out = dict(result) + out["added_list"] = [f"libcuda+0x{x:x} {b[x]['name']} ({b[x]['size']}B)" + for x in added[:args.max_list]] + out["removed_list"] = [f"libcuda+0x{x:x} {a[x]['name']} ({a[x]['size']}B)" + for x in removed[:args.max_list]] + text = json.dumps(out, indent=1) + else: + lines = [] + lines.append(f"corpus diff: {la} -> {lb}") + lines.append("=" * 72) + lines.append(f"functions : A={len(a)} B={len(b)} " + f"(named A={len(named_a)} B={len(named_b)}; " + f"anonymous A={anon_a} B={anon_b})") + lines.append(f"address shift : 0x{shift:x} " + f"(dominant for {dominant*100:.1f}% of name matches)") + lines.append(f"matched : {matched_total} " + f"({len(matched_named)} by name, " + f"{len(shifted_ok)} anonymous at shifted address)") + lines.append(f"match rate : A {result['match_rate_a']*100:.1f}% " + f"B {result['match_rate_b']*100:.1f}% " + f"(name map transfer {result['named_transfer_rate']*100:.1f}%)") + lines.append(f"moved : {len(moved)} named functions changed " + f"address") + lines.append(f"resized : {len(resized)} named + " + f"{len(shifted_resized)} anonymous matched functions " + f"changed size") + lines.append(f"added : {len(added)}") + lines.append(f"removed : {len(removed)}") + lines.append("") + if moved: + lines.append(f"-- moved named (first {args.max_list}) --") + for n, fa, fb in moved[:args.max_list]: + lines.append(f" {n}: libcuda+0x{fa:x} -> libcuda+0x{fb:x} " + f"({fb-fa:+#x})") + if resized: + lines.append(f"-- resized named (first {args.max_list}) --") + for n, fa, fb in resized[:args.max_list]: + lines.append(f" {n}: libcuda+0x{fa:x} {a[fa]['size']}B -> " + f"{b[fb]['size']}B") + if shifted_resized: + lines.append(f"-- resized anonymous (first {args.max_list}) --") + for fa, fb in shifted_resized[:args.max_list]: + lines.append(f" libcuda+0x{fa:x} {a[fa]['size']}B -> " + f"libcuda+0x{fb:x} {b[fb]['size']}B") + if added: + lines.append(f"-- added (first {args.max_list}) --") + lines.extend(" " + x for x in + hexlist(added)) + if removed: + lines.append(f"-- removed (first {args.max_list}) --") + lines.extend(" " + x for x in + hexlist(removed)) + text = "\n".join(lines) + + if args.out: + with open(args.out, "w") as f: + f.write(text + "\n") + else: + print(text) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/userspace/cuda/tools/ghidra/AnalyzeFunctions.java b/userspace/cuda/tools/ghidra/AnalyzeFunctions.java new file mode 100644 index 0000000..e0ae6cb --- /dev/null +++ b/userspace/cuda/tools/ghidra/AnalyzeFunctions.java @@ -0,0 +1,343 @@ +// SPDX-License-Identifier: BSD-3-Clause +// +// AnalyzeFunctions.java - Ghidra post-script: run the "Decompiler Parameter ID" +// analyzer over a caller-supplied list of functions, using the analyzer's +// AddressSetView entry point instead of a whole-program analysis pass. +// +// This is the targeted type-recovery pass for the CUDA userspace driver: the +// full auto-analysis deliberately disables the parameter-ID analyzer because it +// decompiles every function and is very expensive on a 29k-function stripped +// binary. Re-running it over a curated address set (crash path, engine-record +// path, handle pool, RM alloc, resolved wrapper implementations, the largest +// .text functions) recovers parameter counts/types and local variables where +// the decompiler can infer them, without touching the rest of the program. +// +// Usage: +// analyzeHeadless <proj_dir> <proj_name> -process <program> -noanalysis \ +// -scriptPath tools/ghidra \ +// -postScript AnalyzeFunctions.java <addresses_file> \ +// [--timeout <secs>] [--out <tsv>] [--list] [--batch <n>] +// +// <addresses_file> one entry per line: "0x<file_vaddr> [# comment]". +// Addresses are FILE virtual addresses (libcuda+0x...). +// Blank lines and lines starting with '#' are ignored. +// --timeout <secs> per-function decompiler timeout for the analyzer +// (default 60, matching the analyzer's own default). +// --out <tsv> write the before/after report to this file as well. +// --list only resolve the addresses to functions and print them; +// do not run the analyzer (dry run). +// --batch <n> run the analyzer in batches of n functions (default 0 = +// one batch). Useful to checkpoint very large sets. +// +// The script translates file vaddr -> program address via the program's image +// base (0x100000 for this binary), exactly like DecompileAt.java. If an address +// is not inside a known function it is disassembled and turned into one. +// +// The analyzer class was renamed between Ghidra versions +// (DecompilerParameterIdAnalyzer -> DecompilerFunctionAnalyzer); the script +// looks up both and invokes the AddressSetView-based entry point +// added(AddressSetView, TaskMonitor) when present, falling back to +// added(Program, AddressSetView, TaskMonitor, MessageLog). + +import ghidra.app.script.GhidraScript; +import ghidra.app.services.Analyzer; +import ghidra.app.util.importer.MessageLog; +import ghidra.framework.options.Options; +import ghidra.program.model.address.Address; +import ghidra.program.model.address.AddressSet; +import ghidra.program.model.address.AddressSetView; +import ghidra.program.model.listing.Function; +import ghidra.program.model.listing.FunctionIterator; +import ghidra.program.model.listing.Listing; +import ghidra.program.model.listing.Variable; +import ghidra.util.classfinder.ClassSearcher; +import ghidra.util.task.TaskMonitor; + +import java.io.BufferedReader; +import java.io.FileReader; +import java.io.FileWriter; +import java.io.PrintWriter; +import java.lang.reflect.InvocationTargetException; +import java.lang.reflect.Method; +import java.util.ArrayList; +import java.util.List; + +public class AnalyzeFunctions extends GhidraScript { + + private static final String[] ANALYZER_CLASSES = { + "ghidra.app.plugin.core.analysis.DecompilerFunctionAnalyzer", + "ghidra.app.plugin.core.analysis.DecompilerParameterIdAnalyzer", + }; + + private static final String OPTION_TIMEOUT = "Analysis Decompiler Timeout (sec)"; + + private static class Target { + long fileAddr; + Address entry; + String name; + int paramsBefore; + int localsBefore; + int paramsAfter = -1; + int localsAfter = -1; + } + + @Override + public void run() throws Exception { + String[] args = getScriptArgs(); + if (args.length < 1) { + println("usage: AnalyzeFunctions.java <addresses_file> [--timeout secs] " + + "[--out tsv] [--list] [--batch n]"); + return; + } + + String addrFile = args[0]; + int timeoutSecs = 60; + String outFile = null; + boolean listOnly = false; + int batchSize = 0; + for (int i = 1; i < args.length; i++) { + switch (args[i]) { + case "--timeout": + timeoutSecs = Integer.parseInt(args[++i]); + break; + case "--out": + outFile = args[++i]; + break; + case "--list": + listOnly = true; + break; + case "--batch": + batchSize = Integer.parseInt(args[++i]); + break; + default: + println("unknown argument: " + args[i]); + return; + } + } + + long imageBase = currentProgram.getImageBase().getOffset(); + List<Long> fileAddrs = readAddresses(addrFile); + println("AnalyzeFunctions: " + fileAddrs.size() + " address(es), image base 0x" + + Long.toHexString(imageBase)); + + List<Target> targets = resolveTargets(fileAddrs, imageBase); + println("AnalyzeFunctions: resolved " + targets.size() + " function(s)"); + + if (listOnly) { + for (Target t : targets) { + println(String.format(" libcuda+0x%06x %s program=0x%x params=%d locals=%d", + t.fileAddr, t.name, t.entry.getOffset(), t.paramsBefore, + t.localsBefore)); + } + return; + } + + Analyzer analyzer = newAnalyzer(); + if (analyzer == null) { + println("ERROR: no Decompiler Parameter ID analyzer class found"); + return; + } + println("AnalyzeFunctions: analyzer = " + analyzer.getClass().getName() + + ", timeout = " + timeoutSecs + "s"); + + Options opts = currentProgram.getOptions("Analyzers"); + analyzer.registerOptions(opts, currentProgram); + if (opts.contains(OPTION_TIMEOUT)) { + opts.setInt(OPTION_TIMEOUT, timeoutSecs); + } + analyzer.optionsChanged(opts, currentProgram); + + int txId = currentProgram.startTransaction("AnalyzeFunctions parameter ID"); + boolean commit = false; + try { + runAnalyzer(analyzer, targets, batchSize); + commit = true; + } + finally { + currentProgram.endTransaction(txId, commit); + } + + snapshotAfter(targets, imageBase); + report(targets, outFile); + } + + private List<Long> readAddresses(String path) throws Exception { + List<Long> out = new ArrayList<>(); + try (BufferedReader r = new BufferedReader(new FileReader(path))) { + String line; + while ((line = r.readLine()) != null) { + line = line.trim(); + if (line.isEmpty() || line.startsWith("#")) { + continue; + } + String tok = line.split("\\s+")[0]; + if (tok.startsWith("0x") || tok.startsWith("0X")) { + tok = tok.substring(2); + } + try { + out.add(Long.parseUnsignedLong(tok, 16)); + } + catch (NumberFormatException e) { + println("WARN: bad address line: " + line); + } + } + } + return out; + } + + private List<Target> resolveTargets(List<Long> fileAddrs, long imageBase) { + List<Target> targets = new ArrayList<>(); + Listing listing = currentProgram.getListing(); + for (long fileAddr : fileAddrs) { + Address addr = toAddr(fileAddr + imageBase); + Function f = getFunctionContaining(addr); + if (f == null) { + FunctionIterator it = listing.getFunctions(addr, false); + if (it.hasNext()) { + f = it.next(); + } + } + if (f == null) { + disassemble(addr); + f = createFunction(addr, null); + } + if (f == null) { + println("WARN: no function at libcuda+0x" + Long.toHexString(fileAddr)); + continue; + } + Target t = new Target(); + t.fileAddr = fileAddr; + t.entry = f.getEntryPoint(); + t.name = f.getName(); + t.paramsBefore = f.getParameterCount(); + t.localsBefore = f.getLocalVariables().length; + targets.add(t); + } + return targets; + } + + private Analyzer newAnalyzer() throws Exception { + // Preferred path: the plugin service discovers analyzers via ClassSearcher, + // which also works when the analyzer lives in a module that is not visible + // to the script's class loader (Class.forName then fails). + try { + ClassSearcher.search(monitor); + for (Class<? extends Analyzer> c : ClassSearcher.getClasses(Analyzer.class)) { + Object o = c.getDeclaredConstructor().newInstance(); + Analyzer a = (Analyzer) o; + if ("Decompiler Parameter ID".equals(a.getName())) { + return a; + } + } + } + catch (Exception e) { + println("WARN: ClassSearcher lookup failed: " + e); + } + // Fallback: direct reflection over the known class names. + for (String cn : ANALYZER_CLASSES) { + try { + Class<?> c = Class.forName(cn); + Object o = c.getDeclaredConstructor().newInstance(); + if (o instanceof Analyzer) { + return (Analyzer) o; + } + } + catch (ClassNotFoundException e) { + // try next candidate + } + } + return null; + } + + private void runAnalyzer(Analyzer analyzer, List<Target> targets, int batchSize) + throws Exception { + Method addedSet = findMethod(analyzer.getClass(), "added", + AddressSetView.class, TaskMonitor.class); + Method addedProgram = findMethod(analyzer.getClass(), "added", + ghidra.program.model.listing.Program.class, AddressSetView.class, + TaskMonitor.class, MessageLog.class); + if (addedSet == null && addedProgram == null) { + throw new NoSuchMethodException("no supported added(...) entry point on " + + analyzer.getClass().getName()); + } + + int n = (batchSize > 0) ? batchSize : targets.size(); + for (int start = 0; start < targets.size(); start += n) { + int end = Math.min(start + n, targets.size()); + AddressSet set = new AddressSet(); + for (int i = start; i < end; i++) { + set.addRange(targets.get(i).entry, targets.get(i).entry); + } + println(String.format("AnalyzeFunctions: batch %d..%d (%d functions)", + start, end - 1, end - start)); + try { + if (addedSet != null) { + addedSet.invoke(analyzer, set, monitor); + } + else { + addedProgram.invoke(analyzer, currentProgram, set, monitor, + new MessageLog()); + } + } + catch (InvocationTargetException e) { + Throwable cause = e.getCause(); + println("WARN: analyzer threw for batch " + start + ".." + (end - 1) + + ": " + (cause == null ? e : cause)); + } + } + } + + private Method findMethod(Class<?> c, String name, Class<?>... params) { + for (Class<?> k = c; k != null; k = k.getSuperclass()) { + try { + Method m = k.getDeclaredMethod(name, params); + m.setAccessible(true); + return m; + } + catch (NoSuchMethodException e) { + // keep walking up the hierarchy + } + } + return null; + } + + private void snapshotAfter(List<Target> targets, long imageBase) { + for (Target t : targets) { + Function f = getFunctionAt(t.entry); + if (f == null) { + continue; + } + t.paramsAfter = f.getParameterCount(); + Variable[] locals = f.getLocalVariables(); + t.localsAfter = locals.length; + if (t.name.startsWith("FUN_") && !f.getName().equals(t.name)) { + t.name = f.getName(); + } + } + } + + private void report(List<Target> targets, String outFile) throws Exception { + StringBuilder sb = new StringBuilder(); + sb.append("# libcuda 570.144 targeted Decompiler Parameter ID pass\n"); + sb.append("# file_vaddr\tname\tparams_before\tparams_after\tlocals_before\tlocals_after\n"); + int improved = 0; + for (Target t : targets) { + boolean gain = t.paramsAfter > t.paramsBefore || t.localsAfter > t.localsBefore; + if (gain) { + improved++; + } + sb.append(String.format("0x%06x\t%s\t%d\t%d\t%d\t%d%s\n", + t.fileAddr, t.name, t.paramsBefore, t.paramsAfter, + t.localsBefore, t.localsAfter, gain ? "\tGAIN" : "")); + } + sb.append(String.format("# %d/%d functions gained parameters or locals\n", + improved, targets.size())); + String text = sb.toString(); + println(text); + if (outFile != null) { + try (PrintWriter w = new PrintWriter(new FileWriter(outFile))) { + w.print(text); + } + } + } +} diff --git a/userspace/cuda/tools/name_candidates.py b/userspace/cuda/tools/name_candidates.py new file mode 100755 index 0000000..7a36a7c --- /dev/null +++ b/userspace/cuda/tools/name_candidates.py @@ -0,0 +1,224 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: BSD-3-Clause +# +# name_candidates.py - propose API-implementation names for the hand-written +# functions called by the generated CUDA API wrappers. +# +# Background: every exported cu* entry is a generated stub that jumps through +# an interposition slot to a generated wrapper (tracing/tool dispatch); the +# wrapper then calls the hand-written implementation of that API. Export +# families have several variants of one API (cuMemAlloc_v2, cuMemAlloc_ptsz, +# cuMemAlloc_ptds, ...), all of which must reach the same implementation. +# +# Method (robust, structure-only): +# 1. Group exported wrappers by BASE NAME: strip the variant suffixes +# `_v<N>`, `_ptsz`, `_ptds` (in any order/combination) from the export +# name; all exports in a group are variants of one API. +# 2. The implementation of a base name is a DIRECT CALLEE shared by ALL of +# its variants' wrappers, i.e. the intersection of the wrappers' direct +# callee sets. +# 3. Filter out global helpers: a candidate is dropped when it is called by +# >= --max-outside wrappers OUTSIDE its group (trace/init/lock helpers +# are called by dozens to hundreds of unrelated wrappers). Wrappers are +# never candidates for themselves. +# 4. Confidence: +# high - exactly one candidate survives, and either the group has +# >= 2 variants (the survivor is shared by all of them) or the +# survivor is a unique callee (called by exactly one wrapper +# in the whole library); +# medium - exactly one survivor but it is also called by another +# wrapper (possible shared implementation), or several +# survivors remain (the best-ranked one is proposed). +# Ranking among several survivors: fewest outside wrappers first, then +# largest direct-callee count (implementations do work; helpers do not), +# then lowest address, for a deterministic result. +# +# Inputs: +# <exports.tsv> elf-index exports table: addr file_addr size name +# <callgraph.tsv> elf-index callgraph: caller callee caller_file callee_file +# <names.tsv> name map: file_vaddr<TAB>name (wrapper_<export> entries +# provide the wrapper address of each export) +# <out.tsv> output (comments + one row per proposal) +# +# Output rows (TSV): +# file_vaddr proposed_name confidence evidence status +# `proposed_name` is `impl_<base>`; `evidence` carries the wrapper group, +# candidate callees before/after filtering and the outside-wrapper count; +# `status` is always `unverified` (the tool proposes, it does not prove). +# All addresses are FILE virtual addresses (libcuda+0x...). +# +# Usage: +# name_candidates.py <exports.tsv> <callgraph.tsv> <names.tsv> <out.tsv> +# [--max-outside N] [--quiet] + +import argparse +import re +import sys +from collections import defaultdict + + +SUFFIX_RE = re.compile(r"(_ptsz|_ptds)$|_v\d+$") + + +def base_name(name): + """Strip CUDA variant suffixes (_v<N>, _ptsz, _ptds) from an export name.""" + prev = None + while prev != name: + prev = name + name = SUFFIX_RE.sub("", name) + return name + + +def load_exports(path): + rows = [] + with open(path) as f: + f.readline() + for line in f: + p = line.rstrip("\n").split("\t") + if len(p) < 4: + continue + try: + rows.append(dict(addr=int(p[0], 16), file=int(p[1], 16), + size=int(p[2]), name=p[3])) + except ValueError: + continue + return rows + + +def load_names(path): + names = {} + with open(path) as f: + for line in f: + line = line.strip() + if not line or line.startswith("#"): + continue + p = line.split("\t") + if len(p) >= 2: + try: + names[int(p[0], 16)] = p[1] + except ValueError: + pass + return names + + +def load_callgraph(path): + callees = defaultdict(set) + callers = defaultdict(set) + with open(path) as f: + f.readline() + for line in f: + p = line.rstrip("\n").split("\t") + if len(p) < 4: + continue + try: + cf, kf = int(p[2], 16), int(p[3], 16) + except ValueError: + continue + callees[cf].add(kf) + callers[kf].add(cf) + return callees, callers + + +def hx(a): + return f"0x{a:x}" + + +def main(): + ap = argparse.ArgumentParser( + description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter) + ap.add_argument("exports") + ap.add_argument("callgraph") + ap.add_argument("names") + ap.add_argument("out") + ap.add_argument("--max-outside", type=int, default=2, + help="drop a candidate called by >= N wrappers outside " + "its base-name group (default 2)") + ap.add_argument("--quiet", action="store_true") + args = ap.parse_args() + + exports = load_exports(args.exports) + names = load_names(args.names) + callees, callers = load_callgraph(args.callgraph) + + wrapper_by_export = {} + for a, n in names.items(): + if n.startswith("wrapper_"): + wrapper_by_export[n[len("wrapper_"):]] = a + wrapper_set = set(wrapper_by_export.values()) + + groups = defaultdict(list) + for e in exports: + w = wrapper_by_export.get(e["name"]) + if w is not None: + groups[base_name(e["name"])].append((e["name"], w)) + + rows = [] + n_high = n_med = 0 + for group in sorted(groups): + variants = groups[group] + if not variants: + continue + variant_wrappers = {w for _, w in variants} + inter = None + for _, w in variants: + cs = set(callees.get(w, ())) + inter = cs if inter is None else (inter & cs) + if not inter: + continue + before = sorted(inter) + survivors = [] + for c in before: + if c in wrapper_set: + continue + outside = [w for w in callers.get(c, ()) + if w in wrapper_set and w not in variant_wrappers] + if len(outside) >= args.max_outside: + continue + survivors.append((len(outside), c)) + if not survivors: + continue + # rank: fewest outside wrappers, then most direct callees, then address + survivors.sort(key=lambda t: (t[0], -len(callees.get(t[1], ())), t[1])) + best_outside, best = survivors[0] + after = [c for _, c in survivors] + if len(survivors) == 1 and (len(variants) >= 2 or + len(callers.get(best, ()) & wrapper_set) == 1): + confidence = "high" + n_high += 1 + else: + confidence = "medium" + n_med += 1 + evidence = ( + f"group={group} variants=" + ",".join(n for n, _ in variants) + + "; before=" + ",".join(hx(c) for c in before) + + "; after=" + ",".join(hx(c) for c in after) + + f"; outside_wrappers={best_outside}") + rows.append((best, f"impl_{group}", confidence, evidence)) + + rows.sort() + with open(args.out, "w") as f: + f.write("# Open-IOV CUDA API-implementation candidates (our own work product; tracked).\n") + f.write("# Generated by tools/name_candidates.py from the elf-index exports,\n") + f.write("# callgraph and the version's names.tsv; see that script for the method.\n") + f.write("#\n") + f.write("# file_vaddr<TAB>proposed_name<TAB>confidence<TAB>evidence<TAB>status\n") + f.write("# confidence: high = exactly one filtered candidate, shared by >=2 variants\n") + f.write("# or a unique callee; medium = one survivor also called by another\n") + f.write("# wrapper, or several survivors (best-ranked proposed).\n") + f.write("# status: unverified (tool output). Verification = read the body in the\n") + f.write("# decompiled corpus and mark verified/corrected here; only verified\n") + f.write("# names may be copied into names.tsv.\n") + f.write("#\n") + for a, name, conf, ev in rows: + f.write(f"0x{a:08x}\t{name}\t{conf}\t{ev}\tunverified\n") + + if not args.quiet: + print(f"[name_candidates] groups={len(groups)} proposals={len(rows)} " + f"(high={n_high} medium={n_med}) -> {args.out}", + file=sys.stderr) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/userspace/cuda/tools/subsystem_digest.py b/userspace/cuda/tools/subsystem_digest.py new file mode 100755 index 0000000..68f0d09 --- /dev/null +++ b/userspace/cuda/tools/subsystem_digest.py @@ -0,0 +1,190 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: BSD-3-Clause +# +# subsystem_digest.py - per-subsystem call-tree digest from the tracked +# subsystem annotations plus the callgraph index. +# +# For every subsystem in subsystems.tsv it prints: +# * members (address, name, role) - the annotated functions; +# * entry functions - members that are not called by another member of the +# same subsystem (the roots of the subsystem's call tree); +# * a depth-limited call tree from each entry, with the direct callees of a +# node grouped by their subsystem; +# * a cross-subsystem edge summary (how many calls each subsystem makes into +# every other one) - the fastest way to see the layering. +# +# The digest is a navigation aid: it shows where to start reading the +# decompiled corpus for a subsystem and which other subsystems it touches, +# without opening Ghidra. Addresses are FILE virtual addresses (libcuda+0x...). +# +# Usage: +# subsystem_digest.py <subsystems.tsv> <callgraph.tsv> [--out FILE] +# [--depth N] [--max-per-node N] [--subsystem NAME] +# +# <subsystems.tsv> format (comments allowed): +# file_vaddr<TAB>name<TAB>subsystem<TAB>role<TAB>evidence +# <callgraph.tsv> format (header): +# caller<TAB>callee<TAB>caller_file<TAB>callee_file +# +# SPDX-License-Identifier: BSD-3-Clause + +import argparse +import sys +from collections import defaultdict + + +def load_subsystems(path): + members = {} + with open(path) as f: + for line in f: + line = line.strip() + if not line or line.startswith("#"): + continue + p = line.split("\t") + if len(p) < 4: + continue + try: + addr = int(p[0], 16) + except ValueError: + continue + members[addr] = dict(name=p[1], subsystem=p[2], + role=p[3], + evidence=p[4] if len(p) > 4 else "") + return members + + +def load_callgraph(path): + callees = defaultdict(set) + callers = defaultdict(set) + with open(path) as f: + f.readline() + for line in f: + p = line.rstrip("\n").split("\t") + if len(p) < 4: + continue + try: + cf, kf = int(p[2], 16), int(p[3], 16) + except ValueError: + continue + callees[cf].add(kf) + callers[kf].add(cf) + return callees, callers + + +def main(): + ap = argparse.ArgumentParser( + description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter) + ap.add_argument("subsystems") + ap.add_argument("callgraph") + ap.add_argument("--out", default=None, + help="write the digest here (default: stdout)") + ap.add_argument("--depth", type=int, default=2, + help="call-tree depth below each entry (default 2)") + ap.add_argument("--max-per-node", type=int, default=8, + help="max callees shown per node (default 8)") + ap.add_argument("--subsystem", default=None, + help="only digest this subsystem") + args = ap.parse_args() + + members = load_subsystems(args.subsystems) + callees, callers = load_callgraph(args.callgraph) + + by_sub = defaultdict(dict) + for addr, m in members.items(): + by_sub[m["subsystem"]][addr] = m + + def name_of(a): + m = members.get(a) + if m: + return f"{m['name']} (libcuda+0x{a:x})" + return f"libcuda+0x{a:x}" + + def sub_of(a): + m = members.get(a) + return m["subsystem"] if m else "unknown" + + out = [] + out.append("# Subsystem call-tree digest (generated by tools/subsystem_digest.py)") + out.append(f"# subsystems.tsv: {args.subsystems}") + out.append(f"# callgraph.tsv : {args.callgraph}") + out.append("# Addresses are FILE virtual addresses (libcuda+0x...).") + out.append("# entry = annotated member with no caller inside its own subsystem.") + out.append("") + + # cross-subsystem edge counts + cross = defaultdict(int) + for addr, m in members.items(): + for c in callees.get(addr, ()): + cross[(m["subsystem"], sub_of(c))] += 1 + + selected = sorted(by_sub) + if args.subsystem: + selected = [s for s in selected if s == args.subsystem] + if not selected: + print(f"error: subsystem {args.subsystem!r} not in {args.subsystems}", + file=sys.stderr) + return 1 + + for sub in selected: + ms = by_sub[sub] + out.append("=" * 72) + out.append(f"## {sub} ({len(ms)} annotated members)") + out.append("=" * 72) + out.append("members:") + for a in sorted(ms): + m = ms[a] + out.append(f" libcuda+0x{a:x}\t{m['name']}\t{m['role']}") + entries = sorted(a for a in ms + if not (callers.get(a, set()) & set(ms))) + if not entries: + entries = sorted(ms) + out.append("entries: (no entry: every member is called from inside " + "the subsystem; showing all)") + else: + out.append("entries:") + for a in entries: + out.append(f" {name_of(a)}: {ms[a]['role']}") + out.append("call tree:") + + def walk(a, depth, indent, seen): + kids = sorted(callees.get(a, ())) + if not kids or depth <= 0: + return + groups = defaultdict(list) + for k in kids: + groups[sub_of(k)].append(k) + for g in sorted(groups): + shown = groups[g][:args.max_per_node] + more = f" (+{len(groups[g]) - args.max_per_node} more)" \ + if len(groups[g]) > args.max_per_node else "" + out.append(f"{indent}-> {g}: " + + "; ".join(name_of(k) for k in shown) + more) + for k in kids: + if k in seen: + continue + seen.add(k) + walk(k, depth - 1, indent + " ", seen) + + for a in entries: + out.append(f" {name_of(a)}") + walk(a, args.depth, " ", {a}) + out.append("") + + out.append("=" * 72) + out.append("## cross-subsystem calls (edges from -> to)") + out.append("=" * 72) + for (src, dst) in sorted(cross): + out.append(f" {src:<16} -> {dst:<16} {cross[(src, dst)]}") + + text = "\n".join(out) + if args.out: + with open(args.out, "w") as f: + f.write(text + "\n") + else: + print(text) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) |