diff --git a/jouleclaw/SPEC.md b/jouleclaw/SPEC.md index 572d6c552..fe7018e92 100644 --- a/jouleclaw/SPEC.md +++ b/jouleclaw/SPEC.md @@ -119,6 +119,17 @@ Realistic floors: counter is declined there and the host is estimated (`Estimator`). Marketing claims of microjoule precision on Apple Silicon are sales talk; the spec is honest about this. +- Apple Silicon, per process (macOS 27.0.1, unprivileged): the + process's own CPU energy (`proc_pid_rusage`, nJ, live) is + `HwShunt`. Its GPU energy is `HwShunt` in two cases only: an exact + zero when its own GPU time (I/O Registry) did not move, and the + IOReport `GPU Energy` rail (nJ, live) when it had at least 99% of + every GPU client's GPU time over the span. A shared GPU gives no + figure: the rail **MUST NOT** be split by GPU time, which is not + energy-proportional (measured: ~1.9 W vs ~17 W per GPU-second for + two jobs). Without a GPU figure, the process's CPU figure alone is + `Partial`. No live source covers the ANE: a figure for work that + used it is `Partial`. - Consumer AMD GPU, ARM PMU: no usable counter. Use a calibrated `Estimator` and surface the wider tolerance band. diff --git a/jouleclaw/jouleclaw-rs/Cargo.lock b/jouleclaw/jouleclaw-rs/Cargo.lock index c302a3a6a..6e906c56c 100644 --- a/jouleclaw/jouleclaw-rs/Cargo.lock +++ b/jouleclaw/jouleclaw-rs/Cargo.lock @@ -3630,6 +3630,9 @@ dependencies = [ "jouleclaw-core", "libc", "nvml-wrapper", + "objc2", + "objc2-foundation", + "objc2-metal", "serde", "serde_json", "thiserror 2.0.20", @@ -4705,6 +4708,7 @@ dependencies = [ "jouleclaw-backend-apple", "jouleclaw-backend-reference", "jouleclaw-core", + "jouleclaw-energy", "jouleclaw-loader-gguf", "jouleclaw-topology", "serde", diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/gpu_load.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/gpu_load.rs new file mode 100644 index 000000000..e90cbc196 --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/gpu_load.rs @@ -0,0 +1,138 @@ +//! A short, known GPU load from this process: the work the per-process GPU +//! energy counter's liveness probe runs between its samples +//! (`jouleclaw_energy::process_gpu::gpu_rail_liveness`). That crate holds no +//! Metal, so the Metal-owning caller supplies the load: +//! +//! ```ignore +//! let meter = jouleclaw_energy::span::ProcessSpanMeter::new(jouleclaw_backend_apple::gpu_liveness_load); +//! ``` +//! +//! The probe asks the GPU energy rail and this process's own GPU time to +//! advance under it, so the load must run on the GPU (a CPU burn moves +//! neither) and must be waited on. + +/// The GPU time one call aims to run: the probe's tick. +pub const GPU_LOAD_TICK: std::time::Duration = std::time::Duration::from_millis(25); + +/// Put about [`GPU_LOAD_TICK`] of compute work on the GPU from this process — +/// command buffers of a small FMA kernel, each committed and waited on, until +/// their GPU time sums to the tick, or 500 ms of wall time has passed on a GPU +/// too busy to grant it. Does nothing where there is no Metal device (and off +/// macOS). +pub fn gpu_liveness_load() { + #[cfg(target_os = "macos")] + if let Some(b) = inner::burner() { + b.burn(GPU_LOAD_TICK, std::time::Duration::from_millis(500)); + } +} + +#[cfg(target_os = "macos")] +mod inner { + use std::ptr::NonNull; + use std::sync::{Mutex, OnceLock}; + use std::time::{Duration, Instant}; + + use objc2::rc::Retained; + use objc2::runtime::ProtocolObject; + use objc2_foundation::NSString; + use objc2_metal::{ + MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue, MTLComputeCommandEncoder, + MTLComputePipelineState, MTLCreateSystemDefaultDevice, MTLDevice, MTLLibrary, MTLResourceOptions, MTLSize, + }; + + const BURN_MSL: &str = r#" +#include +using namespace metal; +kernel void liveness_burn(device float* b [[buffer(0)]], constant uint& n [[buffer(1)]], uint i [[thread_position_in_grid]]) { + float x = b[i]; + for (uint k = 0; k < n; k++) { x = fma(x, 1.0000001f, 0.0000001f); } + b[i] = x; +} +"#; + const THREADS: usize = 1 << 18; + + pub(super) struct Burner { + queue: Retained>, + pso: Retained>, + buf: Retained>, + /// Loop count per thread; sized at first use to ~5 ms a command + /// buffer. Behind a lock: one thread dispatches at a time. + iters: Mutex, + } + + // SAFETY: Metal documents MTLCommandQueue, compute pipeline states and + // buffers as thread-safe; the loop count is behind a mutex, and that + // mutex also serialises dispatch. + unsafe impl Send for Burner {} + unsafe impl Sync for Burner {} + + pub(super) fn burner() -> Option<&'static Burner> { + static B: OnceLock> = OnceLock::new(); + B.get_or_init(|| { + let device = MTLCreateSystemDefaultDevice()?; + let lib = device.newLibraryWithSource_options_error(&NSString::from_str(BURN_MSL), None).ok()?; + let f = lib.newFunctionWithName(&NSString::from_str("liveness_burn"))?; + let pso = device.newComputePipelineStateWithFunction_error(&f).ok()?; + let buf = device.newBufferWithLength_options(THREADS * 4, MTLResourceOptions::StorageModeShared)?; + let queue = device.newCommandQueue()?; + let b = Burner { queue, pso, buf, iters: Mutex::new(64) }; + let t = (0..3).filter_map(|_| b.dispatch(64)).fold(f64::MAX, f64::min).max(1e-5); + *b.iters.lock().ok()? = ((64.0 * 0.005 / t) as u32).clamp(64, 1 << 20); + Some(b) + }) + .as_ref() + } + + impl Burner { + /// One command buffer of `iters` loops a thread, committed and waited + /// on: its GPU time in seconds, or `None` if Metal refused it. + fn dispatch(&self, iters: u32) -> Option { + let cb = self.queue.commandBuffer()?; + let enc = cb.computeCommandEncoder()?; + enc.setComputePipelineState(&self.pso); + // SAFETY: buffer 0 holds THREADS floats and buffer 1 one u32, as + // the kernel declares; `iters` outlives the call, which copies it. + unsafe { + enc.setBuffer_offset_atIndex(Some(&self.buf), 0, 0); + enc.setBytes_length_atIndex(NonNull::from(&iters).cast(), 4, 1); + } + let tg = self.pso.maxTotalThreadsPerThreadgroup().min(THREADS); + enc.dispatchThreads_threadsPerThreadgroup( + MTLSize { width: THREADS, height: 1, depth: 1 }, + MTLSize { width: tg, height: 1, depth: 1 }, + ); + enc.endEncoding(); + cb.commit(); + cb.waitUntilCompleted(); + Some((cb.GPUEndTime() - cb.GPUStartTime()).max(0.0)) + } + + /// Dispatch until `gpu` of GPU time has run, or `wall` has passed. + pub(super) fn burn(&self, gpu: Duration, wall: Duration) { + let Ok(iters) = self.iters.lock() else { return }; + let t = Instant::now(); + let mut ran = 0.0; + while ran < gpu.as_secs_f64() && t.elapsed() < wall { + match self.dispatch(*iters) { + Some(s) => ran += s, + None => return, + } + } + } + } +} + +#[cfg(all(test, target_os = "macos"))] +mod tests { + /// The load runs, and takes at least its tick of wall time (its GPU time + /// is at least the tick unless the 500 ms wall cap cut it short). + #[test] + fn the_load_runs_on_the_gpu() { + assert!(super::inner::burner().is_some(), "no Metal device"); + let t = std::time::Instant::now(); + super::gpu_liveness_load(); + let took = t.elapsed(); + assert!(took >= super::GPU_LOAD_TICK, "{took:?}"); + assert!(took < std::time::Duration::from_secs(5), "{took:?}"); + } +} diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs index be02ac4a2..ddc1a4e31 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs @@ -21,9 +21,12 @@ //! falls back to the reference backend. mod accelerate; +mod gpu_load; mod metal_compute; pub mod mps_matmul; +pub use gpu_load::{gpu_liveness_load, GPU_LOAD_TICK}; + pub use accelerate::{sgemm_f32, sgemm_f32_bt, AccelerateMatMul}; pub use metal_compute::MetalMatMul; pub use mps_matmul::{ diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/Cargo.toml b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/Cargo.toml index dc5cd665d..371d63d91 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/Cargo.toml +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/Cargo.toml @@ -29,6 +29,14 @@ serde = ["dep:serde", "dep:serde_json"] # downstream crates / binaries enable what their target supports. rapl = [] # Intel/AMD x86 — Linux sysfs powercap (HwShunt while its counter is live) nvml = ["dep:nvml-wrapper", "dep:tracing"] # NVIDIA discrete — NVML cumulative-energy counter (HwShunt while live), power-draw integral (ModelBased) otherwise -ioreport = [] # Apple Silicon M-series — real IOReport energy counter (HwShunt while every CPU rail is live), load-avg estimator fallback +ioreport = [] # Apple Silicon M-series — real IOReport energy counter (HwShunt while every CPU rail is live), load-avg estimator fallback; and the per-process GPU counter (GPU rail × this process's GPU-time share) ina3221 = [] # NVIDIA Jetson (Orin / Thor) — sysfs i2c shunt reader hwmon = [] # Linux generic — for AMD Strix Halo / Ryzen AI Max iGPU + others + +# The per-process GPU counter's hardware test loads the GPU itself (this crate +# holds no Metal; its consumers supply the load). Same Metal bindings as +# `jouleclaw-backend-apple`. +[target.'cfg(target_os = "macos")'.dev-dependencies] +objc2 = "0.6" +objc2-foundation = "0.3" +objc2-metal = "0.3" diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/MEASURED_ENERGY.md b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/MEASURED_ENERGY.md index 6bfb60c9a..81516b19e 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/MEASURED_ENERGY.md +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/MEASURED_ENERGY.md @@ -61,10 +61,57 @@ Notes: `intel-rapl:*/energy_uj` is root-readable only (PLATYPUS mitigation); the counter reads it when run with sufficient privilege. +## Apple Silicon, per process (macOS 27.0.1, unprivileged) + +For work that runs in this process — its threads, and GPU work it dispatches — +two live sources replace the frozen rails: + +| Counter | Source | Domain | Provenance | +|---|---|---|---| +| `process-cpu` ([`process`](src/process.rs)) | `proc_pid_rusage` `ri_energy_nj` | this process's CPU | **HwShunt** while live | +| `ProcessGpuCounter` ([`process_gpu`](src/process_gpu.rs), `ioreport`) | this process's `accumulatedGPUTime` (I/O Registry `IOAccelerator` user clients); IOReport `GPU Energy` (nJ) | this process's GPU | **HwShunt**: 0 µJ exact with no GPU time; the rail at ≥ 99% of every client's GPU time; else no figure | + +The GPU rail is machine-wide and is never split by GPU time: energy is not +proportional to it. On this host a GGUF decode drew ~1.9 W per GPU-second alone +while another session's job drew ~17 W, so a time-share split charged the decode +3–15× its sole-user GPU energy; and the decode's own GPU time grew 5.4× (1.83 s +to 9.81 s for the same 63 tokens) beside a compute burner. A shared GPU gives no +figure (`GpuShared`), nor does a GPU client that vanished mid-span while this +process used the GPU, nor a rail that held while the GPU was busy. + +[`span::ProcessSpanMeter`](src/span.rs) adds the two: `Whole` (DRAM excluded, +as in a RAPL package figure), `Partial` for work declared to use the ANE, which +has no live unprivileged source. Where the GPU has no figure the span is the CPU +figure alone, `HwShunt`, `Partial`, with why; only a CPU counter with no figure +falls back to the host meter. The GPU liveness probe needs a GPU load from the +caller (`jouleclaw_backend_apple::gpu_liveness_load`). + +Measured on an M5 Max (2026-10-04): + +- `tests/process_gpu.rs`, quiet GPU: a 1.5 s compute burn read 39.860 J, this + process 1603.6 ms of 1603.6 ms GPU time, equal to the test's own IOReport + read and `ioreg`'s times. Beside a burning child: `GpuShared` (491.3 of + 1179.8 ms); spans with no GPU work of this process's were 0 µJ exactly, beside + the child and across its exit. +- `gguf_energy` (jouleclaw-runtime example), tinyllama-1.1b Q8_0 with + `Q8_PACKED=1` (decode GEMVs on the GPU), 63 decode tokens: + - sole GPU user: **0.128 J/token** = CPU 0.069 + GPU 0.060, `HwShunt`, + `Whole` (1830.0 ms of 1836.1 ms GPU time, 99.67%; 3.77 J rail), + 71.5 ms/token; + - beside a compute burner: CPU alone 0.053 J/token, `HwShunt`, `Partial`: + the GPU was shared (9.81 s of 13.98 s of GPU time), 225 ms/token. +- The rail at times holds 300–600 ms under load and catches up in one step, + and on a quiet GPU it did not change at all for 4 s, so a snapshot reads it + on its next change, or as it stands once no client accrues GPU time. On this + host 18% of 5 s windows (24% of 10 s) saw a GPU client vanish. + ## Not yet a hardware reading -- **Apple Silicon CPU energy on macOS 27.0.1**, for an unprivileged process — - the IOReport CPU rails are frozen (above), so the load-avg estimator meters it. +- **Apple Silicon host-wide CPU energy on macOS 27.0.1**, for an unprivileged + process — the IOReport CPU rails are frozen (above), so the load-avg estimator + meters work outside this process. +- **Apple Neural Engine energy**, unprivileged: `ANE0` is frozen and ANE work + runs in `aned`; the `H1xANELoadBalancer` clients carry no time or energy. - **x86 CPU energy on Windows** — RAPL MSRs need a kernel driver (e.g. WinRing0) on Windows; the NVML GPU path is the Windows `HwShunt` route. (CPU-side is an `Estimator` there for now.) diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_gpu.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_gpu.rs new file mode 100644 index 000000000..ec26f8569 --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_gpu.rs @@ -0,0 +1,259 @@ +//! Every GPU client's accumulated GPU time, from the I/O Registry (macOS): the +//! time half of the per-process GPU counter (`crate::process_gpu`), whose energy +//! half is the IOReport GPU rail (`crate::apple_ioreport::GpuRails`). +//! +//! Each `IOAccelerator` service (on Apple Silicon one, `AGXAccelerator…`) has +//! one user client per connection a process opened to the GPU, as a child in +//! the `IOService` plane. Each carries `IOUserClientCreator` (`"pid N, name"`) +//! and `AppUsage`, an array with one dictionary per API context, whose +//! `accumulatedGPUTime` is the GPU time that context has spent, in ns. This is +//! public IOKit API, read unprivileged; it is what `ioreg -r -c IOAccelerator -l` +//! prints. A process can hold several clients. A client vanishes from the +//! registry, and its accumulated time with it, when its process closes it or +//! exits. + +use std::ffi::{c_char, c_long, c_void}; + +type CFTypeRef = *const c_void; +type CFStringRef = *const c_void; +type CFArrayRef = *const c_void; +type CFDictionaryRef = *const c_void; +type CFMutableDictionaryRef = *mut c_void; +type CFTypeID = usize; +type IoObject = libc::mach_port_t; + +const K_CFSTRING_ENCODING_UTF8: u32 = 0x0800_0100; +/// `kCFNumberSInt64Type`. +const K_CFNUMBER_SINT64: c_long = 4; +/// `kIOMainPortDefault` is `MACH_PORT_NULL`. +const MAIN_PORT_DEFAULT: libc::mach_port_t = 0; +const KERN_SUCCESS: libc::kern_return_t = 0; + +unsafe extern "C" { + fn IOServiceMatching(name: *const c_char) -> CFMutableDictionaryRef; + fn IOServiceGetMatchingServices( + main_port: libc::mach_port_t, + matching: CFDictionaryRef, + existing: *mut IoObject, + ) -> libc::kern_return_t; + fn IOIteratorNext(iterator: IoObject) -> IoObject; + fn IOObjectRelease(object: IoObject) -> libc::kern_return_t; + fn IORegistryEntryGetChildIterator( + entry: IoObject, + plane: *const c_char, + iterator: *mut IoObject, + ) -> libc::kern_return_t; + fn IORegistryEntryGetRegistryEntryID(entry: IoObject, id: *mut u64) -> libc::kern_return_t; + fn IORegistryEntryCreateCFProperty( + entry: IoObject, + key: CFStringRef, + allocator: *const c_void, + options: u32, + ) -> CFTypeRef; + + fn CFRelease(cf: CFTypeRef); + fn CFGetTypeID(cf: CFTypeRef) -> CFTypeID; + fn CFStringGetTypeID() -> CFTypeID; + fn CFArrayGetTypeID() -> CFTypeID; + fn CFDictionaryGetTypeID() -> CFTypeID; + fn CFNumberGetTypeID() -> CFTypeID; + fn CFArrayGetCount(the_array: CFArrayRef) -> c_long; + fn CFArrayGetValueAtIndex(the_array: CFArrayRef, idx: c_long) -> CFTypeRef; + fn CFDictionaryGetValue(the_dict: CFDictionaryRef, key: *const c_void) -> *const c_void; + fn CFNumberGetValue(number: CFTypeRef, the_type: c_long, value_ptr: *mut c_void) -> bool; + fn CFStringCreateWithCString(allocator: *mut c_void, cstr: *const c_char, encoding: u32) -> CFStringRef; + fn CFStringGetCString(the_string: CFStringRef, buffer: *mut c_char, buffer_size: c_long, encoding: u32) -> bool; +} + +// IOKit and CoreFoundation are linked by `apple_ioreport`, compiled under the +// same configuration as this module. + +/// One GPU client as the registry shows it at one instant. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub(crate) struct RawClient { + /// The client's registry entry ID: unique for the life of the boot, never + /// reused, so it tells a client apart from a later one of the same process. + pub(crate) id: u64, + /// The process that opened it (`IOUserClientCreator`). + pub(crate) pid: i32, + /// Its `accumulatedGPUTime`, summed over its `AppUsage` entries, in ns. + pub(crate) gpu_ns: u64, +} + +/// An owned CF object, released on drop. +struct Cf(CFTypeRef); + +impl Drop for Cf { + fn drop(&mut self) { + if !self.0.is_null() { + // SAFETY: `self.0` is a CF object this process owns (a Create/Copy + // result), released once. + unsafe { CFRelease(self.0) } + } + } +} + +/// An owned I/O Kit object, released on drop. +struct Io(IoObject); + +impl Drop for Io { + fn drop(&mut self) { + if self.0 != 0 { + // SAFETY: `self.0` is an object reference this process owns, + // released once. + unsafe { IOObjectRelease(self.0) }; + } + } +} + +impl Io { + /// The next object of an iterator, or `None` at its end. + fn next_of(iterator: &Io) -> Option { + // SAFETY: `iterator` is a live iterator this process owns. + let o = unsafe { IOIteratorNext(iterator.0) }; + (o != 0).then_some(Io(o)) + } +} + +fn cfstr(s: &std::ffi::CStr) -> Option { + // SAFETY: `s` is a NUL-terminated C string for the duration of the call. + let r = unsafe { CFStringCreateWithCString(std::ptr::null_mut(), s.as_ptr(), K_CFSTRING_ENCODING_UTF8) }; + (!r.is_null()).then_some(Cf(r)) +} + +/// `entry`'s property `key`, owned, if it is set and of the CF type `type_id`. +fn property(entry: &Io, key: &Cf, type_id: CFTypeID) -> Option { + // SAFETY: `entry` is a live registry entry and `key` a live CFString; the + // result is a new reference this process owns (wrapped at once). + let v = Cf(unsafe { IORegistryEntryCreateCFProperty(entry.0, key.0, std::ptr::null(), 0) }); + // SAFETY: `v.0` is a live CF object when non-null. + (!v.0.is_null() && unsafe { CFGetTypeID(v.0) } == type_id).then_some(v) +} + +/// A CFString's text, borrowed from it. +fn text(s: &Cf) -> Option { + let mut buf = [0 as c_char; 256]; + // SAFETY: `s.0` is a live CFString and `buf` a writable buffer of the size + // passed. + let ok = unsafe { CFStringGetCString(s.0, buf.as_mut_ptr(), buf.len() as c_long, K_CFSTRING_ENCODING_UTF8) }; + // SAFETY: on success the buffer holds a NUL-terminated string. + ok.then(|| unsafe { std::ffi::CStr::from_ptr(buf.as_ptr()) }.to_string_lossy().into_owned()) +} + +/// The pid in an `IOUserClientCreator` value, `"pid 172, WindowServer"`. +pub(crate) fn creator_pid(creator: &str) -> Option { + let rest = creator.strip_prefix("pid ")?; + let digits = rest.split(',').next()?.trim(); + digits.parse().ok() +} + +/// The `accumulatedGPUTime` of every `AppUsage` entry, summed, in ns; `None` +/// if an entry is not a dictionary or its time is not a 64-bit integer. +fn app_usage_ns(usage: &Cf, time_key: &Cf) -> Option { + let mut sum: u64 = 0; + // SAFETY: `usage` is a live CFArray (its type was checked); its values and + // their values are borrowed from it for the duration of this call. + unsafe { + for i in 0..CFArrayGetCount(usage.0) { + let entry = CFArrayGetValueAtIndex(usage.0, i); + if entry.is_null() || CFGetTypeID(entry) != CFDictionaryGetTypeID() { + return None; + } + let t = CFDictionaryGetValue(entry, time_key.0); + if t.is_null() { + // An entry that has not recorded GPU time yet. + continue; + } + if CFGetTypeID(t) != CFNumberGetTypeID() { + return None; + } + let mut ns: i64 = 0; + if !CFNumberGetValue(t, K_CFNUMBER_SINT64, &mut ns as *mut i64 as *mut c_void) || ns < 0 { + return None; + } + sum = sum.checked_add(ns as u64)?; + } + } + Some(sum) +} + +/// Every GPU client in the registry now: each `IOAccelerator`'s children in +/// the `IOService` plane that carry an `IOUserClientCreator`, with their +/// accumulated GPU time (a client with no `AppUsage` has spent none). `None` +/// if the registry could not be walked, or a client's creator or usage could +/// not be read: a snapshot that silently left a client out would hand its GPU +/// time to everyone else. +pub(crate) fn gpu_clients() -> Option> { + let creator_key = cfstr(c"IOUserClientCreator")?; + let usage_key = cfstr(c"AppUsage")?; + let time_key = cfstr(c"accumulatedGPUTime")?; + // SAFETY: plain type-ID getters. + let (string_t, array_t) = unsafe { (CFStringGetTypeID(), CFArrayGetTypeID()) }; + + // SAFETY: `IOServiceMatching` returns a new dictionary, which + // `IOServiceGetMatchingServices` consumes whatever it returns. + let services = unsafe { + let matching = IOServiceMatching(c"IOAccelerator".as_ptr()); + if matching.is_null() { + return None; + } + let mut it: IoObject = 0; + if IOServiceGetMatchingServices(MAIN_PORT_DEFAULT, matching, &mut it) != KERN_SUCCESS { + return None; + } + Io(it) + }; + let mut out = Vec::new(); + while let Some(accelerator) = Io::next_of(&services) { + let mut it: IoObject = 0; + // SAFETY: `accelerator` is a live registry entry; the iterator is ours. + if unsafe { IORegistryEntryGetChildIterator(accelerator.0, c"IOService".as_ptr(), &mut it) } != KERN_SUCCESS { + return None; + } + let children = Io(it); + while let Some(child) = Io::next_of(&children) { + let Some(creator) = property(&child, &creator_key, string_t) else { + // Not a user client. + continue; + }; + let pid = creator_pid(&text(&creator)?)?; + let gpu_ns = match property(&child, &usage_key, array_t) { + Some(usage) => app_usage_ns(&usage, &time_key)?, + None => 0, + }; + let mut id: u64 = 0; + // SAFETY: `child` is a live registry entry; `id` is writable. + if unsafe { IORegistryEntryGetRegistryEntryID(child.0, &mut id) } != KERN_SUCCESS { + return None; + } + out.push(RawClient { id, pid, gpu_ns }); + } + } + Some(out) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn the_creator_names_its_pid() { + assert_eq!(creator_pid("pid 172, WindowServer"), Some(172)); + assert_eq!(creator_pid("pid 99917, Google Chrome He"), Some(99_917)); + assert_eq!(creator_pid("pid 5"), Some(5)); + assert_eq!(creator_pid("172, WindowServer"), None); + assert_eq!(creator_pid("pid x, y"), None); + } + + /// On this host the registry walks, and every client is a distinct entry. + #[test] + fn the_registry_lists_gpu_clients_once_each() { + let clients = gpu_clients().expect("the I/O Registry's GPU clients"); + let mut ids: Vec = clients.iter().map(|c| c.id).collect(); + ids.sort_unstable(); + ids.dedup(); + eprintln!("gpu clients: {} (distinct ids {})", clients.len(), ids.len()); + assert!(!clients.is_empty(), "no GPU client at all: WindowServer holds some on any Mac with a display"); + assert_eq!(ids.len(), clients.len()); + } +} diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_ioreport.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_ioreport.rs index 9f6144367..a80744e7c 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_ioreport.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/apple_ioreport.rs @@ -304,43 +304,7 @@ impl AppleIoReportCounter { return opens.then(State::inert); } GROUP_OPENS.fetch_add(1, Ordering::SeqCst); - unsafe { - let group = CFStringCreateWithCString( - std::ptr::null_mut(), - c"Energy Model".as_ptr() as *const _, - K_CFSTRING_ENCODING_UTF8, - ); - if group.is_null() { - return None; - } - let desired = IOReportCopyChannelsInGroup(group, std::ptr::null(), 0, 0, 0); - CFRelease(group); - if desired.is_null() { - return None; - } - let mut subscribed: CFMutableDictionaryRef = std::ptr::null_mut(); - let sub = IOReportCreateSubscription( - std::ptr::null(), - desired, - &mut subscribed as *mut _, - 0, - std::ptr::null(), - ); - CFRelease(desired as CFTypeRef); - if sub.is_null() { - return None; - } - if subscribed.is_null() { - CFRelease(sub); - return None; - } - Some(State { - sub, - channels: subscribed, - prior_sample: None, - cumulative_uj: 0, - }) - } + subscribe_energy_model() } /// Whether this counter meters: IOReport opened **and** the host's CPU @@ -356,6 +320,50 @@ impl AppleIoReportCounter { } } +/// A subscription to the `Energy Model` group alone, or `None` if it cannot be +/// subscribed: the IOReport calls behind [`AppleIoReportCounter::open_energy_model`] +/// (the CPU-rail probe's opener, which counts and can script them) and +/// [`GpuRails::open`] (the GPU rail reader, which does neither). +fn subscribe_energy_model() -> Option { + unsafe { + let group = CFStringCreateWithCString( + std::ptr::null_mut(), + c"Energy Model".as_ptr() as *const _, + K_CFSTRING_ENCODING_UTF8, + ); + if group.is_null() { + return None; + } + let desired = IOReportCopyChannelsInGroup(group, std::ptr::null(), 0, 0, 0); + CFRelease(group); + if desired.is_null() { + return None; + } + let mut subscribed: CFMutableDictionaryRef = std::ptr::null_mut(); + let sub = IOReportCreateSubscription( + std::ptr::null(), + desired, + &mut subscribed as *mut _, + 0, + std::ptr::null(), + ); + CFRelease(desired as CFTypeRef); + if sub.is_null() { + return None; + } + if subscribed.is_null() { + CFRelease(sub); + return None; + } + Some(State { + sub, + channels: subscribed, + prior_sample: None, + cumulative_uj: 0, + }) + } +} + /// Whether a host IOReport status lets the counter meter: `Live`, and nothing /// else. `Unavailable` saw nothing. `CpuRailAbsent` has no CPU rail to meter. /// `CpuRailFrozen` would pass GPU-only energy, or zero, off as the energy of @@ -839,6 +847,101 @@ fn is_cpu_energy_rail(name: &str) -> bool { strip_die_prefix(name).eq_ignore_ascii_case("CPU Energy") } +// ── GPU rail (for `crate::process_gpu`) ───────────────────────────────────── + +/// Whether a channel is a GPU energy aggregate: `GPU Energy`, optionally +/// `DIE__`-prefixed. Not `GPU0` and its other summands: on macOS 27.0.1 +/// those are mJ rails, frozen for an unprivileged reader like every mJ rail, +/// while `GPU Energy` (nJ) advances. +fn is_gpu_energy_rail(name: &str) -> bool { + strip_die_prefix(name).eq_ignore_ascii_case("GPU Energy") +} + +/// A raw IOReport energy value in nanojoules, by its channel's unit label; +/// `None` for a negative value or a unit this does not know. Never guessed: a +/// rail in an unknown unit is no rail. +fn energy_to_nj(raw: i64, unit: &str) -> Option { + let raw = u64::try_from(raw).ok()?; + match unit.trim().to_ascii_lowercase().as_str() { + "pj" => Some(raw / 1_000), + "nj" => Some(raw), + "uj" | "µj" => raw.checked_mul(1_000), + "mj" => raw.checked_mul(1_000_000), + "j" => raw.checked_mul(1_000_000_000), + _ => None, + } +} + +/// The GPU energy rail(s) of the `Energy Model` group, read at their absolute +/// values: what the per-process GPU counter (`crate::process_gpu`) apportions. +/// Its own subscription to the group, not counted in [`GROUP_OPENS`] (which +/// counts the CPU-rail probe's) and never scripted. +pub(crate) struct GpuRails { + state: State, +} + +// The raw CF pointers in `State` are used by one thread at a time: the counter +// that owns a `GpuRails` holds it behind a `Mutex`. +unsafe impl Send for GpuRails {} + +impl GpuRails { + /// Subscribe to the `Energy Model` group, or `None` if it cannot be + /// subscribed. + pub(crate) fn open() -> Option { + subscribe_energy_model().map(|state| Self { state }) + } + + /// Every GPU energy rail in one sample, in sample order, as (name, absolute + /// nJ). `None` if the sample failed, held no GPU rail, or a GPU rail's unit + /// is not one [`energy_to_nj`] knows. + pub(crate) fn sample_nj(&self) -> Option> { + // SAFETY: `self.state` holds a live subscription and its channel + // dictionary (released only when it drops). The sample is released on + // every path; what is got from it is borrowed. + unsafe { + let smp = IOReportCreateSamples(self.state.sub, self.state.channels, std::ptr::null()); + if smp.is_null() { + return None; + } + let key = CFStringCreateWithCString( + std::ptr::null_mut(), + c"IOReportChannels".as_ptr() as *const _, + K_CFSTRING_ENCODING_UTF8, + ); + let arr = if key.is_null() { + std::ptr::null() + } else { + let arr = CFDictionaryGetValue(smp, key as *const _) as CFArrayRef; + CFRelease(key as CFTypeRef); + arr + }; + if arr.is_null() { + CFRelease(smp); + return None; + } + let mut rails = Vec::new(); + let mut unknown_unit = false; + for i in 0..CFArrayGetCount(arr) { + let item = CFArrayGetValueAtIndex(arr, i); + if item.is_null() || !cfstring_eq_ic(IOReportChannelGetGroup(item), "Energy Model") { + continue; + } + let Some(name) = cfstring_to_string(IOReportChannelGetChannelName(item)) else { continue }; + if !is_gpu_energy_rail(&name) { + continue; + } + let unit = cfstring_to_string(IOReportChannelGetUnitLabel(item)).unwrap_or_default(); + match energy_to_nj(IOReportSimpleGetIntegerValue(item, 0) as i64, &unit) { + Some(nj) => rails.push((name, nj)), + None => unknown_unit = true, + } + } + CFRelease(smp); + (!rails.is_empty() && !unknown_unit).then_some(rails) + } + } +} + /// The grade of every reading this counter produces: a hardware reading, the /// grade this crate gives the other vendor energy counters (`RaplCounter`, /// `NvmlCounter`). A vendor counter's reading is a real reading; its trust gap is @@ -1268,6 +1371,21 @@ mod tests { assert_eq!(energy_to_uj(5_000, ""), 5); } + /// The GPU rail reader scales by the unit label and refuses one it does + /// not know (never the nJ guess the summed counter falls back to), and + /// takes the `GPU Energy` aggregate alone. + #[test] + fn the_gpu_rail_is_the_aggregate_in_a_known_unit() { + assert_eq!(energy_to_nj(7, "nJ"), Some(7)); + assert_eq!(energy_to_nj(7, "mJ"), Some(7_000_000)); + assert_eq!(energy_to_nj(7, "uJ"), Some(7_000)); + assert_eq!(energy_to_nj(7_000, "pJ"), Some(7)); + assert_eq!(energy_to_nj(7, ""), None); + assert_eq!(energy_to_nj(-7, "nJ"), None); + assert!(is_gpu_energy_rail("GPU Energy") && is_gpu_energy_rail("DIE_1_GPU Energy")); + assert!(!is_gpu_energy_rail("GPU0") && !is_gpu_energy_rail("GPU SRAM") && !is_gpu_energy_rail("CPU Energy")); + } + /// The counter sums each aggregate once and none of its summands, on the /// real Energy Model channel names of both chips. #[test] @@ -2407,11 +2525,16 @@ mod tests { /// merged, and the outside channel never at all. #[test] fn the_production_status_of_two_dies_stepping_apart_is_frozen_and_each_die_reaches_the_rule_at_its_own_value() { + // Steps at samples 2 and 4: the probe always takes at least six + // samples (the rule's five intervals), so both steps are followed by + // another sample however slowly a loaded host runs its ticks. (They + // were at 3 and 6, and a CI run on a loaded host stopped at six + // samples, before the second step was followed.) fn die_0(i: usize) -> i64 { - 9_494_336 + if i >= 3 { 1_344_869 } else { 0 } + 9_494_336 + if i >= 2 { 1_344_869 } else { 0 } } fn die_1(i: usize) -> i64 { - 10_212_656 + if i >= 6 { 1_391_637 } else { 0 } + 10_212_656 + if i >= 4 { 1_391_637 } else { 0 } } let (status, rec, _) = run_scripted(|i| { let n = i as i64; @@ -2425,7 +2548,7 @@ mod tests { ]) }); assert_eq!(status, IoReportStatus::CpuRailFrozen, "two dies' separate steps, or an outside channel, made a live rail"); - assert!(rec.fed.len() >= 8, "DIE_1's step at sample 6 was never followed by another sample: {}", rec.fed.len()); + assert!(rec.fed.len() >= 6, "DIE_1's step at sample 4 was never followed by another sample: {}", rec.fed.len()); for (i, fed) in rec.fed.iter().enumerate() { let rails = cpu_rails(fed); let want = [(("DIE_0_CPU Energy".to_string(), 0), die_0(i)), (("DIE_1_CPU Energy".to_string(), 0), die_1(i))]; diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/lib.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/lib.rs index 819bc1292..c306fe7d3 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/lib.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/lib.rs @@ -37,6 +37,14 @@ //! wherever it runs, a subprocess included); [`meter::detect_in_process_meter`] //! prefers it, for work that runs on this process's own threads. //! +//! On Apple Silicon (macOS, the `ioreport` feature) it also reads this +//! process's own GPU energy (`process_gpu::ProcessGpuCounter`, domain +//! [`EnergyDomain::ProcessGpu`]): an exact zero where the process spent no GPU +//! time, the machine-wide IOReport GPU rail where it was effectively the GPU's +//! sole user (both `HwShunt`), and no figure where the GPU was shared. +//! [`span::ProcessSpanMeter`] meters a span of in-process work on this +//! process's CPU and GPU together, with its [`Coverage`]. +//! //! The thermodynamic circuit breaker is only as honest as the **worst** //! counter in a request's span. The breaker MUST surface //! [`EnergyCounter::resolution_uj`] and [`EnergyCounter::min_window_ns`] so @@ -71,6 +79,15 @@ pub mod apple; #[cfg(all(feature = "ioreport", target_os = "macos"))] pub mod apple_ioreport; +#[cfg(all(feature = "ioreport", target_os = "macos"))] +mod apple_gpu; + +// This process's own GPU energy (Apple Silicon), and a span of in-process work +// on its CPU and GPU together. Always compiled; the GPU counter is available +// only on macOS with the `ioreport` feature. (Their docs are their own `//!`.) +pub mod process_gpu; +pub mod span; + /// This process's own CPU energy (macOS `proc_pid_rusage`): public API, no /// feature needed. Not part of [`meter::detect_meter`]; see /// [`meter::detect_in_process_meter`]. @@ -228,6 +245,15 @@ pub fn process_cpu_probe_runs() -> usize { } } +/// How many times this process has run the real GPU liveness probe +/// ([`process_gpu::gpu_rail_liveness`]), on any thread. For the test that pins +/// that probe to one run per process; not part of the metering API. Always 0 +/// without the backend. +#[doc(hidden)] +pub fn process_gpu_probe_runs() -> usize { + process_gpu::probe_runs() +} + /// Hardware energy domain a counter reports on. #[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)] #[cfg_attr(feature = "serde", derive(Serialize, Deserialize))] @@ -256,6 +282,33 @@ pub enum EnergyDomain { /// figure over this domain covers work that ran on this process's threads, /// and nothing else. ProcessCpu, + /// The GPU energy of **this process** over a span + /// ([`process_gpu::ProcessGpuCounter`]): zero where it spent no GPU time, + /// the machine-wide GPU energy rail where it had at least 99% of every GPU + /// client's GPU time over the span. Not the CPU that encoded the work, not + /// the ANE, not DRAM. + ProcessGpu, +} + +/// How much of a unit of work a figure accounts for (the family's coverage +/// field: JCP's `Coverage`, with the same two values). +/// +/// A figure **MUST NOT** claim coverage of a domain its counter does not read. +/// A figure over this process's CPU and GPU ([`span::ProcessSpanMeter`]) is +/// `Whole` for work that ran there; DRAM is outside it, as it is outside a +/// RAPL package figure. Work its caller declares used the Neural Engine is +/// `Partial`: there is no live unprivileged ANE source on macOS 27.0.1. So is +/// a span's CPU figure alone, where its GPU has none (the GPU was shared, say). +#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Hash)] +#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))] +pub enum Coverage { + /// The figure accounts for the whole unit of work. + #[default] + Whole, + /// The figure accounts for only part of it. The remainder is + /// **unmeasured**, which is not zero, and a `Partial` figure MUST NOT be + /// presented as a task total. + Partial, } /// A single energy reading with full provenance. @@ -296,6 +349,38 @@ pub enum EnergyError { /// not see the work: that is not a reading, and never a zero. #[error("the counter did not advance across the work: not a reading")] NotAdvanced, + /// A counter's reading at a window's boundary was not current: the IOReport + /// GPU rail held while the GPU was busy (it can hold for hundreds of + /// milliseconds and then catch up in one step), so a window bounded by that + /// reading is short or long by an unknown amount. + #[error("the counter's reading at the window's boundary was not current: it held while the GPU was busy")] + NotCurrent, + /// This process used the GPU over a span, but had less than + /// [`process_gpu::SOLE_USER_PERCENT`] of every GPU client's GPU time: the + /// GPU was shared, and its machine-wide rail cannot be split by GPU time. + #[error( + "the GPU was shared: this process had {self_ns} ns of {all_ns} ns of GPU time; GPU energy is not \ + proportional to GPU time (measured: ~1.9 W vs ~17 W per GPU-second for two jobs), so the rail is not split" + )] + GpuShared { + /// This process's GPU time over the span, ns. + self_ns: u64, + /// Every client's GPU time over the span, ns. + all_ns: u64, + }, + /// A GPU client present when a span started was gone when it ended (its + /// process closed it or exited), or its accumulated GPU time went back. + /// If it was this process's, its own GPU time over the span is unknown; if + /// another's, every client's is, and so whether this process had the GPU + /// to itself: no figure. (A process that spent no GPU time has an exact + /// zero, whatever other clients did.) + #[error("GPU client {client:#x} (pid {pid}) was gone or reset at the span's end: every client's GPU time over it is unknown")] + GpuClientVanished { + /// The process that held the client. + pid: i32, + /// The client's I/O Registry entry ID. + client: u64, + }, /// Underlying platform error. #[error("platform error: {0}")] Platform(String), diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process.rs index d36d3be11..8502f1f63 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process.rs @@ -96,7 +96,7 @@ fn rusage_energy_nj(pid: libc::pid_t) -> Option { } /// This process's cumulative CPU energy, in nanojoules. -fn energy_nj() -> Option { +pub(crate) fn energy_nj() -> Option { // SAFETY: `getpid` has no preconditions and cannot fail. rusage_energy_nj(unsafe { libc::getpid() }) } @@ -132,7 +132,7 @@ fn probe_reader(read: impl Fn() -> Option) -> CounterLiveness { /// is not a zero; a counter that went backwards wrapped or was reset; and one /// that did not advance by a whole µJ — the resolution of the figure — did not /// see the work, which runs on more than zero energy. -fn window_uj(before_nj: Option, after_nj: Option) -> Result { +pub(crate) fn window_uj(before_nj: Option, after_nj: Option) -> Result { let (Some(before), Some(after)) = (before_nj, after_nj) else { return Err(EnergyError::Platform("proc_pid_rusage failed".to_string())); }; diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process_gpu.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process_gpu.rs new file mode 100644 index 000000000..c56ff511a --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/process_gpu.rs @@ -0,0 +1,668 @@ +//! This process's own GPU energy on Apple Silicon, [`EnergyDomain::ProcessGpu`] +//! and nothing else: an exact zero when it spent no GPU time over a span, the +//! machine-wide GPU energy rail when it was effectively the GPU's sole user, and +//! no figure otherwise. +//! +//! **The two sources** (measured on an M5 Max, macOS 27.0.1, unprivileged): +//! +//! - **The GPU rail.** IOReport's `Energy Model` group, channel `GPU Energy` +//! (nJ), advances live for an unprivileged reader under a Metal compute load, +//! where every mJ rail (`CPU Energy`, `GPU0`, `ANE0`, `DRAM0`, …) is frozen. +//! It is machine-wide: every client of the GPU on the host, display +//! compositing among them, and the GPU's idle and static draw. +//! - **Per-client GPU time.** The I/O Registry's `IOAccelerator` user clients, +//! one per connection a process opened to the GPU, each with its creator's +//! pid and an `accumulatedGPUTime` (ns), public IOKit API read unprivileged +//! (`crate::apple_gpu`). A process can hold several clients; they are summed. +//! +//! **The figure** ([`apportion`]). At a span's start and end the counter +//! snapshots the rail and every client's time ([`GpuSnapshot`]). +//! +//! - **This process spent no GPU time** (its own clients' time did not move; +//! a client of its own new at the end counts its whole time): the figure is +//! **0 µJ, exact**, read from the driver's accounting of this process's GPU +//! work. Nothing else is needed for it — not the rail, not any other client. +//! - **It was effectively the sole user** — at least [`SOLE_USER_PERCENT`] +//! (99%) of every client's GPU time over the span: the figure is the rail +//! over the span × `Δt_self / Δt_all`, within 1% of the rail's reading over a +//! span this process owned, the GPU's idle and static draw included as RAPL's +//! package figure includes its idle draw. +//! - **The GPU was shared**: no figure ([`EnergyError::GpuShared`]). The rail +//! cannot be split by GPU time: energy is not proportional to it. Measured: a +//! GGUF decode drew ~1.9 W per GPU-second alone, beside a job drawing ~17 W; +//! a time-share split charged the decode 3–15× its sole-user GPU energy. +//! +//! Every figure is a hardware reading ([`Provenance::HwShunt`]); there is no +//! apportioned grade. +//! +//! **No figure** also when this process's GPU time advanced and: +//! - a client present at the start is gone at the end, or its time went back +//! ([`EnergyError::GpuClientVanished`]) — for one of this process's own +//! clients, whatever else happened: its own time is unknown; for another +//! process's, because `Δt_all`, so the share, is unknown. A client that came +//! and went inside the span is invisible to both snapshots; +//! - the rail's reading at a boundary was not current +//! ([`EnergyError::NotCurrent`]): the rail at times holds for hundreds of +//! milliseconds and then catches up in one step, so each snapshot reads it on +//! its next change — or, where no GPU client accrues any GPU time for +//! [`GPU_IDLE_WINDOW`], as it stands, an idle GPU having next to nothing to +//! publish — and one that got neither within [`RAIL_EDGE_WAIT`] is not current; +//! - the rail did not advance ([`EnergyError::NotAdvanced`]): it did not see +//! the work; or it went back ([`EnergyError::CounterWrap`]); or a snapshot +//! failed. +//! +//! **Liveness.** The rail is a hardware reading only while it is demonstrably +//! live for the GPU domain, and so is this process's GPU time: the family's +//! rule ([`crate::CounterLiveness`]) — every channel advancing in at least five +//! intervals, four in every five, over at least 300 ms — over the GPU rail(s) +//! and this process's own GPU time, while the probe loads **the GPU**. Burning +//! CPU does not move either. This crate holds no Metal, so the load is the +//! caller's: a closure that dispatches a few tens of milliseconds of GPU work +//! and waits for it ([`gpu_rail_liveness`]). Probed once per process, with the +//! first caller's load, and cached. A counter that is not live meters nothing. +//! +//! **Not covered.** The Neural Engine has no unprivileged live source on macOS +//! 27.0.1 (`ANE0` is frozen; ANE work runs in the `aned` daemon, whose clients +//! carry no time or energy). DRAM is not in the GPU rail. See +//! [`crate::span`] for a span's CPU and GPU together, and its coverage. +//! +//! Off macOS, or without the `ioreport` feature, the counter is unavailable and +//! the probe never runs. + +use std::time::Instant; + +use crate::{CounterLiveness, EnergyDomain, EnergyError, Provenance}; + +/// The least share of every client's GPU time over a span, in percent, that +/// this process must have had for its GPU figure to be a hardware reading +/// ([`Provenance::HwShunt`]). +pub const SOLE_USER_PERCENT: u64 = 99; + +/// One GPU client in a [`GpuSnapshot`]. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub struct GpuClient { + /// The client's I/O Registry entry ID: never reused while the host is up. + pub id: u64, + /// The process that opened it. + pub pid: i32, + /// Its accumulated GPU time, ns. + pub gpu_ns: u64, +} + +/// The GPU rail(s) and every GPU client's accumulated GPU time, at one instant. +#[derive(Debug, Clone, PartialEq, Eq)] +pub struct GpuSnapshot { + /// Each GPU energy rail, in sample order: (channel name, absolute nJ). + pub rails_nj: Vec<(String, u64)>, + /// Every GPU client in the registry. + pub clients: Vec, + /// Whether the rails' reading is current: read just after they + /// republished, or while no GPU client accrued GPU time for + /// [`GPU_IDLE_WINDOW`] (an idle GPU has next to nothing to publish), with + /// the clients read right after. `false` when the rail held for + /// [`RAIL_EDGE_WAIT`] while the GPU was busy; such a snapshot bounds no + /// figure. + pub rail_current: bool, + /// When it was taken. + pub at: Instant, +} + +/// How long a snapshot waits for the GPU rail to republish before it reads +/// the clients. The rail is republished about every 30 ms, but at times holds +/// for 300-600 ms and then catches up in one step (M5 Max, macOS 27.0.1, a +/// heavily loaded host: 7 of 37 consecutive ~315 ms windows read zero, each +/// followed by a window of about twice the usual energy). A reading taken +/// mid-hold is short by what the hold has not yet published, so a snapshot +/// reads the rail on its next change, and the clients' times right after. +pub const RAIL_EDGE_WAIT: std::time::Duration = std::time::Duration::from_millis(1000); + +/// How long no GPU client may accrue GPU time for a snapshot to take a held +/// rail as current. On a quiet GPU the rail does not republish at all (an M5 +/// Max, macOS 27.0.1: no change in 4 s of 3 ms samples with no GPU work); an +/// idle GPU has at most its last republish interval of work, and its idle +/// draw, unpublished. +pub const GPU_IDLE_WINDOW: std::time::Duration = std::time::Duration::from_millis(100); + +/// This process's GPU energy over a span: a hardware reading +/// ([`Provenance::HwShunt`], [`Self::provenance`]), or an exact zero. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub struct GpuShare { + /// This process's GPU energy, µJ: 0 when it spent no GPU time; otherwise + /// the rail's advance × `self_ns / all_ns`, with that share at least + /// [`SOLE_USER_PERCENT`]. + pub uj: u64, + /// This process's GPU time over the span, ns (all its clients). + pub self_ns: u64, + /// The GPU rail's advance over the span, µJ: every client's GPU energy, and + /// the GPU's idle and static draw. `None` for an exact zero, which needs + /// no rail. + pub rail_uj: Option, + /// Every client's GPU time over the span, ns (this process's included). + /// `None` for an exact zero, which needs no other client. + pub all_ns: Option, +} + +impl GpuShare { + /// The grade of every GPU figure: a hardware reading. A zero is read from + /// the driver's own accounting of this process's GPU time; a non-zero + /// figure is the rail over a span this process all but owned. + pub fn provenance(&self) -> Provenance { + Provenance::HwShunt + } + + /// This process's share of every client's GPU time, in parts per million; + /// 0 for an exact zero. + pub fn share_ppm(&self) -> u64 { + match self.all_ns { + Some(all) if all > 0 => (u128::from(self.self_ns) * 1_000_000 / u128::from(all)) as u64, + _ => 0, + } + } + + /// The domain of [`Self::uj`]. + pub fn domain(&self) -> EnergyDomain { + EnergyDomain::ProcessGpu + } +} + +/// Each client's GPU time over the span, by registry ID: a client present at +/// both ends its advance, one new at the end its whole time. A client present +/// at the start and gone at the end, or whose time went back, is +/// [`EnergyError::GpuClientVanished`]. Only the clients `pick` selects are +/// summed or checked. +fn spent(start: &GpuSnapshot, end: &GpuSnapshot, pick: impl Fn(&GpuClient) -> bool) -> Result { + let mut ns: u64 = 0; + for c in start.clients.iter().filter(|c| pick(c)) { + let delta = end + .clients + .iter() + .find(|e| e.id == c.id) + .and_then(|e| e.gpu_ns.checked_sub(c.gpu_ns)) + .ok_or(EnergyError::GpuClientVanished { pid: c.pid, client: c.id })?; + ns = ns.saturating_add(delta); + } + for e in end.clients.iter().filter(|e| pick(e) && !start.clients.iter().any(|c| c.id == e.id)) { + ns = ns.saturating_add(e.gpu_ns); + } + Ok(ns) +} + +/// The span's GPU figure for process `pid`, from its start and end snapshots: +/// see the module docs for the rule and when there is none. Pure arithmetic +/// over the snapshots, so the tests drive every case. +/// +/// This process's own GPU time comes first, from its own clients alone (one of +/// them gone is [`EnergyError::GpuClientVanished`]). If it spent none, the +/// figure is an exact zero, whatever the rail or any other client did. Then +/// every other client must be accounted for, this process must have had at +/// least [`SOLE_USER_PERCENT`] of all their time ([`EnergyError::GpuShared`]), +/// and the rail must be current at both ends and have advanced. +pub fn apportion(start: &GpuSnapshot, end: &GpuSnapshot, pid: i32) -> Result { + let self_ns = spent(start, end, |c| c.pid == pid)?; + if self_ns == 0 { + return Ok(GpuShare { uj: 0, self_ns: 0, rail_uj: None, all_ns: None }); + } + + let all_ns = self_ns.saturating_add(spent(start, end, |c| c.pid != pid)?); + if u128::from(self_ns) * 100 < u128::from(all_ns) * u128::from(SOLE_USER_PERCENT) { + return Err(EnergyError::GpuShared { self_ns, all_ns }); + } + + // A rail read mid-hold is not current: no figure is bounded by it. + if !start.rail_current || !end.rail_current { + return Err(EnergyError::NotCurrent); + } + // The rail: each rail's own advance, matched by name and ordinal. + if start.rails_nj.is_empty() || start.rails_nj.len() != end.rails_nj.len() { + return Err(EnergyError::Platform("the GPU rails differ between the span's snapshots".to_string())); + } + let mut rail_nj: u64 = 0; + for ((n0, e0), (n1, e1)) in start.rails_nj.iter().zip(&end.rails_nj) { + if n0 != n1 { + return Err(EnergyError::Platform("the GPU rails differ between the span's snapshots".to_string())); + } + rail_nj = rail_nj.saturating_add(e1.checked_sub(*e0).ok_or(EnergyError::CounterWrap)?); + } + if rail_nj == 0 { + return Err(EnergyError::NotAdvanced); + } + let uj = (u128::from(rail_nj) * u128::from(self_ns) / u128::from(all_ns) / 1_000) as u64; + Ok(GpuShare { uj, self_ns, rail_uj: Some(rail_nj / 1_000), all_ns: Some(all_ns) }) +} + +#[cfg(all(feature = "ioreport", target_os = "macos"))] +mod backend { + use std::sync::atomic::{AtomicUsize, Ordering}; + use std::sync::{Mutex, OnceLock}; + use std::time::Instant; + + use super::{GpuClient, GpuSnapshot}; + use crate::apple_ioreport::GpuRails; + use crate::CounterLiveness; + + pub(crate) static REAL_RUNS: AtomicUsize = AtomicUsize::new(0); + + pub(super) fn this_pid() -> i32 { + // SAFETY: `getpid` has no preconditions and cannot fail. + unsafe { libc::getpid() } + } + + /// Every client now, or `None` if the registry could not be read. + fn clients() -> Option> { + crate::apple_gpu::gpu_clients() + .map(|raw| raw.into_iter().map(|c| GpuClient { id: c.id, pid: c.pid, gpu_ns: c.gpu_ns }).collect()) + } + + /// Every client's GPU time, summed. + fn total_ns(clients: &[GpuClient]) -> u64 { + clients.iter().map(|c| c.gpu_ns).fold(0, u64::saturating_add) + } + + /// The rails on their next change, or as they stand once no client has + /// accrued GPU time for [`super::GPU_IDLE_WINDOW`] (current either way), or + /// after [`super::RAIL_EDGE_WAIT`] (not current); then every client. + pub(super) fn snapshot(rails: &GpuRails) -> Option { + let t = Instant::now(); + let first = rails.sample_nj()?; + let mut idle_since = (Instant::now(), total_ns(&clients()?)); + let (rails_nj, rail_current) = loop { + std::thread::sleep(std::time::Duration::from_millis(2)); + let now = rails.sample_nj()?; + if now != first { + break (now, true); + } + if idle_since.0.elapsed() >= super::GPU_IDLE_WINDOW { + let total = total_ns(&clients()?); + if total == idle_since.1 { + break (now, true); + } + idle_since = (Instant::now(), total); + } + if t.elapsed() >= super::RAIL_EDGE_WAIT { + break (now, false); + } + }; + let clients = clients()?; + Some(GpuSnapshot { rails_nj, clients, rail_current, at: Instant::now() }) + } + + /// The real probe: the shared rule over the GPU rail(s) and this process's + /// own GPU time, with `load` between samples. Uncached; the cache is + /// [`cached_liveness`]. + pub(crate) fn probe(load: impl FnMut()) -> CounterLiveness { + let epoch = Instant::now(); + crate::liveness::probe_with(open_probe_sampler, load, move || epoch.elapsed(), crate::liveness::LIVENESS_BUDGET) + } + + /// The probe's opener: a GPU rail subscription and a sampler that hands the + /// rule each rail's nJ and, last, this process's summed GPU time in ns. + fn open_probe_sampler() -> Option<(usize, impl FnMut() -> Option>)> { + let rails = GpuRails::open()?; + let channels = rails.sample_nj()?.len() + 1; + let pid = this_pid(); + Some((channels, move || { + let mut sample: Vec = rails.sample_nj()?.into_iter().map(|(_, nj)| nj).collect(); + let own: u64 = clients()?.iter().filter(|c| c.pid == pid).map(|c| c.gpu_ns).sum(); + sample.push(own); + Some(sample) + })) + } + + pub(super) fn cached_liveness(load: impl FnMut()) -> CounterLiveness { + *LIVENESS.get_or_init(|| { + REAL_RUNS.fetch_add(1, Ordering::SeqCst); + probe(load) + }) + } + + pub(super) fn cached() -> Option { + LIVENESS.get().copied() + } + + static LIVENESS: OnceLock = OnceLock::new(); + + /// The counter's own subscription, behind a lock: IOReport subscriptions + /// are used by one thread at a time. + pub(super) type Rails = Mutex; + + pub(super) fn open_rails() -> Option { + GpuRails::open().map(Mutex::new) + } +} + +/// This host's GPU rail liveness, probed once per process and cached. The +/// first call, on whichever thread, runs the probe on that thread with `load` +/// as the work between samples: `load` must put a few tens of milliseconds of +/// work on the GPU **from this process** and wait for it (a Metal compute +/// dispatch, committed and waited on), since the probe also asks this process's +/// own GPU time to advance. The probe takes about 0.3 s where both are live, +/// about 1 s where they are not. Every later call ignores its `load` and +/// returns the cached answer. [`CounterLiveness::Unavailable`] without the +/// backend (off macOS, or without `ioreport`), where `load` is never called. +pub fn gpu_rail_liveness(load: impl FnMut()) -> CounterLiveness { + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + backend::cached_liveness(load) + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + let _ = load; + CounterLiveness::Unavailable + } +} + +/// The cached GPU rail liveness, if this process has probed it. +pub fn gpu_rail_liveness_cached() -> Option { + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + backend::cached() + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + None + } +} + +/// How many times this process has run the real GPU liveness probe. Read +/// through [`crate::process_gpu_probe_runs`]. +pub(crate) fn probe_runs() -> usize { + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + backend::REAL_RUNS.load(std::sync::atomic::Ordering::SeqCst) + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + 0 + } +} + +/// One uncached run of the real GPU liveness probe with `load`, for the +/// hardware tests (a CPU burn must not make the GPU counter live); not part of +/// the metering API. `Unavailable` without the backend. +#[doc(hidden)] +pub fn probe_gpu_liveness_uncached(load: impl FnMut()) -> CounterLiveness { + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + backend::probe(load) + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + let _ = load; + CounterLiveness::Unavailable + } +} + +/// Whether a liveness verdict lets the counter read: `Live`, and nothing else. +fn admits(liveness: CounterLiveness) -> bool { + match liveness { + CounterLiveness::Live => true, + CounterLiveness::Unavailable | CounterLiveness::Frozen | CounterLiveness::Implausible => false, + } +} + +/// This process's GPU energy counter (Apple Silicon): see the module docs. +pub struct ProcessGpuCounter { + liveness: CounterLiveness, + #[cfg(all(feature = "ioreport", target_os = "macos"))] + rails: Option, +} + +impl ProcessGpuCounter { + /// The counter, available only when the GPU rail and this process's GPU + /// time are live ([`gpu_rail_liveness`]; the first call in a process + /// probes with `load` on the calling thread). + pub fn new(load: impl FnMut()) -> Self { + Self::given(gpu_rail_liveness(load)) + } + + /// The counter for a host whose GPU liveness probe said `liveness`: a + /// subscription of its own to the rail when that admits it, else none. + fn given(liveness: CounterLiveness) -> Self { + if !admits(liveness) { + return Self::declined(liveness); + } + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + match backend::open_rails() { + Some(rails) => Self { liveness, rails: Some(rails) }, + None => Self::declined(CounterLiveness::Unavailable), + } + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + Self::declined(CounterLiveness::Unavailable) + } + } + + fn declined(liveness: CounterLiveness) -> Self { + Self { + liveness, + #[cfg(all(feature = "ioreport", target_os = "macos"))] + rails: None, + } + } + + /// Whether the counter meters: its sources were observed live. + pub fn is_available(&self) -> bool { + admits(self.liveness) && self.has_rails() + } + + fn has_rails(&self) -> bool { + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + self.rails.is_some() + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + false + } + } + + /// Why the counter is, or is not, available. + pub fn liveness(&self) -> CounterLiveness { + self.liveness + } + + /// The rail(s) and every client's GPU time now; [`EnergyError::NoCounter`] + /// when the counter is not available. + pub fn snapshot(&self) -> Result { + if !self.is_available() { + return Err(EnergyError::NoCounter(EnergyDomain::ProcessGpu)); + } + #[cfg(all(feature = "ioreport", target_os = "macos"))] + { + let rails = self.rails.as_ref().ok_or(EnergyError::NoCounter(EnergyDomain::ProcessGpu))?; + let rails = rails.lock().map_err(|_| EnergyError::Platform("GPU rail mutex poisoned".to_string()))?; + backend::snapshot(&rails) + .ok_or_else(|| EnergyError::Platform("a GPU rail or registry snapshot failed".to_string())) + } + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + { + Err(EnergyError::NoCounter(EnergyDomain::ProcessGpu)) + } + } + + /// This process's GPU figure between two of this counter's snapshots + /// ([`apportion`] for this process's pid). + pub fn share(&self, start: &GpuSnapshot, end: &GpuSnapshot) -> Result { + apportion(start, end, std::process::id() as i32) + } + + /// Run `work` and meter this process's GPU energy over it, or say why there + /// is no figure. `work` runs whether or not the counter is available. + pub fn measure(&self, work: impl FnOnce() -> T) -> (T, Result) { + let start = self.snapshot(); + let out = work(); + let figure = start.and_then(|start| self.snapshot().and_then(|end| self.share(&start, &end))); + (out, figure) + } + + /// Smallest meaningful delta: the rail is in nJ, figures are whole µJ. + pub fn resolution_uj(&self) -> u64 { + 1 + } + + /// Minimum window. Each snapshot reads the rail on a republish, so a + /// span's rail figure is off by at most about one republish interval + /// (~30 ms of GPU energy) at each end; over a second that is a few + /// percent. Meter whole decodes, not single tokens. + pub fn min_window_ns(&self) -> u64 { + 1_000_000_000 + } +} + +#[cfg(test)] +mod tests { + use super::*; + + const ME: i32 = 4242; + + fn client(id: u64, pid: i32, gpu_ns: u64) -> GpuClient { + GpuClient { id, pid, gpu_ns } + } + + fn snap(rail_nj: u64, clients: &[GpuClient]) -> GpuSnapshot { + GpuSnapshot { rails_nj: vec![("GPU Energy".to_string(), rail_nj)], clients: clients.to_vec(), rail_current: true, at: Instant::now() } + } + + /// The measured sole-user case: 2.707 s of 2.712 s against a 40.9 J rail. + #[test] + fn a_sole_user_is_a_hardware_reading_of_the_rail() { + let start = snap(1_000_000_000, &[client(1, ME, 5_000), client(2, 172, 9_000_000)]); + let end = snap(1_000_000_000 + 40_900_000_000, &[client(1, ME, 5_000 + 2_707_000_000), client(2, 172, 9_000_000 + 5_000_000)]); + let s = apportion(&start, &end, ME).expect("a figure"); + assert_eq!((s.self_ns, s.all_ns, s.rail_uj), (2_707_000_000, Some(2_712_000_000), Some(40_900_000))); + assert_eq!(s.uj, (40_900_000_000u128 * 2_707_000_000 / 2_712_000_000 / 1_000) as u64); + assert_eq!(s.provenance(), Provenance::HwShunt, "{s:?}"); + assert_eq!(s.share_ppm(), 998_156); + } + + /// The measured contended case: three processes at 20.4%, 25.2%, 54.1%. + /// None gets a figure: the GPU was shared, and the error says by how much. + #[test] + fn a_shared_gpu_gives_no_figure() { + let start = snap(0, &[client(1, 1, 0), client(2, 2, 0), client(3, 3, 0)]); + let end = snap(30_000_000_000, &[client(1, 1, 665_000_000), client(2, 2, 822_000_000), client(3, 3, 1_764_000_000)]); + for (pid, mine) in [(1, 665_000_000), (2, 822_000_000), (3, 1_764_000_000)] { + match apportion(&start, &end, pid) { + Err(EnergyError::GpuShared { self_ns, all_ns }) => assert_eq!((self_ns, all_ns), (mine, 3_251_000_000)), + other => panic!("pid {pid}: {other:?}"), + } + } + let why = EnergyError::GpuShared { self_ns: 1, all_ns: 9 }.to_string(); + assert!(why.contains("shared") && why.contains("not proportional") && why.contains("1.9 W"), "{why}"); + } + + /// The figure needs 99%; below it the GPU was shared, whatever the rail. + #[test] + fn a_figure_needs_ninety_nine_percent() { + let start = snap(0, &[client(1, ME, 0), client(2, 7, 0)]); + let at = |mine: u64, theirs: u64| apportion(&start, &snap(1_000_000, &[client(1, ME, mine), client(2, 7, theirs)]), ME); + assert_eq!(at(99, 1).map(|s| s.uj).ok(), Some(990)); + assert!(matches!(at(98_999, 1_001), Err(EnergyError::GpuShared { self_ns: 98_999, all_ns: 100_000 }))); + assert_eq!(at(1, 0).map(|s| s.uj).ok(), Some(1_000)); + assert_eq!(at(0, 0).map(|s| (s.uj, s.all_ns)).ok(), Some((0, None))); + } + + /// One of this process's own clients gone, or gone back, leaves its own + /// time unknown: no figure. Another's gone leaves the share unknown: no + /// figure while this process used the GPU. + #[test] + fn a_vanished_client_leaves_no_figure_while_this_process_used_the_gpu() { + let start = snap(0, &[client(1, ME, 0), client(2, 7, 50)]); + let other_gone = snap(5_000_000, &[client(1, ME, 1_000_000)]); + assert!(matches!(apportion(&start, &other_gone, ME), Err(EnergyError::GpuClientVanished { pid: 7, client: 2 }))); + let other_back = snap(5_000_000, &[client(1, ME, 1_000_000), client(2, 7, 49)]); + assert!(matches!(apportion(&start, &other_back, ME), Err(EnergyError::GpuClientVanished { pid: 7, client: 2 }))); + let mine_gone = snap(5_000_000, &[client(2, 7, 50)]); + assert!(matches!(apportion(&start, &mine_gone, ME), Err(EnergyError::GpuClientVanished { pid: ME, client: 1 }))); + let mine_back = GpuSnapshot { clients: vec![client(1, ME, 10), client(2, 7, 50)], ..start.clone() }; + let mine_back_end = snap(5_000_000, &[client(1, ME, 9), client(2, 7, 50)]); + assert!(matches!(apportion(&mine_back, &mine_back_end, ME), Err(EnergyError::GpuClientVanished { pid: ME, client: 1 }))); + } + + /// No GPU time of this process's is an exact zero, a hardware reading: + /// under contention, beside a vanished client, a rail that went back, + /// held, or is not current. None of those can change a zero. + #[test] + fn no_gpu_time_is_an_exact_zero_whatever_else_happened() { + let start = snap(1_000, &[client(1, ME, 10), client(2, 7, 10), client(3, 8, 0)]); + let zero = |s: &GpuSnapshot, e: &GpuSnapshot| { + let z = apportion(s, e, ME).expect("an exact zero"); + assert_eq!((z.uj, z.self_ns, z.rail_uj, z.all_ns, z.provenance()), (0, 0, None, None, Provenance::HwShunt)); + }; + zero(&start, &snap(9_000_000, &[client(1, ME, 10), client(2, 7, 500_000), client(3, 8, 0)])); + zero(&start, &snap(9_000_000, &[client(1, ME, 10)])); + zero(&start, &snap(999, &[client(1, ME, 10), client(2, 7, 5)])); + zero(&start, &GpuSnapshot { rail_current: false, ..snap(9_000, &[client(1, ME, 10)]) }); + zero(&GpuSnapshot { rails_nj: vec![], ..start.clone() }, &snap(9_000, &[client(1, ME, 10)])); + let no_client_of_mine = snap(1_000, &[client(2, 7, 10)]); + zero(&no_client_of_mine, &snap(1_000, &[client(2, 7, 99)])); + } + + /// A client new at the end spent its whole time inside the span; several + /// clients of one process are summed. + #[test] + fn new_clients_count_their_whole_time_and_a_process_sums_its_clients() { + let start = snap(0, &[client(1, ME, 100)]); + let end = snap(10_000_000, &[client(1, ME, 300), client(5, ME, 200), client(6, 9, 2)]); + let s = apportion(&start, &end, ME).unwrap(); + assert_eq!((s.self_ns, s.all_ns), (400, Some(402))); + assert_eq!(s.uj, 9_950); + let shared = snap(10_000_000, &[client(1, ME, 300), client(5, ME, 200), client(6, 9, 600)]); + assert!(matches!(apportion(&start, &shared, ME), Err(EnergyError::GpuShared { self_ns: 400, all_ns: 1_000 }))); + let only_new = snap(10_000_000, &[client(1, ME, 100), client(5, ME, 7)]); + assert_eq!(apportion(&start, &only_new, ME).map(|s| s.self_ns).ok(), Some(7)); + } + + /// This process's time advancing beside a rail that held is no figure; a + /// rail going back is a wrap; a rail not current at either end is none. + #[test] + fn a_sole_users_figure_needs_a_current_advancing_rail() { + let start = snap(1_000, &[client(1, ME, 10)]); + let used = |rail| snap(rail, &[client(1, ME, 900)]); + assert!(matches!(apportion(&start, &used(1_000), ME), Err(EnergyError::NotAdvanced))); + assert!(matches!(apportion(&start, &used(999), ME), Err(EnergyError::CounterWrap))); + assert_eq!(apportion(&start, &used(9_000_000), ME).map(|s| s.uj).ok(), Some(8_999)); + let held = |s: &GpuSnapshot| GpuSnapshot { rail_current: false, ..s.clone() }; + assert!(matches!(apportion(&held(&start), &used(9_000_000), ME), Err(EnergyError::NotCurrent))); + assert!(matches!(apportion(&start, &held(&used(9_000_000)), ME), Err(EnergyError::NotCurrent))); + } + + /// Two dies' rails each advance on their own and are summed; a rail missing + /// from one snapshot is no figure. + #[test] + fn every_rail_is_summed_and_must_be_in_both_snapshots() { + let rails = |a: u64, b: u64| vec![("DIE_0_GPU Energy".to_string(), a), ("DIE_1_GPU Energy".to_string(), b)]; + let start = GpuSnapshot { rails_nj: rails(10, 20), clients: vec![client(1, ME, 0)], rail_current: true, at: Instant::now() }; + let end = GpuSnapshot { rails_nj: rails(2_010, 3_020), clients: vec![client(1, ME, 9)], rail_current: true, at: Instant::now() }; + assert_eq!(apportion(&start, &end, ME).unwrap().uj, 5); + let one = GpuSnapshot { rails_nj: rails(2_010, 3_020)[..1].to_vec(), ..end.clone() }; + assert!(matches!(apportion(&start, &one, ME), Err(EnergyError::Platform(_)))); + } + + /// A counter given any verdict but `Live` meters nothing, and its work + /// still runs. + #[test] + fn a_counter_that_is_not_live_meters_nothing() { + for liveness in [CounterLiveness::Frozen, CounterLiveness::Implausible, CounterLiveness::Unavailable] { + let c = ProcessGpuCounter::given(liveness); + assert!(!c.is_available(), "{liveness:?}"); + assert!(matches!(c.snapshot(), Err(EnergyError::NoCounter(EnergyDomain::ProcessGpu)))); + let (out, figure) = c.measure(|| 7); + assert_eq!(out, 7); + assert!(matches!(figure, Err(EnergyError::NoCounter(EnergyDomain::ProcessGpu))), "{liveness:?}"); + } + } + + /// Without the backend nothing probes, and the load is never run. + #[test] + #[cfg(not(all(feature = "ioreport", target_os = "macos")))] + fn without_the_backend_the_counter_is_unavailable() { + let mut ran = false; + assert_eq!(gpu_rail_liveness(|| ran = true), CounterLiveness::Unavailable); + assert!(!ran); + assert!(!ProcessGpuCounter::new(|| {}).is_available()); + assert_eq!(probe_runs(), 0); + } +} diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/span.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/span.rs new file mode 100644 index 000000000..9998d45bb --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/src/span.rs @@ -0,0 +1,371 @@ +//! A span of in-process work — a Metal decode loop, a GPU kernel this process +//! dispatches and waits on — metered on this process's CPU and GPU together. +//! +//! The figure is this process's CPU energy (`process::ProcessCpuCounter`, +//! [`Provenance::HwShunt`] while live) plus its GPU energy +//! ([`ProcessGpuCounter`]: an exact zero where it spent no GPU time, the GPU +//! rail where it had the GPU to itself, `HwShunt` both) — +//! [`SpanSource::Process`]. Its coverage is this process's CPU and GPU +//! ([`Coverage::Whole`]; DRAM is outside it, as it is outside a RAPL package +//! figure), or [`Coverage::Partial`] for work its caller declares used the +//! Neural Engine, which has no live unprivileged source. +//! +//! Where the GPU counter puts no figure on the span — the GPU was shared, a GPU +//! client vanished, the rail was not current or did not advance, the counter +//! is not live — the span is **its CPU figure alone**: `HwShunt`, +//! [`Coverage::Partial`], and why the GPU has none ([`SpanSource::CpuOnly`]). +//! Only where the CPU counter itself gives no figure does the span fall back to +//! the host meter ([`crate::meter::detect_meter`]: on macOS 27.0.1 the +//! load-average estimator, [`Provenance::Estimator`]), recording why +//! ([`SpanSource::Fallback`]). +//! +//! The CPU half reads every thread of this process, so a span counts whatever +//! else this process runs at the same time; the GPU half every GPU client of +//! this process. + +use std::time::Instant; + +use crate::meter::{detect_meter, CompositeMeter}; +use crate::process_gpu::{GpuShare, GpuSnapshot, ProcessGpuCounter}; +use crate::{Coverage, EnergyError, EnergyReading, Provenance}; + +/// What a span of work ran on, as its caller declares it. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub enum SpanUses { + /// This process's CPU and GPU, and nothing else. + CpuGpu, + /// This process's CPU and GPU, and the Neural Engine (Core ML, say): the + /// figure is [`Coverage::Partial`], the ANE's energy unmeasured. + CpuGpuAne, +} + +/// Where a span's figure came from. +#[derive(Debug, Clone, PartialEq, Eq)] +pub enum SpanSource { + /// This process's own counters, CPU and GPU. + Process { + /// This process's CPU energy, µJ ([`crate::EnergyDomain::ProcessCpu`], + /// a hardware reading). + cpu_uj: u64, + /// This process's GPU energy ([`crate::EnergyDomain::ProcessGpu`]). + gpu: GpuShare, + }, + /// This process's CPU counter alone: the GPU counter gave no figure. The + /// span's figure is [`Coverage::Partial`]: whatever GPU energy the work + /// drew is unmeasured, which is not zero. + CpuOnly { + /// This process's CPU energy, µJ. + cpu_uj: u64, + /// Why the GPU has no figure. + why: String, + }, + /// The host meter, because this process's CPU counter put no figure on + /// the span. + Fallback { + /// The fallback meter's name (`"apple-silicon"`, `"estimation"`, …). + meter: &'static str, + /// Why this process's counters gave no figure. + why: String, + }, +} + +/// The energy of a span of in-process work. +#[derive(Debug, Clone, PartialEq, Eq)] +pub struct SpanEnergy { + /// The span's energy, µJ. + pub uj: u64, + /// The weakest grade of the readings summed. + pub provenance: Provenance, + /// How much of the work the figure accounts for. + pub coverage: Coverage, + /// Where the figure came from. + pub source: SpanSource, +} + +impl SpanEnergy { + /// The CPU and GPU parts, when this process's counters metered both. + pub fn parts(&self) -> Option<(u64, &GpuShare)> { + match &self.source { + SpanSource::Process { cpu_uj, gpu } => Some((*cpu_uj, gpu)), + SpanSource::CpuOnly { .. } | SpanSource::Fallback { .. } => None, + } + } + + /// This process's CPU energy, when its CPU counter metered the span. + pub fn cpu_uj(&self) -> Option { + match &self.source { + SpanSource::Process { cpu_uj, .. } | SpanSource::CpuOnly { cpu_uj, .. } => Some(*cpu_uj), + SpanSource::Fallback { .. } => None, + } + } +} + +/// The weaker of two grades: `Estimator` < `ModelBased` < `HwShunt`. +pub fn weaker(a: Provenance, b: Provenance) -> Provenance { + fn rank(p: Provenance) -> u8 { + match p { + Provenance::HwShunt => 2, + Provenance::ModelBased => 1, + Provenance::Estimator => 0, + } + } + if rank(a) <= rank(b) { a } else { b } +} + +/// The readings taken at a span's start ([`ProcessSpanMeter::start`]). +pub struct SpanStart { + cpu_nj: Option, + gpu: Result, + fallback: Result, + at: Instant, +} + +impl SpanStart { + /// When the span started. + pub fn at(&self) -> Instant { + self.at + } +} + +/// The combination rule, over the two counters' outcomes and the fallback's +/// readings: the counters' sum when both gave a figure; the CPU's alone, +/// `Partial`, with why, when the GPU gave none; else the fallback's window with +/// why. Pure, so the tests drive every case. +pub(crate) fn combine( + cpu: Result, + gpu: Result, + uses: SpanUses, + fallback_meter: &'static str, + fallback: (Result, Result), +) -> Result { + match (cpu, gpu) { + (Ok(cpu_uj), Ok(gpu)) => Ok(SpanEnergy { + uj: cpu_uj.saturating_add(gpu.uj), + provenance: weaker(Provenance::HwShunt, gpu.provenance()), + coverage: match uses { + SpanUses::CpuGpu => Coverage::Whole, + SpanUses::CpuGpuAne => Coverage::Partial, + }, + source: SpanSource::Process { cpu_uj, gpu }, + }), + (Ok(cpu_uj), Err(g)) => Ok(SpanEnergy { + uj: cpu_uj, + provenance: Provenance::HwShunt, + coverage: Coverage::Partial, + source: SpanSource::CpuOnly { cpu_uj, why: format!("process GPU: {g}") }, + }), + (Err(c), gpu) => { + let why = match gpu { + Err(g) => format!("process CPU: {c}; process GPU: {g}"), + Ok(_) => format!("process CPU: {c}"), + }; + let (r0, r1) = fallback; + let window = r0.and_then(|r0| { + r1.and_then(|r1| match r1.uj.checked_sub(r0.uj) { + None => Err(EnergyError::CounterWrap), + Some(0) => Err(EnergyError::NotAdvanced), + Some(uj) => Ok((uj, weaker(r0.provenance, r1.provenance))), + }) + }); + match window { + Ok((uj, provenance)) => Ok(SpanEnergy { + uj, + provenance, + coverage: Coverage::Whole, + source: SpanSource::Fallback { meter: fallback_meter, why }, + }), + Err(e) => Err(EnergyError::Platform(format!("{why}; and the fallback meter `{fallback_meter}`: {e}"))), + } + } + } +} + +/// Meters spans of in-process work on this process's CPU and GPU together, +/// falling back to the host meter where they give no figure. See the module +/// docs. +pub struct ProcessSpanMeter { + #[cfg(target_os = "macos")] + cpu: crate::process::ProcessCpuCounter, + gpu: ProcessGpuCounter, + fallback: CompositeMeter, +} + +impl ProcessSpanMeter { + /// The meter. The first one in a process runs the CPU counter's liveness + /// probe (about 0.3 s of this thread's CPU), the GPU counter's with + /// `gpu_load` ([`crate::process_gpu::gpu_rail_liveness`]: `gpu_load` must + /// put a few tens of milliseconds of work on the GPU from this process and + /// wait for it) and the host meter's; later ones reuse the answers. + pub fn new(gpu_load: impl FnMut()) -> Self { + Self { + #[cfg(target_os = "macos")] + cpu: crate::process::ProcessCpuCounter::new(), + gpu: ProcessGpuCounter::new(gpu_load), + fallback: detect_meter(), + } + } + + /// Whether this process's own counters meter (both live). When `false`, + /// every span falls back to the host meter. + pub fn is_metered(&self) -> bool { + self.cpu_available() && self.gpu.is_available() + } + + fn cpu_available(&self) -> bool { + #[cfg(target_os = "macos")] + { + self.cpu.is_available() + } + #[cfg(not(target_os = "macos"))] + { + false + } + } + + /// The GPU counter. + pub fn gpu(&self) -> &ProcessGpuCounter { + &self.gpu + } + + /// The fallback meter's name. + pub fn fallback_name(&self) -> &'static str { + self.fallback.active_name() + } + + /// Take the span's starting readings. The GPU snapshot comes first and the + /// CPU reading last (and the reverse in [`Self::finish`]), so the CPU + /// time the snapshots spend — a wait for the GPU rail to republish, and + /// IOReport samples — stays outside the CPU window. + pub fn start(&self) -> SpanStart { + let fallback = self.fallback.read().map_err(|e| EnergyError::Platform(e.to_string())); + let gpu = self.gpu.snapshot(); + SpanStart { cpu_nj: self.cpu_nj(), gpu, fallback, at: Instant::now() } + } + + /// End the span begun at `start`: its energy, or why there is none. + pub fn finish(&self, start: SpanStart, uses: SpanUses) -> Result { + let cpu_end = self.cpu_nj(); + let gpu_end = self.gpu.snapshot(); + let fallback_end = self.fallback.read().map_err(|e| EnergyError::Platform(e.to_string())); + let cpu = self.cpu_window(start.cpu_nj, cpu_end); + let gpu = start.gpu.and_then(|s| gpu_end.and_then(|e| self.gpu.share(&s, &e))); + combine(cpu, gpu, uses, self.fallback.active_name(), (start.fallback, fallback_end)) + } + + /// Run `work`, which ran on what `uses` says, and meter it. + pub fn measure(&self, uses: SpanUses, work: impl FnOnce() -> T) -> (T, Result) { + let start = self.start(); + let out = work(); + (out, self.finish(start, uses)) + } + + fn cpu_nj(&self) -> Option { + #[cfg(target_os = "macos")] + { + self.cpu.is_available().then(crate::process::energy_nj).flatten() + } + #[cfg(not(target_os = "macos"))] + { + None + } + } + + fn cpu_window(&self, start: Option, end: Option) -> Result { + #[cfg(target_os = "macos")] + { + if !self.cpu.is_available() { + return Err(EnergyError::NoCounter(crate::EnergyDomain::ProcessCpu)); + } + crate::process::window_uj(start, end) + } + #[cfg(not(target_os = "macos"))] + { + let _ = (start, end); + Err(EnergyError::NoCounter(crate::EnergyDomain::ProcessCpu)) + } + } +} + +#[cfg(test)] +mod tests { + use super::*; + use crate::EnergyDomain; + + fn share(uj: u64) -> GpuShare { + GpuShare { uj, self_ns: 1, rail_uj: Some(uj), all_ns: Some(1) } + } + + fn est(uj: u64) -> Result { + Ok(EnergyReading { uj, timestamp_ns: 0, domain: EnergyDomain::SocTotal, provenance: Provenance::Estimator }) + } + + /// Both counters: the sum, a hardware reading, `Whole` for CPU and GPU + /// work and `Partial` for work that used the ANE. An exact GPU zero keeps + /// the CPU's grade. + #[test] + fn both_counters_sum_as_a_hardware_reading() { + let sole = combine(Ok(300), Ok(share(700)), SpanUses::CpuGpu, "est", (est(0), est(9))).unwrap(); + assert_eq!((sole.uj, sole.provenance, sole.coverage), (1_000, Provenance::HwShunt, Coverage::Whole)); + assert_eq!(sole.parts().map(|(c, g)| (c, g.uj)), Some((300, 700))); + let zero = GpuShare { uj: 0, self_ns: 0, rail_uj: None, all_ns: None }; + let cpu_work = combine(Ok(300), Ok(zero), SpanUses::CpuGpu, "est", (est(0), est(9))).unwrap(); + assert_eq!((cpu_work.uj, cpu_work.provenance, cpu_work.coverage), (300, Provenance::HwShunt, Coverage::Whole)); + let ane = combine(Ok(300), Ok(share(700)), SpanUses::CpuGpuAne, "est", (est(0), est(9))).unwrap(); + assert_eq!((ane.provenance, ane.coverage), (Provenance::HwShunt, Coverage::Partial)); + } + + /// No GPU figure — a shared GPU, a vanished client, a counter not live — + /// leaves the CPU figure alone: a hardware reading, `Partial`, saying why. + /// The fallback meter is not consulted, even where it has a figure. + #[test] + fn no_gpu_figure_is_the_cpu_alone_and_partial() { + for (g, says) in [ + (EnergyError::GpuShared { self_ns: 115, all_ns: 1_000 }, "shared"), + (EnergyError::GpuClientVanished { pid: 7, client: 0x10 }, "pid 7"), + (EnergyError::NoCounter(EnergyDomain::ProcessGpu), "ProcessGpu"), + (EnergyError::NotCurrent, "not current"), + ] { + let e = combine(Ok(300), Err(g), SpanUses::CpuGpu, "apple-silicon", (est(1_000), est(5_000))).unwrap(); + assert_eq!((e.uj, e.provenance, e.coverage), (300, Provenance::HwShunt, Coverage::Partial)); + match &e.source { + SpanSource::CpuOnly { cpu_uj, why } => { + assert_eq!(*cpu_uj, 300); + assert!(why.contains("process GPU") && why.contains(says), "{why}"); + } + other => panic!("{other:?}"), + } + assert!(e.parts().is_none()); + assert_eq!(e.cpu_uj(), Some(300)); + } + let shared = combine(Ok(300), Err(EnergyError::GpuShared { self_ns: 115, all_ns: 1_000 }), SpanUses::CpuGpu, "x", (est(1), est(2))).unwrap(); + assert!(matches!(&shared.source, SpanSource::CpuOnly { why, .. } if why.contains("115 ns of 1000 ns")), "{shared:?}"); + } + + /// Only a CPU counter with no figure sends the span to the fallback meter, + /// with why; no fallback figure either is an error naming both. + #[test] + fn no_cpu_figure_falls_back_and_says_why() { + let e = combine(Err(EnergyError::NotAdvanced), Ok(share(1)), SpanUses::CpuGpu, "apple-silicon", (est(1_000), est(5_000))).unwrap(); + assert_eq!((e.uj, e.provenance, e.coverage), (4_000, Provenance::Estimator, Coverage::Whole)); + match &e.source { + SpanSource::Fallback { meter, why } => { + assert_eq!(*meter, "apple-silicon"); + assert!(why.contains("process CPU"), "{why}"); + } + other => panic!("{other:?}"), + } + assert!(e.parts().is_none() && e.cpu_uj().is_none()); + let both = combine(Err(EnergyError::NotAdvanced), Err(EnergyError::NotCurrent), SpanUses::CpuGpu, "x", (est(1), est(2))).unwrap(); + assert!(matches!(&both.source, SpanSource::Fallback { why, .. } if why.contains("process GPU")), "{both:?}"); + let none = combine(Err(EnergyError::NotAdvanced), Err(EnergyError::NotCurrent), SpanUses::CpuGpu, "x", (est(5), est(5))).unwrap_err(); + assert!(none.to_string().contains("process CPU") && none.to_string().contains("`x`"), "{none}"); + } + + #[test] + fn the_weaker_grade() { + use Provenance::*; + assert_eq!(weaker(HwShunt, ModelBased), ModelBased); + assert_eq!(weaker(Estimator, HwShunt), Estimator); + assert_eq!(weaker(HwShunt, HwShunt), HwShunt); + } +} diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/tests/process_gpu.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/tests/process_gpu.rs new file mode 100644 index 000000000..4f38760cc --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-energy/tests/process_gpu.rs @@ -0,0 +1,624 @@ +//! This process's own GPU energy counter on real hardware, asserted in a process +//! of its own against oracles the test reads itself. +//! +//! The oracles take independent paths to the counter's two sources: the GPU +//! rail through this file's own IOReport calls, and every client's GPU time +//! through `ioreg -r -c IOAccelerator -l` (another process, parsing its text). +//! A third is the test's own Metal command buffers' GPU start and end times. +//! The GPU is shared with whatever else the host runs, so every assertion holds +//! under contention: the oracle's reads sit inside the counter's span, and the +//! counter must agree with them on this process's GPU time exactly, and on every +//! client's to within what the GPU can accrue in the gap between them. +//! +//! One test function, its checks in sequence on one thread; it also serves as +//! the GPU-burning child process the contention and vanished-client checks +//! spawn (run with `JOULECLAW_GPU_BURN_CHILD` set, it burns and returns). Off +//! macOS, or without `ioreport`, it is not compiled; the crate's unit tests +//! assert that the counter is unavailable there. + +#![cfg(all(target_os = "macos", feature = "ioreport"))] + +use std::collections::BTreeMap; +use std::ffi::{c_char, c_long, c_void}; +use std::io::{BufRead, BufReader}; +use std::process::{Child, Command, Stdio}; +use std::ptr::NonNull; +use std::time::{Duration, Instant}; + +use jouleclaw_energy::process_gpu::{ + gpu_rail_liveness_cached, probe_gpu_liveness_uncached, GpuShare, GpuSnapshot, ProcessGpuCounter, +}; +use jouleclaw_energy::span::{ProcessSpanMeter, SpanSource, SpanUses}; +use jouleclaw_energy::{CounterLiveness, Coverage, EnergyDomain, EnergyError, Provenance}; +use objc2::rc::Retained; +use objc2::runtime::ProtocolObject; +use objc2_foundation::NSString; +use objc2_metal::{ + MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue, MTLComputeCommandEncoder, + MTLComputePipelineState, MTLCreateSystemDefaultDevice, MTLDevice, MTLLibrary, MTLResourceOptions, MTLSize, +}; + +const CHILD_ENV: &str = "JOULECLAW_GPU_BURN_CHILD"; +const READY: &str = "GPU-BURN-CHILD-READY"; + +// ── The GPU load: the test's own Metal ────────────────────────────────────── + +const BURN_MSL: &str = "#include \nusing namespace metal;\n\ +kernel void burn(device float *b [[buffer(0)]], constant uint &n [[buffer(1)]], uint i [[thread_position_in_grid]]) {\n\ + float x = b[i]; for (uint k = 0; k < n; k++) { x = fma(x, 1.0000001f, 0.0000001f); } b[i] = x; }\n"; +const THREADS: usize = 1 << 20; + +/// A compute kernel that keeps the GPU busy, one command buffer at a time. +struct Burner { + queue: Retained>, + pso: Retained>, + buf: Retained>, + iters: u32, +} + +impl Burner { + /// Compile the kernel, then size one command buffer to about 8 ms of GPU + /// time. + fn new() -> Self { + let device = MTLCreateSystemDefaultDevice().expect("a Metal device"); + let lib = device.newLibraryWithSource_options_error(&NSString::from_str(BURN_MSL), None).expect("compile"); + let f = lib.newFunctionWithName(&NSString::from_str("burn")).expect("the kernel"); + let pso = device.newComputePipelineStateWithFunction_error(&f).expect("a pipeline"); + let buf = device.newBufferWithLength_options(THREADS * 4, MTLResourceOptions::StorageModeShared).expect("a buffer"); + let queue = device.newCommandQueue().expect("a queue"); + let mut b = Self { queue, pso, buf, iters: 64 }; + let t = (0..3).map(|_| b.dispatch()).fold(f64::MAX, f64::min).max(1e-5); + b.iters = ((64.0 * 0.008 / t) as u32).clamp(64, 1 << 20); + b + } + + /// One command buffer, committed and waited on: its GPU time (end minus + /// start on the GPU's clock), in seconds. + fn dispatch(&self) -> f64 { + let cb = self.queue.commandBuffer().expect("a command buffer"); + let enc = cb.computeCommandEncoder().expect("an encoder"); + enc.setComputePipelineState(&self.pso); + // SAFETY: buffer 0 is THREADS floats, buffer 1 one u32, as the kernel + // declares; `iters` outlives the call, which copies it. + unsafe { + enc.setBuffer_offset_atIndex(Some(&self.buf), 0, 0); + enc.setBytes_length_atIndex(NonNull::from(&self.iters).cast(), 4, 1); + } + let tg = self.pso.maxTotalThreadsPerThreadgroup().min(THREADS); + enc.dispatchThreads_threadsPerThreadgroup( + MTLSize { width: THREADS, height: 1, depth: 1 }, + MTLSize { width: tg, height: 1, depth: 1 }, + ); + enc.endEncoding(); + cb.commit(); + cb.waitUntilCompleted(); + cb.GPUEndTime() - cb.GPUStartTime() + } + + /// Dispatch until `gpu` of command-buffer GPU time has run, or `wall` has + /// passed: the GPU time run, in seconds. + fn burn(&self, gpu: Duration, wall: Duration) -> f64 { + self.burn_counted(gpu, wall).0 + } + + /// [`Self::burn`], and how many command buffers it took. + fn burn_counted(&self, gpu: Duration, wall: Duration) -> (f64, u32) { + let t = Instant::now(); + let (mut ran, mut n) = (0.0, 0); + while ran < gpu.as_secs_f64() && t.elapsed() < wall { + ran += self.dispatch(); + n += 1; + } + (ran, n) + } +} + +/// The liveness probe's load: about 25 ms of this process's GPU time. +fn gpu_tick(b: &Burner) { + b.burn(Duration::from_millis(25), Duration::from_millis(500)); +} + +/// A CPU load of about 25 ms of this thread's CPU time. +fn cpu_tick() { + let t = Instant::now(); + let mut x = 1u64; + while t.elapsed() < Duration::from_millis(25) { + for _ in 0..10_000 { + x = std::hint::black_box(x.wrapping_mul(6_364_136_223_846_793_005).wrapping_add(1)); + } + } +} + +// ── Oracle 1: the GPU rail, through this file's own IOReport calls ───────── + +type CFRef = *const c_void; +unsafe extern "C" { + fn IOReportCopyChannelsInGroup(g: CFRef, sg: CFRef, a: u64, b: u64, c: u64) -> *mut c_void; + fn IOReportCreateSubscription(a: CFRef, ch: *mut c_void, out: *mut *mut c_void, id: u64, o: CFRef) -> CFRef; + fn IOReportCreateSamples(sub: CFRef, ch: *mut c_void, o: CFRef) -> CFRef; + fn IOReportChannelGetChannelName(item: CFRef) -> CFRef; + fn IOReportChannelGetUnitLabel(item: CFRef) -> CFRef; + fn IOReportSimpleGetIntegerValue(item: CFRef, i: i32) -> i64; + fn CFStringCreateWithCString(a: *mut c_void, s: *const c_char, enc: u32) -> CFRef; + fn CFStringGetCString(s: CFRef, buf: *mut c_char, n: c_long, enc: u32) -> bool; + fn CFDictionaryGetValue(d: CFRef, k: CFRef) -> CFRef; + fn CFArrayGetCount(a: CFRef) -> c_long; + fn CFArrayGetValueAtIndex(a: CFRef, i: c_long) -> CFRef; + fn CFRelease(r: CFRef); +} +#[link(name = "IOReport", kind = "dylib")] +#[link(name = "CoreFoundation", kind = "framework")] +unsafe extern "C" {} + +fn cf_text(s: CFRef) -> String { + let mut buf = [0 as c_char; 256]; + // SAFETY: `s` is a live CFString or null (then nothing is read). + if s.is_null() || !unsafe { CFStringGetCString(s, buf.as_mut_ptr(), 256, 0x0800_0100) } { + return String::new(); + } + // SAFETY: on success the buffer holds a NUL-terminated string. + unsafe { std::ffi::CStr::from_ptr(buf.as_ptr()) }.to_string_lossy().into_owned() +} + +/// The test's own subscription to the `Energy Model` group. +struct OracleRail { + sub: CFRef, + channels: *mut c_void, +} + +impl OracleRail { + fn open() -> Self { + // SAFETY: the IOReport calls `ioreg`/`powermetrics` make; every object + // is kept for the test's life or released. + unsafe { + let group = CFStringCreateWithCString(std::ptr::null_mut(), c"Energy Model".as_ptr(), 0x0800_0100); + let wanted = IOReportCopyChannelsInGroup(group, std::ptr::null(), 0, 0, 0); + let mut channels: *mut c_void = std::ptr::null_mut(); + let sub = IOReportCreateSubscription(std::ptr::null(), wanted, &mut channels, 0, std::ptr::null()); + assert!(!sub.is_null() && !channels.is_null(), "the oracle's Energy Model subscription"); + Self { sub, channels } + } + } + + /// The `GPU Energy` channel's absolute value, in nJ (it is the only unit it + /// has been seen in; anything else fails the test). + fn gpu_nj(&self) -> u64 { + // SAFETY: a live subscription; the sample is released. + unsafe { + let s = IOReportCreateSamples(self.sub, self.channels, std::ptr::null()); + assert!(!s.is_null(), "the oracle's sample"); + let key = CFStringCreateWithCString(std::ptr::null_mut(), c"IOReportChannels".as_ptr(), 0x0800_0100); + let arr = CFDictionaryGetValue(s, key); + CFRelease(key); + let mut found = None; + for i in 0..CFArrayGetCount(arr) { + let ch = CFArrayGetValueAtIndex(arr, i); + if cf_text(IOReportChannelGetChannelName(ch)) == "GPU Energy" { + assert_eq!(cf_text(IOReportChannelGetUnitLabel(ch)), "nJ", "the GPU rail's unit"); + found = Some(IOReportSimpleGetIntegerValue(ch, 0) as u64); + } + } + CFRelease(s); + found.expect("a GPU Energy channel") + } + } +} + +// ── Oracle 2: every client's GPU time, from `ioreg`'s text ───────────────── + +/// Every GPU client `ioreg` lists: registry id → (pid, summed +/// `accumulatedGPUTime` ns). +fn ioreg_clients() -> BTreeMap { + let out = Command::new("/usr/sbin/ioreg").args(["-r", "-c", "IOAccelerator", "-l", "-w0"]).output().expect("ioreg"); + assert!(out.status.success(), "ioreg failed"); + let text = String::from_utf8_lossy(&out.stdout); + let mut clients = BTreeMap::new(); + for block in text.split("+-o ").skip(1) { + let Some(creator) = block.split("\"IOUserClientCreator\" = \"pid ").nth(1) else { continue }; + let pid: i32 = creator.split(',').next().unwrap().trim().parse().expect("a pid"); + let id_hex = block.split(", id 0x").nth(1).and_then(|s| s.split(',').next()).expect("an entry id"); + let id = u64::from_str_radix(id_hex, 16).expect("a hex id"); + let ns: u64 = block + .lines() + .find(|l| l.contains("\"AppUsage\" = ")) + .map(|l| { + l.split("\"accumulatedGPUTime\"=").skip(1).map(|s| { + s.chars().take_while(char::is_ascii_digit).collect::().parse::().expect("a time") + }).sum() + }) + .unwrap_or(0); + clients.insert(id, (pid, ns)); + } + assert!(!clients.is_empty(), "ioreg listed no GPU client"); + clients +} + +/// One oracle reading: the rail and every client. +struct Oracle { + rail_nj: u64, + /// When the rail was read: on its first change after the oracle began + /// watching it, so the reading is current (the rail can hold for hundreds + /// of milliseconds and then catch up in one step). + rail_at: Instant, + clients: BTreeMap, + /// When `ioreg` was started and when it had finished: its readings lie + /// between. + clients_from: Instant, + clients_to: Instant, +} + +fn oracle(rail: &OracleRail) -> Oracle { + // The rail on its next change; or, if it holds while `ioreg` shows no + // client accruing GPU time between two of its listings, as it stands (an + // idle GPU has next to nothing to publish). + let first = rail.gpu_nj(); + let waited = Instant::now(); + let mut listed: Option<(Instant, u64)> = None; + let rail_nj = loop { + std::thread::sleep(Duration::from_millis(1)); + let now = rail.gpu_nj(); + if now != first { + break now; + } + if waited.elapsed() >= Duration::from_millis(150) { + let total: u64 = ioreg_clients().values().map(|c| c.1).sum(); + match listed { + Some((at, before)) if at.elapsed() >= Duration::from_millis(100) && before == total => break rail.gpu_nj(), + Some((at, _)) if at.elapsed() < Duration::from_millis(100) => {} + _ => listed = Some((Instant::now(), total)), + } + } + assert!(waited.elapsed() < Duration::from_secs(3), "the oracle's GPU rail held for 3 s while the GPU was busy"); + }; + let rail_at = Instant::now(); + let clients_from = Instant::now(); + let clients = ioreg_clients(); + Oracle { rail_nj, rail_at, clients, clients_from, clients_to: Instant::now() } +} + +/// The oracle's (self_ns, all_ns) between two readings, or `None` if a client +/// it saw at the start was gone at the end. +/// The oracle's (pid's, every client's) GPU time between two readings, over +/// the clients in the second (one new there counting its whole time), and the +/// IDs of the clients it saw at the first and not the second. `None` if a +/// client's time went back. +fn oracle_times(a: &Oracle, b: &Oracle, pid: i32) -> Option<(u64, u64, Vec)> { + let (mut me, mut all) = (0u64, 0u64); + for (id, (p, ns)) in &b.clients { + let d = ns.checked_sub(a.clients.get(id).map_or(0, |c| c.1))?; + all += d; + if *p == pid { + me += d; + } + } + let gone = a.clients.keys().filter(|id| !b.clients.contains_key(id)).copied().collect(); + Some((me, all, gone)) +} + +// ── A GPU-burning child process ───────────────────────────────────────────── + +/// This test binary again, as a child that burns the GPU for `secs` and says +/// so once its first command buffer has run. +fn spawn_burner(secs: u32) -> Child { + let mut child = Command::new(std::env::current_exe().expect("this test binary")) + .args(["--exact", "a_gpu_span_is_a_sole_users_rail_or_an_exact_zero_or_no_figure", "--nocapture", "--test-threads=1"]) + .env(CHILD_ENV, secs.to_string()) + .stdout(Stdio::piped()) + .stderr(Stdio::null()) + .spawn() + .expect("spawn the burner child"); + let mut lines = BufReader::new(child.stdout.take().expect("its stdout")).lines(); + loop { + let line = lines.next().expect("the child exited before it was ready").expect("its stdout"); + // libtest prints the test's name on the same line, before it. + if line.contains(READY) { + break; + } + } + child +} + +/// Kill and reap a child, then wait until `ioreg` no longer lists any client +/// of its pid. +fn kill_and_await_gone(mut child: Child) { + let pid = child.id() as i32; + child.kill().expect("kill the child"); + child.wait().expect("reap the child"); + let t = Instant::now(); + while ioreg_clients().values().any(|(p, _)| *p == pid) { + assert!(t.elapsed() < Duration::from_secs(5), "the dead child's GPU client is still listed"); + std::thread::sleep(Duration::from_millis(50)); + } +} + +// ── The checks ────────────────────────────────────────────────────────────── + +/// The counter's span over `work`, bracketing the oracle's own readings: the +/// counter's start snapshot, the oracle, the work, the oracle, the counter's +/// end snapshot. +fn bracketed(counter: &ProcessGpuCounter, rail: &OracleRail, work: impl FnOnce()) -> (GpuSnapshot, Oracle, Oracle, GpuSnapshot) { + let s = counter.snapshot().expect("a start snapshot"); + let o0 = oracle(rail); + work(); + let o1 = oracle(rail); + let e = counter.snapshot().expect("an end snapshot"); + (s, o0, o1, e) +} + +/// Retry a span whose GPU figure was refused because another process's client +/// vanished — some process on this busy host closing its GPU connection — +/// after checking that the refusal names a client that really is gone. Any +/// other outcome is returned; vanishing on every try fails the test. +fn until_no_client_vanished( + what: &str, + mut span: impl FnMut() -> (Result, GpuSnapshot, GpuSnapshot, X), +) -> (Result, GpuSnapshot, GpuSnapshot, X) { + for attempt in 1..=6 { + let (r, s, e, x) = span(); + match r { + Err(EnergyError::GpuClientVanished { pid, client }) => { + assert!(s.clients.iter().any(|c| c.id == client && c.pid == pid), "{what}: it named a client not at the start"); + let back = e.clients.iter().find(|c| c.id == client).map(|c| c.gpu_ns); + let before = s.clients.iter().find(|c| c.id == client).map(|c| c.gpu_ns); + assert!(back.is_none() || back < before, "{what}: it named a client still there and not reset"); + eprintln!("process-gpu: {what}: attempt {attempt}: client {client:#x} of pid {pid} vanished; again"); + } + r => return (r, s, e, x), + } + } + panic!("{what}: a client vanished in every one of 6 spans"); +} + +fn ms(ns: u64) -> f64 { + ns as f64 / 1e6 +} + +fn line(what: &str, r: &Result) { + match r { + Ok(g) => println!( + "process-gpu: {what}: {:.3} J; this process {:.1} ms of {} GPU time ({:.2}%), rail {}, {:?}", + g.uj as f64 / 1e6, + ms(g.self_ns), + g.all_ns.map_or("(not read)".to_string(), |a| format!("{:.1} ms", ms(a))), + g.share_ppm() as f64 / 1e4, + g.rail_uj.map_or("(not read)".to_string(), |r| format!("{:.3} J", r as f64 / 1e6)), + g.provenance() + ), + Err(e) => println!("process-gpu: {what}: no figure: {e}"), + } +} + +/// The counter's own and every client's GPU time over a span, against the +/// oracle's readings inside it: this process's agrees exactly (it ran no GPU +/// work in the gaps), and every client's is at least the oracle's and at most +/// that plus what the gaps can hold. Returns the oracle's (self, all). +fn agrees_with_oracle(what: &str, self_ns: u64, all_ns: u64, s: &GpuSnapshot, e: &GpuSnapshot, o0: &Oracle, o1: &Oracle) -> (u64, u64) { + let me = std::process::id() as i32; + let (o_me, o_all, gone) = oracle_times(o0, o1, me).expect("no client's time went back in the oracle's window"); + // A client the oracle saw vanish inside a span the counter took a figure + // or a share over came after the counter's start (one there at the start + // and gone at the end would have refused the span): the counter has none + // of its time, and the oracle's sum leaves it out too. + assert!(gone.iter().all(|id| !s.clients.iter().any(|c| c.id == *id)), "{what}: the oracle saw a client vanish that the counter's start held"); + // How much longer the counter's window is than the oracle's, at most: the + // counter read the clients within 20 ms before each snapshot's `at`. + let outer = e.at - s.at + Duration::from_millis(20); + let gap_ns = outer.saturating_sub(o1.clients_from.saturating_duration_since(o0.clients_to)).as_nanos() as u64; + println!( + "process-gpu: {what}: oracle: this process {:.1} ms of {:.1} ms GPU time; outer window {:.0} ms, gap {:.0} ms", + ms(o_me), + ms(o_all), + (e.at - s.at).as_secs_f64() * 1e3, + ms(gap_ns) + ); + assert_eq!(self_ns, o_me, "{what}: this process's GPU time, counter vs ioreg"); + assert!( + all_ns >= o_all && all_ns <= o_all + gap_ns + gap_ns / 10 + 5_000_000, + "{what}: every client's GPU time: counter {} ms, oracle {} ms, gap {} ms", + ms(all_ns), + ms(o_all), + ms(gap_ns) + ); + (o_me, o_all) +} + +/// A GPU-free span: an exact zero, a hardware reading, with nothing else read. +fn assert_exact_zero(what: &str, r: &Result) { + line(what, r); + let g = r.as_ref().unwrap_or_else(|e| panic!("{what}: no GPU work of this process's is an exact zero, not {e}")); + assert_eq!((g.uj, g.self_ns, g.rail_uj, g.all_ns, g.provenance()), (0, 0, None, None, Provenance::HwShunt), "{what}"); +} + +#[test] +fn a_gpu_span_is_a_sole_users_rail_or_an_exact_zero_or_no_figure() { + if let Ok(secs) = std::env::var(CHILD_ENV) { + // The child: burn, say so after the first command buffer, and keep on. + let b = Burner::new(); + b.dispatch(); + println!("{READY}"); + b.burn(Duration::from_secs(secs.parse().expect("seconds")), Duration::from_secs(60)); + return; + } + let rail = OracleRail::open(); + + // 1. Live, probed once per process with the first caller's GPU load, + // however many threads ask; a CPU load never makes it live. + let askers: Vec<_> = (0..4) + .map(|_| { + std::thread::spawn(|| { + let b = Burner::new(); + let c = ProcessGpuCounter::new(|| gpu_tick(&b)); + (c.liveness(), c.is_available()) + }) + }) + .collect(); + let seen: Vec<_> = askers.into_iter().map(|t| t.join().expect("an asker")).collect(); + let runs = jouleclaw_energy::process_gpu_probe_runs(); + let cpu_only = probe_gpu_liveness_uncached(cpu_tick); + println!("process-gpu: liveness seen by 4 threads={seen:?}, probe runs={runs}; under a CPU load only: {cpu_only:?}"); + assert_eq!(runs, 1, "the GPU liveness probe did not run exactly once"); + assert!(seen.iter().all(|s| *s == (CounterLiveness::Live, true)), "{seen:?}"); + assert_eq!(gpu_rail_liveness_cached(), Some(CounterLiveness::Live)); + assert_eq!(cpu_only, CounterLiveness::Frozen, "a CPU burn made this process's GPU sources look live"); + let burner = Burner::new(); + let counter = ProcessGpuCounter::new(|| panic!("the cached liveness probed again")); + assert!(counter.is_available()); + let span_meter = ProcessSpanMeter::new(|| panic!("the cached liveness probed again")); + assert!(span_meter.is_metered()); + + // 2. A GPU burn of ~1.5 s, against both oracles: the rail, as the sole + // user; `GpuShared`, with the shares the oracle sees, otherwise. + let (mut ran, mut buffers) = (0.0, 0u32); + let (r, s, e, (o0, o1)) = until_no_client_vanished("1.5 s burn", || { + let (s, o0, o1, e) = bracketed(&counter, &rail, || (ran, buffers) = burner.burn_counted(Duration::from_millis(1500), Duration::from_secs(20))); + (counter.share(&s, &e), s, e, (o0, o1)) + }); + line("1.5 s GPU burn", &r); + println!("process-gpu: 1.5 s GPU burn: {buffers} command buffers ran {:.1} ms on the GPU", ran * 1e3); + let (self_ns, all_ns) = match &r { + Ok(g) => (g.self_ns, g.all_ns.expect("a non-zero figure reads every client")), + Err(EnergyError::GpuShared { self_ns, all_ns }) => (*self_ns, *all_ns), + Err(other) => panic!("1.5 s burn: {other}"), + }; + agrees_with_oracle("1.5 s GPU burn", self_ns, all_ns, &s, &e, &o0, &o1); + // This process's GPU time runs a little past its command buffers' own GPU + // spans (as the sole user, 1618.3 ms for 1500.9 ms of spans over ~190 + // buffers: ~0.6 ms a buffer), so the ceiling allows 2 ms a buffer. Under + // contention the spans hold other clients' time slices too (1.50 s of + // spans for 164 ms of this process's time), so the floor is low. + assert!(self_ns >= 20_000_000, "{} ms of GPU time for a 1.5 s burn", ms(self_ns)); + assert!(self_ns as f64 <= ran * 1e9 + f64::from(buffers) * 2e6 + 5e6, "more GPU time than {buffers} command buffers ran"); + let sole = u128::from(self_ns) * 100 >= u128::from(all_ns) * 99; + match &r { + Ok(g) => { + assert!(sole, "a figure below 99% of the GPU time: {g:?}"); + // The rail: at least the oracle's inner window, at most that plus + // 120 W over the gap and a republish interval at each end. + let rail_uj = g.rail_uj.expect("a non-zero figure reads the rail"); + let o_rail_uj = (o1.rail_nj - o0.rail_nj) / 1_000; + let rail_gap_ns = (e.at - s.at + Duration::from_millis(20)).saturating_sub(o1.rail_at - o0.rail_at).as_nanos() as u64 + 100_000_000; + println!("process-gpu: 1.5 s GPU burn: rail {:.3} J, oracle's inner {:.3} J", rail_uj as f64 / 1e6, o_rail_uj as f64 / 1e6); + assert!(rail_uj + 1 >= o_rail_uj && rail_uj <= o_rail_uj + rail_gap_ns * 120 / 1_000, "rail {rail_uj} µJ vs the oracle's {o_rail_uj} µJ"); + let expected = (u128::from(rail_uj) * u128::from(g.self_ns) / u128::from(all_ns)) as u64; + assert!(g.uj.abs_diff(expected) <= 1 && g.uj > 0, "the figure {} µJ is not the rail times the share, {expected} µJ", g.uj); + assert_eq!(g.domain(), EnergyDomain::ProcessGpu); + } + Err(_) => assert!(!sole, "GpuShared at or above 99% of the GPU time: {r:?}"), + } + + // 3. A span with no GPU work of this process's is an exact zero — every + // time, whatever else on the host uses the GPU or closes a client. + for k in 0..3 { + let s = counter.snapshot().expect("a snapshot"); + std::thread::sleep(Duration::from_millis(300)); + let e = counter.snapshot().expect("a snapshot"); + assert_exact_zero(&format!("300 ms idle #{k}"), &counter.share(&s, &e)); + } + + // 4. Contended: a child burns the GPU through this process's spans. + let child = spawn_burner(60); + let child_pid = child.id() as i32; + // a. No GPU work here beside it: still an exact zero, while the oracle + // sees the child's GPU time in the window. + let (s, o0, o1, e) = bracketed(&counter, &rail, || std::thread::sleep(Duration::from_millis(400))); + let (child_ns, _, _) = oracle_times(&o0, &o1, child_pid).expect("no client's time went back in the oracle's window"); + assert_exact_zero(&format!("400 ms idle beside a child that ran {:.1} ms on the GPU", ms(child_ns)), &counter.share(&s, &e)); + assert!(child_ns >= 20_000_000, "the child did not contend: {} ms", ms(child_ns)); + // b. GPU work here beside it: `GpuShared`, at the oracle's times. + let (r, s, e, (o0, o1)) = until_no_client_vanished("contended 1 s burn", || { + let (s, o0, o1, e) = bracketed(&counter, &rail, || { + burner.burn(Duration::from_millis(1000), Duration::from_secs(20)); + }); + (counter.share(&s, &e), s, e, (o0, o1)) + }); + line("1 s GPU burn beside a burning child", &r); + let Err(EnergyError::GpuShared { self_ns, all_ns }) = r else { panic!("a shared GPU gave {r:?}") }; + agrees_with_oracle("1 s GPU burn beside a burning child", self_ns, all_ns, &s, &e, &o0, &o1); + let (child_ns, _, _) = oracle_times(&o0, &o1, child_pid).expect("no client's time went back in the oracle's window"); + println!("process-gpu: oracle: the child ran {:.1} ms on the GPU in that window", ms(child_ns)); + assert!(child_ns >= 20_000_000 && self_ns > 0, "the child did not contend: {} ms", ms(child_ns)); + assert!(u128::from(self_ns) * 100 < u128::from(all_ns) * 99); + // c. The span meter there: the CPU alone, `Partial`, saying why; and a + // CPU-only span beside the child keeps its whole figure. + let shared = (0..6) + .map(|_| span_meter.measure(SpanUses::CpuGpu, || burner.burn(Duration::from_millis(500), Duration::from_secs(10))).1.expect("a figure")) + .find(|r| !matches!(&r.source, SpanSource::CpuOnly { why, .. } if why.contains("gone or reset"))) + .expect("a client vanished in every one of 6 spans"); + println!("process-gpu: span meter, GPU burn beside the child: {shared:?}"); + match &shared.source { + SpanSource::CpuOnly { cpu_uj, why } => { + assert!(why.contains("GPU was shared") && why.contains("ns of"), "{why}"); + assert_eq!((shared.uj, shared.provenance, shared.coverage), (*cpu_uj, Provenance::HwShunt, Coverage::Partial)); + assert!(*cpu_uj > 0); + } + other => panic!("a span on a shared GPU was not the CPU alone: {other:?}"), + } + let (_, cpu_span) = span_meter.measure(SpanUses::CpuGpu, cpu_tick); + let cpu_span = cpu_span.expect("a figure"); + println!("process-gpu: span meter, CPU-only work beside the child: {cpu_span:?}"); + let (cpu_uj, gpu) = cpu_span.parts().expect("a CPU-only span is metered whole"); + assert_eq!((gpu.uj, gpu.self_ns, cpu_span.provenance, cpu_span.coverage), (0, 0, Provenance::HwShunt, Coverage::Whole)); + assert!(cpu_uj > 0 && cpu_span.uj == cpu_uj); + + // 5. The child exits inside a span. With no GPU work here: an exact zero. + // With GPU work here: no GPU figure (`GpuClientVanished`), and the span + // meter gives the CPU alone, `Partial`, saying why. + let second = spawn_burner(60); + let second_pid = second.id() as i32; + let s = counter.snapshot().expect("a snapshot"); + assert!(s.clients.iter().any(|c| c.pid == child_pid), "the child's client is not in the start snapshot"); + kill_and_await_gone(child); + let e = counter.snapshot().expect("a snapshot"); + assert!(!e.clients.iter().any(|c| c.pid == child_pid), "the dead child's client is in the end snapshot"); + assert_exact_zero("idle across the child's exit", &counter.share(&s, &e)); + let span = span_meter.start(); + let s = counter.snapshot().expect("a snapshot"); + assert!(s.clients.iter().any(|c| c.pid == second_pid), "the second child's client is not in the start snapshot"); + burner.burn(Duration::from_millis(200), Duration::from_secs(10)); + kill_and_await_gone(second); + let e = counter.snapshot().expect("a snapshot"); + let refused = counter.share(&s, &e); + line("GPU burn across the second child's exit", &refused); + assert!(matches!(refused, Err(EnergyError::GpuClientVanished { pid, .. }) if pid == second_pid), "{refused:?}"); + let across = span_meter.finish(span, SpanUses::CpuGpu).expect("a figure"); + println!("process-gpu: span meter across it: {across:?}"); + match &across.source { + SpanSource::CpuOnly { why, .. } => assert!(why.contains("process GPU") && why.contains("gone or reset"), "{why}"), + other => panic!("a span across a vanished client was not the CPU alone: {other:?}"), + } + assert_eq!((across.provenance, across.coverage), (Provenance::HwShunt, Coverage::Partial)); + + // 6. The span meter on a GPU burn with the children gone: CPU + GPU, + // `Whole` (`Partial` when the caller says the work used the ANE), as the + // sole user; the CPU alone, `Partial`, if the host's other GPU work + // shared the GPU. + for uses in [SpanUses::CpuGpu, SpanUses::CpuGpuAne] { + let r = (0..6) + .map(|_| span_meter.measure(uses, || burner.burn(Duration::from_millis(500), Duration::from_secs(10))).1.expect("a figure")) + .find(|r| !matches!(&r.source, SpanSource::CpuOnly { why, .. } if why.contains("gone or reset"))) + .expect("a client vanished in every one of 6 spans"); + match &r.source { + SpanSource::Process { cpu_uj, gpu } => { + println!( + "process-gpu: span meter, {uses:?}: {:.3} J = CPU {:.3} J + GPU {:.3} J ({:.2}% of the GPU time), {:?}, {:?}", + r.uj as f64 / 1e6, + *cpu_uj as f64 / 1e6, + gpu.uj as f64 / 1e6, + gpu.share_ppm() as f64 / 1e4, + r.provenance, + r.coverage + ); + assert_eq!(r.uj, cpu_uj + gpu.uj); + assert!(*cpu_uj > 0 && gpu.uj > 0 && gpu.share_ppm() >= 990_000, "{r:?}"); + assert_eq!(r.provenance, Provenance::HwShunt); + assert_eq!(r.coverage, if uses == SpanUses::CpuGpu { Coverage::Whole } else { Coverage::Partial }); + } + SpanSource::CpuOnly { cpu_uj, why } => { + println!("process-gpu: span meter, {uses:?}: CPU alone {:.3} J, Partial: {why}", *cpu_uj as f64 / 1e6); + assert!(why.contains("GPU was shared"), "{why}"); + assert_eq!((r.uj, r.provenance, r.coverage), (*cpu_uj, Provenance::HwShunt, Coverage::Partial)); + } + SpanSource::Fallback { .. } => panic!("the CPU counter is live, so no span falls back: {r:?}"), + } + } +} diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-harness-econ/src/lib.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-harness-econ/src/lib.rs index b79648e7b..6cb27ca4a 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-harness-econ/src/lib.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-harness-econ/src/lib.rs @@ -55,7 +55,8 @@ use jouleclaw_cascade::{ Answer, AnswerError, ExecutionTrace, L4ModelId, Query, QueryInput, Runtime, Tier, TierEstimate, TierId, VerificationStatus, }; -use jouleclaw_energy::{EnergyReading, Provenance}; +use jouleclaw_energy::span::SpanEnergy; +use jouleclaw_energy::{Coverage, EnergyError, EnergyReading, Provenance}; use jouleclaw_skill::{InMemorySkillStore, SharedSkills, SkillInducer, SkillStore}; use serde::{Deserialize, Serialize}; @@ -541,6 +542,28 @@ pub fn window_joules( } } +/// The joules `jouleclaw_energy::span::ProcessSpanMeter` put on one window of +/// in-process work — a model dispatch whose decode runs on this process's CPU and +/// its Metal GPU — at the figure's grade: `HwShunt` for this process's own CPU and +/// GPU counters, `Estimator` where the span fell back to the host meter. Or why +/// there is no figure: the meter gave none, or its figure is `Partial` — the CPU +/// alone because the GPU had no figure (it was shared, say), or work that used the +/// Neural Engine, whose energy no live source reads — which MUST NOT be billed as +/// the dispatch's total. +pub fn span_joules(span: Result) -> Result<(f64, Provenance), String> { + let span = span.map_err(|e| format!("the span meter put no figure on the window: {e}"))?; + if span.coverage != Coverage::Whole { + return Err(format!( + "the span's figure ({} µJ) covers only part of the work (coverage {:?}): it is not the dispatch's total", + span.uj, span.coverage + )); + } + match span.uj { + 0 => Err("the span meter registered no energy across the window".to_string()), + uj => Ok((uj as f64 / 1e6, span.provenance)), + } +} + /// How a report says a metered figure at `provenance` was obtained: `"measured"` for a /// reading (a hardware shunt, a live vendor energy counter, or a vendor energy model), /// `"estimated"` for an estimate. The local-model examples print each dispatch's joules @@ -589,6 +612,36 @@ mod tests { assert!(why.contains("registered no energy"), "{why}"); } + /// A span's figure is billed whole only when it covers the whole dispatch, at + /// its own grade; no figure, or a `Partial` one, is no dispatch total. + #[test] + fn a_span_is_billed_only_when_it_covers_the_whole_dispatch() { + use jouleclaw_energy::process_gpu::GpuShare; + use jouleclaw_energy::span::SpanSource; + let gpu = GpuShare { uj: 600_000, self_ns: 10, rail_uj: Some(600_000), all_ns: Some(10) }; + let span = |coverage| SpanEnergy { + uj: 1_000_000, + provenance: Provenance::HwShunt, + coverage, + source: SpanSource::Process { cpu_uj: 400_000, gpu }, + }; + assert_eq!(span_joules(Ok(span(Coverage::Whole))), Ok((1.0, Provenance::HwShunt))); + let why = span_joules(Ok(span(Coverage::Partial))).unwrap_err(); + assert!(why.contains("only part"), "{why}"); + // The CPU alone, where the GPU was shared: never a dispatch total. + let cpu_only = SpanEnergy { + uj: 400_000, + provenance: Provenance::HwShunt, + coverage: Coverage::Partial, + source: SpanSource::CpuOnly { cpu_uj: 400_000, why: "process GPU: the GPU was shared".into() }, + }; + let why = span_joules(Ok(cpu_only)).unwrap_err(); + assert!(why.contains("only part"), "{why}"); + let why = span_joules(Err(EnergyError::GpuClientVanished { pid: 7, client: 1 })).unwrap_err(); + assert!(why.contains("no figure"), "{why}"); + assert_eq!(joules_obtained(Provenance::ModelBased), "measured"); + } + #[test] fn an_estimate_is_never_reported_as_measured() { assert_eq!(joules_obtained(Provenance::Estimator), "estimated"); diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/energy_autotune.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/energy_autotune.rs index 8c0248d58..992a9a1ab 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/energy_autotune.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/energy_autotune.rs @@ -193,8 +193,30 @@ impl Metered { domain: format!("{:?}", meter.domain()), } } + + /// The record for costs measured on this process's own CPU and GPU counters + /// (`jouleclaw_energy::span::ProcessSpanMeter`), whose weakest span was graded + /// `grade` — `HwShunt`, as every span metered whole by those counters is; a + /// `Partial` span (the CPU alone, the GPU having no figure) never reaches a + /// record (see `tune_shape_live_process`). + pub fn of_process_spans(grade: jouleclaw_energy::Provenance) -> Self { + use jouleclaw_energy::Provenance as P; + let grade = match grade { + P::HwShunt => "HwShunt", + P::ModelBased => "ModelBased", + P::Estimator => "Estimator", + }; + Metered { + meter: PROCESS_SPAN_METER.to_string(), + energy_provenance: grade.to_string(), + domain: "ProcessCpu+ProcessGpu".to_string(), + } + } } +/// The meter name a cost measured on this process's CPU and GPU counters records. +pub const PROCESS_SPAN_METER: &str = "process-cpu+gpu"; + /// A learned, deterministic map from op-shape to the backend that runs it in the /// fewest joules. Deterministic (a `BTreeMap`) so a persisted or sealed profile is /// byte-stable. The dispatcher consults it; a shape with no entry falls back to @@ -602,6 +624,24 @@ mod tests { } } + /// Costs measured on this process's CPU and GPU counters are recorded under their + /// own meter name and both domains, at the weakest span's grade. + #[cfg(feature = "fabric-meter")] + #[test] + fn process_span_costs_are_recorded_with_both_domains_at_their_grade() { + use jouleclaw_energy::Provenance as P; + let m = Metered::of_process_spans(P::HwShunt); + assert_eq!( + m, + Metered { + meter: "process-cpu+gpu".into(), + energy_provenance: "HwShunt".into(), + domain: "ProcessCpu+ProcessGpu".into() + } + ); + assert_eq!(Metered::of_process_spans(P::Estimator).energy_provenance, "Estimator"); + } + /// A live IOReport counter's costs are recorded, and sealed, as hardware readings /// over the SoC — the counter's own name, grade and domain; an estimator's as /// estimates. diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/fabric_dispatch.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/fabric_dispatch.rs index 1b2e40810..408827034 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/fabric_dispatch.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-loader-gguf/src/fabric_dispatch.rs @@ -443,6 +443,10 @@ mod enabled { /// advance, or either backend's window draws no energy above idle /// ([`crate::marginal_energy`]). An unmeasured backend never wins, or loses, a /// tune-off against a measured one. + /// + /// Where the IOReport counter is unavailable (macOS 27.0.1: its CPU rails are + /// frozen for an unprivileged reader), each backend is metered instead on this + /// process's own CPU and GPU counters ([`tune_shape_live_process`]). #[cfg(feature = "fabric-meter")] pub fn tune_shape_live( a: &[f32], @@ -459,7 +463,7 @@ mod enabled { let counter = jouleclaw_energy::apple_ioreport::AppleIoReportCounter::new(); if !counter.is_available() { - return None; + return tune_shape_live_process(a, w, m, n, k, window_secs); } let read = || counter.read().ok().map(|r| r.uj); // Idle baseline, so we charge each kernel its *marginal* energy, not the @@ -501,6 +505,69 @@ mod enabled { Some(cost) } + /// [`tune_shape_live`] on this process's own CPU and GPU counters + /// (`jouleclaw_energy::span::ProcessSpanMeter`): each backend's `window_secs` + /// loop is one span, its energy this process's CPU energy plus its GPU energy + /// (an exact zero for the CPU candidate; the GPU rail for the GPU candidate, + /// where this process had the GPU to itself), per op. No idle subtraction: both + /// counters are this process's alone. Only a span metered **whole** competes: + /// one that is `Partial` — the CPU alone, because the GPU was shared or a GPU + /// client vanished mid-span — or fell back to the host meter voids the tune, + /// `None`, recording nothing, as does a meter that is not live. Ranking a CPU + /// candidate's whole figure against a GPU candidate's CPU-only one would pick + /// the GPU for energy it never counted. + #[cfg(feature = "fabric-meter")] + pub fn tune_shape_live_process( + a: &[f32], + w: &[u8], + m: usize, + n: usize, + k: usize, + window_secs: f64, + ) -> Option { + use crate::energy_autotune::{record_global_metered, tune_measured, Backend, Cost, Metered, ShapeKey}; + use crate::marginal_energy::per_op_uj; + use jouleclaw_energy::span::{weaker, ProcessSpanMeter, SpanUses}; + use jouleclaw_energy::Provenance; + use std::time::Instant; + + let meter = ProcessSpanMeter::new(jouleclaw_backend_apple::gpu_liveness_load); + if !meter.is_metered() { + return None; + } + let mut grade = Provenance::HwShunt; + let mut out = vec![0f32; m * n]; + let measure = |backend: Backend| -> Option<(u64, u64)> { + let run = |out: &mut [f32]| match backend { + Backend::Cpu => matmul_bf16_amx(a, w, &[], m, n, k, out), + Backend::Gpu => matmul_bf16_gpu(a, w, &[], m, n, k, out), + }; + if !run(&mut out) { + return Some((u64::MAX, u64::MAX)); + } + let span = meter.start(); + let t0 = Instant::now(); + let mut ops = 0u64; + while t0.elapsed().as_secs_f64() < window_secs { + run(&mut out); + ops += 1; + } + let dt = t0.elapsed().as_secs_f64(); + let e = meter.finish(span, SpanUses::CpuGpu).ok()?; + if e.coverage != jouleclaw_energy::Coverage::Whole { + return None; + } + e.parts()?; + grade = weaker(grade, e.provenance); + let ops = ops.max(1); + let uj = per_op_uj(e.uj as f64, ops).ok()?; + Some((uj, (dt / ops as f64 * 1e9) as u64)) + }; + let cost: Cost = tune_measured(&[Backend::Cpu, Backend::Gpu], 1, measure)?; + record_global_metered(ShapeKey::new(m, n, k), cost, Metered::of_process_spans(grade)); + Some(cost) + } + /// Fused fabric dispatch: when the fabric is on, the GPU /// member is certified, and EVERY job's weight is (or /// becomes) F16-resident, all jobs encode into ONE @@ -1277,7 +1344,7 @@ pub use enabled::{ /// The live energy-tuner is only present when the meter is compiled in. #[cfg(feature = "fabric-meter")] -pub use enabled::tune_shape_live; +pub use enabled::{tune_shape_live, tune_shape_live_process}; /// One op in a fused dispatch window (no-fabric stub). #[cfg(not(feature = "fabric"))] @@ -1397,9 +1464,12 @@ mod tests { /// shape only when the meter measured every backend's window, never a zero, and /// what it returns is exactly what it recorded. Where the counter is declined — /// macOS 27.0.1 on an M5 Max and an M3 Ultra, whose CPU energy rail does not - /// advance — it returns `None` and the global profile gains nothing for the shape. - /// It used to fall to `(ΔE − P_idle·Δt).max(0) = 0` µJ per op on a frozen rail and - /// record that zero, which then beat every later measurement. + /// advance — the IOReport path measures nothing (it used to fall to + /// `(ΔE − P_idle·Δt).max(0) = 0` µJ per op on a frozen rail and record that zero, + /// which then beat every later measurement); the tune is metered on this + /// process's own CPU and GPU counters instead, and recorded under their name as + /// `HwShunt` — or, where a span was not metered whole (the GPU shared with + /// another session), returns `None` and records nothing. #[cfg(feature = "fabric-meter")] #[test] fn a_live_tune_records_only_what_the_meter_measured() { @@ -1411,16 +1481,26 @@ mod tests { let w: Vec = (0..n * k) .flat_map(|i| crate::matmul_bf16::f32_to_bf16_rne(((i % 7) as f32 - 3.0) * 0.01).to_le_bytes()) .collect(); - let tuned = super::tune_shape_live(&a, &w, m, n, k, 0.05); + let tuned = super::tune_shape_live(&a, &w, m, n, k, 0.3); let key = ShapeKey::new(m, n, k); let recorded = global_snapshot().into_iter().find(|(s, _)| *s == key).map(|(_, c)| c); - eprintln!("tune_shape_live on this host: ioreport={status:?} tuned={tuned:?} recorded={recorded:?}"); + let metered = crate::energy_autotune::global_snapshot_metered() + .into_iter() + .find(|(s, _, _)| *s == key) + .and_then(|(_, _, m)| m); + eprintln!("tune_shape_live on this host: ioreport={status:?} tuned={tuned:?} recorded={recorded:?} by {metered:?}"); assert_eq!(tuned, recorded, "a live tune records exactly what it returns"); if let Some(cost) = recorded { assert!(cost.uj > 0 && cost.uj != u64::MAX, "recorded {cost:?}"); } if status != jouleclaw_energy::IoReportStatus::Live { - assert_eq!(tuned, None, "the IOReport counter is declined ({status:?}), so nothing is measured"); + if tuned.is_some() { + let metered = metered.expect("a recorded tune records its meter"); + assert_eq!(metered.meter, crate::energy_autotune::PROCESS_SPAN_METER, "{metered:?}"); + assert_eq!(metered.energy_provenance, "HwShunt", "{metered:?}"); + } else { + assert!(metered.is_none()); + } } } diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/Cargo.toml b/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/Cargo.toml index b06068156..83fcd0c61 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/Cargo.toml +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/Cargo.toml @@ -21,3 +21,8 @@ bincode = "1" # Enables the loader's device-resident decode driver (Apple Metal): # fast GGUF-quantized decode via DevstralResident + hybrid prefill. fabric = ["jouleclaw-loader-gguf/fabric"] + +[dev-dependencies] +# The `gguf_energy` example meters a GGUF decode on this process's CPU and GPU +# (`ProcessSpanMeter`, Apple Silicon IOReport GPU rail apportioned by GPU time). +jouleclaw-energy = { path = "../jouleclaw-energy", features = ["ioreport"] } diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/examples/gguf_energy.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/examples/gguf_energy.rs new file mode 100644 index 000000000..b56ef96dc --- /dev/null +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-runtime/examples/gguf_energy.rs @@ -0,0 +1,131 @@ +//! Generate from a GGUF and meter the decode on this process's CPU and GPU: +//! joules per token, CPU and GPU apart, with the GPU figure's grade and this +//! process's share of the GPU's time. +//! +//! ```text +//! Q8_PACKED=1 cargo run --release -p jouleclaw-runtime --example gguf_energy -- \ +//! /path/model.gguf "prompt" [n_tokens] +//! ``` +//! +//! `Q8_PACKED=1` (Q8_0 weights), `Q4K_PACKED=1` / `Q6K_PACKED=1` (Q4_K / Q6_K) +//! keep the weights packed and run the decode GEMVs on the GPU +//! (`jouleclaw-backend-apple`); without them the decode runs on the CPU. +//! +//! The meter is `jouleclaw_energy::span::ProcessSpanMeter`: this process's CPU +//! energy (`process-cpu`, a hardware reading) plus its GPU energy — the IOReport +//! GPU rail when this process had at least 99% of every GPU client's time over +//! the span, `HwShunt` both. Where the GPU was shared (or a GPU client of +//! another process exited during the span) the GPU has no figure, and the span +//! is the CPU's alone, `Partial`, with why. Prefill (prompt through the first +//! token) and decode (every later token) are metered as two spans. + +use jouleclaw_energy::span::{ProcessSpanMeter, SpanEnergy, SpanSource, SpanUses}; +use jouleclaw_energy::EnergyError; +use jouleclaw_loader_gguf::read_gguf_file; +use jouleclaw_loader_gguf::sample::SamplingConfig; +use jouleclaw_loader_gguf::tokenizer::Vocab; +use jouleclaw_runtime::generate::{GenerateConfig, KvCacheKind, TokenizerKind}; +use jouleclaw_runtime::streaming::Conversation; +use jouleclaw_runtime::Runtime; +use std::time::Instant; + +fn report(what: &str, tokens: usize, secs: f64, e: &Result) { + let n = tokens.max(1) as f64; + println!("{what}: {tokens} token(s) in {:.2} s ({:.1} ms/token)", secs, secs * 1e3 / n); + match e { + Ok(e) => { + println!( + " total {:.4} J/token ({:.3} J), {:?}, coverage {:?}", + e.uj as f64 / 1e6 / n, + e.uj as f64 / 1e6, + e.provenance, + e.coverage + ); + match &e.source { + SpanSource::Process { cpu_uj, gpu } => match (gpu.all_ns, gpu.rail_uj) { + (Some(all_ns), Some(rail_uj)) => println!( + " CPU {:.4} J/token (HwShunt) | GPU {:.4} J/token ({:?}): this process {:.1} ms of {:.1} ms \ + GPU time ({:.2}%) of a {:.3} J GPU rail", + *cpu_uj as f64 / 1e6 / n, + gpu.uj as f64 / 1e6 / n, + gpu.provenance(), + gpu.self_ns as f64 / 1e6, + all_ns as f64 / 1e6, + gpu.share_ppm() as f64 / 1e4, + rail_uj as f64 / 1e6 + ), + _ => println!( + " CPU {:.4} J/token (HwShunt) | GPU 0 J/token, exact: this process spent no GPU time", + *cpu_uj as f64 / 1e6 / n + ), + }, + SpanSource::CpuOnly { cpu_uj, why } => println!( + " CPU alone {:.4} J/token (HwShunt), Partial: no GPU figure, because {why}", + *cpu_uj as f64 / 1e6 / n + ), + SpanSource::Fallback { meter, why } => println!(" fallback meter `{meter}`, because {why}"), + } + } + Err(e) => println!(" no figure: {e}"), + } +} + +fn main() { + let args: Vec = std::env::args().collect(); + if args.len() < 2 { + eprintln!("usage: gguf_energy [prompt] [n_tokens]"); + std::process::exit(1); + } + let path = &args[1]; + let prompt = args.get(2).map(|s| s.as_str()).unwrap_or("The capital of France is"); + let n: usize = args.get(3).and_then(|s| s.parse().ok()).unwrap_or(64); + + let t0 = Instant::now(); + let model = read_gguf_file(path).expect("read_gguf_file"); + let vocab = Vocab::from_gguf(&model).expect("Vocab::from_gguf"); + let arch = model.metadata.get("general.architecture").and_then(|v| v.as_string()).unwrap_or("?").to_string(); + eprintln!("arch={arch} tensors={} | loaded in {:?}", model.tensors.len(), t0.elapsed()); + + let meter = ProcessSpanMeter::new(jouleclaw_backend_apple::gpu_liveness_load); + println!( + "meter: process counters {} (GPU liveness {:?}); fallback `{}`", + if meter.is_metered() { "live" } else { "NOT live" }, + meter.gpu().liveness(), + meter.fallback_name() + ); + + let mut conv = Conversation::with_runtime(&model, &vocab, (n + 64).max(128), Runtime::boot()) + .expect("Conversation::with_runtime"); + let cfg = GenerateConfig { + max_new_tokens: n, + add_bos: true, + tokenizer_kind: TokenizerKind::Auto, + cache_kind: KvCacheKind::InPlace, + sampling: SamplingConfig::greedy(), + max_seq: Some((n + 64).max(128)), + stop_strings: vec![], + }; + + let mut prefill_start = Some(meter.start()); + let t = Instant::now(); + let mut tokens: Vec = Vec::new(); + let mut prefill = None; + let mut decode = None; + for st in conv.extend(prompt, &cfg).expect("extend") { + tokens.push(st.expect("decode step").id); + if let Some(start) = prefill_start.take() { + let secs = t.elapsed().as_secs_f64(); + prefill = Some((secs, meter.finish(start, SpanUses::CpuGpu))); + decode = Some((Instant::now(), meter.start())); + } + } + let decode_figure = decode.map(|(t, s)| (t.elapsed().as_secs_f64(), meter.finish(s, SpanUses::CpuGpu))); + + println!("\n>>> {}{}\n", prompt, vocab.decode_bpe(&tokens)); + if let Some((secs, e)) = &prefill { + report("prefill (prompt + first token)", 1, *secs, e); + } + if let Some((secs, e)) = &decode_figure { + report("decode", tokens.len().saturating_sub(1), *secs, e); + } +} diff --git a/mesh/mesh-rs/crates/mesh-meter/src/lib.rs b/mesh/mesh-rs/crates/mesh-meter/src/lib.rs index 02357e2af..131c52a87 100644 --- a/mesh/mesh-rs/crates/mesh-meter/src/lib.rs +++ b/mesh/mesh-rs/crates/mesh-meter/src/lib.rs @@ -307,8 +307,9 @@ mod hw { /// or an external rail whose scope this crate cannot know. Nor does one /// process's own CPU energy (`ProcessCpu`): this bridge cannot see whether a /// request ran on that process's threads alone, and it misses any GPU work - /// and any other process serving it. The match is exhaustive, so a new domain - /// has to be decided here. + /// and any other process serving it. Nor does one process's own GPU energy + /// (`ProcessGpu`), for the same reasons and because its CPU is not in it. + /// The match is exhaustive, so a new domain has to be decided here. pub fn covers_whole_request(domain: EnergyDomain) -> bool { match domain { EnergyDomain::SocTotal | EnergyDomain::CpuPkg => true, @@ -318,7 +319,8 @@ mod hw { | EnergyDomain::Npu | EnergyDomain::Ane | EnergyDomain::ExternalRail(_) - | EnergyDomain::ProcessCpu => false, + | EnergyDomain::ProcessCpu + | EnergyDomain::ProcessGpu => false, } } @@ -637,6 +639,7 @@ mod hw { (Ane, false), (ExternalRail(0), false), (ProcessCpu, false), + (ProcessGpu, false), ] { assert_eq!(covers_whole_request(d), whole, "{d:?}"); }