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

cleanup: remove dead graph_adam buffers + q_mean_ema_ptr parameter

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-17 02:10:30 +02:00
parent 11b1a1ca9f
commit ca50bbf837
2 changed files with 4 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
crates/ml/src/cuda_pipeline/c51_grad_kernel.cu
View File

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

17
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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<f32>,
/// Adam first moment for graph_params [60]. Reserved for future graph-level training.
graph_adam_m: CudaSlice<f32>,
/// Adam second moment for graph_params [60]. Reserved for future graph-level training.
graph_adam_v: CudaSlice<f32>,
/// 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::<f32>(60)
.map_err(|e| MLError::ModelError(format!("alloc graph_adam_m: {e}")))?;
let graph_adam_v = stream.alloc_zeros::<f32>(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),