From 185add7dc8211592f24c9f8f86e9784efb8e0861 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 24 May 2026 14:20:05 +0200 Subject: [PATCH] =?UTF-8?q?feat(rl):=20adaptive=20RATIO=20+=20EWMA=20V=5FM?= =?UTF-8?q?IN/V=5FMAX=20+=20=CE=BB=5Fdistill=20bump?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit wwcsz analysis identified atom-resolution starvation + asymmetric RATIO mismatch as the empirical ceiling on win rate (38.56% vs break-even 45.3%). Three coupled fixes shipped in one pass per "no deferrals": (a) Adaptive RATIO from observed |loss|/|win| EMAs: - apply_reward_scale tracks max(-scaled, 0) per step (slot 489) - rl_reward_clamp_controller maintains neg_max_ema (slot 490, sparse-aware like pos_max_ema) - RATIO = clamp(MIN=1.0, neg_ema/pos_ema, MAX=3.0); writes to slot 481 - Removes built-in 3:1 loss-aversion bias when reality is symmetric (wwcsz: actual avg|loss|/avg(win) = 0.83). Floor 1.0 prevents inverted asymmetry; ceiling 3.0 preserves original loss-aversion as the worst case. (b) C51 V_MAX/V_MIN: ratchet → slow EWMA (α=0.001, half-life ~700 steps): - Static ratchet wasted atom resolution on rare tails — wwcsz had V_MIN=-60, V_MAX=20 but realized rewards mostly in [-5, +5] (Δz=4 vs typical reward magnitude 1-5) - Slow EWMA lets atom span shrink toward active reward range, gaining resolution where data lives. Floors at [-1, +1] preserve original C51 baseline as the worst case. - Slow α gives Q's atom mapping time to be valid across encoder/head co-adaptation (vs aggressive EWMA which would invalidate Q's learned distribution every step) (c) Q→π distillation λ bumped 0.01 → 0.05: - wwcsz showed KL dropped 2.10 → 0.30 with λ=0.01 — Q signal landing but conservatively. Bump tests whether stronger Q pull translates to better policy → better R/done. Diag exposes neg_scaled_max + neg_scaled_max_ema so the RATIO adaptation chain is observable. apply_reward_scale shared_mem doubled from 2× to 3× block × f32 to fit the three parallel reductions (abs, pos, neg). Companion to investigation (e) — n_rollout_steps controller was suspected of misalignment (256-8192 vs trade_duration ≈ 6 steps) but turned out to be a K-loop param, not used in Bellman target. 1-step Bellman with γ-bootstrap is the actual mechanism; closed without code change. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/apply_reward_scale.cu | 49 +++++--- .../cuda/rl_reward_clamp_controller.cu | 114 +++++++++++++----- crates/ml-alpha/examples/alpha_rl_train.rs | 6 + crates/ml-alpha/src/rl/isv_slots.rs | 26 +++- crates/ml-alpha/src/trainer/integrated.rs | 11 +- 5 files changed, 152 insertions(+), 54 deletions(-) diff --git a/crates/ml-alpha/cuda/apply_reward_scale.cu b/crates/ml-alpha/cuda/apply_reward_scale.cu index 72634e680..f3983e977 100644 --- a/crates/ml-alpha/cuda/apply_reward_scale.cu +++ b/crates/ml-alpha/cuda/apply_reward_scale.cu @@ -70,13 +70,21 @@ // reach the C51 atom span properly. #define RL_POS_SCALED_REWARD_MAX_INDEX 478 +// wwcsz followup 2026-05-24: also track per-step max(-scaled, 0) — the +// loss-side tail — so the RATIO (LOSS/WIN) can become adaptive from +// observed loss/win magnitudes rather than being hardcoded 3:1. wwcsz +// data showed avg_loss=-$4.36 vs avg_win=+$5.27 (actual ratio 0.83); +// the 3:1 built-in bias was over-emphasising loss-aversion vs reality. +#define RL_NEG_SCALED_REWARD_MAX_INDEX 489 + // Single block; grid-stride for b_size > BLOCK_DIM. Caller launches // with block_dim = min(b_size, 256), grid_dim = 1, shared bytes = // block_dim * sizeof(float). -// Shared-mem layout: caller passes `shared_mem_bytes = 2 * block_dim * -// sizeof(float)`. Lower half [0..block_dim) tracks max(|scaled|) for -// the diag slot; upper half [block_dim..2*block_dim) tracks max(positive -// scaled, 0) for the adaptive clamp controller's input signal. +// Shared-mem layout: caller passes `shared_mem_bytes = 3 * block_dim * +// sizeof(float)`. Three parallel reductions: +// [0..block_dim) s_abs — max(|scaled|) for diag slot +// [block_dim..2*block_dim) s_pos — max(positive scaled, 0) +// [2*block_dim..3*block_dim) s_neg — max(-scaled, 0) for adaptive RATIO extern "C" __global__ void apply_reward_scale( float* __restrict__ isv, // ISV bus IN/OUT float* __restrict__ rewards, // [b_size] IN/OUT (scaled+clamped) @@ -85,6 +93,7 @@ extern "C" __global__ void apply_reward_scale( extern __shared__ float smem[]; float* s_abs = smem; float* s_pos = smem + blockDim.x; + float* s_neg = smem + 2 * blockDim.x; const int tid = threadIdx.x; const float scale = isv[RL_REWARD_SCALE_INDEX]; // ISV-driven clamp bounds (was hardcoded 1.0 / 3.0). Reading per @@ -93,11 +102,13 @@ extern "C" __global__ void apply_reward_scale( const float clamp_win = isv[RL_REWARD_CLAMP_WIN_INDEX]; const float clamp_loss = isv[RL_REWARD_CLAMP_LOSS_INDEX]; - // Pass 1 — scale, clamp, accumulate per-thread max |scaled| AND max - // max(positive scaled, 0). Two independent local maxes tracked in - // a single grid-stride loop so the kernel's HBM read pass is shared. + // Pass 1 — scale, clamp, accumulate per-thread max |scaled|, max + // max(positive scaled, 0), AND max(-scaled, 0). Three independent + // local maxes tracked in a single grid-stride loop so the kernel's + // HBM read pass is shared. float local_abs = 0.0f; float local_pos = 0.0f; + float local_neg = 0.0f; for (int b = tid; b < b_size; b += blockDim.x) { const float raw = rewards[b]; const float scaled = raw * scale; @@ -109,14 +120,17 @@ extern "C" __global__ void apply_reward_scale( // Track positive tail BEFORE clamping for the adaptive clamp // controller — winning-trade signal that the current static - // WIN=1.0 bound clips. Negative-tail (losses) tracked - // implicitly via max(|scaled|) at the diag slot since losses - // dominate the abs-distribution; the controller assumes the - // documented loss-aversion ratio (LOSS = ratio × WIN) so we - // don't need a separate negative-tail tracker. + // WIN=1.0 bound clips. const float p = fmaxf(scaled, 0.0f); if (p > local_pos) local_pos = p; + // Track negative tail magnitude (max(-scaled, 0)) for adaptive + // RATIO calculation per wwcsz followup. Separate from |scaled| + // because we need positive-only and negative-only EMAs to + // compute observed loss/win magnitude ratio. + const float n = fmaxf(-scaled, 0.0f); + if (n > local_neg) local_neg = n; + // Asymmetric clamp: [-clamp_loss, +clamp_win] from ISV. // Defaults preserve loss aversion (loss > win) while bounding // V regression target. @@ -126,10 +140,11 @@ extern "C" __global__ void apply_reward_scale( } s_abs[tid] = local_abs; s_pos[tid] = local_pos; + s_neg[tid] = local_neg; __syncthreads(); // Pass 2 — block tree reduce (no atomics per pearl_no_atomicadd). - // Both reductions advance together to share the __syncthreads barrier. + // Three reductions advance together to share the __syncthreads barrier. for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { if (tid < stride) { const float a_a = s_abs[tid]; @@ -138,15 +153,19 @@ extern "C" __global__ void apply_reward_scale( const float p_a = s_pos[tid]; const float p_b = s_pos[tid + stride]; s_pos[tid] = p_a > p_b ? p_a : p_b; + const float n_a = s_neg[tid]; + const float n_b = s_neg[tid + stride]; + s_neg[tid] = n_a > n_b ? n_a : n_b; } __syncthreads(); } - // Thread 0 publishes per-step maxes to ISV. Both are POINT + // Thread 0 publishes per-step maxes to ISV. All are POINT // measurements; the controller maintains the EMA over the positive - // signal in its own slot. + // and negative signals in its own slots. if (tid == 0) { isv[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX] = s_abs[0]; isv[RL_POS_SCALED_REWARD_MAX_INDEX] = s_pos[0]; + isv[RL_NEG_SCALED_REWARD_MAX_INDEX] = s_neg[0]; } } diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index db8fb139f..1a8640e39 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -68,15 +68,31 @@ // MARGIN adaptation slots (audit 2026-05-24 follow-up). #define RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX 482 #define RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX 483 -// C51 atom span ratchet slots (audit 2026-05-24 second follow-up). -// Coupling: V_MAX tracks max(V_MAX_prev, max(1.0, WIN_clamp)); -// V_MIN tracks min(V_MIN_prev, min(-1.0, -LOSS_clamp)). Ratchet -// (monotone-grow) prevents destabilising Q's learned distribution -// when WIN_clamp temporarily shrinks. +// C51 atom span slots (audit 2026-05-24 second follow-up). +// Initial design: ratchet (monotone-grow). wwcsz followup +// 2026-05-24: replaced with slow EWMA (α=0.001, half-life ~700 +// steps) to focus atom resolution on the ACTIVE reward range. +// Static ratchet wasted resolution on rare tails — with avg +// rewards in [-5, +5] but atom span at [-60, +20], Δz=4 meant Q +// couldn't differentiate "slightly winning" from "slightly losing" +// actions. Slow EWMA lets the span shrink toward the current +// active range while floors at [-1, +1] preserve the original C51 +// baseline. The slow α gives Q's atom mapping time to be valid +// across encoder/head co-adaptation. #define RL_C51_V_MAX_INDEX 484 #define RL_C51_V_MIN_INDEX 485 +// Adaptive RATIO slots (wwcsz followup) — negative tail signal that +// drives RATIO = LOSS/WIN bound asymmetry. The static RATIO=3.0 +// over-emphasised loss-aversion vs observed |loss|/|win| ≈ 0.83. +#define RL_NEG_SCALED_REWARD_MAX_INDEX 489 +#define RL_NEG_SCALED_REWARD_MAX_EMA_INDEX 490 + #define MIN_WIN 1.0f +#define MIN_RATIO 1.0f // no inverted asymmetry (loss never < win) +#define MAX_RATIO 3.0f // no worse than original loss-aversion +#define V_BOUND_FLOOR 1.0f // |V_MIN| and V_MAX both ≥ 1.0 +#define V_BOUND_EWMA_ALPHA 0.001f // slow — half-life ~700 steps #define WIENER_ALPHA_FLOOR 0.4f // MARGIN Schulman bounded-step bounds + adjust rate per @@ -178,43 +194,77 @@ extern "C" __global__ void rl_reward_clamp_controller( isv[RL_REWARD_CLAMP_MARGIN_INDEX] = margin; } - // Step 4: compute adaptive WIN bound: MARGIN × EMA, floored. - // No upper cap per the header note — the MARGIN ∈ [1, 5] × EMA - // composition is the structural bound; downstream C51 projection - // saturates beyond V_MAX without needing this kernel to enforce. - const float win_eff = fmaxf(MIN_WIN, margin * ema_new); + // ── Adaptive RATIO from observed neg/pos EMAs (wwcsz followup). ── + // + // The static RATIO=3.0 baked 3:1 loss-aversion into Q's value + // representation. wwcsz showed observed |loss|/|win| ≈ 0.83, + // making Q over-cautious. Track per-step max(-scaled, 0) via a + // sparse-aware EMA (same discipline as pos_max_ema) and compute + // RATIO = clamp(MIN=1.0, neg_ema/pos_ema, MAX=3.0). Floor 1.0 + // prevents inverted asymmetry; ceiling 3.0 preserves original + // loss-aversion as the worst case. + const float neg_max = isv[RL_NEG_SCALED_REWARD_MAX_INDEX]; + const float neg_prev = isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX]; + float neg_new = neg_prev; + if (neg_max > 0.0f) { + if (neg_prev == 0.0f) { + neg_new = neg_max; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + neg_new = (1.0f - a) * neg_prev + a * neg_max; + } + isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX] = neg_new; + } + // Adaptive RATIO. Defer if EMAs not warmed up yet — keep + // statically-seeded RATIO=3.0 until both EMAs are non-zero. + if (ema_new > 0.0f && neg_new > 0.0f) { + const float ratio_target = neg_new / ema_new; + const float ratio_clamped = fmaxf(MIN_RATIO, + fminf(MAX_RATIO, ratio_target)); + isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped; + } - // LOSS = RATIO × WIN preserves asymmetry. RATIO seeded to 3.0. + // Step 4: compute adaptive WIN/LOSS bounds: MARGIN × EMA, floored. + // No upper cap on WIN per the header note — the MARGIN ∈ [1, 5] + // × EMA composition is the structural bound. LOSS = RATIO × WIN + // with RATIO now adaptive. + const float win_eff = fmaxf(MIN_WIN, margin * ema_new); const float ratio = isv[RL_REWARD_CLAMP_RATIO_INDEX]; const float loss_eff = ratio * win_eff; isv[RL_REWARD_CLAMP_WIN_INDEX] = win_eff; isv[RL_REWARD_CLAMP_LOSS_INDEX] = loss_eff; - // ── C51 atom span ratchet (audit follow-up). ────────────────── + // ── C51 atom span — slow EWMA (was ratchet). ────────────────── // - // V_MAX/V_MIN track the reward clamp with monotone-grow semantics - // — once we've admitted a winning trade of magnitude X, the C51 - // atom mapping permanently includes X in its support. This lets - // bellman_target_projection's clip pass through wins > 1.0 to - // distinct atoms instead of saturating the top atom. - // - // Floors: V_MAX ≥ 1.0, V_MIN ≤ -1.0 preserve the original C51 - // [-1, +1] support as a hard minimum (Q's distribution can't - // become coarser than the design baseline). + // wwcsz followup: replace monotone-grow with slow EWMA so atom + // resolution focuses on the active reward range. Static ratchet + // wasted Δz on rare tails (V_MAX=20 with rewards mostly in + // [-5, +5] → Δz=2 wasted on extremes). Slow α=0.001 + // (half-life ~700 steps) gives Q's atom mapping time to be valid + // across encoder/head co-adaptation. Floors at [-1, +1] + // preserve the original C51 baseline as the worst case. const float v_max_prev = isv[RL_C51_V_MAX_INDEX]; const float v_min_prev = isv[RL_C51_V_MIN_INDEX]; - const float v_max_floor = 1.0f; - const float v_min_floor = -1.0f; - // Ratchet: target = max(floor, win); kept = max(prev, target). - const float v_max_target = fmaxf(v_max_floor, win_eff); - const float v_min_target = fminf(v_min_floor, -loss_eff); - // Bootstrap on sentinel 0 — first emit replaces directly with - // target (per pearl_first_observation_bootstrap). - const float v_max_new = (v_max_prev == 0.0f) ? v_max_target - : fmaxf(v_max_prev, v_max_target); - const float v_min_new = (v_min_prev == 0.0f) ? v_min_target - : fminf(v_min_prev, v_min_target); + const float v_max_target = fmaxf(V_BOUND_FLOOR, win_eff); + const float v_min_target = fminf(-V_BOUND_FLOOR, -loss_eff); + // Bootstrap on sentinel 0 — first emit replaces directly (per + // pearl_first_observation_bootstrap). + float v_max_new, v_min_new; + if (v_max_prev == 0.0f) { + v_max_new = v_max_target; + } else { + v_max_new = (1.0f - V_BOUND_EWMA_ALPHA) * v_max_prev + + V_BOUND_EWMA_ALPHA * v_max_target; + v_max_new = fmaxf(V_BOUND_FLOOR, v_max_new); + } + if (v_min_prev == 0.0f) { + v_min_new = v_min_target; + } else { + v_min_new = (1.0f - V_BOUND_EWMA_ALPHA) * v_min_prev + + V_BOUND_EWMA_ALPHA * v_min_target; + v_min_new = fminf(-V_BOUND_FLOOR, v_min_new); + } isv[RL_C51_V_MAX_INDEX] = v_max_new; isv[RL_C51_V_MIN_INDEX] = v_min_new; } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 22a637f7c..4352dbe29 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -69,6 +69,7 @@ use ml_alpha::rl::isv_slots::{ RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX, RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX, RL_C51_V_MAX_INDEX, RL_C51_V_MIN_INDEX, RL_Q_DISTILL_LAMBDA_INDEX, RL_Q_DISTILL_TEMPERATURE_INDEX, RL_Q_DISTILL_KL_EMA_INDEX, + RL_NEG_SCALED_REWARD_MAX_INDEX, RL_NEG_SCALED_REWARD_MAX_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PI_GRAD_NORM_EMA_INDEX, RL_PPO_CLIP_INDEX, RL_PPO_LOG_RATIO_ABS_MAX_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX, @@ -689,6 +690,11 @@ fn main() -> Result<()> { // successfully chasing Q's preferences. "q_distill_kl_ema": isv[RL_Q_DISTILL_KL_EMA_INDEX], + // Negative-tail tracking — input to adaptive RATIO. + "neg_scaled_max": + isv[RL_NEG_SCALED_REWARD_MAX_INDEX], + "neg_scaled_max_ema": + isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX], }, // audit — PPO importance-ratio clamp diagnostics. // diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 18dab67fa..c5bdea2e9 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -42,8 +42,9 @@ //! | 482-483 | 2 MARGIN-adaptation slots (clip-rate EMA + target) | rl_reward_clamp_controller (Schulman step on MARGIN) | //! | 484-485 | 2 C51 atom-span ratchet slots (V_MAX + V_MIN) | rl_reward_clamp_controller (ratchet) + rl_atom_support_update (writes atom_supports_d) | //! | 486-488 | 3 Q→π distillation slots (λ + τ + KL_ema) | rl_isv_write (seed) + rl_q_pi_distill_grad (ADDS to pi_grad_logits + writes KL_ema) | +//! | 489-490 | 2 negative-tail tracking slots (raw + EMA) for adaptive RATIO | apply_reward_scale + rl_reward_clamp_controller | //! -//! Total: 89 slots; `RL_SLOTS_END = 489`. +//! Total: 91 slots; `RL_SLOTS_END = 491`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -684,6 +685,27 @@ pub const RL_Q_DISTILL_TEMPERATURE_INDEX: usize = 487; /// (KL should drop over training). pub const RL_Q_DISTILL_KL_EMA_INDEX: usize = 488; +// ─── Adaptive RATIO from observed loss/win magnitudes (wwcsz followup) ── +// +// wwcsz showed avg_loss=-$4.36 vs avg_win=+$5.27 (actual |loss|/|win| +// = 0.83). The static RATIO=3.0 baked into the reward clamp + +// C51 atom span was over-emphasising loss-aversion vs reality. +// Make RATIO adaptive from observed negative/positive tail EMAs, +// floored at 1.0 (no inverted asymmetry) and ceiled at 3.0 +// (no worse than original loss-aversion). + +/// Per-step max(-scaled, 0) — POINT measurement published by +/// `apply_reward_scale` for the adaptive RATIO controller. +/// Overwritten every step. +pub const RL_NEG_SCALED_REWARD_MAX_INDEX: usize = 489; + +/// EMA of `RL_NEG_SCALED_REWARD_MAX_INDEX` — maintained by +/// `rl_reward_clamp_controller` via Wiener-α blend (sparse-aware: +/// only updates when neg_max > 0). The controller computes +/// RATIO = clamp(1.0, neg_ema / pos_ema, 3.0) from this and the +/// existing positive-tail EMA. +pub const RL_NEG_SCALED_REWARD_MAX_EMA_INDEX: usize = 490; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 489; +pub const RL_SLOTS_END: usize = 491; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 7c460aba6..9d33ec90f 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1256,9 +1256,10 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX, 0.05), (crate::rl::isv_slots::RL_C51_V_MAX_INDEX, 1.0), (crate::rl::isv_slots::RL_C51_V_MIN_INDEX, -1.0), - // Q→π distillation (vj5f6 followup) — λ small so PPO - // dominates; τ=1.0 for canonical softmax(E_Q). - (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_INDEX, 0.01), + // Q→π distillation (vj5f6 followup) — λ raised to 0.05 + // (wwcsz followup) so Q's now-calibrated signal more + // strongly drives π updates; τ=1.0 canonical Boltzmann. + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_INDEX, 0.05), (crate::rl::isv_slots::RL_Q_DISTILL_TEMPERATURE_INDEX, 1.0), (crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01), (crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99), @@ -1753,7 +1754,7 @@ impl IntegratedTrainer { let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (block_x, 1, 1), - shared_mem_bytes: 2 * block_x * (std::mem::size_of::() as u32), + shared_mem_bytes: 3 * block_x * (std::mem::size_of::() as u32), }; let b_size_i = b_size as i32; let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn); @@ -3134,7 +3135,7 @@ impl IntegratedTrainer { let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (block_x, 1, 1), - shared_mem_bytes: 2 * block_x * (std::mem::size_of::() as u32), + shared_mem_bytes: 3 * block_x * (std::mem::size_of::() as u32), }; let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn); launch