diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 7a47af43c..1130cc335 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -25,7 +25,6 @@ extern "C" __global__ void c51_grad_kernel( float entropy_coeff, const float* __restrict__ branch_scales, /* [B, 4] per-sample per-branch gradient scale */ const float* __restrict__ per_sample_support, /* [B, 4, 3] stride-12 per-sample per-branch [v_min, v_max, delta_z] (Phase 2d) */ - const float* __restrict__ liquid_mod, /* [4] pinned device-mapped per-branch modulators */ const float* __restrict__ atom_positions, /* [4, num_atoms] adaptive positions. NULL = linear. */ const float* __restrict__ q_mean_ema_ptr, /* [1] pinned device-mapped Q-mean EMA */ /* Task 2.X adaptive magnitude fix + Task 2.Y adaptive direction fix — @@ -235,8 +234,7 @@ extern "C" __global__ void c51_grad_kernel( z_norm = 2.0f * (float)j / fmaxf((float)(num_atoms - 1), 1.0f) - 1.0f; delta_z = per_sample_support[support_base + 2]; } - float velocity_mod = liquid_mod[d]; - float spread_scale = inv_batch * delta_z * velocity_mod; + float spread_scale = inv_batch * delta_z; /* Asymmetric spread: challenger (adjacent to taken) gets pushed UP. * Creates a two-horse race instead of one-horse monopoly. diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index f7673d9c5..481ff3bb9 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -4518,86 +4518,6 @@ extern "C" __global__ void qlstm_step( } /* ================================================================== */ -/* Kernel: liquid_tau_rk4_step */ -/* ================================================================== */ - -/** - * Liquid tau RK4/Euler adaptive ODE step — per-branch dynamics. - * - * For each branch d (0..3): - * velocity = q_gap[d] - q_gap_ema[d] - * tau = tau_min + (tau_max - tau_min) * sigmoid(velocity / delta_z) - * - * if context_norm > rk4_threshold: - * RK4 step: x_new = x + (k1 + 2*k2 + 2*k3 + k4) / 6 - * else: - * Euler step: x_new = x + f(x) - * - * where f(x) = (1/tau) * (1.0 - x) - * liquid_mod[d] = clamp(x_new, 0.1, 2.0) - * - * Also updates q_gap_ema with adaptive alpha. - * - * One thread, one block. Grid: (1,1,1), Block: (1,1,1). - */ -extern "C" __global__ void liquid_tau_rk4_step( - const float* __restrict__ branch_q_gaps, /* [4] current per-branch Q-gaps */ - float* __restrict__ q_gap_ema, /* [4] persistent EMA (read-modify-write) */ - float* __restrict__ liquid_mod, /* [4] output modulation scalars */ - const float* __restrict__ context, /* [8] from qlstm_step (or zeros if not ready) */ - float delta_z_approx, /* eval_q_std_ema */ - float rk4_threshold /* 0.5 default */ -) { - if (threadIdx.x != 0 || blockIdx.x != 0) return; - - float tau_min = 0.01f; - float tau_max = 1.0f; - - /* Compute context norm to decide Euler vs RK4 */ - float ctx_norm = 0.0f; - for (int i = 0; i < 8; i++) ctx_norm += context[i] * context[i]; - ctx_norm = sqrtf(ctx_norm); - int use_rk4 = (ctx_norm > rk4_threshold) ? 1 : 0; - - for (int d = 0; d < 4; d++) { - float q_gap_d = branch_q_gaps[d]; - - /* Adaptive EMA update */ - float ema = q_gap_ema[d]; - if (ema < 1e-12f) { - ema = q_gap_d; - } else { - float err = fabsf(q_gap_d - ema); - float alpha = fminf(fmaxf(err / (err + fmaxf(ema, 0.001f)), 0.01f), 0.3f); - ema = (1.0f - alpha) * ema + alpha * q_gap_d; - } - q_gap_ema[d] = ema; - - float velocity = q_gap_d - ema; - float sig_arg = velocity / fmaxf(delta_z_approx, 0.01f); - float sig = 1.0f / (1.0f + expf(-sig_arg)); - float tau = tau_min + (tau_max - tau_min) * sig; - float inv_tau = 1.0f / tau; - - float x = liquid_mod[d]; - float x_new; - - if (use_rk4) { - /* RK4 step: f(x) = inv_tau * (1.0 - x) */ - float k1 = inv_tau * (1.0f - x); - float k2 = inv_tau * (1.0f - (x + 0.5f * k1)); - float k3 = inv_tau * (1.0f - (x + 0.5f * k2)); - float k4 = inv_tau * (1.0f - (x + k3)); - x_new = x + (k1 + 2.0f * k2 + 2.0f * k3 + k4) / 6.0f; - } else { - /* Euler step */ - x_new = x + inv_tau * (1.0f - x); - } - - liquid_mod[d] = fminf(fmaxf(x_new, 0.1f), 2.0f); - } -} - /* ================================================================== */ /* Denoise cuBLAS helper kernels */ /* ================================================================== */ diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index f1686b4f0..290b37cf0 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1230,8 +1230,8 @@ pub struct GpuDqnTrainer { sel_t_pinned: *mut i32, sel_t_dev_ptr: u64, - // ── Liquid tau ── - /// [4] per-branch Q-gap EMA — device buffer (read-modify-write by liquid_tau_rk4_step). + // ── Per-branch Q-gap EMA (trajectory backtracking state) ── + /// [4] per-branch Q-gap EMA — device buffer (read-modify-write for backtracking restore). per_branch_q_gap_ema_buf: CudaSlice, /// Last computed per-branch Q-gaps from reduce_current_q_stats. last_per_branch_q_gaps: [f32; 4], @@ -1242,11 +1242,6 @@ pub struct GpuDqnTrainer { q_mag_means_cached: [f32; 3], /// EMA of atom utilization [0,1] for adaptive entropy regularization. utilization_ema: f32, - /// [4] liquid modulation scalars — device buffer (written by liquid_tau_rk4_step). - liquid_mod_buf: CudaSlice, - /// Kernel handle for liquid_tau_rk4_step. - liquid_tau_rk4_kernel: CudaFunction, - // ── VSN + GLU scratch ── vsn_masked_buf: CudaSlice, // [B, SH2] vsn_kernel: CudaFunction, @@ -3218,33 +3213,6 @@ impl GpuDqnTrainer { /// Update per-branch liquid tau modulation — GPU kernel (RK4/Euler adaptive ODE). /// /// Reads per_branch_q_gaps (pinned device-mapped) and qlstm_context from device. - /// Updates per_branch_q_gap_ema_buf and liquid_mod_buf on device. - /// Zero CPU compute. - pub fn update_liquid_tau(&mut self, delta_z_approx: f32) -> Result<(), MLError> { - let gaps_ptr = self.per_branch_q_gaps_dev_ptr; - let ema_ptr = self.per_branch_q_gap_ema_buf.raw_ptr(); - let mod_ptr = self.liquid_mod_buf.raw_ptr(); - let ctx_ptr = self.qlstm_context.raw_ptr(); - let rk4_thresh = 0.5_f32; - unsafe { - self.stream - .launch_builder(&self.liquid_tau_rk4_kernel) - .arg(&gaps_ptr) - .arg(&ema_ptr) - .arg(&mod_ptr) - .arg(&ctx_ptr) - .arg(&delta_z_approx) - .arg(&rk4_thresh) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }) - .map_err(|e| MLError::ModelError(format!("liquid_tau_rk4: {e}")))?; - } - Ok(()) - } - /// Update eval v_range from observed per-branch Q-value statistics. /// Called at epoch boundary after `reduce_current_q_stats_per_branch`. /// @@ -3354,13 +3322,6 @@ impl GpuDqnTrainer { } } - // GPU kernel: liquid tau RK4/Euler adaptive ODE (zero CPU compute). - // Uses the direction-branch (b0) Q-std as the global delta_z proxy — - // matches the single-EMA path before per-branch promotion. - let delta_z = self.eval_q_std_ema[0].max(0.01); - if let Err(e) = self.update_liquid_tau(delta_z) { - tracing::warn!("liquid_tau_rk4 kernel failed: {e}"); - } let _ = per_branch_q_gaps; // already on device via per_branch_q_gaps_dev_ptr } pub fn set_per_sample_support_ptr(&mut self, ptr: u64) { @@ -7532,21 +7493,9 @@ impl GpuDqnTrainer { dev_ptr }; - // ── Liquid tau: device buffers (EMA + modulation scalars) ──────── + // ── Per-branch Q-gap EMA (trajectory backtracking state) ──────── let per_branch_q_gap_ema_buf = stream.alloc_zeros::(4) // [4] .map_err(|e| MLError::ModelError(format!("alloc per_branch_q_gap_ema_buf: {e}")))?; - // Initialize liquid_mod_buf to 1.0 (identity modulation). - let liquid_mod_buf = stream.alloc_zeros::(4) // [4] - .map_err(|e| MLError::ModelError(format!("alloc liquid_mod_buf: {e}")))?; - unsafe { - let liquid_mod_init = [1.0_f32, 1.0_f32, 1.0_f32, 1.0_f32]; - cudarc::driver::sys::cuMemcpyHtoDAsync_v2( - liquid_mod_buf.raw_ptr(), - liquid_mod_init.as_ptr().cast(), - 4 * std::mem::size_of::(), - stream.cu_stream(), - ); - } // ── VSN + GLU scratch buffers ───────────────────────────────── let vsn_masked_buf = alloc_f32(&stream, b * config.shared_h2, "vsn_masked_buf")?; @@ -8683,11 +8632,6 @@ impl GpuDqnTrainer { dp }; info!("GpuDqnTrainer: qlstm_step + qlstm_train_step kernels loaded (528 params, xLSTM mLSTM cell, scale={:.4})", qlstm_xavier_scale); - let cpbi_module_liquid = stream.context().load_cubin(EXPECTED_Q_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("cpbi cubin (liquid_tau): {e}")))?; - let liquid_tau_rk4_kernel = cpbi_module_liquid.load_function("liquid_tau_rk4_step") - .map_err(|e| MLError::ModelError(format!("liquid_tau_rk4_step load: {e}")))?; - info!("GpuDqnTrainer: liquid_tau_rk4_step kernel loaded (RK4/Euler adaptive ODE, 4 branches)"); // ── Mamba2 temporal scan: rolling history + selective SSM ────────── // W_A and W_B project from h_history width (SH2+OFI_EMBED_DIM) to STATE_D. @@ -9113,8 +9057,6 @@ impl GpuDqnTrainer { last_per_branch_q_gaps: [0.0; 4], q_mag_means_cached: [0.0_f32; 3], utilization_ema: 1.0, - liquid_mod_buf, - liquid_tau_rk4_kernel, vsn_masked_buf, vsn_kernel, glu_gate_pre_buf, @@ -11943,7 +11885,7 @@ impl GpuDqnTrainer { } // Write per-branch Q-gaps to pinned device-mapped buffer — no HtoD copy. - // GPU reads via dev_ptr (qlstm_step + liquid_tau_rk4_step kernels). + // GPU reads via dev_ptr (qlstm_step kernel). unsafe { std::ptr::copy_nonoverlapping( self.last_per_branch_q_gaps.as_ptr(), @@ -13977,7 +13919,6 @@ impl GpuDqnTrainer { let actions_buf_ptr = self.actions_buf.raw_ptr(); let d_value_logits_buf_ptr = self.d_value_logits_buf.raw_ptr(); let d_adv_logits_buf_ptr = self.d_adv_logits_buf.raw_ptr(); - let liquid_mod_buf_ptr = self.liquid_mod_buf.raw_ptr(); let atom_positions_buf_ptr = self.atom_positions_buf.raw_ptr(); // Backward-standardization fix (2026-04-23): thread the magnitude- @@ -14009,9 +13950,8 @@ impl GpuDqnTrainer { .arg(&entropy_coeff) .arg(&self.branch_scales_ptr) .arg(&self.per_sample_support_ptr) - .arg(&liquid_mod_buf_ptr) .arg(&atom_positions_buf_ptr) - .arg(&self.q_mean_ema_dev_ptr) // unused but preserves 19-param graph node structure on Hopper + .arg(&self.q_mean_ema_dev_ptr) // unused; preserves param-count graph node structure on Hopper // Task 2.X adaptive magnitude fix — ISV signal bus; kernel reads // slots [13] (q_mag_spread_ema) and [14] (q_dir_spread_ema) to // derive the magnitude-branch adaptive bin weight. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 45a16b11b..98169a77d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -157,7 +157,7 @@ | Module / kernel | Consumer path | Classification | Notes | Action | |---|---|---|---|---| | `mamba/mod.rs` + `mamba/trainable_adapter.rs` | `trainers/mamba2.rs`, `model_factory.rs`, `inference_validator.rs`; Mamba2 temporal kernel wired fully in `gpu_dqn_trainer.rs` | Wired (DQN trunk) + OUT-of-DQN-scope (supervised `trainers/mamba2.rs`) | Forward and backward both implemented in mamba2_temporal_kernel.cu | — | -| `liquid/mod.rs` + `liquid/adapter.rs` | `trainers/dqn/fused_training.rs` (features), `gpu_dqn_trainer.rs` (`liquid_tau_rk4_kernel`); also `trainers/liquid.rs` supervised path | Partial | Liquid tau ODE modulation wired in DQN kernel; `LiquidTrainableAdapter` from supervised path not called in DQN backward | Plan 2 D.7 audits full backward wiring | +| `liquid/mod.rs` + `liquid/adapter.rs` | `trainers/liquid.rs` supervised path only | Deleted (DQN) | D.7 audit (2026-04-24): `liquid_tau_rk4_step` kernel was a mathematical identity — ODE f(x)=(1/tau)*(1-x) has fixed point x=1.0; initialised to 1.0, it never deviated. `liquid_mod_buf` always held [1.0,1.0,1.0,1.0], so `velocity_mod` in c51_grad_kernel multiplied spread_scale by 1.0 unconditionally. No ISV slot existed (violates §4.C.6). Kernel, buf, and call removed. `LiquidTrainableAdapter` (supervised path) unchanged. | | `tft/mod.rs` + `tft/trainable_adapter.rs` | `trainers/tft/trainer.rs` has full forward + `backward_step()`; `TrainableTFT::backward()` returns error with explicit message that backward is via `TFTTrainer::train()` | Partial | TFT backward available via dedicated trainer; `UnifiedTrainable::backward()` slot returns informational error (not a stub — documented routing) | Plan 4 E.1/E.3 — DQN concept adoption; full backward path tracked in task #76 | ## Data Pipeline Modules diff --git a/docs/ml-supervised-to-dqn-concept-audit.md b/docs/ml-supervised-to-dqn-concept-audit.md index 7d4f5bd7f..64fb63f37 100644 --- a/docs/ml-supervised-to-dqn-concept-audit.md +++ b/docs/ml-supervised-to-dqn-concept-audit.md @@ -10,7 +10,7 @@ Authoritative decisions on which ml-supervised concepts the DQN adopts. Per Inva | TLOB attention encoder | ml-supervised::tlob | IN | D.8 (Plan 2): wire to DQN trunk | pending | | Mamba2 forward | ml-supervised::mamba2 | WIRED | already in DQN trunk | — | | Mamba2 backward | ml-supervised::mamba2 | IN | D.1 (Plan 2): complete | pending | -| Liquid Time-constant | ml-supervised::liquid | AUDIT-THEN-WIRE-OR-DELETE | D.7 (Plan 2): audit decides | pending | +| Liquid Time-constant | ml-supervised::liquid | OUT (deleted from DQN) | D.7 (2026-04-24): ODE f(x)=(1/tau)*(1-x) is identity at fixed point x=1.0; initialised to 1.0 and never moved. c51_grad spread_scale multiplied by 1.0 unconditionally. No ISV slot (§4.C.6 violation). Kernel + buf deleted. Supervised LiquidTrainableAdapter unchanged. | complete | | xLSTM | ml-supervised::xlstm | OUT | Redundant with Mamba2 + TLOB; no DQN wiring | — | | KAN | ml-supervised::kan | OUT | Function-approx not a bottleneck; no DQN wiring | — | | Encoder-decoder separation | ml-supervised::tft | IN | E.4: extend D.3 horizon-decomposed V to full trunk/head separation | pending |