From 924aaa775eea2d20599b641fc7f6b981e504f963 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 14:58:58 +0200 Subject: [PATCH] =?UTF-8?q?fix:=20restore=20c51=5Fgrad=2019th=20param=20+?= =?UTF-8?q?=20HER=20(32,1,1)=20=E2=80=94=20match=20working=20SHA=20graph?= =?UTF-8?q?=5Fmega?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit graph_mega on Hopper requires EXACT node structure match with working SHA 11b1a1ca9. Two remaining differences caused the hang: - c51_grad: 18→19 params (restored unused q_mean_ema_ptr) - HER relabel: (1,1,1)→(32,1,1) (restored warp-parallel launch) Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml/src/cuda_pipeline/c51_grad_kernel.cu | 3 ++- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) 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),