From 21b8ebc9db1d0ae649d6406087ea22b54e3b5775 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 23 Sep 2026 12:11:04 -0400 Subject: 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). --- userspace/cuda/570.144/.gitignore | 10 +- userspace/cuda/570.144/README.md | 13 + userspace/cuda/570.144/annotations/README.md | 35 + userspace/cuda/570.144/annotations/names.tsv | 1304 ++++++++++++++++++++ userspace/cuda/570.144/annotations/subsystems.tsv | 37 + .../cuda/570.144/docs/ANALYSIS-cuctxcreate-segv.md | 13 + userspace/cuda/READ-ME-FIRST.md | 151 +++ userspace/cuda/README.md | 124 ++ userspace/cuda/docs/01-architecture.md | 302 +++++ userspace/cuda/docs/02-build-and-execution.md | 190 +++ userspace/cuda/docs/03-driver-and-gpu-interface.md | 276 +++++ .../cuda/docs/04-virtualization-and-emulation.md | 244 ++++ userspace/cuda/docs/05-570.144-decompilation.md | 315 +++++ userspace/cuda/docs/06-tooling.md | 165 +++ userspace/cuda/docs/07-version-upgrade-playbook.md | 98 ++ .../cuda/docs/08-virtualization-approaches.md | 67 + userspace/cuda/tools/annotate_corpus.py | 155 +++ userspace/cuda/tools/callwalk.py | 133 ++ userspace/cuda/tools/cudatool.py | 515 ++++++++ userspace/cuda/tools/ghidra/ApplyNameMap.java | 86 ++ userspace/cuda/tools/ghidra/BatchDecompile.java | 141 +++ userspace/cuda/tools/ghidra/ConfigureAnalysis.java | 84 ++ userspace/cuda/tools/ghidra/DecompileAt.java | 86 ++ .../cuda/tools/ghidra/ExportFunctionIndex.java | 192 +++ userspace/cuda/tools/ghidra/PrintProgramInfo.java | 38 + userspace/cuda/tools/ghidra_batch_decompile.sh | 83 ++ userspace/cuda/tools/ghidra_decompile_at.sh | 29 + userspace/cuda/tools/ghidra_export_index.sh | 27 + userspace/cuda/tools/ghidra_full_analysis.sh | 49 + userspace/cuda/tools/ghidra_index_merge.sh | 18 + 30 files changed, 4979 insertions(+), 1 deletion(-) create mode 100644 userspace/cuda/570.144/annotations/README.md create mode 100644 userspace/cuda/570.144/annotations/names.tsv create mode 100644 userspace/cuda/570.144/annotations/subsystems.tsv create mode 100644 userspace/cuda/READ-ME-FIRST.md create mode 100644 userspace/cuda/README.md create mode 100644 userspace/cuda/docs/01-architecture.md create mode 100644 userspace/cuda/docs/02-build-and-execution.md create mode 100644 userspace/cuda/docs/03-driver-and-gpu-interface.md create mode 100644 userspace/cuda/docs/04-virtualization-and-emulation.md create mode 100644 userspace/cuda/docs/05-570.144-decompilation.md create mode 100644 userspace/cuda/docs/06-tooling.md create mode 100644 userspace/cuda/docs/07-version-upgrade-playbook.md create mode 100644 userspace/cuda/docs/08-virtualization-approaches.md create mode 100755 userspace/cuda/tools/annotate_corpus.py create mode 100755 userspace/cuda/tools/callwalk.py create mode 100755 userspace/cuda/tools/cudatool.py create mode 100644 userspace/cuda/tools/ghidra/ApplyNameMap.java create mode 100644 userspace/cuda/tools/ghidra/BatchDecompile.java create mode 100644 userspace/cuda/tools/ghidra/ConfigureAnalysis.java create mode 100644 userspace/cuda/tools/ghidra/DecompileAt.java create mode 100644 userspace/cuda/tools/ghidra/ExportFunctionIndex.java create mode 100644 userspace/cuda/tools/ghidra/PrintProgramInfo.java create mode 100755 userspace/cuda/tools/ghidra_batch_decompile.sh create mode 100755 userspace/cuda/tools/ghidra_decompile_at.sh create mode 100755 userspace/cuda/tools/ghidra_export_index.sh create mode 100755 userspace/cuda/tools/ghidra_full_analysis.sh create mode 100755 userspace/cuda/tools/ghidra_index_merge.sh 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_vaddrnamesource` | symbol map to apply to a Ghidra project (or any address-keyed analysis). `source` is `export`, `wrapper`, `verified`, or `todo`. | +| `subsystems.tsv` | `file_vaddrnamesubsystemroleevidence` | 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_`. +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: \t\t +# 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_vaddrnamesubsystemroleevidence +# 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 +``` +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 +``` +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.`): 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__.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__.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 # size, exports, ABI note +tools/cudatool.py exports # symbol map +tools/cudatool.py functions # FDE function map +tools/cudatool.py delta # interface-level diff +tools/ghidra_full_analysis.sh # 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 file_addr=libcuda+0x` so +citations in docs and gdb sessions use the same numbers. + +## 5. Known limits + +* Names are `FUN_` 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.` 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 # build id, sections, exports, ABI note +tools/cudatool.py exports # addr size name (file vaddrs) +tools/cudatool.py functions # function ranges from .eh_frame_hdr +tools/cudatool.py nearest # function start + nearest export +tools/cudatool.py crash # dmesg -> file vaddr +tools/cudatool.py strings [--grep RE] +tools/cudatool.py note # .note.cuda.fwd_compatibility JSON +tools/cudatool.py delta # 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 + +# 2. export the index (fast, read-only) +tools/ghidra_export_index.sh + +# 3. decompile everything (parallel, resumable; safe to re-run) +tools/ghidra_batch_decompile.sh \ + [chunk_hex] [jobs] [timeout_s] +tools/ghidra_index_merge.sh +``` + +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 + `/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 +``` + +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_` 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 /callgraph.tsv callees 0x31b7f0 2 10 +tools/callwalk.py /callgraph.tsv callers 0x453ef0 +tools/callwalk.py /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 0x45a670 out.c + +# corpus search by file address (headers are annotated) +grep -n "libcuda+0x45a670" /buckets/*.c + +# index lookups +awk -F'\t' '$2=="0x45a670"' /decompiled_index.tsv +grep -P "^0x0055a670\t" /callgraph.tsv # direct callers/callees + +# exports/strings +grep cuLaunchKernel /exports.tsv +``` + +## 4. Version deltas — `cudatool.py delta` + +``` +tools/cudatool.py delta +``` + +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 -process -noanalysis \ + -scriptPath tools/ghidra -postScript ApplyNameMap.java [--dry-run] +``` + +The map is `file_vaddrname` 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.` (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//` — 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 +tools/cudatool.py delta > run/delta.txt +tools/cudatool.py crash # 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 /run/ghidra-full cuda-full +tools/ghidra_export_index.sh /run/ghidra-full cuda-full /run/analysis/elf-index +tools/ghidra_batch_decompile.sh /run/ghidra-full cuda-full /run/analysis/decompiled \ + 0x100000 5 120 +tools/ghidra_index_merge.sh /run/analysis/decompiled +``` + +Get ``/`` from `cudatool.py info`. Confirm the image base +with `PrintProgramInfo.java` before citing addresses. + +## 3. Transfer names and knowledge + +``` +cp /annotations/names.tsv /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 cuda-full -process -noanalysis \ + -scriptPath tools/ghidra -postScript ApplyNameMap.java /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--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_ references inside the +# function bodies, so a renamed function is readable everywhere. +# +# Inputs (produced by the toolchain): +# /buckets/bucket_*.c raw decompiled C +# /decompiled_index.tsv addr file_addr size name file line status +# /callgraph.tsv caller callee caller_file callee_file +# file_vaddrname +# +# Usage: +# annotate_corpus.py +# +# 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[^\n]*)\n \* addr=(?P0x[0-9a-f]+) " + r"file_addr=libcuda\+(?P0x[0-9a-f]+) size=(?P\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 callees [depth] [max_per_node] +# callwalk.py callers [depth] [max_per_node] +# callwalk.py chain # 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_vaddrname) 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 ELF/version summary +# exports [--json] exported symbols (file vaddrs) +# functions [--json] function ranges from .eh_frame_hdr +# nearest nearest preceding function/export for a vaddr +# crash dmesg mapping base + ip -> file vaddr + symbol +# strings [--grep RE] defined strings with file offsets +# delta interface-level diff between two versions +# wrappers export stub -> generated wrapper target +# note 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( + "> 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(" 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("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 : ") + if prev: + print(f"nearest export : {prev['name']} + 0x{target - prev['value']:x} " + f"(export at 0x{prev['value']:x})") + else: + print("nearest export : ") + + +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): +// [] +// Lines starting with '#' and empty lines are ignored. Addresses are file +// virtual addresses ("libcuda+0x..."), not Ghidra image-base offsets. +// +// Usage (analyzeHeadless): +// -process -noanalysis \ +// -postScript ApplyNameMap.java [--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 [--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 -readOnly -noanalysis \ +// -postScript BatchDecompile.java [timeoutSecs] +// +// Outputs: +// /buckets/bucket__.c decompiled C, functions in order +// /index/bucket__.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 [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 -process -readOnly -noanalysis \ +// -postScript DecompileAt.java [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 [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 \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 -noanalysis \ +// -postScript ExportFunctionIndex.java +// +// Outputs (all TSV, one row per item, with a header row): +// /functions.tsv addr size name namespace is_exported +// is_external is_thunk has_body entry_bytes +// /exports.tsv addr size name +// /imports.tsv addr name library +// /blocks.tsv name start end size r/w/x +// /callgraph.tsv caller_addr callee_addr call_sites +// /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 "); + 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 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 -process -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 /wk/ and +# round-robins the chunks across the copies. Copies are kept (they make +# resuming cheap) and can be removed with: rm -rf /wk +# +# Usage: +# ghidra_batch_decompile.sh \ +# [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 +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 [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 "$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 [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 [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 +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_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 [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 [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 +set -euo pipefail +OUTDIR="${1:?usage: $0 }" + +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" -- cgit v1.2.3