userspace/cuda: add CUDA userspace driver analysis, tooling and name maps

Add the analysis of the NVIDIA CUDA userspace driver used by Open-IOV VMs:

* docs/01-08: architecture (layers, object model, threads, sync, memory,
  channels, launch), build and execution model, the driver/kernel/GPU
  interface contract, virtualization approaches, the 570.144 decompilation
  map with the corrected libcuda crash path and the current engine-topology
  blocker, tooling, and a version-upgrade playbook.
* tools/: version-agnostic toolchain - stdlib ELF/interface toolkit
  (exports, unwind-table function map, wrapper resolution, version delta,
  crash address decode), callgraph walker, corpus annotator, and Ghidra
  analysis/decompilation/name-map scripts with resumable parallel drivers.
* 570.144/annotations/: our name map (exports + resolved wrappers + verified
  internals) and subsystem index, tracked so naming work accumulates.
* READ-ME-FIRST.md: reading guide with the fast lookup recipes.

The full decompiled corpus and annotated corpus stay untracked (vendor-
derived, gitignored under the version run/ tree).

30 files changed, 4979 insertions, 1 deletions

diff --git a/userspace/cuda/570.144/.gitignore b/userspace/cuda/570.144/.gitignore
index c5997d3..2427d64 100644
--- a/userspace/cuda/570.144/.gitignore
+++ b/userspace/cuda/570.144/.gitignore
@@ -11,8 +11,16 @@ run/*.txt
 run/*.asm
 run/*.json
 run/*.log
+run/*.out
 
-# Ghidra projects
+# Ghidra projects (single-function and full-analysis)
 run/*.rep/
 run/*.gpr
 run/ghidra-*/
+run/ghidra-full/
+
+# Full-corpus analysis outputs (index TSVs, decompiled buckets, project copies)
+run/analysis/
+
+# Analysis driver logs
+run/analysis-*.log
diff --git a/userspace/cuda/570.144/README.md b/userspace/cuda/570.144/README.md
index f32b762..e2e56fd 100644
--- a/userspace/cuda/570.144/README.md
+++ b/userspace/cuda/570.144/README.md
@@ -2,6 +2,19 @@
 
 Tools and notes for **debugging CUDA execution inside Open-IOV guests**.
 
+> **Full analysis now lives one level up.** See `../README.md` and
+> `../docs/` for the architecture documents, the complete 570.144
+> decompilation (29,254 functions, annotated corpus under
+> `run/analysis/decompiled-annotated/`), the corrected crash-path analysis
+> (`../docs/05-570.144-decompilation.md`) and the reusable tooling
+> (`../tools/`).
+>
+> **Address warning:** Ghidra loads this binary at image base `0x100000`.
+> `scripts/ghidra_decompile_project.sh` passed file vaddrs straight to Ghidra,
+> so decompilations produced with it (e.g. `run/helper_453ef0.c`) are of
+> addresses 0x100000 too low. Use `../tools/ghidra_decompile_at.sh`, which
+> translates file vaddr → Ghidra address automatically.
+
 ## Purpose (read this first)
 
 These tools exist to make virtualized CUDA *work correctly and fast* in
diff --git a/userspace/cuda/570.144/annotations/README.md b/userspace/cuda/570.144/annotations/README.md
new file mode 100644
index 0000000..868c503
--- /dev/null
+++ b/userspace/cuda/570.144/annotations/README.md
@@ -0,0 +1,35 @@
+# Open-IOV CUDA analysis annotations (tracked)
+
+These files are our own work product: names, roles and evidence for functions
+in the CUDA userspace driver. They are tracked in git so that analysis effort
+accumulates across driver versions.
+
+## Files
+
+| 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`. |
+| `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. |
+
+Addresses are **file virtual addresses** (`libcuda+0x...`). Apply the map with
+`tools/ghidra/ApplyNameMap.java`; see `docs/06-tooling.md`.
+
+## How names.tsv is generated
+
+1. **Exports** — exact, from `.dynsym` (`tools/cudatool.py exports`).
+2. **Wrappers** — the generated `cu*` stubs jump through an interposition slot
+   whose initial value is a relocation to the generated wrapper function.
+   `tools/cudatool.py wrappers` resolves all of them; they are named
+   `wrapper_<export>`.
+3. **Verified internals** — functions read and understood during the
+   decompilation effort. Each is added by hand with evidence recorded in
+   `subsystems.tsv` and the public docs.
+4. **Transfers** — when a new driver version is analysed, apply the previous
+   map, check how many names land (`ApplyNameMap.java --dry-run`), fix the
+   misses, and record the delta for the version.
+
+## Rules
+
+* Never guess silently. A name without evidence is worse than `FUN_...`; put
+  open questions in `subsystems.tsv` with `evidence=todo` instead.
+* Names must describe behaviour, not vendor internals we cannot justify.
diff --git a/userspace/cuda/570.144/annotations/names.tsv b/userspace/cuda/570.144/annotations/names.tsv
new file mode 100644
index 0000000..403bef3
--- /dev/null
+++ b/userspace/cuda/570.144/annotations/names.tsv
@@ -0,0 +1,1304 @@
+# Open-IOV CUDA 570.144 name map (our own work product; tracked).
+# Format: <file_vaddr>\t<name>\t<source>
+# Addresses are FILE virtual addresses (libcuda+0x...). Apply with
+# tools/ghidra/ApplyNameMap.java; see annotations/README.md.
+#
+# --- exports (exact, from .dynsym) ---
+0x001dc040	cuMemGetAttribute	export
+0x001dc050	cuMemGetAttribute_v2	export
+0x002414b0	cudbgReportDriverApiError	export
+0x002414c0	cudbgReportDriverInternalError	export
+0x00249c70	cudbgGetAPI	export
+0x00249e40	cudbgGetAPIVersion	export
+0x0024b1f0	cudbgApiAttach	export
+0x0024b2a0	cudbgApiDetach	export
+0x0025bbd0	cudbgApiInit	export
+0x0026e470	cuEGLApiInit	export
+0x00310020	cuGetErrorString	export
+0x00310040	cuGetErrorName	export
+0x00310060	cuInit	export
+0x00310080	cuDriverGetVersion	export
+0x003100a0	cuDeviceGet	export
+0x003100c0	cuDeviceGetCount	export
+0x003100e0	cuDeviceGetName	export
+0x00310100	cuDeviceGetUuid	export
+0x00310120	cuDeviceGetUuid_v2	export
+0x00310140	cuDeviceGetLuid	export
+0x00310160	cuDeviceTotalMem_v2	export
+0x00310180	cuDeviceGetTexture1DLinearMaxWidth	export
+0x003101a0	cuDeviceGetAttribute	export
+0x003101c0	cuDeviceGetNvSciSyncAttributes	export
+0x003101e0	cuDeviceSetMemPool	export
+0x00310200	cuDeviceGetMemPool	export
+0x00310220	cuDeviceGetDefaultMemPool	export
+0x00310240	cuDeviceGetExecAffinitySupport	export
+0x00310260	cuFlushGPUDirectRDMAWrites	export
+0x00310280	cuDeviceGetProperties	export
+0x003102a0	cuDeviceComputeCapability	export
+0x003102c0	cuDevicePrimaryCtxRetain	export
+0x003102e0	cuDevicePrimaryCtxRelease_v2	export
+0x00310300	cuDevicePrimaryCtxSetFlags_v2	export
+0x00310320	cuDevicePrimaryCtxGetState	export
+0x00310340	cuDevicePrimaryCtxReset_v2	export
+0x00310360	cuCtxCreate_v2	export
+0x00310380	cuCtxCreate_v3	export
+0x003103a0	cuCtxCreate_v4	export
+0x003103c0	cuCtxDestroy_v2	export
+0x003103e0	cuCtxPushCurrent_v2	export
+0x00310400	cuCtxPopCurrent_v2	export
+0x00310420	cuCtxSetCurrent	export
+0x00310440	cuCtxGetCurrent	export
+0x00310460	cuCtxGetDevice	export
+0x00310480	cuCtxGetFlags	export
+0x003104a0	cuCtxSetFlags	export
+0x003104c0	cuCtxGetId	export
+0x003104e0	cuCtxSynchronize	export
+0x00310500	cuCtxSetLimit	export
+0x00310520	cuCtxGetLimit	export
+0x00310540	cuCtxGetCacheConfig	export
+0x00310560	cuCtxSetCacheConfig	export
+0x00310580	cuCtxGetApiVersion	export
+0x003105a0	cuCtxGetStreamPriorityRange	export
+0x003105c0	cuCtxResetPersistingL2Cache	export
+0x003105e0	cuCtxGetExecAffinity	export
+0x00310600	cuCtxRecordEvent	export
+0x00310620	cuCtxWaitEvent	export
+0x00310640	cuCtxAttach	export
+0x00310660	cuCtxDetach	export
+0x00310680	cuCtxGetSharedMemConfig	export
+0x003106a0	cuCtxSetSharedMemConfig	export
+0x003106c0	cuModuleLoad	export
+0x003106e0	cuModuleLoadData	export
+0x00310700	cuModuleLoadDataEx	export
+0x00310720	cuModuleLoadFatBinary	export
+0x00310740	cuModuleUnload	export
+0x00310760	cuModuleGetLoadingMode	export
+0x00310780	cuModuleGetFunction	export
+0x003107a0	cuModuleGetFunctionCount	export
+0x003107c0	cuModuleEnumerateFunctions	export
+0x003107e0	cuModuleGetGlobal_v2	export
+0x00310800	cuLinkCreate_v2	export
+0x00310820	cuLinkAddData_v2	export
+0x00310840	cuLinkAddFile_v2	export
+0x00310860	cuLinkComplete	export
+0x00310880	cuLinkDestroy	export
+0x003108a0	cuModuleGetTexRef	export
+0x003108c0	cuModuleGetSurfRef	export
+0x003108e0	cuLibraryLoadData	export
+0x00310900	cuLibraryLoadFromFile	export
+0x00310920	cuLibraryUnload	export
+0x00310940	cuLibraryGetKernel	export
+0x00310960	cuLibraryGetKernelCount	export
+0x00310980	cuLibraryEnumerateKernels	export
+0x003109a0	cuLibraryGetModule	export
+0x003109c0	cuKernelGetFunction	export
+0x003109e0	cuKernelGetLibrary	export
+0x00310a00	cuLibraryGetGlobal	export
+0x00310a20	cuLibraryGetManaged	export
+0x00310a40	cuLibraryGetUnifiedFunction	export
+0x00310a60	cuKernelGetAttribute	export
+0x00310a80	cuKernelSetAttribute	export
+0x00310aa0	cuKernelSetCacheConfig	export
+0x00310ac0	cuKernelGetName	export
+0x00310ae0	cuKernelGetParamInfo	export
+0x00310b00	cuMemGetInfo_v2	export
+0x00310b20	cuMemAlloc_v2	export
+0x00310b40	cuMemAllocPitch_v2	export
+0x00310b60	cuMemFree_v2	export
+0x00310b80	cuMemGetAddressRange_v2	export
+0x00310ba0	cuMemAllocHost_v2	export
+0x00310bc0	cuMemFreeHost	export
+0x00310be0	cuMemHostAlloc	export
+0x00330590	cuMemHostGetDevicePointer_v2	export
+0x003305b0	cuMemHostGetFlags	export
+0x003305d0	cuMemAllocManaged	export
+0x003305f0	cuDeviceRegisterAsyncNotification	export
+0x00330610	cuDeviceUnregisterAsyncNotification	export
+0x00330630	cuDeviceGetByPCIBusId	export
+0x00330650	cuDeviceGetPCIBusId	export
+0x00330670	cuIpcGetEventHandle	export
+0x00330690	cuIpcOpenEventHandle	export
+0x003306b0	cuIpcGetMemHandle	export
+0x003306d0	cuIpcOpenMemHandle_v2	export
+0x003306f0	cuIpcCloseMemHandle	export
+0x00330710	cuMemHostRegister_v2	export
+0x00330730	cuMemHostUnregister	export
+0x00330750	cuMemcpy_ptds	export
+0x00330770	cuMemcpyPeer_ptds	export
+0x00330790	cuMemcpyHtoD_v2_ptds	export
+0x003307b0	cuMemcpyDtoH_v2_ptds	export
+0x003307d0	cuMemcpyDtoD_v2_ptds	export
+0x003307f0	cuMemcpyDtoA_v2_ptds	export
+0x00330810	cuMemcpyAtoD_v2_ptds	export
+0x00330830	cuMemcpyHtoA_v2_ptds	export
+0x00330850	cuMemcpyAtoH_v2_ptds	export
+0x00330870	cuMemcpyAtoA_v2_ptds	export
+0x00330890	cuMemcpy2D_v2_ptds	export
+0x003308b0	cuMemcpy2DUnaligned_v2_ptds	export
+0x003308d0	cuMemcpy3D_v2_ptds	export
+0x003308f0	cuMemcpy3DPeer_ptds	export
+0x00330910	cuMemcpyAsync_ptsz	export
+0x00330930	cuMemcpyPeerAsync_ptsz	export
+0x00330950	cuMemcpyHtoDAsync_v2_ptsz	export
+0x00330970	cuMemcpyDtoHAsync_v2_ptsz	export
+0x00330990	cuMemcpyDtoDAsync_v2_ptsz	export
+0x003309b0	cuMemcpyHtoAAsync_v2_ptsz	export
+0x003309d0	cuMemcpyAtoHAsync_v2_ptsz	export
+0x003309f0	cuMemcpy2DAsync_v2_ptsz	export
+0x00330a10	cuMemcpy3DAsync_v2_ptsz	export
+0x00330a30	cuMemcpy3DPeerAsync_ptsz	export
+0x00330a50	cuMemcpyBatchAsync_ptsz	export
+0x00330a70	cuMemcpy3DBatchAsync_ptsz	export
+0x00330a90	cuMemsetD8_v2_ptds	export
+0x00330ac0	cuMemsetD16_v2_ptds	export
+0x00330ae0	cuMemsetD32_v2_ptds	export
+0x00330b00	cuMemsetD2D8_v2_ptds	export
+0x00330b20	cuMemsetD2D16_v2_ptds	export
+0x00330b40	cuMemsetD2D32_v2_ptds	export
+0x00330b60	cuMemsetD8Async_ptsz	export
+0x00330b90	cuMemsetD16Async_ptsz	export
+0x00330bb0	cuMemsetD32Async_ptsz	export
+0x00330bd0	cuMemsetD2D8Async_ptsz	export
+0x00330bf0	cuMemsetD2D16Async_ptsz	export
+0x00330c10	cuMemsetD2D32Async_ptsz	export
+0x00330c30	cuArrayCreate_v2	export
+0x00330c50	cuArrayGetDescriptor_v2	export
+0x00330c70	cuArrayGetSparseProperties	export
+0x00330c90	cuMipmappedArrayGetSparseProperties	export
+0x00330cb0	cuArrayGetMemoryRequirements	export
+0x00330cd0	cuMipmappedArrayGetMemoryRequirements	export
+0x00330cf0	cuArrayGetPlane	export
+0x00330d10	cuArrayDestroy	export
+0x00330d30	cuArray3DCreate_v2	export
+0x00330d50	cuArray3DGetDescriptor_v2	export
+0x00330d70	cuMipmappedArrayCreate	export
+0x00330d90	cuMipmappedArrayGetLevel	export
+0x00330db0	cuMipmappedArrayDestroy	export
+0x00330dd0	cuMemGetHandleForAddressRange	export
+0x00330df0	cuMemBatchDecompressAsync_ptsz	export
+0x00330e10	cuMemAddressReserve	export
+0x00330e30	cuMemAddressFree	export
+0x00330e50	cuMemCreate	export
+0x00330e70	cuMemRelease	export
+0x00330e90	cuMemMap	export
+0x00330eb0	cuMemMapArrayAsync_ptsz	export
+0x00330ed0	cuMemUnmap	export
+0x00330ef0	cuMemSetAccess	export
+0x00330f10	cuMemGetAccess	export
+0x00330f30	cuMemExportToShareableHandle	export
+0x00330f50	cuMemImportFromShareableHandle	export
+0x00330f70	cuMemGetAllocationGranularity	export
+0x00330f90	cuMemGetAllocationPropertiesFromHandle	export
+0x00330fb0	cuMemRetainAllocationHandle	export
+0x00330fd0	cuMemFreeAsync_ptsz	export
+0x00330ff0	cuMemAllocAsync_ptsz	export
+0x00331010	cuMemPoolTrimTo	export
+0x00331030	cuMemPoolSetAttribute	export
+0x00331050	cuMemPoolGetAttribute	export
+0x00331070	cuMemPoolSetAccess	export
+0x00331090	cuMemPoolGetAccess	export
+0x003310b0	cuMemPoolCreate	export
+0x003310d0	cuMemPoolDestroy	export
+0x003310f0	cuMemAllocFromPoolAsync_ptsz	export
+0x00331110	cuMemPoolExportToShareableHandle	export
+0x00331130	cuMemPoolImportFromShareableHandle	export
+0x00331150	cuMemPoolExportPointer	export
+0x00331170	cuMemPoolImportPointer	export
+0x00331190	cuMulticastCreate	export
+0x003311b0	cuMulticastAddDevice	export
+0x003311d0	cuMulticastBindMem	export
+0x003311f0	cuMulticastBindAddr	export
+0x00331210	cuMulticastUnbind	export
+0x00331230	cuMulticastGetGranularity	export
+0x00331250	cuPointerGetAttribute	export
+0x00331270	cuMemPrefetchAsync_ptsz	export
+0x00331290	cuMemPrefetchAsync_v2_ptsz	export
+0x003312b0	cuMemAdvise	export
+0x003312d0	cuMemAdvise_v2	export
+0x003312f0	cuMemRangeGetAttribute	export
+0x00331310	cuMemRangeGetAttributes	export
+0x00331330	cuPointerSetAttribute	export
+0x00331350	cuPointerGetAttributes	export
+0x00331370	cuStreamCreate	export
+0x00331390	cuStreamCreateWithPriority	export
+0x003313b0	cuStreamGetPriority_ptsz	export
+0x003313d0	cuStreamGetDevice_ptsz	export
+0x003313f0	cuStreamGetFlags_ptsz	export
+0x00331410	cuStreamGetId_ptsz	export
+0x00331430	cuStreamGetCtx_ptsz	export
+0x00331450	cuStreamGetCtx_v2_ptsz	export
+0x00331470	cuStreamWaitEvent_ptsz	export
+0x00331490	cuStreamAddCallback_ptsz	export
+0x003314b0	cuStreamBeginCapture_v2_ptsz	export
+0x003314d0	cuStreamBeginCaptureToGraph_ptsz	export
+0x003314f0	cuThreadExchangeStreamCaptureMode	export
+0x00331510	cuStreamEndCapture_ptsz	export
+0x00331530	cuStreamIsCapturing_ptsz	export
+0x00331550	cuStreamGetCaptureInfo_v2_ptsz	export
+0x00331570	cuStreamGetCaptureInfo_v3_ptsz	export
+0x00331590	cuStreamUpdateCaptureDependencies_ptsz	export
+0x003315b0	cuStreamUpdateCaptureDependencies_v2_ptsz	export
+0x003315d0	cuStreamAttachMemAsync_ptsz	export
+0x003315f0	cuStreamQuery_ptsz	export
+0x00331610	cuStreamSynchronize_ptsz	export
+0x00331630	cuStreamDestroy_v2	export
+0x00331650	cuStreamCopyAttributes_ptsz	export
+0x00331670	cuStreamGetAttribute_ptsz	export
+0x00331690	cuStreamSetAttribute_ptsz	export
+0x003316b0	cuEventCreate	export
+0x003316d0	cuEventRecord_ptsz	export
+0x003316f0	cuEventRecordWithFlags_ptsz	export
+0x00331710	cuEventQuery	export
+0x00331730	cuEventSynchronize	export
+0x00331750	cuEventDestroy_v2	export
+0x00331770	cuEventElapsedTime	export
+0x00331790	cuEventElapsedTime_v2	export
+0x003317b0	cuImportExternalMemory	export
+0x003317d0	cuExternalMemoryGetMappedBuffer	export
+0x003317f0	cuExternalMemoryGetMappedMipmappedArray	export
+0x00331810	cuDestroyExternalMemory	export
+0x00331830	cuImportExternalSemaphore	export
+0x00331850	cuSignalExternalSemaphoresAsync_ptsz	export
+0x00331870	cuWaitExternalSemaphoresAsync_ptsz	export
+0x00331890	cuDestroyExternalSemaphore	export
+0x003318b0	cuStreamWaitValue32_v2_ptsz	export
+0x003318d0	cuStreamWaitValue64_v2_ptsz	export
+0x003318f0	cuStreamWriteValue32_v2_ptsz	export
+0x00331910	cuStreamWriteValue64_v2_ptsz	export
+0x00331930	cuStreamBatchMemOp_v2_ptsz	export
+0x00331950	cuFuncGetAttribute	export
+0x00331970	cuFuncSetAttribute	export
+0x00331990	cuFuncSetCacheConfig	export
+0x003319b0	cuFuncGetModule	export
+0x003319d0	cuFuncGetName	export
+0x003319f0	cuFuncGetParamInfo	export
+0x00331a10	cuFuncIsLoaded	export
+0x00331a30	cuFuncLoad	export
+0x00331a50	cuLaunchKernel_ptsz	export
+0x00331a70	cuLaunchKernelEx_ptsz	export
+0x00331a90	cuLaunchCooperativeKernel_ptsz	export
+0x00331ab0	cuLaunchCooperativeKernelMultiDevice	export
+0x00331ad0	cuLaunchHostFunc_ptsz	export
+0x00331af0	cuFuncSetBlockShape	export
+0x00331b10	cuFuncSetSharedSize	export
+0x00331b30	cuParamSetSize	export
+0x00331b50	cuParamSeti	export
+0x00331b70	cuParamSetf	export
+0x00331b90	cuParamSetv	export
+0x00331bb0	cuLaunch	export
+0x00331bd0	cuLaunchGrid	export
+0x00331bf0	cuLaunchGridAsync	export
+0x00331c10	cuParamSetTexRef	export
+0x00331c30	cuFuncSetSharedMemConfig	export
+0x00331c50	cuGraphCreate	export
+0x00331c70	cuGraphAddKernelNode_v2	export
+0x0034f650	cuGraphKernelNodeGetParams_v2	export
+0x0034f670	cuGraphKernelNodeSetParams_v2	export
+0x0034f690	cuGraphAddMemcpyNode	export
+0x0034f6b0	cuGraphMemcpyNodeGetParams	export
+0x0034f6d0	cuGraphMemcpyNodeSetParams	export
+0x0034f6f0	cuGraphAddMemsetNode	export
+0x0034f710	cuGraphMemsetNodeGetParams	export
+0x0034f730	cuGraphMemsetNodeSetParams	export
+0x0034f750	cuGraphAddHostNode	export
+0x0034f770	cuGraphHostNodeGetParams	export
+0x0034f790	cuGraphHostNodeSetParams	export
+0x0034f7b0	cuGraphAddChildGraphNode	export
+0x0034f7d0	cuGraphChildGraphNodeGetGraph	export
+0x0034f7f0	cuGraphAddEmptyNode	export
+0x0034f810	cuGraphAddEventRecordNode	export
+0x0034f830	cuGraphEventRecordNodeGetEvent	export
+0x0034f850	cuGraphEventRecordNodeSetEvent	export
+0x0034f870	cuGraphAddEventWaitNode	export
+0x0034f890	cuGraphEventWaitNodeGetEvent	export
+0x0034f8b0	cuGraphEventWaitNodeSetEvent	export
+0x0034f8d0	cuGraphAddExternalSemaphoresSignalNode	export
+0x0034f8f0	cuGraphExternalSemaphoresSignalNodeGetParams	export
+0x0034f910	cuGraphExternalSemaphoresSignalNodeSetParams	export
+0x0034f930	cuGraphAddExternalSemaphoresWaitNode	export
+0x0034f950	cuGraphExternalSemaphoresWaitNodeGetParams	export
+0x0034f970	cuGraphExternalSemaphoresWaitNodeSetParams	export
+0x0034f990	cuGraphAddBatchMemOpNode	export
+0x0034f9b0	cuGraphBatchMemOpNodeGetParams	export
+0x0034f9d0	cuGraphBatchMemOpNodeSetParams	export
+0x0034f9f0	cuGraphExecBatchMemOpNodeSetParams	export
+0x0034fa10	cuGraphAddMemAllocNode	export
+0x0034fa30	cuGraphMemAllocNodeGetParams	export
+0x0034fa50	cuGraphAddMemFreeNode	export
+0x0034fa70	cuGraphMemFreeNodeGetParams	export
+0x0034fa90	cuDeviceGraphMemTrim	export
+0x0034fab0	cuDeviceGetGraphMemAttribute	export
+0x0034fad0	cuDeviceSetGraphMemAttribute	export
+0x0034faf0	cuGraphClone	export
+0x0034fb10	cuGraphNodeFindInClone	export
+0x0034fb30	cuGraphNodeGetType	export
+0x0034fb50	cuGraphGetNodes	export
+0x0034fb70	cuGraphGetRootNodes	export
+0x0034fb90	cuGraphGetEdges	export
+0x0034fbb0	cuGraphGetEdges_v2	export
+0x0034fbd0	cuGraphNodeGetDependencies	export
+0x0034fbf0	cuGraphNodeGetDependencies_v2	export
+0x0034fc10	cuGraphNodeGetDependentNodes	export
+0x0034fc30	cuGraphNodeGetDependentNodes_v2	export
+0x0034fc50	cuGraphAddDependencies	export
+0x0034fc70	cuGraphAddDependencies_v2	export
+0x0034fc90	cuGraphRemoveDependencies	export
+0x0034fcb0	cuGraphRemoveDependencies_v2	export
+0x0034fcd0	cuGraphDestroyNode	export
+0x0034fcf0	cuGraphInstantiateWithFlags	export
+0x0034fd10	cuGraphInstantiateWithParams_ptsz	export
+0x0034fd30	cuGraphExecGetFlags	export
+0x0034fd50	cuGraphExecKernelNodeSetParams_v2	export
+0x0034fd70	cuGraphExecMemcpyNodeSetParams	export
+0x0034fd90	cuGraphExecMemsetNodeSetParams	export
+0x0034fdb0	cuGraphExecHostNodeSetParams	export
+0x0034fdd0	cuGraphExecChildGraphNodeSetParams	export
+0x0034fdf0	cuGraphExecEventRecordNodeSetEvent	export
+0x0034fe10	cuGraphExecEventWaitNodeSetEvent	export
+0x0034fe30	cuGraphExecExternalSemaphoresSignalNodeSetParams	export
+0x0034fe50	cuGraphExecExternalSemaphoresWaitNodeSetParams	export
+0x0034fe70	cuGraphNodeSetEnabled	export
+0x0034fe90	cuGraphNodeGetEnabled	export
+0x0034feb0	cuGraphUpload_ptsz	export
+0x0034fed0	cuGraphLaunch_ptsz	export
+0x0034fef0	cuGraphExecDestroy	export
+0x0034ff10	cuGraphDestroy	export
+0x0034ff30	cuGraphExecUpdate_v2	export
+0x0034ff50	cuGraphKernelNodeCopyAttributes	export
+0x0034ff70	cuGraphKernelNodeGetAttribute	export
+0x0034ff90	cuGraphKernelNodeSetAttribute	export
+0x0034ffb0	cuGraphDebugDotPrint	export
+0x0034ffd0	cuUserObjectCreate	export
+0x0034fff0	cuUserObjectRetain	export
+0x00350010	cuUserObjectRelease	export
+0x00350030	cuGraphRetainUserObject	export
+0x00350050	cuGraphReleaseUserObject	export
+0x00350070	cuGraphAddNode	export
+0x00350090	cuGraphAddNode_v2	export
+0x003500b0	cuGraphNodeSetParams	export
+0x003500d0	cuGraphExecNodeSetParams	export
+0x003500f0	cuGraphConditionalHandleCreate	export
+0x00350110	cuOccupancyMaxActiveBlocksPerMultiprocessor	export
+0x00350130	cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags	export
+0x00350150	cuOccupancyMaxPotentialBlockSize	export
+0x00350170	cuOccupancyMaxPotentialBlockSizeWithFlags	export
+0x00350190	cuOccupancyAvailableDynamicSMemPerBlock	export
+0x003501b0	cuOccupancyMaxPotentialClusterSize	export
+0x003501d0	cuOccupancyMaxActiveClusters	export
+0x003501f0	cuTexRefSetArray	export
+0x00350210	cuTexRefSetMipmappedArray	export
+0x00350230	cuTexRefSetAddress_v2	export
+0x00350250	cuTexRefSetAddress2D_v3	export
+0x00350270	cuTexRefSetFormat	export
+0x00350290	cuTexRefSetAddressMode	export
+0x003502b0	cuTexRefSetFilterMode	export
+0x003502d0	cuTexRefSetMipmapFilterMode	export
+0x003502f0	cuTexRefSetMipmapLevelBias	export
+0x00350310	cuTexRefSetMipmapLevelClamp	export
+0x00350330	cuTexRefSetMaxAnisotropy	export
+0x00350350	cuTexRefSetBorderColor	export
+0x00350370	cuTexRefSetFlags	export
+0x00350390	cuTexRefGetAddress_v2	export
+0x003503b0	cuTexRefGetArray	export
+0x003503d0	cuTexRefGetMipmappedArray	export
+0x003503f0	cuTexRefGetAddressMode	export
+0x00350410	cuTexRefGetFilterMode	export
+0x00350430	cuTexRefGetFormat	export
+0x00350450	cuTexRefGetMipmapFilterMode	export
+0x00350470	cuTexRefGetMipmapLevelBias	export
+0x00350490	cuTexRefGetMipmapLevelClamp	export
+0x003504b0	cuTexRefGetMaxAnisotropy	export
+0x003504d0	cuTexRefGetBorderColor	export
+0x003504f0	cuTexRefGetFlags	export
+0x00350510	cuTexRefCreate	export
+0x00350530	cuTexRefDestroy	export
+0x00350550	cuSurfRefSetArray	export
+0x00350570	cuSurfRefGetArray	export
+0x00350590	cuTexObjectCreate	export
+0x003505b0	cuTexObjectDestroy	export
+0x003505d0	cuTexObjectGetResourceDesc	export
+0x003505f0	cuTexObjectGetTextureDesc	export
+0x00350610	cuTexObjectGetResourceViewDesc	export
+0x00350630	cuSurfObjectCreate	export
+0x00350650	cuSurfObjectDestroy	export
+0x00350670	cuSurfObjectGetResourceDesc	export
+0x00350690	cuTensorMapEncodeTiled	export
+0x003506b0	cuTensorMapEncodeIm2col	export
+0x003506d0	cuTensorMapEncodeIm2colWide	export
+0x003506f0	cuTensorMapReplaceAddress	export
+0x00350710	cuDeviceCanAccessPeer	export
+0x00350730	cuCtxEnablePeerAccess	export
+0x00350750	cuCtxDisablePeerAccess	export
+0x00350770	cuDeviceGetP2PAttribute	export
+0x00350790	cuGraphicsUnregisterResource	export
+0x003507b0	cuGraphicsSubResourceGetMappedArray	export
+0x003507d0	cuGraphicsResourceGetMappedMipmappedArray	export
+0x003507f0	cuGraphicsResourceGetMappedPointer_v2	export
+0x00350810	cuGraphicsResourceSetMapFlags_v2	export
+0x00350830	cuGraphicsMapResources_ptsz	export
+0x00350850	cuGraphicsUnmapResources_ptsz	export
+0x00350870	cuGetProcAddress_v2	export
+0x00350890	cuCoredumpGetAttribute	export
+0x003508b0	cuCoredumpGetAttributeGlobal	export
+0x003508d0	cuCoredumpSetAttribute	export
+0x003508f0	cuCoredumpSetAttributeGlobal	export
+0x00350910	cuGetExportTable	export
+0x00350930	cuGreenCtxCreate	export
+0x00350950	cuGreenCtxDestroy	export
+0x00350970	cuCtxFromGreenCtx	export
+0x00350990	cuDeviceGetDevResource	export
+0x003509b0	cuCtxGetDevResource	export
+0x003509d0	cuGreenCtxGetDevResource	export
+0x003509f0	cuDevSmResourceSplitByCount	export
+0x00350a10	cuDevResourceGenerateDesc	export
+0x00350a30	cuGreenCtxRecordEvent	export
+0x00350a50	cuGreenCtxWaitEvent	export
+0x00350a70	cuStreamGetGreenCtx	export
+0x00350a90	cuGreenCtxStreamCreate	export
+0x00350ab0	cuMemHostRegister	export
+0x00350ad0	cuGraphicsResourceSetMapFlags	export
+0x00350af0	cuLinkCreate	export
+0x00350b10	cuLinkAddData	export
+0x00350b30	cuLinkAddFile	export
+0x00350b50	cuTexRefSetAddress2D_v2	export
+0x00350b70	cuDeviceTotalMem	export
+0x00350b90	cuCtxCreate	export
+0x00350bb0	cuModuleGetGlobal	export
+0x0036e7b0	cuMemGetInfo	export
+0x0036e7d0	cuMemAlloc	export
+0x0036e7f0	cuMemAllocPitch	export
+0x0036e810	cuMemFree	export
+0x0036e830	cuMemGetAddressRange	export
+0x0036e850	cuMemAllocHost	export
+0x0036e870	cuMemHostGetDevicePointer	export
+0x0036e890	cuMemcpyHtoD	export
+0x0036e8b0	cuMemcpyDtoH	export
+0x0036e8d0	cuMemcpyDtoD	export
+0x0036e8f0	cuMemcpyDtoA	export
+0x0036e910	cuMemcpyAtoD	export
+0x0036e930	cuMemcpyHtoA	export
+0x0036e950	cuMemcpyAtoH	export
+0x0036e970	cuMemcpyAtoA	export
+0x0036e990	cuMemcpyHtoAAsync	export
+0x0036e9b0	cuMemcpyAtoHAsync	export
+0x0036e9d0	cuMemcpy2D	export
+0x0036e9f0	cuMemcpy2DUnaligned	export
+0x0036ea10	cuMemcpy3D	export
+0x0036ea30	cuMemcpyHtoDAsync	export
+0x0036ea50	cuMemcpyDtoHAsync	export
+0x0036ea70	cuMemcpyDtoDAsync	export
+0x0036ea90	cuMemcpy2DAsync	export
+0x0036eab0	cuMemcpy3DAsync	export
+0x0036ead0	cuMemsetD8	export
+0x0036eb00	cuMemsetD16	export
+0x0036eb20	cuMemsetD32	export
+0x0036eb40	cuMemsetD2D8	export
+0x0036eb60	cuMemsetD2D16	export
+0x0036eb80	cuMemsetD2D32	export
+0x0036eba0	cuArrayCreate	export
+0x0036ebc0	cuArrayGetDescriptor	export
+0x0036ebe0	cuArray3DCreate	export
+0x0036ec00	cuArray3DGetDescriptor	export
+0x0036ec20	cuTexRefSetAddress	export
+0x0036ec40	cuTexRefSetAddress2D	export
+0x0036ec60	cuTexRefGetAddress	export
+0x0036ec80	cuGraphicsResourceGetMappedPointer	export
+0x0036eca0	cuCtxDestroy	export
+0x0036ecc0	cuCtxPopCurrent	export
+0x0036ece0	cuCtxPushCurrent	export
+0x0036ed00	cuStreamDestroy	export
+0x0036ed20	cuEventDestroy	export
+0x0036ed40	cuDevicePrimaryCtxRelease	export
+0x0036ed60	cuDevicePrimaryCtxReset	export
+0x0036ed80	cuDevicePrimaryCtxSetFlags	export
+0x0036eda0	cuMemcpyHtoD_v2	export
+0x0036edc0	cuMemcpyDtoH_v2	export
+0x0036ede0	cuMemcpyDtoD_v2	export
+0x0036ee00	cuMemcpyDtoA_v2	export
+0x0036ee20	cuMemcpyAtoD_v2	export
+0x0036ee40	cuMemcpyHtoA_v2	export
+0x0036ee60	cuMemcpyAtoH_v2	export
+0x0036ee80	cuMemcpyAtoA_v2	export
+0x0036eea0	cuMemcpyHtoAAsync_v2	export
+0x0036eec0	cuMemcpyAtoHAsync_v2	export
+0x0036eee0	cuMemcpy2D_v2	export
+0x0036ef00	cuMemcpy2DUnaligned_v2	export
+0x0036ef20	cuMemcpy3D_v2	export
+0x0036ef40	cuMemcpyHtoDAsync_v2	export
+0x0036ef60	cuMemcpyDtoHAsync_v2	export
+0x0036ef80	cuMemcpyDtoDAsync_v2	export
+0x0036efa0	cuMemcpy2DAsync_v2	export
+0x0036efc0	cuMemcpy3DAsync_v2	export
+0x0036efe0	cuMemsetD8_v2	export
+0x0036f010	cuMemsetD16_v2	export
+0x0036f030	cuMemsetD32_v2	export
+0x0036f050	cuMemsetD2D8_v2	export
+0x0036f070	cuMemsetD2D16_v2	export
+0x0036f090	cuMemsetD2D32_v2	export
+0x0036f0b0	cuMemcpy	export
+0x0036f0d0	cuMemcpyAsync	export
+0x0036f0f0	cuMemcpyPeer	export
+0x0036f110	cuMemcpyPeerAsync	export
+0x0036f130	cuMemcpy3DPeer	export
+0x0036f150	cuMemcpy3DPeerAsync	export
+0x0036f170	cuMemcpyBatchAsync	export
+0x0036f190	cuMemcpy3DBatchAsync	export
+0x0036f1b0	cuMemsetD8Async	export
+0x0036f1e0	cuMemsetD16Async	export
+0x0036f200	cuMemsetD32Async	export
+0x0036f220	cuMemsetD2D8Async	export
+0x0036f240	cuMemsetD2D16Async	export
+0x0036f260	cuMemsetD2D32Async	export
+0x0036f280	cuStreamGetPriority	export
+0x0036f2a0	cuStreamGetId	export
+0x0036f2c0	cuStreamGetFlags	export
+0x0036f2e0	cuStreamGetDevice	export
+0x0036f300	cuStreamGetCtx	export
+0x0036f320	cuStreamGetCtx_v2	export
+0x0036f340	cuStreamWaitEvent	export
+0x0036f360	cuStreamAddCallback	export
+0x0036f380	cuStreamAttachMemAsync	export
+0x0036f3a0	cuStreamQuery	export
+0x0036f3c0	cuStreamSynchronize	export
+0x0036f3e0	cuEventRecord	export
+0x0036f400	cuEventRecordWithFlags	export
+0x0036f420	cuLaunchKernel	export
+0x0036f440	cuLaunchKernelEx	export
+0x0036f460	cuLaunchHostFunc	export
+0x0036f480	cuGraphicsMapResources	export
+0x0036f4a0	cuGraphicsUnmapResources	export
+0x0036f4c0	cuStreamWriteValue32	export
+0x0036f4e0	cuStreamWaitValue32	export
+0x0036f500	cuStreamWriteValue64	export
+0x0036f520	cuStreamWaitValue64	export
+0x0036f540	cuStreamBatchMemOp	export
+0x0036f560	cuStreamWriteValue32_ptsz	export
+0x0036f580	cuStreamWaitValue32_ptsz	export
+0x0036f5a0	cuStreamWriteValue64_ptsz	export
+0x0036f5c0	cuStreamWaitValue64_ptsz	export
+0x0036f5e0	cuStreamBatchMemOp_ptsz	export
+0x0036f600	cuStreamWriteValue32_v2	export
+0x0036f620	cuStreamWaitValue32_v2	export
+0x0036f640	cuStreamWriteValue64_v2	export
+0x0036f660	cuStreamWaitValue64_v2	export
+0x0036f680	cuStreamBatchMemOp_v2	export
+0x0036f6a0	cuMemPrefetchAsync	export
+0x0036f6c0	cuMemPrefetchAsync_v2	export
+0x0036f6e0	cuLaunchCooperativeKernel	export
+0x0036f700	cuSignalExternalSemaphoresAsync	export
+0x0036f720	cuWaitExternalSemaphoresAsync	export
+0x0036f740	cuStreamBeginCapture	export
+0x0036f760	cuStreamBeginCapture_ptsz	export
+0x0036f780	cuStreamBeginCapture_v2	export
+0x0036f7a0	cuStreamBeginCaptureToGraph	export
+0x0036f7c0	cuStreamEndCapture	export
+0x0036f7e0	cuStreamIsCapturing	export
+0x0036f800	cuStreamGetCaptureInfo	export
+0x0036f820	cuStreamGetCaptureInfo_ptsz	export
+0x0036f840	cuStreamGetCaptureInfo_v2	export
+0x0036f860	cuStreamGetCaptureInfo_v3	export
+0x0036f880	cuGraphAddKernelNode	export
+0x0036f8a0	cuGraphKernelNodeGetParams	export
+0x0036f8c0	cuGraphKernelNodeSetParams	export
+0x0036f8e0	cuGraphExecKernelNodeSetParams	export
+0x0036f900	cuGraphInstantiateWithParams	export
+0x0036f920	cuGraphExecUpdate	export
+0x0036f940	cuGraphUpload	export
+0x0036f960	cuGraphLaunch	export
+0x0036f980	cuStreamCopyAttributes	export
+0x0036f9a0	cuStreamGetAttribute	export
+0x0036f9c0	cuStreamSetAttribute	export
+0x0036f9e0	cuIpcOpenMemHandle	export
+0x0036fa00	cuGraphInstantiate	export
+0x0036fa20	cuGraphInstantiate_v2	export
+0x0036fa40	cuMemMapArrayAsync	export
+0x0036fa60	cuMemFreeAsync	export
+0x0036fa80	cuMemAllocAsync	export
+0x0036faa0	cuMemAllocFromPoolAsync	export
+0x0036fac0	cuStreamUpdateCaptureDependencies	export
+0x0036fae0	cuStreamUpdateCaptureDependencies_v2	export
+0x0036fb00	cuMemBatchDecompressAsync	export
+0x0036fb20	cuGetProcAddress	export
+0x0036fb40	cuCheckpointProcessGetRestoreThreadId	export
+0x0036fb60	cuCheckpointProcessGetState	export
+0x0036fb80	cuCheckpointProcessLock	export
+0x0036fba0	cuCheckpointProcessCheckpoint	export
+0x0036fbc0	cuCheckpointProcessRestore	export
+0x0036fbe0	cuCheckpointProcessUnlock	export
+0x0036fc00	cuGraphicsEGLRegisterImage	export
+0x0036fc20	cuEGLStreamConsumerConnect	export
+0x0036fc40	cuEGLStreamConsumerConnectWithFlags	export
+0x0036fc60	cuEGLStreamConsumerDisconnect	export
+0x0036fc80	cuEGLStreamConsumerAcquireFrame	export
+0x0036fca0	cuEGLStreamConsumerReleaseFrame	export
+0x0036fcc0	cuEGLStreamProducerConnect	export
+0x0036fce0	cuEGLStreamProducerDisconnect	export
+0x0036fd00	cuEGLStreamProducerPresentFrame	export
+0x0036fd20	cuEGLStreamProducerReturnFrame	export
+0x0036fd40	cuGraphicsResourceGetMappedEglFrame	export
+0x0036fd60	cuGraphicsGLRegisterBuffer	export
+0x0036fd80	cuGraphicsGLRegisterImage	export
+0x0036fda0	cuGLGetDevices_v2	export
+0x0036fdc0	cuGLCtxCreate_v2	export
+0x0036fde0	cuGLInit	export
+0x0036fe00	cuGLRegisterBufferObject	export
+0x0036fe20	cuGLMapBufferObject_v2_ptds	export
+0x0036fe40	cuGLUnmapBufferObject	export
+0x0036fe60	cuGLUnregisterBufferObject	export
+0x0036fe80	cuGLSetBufferObjectMapFlags	export
+0x0036fea0	cuGLMapBufferObjectAsync_v2_ptsz	export
+0x0036fec0	cuGLUnmapBufferObjectAsync	export
+0x0036fee0	cuGLGetDevices	export
+0x0036ff00	cuGLMapBufferObject_v2	export
+0x0036ff20	cuGLMapBufferObjectAsync_v2	export
+0x0038d820	cuGLCtxCreate	export
+0x0038d840	cuGLMapBufferObject	export
+0x0038d860	cuGLMapBufferObjectAsync	export
+0x0038d880	cuProfilerInitialize	export
+0x0038d8a0	cuProfilerStart	export
+0x0038d8c0	cuProfilerStop	export
+0x0038d8e0	cuVDPAUGetDevice	export
+0x0038d900	cuVDPAUCtxCreate_v2	export
+0x0038d920	cuGraphicsVDPAURegisterVideoSurface	export
+0x0038d940	cuGraphicsVDPAURegisterOutputSurface	export
+0x0038d960	cuVDPAUCtxCreate	export
+0x00487710	cudbgMain	export
+#
+# --- generated wrappers (initial interposition slot targets) ---
+0x0036e2c0	wrapper_cuModuleGetGlobal	wrapper
+0x003225e0	wrapper_cuCtxGetFlags	wrapper
+0x00371570	wrapper_cuMemGetInfo	wrapper
+0x0034b020	wrapper_cuFuncGetAttribute	wrapper
+0x0037e6a0	wrapper_cuMemcpyAsync	wrapper
+0x0037c8b0	wrapper_cuMemcpyDtoA_v2	wrapper
+0x00375390	wrapper_cuMemsetD32Async	wrapper
+0x00385b60	wrapper_cuGraphExecKernelNodeSetParams	wrapper
+0x003240f0	wrapper_cuCtxSetCacheConfig	wrapper
+0x0034c8d0	wrapper_cuLaunchKernelEx_ptsz	wrapper
+0x00344640	wrapper_cuPointerSetAttribute	wrapper
+0x0034a240	wrapper_cuDestroyExternalSemaphore	wrapper
+0x003270a0	wrapper_cuLibraryLoadFromFile	wrapper
+0x00368b70	wrapper_cuDeviceCanAccessPeer	wrapper
+0x00359120	wrapper_cuGraphAddExternalSemaphoresWaitNode	wrapper
+0x00353810	wrapper_cuOccupancyMaxActiveClusters	wrapper
+0x0035f6a0	wrapper_cuGraphExecUpdate_v2	wrapper
+0x00340590	wrapper_cuThreadExchangeStreamCaptureMode	wrapper
+0x00358620	wrapper_cuGraphAddEventRecordNode	wrapper
+0x0035c060	wrapper_cuGraphGetEdges	wrapper
+0x0034da70	wrapper_cuFuncSetBlockShape	wrapper
+0x0034ceb0	wrapper_cuLaunchCooperativeKernelMultiDevice	wrapper
+0x00366510	wrapper_cuTensorMapEncodeTiled	wrapper
+0x0031f700	wrapper_cuCtxGetExecAffinity	wrapper
+0x00369cc0	wrapper_cuGraphicsResourceGetMappedPointer_v2	wrapper
+0x0035e3c0	wrapper_cuGraphExecChildGraphNodeSetParams	wrapper
+0x003939c0	wrapper_cuGLMapBufferObjectAsync	wrapper
+0x003324b0	wrapper_cuFuncGetParamInfo	wrapper
+0x003369f0	wrapper_cuMemcpyAtoA_v2_ptds	wrapper
+0x00342000	wrapper_cuMemPoolExportToShareableHandle	wrapper
+0x003682a0	wrapper_cuSurfObjectGetResourceDesc	wrapper
+0x0031d160	wrapper_cuModuleGetLoadingMode	wrapper
+0x003678e0	wrapper_cuTexObjectGetResourceDesc	wrapper
+0x0033f640	wrapper_cuMemUnmap	wrapper
+0x003275f0	wrapper_cuLibraryGetKernel	wrapper
+0x0038cab0	wrapper_cuGLUnmapBufferObject	wrapper
+0x00389830	wrapper_cuEGLStreamConsumerConnectWithFlags	wrapper
+0x0036c810	wrapper_cuGreenCtxRecordEvent	wrapper
+0x0031acd0	wrapper_cuModuleEnumerateFunctions	wrapper
+0x00349e40	wrapper_cuSignalExternalSemaphoresAsync_ptsz	wrapper
+0x0031d570	wrapper_cuMemGetInfo_v2	wrapper
+0x00364330	wrapper_cuTexRefSetBorderColor	wrapper
+0x00394b80	wrapper_cuGraphicsVDPAURegisterOutputSurface	wrapper
+0x0033b7a0	wrapper_cuMipmappedArrayGetLevel	wrapper
+0x00341620	wrapper_cuMemPoolGetAttribute	wrapper
+0x0031b1f0	wrapper_cuKernelGetName	wrapper
+0x0037a2f0	wrapper_cuStreamGetCtx	wrapper
+0x00323e70	wrapper_cuCtxGetSharedMemConfig	wrapper
+0x0038b170	wrapper_cuGraphicsGLRegisterBuffer	wrapper
+0x00348360	wrapper_cuEventQuery	wrapper
+0x003598e0	wrapper_cuGraphBatchMemOpNodeSetParams	wrapper
+0x00371ae0	wrapper_cuMemAlloc	wrapper
+0x00383f60	wrapper_cuSignalExternalSemaphoresAsync	wrapper
+0x0035f460	wrapper_cuGraphDestroy	wrapper
+0x00381e60	wrapper_cuStreamBatchMemOp_v2	wrapper
+0x00368510	wrapper_cuTexObjectDestroy	wrapper
+0x00343cd0	wrapper_cuMemPrefetchAsync_ptsz	wrapper
+0x00337e60	wrapper_cuMemcpyAtoHAsync_v2_ptsz	wrapper
+0x00374bd0	wrapper_cuMemcpy2DAsync	wrapper
+0x00331c90	wrapper_cuMemRangeGetAttributes	wrapper
+0x00363300	wrapper_cuTexRefSetAddressMode	wrapper
+0x00322ef0	wrapper_cuModuleGetSurfRef	wrapper
+0x0035aaf0	wrapper_cuGraphNodeGetEnabled	wrapper
+0x0035e680	wrapper_cuGraphNodeSetEnabled	wrapper
+0x00349270	wrapper_cuImportExternalMemory	wrapper
+0x0035cf40	wrapper_cuGraphNodeGetDependentNodes_v2	wrapper
+0x00369700	wrapper_cuGraphicsSubResourceGetMappedArray	wrapper
+0x0033a9c0	wrapper_cuArray3DCreate_v2	wrapper
+0x00384160	wrapper_cuWaitExternalSemaphoresAsync	wrapper
+0x0035b360	wrapper_cuGraphAddMemFreeNode	wrapper
+0x00323230	wrapper_cuCtxSynchronize	wrapper
+0x00336600	wrapper_cuMemcpyHtoA_v2_ptds	wrapper
+0x0037b860	wrapper_cuEventDestroy	wrapper
+0x0036a060	wrapper_cuGraphicsResourceSetMapFlags_v2	wrapper
+0x00358f30	wrapper_cuGraphExternalSemaphoresSignalNodeSetParams	wrapper
+0x0033b0a0	wrapper_cuMipmappedArrayGetMemoryRequirements	wrapper
+0x003265a0	wrapper_cuLinkAddFile_v2	wrapper
+0x00363db0	wrapper_cuTexRefSetMipmapLevelClamp	wrapper
+0x0031a850	wrapper_cuDriverGetVersion	wrapper
+0x0037af90	wrapper_cuStreamCopyAttributes	wrapper
+0x00388cc0	wrapper_cuCheckpointProcessLock	wrapper
+0x003322d0	wrapper_cuFuncGetName	wrapper
+0x0037b5b0	wrapper_cuStreamDestroy	wrapper
+0x0034a4b0	wrapper_cuStreamBatchMemOp_v2_ptsz	wrapper
+0x00357cd0	wrapper_cuGraphAddHostNode	wrapper
+0x003438e0	wrapper_cuMulticastBindAddr	wrapper
+0x0035b050	wrapper_cuGraphAddMemAllocNode	wrapper
+0x003693f0	wrapper_cuDeviceGetP2PAttribute	wrapper
+0x0038c830	wrapper_cuGLMapBufferObjectAsync_v2	wrapper
+0x0035eb30	wrapper_cuGraphUpload_ptsz	wrapper
+0x00379c30	wrapper_cuStreamGetId	wrapper
+0x00324540	wrapper_cuCtxResetPersistingL2Cache	wrapper
+0x0031bc10	wrapper_cuDeviceGetCount	wrapper
+0x0035df50	wrapper_cuGraphExecKernelNodeSetParams_v2	wrapper
+0x00388f50	wrapper_cuCheckpointProcessCheckpoint	wrapper
+0x003465b0	wrapper_cuStreamUpdateCaptureDependencies_ptsz	wrapper
+0x0036d650	wrapper_cuLinkAddData	wrapper
+0x003398b0	wrapper_cuMemsetD2D32Async_ptsz	wrapper
+0x00341ab0	wrapper_cuMemPoolCreate	wrapper
+0x00335a80	wrapper_cuMemcpyPeer_ptds	wrapper
+0x003206d0	wrapper_cuCtxCreate_v2	wrapper
+0x00346860	wrapper_cuStreamEndCapture_ptsz	wrapper
+0x00320a70	wrapper_cuCtxCreate_v3	wrapper
+0x00320f40	wrapper_cuCtxCreate_v4	wrapper
+0x003608c0	wrapper_cuGraphReleaseUserObject	wrapper
+0x0035f1d0	wrapper_cuGraphExecDestroy	wrapper
+0x00393230	wrapper_cuVDPAUCtxCreate	wrapper
+0x00379670	wrapper_cuGraphicsUnmapResources	wrapper
+0x00348ec0	wrapper_cuEventElapsedTime_v2	wrapper
+0x0035dbf0	wrapper_cuGraphExecGetFlags	wrapper
+0x0031f9a0	wrapper_cuCtxAttach	wrapper
+0x0031ff10	wrapper_cuDevicePrimaryCtxRelease_v2	wrapper
+0x00364b00	wrapper_cuTexRefGetMipmappedArray	wrapper
+0x0036a2f0	wrapper_cuGraphicsMapResources_ptsz	wrapper
+0x003816d0	wrapper_cuLaunchCooperativeKernel	wrapper
+0x00365ab0	wrapper_cuTexRefGetMaxAnisotropy	wrapper
+0x00337270	wrapper_cuMemcpyAsync_ptsz	wrapper
+0x0033e540	wrapper_cuMemAddressReserve	wrapper
+0x0035d7b0	wrapper_cuGraphInstantiateWithFlags	wrapper
+0x0032a230	wrapper_cuMemAllocHost_v2	wrapper
+0x00369090	wrapper_cuCtxDisablePeerAccess	wrapper
+0x00368db0	wrapper_cuCtxEnablePeerAccess	wrapper
+0x003603b0	wrapper_cuUserObjectRetain	wrapper
+0x003544f0	wrapper_cuDevSmResourceSplitByCount	wrapper
+0x0033e070	wrapper_cuMemBatchDecompressAsync_ptsz	wrapper
+0x0031aa10	wrapper_cuModuleGetFunctionCount	wrapper
+0x0035cc40	wrapper_cuGraphNodeGetDependentNodes	wrapper
+0x00342fb0	wrapper_cuMulticastUnbind	wrapper
+0x00353bd0	wrapper_cuTexRefSetMipmappedArray	wrapper
+0x0034dce0	wrapper_cuFuncSetSharedSize	wrapper
+0x0037e890	wrapper_cuMemcpyPeer	wrapper
+0x0035fef0	wrapper_cuGraphDebugDotPrint	wrapper
+0x0036b2d0	wrapper_cuCoredumpGetAttributeGlobal	wrapper
+0x00325460	wrapper_cuModuleUnload	wrapper
+0x00342bc0	wrapper_cuMemPoolImportPointer	wrapper
+0x0037dd60	wrapper_cuMemcpyDtoHAsync_v2	wrapper
+0x003449a0	wrapper_cuPointerGetAttributes	wrapper
+0x003358a0	wrapper_cuMemcpy_ptds	wrapper
+0x0031ef50	wrapper_cuFlushGPUDirectRDMAWrites	wrapper
+0x0037e300	wrapper_cuMemcpy3DAsync_v2	wrapper
+0x0037e140	wrapper_cuMemcpy2DAsync_v2	wrapper
+0x00326250	wrapper_cuLinkAddData_v2	wrapper
+0x00388600	wrapper_cuGetProcAddress	wrapper
+0x00353db0	wrapper_cuSurfRefSetArray	wrapper
+0x0032a5d0	wrapper_cuMemFreeHost	wrapper
+0x00365030	wrapper_cuTexRefGetFormat	wrapper
+0x0036d190	wrapper_cuMemHostRegister	wrapper
+0x003822c0	wrapper_cuStreamWaitValue32	wrapper
+0x00329830	wrapper_cuMemAlloc_v2	wrapper
+0x00343520	wrapper_cuMulticastCreate	wrapper
+0x0037ce80	wrapper_cuMemcpyAtoH_v2	wrapper
+0x003514e0	wrapper_cuGraphEventWaitNodeGetEvent	wrapper
+0x00338070	wrapper_cuMemcpy2DAsync_v2_ptsz	wrapper
+0x00360b00	wrapper_cuGraphNodeSetParams	wrapper
+0x00362b70	wrapper_cuTexRefSetAddress2D_v2	wrapper
+0x003626c0	wrapper_cuTexRefSetAddress2D_v3	wrapper
+0x00355070	wrapper_cuGraphMemcpyNodeGetParams	wrapper
+0x003295d0	wrapper_cuKernelSetCacheConfig	wrapper
+0x00367c60	wrapper_cuTexObjectGetTextureDesc	wrapper
+0x0035b8e0	wrapper_cuDeviceGetGraphMemAttribute	wrapper
+0x00346f90	wrapper_cuStreamAttachMemAsync_ptsz	wrapper
+0x0035e8d0	wrapper_cuStreamGetGreenCtx	wrapper
+0x00337c60	wrapper_cuMemcpyHtoAAsync_v2_ptsz	wrapper
+0x003724c0	wrapper_cuMemAllocHost	wrapper
+0x00325090	wrapper_cuModuleLoadDataEx	wrapper
+0x003407f0	wrapper_cuMemAllocAsync_ptsz	wrapper
+0x00336d90	wrapper_cuMemcpy2DUnaligned_v2_ptds	wrapper
+0x00337a70	wrapper_cuMemcpyDtoDAsync_v2_ptsz	wrapper
+0x003252a0	wrapper_cuModuleLoadFatBinary	wrapper
+0x00381a40	wrapper_cuStreamBatchMemOp	wrapper
+0x0033ae60	wrapper_cuArrayGetMemoryRequirements	wrapper
+0x003367f0	wrapper_cuMemcpyAtoH_v2_ptds	wrapper
+0x00333280	wrapper_cuDeviceUnregisterAsyncNotification	wrapper
+0x00370310	wrapper_cuArrayGetDescriptor	wrapper
+0x00358960	wrapper_cuGraphAddEventWaitNode	wrapper
+0x0036bb70	wrapper_cuGreenCtxCreate	wrapper
+0x00341070	wrapper_cuMemAllocFromPoolAsync_ptsz	wrapper
+0x00343af0	wrapper_cuPointerGetAttribute	wrapper
+0x00382050	wrapper_cuStreamWriteValue32	wrapper
+0x00384e10	wrapper_cuStreamBeginCaptureToGraph	wrapper
+0x0034e480	wrapper_cuParamSetTexRef	wrapper
+0x0038b4e0	wrapper_cuGraphicsGLRegisterImage	wrapper
+0x0033c720	wrapper_cuMemGetHandleForAddressRange	wrapper
+0x003858a0	wrapper_cuGraphKernelNodeSetParams	wrapper
+0x00373480	wrapper_cuMemcpyHtoA	wrapper
+0x0034adc0	wrapper_cuStreamWriteValue64_v2_ptsz	wrapper
+0x0038ab50	wrapper_cuEGLStreamProducerReturnFrame	wrapper
+0x00372b30	wrapper_cuMemcpyHtoD	wrapper
+0x00373db0	wrapper_cuMemcpyAtoHAsync	wrapper
+0x0036acb0	wrapper_cuGetProcAddress_v2	wrapper
+0x00382a10	wrapper_cuStreamWriteValue32_ptsz	wrapper
+0x00334c40	wrapper_cuIpcGetMemHandle	wrapper
+0x0035d2b0	wrapper_cuGetExportTable	wrapper
+0x00389100	wrapper_cuCheckpointProcessRestore	wrapper
+0x00363b20	wrapper_cuTexRefSetMipmapLevelBias	wrapper
+0x00353f90	wrapper_cuTexObjectCreate	wrapper
+0x00371f00	wrapper_cuMemFree	wrapper
+0x00394800	wrapper_cuGraphicsVDPAURegisterVideoSurface	wrapper
+0x00338b90	wrapper_cuMemcpy3DBatchAsync_ptsz	wrapper
+0x00352bf0	wrapper_cuGraphExecMemcpyNodeSetParams	wrapper
+0x00342800	wrapper_cuMemPoolExportPointer	wrapper
+0x00383160	wrapper_cuStreamWaitValue64_ptsz	wrapper
+0x00335c80	wrapper_cuMemcpyHtoD_v2_ptds	wrapper
+0x0036d9a0	wrapper_cuLinkAddFile	wrapper
+0x00386580	wrapper_cuGraphExecUpdate	wrapper
+0x00379e80	wrapper_cuStreamGetFlags	wrapper
+0x0035b6a0	wrapper_cuDeviceGraphMemTrim	wrapper
+0x003217d0	wrapper_cuCtxDetach	wrapper
+0x0035fc30	wrapper_cuGraphKernelNodeSetAttribute	wrapper
+0x003699c0	wrapper_cuGraphicsResourceGetMappedMipmappedArray	wrapper
+0x0034cb40	wrapper_cuLaunchCooperativeKernel_ptsz	wrapper
+0x0037caa0	wrapper_cuMemcpyAtoD_v2	wrapper
+0x0038c350	wrapper_cuGLMapBufferObjectAsync_v2_ptsz	wrapper
+0x00377e40	wrapper_cuTexRefSetAddress2D	wrapper
+0x0038ce20	wrapper_cuGLUnmapBufferObjectAsync	wrapper
+0x0033b550	wrapper_cuArray3DGetDescriptor_v2	wrapper
+0x00324310	wrapper_cuCtxSetSharedMemConfig	wrapper
+0x0037a560	wrapper_cuStreamGetCtx_v2	wrapper
+0x00384b60	wrapper_cuStreamUpdateCaptureDependencies	wrapper
+0x0036a9f0	wrapper_cuGraphicsResourceSetMapFlags	wrapper
+0x0031ccb0	wrapper_cuCtxGetId	wrapper
+0x00355430	wrapper_cuGraphChildGraphNodeGetGraph	wrapper
+0x00355b30	wrapper_cuCtxFromGreenCtx	wrapper
+0x0036be60	wrapper_cuGreenCtxDestroy	wrapper
+0x00374d90	wrapper_cuMemcpy3DAsync	wrapper
+0x00360f80	wrapper_cuGraphExecNodeSetParams	wrapper
+0x0031bdf0	wrapper_cuDeviceGetName	wrapper
+0x00326d30	wrapper_cuLibraryLoadData	wrapper
+0x00335520	wrapper_cuMemHostRegister_v2	wrapper
+0x00327c10	wrapper_cuLibraryGetKernelCount	wrapper
+0x00334180	wrapper_cuDeviceGetPCIBusId	wrapper
+0x00323bf0	wrapper_cuCtxGetCacheConfig	wrapper
+0x00355da0	wrapper_cuGraphAddMemcpyNode	wrapper
+0x003268c0	wrapper_cuLinkComplete	wrapper
+0x00336220	wrapper_cuMemcpyDtoA_v2_ptds	wrapper
+0x00345130	wrapper_cuStreamWaitEvent_ptsz	wrapper
+0x00368810	wrapper_cuSurfObjectDestroy	wrapper
+0x0033e900	wrapper_cuMemAddressFree	wrapper
+0x00354d30	wrapper_cuGraphKernelNodeCopyAttributes	wrapper
+0x00356980	wrapper_cuGraphAddMemsetNode	wrapper
+0x003730a0	wrapper_cuMemcpyDtoA	wrapper
+0x00334720	wrapper_cuIpcOpenEventHandle	wrapper
+0x00372ed0	wrapper_cuMemcpyDtoD	wrapper
+0x0031dae0	wrapper_cuDeviceGetTexture1DLinearMaxWidth	wrapper
+0x0033f800	wrapper_cuMemSetAccess	wrapper
+0x0038a2e0	wrapper_cuEGLStreamConsumerReleaseFrame	wrapper
+0x00370990	wrapper_cuStreamGetCaptureInfo_v2	wrapper
+0x00383890	wrapper_cuStreamWriteValue64_v2	wrapper
+0x00349830	wrapper_cuExternalMemoryGetMappedMipmappedArray	wrapper
+0x00374120	wrapper_cuMemcpy2D	wrapper
+0x00370bd0	wrapper_cuStreamGetCaptureInfo_v3	wrapper
+0x00388130	wrapper_cuMemBatchDecompressAsync	wrapper
+0x0033d2d0	wrapper_cuStreamGetCtx_ptsz	wrapper
+0x00372d00	wrapper_cuMemcpyDtoH	wrapper
+0x003460d0	wrapper_cuStreamBeginCaptureToGraph_ptsz	wrapper
+0x003618f0	wrapper_cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags	wrapper
+0x00378550	wrapper_cuGraphicsResourceGetMappedPointer	wrapper
+0x0034ab60	wrapper_cuStreamWriteValue32_v2_ptsz	wrapper
+0x00363890	wrapper_cuTexRefSetMipmapFilterMode	wrapper
+0x0034e0b0	wrapper_cuLaunchGrid	wrapper
+0x00358ca0	wrapper_cuGraphAddExternalSemaphoresSignalNode	wrapper
+0x0031df80	wrapper_cuDeviceGetProperties	wrapper
+0x0037d080	wrapper_cuMemcpyAtoA_v2	wrapper
+0x003582b0	wrapper_cuGraphAddEmptyNode	wrapper
+0x0038bf50	wrapper_cuGLRegisterBufferObject	wrapper
+0x00357ad0	wrapper_cuGraphMemsetNodeSetParams	wrapper
+0x00384360	wrapper_cuStreamBeginCapture	wrapper
+0x00357800	wrapper_cuDeviceGetDevResource	wrapper
+0x003760b0	wrapper_cuMemsetD32	wrapper
+0x00332f80	wrapper_cuDeviceRegisterAsyncNotification	wrapper
+0x00392b60	wrapper_cuGLCtxCreate	wrapper
+0x00380100	wrapper_cuLaunchHostFunc	wrapper
+0x00365d30	wrapper_cuTexRefGetBorderColor	wrapper
+0x0034b610	wrapper_cuFuncIsLoaded	wrapper
+0x0037aa70	wrapper_cuStreamGetCaptureInfo	wrapper
+0x003827a0	wrapper_cuStreamWaitValue64	wrapper
+0x0033fde0	wrapper_cuMemImportFromShareableHandle	wrapper
+0x00367160	wrapper_cuTensorMapReplaceAddress	wrapper
+0x0031c790	wrapper_cuDeviceGetExecAffinitySupport	wrapper
+0x00380dd0	wrapper_cuEventRecord	wrapper
+0x0033ee20	wrapper_cuMemRelease	wrapper
+0x0037df50	wrapper_cuMemcpyDtoDAsync_v2	wrapper
+0x00366240	wrapper_cuTexRefGetAddressMode	wrapper
+0x00351270	wrapper_cuGraphEventRecordNodeSetEvent	wrapper
+0x0038d340	wrapper_cuGLSetBufferObjectMapFlags	wrapper
+0x0036b6a0	wrapper_cuCoredumpSetAttribute	wrapper
+0x00321500	wrapper_cuCtxDestroy_v2	wrapper
+0x00380840	wrapper_cuStreamQuery	wrapper
+0x0034b890	wrapper_cuFuncLoad	wrapper
+0x00380510	wrapper_cuStreamAttachMemAsync	wrapper
+0x00373870	wrapper_cuMemcpyAtoA	wrapper
+0x00335700	wrapper_cuMemHostUnregister	wrapper
+0x0035a690	wrapper_cuGraphExecExternalSemaphoresSignalNodeSetParams	wrapper
+0x00389da0	wrapper_cuEGLStreamProducerDisconnect	wrapper
+0x003529e0	wrapper_cuGraphRemoveDependencies_v2	wrapper
+0x003527f0	wrapper_cuGraphRemoveDependencies	wrapper
+0x00323790	wrapper_cuCtxSetLimit	wrapper
+0x0036cad0	wrapper_cuGreenCtxWaitEvent	wrapper
+0x00374460	wrapper_cuMemcpy3D	wrapper
+0x00373290	wrapper_cuMemcpyAtoD	wrapper
+0x0035da00	wrapper_cuGraphInstantiateWithParams_ptsz	wrapper
+0x00377370	wrapper_cuMemsetD2D32_v2	wrapper
+0x00373670	wrapper_cuMemcpyAtoH	wrapper
+0x00333720	wrapper_cuMemExportToShareableHandle	wrapper
+0x003284c0	wrapper_cuKernelGetFunction	wrapper
+0x0037ad00	wrapper_cuStreamGetCaptureInfo_ptsz	wrapper
+0x0038a630	wrapper_cuEGLStreamProducerConnect	wrapper
+0x0038a0f0	wrapper_cuEGLStreamConsumerAcquireFrame	wrapper
+0x00382530	wrapper_cuStreamWriteValue64	wrapper
+0x0031fc50	wrapper_cuDevicePrimaryCtxRetain	wrapper
+0x00367620	wrapper_cuTexRefDestroy	wrapper
+0x00332da0	wrapper_cuMemAllocManaged	wrapper
+0x00321df0	wrapper_cuCtxPopCurrent_v2	wrapper
+0x003418d0	wrapper_cuMemPoolSetAccess	wrapper
+0x0037f830	wrapper_cuStreamWaitEvent	wrapper
+0x0035ee80	wrapper_cuGraphLaunch_ptsz	wrapper
+0x003539f0	wrapper_cuTexRefSetArray	wrapper
+0x003929a0	wrapper_cuProfilerInitialize	wrapper
+0x0034e290	wrapper_cuLaunchGridAsync	wrapper
+0x0037bad0	wrapper_cuDevicePrimaryCtxRelease	wrapper
+0x00337880	wrapper_cuMemcpyDtoHAsync_v2_ptsz	wrapper
+0x0037f610	wrapper_cuMemcpy3DBatchAsync	wrapper
+0x00383af0	wrapper_cuStreamWaitValue64_v2	wrapper
+0x0034bde0	wrapper_cuFuncSetCacheConfig	wrapper
+0x00385e50	wrapper_cuGraphInstantiateWithParams	wrapper
+0x00387660	wrapper_cuMemMapArrayAsync	wrapper
+0x0034b3b0	wrapper_cuFuncGetModule	wrapper
+0x003792f0	wrapper_cuGraphicsMapResources	wrapper
+0x0038c5d0	wrapper_cuGLMapBufferObject_v2	wrapper
+0x003668e0	wrapper_cuTensorMapEncodeIm2col	wrapper
+0x00349430	wrapper_cuExternalMemoryGetMappedBuffer	wrapper
+0x00332be0	wrapper_cuMemHostGetFlags	wrapper
+0x00336410	wrapper_cuMemcpyAtoD_v2_ptds	wrapper
+0x0038b960	wrapper_cuGLGetDevices	wrapper
+0x00328d40	wrapper_cuLibraryGetUnifiedFunction	wrapper
+0x00386f40	wrapper_cuStreamSetAttribute	wrapper
+0x0038c0f0	wrapper_cuGLMapBufferObject_v2_ptds	wrapper
+0x0037fd10	wrapper_cuStreamAddCallback	wrapper
+0x003564c0	wrapper_cuSurfRefGetArray	wrapper
+0x0034e970	wrapper_cuGraphAddKernelNode_v2	wrapper
+0x003775a0	wrapper_cuArrayCreate	wrapper
+0x0033cc10	wrapper_cuStreamGetDevice_ptsz	wrapper
+0x003747f0	wrapper_cuMemcpyDtoHAsync	wrapper
+0x0037a0d0	wrapper_cuStreamGetDevice	wrapper
+0x003848a0	wrapper_cuStreamBeginCapture_v2	wrapper
+0x00351d60	wrapper_cuGraphBatchMemOpNodeGetParams	wrapper
+0x003343c0	wrapper_cuIpcGetEventHandle	wrapper
+0x0036c5e0	wrapper_cuDevResourceGenerateDesc	wrapper
+0x00388a20	wrapper_cuCheckpointProcessGetState	wrapper
+0x00393760	wrapper_cuGLMapBufferObject	wrapper
+0x003383f0	wrapper_cuMemcpy3DPeerAsync_ptsz	wrapper
+0x00350dd0	wrapper_cuGraphHostNodeSetParams	wrapper
+0x00344f00	wrapper_cuStreamCreateWithPriority	wrapper
+0x00344d00	wrapper_cuStreamCreate	wrapper
+0x0033bf70	wrapper_cuArrayDestroy	wrapper
+0x0035ad40	wrapper_cuGraphExecBatchMemOpNodeSetParams	wrapper
+0x003799f0	wrapper_cuStreamGetPriority	wrapper
+0x003385b0	wrapper_cuMemcpyBatchAsync_ptsz	wrapper
+0x003432c0	wrapper_cuMulticastGetGranularity	wrapper
+0x0035c8c0	wrapper_cuGraphNodeGetDependencies_v2	wrapper
+0x00339ef0	wrapper_cuMemsetD32_v2_ptds	wrapper
+0x00320410	wrapper_cuDevicePrimaryCtxSetFlags_v2	wrapper
+0x003292b0	wrapper_cuKernelSetAttribute	wrapper
+0x0037d690	wrapper_cuMemcpy2D_v2	wrapper
+0x00361ba0	wrapper_cuOccupancyMaxPotentialBlockSize	wrapper
+0x00372680	wrapper_cuMemHostGetDevicePointer	wrapper
+0x00364db0	wrapper_cuTexRefGetFilterMode	wrapper
+0x00362fb0	wrapper_cuTexRefSetFormat	wrapper
+0x0033ac20	wrapper_cuArrayGetDescriptor_v2	wrapper
+0x0031b430	wrapper_cuMemAllocPitch_v2	wrapper
+0x00323520	wrapper_cuCtxGetStreamPriorityRange	wrapper
+0x0035bdd0	wrapper_cuGraphClone	wrapper
+0x0031ecb0	wrapper_cuDeviceGetDefaultMemPool	wrapper
+0x00359bd0	wrapper_cuUserObjectCreate	wrapper
+0x00356750	wrapper_cuGraphMemcpyNodeSetParams	wrapper
+0x0038d500	wrapper_cuMemPrefetchAsync_v2	wrapper
+0x0034a900	wrapper_cuStreamWaitValue64_v2_ptsz	wrapper
+0x00342da0	wrapper_cuMulticastAddDevice	wrapper
+0x0033c9d0	wrapper_cuStreamGetPriority_ptsz	wrapper
+0x0035f970	wrapper_cuGraphKernelNodeGetAttribute	wrapper
+0x0031b9b0	wrapper_cuDeviceGet	wrapper
+0x00335e60	wrapper_cuMemcpyDtoH_v2_ptds	wrapper
+0x00360650	wrapper_cuGraphRetainUserObject	wrapper
+0x0031c2c0	wrapper_cuDeviceGetUuid_v2	wrapper
+0x0037ee60	wrapper_cuMemcpy3DPeerAsync	wrapper
+0x00336f30	wrapper_cuMemcpy3D_v2_ptds	wrapper
+0x00336bf0	wrapper_cuMemcpy2D_v2_ptds	wrapper
+0x00321a60	wrapper_cuCtxPushCurrent_v2	wrapper
+0x0038b6d0	wrapper_cuGLGetDevices_v2	wrapper
+0x00337690	wrapper_cuMemcpyHtoDAsync_v2_ptsz	wrapper
+0x0033f390	wrapper_cuMemMapArrayAsync_ptsz	wrapper
+0x00347dc0	wrapper_cuEventCreate	wrapper
+0x00351b50	wrapper_cuGraphExternalSemaphoresWaitNodeGetParams	wrapper
+0x00322bb0	wrapper_cuModuleGetTexRef	wrapper
+0x00374600	wrapper_cuMemcpyHtoDAsync	wrapper
+0x00370550	wrapper_cuArray3DGetDescriptor	wrapper
+0x00333560	wrapper_cuDeviceGetByPCIBusId	wrapper
+0x0038af50	wrapper_cuGLInit	wrapper
+0x00361f20	wrapper_cuOccupancyMaxPotentialBlockSizeWithFlags	wrapper
+0x0034c1a0	wrapper_cuFuncSetSharedMemConfig	wrapper
+0x00359f40	wrapper_cuGraphExecHostNodeSetParams	wrapper
+0x0034efc0	wrapper_cuMemAdvise_v2	wrapper
+0x0031dd40	wrapper_cuDeviceGetAttribute	wrapper
+0x00329e00	wrapper_cuMemGetAddressRange_v2	wrapper
+0x0032a3f0	wrapper_cuMemHostAlloc	wrapper
+0x00383d50	wrapper_cuMemPrefetchAsync	wrapper
+0x0033d7c0	wrapper_cuStreamIsCapturing_ptsz	wrapper
+0x0031a630	wrapper_cuGetErrorName	wrapper
+0x0036c340	wrapper_cuGreenCtxGetDevResource	wrapper
+0x00327f30	wrapper_cuLibraryEnumerateKernels	wrapper
+0x00347290	wrapper_cuStreamQuery_ptsz	wrapper
+0x003561d0	wrapper_cuTexRefGetAddress_v2	wrapper
+0x00387e30	wrapper_cuMemAllocFromPoolAsync	wrapper
+0x00346d30	wrapper_cuStreamGetCaptureInfo_v3_ptsz	wrapper
+0x00326b60	wrapper_cuLinkDestroy	wrapper
+0x0034d540	wrapper_cuParamSetf	wrapper
+0x00324d10	wrapper_cuModuleLoad	wrapper
+0x003707a0	wrapper_cuEventRecordWithFlags	wrapper
+0x0034d2d0	wrapper_cuParamSeti	wrapper
+0x003852f0	wrapper_cuStreamEndCapture	wrapper
+0x003645e0	wrapper_cuTexRefSetFlags	wrapper
+0x0036ff40	wrapper_cuMemAllocPitch	wrapper
+0x00357470	wrapper_cuGraphGetRootNodes	wrapper
+0x0033c270	wrapper_cuMipmappedArrayCreate	wrapper
+0x00357100	wrapper_cuGraphGetNodes	wrapper
+0x00347590	wrapper_cuStreamSynchronize_ptsz	wrapper
+0x0033ec30	wrapper_cuMemCreate	wrapper
+0x00328250	wrapper_cuLibraryGetModule	wrapper
+0x00353630	wrapper_cuOccupancyMaxPotentialClusterSize	wrapper
+0x00353200	wrapper_cuGraphAddNode_v2	wrapper
+0x00354190	wrapper_cuSurfObjectCreate	wrapper
+0x0034d7b0	wrapper_cuParamSetv	wrapper
+0x00349c80	wrapper_cuImportExternalSemaphore	wrapper
+0x0033a780	wrapper_cuArrayCreate_v2	wrapper
+0x00345e10	wrapper_cuStreamBeginCapture_v2_ptsz	wrapper
+0x0031e730	wrapper_cuDeviceSetMemPool	wrapper
+0x0033bcb0	wrapper_cuMipmappedArrayGetSparseProperties	wrapper
+0x00325a70	wrapper_cuModuleGetGlobal_v2	wrapper
+0x0034a6a0	wrapper_cuStreamWaitValue32_v2_ptsz	wrapper
+0x00355920	wrapper_cuGraphNodeGetType	wrapper
+0x0031e280	wrapper_cuDeviceComputeCapability	wrapper
+0x003320e0	wrapper_cuEventRecordWithFlags_ptsz	wrapper
+0x00337460	wrapper_cuMemcpyPeerAsync_ptsz	wrapper
+0x003657d0	wrapper_cuTexRefGetMipmapLevelClamp	wrapper
+0x0034eca0	wrapper_cuMemPrefetchAsync_v2_ptsz	wrapper
+0x00360100	wrapper_cuUserObjectRelease	wrapper
+0x0031cf20	wrapper_cuCtxGetApiVersion	wrapper
+0x00353430	wrapper_cuOccupancyAvailableDynamicSMemPerBlock	wrapper
+0x00347820	wrapper_cuStreamDestroy_v2	wrapper
+0x003595a0	wrapper_cuGraphAddBatchMemOpNode	wrapper
+0x003635e0	wrapper_cuTexRefSetFilterMode	wrapper
+0x0033a550	wrapper_cuMemsetD2D32_v2_ptds	wrapper
+0x0038d1a0	wrapper_cuGLUnregisterBufferObject	wrapper
+0x0036b8c0	wrapper_cuCoredumpSetAttributeGlobal	wrapper
+0x003833d0	wrapper_cuStreamWriteValue32_v2	wrapper
+0x00351080	wrapper_cuGraphEventRecordNodeGetEvent	wrapper
+0x00354a50	wrapper_cuGraphKernelNodeSetParams_v2	wrapper
+0x00376d10	wrapper_cuMemsetD32_v2	wrapper
+0x003391d0	wrapper_cuMemsetD32Async_ptsz	wrapper
+0x0036a670	wrapper_cuGraphicsUnmapResources_ptsz	wrapper
+0x00371060	wrapper_cuStreamUpdateCaptureDependencies_v2	wrapper
+0x0033ffc0	wrapper_cuMemRetainAllocationHandle	wrapper
+0x00354350	wrapper_cuGraphicsUnregisterResource	wrapper
+0x0035a450	wrapper_cuGraphExecEventWaitNodeSetEvent	wrapper
+0x0035bb60	wrapper_cuDeviceSetGraphMemAttribute	wrapper
+0x003239d0	wrapper_cuCtxGetLimit	wrapper
+0x0037d280	wrapper_cuMemcpyHtoAAsync_v2	wrapper
+0x00340dc0	wrapper_cuMemPoolTrimTo	wrapper
+0x00346af0	wrapper_cuStreamGetCaptureInfo_v2_ptsz	wrapper
+0x00348b00	wrapper_cuEventElapsedTime	wrapper
+0x00332730	wrapper_cuMemHostGetDevicePointer_v2	wrapper
+0x003228d0	wrapper_cuCtxSetFlags	wrapper
+0x00382ef0	wrapper_cuStreamWriteValue64_ptsz	wrapper
+0x00381460	wrapper_cuLaunchKernelEx	wrapper
+0x0035c5b0	wrapper_cuGraphNodeGetDependencies	wrapper
+0x0031d320	wrapper_cuDeviceTotalMem_v2	wrapper
+0x0036d360	wrapper_cuLinkCreate	wrapper
+0x00345a00	wrapper_cuLaunchHostFunc_ptsz	wrapper
+0x00324a60	wrapper_cuCtxWaitEvent	wrapper
+0x00393e70	wrapper_cuProfilerStop	wrapper
+0x0034bae0	wrapper_cuFuncSetAttribute	wrapper
+0x00377a20	wrapper_cuTexRefSetAddress	wrapper
+0x00371280	wrapper_cuTexRefGetAddress	wrapper
+0x0031afa0	wrapper_cuKernelGetLibrary	wrapper
+0x0031c060	wrapper_cuDeviceGetUuid	wrapper
+0x003788f0	wrapper_cuCtxDestroy	wrapper
+0x00381c50	wrapper_cuStreamBatchMemOp_ptsz	wrapper
+0x00364870	wrapper_cuTexRefGetArray	wrapper
+0x00386040	wrapper_cuGraphInstantiate	wrapper
+0x0037e4c0	wrapper_cuMemcpy	wrapper
+0x0036cd20	wrapper_cuGreenCtxStreamCreate	wrapper
+0x0033d540	wrapper_cuStreamGetCtx_v2_ptsz	wrapper
+0x003940a0	wrapper_cuVDPAUGetDevice	wrapper
+0x00345610	wrapper_cuStreamAddCallback_ptsz	wrapper
+0x00375a70	wrapper_cuMemsetD2D32Async	wrapper
+0x0037ecc0	wrapper_cuMemcpy3DPeer	wrapper
+0x00348040	wrapper_cuEventRecord_ptsz	wrapper
+0x00351940	wrapper_cuGraphExternalSemaphoresSignalNodeGetParams	wrapper
+0x0031f460	wrapper_cuCtxGetDevice	wrapper
+0x00389a60	wrapper_cuEGLStreamConsumerDisconnect	wrapper
+0x00347ab0	wrapper_cuStreamSetAttribute_ptsz	wrapper
+0x00357fa0	wrapper_cuGraphAddChildGraphNode	wrapper
+0x00392ea0	wrapper_cuVDPAUCtxCreate_v2	wrapper
+0x00366cf0	wrapper_cuTensorMapEncodeIm2colWide	wrapper
+0x00351f70	wrapper_cuGraphMemAllocNodeGetParams	wrapper
+0x003622a0	wrapper_cuTexRefSetAddress_v2	wrapper
+0x0033b2d0	wrapper_cuArrayGetPlane	wrapper
+0x00378f70	wrapper_cuCtxPushCurrent	wrapper
+0x00361390	wrapper_cuGraphConditionalHandleCreate	wrapper
+0x00367fa0	wrapper_cuTexObjectGetResourceViewDesc	wrapper
+0x0036c0d0	wrapper_cuCtxGetDevResource	wrapper
+0x003862e0	wrapper_cuGraphInstantiate_v2	wrapper
+0x0037d9d0	wrapper_cuMemcpy3D_v2	wrapper
+0x00328770	wrapper_cuLibraryGetGlobal	wrapper
+0x0037c050	wrapper_cuDevicePrimaryCtxSetFlags	wrapper
+0x00365540	wrapper_cuTexRefGetMipmapLevelBias	wrapper
+0x00349a10	wrapper_cuDestroyExternalMemory	wrapper
+0x0031ca10	wrapper_cuDevicePrimaryCtxGetState	wrapper
+0x00361640	wrapper_cuOccupancyMaxActiveBlocksPerMultiprocessor	wrapper
+0x00341370	wrapper_cuMemPoolSetAttribute	wrapper
+0x0033f180	wrapper_cuMemMap	wrapper
+0x00383630	wrapper_cuStreamWaitValue32_v2	wrapper
+0x00329000	wrapper_cuKernelGetAttribute	wrapper
+0x003370d0	wrapper_cuMemcpy3DPeer_ptds	wrapper
+0x00387910	wrapper_cuMemFreeAsync	wrapper
+0x003810f0	wrapper_cuLaunchKernel	wrapper
+0x00336040	wrapper_cuMemcpyDtoD_v2_ptds	wrapper
+0x0037c310	wrapper_cuMemcpyHtoD_v2	wrapper
+0x0038a920	wrapper_cuEGLStreamProducerPresentFrame	wrapper
+0x00352ff0	wrapper_cuGraphAddNode	wrapper
+0x00384600	wrapper_cuStreamBeginCapture_ptsz	wrapper
+0x0033fb10	wrapper_cuMemGetAccess	wrapper
+0x00338230	wrapper_cuMemcpy3DAsync_v2_ptsz	wrapper
+0x00340310	wrapper_cuMemFreeAsync_ptsz	wrapper
+0x00333d40	wrapper_cuMemGetAllocationPropertiesFromHandle	wrapper
+0x0034df10	wrapper_cuLaunch	wrapper
+0x00364080	wrapper_cuTexRefSetMaxAnisotropy	wrapper
+0x0038bc10	wrapper_cuGLCtxCreate_v2	wrapper
+0x00331ec0	wrapper_cuStreamUpdateCaptureDependencies_v2_ptsz	wrapper
+0x003442d0	wrapper_cuMemRangeGetAttribute	wrapper
+0x00328a70	wrapper_cuKernelGetParamInfo	wrapper
+0x003749e0	wrapper_cuMemcpyDtoDAsync	wrapper
+0x0035a8c0	wrapper_cuGraphExecExternalSemaphoresWaitNodeSetParams	wrapper
+0x003525e0	wrapper_cuGraphAddDependencies_v2	wrapper
+0x00389640	wrapper_cuEGLStreamConsumerConnect	wrapper
+0x0037c4f0	wrapper_cuMemcpyDtoH_v2	wrapper
+0x0037d830	wrapper_cuMemcpy2DUnaligned_v2	wrapper
+0x0034c560	wrapper_cuLaunchKernel_ptsz	wrapper
+0x00356e20	wrapper_cuGraphMemsetNodeGetParams	wrapper
+0x0037f020	wrapper_cuMemcpyBatchAsync	wrapper
+0x00386bf0	wrapper_cuGraphLaunch	wrapper
+0x00393c40	wrapper_cuProfilerStart	wrapper
+0x00373a70	wrapper_cuMemcpyHtoAAsync	wrapper
+0x00348610	wrapper_cuEventSynchronize	wrapper
+0x003892b0	wrapper_cuCheckpointProcessUnlock	wrapper
+0x0037a7e0	wrapper_cuStreamIsCapturing	wrapper
+0x00333a70	wrapper_cuMemGetAllocationGranularity	wrapper
+0x003868a0	wrapper_cuGraphUpload	wrapper
+0x00382c80	wrapper_cuStreamWaitValue32_ptsz	wrapper
+0x003273c0	wrapper_cuLibraryUnload	wrapper
+0x00378c00	wrapper_cuCtxPopCurrent	wrapper
+0x00355670	wrapper_cuGraphNodeFindInClone	wrapper
+0x00378250	wrapper_cuGraphicsResourceGetMappedEglFrame	wrapper
+0x00376710	wrapper_cuMemsetD2D32	wrapper
+0x0037cc90	wrapper_cuMemcpyHtoA_v2	wrapper
+0x00325760	wrapper_cuModuleGetFunction	wrapper
+0x003523f0	wrapper_cuGraphAddDependencies	wrapper
+0x0037d480	wrapper_cuMemcpyAtoHAsync_v2	wrapper
+0x00352200	wrapper_cuGraphMemFreeNodeGetParams	wrapper
+0x00385580	wrapper_cuGraphAddKernelNode	wrapper
+0x003488b0	wrapper_cuEventDestroy_v2	wrapper
+0x0036dcc0	wrapper_cuDeviceTotalMem	wrapper
+0x003201d0	wrapper_cuDevicePrimaryCtxReset_v2	wrapper
+0x0035a210	wrapper_cuGraphExecEventRecordNodeSetEvent	wrapper
+0x0031ea00	wrapper_cuDeviceGetMemPool	wrapper
+0x00352df0	wrapper_cuGraphExecMemsetNodeSetParams	wrapper
+0x003423c0	wrapper_cuMemPoolImportFromShareableHandle	wrapper
+0x00389460	wrapper_cuGraphicsEGLRegisterImage	wrapper
+0x0033b9f0	wrapper_cuArrayGetSparseProperties	wrapper
+0x003777d0	wrapper_cuArray3DCreate	wrapper
+0x00387250	wrapper_cuIpcOpenMemHandle	wrapper
+0x00340a90	wrapper_cuMemPoolGetAccess	wrapper
+0x0031b7f0	wrapper_cuInit	wrapper
+0x003720a0	wrapper_cuMemGetAddressRange	wrapper
+0x00365fc0	wrapper_cuTexRefGetFlags	wrapper
+0x0035c2f0	wrapper_cuGraphGetEdges_v2	wrapper
+0x0036af60	wrapper_cuCoredumpGetAttribute	wrapper
+0x0033c450	wrapper_cuMipmappedArrayDestroy	wrapper
+0x0037ea90	wrapper_cuMemcpyPeerAsync	wrapper
+0x00322190	wrapper_cuCtxSetCurrent	wrapper
+0x0037c6d0	wrapper_cuMemcpyDtoD_v2	wrapper
+0x0031f1f0	wrapper_cuCtxGetCurrent	wrapper
+0x0036df20	wrapper_cuCtxCreate	wrapper
+0x00324770	wrapper_cuCtxRecordEvent	wrapper
+0x0034e730	wrapper_cuGraphCreate	wrapper
+0x0037db70	wrapper_cuMemcpyHtoDAsync_v2	wrapper
+0x003593b0	wrapper_cuGraphExternalSemaphoresWaitNodeSetParams	wrapper
+0x0033da50	wrapper_cuStreamCopyAttributes_ptsz	wrapper
+0x00388860	wrapper_cuCheckpointProcessGetRestoreThreadId	wrapper
+0x00334f30	wrapper_cuIpcOpenMemHandle_v2	wrapper
+0x00370e30	wrapper_cuGraphKernelNodeGetParams	wrapper
+0x00329c60	wrapper_cuMemFree_v2	wrapper
+0x0033ce30	wrapper_cuStreamGetFlags_ptsz	wrapper
+0x0031c560	wrapper_cuDeviceGetLuid	wrapper
+0x003742c0	wrapper_cuMemcpy2DUnaligned	wrapper
+0x0034f300	wrapper_cuMulticastBindMem	wrapper
+0x0031a410	wrapper_cuGetErrorString	wrapper
+0x0037bdb0	wrapper_cuDevicePrimaryCtxReset	wrapper
+0x00343ee0	wrapper_cuMemAdvise	wrapper
+0x00387b90	wrapper_cuMemAllocAsync	wrapper
+0x0033dd70	wrapper_cuStreamGetAttribute_ptsz	wrapper
+0x0031e4f0	wrapper_cuDeviceGetNvSciSyncAttributes	wrapper
+0x003516d0	wrapper_cuGraphEventWaitNodeSetEvent	wrapper
+0x00350bd0	wrapper_cuGraphHostNodeGetParams	wrapper
+0x00325f60	wrapper_cuLinkCreate_v2	wrapper
+0x003673f0	wrapper_cuTexRefCreate	wrapper
+0x00324ed0	wrapper_cuModuleLoadData	wrapper
+0x00354730	wrapper_cuGraphKernelNodeGetParams_v2	wrapper
+0x00380b40	wrapper_cuStreamSynchronize	wrapper
+0x0037b2b0	wrapper_cuStreamGetAttribute	wrapper
+0x0034d080	wrapper_cuParamSetSize	wrapper
+0x0033d080	wrapper_cuStreamGetId_ptsz	wrapper
+0x0034a040	wrapper_cuWaitExternalSemaphoresAsync_ptsz	wrapper
+0x003278e0	wrapper_cuLibraryGetManaged	wrapper
+0x00335380	wrapper_cuIpcCloseMemHandle	wrapper
+0x003652c0	wrapper_cuTexRefGetMipmapFilterMode	wrapper
+0x0035d590	wrapper_cuGraphDestroyNode	wrapper
+0x00341d90	wrapper_cuMemPoolDestroy	wrapper
+#
+# --- internals verified during the 2026-09 decompilation (see docs/05) ---
+0x3a9c20	cudaHandlePoolAlloc	verified
+0x3a2f50	cudaHandlePoolFree	verified
+0x453ef0	cudaEngineContextRegisterHelper	verified
+0x45a670	cudaDeviceInitEngineContexts	verified
+0x4676f0	cudaEngineContextCreateHook	verified
+0x481360	rmEngineContextAlloc	verified
+0x453cc0	cudaEngineContextCleanup	verified
diff --git a/userspace/cuda/570.144/annotations/subsystems.tsv b/userspace/cuda/570.144/annotations/subsystems.tsv
new file mode 100644
index 0000000..5a512bd
--- /dev/null
+++ b/userspace/cuda/570.144/annotations/subsystems.tsv
@@ -0,0 +1,37 @@
+# Open-IOV CUDA 570.144 subsystem index (tracked)
+#
+# Format: file_vaddr<TAB>name<TAB>subsystem<TAB>role<TAB>evidence
+# Addresses are FILE virtual addresses (libcuda+0x...). See README.md.
+#
+# subsystem values: api | engine | engine-engines | mem | stream | launch | sync |
+#                   module | rm | uvm | jit | table | unknown
+# evidence: decompiled (address inspected in the 570.144 corpus), trace (live),
+#           relocation (derived from the stub), todo (open question)
+#
+# --- public API entry stubs (generated) ---
+0x310060	cuInit	api	exported entry stub; checks destroyed-state then jumps through interposition slot	relocation
+0x310360	cuCtxCreate_v2	api	exported entry stub	relocation
+0x310b20	cuMemAlloc_v2	api	exported entry stub	relocation
+0x3106e0	cuModuleLoadData	api	exported entry stub	relocation
+0x36f420	cuLaunchKernel	api	exported entry stub	relocation
+0x36f3c0	cuStreamSynchronize	api	exported entry stub	relocation
+# --- generated wrappers (tracing/tool dispatch) ---
+0x31b7f0	wrapper_cuInit	api	initial interposition target for cuInit	relocation
+0x3206d0	wrapper_cuCtxCreate_v2	api	initial interposition target for cuCtxCreate_v2	relocation
+0x329830	wrapper_cuMemAlloc_v2	api	initial interposition target for cuMemAlloc_v2	relocation
+0x324ed0	wrapper_cuModuleLoadData	api	initial interposition target for cuModuleLoadData	relocation
+0x3810f0	wrapper_cuLaunchKernel	api	initial interposition target for cuLaunchKernel	relocation
+0x380b40	wrapper_cuStreamSynchronize	api	initial interposition target for cuStreamSynchronize	relocation
+# --- 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
+0x453cc0	cudaEngineContextCleanup	engine-engines	failure-path cleanup of an engine context; runs on the rolled-back context (rc=999)	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
+# --- 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
diff --git a/userspace/cuda/570.144/docs/ANALYSIS-cuctxcreate-segv.md b/userspace/cuda/570.144/docs/ANALYSIS-cuctxcreate-segv.md
index 9a15b51..db38802 100644
--- a/userspace/cuda/570.144/docs/ANALYSIS-cuctxcreate-segv.md
+++ b/userspace/cuda/570.144/docs/ANALYSIS-cuctxcreate-segv.md
@@ -1,5 +1,18 @@
 # libcuda 570.144 `cuCtxCreate` segfault - analysis (2026-09-23)
 
+> **Status: historical (2026-09-23 PART 9).** The segfault described here was
+> fixed in PART 10 (the RM allocation reply now echoes the request parameters,
+> which prevents a 16-byte class-params copy-back from zeroing the caller's
+> saved frame pointer). The current blocker is the missing async copy engine in
+> the emulated engine list (`cuCtxCreate` returns `999`). For the corrected,
+> current analysis see
+> [`../../docs/05-570.144-decompilation.md`](../../docs/05-570.144-decompilation.md)
+> and [`../../READ-ME-FIRST.md`](../../READ-ME-FIRST.md). The runtime
+> observations below (gdb traces, ioctl sequences) remain valid, and note that
+> the decompiled files it references were produced at the wrong Ghidra image
+> base (see PART 9's own address note and `../../docs/05` section 1.1).
+
+
 First analysis note for the CUDA userspace driver under Open-IOV. Companion to
 `../README.md` (tools) and the host/guest A/B harness (`tools/cuda-ab/`).
 
diff --git a/userspace/cuda/READ-ME-FIRST.md b/userspace/cuda/READ-ME-FIRST.md
new file mode 100644
index 0000000..08f0b28
--- /dev/null
+++ b/userspace/cuda/READ-ME-FIRST.md
@@ -0,0 +1,151 @@
+# READ-ME-FIRST — CUDA analysis guide for agents
+
+You are being pointed here because you are working on CUDA bring-up,
+debugging or emulation for Open-IOV and there is now a body of analysis that
+will save you hours. This file tells you **what to read, in what order, and
+why it matters**, plus the fastest recipes for the questions you are likely to
+have.
+
+Everything referenced here lives under `open-iov/userspace/cuda/` and was
+written by the CUDA-decompilation analysis (2026-09). Nothing in this area
+modifies the guest, the kernel module or the emulation; it is analysis plus
+tooling.
+
+---
+
+## 0. Read these two things first (5 minutes, they change how you debug)
+
+1. **`docs/05-570.144-decompilation.md` §1.1 — the address convention.**
+   Ghidra loads these binaries at **image base `0x100000`**. The original
+   per-version helper passed file vaddrs straight to Ghidra, so the old
+   decompiled files in `570.144/run/` (`helper_453ef0.c`, `fn_0x481360.c`, …)
+   are **the wrong functions** (0x100000 too low). The runtime observations in
+   `570.144/docs/ANALYSIS-cuctxcreate-segv.md` (gdb traces, live table values,
+   ioctl traces) are still valid; any conclusion drawn from its decompiled
+   code must be re-derived. The new tooling translates addresses for you:
+   `tools/ghidra_decompile_at.sh` and the annotated corpus.
+2. **`docs/05` §3 — the `cuCtxCreate` path: crash solved, engine topology
+   blocking.** The old fault at `libcuda+0x453f90` (helper at `+0x453ef0`) is
+   solved: the alloc reply must **echo the request's params**, because the
+   driver copies 16 bytes of class params back over the caller's 8-byte local
+   (the hook's saved `rbp`) — a zeroed reply is what made `rbp=0` (§3.5). The
+   current blocker is one step later: the class-filtered engine list is
+   missing the async copy engine `COPY14`, so `cuCtxCreate` returns `rc=999`
+   (§3.6). §3.3 gives the engine-record layout (`base+0x2fec`,
+   `+0x300c+(idx-1)*0x14`, `+0x350c`, `+0x351c+(idx-0x16)*0x14`; fields
+   +0/+8/+0xc) and §3.7 lists exactly what the emulation should check first.
+
+If you read nothing else, read those two.
+
+## 1. The five-document core (30–45 minutes total)
+
+Read in this order; each builds on the previous one.
+
+| # | Document | Time | What you gain |
+|---|---|---|---|
+| 1 | `docs/01-architecture.md` | 15 min | The mental model: layers (API → engine → driver-model → hardware), context/device/stream/channel/memory objects, the marker + tracking-semaphore sync model, launch/QMD flow, UVM integration, threading. This is the vocabulary the rest of the docs use. |
+| 2 | `docs/03-driver-and-gpu-interface.md` | 15 min | The emulation contract: device nodes and every ioctl CUDA issues, RM object classes, control families, memory/VA/mapping, channels/GPFIFO/USERD/doorbells, semaphores/notifiers/interrupts, UVM, GSP-RM. Ends with an 8-point emulation checklist. |
+| 3 | `docs/05-570.144-decompilation.md` | 10 min | The actual binary: function counts, address conventions, API stub → wrapper → implementation mechanism, the `cuCtxCreate` crash-path resolution and current engine-topology blocker, and how to look anything up in the corpus. |
+| 4 | `docs/04-virtualization-and-emulation.md` | 15 min | Design doc: how NVIDIA's vGPU split works, how Open-IOV differs (emulate the interface, execute the work), and the anti-fragility/performance/portability/security principles that constrain any fix you make. |
+| 5 | `docs/02-build-and-execution.md` | 10 min | How the library is built (generated API wrappers, entry tables, embedded device code) and how it executes in a process (init sequence, threads, interposition, `/dev` access, why LD_PRELOAD can't see its syscalls). |
+
+## 2. Reference material (use when you need it)
+
+* `docs/06-tooling.md` — every tool, its usage and its cost. The recipes below
+  are extracts.
+* `docs/07-version-upgrade-playbook.md` — how to redo the analysis when the
+  driver/CUDA/GSP changes. Read §1 if you only need the cheap `delta` triage.
+* `docs/08-virtualization-approaches.md` — one page on the plausible execution
+  approaches (software emulation vs host-RM proxy vs direct submission vs
+  hybrid), their trade-offs, and the honest current stage of the project.
+* `570.144/annotations/README.md` + `names.tsv` + `subsystems.tsv` — our name
+  map (659 exports, 628 resolved wrappers, verified internals) and subsystem
+  index. The map is applied to the annotated corpus already.
+
+## 3. Fast recipes (the 90% cases)
+
+**"What does the driver do with the reply we just gave it?"**
+→ Find the ioctl/control in the guest trace, then locate the consuming code:
+`docs/03` maps controls to subsystems; `docs/05` §4 shows corpus searches.
+The annotated corpus headers list callers/callees with names, so you can walk
+the path.
+
+**"Where is `libcuda+0xNNNNNN`?"**
+```
+tools/cudatool.py nearest 570.144/run/libcuda.so.570.144 0xNNNNNN
+tools/ghidra_decompile_at.sh 570.144/run/ghidra-full cuda-full \
+    libcuda.so.570.144 0xNNNNNN /tmp/fn.c
+grep -n "libcuda+0xNNNNNN" 570.144/run/analysis/decompiled-annotated/buckets/*.c
+```
+
+**"A guest segfault gives me `lib[base+size] ip` — what is it?"**
+```
+tools/cudatool.py crash 570.144/run/libcuda.so.570.144 <base> <ip>
+```
+It prints the file vaddr, containing function start, and nearest export.
+
+**"Which functions call X / does X call?"**
+```
+tools/callwalk.py 570.144/run/analysis/elf-index/callgraph.tsv callers 0xNNNNNN
+tools/callwalk.py 570.144/run/analysis/elf-index/callgraph.tsv callees 0xNNNNNN 2
+tools/callwalk.py 570.144/run/analysis/elf-index/callgraph.tsv chain 0xA 0xB
+```
+
+**"What changed between driver versions?"**
+```
+tools/cudatool.py delta <old libcuda> <new libcuda>
+```
+Same-generation builds (e.g. 570.144 vs 570.211.01) show the same 659 exports
+with small moves — the name map transfers; new/removed exports flag feature
+changes.
+
+**"I need the whole picture of a subsystem."**
+→ `docs/01` for the model, then `570.144/annotations/subsystems.tsv` for the
+known addresses, then the annotated corpus. `docs/05` §5 lists what the
+decompiler does *not* recover (types, indirect targets).
+
+## 4. What is where on disk
+
+```
+userspace/cuda/
+├── READ-ME-FIRST.md                 ← you are here
+├── README.md                        index + honest status
+├── WORKING-HERE.md                  coordination rules (run/ is shared scratch)
+├── docs/01..07                      public analysis (tracked)
+├── tools/                           reusable tooling (tracked)
+├── LOCAL-SOURCES.md               local source map (gitignored)
+└── 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
+```
+
+## 5. Rules while you work in this area
+
+* **Do not modify anything outside `userspace/cuda/` from analysis work.**
+  Emulation fixes go through the kernel module and the main handoff.
+* **Do not clean `run/`** — it holds hours of analysis. You may
+  `rm -rf 570.144/run/analysis/decompiled/wk` to reclaim ~2 GB (per-chunk
+  project copies; the batch decompiler recreates them).
+* **Decompiled output is never committed.** It is vendor-derived; the
+  gitignores in `open-iov/` and `570.144/` cover the binary, projects and
+  corpus. Tracked artifacts are our own words, names, tables and tools.
+* **Write the public description first.** Tracked docs describe components
+  and behaviour; keep implementation details in the code and its comments.
+* **When you fix the emulation**, update `docs/03` (contract) and, if the
+  interface understanding changed, `docs/05`; record the delta for the
+  version.
+
+## 6. The shortest possible version
+
+1. `docs/05` §1.1 and §3 — address fix + the `cuCtxCreate` path (crash
+   solved; engine-topology blocker) (read now).
+2. `docs/01` and `docs/03` — the model and the contract.
+3. `docs/04` — the design principles your fix must respect.
+4. `docs/06` — the recipes; then `tools/`.
diff --git a/userspace/cuda/README.md b/userspace/cuda/README.md
new file mode 100644
index 0000000..e949d0e
--- /dev/null
+++ b/userspace/cuda/README.md
@@ -0,0 +1,124 @@
+# CUDA userspace driver analysis (Open-IOV)
+
+This directory contains the Open-IOV analysis of the NVIDIA CUDA userspace
+driver (`libcuda.so.<version>`): what it is, how it is built, what it asks of
+the kernel driver and the GPU, how that maps onto Open-IOV's emulation and
+passthrough model, and the tooling to repeat the analysis on future driver
+versions.
+
+It exists to make virtualized CUDA **work correctly and fast** in Open-IOV
+VMs and to keep working across driver/firmware/hardware changes. It is **not**
+a CUDA reimplementation and not a way to produce an open-source CUDA: we
+disassemble and decompile the vendor userspace driver to understand the
+interface it expects, and we publish our own descriptions, tables and tools.
+Decompiled output is never committed.
+
+## Start here
+
+> **New here (agent or human)? Read [`READ-ME-FIRST.md`](READ-ME-FIRST.md)
+> first** — it is a reading guide: what to read in what order, what each
+> document gives you, fast lookup recipes, and the two findings that change
+> how you debug (Ghidra image base and the `cuCtxCreate` path — crash solved,
+> engine topology blocking).
+
+| Document | What it covers |
+|---|---|
+| [`docs/01-architecture.md`](docs/01-architecture.md) | the CUDA internal architecture: layers, object model, threads, memory, channels, launches, sync |
+| [`docs/02-build-and-execution.md`](docs/02-build-and-execution.md) | how the driver is compiled and how it executes in a process (exports, generated wrappers, entry tables, embedded device code) |
+| [`docs/03-driver-and-gpu-interface.md`](docs/03-driver-and-gpu-interface.md) | everything the driver touches below itself: device nodes, RM controls, memory, channels, doorbells, UVM, interrupts |
+| [`docs/04-virtualization-and-emulation.md`](docs/04-virtualization-and-emulation.md) | NVIDIA's vGPU split, how Open-IOV differs, anti-fragility, performance, portability, security |
+| [`docs/05-570.144-decompilation.md`](docs/05-570.144-decompilation.md) | the 570.144 binary map, key functions, call chains, the solved `cuCtxCreate` crash, and the current engine-topology blocker |
+| [`docs/06-tooling.md`](docs/06-tooling.md) | the reusable toolchain: ELF toolkit, Ghidra analysis/decompilation, name transfer, version deltas |
+| [`docs/07-version-upgrade-playbook.md`](docs/07-version-upgrade-playbook.md) | step-by-step method to re-run the analysis on a new driver version |
+| [`docs/08-virtualization-approaches.md`](docs/08-virtualization-approaches.md) | one-page comparison of execution-forwarding options (software emulation, host-RM proxy, direct submission, hybrid) and where the project stands |
+
+Per-version material lives in version directories (`570.144/`), including the
+crash-focused debug helpers (`570.144/scripts/`) and the original
+`cuCtxCreate` write-up (`570.144/docs/ANALYSIS-cuctxcreate-segv.md`).
+
+## Status (honest)
+
+* **570.144 decompilation: complete for the whole library.** A full Ghidra
+  auto-analysis produced 29,494 functions, 659 exported functions, 77,742 call edges and
+  a complete decompiled C corpus; the corpus is machine-generated analysis
+  output and is therefore kept out of git (see "Working with decompiled
+  output" below).
+* **Name recovery: partial.** Exports are exact; internal function names are
+  recovered from call structure, string usage and cross-version matching. The
+  name map is deliberately a separate, reviewable artifact
+  (`570.144/annotations/names.tsv`) so it accumulates across versions.
+* **Crash path (`cuCtxCreate`): solved.** The alloc reply must echo the
+  request's params: the guest driver copies 16 bytes of class params back over
+  the caller's 8-byte local, which zeroed the engine hook's saved `rbp` and
+  faulted at `libcuda+0x453f90`. Echoing the request (only `status` forced to
+  `NV_OK`) makes that copy-back idempotent; `cuCtxCreate` no longer segfaults
+  (see `docs/05-...` §3.5).
+* **Current blocker: engine topology.** `cuCtxCreate` now returns `rc=999`
+  (`CUDA_ERROR_UNKNOWN`); the guest's class-filtered engine list has 6 engines
+  (GR + CE0..CE4) where the oracle has 7 — the async copy engine `COPY14`
+  (descriptor `0x793ceb0e`, instance 14, client type `0x38`) is missing, so the
+  walker reads a zeroed secondary record, the alloc is rejected (class `0x41`,
+  st `0x19`) and the context is rolled back (`docs/05` §3.6).
+* **Open-IOV relevance: documented.** The emulation boundary, the controls the
+  driver issues, and the anti-fragility strategy are in `docs/03`/`docs/04`.
+
+## Layout
+
+```
+userspace/cuda/
+├── README.md                 this file (index + status)
+├── READ-ME-FIRST.md          reading guide for agents/debuggers
+├── WORKING-HERE.md           coordination note for parallel agents
+├── docs/                     public analysis documents
+├── tools/                    reusable, version-agnostic tooling
+│   ├── cudatool.py           ELF/export/FDE/wrapper/version-delta toolkit
+│   ├── callwalk.py           callgraph callee/caller/path walker
+│   ├── annotate_corpus.py    annotate + rename the decompiled corpus
+│   ├── ghidra_full_analysis.sh
+│   ├── ghidra_export_index.sh
+│   ├── ghidra_batch_decompile.sh
+│   ├── ghidra_index_merge.sh
+│   ├── ghidra_decompile_at.sh
+│   └── ghidra/               Ghidra Java scripts (analysis, decompile, name map)
+└── 570.144/
+    ├── README.md             version-specific debug tooling (existing)
+    ├── docs/                 version-specific write-ups
+    ├── scripts/              version-specific helpers
+    ├── ghidra/               version-specific Ghidra script
+    ├── annotations/          our name map + subsystem index (tracked)
+    └── run/                  analysis scratch: extracted binary, Ghidra projects,
+                              decompiled corpus (all gitignored)
+```
+
+## Working with decompiled output
+
+The decompiled corpus for 570.144 is under
+`570.144/run/analysis/decompiled/` (gitignored):
+
+* `buckets/bucket_<start>_<end>.c` — decompiled C for functions in that
+  program-address range, in address order, each with a header comment giving
+  the function address and `libcuda+0x...` file offset;
+* `index/bucket_<start>_<end>.tsv` + merged `decompiled_index.tsv` —
+  `addr, file_addr, size, name, file, line, status` for every function.
+
+Search the corpus by address or by name:
+
+```
+grep -n "libcuda+0x453ef0" 570.144/run/analysis/decompiled/buckets/*.c
+awk -F'\t' '$4=="FUN_0055a670"' 570.144/run/analysis/decompiled/decompiled_index.tsv
+```
+
+**Address convention (important).** Public docs cite *file virtual addresses*
+(`libcuda+0x...`). Ghidra loads these binaries at image base `0x100000`, so
+program addresses inside Ghidra/Ghidra output are `file vaddr + 0x100000`.
+The tooling translates automatically; if you use raw Ghidra scripts, check
+with `tools/ghidra/PrintProgramInfo.java`. An earlier analysis in this tree
+did not account for this and decompiled addresses 0x100000 too low — see
+`docs/05-570.144-decompilation.md`.
+
+## Rules
+
+* Tracked docs describe components and behaviour only.
+* Decompiled/disassembled vendor output is never committed (see
+  `.gitignore`); we cite `libcuda+0x...` offsets and describe behaviour.
+* Every finding lands in `docs/` (structured) and, when deep, `wiki/`.
diff --git a/userspace/cuda/docs/01-architecture.md b/userspace/cuda/docs/01-architecture.md
new file mode 100644
index 0000000..2cde42f
--- /dev/null
+++ b/userspace/cuda/docs/01-architecture.md
@@ -0,0 +1,302 @@
+# 01 — CUDA userspace driver architecture
+
+This document describes the *userspace* CUDA driver (`libcuda.so`), the layer
+that applications and language runtimes talk to. Kernel drivers, GSP firmware
+and the GPU are covered only as the interfaces this library consumes.
+
+Everything below is established from decompilation of shipped binaries,
+open-source study and live traces. Addresses are file virtual addresses
+(`libcuda+0x...`) for driver **570.144** (CUDA 12.8, build id
+`c0e1e77ea249944ac4fd43faebed9498a729606a`, md5
+`ed5b2fc0044df6052df1d5113754c499`).
+
+---
+
+## 1. What the library is
+
+`libcuda.so.1.1` is a single ~71 MB, fully stripped position-independent
+shared object. For 570.144:
+
+| Property | Value |
+|---|---|
+| Exported functions | 659 (public CUDA driver API plus internal/legacy entry points) |
+| Imported functions | 243 (libc/libpthread/libdl/librt only — no other NVIDIA userspace dependency for the CUDA path) |
+| `.text` | `0x179ec0 .. 0xe94c63` (13.7 MB) |
+| Functions found by full analysis | 29,494 (28,925 have unwind tables) |
+| Static data | `.rodata` 53 MB (error strings, tables, embedded device code), `.data`/`.bss` ~1.4 MB |
+| Build id | `c0e1e7…`; a matching `libcuda.so.1.1.debug` is referenced but not shipped |
+| CUDA version note | `.note.cuda.fwd_compatibility` says CUDA 12.8, supported kernel-driver ABIs `[470, 535, 550, 560, 565, 570]` |
+
+Two entry surfaces coexist in the same library:
+
+1. **The public driver API** — `cu*` symbols (`cuInit libcuda+0x310060`,
+   `cuCtxCreate_v2 libcuda+0x310360`, `cuMemAlloc_v2 libcuda+0x310b20`,
+   `cuLaunchKernel libcuda+0x36f420`, …).
+2. **The entry-table API** — `cuGetExportTable` (`libcuda+0x350910`),
+   a UUID-keyed lookup that hands internal function tables to the CUDA
+   runtime (`libcudart`), tooling (profiler/debugger/memcheck) and
+   interposition libraries. These tables are versioned by a leading
+   `struct_size` field and are append-only.
+
+The library embeds device code (cubin images for memcpy/memset kernels,
+graph scheduler, syscall trampolines and preemption-handler assembly) in its
+read-only data; at initialization it uploads them through the normal module
+machinery.
+
+### 1.1 The CUDA runtime (`libcudart`) and the driver
+
+The CUDA *runtime* library (`libcudart`) is a separate userspace library that
+applications normally link against; it implements the `cuda*` API (as opposed
+to the driver's `cu*` API) and is statically linked in many test programs
+(including the project's `vector_add` A/B binary). There is no kernel
+interface for it: everything it does bottoms out in `libcuda`. Two facts
+matter for analysis and debugging:
+
+* **`libcudart` calls the driver through the entry tables as well as the
+  public API.** Fatbin loading, primary-context initialization and some
+  launch paths use the UUID tables (see §2.3) for speed and to pass internal
+  parameters the public API cannot express. So when a guest misbehaves in
+  `cuda*` calls, the relevant driver code may be reached through a different
+  door than the public `cu*` entry point.
+* **`libcudart` version and driver version are negotiated.** The driver
+  reports its CUDA version (`cuDriverGetVersion`) and exposes versioned
+  entry points; the runtime selects compatible paths. A guest crash can
+  therefore involve a runtime→driver version negotiation rather than a
+  single API call.
+
+In the Open-IOV guest the runtime is **stock and not our concern to modify**;
+it is, however, part of the A/B surface: the same statically linked binary
+run on the host (real GPU, apt driver) and in the guest isolates whether a
+difference is in the runtime or below it.
+
+## 2. Layer map
+
+The library is strictly layered. Each layer has one job, and the boundary
+between layers is where the interesting interfaces live:
+
+```
+application / libcudart
+        │  public cu* API  or  UUID entry tables
+        ▼
+public API wrappers            generated: per-thread/global checks, tracing,
+        │                      interposition hook table, errors → CUresult
+        ▼
+API layer                      hand-written per domain: TLS context stack,
+        │                      init/state validation, argument checks, one
+        │                      context lock, one call into the engine
+        ▼
+CUDA internal engine            the actual driver: contexts, devices, memory,
+        │                      streams, events, modules, graphs, sync
+        ├───────────────┬───────────────────────────┬───────────────┐
+        ▼               ▼                           ▼               ▼
+driver-model layer    hardware layer            UVM userspace    JIT/ELF
+(channels, allocs,    (pushbuffer/QMD/method    (managed memory, (PTX→SASS,
+ controls; RM client)  encoding per chip)        page migration)  cubin loading)
+        │                       │                     │
+        ▼                       ▼                     ▼
+/dev/nvidia* ioctls      GPU queues/doorbells    /dev/nvidia-uvm
+        │
+        ▼
+kernel driver + GSP-RM + GPU
+```
+
+The names of the layers are descriptive working names.
+
+### 2.1 API layer
+
+The `cu*` exports are **generated**, not hand-written. A master export list
+(name, platform flags, the CUDA version in which the symbol was introduced)
+is turned into:
+
+* the exported symbol table (`cuInit`, `cuCtxCreate`, `cuCtxCreate_v2`,
+  `_ptsz`/`_ptds` per-thread-stream variants, …);
+* a wrapper per function that performs the "driver destroyed" check, optional
+  API serialization, then dispatches through a **global interposition table**
+  so tools can replace any entry point before initialization;
+* a table of tool callbacks (enter/exit events keyed by stable numeric IDs)
+  consumed by the profiling/debugging entry tables;
+* the `cuGetProcAddress` lookup tables, which resolve
+  `(name, cudaVersion, flags)` to the right symbol variant.
+
+The hand-written API layer below the wrappers is uniform:
+
+1. read the per-thread context stack (the TLS "current context");
+2. verify the driver is initialized (`CUDA_ERROR_NOT_INITIALIZED` /
+   `_DEINITIALIZED`) and that the object handles are valid;
+3. validate arguments and version gates;
+4. take exactly one context lock (read or write; recursive RW lock);
+5. call one engine function;
+6. return a `CUresult`; sticky asynchronous errors are checked here too.
+
+This uniformity is what makes the API layer cheap to map: the semantic work
+always happens one layer down, in the engine.
+
+### 2.2 CUDA internal engine
+
+The engine owns all driver state. Its main object graph:
+
+* **Process globals** — initialization state, the single kernel-client handle
+  ("RM client") for the process, device list, handle pools, caches, per-thread
+  context stack.
+* **Device** — chip/arch identity, capability bitfields, limits, hardware
+  layer vtable, driver-model vtable, primary context. One per visible GPU.
+* **Context** — an application's CUDA context: its own memory manager, channel
+  manager, stream manager, semaphore pools, QMD pools, trap handler, ISR
+  thread, profiler state. A context is a *userspace* object; at the kernel
+  level it materializes as channel groups + channels + a VA space, not as a
+  distinct kernel object.
+* **Stream** — an ordered work queue. Each stream owns a small pool of QMDs
+  (launch descriptors) and a marker (see §2.5).
+* **Module/Function** — a loaded cubin image and its kernels; functions carry
+  the compiled metadata needed to build QMDs and constant banks.
+* **Memory objects** — a hierarchy: *memblock* (an allocation from the
+  kernel driver, possibly VA-reserved and mapped), *memobj* (a suballocation
+  of a memblock or a standalone object), plus the pointer/VA heaps that
+  resolve device pointers and unified addresses back to objects.
+* **Events, graphs, libraries** — user-visible objects layered on the same
+  marker/memory machinery.
+
+### 2.3 Driver-model layer
+
+Everything that depends on *which* kernel driver is present goes through a
+function table selected at initialization:
+
+* the **RM backend** (Liberty/vGPU, Linux and Windows datacenter) — the one
+  relevant to Open-IOV;
+* a SoC backend for Tegra-class `nvrm_gpu` stacks;
+* an Amodel backend (in-process simulator);
+* an MPS backend (multi-process service client/server);
+* a WDDM backend (Windows).
+
+The RM backend translates engine operations into *RM controls* and object
+allocations over `/dev/nvidiactl`/`/dev/nvidiaN`. It also owns the
+**driver-ABI adaptation** layer: the userspace driver can talk to kernel
+drivers from several ABI generations (570.144 advertises 470…570); one
+adaptation instance per ABI is compiled into the library and the matching one
+is selected at `cuInit` by brute-force version negotiation.
+
+### 2.4 Hardware layer
+
+Chip-specific behavior — method encodings, QMD layout, semaphore release
+encodings, memcpy engine commands, doorbell format — lives behind a
+per-architecture vtable selected from the device's compute capability. The
+userspace driver *builds* the GPU command streams itself; the kernel driver
+mainly allocates objects, maps memory, schedules channels and services
+interrupts.
+
+### 2.5 Synchronization model (the core mental model)
+
+Three primitives explain almost all CUDA synchronization:
+
+1. **Tracking semaphore** — each channel has a GPU-visible memory location
+   whose value counts completed pushes. The CPU tracks three counters:
+   issued-by-CPU, issued-to-GPU and finished-by-GPU. Channel completion is
+   "semaphore reached value N".
+2. **Marker** — a small set of entries (tracking-semaphore values, QMD launch
+   IDs, CPU semaphores, UVM semaphores) that together describe "all work
+   before this point". Every stream operation appends to / waits on markers;
+   stream sync, event record and stream-wait-event are marker operations.
+3. **QMD semaphore** — each launch descriptor (QMD) carries its own completion
+   semaphore, released by the GPU when the grid completes. Markers reference
+   `(QMD, launch id)` and the driver compares the QMD semaphore payload.
+
+Blocking waits are implemented with an ISR thread per context plus a DPC
+queue; a "blocking sync" fast path lets the driver park a thread on a GPU
+event rather than spin.
+
+### 2.6 Memory model
+
+* Device pointers come from a per-context address-space manager with four
+  views: 32-bit/40-bit device pointers, full device VA, host VA and unified
+  VA (UVA). Unified addresses are resolved by a process-global registry, so
+  any context can look up a pointer allocated by another.
+* Allocation is a two-step design: reserve VA, then back it. Backing comes
+  from kernel allocations (vidmem/sysmem), imported handles (IPC), host
+  registrations, or the UVM driver (managed memory).
+* Managed memory and pageable-memory access are delegated to the UVM
+  userspace library/driver, with CUDA maintaining a stream-dependency DAG so
+  UVM can migrate pages correctly (see §2.7).
+* Copies are dispatched by a policy engine: host↔host in the CPU, small
+  host→device inline in the compute pushbuffer, otherwise staged through
+  pinned bounce buffers or sent to a copy engine channel. The choice depends
+  on size, alignment, engine availability and whether UVM is active.
+
+### 2.7 UVM integration
+
+Managed memory is a collaboration: CUDA owns streams and knows the dependency
+order; UVM owns page placement and migration. CUDA registers each channel and
+each managed allocation with UVM, tracks stream dependencies in a DAG
+(`stream → node → event`), and converts user synchronizations
+(`cuStreamSynchronize`, event waits) into UVM DAG syncs. If UVM is
+unavailable or unlicensed (vGPU profiles without UVM), the driver falls back
+to zero-copy-style behavior.
+
+### 2.8 JIT and module loading
+
+`cuModuleLoad*` accepts fatbins, PTX, cubins and (newer) libraries. The path
+is: extract the right device image for the target ISA; JIT-compile PTX if
+needed (the embedded PTX compiler is loaded as a separate library); parse and
+relocate the ELF; register kernels; apply workarounds; make constant banks and
+texture/surface bindings; then upload through the channel machinery. Kernel
+launch subsequently builds a QMD referencing the loaded code and constant
+banks, and submits it with a PCAS ("prepare/switch") method stream on a
+compute channel.
+
+## 3. Execution model inside the process
+
+* **Threads.** Every API call runs on the caller thread. The ISR thread per
+  context services interrupts/DPCs; a per-context boost thread handles
+  drop-to-idle clock management. Lock ordering is globally defined and
+  enforced in debug builds, which is why deadlock analysis needs the lock
+  order list rather than guesswork.
+* **Context stack.** The "current context" is per thread, with a stack
+  (`cuCtxPushCurrent`/`PopCurrent`). Most calls implicitly use the top of the
+  stack; the primary context is retained on first use.
+* **Lazy initialization.** `cuInit` initializes globals and enumerates
+  devices but creates no context. Device enumeration itself performs RM
+  queries; the first context creates the bulk of the resources (channels,
+  semaphores, QMD pools, ISR thread, UVM registration).
+* **Async errors.** GPU faults surface asynchronously: the driver checks
+  channel error notifiers when synchronizing or when a marker reports error,
+  and returns sticky errors from subsequent API calls.
+
+## 4. The 570.144 binary in one screen
+
+| Object | Export | File vaddr |
+|---|---|---|
+| initialization | `cuInit` | `0x310060` |
+| device enumeration | `cuDeviceGet` / `cuDeviceGetCount` | `0x3100a0` / `0x3100c0` |
+| context | `cuCtxCreate_v2` / `_v3` | `0x310360` / `0x310380` |
+| module load | `cuModuleLoadData` / `Ex` | `0x3106e0` / `0x310700` |
+| library API | `cuLibraryLoadData` | `0x3108e0` |
+| memory | `cuMemAlloc_v2` / `cuMemFree_v2` | `0x310b20` / `0x310b60` |
+| VMM API | `cuMemAddressReserve` / `cuMemCreate` / `cuMemMap` / `cuMemSetAccess` | `0x330e10`…`0x330ef0` |
+| stream/event | `cuStreamCreate` / `cuStreamSynchronize` / `cuEventRecord` | `0x331370` / `0x36f3c0` / `0x36f3e0` |
+| launch | `cuLaunchKernel` / `_ptsz` | `0x36f420` / `0x331a50` |
+| copy | `cuMemcpyHtoD_v2` / `cuMemcpyDtoH_v2` | `0x36eda0` / `0x36edc0` |
+| async alloc | `cuMemAllocAsync` | `0x36fa80` |
+| dispatch | `cuGetProcAddress` / `cuGetExportTable` | `0x36fb20` / `0x350910` |
+
+The exports are only 24–30-byte stubs; the work happens in the internal
+functions reached from them. Full call chains for the core paths (init,
+device attach, context create, memory, module load, launch, sync, export
+tables) are in `docs/05-570.144-decompilation.md` and the machine-readable
+index `570.144/run/analysis/elf-index/callgraph.tsv`.
+
+## 5. Why this architecture is good news for Open-IOV
+
+1. **One narrow kernel interface.** Everything below the engine funnels
+   through the RM backend: object allocations, controls, memory mappings and
+   the `/dev/nvidia*` ioctls. That is the surface Open-IOV emulates — not the
+   whole API.
+2. **The driver builds its own command streams.** Kernel launches, copies and
+   syncs are encoded by *userspace* into pushbuffers, GPFIFO entries and QMDs.
+   The kernel driver does not generate commands that must be faithfully
+   emulated; it allocates, maps, schedules and services interrupts.
+3. **Stable, versioned side channels.** Entry tables, `cuGetProcAddress` and
+   the driver-ABI adaptation layer are explicitly versioned, so a version
+   change moves details, not the shape.
+4. **The object model is explicit and enumerable.** The exact kernel objects a
+   context needs are visible in the decompiled call order (see `docs/05`),
+   which is what makes a host-side emulation able to satisfy the driver
+   without touching it.
diff --git a/userspace/cuda/docs/02-build-and-execution.md b/userspace/cuda/docs/02-build-and-execution.md
new file mode 100644
index 0000000..9fc2f02
--- /dev/null
+++ b/userspace/cuda/docs/02-build-and-execution.md
@@ -0,0 +1,190 @@
+# 02 — How the CUDA userspace driver is built and how it executes
+
+This document explains the compilation model behind `libcuda.so` and the
+runtime model of the resulting binary: what is generated, what is compiled in,
+what the loader does, and what the driver does from process start to first
+kernel launch. It is written to be re-derivable on any driver version with the
+tooling in `tools/`.
+
+---
+
+## 1. Build model
+
+### 1.1 One library, many components
+
+The driver is one shared object assembled from several thousand translation
+units. In architecture terms the inputs are:
+
+| Input group | Nature | Notes |
+|---|---|---|
+| API wrappers | **generated** C | one exported function per API name/version variant; adds destroyed-state check, optional serialization, tracing/interposition dispatch |
+| API implementation | hand-written C | one module per API domain |
+| CUDA internal engine | hand-written C | the bulk of the driver |
+| Driver-model backends | hand-written C/C++ | RM/SoC/simulator/MPS/WDDM behind one function-table interface |
+| Hardware layer | hand-written C | per-architecture method/QMD encoders |
+| Device code | **assembly and CUDA C** compiled to cubins, embedded as byte arrays | memcpy/memset kernels, syscall trampolines, trap/exit handlers, preemption, graph scheduler |
+| Export lists | text | symbol name + platform flags + the CUDA version the symbol appeared in |
+| ABI adaptation | **one compiled instance per kernel-driver ABI** | same source compiled repeatedly with different ABI headers and symbol localization |
+| Generators | Perl/Python/Java | produce wrappers, proc-address tables, entry tables, interposition tables, error tables, abstraction-layer headers |
+
+Two consequences matter for analysis:
+
+* **The public ABI is data-driven.** Adding an API function is mostly an entry
+  in a list plus a hand-written implementation; the exported symbol,
+  the tracing callback ID, the `cuGetProcAddress` entry and the interposition
+  slot are generated from that list. A new driver version usually moves this
+  data, not the architecture.
+* **Versioning is explicit.** Symbols carry `_vN` suffixes and a
+  "since CUDA X.Y" attribute; per-thread default stream variants (`_ptsz`)
+  exist for many calls; `cuGetProcAddress` selects by requested version.
+
+### 1.2 Embedded device code
+
+The library carries its own GPU code. Device assembly sources are preprocessed,
+assembled per architecture, converted to C arrays and linked into the library;
+CUDA C device sources are compiled with the toolchain's device compiler. At
+runtime the driver treats them like user modules: it extracts the image for
+the current architecture and loads it through its own module machinery. This
+is why a stripped `libcuda.so` still contains many small ELF images and
+`.nv.*` section names in `.rodata` — those are cubins, not host code.
+
+### 1.3 Build flavors
+
+* `release` separates debug info into a `.debug` file (which NVIDIA does not
+  ship in the driver packages we analyse, hence the fully stripped binary).
+* `develop`/`debug` add logging and assertions; some entry tables (test
+  interfaces) are present only outside public-release builds.
+* A "safety" configuration removes the tooling dispatch entirely (exports call
+  the implementation directly).
+* The kernel-module build of the same sources exists but is not what runs in
+  a guest.
+
+### 1.4 What the produced binary looks like
+
+For 570.144 (`cudatool.py info`):
+
+```
+build id        : c0e1e77ea249944ac4fd43faebed9498a729606a
+exec segment base: 0x179000
+.text           : 0x179ec0 .. 0xe94c64 (13,741,476 bytes)
+exports         : 659        imports: 243 (libc/libpthread/libdl/librt)
+cuda note       : CUDA 12.8, driver ABI [470,535,550,560,565,570], device [1,2,7..14]
+debuglink       : libcuda.so.1.1.debug (not shipped)
+```
+
+Sections of interest:
+
+| Section | Why it matters |
+|---|---|
+| `.dynsym`/`.dynstr` | the 659 exports and 243 imports; no local symbols (stripped) |
+| `.text` | all host code |
+| `.rodata` (53 MB) | error strings, dispatch/version tables, **embedded cubins** |
+| `.data.rel.ro` (~900 KB) | relocated pointer tables (entry tables, vtables, function tables) |
+| `.eh_frame`/`.eh_frame_hdr` | unwind info; also the best stripped-binary function map (28,925 FDEs) |
+| `.note.cuda.fwd_compatibility` | CUDA version + supported kernel-driver ABI list |
+| `.gnu_debuglink` | name/CRC of the (absent) debug file |
+
+## 2. Execution model
+
+### 2.1 Load and initialization
+
+1. The dynamic loader maps the four PT_LOAD segments and relocates
+   `.rela.dyn`/`.rela.plt`. The library has no NVIDIA userspace dependencies
+   for the CUDA path; everything below it is a device file.
+2. C++-style constructors registered in `.init_array` run. They set up global
+   mutexes and install an `atexit`/destructor that marks the driver destroyed.
+3. The first `cuInit` runs the initialization state machine:
+   * initialize TLS and globals; select the driver-model backend;
+   * open the kernel driver's control device and create the process-wide
+     **kernel client** (one per process, not per context);
+   * enumerate and attach GPUs, build the device list (chip, caps, limits,
+     engine lists, memory sizes, BAR info);
+   * initialize the UVM userspace layer if the platform supports it;
+   * load embedded device code caches lazily.
+4. Subsequent API calls use the per-thread context stack; the first
+   `cuCtxCreate` builds the heavy per-context state.
+
+An `LD_PRELOAD` shim cannot observe the driver's kernel traffic: the RM client
+issues **raw `syscall(SYS_ioctl)`**, not libc `ioctl`. Instrumentation is done
+with `ptrace` (see `tools/guest-diag/ioctl_trace.c` in the parent tree) or
+with kernel-side tracing.
+
+### 2.2 Interposition and tooling dispatch
+
+Every exported call goes through a global function table before reaching the
+implementation. A tool library can register an interposition layer *before*
+`cuInit` and receive the original pointers to chain to. This is the sanctioned
+mechanism used by profilers/debuggers, and it is also a clean hook point for
+A/B instrumentation in the guest.
+
+### 2.3 Entry tables
+
+`cuGetExportTable(uuid)` returns pointers to internal tables for:
+
+* the CUDA runtime library (fatbin loading, primary-context initialization,
+  hidden module-load paths, kernel launch with internal parameters);
+* tooling (profiling, debugging, memory checking, launch tracking);
+* debugger support (context-local storage, trap handlers);
+* interop and special platforms (video, NVN, OpenCL bridges).
+
+Tables are identified by 128-bit UUIDs and start with a size field, so old
+clients and new drivers coexist. This is a key fact for emulation: the guest
+runtime/tools may call *through these tables* instead of the public API, and
+all of it still resolves to the same internal engine.
+
+### 2.4 Threading and process model
+
+* API calls are caller-threaded; each context has a recursive RW lock.
+* One ISR thread per context handles interrupts/DPC callbacks.
+* One boost thread per context handles drop-to-idle clock policy.
+* MPS (multi-process service) inserts a client/server IPC layer *below* the
+  normal API: clients share the server's channels; the wire format is a
+  private socket protocol, not the public API.
+* Inter-process memory sharing uses kernel-driver exported handles and
+  (optionally) POSIX file descriptors; CUDA's own IPC mem handles are an RM
+  object plus a key.
+
+### 2.5 What the driver touches at runtime (overview)
+
+| Resource | Access | Purpose |
+|---|---|---|
+| `/dev/nvidiactl` | ioctl | control device: version negotiation, client creation, GPU attach, OS events, exported-object fds |
+| `/dev/nvidiaN` | open/ioctl/mmap | per-GPU: object alloc/free, controls, memory alloc/map, interrupts |
+| `/dev/nvidia-uvm` | open/ioctl | managed memory, page migration, channel registration (via UVM userspace API) |
+| BAR regions | mmap | user-mode doorbells (USERD), interrupt tree/usermode region, framebuffer apertures |
+| sysmem | mmap/alloc | command buffers, semaphores, page tables (when the driver owns them) |
+| shared memory / sockets | IPC | MPS, IPC events/memory |
+
+The exact ioctl vocabulary observed from the 570.144 guest is enumerated in
+`docs/03-driver-and-gpu-interface.md`.
+
+## 3. Running in a guest (what Open-IOV sees)
+
+* The library is **unmodified** in the guest; it links against the guest's
+  kernel driver (`nvidia.ko` + `nvidia-uvm.ko`) exactly as on bare metal.
+* The kernel driver is also unmodified; it talks to the emulated PCI function
+  and the GSP-RM interface that Open-IOV emulates.
+* Therefore, from libcuda's point of view, the guest is a normal NVIDIA
+  machine whose kernel driver has certain capabilities and whose GPU answers
+  RM controls. The crash class we chase (`cuCtxCreate` segfault) is a
+  *userspace* consequence of *kernel/GSP* replies that differ from bare metal.
+* The library's own state is therefore the best available oracle: when the
+  driver misbehaves, comparing its state (or its control traffic) against a
+  bare-metal A/B run localizes the emulation gap. The A/B harness in
+  `tools/cuda-ab/` does this at the workload level; `docs/05` does it at the
+  function level.
+
+## 4. Reproducing the build/execution analysis for a new version
+
+```
+tools/cudatool.py info       <libcuda.so.NEW>     # size, exports, ABI note
+tools/cudatool.py exports    <libcuda.so.NEW>     # symbol map
+tools/cudatool.py functions  <libcuda.so.NEW>     # FDE function map
+tools/cudatool.py delta <old> <new>               # interface-level diff
+tools/ghidra_full_analysis.sh <libcuda.so.NEW>    # full auto-analysis
+tools/ghidra_export_index.sh ...                  # function/call/string index
+tools/ghidra_batch_decompile.sh ...               # decompiled corpus
+```
+
+See `docs/06-tooling.md` for details and `docs/07-version-upgrade-playbook.md`
+for the full procedure.
diff --git a/userspace/cuda/docs/03-driver-and-gpu-interface.md b/userspace/cuda/docs/03-driver-and-gpu-interface.md
new file mode 100644
index 0000000..4c4caa2
--- /dev/null
+++ b/userspace/cuda/docs/03-driver-and-gpu-interface.md
@@ -0,0 +1,276 @@
+# 03 — What the CUDA userspace driver asks of the kernel and the GPU
+
+This document is the emulation contract: everything `libcuda` needs from
+below itself. It is the checklist Open-IOV's host-side emulation has to
+satisfy for unmodified CUDA to run, and the map used when a guest call fails.
+
+Sources: live ioctl traces from the guest (ptrace tracer), host-side GSP-RM
+emulation logs, and decompilation of the RM backend in `libcuda` 570.144. The
+ioctl numbers below were observed on the wire; the object classes and control
+IDs are quoted from the trace and from public CUDA/RM headers.
+
+---
+
+## 1. Device nodes and the ioctl vocabulary
+
+| Node | Used for |
+|---|---|
+| `/dev/nvidiactl` | process-wide control: version negotiation, client creation, GPU attach, OS events, exported-object file descriptors |
+| `/dev/nvidiaN` | per-GPU: object alloc/free, RM controls, memory alloc/map, event allocation |
+| `/dev/nvidia-uvm` | managed memory and page migration (UVM userspace API) |
+| `/dev/nvidia-vgpuN` | vGPU plugin node (only inside NVIDIA's own vGPU stack; not used by stock CUDA paths) |
+
+Control-device ioctls (`_IOWR('F', cmd, …)` with magic `'F'`):
+
+| cmd | Name | Purpose |
+|---|---|---|
+| 200 (`0xc8`) | `NV_ESC_CARD_INFO` | enumerate GPUs and minors |
+| 201 (`0xc9`) | `NV_ESC_REGISTER_FD` | bind a per-GPU fd to the control fd |
+| 206/207 | `NV_ESC_ALLOC_OS_EVENT` / `FREE_OS_EVENT` | OS event objects (blocking waits, GPU events) |
+| 209 (`0xd1`) | `NV_ESC_STATUS_CODE` | driver status on open failure |
+| 210 (`0xd2`) | `NV_ESC_CHECK_VERSION_STR` | RM/driver version negotiation |
+| 211 (`0xd3`) | `NV_ESC_IOCTL_XFER_CMD` | inter-GPU ioctl forwarding |
+| 212 (`0xd4`) | `NV_ESC_ATTACH_GPUS_TO_FD` | attach GPUs to an fd (exported objects) |
+| 213 (`0xd5`) | `NV_ESC_QUERY_DEVICE_INTR` | interrupt query |
+| 214 (`0xd6`) | `NV_ESC_SYS_PARAMS` | system parameters |
+
+RM escapes (same magic, used on control and per-GPU fds):
+
+| cmd | Name | Purpose |
+|---|---|---|
+| 0x27 | `NV_ESC_RM_ALLOC_MEMORY` | allocate memory (older path; NVOS02) |
+| 0x28 | `NV_ESC_RM_ALLOC_OBJECT` | allocate object (NVOS05) |
+| 0x29 | `NV_ESC_RM_FREE` | free object (NVOS00) |
+| 0x2a | `NV_ESC_RM_CONTROL` | issue an RM control (NVOS54) |
+| 0x2b | `NV_ESC_RM_ALLOC` | allocate object (NVOS21/NVOS64) — **the main alloc path** |
+| 0x34 | `NV_ESC_RM_DUP_OBJECT` | duplicate/import an object |
+| 0x35 | `NV_ESC_RM_SHARE` | share object |
+| 0x41 | `NV_ESC_RM_IDLE_CHANNELS` | idle channels |
+| 0x4a | `NV_ESC_RM_VID_HEAP_CONTROL` | heap allocation (NVOS32) |
+| 0x4e / 0x4f | `NV_ESC_RM_MAP_MEMORY` / `UNMAP_MEMORY` | map memory into the process (NVOS33/NVOS34) |
+| 0x57 / 0x58 | `NV_ESC_RM_MAP_MEMORY_DMA` / `UNMAP_MEMORY_DMA` | map memory into a device VA space (NVOS46/NVOS47) |
+| 0x5c / 0x5d | `NV_ESC_RM_EXPORT_OBJECT_TO_FD` / `IMPORT_OBJECT_FROM_FD` | cross-process object sharing |
+| 0x52 | `NV_ESC_RM_GET_EVENT_DATA` | read event payload (NVOS41) |
+| 0x32/0x33 | `NV_ESC_RM_CONFIG_GET/SET` | registry-style configuration |
+
+In a trace, the ioctl request word encodes the escape and the size; e.g.
+`_IOWR('F',0x2b,48)` = `NV_ESC_RM_ALLOC` with a 48-byte payload. Seeing this
+request in a crash register dump (`rsi=0xc030462b`) identifies "driver was
+allocating an RM object" — that is exactly what happens at the `cuCtxCreate`
+crash (see `docs/05`).
+
+## 2. The RM object model as CUDA uses it
+
+Every RM call names a **class**. The classes CUDA allocates, in the order
+observed for 570.144 on an emulated TU104:
+
+| Class | Name (public RM header) | When |
+|---|---|---|
+| `0x0` | `NV01_ROOT` | process client (once, at `cuInit`) |
+| `0x80` | `NV01_DEVICE_0` | per GPU, at attach |
+| `0x2080` | `NV20_SUBDEVICE_0` | per GPU/subdevice |
+| `0x40` | `NV01_EVENT` | events (several) |
+| `0xa06c` | `KEPLER_CHANNEL_GROUP_A` | context channel group (TSG) |
+| `0x9067` | channel-group/runlist object | context |
+| `0x3e` | `NV01_MEMORY_*` | memory objects |
+| `0xc46f` | `TURING_CHANNEL_GPFIFO_A` | the channel itself |
+| `0xc5c0` | `TURING_COMPUTE_A` | compute engine object bound to the channel |
+| `0xc5b5` | (compute class, second engine) | used by the host A/B run for a CE-style object |
+| `0xa080` | `KEPLER_DEVICE_VGPU` | vGPU capability query (vGPU only) |
+| `0x90f1`-family | VA space (`FERMI_VASPACE_A` etc.) | memory model |
+
+`cuCtxCreate`'s observed ioctl sequence (guest, 570.144):
+
+```
+NV_ESC_RM_ALLOC class=0x0      (NV01_ROOT)
+NV_ESC_RM_ALLOC class=0x80     (NV01_DEVICE_0)
+NV_ESC_RM_ALLOC class=0x2080   (NV20_SUBDEVICE_0)
+NV_ESC_RM_ALLOC class=0x40     (events)
+NV_ESC_RM_ALLOC class=0xa06c   (KEPLER_CHANNEL_GROUP_A)
+NV_ESC_RM_ALLOC class=0x9067   (channel group object)
+NV_ESC_RM_ALLOC class=0x3e     (memory)
+NV_ESC_RM_ALLOC class=0xc46f   (TURING_CHANNEL_GPFIFO_A)      <- channel
+NV_ESC_RM_CONTROL cmd=0x80170d (NV0080_CTRL_FIFO_GET_CHANNELLIST)
+NV_ESC_RM_ALLOC class=0xc5c0   (TURING_COMPUTE_A)             <- last alloc
+[engine registration attempt; pre-fix this segfaulted — now rc=999, see docs/05 §3.5-3.6]
+```
+
+The first two allocs (`NV01_ROOT`, `NV01_DEVICE_0`) actually happen at
+`cuInit`/device attach; the trace window shows them repeated per client setup.
+
+## 3. RM controls the CUDA path uses
+
+Controls are `NV_ESC_RM_CONTROL` with a 32-bit command ID and a fixed-size
+parameter block. Families relevant to CUDA bring-up (IDs in the low 32 bits
+are `(class << 16) | cmd`; the list below is from traces and decompiled call
+sites, not exhaustive):
+
+* **Device/GPU info** — `NV0080_CTRL_CMD_GPU_GET_VIRTUALIZATION_MODE`
+  (`0x00800089`), `NV2080_CTRL_CMD_GPU_GET_INFO`, `GPU_GET_ENGINES`,
+  `GPU_GET_CLASSLIST`, `GPU_GET_SIMULATION_INFO`, `FB_GET_INFO`,
+  `BUS_GET_INFO`, `BUS_GET_PCI_INFO`, `BUS_GET_PCI_BAR_INFO`,
+  `MC_GET_ARCH_INFO`, `GR_GET_INFO`, `GR_GET_GPC_MASK`, `GR_GET_TPC_MASK`,
+  `GR_GET_GLOBAL_SM_ORDER`, `FB_GET_LTC_INFO_FOR_FBP`, `CE_GET_CAPS`,
+  `GSP_GET_FEATURES`, `GPU_GET_SW_FEATURES`, `GPU_GET_MEMOP_ENABLE`,
+  `FB_GET_CAPS`, `GR_GET_CAPS_V2`.
+* **FIFO/channels** — `NV0080_CTRL_CMD_FIFO_GET_CHANNELLIST` (`0x80170d`),
+  `NV0080_CTRL_CMD_FIFO_GET_CAPS`, `NVA06F/NVA06C_CTRL_CMD_GPFIFO_SCHEDULE`,
+  `SET_ERROR_NOTIFIER`, `NVA06C_CTRL_CMD_SET/GET_TIMESLICE`,
+  `NV2080_CTRL_CMD_FIFO_DISABLE_CHANNELS`, per-class
+  `GET_CLASS_ENGINEID` (e.g. `NV906F`, `NVA06F`-family),
+  `NVC36F_CTRL_CMD_GPFIFO_GET_WORK_SUBMIT_TOKEN`.
+* **GR/context** — `GR_CTXSW_PM_MODE`, `GR_CTXSW_SMPC_MODE`,
+  `GR_SET_CTXSW_PREEMPTION_MODE`, `GR_PC_SAMPLING_MODE`,
+  `GR_GET_CTX_BUFFER_SIZE`, `FLCN_GET_CTX_BUFFER_SIZE`.
+* **Memory** — `FB_GET_INFO` (BAR1 info), `NV0000_CTRL_CMD_OS_UNIX_*`
+  (flush user cache, exported objects), `NV0041_CTRL_CMD_GET_SURFACE_INFO`,
+  `NV0000_CTRL_CMD_CLIENT_GET_ADDR_SPACE_TYPE`.
+* **Events/interrupts** — `NV90E6_CTRL_CMD_MASTER_GET_VIRTUAL_FUNCTION_ERROR_CONT_INTR_MASK`
+  (interrupt-tree setup), `NV83DE_CTRL_CMD_DEBUG_*` (SM error states, debug
+  mode), `RC_ENABLE/DISABLE/SOFT_DISABLE_WATCHDOG`, `RELEASE_WATCHDOG_REQUESTS`.
+* **vGPU/licensing** — `NVA080_CTRL_CMD_VGPU_GET_CONFIG` (`0xa0800301`),
+  `NVA080_CTRL_CMD_VGPU_SET_LICENSE_INFO`,
+  `NV2080_CTRL_CMD_GPU_GET_LICENSABLE_FEATURES`,
+  `GPU_GET_GRID_UNLICENSED_STATE_MACHINE_INFO` (`0x2080018c`).
+* **Perf/clocks** — `GPU_GET_COMPUTE_POLICY_CONFIG`, boost/sync-boost
+  controls, `NV2080_CTRL_CMD_GPU_ACQUIRE/RELEASE_COMPUTE_MODE_RESERVATION`.
+
+Traces show the driver tolerating unknown-but-successful controls: it only
+fails when a *specific* reply field is consumed and wrong. This is why the
+emulation can start with generic success and then bisect the consumed fields
+(documented in `docs/04` as the "reply-fidelity ladder").
+
+## 4. Memory: allocation, mapping, VA
+
+CUDA's memory path uses three layers:
+
+1. **VA reservation** — `NV_ESC_RM_VID_HEAP_CONTROL` with a virtual
+   `NVOS32_FUNCTION_ALLOC_SIZE` request (or no reservation at all when UVM
+   owns the VA space).
+2. **Backing allocation** — vidmem/sysmem `NVOS32` allocation with attributes:
+   location (`VIDMEM`/`PCI`/`SYSTEM`), page size, coherency, compression,
+   protection, fixed physical address (if requested), or pinning an existing
+   host address (`NV01_MEMORY_SYSTEM_OS_DESCRIPTOR`).
+3. **Mapping** — `NV_ESC_RM_MAP_MEMORY` (`NVOS33`) to map into the process
+   (for host-visible/CPU access, flags include MAP_FIXED, RESERVE_ON_UNMAP,
+   DIRECT for incoherent sysmem) and `NV_ESC_RM_MAP_MEMORY_DMA` (`NVOS46`) to
+   map into a device VA space (flags include RW/RO, snoop, fixed DMA offset,
+   page size, peer flags).
+
+Related classes: `FERMI_VASPACE_A` (VA space; one is shared with UVM, one
+legacy, optionally one for fabric-local aperture), `NV01_MEMORY_FABRIC*` /
+`NV_MEMORY_FABRIC*` (NVLink/fabric), `NV01_MEMORY_SYSTEM_OS_DESCRIPTOR`
+(pinned host memory), `NV01_MEMORY_LOCAL_USER` (vidmem).
+
+For Open-IOV the practical summary is: **all guest memory that CUDA uses is
+either an RM allocation backed by emulated VRAM/BAR1, or pinned guest RAM that
+the GPU must be able to DMA to.** The project's current model backs VRAM with
+host RAM and executes push-buffer work in the host (see the parent handoff).
+
+## 5. Channels, GPFIFO, USERD and doorbells
+
+A CUDA channel is a GPFIFO (array of pushbuffer descriptors) plus one or more
+engine objects:
+
+* GPFIFO memory: allocated by CUDA (vidmem or sysmem), passed to the kernel
+  driver at channel creation with its offset and entry count, together with
+  an error notifier object.
+* **USERD**: on Volta and later, CUDA allocates the user-mode doorbell/control
+  area itself and passes its handles to the channel allocation; older
+  generations map the kernel-provided USERD. On Hopper, a second BAR1-mapped
+  usermode region is used for doorbells.
+* **Doorbell**: submission writes the GPFIFO put pointer into USERD and rings
+  a doorbell token (obtained via `NVC36F_CTRL_CMD_GPFIFO_GET_WORK_SUBMIT_TOKEN`
+  on newer ABIs, or a static token on older ones) at a token/offset the
+  channel allocation returned.
+* **Pushbuffers**: command streams in vidmem or sysmem; a push is a sequence
+  of methods (`SET_OBJECT`, engine methods, semaphore releases) terminated by
+  a tracking-semaphore release and a new GPFIFO entry.
+* **Context share (Volta+)**: subcontexts are created by allocating a
+  `FERMI_CONTEXT_SHARE_A` object with the VA space handle and subcontext id;
+  channels then reference it instead of the VA space directly.
+
+Open-IOV relevance: the host emulation sees the GPFIFO and pushbuffer memory
+the guest writes. The project already decodes push buffers as method streams
+and executes memory ops (`kernel/gsp_sema_service.c`); this section defines
+the contract those streams follow.
+
+## 6. Synchronization objects below the API
+
+* **Semaphores** — plain 32/64-bit memory values, allocated by CUDA in
+  GPU-visible memory. The GPU releases them (semaphore release methods or
+  inline memcpy of the payload), the CPU polls or waits. Tracking semaphores
+  count completed pushes; QMD semaphores count completed grids.
+* **Notifiers** — 4 KB error-notifier memory per channel; the kernel driver
+  writes a notification record on channel errors, which CUDA reads when
+  synchronizing to map to `CUDA_ERROR_LAUNCH_*`.
+* **OS events** — kernel objects (`NV_ESC_ALLOC_OS_EVENT`) used to wake
+  blocked threads; GPU events bind an engine event to an OS event so an
+  interrupt can signal user space.
+* **Interrupts** — the kernel driver owns the GPU interrupt; CUDA's ISR
+  thread is woken through the OS event/monitored-fence mechanism and then
+  walks channel state, updates blocking-sync accounting and runs DPCs.
+* **Trap handler** — for debuggers and `device-side assert`, CUDA installs a
+  trap handler (embedded device code) and reads a trace buffer in GPU memory.
+
+## 7. UVM interface
+
+When UVM is available, the userspace driver calls the UVM API (through
+`/dev/nvidia-uvm`):
+
+* set the expected kernel driver version (forward compatibility);
+* initialize UVM (sharing mode for MPS, HMM/pageable-memory options);
+* register each GPU and each VA space;
+* register each channel with its platform parameters (the RM handles and
+  context-buffer VA/size the kernel driver provided);
+* allocate/register managed allocations, prefetch/advice calls;
+* semaphores: UVM semaphore acquire/release are pushed into the same streams.
+
+If UVM is unavailable or the vGPU profile lacks UVM/fault capability, CUDA
+disables kernel-mode UVM and degrades to zero-copy behavior.
+
+## 8. GSP-RM in the picture
+
+Modern drivers run the resource manager on the GSP (a RISC-V core on the
+GPU). From the userspace driver's perspective this changes nothing: it still
+issues RM ioctls. Inside the kernel driver, controls become GSP-RM RPCs; in
+the Open-IOV emulation, the host answers those RPCs. The host log for the
+`cuCtxCreate` window shows, in order:
+
+```
+CE_GET_FAULT_METHOD_BUFFER_SIZE -> 12288
+GPU_PROMOTE_CTX (0x2080012b, 560 B) -> generic SUCCESS
+TURING_CHANNEL_GPFIFO_A alloc, TURING_COMPUTE_A alloc
+```
+
+The userspace failure now surfaces after the engine registration attempt: with
+the alloc reply fixed (`docs/05` §3.5) the walker reads a zeroed secondary
+engine record and `cuCtxCreate` returns `rc=999` (`docs/05` §3.6). The
+function-level path is in `docs/05`; the emulation-side question is what input
+builds the driver's class-filtered engine list.
+
+## 9. Emulation checklist (derived from the above)
+
+For an unmodified CUDA 12.x userspace driver in a guest to reach a working
+context, the host emulation must provide, at minimum:
+
+1. **Control device semantics** — version negotiation, client creation, GPU
+   attach, OS events, per-GPU fds.
+2. **Object allocation** for the classes in §2, with plausible sizes/handles,
+   and successful `NV_ESC_RM_FREE`.
+3. **Device/engine topology** — a consistent engine list and class list
+   (`GPU_GET_ENGINES`, `GPU_GET_CLASSLIST`, static info) matching what the
+   caps/controls claim; per-engine records must be self-consistent (see the
+   engine-descriptor discussion in `docs/05`).
+4. **Memory** — VA spaces, vidmem/sysmem allocations with the requested
+   attributes, device and host mappings, BAR1 info.
+5. **Channels** — GPFIFO alloc with notifier and USERD handles, engine object
+   alloc, scheduling, channel ids, doorbell tokens/offsets.
+6. **Interrupt/event plumbing** — interrupt tree or usermode region, OS
+   events, error notifiers.
+7. **Reply fidelity for consumed controls** — the small set of controls whose
+   *values* (not just status) the driver reads: memory sizes, engine ids,
+   class engine ids, BAR1 info, channel ids, version info, UVM caps.
+8. **Push-buffer execution** — the GPU work encoded by userspace must actually
+   happen (memsets, copies, semaphore releases, kernel grids) or the driver
+   will hang in synchronization rather than fail fast.
diff --git a/userspace/cuda/docs/04-virtualization-and-emulation.md b/userspace/cuda/docs/04-virtualization-and-emulation.md
new file mode 100644
index 0000000..89adc35
--- /dev/null
+++ b/userspace/cuda/docs/04-virtualization-and-emulation.md
@@ -0,0 +1,244 @@
+# 04 — Virtualizing CUDA: NVIDIA's vGPU model, Open-IOV's model, and how to stay unbreakable
+
+This is the design document for CUDA virtualization in Open-IOV. It answers
+four questions:
+
+1. How does NVIDIA virtualize CUDA in its own vGPU product?
+2. How does Open-IOV's model differ, and why?
+3. How do we stay working across CUDA/driver/GSP/firmware/hardware changes
+   (anti-fragility)?
+4. How do we get performance, portability, security and correctness?
+
+Companion documents: `docs/01-architecture.md` (what the driver is),
+`docs/03-driver-and-gpu-interface.md` (the interface we must satisfy).
+
+---
+
+## 1. NVIDIA's vGPU architecture (behavioural summary)
+
+NVIDIA's vGPU stack splits the driver across the VM boundary. The **guest is
+fully unmodified**: stock `libcuda`, stock `nvidia.ko`, stock `nvidia-uvm.ko`.
+Virtualization is hidden below the RM client interface, using a mediated PCI
+function and a host-side device model:
+
+```
+guest userspace      libcuda (unmodified)  ── /dev/nvidia* ioctls ──┐
+guest kernel         nvidia.ko, nvidia-uvm.ko (unmodified)         │
+                     │                                             │
+                     │  emulated PCI function: config space,       │
+                     │  BAR0/BAR1 MMIO, doorbells, MSI-X, VGA      │
+                     ▼                                             │
+hypervisor           mediated device (VFIO/mdev on KVM, VMBus on    │
+                     Hyper-V, device model on Xen/VMware)          │
+                     │                                             │
+host kernel          vGPU kernel module (VFIO side)                │
+                     │  ioctls: create/destroy, pin/unpin pages,   │
+                     │  reg access, mmio mappings, interrupt       │
+                     │  injection, migration state                 │
+                     ▼                                             │
+host userspace       vGPU manager daemon + plugin (VMIOP)  ◄───────┘
+                     │  NvRm API
+                     ▼
+host kernel + GPU    host nvidia.ko / GSP-RM / physical GPU
+```
+
+The pieces that make this work:
+
+* **A mediated PCI function.** The guest sees a GPU with a *vGPU profile*
+  (a slice of framebuffer, a set of engines, display heads, encoder capacity).
+  Config space, BAR sizes, MMIO registers, doorbells and MSI-X are emulated or
+  forwarded; the profile is what licenses and limits the VM.
+* **A guest-host wire protocol (VMIOP) between the guest kernel driver and the
+  host plugin.** The guest kernel driver's MMIO/register accesses arrive at
+  the host plugin as emulation callbacks (space, offset, width, data,
+  cacheability). The plugin either satisfies them from backing memory (with
+  per-register callback gating) or forwards them into host RM.
+* **Guest memory mapping and pinning.** Guest pages are pinned and mapped into
+  the plugin so the physical GPU can DMA into guest RAM; the guest framebuffer
+  is carved from physical VRAM and mapped into the guest's physical address
+  space.
+* **Interrupt injection.** Device interrupts are delivered to the guest as
+  MSI/MSI-X through the hypervisor.
+* **A capability negotiation on the CUDA side.** The CUDA userspace driver
+  detects vGPU (`GPU_GET_VIRTUALIZATION_MODE == VGX`) and queries a vGPU
+  capability block: is CUDA enabled for this profile, is UVM available, are
+  replayable faults available, is GPU-direct RDMA available, is TSG timeslice
+  override available. It then *restricts itself*: no UVM if the profile lacks
+  it, no debugger if disabled, no checkpoint/CRIU, CILP instead of MBP for
+  preemption, and it enforces the GRID license state (unlicensed VMs get
+  `CUDA_ERROR_DEVICE_NOT_LICENSED` or a sleep/timebomb behavior).
+* **Real objects on the host.** Channels, VA spaces, memory objects, engine
+  contexts and command execution are *real* host RM objects bound to real
+  channels; command streams are passed through and executed by the physical
+  GPU. What is virtualized is the *view* (profiles, MMIO, interrupts,
+  migration), not the execution.
+
+Two design consequences worth naming:
+
+* The guest driver is never patched; the virtualization is data (capabilities,
+  object parameters, emulated registers) plus a host device model.
+* The split is at the *kernel-driver ↔ hardware* boundary. Everything above
+  it (CUDA userspace, RM API semantics) is the stock stack.
+
+The exact plugin interfaces, capability bits, object classes and migration
+stages are vendor-internal and not needed to build an interoperable
+alternative.
+
+## 2. Open-IOV's model: emulate the interface, pass through the work
+
+Open-IOV makes the same fundamental choice as NVIDIA — **the guest stays
+stock** — but implements the host side differently:
+
+| Concern | NVIDIA vGPU | Open-IOV |
+|---|---|---|
+| Guest stack | stock | stock (identical constraint) |
+| Device presented | mediated PCI function with a vGPU profile, backed by host RM | emulated PCI function (mdev/VFIO) backed by the host module + GSP-RM emulation |
+| Where commands execute | physical GPU via host RM channels | host executes the work the guest encodes (e.g. memsets/copies/semaphore releases) or passes through, depending on the resource |
+| Host-side control plane | vGPU manager daemon + VMIOP plugin + host RM | `open-iov-vfio.ko` emulating the GSP-RM interface and device model directly in the kernel |
+| Licensing | GRID license server / vGPU licensing; CUDA gated on license state | no GRID licensing; the driver's vGPU detection must be answered in a way that does not gate CUDA (or the guest must not be told it is a licensed-vGPU-only environment) |
+| CUDA userspace | unmodified, capability-limited by vGPU profile | unmodified, no profile limits imposed |
+| Migration | vGPU-specific save/restore protocol with dirty-page tracking | 4-phase migration in the host module (project feature) |
+
+The key architectural difference: **Open-IOV does not need to virtualize the
+GPU's register interface at all if it can emulate the GSP-RM interface the
+kernel driver consumes.** The kernel driver asks the GSP (via RPCs) for
+objects, memory, channels, engines and caps; if the emulation answers those
+RPCs coherently, the driver builds its own command streams and the guest CUDA
+stack works. This moves the emulation boundary *up* from MMIO registers to
+RPC messages, which is far more stable: RPCs are versioned and typed, MMIO
+register maps are not.
+
+Where pass-through is possible (real GPU assigned, SR-IOV, or a future
+split-execution mode), the same userspace stack runs unchanged; the host-side
+work is the only difference. That is the portability goal: the analysis in
+`docs/01`–`docs/03` describes the *interface*, and both emulation and
+passthrough are implementations of it.
+
+## 3. Anti-fragility: how we survive version churn
+
+CUDA, the kernel driver, GSP firmware and GPU architectures all change on
+independent schedules. Our analysis and implementation are designed so that
+change moves *data*, not *architecture*:
+
+1. **Anchor on the stable interfaces.** The kernel-driver interface (RM
+   objects, controls, ioctls) and the userspace driver's internal layering
+   have been stable for many years; the driver-ABI adaptation layer inside
+   `libcuda` exists precisely because NVIDIA itself maintains that stability.
+   We track the interface, not the implementation.
+2. **Prefer behaviour over constants.** Every numeric ID we depend on
+   (classes, controls, capability bits) is versioned in config
+   (`config/versions/*.yaml`, `config/hardware/*.yaml`), with the analysis
+   tooling able to re-derive them from a new binary (`cudatool.py delta`,
+   `docs/07`).
+3. **Never require driver patches.** No guest modification, no userspace
+   shim, no `LD_PRELOAD`. Anything that would break on a driver update is
+   rejected by design. The only supported hook points are the ones NVIDIA
+   itself keeps stable: entry tables, `cuGetProcAddress` versioning, the RM
+   ABI.
+4. **Emulate replies, not side effects.** The host emulation answers RPCs with
+   *coherent data*; it does not fake hardware state that the driver might
+   later cross-check. Where a value is consumed (sizes, IDs, caps), we must
+   supply a consistent value — the crash analysis in `docs/05` is exactly a
+   case of inconsistent data reaching a userspace consumer.
+5. **Name transfer instead of re-analysis.** Function names and structure
+   knowledge are stored in a reviewable name map (`570.144/annotations/`) and
+   can be re-applied to a new version in seconds, then corrected where
+   signatures moved. Analysis effort accumulates across versions.
+6. **Decompose by failure mode.** The analysis is organized by what can break:
+   * *API layer* — new/changed entry points (`delta` detects them);
+   * *engine* — new object types, new control flows;
+   * *driver-model* — new controls/ABI versions;
+   * *hardware layer* — new chips and QMD/method encodings;
+   * *device code* — new syscalls/trampolines embedded in the library;
+   * *UVM* — page-migration and channel registration changes.
+7. **Keep a reference oracle.** The same CUDA workload runs on bare metal
+   (host A/B harness) and in the guest; divergence localizes the break to the
+   emulation, and the ioctl/control traces localize it further. This is the
+   fastest known way to find an emulation gap without understanding the whole
+   stack.
+
+## 4. Performance
+
+The performance budget for virtualized CUDA is dominated by three things, in
+order:
+
+1. **Push-buffer execution.** Kernel launches, copies and syncs are encoded by
+   guest userspace into memory. The host must execute that work efficiently.
+   The project's design decodes pushes as method streams and executes them
+   (memsets/copies/semaphore releases) rather than round-tripping each method.
+   Fast path: operate directly on the mapped memory; no per-method syscall.
+2. **Synchronization latency.** Every launch ends in a semaphore release; every
+   `cuStreamSynchronize` waits on a marker. Host-side completion must be
+   signaled promptly (ISR/MSI-X injection latency matters more than raw
+   throughput). The userspace driver's blocking-sync path can park threads on
+   OS events, so delivering the completion interrupt promptly avoids spinning.
+3. **Control-plane latency (setup, not steady state).** Context creation,
+   module load, first launch are control-heavy; these happen once. Optimizing
+   them is about correctness first (see `docs/05`), not microseconds.
+
+Design rules that follow:
+
+* **Fast path in the host kernel**: the push decoder and semaphore service run
+  in the kernel module, with no daemon in the loop for steady-state work.
+* **Batching**: process a whole push (or a whole page of pushes) in one pass,
+  keeping decode state across methods; avoid per-method wakeups.
+* **No copy where mapping suffices**: guest memory the GPU reads/writes is
+  pinned and mapped; the emulation should operate on it in place.
+* **Measure against the A/B oracle**: `tools/cuda-ab/` gives host-vs-guest
+  numbers for the same binary; a regression is visible immediately.
+* **Latency instrumentation**: RPC round-trip and interrupt-injection timing
+  are traceable (`trace_printk` in the module) so hot spots are found with
+  data.
+
+## 5. Portability
+
+The same analysis supports:
+
+* **Multiple driver versions** (550/560/570/580+, Windows and Linux): the
+  versioned interfaces are data-driven; the name map and config carry the
+  per-version facts.
+* **Multiple GPU architectures** (Turing → Blackwell): the hardware layer
+  (QMD/method encodings) is selected by compute capability inside the driver;
+  the host must only understand the *effects* (copy/memset/semaphore/grid),
+  not every encoding. Where a new architecture introduces a genuinely new
+  effect, the decoder learns it, not the whole stack.
+* **Multiple hypervisors**: the emulation boundary is VFIO/mdev today, but the
+  interface analysis is hypervisor-neutral (the same stock guest runs under
+  libvirt, / or raw QEMU).
+* **Emulation and passthrough**: as noted in §2, both are implementations of
+  the same interface; a passthrough deployment validates the guest stack and
+  narrows the emulation's responsibility to the control plane.
+
+## 6. Security and correctness
+
+* **Isolation.** The guest's CUDA objects must map to host objects with the
+  guest's identity and limits; no guest handle may name a host object it does
+  not own. Handles are per-client; the emulation must not let a guest forge
+  handles or reuse freed ones (handle generation).
+* **Memory safety.** Every guest-provided address/size (push buffers, GPFIFO
+  entries, semaphore addresses, VA ranges) is untrusted input to the host
+  decoder. Validate ranges against the mapped regions, bound loops, and never
+  index arrays with guest-controlled values.
+* **DMA containment.** With an IOMMU, guest memory used by the GPU is mapped
+  explicitly; the emulation must not grant broader DMA windows than the
+  allocation requires.
+* **Interrupt/event safety.** Event objects are per-client; signalling must be
+  routed to the right VM only.
+* **Correctness before speed.** A wrong reply is worse than a slow one: it
+  corrupts driver state (as in `docs/05`). The reply-fidelity ladder is:
+  status-only → coherent zeroed params → coherent values for consumed fields →
+  exact values. Every rung should be validated by the A/B harness before
+  moving on.
+* **Auditability.** All host-side emulation decisions are logged (control ID,
+  size, status, consuming context) so a guest failure can be reconstructed
+  from the host log alone.
+
+## 7. The practical path (where this analysis plugs in)
+
+1. Use `docs/03` as the contract checklist for the emulation.
+2. Use `docs/05` + the decompiled corpus to answer "what does the driver do
+   with the reply we just gave it" for any failure.
+3. Use `docs/06`/`docs/07` to re-run the analysis when the driver, CUDA
+   runtime, GSP firmware or hardware changes.
+4. Keep the guest stock; keep the interface emulated; keep the work executed.
+   Everything in this directory is organized around that sentence.
diff --git a/userspace/cuda/docs/05-570.144-decompilation.md b/userspace/cuda/docs/05-570.144-decompilation.md
new file mode 100644
index 0000000..84b377e
--- /dev/null
+++ b/userspace/cuda/docs/05-570.144-decompilation.md
@@ -0,0 +1,315 @@
+# 05 — 570.144 decompilation: map, key paths, and the `cuCtxCreate` bring-up
+
+This document is the version-specific result of the comprehensive
+decompilation of `libcuda.so.570.144` and the reference for anyone debugging
+the guest's CUDA bring-up. Addresses are **file virtual addresses**
+(`libcuda+0x...`).
+
+Artifacts behind it (gitignored, local):
+
+* `run/libcuda.so.570.144` — the guest's library, md5
+  `ed5b2fc0044df6052df1d5113754c499`;
+* `run/ghidra-full/` — the analyzed Ghidra project (29,494 functions);
+* `run/analysis/elf-index/` — `functions.tsv`, `exports.tsv`,
+  `callgraph.tsv` (77,742 edges), `strings.tsv` (102,324 strings),
+  `blocks.tsv`;
+* `run/analysis/decompiled/` — the full decompiled corpus
+  (`buckets/*.c`, `index/*.tsv`).
+
+Reproduce with `tools/ghidra_full_analysis.sh`, `tools/ghidra_export_index.sh`
+and `tools/ghidra_batch_decompile.sh` (see `docs/06-tooling.md`).
+
+---
+
+## 1. Analysis result
+
+| Metric | Value |
+|---|---|
+| Functions (full auto-analysis) | 29,494 |
+| Functions with unwind info (FDEs) | 28,925 |
+| Exported symbols | 659 (plus 31 non-function exports) |
+| Call edges | 77,742 |
+| Defined strings | 102,324 |
+| Functions decompiled | 29,254 (all non-external functions in the executable range; 240 external imports) |
+| Decompilation failures | 0 |
+
+The analysis used Ghidra 12.1.2 with the two optional identifier analyzers
+disabled (parameter-ID and function-ID); everything else (disassembly,
+aggressive instruction finding, exception handlers, stack, reference
+analyzers) was on. This is the right trade-off for a stripped binary: it
+finishes in minutes and still finds essentially every function.
+
+### 1.1 Address convention (read this before using any offset)
+
+Ghidra loads these shared objects at **image base `0x100000`**, so:
+
+```
+Ghidra program address = file virtual address + 0x100000
+```
+
+The tools translate automatically. The earlier analysis in this tree did not:
+`run/helper_453ef0.c` and friends were produced by passing a file vaddr
+(`0x453ef0`) straight to Ghidra, which decompiled `libcuda+0x353ef0` — a
+*different function*. Every conclusion in the old write-up that depended on
+decompiled code was therefore about the wrong function. The runtime
+observations (gdb register traces, live table values, ioctl traces) remain
+valid and are used below.
+
+Correct addresses for the old write-up's landmarks:
+
+| What the old note called | File vaddr | Ghidra address | Correct identity |
+|---|---|---|---|
+| `FUN_00453ef0` (helper) | `0x453ef0` | `0x553ef0` | engine-context registration helper (per engine) |
+| crash site | `0x453f90` | `0x553f90` | continuation inside the helper after the hook call |
+| `FUN_0045a670` (walker) | `0x45a670` | `0x55a670` | device engine/context walker (calls the helper twice) |
+| `FUN_004676f0` (hook) | `0x4676f0` | `0x5676f0` | hook reached through `[table+0x40]+0x530` |
+| hook callee | `0x481360` | `0x581360` | engine context create/register (issues `NV_ESC_RM_ALLOC`) |
+| `FUN_003a9c20` | `0x3a9c20` | `0x4a9c20` | handle-pool allocation called before the hook |
+
+## 2. The public API stubs
+
+Every exported `cu*` is a 24–30 byte stub:
+
+1. compare the global initialization word against the "destroyed" sentinel
+   `0x321CBA00`; if equal return `4` (`CUDA_ERROR_DEINITIALIZED`);
+2. otherwise `jmp [slot]` — an indirect jump through the interposition table.
+
+Initial slot values (from relocations) give the generated wrapper addresses:
+
+| Export (file vaddr) | Wrapper target |
+|---|---|
+| `cuInit` `0x310060` | `0x31b7f0` |
+| `cuCtxCreate_v2` `0x310360` | `0x3206d0` |
+| `cuModuleLoadData` `0x3106e0` | `0x324ed0` |
+| `cuMemAlloc_v2` `0x310b20` | `0x329830` |
+| `cuStreamSynchronize` `0x36f3c0` | `0x380b40` |
+| `cuLaunchKernel` `0x36f420` | `0x3810f0` |
+
+This is the mechanism described in `docs/01` §2.2/§2.3: the wrapper does the
+tracing/tool callbacks and then calls the hand-written implementation. When
+chasing a call, follow the slot; when a tool interposes, the slot changes.
+
+## 3. The `cuCtxCreate` path: crash solved, engine topology blocking
+
+The 2026-09 crash was deterministic: on a fresh guest `cuCtxCreate` segfaulted
+at `libcuda+0x453f90`. That is **solved** (§3.5). The call now returns
+`rc=999` (`CUDA_ERROR_UNKNOWN`), and the active blocker is engine topology —
+the emulated engine list is missing the async copy engine (§3.6).
+
+Fault facts from the crash period (live traces, corrected decompilation):
+
+```
+guest dmesg : segfault at ffffffffffffffc8  (NULL-0x38)
+              ip in libcuda.so.570.144, mapping base 0x735cfcb79000
+              file vaddr = 0x453f90
+gdb         : rbp=0, rax=0, rsi=0xc030462b (_IOWR('F',0x2b,48) = NV_ESC_RM_ALLOC)
+```
+
+### 3.1 The faulting instruction
+
+```
+libcuda+0x453ef0  push rbp; mov rbp,rsp; ...            <- helper entry
+...
+libcuda+0x453f5f  call qword ptr [rax+0x530]            <- engine hook
+libcuda+0x453f65  mov r12d,eax
+libcuda+0x453f68  test eax,eax
+libcuda+0x453f6a  je  0x453f90
+...
+libcuda+0x453f90  mov rax,[rbp-0x38]                    <- FAULT (rbp==0)
+```
+
+The helper saved a table pointer at `[rbp-0x38]` in its prologue and reloads
+it after the hook returns. With `rbp==0` this reads `0xffffffffffffffc8`,
+exactly the fault address.
+
+### 3.2 What the helper does (behaviour)
+
+`libcuda+0x453ef0(device, engine_desc)`:
+
+1. `idx = engine_desc->[0]`, with `1` normalized to `0` (the GR engine);
+2. read `table = *(*(*device + 8) + 0x40)` and save it;
+3. allocate a handle from the global handle pool (`libcuda+0x3a9c20`,
+   a bitmap allocator) → `engine_id`;
+4. if the lookup succeeded, call hook `[*(table2+0x40)+0x530]`
+   (`libcuda+0x4676f0`) with `(global, *(device+0x9488), engine_id,
+   engine_desc)` — this is the engine create/register call that ultimately
+   issues the `NV_ESC_RM_ALLOC` seen in `rsi` at the fault;
+5. if the hook returns 0, reload the saved table, read
+   `class = *(table+0xd40)`, and call a second hook
+   `[*(table2+0x40)+0x250]` with `(global, *(device+0x9488), engine_id,
+   class, device + idx*4 + 0x68)`;
+6. on success record `device[idx*4 + 0x50] = engine_id` and
+   `device[0x9488][idx*8 + 0x88] = engine_desc`;
+7. on failure, clean up (`libcuda+0x453cc0`, `libcuda+0x3a2f50`).
+
+So the helper is "register one engine with the driver": look up the engine,
+create its context (hook 1), then bind/allocate the engine object (hook 2),
+then publish the result into per-engine arrays.
+
+### 3.3 The engine walker
+
+`libcuda+0x45a670` walks the device's engine records and calls the helper at
+`libcuda+0x45afc1` and `libcuda+0x45b02b` (per engine, and once more for a
+secondary engine at `[device-table + 0x75c]`). The descriptor is selected
+from 20-byte (0x14) records:
+
+```
+engine index 0        -> table + 0x2fec
+engine index 1..0xa   -> table + 0x300c + (index-1)*0x14
+engine index 0xb      -> table + 0x350c
+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.
+
+### 3.4 The engine create/register function
+
+`libcuda+0x481360` (called by the hook) is the function that actually asks RM
+for an engine context. Behaviour (condensed):
+
+* class `0x80` (device): find an existing context for the class, create one,
+  then issue an RM alloc (`NV_ESC_RM_ALLOC`, request `0xc030462b`, 48-byte
+  parameter block) to bind it;
+* classes `0x5`/`0x79`: use a per-class object from a global table, with a
+  bounds-checked index;
+* class `0x2080` (subdevice): iterate a global array of 32 engine
+  descriptors, matching `[context+0x28]`, and for a match build a 0x28-byte
+  RM control (`0x202`) and alloc;
+* default: generic RM alloc via `libcuda+0x480d00`.
+
+Failure returns are small status codes (`0x1f`, `0x22`, `0x26`, `0x36`,
+`0x3d`, `0x56`, `0x59`, …) which the hook then routes into the driver's error
+reporting (`libcuda+0x387bf0`).
+
+### 3.5 Resolution: the alloc-params copy-back
+
+The engine-registration hook reaches `libcuda+0x481360`, which for the
+`TURING_COMPUTE_A` class (`0xc5c0`) issues `NV_ESC_RM_ALLOC` (ioctl `0x2b`,
+NVOS64 form) with `pAllocParms` pointing at the caller's **own 8-byte stack
+local** and `paramsSize = 0`.
+
+The guest driver does not honour the caller's `paramsSize`: for class `0xc5c0`
+its resource table declares 16-byte class params (`NV_GR_ALLOCATION_PARAMETERS`),
+so it
+
+1. reads 16 bytes starting at that 8-byte local — the upper 8 bytes are the
+   hook's saved `rbp` — and sends them as the alloc parameters, and
+2. on the reply, copies 16 bytes of the reply's params back over the same
+   buffer.
+
+With the old generic (zeroed) alloc reply, step 2 zeroed the hook's saved
+`rbp`. The hook's `leave` then loaded `rbp = 0`, and `mov rax,[rbp-0x38]` at
+`libcuda+0x453f90` faulted at `0xffffffffffffffc8`.
+
+Live gdb intervals confirmed the writer: the hook's saved-`rbp` slot was intact
+at the pre-ioctl site (`libcuda+0x481448`) and zeroed at the post-ioctl site
+(`libcuda+0x48146a`), and hardware watchpoints never saw an in-process store —
+the store is the kernel's `copy_to_user` inside the ioctl.
+
+**Fix (host module, live-verified):** the alloc reply now echoes the request's
+params back, with only `status` forced to `NV_OK`. The driver's class-size
+copy-back then writes back exactly the bytes it read — including the caller's
+saved-`rbp` slot — and the `libcuda+0x453f90` fault disappears. This is also
+what the real GSP-RM does for these in/out params: the host runs the same
+libcuda code and survives for the same reason. It is the "coherent params"
+rung of the reply-fidelity ladder (`docs/04` §6).
+
+**Historical: why `rbp` was zero (solved).** Earlier revisions of this document
+listed three candidate writers of the hook's saved-`rbp` slot; the writer was
+the alloc reply's class-size copy-back described above.
+
+### 3.6 Current blocker (2026-09-23): the missing async copy engine
+
+With the copy-back fixed, the walker completes but `cuCtxCreate` returns
+`rc=999` (`CUDA_ERROR_UNKNOWN`); `cuMemAlloc` and later calls are not reached.
+Measured with the guest ioctl trace plus a host A/B oracle (the same libcuda
+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
+  `{class, engDesc}` pair table pairs OBJCE instances `0..7` with class
+  `0xcab5` and instances `0x0b..0x13` (including `0x0e` = 14) with class
+  `0xc9b5`.
+* The driver filters the client engine table through its per-(class, instance)
+  class database (`classDB`) before the engine walker (§3.3) sees it. With no
+  classDB entry for descriptor `0x793ceb0e` the engine is silently skipped and
+  libcuda never learns it exists.
+* The walker then registers a secondary engine from a record that is zeroed in
+  the guest (host: `t=2 cls=0xc5b5`; guest: `t=0 cls=0`, dumped live with gdb).
+  The zero class makes libcuda issue `NV_ESC_RM_ALLOC` for class `0`; the
+  ioctl escape path rewrites root classes to `NV01_ROOT_CLIENT` (`0x41`), so
+  the trace shows the single failing call `class=0x41 st=0x19` immediately
+  after the compute-object alloc.
+* libcuda rolls back — it frees the channel and the compute object — and
+  returns 999. The guest kernel still oopses in
+  `uvm_pmm_devmem_page_to_va_space` while tearing down the failed context.
+
+A first emulator-side attempt (advertise the `COPY14` capability bit in the
+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.
+* **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).
+* 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
+  immediately consumes the reply's params and (for other classes) engine ids.
+
+## 4. Navigating the decompiled corpus
+
+```
+# by address (file vaddr):
+grep -n "libcuda+0x45a670" 570.144/run/analysis/decompiled/buckets/*.c
+
+# by Ghidra name:
+awk -F'\t' '$4=="FUN_0055a670"' 570.144/run/analysis/decompiled/decompiled_index.tsv
+
+# callers of a function (program addresses):
+grep -P "^0x0055a670\t" 570.144/run/analysis/elf-index/callgraph.tsv
+
+# exports and their sizes:
+570.144/run/analysis/elf-index/exports.tsv
+
+# strings around an address (error messages, table names):
+awk -F'\t' '$1>="0x..." && $1<="0x..."' 570.144/run/analysis/elf-index/strings.tsv
+```
+
+Every bucket header carries `addr=0x<program> file_addr=libcuda+0x<file>` so
+citations in docs and gdb sessions use the same numbers.
+
+## 5. Known limits
+
+* Names are `FUN_<addr>` unless recovered; `570.144/annotations/names.tsv`
+  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`.
+* 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.
+* Embedded device code (cubins) is present in `.rodata` and is *not* host
+  code; the corpus only covers `.text`.
diff --git a/userspace/cuda/docs/06-tooling.md b/userspace/cuda/docs/06-tooling.md
new file mode 100644
index 0000000..85d1c8c
--- /dev/null
+++ b/userspace/cuda/docs/06-tooling.md
@@ -0,0 +1,165 @@
+# 06 — Tooling: how to analyse a CUDA userspace driver
+
+The toolchain in `tools/` is version-agnostic. Point it at any
+`libcuda.so.<version>` and it produces the same artifacts as for 570.144. All
+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
+├── 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
+├── ghidra_batch_decompile.sh       parallel, resumable full decompilation
+├── ghidra_decompile_at.sh          decompile one function by file vaddr
+├── ghidra_index_merge.sh           merge batch indexes
+└── ghidra/
+    ├── ConfigureAnalysis.java      analyzer tuning for big stripped binaries
+    ├── ExportFunctionIndex.java    index exporter
+    ├── BatchDecompile.java         chunk decompiler
+    ├── DecompileAt.java            single function (file-vaddr aware)
+    ├── ApplyNameMap.java           apply/transfer a name map
+    └── PrintProgramInfo.java       address convention + stats
+```
+
+## 1. Fast facts without Ghidra — `cudatool.py`
+
+```
+tools/cudatool.py info      <lib>          # build id, sections, exports, ABI note
+tools/cudatool.py exports   <lib>          # addr size name (file vaddrs)
+tools/cudatool.py functions <lib>          # function ranges from .eh_frame_hdr
+tools/cudatool.py nearest   <lib> <addr>   # function start + nearest export
+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
+```
+
+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.
+
+### Address conventions
+
+* `cudatool.py` prints **file virtual addresses** — the `libcuda+0x...` used
+  everywhere in this directory.
+* Ghidra may load the same binary at a non-zero image base. For 570.144 it is
+  `0x100000`; check with `PrintProgramInfo.java`. The Ghidra wrappers here
+  translate file vaddr → program address automatically; raw Ghidra scripts do
+  not.
+
+## 2. Full analysis — Ghidra
+
+```
+# 1. import + analyze (keep the project; can take minutes to hours)
+tools/ghidra_full_analysis.sh <lib> <project_dir> <project_name>
+
+# 2. export the index (fast, read-only)
+tools/ghidra_export_index.sh <project_dir> <project_name> <program> <outdir>
+
+# 3. decompile everything (parallel, resumable; safe to re-run)
+tools/ghidra_batch_decompile.sh <project_dir> <project_name> <program> <outdir> \
+    <start_hex> <end_hex> [chunk_hex] [jobs] [timeout_s]
+tools/ghidra_index_merge.sh <outdir>
+```
+
+Notes:
+
+* `ConfigureAnalysis.java` disables the two expensive identifier analyzers
+  (parameter-ID, function-ID) and keeps everything needed for a stripped
+  binary. Full analysis of 71 MB then takes minutes, not hours.
+* `ghidra_batch_decompile.sh` gives **every chunk its own project copy** under
+  `<outdir>/wk/` because Ghidra takes an exclusive project lock even in
+  read-only mode; completed chunks are skipped on re-run. Delete `wk/` to
+  reclaim disk.
+* The program name passed to Ghidra is the binary's file name
+  (`libcuda.so.570.144`).
+* Do **not** pass `-analysisTimeoutPerFile 0`: in this Ghidra build that means
+  "abort analysis immediately" (the original per-version helper in
+  `570.144/scripts/` did this and produced an analysed-looking project with no
+  analysis at all).
+
+## 2b. Annotating the corpus
+
+```
+tools/annotate_corpus.py <outdir> <indexdir> <names.tsv> <annotated_outdir>
+```
+
+Rewrites every function header with the recovered name, file address, size and
+direct callers/callees (named where the map has them), and substitutes
+recovered names for `FUN_<addr>` references in the bodies. Runs in seconds for
+30k functions; the output is gitignored like the raw corpus.
+
+## 2c. Walking the callgraph without Ghidra
+
+```
+tools/callwalk.py <indexdir>/callgraph.tsv callees 0x31b7f0 2 10
+tools/callwalk.py <indexdir>/callgraph.tsv callers 0x453ef0
+tools/callwalk.py <indexdir>/callgraph.tsv chain 0x31b7f0 0x2cb940
+```
+
+Addresses are file vaddrs; `names.tsv` next to the callgraph is used for
+labels.
+
+## 3. Looking things up
+
+```
+# one function by file vaddr
+tools/ghidra_decompile_at.sh <project_dir> <project_name> <program> 0x45a670 out.c
+
+# corpus search by file address (headers are annotated)
+grep -n "libcuda+0x45a670" <outdir>/buckets/*.c
+
+# index lookups
+awk -F'\t' '$2=="0x45a670"' <outdir>/decompiled_index.tsv
+grep -P "^0x0055a670\t" <indexdir>/callgraph.tsv     # direct callers/callees
+
+# exports/strings
+grep cuLaunchKernel <indexdir>/exports.tsv
+```
+
+## 4. Version deltas — `cudatool.py delta`
+
+```
+tools/cudatool.py delta <libA> <libB>
+```
+
+Reports, at the interface level: added/removed/moved/resized exports, import
+count changes, `.text` growth, build-id and the CUDA/ABI note. This is the
+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).
+
+## 5. Name transfer — `ApplyNameMap.java`
+
+```
+tools/ghidra_decompile_at.sh-free invocation:
+  analyzeHeadless <project_dir> <project_name> -process <program> -noanalysis \
+      -scriptPath tools/ghidra -postScript ApplyNameMap.java <names.tsv> [--dry-run]
+```
+
+The map is `file_vaddr<TAB>name` per line (comments allowed). Name maps live
+in version directories (`570.144/annotations/names.tsv`) and are tracked, so
+naming work accumulates: apply the previous version's map to a new binary,
+then fix the misses. `--dry-run` reports how many names would land.
+
+## 6. Writing new Ghidra scripts
+
+Follow the existing pattern: a `GhidraScript` subclass in `tools/ghidra/`,
+args via `getScriptArgs()`, and file-vaddr ↔ program-address translation via
+`currentProgram.getImageBase()` (see `DecompileAt.java`). Keep outputs under
+the version's gitignored `run/` tree unless the artifact is our own work
+product (name maps, indexes of our findings), which may be tracked.
+
+## 7. Cost guide (measured on 570.144, 16 cores)
+
+| Step | Time | Disk |
+|---|---|---|
+| import alone | ~1 min | 324 MB project |
+| full analysis | ~10–15 min | same project |
+| index export | ~1–2 min | ~25 MB TSV |
+| full decompilation (5 workers) | ~20–30 min | ~3 GB C + 2 GB project copies |
+
+Notes for a rerun: set `GHIDRA_HEADLESS_MAXMEM` (default 12G for analysis,
+4G per batch job) and keep total batch jobs ≤ 5–6 on a 32 GB machine.
diff --git a/userspace/cuda/docs/07-version-upgrade-playbook.md b/userspace/cuda/docs/07-version-upgrade-playbook.md
new file mode 100644
index 0000000..9e53514
--- /dev/null
+++ b/userspace/cuda/docs/07-version-upgrade-playbook.md
@@ -0,0 +1,98 @@
+# 07 — Version-upgrade playbook: re-running the CUDA analysis
+
+When the guest driver, CUDA runtime, GSP firmware or GPU generation changes,
+this is the procedure that refreshes the analysis with minimal effort and
+keeps the name map and interface knowledge accumulating.
+
+## 0. Before you start
+
+* You need the new `libcuda.so.<version>` (from the guest with
+  `570.144/scripts/fetch_from_guest.sh`, or from the driver package with
+  `--extract-only`).
+* Work in a new version directory `userspace/cuda/<version>/` — never
+  overwrite an analysed version. Copy the structure from `570.144/`
+  (`scripts/`, `ghidra/`, `annotations/`) and create `run/`.
+* Update `WORKING-HERE.md` so parallel agents know which tree is being
+  analysed.
+
+## 1. Cheap triage (minutes)
+
+```
+tools/cudatool.py info    <lib>
+tools/cudatool.py delta   <previous_lib> <lib> > run/delta.txt
+tools/cudatool.py crash   <lib> <map_base> <ip>     # if a guest crash is known
+```
+
+Read `delta.txt`:
+
+* **no export moves, small text growth** → patch release: most names and
+  offsets from the previous version transfer; spot-check the subsystems
+  touched by the release notes.
+* **new exports** → feature release: identify the new API surface, map it to
+  engine functions via decompilation of the new stubs.
+* **many moves** → reshuffle (toolchain/linker change): plan a full re-analysis
+  and re-apply the name map, expecting to fix a fraction.
+
+## 2. Full analysis (tens of minutes)
+
+```
+tools/ghidra_full_analysis.sh <lib> <version_dir>/run/ghidra-full cuda-full
+tools/ghidra_export_index.sh  <version_dir>/run/ghidra-full cuda-full <program> <version_dir>/run/analysis/elf-index
+tools/ghidra_batch_decompile.sh <version_dir>/run/ghidra-full cuda-full <program> <version_dir>/run/analysis/decompiled \
+    <text_start> <text_end> 0x100000 5 120
+tools/ghidra_index_merge.sh <version_dir>/run/analysis/decompiled
+```
+
+Get `<text_start>`/`<text_end>` from `cudatool.py info`. Confirm the image base
+with `PrintProgramInfo.java` before citing addresses.
+
+## 3. Transfer names and knowledge
+
+```
+cp <previous>/annotations/names.tsv <new>/annotations/names.tsv
+# translate addresses: the map is file-vaddr keyed, so for unchanged
+# functions the offsets are usually identical; where the delta shows moves,
+# adjust from the decompiled corpus.
+analyzeHeadless <proj> cuda-full -process <program> -noanalysis \
+    -scriptPath tools/ghidra -postScript ApplyNameMap.java <new>/annotations/names.tsv --dry-run
+# then without --dry-run
+```
+
+Also transfer the *subsystem index* (`annotations/subsystems.tsv`: address,
+our name, subsystem, evidence) and re-validate the entries that the delta
+marked as moved.
+
+## 4. Validate against the guest
+
+1. Load the stack and run the guest smoke test (`tools/test/cuda_drv_test.c`
+   compiled as `/tmp/cdt`), with the ioctl tracer
+   (`tools/guest-diag/ioctl_trace.c`) attached.
+2. Compare the observed ioctl/control sequence with the previous version's
+   `docs/05` sequence; new or reordered controls are the interesting deltas.
+3. Run the host/guest A/B harness (`tools/cuda-ab/ab.sh all`) for a numerical
+   correctness check.
+4. If the guest fails, use the crash decoder + corpus lookup (§1, §2) to get
+   from a dmesg `ip` to the exact function and line in the decompiled corpus,
+   then to the control reply the driver consumed.
+
+## 5. Update the documents
+
+* `docs/05-<version>-decompilation.md` — new version-specific map and crash
+  notes (copy the structure from `docs/05-570.144-decompilation.md`).
+* `annotations/names.tsv` and `annotations/subsystems.tsv` — the deltas.
+* `docs/04-virtualization-and-emulation.md` — only if the *interface* changed
+  (new controls, new capability bits, new object classes).
+* `README.md` status table.
+
+## 6. What not to do
+
+* Do not patch the guest driver or add `LD_PRELOAD` shims to make things work;
+  that breaks the anti-fragility rule (`docs/04` §3). If a reply must change,
+  change the emulation.
+* Do not commit decompiled output, extracted binaries or Ghidra projects; they
+  are gitignored for a reason (vendor-derived, large).
+* Do not name non-public source locations in tracked docs; describe behaviour
+  and cite `libcuda+0x...` offsets.
+* Do not run a second analysis in the same Ghidra project while another
+  process has it open (exclusive lock, even read-only); use the per-chunk
+  project copies in `ghidra_batch_decompile.sh`.
diff --git a/userspace/cuda/docs/08-virtualization-approaches.md b/userspace/cuda/docs/08-virtualization-approaches.md
new file mode 100644
index 0000000..3523ad1
--- /dev/null
+++ b/userspace/cuda/docs/08-virtualization-approaches.md
@@ -0,0 +1,67 @@
+# 08 — Virtualization approaches for CUDA: options, trade-offs, and where we are
+
+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`,
+`docs/04-implementation/02-practical-challenges-and-adversarial-interop.md`).
+
+## 1. The options
+
+| # | Approach | What runs where | Benefits | Drawbacks / what we must own |
+|---|---|---|---|---|
+| **0** | **Full software emulation** — execute every command in the host emulator (memsets, copies, semaphores, and in principle grids) | Everything in the host module | No physical GPU needed at all; perfectly isolated; simplest correctness story; portable | Orders of magnitude too slow for compute; only viable for state/memory operations; never the whole answer |
+| **A** | **Host RM/GSP proxy** — guest push → translate → host driver's RM API → real channel → GPU | Control and work both through the host driver | Reuses a mature driver's knowledge (page tables, scheduling, faults); fastest path to *correct*; no direct hardware ownership | Extra hop (latency/throughput cost); depends on host driver interfaces staying usable; sharing the host GPU with the host driver needs a policy |
+| **B** | **Direct engine submission** — guest push → translate → our own channel/push construction → GPU | Control emulated, work submitted by us | Lowest latency, no host-driver dependency, full control of scheduling/isolation | We must own page tables, channel/runlist setup, cache invalidation, completion, fault handling — the largest engineering burden |
+| **C** | **Hybrid** — management via A, hot path via B | Split by operation class | Pragmatic: correctness now, performance where it costs | Two data paths to keep consistent; the boundary must be an interface from day one |
+| **D** | **Hardware partitioning (SR-IOV/VF, MIG)** | Guest on its own partition/VF | Strong isolation, near-native performance | Not available on consumer GPUs; still needs a device model for the VF; complementary, not a replacement for A/B/C |
+| **E** | **Guest-assisted (patch driver / shim / API interception)** | Guest modified | Easy performance wins in the short term | Breaks the "stock guest" rule, the anti-fragility rules and the legal/clean-room stance — **rejected** |
+
+**Fast paths are orthogonal to the choice:** pure state operations (semaphores,
+fences, some memops) should be satisfied directly in the emulation regardless
+of A/B/C — they are microseconds in hardware and would lose by round-tripping.
+
+## 2. Design implications
+
+* Where the boundary sits decides what we must own: A keeps ownership in the
+  host driver, B moves it into Open-IOV. Either way the **interface** the
+  guest sees is identical — that is what makes the choice reversible.
+* Isolation and correctness do not depend on the path, but **the failure
+  modes do**: A fails as host-driver errors and sharing conflicts, B fails as
+  hardware faults and page-table/scheduling bugs we introduced.
+* Performance targets (launch overhead, copy bandwidth, completion latency)
+  can only be met by forwarding *real work*; software execution is for
+  state/memory operations and as a correctness bridge, not the data plane.
+* The work path must therefore be abstracted behind one interface now, so A
+  can be swapped for B (or C) without touching the control plane.
+
+## 3. What Open-IOV is pursuing
+
+**Stated architecture:** *emulate the interface, pass through the work, never
+patch anyone* — a mediated-passthrough vGPU where PCI config/BARs, the GSP boot
+handshake, the GSP-RM protocol and interrupt routing are emulated in
+`open-iov-vfio.ko`, while real work is meant to execute on the physical GPU.
+The project's own decision procedure is explicit: **prototype A first for
+correctness, measure, then move the hot path to B if the budgets demand it**
+(option C in practice), with the work path abstracted from day one.
+
+**Current reality (2026-09, honest):** the *control plane* follows the
+architecture — GSP-RM emulation, RPC handling, BAR/PCI/MSI-X emulation and
+object/control replies are implemented, and the guest's stock driver +
+`libcuda` now get through `cuInit`, `cuDeviceGetCount`/`cuDeviceGet` and into
+`cuCtxCreate`. The old `cuCtxCreate` segfault is solved (the alloc reply must
+echo the request's params — `docs/05` §3.5); context creation is blocked one
+step later by engine topology: the class-filtered engine list is missing the
+async copy engine, so the walker reads a zeroed secondary record and the call
+returns `rc=999` (`docs/05` §3.6). The *data plane* is at the bridge stage: the
+host module decodes the guest's push buffers as method streams and **executes
+the small work it understands itself** (memsets, copies, semaphore releases)
+instead of forwarding it to the physical GPU. No CUDA kernel has been
+forwarded yet. So: yes, we are pursuing the intended architecture, but we are
+in the "emulate + software-execute state/memory ops" phase of it, with the
+A→B/C forwarding decision still ahead.
+
+**Risks to guard now:** (1) software execution silently becoming the data
+plane instead of a bridge; (2) the work-path interface being skipped because
+the direct path "works for now"; (3) tuning the emulation to one driver
+version instead of the interface. The anti-fragility, performance and security
+rules in `docs/04` exist to prevent exactly this drift.
diff --git a/userspace/cuda/tools/annotate_corpus.py b/userspace/cuda/tools/annotate_corpus.py
new file mode 100755
index 0000000..972f923
--- /dev/null
+++ b/userspace/cuda/tools/annotate_corpus.py
@@ -0,0 +1,155 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# annotate_corpus.py - turn the raw decompiled corpus into an annotated,
+# human-navigable corpus.
+#
+# For every function in the decompiled buckets it rewrites the header comment
+# with:
+#   * the recovered name (from a names.tsv map) or the Ghidra FUN_ name;
+#   * file virtual address and size;
+#   * direct callers and callees (from callgraph.tsv), labelled with names;
+#   * a "seen in" marker for functions referenced by the docs (optional).
+# It also substitutes recovered names for FUN_<addr> references inside the
+# function bodies, so a renamed function is readable everywhere.
+#
+# Inputs (produced by the toolchain):
+#   <outdir>/buckets/bucket_*.c          raw decompiled C
+#   <outdir>/decompiled_index.tsv        addr file_addr size name file line status
+#   <indexdir>/callgraph.tsv             caller callee caller_file callee_file
+#   <names.tsv>                          file_vaddr<TAB>name
+#
+# Usage:
+#   annotate_corpus.py <outdir> <indexdir> <names.tsv> <annotated_outdir>
+#
+# The output directory is gitignored like the raw corpus; this tool is tracked
+# and reusable. Addresses are file virtual addresses in all inputs; the corpus
+# uses Ghidra program addresses, so the tool translates via the index's
+# addr/file_addr columns (no hard-coded image base).
+
+import os
+import re
+import sys
+from collections import defaultdict
+
+
+def load_index(path):
+    rows = {}
+    with open(path) as f:
+        header = f.readline()
+        for line in f:
+            p = line.rstrip("\n").split("\t")
+            if len(p) < 6:
+                continue
+            prog = int(p[0], 16)
+            rows[prog] = dict(file=int(p[1], 16), size=int(p[2]),
+                              name=p[3], file_name=p[4], line=int(p[5]))
+    return rows
+
+
+def load_callgraph(path):
+    callees = defaultdict(list)
+    callers = defaultdict(list)
+    with open(path) as f:
+        f.readline()
+        for line in f:
+            p = line.rstrip("\n").split("\t")
+            if len(p) < 4:
+                continue
+            cf, kf = int(p[2], 16), int(p[3], 16)
+            callees[cf].append(kf)
+            callers[kf].append(cf)
+    return callees, callers
+
+
+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 label(file_addr, names, index):
+    name = names.get(file_addr)
+    if name:
+        return f"{name} (libcuda+0x{file_addr:x})"
+    return f"libcuda+0x{file_addr:x}"
+
+
+def main():
+    if len(sys.argv) < 5:
+        print(__doc__)
+        sys.exit(2)
+    outdir, indexdir, namemap, out = sys.argv[1:5]
+    index = load_index(os.path.join(outdir, "decompiled_index.tsv"))
+    callees, callers = load_callgraph(os.path.join(indexdir, "callgraph.tsv"))
+    names = load_names(namemap)
+    # program address -> recovered name, for body substitution
+    prog_names = {p: names[r["file"]] for p, r in index.items() if r["file"] in names}
+
+    os.makedirs(os.path.join(out, "buckets"), exist_ok=True)
+    n = 0
+    for bucket in sorted(os.listdir(os.path.join(outdir, "buckets"))):
+        if not bucket.endswith(".c"):
+            continue
+        src = os.path.join(outdir, "buckets", bucket)
+        dst = os.path.join(out, "buckets", bucket)
+        with open(src) as f:
+            text = f.read()
+
+        def annotate(m):
+            nonlocal n
+            prog = int(m.group("addr"), 16)
+            file_addr = int(m.group("file"), 16)
+            size = int(m.group("size"))
+            rec = index.get(prog, {})
+            name = names.get(file_addr, rec.get("name", f"FUN_{prog:08x}"))
+            cl = sorted(set(callees.get(file_addr, [])))
+            cr = sorted(set(callers.get(file_addr, [])))
+            lines = [
+                "/* ------------------------------------------------------------------",
+                f" * {name}",
+                f" * file_addr=libcuda+0x{file_addr:x} size={size}",
+            ]
+            if cr:
+                shown = [label(a, names, index) for a in cr[:8]]
+                more = f" (+{len(cr)-8} more)" if len(cr) > 8 else ""
+                lines.append(" * called by: " + "; ".join(shown) + more)
+            if cl:
+                shown = [label(a, names, index) for a in cl[:8]]
+                more = f" (+{len(cl)-8} more)" if len(cl) > 8 else ""
+                lines.append(" * calls:     " + "; ".join(shown) + more)
+            lines.append(" *----------------------------------------------------------------- */")
+            n += 1
+            return "\n" + "\n".join(lines) + "\n"
+
+        text = re.sub(
+            r"\n/\* -+\n \* (?P<name>[^\n]*)\n \* addr=(?P<addr>0x[0-9a-f]+) "
+            r"file_addr=libcuda\+(?P<file>0x[0-9a-f]+) size=(?P<size>\d+)\n"
+            r" \*-+ \*/",
+            annotate, text)
+
+        if prog_names:
+            def sub(m):
+                a = int(m.group(1), 16)
+                return prog_names.get(a, m.group(0))
+            text = re.sub(r"\bFUN_([0-9a-f]{8})\b", sub, text)
+
+        with open(dst, "w") as f:
+            f.write(text)
+
+    print(f"[annotate_corpus] annotated {n} functions -> {out}/buckets")
+    print(f"[annotate_corpus] names applied: {len(names)}")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/userspace/cuda/tools/callwalk.py b/userspace/cuda/tools/callwalk.py
new file mode 100755
index 0000000..15feb45
--- /dev/null
+++ b/userspace/cuda/tools/callwalk.py
@@ -0,0 +1,133 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# callwalk.py - walk the Open-IOV callgraph index and print callee trees /
+# caller trees around a function. This is the quick way to go from an export
+# wrapper to the internal implementation chain without opening Ghidra.
+#
+# Usage:
+#   callwalk.py <callgraph.tsv> callees <hex_file_vaddr> [depth] [max_per_node]
+#   callwalk.py <callgraph.tsv> callers <hex_file_vaddr> [depth] [max_per_node]
+#   callwalk.py <callgraph.tsv> chain   <hex_from> <hex_to>   # shortest path
+#
+# Addresses are FILE virtual addresses (libcuda+0x...); the callgraph columns
+# `caller_file`/`callee_file` are used, so Ghidra image-base offsets do not
+# matter. `names.tsv` (file_vaddr<TAB>name) is picked up next to the callgraph
+# when present, and used to label nodes.
+#
+# SPDX-License-Identifier: BSD-3-Clause
+
+import sys
+from collections import defaultdict, deque
+
+
+def load_graph(path):
+    callees = defaultdict(list)
+    callers = defaultdict(list)
+    with open(path) as f:
+        header = f.readline()
+        for line in f:
+            parts = line.rstrip("\n").split("\t")
+            if len(parts) < 4:
+                continue
+            cf, kf = int(parts[2], 16), int(parts[3], 16)
+            callees[cf].append(kf)
+            callers[kf].append(cf)
+    return callees, callers
+
+
+def load_names(path):
+    names = {}
+    try:
+        with open(path) as f:
+            for line in f:
+                line = line.strip()
+                if not line or line.startswith("#"):
+                    continue
+                parts = line.split("\t")
+                if len(parts) >= 2:
+                    try:
+                        names[int(parts[0], 16)] = parts[1]
+                    except ValueError:
+                        pass
+    except FileNotFoundError:
+        pass
+    return names
+
+
+def label(a, names):
+    return f"0x{a:08x}({names[a]})" if a in names else f"0x{a:08x}"
+
+
+def walk(adj, start, depth, limit):
+    seen = {start}
+    frontier = [start]
+    for d in range(depth):
+        nxt = []
+        for node in frontier:
+            kids = sorted(set(adj.get(node, [])))
+            if not kids:
+                continue
+            print("  " * (d + 1) + " | ".join(label(k, NAMES) for k in kids[:limit])
+                  + (" ..." if len(kids) > limit else ""))
+            for k in kids:
+                if k not in seen:
+                    seen.add(k)
+                    nxt.append(k)
+        frontier = nxt
+        if not frontier:
+            break
+
+
+def shortest_path(callees, start, goal):
+    if start == goal:
+        return [start]
+    prev = {start: None}
+    q = deque([start])
+    while q:
+        node = q.popleft()
+        for k in callees.get(node, []):
+            if k in prev:
+                continue
+            prev[k] = node
+            if k == goal:
+                path = [goal]
+                while prev[path[-1]] is not None:
+                    path.append(prev[path[-1]])
+                return list(reversed(path))
+            q.append(k)
+    return None
+
+
+def main():
+    if len(sys.argv) < 4:
+        print(__doc__)
+        sys.exit(2)
+    graph = sys.argv[1]
+    mode = sys.argv[2]
+    addr = int(sys.argv[3], 16)
+    depth = int(sys.argv[4]) if len(sys.argv) > 4 else 4
+    limit = int(sys.argv[5]) if len(sys.argv) > 5 else 8
+
+    global NAMES
+    callees, callers = load_graph(graph)
+    NAMES = load_names(graph.rsplit("/", 1)[0] + "/names.tsv")
+
+    if mode == "chain":
+        goal = int(sys.argv[4], 16)
+        path = shortest_path(callees, addr, goal)
+        if not path:
+            print("no direct path found")
+        else:
+            for a in path:
+                print(label(a, NAMES))
+        return
+
+    print(f"{mode} of {label(addr, NAMES)}:")
+    walk(callees if mode == "callees" else callers, addr, depth, limit)
+
+
+NAMES = {}
+
+if __name__ == "__main__":
+    main()
diff --git a/userspace/cuda/tools/cudatool.py b/userspace/cuda/tools/cudatool.py
new file mode 100755
index 0000000..2e4c1c6
--- /dev/null
+++ b/userspace/cuda/tools/cudatool.py
@@ -0,0 +1,515 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# cudatool.py - Open-IOV CUDA userspace driver analysis toolkit.
+#
+# A dependency-free (stdlib-only) ELF toolkit for the NVIDIA CUDA userspace
+# driver, usable on every driver version without a Ghidra project. It gives
+# the fast facts that do not require decompilation: exports, function ranges
+# derived from the exception-handling tables, the CUDA forward-compatibility
+# note, and a dmesg/crash-address decoder. It is also the version-delta tool:
+# `delta` compares two driver versions and reports exactly what changed at
+# the interface level, which is what keeps the Open-IOV side anti-fragile
+# across driver upgrades.
+#
+# Commands
+#   info <lib>                     ELF/version summary
+#   exports <lib> [--json]         exported symbols (file vaddrs)
+#   functions <lib> [--json]       function ranges from .eh_frame_hdr
+#   nearest <lib> <hex_addr>       nearest preceding function/export for a vaddr
+#   crash <lib> <map_base> <ip>    dmesg mapping base + ip -> file vaddr + symbol
+#   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
+#   note <lib>                     CUDA forward-compatibility note (JSON)
+#
+# Address convention: every address printed by this tool is a FILE VIRTUAL
+# ADDRESS, i.e. the `libcuda+0x...` numbers used in the project docs. Ghidra
+# may load the same file at a non-zero image base (0x100000 for these builds);
+# use tools/ghidra_decompile_at.sh which translates file vaddrs automatically.
+#
+# License note: this tool only *reads* the vendor binary for interoperability
+# analysis; it contains no vendor code and is published under BSD-3-Clause.
+
+import argparse
+import json
+import re
+import struct
+import sys
+
+ELF_MAGIC = b"\x7fELF"
+
+
+class ELFError(Exception):
+    pass
+
+
+class ELF:
+    """Minimal 64-bit little-endian ELF reader."""
+
+    def __init__(self, path):
+        self.path = path
+        with open(path, "rb") as f:
+            self.data = f.read()
+        if self.data[:4] != ELF_MAGIC:
+            raise ELFError(f"{path}: not an ELF file")
+        if self.data[4] != 2:
+            raise ELFError(f"{path}: not ELF64")
+        if self.data[5] != 1:
+            raise ELFError(f"{path}: not little-endian")
+        (self.e_type, self.e_machine, _v, self.e_entry, self.e_phoff, self.e_shoff,
+         self.e_flags, self.e_ehsize, self.e_phentsize, self.e_phnum,
+         self.e_shentsize, self.e_shnum, self.e_shstrndx) = struct.unpack_from(
+            "<HHIQQQIHHHHHH", self.data, 16)
+
+        self.program_headers = []
+        for i in range(self.e_phnum):
+            off = self.e_phoff + i * self.e_phentsize
+            (p_type, p_flags, p_offset, p_vaddr, p_paddr, p_filesz, p_memsz,
+             p_align) = struct.unpack_from("<IIQQQQQQ", self.data, off)
+            self.program_headers.append(dict(
+                type=p_type, flags=p_flags, offset=p_offset, vaddr=p_vaddr,
+                paddr=p_paddr, filesz=p_filesz, memsz=p_memsz, align=p_align,
+                index=i))
+
+        self.sections = []
+        self.section_by_name = {}
+        for i in range(self.e_shnum):
+            off = self.e_shoff + i * self.e_shentsize
+            (sh_name, sh_type, sh_flags, sh_addr, sh_offset, sh_size,
+             sh_link, sh_info, sh_addralign, sh_entsize) = struct.unpack_from(
+                "<IIQQQQIIQQ", self.data, off)
+            self.sections.append(dict(
+                index=i, name_off=sh_name, type=sh_type, flags=sh_flags,
+                addr=sh_addr, offset=sh_offset, size=sh_size, link=sh_link,
+                info=sh_info, addralign=sh_addralign, entsize=sh_entsize))
+        if self.e_shstrndx:
+            shstr = self.sections[self.e_shstrndx]
+            base = shstr["offset"]
+            for s in self.sections:
+                end = self.data.find(b"\0", base + s["name_off"])
+                s["name"] = self.data[base + s["name_off"]:end].decode(
+                    "utf-8", "replace")
+                self.section_by_name[s["name"]] = s
+
+    # ---- helpers ---------------------------------------------------------
+
+    def section(self, name):
+        return self.section_by_name.get(name)
+
+    def exec_segment_base(self):
+        """Base address reported by dmesg/libcuda mappings: first executable
+        PT_LOAD vaddr aligned as Linux maps it (page-aligned start)."""
+        for ph in self.program_headers:
+            if ph["type"] == 1 and (ph["flags"] & 1):  # PT_LOAD | PF_X
+                return ph["vaddr"] & ~0xFFF
+        raise ELFError("no executable PT_LOAD segment")
+
+    def read_string(self, offset):
+        end = self.data.find(b"\0", offset)
+        return self.data[offset:end].decode("utf-8", "replace")
+
+    def build_id(self):
+        for s in self.sections:
+            if s["name"] == ".note.gnu.build-id":
+                off = s["offset"]
+                namesz, descsz, ntype = struct.unpack_from("<III", self.data, off)
+                name = self.data[off + 12:off + 12 + namesz].rstrip(b"\0")
+                if name == b"GNU" and ntype == 3:
+                    desc = self.data[off + 12 + ((namesz + 3) & ~3):
+                                      off + 12 + ((namesz + 3) & ~3) + descsz]
+                    return desc.hex()
+        return None
+
+    # ---- symbols ---------------------------------------------------------
+
+    def dynamic_symbols(self):
+        symtab = self.section(".dynsym")
+        if not symtab:
+            return []
+        strtab = self.sections[symtab["link"]]
+        out = []
+        n = symtab["size"] // 24
+        for i in range(n):
+            off = symtab["offset"] + i * 24
+            st_name, st_info, st_other, st_shndx, st_value, st_size = \
+                struct.unpack_from("<IBBHQQ", self.data, off)
+            name = self.read_string(strtab["offset"] + st_name)
+            out.append(dict(
+                name=name,
+                value=st_value,
+                size=st_size,
+                type=st_info & 0xF,
+                bind=st_info >> 4,
+                shndx=st_shndx,
+                defined=st_shndx != 0,
+            ))
+        return out
+
+    def exports(self):
+        return [s for s in self.dynamic_symbols()
+                if s["defined"] and s["type"] == 2 and s["value"]]
+
+    def imports(self):
+        return [s for s in self.dynamic_symbols()
+                if not s["defined"] and s["name"]]
+
+    # ---- exception handling function table -------------------------------
+
+    def function_ranges(self):
+        """Return [(start, end)] from .eh_frame_hdr.
+
+        .eh_frame_hdr layout (DWARF .eh_frame_hdr, version 1):
+          u8 version; u8 eh_frame_ptr_enc; u8 fde_count_enc; u8 table_enc;
+          encoded eh_frame_ptr; encoded fde_count;
+          table[count] { encoded pc_begin; encoded fde_addr; }
+
+        GCC emits: eh_frame_ptr_enc = DW_EH_PE_pcrel|sdata4 (0x1b),
+        fde_count_enc = udata4 (0x03), table_enc = DW_EH_PE_datarel|sdata4
+        (0x3b). The datarel anchor is the start of .eh_frame_hdr. The parser
+        below handles the encodings generally (pcrel/datarel + 1/2/4/8-byte
+        and ULEB/SLEB), so it also works for other toolchains.
+        """
+        hdr = self.section(".eh_frame_hdr")
+        if not hdr:
+            return []
+        base = hdr["offset"]
+        anchor = hdr["addr"]
+        version = self.data[base]
+        if version != 1:
+            raise ELFError(f"unsupported .eh_frame_hdr version {version}")
+        eh_ptr_enc, fde_count_enc, table_enc = self.data[base + 1:base + 4]
+
+        def enc_size(enc):
+            return {0x0: 0, 0x1: 1, 0x2: 2, 0x3: 4, 0x4: 8,
+                    0x9: 1, 0xA: 2, 0xB: 4, 0xC: 8}.get(enc & 0x0F, 0)
+
+        def decode(enc, ptr, field_addr):
+            fmt = enc & 0x0F
+            if fmt == 0x0:
+                return 0
+            if fmt == 0x1:
+                val, shift = 0, 0
+                while True:
+                    b = self.data[ptr]
+                    ptr += 1
+                    val |= (b & 0x7F) << shift
+                    if not (b & 0x80):
+                        break
+                    shift += 7
+                return val
+            if fmt == 0x2:
+                return struct.unpack_from("<H", self.data, ptr)[0]
+            if fmt == 0x3:
+                return struct.unpack_from("<I", self.data, ptr)[0]
+            if fmt == 0x4:
+                return struct.unpack_from("<Q", self.data, ptr)[0]
+            if fmt == 0x9:
+                val, shift = 0, 0
+                while True:
+                    b = self.data[ptr]
+                    ptr += 1
+                    val |= (b & 0x7F) << shift
+                    shift += 7
+                    if not (b & 0x80):
+                        if b & 0x40:
+                            val -= 1 << shift
+                        break
+                return val
+            if fmt == 0xA:
+                return struct.unpack_from("<h", self.data, ptr)[0]
+            if fmt == 0xB:
+                return struct.unpack_from("<i", self.data, ptr)[0]
+            if fmt == 0xC:
+                return struct.unpack_from("<q", self.data, ptr)[0]
+            raise ELFError(f"unsupported eh encoding 0x{enc:02x}")
+
+        def absolute(enc, ptr):
+            val = decode(enc, ptr, anchor)
+            kind = enc & 0x70
+            if kind == 0x10:      # DW_EH_PE_pcrel: relative to field address
+                val += anchor + (ptr - base)
+            elif kind == 0x30:    # DW_EH_PE_datarel
+                val += anchor
+            return val
+
+        p = base + 4
+        p += enc_size(eh_ptr_enc)          # skip eh_frame_ptr
+        count = decode(fde_count_enc, p, anchor)
+        p += enc_size(fde_count_enc)
+        starts = []
+        for _ in range(count):
+            pc = absolute(table_enc, p)
+            p += enc_size(table_enc)
+            p += enc_size(table_enc)       # skip fde address
+            starts.append(pc)
+        starts.sort()
+        out = []
+        for i, pc in enumerate(starts):
+            end = starts[i + 1] if i + 1 < len(starts) else None
+            out.append((pc, end))
+        return out
+
+    # ---- CUDA forward compatibility note ---------------------------------
+
+    def cuda_note(self):
+        s = self.section(".note.cuda.fwd_compatibility")
+        if not s:
+            return None
+        off = s["offset"]
+        namesz, descsz, ntype = struct.unpack_from("<III", self.data, off)
+        name_off = off + 12
+        name = self.data[name_off:name_off + namesz].rstrip(b"\0").decode()
+        desc_off = name_off + ((namesz + 3) & ~3)
+        desc = self.data[desc_off:desc_off + descsz].rstrip(b"\0").decode()
+        try:
+            parsed = json.loads(desc)
+        except json.JSONDecodeError:
+            parsed = None
+        return dict(owner=name, raw=desc, parsed=parsed)
+
+    def debuglink(self):
+        s = self.section(".gnu_debuglink")
+        if not s:
+            return None
+        name = self.read_string(s["offset"])
+        crc = struct.unpack_from("<I", self.data, s["offset"] + (len(name) + 4) // 4 * 4)[0]
+        return dict(name=name, crc=crc)
+
+    # ---- export stubs -> generated wrapper ---------------------------------
+
+    def relative_relocs(self):
+        """Map address -> R_X86_64_RELATIVE addend (the initial slot value)."""
+        out = {}
+        s = self.section(".rela.dyn")
+        if not s:
+            return out
+        n = s["size"] // 24
+        for i in range(n):
+            off = s["offset"] + i * 24
+            r_offset, r_info, r_addend = struct.unpack_from("<QQq", self.data, off)
+            if (r_info & 0xFFFFFFFF) == 8:  # R_X86_64_RELATIVE
+                out[r_offset] = r_addend
+        return out
+
+    def wrappers(self):
+        """Resolve each exported cu* stub to its initial wrapper target.
+
+        The generated stubs are:
+            cmp dword ptr [rip+d], 0x321cba00   (7 bytes)
+            je  +0xc                            (2 bytes)
+            jmp qword ptr [rip+d]               (6 bytes, disp at stub+14)
+        so the interposition slot is stub + 18 + disp32. Its initial value is
+        set by a R_X86_64_RELATIVE relocation and points at the generated
+        wrapper (tracing/tool dispatch), whose first callee is the hand-written
+        implementation.
+        """
+        relocs = self.relative_relocs()
+        out = []
+        for s in self.exports():
+            if not s["name"].startswith("cu"):
+                continue
+            a = s["value"]
+            # sanity-check the cmp/jmp pattern before trusting the arithmetic
+            if self.data[a] != 0x81 or self.data[a + 10] != 0x74 or \
+                    self.data[a + 12] != 0xFF or self.data[a + 13] != 0x25:
+                continue
+            disp = struct.unpack_from("<i", self.data, a + 14)[0]
+            slot = a + 18 + disp
+            target = relocs.get(slot)
+            out.append(dict(name=s["name"], stub=a, slot=slot, wrapper=target))
+        return out
+
+    def text_bounds(self):
+        t = self.section(".text")
+        if not t:
+            return None
+        return (t["addr"], t["addr"] + t["size"])
+
+
+# ---- commands ------------------------------------------------------------
+
+
+def cmd_info(args):
+    elf = ELF(args.lib)
+    tb = elf.text_bounds()
+    note = elf.cuda_note()
+    print(f"file            : {args.lib}")
+    print(f"build id        : {elf.build_id()}")
+    print(f"e_type          : {elf.e_type} (3=DYN shared object)")
+    print(f"exec segment base: 0x{elf.exec_segment_base():x}")
+    print(f".text           : 0x{tb[0]:x} .. 0x{tb[1]:x} ({tb[1]-tb[0]} bytes)")
+    print(f"exports         : {len(elf.exports())}")
+    print(f"imports         : {len(elf.imports())}")
+    if note and note["parsed"]:
+        p = note["parsed"]
+        print(f"cuda note       : CUDA {p.get('CUDA version')}, "
+              f"driver ABI {p.get('driver')}, device {p.get('device')}")
+    dl = elf.debuglink()
+    if dl:
+        print(f"debuglink       : {dl['name']} (crc 0x{dl['crc']:08x})")
+    for s in elf.sections:
+        print(f"  section {s['name']:<28} addr=0x{s['addr']:08x} "
+              f"off=0x{s['offset']:08x} size={s['size']}")
+
+
+def cmd_exports(args):
+    elf = ELF(args.lib)
+    ex = sorted(elf.exports(), key=lambda s: s["value"])
+    if args.json:
+        print(json.dumps([{k: v for k, v in s.items()} for s in ex], indent=1))
+        return
+    for s in ex:
+        print(f"0x{s['value']:08x} {s['size']:>8} {s['name']}")
+
+
+def cmd_functions(args):
+    elf = ELF(args.lib)
+    fr = elf.function_ranges()
+    if args.json:
+        print(json.dumps([dict(start=hex(a), end=hex(b) if b else None)
+                          for a, b in fr], indent=1))
+        return
+    for a, b in fr:
+        print(f"0x{a:08x} 0x{b:08x}" if b else f"0x{a:08x} -")
+
+
+def cmd_nearest(args):
+    elf = ELF(args.lib)
+    target = int(args.addr, 16)
+    ex = sorted(elf.exports(), key=lambda s: s["value"])
+    prev = None
+    for s in ex:
+        if s["value"] <= target:
+            prev = s
+        else:
+            break
+    fr = elf.function_ranges()
+    fstart = None
+    for a, b in fr:
+        if a <= target and (b is None or target < b):
+            fstart = a
+            break
+    print(f"vaddr            : 0x{target:x}")
+    print(f"function start   : 0x{fstart:x}" if fstart is not None else
+          "function start   : <not covered by .eh_frame_hdr>")
+    if prev:
+        print(f"nearest export   : {prev['name']} + 0x{target - prev['value']:x} "
+              f"(export at 0x{prev['value']:x})")
+    else:
+        print("nearest export   : <none>")
+
+
+def cmd_crash(args):
+    elf = ELF(args.lib)
+    base = int(args.map_base, 16)
+    ip = int(args.ip, 16)
+    vaddr = ip - base + elf.exec_segment_base()
+    print(f"file vaddr       : 0x{vaddr:x}")
+    args.addr = hex(vaddr)
+    cmd_nearest(args)
+
+
+def cmd_strings(args):
+    elf = ELF(args.lib)
+    pat = re.compile(args.grep) if args.grep else None
+    rodata = [s for s in elf.sections if s["type"] == 1 and not (s["flags"] & 4)]
+    seen = set()
+    for s in rodata:
+        blob = elf.data[s["offset"]:s["offset"] + s["size"]]
+        for m in re.finditer(rb"[\x20-\x7e]{5,}", blob):
+            text = m.group().decode()
+            if pat and not pat.search(text):
+                continue
+            key = (s["name"], m.start())
+            if key in seen:
+                continue
+            seen.add(key)
+            print(f"0x{s['addr'] + m.start():08x} {text}")
+
+
+def cmd_delta(args):
+    a, b = ELF(args.libA), ELF(args.libB)
+    ea = {s["name"]: s for s in a.exports()}
+    eb = {s["name"]: s for s in b.exports()}
+    added = sorted(set(eb) - set(ea))
+    removed = sorted(set(ea) - set(eb))
+    moved = []
+    resized = []
+    for n in sorted(set(ea) & set(eb)):
+        if ea[n]["value"] != eb[n]["value"]:
+            moved.append((n, ea[n]["value"], eb[n]["value"]))
+        if ea[n]["size"] != eb[n]["size"]:
+            resized.append((n, ea[n]["size"], eb[n]["size"]))
+    print(f"A: {args.libA}")
+    print(f"   build-id {a.build_id()} exports {len(ea)}")
+    print(f"B: {args.libB}")
+    print(f"   build-id {b.build_id()} exports {len(eb)}")
+    na, nb = a.cuda_note(), b.cuda_note()
+    print(f"note A: {na['raw'] if na else None}")
+    print(f"note B: {nb['raw'] if nb else None}")
+    tb_a, tb_b = a.text_bounds(), b.text_bounds()
+    print(f"text A: 0x{tb_a[0]:x}..0x{tb_a[1]:x} ({tb_a[1]-tb_a[0]})")
+    print(f"text B: 0x{tb_b[0]:x}..0x{tb_b[1]:x} ({tb_b[1]-tb_b[0]})")
+    print(f"exports added   ({len(added)}): " + ", ".join(added[:40]) +
+          (" ..." if len(added) > 40 else ""))
+    print(f"exports removed ({len(removed)}): " + ", ".join(removed[:40]) +
+          (" ..." if len(removed) > 40 else ""))
+    print(f"exports moved   ({len(moved)})" +
+          (": " + ", ".join(f"{n} 0x{x:x}->0x{y:x}" for n, x, y in moved[:10])
+           if moved else ""))
+    print(f"exports resized ({len(resized)})" +
+          (": " + ", ".join(f"{n} {x}->{y}" for n, x, y in resized[:10])
+           if resized else ""))
+    print(f"imports A={len(a.imports())} B={len(b.imports())}")
+
+
+def cmd_wrappers(args):
+    elf = ELF(args.lib)
+    ws = elf.wrappers()
+    if args.json:
+        print(json.dumps(ws, indent=1))
+        return
+    for w in ws:
+        if w["wrapper"] is None:
+            continue
+        print(f"0x{w['stub']:08x} -> 0x{w['wrapper']:08x} {w['name']}")
+
+
+def cmd_note(args):
+    elf = ELF(args.lib)
+    print(json.dumps(elf.cuda_note(), indent=2))
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__,
+                                 formatter_class=argparse.RawDescriptionHelpFormatter)
+    sub = ap.add_subparsers(dest="cmd", required=True)
+
+    p = sub.add_parser("info"); p.add_argument("lib"); p.set_defaults(fn=cmd_info)
+    p = sub.add_parser("exports"); p.add_argument("lib")
+    p.add_argument("--json", action="store_true"); p.set_defaults(fn=cmd_exports)
+    p = sub.add_parser("functions"); p.add_argument("lib")
+    p.add_argument("--json", action="store_true"); p.set_defaults(fn=cmd_functions)
+    p = sub.add_parser("nearest"); p.add_argument("lib"); p.add_argument("addr")
+    p.set_defaults(fn=cmd_nearest)
+    p = sub.add_parser("crash"); p.add_argument("lib")
+    p.add_argument("map_base"); p.add_argument("ip"); p.set_defaults(fn=cmd_crash)
+    p = sub.add_parser("strings"); p.add_argument("lib")
+    p.add_argument("--grep"); p.set_defaults(fn=cmd_strings)
+    p = sub.add_parser("delta"); p.add_argument("libA"); p.add_argument("libB")
+    p.set_defaults(fn=cmd_delta)
+    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)
+
+    args = ap.parse_args()
+    try:
+        args.fn(args)
+    except ELFError as e:
+        print(f"error: {e}", file=sys.stderr)
+        sys.exit(1)
+
+
+if __name__ == "__main__":
+    main()
diff --git a/userspace/cuda/tools/ghidra/ApplyNameMap.java b/userspace/cuda/tools/ghidra/ApplyNameMap.java
new file mode 100644
index 0000000..b8be68f
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/ApplyNameMap.java
@@ -0,0 +1,86 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// ApplyNameMap.java - Ghidra post-script: apply a symbol name map to a
+// program. This is the version-transfer hook of the Open-IOV CUDA tooling:
+// once a name map exists for version N (derived from public documentation,
+// decompiled behaviour or an earlier version), it can be applied to version
+// N+1 in seconds, so analysis effort accumulates instead of restarting.
+//
+// Input format (TSV, no header required):
+//   <hex-addr> <name> [<optional namespace>]
+// Lines starting with '#' and empty lines are ignored. Addresses are file
+// virtual addresses ("libcuda+0x..."), not Ghidra image-base offsets.
+//
+// Usage (analyzeHeadless):
+//   -process <program> -noanalysis \
+//     -postScript ApplyNameMap.java <namemap.tsv> [--dry-run]
+//
+// The script prints how many names were applied, and leaves a
+// "cuda-names-applied" user property on the program with the map path and
+// count so a later session can see it happened.
+
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.Listing;
+import ghidra.program.model.symbol.SourceType;
+
+import java.io.BufferedReader;
+import java.io.FileReader;
+
+public class ApplyNameMap extends GhidraScript {
+
+    @Override
+    public void run() throws Exception {
+        String[] args = getScriptArgs();
+        if (args.length < 1) {
+            println("usage: ApplyNameMap.java <namemap.tsv> [--dry-run]");
+            return;
+        }
+        boolean dryRun = args.length > 1 && args[1].equals("--dry-run");
+        Listing listing = currentProgram.getListing();
+        long imageBase = currentProgram.getImageBase().getOffset();
+        int applied = 0, missing = 0, lines = 0;
+
+        try (BufferedReader r = new BufferedReader(new FileReader(args[0]))) {
+            String line;
+            while ((line = r.readLine()) != null && !monitor.isCancelled()) {
+                line = line.trim();
+                if (line.isEmpty() || line.startsWith("#")) continue;
+                lines++;
+                String[] parts = line.split("[ \t]+");
+                if (parts.length < 2) continue;
+                long addr;
+                try {
+                    addr = Long.parseUnsignedLong(parts[0].replaceFirst("^0x", ""), 16);
+                } catch (NumberFormatException e) {
+                    continue;
+                }
+                String name = parts[1];
+                // Map keys are file virtual addresses (libcuda+0x...).
+                Address a = toAddr(addr + imageBase);
+                Function f = listing.getFunctionAt(a);
+                if (f == null) {
+                    missing++;
+                    continue;
+                }
+                if (!dryRun) {
+                    try {
+                        f.setName(name, SourceType.USER_DEFINED);
+                    } catch (Exception e) {
+                        println("  cannot rename 0x" + Long.toHexString(addr)
+                                + " -> " + name + ": " + e.getMessage());
+                        continue;
+                    }
+                }
+                applied++;
+            }
+        }
+        if (!dryRun) {
+            currentProgram.getOptions("OpenIOV")
+                    .setString("cuda-names-applied", args[0] + " (" + applied + " names)");
+        }
+        println(String.format("[ApplyNameMap] lines=%d applied=%s missing=%d%s",
+                lines, String.valueOf(applied), missing, dryRun ? " (dry run)" : ""));
+    }
+}
diff --git a/userspace/cuda/tools/ghidra/BatchDecompile.java b/userspace/cuda/tools/ghidra/BatchDecompile.java
new file mode 100644
index 0000000..0319681
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/BatchDecompile.java
@@ -0,0 +1,141 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// BatchDecompile.java - Ghidra post-script: decompile every function whose
+// entry address lies in [start, end) and write the C output plus a
+// machine-readable index.
+//
+// Designed for a very large stripped vendor library (~100k functions, 70 MB).
+// The caller splits the code address range into chunks and runs one
+// analyzeHeadless process per chunk (all with -readOnly against the same
+// analyzed project), which parallelizes and makes the run resumable: a chunk
+// whose output already exists is skipped.
+//
+// Usage (analyzeHeadless, on an already analyzed program):
+//   -process <program> -readOnly -noanalysis \
+//     -postScript BatchDecompile.java <startAddrHex> <endAddrHex> <outdir> [timeoutSecs]
+//
+// Outputs:
+//   <outdir>/buckets/bucket_<start>_<end>.c      decompiled C, functions in order
+//   <outdir>/index/bucket_<start>_<end>.tsv      addr  size  name  file  line  status
+//
+// Decompiled output is vendor-derived analysis output: it stays in the
+// version's gitignored run/ tree and is NEVER committed.
+
+import ghidra.app.script.GhidraScript;
+import ghidra.app.decompiler.DecompInterface;
+import ghidra.app.decompiler.DecompileResults;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionIterator;
+import ghidra.program.model.listing.Listing;
+
+import java.io.BufferedWriter;
+import java.io.File;
+import java.io.FileWriter;
+import java.io.PrintWriter;
+
+public class BatchDecompile extends GhidraScript {
+
+    private static int countLines(String s) {
+        int n = 0;
+        for (int i = 0; i < s.length(); i++) {
+            if (s.charAt(i) == '\n') n++;
+        }
+        return n;
+    }
+
+    @Override
+    public void run() throws Exception {
+        String[] args = getScriptArgs();
+        if (args.length < 3) {
+            println("usage: BatchDecompile.java <startAddrHex> <endAddrHex> <outdir> [timeoutSecs]");
+            return;
+        }
+        long start = Long.parseUnsignedLong(args[0].replaceFirst("^0x", ""), 16);
+        long end = Long.parseUnsignedLong(args[1].replaceFirst("^0x", ""), 16);
+        String outdir = args[2];
+        int timeout = args.length > 3 ? Integer.parseInt(args[3]) : 120;
+
+        File buckets = new File(outdir, "buckets");
+        File indexDir = new File(outdir, "index");
+        buckets.mkdirs();
+        indexDir.mkdirs();
+        String tag = String.format("%08x_%08x", start, end);
+        File outFile = new File(buckets, "bucket_" + tag + ".c");
+        File idxFile = new File(indexDir, "bucket_" + tag + ".tsv");
+
+        if (outFile.isFile() && idxFile.isFile()) {
+            println("[BatchDecompile] chunk " + tag + " already done, skipping");
+            return;
+        }
+
+        Listing listing = currentProgram.getListing();
+        long imageBase = currentProgram.getImageBase().getOffset();
+        DecompInterface dif = new DecompInterface();
+        dif.toggleCCode(true);
+        dif.toggleSyntaxTree(false);
+        dif.openProgram(currentProgram);
+
+        int n = 0, fail = 0, skip = 0, lines = 0;
+
+        try (PrintWriter idx = new PrintWriter(new BufferedWriter(new FileWriter(idxFile)));
+             PrintWriter out = new PrintWriter(new BufferedWriter(new FileWriter(outFile)))) {
+
+            String banner = "/* CUDA userspace driver decompilation - Open-IOV project\n"
+                    + " * program: " + currentProgram.getName() + "\n"
+                    + " * range:   [0x" + Long.toHexString(start) + ", 0x"
+                    + Long.toHexString(end) + ")\n"
+                    + " * Machine-generated analysis output; not part of the public tree.\n"
+                    + " */\n";
+            out.print(banner);
+            lines += countLines(banner);
+            idx.println("addr\tfile_addr\tsize\tname\tfile\tline\tstatus");
+
+            FunctionIterator fit = listing.getFunctions(true);
+            while (fit.hasNext() && !monitor.isCancelled()) {
+                Function f = fit.next();
+                if (f.isExternal()) continue;
+                long addr = f.getEntryPoint().getOffset();
+                if (addr < start || addr >= end) continue;
+
+                long size = f.getBody().getNumAddresses();
+                String header = String.format(
+                        "\n/* ------------------------------------------------------------------"
+                        + "----------\n"
+                        + " * %s\n"
+                        + " * addr=0x%08x file_addr=libcuda+0x%x size=%d\n"
+                        + " *-------------------------------------------------------------------"
+                        + "-------- */\n",
+                        f.getName(), addr, addr - imageBase, size);
+                out.print(header);
+                lines += countLines(header);
+
+                String status = "ok";
+                String code;
+                if (size > 200000) {
+                    status = "skipped-huge";
+                    skip++;
+                    code = "/* skipped: body too large (" + size + " bytes) */\n";
+                } else {
+                    DecompileResults res = dif.decompileFunction(f, timeout, monitor);
+                    if (res != null && res.decompileCompleted()) {
+                        code = res.getDecompiledFunction().getC();
+                    } else {
+                        status = "failed";
+                        fail++;
+                        code = "/* decompilation failed */\n";
+                    }
+                }
+                out.print(code);
+                out.print("\n");
+                lines += countLines(code) + 1;
+                idx.printf("0x%08x\t0x%08x\t%d\t%s\t%s\t%d\t%s%n",
+                        addr, addr - imageBase, size, f.getName(),
+                        outFile.getName(), lines - countLines(code), status);
+                n++;
+            }
+        }
+
+        println(String.format("[BatchDecompile] %s: functions=%d failed=%d skipped=%d -> %s",
+                tag, n, fail, skip, outFile.getAbsolutePath()));
+    }
+}
diff --git a/userspace/cuda/tools/ghidra/ConfigureAnalysis.java b/userspace/cuda/tools/ghidra/ConfigureAnalysis.java
new file mode 100644
index 0000000..71bb75e
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/ConfigureAnalysis.java
@@ -0,0 +1,84 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// ConfigureAnalysis.java - Ghidra pre-script used by the Open-IOV CUDA
+// tooling to make headless auto-analysis of a large stripped vendor library
+// tractable and reproducible.
+//
+// The CUDA userspace driver is a ~70 MB, fully stripped x86-64 shared object
+// with roughly a hundred thousand functions. The default analyzer set spends
+// the majority of its time in the two optional "identifier" analyzers, which
+// are not needed to obtain readable decompilation:
+//
+//   * Decompiler Parameter ID  - decompiles every function to recover
+//                                parameter types (very expensive, no value
+//                                for a stripped binary where we have no
+//                                type definitions for internal structs)
+//   * Function ID (FID)        - matches function bodies against .fidb
+//                                signature databases (liberal, and we have
+//                                no curated database for this toolchain)
+//
+// Everything else (disassembly, function start search, aggressive instruction
+// finder, exception handling, stack analysis, reference analyzers, ...) is
+// kept, because function discovery quality on a stripped binary depends on
+// them.
+//
+// Usage (analyzeHeadless):
+//   -preScript ConfigureAnalysis.java [verbose]
+//
+// Re-runnable: the same settings are applied on every analysis of a program.
+
+import ghidra.app.script.GhidraScript;
+import ghidra.framework.options.Options;
+
+public class ConfigureAnalysis extends GhidraScript {
+
+    private static final String[] DISABLE = {
+        "Decompiler Parameter ID",
+        "Function ID",
+    };
+
+    private static final String[] ENABLE = {
+        "Aggressive Instruction Finder",
+        "GCC Exception Handlers",
+        "Function Start Search",
+        "Shared Return Calls",
+        "Non-Returning Functions - Known",
+        "Non-Returning Functions - Discovered",
+    };
+
+    @Override
+    public void run() throws Exception {
+        boolean verbose = getScriptArgs().length > 0;
+        Options opts = currentProgram.getOptions("Analyzers");
+
+        if (verbose) {
+            println("[ConfigureAnalysis] analyzer options before:");
+            for (String name : opts.getOptionNames()) {
+                println("  " + name + " = " + opts.getObject(name, null));
+            }
+        }
+
+        for (String name : DISABLE) {
+            if (opts.contains(name)) {
+                opts.setBoolean(name, false);
+                println("[ConfigureAnalysis] disabled: " + name);
+            } else {
+                println("[ConfigureAnalysis] WARN not found: " + name);
+            }
+        }
+
+        for (String name : ENABLE) {
+            if (opts.contains(name)) {
+                opts.setBoolean(name, true);
+                println("[ConfigureAnalysis] enabled: " + name);
+            }
+        }
+
+        if (verbose) {
+            println("[ConfigureAnalysis] analyzer options after:");
+            for (String name : opts.getOptionNames()) {
+                println("  " + name + " = " + opts.getObject(name, null));
+            }
+        }
+    }
+}
diff --git a/userspace/cuda/tools/ghidra/DecompileAt.java b/userspace/cuda/tools/ghidra/DecompileAt.java
new file mode 100644
index 0000000..ed8cf5d
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/DecompileAt.java
@@ -0,0 +1,86 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// DecompileAt.java - Ghidra post-script: decompile the function containing a
+// given file virtual address. Version-agnostic copy of the per-version
+// helper used with ghidra_decompile_at.sh.
+//
+// Usage:
+//   analyzeHeadless <proj> <name> -process <program> -readOnly -noanalysis \
+//     -postScript DecompileAt.java <hex-vaddr> [outfile]
+//
+// If the address does not fall inside a known function, the script tries to
+// disassemble at the address; if the target is outside any function body it
+// prints a comment with the nearest preceding function instead of failing.
+
+import ghidra.app.script.GhidraScript;
+import ghidra.app.decompiler.DecompInterface;
+import ghidra.app.decompiler.DecompileResults;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionIterator;
+import ghidra.program.model.listing.Listing;
+
+import java.io.FileWriter;
+
+public class DecompileAt extends GhidraScript {
+    @Override
+    public void run() throws Exception {
+        String[] args = getScriptArgs();
+        if (args.length < 1) {
+            println("usage: DecompileAt.java <hex-vaddr> [outfile]");
+            return;
+        }
+        long fileAddr = Long.parseUnsignedLong(args[0].replaceFirst("^0x", ""), 16);
+        // Arguments are FILE virtual addresses (libcuda+0x...). Ghidra may have
+        // loaded the program at a non-zero image base (0x100000 in this build),
+        // so translate before resolving. If image base is 0 this is a no-op.
+        long imageBase = currentProgram.getImageBase().getOffset();
+        Address addr = toAddr(fileAddr + imageBase);
+        long target = addr.getOffset();
+
+        Function f = getFunctionContaining(addr);
+        if (f == null) {
+            Listing listing = currentProgram.getListing();
+            FunctionIterator it = listing.getFunctions(addr, false);
+            if (it.hasNext()) f = it.next();
+        }
+        if (f == null) {
+            disassemble(addr);
+            f = createFunction(addr, null);
+        }
+        StringBuilder sb = new StringBuilder();
+        if (f == null) {
+            sb.append("// no function at 0x").append(Long.toHexString(target)).append("\n");
+        } else {
+            DecompInterface dif = new DecompInterface();
+            dif.toggleCCode(true);
+            dif.toggleSyntaxTree(true);
+            dif.openProgram(currentProgram);
+            DecompileResults res = dif.decompileFunction(f, 300, monitor);
+            sb.append("// function ").append(f.getName())
+              .append(" at libcuda+0x")
+              .append(Long.toHexString(f.getEntryPoint().getOffset() - imageBase))
+              .append(" (program 0x").append(Long.toHexString(f.getEntryPoint().getOffset()))
+              .append(", query file 0x").append(Long.toHexString(fileAddr)).append(")\n");
+            if (res != null && res.decompileCompleted()) {
+                sb.append("#include <stdint.h>\n")
+                  .append("typedef unsigned char undefined1;\n")
+                  .append("typedef unsigned short undefined2;\n")
+                  .append("typedef unsigned int undefined4;\n")
+                  .append("typedef unsigned long long undefined8;\n")
+                  .append("typedef long long longlong;\n")
+                  .append("typedef unsigned long long ulonglong;\n\n")
+                  .append(res.getDecompiledFunction().getC());
+            } else {
+                sb.append("// decompilation failed\n");
+            }
+        }
+        String text = sb.toString();
+        if (args.length > 1) {
+            try (FileWriter w = new FileWriter(args[1])) {
+                w.write(text);
+            }
+        }
+        println(text);
+    }
+}
diff --git a/userspace/cuda/tools/ghidra/ExportFunctionIndex.java b/userspace/cuda/tools/ghidra/ExportFunctionIndex.java
new file mode 100644
index 0000000..c36f98f
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/ExportFunctionIndex.java
@@ -0,0 +1,192 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// ExportFunctionIndex.java - Ghidra post-script: dump a complete inventory of
+// the analyzed program: functions, symbols (exports/imports), memory blocks
+// and per-function references to interesting callees, plus the address of
+// every function that is reachable from the exported CUDA entry points.
+//
+// This is the primary "map" artifact for a CUDA userspace driver version: it
+// lets later tooling and analysis answer "what is at libcuda+0xNNNNNN" and
+// "which internal function implements exported cuFoo".
+//
+// Usage (analyzeHeadless, on an already analyzed program):
+//   -process <program> -noanalysis \
+//     -postScript ExportFunctionIndex.java <outdir>
+//
+// Outputs (all TSV, one row per item, with a header row):
+//   <outdir>/functions.tsv    addr  size  name  namespace  is_exported
+//                             is_external  is_thunk  has_body  entry_bytes
+//   <outdir>/exports.tsv      addr  size  name
+//   <outdir>/imports.tsv      addr  name  library
+//   <outdir>/blocks.tsv       name  start  end  size  r/w/x
+//   <outdir>/callgraph.tsv    caller_addr  callee_addr  call_sites
+//   <outdir>/strings.tsv      addr  length  string (first 200 chars)
+//
+// The callgraph is produced from Ghidra's reference model for call
+// instructions inside each function body.
+
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.address.AddressSetView;
+import ghidra.program.model.listing.Data;
+import ghidra.program.model.listing.DataIterator;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionIterator;
+import ghidra.program.model.listing.Listing;
+import ghidra.program.model.mem.MemoryBlock;
+import ghidra.program.model.symbol.Reference;
+import ghidra.program.model.symbol.ReferenceIterator;
+import ghidra.program.model.symbol.Symbol;
+import ghidra.program.model.symbol.SymbolIterator;
+import ghidra.program.model.symbol.SymbolTable;
+
+import java.io.BufferedWriter;
+import java.io.FileWriter;
+import java.io.PrintWriter;
+import java.util.HashMap;
+import java.util.Map;
+
+public class ExportFunctionIndex extends GhidraScript {
+
+    private PrintWriter open(String dir, String name) throws Exception {
+        return new PrintWriter(new BufferedWriter(new FileWriter(dir + "/" + name)));
+    }
+
+    @Override
+    public void run() throws Exception {
+        String[] args = getScriptArgs();
+        if (args.length < 1) {
+            println("usage: ExportFunctionIndex.java <outdir>");
+            return;
+        }
+        String outdir = args[0];
+        new java.io.File(outdir).mkdirs();
+
+        Listing listing = currentProgram.getListing();
+        SymbolTable symtab = currentProgram.getSymbolTable();
+        // Ghidra may load a PIE/shared object at a non-zero image base (this
+        // build uses 0x100000). All *file virtual addresses* (what the public
+        // docs cite as libcuda+0x...) are program addresses minus this base.
+        long imageBase = currentProgram.getImageBase().getOffset();
+
+        // ---- functions -------------------------------------------------------
+        int nfunc = 0, nexport = 0, nexternal = 0;
+        Map<Long, String> exported = new HashMap<>();
+        try (PrintWriter w = open(outdir, "functions.tsv")) {
+            w.println("addr\tfile_addr\tsize\tname\tnamespace\texported\texternal\tthunk\tbody");
+            FunctionIterator fit = listing.getFunctions(true);
+            while (fit.hasNext() && !monitor.isCancelled()) {
+                Function f = fit.next();
+                Address ep = f.getEntryPoint();
+                boolean isExt = f.isExternal();
+                boolean isThunk = f.isThunk();
+                String ns = f.getParentNamespace() != null
+                        ? f.getParentNamespace().getName(true) : "";
+                String name = f.getName();
+                long size = f.getBody().getNumAddresses();
+                w.printf("0x%08x\t0x%08x\t%d\t%s\t%s\t%s\t%s\t%s\t%s%n",
+                        ep.getOffset(), ep.getOffset() - imageBase, size, name, ns,
+                        isExt ? "0" : (symtab.isExternalEntryPoint(ep) ? "1" : "0"),
+                        isExt ? "1" : "0", isThunk ? "1" : "0",
+                        size > 0 ? "1" : "0");
+                nfunc++;
+                if (!isExt) {
+                    for (Symbol s : symtab.getSymbols(ep)) {
+                        if (s.isExternalEntryPoint() || s.isGlobal()) {
+                            exported.put(ep.getOffset(), s.getName(true));
+                        }
+                    }
+                }
+                if (isExt) nexternal++;
+            }
+        }
+
+        // ---- exports: symbols flagged external-entry / global in .text ------
+        // The loader marks ELF .dynsym definitions; enumerating them via the
+        // symbol table is exact and independent of the function iterator.
+        try (PrintWriter w = open(outdir, "exports.tsv")) {
+            w.println("addr\tfile_addr\tsize\tname");
+            SymbolIterator sit = symtab.getAllSymbols(true);
+            while (sit.hasNext() && !monitor.isCancelled()) {
+                Symbol s = sit.next();
+                Address a = s.getAddress();
+                if (a == null || !a.isMemoryAddress()) continue;
+                if (!symtab.isExternalEntryPoint(a)) continue;
+                Function f = listing.getFunctionAt(a);
+                long size = f != null ? f.getBody().getNumAddresses() : 0;
+                w.printf("0x%08x\t0x%08x\t%d\t%s%n", a.getOffset(), a.getOffset() - imageBase, size, s.getName(true));
+                nexport++;
+            }
+        }
+
+        // ---- imports: external functions and the library they come from ----
+        try (PrintWriter w = open(outdir, "imports.tsv")) {
+            w.println("addr\tfile_addr\tname\tlibrary");
+            FunctionIterator fit = listing.getFunctions(true);
+            while (fit.hasNext() && !monitor.isCancelled()) {
+                Function f = fit.next();
+                if (!f.isExternal()) continue;
+                String lib = f.getParentNamespace() != null
+                        ? f.getParentNamespace().getName(true) : "";
+                w.printf("0x%08x\t0x%08x\t%s\t%s%n",
+                        f.getEntryPoint().getOffset(),
+                        f.getEntryPoint().getOffset() - imageBase, f.getName(), lib);
+            }
+        }
+
+        // ---- memory blocks ---------------------------------------------------
+        try (PrintWriter w = open(outdir, "blocks.tsv")) {
+            w.println("name\tstart\tend\tsize\tr\tw\tx\tinitialized");
+            for (MemoryBlock b : currentProgram.getMemory().getBlocks()) {
+                w.printf("%s\t0x%08x\t0x%08x\t%d\t%s\t%s\t%s\t%s%n",
+                        b.getName(), b.getStart().getOffset(), b.getEnd().getOffset(),
+                        b.getSize(), b.isRead() ? "1" : "0", b.isWrite() ? "1" : "0",
+                        b.isExecute() ? "1" : "0", b.isInitialized() ? "1" : "0");
+            }
+        }
+
+        // ---- call graph ------------------------------------------------------
+        // Function.getCalledFunctions() is the reliable API: it resolves both
+        // direct calls and call-through references recorded by the analyzers,
+        // independent of how the reference manager indexes them.
+        int edges = 0;
+        try (PrintWriter w = open(outdir, "callgraph.tsv")) {
+            w.println("caller\tcallee\tcaller_file\tcallee_file");
+            FunctionIterator fit = listing.getFunctions(true);
+            while (fit.hasNext() && !monitor.isCancelled()) {
+                Function f = fit.next();
+                if (f.isExternal()) continue;
+                for (Function callee : f.getCalledFunctions(monitor)) {
+                    w.printf("0x%08x\t0x%08x\t0x%08x\t0x%08x%n",
+                            f.getEntryPoint().getOffset(),
+                            callee.getEntryPoint().getOffset(),
+                            f.getEntryPoint().getOffset() - imageBase,
+                            callee.getEntryPoint().getOffset() - imageBase);
+                    edges++;
+                }
+            }
+        }
+
+        // ---- defined strings -------------------------------------------------
+        int nstr = 0;
+        try (PrintWriter w = open(outdir, "strings.tsv")) {
+            w.println("addr\tfile_addr\tlength\tstring");
+            DataIterator dit = listing.getDefinedData(true);
+            while (dit.hasNext() && !monitor.isCancelled()) {
+                Data d = dit.next();
+                if (!d.hasStringValue()) continue;
+                Object v = d.getValue();
+                if (v == null) continue;
+                String s = v.toString().replace("\t", " ").replace("\n", "\\n");
+                if (s.length() > 200) s = s.substring(0, 200);
+                w.printf("0x%08x\t0x%08x\t%d\t%s%n", d.getAddress().getOffset(),
+                        d.getAddress().getOffset() - imageBase, d.getLength(), s);
+                nstr++;
+            }
+        }
+
+        println(String.format(
+                "[ExportFunctionIndex] functions=%d exports=%d imports=%d callgraph_edges=%d strings=%d",
+                nfunc, nexport, nexternal, edges, nstr));
+    }
+}
diff --git a/userspace/cuda/tools/ghidra/PrintProgramInfo.java b/userspace/cuda/tools/ghidra/PrintProgramInfo.java
new file mode 100644
index 0000000..86a31b3
--- /dev/null
+++ b/userspace/cuda/tools/ghidra/PrintProgramInfo.java
@@ -0,0 +1,38 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// PrintProgramInfo.java - Ghidra post-script: print the program's image base,
+// executable block ranges and function count. Used to establish the address
+// convention of a project before running decompilation tools.
+//
+// Usage:
+//   analyzeHeadless <proj> <name> -process <program> -readOnly -noanalysis \
+//     -postScript PrintProgramInfo.java
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionIterator;
+import ghidra.program.model.mem.MemoryBlock;
+
+public class PrintProgramInfo extends GhidraScript {
+    @Override
+    public void run() throws Exception {
+        println("program        = " + currentProgram.getName());
+        println("language       = " + currentProgram.getLanguageID());
+        println("image base     = 0x" + Long.toHexString(currentProgram.getImageBase().getOffset()));
+        println("executable blocks:");
+        for (MemoryBlock b : currentProgram.getMemory().getBlocks()) {
+            if (b.isExecute()) {
+                println("  0x" + Long.toHexString(b.getStart().getOffset())
+                        + " - 0x" + Long.toHexString(b.getEnd().getOffset())
+                        + "  " + b.getName());
+            }
+        }
+        int n = 0;
+        FunctionIterator fit = currentProgram.getListing().getFunctions(true);
+        while (fit.hasNext() && !monitor.isCancelled()) {
+            Function f = fit.next();
+            if (!f.isExternal()) n++;
+        }
+        println("functions      = " + n);
+    }
+}
diff --git a/userspace/cuda/tools/ghidra_batch_decompile.sh b/userspace/cuda/tools/ghidra_batch_decompile.sh
new file mode 100755
index 0000000..bce8cb9
--- /dev/null
+++ b/userspace/cuda/tools/ghidra_batch_decompile.sh
@@ -0,0 +1,83 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ghidra_batch_decompile.sh - decompile the whole code range of an analyzed
+# CUDA program in parallel. Chunks already present are skipped, so the command
+# is safe to re-run and to interrupt.
+#
+# Ghidra takes an exclusive project lock even in -readOnly mode, so a single
+# project cannot be opened by several processes. To parallelize, this script
+# makes one copy of the analyzed project per worker under <outdir>/wk/ and
+# round-robins the chunks across the copies. Copies are kept (they make
+# resuming cheap) and can be removed with: rm -rf <outdir>/wk
+#
+# Usage:
+#   ghidra_batch_decompile.sh <project_dir> <project_name> <program> <outdir> \
+#                             <start_hex> <end_hex> [chunk_size_hex] [jobs] [timeout_s]
+#
+# Defaults: chunk 0x100000, jobs 5, timeout 120.
+#   env: GHIDRA_HOME, GHIDRA_HEADLESS_MAXMEM (default 4G for batch jobs), MAX_CPU
+#
+# After it completes, merge per-chunk indexes with:
+#   ghidra_index_merge.sh <outdir>
+set -euo pipefail
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+GHIDRA_HOME="${GHIDRA_HOME:-/snap/ghidra/current/ghidra}"
+ANALYZE="$GHIDRA_HOME/support/analyzeHeadless"
+export GHIDRA_HEADLESS_MAXMEM="${GHIDRA_HEADLESS_MAXMEM:-4G}"
+
+PROJECT_DIR="${1:?usage: $0 <project_dir> <project_name> <program> <outdir> <start_hex> <end_hex> [chunk_hex] [jobs] [timeout_s]}"
+PROJECT_NAME="${2:?}"
+PROGRAM="${3:?}"
+OUTDIR="${4:?}"
+START="${5:?}"
+END="${6:?}"
+CHUNK="${7:-0x100000}"
+JOBS="${8:-5}"
+TIMEOUT="${9:-120}"
+
+mkdir -p "$OUTDIR/buckets" "$OUTDIR/index" "$OUTDIR/logs" "$OUTDIR/wk"
+
+START_D=$((START)); END_D=$((END)); CHUNK_D=$((CHUNK))
+
+# Build the chunk list first: every chunk gets its OWN project copy, so two
+# concurrent jobs can never race on the same Ghidra project lock.
+ranges=""
+n=0
+for ((a=START_D; a<END_D; a+=CHUNK_D)); do
+    b=$((a+CHUNK_D)); [ "$b" -gt "$END_D" ] && b=$END_D
+    ranges+=$(printf "0x%x 0x%x %d\n" "$a" "$b" "$n")$'\n'
+    n=$((n+1))
+done
+
+# Create project copies for the chunks that still need work.
+while read -r _a _b w; do
+    [ -z "$w" ] && continue
+    out="$OUTDIR/buckets/bucket_${_a#0x}_${_b#0x}.c"
+    idx="$OUTDIR/index/bucket_${_a#0x}_${_b#0x}.tsv"
+    if [ -f "$out" ] && [ -f "$idx" ]; then continue; fi
+    wdir="$OUTDIR/wk"
+    if [ ! -d "$wdir/$PROJECT_NAME.w$w.rep" ]; then
+        echo "[*] creating project copy $w"
+        cp -a "$PROJECT_DIR/$PROJECT_NAME.gpr" "$wdir/$PROJECT_NAME.w$w.gpr"
+        cp -a "$PROJECT_DIR/$PROJECT_NAME.rep" "$wdir/$PROJECT_NAME.w$w.rep"
+    fi
+done <<< "$ranges"
+
+echo "[*] decompiling $PROGRAM in $n chunks across $JOBS workers"
+echo "$ranges" | xargs -P "$JOBS" -n 3 bash -c '
+    set -e
+    a="$1"; b="$2"; w="$3"
+    tag="${a#0x}_${b#0x}"
+    log="'$OUTDIR'/logs/decompile_${tag}.log"
+    GHIDRA_HEADLESS_MAXMEM="'$GHIDRA_HEADLESS_MAXMEM'" "'$ANALYZE'" "'$OUTDIR'/wk" "'$PROJECT_NAME'.w$w" \
+        -process "'$PROGRAM'" -readOnly -noanalysis \
+        -scriptPath "'$HERE'/ghidra" \
+        -postScript BatchDecompile.java "$a" "$b" "'$OUTDIR'" "'$TIMEOUT'" \
+        > "$log" 2>&1 || { echo "[FAIL] $a-$b (see $log)"; exit 1; }
+    tail -1 "$log"
+' _
+
+echo "[ok] batch decompilation done -> $OUTDIR/buckets"
+echo "     merge indexes: $HERE/ghidra_index_merge.sh $OUTDIR"
diff --git a/userspace/cuda/tools/ghidra_decompile_at.sh b/userspace/cuda/tools/ghidra_decompile_at.sh
new file mode 100755
index 0000000..1fa9035
--- /dev/null
+++ b/userspace/cuda/tools/ghidra_decompile_at.sh
@@ -0,0 +1,29 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ghidra_decompile_at.sh - decompile the function containing a file virtual
+# address from an already analyzed project. Version-agnostic replacement for
+# the per-version helper (same idea, any project/binary).
+#
+# Usage: ghidra_decompile_at.sh <project_dir> <project_name> <program> <hex_vaddr> [outfile]
+set -euo pipefail
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+GHIDRA_HOME="${GHIDRA_HOME:-/snap/ghidra/current/ghidra}"
+ANALYZE="$GHIDRA_HOME/support/analyzeHeadless"
+export GHIDRA_HEADLESS_MAXMEM="${GHIDRA_HEADLESS_MAXMEM:-12G}"
+
+PROJECT_DIR="${1:?usage: $0 <project_dir> <project_name> <program> <hex_vaddr> [outfile]}"
+PROJECT_NAME="${2:?}"
+PROGRAM="${3:?}"
+VADDR="${4:?}"
+OUT="${5:-$PROJECT_DIR/fn_${VADDR}.c}"
+
+"$ANALYZE" "$PROJECT_DIR" "$PROJECT_NAME" \
+    -process "$PROGRAM" -readOnly -noanalysis \
+    -scriptPath "$HERE/ghidra" \
+    -postScript DecompileAt.java "$VADDR" "$OUT" \
+    2>&1 | tail -3
+
+[ -f "$OUT" ] || { echo "[FAIL] no output"; exit 1; }
+echo "[ok] $OUT"
diff --git a/userspace/cuda/tools/ghidra_export_index.sh b/userspace/cuda/tools/ghidra_export_index.sh
new file mode 100755
index 0000000..3709aae
--- /dev/null
+++ b/userspace/cuda/tools/ghidra_export_index.sh
@@ -0,0 +1,27 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ghidra_export_index.sh - export the complete function/symbol/callgraph/
+# string index of an analyzed CUDA program.
+#
+# Usage: ghidra_export_index.sh <project_dir> <project_name> <program> <outdir>
+set -euo pipefail
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+GHIDRA_HOME="${GHIDRA_HOME:-/snap/ghidra/current/ghidra}"
+ANALYZE="$GHIDRA_HOME/support/analyzeHeadless"
+export GHIDRA_HEADLESS_MAXMEM="${GHIDRA_HEADLESS_MAXMEM:-12G}"
+
+PROJECT_DIR="${1:?usage: $0 <project_dir> <project_name> <program> <outdir>}"
+PROJECT_NAME="${2:?}"
+PROGRAM="${3:?}"
+OUTDIR="${4:?}"
+mkdir -p "$OUTDIR"
+
+"$ANALYZE" "$PROJECT_DIR" "$PROJECT_NAME" \
+    -process "$PROGRAM" \
+    -readOnly -noanalysis \
+    -scriptPath "$HERE/ghidra" \
+    -postScript ExportFunctionIndex.java "$OUTDIR"
+
+echo "[ok] index written to $OUTDIR"
diff --git a/userspace/cuda/tools/ghidra_full_analysis.sh b/userspace/cuda/tools/ghidra_full_analysis.sh
new file mode 100755
index 0000000..a143a80
--- /dev/null
+++ b/userspace/cuda/tools/ghidra_full_analysis.sh
@@ -0,0 +1,49 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ghidra_full_analysis.sh - import a CUDA userspace library into a fresh
+# Ghidra project and run a full auto-analysis tuned for a large stripped
+# vendor binary.
+#
+# This is the fixed version of the original per-version helper: passing
+# `-analysisTimeoutPerFile 0` made this Ghidra build abort analysis after
+# 0 seconds (see run/ghidra-analyze.log: "Analysis timed out at 0 seconds"),
+# so the old project contains no analysis at all.
+#
+# Usage: ghidra_full_analysis.sh <binary> [project_dir] [project_name]
+#   env: GHIDRA_HOME          (default /snap/ghidra/current/ghidra)
+#        GHIDRA_HEADLESS_MAXMEM (default 12G)
+#
+# The project is kept so the expensive analysis runs exactly once; all later
+# decompilation uses ghidra_batch_decompile.sh / ghidra_decompile_at.sh with
+# -noanalysis against the same project.
+set -euo pipefail
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+GHIDRA_HOME="${GHIDRA_HOME:-/snap/ghidra/current/ghidra}"
+ANALYZE="$GHIDRA_HOME/support/analyzeHeadless"
+export GHIDRA_HEADLESS_MAXMEM="${GHIDRA_HEADLESS_MAXMEM:-12G}"
+
+BIN="${1:?usage: $0 <binary> [project_dir] [project_name]}"
+PROJECT_DIR="${2:-$(dirname "$BIN")/ghidra-full}"
+PROJECT_NAME="${3:-cuda-full}"
+
+[ -x "$ANALYZE" ] || { echo "analyzeHeadless not found at $ANALYZE"; exit 1; }
+[ -f "$BIN" ]     || { echo "no such binary: $BIN"; exit 1; }
+mkdir -p "$PROJECT_DIR"
+
+echo "[*] importing + analyzing $BIN"
+echo "    project: $PROJECT_DIR/$PROJECT_NAME"
+echo "    heap:    $GHIDRA_HEADLESS_MAXMEM"
+echo "    (Decompiler Parameter ID and Function ID are disabled by"
+echo "     ConfigureAnalysis.java; everything else stays on.)"
+
+"$ANALYZE" "$PROJECT_DIR" "$PROJECT_NAME" \
+    -import "$BIN" \
+    -overwrite \
+    -scriptPath "$HERE/ghidra" \
+    -preScript ConfigureAnalysis.java \
+    -analysisTimeoutPerFile 14400 \
+    -max-cpu "${MAX_CPU:-8}"
+
+echo "[ok] analysis complete: $PROJECT_DIR/$PROJECT_NAME"
diff --git a/userspace/cuda/tools/ghidra_index_merge.sh b/userspace/cuda/tools/ghidra_index_merge.sh
new file mode 100755
index 0000000..44479e9
--- /dev/null
+++ b/userspace/cuda/tools/ghidra_index_merge.sh
@@ -0,0 +1,18 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# ghidra_index_merge.sh - merge the per-chunk decompiled_index.tsv files from
+# ghidra_batch_decompile.sh into one sorted index and print coverage stats.
+#
+# Usage: ghidra_index_merge.sh <outdir>
+set -euo pipefail
+OUTDIR="${1:?usage: $0 <outdir>}"
+
+head -1 "$OUTDIR"/index/bucket_*.tsv | head -1 > "$OUTDIR/decompiled_index.tsv"
+cat "$OUTDIR"/index/bucket_*.tsv | grep -v '^addr' | sort -t$'\t' -k1,1 >> "$OUTDIR/decompiled_index.tsv"
+
+total=$(grep -c . "$OUTDIR/decompiled_index.tsv" || true)
+failed=$(grep -c $'\tfailed$' "$OUTDIR/decompiled_index.tsv" || true)
+skipped=$(grep -c 'skipped-huge' "$OUTDIR/decompiled_index.tsv" || true)
+echo "[index] rows=$((total-1)) failed=$failed skipped-huge=$skipped"
+echo "[index] $OUTDIR/decompiled_index.tsv"