diff --git a/crates/onnx-runtime-ep-cuda/src/lib.rs b/crates/onnx-runtime-ep-cuda/src/lib.rs index 89c16ccd5e..6f76f34c82 100644 --- a/crates/onnx-runtime-ep-cuda/src/lib.rs +++ b/crates/onnx-runtime-ep-cuda/src/lib.rs @@ -89,8 +89,8 @@ pub use kernels::{CUDA_COVERED_OPS, build_cuda_registry, build_cuda_registry_wit pub use provider::CudaExecutionProvider; pub use weight_paging::{ CudaResidencyStats, CudaWeightPage, CudaWeightPager, CudaWeightResidency, DeviceOffloadPolicy, - GlobalOffloadStats, WEIGHT_OFFLOAD_DEVICE_BYTES_ENV, WEIGHT_OFFLOAD_ENV, global_offload_stats, - reset_global_offload_stats, + GlobalOffloadStats, WEIGHT_OFFLOAD_ASYNC_PAGEIN_ENV, WEIGHT_OFFLOAD_DEVICE_BYTES_ENV, + WEIGHT_OFFLOAD_ENV, global_offload_stats, reset_global_offload_stats, }; /// Number of additional u32 words required by the CUDA device argmax result buffer. diff --git a/crates/onnx-runtime-ep-cuda/src/provider.rs b/crates/onnx-runtime-ep-cuda/src/provider.rs index 2e5bc007dd..bd86c3fc74 100644 --- a/crates/onnx-runtime-ep-cuda/src/provider.rs +++ b/crates/onnx-runtime-ep-cuda/src/provider.rs @@ -81,7 +81,10 @@ impl CudaExecutionProvider { let budget = offload_policy .device_budget_bytes .unwrap_or(DEFAULT_DEVICE_OFFLOAD_BUDGET_BYTES); - Arc::new(CudaWeightResidency::new(runtime.clone(), budget)) + Arc::new( + CudaWeightResidency::new(runtime.clone(), budget) + .with_async_pagein(offload_policy.async_pagein), + ) }); Ok(Self { device: DeviceId::cuda(ordinal), @@ -812,4 +815,185 @@ extern "C" __global__ void copy_out(const float* in, float* out, unsigned long l ep.deallocate(dst).unwrap(); ep.deallocate(out).unwrap(); } + + // Anti-regression lock for the async, fence-ordered weight page-in (#87 first + // increment). Both arms drive the transfer and compute streams through the + // *same* primitive chain `CudaWeightPage::upload_async` composes internally + // (`htod_async` + `record_copy_fence`), differing only in how the + // compute-stream consumer is ordered relative to the transfer: + // + // * Positive (real page-in ordering): a spin-delay holds the H2D copy + // pending on the transfer stream, then `compute_wait_fence` orders the + // compute-stream consumer after it, so the consumer reads the fully + // paged-in bytes. Deleting `compute_wait_fence` leaves the consumer to + // read the pre-copy POISON, so the lock is non-vacuous. + // * Negative (deterministic poison control): the transfer is event-ordered + // strictly *after* the consumer (`record_compute_fence` + `copy_wait_fence`), + // so with no `compute_wait_fence` the consumer provably reads pre-transfer + // POISON. This proves the compute-side wait is load-bearing without a + // wall-clock race — an earlier revision raced the consumer against a + // spin-delayed copy, which the parallel, captured `cargo test` invocation + // flaked whenever GPU contention delayed the consumer kernel past the copy. + // + // Every device/pinned allocation is hoisted out of the timing window, so no + // synchronizing `cuMemAlloc`/`cuMemHostAlloc` can drain the copy-stream + // spin-delay: the delay→async-copy→fence→consume window is the only thing the + // positive arm's ordering depends on. A trailing `upload_async` byte-parity + // check keeps the real allocate+stage+copy+fence entry point under test. + #[test] + fn async_pagein_fence_orders_weight_page_in_consumer() { + use cudarc::driver::{LaunchConfig, PushKernelArg}; + + const MODULE: &str = "cuda_ep_async_pagein_test"; + const SOURCE: &str = r#" +extern "C" __global__ void spin_delay(long long spin) { + long long start = clock64(); + while (clock64() - start < spin) { } +} +extern "C" __global__ void copy_out(const float* in, float* out, unsigned long long n) { + unsigned long long i = (unsigned long long)blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n) return; + out[i] = in[i]; +} +"#; + let Ok(ep) = CudaExecutionProvider::initialized(0) else { + eprintln!("skipping async page-in fence test: CUDA EP unavailable"); + return; + }; + let runtime = ep.runtime().clone(); + let spin_delay = runtime + .nvrtc_function(MODULE, SOURCE, "spin_delay") + .unwrap(); + let copy_out = runtime.nvrtc_function(MODULE, SOURCE, "copy_out").unwrap(); + + let n = 4096usize; + let bytes = n * std::mem::size_of::(); + let n_u64 = n as u64; + let payload: Vec = (0..n).map(|i| 5.0 + (i % 13) as f32).collect(); + let payload_bytes = + unsafe { std::slice::from_raw_parts(payload.as_ptr().cast::(), bytes) }; + + // Hoist every device/pinned allocation OUT of the per-iteration timing + // window: a synchronizing `cuMemAlloc`/`cuMemHostAlloc` between the + // spin-delay and the consumer would drain the delay and let ordering lean + // on the alloc instead of the fence. All buffers are reused each iteration. + let poison = vec![-777.0f32; n]; + let poison_bytes = + unsafe { std::slice::from_raw_parts(poison.as_ptr().cast::(), bytes) }; + let pos_dst = ep.allocate(bytes, 256).unwrap(); + let neg_dst = ep.allocate(bytes, 256).unwrap(); + let out = ep.allocate(bytes, 256).unwrap(); + let pos_dst_p = cuptr(pos_dst.as_ptr()); + let neg_dst_p = cuptr(neg_dst.as_ptr()); + let out_p = cuptr(out.as_ptr()); + let mut staging = runtime.alloc_pinned(bytes).unwrap(); + staging.as_mut_slice().copy_from_slice(payload_bytes); + let spin: i64 = 8_000_000; + + for _ in 0..8 { + // ── Positive: the real page-in ordering. Poison the destination, hold + // the H2D copy pending behind a spin-delay, then order the + // compute-stream consumer after the transfer with `compute_wait_fence` + // (the exact `htod_async` + `record_copy_fence` chain `upload_async` + // composes). With the fence the consumer reads the paged-in payload; + // delete the fence and it reads the poison below. + unsafe { runtime.htod(poison_bytes, pos_dst_p) }.unwrap(); + runtime.synchronize().unwrap(); + + let mut delay = runtime.copy_stream().launch_builder(&spin_delay); + delay.arg(&spin); + unsafe { delay.launch(LaunchConfig::for_num_elems(1)).unwrap() }; + + unsafe { runtime.htod_async(staging.as_slice(), pos_dst_p) }.unwrap(); + let fence = runtime.record_copy_fence().unwrap(); + runtime.compute_wait_fence(fence).unwrap(); + + let mut consume = runtime.stream().launch_builder(©_out); + consume.arg(&pos_dst_p).arg(&out_p).arg(&n_u64); + unsafe { + consume + .launch(LaunchConfig::for_num_elems(n as u32)) + .unwrap() + }; + let mut got = vec![0.0f32; n]; + let got_bytes = + unsafe { std::slice::from_raw_parts_mut(got.as_mut_ptr().cast::(), bytes) }; + unsafe { runtime.dtoh(got_bytes, out_p) }.unwrap(); + runtime.sync_copy_stream().unwrap(); + assert_eq!( + got, payload, + "async page-in consumer read stale bytes — compute_wait_fence did \ + not order the transfer before the compute-stream read" + ); + + // ── Negative (deterministic poison control): event-order the transfer + // strictly AFTER the consumer, so with NO `compute_wait_fence` the + // consumer provably reads pre-transfer poison. The `copy_wait_fence` + // on a compute-stream fence removes all wall-clock racing — the + // outcome never depends on the consumer winning against a delayed copy. + unsafe { runtime.htod(poison_bytes, neg_dst_p) }.unwrap(); + runtime.synchronize().unwrap(); + + let mut consume = runtime.stream().launch_builder(©_out); + consume.arg(&neg_dst_p).arg(&out_p).arg(&n_u64); + unsafe { + consume + .launch(LaunchConfig::for_num_elems(n as u32)) + .unwrap() + }; + // Hold the transfer until the consumer above has read `neg_dst`. + let consumer_fence = runtime.record_compute_fence().unwrap(); + runtime.copy_wait_fence(consumer_fence).unwrap(); + unsafe { runtime.htod_async(staging.as_slice(), neg_dst_p) }.unwrap(); + let _unused_fence = runtime.record_copy_fence().unwrap(); + + let mut raced = vec![0.0f32; n]; + let raced_bytes = + unsafe { std::slice::from_raw_parts_mut(raced.as_mut_ptr().cast::(), bytes) }; + unsafe { runtime.dtoh(raced_bytes, out_p) }.unwrap(); + // Drain the transfer (which lands after the consumer) before the next + // iteration reuses `neg_dst` / `staging`. + runtime.sync_copy_stream().unwrap(); + assert_eq!( + raced, poison, + "un-ordered async page-in consumer did NOT read poison — the \ + compute-stream wait is not load-bearing, so this lock proves nothing" + ); + + // ── Real `upload_async` entry point: allocate + stage + async-copy + + // fence, then a fenced consumer must observe the byte-identical + // payload. Keeps the production API (not just its primitive chain) + // under regression cover. + let (page, page_fence) = crate::weight_paging::CudaWeightPage::upload_async( + &runtime, + DataType::Float32, + vec![n], + payload_bytes, + ) + .unwrap(); + runtime.compute_wait_fence(page_fence).unwrap(); + let page_p = cuptr(page.device_ptr()); + let mut consume = runtime.stream().launch_builder(©_out); + consume.arg(&page_p).arg(&out_p).arg(&n_u64); + unsafe { + consume + .launch(LaunchConfig::for_num_elems(n as u32)) + .unwrap() + }; + let mut paged = vec![0.0f32; n]; + let paged_bytes = + unsafe { std::slice::from_raw_parts_mut(paged.as_mut_ptr().cast::(), bytes) }; + unsafe { runtime.dtoh(paged_bytes, out_p) }.unwrap(); + assert_eq!( + paged, payload, + "upload_async page-in read stale bytes — the returned copy fence \ + did not order the transfer before the compute-stream read" + ); + drop(page); + } + + ep.deallocate(pos_dst).unwrap(); + ep.deallocate(neg_dst).unwrap(); + ep.deallocate(out).unwrap(); + } } diff --git a/crates/onnx-runtime-ep-cuda/src/weight_paging.rs b/crates/onnx-runtime-ep-cuda/src/weight_paging.rs index 246f97ee3a..e30f9e25f9 100644 --- a/crates/onnx-runtime-ep-cuda/src/weight_paging.rs +++ b/crates/onnx-runtime-ep-cuda/src/weight_paging.rs @@ -26,7 +26,7 @@ use onnx_runtime_ep_api::{ }; use onnx_runtime_ir::DataType; -use crate::runtime::{CudaRuntime, raw_ptr}; +use crate::runtime::{CudaRuntime, PinnedStaging, raw_ptr}; /// Process-global weight-offload activity counters. Every [`CudaWeightResidency`] /// updates these in addition to its own instance stats, so an end-to-end decode @@ -71,14 +71,36 @@ pub const WEIGHT_OFFLOAD_ENV: &str = onnx_runtime_ep_cpu::WEIGHT_OFFLOAD_ENV; /// manager is constructed with a caller-chosen default. pub const WEIGHT_OFFLOAD_DEVICE_BYTES_ENV: &str = "ONNX_GENAI_WEIGHT_OFFLOAD_DEVICE_BYTES"; +/// Sub-knob (default ON) selecting the asynchronous, fence-ordered residency +/// page-in over the legacy synchronous `cuMemcpyHtoD`. Set to exactly `0` to +/// force the old synchronous page-in — used as the A/B "before" arm to measure +/// the overlap win, and as a kill-switch. It only has any effect when weight +/// offload itself is enabled (`ONNX_GENAI_WEIGHT_OFFLOAD=1`); with offload off +/// the resident fast path is untouched and byte-identical regardless. +pub const WEIGHT_OFFLOAD_ASYNC_PAGEIN_ENV: &str = "ONNX_GENAI_WEIGHT_OFFLOAD_ASYNC_PAGEIN"; + /// Whether/how the CUDA EP should page offloaded weights into a bounded VRAM /// residency cache. Disabled by default so the resident fast path is untouched /// and byte-identical. -#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +#[derive(Clone, Copy, Debug, PartialEq, Eq)] pub struct DeviceOffloadPolicy { pub enabled: bool, /// Explicit VRAM budget in bytes, if the operator pinned one. pub device_budget_bytes: Option, + /// Use the asynchronous, fence-ordered page-in (default `true`). Disabled + /// only when `ONNX_GENAI_WEIGHT_OFFLOAD_ASYNC_PAGEIN=0`, which restores the + /// synchronous page-in for A/B measurement or as a kill-switch. + pub async_pagein: bool, +} + +impl Default for DeviceOffloadPolicy { + fn default() -> Self { + Self { + enabled: false, + device_budget_bytes: None, + async_pagein: true, + } + } } impl DeviceOffloadPolicy { @@ -88,9 +110,13 @@ impl DeviceOffloadPolicy { let device_budget_bytes = std::env::var(WEIGHT_OFFLOAD_DEVICE_BYTES_ENV) .ok() .and_then(|value| parse_budget_bytes(&value)); + // Async page-in defaults ON; only an explicit `0` disables it. + let async_pagein = + std::env::var_os(WEIGHT_OFFLOAD_ASYNC_PAGEIN_ENV).is_none_or(|value| value != "0"); Self { enabled, device_budget_bytes, + async_pagein, } } } @@ -134,6 +160,11 @@ pub struct CudaWeightPage { len: usize, dtype: DataType, shape: Vec, + /// Owned pinned host staging that backed an asynchronous page-in, kept alive + /// so an in-flight `cuMemcpyHtoDAsync` never reads freed host memory. `None` + /// for pages uploaded synchronously (the source is fully consumed before + /// `htod` returns). Freed on drop — long after the copy has completed. + staging: Option, } impl CudaWeightPage { @@ -158,6 +189,7 @@ impl CudaWeightPage { len: bytes.len(), dtype, shape, + staging: None, }; // SAFETY: `ptr` owns `bytes.len()` bytes; `page`'s Drop frees it if the // copy below fails. @@ -166,6 +198,59 @@ impl CudaWeightPage { Ok(page) } + /// Asynchronous, overlap-friendly variant of [`Self::upload`]: allocate a + /// VRAM page and enqueue a host→device copy of `bytes` on the runtime's + /// dedicated transfer stream, returning the page plus a **copy fence** the + /// caller MUST order the consuming compute work after (via + /// [`CudaRuntime::compute_wait_fence`]). Because the copy is asynchronous, + /// the source bytes are first staged into a page-locked buffer that the + /// returned [`CudaWeightPage`] owns, so the transfer keeps reading valid host + /// memory until it completes — the staging outlives the page's whole life, + /// which itself outlives any consumer ordered after the fence. Frees the VRAM + /// (and staging) on any failure. + pub fn upload_async( + runtime: &Arc, + dtype: DataType, + shape: Vec, + bytes: &[u8], + ) -> Result<(Self, u64), WeightHandleError> { + if bytes.is_empty() { + return Err(WeightHandleError::MissingRegions); + } + let ptr = runtime + .alloc_raw(bytes.len()) + .map_err(|error| WeightHandleError::DeviceBinding(format!("VRAM alloc: {error}")))?; + let mut staging = runtime + .alloc_pinned(bytes.len()) + .map_err(|error| WeightHandleError::DeviceBinding(format!("pinned alloc: {error}")))?; + staging.as_mut_slice().copy_from_slice(bytes); + // Own both the VRAM `ptr` and the pinned `staging` before enqueuing the + // copy, so any error below drops `page` and frees both exactly once. + let page = Self { + runtime: Arc::clone(runtime), + ptr, + len: bytes.len(), + dtype, + shape, + staging: Some(staging), + }; + // SAFETY: `dst` (`ptr`) owns `len` bytes; the async source is the pinned + // staging `page` owns, which outlives the transfer (freed only on the + // page's Drop, well after the copy fence has been awaited). + let staged = page + .staging + .as_ref() + .expect("staging was just set") + .as_slice(); + unsafe { runtime.htod_async(staged, ptr) }.map_err(|error| { + WeightHandleError::DeviceBinding(format!("async H2D copy: {error}")) + })?; + let fence = runtime + .record_copy_fence() + .map_err(|error| WeightHandleError::DeviceBinding(format!("copy fence: {error}")))?; + Ok((page, fence)) + } + /// Opaque device pointer to the paged bytes, for a kernel `TensorView`. pub fn device_ptr(&self) -> *const std::ffi::c_void { raw_ptr(self.ptr) @@ -251,6 +336,7 @@ impl LazyDeviceWeightBinder for CudaWeightPager<'_, S> { len: total, dtype: weight.dtype, shape: weight.shape.clone(), + staging: None, }; let mut offset: usize = 0; @@ -292,8 +378,20 @@ impl LazyDeviceWeightBinder for CudaWeightPager<'_, S> { /// /// A single weight larger than the whole budget is still paged in (correctness /// beats the budget); the cache simply runs transiently over budget for it. +/// +/// Page-in defaults to the asynchronous, fence-ordered path +/// ([`CudaWeightPage::upload_async`]): the H2D copy is enqueued on the runtime's +/// dedicated transfer stream so it overlaps the in-flight compute, and the +/// consuming compute stream is ordered after it with a completion fence. The +/// full compute-stream drain is now taken **only when admitting must evict** (so +/// eviction never frees a page a prior kernel still reads); a page-in that fits +/// the budget no longer host-blocks, which is what lets the transfer overlap. +/// Set `async_pagein = false` (env `ONNX_GENAI_WEIGHT_OFFLOAD_ASYNC_PAGEIN=0`) to +/// restore the legacy synchronous page-in for A/B comparison. pub struct CudaWeightResidency { runtime: Arc, + /// Whether misses page in asynchronously (see [`CudaWeightPage::upload_async`]). + async_pagein: bool, inner: Mutex, } @@ -310,10 +408,12 @@ struct ResidencyInner { } impl CudaWeightResidency { - /// Build a residency cache with an explicit VRAM `budget_bytes`. + /// Build a residency cache with an explicit VRAM `budget_bytes`. Page-in is + /// asynchronous by default; chain [`Self::with_async_pagein`] to override. pub fn new(runtime: Arc, budget_bytes: u64) -> Self { Self { runtime, + async_pagein: true, inner: Mutex::new(ResidencyInner { budget: budget_bytes, resident_bytes: 0, @@ -327,6 +427,12 @@ impl CudaWeightResidency { } } + /// Select the asynchronous (default `true`) vs synchronous page-in path. + pub fn with_async_pagein(mut self, async_pagein: bool) -> Self { + self.async_pagein = async_pagein; + self + } + /// Return the device page for `key`, paging it in from `source` on a miss and /// evicting LRU pages to respect the budget. The returned [`Arc`] keeps the /// page resident for as long as the caller holds it. @@ -360,62 +466,107 @@ impl CudaWeightResidency { return Ok(hit); } let resident = weight.materialize()?; - let page = Arc::new(CudaWeightPage::upload( - &self.runtime, - resident.dtype, - resident.shape.clone(), - resident.bytes(), - )?); - self.admit(key, page) + if self.async_pagein { + // Async, fence-ordered page-in: enqueue the H2D on the transfer stream + // (overlapping the in-flight compute), admit, then order the compute + // stream after the transfer's completion event so the consuming kernel + // waits only on the copy — never a full host sync. + let (page, copy_fence) = CudaWeightPage::upload_async( + &self.runtime, + resident.dtype, + resident.shape.clone(), + resident.bytes(), + )?; + let admitted = self.admit(key, Arc::new(page))?; + self.runtime + .compute_wait_fence(copy_fence) + .map_err(|error| { + WeightHandleError::DeviceBinding(format!("copy fence wait: {error}")) + })?; + Ok(admitted) + } else { + // Legacy synchronous page-in (A/B "before" arm / kill-switch): the + // blocking `htod` serializes the transfer with compute. + let page = Arc::new(CudaWeightPage::upload( + &self.runtime, + resident.dtype, + resident.shape.clone(), + resident.bytes(), + )?); + self.admit(key, page) + } } /// Look up `key`, marking it most-recently-used and counting a hit. fn get_hit(&self, key: u64) -> Option> { let mut inner = self.lock(); if let Some(page) = inner.pages.get(&key).cloned() { - inner.touch(key); - inner.hits += 1; - GLOBAL_HITS.fetch_add(1, Ordering::Relaxed); + inner.record_hit(key); Some(page) } else { None } } + /// Drain the transfer stream so a just-uploaded page we are about to drop + /// cannot have an in-flight async copy still reading its (freeing) VRAM or + /// pinned staging. A no-op for synchronously-uploaded pages (idle copy stream). + fn drain_copy_stream(&self) -> Result<(), WeightHandleError> { + self.runtime.sync_copy_stream().map_err(|error| { + WeightHandleError::DeviceBinding(format!("transfer stream sync: {error}")) + }) + } + /// Insert a freshly paged-in `page` under `key`, evicting LRU pages to fit the - /// budget. Synchronizes the compute stream first so no in-flight kernel still - /// references an about-to-be-freed page's VRAM. (Offload is mutually exclusive - /// with CUDA graph capture, so this sync is never capture-illegal.) + /// budget. + /// + /// The compute stream is drained **only when admitting must evict**, so no + /// in-flight kernel still references an about-to-be-freed page's VRAM (the + /// original WAR/reuse guarantee). A page-in that fits the budget frees nothing + /// and therefore skips the sync, letting its async transfer overlap the + /// current compute. (Offload is mutually exclusive with CUDA graph capture, so + /// this sync is never capture-illegal.) fn admit( &self, key: u64, page: Arc, ) -> Result, WeightHandleError> { let bytes = page.len() as u64; - // Weight offload and CUDA graph capture are mutually exclusive: the decode - // session declines capture whenever offload is enabled (see - // `resolve_graph_capture_enabled`), so paging never runs during capture. - // The synchronize below — and the alloc/copy/free it guards against — are - // themselves capture-illegal, so we can unconditionally sync here. + { + let mut inner = self.lock(); + // A concurrent caller may have populated `key` while we paged in; + // prefer the already-resident page and drop ours (its Drop frees the + // VRAM + staging — drained first so it never races an in-flight copy). + if let Some(existing) = inner.pages.get(&key).cloned() { + inner.record_hit(key); + drop(inner); + self.drain_copy_stream()?; + return Ok(existing); + } + // Fits without eviction: nothing is freed, so no consumer drain is + // needed. Admit directly so the async transfer keeps overlapping. + if inner.resident_bytes.saturating_add(bytes) <= inner.budget { + inner.insert_page(key, Arc::clone(&page), bytes); + return Ok(page); + } + } + // Eviction required: drain in-flight consumers before freeing any page. + // Weight offload and CUDA graph capture are mutually exclusive (the decode + // session declines capture whenever offload is enabled), so this + // synchronize is never capture-illegal. self.runtime .synchronize() .map_err(|error| WeightHandleError::DeviceBinding(format!("stream sync: {error}")))?; let mut inner = self.lock(); - // A concurrent caller may have populated `key` while we paged in; prefer - // the already-resident page and drop ours (its Drop frees the VRAM). + // Re-check after releasing the lock for the sync. if let Some(existing) = inner.pages.get(&key).cloned() { - inner.touch(key); - inner.hits += 1; - GLOBAL_HITS.fetch_add(1, Ordering::Relaxed); + inner.record_hit(key); + drop(inner); + self.drain_copy_stream()?; return Ok(existing); } inner.evict_to_fit(bytes); - inner.pages.insert(key, Arc::clone(&page)); - inner.order.push(key); - inner.resident_bytes += bytes; - inner.peak_resident_bytes = inner.peak_resident_bytes.max(inner.resident_bytes); - inner.page_ins += 1; - GLOBAL_PAGE_INS.fetch_add(1, Ordering::Relaxed); + inner.insert_page(key, Arc::clone(&page), bytes); Ok(page) } @@ -458,6 +609,25 @@ impl ResidencyInner { } } + /// Record a cache hit for `key`: mark it most-recently-used and bump the + /// per-instance and process-global hit counters. + fn record_hit(&mut self, key: u64) { + self.touch(key); + self.hits += 1; + GLOBAL_HITS.fetch_add(1, Ordering::Relaxed); + } + + /// Insert a freshly paged-in `page` of `bytes` under `key`, updating the LRU + /// order, residency accounting, and the page-in counters. + fn insert_page(&mut self, key: u64, page: Arc, bytes: u64) { + self.pages.insert(key, page); + self.order.push(key); + self.resident_bytes += bytes; + self.peak_resident_bytes = self.peak_resident_bytes.max(self.resident_bytes); + self.page_ins += 1; + GLOBAL_PAGE_INS.fetch_add(1, Ordering::Relaxed); + } + /// Evict least-recently-used, currently-unreferenced pages until admitting /// `incoming` bytes fits the budget (best effort; stops when nothing more is /// evictable).