gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

refactor(gpu-log): rename feature cuda-diag-log → kernel-step-trace

Browse Source 
Per user direction: feature-gating is appropriate for per-step diagnostic
logging (real perf/memory cost) but the name should be specific to the
mechanism, not a generic "diag-log" catch-all. kernel-step-trace
describes what the feature provides — per-step records written to the
ring by kernels.

Pure rename, no behavioral changes. Touched: Cargo.toml feature decl,
lib.rs cfg attribute, perception.rs (~25 cfg attributes including
not(feature) pairs), gpu_log.rs + smoothness_lambda_controller.cu doc
comments, gpu_log_ring_invariants.rs file-level cfg + ignore-attribute
text.
This commit is contained in:
jgrusewski
2026-05-21 01:44:29 +02:00
parent adc3506d3e
commit fa41b39ca4
6 changed files with 37 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
crates/ml-alpha/Cargo.toml
View File

@@ -14,7 +14,7 @@ publish = false
[features]
default = ["cuda"]
cuda = []
cuda-diag-log = []
kernel-step-trace = []
[dependencies]
ml-core = { workspace = true }

4
crates/ml-alpha/cuda/smoothness_lambda_controller.cu
View File

@@ -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;

2
crates/ml-alpha/src/gpu_log.rs
View File

@@ -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

2
crates/ml-alpha/src/lib.rs
View File

@@ -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;

60
crates/ml-alpha/src/trainer/perception.rs
View File

@@ -316,28 +316,28 @@ pub struct PerceptionTrainer {
/// Cached handle for `smoothness_lambda_controller`.
smoothness_controller_fn: CudaFunction,
_smoothness_controller_module: Arc<CudaModule>,
// ── 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<crate::gpu_log::LogRing>,
#[cfg(feature = "cuda-diag-log")]
#[cfg(feature = "kernel-step-trace")]
log_step_counter_d: CudaSlice<i32>,
/// 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<crate::pinned_mem::MappedRecordBuffer<i32>>,
#[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<CudaModule>,
#[cfg(feature = "cuda-diag-log")]
#[cfg(feature = "kernel-step-trace")]
log_drain_stop: std::sync::Arc<std::sync::atomic::AtomicBool>,
#[cfg(feature = "cuda-diag-log")]
#[cfg(feature = "kernel-step-trace")]
log_drain_handle: Option<tokio::task::JoinHandle<()>>,
// 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::<i32>(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::<i32>::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<CudaStream>, src: &CudaSlice<f32>) -> Result<Vec<f32>>
// 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;

4
crates/ml-alpha/tests/gpu_log_ring_invariants.rs
View File

@@ -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};