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),