Struct AtlasCudaBackend 

pub struct AtlasCudaBackend { /* private fields */ }

Production GPU backend wrapping AtlasRegistry + raw CUDA driver API.

Owns this model’s kernel modules. The registry used to be a process singleton reached through AtlasRegistry::get(); it is now loaded per model and propagated from here, so a swapped-in model cannot run the previous model’s kernels. Dropping the last backend unloads them.

Implementations

impl AtlasCudaBackend

pub fn live_bytes(&self) -> usize

Total live device bytes this backend has allocated and not freed.

pub fn alloc_report(&self, top_n: usize, min_mb: usize) -> String

Human-readable attribution of live device memory, biggest site first.

Aggregated by allocating call site rather than by pointer: one site looping over 48 SSM layers is one line reading 9.7 GB across 48 allocations, which is the shape that makes an over-sized pool obvious. Sites below min_mb are folded into a remainder line so the report stays readable while still summing to the true total.

impl AtlasCudaBackend

pub fn new( ordinal: usize, ptx_modules: &[(&'static str, &'static [u8])], ) -> Result<Self>

Initialize the CUDA backend on the given GPU ordinal.

Loads the provided PTX modules for THIS model. Use atlas_kernels::ptx_for_model() or ptx_modules() to obtain the correct module set. Each call produces an independent module set — the CUDA context and stream are shared, nothing else is.

pub fn poison_redzones(&self, lo: usize, hi: usize) -> Result<()>

Re-poison every guard band: 0xEE for zones whose creation index is in [lo, hi), 0x00 for all the others.

The BISECTION half of the A55 red-zone hunt. The zones themselves never move, so every call leaves the device heap byte-for-byte identical and only the CONTENTS of the guard bands change — which is exactly the variable the read detector proved matters. Narrowing [lo, hi) until the completion flips names the allocation being read past.

pub fn scan_redzones(&self) -> Result<usize>

Read every guard band back and report the ones that no longer hold the fill byte.

Returns the number of violated zones. Each violation is logged with the allocation’s creation index, its size, and the first byte of the pad that changed — the size is what identifies the buffer (cross-reference the arena sizes in BufferSizes), and the offset is how far past the end the writer reached.

RE-FILLS every violated zone before returning, so a repeat offender is reported once per scan rather than once and then forever.

pub fn sweep_unreleased(&self) -> usize

Free every allocation this backend made and nobody released.

The backstop for allocations no ModelResource covers — chiefly the loaders’ fused weights, which are owned by layer structs rather than by any pool. Returns how many were reclaimed; since 2026-08-19 the ledger also carries each one’s size and call site, so the sweep can say how many BYTES had no owner and name the sites they came from instead of only counting them. A non-zero count after a clean teardown is a leak, and the log line now points at the code that made it.

Runs LAST in teardown, after every ModelResource::release, so it only ever sees what those missed — and each free here has already been removed from the ledger by forget_alloc, so it cannot double-free.

pub fn registry(&self) -> &Arc<AtlasRegistry>

Trait Implementations

impl Drop for AtlasCudaBackend

Last-resort reclamation for a backend that never reached model teardown.

A load that FAILS part-way leaves whatever it had already allocated on the ledger, and no Model is ever built to tear down. On a hot-swap that memory is not merely leaked, it is actively harmful: the outgoing model is already gone, and the restore then loads into a budget the dead attempt is still holding. That is not hypothetical — a 35B swap failed at kernel selection and the 27B restore died with “only 14.08 GB remains but 17.38 GB is needed”, leaving the server with no model at all.

On the normal path this frees nothing: Model::teardown drains the ledger first, so the sweep finds an empty set. Freeing here is the safe case described in atlas_core::scope — nothing is allocating against a backend that is being dropped.

fn drop(&mut self)

Executes the destructor for this type.

impl GpuBackend for AtlasCudaBackend

fn alloc(&self, bytes: usize) -> Result<DevicePtr>

Allocate bytes of device memory.

fn scan_redzones(&self) -> Result<usize>

A55 diagnostic: read every allocation’s trailing guard band back and report the ones a kernel wrote past. Returns the violation count. Ok(0) when ATLAS_REDZONE is unset or the backend has no red zones — every backend but CUDA.

fn poison_redzones(&self, lo: usize, hi: usize) -> Result<()>

A55 bisection: poison guard bands [lo, hi) with 0xEE and the rest with 0x00. Layout-preserving by construction — nothing is allocated, moved or resized.

fn alloc_managed(&self, bytes: usize) -> Result<DevicePtr>

Allocate managed (unified) memory. On GB10, this allows over-subscribing physical GPU memory — Linux pages overflow to NVMe swap automatically. Managed memory is slower than device memory but avoids OOM.

fn free(&self, ptr: DevicePtr) -> Result<()>

Free device memory.

fn live_bytes(&self) -> Option<usize>

Live device bytes this backend has allocated and not freed, if it tracks them. None for backends with no ledger (mock/CPU).

fn alloc_report(&self, top_n: usize, min_mb: usize) -> Option<String>

Attribution of live device memory by allocating call site, biggest first. None for backends with no ledger.

fn sweep_unreleased(&self) -> usize

Free every allocation this backend made that nobody released, and report how many there were.

fn copy_h2d(&self, src: &[u8], dst: DevicePtr) -> Result<()>

Copy from host to device.

fn copy_d2h(&self, src: DevicePtr, dst: &mut [u8]) -> Result<()>

Copy from device to host.

fn copy_d2h_on_stream( &self, src: DevicePtr, dst: &mut [u8], stream: u64, ) -> Result<()>

Synchronous device-to-host copy ordered after work on stream.

fn copy_d2h_async( &self, src: DevicePtr, dst: &mut [u8], stream: u64, ) -> Result<()>

Async device-to-host copy (no stream synchronization).

fn copy_d2d(&self, src: DevicePtr, dst: DevicePtr, bytes: usize) -> Result<()>

Copy device to device.

fn launch( &self, func: KernelHandle, grid: [u32; 3], block: [u32; 3], shared_mem: u32, stream: u64, params: &mut [*mut c_void], ) -> Result<()>

Launch a kernel on the given CUDA stream.

fn stream_is_capturing(&self, stream: u64) -> bool

Whether stream is inside an active CUDA-graph capture. Telemetry taps MUST check this before any sync/D2H on a potentially-captured stream — those calls invalidate the capture (CUDA 901) and wedge the serve. Default false (backends without capture, or without a query API, never capture through this trait’s eager paths).

fn synchronize(&self, stream: u64) -> Result<()>

Synchronize a CUDA stream (blocks until all work completes).

fn default_stream(&self) -> u64

Get the default stream handle.

fn op_cache(&self) -> &OpCache

This backend’s memoized kernel handles and scratch allocations.

fn debug_sync_kernels(&self) -> bool

Synchronise the stream after every kernel launch, so an asynchronous illegal-address fault is reported at the kernel that caused it rather than at a later sync. Resolved once when the backend is built; read on the launch path, which is why it is not a per-launch getenv.

fn kernel_registry(&self) -> Option<Arc<AtlasRegistry>>

This backend’s model-scoped kernel modules, for the few callers that need the registry itself rather than a kernel handle — resolving a __device__ symbol, for instance. None on backends that have no such concept, which is why it is an accessor rather than a downcast.

fn kernel(&self, module: &str, func_name: &str) -> Result<KernelHandle>

Look up a kernel function by module and function name.

fn copy_h2d_async(&self, src: &[u8], dst: DevicePtr, stream: u64) -> Result<()>

Async host-to-device copy: src may be dropped or overwritten the moment this returns.

fn copy_h2d_async_retained( &self, src: &[u8], dst: DevicePtr, stream: u64, ) -> Result<()>

Async host-to-device copy for a source the CALLER keeps alive.

fn copy_d2d_async( &self, src: DevicePtr, dst: DevicePtr, bytes: usize, stream: u64, ) -> Result<()>

Async device-to-device copy (no stream synchronization).

fn copy_d2d_2d_async( &self, src: DevicePtr, src_pitch: usize, dst: DevicePtr, dst_pitch: usize, width_bytes: usize, height: usize, stream: u64, ) -> Result<()>

Strided device-to-device 2D (pitched) copy: height rows of width_bytes, source rows spaced by src_pitch, dest rows by dst_pitch. Default = per-row copy_d2d_async loop; the CUDA backend overrides with ONE cudaMemcpy2DAsync (replaces the per-token Z-copy loop = up to num_tokens×num_ssm_layers launches/forward).

fn begin_capture(&self, stream: u64) -> Result<()>

Begin capturing CUDA operations on stream into a graph.

fn end_capture(&self, stream: u64) -> Result<GraphHandle>

End capture and return an instantiated graph ready for replay.

fn abort_capture_if_active(&self, stream: u64)

Best-effort: if stream is mid graph-capture, end that capture so the stream returns to normal mode (discarding any partial graph). Call this on an error path that unwound out of a begin_capture/end_capture region (e.g. a fold refuse bailed mid-capture) — otherwise the stream is left recording and every subsequent op fails with STREAM_CAPTURE_UNSUPPORTED, bricking the server. No-op if not capturing.

fn launch_graph(&self, graph: GraphHandle, stream: u64) -> Result<()>

Replay all operations captured in the graph on stream.

fn destroy_graph(&self, graph: GraphHandle) -> Result<()>

Destroy an instantiated graph, freeing resources.

fn memset(&self, ptr: DevicePtr, value: u8, bytes: usize) -> Result<()>

Set device memory to a byte value (synchronous — waits for completion).

fn memset_async( &self, ptr: DevicePtr, value: u8, bytes: usize, stream: u64, ) -> Result<()>

Set device memory to a byte value on the given stream (async — does not wait).

fn total_memory(&self) -> Result<usize>

Total device memory in bytes.

fn free_memory(&self) -> Result<usize>

Free device memory in bytes.

fn device_free_memory(&self) -> Result<usize>

Free device memory as the DRIVER reports it, with no host leg.

fn live_alloc_count(&self) -> usize

Live (allocated, not yet freed) device allocations on this backend.

fn sm_count(&self) -> Result<u32>

Number of streaming multiprocessors (CUDA SMs / HIP CUs) on the device.

fn create_stream(&self) -> Result<u64>

Create a new CUDA stream (for overlapping work).

fn bind_to_thread(&self) -> Result<()>

Bind the CUDA context to the current thread.

fn create_event(&self) -> Result<u64>

Create a CUDA event (for inter-stream synchronization).

fn record_event(&self, event: u64, stream: u64) -> Result<()>

Record an event on a stream (marks a point in the stream’s work).

fn stream_wait_event(&self, stream: u64, event: u64) -> Result<()>

Make a stream wait for an event (GPU-side sync, CPU does not block).

fn event_synchronize(&self, event: u64) -> Result<()>

Block the calling host thread until all work already recorded against the event — e.g. an async D2H copy issued on the graph stream followed by record_event, then event_synchronize right before the host dereferences the destination pinned buffer. Cheaper than synchronize(stream) when the stream has work beyond the event you care about: this only waits for the recorded point, not for everything subsequently enqueued.

fn destroy_event(&self, event: u64) -> Result<()>

Destroy an event.

fn host_ptr_to_device(&self, host: *mut u8) -> Result<DevicePtr>

Device-side alias of a page-locked host pointer from Self::alloc_host_pinned (cuMemHostGetDevicePointer). On UMA parts (GB10) this lets a KERNEL write results directly into host-visible memory, eliminating the copy-engine op for tiny readbacks entirely. Default: unsupported.

fn alloc_host_pinned(&self, bytes: usize) -> Result<*mut u8>

Allocate page-locked (pinned) host memory for efficient async H2D.

fn free_host_pinned(&self, ptr: *mut u8, _bytes: usize) -> Result<()>

Free page-locked host memory previously allocated by alloc_host_pinned.

fn launch_typed( &self, func: KernelHandle, grid: [u32; 3], block: [u32; 3], shared_mem: u32, stream: u64, args: &[KernelArg<'_>], ) -> Result<()>

Typed-args kernel launch.