NanoVDB: chain DeviceBuffer::recordUse and expose orderAfterPriorUses (CUDA)

[swahtz]

Summary

Two small, coupled changes to cuda::DeviceBuffer's use tracking, prompted by review discussion on #2225's recordUse bindings.

1. recordUse now chains across streams (bug fix). The buffer tracks one event per device, and cudaEventRecord moves an event. So with uses recorded on streams A then B, B's record discarded the only coverage of A's in-flight work, and the device free (destructor / clear() / move-assignment) could run while A was still using the buffer. This violated the invariant already documented on mEvents ("every use waits on this event before issuing work and re-records it afterwards") — the internal upload/download paths follow it, but the public recordUse skipped the wait half. recordUse now issues cudaStreamWaitEvent(stream, event) before re-recording, so the single event transitively covers every recorded use. Side effect (documented): work subsequently issued on the recording stream is also ordered after the previously recorded use — for a shared buffer that is the expected visibility ordering.

2. orderAfterPriorUses is now public (API). It is the consume-side companion of recordUse: an external consumer (e.g. the Python bindings' __cuda_array_interface__ / DLPack exports in #2225) calls it with its own stream before reading, so it cannot observe a partially-written buffer after an asynchronous upload. It was previously private, leaving no public way to order a consumer stream against the tracked uses.

Testing

• New TestNanoVDBCUDA.DeviceBufferChainedRecordUse regression test, modeled on the existing DeviceBuffer*FreeOrdering harness: a non-blocking stream is parked behind a busy-wait with a recorded late write, a second idle stream records a use after it, and the buffer is destroyed. Verified A/B: fails before the fix ("block recycled: true", the recycled allocation is clobbered by the late write) and passes after.
• Full nanovdb_test_cuda suite passes on a Blackwell (sm_120) GPU (the one pre-existing exception, UnifiedBuffer_IO, requires a fixture written by the host suite's GridCountAndIndex and passes when run after it — unrelated to this change).

🤖 Generated with Claude Code

[Copilot]

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

[harrism]

Solid fix — the event-move hazard is real, the chaining mechanism is correct (verified the never-recorded-event no-op, cross-device wait legality, first-creation branch, and per-device chain independence), and the A/B-verified regression test is exactly the right discipline. Three comments inline: one design trade worth acknowledging (concurrent-reader serialization), one test-robustness suggestion (vacuous pass when the block isn't recycled), and an optional style nit.

[harrism]

The fix is right for the motivating case (write on A, then record on B), but note it also changes recordUse from an observer into an ordering mutator for concurrent readers: two readers recording uses on streams A then B were previously independent, and now B serializes behind A's capture point. The single-event design can't distinguish read‖read from write→read, and the doc note's "later consumers observing earlier writes is the expected ordering" only covers the latter. Correctness over concurrency is the right trade for a tracking convenience — but do the #2225 bindings (or other consumers) fan out concurrent readers today? If so it might deserve a sentence in the doc note; if a profile ever surfaces this, the fix would be a read/write-separated or per-record event scheme rather than a revert.

[swahtz]

Agreed on the trade, and to your question: no — the #2225 bindings never call recordUse internally. The CAI/DLPack exports call orderAfterPriorUses (the consume-side wait), which doesn't record and so doesn't serialize anything; recordUse is only ever user-invoked from Python. So the read‖read serialization is strictly opt-in today, and the cost is one event-wait on the recording stream. I've extended the doc note in 5237e3e to state the concurrent-reader serialization explicitly and name the upgrade path (read/write-separated or per-record events, not a revert).

[harrism]

The detection here depends on the freed block being recycled into victim; when the allocator doesn't recycle (different pool state, or a device without pool support), clobbered is 0 and the test passes without having tested the chain. It reports recycled in the failure message but never requires it — so silence is indistinguishable from success. Inherited from the FreeOrdering harness family, so arguably out of scope here, but a one-liner like if (!recycled) GTEST_SKIP() << "allocator did not recycle the block; ordering not exercised"; would make a vacuous run visible instead of green.

[swahtz]

Adopted in 5237e3e, with one check first: I wanted to be sure non-recycling couldn't be a consequence of the fix (in which case the skip would fire precisely on success). Probed on sm_120 across repeated trials post-fix: recycled=1, stillPending=1 every time — the stream-ordered pool hands the same block back with the dependency attached rather than refusing to recycle, so on a working platform the skip never fires and the clobber detection is genuinely exercised. On a non-recycling platform it converts a silent vacuous green into a visible SKIP. I scoped it to the new test; happy to apply the same guard to the two inherited FreeOrdering tests as a follow-up if you'd like.

[harrism]

Optional style nit: this publishes the method by carving a public:/private: island inside the private section. Moving the method down to the public block keeps the access sections contiguous, which is the pattern elsewhere in the file.

[swahtz]

Done in 5237e3e — moved into the public section directly above recordUse, so the produce/consume halves of the tracking API sit together and the access sections stay contiguous.