diff --git a/crates/ml-alpha/Cargo.toml b/crates/ml-alpha/Cargo.toml index 53fd024de..36548ef45 100644 --- a/crates/ml-alpha/Cargo.toml +++ b/crates/ml-alpha/Cargo.toml @@ -14,7 +14,7 @@ publish = false [features] default = ["cuda"] cuda = [] -cuda-diag-log = [] +kernel-step-trace = [] [dependencies] ml-core = { workspace = true } diff --git a/crates/ml-alpha/cuda/smoothness_lambda_controller.cu b/crates/ml-alpha/cuda/smoothness_lambda_controller.cu index 541ca37ae..0a58e75a2 100644 --- a/crates/ml-alpha/cuda/smoothness_lambda_controller.cu +++ b/crates/ml-alpha/cuda/smoothness_lambda_controller.cu @@ -20,7 +20,7 @@ // no host branching; thread-id gating only. // // GPU log ring producer: emits three records per call when a non-null -// `g_log_ring` is passed (cuda-diag-log feature on the Rust side): +// `g_log_ring` is passed (kernel-step-trace feature on the Rust side): // RT_INPUT — raw_per_h[5] + jitter_in[5] (pre-EMA state) // RT_STATE — jitter_ema[5] + target[5] (post-EMA state + derived target) // RT_OUTPUT — excess[5] + lambda_out[5] (control signal + emitted λ) @@ -49,7 +49,7 @@ extern "C" __global__ void smoothness_lambda_controller( int* __restrict__ first_obs, // [1] in/out — sentinel float base_lambda, // scalar — amplitude knob float* __restrict__ lambda_out, // [5] output — λ for next step - LogRing* g_log_ring, // nullable — log ring (cuda-diag-log) + LogRing* g_log_ring, // nullable — log ring (kernel-step-trace) const int* g_step_counter // nullable — device step counter ) { const int h = threadIdx.x; diff --git a/crates/ml-alpha/src/gpu_log.rs b/crates/ml-alpha/src/gpu_log.rs index 68082fdc1..d4d88bf3e 100644 --- a/crates/ml-alpha/src/gpu_log.rs +++ b/crates/ml-alpha/src/gpu_log.rs @@ -2,7 +2,7 @@ //! //! See `docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md`. //! -//! Behaviour: when the `cuda-diag-log` feature is enabled, the trainer +//! Behaviour: when the `kernel-step-trace` feature is enabled, the trainer //! allocates a mapped-pinned ring buffer + a 1-element device step counter, //! passes their pointers into logging-capable kernels, and starts a tokio //! drain task that polls the ring at 500 ms cadence and emits structured diff --git a/crates/ml-alpha/src/lib.rs b/crates/ml-alpha/src/lib.rs index 8839602be..6e9be402e 100644 --- a/crates/ml-alpha/src/lib.rs +++ b/crates/ml-alpha/src/lib.rs @@ -32,7 +32,7 @@ pub mod heads; pub mod isv; pub mod pinned; pub mod pinned_mem; -#[cfg(feature = "cuda-diag-log")] +#[cfg(feature = "kernel-step-trace")] pub mod gpu_log; pub mod trainer; diff --git a/crates/ml-alpha/src/trainer/perception.rs b/crates/ml-alpha/src/trainer/perception.rs index 622c00f5f..4ea3ea787 100644 --- a/crates/ml-alpha/src/trainer/perception.rs +++ b/crates/ml-alpha/src/trainer/perception.rs @@ -316,28 +316,28 @@ pub struct PerceptionTrainer { /// Cached handle for `smoothness_lambda_controller`. smoothness_controller_fn: CudaFunction, _smoothness_controller_module: Arc, - // ── GPU log ring (feature: cuda-diag-log) ── + // ── GPU log ring (feature: kernel-step-trace) ── // When enabled, the trainer allocates a mapped-pinned ring + a device // step counter, passes their pointers into logging-capable kernels // (currently `smoothness_lambda_controller`), and spawns a tokio drain // task that polls the ring at 500 ms cadence and emits structured // tracing events. See `docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md`. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_ring: std::sync::Arc, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_step_counter_d: CudaSlice, /// Mapped-pinned host shadow of `log_step_counter_d`. Updated via /// memcpy_dtod_async inside the captured region; read by the drain /// task each tick. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_step_counter_host_shadow: std::sync::Arc>, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] gpu_log_tick_fn: CudaFunction, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] _gpu_log_module: Arc, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_drain_stop: std::sync::Arc, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_drain_handle: Option>, // X6: LN_b weights (formerly ln_gain_d / ln_bias_d) moved to // `self.trunk.ln_b_gain_d` / `ln_b_bias_d`. LN_b is the LayerNorm @@ -1129,24 +1129,24 @@ impl PerceptionTrainer { let smoothness_jitter_first_obs_d = stream.alloc_zeros::(1) .context("smoothness_jitter_first_obs_d alloc")?; - // ── GPU log ring allocation (feature: cuda-diag-log) ── + // ── GPU log ring allocation (feature: kernel-step-trace) ── // Per `feedback_no_htod_htoh_only_mapped_pinned`: ring is mapped- // pinned; step counter shadow is mapped-pinned; the on-stream // DtoD memcpy of the device step counter into the host shadow // is captured in the graph, so the drainer sees the latest // counter without an extra sync. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_ring = std::sync::Arc::new(crate::gpu_log::LogRing::alloc()?); - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_step_counter_d = crate::gpu_log::alloc_step_counter(&stream)?; - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_step_counter_host_shadow = { let buf = unsafe { crate::pinned_mem::MappedRecordBuffer::::new(1) } .map_err(|e| anyhow::anyhow!("log step counter shadow: {e}"))?; buf.write_record(0, 0); std::sync::Arc::new(buf) }; - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let (gpu_log_tick_fn, gpu_log_module) = { const GPU_LOG_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/gpu_log_ring.cubin")); let module = ctx.load_cubin(GPU_LOG_CUBIN.to_vec()) @@ -1154,7 +1154,7 @@ impl PerceptionTrainer { let f = module.load_function("gpu_log_tick").context("load gpu_log_tick")?; (f, module) }; - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_drain_stop = std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)); // The drain task uses `tokio::spawn`, which requires a runtime // handle on the current thread. In production the trainer runs @@ -1163,7 +1163,7 @@ impl PerceptionTrainer { // (records still land in the ring; the test simply doesn't read // them through the structured tracing path). This is the // canonical pattern for opt-in async telemetry. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_drain_handle = if tokio::runtime::Handle::try_current().is_ok() { Some(crate::gpu_log::spawn_drain_task( log_ring.clone(), @@ -1198,19 +1198,19 @@ impl PerceptionTrainer { smoothness_jitter_first_obs_d, smoothness_controller_fn, _smoothness_controller_module: smoothness_controller_module, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_ring, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_step_counter_d, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_step_counter_host_shadow, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] gpu_log_tick_fn, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] _gpu_log_module: gpu_log_module, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_drain_stop, - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] log_drain_handle, // LayerNorm state. `ln_gain_d` / `ln_bias_d` constructed // above via `upload` (gain=1, bias=0). LN_HIDDEN matches @@ -1668,13 +1668,13 @@ impl PerceptionTrainer { k_seq: usize, total_snaps: usize, ) -> Result<()> { - // ── GPU log ring tick (feature: cuda-diag-log) ── + // ── GPU log ring tick (feature: kernel-step-trace) ── // Runs FIRST inside the captured graph so every downstream // logging-capable kernel sees the new step counter. Single- // thread/single-block kernel: no host malloc, no host branch // (per `pearl_cudarc_disable_event_tracking_for_graph_capture` // and `pearl_no_host_branches_in_captured_graph`). - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] { let tick_cfg = LaunchConfig { grid_dim: (1, 1, 1), @@ -2157,16 +2157,16 @@ impl PerceptionTrainer { // GPU log ring pointers. With feature off, both are 0 (null) and // the kernel's `log_record()` short-circuits at the nullptr guard // — behaviorally identical to the pre-log kernel. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_ring_ptr: cudarc::driver::sys::CUdeviceptr = self.log_ring.dev_ptr(); - #[cfg(not(feature = "cuda-diag-log"))] + #[cfg(not(feature = "kernel-step-trace"))] let log_ring_ptr: cudarc::driver::sys::CUdeviceptr = 0; - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] let log_step_ptr: cudarc::driver::sys::CUdeviceptr = { let (p, _g) = self.log_step_counter_d.device_ptr(&self.stream); p }; - #[cfg(not(feature = "cuda-diag-log"))] + #[cfg(not(feature = "kernel-step-trace"))] let log_step_ptr: cudarc::driver::sys::CUdeviceptr = 0; { let mut launch = self.stream.launch_builder(&self.smoothness_controller_fn); @@ -2755,7 +2755,7 @@ impl PerceptionTrainer { // how many steps have completed since its last drain. Captured // inside the same graph as the other shadow memcpys, so a single // end-of-step sync makes all of them visible. - #[cfg(feature = "cuda-diag-log")] + #[cfg(feature = "kernel-step-trace")] unsafe { let (src_ptr, _g) = self.log_step_counter_d.device_ptr(&self.stream); cudarc::driver::result::memcpy_dtod_async( @@ -4179,7 +4179,7 @@ fn download(stream: &Arc, src: &CudaSlice) -> Result> // runtime / hanging shutdown. Drop cannot be async, so we cannot await // the handle — abort is the canonical pattern for tokio JoinHandle in // non-async destructors. -#[cfg(feature = "cuda-diag-log")] +#[cfg(feature = "kernel-step-trace")] impl Drop for PerceptionTrainer { fn drop(&mut self) { use std::sync::atomic::Ordering; diff --git a/crates/ml-alpha/tests/gpu_log_ring_invariants.rs b/crates/ml-alpha/tests/gpu_log_ring_invariants.rs index 3ee5afe2d..49eb195e3 100644 --- a/crates/ml-alpha/tests/gpu_log_ring_invariants.rs +++ b/crates/ml-alpha/tests/gpu_log_ring_invariants.rs @@ -10,7 +10,7 @@ //! Test #7 is the GPU oracle path: launches the tick kernel via cudarc, //! confirms the device-side step counter increments. -#![cfg(feature = "cuda-diag-log")] +#![cfg(feature = "kernel-step-trace")] use std::sync::Arc; use std::sync::atomic::{AtomicBool, Ordering}; @@ -250,7 +250,7 @@ async fn decoder_dispatch_unknown_falls_through() { /// `perception_overfit` smoke once a kernel wires `log_record()` into /// the ring (Task 7). #[test] -#[ignore = "requires CUDA — run via `cargo test --features cuda-diag-log -- --ignored`"] +#[ignore = "requires CUDA — run via `cargo test --features kernel-step-trace -- --ignored`"] fn kernel_writes_visible_to_host_smoke() { use cudarc::driver::{LaunchConfig, PushKernelArg};