You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
hipGraphCreate hands the application a pointer to our heap GraphHandle, not a native ihipGraph*. The real graph lives inside it as native_graph. Same for NodeHandle::native_node and ExecState::native_exec.
Every difficulty in redline-hipgraph descends from that one decision:
Any hipGraph entry point we do not export receives our heap pointer and hands it to HIP, which dereferences it as its own type — silent memory corruption in an application that worked before the preload. This is what the 72 shims in hipgraph: shim the 72 unexported entry points that took our handles #5 exist to prevent.
The registry is HashSet<usize> — bare addresses. is_graph/graph_handle/exec_handle/node_snapshot test membership under a mutex, drop it, then dereference the pointer as a fabricated &'static. A concurrent hipGraphDestroy frees it mid-call. Address reuse gives ABA.
The failure mode is inverted from what an interposer under someone else's application should have. Today an entry point we have not thought of corrupts; it should lose acceleration.
Proposal
Return the real native handle to the application; keep the retained PM4 plan in a side table keyed by that native pointer.
Type confusion becomes structurally impossible — the app never holds our pointer.
Unknown entry points degrade to lost acceleration.
The lifetime race dissolves: a side-table miss is a clean miss, not a UAF. No Arc registry needed.
Cost is one hash lookup per intercepted call, against a HIP call. Replay — the hot path — is untouched.
Measured blast radius
Bucketing all 72 shims from #5 (1d2a952, 4017 lines under src/shims/):
bucket
count
fate
DELETE — pure handle translator
23
removed entirely
KEEP — mutates, must invalidate our plan
32
shrinks to side-table lookup + flag
KEEP-CREATE — creates an object we must record
17
shrinks to side-table insert
~620 shim lines delete outright, plus all five translation helpers (~104 lines); the 49 survivors get materially thinner. Version script drops 96 → 73 symbols.
src/shims/introspect.rs deletes entirely. Three of the four symbols llama.cpp remaps via ggml-cuda/vendors/hip.h (GetNodes, NodeGetType, KernelNodeGetParams) would need no shim at all — only KernelNodeSetParams stays. /cc @pwilkin
The one real obstacle
Stream capture. hipStreamBeginCapture does allocate_graph(0) — our graph exists with no native twin, capture-time nodes get native_node = 0, and native identity attaches only at hipStreamEndCapture.
It resolves: during capture the application never holds the graph handle. It first becomes visible at EndCapture, by which point the native graph exists. Provisional state stays keyed by stream (which global().captures already does), and native identity is published at EndCapture — where reconcile_captured_native_nodes already treats the native graph as ground truth.
The other zero-shadow source is degenerate, not a design feature: native_graph == 0 outside capture only happens when real_symbol("hipGraphCreate") returns None, or when a failed native create was swallowed by unwrap_or(0).
Counter-risk
A future unshimmed API that mutates a graph would leave our plan stale and replay something wrong. The answer is only available because of the redesign: with the native graph always authoritative, validate our modelled topology against hipGraphGetNodes at instantiate time and force native on mismatch. Once per instantiate, never per replay.
Sequencing
Merge hipgraph: shim the 72 unexported entry points that took our handles #5 first — it removes a deterministic corruption today. Not wasted: the 72/72 signature audit against hip_runtime_api.h and the readelf-derived version-node mapping carry over, and 49 of 72 shims survive in thinner form.
Redesign as one PR that deletes more than it adds.
Then the fmt/clippy paydown and promoting those gates to required.
Rejected alternative
An Arc registry (HashMap<usize, Arc<GraphHandle>>) closes the UAF and both residual races in about a day, with small call-site churn. Rejected as the primary plan: it buys safety for the design that generates the work, and leaves every future ROCm entry point a corruption risk. Keep it as a stopgap only if #5 must be sound now and the redesign slips.
Testability
None of this needs a GPU. The side table is pure host-side state, so a multithreaded create/lookup/destroy stress test runs in the CI added in #4 — turning "we reasoned about the lifetime" into "the machine checks it on every PR."
Not fixed by this
An application that destroys a graph while another thread uses it stays broken — identically to how it breaks without the preload. That is the correct boundary.
Full design note with file:line evidence: docs/investigations/2026-08-01-hipgraph-handle-ownership-redesign.md (lands with the redesign branch).
Problem
hipGraphCreatehands the application a pointer to our heapGraphHandle, not a nativeihipGraph*. The real graph lives inside it asnative_graph. Same forNodeHandle::native_nodeandExecState::native_exec.Every difficulty in
redline-hipgraphdescends from that one decision:HashSet<usize>— bare addresses.is_graph/graph_handle/exec_handle/node_snapshottest membership under a mutex, drop it, then dereference the pointer as a fabricated&'static. A concurrenthipGraphDestroyfrees it mid-call. Address reuse gives ABA.The failure mode is inverted from what an interposer under someone else's application should have. Today an entry point we have not thought of corrupts; it should lose acceleration.
Proposal
Return the real native handle to the application; keep the retained PM4 plan in a side table keyed by that native pointer.
Arcregistry needed.Measured blast radius
Bucketing all 72 shims from #5 (
1d2a952, 4017 lines undersrc/shims/):~620 shim lines delete outright, plus all five translation helpers (~104 lines); the 49 survivors get materially thinner. Version script drops 96 → 73 symbols.
src/shims/introspect.rsdeletes entirely. Three of the four symbols llama.cpp remaps viaggml-cuda/vendors/hip.h(GetNodes,NodeGetType,KernelNodeGetParams) would need no shim at all — onlyKernelNodeSetParamsstays. /cc @pwilkinThe one real obstacle
Stream capture.
hipStreamBeginCapturedoesallocate_graph(0)— our graph exists with no native twin, capture-time nodes getnative_node = 0, and native identity attaches only athipStreamEndCapture.It resolves: during capture the application never holds the graph handle. It first becomes visible at
EndCapture, by which point the native graph exists. Provisional state stays keyed by stream (whichglobal().capturesalready does), and native identity is published atEndCapture— wherereconcile_captured_native_nodesalready treats the native graph as ground truth.The other zero-shadow source is degenerate, not a design feature:
native_graph == 0outside capture only happens whenreal_symbol("hipGraphCreate")returnsNone, or when a failed native create was swallowed byunwrap_or(0).Counter-risk
A future unshimmed API that mutates a graph would leave our plan stale and replay something wrong. The answer is only available because of the redesign: with the native graph always authoritative, validate our modelled topology against
hipGraphGetNodesat instantiate time and force native on mismatch. Once per instantiate, never per replay.Sequencing
hip_runtime_api.hand thereadelf-derived version-node mapping carry over, and 49 of 72 shims survive in thinner form.Rejected alternative
An
Arcregistry (HashMap<usize, Arc<GraphHandle>>) closes the UAF and both residual races in about a day, with small call-site churn. Rejected as the primary plan: it buys safety for the design that generates the work, and leaves every future ROCm entry point a corruption risk. Keep it as a stopgap only if #5 must be sound now and the redesign slips.Testability
None of this needs a GPU. The side table is pure host-side state, so a multithreaded create/lookup/destroy stress test runs in the CI added in #4 — turning "we reasoned about the lifetime" into "the machine checks it on every PR."
Not fixed by this
An application that destroys a graph while another thread uses it stays broken — identically to how it breaks without the preload. That is the correct boundary.
Full design note with file:line evidence:
docs/investigations/2026-08-01-hipgraph-handle-ownership-redesign.md(lands with the redesign branch).