feat(rl): Phase 4.4 — ISV-adaptive V blend controller

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 <noreply@anthropic.com>

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

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];
+}

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

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;

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<CudaModule>,
     rl_ensemble_action_value_fn: CudaFunction,
+    // Phase 4.4 adaptive V blend.
+    _rl_v_blend_module: Arc<CudaModule>,
+    rl_v_blend_fn: CudaFunction,
+    _rl_v_blend_alpha_controller_module: Arc<CudaModule>,
+    rl_v_blend_alpha_controller_fn: CudaFunction,
+    /// Phase 4.4 blended V baselines fed to compute_advantage_return.
+    pub v_blended_d: CudaSlice<f32>,
+    pub v_blended_tp1_d: CudaSlice<f32>,
     // λ_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::<f32>(b_size)
             .context("alloc dueling_v_tp1_d")?;
+        // Phase 4.4 adaptive V blend output buffers.
+        let v_blended_d = stream
+            .alloc_zeros::<f32>(b_size)
+            .context("alloc v_blended_d")?;
+        let v_blended_tp1_d = stream
+            .alloc_zeros::<f32>(b_size)
+            .context("alloc v_blended_tp1_d")?;
         let dueling_a_d = stream
             .alloc_zeros::<f32>(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);