diff options
| author | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-09-23 13:48:00 -0400 |
|---|---|---|
| committer | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-09-23 13:48:00 -0400 |
| commit | c587eac377c1c5b720142edadebaace5cb1b17c0 (patch) | |
| tree | c381be95dd02151022f58c3e84c6eeb0dfe91b25 /userspace | |
| parent | bd6c3ee69d7157e2dabf9350be52367c6a6e3c42 (diff) | |
| download | open-iov-c587eac377c1c5b720142edadebaace5cb1b17c0.tar.gz open-iov-c587eac377c1c5b720142edadebaace5cb1b17c0.zip | |
userspace/cuda: complete the analysis pass
* Name recovery: name_candidates.py proposes the hand-written API-layer
implementation for each exported wrapper by grouping version variants and
intersecting their callees; 180 high-confidence proposals recorded in
annotations/name-candidates.tsv, 29 of them body-verified and promoted to
names.tsv with evidence in subsystems.tsv.
* Engine records: identified the 0x14-byte per-engine descriptor, mapped its
fields (kind, engine type, instance, class id, CE caps) and the existence
callback, and documented how the missing async copy engine (COPY14) leaves
the secondary record zeroed and blocks context creation (docs/05 section
3.7, public-safe).
* Targeted type recovery: AnalyzeFunctions.java runs the parameter-ID
analyzer over a curated list of 80 functions; 16 improved materially and
the typed decompilations are kept under the version run tree.
* Cross-version: cudatool.py fdematch measures how names transfer between
patch versions (570.144 -> 570.211.01: exports 100%, verified internals
~70%, generated wrappers must be re-resolved by relocation, never by
address arithmetic); documented in the upgrade playbook.
* Corpus tooling: diff_corpus.py (version-to-version corpus diff) and
subsystem_digest.py (per-subsystem call trees).
* Deprecate the old per-version helpers whose address handling and analysis
timeout produced wrong results.
Generated corpora and typed outputs remain untracked under the version run
tree.
Diffstat (limited to 'userspace')
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()) |