diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index d719697b2..8ffb6883c 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -913,7 +913,7 @@ impl GpuDqnTrainer { /// Reference to the states buffer on GPU. /// - /// Shape: `[B, STATE_DIM]` — contains the batch's states after `upload_batch_gpu()`. + /// Shape: `[B, STATE_DIM]` — contains the batch's states after batch upload. /// Used by IQL value network to read states without Candle tensor intermediaries. pub fn states_buf(&self) -> &CudaSlice { &self.states_buf @@ -3296,15 +3296,9 @@ impl GpuDqnTrainer { } } - // Single DtoH readback at the end — read mean sensitivity - let n_features = market_dim.min(14); - let mut host_sens = vec![0.0_f32; market_dim]; - self.stream.synchronize() - .map_err(|e| MLError::ModelError(format!("causal final sync: {e}")))?; - self.stream.memcpy_dtoh(&self.causal_sensitivity_buf, &mut host_sens[..market_dim]) - .map_err(|e| MLError::ModelError(format!("causal readback: {e}")))?; - let total: f32 = host_sens[..n_features].iter().sum(); - Ok(total / n_features as f32) + // Sensitivity stays on GPU — no readback needed. + // The caller discards the return value (only error handling). + Ok(0.0) } /// #32 Gradient Vaccine: project out overfitting gradient components. @@ -3339,8 +3333,9 @@ impl GpuDqnTrainer { self.stream.memset_zeros(&mut self.d_adv_logits_buf) .map_err(|e| MLError::ModelError(format!("vaccine zero d_adv: {e}")))?; - // Step 3: Upload vaccine batch + forward+backward (NON-graph path) - self.upload_batch_gpu(vaccine_batch)?; + // Step 3: Upload vaccine batch via indirect pointers + forward+backward + self.upload_batch_ptrs_6(vaccine_batch); + self.submit_indirect_upload_ops()?; // Forward pass (cuBLAS, outside graph) self.launch_cublas_forward()?; self.launch_curiosity_inference()?; @@ -4573,78 +4568,6 @@ impl GpuDqnTrainer { /// Upload batch data from GPU tensors — pure CudaSlice, zero Candle temporaries. - /// - /// GpuBatch layout from GpuReplayBuffer: - /// - states/next_states: BF16 (stored in BF16 ring buffer) - /// - rewards/dones: BF16 GpuTensor - /// - weights: CudaSlice (IS-weights kept as f32 to avoid bf16 overflow → Inf → NaN) - /// - actions: U32 (stored in U32 ring buffer) - /// - /// No Candle `to_dtype()` temporaries — eliminates Drop/event conflicts with forked stream. - fn upload_batch_gpu( - &mut self, - gpu_batch: &crate::dqn::replay_buffer_type::GpuBatch, - ) -> Result<(), MLError> { - let b = self.config.batch_size; - - let f32_size = std::mem::size_of::(); - - // States + next_states: contiguous [B, SD] in GpuBatch → padded [B, pad128(SD)] - self.launch_pad_states( - self.states_buf.raw_ptr(), - gpu_batch.states.data().raw_ptr(), - b, - )?; - self.launch_pad_states( - self.next_states_buf.raw_ptr(), - gpu_batch.next_states.data().raw_ptr(), - b, - )?; - - // Rewards, dones: f32 DtoD (no bf16 NaN risk) - dtod_copy( - self.rewards_buf.raw_ptr(), - gpu_batch.rewards.raw_ptr(), - b * f32_size, &self.stream, 2, "rewards", - )?; - dtod_copy( - self.dones_buf.raw_ptr(), - gpu_batch.dones.raw_ptr(), - b * f32_size, &self.stream, 3, "dones", - )?; - - // IS-weights: f32 DtoD (bf16 overflows to Inf for PER weights) - dtod_copy( - self.is_weights_buf.raw_ptr(), - gpu_batch.weights.raw_ptr(), - b * f32_size, &self.stream, 4, "is_weights", - )?; - - // Actions: CudaSlice → i32 buf (async DtoD, zero CPU sync) - dtod_copy( - self.actions_buf.raw_ptr(), gpu_batch.actions.raw_ptr(), - b * std::mem::size_of::(), &self.stream, 4, "actions", - )?; - - Ok(()) - } - - // ═══════════════════════════════════════════════════════════════════ - // Kernel launch methods - // ═══════════════════════════════════════════════════════════════════ - - /// Launch the cuBLAS batched forward pass (online + target networks). - /// - /// Replaces `launch_forward_loss` for the forward phase only. - /// The C51 distributional loss is still computed by the existing NVRTC kernel - /// which reads the cuBLAS output buffers (`on_v_logits_buf`, `on_b_logits_buf`, - /// `tg_v_logits_buf`, `tg_b_logits_buf`) together with the saved activations. - /// - /// This raises GPU occupancy from ~1.56% to >60% on H100 by replacing the - /// 1-warp-per-sample shared-memory-bound kernel with cuBLAS DGEMM which - /// uses tensor cores and efficiently tiles across the full SM. - /// - /// Returns `Ok(false)` if cuBLAS is not initialized (caller falls back to BF16 kernel). fn launch_cublas_forward(&self) -> Result<(), MLError> { let cublas = &self.cublas_forward; diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index e94920bcf..d8ebd9f55 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -889,12 +889,8 @@ impl FusedTrainingCtx { // ── Step 5d2: Causal Intervention (#34) — per-feature sensitivity ── // #34 Causal intervention: always active (one production path) - { - let step = self.steps_since_varmap_sync; - if let Err(e) = self.trainer.run_causal_intervention(step) { - tracing::warn!("Causal intervention failed (non-fatal): {e}"); - } - } + // Causal intervention removed — 14 cuBLAS forwards every 100 steps + // with no readback or training decision based on the result. Pure waste. // ── Step 5e: Gradient Vaccine (#32) — project out overfitting gradients ── // If a vaccine batch is provided, compute its gradient and project out