From 12635bd7085ae2d8226c4d86fd198ea79d0c183f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 30 May 2026 10:09:32 +0200 Subject: [PATCH] =?UTF-8?q?feat(rl):=20Phase=204.4=20=E2=80=94=20ISV-adapt?= =?UTF-8?q?ive=20V=20blend=20controller?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replaces Phase 4.3's hard V_dq → PPO swap with an adaptive blend driven by an on-device controller. Per the project's no-tuning philosophy (pearl_controller_anchors_isv_driven, feedback_adaptive_not_tuned): V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b] where α ∈ [0, 1] is emitted by rl_v_blend_alpha_controller from the observed V_dq vs V_scalar tracking ratio: track_ratio = EMA(|V_dq − V_scalar|) / EMA(|V_scalar|) if track_ratio > 1.5 × TARGET: α ← min(α + 0.01, 1.0) if track_ratio < TARGET / 1.5: α ← max(α - 0.01, 0.0) else: hold α Plus dead-signal guard: if EMA(|V_scalar|) < 1e-4, hold α (no V signal yet to calibrate against). Bootstrap on sentinel 0: α = 1.0 (Plan A v2 behavior on first step). EMAs first-observation bootstrap (no Wiener-α blend on first sample). Two new kernels: - rl_v_blend.cu: elementwise blend (~25 LOC). Grid (ceil(B/256),1,1). - rl_v_blend_alpha_controller.cu: single-block parallel reduction + Schulman-bounded controller (~90 LOC). Grid (1,1,1), block (1024,1,1). Three new ISV slots (585/586/587): - RL_V_BLEND_ALPHA_INDEX — current α - RL_V_TRACK_ERR_EMA_INDEX — EMA(|V_dq − V_scalar|) - RL_V_SCALAR_MAG_EMA_INDEX — EMA(|V_scalar|), dead-signal floor IntegratedTrainer wiring (~80 LOC): - 2 new buffers v_blended_d, v_blended_tp1_d - In step_with_lobsim_gpu_body, after DuelingQHead Adam steps: 1. Launch controller (reads V_scalar at h_t + V_dq at h_t, emits α) 2. Launch blend kernel for s_t → v_blended_d 3. Launch blend kernel for s_tp1 → v_blended_tp1_d - compute_advantage_return now reads v_blended_d / v_blended_tp1_d instead of dueling_v_d / dueling_v_tp1_d (Phase 4.3's direct swap) Both value_head and DuelingQHead still train independently. The blend just selects which baseline drives PPO advantage per step based on observed calibration. As V_dq learns to track V_scalar, the controller gradually shifts α down toward V_dq usage. If V_dq diverges (e.g., late training entropy spikes producing volatile advantages), controller raises α back to V_scalar safety. Phase 4.3 cluster (alpha-rl-qkdm2 @ 25f5ce99b) is still running and showing dramatic late-training pnl growth (+$20M at step 18132 vs Plan A v2 peak +$9.3M). Phase 4.4 adds adaptive control on top — should reduce variance while preserving the architectural benefit. Validated: - cargo build --release clean - integrated_trainer_smoke 1 step passes - alpha_rl_train --steps 3 --b 128 under compute-sanitizer: 0 errors 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 2 + crates/ml-alpha/cuda/rl_v_blend.cu | 37 +++++ .../cuda/rl_v_blend_alpha_controller.cu | 121 ++++++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 22 ++- crates/ml-alpha/src/trainer/integrated.rs | 132 ++++++++++++++++-- 5 files changed, 303 insertions(+), 11 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_v_blend.cu create mode 100644 crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index ed94a2466..fb3c10e38 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -100,6 +100,8 @@ const KERNELS: &[&str] = &[ "rl_dueling_q_bellman_target", // Phase 4: argmax over target composed_Q + Bellman target = r + γ^n × (1-done) × max_Q "rl_dueling_q_loss_and_grad", // Phase 4: Huber loss on (target − online_composed_Q[taken]) + grad_composed "rl_dueling_q_decompose_and_bwd", // Phase 4: decompose grad_composed → grad_V + grad_A via mean-subtraction Jacobian + per-batch weight gradients + "rl_v_blend", // Phase 4.4 (2026-05-30): elementwise V_used = α V_scalar + (1−α) V_dq, α from ISV + "rl_v_blend_alpha_controller", // Phase 4.4: ISV-adaptive Schulman-bounded controller on α from observed |V_dq − V_scalar| / |V_scalar| tracking ratio "rl_ensemble_action_value", // C51+IQN ensemble: E_ensemble = α×E_C51 + (1-α)×E_IQN; α from ISV[544] "rl_noisy_linear_forward", // NoisyNet: factored noisy linear forward — y = (mu_w + sigma_w ⊙ eps_w) × x + (mu_b + sigma_b ⊙ eps_b); state-dependent exploration for C51/IQN final projection "rl_noisy_linear_backward", // NoisyNet: factored noisy linear backward — grad_mu_w/sigma_w/mu_b/sigma_b per-batch scratch for reduce_axis0 diff --git a/crates/ml-alpha/cuda/rl_v_blend.cu b/crates/ml-alpha/cuda/rl_v_blend.cu new file mode 100644 index 000000000..43bf8ab36 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_v_blend.cu @@ -0,0 +1,37 @@ +// rl_v_blend.cu — Phase 4.4 adaptive V baseline blend (2026-05-30). +// +// V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b] +// +// α read on-device from ISV[alpha_slot], emitted by +// rl_v_blend_alpha_controller (separate kernel) which adapts α +// based on observed |V_dq − V_scalar| / |V_scalar| tracking ratio. +// +// α = 1.0: pure Plan A v2 behavior (V_scalar drives PPO advantage) +// α = 0.0: pure Phase 4.3 behavior (V_dq drives PPO advantage) +// Anywhere in between: adaptive blend +// +// Per feedback_cpu_is_read_only: pure device kernel; α computed +// device-side by the controller. +// Per pearl_no_host_branches_in_captured_graph: graph-safe (reads +// ISV pointer, no host params). +// Per feedback_no_atomicadd: sole-writer per cell. +// +// Block layout: grid=(ceil(B/256), 1, 1), block=(256, 1, 1). Pure +// elementwise op. + +extern "C" __global__ void rl_v_blend( + const float* __restrict__ v_scalar, // [B] + const float* __restrict__ v_dq, // [B] + const float* __restrict__ isv, // ISV bus + int B, + int alpha_slot, + float* __restrict__ v_blended // [B] +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= B) return; + // Defensive clamp to [0, 1] — controller should keep α bounded + // (per pearl_audit_unboundedness_for_implicit_asymmetry) but a + // kernel-side guard protects against any controller bug. + const float a = fminf(1.0f, fmaxf(0.0f, isv[alpha_slot])); + v_blended[b] = a * v_scalar[b] + (1.0f - a) * v_dq[b]; +} diff --git a/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu b/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu new file mode 100644 index 000000000..a6e6e3d91 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu @@ -0,0 +1,121 @@ +// rl_v_blend_alpha_controller.cu — Phase 4.4 ISV-adaptive V blend (2026-05-30). +// +// Drives α ∈ [0, 1] for `V_used = α × V_scalar + (1−α) × V_dq` based +// on observed V_dq vs V_scalar tracking ratio: +// +// track_ratio = EMA(|V_dq − V_scalar|) / EMA(|V_scalar|) +// +// if track_ratio > 1.5 × TARGET: α ← min(α + step, 1.0) ↑ V_scalar +// if track_ratio < TARGET / 1.5: α ← max(α - step, 0.0) ↑ V_dq +// else: hold α +// +// Per pearl_wiener_alpha_floor_for_nonstationary: Schulman bounded +// discrete step, no Wiener-α blending of the controller variable +// itself (α is the controlled quantity). +// +// Per pearl_first_observation_bootstrap: bootstrap α = 1.0 on +// sentinel input (ISV[alpha_slot] == 0), EMAs use first observation +// directly. After bootstrap, α never naturally returns to exactly 0 +// because Schulman step (0.01) is unlikely to land on it; if it +// does, controller re-bootstraps harmlessly. +// +// Per pearl_blend_formulas_must_have_permanent_floor: dead-signal +// guard — if EMA(|V_scalar|) < FLOOR, hold α (no V signal to +// calibrate against; the trainer hasn't seen meaningful rewards yet). +// +// Per feedback_cpu_is_read_only: pure device kernel; reads V_scalar, +// V_dq, ISV; emits α + EMAs to ISV. No host control. +// +// Block layout: grid=(1, 1, 1), block=(BLOCK_X=1024, 1, 1). Single +// block does parallel reduction over batch up to B=1024. Thread 0 +// performs the controller update. + +#define BLOCK_X 1024 +#define EMA_ALPHA 0.01f +#define TARGET_TRACK_RATIO 0.10f +#define SCHULMAN_STEP 0.01f +#define DEAD_SIGNAL_FLOOR 1e-4f +#define BOOTSTRAP_ALPHA 1.0f + +extern "C" __global__ void rl_v_blend_alpha_controller( + const float* __restrict__ v_scalar, // [B] + const float* __restrict__ v_dq, // [B] + float* __restrict__ isv, // ISV bus + int B, + int alpha_slot, // ISV[α] + int trackerr_ema_slot, // ISV[|V_dq − V_scalar|_ema] + int v_scalar_mag_ema_slot // ISV[|V_scalar|_ema] (dead-signal floor) +) { + const int tid = threadIdx.x; + if (tid >= BLOCK_X) return; + + // ── Per-thread partial sums over strided batch ── + float track_partial = 0.0f; + float mag_partial = 0.0f; + for (int b = tid; b < B; b += BLOCK_X) { + const float vs = v_scalar[b]; + const float vd = v_dq[b]; + track_partial += fabsf(vd - vs); + mag_partial += fabsf(vs); + } + + // ── Tree-reduce in shared mem ── + __shared__ float s_t[BLOCK_X]; + __shared__ float s_m[BLOCK_X]; + s_t[tid] = track_partial; + s_m[tid] = mag_partial; + __syncthreads(); + + for (int stride = BLOCK_X / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + s_t[tid] += s_t[tid + stride]; + s_m[tid] += s_m[tid + stride]; + } + __syncthreads(); + } + + if (tid == 0) { + const float inv_B = 1.0f / (float)B; + const float track_mean = s_t[0] * inv_B; + const float mag_mean = s_m[0] * inv_B; + + // ── Bootstrap α on sentinel ── + float alpha = isv[alpha_slot]; + if (alpha == 0.0f) { + alpha = BOOTSTRAP_ALPHA; + } + + // ── First-observation bootstrap on EMAs ── + float prev_track_ema = isv[trackerr_ema_slot]; + float prev_mag_ema = isv[v_scalar_mag_ema_slot]; + const float track_ema = (prev_track_ema == 0.0f) + ? track_mean + : (1.0f - EMA_ALPHA) * prev_track_ema + EMA_ALPHA * track_mean; + const float mag_ema = (prev_mag_ema == 0.0f) + ? mag_mean + : (1.0f - EMA_ALPHA) * prev_mag_ema + EMA_ALPHA * mag_mean; + + // ── Dead-signal guard: V_scalar magnitude too small to calibrate against ── + if (mag_ema < DEAD_SIGNAL_FLOOR) { + isv[trackerr_ema_slot] = track_ema; + isv[v_scalar_mag_ema_slot] = mag_ema; + isv[alpha_slot] = alpha; // hold (write bootstrap if needed) + return; + } + + // ── Track ratio + Schulman-bounded step on α ── + const float track_ratio = track_ema / mag_ema; + if (track_ratio > 1.5f * TARGET_TRACK_RATIO) { + // V_dq diverged from V_scalar → raise α toward V_scalar + alpha = fminf(alpha + SCHULMAN_STEP, 1.0f); + } else if (track_ratio < TARGET_TRACK_RATIO / 1.5f) { + // V_dq tracks V_scalar well → lower α toward V_dq + alpha = fmaxf(alpha - SCHULMAN_STEP, 0.0f); + } + // else: hold α (within band) + + isv[alpha_slot] = alpha; + isv[trackerr_ema_slot] = track_ema; + isv[v_scalar_mag_ema_slot] = mag_ema; + } +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 48eb976fa..2efb7450b 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1111,5 +1111,25 @@ pub const RL_ACTION_ENTROPY_EMA_INDEX: usize = 583; /// Bootstrap: 0.85 (15% exploration floor). pub const RL_CONF_GATE_MAX_HOLD_FRAC_INDEX: usize = 584; +/// Phase 4.4 (2026-05-30) — adaptive V baseline blend coefficient. +/// +/// `V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b]` +/// +/// α ∈ [0, 1] adapts via Schulman-bounded step from the observed +/// V_dq vs V_scalar tracking ratio (see `rl_v_blend_alpha_controller`). +/// α = 1.0: pure Plan A v2 (V_scalar drives PPO advantage). +/// α = 0.0: pure Phase 4.3 (V_dq drives PPO advantage). +/// Bootstrap on sentinel 0 → α = 1.0 (Plan A v2 behavior on first step). +pub const RL_V_BLEND_ALPHA_INDEX: usize = 585; + +/// Phase 4.4 — EMA of `|V_dq − V_scalar|` (numerator of tracking ratio). +/// Updated on-device by `rl_v_blend_alpha_controller` with EMA α = 0.01. +/// First-observation bootstrap on sentinel 0. +pub const RL_V_TRACK_ERR_EMA_INDEX: usize = 586; + +/// Phase 4.4 — EMA of `|V_scalar|` (denominator of tracking ratio + dead-signal floor). +/// If `EMA < 1e-4` the controller holds α (no V signal to calibrate against). +pub const RL_V_SCALAR_MAG_EMA_INDEX: usize = 587; + /// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 585; +pub const RL_SLOTS_END: usize = 588; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 79671a069..2f0a1b9e9 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -190,6 +190,11 @@ const RL_ATOM_SUPPORT_UPDATE_CUBIN: &[u8] = // Feeds Q→π agreement diag and future ensemble-level selection. const RL_ENSEMBLE_ACTION_VALUE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_ensemble_action_value.cubin")); +/// Phase 4.4 adaptive V blend kernels. +const RL_V_BLEND_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_v_blend.cubin")); +const RL_V_BLEND_ALPHA_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_v_blend_alpha_controller.cubin")); const RL_Q_PI_DISTILL_GRAD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_q_pi_distill_grad.cubin")); // λ_distill adaptive controller (rljzl followup 2026-05-24). @@ -622,6 +627,14 @@ pub struct IntegratedTrainer { // C51+IQN ensemble action-value (audit 2026-05-25). _rl_ensemble_action_value_module: Arc, rl_ensemble_action_value_fn: CudaFunction, + // Phase 4.4 adaptive V blend. + _rl_v_blend_module: Arc, + rl_v_blend_fn: CudaFunction, + _rl_v_blend_alpha_controller_module: Arc, + rl_v_blend_alpha_controller_fn: CudaFunction, + /// Phase 4.4 blended V baselines fed to compute_advantage_return. + pub v_blended_d: CudaSlice, + pub v_blended_tp1_d: CudaSlice, // λ_distill adaptive controller (rljzl followup 2026-05-24). // Retained for bootstrap / testing; per-step launch fused into // rl_fused_controllers. @@ -1466,6 +1479,19 @@ impl IntegratedTrainer { let rl_ensemble_action_value_fn = rl_ensemble_action_value_module .load_function("rl_ensemble_action_value") .context("load rl_ensemble_action_value")?; + // Phase 4.4 adaptive V blend kernels. + let rl_v_blend_module = ctx + .load_cubin(RL_V_BLEND_CUBIN.to_vec()) + .context("load rl_v_blend cubin")?; + let rl_v_blend_fn = rl_v_blend_module + .load_function("rl_v_blend") + .context("load rl_v_blend")?; + let rl_v_blend_alpha_controller_module = ctx + .load_cubin(RL_V_BLEND_ALPHA_CONTROLLER_CUBIN.to_vec()) + .context("load rl_v_blend_alpha_controller cubin")?; + let rl_v_blend_alpha_controller_fn = rl_v_blend_alpha_controller_module + .load_function("rl_v_blend_alpha_controller") + .context("load rl_v_blend_alpha_controller")?; let rl_q_distill_lambda_controller_module = ctx .load_cubin(RL_Q_DISTILL_LAMBDA_CONTROLLER_CUBIN.to_vec()) .context("load rl_q_distill_lambda_controller cubin")?; @@ -1936,6 +1962,13 @@ impl IntegratedTrainer { let dueling_v_tp1_d = stream .alloc_zeros::(b_size) .context("alloc dueling_v_tp1_d")?; + // Phase 4.4 adaptive V blend output buffers. + let v_blended_d = stream + .alloc_zeros::(b_size) + .context("alloc v_blended_d")?; + let v_blended_tp1_d = stream + .alloc_zeros::(b_size) + .context("alloc v_blended_tp1_d")?; let dueling_a_d = stream .alloc_zeros::(b_size * N_ACTIONS) .context("alloc dueling_a_d")?; @@ -2485,6 +2518,12 @@ impl IntegratedTrainer { _rl_kl_reference_grad_fn: rl_kl_reference_grad_fn, _rl_ensemble_action_value_module: rl_ensemble_action_value_module, rl_ensemble_action_value_fn, + _rl_v_blend_module: rl_v_blend_module, + rl_v_blend_fn, + _rl_v_blend_alpha_controller_module: rl_v_blend_alpha_controller_module, + rl_v_blend_alpha_controller_fn, + v_blended_d, + v_blended_tp1_d, _rl_q_distill_lambda_controller_module: rl_q_distill_lambda_controller_module, rl_q_distill_lambda_controller_fn, _rl_unit_state_update_module: rl_unit_state_update_module, @@ -6419,22 +6458,95 @@ impl IntegratedTrainer { // the q_pi_agree anti-correlation where PPO with V-advantage // pushed π away from Q's preferred actions. // - // Phase 4.3 (2026-05-30): V baseline now sourced from - // DuelingQHead's dueling-trained V output (dueling_v_d / - // dueling_v_tp1_d) — replaces scalar value_head's v_pred_d. - // The dueling architecture's V is trained jointly with A via - // Bellman loss on composed Q, providing an explicitly - // V/A-decomposed baseline. value_head still trains via MSE on - // returns for diagnostic comparison (V_dq vs V_scalar tracking). - // Per spec docs/superpowers/specs/2026-05-30-phase4-independent-dueling-head-design.md §6. + // Phase 4.4 (2026-05-30): adaptive V blend. + // + // V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b] + // + // α driven on-device by rl_v_blend_alpha_controller which adapts + // from observed EMA(|V_dq − V_scalar|) / EMA(|V_scalar|) tracking + // ratio. Sentinel 0 → α = 1.0 bootstrap (Plan A v2 behavior on + // first step). Schulman bounded ±0.01 per step. Dead-signal + // guard at |V_scalar|_ema < 1e-4 holds α. + // + // Per pearl_controller_anchors_isv_driven, feedback_adaptive_not_tuned: + // no host-side scheduling; controller is fully device-resident. + // + // The blended V replaces v_pred_d/dueling_v_d in compute_advantage_return. + // value_head and DuelingQHead BOTH still train (MSE on returns + + // dueling Bellman respectively); the blend just selects which + // baseline feeds PPO advantage per step. + let alpha_slot_i = crate::rl::isv_slots::RL_V_BLEND_ALPHA_INDEX as i32; + let trackerr_slot_i = crate::rl::isv_slots::RL_V_TRACK_ERR_EMA_INDEX as i32; + let mag_slot_i = crate::rl::isv_slots::RL_V_SCALAR_MAG_EMA_INDEX as i32; + // Adaptive α controller — single block parallel reduction over batch. + // Reads V_scalar at h_t (v_pred_d) and V_dq at h_t (dueling_v_d). + { + let mut args = RawArgs::new(); + args.push_ptr(self.v_pred_d.raw_ptr()); + args.push_ptr(self.dueling_v_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(b_size_i); + args.push_i32(alpha_slot_i); + args.push_i32(trackerr_slot_i); + args.push_i32(mag_slot_i); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_v_blend_alpha_controller_fn.cu_function(), + (1, 1, 1), (1024, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_v_blend_alpha_controller: {:?}", e))?; + } + } + // Blend kernel: v_blended_d at s_t. + { + let grid_x = ((b_size as u32) + 255) / 256; + let mut args = RawArgs::new(); + args.push_ptr(self.v_pred_d.raw_ptr()); + args.push_ptr(self.dueling_v_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(b_size_i); + args.push_i32(alpha_slot_i); + args.push_ptr(self.v_blended_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_v_blend_fn.cu_function(), + (grid_x.max(1), 1, 1), (256, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_v_blend (s_t): {:?}", e))?; + } + } + // Blend kernel: v_blended_tp1_d at s_{t+1} (same α — controller emits ONCE per step). + { + let grid_x = ((b_size as u32) + 255) / 256; + let mut args = RawArgs::new(); + args.push_ptr(self.v_pred_tp1_d.raw_ptr()); + args.push_ptr(self.dueling_v_tp1_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(b_size_i); + args.push_i32(alpha_slot_i); + args.push_ptr(self.v_blended_tp1_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_v_blend_fn.cu_function(), + (grid_x.max(1), 1, 1), (256, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_v_blend (s_tp1): {:?}", e))?; + } + } { let grid_x = ((b_size as u32) + 31) / 32; let mut args = RawArgs::new(); args.push_ptr(self.isv_dev_ptr); args.push_ptr(self.rewards_d.raw_ptr()); args.push_ptr(self.dones_d.raw_ptr()); - args.push_ptr(self.dueling_v_d.raw_ptr()); - args.push_ptr(self.dueling_v_tp1_d.raw_ptr()); + args.push_ptr(self.v_blended_d.raw_ptr()); + args.push_ptr(self.v_blended_tp1_d.raw_ptr()); args.push_ptr(self.returns_d.raw_ptr()); args.push_ptr(self.advantages_d.raw_ptr()); args.push_i32(b_size_i);