From ca50bbf8371f5b7875f6e1cf1cf2d08c7b706017 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 02:10:30 +0200 Subject: [PATCH] cleanup: remove dead graph_adam buffers + q_mean_ema_ptr parameter MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit graph_adam_m/v/step were allocated but never used — graph_params are intentionally fixed at near-identity initialization. q_mean_ema_ptr removed from c51_grad_kernel (ISV replaced Q-drift penalty, no ABI to maintain for compiled cubins). Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml/src/cuda_pipeline/c51_grad_kernel.cu | 8 +------- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 17 +++-------------- 2 files changed, 4 insertions(+), 21 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 6a3f4dacd..15a5b03fa 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -26,8 +26,7 @@ extern "C" __global__ void c51_grad_kernel( const float* __restrict__ branch_scales, /* [B, 4] per-sample per-branch gradient scale */ const float* __restrict__ per_sample_support, /* [B, 3] per-sample [v_min, v_max, delta_z] */ 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 */ + const float* __restrict__ atom_positions) /* [4, num_atoms] adaptive positions. NULL = linear. */ { int tid = blockIdx.x * blockDim.x + threadIdx.x; int total_elems = batch_size * num_atoms; @@ -133,11 +132,6 @@ extern "C" __global__ void c51_grad_kernel( branch_base += batch_size * A_d * num_atoms; } - /* Q-mean drift: handled adaptively by E1 enrichment (Q-value reality check). - * Hardcoded penalty removed — E1 computes bias correction from actual eval - * performance and adapts to changing reward scales automatically. */ - (void)q_mean_ema_ptr; /* parameter kept for ABI compatibility */ - /* Single deterministic write — no atomicAdd */ d_value_logits[b * num_atoms + j] = d_val_sum; } diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index ca52ca5a7..498160de3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1226,13 +1226,10 @@ pub struct GpuDqnTrainer { // No gradient backward — the 4 edges encode structural dependencies (dir→mag, etc). // Training these would require a graph-level loss, deferred to future work. /// Edge parameters: 4 edges × 15 params (W_gate[6] + W_msg[9]) = 60. + /// Fixed at near-identity initialization — not trained. The 4 edges encode + /// structural dependencies (dir→mag, etc.) via domain knowledge; training + /// would require a graph-level loss (deferred to future work). graph_params: CudaSlice, - /// Adam first moment for graph_params [60]. Reserved for future graph-level training. - graph_adam_m: CudaSlice, - /// Adam second moment for graph_params [60]. Reserved for future graph-level training. - graph_adam_v: CudaSlice, - /// Adam step counter for graph message pass optimizer. Reserved for future graph-level training. - graph_adam_step: i32, /// branch_graph_message_pass kernel handle. graph_msg_kernel: CudaFunction, @@ -5190,10 +5187,6 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("cpbi cubin (graph_msg): {e}")))?; let graph_msg_kernel = cpbi_module_graph.load_function("branch_graph_message_pass") .map_err(|e| MLError::ModelError(format!("branch_graph_message_pass load: {e}")))?; - let graph_adam_m = stream.alloc_zeros::(60) - .map_err(|e| MLError::ModelError(format!("alloc graph_adam_m: {e}")))?; - let graph_adam_v = stream.alloc_zeros::(60) - .map_err(|e| MLError::ModelError(format!("alloc graph_adam_v: {e}")))?; // Initialize graph_params: W_gate biases near zero, W_msg near identity (0.1 on diagonal). // Layout per edge (15 floats): W_gate[6] then W_msg[9 = 3×3 row-major]. let mut graph_params_host = [0.0_f32; 60]; @@ -5810,9 +5803,6 @@ impl GpuDqnTrainer { q_mean_ema_pinned, q_mean_ema_dev_ptr, graph_params, - graph_adam_m, - graph_adam_v, - graph_adam_step: 0, graph_msg_kernel, denoise_params, denoise_adam_m, @@ -8666,7 +8656,6 @@ impl GpuDqnTrainer { .arg(&self.per_sample_support_ptr) .arg(&self.liquid_mod_buf.raw_ptr()) .arg(&self.atom_positions_buf) - .arg(&self.q_mean_ema_dev_ptr) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1),