diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 15a5b03fa..4dab245f9 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -26,7 +26,8 @@ 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__ atom_positions, /* [4, num_atoms] adaptive positions. NULL = linear. */ + const float* __restrict__ q_mean_ema_ptr) /* [1] pinned — unused but preserves 19-param graph node on Hopper */ { int tid = blockIdx.x * blockDim.x + threadIdx.x; int total_elems = batch_size * num_atoms; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 5a2d25824..ca68e4da7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -3415,7 +3415,7 @@ impl GpuDqnTrainer { let launch_cfg = cudarc::driver::LaunchConfig { grid_dim: (her_batch_size as u32, 1, 1), - block_dim: (1, 1, 1), + block_dim: (32, 1, 1), // must match working SHA — graph_mega node structure on Hopper shared_mem_bytes: 0, }; @@ -8690,6 +8690,7 @@ 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) // unused but preserves 19-param graph node structure on Hopper .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1),