TrainerRank reusable-cache admission can leave insufficient physical CUDA memory for backward library allocations

[bradhilton]

Current status — September 17, 2026: partially addressed; still open.

Brad approved PR #900, merged as c009557e2f2488a83551d29c3253b708edef322c. Admission now uses physical free CUDA memory and can release unused allocator cache under a measured cost budget. The scoped two-H200 qualification and source/CI checks passed before merge. This does not establish full-training overhead, sufficient memory throughout arbitrary backwards, or adoption by every Caladan experiment/image.

The earlier PR #888 is a separate, still-held proposal to release cache at a low-headroom handoff to backward. It is not adopted by #900, and its process-wide effects and throughput tradeoffs remain relevant. It should not be merged or relabeled as superseded without reassessing that remaining boundary against the current admission policy.

Active work with #848: qualify the current code on the retained workload and resolve underestimated backward demand. Keep the library-allocation problem separate from the repeated-gradient variability tracked by #902. All recent completed diagnostic resources are cleaned up. Representative end-to-end overhead and broader backward-memory safety remain unqualified.

The historical September 12 status below is preserved; its description of the then-current candidate is superseded by this update.

---

September 12 update: the earlier statement that conditional/full-step and multi-device behavior are entirely unqualified is superseded by bounded native evidence. Three complete updates on one H200 exercised four actual conditional releases; a separate two-H200 run verified that selected-device release also trims the other device's unused cache while preserving the checked live values/gradients and caller device/streams. All resources were independently retired. Exact retained evidence and limits: /home/brad/.local/share/schulman/art870-two-device-gap-inventory-20260912-capture/REPORT.md.

ART #888 remains OPEN/DRAFT at1763a46ac9b4f465e38440baf8cc3e7922da69a3, source-correctness reviews and CI green, adoption HELD. The 3% trigger is best effort, not a guaranteed reserve. Process-wide cache retirement/waits are significant behavior requiring Brad's decision. Representative throughput/refill costs, arbitrary concurrent users, and broader distributed/negative cases remain unqualified. Do not repeat the completed two-device witness merely to restate the already observed process-wide effect.

---

Earlier record (preserved; status above is current):

Status reconciliation — September 11, 2026 (Schulman)

Confirmed allocation failure; local mitigation only. Unconditional cache release let the pinned first-backward diagnostic finish, but the proposed conditional trigger, complete optimizer behavior and multi-device side effects remain unqualified. torch.cuda.empty_cache can trim caches across devices in the process; a selected-device context does not isolate that effect.

Planned lane: Schulman and subagents, queued with #848/#256. Complete the conditional/multi-device and full-step qualification before proposing a generic fix. The held local patch is not merged, does not modify art.megatron, and is not covered by a no-significant-behavior assumption.

---

Historical report (preserved):

Type: confirmed backward allocation failure with a controlled experiment-local mitigation; generic policy remains unresolved.
Agent owner: Schulman, related to #848.

A rank-1 Qwen3.6-35B-A3B eight-pair dynamics batch on one H200 passed its forwards, then failed in micro_batch_loss.backward with CUBLAS_STATUS_ALLOC_FAILED from cublasCreate(handle). GeneratorExit was secondary iterator teardown. The exact internal cuBLAS request size was not measured.

At the first gradient admission, live allocated/reserved memory was about 74.662/147.394 decimal GB with only 0.034734 GB physically free. The planner treated reusable allocator slack as available; its 3% capacity subtraction did not physically reserve CUDA memory for external library allocations. At the caller boundary, an unconditional torch.cuda.empty_cache() released 42.754638 GB without changing live allocation. The matched release-only replay completed all five backward waves and the optimizer; the observe-only replay failed. Pre-intervention plans/budgets and corresponding forward counters matched exactly.

The qualified mitigation is experiment-local cache release at the start of every gradient caller, after receiving its microbatch. It is not equivalent to a conditional 3% policy, a pre-yield trainer hook, or handle warmup. Later nonzero-generator-gradient updates also succeeded, but larger traces and full validation remain unqualified. Cold forward estimate underprediction is a separate problem.

Investigate a bounded generic physical-headroom/library-allocation contract with mixed-grad plans, device/stream behavior and normal downstream workloads. Prefer trainer_rank-local work; no art.megatron change is proposed. Do not silently globalize the workaround based on one fixture.

Evidence: /home/brad/.local/share/schulman/retail49-memory-failure-20260909/MORNING-ALLOCATOR.md and replay-{observe,release}-comparison/. Local Caladan mitigation 431e424 and timing-only follow-up 177a9fce. Raw failures and cleanup receipts are preserved.

[bradhilton]

The new three-arm cold eight-pair comparison is complete and all owned resources are independently absent. Observe-only again fails at first backward with cuBLAS handle allocation, with41,615,360 bytes physically free. Both caller-release and trainer pre-yield release complete all five backwards plus optimizer with matching finite metrics and byte-identical auxiliary head outputs; first G gradient is zero by the initialized objective.

Both placements release42,754,637,824 bytes of reserved cache and raise physical free memory to about42.8GB. There is no demonstrated pre-yield advantage on this fixture, and the tested endpoints do not determine a safe generic threshold or establish a general library-allocation contract.

Terminal harness failures remain preserved: the caller arm completes its update, then a standard JSON boundary rejects the cal.run dict proxy; the pre-yield arm completes, then trainer.close exceeds60s. A CPU reproduction shows explicit dict(result) is appropriate at that JSON serialization boundary. External UID/service/process cleanup succeeds for all three arms. These are not clean exit0 runs.

Evidence: /home/brad/.local/share/schulman/art-headroom-arms-20260909d/REPORT.md, SHA256 d11ac00d7bdce959ef4a56636c20935bd9aff7317ec2753528097aea262a1831; sealed evidence manifest 35f88e971b5daa8092782458fe793cc3316ab8ad80a9c116e1fa9271fc9d5d2a. The existing049 caller release remains in main. Generic trainer_rank policy remains unresolved; no art.megatron change or new launch.

[bradhilton]

The new bounded V7 observe run reached the exact first backward and preserved a finalized13,362,126-byte Nsight report (SHA25673948f6b4d55da54da4fcfe42be029a8bcf11623246b846a0c197750a7985001). Independent read-only SQLite analysis finds a concrete allocation failure: cudaMalloc_v3020 return2, correlation1293703, wholly inside cublasCreate_v2 on the same worker thread. The owning process is actorPID9568; the worker is TID37872, pt_autograd_0. It falls50.294ms after admission3's backward_start and5.575ms before its RuntimeError boundary. The main telemetry thread9568 is distinct; the association uses exact process namespace, same-worker CUDA/cuBLAS nesting and checked monotonic-clock alignment, not an invented CUDA admission label.

Admission3 reported required7,976,316,880bytes against available68,242,798,285bytes. At handoff, PyTorch allocated92,197,523,968bytes/reserved147,394,134,016bytes while physical CUDA free was14,352,384bytes. No backward completed. This supports physical allocation pressure during cuBLAS creation on this exact fixture, beyond the exception classifier alone. The failed request size, callchain and cuBLAS return-status field are absent from this export; some CUDA/NVTX events may be missing, and hardware tracing fell back to software. The exported memory table has allocation records only, so no complete allocation/free ledger or generic safe-reserve threshold is claimed.

Report /home/brad/.local/share/schulman/art870-v7-trace-independent-20260910/REPORT.md SHA252165e49246d0f8a613e74705ef87090c5857201749ae90d1693c7664922cc5; manifest23e50cf1ecd73e69edf4680c5f3ec86cf4263a76fa7040dcf159026519e0ebf1. All10 original/copy hashes, SQLite integrity, owner-token/transfer and timestamp/identity joins pass. No live-resource operation was performed for this analysis. A same-source release-mode intervention is the next bounded comparison under root registration; it would qualify that intervention on this fixture, not establish a general reserve policy.

[bradhilton]

#870 release-d: per-file trace cap censored the run

The original retained companion exception is Trace file byte budget exceeded, from snapshot()'s regular-file st_size > max_file_bytes check.
It is not a CUDA OOM, an inferred safety margin, or a 90% threshold. The configured
per-file cap is 1,073,741,824 bytes. The exact offending relative path and size
were not included in the failure response, so they cannot be reconstructed here.

trace-guard.json is the last successful poll at 3:42:33 AM Mountain:
1,027,608,576 largest-file bytes, 1,329,111,040 aggregate accounted bytes and
1,071 entries. The next poll threw before either observation receipt was updated.
The successful snapshot's remaining space does not describe the failing poll.
The original companion traceback joins the actual launched V9 source through
companion.py:194 → trace_guard.check:470 → remote_snapshot:428 → snapshot.

Identity: source manifest 0f89ce6047d3436d26f0194a8cdce0fb07d760c0aff048c0528ce773e4f5bc84,
release mode, root registration260707, exact pod UID
07fb3420-0c0d-406b-a968-9e1c4d05bbbe, actor PID6986, ART7ea4db7b implementation
044c9953…, Caladan99441ee1 and immutable ecf9 image. The full input manifest
records original paths, sizes and hashes without copying raw driver logs or
touching the frozen run. Uploaded/runtime guard bytes equal the launched source.

Phase evidence is limited:

Durable boundary	Mountain time
Root launch receipt	3:30:25 AM
Ready actor/image and actor-source receipt	3:36:39 AM
Collector installation receipt, actor6986	3:36:40 AM
Measured-kernel receipt after driver load/prefetch return	3:37:06 AM
Last successful trace poll	3:42:33 AM
Companion begins cleanup after original exception	3:42:38 AM
Supervisor records driver exit−9	3:42:39 AM
Companion records verified cleanup	3:42:48 AM

Receipt times above are durable file timestamps where no event timestamp exists;
they are not precise actor-operation timing. kernel.json is written immediately
before the queued measurement await. It proves the driver reached that boundary,
not that the actor entered or finished the function. No retained local record
proves prepass completion, capture_start, gradient handoff/cache release, or
backward start/completion. No first-backward.json, failure/closing terminal or
completed entrypoint receipt exists. Cicero confirms this V9 composition has
terminal-only actor telemetry transfer; no separate actor capture is owned by
him. An unrelated successful observe trace cannot fill this attempt's missing
phase join. No model OOM is present in the retained evidence; its absence is not
proof of what the forcibly stopped actor was doing.

The retained canonical cleanup receipt reports no remaining pods, services or
processes, sky_down_exit=0, and no cleanup errors. Root independently owns that
resource verification. This auditor made no Kubernetes, actor, queue or resource
calls and did not independently relaunch/revalidate a live object.

An unchanged-source CPU witness uses tiny owned scratch files and mocked GETs.
A 900-byte file passes a 1,024-byte cap; growth to 1,025 bytes raises the exact
error and leaves the previous durable guard receipt byte-identical. The scratch
was removed. It demonstrates the stale-success-receipt behavior without creating
a GPU workload, large file, Kubernetes object or source patch.

Smallest defensible next counterfactual

First make the existing bounded observation path survive another censor:
persist primitive collector events incrementally under the exact owner, and have
the existing companion read bounded summaries. At a file-budget failure, retain
the owned relative path, size, cap, partial inventory and failure timestamp
before terminating. Preserve error identity and all caps, UID/birth checks,
deadlines and canonical cleanup. This needs no new actor queue or supervisor.
The existing collector already limits records to256 and each record to16KiB;
any incremental response must also fit the monitor's64KiB response bound.

Then consider a new paired observe/release capture range from native gradient
handoff through the first backward, keeping the 1GiB/2GiB caps and existing
CUDA/cuBLAS trace switches. Start profiling in a source-gated instance wrapper
at entry to the existing _test_gradient_handoff_cache, before its before
memory snapshot. Starting inside _telemetry_phase would be too late: that
snapshot already exists, and profiler startup could contaminate the reported
before/after cache-release difference. Stop after first backward under the
existing ownership/error handling. Prepass and gradient forward still execute
unchanged, and original candidate admission remains captured; their detailed
CUDA trace would be excluded.

This is a scope-reduction hypothesis, not proof it will fit: the failed file's
type/contents are unknown and backward alone may still exceed the cap. It requires
CPU ordering/source/ownership/partial-failure tests and a fresh registered pair
with matched input/source/seed and actual admission/layout/state checks. A new
shorter-range release cannot be treated as an identical-instrumentation pair
with old observe-e. No optimizer, publication, arbitrary warmup or cap increase
is needed for the proposed diagnostic.

Dropping --cuda-trace-all-apis=true is another possible reduction, but not yet a
coverage proof for the required allocation/library APIs. NVIDIA documents added
overhead from tracing every CUDA API and from memory-usage collection; exact
staged CLI help confirms those flags. Prefer the narrower range first, retaining
the target release/backward evidence. NVIDIA Nsight Systems User Guide

No implementation or retry is authorized by this report. Root retains scope,
independent review, source sealing, capacity and launch decisions.

[bradhilton]

A separate prepass allocation failure is now documented for follow-up under the physical-headroom investigation. Fresh residual attempt J completed both no-grad forwards and all eight captured/tensorized trajectories, then failed in the unchanged rank.dp_reduce(statistics) before any gradient forward or successful update. NCCL reports Failed to CUDA calloc async 24 bytes; this message alone does not establish CUDA OOM, the failed allocation’s underlying cause, or a generic memory-policy defect.

The preceding caller boundary records 75,600,376,320 bytes allocated, 147,463,340,032 reserved and 381,353,984 physically free. The larger admitted prepass used 228,142 packed / 928,901 logical tokens across seven pairs and 79 requests; its original required/available check was 68,177,016,880 / 69,433,447,117 bytes. Measured incremental forward peak was 46,946,880,000 bytes, below that estimate. These measurements do not demonstrate forward underprediction. Caller-owned last-microbatch tensors were still live.

The evidence is frozen at /home/brad/.local/share/schulman/research-j-prepass-failure-20260910-harvey, manifest 01d2dff0a5bf13a8b393317858309e0f3d6272091c170c11e83641e869f3360b. Exact run/pod/source/queue joins, original checks, all raw captures and errors remain private. Root independently verified the three process births/groups, pod and two service UIDs, and Sky owner absent after cleanup.

A fresh, unarmed local L candidate adds exactly one torch.cuda.empty_cache() immediately before the unchanged reduction. It retains all caller tensors and collectives and restores byte-for-byte to J after deleting that line. Independent review is in progress; 22 root CPU source/native/binder controls pass. It is an unqualified local mitigation, not a generic ART change. Fresh generated inputs will differ, so a successful future L would not be an exact replay of J. The cuBLAS first-backward trace and its separate observe/release experiments remain distinct evidence.

[bradhilton]

Schulman, September 10 morning update. This issue remains open.

The completed V11 observe-g trace joins a failed cudaMalloc_v3020 return 2 inside cublasCreate_v2 on the same actor and backward thread. The full 3,074,756-byte Nsight report has SHA256 3229d2f4358fde4769f930234f0d45968dba38891ac407fe0d634a2478433411; all thirteen primitive events were returned. Original admission: predicted 7,976,316,880 bytes against 68,242,798,285 available. Handoff: 92,197,523,968 allocated, 147,394,134,016 reserved, only 14,352,384 physically free. The failed allocation's requested byte count is unknown. The trace records an early failure and does not qualify the five-second capture-stop path.

The earlier cache-release arm freed 42,754,637,824 bytes of reserved memory with allocated tensors unchanged, but was censored by trace/disk guards before any backward-end result. That is neither a successful backward nor an OOM result. Same-source V11 release-f then correctly refused before driver launch because free disk was below its unchanged 23-GiB starting requirement. Release-g was registered but never granted/launched; release-h remains unarmed. By the time staging and disk gates cleared, the remaining work window was shorter than the observed V11 launch-to-backward time of 610.86 seconds, before accounting for release completion and drain. This empirical timing is not a universal minimum. No deadline or budget was relaxed.

All allocated attempts have independent exact pod/service UID, PID-birth/process-group and hosted Sky cleanup receipts. The generic physical-headroom policy remains unqualified; no generic reserve threshold was adopted. The separate prepass reduction failure (NCCL allocation of 24 bytes) is preserved as another boundary, without asserting that it proves the same cause.

Evidence: /home/brad/.local/share/schulman/art870-v11-observe-audit-20260910-carver/REPORT.md, SHA256 16570fa9ef434fab4b339d7d42ed7e7b3e6a932aafb785dffe3048fcf4dddb2c; older release-censoring and raw evidence remain under the corresponding owned attempt roots. Raw traces/logs remain private. Next useful step is a fresh sufficiently bounded same-source release qualification, followed by a separately reviewed generic fix.

[bradhilton]

Schulman — September10 matched V11 release result. This issue remains open.

The fresh same-source release-j completed the first backward after the exact original admitted candidate and before-state matched observe-g. Candidate3:98,647 packed /345,213 logical tokens; required7,976,316,880 B, available68,242,798,285 B. Handoff before:92,197,523,968 allocated /147,394,134,016 reserved /14,352,384 physically free. One native pre-yield cache release freed42,754,637,824 B with allocated unchanged; physical free became42,768,990,208 B. The backward completed169.115 seconds later with no error.

The finalized6,922,681-byte raw trace (7af247c7407348b05fd27d99ab917647216c3b62e5687531805d58f31b05c107) exports cleanly. On the same release actor/autograd worker, cublasCreate_v2 encloses three successful cudaMalloc_v3020 calls; exact allocation-event joins record1,024 B,131,072 B and67,108,864 B. Observe-g had a failed CUDA allocation inside that library call. All175,876 recorded release CUDA API calls return0. This trace covers the initial5-second library-allocation window; primitive telemetry continues through backward. The checked profiler-stop path succeeds after the automatic5-second stop, followed by finalized trace drain. Nsight's incomplete-event warnings remain preserved; this is not an exhaustive full-backward trace or throughput qualification.

Lifecycle remains a separate failure: after result/trace returned, close removed the pod and a companion read reported404. Companion then stopped its supervisor/driver and verified cleanup. Root independently confirmed the exact pod, two service UIDs, three PID births/groups and Sky owner absent. This is successful bounded diagnostic work with a shutdown harness failure, not a clean end-to-end run.

Source/physics remained exact V11 45384beee0fbe5ac44553f8987f3d724ae756c847a4eb4447c238c61d6354ffd, ART7ea4db7b/Cal99441ee1/imageecf9, seed0. A fresh reviewed configuration used the existing zero shared-disk-delta-budget mode after release-i's global filesystem-growth guard fired;20GiB sampled free floor,16GiB sampled RSS,trace1GiB/file+2GiB aggregate and20+15-minute lifecycle remained. This is not a per-owner disk quota. Earlier attempts remain unchanged.

Evidence: /home/brad/.local/share/schulman/art870-v11-release-j-audit-20260910-carver/REPORT.md, SHA256 6aac97680648de11784618a8eb1ab03dcce6724ffaaa28318907735bd2c7bf58; manifest 633e4b4e4a7243aa93d796ab4e97897539e851f20faf23f34e9ed477207b7b58. Raw traces/logs remain private.

This qualifies the release mechanism for the matched fixture, not a generic reserve amount: the failed observe allocation's requested size and minimum sufficient headroom remain unknown. The existing private trainer_rank-only conditional candidate dc04dcf0 is being reassessed as a best-effort materialization of the planner's existing3% reserve. Its thresholds, other boundaries/backends and performance still require qualification; no shared policy, pin, merge or deployment has been adopted. Separate prepass/NCCL and forward-admission issues remain distinct.

[bradhilton]

Schulman — correction to the earlier process-cleanup scope

Today's measured research-directory cleanup found five detached native Sky command sessions (15 local processes) outside the groups recorded by the original qualification drivers and companions. This includes the release-j / observe-g roots discussed here. All recorded process-group/UID checks remain evidence of those exact checks; they did not establish absence of every detached client. The earlier broader local-process cleanup claim was therefore incomplete.

The births/cwds/executables/command hashes join the expired registrations and frozen native command recipes. Two OS snapshots showed continuing CPU/I/O, but no request counts or backend operation/resource attribution follows. Root retired the verified clients with identity-bound pidfds; independent verification at 3:37:33 PM Mountain on September10 found every original PID/birth and all five groups absent. No Sky/API request, allocation or deployment was issued. Current backend state was not re-queried; historical resource checks must not be represented as fresh checks.

The distinct current-source containment gap and required CPU reproduction are now documented in OpenPipe/caladan#336. This is separate from OpenPipe/caladan#298's same-group early-leader-exit correction and OpenPipe/caladan#318/#331's normal-teardown transition. A new-session command can survive other forced driver failures even when normal teardown is corrected.

Evidence: /home/brad/.local/share/schulman/orphan-sky-clients-audit-20260910/REPORT.md (SHA256 029bd1ffdd2586ac5370bdc174c15057acfe32e8d0b3060557b9e9128d230108), exact original/retirement receipts and independent post-retirement-verification.json; root signal journal in sibling orphan-sky-clients-retirement-20260910/.

[bradhilton]

Schulman: ART #870 remains open. The current reviewable private candidate is a9552fa atop accepted/current main 66f644d; dc04dcf0 in the last issue comment is superseded by this native-allocator-only composition. It adds 37 runtime lines entirely within art.trainer_rank and no public API or art.megatron change. Existing author and independent CPU/source reviews are clear, and I verified their exact sources and receipts; no additional source defect was found.

The matched V11 observe/release evidence proves that freeing unused cache allowed the pinned fixture's first backward to finish. It tested an unconditional diagnostic hook, not the proposed conditional 3% trigger. The remaining adoption boundary is concrete: PyTorch native empty_cache trims caches across devices in the process, so selected-device context does not confine its side effects. Actual conditional GPU behavior, multi-device effects, complete training/optimizer behavior and warm/query/refill cost remain unqualified. Observer compatibility must also be reviewed against the changed source hash.

Held patch/PR draft: /home/brad/.local/share/schulman/art870-accepted-composition-20260910-carver/. Independent review: /home/brad/.local/share/schulman/art870-accepted-review-20260910-harvey/REVIEW.md. Current audit and minimum qualification boundaries: /home/brad/.local/share/schulman/art870-candidate-status-audit-20260910/REPORT.md. No GPU/provider calls, trace rehydration, source edits, rollout or policy decision occurred. #848 estimator/profile work stays with Pasteur; the no-grad NCCL prepass failure remains a distinct unqualified boundary.