diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
index 8159c1a88..a439574e4 100644
--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -1,7 +1,7 @@
 // Vectorized backtest environment step kernel.
 // One thread per walk-forward window. Each thread steps sequentially.
 //
-// Portfolio state layout per window [8 floats]:
+// Portfolio state layout per window [8 bf16]:
 //   [0] value        - current portfolio value
 //   [1] position     - current position size (-1.0 to +1.0)
 //   [2] cash         - cash balance
@@ -21,9 +21,9 @@
 // cum_return and step_count are intentionally zeroed — the breached episode's
 // metrics are captured in done_flags and step_returns before this reset.
 __device__ void handle_capital_floor_breach(
-    float* portfolio_state, int ps,
-    float new_capital, float step_ret,
-    float* step_rewards, float* step_returns, int* actions_history, int* done_flags,
+    __nv_bfloat16* portfolio_state, int ps,
+    __nv_bfloat16 new_capital, __nv_bfloat16 step_ret,
+    __nv_bfloat16* step_rewards, __nv_bfloat16* step_returns, int* actions_history, int* done_flags,
     int w, int max_len, int current_step, int b0_size
 ) {
     step_rewards[w] = step_ret;
@@ -31,13 +31,13 @@ __device__ void handle_capital_floor_breach(
     actions_history[w * max_len + current_step] = b0_size / 2;  // Flat
     // Full episode restart — metrics captured before this reset
     portfolio_state[ps + 0] = new_capital;
-    portfolio_state[ps + 1] = 0.0f;
+    portfolio_state[ps + 1] = bf16_zero();
     portfolio_state[ps + 2] = new_capital;
-    portfolio_state[ps + 3] = 0.0f;
+    portfolio_state[ps + 3] = bf16_zero();
     portfolio_state[ps + 4] = new_capital;
-    portfolio_state[ps + 5] = 0.0f;
-    portfolio_state[ps + 6] = 0.0f;
-    portfolio_state[ps + 7] = 0.0f;
+    portfolio_state[ps + 5] = bf16_zero();
+    portfolio_state[ps + 6] = bf16_zero();
+    portfolio_state[ps + 7] = bf16_zero();
     done_flags[w] = 1;
 }
 
@@ -50,15 +50,15 @@ extern "C" __global__ void backtest_env_step(
     const int*   __restrict__ actions,      // [n_windows] factored action index
 
     // Portfolio state (read-write, persistent across steps)
-    float*       portfolio_state,           // [n_windows * PORTFOLIO_STATE_SIZE]
+    __nv_bfloat16*       portfolio_state,           // [n_windows * PORTFOLIO_STATE_SIZE]
 
     // Step outputs
-    float*       step_rewards,              // [n_windows]
-    float*       step_returns,              // [n_windows * max_len] (accumulated)
+    __nv_bfloat16*       step_rewards,              // [n_windows]
+    __nv_bfloat16*       step_returns,              // [n_windows * max_len] (accumulated)
     int*         done_flags,                // [n_windows]
     int*         actions_history,           // [n_windows * max_len] (accumulated, for metrics)
 
-    // Config
+    // Config — host scalars stay float, convert on first use
     int          n_windows,
     int          max_len,
     float        max_position,
@@ -95,34 +95,38 @@ extern "C" __global__ void backtest_env_step(
         return;
     }
 
-    // Read current prices
+    // Read current prices — native bf16
     int price_base = (w * max_len + current_step) * 4;
-    float open  = (float)prices[price_base + 0];
-    float high  = (float)prices[price_base + 1];
-    float low   = (float)prices[price_base + 2];
-    float close = (float)prices[price_base + 3];
+    __nv_bfloat16 open  = prices[price_base + 0];
+    __nv_bfloat16 high  = prices[price_base + 1];
+    __nv_bfloat16 low   = prices[price_base + 2];
+    __nv_bfloat16 close = prices[price_base + 3];
 
-    // Read portfolio state from shared memory tile
-    float value      = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 0];
-    float position   = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 1];
-    float cash       = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 2];
-    float entry_price = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 3];
-    float max_equity = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 4];
-    float hold_time  = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 5];
-    float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6];
+    // Read portfolio state from shared memory tile — native bf16
+    __nv_bfloat16 value      = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 0];
+    __nv_bfloat16 position   = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 1];
+    __nv_bfloat16 cash       = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 2];
+    __nv_bfloat16 entry_price = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 3];
+    __nv_bfloat16 max_equity = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 4];
+    __nv_bfloat16 hold_time  = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 5];
+    __nv_bfloat16 cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6];
 
     // ── Capital floor circuit breaker (shared: trade_physics.cuh) ────────
-    if (check_capital_floor(value, max_equity)) {
+    // trade_physics expects float — promote for the call, then back to bf16
+    if (check_capital_floor(__bfloat162float(value), __bfloat162float(max_equity))) {
         // Force flat: close any open position at current price (notional model)
-        if (fabsf(position) > 0.001f && close > 0.0f) {
-            float exit_cost = compute_tx_cost(position, close, tx_cost_bps, spread_cost,
+        if (__bfloat162float(bf16_fabs(position)) > 0.001f && close > bf16_zero()) {
+            float exit_cost = compute_tx_cost(__bfloat162float(position), __bfloat162float(close),
+                                              tx_cost_bps, spread_cost,
                                               max_position, 0, -1.0f);
-            cash += position * close;   // sell position at market (notional)
-            cash -= exit_cost;
-            position = 0.0f;
+            cash = cash + position * close;   // sell position at market (notional)
+            cash = cash - bf16(exit_cost);
+            position = bf16_zero();
         }
-        float liq_value = cash;
-        float liq_ret = (value > 0.01f) ? (liq_value - value) / value : 0.0f;
+        __nv_bfloat16 liq_value = cash;
+        __nv_bfloat16 liq_ret = (value > bf16(0.01f))
+            ? (liq_value - value) / value
+            : bf16_zero();
         handle_capital_floor_breach(portfolio_state, ps, liq_value, liq_ret,
                                     step_rewards, step_returns, actions_history, done_flags,
                                     w, max_len, current_step, b0_size);
@@ -132,22 +136,27 @@ extern "C" __global__ void backtest_env_step(
     // ── Decode action (shared: trade_physics.cuh) ────────────────────────
     int action_val = actions[w];
     int exposure_idx = decode_exposure_index(action_val, b0_size, b1_size, b2_size);
-    float target_exposure = compute_target_position(exposure_idx, b0_size, max_position);
+    // trade_physics compute_target_position returns float
+    __nv_bfloat16 target_exposure = bf16(compute_target_position(exposure_idx, b0_size, max_position));
     int order_type_idx = decode_order_type(action_val, b1_size, b2_size);
 
     // ── Margin-aware position cap (shared: trade_physics.cuh) ────────────
-    // Prevents overleveraging when equity is depleted — a depleted account
-    // can't hold the same position as a full account.
     // margin = price * multiplier * margin_pct (e.g. 5000 * 50 * 0.06 = $15K for ES)
-    float margin_per_contract = close * contract_multiplier * margin_pct;
-    target_exposure = apply_margin_cap(target_exposure, value, margin_per_contract, max_equity);
+    __nv_bfloat16 margin_per_contract = close * bf16(contract_multiplier) * bf16(margin_pct);
+    target_exposure = bf16(apply_margin_cap(__bfloat162float(target_exposure),
+                                            __bfloat162float(value),
+                                            __bfloat162float(margin_per_contract),
+                                            __bfloat162float(max_equity)));
 
     // Suppress unused variable warnings
     (void)open; (void)high; (void)low;
 
     // ── Hold enforcement (shared: trade_physics.cuh) ─────────────────────
     int is_last_bar = (current_step >= wlen - 1) ? 1 : 0;
-    target_exposure = enforce_hold(position, target_exposure, hold_time, min_hold_bars, is_last_bar);
+    target_exposure = bf16(enforce_hold(__bfloat162float(position),
+                                        __bfloat162float(target_exposure),
+                                        __bfloat162float(hold_time),
+                                        min_hold_bars, is_last_bar));
 
     // ── Trailing stop (shared: trade_physics.cuh) ────────────────────────
     // Exit when profit retreats from peak. Uses 0.5% base distance.
@@ -156,43 +165,58 @@ extern "C" __global__ void backtest_env_step(
     {
         // Portfolio-weighted trade return: unrealized P&L / equity
         // Matches training kernel's unrealized_trade_pnl computation
-        float trade_ret = 0.0f;
-        if (fabsf(position) > 0.001f && entry_price > 0.0f && value > 1.0f) {
-            float unrealized = position * (close - entry_price);
+        __nv_bfloat16 trade_ret = bf16_zero();
+        if (__bfloat162float(bf16_fabs(position)) > 0.001f
+            && entry_price > bf16_zero()
+            && value > bf16_one()) {
+            __nv_bfloat16 unrealized = position * (close - entry_price);
             trade_ret = unrealized / value;
         }
-        if (check_trailing_stop(hold_time, min_hold_bars, max_equity,
-                                value,  /* prev_equity: state[0] = previous step's final equity, same as training's ps[9] */
-                                trade_ret, 0.005f, 1.0f, 1.0f)) {
-            target_exposure = 0.0f;  // Force flat — trailing stop triggered
+        if (check_trailing_stop(__bfloat162float(hold_time), min_hold_bars,
+                                __bfloat162float(max_equity),
+                                __bfloat162float(value),
+                                __bfloat162float(trade_ret), 0.005f, 1.0f, 1.0f)) {
+            target_exposure = bf16_zero();  // Force flat — trailing stop triggered
         }
     }
 
     // ── Execute trade (shared: trade_physics.cuh) ────────────────────────
-    float prev_position = position;
-    float tx_cost = execute_trade(&position, &cash, target_exposure, close,
-                                  tx_cost_bps, spread_cost, max_position,
-                                  order_type_idx, -1.0f);
+    __nv_bfloat16 prev_position = position;
+    {
+        float f_position = __bfloat162float(position);
+        float f_cash     = __bfloat162float(cash);
+        float tx_cost = execute_trade(&f_position, &f_cash,
+                                      __bfloat162float(target_exposure),
+                                      __bfloat162float(close),
+                                      tx_cost_bps, spread_cost, max_position,
+                                      order_type_idx, -1.0f);
+        position = bf16(f_position);
+        cash     = bf16(f_cash);
+        (void)tx_cost;
+    }
 
     // Update entry_price only on NEW trade entry or reversal (for trailing stop)
-    if (fabsf(prev_position) < 0.001f && fabsf(position) > 0.001f) {
+    if (__bfloat162float(bf16_fabs(prev_position)) < 0.001f
+        && __bfloat162float(bf16_fabs(position)) > 0.001f) {
         entry_price = close;  // new entry from flat
-    } else if (prev_position * position < 0.0f) {
+    } else if (__bfloat162float(prev_position) * __bfloat162float(position) < 0.0f) {
         entry_price = close;  // reversal
-    } else if (fabsf(position) < 0.001f) {
-        entry_price = 0.0f;   // went flat
+    } else if (__bfloat162float(bf16_fabs(position)) < 0.001f) {
+        entry_price = bf16_zero();   // went flat
     }
     // On same-direction scaling (L50->L100), keep original entry_price
 
     // ── Hold time tracking (shared: trade_physics.cuh) ───────────────────
-    hold_time = update_hold_time(prev_position, position, hold_time);
+    hold_time = bf16(update_hold_time(__bfloat162float(prev_position),
+                                      __bfloat162float(position),
+                                      __bfloat162float(hold_time)));
 
     // Mark-to-market current position (notional model: equity = cash + position * price)
-    float new_value = cash + position * close;
+    __nv_bfloat16 new_value = cash + position * close;
 
     // Update max_equity BEFORE floor check — prevents stale peak from
     // missing breaches or triggering false ones after profitable trades.
-    max_equity = fmaxf(max_equity, new_value);
+    max_equity = bf16_fmax(max_equity, new_value);
 
     // ── Post-trade floor check: catch intra-step breaches ────────────
     // The pre-trade check (top of kernel) catches breaches from the
@@ -200,18 +224,21 @@ extern "C" __global__ void backtest_env_step(
     // or transaction costs. Without it, max_dd overshoots the 75% floor
     // because the metrics kernel records the breached value before the
     // next step's pre-trade check fires.
-    if (check_capital_floor(new_value, max_equity)) {
+    if (check_capital_floor(__bfloat162float(new_value), __bfloat162float(max_equity))) {
         // Emergency liquidation: close position at current price (notional model)
-        if (fabsf(position) > 0.001f) {
-            float exit_cost = compute_tx_cost(position, close, tx_cost_bps, spread_cost,
-                                              max_position, 0, -1.0f);
-            cash += position * close;   // sell position at market (notional)
-            cash -= exit_cost;
-            position = 0.0f;
+        if (__bfloat162float(bf16_fabs(position)) > 0.001f) {
+            float f_exit_cost = compute_tx_cost(__bfloat162float(position), __bfloat162float(close),
+                                                tx_cost_bps, spread_cost,
+                                                max_position, 0, -1.0f);
+            cash = cash + position * close;   // sell position at market (notional)
+            cash = cash - bf16(f_exit_cost);
+            position = bf16_zero();
             new_value = cash;           // flat: equity = cash
-            entry_price = 0.0f;
+            entry_price = bf16_zero();
         }
-        float step_ret = (value > 0.01f) ? (new_value - value) / value : 0.0f;
+        __nv_bfloat16 step_ret = (value > bf16(0.01f))
+            ? (new_value - value) / value
+            : bf16_zero();
         handle_capital_floor_breach(portfolio_state, ps, new_value, step_ret,
                                     step_rewards, step_returns, actions_history, done_flags,
                                     w, max_len, current_step, b0_size);
@@ -219,9 +246,11 @@ extern "C" __global__ void backtest_env_step(
     }
 
     // Step return
-    float step_ret = (value > 0.0f) ? (new_value - value) / value : 0.0f;
-    float new_cum_return = cum_return + step_ret;
-    float new_max = fmaxf(max_equity, new_value);
+    __nv_bfloat16 step_ret = (value > bf16_zero())
+        ? (new_value - value) / value
+        : bf16_zero();
+    __nv_bfloat16 new_cum_return = cum_return + step_ret;
+    __nv_bfloat16 new_max = bf16_fmax(max_equity, new_value);
 
     // Write portfolio state
     portfolio_state[ps + 0] = new_value;
@@ -231,7 +260,7 @@ extern "C" __global__ void backtest_env_step(
     portfolio_state[ps + 4] = new_max;
     portfolio_state[ps + 5] = hold_time;
     portfolio_state[ps + 6] = new_cum_return;
-    portfolio_state[ps + 7] += 1.0f;
+    portfolio_state[ps + 7] = portfolio_state[ps + 7] + bf16_one();
 
     // Outputs
     step_rewards[w] = step_ret;
@@ -240,8 +269,8 @@ extern "C" __global__ void backtest_env_step(
     // Re-encode position -> exposure_idx -> factored action to ensure metrics
     // kernel counts real position changes, not model-requested actions.
     {
-        float actual_exposure_frac = position / fmaxf(max_position, 0.01f);
-        int actual_exp_idx = (int)roundf((actual_exposure_frac + 1.0f) * 0.5f * (float)(b0_size - 1));
+        float f_actual_frac = __bfloat162float(position) / fmaxf(max_position, 0.01f);
+        int actual_exp_idx = (int)roundf((f_actual_frac + 1.0f) * 0.5f * (float)(b0_size - 1));
         if (actual_exp_idx < 0) actual_exp_idx = 0;
         if (actual_exp_idx >= b0_size) actual_exp_idx = b0_size - 1;
         // Preserve original order/urgency from model action
diff --git a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
index 55eba84cf..237750f9e 100644
--- a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
@@ -7,7 +7,7 @@
 //   [market_dim+8 .. market_dim+8+16)      : 16 multi-timeframe features
 //   [market_dim+24 .. state_dim)           : zero-pad / OFI / alignment
 //
-// Portfolio state layout per window [8 floats]:
+// Portfolio state layout per window [8 bf16]:
 //   [0] value        - current portfolio value
 //   [1] position     - current position size (-1.0 to +1.0)
 //   [2] cash         - cash balance
@@ -18,31 +18,31 @@
 //   [7] step_count   - number of completed steps
 //
 // 8 portfolio features (matches training experience_state_gather):
-//   +0: position                       — raw position [-1, +1]
-//   +1: unrealized_pnl / equity        — how is THIS trade doing?
-//   +2: drawdown                       — (peak - equity) / peak [0, 1]
-//   +3: hold_time / 100.0              — normalized holding duration (step_count as proxy)
-//   +4: realized_pnl / equity          — cumulative return signal
-//   +5: distance_to_floor              — (equity - floor) / equity [0, 1]
-//   +6: trade_return                   — unrealized PnL / equity
-//   +7: cash / equity                  — available capital ratio
+//   +0: position                       -- raw position [-1, +1]
+//   +1: unrealized_pnl / equity        -- how is THIS trade doing?
+//   +2: drawdown                       -- (peak - equity) / peak [0, 1]
+//   +3: hold_time / 100.0              -- normalized holding duration (step_count as proxy)
+//   +4: realized_pnl / equity          -- cumulative return signal
+//   +5: distance_to_floor              -- (equity - floor) / equity [0, 1]
+//   +6: trade_return                   -- unrealized PnL / equity
+//   +7: cash / equity                  -- available capital ratio
 //
-// 16 multi-timeframe features (4 lookbacks × 4 features):
+// 16 multi-timeframe features (4 lookbacks x 4 features):
 //   Lookback windows: 5, 15, 60, 240 bars
 //   Per window: return, volatility, volume_trend, momentum
 
 extern "C" __global__ void gather_states(
     const __nv_bfloat16* __restrict__ features,      // [n_windows * max_len * feat_dim]
     const __nv_bfloat16* __restrict__ portfolio,     // [n_windows * 8]
-    float*       states_out,                 // [n_windows * state_dim]
+    __nv_bfloat16*       states_out,                 // [n_windows * state_dim]
     int          n_windows,
     int          max_len,
     int          feat_dim,
     int          state_dim,
     int          current_step,
-    float        initial_capital,
-    float        spread_cost,
-    int          ofi_dim                     // 0 = no OFI, 8 = standard OFI features
+    float        initial_capital,       // host scalar — convert on first use
+    float        spread_cost,           // host scalar — convert on first use
+    int          ofi_dim                // 0 = no OFI, 8 = standard OFI features
 ) {
     // Shared memory tile for coalesced portfolio reads — 8 BF16 values per thread
     __shared__ __nv_bfloat16 shmem_portfolio[256 * 8];
@@ -65,13 +65,13 @@ extern "C" __global__ void gather_states(
     int feat_base = (w * max_len + current_step) * feat_dim;
     int out_base  = w * state_dim;
 
-    // Read portfolio state from shared memory tile
-    float value      = (float)shmem_portfolio[local_tid * 8 + 0];
-    float position   = (float)shmem_portfolio[local_tid * 8 + 1];
-    float cash       = (float)shmem_portfolio[local_tid * 8 + 2];
-    float entry_price = (float)shmem_portfolio[local_tid * 8 + 3];
-    float max_equity = (float)shmem_portfolio[local_tid * 8 + 4];
-    float step_count = (float)shmem_portfolio[local_tid * 8 + 7];
+    // Read portfolio state from shared memory tile — native bf16
+    __nv_bfloat16 value      = shmem_portfolio[local_tid * 8 + 0];
+    __nv_bfloat16 position   = shmem_portfolio[local_tid * 8 + 1];
+    __nv_bfloat16 cash       = shmem_portfolio[local_tid * 8 + 2];
+    __nv_bfloat16 entry_price = shmem_portfolio[local_tid * 8 + 3];
+    __nv_bfloat16 max_equity = shmem_portfolio[local_tid * 8 + 4];
+    __nv_bfloat16 step_count = shmem_portfolio[local_tid * 8 + 7];
 
     // Determine market_dim: if OFI is present, market features are feat_dim - ofi_dim
     int market_dim = (ofi_dim > 0) ? (feat_dim - ofi_dim) : feat_dim;
@@ -79,51 +79,51 @@ extern "C" __global__ void gather_states(
     // ── 1. Market features [0 .. market_dim) ── //
     int i = 0;
     for (; i + 3 < market_dim; i += 4) {
-        states_out[out_base + i]     = (float)__ldg(&features[feat_base + i]);
-        states_out[out_base + i + 1] = (float)__ldg(&features[feat_base + i + 1]);
-        states_out[out_base + i + 2] = (float)__ldg(&features[feat_base + i + 2]);
-        states_out[out_base + i + 3] = (float)__ldg(&features[feat_base + i + 3]);
+        states_out[out_base + i]     = __ldg(&features[feat_base + i]);
+        states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
+        states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
+        states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
     }
     for (; i < market_dim; i++) {
-        states_out[out_base + i] = (float)__ldg(&features[feat_base + i]);
+        states_out[out_base + i] = __ldg(&features[feat_base + i]);
     }
 
     // ── 2. Portfolio features [market_dim .. market_dim+8) ── //
     // Match training experience_state_gather exactly.
-    float equity = (value > 1.0f) ? value : 1.0f;
+    __nv_bfloat16 equity = (value > bf16_one()) ? value : bf16_one();
 
     // Drawdown from peak
-    float drawdown = (max_equity > 1.0f)
+    __nv_bfloat16 drawdown = (max_equity > bf16_one())
                    ? (max_equity - value) / max_equity
-                   : 0.0f;
-    if (drawdown < 0.0f) drawdown = 0.0f;
+                   : bf16_zero();
+    if (drawdown < bf16_zero()) drawdown = bf16_zero();
 
     // Unrealized PnL: position * (current_close - entry_price)
     // Use feature[0] as close price proxy (same as training kernel)
-    float close_now = (float)__ldg(&features[feat_base]);
-    float unrealized_pnl = (entry_price > 0.0f && position != 0.0f)
+    __nv_bfloat16 close_now = __ldg(&features[feat_base]);
+    __nv_bfloat16 unrealized_pnl = (entry_price > bf16_zero() && position != bf16_zero())
                          ? position * (close_now - entry_price)
-                         : 0.0f;
+                         : bf16_zero();
 
     // Realized PnL: derive from equity state
     // In backtest, realized PnL = value - initial_capital (approximate)
-    float realized_pnl = value - initial_capital;
+    __nv_bfloat16 realized_pnl = value - bf16(initial_capital);
 
     // Trade return (unrealized / equity)
-    float trade_return = unrealized_pnl / equity;
+    __nv_bfloat16 trade_return = unrealized_pnl / equity;
 
     // Capital floor distance (75% of peak)
-    float floor = max_equity * 0.75f;
-    float floor_dist = (equity > floor && equity > 1.0f)
-                     ? (equity - floor) / equity
-                     : 0.0f;
+    __nv_bfloat16 floor_val = max_equity * bf16(0.75f);
+    __nv_bfloat16 floor_dist = (equity > floor_val && equity > bf16_one())
+                     ? (equity - floor_val) / equity
+                     : bf16_zero();
 
     int pf_base = market_dim;
     if (pf_base + 7 < state_dim) {
         states_out[out_base + pf_base + 0] = position;                    // raw position
         states_out[out_base + pf_base + 1] = unrealized_pnl / equity;     // trade P&L signal
         states_out[out_base + pf_base + 2] = drawdown;                    // risk: how deep are we?
-        states_out[out_base + pf_base + 3] = step_count / 100.0f;         // how long in trade?
+        states_out[out_base + pf_base + 3] = step_count / bf16(100.0f);   // how long in trade?
         states_out[out_base + pf_base + 4] = realized_pnl / equity;       // session P&L
         states_out[out_base + pf_base + 5] = floor_dist;                  // distance to game over
         states_out[out_base + pf_base + 6] = trade_return;                // this trade's return
@@ -131,7 +131,7 @@ extern "C" __global__ void gather_states(
     }
 
     // ── 3. Multi-timeframe features [market_dim+8 .. market_dim+8+16) ── //
-    // 4 lookback windows × 4 features = 16 GPU-native features.
+    // 4 lookback windows x 4 features = 16 GPU-native features.
     // Matches training experience_state_gather exactly.
     const int lookbacks[4] = {5, 15, 60, 240};
     int mtf_base = market_dim + 8;
@@ -147,52 +147,62 @@ extern "C" __global__ void gather_states(
         if (past_step >= 0 && slot + 3 < state_dim) {
             int now_feat_off  = (window_bar_base + current_step) * feat_dim;
             int past_feat_off = (window_bar_base + past_step) * feat_dim;
-            float close_cur  = (float)__ldg(&features[now_feat_off]);
-            float close_past = (float)__ldg(&features[past_feat_off]);
+            __nv_bfloat16 close_cur  = __ldg(&features[now_feat_off]);
+            __nv_bfloat16 close_past = __ldg(&features[past_feat_off]);
 
             // Return over N bars
-            float ret = (close_past > 0.0f) ? (close_cur - close_past) / close_past : 0.0f;
-            float scaled_ret = ret * 100.0f;
-            states_out[out_base + slot + 0] = fmaxf(-10.0f, fminf(10.0f, scaled_ret));
+            __nv_bfloat16 ret = (close_past > bf16_zero())
+                ? (close_cur - close_past) / close_past
+                : bf16_zero();
+            __nv_bfloat16 scaled_ret = ret * bf16(100.0f);
+            states_out[out_base + slot + 0] = bf16_fmax(bf16(-10.0f), bf16_fmin(bf16(10.0f), scaled_ret));
 
             // Volatility: scan high/low over window (from close changes)
-            float max_val = close_cur;
-            float min_val = close_cur;
-            float vol_sum = 0.0f;
+            __nv_bfloat16 max_val = close_cur;
+            __nv_bfloat16 min_val = close_cur;
+            __nv_bfloat16 vol_sum = bf16_zero();
             int vol_count = 0;
             for (int j = past_step; j <= current_step; j++) {
                 int j_off = (window_bar_base + j) * feat_dim;
-                float v = (float)__ldg(&features[j_off]);
+                __nv_bfloat16 v = __ldg(&features[j_off]);
                 if (v > max_val) max_val = v;
                 if (v < min_val) min_val = v;
                 // Volume proxy: feature index 4
                 if (feat_dim > 4) {
-                    vol_sum += (float)__ldg(&features[j_off + 4]);
+                    vol_sum = vol_sum + __ldg(&features[j_off + 4]);
                     vol_count++;
                 }
             }
-            float range = (close_cur > 0.0f) ? (max_val - min_val) / close_cur : 0.0f;
-            float scaled_range = range * 100.0f;
-            states_out[out_base + slot + 1] = fmaxf(0.0f, fminf(10.0f, scaled_range));
+            __nv_bfloat16 range = (close_cur > bf16_zero())
+                ? (max_val - min_val) / close_cur
+                : bf16_zero();
+            __nv_bfloat16 scaled_range = range * bf16(100.0f);
+            states_out[out_base + slot + 1] = bf16_fmax(bf16_zero(), bf16_fmin(bf16(10.0f), scaled_range));
 
             // Volume trend: current vs average
-            float avg_vol = (vol_count > 0) ? vol_sum / (float)vol_count : 1.0f;
-            float cur_vol = (feat_dim > 4) ? (float)__ldg(&features[now_feat_off + 4]) : 1.0f;
-            float vol_ratio = (avg_vol > 0.0f) ? cur_vol / avg_vol : 1.0f;
-            states_out[out_base + slot + 2] = fmaxf(0.0f, fminf(5.0f, vol_ratio));
+            __nv_bfloat16 avg_vol = (vol_count > 0)
+                ? vol_sum / bf16((float)vol_count)
+                : bf16_one();
+            __nv_bfloat16 cur_vol = (feat_dim > 4)
+                ? __ldg(&features[now_feat_off + 4])
+                : bf16_one();
+            __nv_bfloat16 vol_ratio = (avg_vol > bf16_zero())
+                ? cur_vol / avg_vol
+                : bf16_one();
+            states_out[out_base + slot + 2] = bf16_fmax(bf16_zero(), bf16_fmin(bf16(5.0f), vol_ratio));
 
             // Momentum: position within range [0=bottom, 1=top]
-            float range_size = max_val - min_val;
-            states_out[out_base + slot + 3] = (range_size > 0.0f)
+            __nv_bfloat16 range_size = max_val - min_val;
+            states_out[out_base + slot + 3] = (range_size > bf16_zero())
                           ? (close_cur - min_val) / range_size
-                          : 0.5f;
+                          : bf16(0.5f);
         } else {
             // Not enough history — zero pad
             if (slot + 3 < state_dim) {
-                states_out[out_base + slot + 0] = 0.0f;
-                states_out[out_base + slot + 1] = 0.0f;
-                states_out[out_base + slot + 2] = 0.0f;
-                states_out[out_base + slot + 3] = 0.0f;
+                states_out[out_base + slot + 0] = bf16_zero();
+                states_out[out_base + slot + 1] = bf16_zero();
+                states_out[out_base + slot + 2] = bf16_zero();
+                states_out[out_base + slot + 3] = bf16_zero();
             }
         }
     }
@@ -202,26 +212,26 @@ extern "C" __global__ void gather_states(
         int ofi_out_base = market_dim + 8 + 16;
         // OFI features are stored after market features in the feature vector
         for (i = 0; i + 3 < ofi_dim; i += 4) {
-            states_out[out_base + ofi_out_base + i]     = (float)__ldg(&features[feat_base + market_dim + i]);
-            states_out[out_base + ofi_out_base + i + 1] = (float)__ldg(&features[feat_base + market_dim + i + 1]);
-            states_out[out_base + ofi_out_base + i + 2] = (float)__ldg(&features[feat_base + market_dim + i + 2]);
-            states_out[out_base + ofi_out_base + i + 3] = (float)__ldg(&features[feat_base + market_dim + i + 3]);
+            states_out[out_base + ofi_out_base + i]     = __ldg(&features[feat_base + market_dim + i]);
+            states_out[out_base + ofi_out_base + i + 1] = __ldg(&features[feat_base + market_dim + i + 1]);
+            states_out[out_base + ofi_out_base + i + 2] = __ldg(&features[feat_base + market_dim + i + 2]);
+            states_out[out_base + ofi_out_base + i + 3] = __ldg(&features[feat_base + market_dim + i + 3]);
         }
         for (; i < ofi_dim; i++) {
-            states_out[out_base + ofi_out_base + i] = (float)__ldg(&features[feat_base + market_dim + i]);
+            states_out[out_base + ofi_out_base + i] = __ldg(&features[feat_base + market_dim + i]);
         }
     }
 
     // ── 5. Zero-pad remaining [filled .. state_dim) ── //
     int filled = market_dim + 8 + 16 + ofi_dim;
     for (i = filled; i + 3 < state_dim; i += 4) {
-        states_out[out_base + i]     = 0.0f;
-        states_out[out_base + i + 1] = 0.0f;
-        states_out[out_base + i + 2] = 0.0f;
-        states_out[out_base + i + 3] = 0.0f;
+        states_out[out_base + i]     = bf16_zero();
+        states_out[out_base + i + 1] = bf16_zero();
+        states_out[out_base + i + 2] = bf16_zero();
+        states_out[out_base + i + 3] = bf16_zero();
     }
     for (; i < state_dim; i++) {
-        states_out[out_base + i] = 0.0f;
+        states_out[out_base + i] = bf16_zero();
     }
 
     // Suppress unused variable warning
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 2cf322518..dfc9d1d88 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -3,7 +3,7 @@
  *
  * Three focused CUDA kernels that handle the non-Q-forward parts of DQN
  * experience collection.  The Q-network forward pass is performed by cuBLAS
- * SGEMM from the Rust host between kernel launches.
+ * GemmEx BF16 from the Rust host between kernel launches.
  *
  * Kernels:
  *   1. experience_state_gather  — assemble batch state tensor for cuBLAS
@@ -13,10 +13,7 @@
  * Design:
  *   - Grid: ceil(N/256), Block: 256.  One thread per episode — trivially parallel.
  *   - No shared memory required.
- *   - Standalone file: does NOT include common_device_functions.cuh.
- *     Constants normally injected via NVRTC are guarded with #ifndef
- *     fallback defaults; production use injects the correct values at
- *     NVRTC compile time.
+ *   - Native __nv_bfloat16 everywhere — zero float storage on GPU.
  *
  * State layout (matches common_device_functions.cuh):
  *   [0 .. market_dim)                      : market features
@@ -32,7 +29,7 @@
  *   [5]  (reserved)      — was pnl_ema, unused by reward v6
  *   [6]  (reserved)      — was pnl_var, unused by reward v6
  *   [7]  peak_equity     — high-water mark (init to initial_capital)
- *   [8]  flat_counter    — consecutive flat steps (float for GPU simplicity)
+ *   [8]  flat_counter    — consecutive flat steps
  *   [9]  prev_equity     — equity at previous step (init to initial_capital)
  *   [10] hold_time       — consecutive steps with position (total, not just losing)
  *   [11] realized_pnl    — cumulative realized PnL
@@ -49,7 +46,7 @@
  */
 
 /* ------------------------------------------------------------------ */
-/* Portfolio stride for experience kernels (12 floats per episode).    */
+/* Portfolio stride for experience kernels (20 bf16 per episode).      */
 /* portfolio_sim_kernel uses its own stride of 8 — do NOT change it.  */
 /* ------------------------------------------------------------------ */
 #define PORTFOLIO_STRIDE 20
@@ -59,16 +56,10 @@
 #include "trade_physics.cuh"
 
 /* ------------------------------------------------------------------ */
-/* BF16 I/O helper — uses the __nv_bfloat16 explicit float conversion  */
-/* operator; keeps this file free of raw bfloat16 intrinsic calls.    */
+/* BF16 support — common_device_functions.cuh is prepended by build.rs */
 /* ------------------------------------------------------------------ */
 #include <cuda_bf16.h>
 
-/** Load a __nv_bfloat16 value from a BF16 buffer into a float. */
-__device__ __forceinline__ float bf16f(__nv_bfloat16 x) {
-    return (float)x;
-}
-
 /* ------------------------------------------------------------------ */
 /* Compile-time constants — overridable via NVRTC #define injection.  */
 /* ------------------------------------------------------------------ */
@@ -92,9 +83,9 @@ __device__ __forceinline__ float bf16f(__nv_bfloat16 x) {
 /**
  * LCG random — returns float in [0, 1).
  *
- * Identical implementation to gpu_random() in common_device_functions.cuh
- * so episodes share the same statistical properties when mixing old and
- * new kernels during the transition period.
+ * RNG state manipulation must stay in 32-bit unsigned int (integer math).
+ * Returns float because [0,1) probability comparisons with epsilon need
+ * more precision than bf16 can offer (bf16 has only 8 mantissa bits).
  */
 __device__ __forceinline__ float lcg_random(unsigned int* state) {
     *state = *state * 1664525u + 1013904223u;
@@ -102,13 +93,13 @@ __device__ __forceinline__ float lcg_random(unsigned int* state) {
 }
 
 /**
- * Argmax over a float array of length n.
+ * Argmax over a __nv_bfloat16 array of length n.
  * Returns the index of the maximum element; ties broken in favour of the
  * lowest index (first maximum found).
  */
-__device__ __forceinline__ int argmax_n(const float* arr, int n) {
+__device__ __forceinline__ int argmax_n(const __nv_bfloat16* arr, int n) {
     int best_idx = 0;
-    float best_val = arr[0];
+    __nv_bfloat16 best_val = arr[0];
     for (int j = 1; j < n; j++) {
         if (arr[j] > best_val) {
             best_val = arr[j];
@@ -127,15 +118,15 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) {
 
 /**
  * Gather market features and portfolio state into a flat batch tensor
- * suitable for cuBLAS SGEMM (Q-network forward pass).
+ * suitable for cuBLAS GemmEx (Q-network forward pass).
  *
  * Grid: ceil(N / 256), Block: 256.  One thread per episode.
  *
- * @param market_features   [total_bars, market_dim]  raw market feature matrix
+ * @param market_features   [total_bars, market_dim]  raw market feature matrix (bf16)
  * @param episode_starts    [N]                        global bar index of episode start
  * @param current_timesteps [N]                        timestep offset within episode (read-only)
- * @param portfolio_states  [N, 3]                     {position, cash, portfolio_value}
- * @param batch_states      [N, state_dim]             output: assembled state batch
+ * @param portfolio_states  [N, 3]                     {position, cash, portfolio_value} (bf16)
+ * @param batch_states      [N, state_dim]             output: assembled state batch (bf16)
  * @param N                                            number of episodes
  * @param total_bars                                   number of bars in market_features
  * @param state_dim                                    full state dimension (tensor-core aligned)
@@ -146,7 +137,7 @@ extern "C" __global__ void experience_state_gather(
     const int*   __restrict__ episode_starts,
     const int*   __restrict__ current_timesteps,
     const __nv_bfloat16* __restrict__ portfolio_states,
-    float*                    batch_states,
+    __nv_bfloat16*                batch_states,
     int N,
     int total_bars,
     int state_dim,
@@ -157,19 +148,19 @@ extern "C" __global__ void experience_state_gather(
 
     int bar_idx = episode_starts[i] + current_timesteps[i];
 
-    float* out = batch_states + (long long)i * state_dim;
+    __nv_bfloat16* out = batch_states + (long long)i * state_dim;
 
     /* Out-of-data: write zero state; cuBLAS will produce Q-values of ~0. */
     if (bar_idx >= total_bars) {
         for (int k = 0; k < state_dim; k++)
-            out[k] = 0.0f;
+            out[k] = bf16_zero();
         return;
     }
 
     /* -- Market features: [0 .. market_dim) -- */
     const __nv_bfloat16* mf_row = market_features + (long long)bar_idx * market_dim;
     for (int k = 0; k < market_dim; k++)
-        out[k] = bf16f(mf_row[k]);
+        out[k] = mf_row[k];
 
     /* -- Portfolio features: [market_dim .. market_dim+8) --
      *
@@ -190,57 +181,58 @@ extern "C" __global__ void experience_state_gather(
      * The network learns the actual scale via the position feature.
      */
     const __nv_bfloat16* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
-    float position           = bf16f(ps[0]);
-    float cash               = bf16f(ps[1]);
-    float portfolio_value    = bf16f(ps[2]);
-    float peak_equity        = bf16f(ps[7]);
-    float prev_equity        = bf16f(ps[9]);
-    float hold_time          = bf16f(ps[10]);
-    float realized_pnl       = bf16f(ps[11]);
-    float entry_price        = bf16f(ps[12]);
-    float trade_start_pnl    = bf16f(ps[13]);
+    __nv_bfloat16 position        = ps[0];
+    __nv_bfloat16 cash            = ps[1];
+    __nv_bfloat16 portfolio_value = ps[2];
+    __nv_bfloat16 peak_equity     = ps[7];
+    __nv_bfloat16 prev_equity     = ps[9];
+    __nv_bfloat16 hold_time       = ps[10];
+    __nv_bfloat16 realized_pnl    = ps[11];
+    __nv_bfloat16 entry_price     = ps[12];
+    __nv_bfloat16 trade_start_pnl = ps[13];
 
-    float equity = (portfolio_value > 1.0f) ? portfolio_value : 1.0f;
-    float drawdown = (peak_equity > 1.0f)
+    __nv_bfloat16 one = bf16_one();
+    __nv_bfloat16 equity = (portfolio_value > one) ? portfolio_value : one;
+    __nv_bfloat16 drawdown = (peak_equity > one)
                    ? (peak_equity - portfolio_value) / peak_equity
-                   : 0.0f;
-    drawdown = (drawdown > 0.0f) ? drawdown : 0.0f;
+                   : bf16_zero();
+    drawdown = (drawdown > bf16_zero()) ? drawdown : bf16_zero();
 
     /* Unrealized P&L: current position mark-to-market minus entry cost */
-    float unrealized_pnl = (entry_price > 0.0f && position != 0.0f)
-                         ? position * (bf16f(market_features[(long long)bar_idx * market_dim]) - entry_price)
-                         : 0.0f;
+    __nv_bfloat16 unrealized_pnl = (entry_price > bf16_zero() && position != bf16_zero())
+                         ? position * (market_features[(long long)bar_idx * market_dim] - entry_price)
+                         : bf16_zero();
 
     /* Trade return since entry */
-    float trade_return = (trade_start_pnl != 0.0f || realized_pnl != 0.0f)
+    __nv_bfloat16 trade_return = (trade_start_pnl != bf16_zero() || realized_pnl != bf16_zero())
                        ? (realized_pnl - trade_start_pnl + unrealized_pnl) / equity
-                       : 0.0f;
+                       : bf16_zero();
 
     /* Capital floor distance (75% of peak = capital floor) */
-    float floor = peak_equity * 0.75f;
-    float floor_dist = (equity > floor && equity > 1.0f)
-                     ? (equity - floor) / equity
-                     : 0.0f;
+    __nv_bfloat16 floor_val = peak_equity * bf16(0.75f);
+    __nv_bfloat16 floor_dist = (equity > floor_val && equity > one)
+                     ? (equity - floor_val) / equity
+                     : bf16_zero();
 
     int portfolio_base = market_dim;
     if (portfolio_base + 7 < state_dim) {
-        out[portfolio_base + 0] = position;                        /* raw position */
-        out[portfolio_base + 1] = unrealized_pnl / equity;         /* trade P&L signal */
-        out[portfolio_base + 2] = drawdown;                        /* risk: how deep are we? */
-        out[portfolio_base + 3] = hold_time / 100.0f;              /* how long in trade? */
-        out[portfolio_base + 4] = realized_pnl / equity;           /* session P&L */
-        out[portfolio_base + 5] = floor_dist;                      /* distance to game over */
-        out[portfolio_base + 6] = trade_return;                    /* this trade's return */
-        out[portfolio_base + 7] = cash / equity;                   /* available capital */
+        out[portfolio_base + 0] = position;                              /* raw position */
+        out[portfolio_base + 1] = unrealized_pnl / equity;              /* trade P&L signal */
+        out[portfolio_base + 2] = drawdown;                              /* risk: how deep are we? */
+        out[portfolio_base + 3] = hold_time / bf16(100.0f);             /* how long in trade? */
+        out[portfolio_base + 4] = realized_pnl / equity;                /* session P&L */
+        out[portfolio_base + 5] = floor_dist;                            /* distance to game over */
+        out[portfolio_base + 6] = trade_return;                          /* this trade's return */
+        out[portfolio_base + 7] = cash / equity;                         /* available capital */
     }
 
     /* -- Multi-timeframe features: [market_dim+8 .. market_dim+8+16) --
      *
-     * 4 lookback windows × 4 features = 16 GPU-native features.
+     * 4 lookback windows x 4 features = 16 GPU-native features.
      * Computed directly from the market data already on GPU — no CPU.
      * The model sees price action at multiple timescales simultaneously.
      *
-     * Lookback windows: 5, 15, 60, 240 bars (≈ 5min, 15min, 1hr, 4hr)
+     * Lookback windows: 5, 15, 60, 240 bars
      * Features per window:
      *   +0: return = (close_now - close_N) / close_N
      *   +1: volatility = max_high - min_low over N bars / close_now
@@ -250,60 +242,58 @@ extern "C" __global__ void experience_state_gather(
     const int lookbacks[4] = {5, 15, 60, 240};
     int mtf_base = market_dim + 8;
     for (int lb = 0; lb < 4; lb++) {
-        int N = lookbacks[lb];
-        int past_idx = bar_idx - N;
+        int Nlb = lookbacks[lb];
+        int past_idx = bar_idx - Nlb;
         int slot = mtf_base + lb * 4;
 
         if (past_idx >= 0 && slot + 3 < state_dim) {
-            /* Close prices: feature[0] of each bar is typically close or a return.
-             * Use raw market_features — index 0 is the first feature per bar. */
             const __nv_bfloat16* now_row = market_features + (long long)bar_idx * market_dim;
             const __nv_bfloat16* past_row = market_features + (long long)past_idx * market_dim;
-            float close_now = bf16f(now_row[0]);
-            float close_past = bf16f(past_row[0]);
+            __nv_bfloat16 close_now = now_row[0];
+            __nv_bfloat16 close_past = past_row[0];
 
             /* Return over N bars */
-            float ret = (close_past > 0.0f) ? (close_now - close_past) / close_past : 0.0f;
-            float scaled_ret = ret * 100.0f;
-            out[slot + 0] = fmaxf(-10.0f, fminf(10.0f, scaled_ret)); /* clamp ±10% */
+            __nv_bfloat16 ret = (close_past > bf16_zero()) ? (close_now - close_past) / close_past : bf16_zero();
+            __nv_bfloat16 scaled_ret = ret * bf16(100.0f);
+            out[slot + 0] = bf16_fmax(bf16(-10.0f), bf16_fmin(bf16(10.0f), scaled_ret)); /* clamp +/-10% */
 
             /* Volatility: scan high/low over window (approx from close changes) */
-            float max_val = close_now;
-            float min_val = close_now;
-            float vol_sum = 0.0f;
+            __nv_bfloat16 max_val = close_now;
+            __nv_bfloat16 min_val = close_now;
+            __nv_bfloat16 vol_sum = bf16_zero();
             int vol_count = 0;
             for (int j = past_idx; j <= bar_idx && j < total_bars; j++) {
-                float v = bf16f(market_features[(long long)j * market_dim]);
+                __nv_bfloat16 v = market_features[(long long)j * market_dim];
                 if (v > max_val) max_val = v;
                 if (v < min_val) min_val = v;
                 /* Volume proxy: use feature index 4 if it exists (commonly volume) */
                 if (market_dim > 4) {
-                    vol_sum += bf16f(market_features[(long long)j * market_dim + 4]);
+                    vol_sum = vol_sum + market_features[(long long)j * market_dim + 4];
                     vol_count++;
                 }
             }
-            float range = (close_now > 0.0f) ? (max_val - min_val) / close_now : 0.0f;
-            float scaled_range = range * 100.0f;
-            out[slot + 1] = fmaxf(0.0f, fminf(10.0f, scaled_range)); /* clamp 0-10% */
+            __nv_bfloat16 range = (close_now > bf16_zero()) ? (max_val - min_val) / close_now : bf16_zero();
+            __nv_bfloat16 scaled_range = range * bf16(100.0f);
+            out[slot + 1] = bf16_fmax(bf16_zero(), bf16_fmin(bf16(10.0f), scaled_range)); /* clamp 0-10% */
 
             /* Volume trend: current vs average */
-            float avg_vol = (vol_count > 0) ? vol_sum / (float)vol_count : 1.0f;
-            float cur_vol = (market_dim > 4) ? bf16f(now_row[4]) : 1.0f;
-            float vol_ratio = (avg_vol > 0.0f) ? cur_vol / avg_vol : 1.0f;
-            out[slot + 2] = fmaxf(0.0f, fminf(5.0f, vol_ratio)); /* clamp 0-5x */
+            __nv_bfloat16 avg_vol = (vol_count > 0) ? vol_sum / bf16((float)vol_count) : bf16_one();
+            __nv_bfloat16 cur_vol = (market_dim > 4) ? now_row[4] : bf16_one();
+            __nv_bfloat16 vol_ratio = (avg_vol > bf16_zero()) ? cur_vol / avg_vol : bf16_one();
+            out[slot + 2] = bf16_fmax(bf16_zero(), bf16_fmin(bf16(5.0f), vol_ratio)); /* clamp 0-5x */
 
             /* Momentum: position within range [0=bottom, 1=top] */
-            float range_size = max_val - min_val;
-            out[slot + 3] = (range_size > 0.0f)
+            __nv_bfloat16 range_size = max_val - min_val;
+            out[slot + 3] = (range_size > bf16_zero())
                           ? (close_now - min_val) / range_size
-                          : 0.5f;
+                          : bf16(0.5f);
         } else {
             /* Not enough history — zero pad */
             if (slot + 3 < state_dim) {
-                out[slot + 0] = 0.0f;
-                out[slot + 1] = 0.0f;
-                out[slot + 2] = 0.0f;
-                out[slot + 3] = 0.0f;
+                out[slot + 0] = bf16_zero();
+                out[slot + 1] = bf16_zero();
+                out[slot + 2] = bf16_zero();
+                out[slot + 3] = bf16_zero();
             }
         }
     }
@@ -316,7 +306,7 @@ extern "C" __global__ void experience_state_gather(
     /* -- Zero-pad remaining dimensions (tensor-core alignment) -- */
     int filled = market_dim + 8 + 16; /* 42 market + 8 portfolio + 16 multi-tf */
     for (int k = filled; k < state_dim; k++)
-        out[k] = 0.0f;
+        out[k] = bf16_zero();
 }
 
 /* ================================================================== */
@@ -338,21 +328,25 @@ extern "C" __global__ void experience_state_gather(
  *
  * Grid: ceil(N / 256), Block: 256.  One thread per episode.
  *
- * @param q_values          [N, q_stride]  Q-values from cuBLAS
+ * @param q_values          [N, q_stride]  Q-values from cuBLAS (bf16)
  * @param out_actions       [N]            selected factored action index (output)
  * @param rng_states        [N]            per-episode LCG RNG counter (updated in place)
+ * @param out_q_gaps        [N]            output: Q-gap per episode for conviction sizing (bf16)
  * @param epsilon                          exploration probability in [0, 1]
  * @param N                                number of episodes
  * @param b0_size                          exposure branch size (default 5)
  * @param b1_size                          order branch size (default 3)
  * @param b2_size                          urgency branch size (default 3)
  * @param q_gap_threshold                  min Q-gap for trade entry (0.0 = disabled)
+ * @param portfolio_states  [N, 20]        read-only for hold enforcement (bf16)
+ * @param min_hold_bars                    minimum bars to hold before switching
+ * @param max_position                     max allowed position (host scalar)
  */
 extern "C" __global__ void experience_action_select(
-    const float*  __restrict__ q_values,
+    const __nv_bfloat16*  __restrict__ q_values,
     int*                       out_actions,
     unsigned int*              rng_states,
-    float*                     out_q_gaps,    /* [N] output: Q-gap per episode for conviction sizing */
+    __nv_bfloat16*             out_q_gaps,    /* [N] output: Q-gap per episode for conviction sizing */
     float epsilon,
     int N,
     int b0_size,
@@ -368,10 +362,10 @@ extern "C" __global__ void experience_action_select(
 
     unsigned int rng = rng_states[i];
 
-    int q_stride      = b0_size + b1_size + b2_size;
-    const float* q_b0 = q_values + (long long)i * q_stride;
-    const float* q_b1 = q_b0 + b0_size;
-    const float* q_b2 = q_b1 + b1_size;
+    int q_stride = b0_size + b1_size + b2_size;
+    const __nv_bfloat16* q_b0 = q_values + (long long)i * q_stride;
+    const __nv_bfloat16* q_b1 = q_b0 + b0_size;
+    const __nv_bfloat16* q_b2 = q_b1 + b1_size;
 
     int a0, a1, a2;
 
@@ -380,24 +374,29 @@ extern "C" __global__ void experience_action_select(
      * Portfolio state: ps[0] = position, ps[10] = hold_time.
      * When portfolio_states is NULL (backtest evaluator), skip hold enforcement. */
     int in_hold = 0;
-    float cur_position = 0.0f;
+    __nv_bfloat16 cur_position = bf16_zero();
     if (min_hold_bars > 0 && portfolio_states != NULL) {
         int ps_base = i * 20;  /* PORTFOLIO_STRIDE = 20 */
-        float hold_time_val = bf16f(portfolio_states[ps_base + 10]);
-        cur_position = bf16f(portfolio_states[ps_base + 0]);
-        in_hold = (hold_time_val > 0.0f && hold_time_val < (float)min_hold_bars);
+        __nv_bfloat16 hold_time_val = portfolio_states[ps_base + 10];
+        cur_position = portfolio_states[ps_base + 0];
+        in_hold = (hold_time_val > bf16_zero() && hold_time_val < bf16((float)min_hold_bars));
     }
 
+    __nv_bfloat16 bf16_max_pos = bf16(max_position);
+
     if (in_hold) {
         /* Compute current exposure index from position */
-        float exposure_frac = cur_position / (max_position > 0.0f ? max_position : 1.0f);
-        int cur_exp = (int)roundf((exposure_frac + 1.0f) * 0.5f * (float)(b0_size - 1));
+        __nv_bfloat16 denom = (bf16_max_pos > bf16_zero()) ? bf16_max_pos : bf16_one();
+        __nv_bfloat16 exposure_frac = cur_position / denom;
+        /* Need float precision for roundf */
+        float ef = __bfloat162float(exposure_frac);
+        int cur_exp = (int)roundf((ef + 1.0f) * 0.5f * (float)(b0_size - 1));
         if (cur_exp < 0) cur_exp = 0;
         if (cur_exp >= b0_size) cur_exp = b0_size - 1;
         a0 = cur_exp;  /* Force current exposure — skip Q-value selection entirely */
         /* Zero Q-gap during hold — conviction sizing is meaningless when forced */
         if (out_q_gaps != NULL) {
-            out_q_gaps[i] = 0.0f;
+            out_q_gaps[i] = bf16_zero();
         }
     } else {
     /* Branch 0: exposure — with Q-gap conviction filter.
@@ -412,9 +411,9 @@ extern "C" __global__ void experience_action_select(
         /* Q-gap filter: force flat when conviction is low */
         if (q_gap_threshold > 0.0f && b0_size > 2) {
             int flat_idx = b0_size / 2; /* center = Flat (index 4 for 9-action) */
-            float q_best = q_b0[a0];
-            float q_flat = q_b0[flat_idx];
-            if (q_best - q_flat < q_gap_threshold) {
+            __nv_bfloat16 q_best = q_b0[a0];
+            __nv_bfloat16 q_flat = q_b0[flat_idx];
+            if (q_best - q_flat < bf16(q_gap_threshold)) {
                 a0 = flat_idx;
             }
         }
@@ -437,7 +436,7 @@ extern "C" __global__ void experience_action_select(
         a2 = argmax_n(q_b2, b2_size);
     }
 
-    /* Compose factored action (Tavakoli et al., 2018, §3.1) */
+    /* Compose factored action (Tavakoli et al., 2018, S3.1) */
     int action_idx = a0 * b1_size * b2_size + a1 * b2_size + a2;
 
     out_actions[i] = action_idx;
@@ -449,10 +448,10 @@ extern "C" __global__ void experience_action_select(
      * Skip when in_hold — already zeroed in the hold branch above. */
     if (out_q_gaps != NULL && !in_hold) {
         int flat_idx = b0_size / 2;
-        float q_best = q_b0[argmax_n(q_b0, b0_size)];
-        float q_flat = q_b0[flat_idx];
-        float gap = q_best - q_flat;
-        out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f;
+        __nv_bfloat16 q_best = q_b0[argmax_n(q_b0, b0_size)];
+        __nv_bfloat16 q_flat = q_b0[flat_idx];
+        __nv_bfloat16 gap = q_best - q_flat;
+        out_q_gaps[i] = (gap > bf16_zero()) ? gap : bf16_zero();
     }
 }
 
@@ -485,50 +484,19 @@ extern "C" __global__ void experience_action_select(
  *   [2] raw_close      — raw close price (for position cost + tx)
  *   [3] raw_next       — raw next-bar close (for mark-to-market PnL)
  *
- * portfolio_states layout: [N, PORTFOLIO_STRIDE=20]  (read-write)
+ * portfolio_states layout: [N, PORTFOLIO_STRIDE=20]  (read-write, bf16)
  *   See file header for field definitions.
- *
- * @param targets            [total_bars, 4]   price data
- * @param episode_starts     [N]               global bar offset per episode
- * @param current_timesteps  [N]               current episode step (incremented)
- * @param actions            [N]               factored action from action_select
- * @param portfolio_states   [N, 20]           full per-episode state (read-write)
- * @param out_states         [N, L, state_dim] replay output: states
- * @param out_actions        [N, L]            replay output: actions
- * @param out_rewards        [N, L]            replay output: rewards
- * @param out_dones          [N, L]            replay output: done flags
- * @param batch_states       [N, state_dim]    current state batch from gather kernel
- * @param max_position                         max allowed contract position (absolute)
- * @param tx_cost_multiplier                   proportional transaction cost rate
- * @param loss_aversion                        asymmetric loss scaling factor (~1.5)
- * @param features           [total_bars, market_dim] market features for regime detection
- * @param market_dim                           MARKET_DIM (42)
- * @param L                                    episode length (replay buffer columns)
- * @param N                                    number of episodes
- * @param total_bars                           length of targets/market_features
- * @param state_dim                            state vector length
- * @param b0_size                              exposure branch size (default 9)
- * @param b1_size                              order branch size (default 3)
- * @param b2_size                              urgency branch size (default 3)
- * @param current_t                            timestep index (0-based) in this batch
- * @param cvar_scales        [N] or NULL       CVaR position scaling
- * @param q_gaps             [N] or NULL       Q-gap conviction scaling
- * @param raw_returns_out    [N, L] or NULL    true per-bar portfolio return (unshaped)
- * @param min_hold_bars                       minimum bars to hold before exiting or reversing
- * @param spread_cost                        bid-ask spread cost per unit (matches backtest)
- * @param contract_multiplier                dollar multiplier per point (50 for ES, 20 for NQ)
- * @param margin_pct                         initial margin as fraction of notional (0.06 = 6%)
  */
 extern "C" __global__ void experience_env_step(
     const __nv_bfloat16* __restrict__ targets,
     const int*   __restrict__ episode_starts,
     int*                      current_timesteps,
     const int*   __restrict__ actions,
-    float*                    portfolio_states,
-    float*                    out_states,
+    __nv_bfloat16*            portfolio_states,
+    __nv_bfloat16*            out_states,
     int*                      out_actions,
-    float*                    out_rewards,
-    float*                    out_dones,
+    __nv_bfloat16*            out_rewards,
+    __nv_bfloat16*            out_dones,
     const __nv_bfloat16* __restrict__ batch_states,
     float max_position,
     float tx_cost_multiplier,
@@ -545,7 +513,7 @@ extern "C" __global__ void experience_env_step(
     int current_t,
     const __nv_bfloat16* __restrict__ cvar_scales,  /* [N] or NULL — CVaR position scaling */
     const __nv_bfloat16* __restrict__ q_gaps,       /* [N] or NULL — Q-gap conviction scaling */
-    float*                    raw_returns_out, /* [N, L] output: true per-bar portfolio return (unshapen) */
+    __nv_bfloat16*            raw_returns_out, /* [N, L] output: true per-bar portfolio return */
     int min_hold_bars,  /* minimum bars to hold before exiting or reversing */
     float spread_cost,  /* bid-ask spread cost per unit (matches backtest) */
     float contract_multiplier,  /* e.g. 50.0 for ES, 20.0 for NQ */
@@ -562,9 +530,9 @@ extern "C" __global__ void experience_env_step(
 
     /* ---- Write current state to replay buffer ---- */
     const __nv_bfloat16* src = batch_states + (long long)i * state_dim;
-    float*       dst = out_states   + out_off * state_dim;
+    __nv_bfloat16*       dst = out_states   + out_off * state_dim;
     for (int k = 0; k < state_dim; k++)
-        dst[k] = bf16f(src[k]);
+        dst[k] = src[k];
 
     /* ---- Write action ---- */
     int action_idx = actions[i];
@@ -572,8 +540,8 @@ extern "C" __global__ void experience_env_step(
 
     /* ---- Out-of-data check ---- */
     if (bar_idx >= total_bars - 1) {
-        out_rewards[out_off] = 0.0f;
-        out_dones[out_off]   = 1.0f;
+        out_rewards[out_off] = bf16_zero();
+        out_dones[out_off]   = bf16_one();
         /* Do not advance timestep: episode is at the data boundary. */
         return;
     }
@@ -583,32 +551,42 @@ extern "C" __global__ void experience_env_step(
      *   [0:1] = preprocessed (log-return normalized) — for Q-network
      *   [2:3] = raw dollar prices — for portfolio simulation P&L + tx costs */
     const __nv_bfloat16* tgt = targets + (long long)bar_idx * 4;
-    float raw_close  = bf16f(tgt[2]);
-    float raw_next   = bf16f(tgt[3]);
+    /* Trade physics (execute_trade, compute_tx_cost, etc.) use float internally
+     * because they involve multi-step accumulation where bf16 precision is
+     * insufficient (e.g. cash -= delta * price; cash -= cost). We convert
+     * prices to float here for the physics section, then store results back
+     * as bf16 at the end. */
+    float raw_close  = __bfloat162float(tgt[2]);
+    float raw_next   = __bfloat162float(tgt[3]);
 
     /* Guard against degenerate prices from data gaps. */
     if (raw_close <= 0.0f) raw_close = 1.0f;
     if (raw_next  <= 0.0f) raw_next  = raw_close;
 
     /* ---- Read full portfolio state (PORTFOLIO_STRIDE=20) ---- */
-    float* ps            = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
-    float position       = ps[0];
-    float cash           = ps[1];
+    /* Portfolio arithmetic uses float accumulators because trade physics
+     * functions (execute_trade, apply_margin_cap, etc.) in trade_physics.cuh
+     * are all float. Converting the entire physics engine to bf16 would
+     * require rewriting the shared header used by backtest_env_kernel too.
+     * We load bf16 → float here, compute, then store float → bf16 at end. */
+    __nv_bfloat16* ps       = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
+    float position           = __bfloat162float(ps[0]);
+    float cash               = __bfloat162float(ps[1]);
     /* ps[2] = portfolio_value (updated at end) */
     /* ps[3:6] reserved (unused by reward v6) */
-    float peak_equity    = ps[7];
-    float flat_counter   = ps[8];
-    float prev_equity    = ps[9];
-    float hold_time      = ps[10];
+    float peak_equity        = __bfloat162float(ps[7]);
+    float flat_counter       = __bfloat162float(ps[8]);
+    float prev_equity        = __bfloat162float(ps[9]);
+    float hold_time          = __bfloat162float(ps[10]);
     /* ps[11] = realized_pnl (cumulative, updated at end) */
-    float entry_price    = ps[12];   /* price when trade was entered */
-    float trade_start_pnl = ps[13]; /* realized_pnl snapshot at trade entry */
-    float win_count      = ps[14];  /* Kelly: number of profitable trade exits */
-    float loss_count     = ps[15];  /* Kelly: number of losing trade exits */
-    float sum_wins       = ps[16];  /* Kelly: cumulative profit from winners */
-    float sum_losses     = ps[17];  /* Kelly: cumulative |loss| from losers */
-    float sum_returns    = ps[18];  /* Kelly: cumulative net returns (for μ) */
-    float sum_sq_returns = ps[19];  /* Kelly: cumulative squared returns (for σ²) */
+    float entry_price        = __bfloat162float(ps[12]);   /* price when trade was entered */
+    float trade_start_pnl    = __bfloat162float(ps[13]); /* realized_pnl snapshot at trade entry */
+    float win_count          = __bfloat162float(ps[14]);  /* Kelly: number of profitable trade exits */
+    float loss_count         = __bfloat162float(ps[15]);  /* Kelly: number of losing trade exits */
+    float sum_wins           = __bfloat162float(ps[16]);  /* Kelly: cumulative profit from winners */
+    float sum_losses         = __bfloat162float(ps[17]);  /* Kelly: cumulative |loss| from losers */
+    float sum_returns        = __bfloat162float(ps[18]);  /* Kelly: cumulative net returns (for mu) */
+    float sum_sq_returns     = __bfloat162float(ps[19]);  /* Kelly: cumulative squared returns (for sigma^2) */
 
     /* Pre-trade capital floor: skip trade execution on blown accounts.
      * When the floor triggers, write done=1 AND reset the portfolio to
@@ -616,28 +594,28 @@ extern "C" __global__ void experience_env_step(
      * the reset, the blown portfolio persists and every subsequent step
      * hits the floor again (producing fake 92% MaxDD from stuck episodes). */
     {
-        float portfolio_val = ps[2];
+        float portfolio_val = __bfloat162float(ps[2]);
         if (check_capital_floor(portfolio_val, peak_equity)) {
-            out_rewards[out_off] = -10.0f;
-            out_dones[out_off] = 1.0f;
+            out_rewards[out_off] = bf16(-10.0f);
+            out_dones[out_off] = bf16_one();
             /* Reset portfolio for next episode (fresh capital) */
             float init_cap = peak_equity;  /* use peak as initial for next episode */
-            ps[0]  = 0.0f;       /* position = flat */
-            ps[1]  = init_cap;   /* cash = initial_capital */
-            ps[2]  = init_cap;   /* portfolio_value */
-            ps[7]  = init_cap;   /* peak_equity */
-            ps[8]  = 0.0f;       /* flat_counter */
-            ps[9]  = init_cap;   /* prev_equity */
-            ps[10] = 0.0f;       /* hold_time */
-            ps[11] = 0.0f;       /* realized_pnl */
-            ps[12] = 0.0f;       /* entry_price */
-            ps[13] = 0.0f;       /* trade_start_pnl */
-            ps[14] = 0.0f;       /* win_count (reset Kelly) */
-            ps[15] = 0.0f;       /* loss_count */
-            ps[16] = 0.0f;       /* sum_wins */
-            ps[17] = 0.0f;       /* sum_losses */
-            ps[18] = 0.0f;       /* sum_returns */
-            ps[19] = 0.0f;       /* sum_sq_returns */
+            ps[0]  = bf16_zero();                   /* position = flat */
+            ps[1]  = __float2bfloat16(init_cap);    /* cash = initial_capital */
+            ps[2]  = __float2bfloat16(init_cap);    /* portfolio_value */
+            ps[7]  = __float2bfloat16(init_cap);    /* peak_equity */
+            ps[8]  = bf16_zero();                   /* flat_counter */
+            ps[9]  = __float2bfloat16(init_cap);    /* prev_equity */
+            ps[10] = bf16_zero();                   /* hold_time */
+            ps[11] = bf16_zero();                   /* realized_pnl */
+            ps[12] = bf16_zero();                   /* entry_price */
+            ps[13] = bf16_zero();                   /* trade_start_pnl */
+            ps[14] = bf16_zero();                   /* win_count (reset Kelly) */
+            ps[15] = bf16_zero();                   /* loss_count */
+            ps[16] = bf16_zero();                   /* sum_wins */
+            ps[17] = bf16_zero();                   /* sum_losses */
+            ps[18] = bf16_zero();                   /* sum_returns */
+            ps[19] = bf16_zero();                   /* sum_sq_returns */
             current_timesteps[i] = 0;  /* reset episode timer */
             return;
         }
@@ -661,7 +639,7 @@ extern "C" __global__ void experience_env_step(
 
     /* CVaR position scaling: scale down in high-risk quantiles */
     if (cvar_scales != NULL) {
-        float cvar_scale = bf16f(cvar_scales[i]);
+        float cvar_scale = __bfloat162float(cvar_scales[i]);
         if (cvar_scale > 0.0f && cvar_scale < 1.0f) {
             target_position *= cvar_scale;
         }
@@ -669,7 +647,7 @@ extern "C" __global__ void experience_env_step(
 
     /* Q-gap conviction: scale position by trading conviction */
     if (q_gaps != NULL) {
-        float q_gap = bf16f(q_gaps[i]);
+        float q_gap = __bfloat162float(q_gaps[i]);
         float conviction = fminf(fmaxf(q_gap * 0.5f, 0.25f), 1.0f);
         target_position *= conviction;
     }
@@ -679,8 +657,8 @@ extern "C" __global__ void experience_env_step(
     {
         float total_trades = win_count + loss_count;
         if (total_trades >= 20.0f) {
-            float sum_wins_val  = ps[16];
-            float sum_losses_val = ps[17];
+            float sum_wins_val  = __bfloat162float(ps[16]);
+            float sum_losses_val = __bfloat162float(ps[17]);
         float avg_win  = sum_wins_val / fmaxf(win_count, 1.0f);
         float avg_loss = sum_losses_val / fmaxf(loss_count, 1.0f);
         float win_rate = win_count / total_trades;
@@ -699,10 +677,10 @@ extern "C" __global__ void experience_env_step(
     /* Margin-aware position cap (shared: trade_physics.cuh).
      * Prevents overleveraging when equity is depleted — a depleted account
      * can't hold the same position size as a full account.
-     * margin ≈ 6% of notional (CME ES initial margin ~$15K per contract). */
+     * margin ~ 6% of notional (CME ES initial margin ~$15K per contract). */
     {
         float margin_per_contract = raw_close * contract_multiplier * margin_pct;
-        float portfolio_val = ps[2];
+        float portfolio_val = __bfloat162float(ps[2]);
         target_position = apply_margin_cap(target_position, portfolio_val, margin_per_contract, peak_equity);
     }
 
@@ -718,7 +696,7 @@ extern "C" __global__ void experience_env_step(
      * proxy for reward normalization — ATR(14) at feature[9] handles that. */
     float cusum_raw = 0.0f;
     if (features != NULL && bar_idx < total_bars && market_dim > 41) {
-        cusum_raw = bf16f(features[(long long)bar_idx * market_dim + 41]);
+        cusum_raw = __bfloat162float(features[(long long)bar_idx * market_dim + 41]);
     }
     float spread_scale = cusum_raw / 0.5f;
     spread_scale = (spread_scale < 0.5f) ? 0.5f : ((spread_scale > 2.0f) ? 2.0f : spread_scale);
@@ -727,15 +705,16 @@ extern "C" __global__ void experience_env_step(
     float pre_trade_position = position;
 
     /* Sign of position BEFORE trade: -1 (short), 0 (flat), +1 (long) */
-    int prev_sign = (ps[0] > 0.001f) ? 1 : ((ps[0] < -0.001f) ? -1 : 0);
+    float ps0_f = __bfloat162float(ps[0]);
+    int prev_sign = (ps0_f > 0.001f) ? 1 : ((ps0_f < -0.001f) ? -1 : 0);
 
-    /* ════════════════════════════════════════════════════════════════════
+    /* ================================================================
      * DYNAMIC TRAILING STOP — regime-adaptive, locks in profits.
      *
      * Runs BEFORE execute_trade (same order as backtest kernel).
      * If triggered, overrides target_position to 0 (force flat).
      * execute_trade then handles the exit normally.
-     * ════════════════════════════════════════════════════════════════════ */
+     * ================================================================ */
     int trail_triggered = 0;
     if (prev_sign != 0 && hold_time > 0.0f) {
         float current_unrealized = pre_trade_position * (raw_close - entry_price);
@@ -754,25 +733,25 @@ extern "C" __global__ void experience_env_step(
         }
     }
 
-    /* ════════════════════════════════════════════════════════════════════
+    /* ================================================================
      * HOLD ENFORCEMENT: block exits and reversals during minimum hold.
      *
      * Runs BEFORE execute_trade (same order as backtest kernel).
      * If hold is violated, override target_position to keep old position.
-     * execute_trade then sees delta≈0 and does nothing.
+     * execute_trade then sees delta~0 and does nothing.
      *
      * Layer 2 safety net — Layer 1 (action masking) prevents most hold
      * violations, but epsilon exploration can still slip through.
-     * ════════════════════════════════════════════════════════════════════ */
+     * ================================================================ */
     int is_last_bar = (bar_idx >= total_bars - 1) ? 1 : 0;
-    float enforced_position = enforce_hold(ps[0], target_position, hold_time, min_hold_bars, is_last_bar);
+    float enforced_position = enforce_hold(ps0_f, target_position, hold_time, min_hold_bars, is_last_bar);
     int hold_violation = (fabsf(enforced_position - target_position) > 0.001f);
 
     if (hold_violation) {
         target_position = enforced_position;  /* keep old position */
 
         /* Fix the stored action to match held exposure (prevent action aliasing) */
-        float prev_exposure_frac = ps[0] / (max_position > 0.0f ? max_position : 1.0f);
+        float prev_exposure_frac = ps0_f / (max_position > 0.0f ? max_position : 1.0f);
         int held_exposure = (int)roundf((prev_exposure_frac + 1.0f) * 0.5f * (float)(b0_size - 1));
         if (held_exposure < 0) held_exposure = 0;
         if (held_exposure >= b0_size) held_exposure = b0_size - 1;
@@ -811,14 +790,14 @@ extern "C" __global__ void experience_env_step(
 
     if (entering_trade) {
         entry_price = raw_close;
-        trade_start_pnl = ps[11];
+        trade_start_pnl = __bfloat162float(ps[11]);
     }
 
     /* On reversal: close old segment, open new one.
      * MUST run AFTER hold enforcement — if hold guard cancelled the reversal,
      * reversing_trade is 0 and this block correctly does nothing. */
     if (reversing_trade) {
-        float closing_pnl = ps[11] + old_pos_pnl - trade_start_pnl;
+        float closing_pnl = __bfloat162float(ps[11]) + old_pos_pnl - trade_start_pnl;
         reversal_return = closing_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
 
         /* Kelly stats for completed segment */
@@ -834,7 +813,7 @@ extern "C" __global__ void experience_env_step(
 
         /* Reset for new segment */
         entry_price = raw_close;
-        trade_start_pnl = ps[11] + old_pos_pnl;
+        trade_start_pnl = __bfloat162float(ps[11]) + old_pos_pnl;
     }
 
     /* Save hold_time BEFORE reset — sparse reward needs the pre-reset value
@@ -852,7 +831,7 @@ extern "C" __global__ void experience_env_step(
         flat_counter = 0.0f;
     }
 
-    /* ════════════════════════════════════════════════════════════════════
+    /* ================================================================
      * REWARD v6: Sparse Trade-Completion Only
      *
      * Validated by ETDQN (Takara et al., 2023) — outperformed standard
@@ -861,8 +840,8 @@ extern "C" __global__ void experience_env_step(
      * Design:
      *   DURING TRADE: reward = 0.0  (ZERO — no noise, no signal)
      *   WHEN FLAT:    reward = 0.0  (no trade, no signal)
-     *   AT EXIT:      reward = SCALE × vol_normalized(trade_return)
-     *   TURNOVER:     reward -= PENALTY × |delta_position|
+     *   AT EXIT:      reward = SCALE * vol_normalized(trade_return)
+     *   TURNOVER:     reward -= PENALTY * |delta_position|
      *
      * Why zero during trade: per-bar ES returns (~0.001) have SNR of 0.001
      * after weighting — mathematically unlearnable. The model trains on
@@ -876,13 +855,13 @@ extern "C" __global__ void experience_env_step(
      * Vol normalization (Zhang 2020): makes rewards comparable across
      * trending vs ranging regimes. The model targets risk-adjusted
      * returns (Sharpe per trade), not raw P&L.
-     * ════════════════════════════════════════════════════════════════════ */
+     * ================================================================ */
 
     float reward = 0.0f;
 
     /* ---- Sparse: trade completion reward (vol-normalized) ---- */
     if (segment_complete && segment_hold_time > 0.0f) {
-        float segment_pnl = ps[11] + raw_pnl - trade_start_pnl;
+        float segment_pnl = __bfloat162float(ps[11]) + raw_pnl - trade_start_pnl;
         float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
 
         /* Kelly statistics for EXIT (reversal Kelly already done above) */
@@ -910,7 +889,7 @@ extern "C" __global__ void experience_env_step(
          * Makes reward comparable across trending and ranging regimes. */
         float atr_norm = 0.0f;
         if (features != NULL && bar_idx < total_bars && market_dim > 9) {
-            atr_norm = bf16f(features[(long long)bar_idx * market_dim + 9]);
+            atr_norm = __bfloat162float(features[(long long)bar_idx * market_dim + 9]);
         }
         float log_atr = atr_norm * 16.0f - 7.0f;
         float atr_pct = expf(log_atr) / fmaxf(raw_close, 1.0f);
@@ -981,30 +960,30 @@ extern "C" __global__ void experience_env_step(
     int done     = (next_bar >= total_bars || check_capital_floor(new_portfolio_value, peak_equity)) ? 1 : 0;
 
     /* ---- Update full portfolio state (PORTFOLIO_STRIDE=20) ---- */
-    ps[0]  = position;
-    ps[1]  = cash;
-    ps[2]  = new_portfolio_value;
+    ps[0]  = __float2bfloat16(position);
+    ps[1]  = __float2bfloat16(cash);
+    ps[2]  = __float2bfloat16(new_portfolio_value);
     /* ps[3:6] reserved — reward v6 does not use DSR/PnL EMA */
-    ps[7]  = peak_equity;
-    ps[8]  = flat_counter;
-    ps[9]  = new_portfolio_value;  /* prev_equity = current equity for next step */
-    ps[10] = hold_time;
+    ps[7]  = __float2bfloat16(peak_equity);
+    ps[8]  = __float2bfloat16(flat_counter);
+    ps[9]  = __float2bfloat16(new_portfolio_value);  /* prev_equity = current equity for next step */
+    ps[10] = __float2bfloat16(hold_time);
     /* realized_pnl accumulates. On reversal bars, use old_pos_pnl
      * (saved before position was overwritten at ps[0] = position).
      * For non-reversal bars, raw_pnl (= current position's PnL) is correct. */
     if (reversing_trade) {
-        ps[11] = ps[11] + old_pos_pnl; /* old position's PnL (saved at line ~700) */
+        ps[11] = __float2bfloat16(__bfloat162float(ps[11]) + old_pos_pnl); /* old position's PnL */
     } else {
-        ps[11] = ps[11] + raw_pnl; /* normal: current position's PnL */
+        ps[11] = __float2bfloat16(__bfloat162float(ps[11]) + raw_pnl); /* normal: current position's PnL */
     }
-    ps[12] = entry_price;          /* preserved across bars of a trade */
-    ps[13] = trade_start_pnl;      /* realized_pnl snapshot at trade entry */
-    ps[14] = win_count;            /* Kelly: accumulated across episode */
-    ps[15] = loss_count;
-    ps[16] = sum_wins;
-    ps[17] = sum_losses;
-    ps[18] = sum_returns;          /* Kelly continuous: Σ returns */
-    ps[19] = sum_sq_returns;       /* Kelly continuous: Σ returns² */
+    ps[12] = __float2bfloat16(entry_price);          /* preserved across bars of a trade */
+    ps[13] = __float2bfloat16(trade_start_pnl);      /* realized_pnl snapshot at trade entry */
+    ps[14] = __float2bfloat16(win_count);            /* Kelly: accumulated across episode */
+    ps[15] = __float2bfloat16(loss_count);
+    ps[16] = __float2bfloat16(sum_wins);
+    ps[17] = __float2bfloat16(sum_losses);
+    ps[18] = __float2bfloat16(sum_returns);          /* Kelly continuous: sum returns */
+    ps[19] = __float2bfloat16(sum_sq_returns);       /* Kelly continuous: sum returns^2 */
 
     /* ---- NO global reward clamp ---- */
     /* Reward v6: sparse trade-completion only signal.
@@ -1015,8 +994,8 @@ extern "C" __global__ void experience_env_step(
     /* Final NaN guard — if reward is NaN/Inf, write 0.0 instead of poisoning
      * the replay buffer. This prevents gradient explosion from propagating. */
     if (isnan(reward) || isinf(reward)) reward = 0.0f;
-    out_rewards[out_off] = reward;
-    out_dones[out_off]   = (float)done;
+    out_rewards[out_off] = __float2bfloat16(reward);
+    out_dones[out_off]   = __float2bfloat16((float)done);
 
     /* ---- Write RAW portfolio return (unshaped) for accurate Sharpe/MaxDD ---- */
     /* True fractional return: (equity_t - equity_{t-1}) / equity_{t-1}.
@@ -1027,7 +1006,7 @@ extern "C" __global__ void experience_env_step(
             ? (new_portfolio_value - prev_equity) / prev_equity
             : 0.0f;
         if (isnan(portfolio_return) || isinf(portfolio_return)) portfolio_return = 0.0f;
-        raw_returns_out[out_off] = portfolio_return;
+        raw_returns_out[out_off] = __float2bfloat16(portfolio_return);
     }
 
     /* ---- Advance timestep counter ---- */
@@ -1038,43 +1017,47 @@ extern "C" __global__ void experience_env_step(
      * Without this, the blown portfolio persists and every subsequent step
      * hits the pre-trade floor check, producing stuck episodes with fake MaxDD. */
     if (done) {
-        float init_cap = peak_equity;  /* use peak as starting capital for next ep */
-        ps[0]  = 0.0f;       /* position = flat */
-        ps[1]  = init_cap;   /* cash */
-        ps[2]  = init_cap;   /* portfolio_value */
-        ps[7]  = init_cap;   /* peak_equity */
-        ps[8]  = 0.0f;       /* flat_counter */
-        ps[9]  = init_cap;   /* prev_equity */
-        ps[10] = 0.0f;       /* hold_time */
-        ps[11] = 0.0f;       /* realized_pnl */
-        ps[12] = 0.0f;       /* entry_price */
-        ps[13] = 0.0f;       /* trade_start_pnl */
-        ps[14] = 0.0f;       /* win/loss/Kelly counters */
-        ps[15] = 0.0f;
-        ps[16] = 0.0f;
-        ps[17] = 0.0f;
-        ps[18] = 0.0f;
-        ps[19] = 0.0f;
+        __nv_bfloat16 init_cap = __float2bfloat16(peak_equity);  /* use peak as starting capital for next ep */
+        ps[0]  = bf16_zero();    /* position = flat */
+        ps[1]  = init_cap;       /* cash */
+        ps[2]  = init_cap;       /* portfolio_value */
+        ps[7]  = init_cap;       /* peak_equity */
+        ps[8]  = bf16_zero();    /* flat_counter */
+        ps[9]  = init_cap;       /* prev_equity */
+        ps[10] = bf16_zero();    /* hold_time */
+        ps[11] = bf16_zero();    /* realized_pnl */
+        ps[12] = bf16_zero();    /* entry_price */
+        ps[13] = bf16_zero();    /* trade_start_pnl */
+        ps[14] = bf16_zero();    /* win/loss/Kelly counters */
+        ps[15] = bf16_zero();
+        ps[16] = bf16_zero();
+        ps[17] = bf16_zero();
+        ps[18] = bf16_zero();
+        ps[19] = bf16_zero();
         current_timesteps[i] = 0;
     }
 }
 
-/* ══════════════════════════════════════════════════════════════════════════
+/* ======================================================================
  * LEGACY: Portfolio simulation kernel (used by GpuPortfolioSimulator)
  *
  * Single-thread sequential processing for hyperopt/backtest paths.
  * NOT part of the timestep-loop experience collection — kept for
  * GpuPortfolioSimulator compatibility in gpu_portfolio.rs.
- * ══════════════════════════════════════════════════════════════════════════ */
+ *
+ * NOTE: portfolio_state, portfolio_out, rewards_out are bf16 on Rust side.
+ * Internal arithmetic uses float accumulators for precision, converts at
+ * boundaries.
+ * ====================================================================== */
 
 /* action_to_exposure and action_to_tx_cost are provided by common_device_functions.cuh */
 
 extern "C" __global__ void portfolio_sim_kernel(
     const __nv_bfloat16* __restrict__ targets,
     const int*   __restrict__ actions,
-    float*       portfolio_state,
-    float*       portfolio_out,
-    float*       rewards_out,
+    __nv_bfloat16* portfolio_state,
+    __nv_bfloat16* portfolio_out,
+    __nv_bfloat16* rewards_out,
     int*         done_out,
     int          batch_start,
     int          batch_size,
@@ -1084,24 +1067,24 @@ extern "C" __global__ void portfolio_sim_kernel(
 ) {
     if (threadIdx.x != 0 || blockIdx.x != 0) return;
 
-    float cash          = portfolio_state[0];
-    float position      = portfolio_state[1];
-    float entry_price   = portfolio_state[2];
-    float initial_cap   = portfolio_state[3];
-    float spread        = portfolio_state[4];
-    float last_price    = portfolio_state[5];
-    float reserve_pct   = portfolio_state[6];
-    float cum_costs     = portfolio_state[7];
+    float cash          = __bfloat162float(portfolio_state[0]);
+    float position      = __bfloat162float(portfolio_state[1]);
+    float entry_price   = __bfloat162float(portfolio_state[2]);
+    float initial_cap   = __bfloat162float(portfolio_state[3]);
+    float spread        = __bfloat162float(portfolio_state[4]);
+    float last_price    = __bfloat162float(portfolio_state[5]);
+    float reserve_pct   = __bfloat162float(portfolio_state[6]);
+    float cum_costs     = __bfloat162float(portfolio_state[7]);
 
     for (int b = 0; b < batch_size; b++) {
         int global_idx = batch_start + b;
         int action_idx = actions[b];
 
         int t_offset = global_idx * 4;
-        float current_close     = bf16f(targets[t_offset + 0]);
-        float next_close        = bf16f(targets[t_offset + 1]);
-        float current_close_raw = bf16f(targets[t_offset + 2]);
-        float next_close_raw    = bf16f(targets[t_offset + 3]);
+        float current_close     = __bfloat162float(targets[t_offset + 0]);
+        float next_close        = __bfloat162float(targets[t_offset + 1]);
+        float current_close_raw = __bfloat162float(targets[t_offset + 2]);
+        float next_close_raw    = __bfloat162float(targets[t_offset + 3]);
 
         float price = (current_close_raw != 0.0f) ? current_close_raw : current_close;
         if (price <= 0.0f) price = 1.0f;
@@ -1176,9 +1159,9 @@ extern "C" __global__ void portfolio_sim_kernel(
         float max_pos_norm = (price > 0.0f) ? initial_cap / price : 1.0f;
         float pos_norm = position / max_pos_norm;
 
-        portfolio_out[b * 3 + 0] = next_norm;
-        portfolio_out[b * 3 + 1] = pos_norm;
-        portfolio_out[b * 3 + 2] = spread;
+        portfolio_out[b * 3 + 0] = __float2bfloat16(next_norm);
+        portfolio_out[b * 3 + 1] = __float2bfloat16(pos_norm);
+        portfolio_out[b * 3 + 2] = __float2bfloat16(spread);
 
         float reward = 0.0f;
         if (current_norm > 0.0f) {
@@ -1190,7 +1173,7 @@ extern "C" __global__ void portfolio_sim_kernel(
             reward -= (abs_pos - 0.8f) * 5.0f * 0.1f;
         }
 
-        rewards_out[b] = reward;
+        rewards_out[b] = __float2bfloat16(reward);
 
         int step = global_idx + 1;
         int time_done = (step % episode_length == 0) ? 1 : 0;
@@ -1205,14 +1188,14 @@ extern "C" __global__ void portfolio_sim_kernel(
         }
     }
 
-    portfolio_state[0] = cash;
-    portfolio_state[1] = position;
-    portfolio_state[2] = entry_price;
-    portfolio_state[3] = initial_cap;
-    portfolio_state[4] = spread;
-    portfolio_state[5] = last_price;
-    portfolio_state[6] = reserve_pct;
-    portfolio_state[7] = cum_costs;
+    portfolio_state[0] = __float2bfloat16(cash);
+    portfolio_state[1] = __float2bfloat16(position);
+    portfolio_state[2] = __float2bfloat16(entry_price);
+    portfolio_state[3] = __float2bfloat16(initial_cap);
+    portfolio_state[4] = __float2bfloat16(spread);
+    portfolio_state[5] = __float2bfloat16(last_price);
+    portfolio_state[6] = __float2bfloat16(reserve_pct);
+    portfolio_state[7] = __float2bfloat16(cum_costs);
 }
 
 /* ================================================================== */
@@ -1232,9 +1215,9 @@ extern "C" __global__ void portfolio_sim_kernel(
  *
  * Grid: ceil(N / 256), Block: 256.  One thread per episode.
  *
- * @param v_logits    [N, NA]                 value head logits
- * @param b_logits    [N, (B0+B1+B2)*NA]      branch advantage logits
- * @param q_values    [N, B0+B1+B2]           output: expected Q per action
+ * @param v_logits    [N, NA]                 value head logits (bf16)
+ * @param b_logits    [N, (B0+B1+B2)*NA]      branch advantage logits (bf16)
+ * @param q_values    [N, B0+B1+B2]           output: expected Q per action (bf16)
  * @param N                                    number of samples
  * @param num_atoms                            C51 atom count (NA)
  * @param b0_size                              exposure branch size
@@ -1246,7 +1229,7 @@ extern "C" __global__ void portfolio_sim_kernel(
 extern "C" __global__ void compute_expected_q(
     const __nv_bfloat16* __restrict__ v_logits,
     const __nv_bfloat16* __restrict__ b_logits,
-    float*                    q_values,
+    __nv_bfloat16*                q_values,
     int N,
     int num_atoms,
     int b0_size,
@@ -1276,28 +1259,29 @@ extern "C" __global__ void compute_expected_q(
 
         /* Combine value + advantage (no mean subtraction for action selection —
          * the argmax is invariant to the shared value baseline). */
-        /* Softmax + expectation over atoms */
+        /* Softmax + expectation over atoms — float accumulation for numerical
+         * stability (softmax is sensitive to precision). */
         float max_logit = -1e30f;
         for (int z = 0; z < num_atoms; z++) {
-            float logit = bf16f(v_row[z]) + bf16f(adv_a[z]);
+            float logit = __bfloat162float(v_row[z]) + __bfloat162float(adv_a[z]);
             if (logit > max_logit) max_logit = logit;
         }
 
         float sum_exp = 0.0f;
         for (int z = 0; z < num_atoms; z++) {
-            float logit = bf16f(v_row[z]) + bf16f(adv_a[z]);
+            float logit = __bfloat162float(v_row[z]) + __bfloat162float(adv_a[z]);
             sum_exp += expf(logit - max_logit);
         }
 
         float expected_q = 0.0f;
         for (int z = 0; z < num_atoms; z++) {
-            float logit = bf16f(v_row[z]) + bf16f(adv_a[z]);
+            float logit = __bfloat162float(v_row[z]) + __bfloat162float(adv_a[z]);
             float prob = expf(logit - max_logit) / sum_exp;
             float z_val = v_min + (float)z * dz;
             expected_q += prob * z_val;
         }
 
-        q_values[(long long)i * total_actions + a] = expected_q;
+        q_values[(long long)i * total_actions + a] = __float2bfloat16(expected_q);
     }
 }
 
@@ -1310,9 +1294,9 @@ extern "C" __global__ void compute_expected_q(
  *
  * For each episode, reads prices from `targets` and ADX from `features` at the
  * current bar position. Computes a simple fast/slow price crossover signal:
- *   - If close > open (bullish bar) AND ADX > 25 → Long (a0 = b0-1)
- *   - If close < open (bearish bar) AND ADX > 25 → Short (a0 = 0)
- *   - Otherwise → no override
+ *   - If close > open (bullish bar) AND ADX > 25 -> Long (a0 = b0-1)
+ *   - If close < open (bearish bar) AND ADX > 25 -> Short (a0 = 0)
+ *   - Otherwise -> no override
  *
  * With probability `expert_ratio`, overrides the Q-network's exposure action.
  * Order and urgency branches are preserved from the Q-network.
@@ -1354,17 +1338,19 @@ extern "C" __global__ void expert_action_override(
     /* Read prices: compute multi-bar momentum as EMA proxy.
      * True EMA requires sequential scan (O(n) per thread). Instead, use
      * 5-bar price momentum + 20-bar momentum divergence as a crossover proxy.
-     * When short-term momentum > long-term momentum AND ADX is strong → trend.
+     * When short-term momentum > long-term momentum AND ADX is strong -> trend.
      * This approximates fast/slow EMA crossover without sequential scan. */
     if (bar < 20) return;  /* Need 20 bars of history */
 
-    float adx = bf16f(features[bar * market_dim + 40]);
-    float cusum = bf16f(features[bar * market_dim + 41]);
+    __nv_bfloat16 adx_bf = features[bar * market_dim + 40];
+    __nv_bfloat16 cusum_bf = features[bar * market_dim + 41];
+    float adx = __bfloat162float(adx_bf);
+    float cusum = __bfloat162float(cusum_bf);
 
     /* 5-bar return (fast proxy) */
-    float close_now = bf16f(targets[bar * 4 + 3]);
-    float close_5   = bf16f(targets[(bar - 5) * 4 + 3]);
-    float close_20  = bf16f(targets[(bar - 20) * 4 + 3]);
+    float close_now = __bfloat162float(targets[bar * 4 + 3]);
+    float close_5   = __bfloat162float(targets[(bar - 5) * 4 + 3]);
+    float close_20  = __bfloat162float(targets[(bar - 20) * 4 + 3]);
     float ret_5  = (close_now - close_5)  / (close_5  + 1e-8f);
     float ret_20 = (close_now - close_20) / (close_20 + 1e-8f);
 
@@ -1387,7 +1373,7 @@ extern "C" __global__ void expert_action_override(
      * portfolio_states[i * PORTFOLIO_STRIDE + 0] = current position (-1 to +1).
      * If expert says Long and we're already Long (position > 0.5), don't waste
      * the override. Same for Short. */
-    float current_pos = bf16f(portfolio_states[i * 20 + 0]);  /* PORTFOLIO_STRIDE=20, pos at idx 0 */
+    float current_pos = __bfloat162float(portfolio_states[i * 20 + 0]);  /* PORTFOLIO_STRIDE=20, pos at idx 0 */
     if (expert_a0 == b0_size - 1 && current_pos > 0.5f) return;  /* Already Long */
     if (expert_a0 == 0 && current_pos < -0.5f) return;           /* Already Short */
 
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 44f46f062..b9bf71c2a 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -3014,8 +3014,8 @@ impl GpuDqnTrainer {
         dtod_copy(next_states_dst, staging_base + byte_offset, next_states_bytes, &self.stream, 1, "upload_scatter")?;
         byte_offset += next_states_bytes as u64;
 
-        // actions: B elements (reinterpreted as f32 in staging, copy raw bytes to i32 buf)
-        let actions_bytes = b * f32_size; // i32 and f32 are both 4 bytes
+        // actions: B elements as i32 (4 bytes each, not bf16)
+        let actions_bytes = b * std::mem::size_of::<i32>();
         let actions_dst = raw_device_ptr_i32(&self.actions_buf, &self.stream);
         dtod_copy(actions_dst, staging_base + byte_offset, actions_bytes, &self.stream, 2, "upload_scatter")?;
         byte_offset += actions_bytes as u64;
diff --git a/crates/ml/src/cuda_pipeline/ppo_experience_kernel.cu b/crates/ml/src/cuda_pipeline/ppo_experience_kernel.cu
index b5c6a6e5e..6268169c2 100644
--- a/crates/ml/src/cuda_pipeline/ppo_experience_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/ppo_experience_kernel.cu
@@ -1,5 +1,5 @@
 /**
- * Zero-Roundtrip PPO Experience Collection Kernel
+ * Zero-Roundtrip PPO Experience Collection Kernel — Native BF16
  *
  * Requires common_device_functions.cuh prepended via NVRTC source concatenation.
  * Launch config: grid=(ceil(N/32),1,1), block=(32,1,1).
@@ -7,6 +7,10 @@
  *
  * Phase A: Forward rollout (L timesteps) — actor, critic, portfolio, rewards
  * Phase B: Backward GAE scan — advantages and returns
+ *
+ * ALL internal computation uses __nv_bfloat16. Host scalar parameters
+ * (gamma, gae_lambda, etc.) arrive as float and are converted once at
+ * kernel entry via bf16().
  */
 
 /* Override NUM_ACTIONS for PPO: use full 45-action factored space */
@@ -28,49 +32,77 @@
 #define MAX_GAE_LEN 500
 
 /* ------------------------------------------------------------------ */
-/*  PPO-Specific Device Functions                                      */
+/*  BF16 matvec overload: bf16 weights, bf16 input, bf16 output        */
 /* ------------------------------------------------------------------ */
 
 /**
- * PPO Actor MLP forward pass.
+ * Matrix-vector multiply: output = W * input + b, with optional LeakyReLU.
+ * Fully native BF16: weights, biases, input, and output are all __nv_bfloat16.
+ * Accumulation in BF16 (matching tensor-core semantics on SM80+).
+ */
+__device__ void matvec_leaky_relu_bf16(
+    const __nv_bfloat16* __restrict__ W,
+    const __nv_bfloat16* __restrict__ b,
+    const __nv_bfloat16* input,
+    __nv_bfloat16* output,
+    int in_dim,
+    int out_dim,
+    int activate
+) {
+    for (int j = 0; j < out_dim; j++) {
+        __nv_bfloat16 acc = b[j];
+        const __nv_bfloat16* row = W + j * in_dim;
+        for (int i = 0; i < in_dim; i++) {
+            acc = acc + row[i] * input[i];
+        }
+        output[j] = activate ? leaky_relu_bf16(acc) : acc;
+    }
+}
+
+/* ------------------------------------------------------------------ */
+/*  PPO-Specific Device Functions (Native BF16)                        */
+/* ------------------------------------------------------------------ */
+
+/**
+ * PPO Actor MLP forward pass — native BF16.
  *
  *   state[54] -> h1[128] (LeakyReLU) -> h2[64] (LeakyReLU) -> logits[45]
  *
- * @param state   Input state vector [STATE_DIM]
+ * @param state   Input state vector [STATE_DIM]  (__nv_bfloat16)
  * @param pw1     Layer 1 weights [ACTOR_H1, STATE_DIM]
  * @param pb1     Layer 1 biases  [ACTOR_H1]
  * @param pw2     Layer 2 weights [ACTOR_H2, ACTOR_H1]
  * @param pb2     Layer 2 biases  [ACTOR_H2]
  * @param pw3     Output weights  [NUM_ACTIONS, ACTOR_H2]
  * @param pb3     Output biases   [NUM_ACTIONS]
- * @param h1      Scratch buffer  [ACTOR_H1]
- * @param h2      Scratch buffer  [ACTOR_H2]
- * @param logits  Output logits   [NUM_ACTIONS]
+ * @param h1      Scratch buffer  [ACTOR_H1]  (__nv_bfloat16)
+ * @param h2      Scratch buffer  [ACTOR_H2]  (__nv_bfloat16)
+ * @param logits  Output logits   [NUM_ACTIONS] (__nv_bfloat16)
  */
 __device__ void ppo_actor_forward(
-    const float* state,
+    const __nv_bfloat16* state,
     const __nv_bfloat16* __restrict__ pw1,   /* [ACTOR_H1, STATE_DIM] */
     const __nv_bfloat16* __restrict__ pb1,   /* [ACTOR_H1]            */
     const __nv_bfloat16* __restrict__ pw2,   /* [ACTOR_H2, ACTOR_H1]  */
     const __nv_bfloat16* __restrict__ pb2,   /* [ACTOR_H2]            */
     const __nv_bfloat16* __restrict__ pw3,   /* [NUM_ACTIONS, ACTOR_H2] */
     const __nv_bfloat16* __restrict__ pb3,   /* [NUM_ACTIONS]         */
-    float* h1,                       /* [ACTOR_H1] scratch    */
-    float* h2,                       /* [ACTOR_H2] scratch    */
-    float* logits                    /* [NUM_ACTIONS] output  */
+    __nv_bfloat16* h1,                       /* [ACTOR_H1] scratch    */
+    __nv_bfloat16* h2,                       /* [ACTOR_H2] scratch    */
+    __nv_bfloat16* logits                    /* [NUM_ACTIONS] output  */
 ) {
     /* Hidden layer 1: state -> h1 with LeakyReLU */
-    matvec_leaky_relu(pw1, pb1, state, h1, STATE_DIM, ACTOR_H1, 1);
+    matvec_leaky_relu_bf16(pw1, pb1, state, h1, STATE_DIM, ACTOR_H1, 1);
 
     /* Hidden layer 2: h1 -> h2 with LeakyReLU */
-    matvec_leaky_relu(pw2, pb2, h1, h2, ACTOR_H1, ACTOR_H2, 1);
+    matvec_leaky_relu_bf16(pw2, pb2, h1, h2, ACTOR_H1, ACTOR_H2, 1);
 
     /* Output layer: h2 -> logits (no activation) */
-    matvec_leaky_relu(pw3, pb3, h2, logits, ACTOR_H2, NUM_ACTIONS, 0);
+    matvec_leaky_relu_bf16(pw3, pb3, h2, logits, ACTOR_H2, NUM_ACTIONS, 0);
 }
 
 /**
- * Stable softmax + categorical sampling.
+ * Stable softmax + categorical sampling — native BF16.
  *
  * 1. Find max logit for numerical stability
  * 2. exp(logit - max) and accumulate sum
@@ -78,44 +110,44 @@ __device__ void ppo_actor_forward(
  * 4. CDF scan with LCG random draw -> action index
  * 5. Compute log(p[action]) for PPO loss
  *
- * @param logits       Input logits [NUM_ACTIONS]
- * @param probs        Scratch + output probabilities [NUM_ACTIONS]
+ * @param logits       Input logits [NUM_ACTIONS] (__nv_bfloat16)
+ * @param probs        Scratch + output probabilities [NUM_ACTIONS] (__nv_bfloat16)
  * @param rng          Pointer to LCG RNG state
  * @param out_action   Output: selected action index
- * @param out_log_prob Output: log probability of selected action
+ * @param out_log_prob Output: log probability of selected action (__nv_bfloat16)
  */
 __device__ void softmax_sample(
-    const float* logits,
-    float* probs,
+    const __nv_bfloat16* logits,
+    __nv_bfloat16* probs,
     unsigned int* rng,
     int* out_action,
-    float* out_log_prob
+    __nv_bfloat16* out_log_prob
 ) {
     /* Step 1: Find max logit for numerical stability */
-    float max_logit = logits[0];
+    __nv_bfloat16 max_logit = logits[0];
     for (int i = 1; i < NUM_ACTIONS; i++) {
-        if (logits[i] > max_logit) max_logit = logits[i];
+        max_logit = bf16_fmax(max_logit, logits[i]);
     }
 
     /* Step 2: exp(logit - max) and accumulate sum */
-    float sum_exp = 0.0f;
+    __nv_bfloat16 sum_exp = bf16_zero();
     for (int i = 0; i < NUM_ACTIONS; i++) {
-        probs[i] = expf(logits[i] - max_logit);
-        sum_exp += probs[i];
+        probs[i] = bf16_exp(logits[i] - max_logit);
+        sum_exp = sum_exp + probs[i];
     }
 
     /* Step 3: Normalize to probabilities */
-    float inv_sum = 1.0f / fmaxf(sum_exp, 1e-8f);
+    __nv_bfloat16 inv_sum = bf16_one() / bf16_fmax(sum_exp, bf16(1e-8f));
     for (int i = 0; i < NUM_ACTIONS; i++) {
-        probs[i] *= inv_sum;
+        probs[i] = probs[i] * inv_sum;
     }
 
     /* Step 4: CDF scan + random draw -> action index */
-    float u = gpu_random(rng);
-    float cdf = 0.0f;
+    __nv_bfloat16 u = bf16(gpu_random(rng));
+    __nv_bfloat16 cdf = bf16_zero();
     int action = NUM_ACTIONS - 1;  /* default to last action */
     for (int i = 0; i < NUM_ACTIONS; i++) {
-        cdf += probs[i];
+        cdf = cdf + probs[i];
         if (u < cdf) {
             action = i;
             break;
@@ -123,28 +155,28 @@ __device__ void softmax_sample(
     }
 
     /* Step 5: log probability for PPO loss */
-    float p = fmaxf(probs[action], 1e-8f);  /* clamp for log safety */
+    __nv_bfloat16 p = bf16_fmax(probs[action], bf16(1e-8f));  /* clamp for log safety */
     *out_action   = action;
-    *out_log_prob = logf(p);
+    *out_log_prob = bf16_log(p);
 }
 
 /**
- * PPO Critic (value network) forward pass — 5-layer deep with ping-pong buffers.
+ * PPO Critic (value network) forward pass — native BF16, 5-layer deep with ping-pong buffers.
  *
  *   state[54] -> 512(LReLU) -> 384(LReLU) -> 256(LReLU) -> 128(LReLU) -> 64(LReLU) -> 1
  *
  * Uses two scratch buffers (scratch_a[512], scratch_b[512]) that alternate
  * between layers to avoid extra memory.
  *
- * @param state      Input state vector [STATE_DIM]
+ * @param state      Input state vector [STATE_DIM] (__nv_bfloat16)
  * @param vw1-vw6    Weight matrices for each layer
  * @param vb1-vb6    Bias vectors for each layer
- * @param scratch_a  Ping-pong scratch buffer A [CRITIC_H1] (512 wide)
- * @param scratch_b  Ping-pong scratch buffer B [CRITIC_H1] (512 wide)
- * @return           Scalar value estimate V(s)
+ * @param scratch_a  Ping-pong scratch buffer A [CRITIC_H1] (512 wide) (__nv_bfloat16)
+ * @param scratch_b  Ping-pong scratch buffer B [CRITIC_H1] (512 wide) (__nv_bfloat16)
+ * @return           Scalar value estimate V(s) (__nv_bfloat16)
  */
-__device__ float ppo_critic_forward(
-    const float* state,
+__device__ __nv_bfloat16 ppo_critic_forward(
+    const __nv_bfloat16* state,
     const __nv_bfloat16* __restrict__ vw1,   /* [CRITIC_H1, STATE_DIM]  = [512, 54]  */
     const __nv_bfloat16* __restrict__ vb1,   /* [CRITIC_H1]             = [512]      */
     const __nv_bfloat16* __restrict__ vw2,   /* [CRITIC_H2, CRITIC_H1]  = [384, 512] */
@@ -157,81 +189,247 @@ __device__ float ppo_critic_forward(
     const __nv_bfloat16* __restrict__ vb5,   /* [CRITIC_H5]             = [64]       */
     const __nv_bfloat16* __restrict__ vw6,   /* [1, CRITIC_H5]          = [1, 64]    */
     const __nv_bfloat16* __restrict__ vb6,   /* [1]                                  */
-    float* scratch_a,                /* [CRITIC_H1] = [512] ping-pong A      */
-    float* scratch_b                 /* [CRITIC_H1] = [512] ping-pong B      */
+    __nv_bfloat16* scratch_a,                /* [CRITIC_H1] = [512] ping-pong A      */
+    __nv_bfloat16* scratch_b                 /* [CRITIC_H1] = [512] ping-pong B      */
 ) {
     /* Layer 1: state[54] -> scratch_a[512] with LeakyReLU */
-    matvec_leaky_relu(vw1, vb1, state, scratch_a, STATE_DIM, CRITIC_H1, 1);
+    matvec_leaky_relu_bf16(vw1, vb1, state, scratch_a, STATE_DIM, CRITIC_H1, 1);
 
     /* Layer 2: scratch_a[512] -> scratch_b[384] with LeakyReLU */
-    matvec_leaky_relu(vw2, vb2, scratch_a, scratch_b, CRITIC_H1, CRITIC_H2, 1);
+    matvec_leaky_relu_bf16(vw2, vb2, scratch_a, scratch_b, CRITIC_H1, CRITIC_H2, 1);
 
     /* Layer 3: scratch_b[384] -> scratch_a[256] with LeakyReLU */
-    matvec_leaky_relu(vw3, vb3, scratch_b, scratch_a, CRITIC_H2, CRITIC_H3, 1);
+    matvec_leaky_relu_bf16(vw3, vb3, scratch_b, scratch_a, CRITIC_H2, CRITIC_H3, 1);
 
     /* Layer 4: scratch_a[256] -> scratch_b[128] with LeakyReLU */
-    matvec_leaky_relu(vw4, vb4, scratch_a, scratch_b, CRITIC_H3, CRITIC_H4, 1);
+    matvec_leaky_relu_bf16(vw4, vb4, scratch_a, scratch_b, CRITIC_H3, CRITIC_H4, 1);
 
     /* Layer 5: scratch_b[128] -> scratch_a[64] with LeakyReLU */
-    matvec_leaky_relu(vw5, vb5, scratch_b, scratch_a, CRITIC_H4, CRITIC_H5, 1);
+    matvec_leaky_relu_bf16(vw5, vb5, scratch_b, scratch_a, CRITIC_H4, CRITIC_H5, 1);
 
     /* Output layer: scratch_a[64] -> scalar (no activation) */
-    float value = (float)vb6[0];
+    __nv_bfloat16 value = vb6[0];
     for (int i = 0; i < CRITIC_H5; i++) {
-        value += (float)vw6[i] * scratch_a[i];
+        value = value + vw6[i] * scratch_a[i];
     }
 
     return value;
 }
 
 /**
- * Generalized Advantage Estimation (GAE) backward scan.
+ * Generalized Advantage Estimation (GAE) backward scan — native BF16.
  *
  * Computes advantages and returns by scanning backwards through
  * the collected rollout data:
  *
  *   for t = L-1 down to 0:
- *       delta = rewards[t] + gamma * values[t+1] * (1-dones[t]) - values[t]
- *       gae   = delta + gamma * lambda * (1-dones[t]) * gae
+ *       delta = rewards[t] + gm * values[t+1] * (1-dones[t]) - values[t]
+ *       gae   = delta + gm * lm * (1-dones[t]) * gae
  *       advantages[t] = gae
  *       returns[t]    = gae + values[t]
  *
- * @param rewards     Per-timestep rewards [L]
- * @param values      Per-timestep value estimates [L+1] (values[L] is bootstrap)
- * @param dones       Per-timestep done flags [L] (1.0 = done, 0.0 = not done)
- * @param advantages  Output advantage estimates [L]
- * @param returns     Output return targets [L]
+ * @param rewards     Per-timestep rewards [L] (__nv_bfloat16)
+ * @param values      Per-timestep value estimates [L+1] (__nv_bfloat16, values[L] is bootstrap)
+ * @param dones       Per-timestep done flags [L] (__nv_bfloat16, 1.0 = done, 0.0 = not done)
+ * @param advantages  Output advantage estimates [L] (__nv_bfloat16)
+ * @param returns     Output return targets [L] (__nv_bfloat16)
  * @param L           Number of timesteps
- * @param gamma       Discount factor
- * @param lambda      GAE lambda parameter
+ * @param gm          Discount factor (bf16)
+ * @param lm          GAE lambda parameter (bf16)
  */
 __device__ void compute_gae_backward(
-    const float* rewards,
-    const float* values,
-    const float* dones,
-    float* advantages,
-    float* returns,
+    const __nv_bfloat16* rewards,
+    const __nv_bfloat16* values,
+    const __nv_bfloat16* dones,
+    __nv_bfloat16* advantages,
+    __nv_bfloat16* returns,
     int L,
-    float gamma,
-    float lambda
+    __nv_bfloat16 gm,
+    __nv_bfloat16 lm
 ) {
-    float gae = 0.0f;
+    __nv_bfloat16 gae = bf16_zero();
 
     for (int t = L - 1; t >= 0; t--) {
-        float not_done = 1.0f - dones[t];
-        float delta = rewards[t] + gamma * values[t + 1] * not_done - values[t];
-        gae = delta + gamma * lambda * not_done * gae;
+        __nv_bfloat16 not_done = bf16_one() - dones[t];
+        __nv_bfloat16 delta = rewards[t] + gm * values[t + 1] * not_done - values[t];
+        gae = delta + gm * lm * not_done * gae;
         advantages[t] = gae;
         returns[t]    = gae + values[t];
     }
 }
 
+/**
+ * BF16 curiosity inference: builds input in bf16, calls through bf16 matvec,
+ * returns bf16 MSE clamped to max_reward.
+ */
+__device__ __nv_bfloat16 curiosity_inference_bf16(
+    const __nv_bfloat16* state,
+    const __nv_bfloat16* next_state,
+    int action_idx,
+    const __nv_bfloat16* __restrict__ w_c1,
+    const __nv_bfloat16* __restrict__ b_c1,
+    const __nv_bfloat16* __restrict__ w_c2,
+    const __nv_bfloat16* __restrict__ b_c2,
+    __nv_bfloat16* scratch,
+    __nv_bfloat16 max_reward
+) {
+    __nv_bfloat16 input[CUR_INPUT];
+    for (int i = 0; i < MARKET_DIM; i++) {
+        input[i] = state[i];
+    }
+
+    /* Action to category one-hot */
+    int category;
+    if (action_idx <= 1)      category = 0;  /* Short100/Short50 */
+    else if (action_idx == 2) category = 1;  /* Flat */
+    else                      category = 2;  /* Long50/Long100 */
+
+    input[MARKET_DIM + 0] = (category == 0) ? bf16_one() : bf16_zero();
+    input[MARKET_DIM + 1] = (category == 1) ? bf16_one() : bf16_zero();
+    input[MARKET_DIM + 2] = (category == 2) ? bf16_one() : bf16_zero();
+
+    /* Hidden layer */
+    matvec_leaky_relu_bf16(w_c1, b_c1, input, scratch, CUR_INPUT, CUR_HIDDEN, 1);
+
+    /* Output layer (no activation) */
+    __nv_bfloat16 pred[CUR_OUTPUT];
+    matvec_leaky_relu_bf16(w_c2, b_c2, scratch, pred, CUR_HIDDEN, CUR_OUTPUT, 0);
+
+    /* MSE against actual next_state features (first MARKET_DIM) */
+    __nv_bfloat16 mse = bf16_zero();
+    for (int i = 0; i < CUR_OUTPUT; i++) {
+        __nv_bfloat16 diff = pred[i] - next_state[i];
+        mse = mse + diff * diff;
+    }
+    mse = mse / bf16((float)CUR_OUTPUT);
+
+    return bf16_fmin(mse, max_reward);
+}
+
+/**
+ * BF16 barrier_init: initializes triple-barrier on position open.
+ * barrier_state[5] and barrier_config[3] are both __nv_bfloat16.
+ */
+__device__ __forceinline__ void barrier_init_bf16(
+    __nv_bfloat16* barrier_state,
+    const __nv_bfloat16* __restrict__ barrier_config,
+    __nv_bfloat16 entry_price,
+    int current_step
+) {
+    __nv_bfloat16 profit_mult = barrier_config[0];
+    __nv_bfloat16 loss_mult   = barrier_config[1];
+    int           max_hold    = (int)__bfloat162float(barrier_config[2]);
+
+    barrier_state[0] = entry_price;
+    barrier_state[1] = entry_price * profit_mult;
+    barrier_state[2] = entry_price * loss_mult;
+    barrier_state[3] = bf16((float)(current_step + max_hold));
+    barrier_state[4] = bf16_zero();
+}
+
+/**
+ * BF16 barrier_check: check triple-barrier conditions.
+ * Returns label: +1 = profit, -1 = loss, 0 = pending/expired.
+ */
+__device__ __forceinline__ int barrier_check_bf16(
+    __nv_bfloat16* barrier_state,
+    __nv_bfloat16 price,
+    int current_step,
+    __nv_bfloat16 position
+) {
+    if (barrier_state[0] <= bf16_zero()) return 0;
+
+    __nv_bfloat16 upper  = barrier_state[1];
+    __nv_bfloat16 lower  = barrier_state[2];
+    int           expiry = (int)__bfloat162float(barrier_state[3]);
+    __nv_bfloat16 sign   = (position >= bf16_zero()) ? bf16_one() : bf16(-1.0f);
+
+    int label = 0;
+    if (sign > bf16_zero()) {
+        if (price >= upper) label = 1;
+        else if (price <= lower) label = -1;
+    } else {
+        if (price <= lower) label = 1;
+        else if (price >= upper) label = -1;
+    }
+
+    if (label == 0 && current_step >= expiry) {
+        __nv_bfloat16 entry = barrier_state[0];
+        __nv_bfloat16 pnl   = (price - entry) * sign;
+        label = (pnl > bf16_zero()) ? 1 : -1;
+    }
+
+    if (label != 0) {
+        barrier_state[4] = bf16((float)label);
+    }
+    return label;
+}
+
+/** BF16 barrier reset. */
+__device__ __forceinline__ void barrier_reset_bf16(__nv_bfloat16* barrier_state) {
+    barrier_state[0] = bf16_zero();
+    barrier_state[1] = bf16_zero();
+    barrier_state[2] = bf16_zero();
+    barrier_state[3] = bf16_zero();
+    barrier_state[4] = bf16_zero();
+}
+
+/**
+ * BF16 diversity entropy penalty.
+ * Returns bf16 penalty value. Window/meta are int arrays (unchanged).
+ */
+__device__ __nv_bfloat16 diversity_entropy_bf16(
+    int* diversity_window,
+    int* diversity_meta,
+    int action_idx
+) {
+    int category = action_idx;
+    if (category < 0) category = 0;
+    if (category >= DQN_NUM_ACTIONS) category = DQN_NUM_ACTIONS - 1;
+
+    int pos   = diversity_meta[0];
+    int count = diversity_meta[1];
+
+    diversity_window[pos] = category;
+    diversity_meta[0] = (pos + 1) % DIVERSITY_WINDOW;
+    if (count < DIVERSITY_WINDOW) {
+        count++;
+        diversity_meta[1] = count;
+    }
+
+    if (count < 2) return bf16_zero();
+
+    int counts[DQN_NUM_ACTIONS];
+    for (int i = 0; i < DQN_NUM_ACTIONS; i++) counts[i] = 0;
+    for (int i = 0; i < count; i++) {
+        int c = diversity_window[i];
+        if (c >= 0 && c < DQN_NUM_ACTIONS) counts[c]++;
+    }
+
+    /* Shannon entropy in bf16 */
+    __nv_bfloat16 entropy = bf16_zero();
+    __nv_bfloat16 inv_n = bf16_one() / bf16((float)count);
+    for (int c = 0; c < DQN_NUM_ACTIONS; c++) {
+        if (counts[c] > 0) {
+            __nv_bfloat16 p = bf16((float)counts[c]) * inv_n;
+            /* log2(p) = log(p) / log(2) */
+            entropy = entropy - p * bf16_log(p) / bf16(0.6931472f);
+        }
+    }
+
+    __nv_bfloat16 max_entropy = bf16_log(bf16((float)DQN_NUM_ACTIONS)) / bf16(0.6931472f);
+    if (entropy < max_entropy) {
+        return bf16_zero() - ((max_entropy - entropy) / max_entropy) * bf16(0.1f);
+    }
+    return bf16_zero();
+}
+
 /* ------------------------------------------------------------------ */
 /*  Main Kernel                                                        */
 /* ------------------------------------------------------------------ */
 
 /**
- * Full PPO experience collection kernel.
+ * Full PPO experience collection kernel — native BF16.
  *
  * Each thread runs one independent episode of L timesteps (Phase A),
  * then performs a backward GAE scan (Phase B).
@@ -272,16 +470,16 @@ extern "C" __global__ void ppo_full_experience_kernel(
     const __nv_bfloat16* __restrict__ cur_w2,  /* [CUR_OUTPUT, CUR_HIDDEN] */
     const __nv_bfloat16* __restrict__ cur_b2,  /* [CUR_OUTPUT]             */
 
-    /* ---- Per-episode mutable state arrays ---- */
-    float* portfolio_states,        /* [N, PORTFOLIO_STATE_SIZE] */
-    float* barrier_states,          /* [N, BARRIER_STATE_SIZE]   */
+    /* ---- Per-episode mutable state arrays (BF16) ---- */
+    __nv_bfloat16* portfolio_states,        /* [N, PORTFOLIO_STATE_SIZE] */
+    __nv_bfloat16* barrier_states,          /* [N, BARRIER_STATE_SIZE]   */
     int*   diversity_windows,       /* [N, DIVERSITY_WINDOW]     */
     int*   diversity_metas,         /* [N, 2]                    */
 
     /* ---- Barrier config (shared) ---- */
     const __nv_bfloat16* __restrict__ barrier_config,  /* [3]: profit_mult, loss_mult, max_bars */
 
-    /* ---- Scalar configs ---- */
+    /* ---- Scalar configs (host floats, converted to bf16 at entry) ---- */
     float max_position,
     int   episode_length,
     int   total_bars,
@@ -298,31 +496,41 @@ extern "C" __global__ void ppo_full_experience_kernel(
     /* ---- RNG states [N] ---- */
     unsigned int* rng_states,
 
-    /* ---- Output arrays ---- */
-    float* out_states,        /* [N, L, STATE_DIM]  */
+    /* ---- Output arrays (BF16) ---- */
+    __nv_bfloat16* out_states,        /* [N, L, STATE_DIM]  */
     int*   out_actions,       /* [N, L]             */
-    float* out_log_probs,     /* [N, L]             */
-    float* out_advantages,    /* [N, L]             */
-    float* out_returns,       /* [N, L]             */
+    __nv_bfloat16* out_log_probs,     /* [N, L]             */
+    __nv_bfloat16* out_advantages,    /* [N, L]             */
+    __nv_bfloat16* out_returns,       /* [N, L]             */
     int*   out_dones          /* [N, L]             */
 ) {
     int tid = blockIdx.x * blockDim.x + threadIdx.x;
     if (tid >= N) return;
 
+    /* ---- Convert host scalar params to bf16 once ---- */
+    __nv_bfloat16 bf_max_position      = bf16(max_position);
+    __nv_bfloat16 bf_gamma             = bf16(gamma);
+    __nv_bfloat16 bf_gae_lambda        = bf16(gae_lambda);
+    __nv_bfloat16 bf_curiosity_max_rew = bf16(curiosity_max_reward);
+    __nv_bfloat16 bf_barrier_scale     = bf16(barrier_scale);
+    __nv_bfloat16 bf_diversity_scale   = bf16(diversity_scale);
+    __nv_bfloat16 bf_curiosity_scale   = bf16(curiosity_scale);
+    __nv_bfloat16 bf_risk_weight       = bf16(risk_weight);
+
     /* ---- Load per-thread portfolio state ---- */
     int ps_off = tid * PORTFOLIO_STATE_SIZE;
-    float cash        = portfolio_states[ps_off + 0];
-    float position    = portfolio_states[ps_off + 1];
-    float entry_price = portfolio_states[ps_off + 2];
-    float initial_cap = portfolio_states[ps_off + 3];
-    float spread      = portfolio_states[ps_off + 4];
-    float last_price  = portfolio_states[ps_off + 5];
-    float reserve_pct = portfolio_states[ps_off + 6];
-    float cum_costs   = portfolio_states[ps_off + 7];
+    __nv_bfloat16 cash        = portfolio_states[ps_off + 0];
+    __nv_bfloat16 position    = portfolio_states[ps_off + 1];
+    __nv_bfloat16 entry_price = portfolio_states[ps_off + 2];
+    __nv_bfloat16 initial_cap = portfolio_states[ps_off + 3];
+    __nv_bfloat16 spread      = portfolio_states[ps_off + 4];
+    __nv_bfloat16 last_price  = portfolio_states[ps_off + 5];
+    __nv_bfloat16 reserve_pct = portfolio_states[ps_off + 6];
+    __nv_bfloat16 cum_costs   = portfolio_states[ps_off + 7];
 
     /* ---- Load per-thread barrier state ---- */
     int bs_off = tid * BARRIER_STATE_SIZE;
-    float barrier_st[BARRIER_STATE_SIZE];
+    __nv_bfloat16 barrier_st[BARRIER_STATE_SIZE];
     for (int i = 0; i < BARRIER_STATE_SIZE; i++)
         barrier_st[i] = barrier_states[bs_off + i];
 
@@ -340,29 +548,23 @@ extern "C" __global__ void ppo_full_experience_kernel(
     unsigned int rng = rng_states[tid];
     int ep_start = episode_starts[tid];
 
-    /* ---- Per-thread scratch buffers ---- */
-    float state[STATE_DIM];
-    float actor_h1[ACTOR_H1];
-    float actor_h2[ACTOR_H2];
-    float logits[NUM_ACTIONS];
-    float probs[NUM_ACTIONS];
-    float critic_a[CRITIC_H1];   /* ping-pong buffer A (512 wide) */
-    float critic_b[CRITIC_H1];   /* ping-pong buffer B (512 wide) */
-    float next_state[STATE_DIM];
-    float cur_scratch[CUR_HIDDEN];
+    /* ---- Per-thread scratch buffers (BF16) ---- */
+    __nv_bfloat16 state[STATE_DIM];
+    __nv_bfloat16 actor_h1[ACTOR_H1];
+    __nv_bfloat16 actor_h2[ACTOR_H2];
+    __nv_bfloat16 logits[NUM_ACTIONS];
+    __nv_bfloat16 probs[NUM_ACTIONS];
+    __nv_bfloat16 critic_a[CRITIC_H1];   /* ping-pong buffer A (512 wide) */
+    __nv_bfloat16 critic_b[CRITIC_H1];   /* ping-pong buffer B (512 wide) */
+    __nv_bfloat16 next_state[STATE_DIM];
+    __nv_bfloat16 cur_scratch[CUR_HIDDEN];
 
-    /* BF16->float conversion buffers for functions that expect float* */
-    float barrier_config_f[3];
-    barrier_config_f[0] = (float)barrier_config[0];
-    barrier_config_f[1] = (float)barrier_config[1];
-    barrier_config_f[2] = (float)barrier_config[2];
-
-    /* GAE accumulation arrays */
-    float gae_values[MAX_GAE_LEN + 1];   /* values[L] is bootstrap */
-    float gae_rewards[MAX_GAE_LEN];
-    float gae_dones[MAX_GAE_LEN];
-    float gae_advantages[MAX_GAE_LEN];
-    float gae_returns[MAX_GAE_LEN];
+    /* GAE accumulation arrays (BF16) */
+    __nv_bfloat16 gae_values[MAX_GAE_LEN + 1];   /* values[L] is bootstrap */
+    __nv_bfloat16 gae_rewards[MAX_GAE_LEN];
+    __nv_bfloat16 gae_dones[MAX_GAE_LEN];
+    __nv_bfloat16 gae_advantages[MAX_GAE_LEN];
+    __nv_bfloat16 gae_returns[MAX_GAE_LEN];
 
     int step_in_episode = 0;
     int actual_L = (L <= MAX_GAE_LEN) ? L : MAX_GAE_LEN;
@@ -378,35 +580,35 @@ extern "C" __global__ void ppo_full_experience_kernel(
         /* Handle out-of-data */
         if (global_bar >= total_bars - 1) {
             for (int i = 0; i < STATE_DIM; i++)
-                out_states[out_off * STATE_DIM + i] = 0.0f;
+                out_states[out_off * STATE_DIM + i] = bf16_zero();
             out_actions[out_off]   = 0;
-            out_log_probs[out_off] = 0.0f;
+            out_log_probs[out_off] = bf16_zero();
             out_dones[out_off]     = 1;
-            gae_values[t]  = 0.0f;
-            gae_rewards[t] = 0.0f;
-            gae_dones[t]   = 1.0f;
+            gae_values[t]  = bf16_zero();
+            gae_rewards[t] = bf16_zero();
+            gae_dones[t]   = bf16_one();
             continue;
         }
 
         /* ---- Step 1: Read 51 market features from global memory ---- */
         int mf_off = global_bar * MARKET_DIM;
         for (int i = 0; i < MARKET_DIM; i++)
-            state[i] = (float)market_features[mf_off + i];
+            state[i] = market_features[mf_off + i];
 
         /* ---- Step 2: Compute 3 portfolio features ---- */
         int t_off = global_bar * 4;
-        float current_close     = (float)targets[t_off + 0];
-        float next_close        = (float)targets[t_off + 1];
-        float current_close_raw = (float)targets[t_off + 2];
-        float next_close_raw    = (float)targets[t_off + 3];
+        __nv_bfloat16 current_close     = targets[t_off + 0];
+        __nv_bfloat16 next_close        = targets[t_off + 1];
+        __nv_bfloat16 current_close_raw = targets[t_off + 2];
+        __nv_bfloat16 next_close_raw    = targets[t_off + 3];
 
-        float price = (current_close_raw != 0.0f) ? current_close_raw : current_close;
-        if (price <= 0.0f) price = 1.0f;
+        __nv_bfloat16 price = (current_close_raw != bf16_zero()) ? current_close_raw : current_close;
+        if (price <= bf16_zero()) price = bf16_one();
 
-        float current_value = cash + position * price;
-        float current_norm  = current_value / initial_cap;
-        float max_pos_norm  = (price > 0.0f) ? initial_cap / price : 1.0f;
-        float pos_norm      = position / max_pos_norm;
+        __nv_bfloat16 current_value = cash + position * price;
+        __nv_bfloat16 current_norm  = current_value / initial_cap;
+        __nv_bfloat16 max_pos_norm  = (price > bf16_zero()) ? initial_cap / price : bf16_one();
+        __nv_bfloat16 pos_norm      = position / max_pos_norm;
 
         state[MARKET_DIM + 0] = current_norm;   /* normalized value    */
         state[MARKET_DIM + 1] = pos_norm;        /* normalized position */
@@ -420,11 +622,11 @@ extern "C" __global__ void ppo_full_experience_kernel(
         );
 
         int action_idx;
-        float log_prob;
+        __nv_bfloat16 log_prob;
         softmax_sample(logits, probs, &rng, &action_idx, &log_prob);
 
         /* ---- Step 4: Critic forward -> value estimate ---- */
-        float value = ppo_critic_forward(
+        __nv_bfloat16 value = ppo_critic_forward(
             state,
             vw1, vb1, vw2, vb2, vw3, vb3,
             vw4, vb4, vw5, vb5, vw6, vb6,
@@ -435,133 +637,133 @@ extern "C" __global__ void ppo_full_experience_kernel(
         gae_values[t] = value;
 
         /* ---- Step 5: Portfolio simulation ---- */
-        float target_exposure  = factored_action_to_exposure(action_idx);
-        float target_position  = target_exposure * max_position;
-        float tx_rate          = ppo_action_to_tx_cost(action_idx);
+        __nv_bfloat16 target_exposure  = bf16(factored_action_to_exposure(action_idx));
+        __nv_bfloat16 target_position  = target_exposure * bf_max_position;
+        __nv_bfloat16 tx_rate          = bf16(ppo_action_to_tx_cost(action_idx));
 
         /* Detect reversal (sign change) */
-        int is_reversal = (position > 0.0f && target_position < 0.0f) ||
-                          (position < 0.0f && target_position > 0.0f);
+        int is_reversal = (position > bf16_zero() && target_position < bf16_zero()) ||
+                          (position < bf16_zero() && target_position > bf16_zero());
 
         if (is_reversal) {
             /* Phase 1: Close current position */
-            float close_cash = position * price;
-            float close_cost = fabsf(position) * price * tx_rate;
-            cash += close_cash - close_cost;
-            cum_costs += close_cost;
+            __nv_bfloat16 close_cash = position * price;
+            __nv_bfloat16 close_cost = bf16_fabs(position) * price * tx_rate;
+            cash = cash + close_cash - close_cost;
+            cum_costs = cum_costs + close_cost;
 
             /* Phase 2: Open opposite position */
-            float reserve = (reserve_pct > 0.0f) ? current_value * (reserve_pct / 100.0f) : 0.0f;
-            float affordable = fmaxf(cash - reserve, 0.0f);
-            float max_contracts = (price > 0.0f) ? affordable / (price * (1.0f + tx_rate)) : 0.0f;
-            max_contracts = floorf(max_contracts);
-            float actual = fminf(max_contracts, fabsf(target_position));
+            __nv_bfloat16 reserve = (reserve_pct > bf16_zero()) ? current_value * (reserve_pct / bf16(100.0f)) : bf16_zero();
+            __nv_bfloat16 affordable = bf16_fmax(cash - reserve, bf16_zero());
+            __nv_bfloat16 max_contracts = (price > bf16_zero()) ? affordable / (price * (bf16_one() + tx_rate)) : bf16_zero();
+            max_contracts = bf16_floor(max_contracts);
+            __nv_bfloat16 actual = bf16_fmin(max_contracts, bf16_fabs(target_position));
 
-            if (actual > 0.0f) {
-                float new_pos = (target_position > 0.0f) ? actual : -actual;
-                float open_cost = actual * price * tx_rate;
-                cash -= new_pos * price + open_cost;
-                cum_costs += open_cost;
+            if (actual > bf16_zero()) {
+                __nv_bfloat16 new_pos = (target_position > bf16_zero()) ? actual : bf16_zero() - actual;
+                __nv_bfloat16 open_cost = actual * price * tx_rate;
+                cash = cash - new_pos * price - open_cost;
+                cum_costs = cum_costs + open_cost;
                 position = new_pos;
                 entry_price = price;
             } else {
-                position = 0.0f;
-                entry_price = 0.0f;
+                position = bf16_zero();
+                entry_price = bf16_zero();
             }
         } else {
             /* Non-reversal: adjust position directly */
-            float delta = target_position - position;
-            if (fabsf(delta) > 0.0f) {
-                float trade_cost = fabsf(delta) * price * tx_rate;
-                cum_costs += trade_cost;
-                cash -= trade_cost;
+            __nv_bfloat16 delta = target_position - position;
+            if (bf16_fabs(delta) > bf16_zero()) {
+                __nv_bfloat16 trade_cost = bf16_fabs(delta) * price * tx_rate;
+                cum_costs = cum_costs + trade_cost;
+                cash = cash - trade_cost;
 
                 /* Cash reserve check for buys */
-                if (delta > 0.0f && reserve_pct > 0.0f) {
-                    float pv = cash + position * price;
-                    float reserve = pv * (reserve_pct / 100.0f);
-                    float buy_cost = delta * price;
+                if (delta > bf16_zero() && reserve_pct > bf16_zero()) {
+                    __nv_bfloat16 pv = cash + position * price;
+                    __nv_bfloat16 reserve = pv * (reserve_pct / bf16(100.0f));
+                    __nv_bfloat16 buy_cost = delta * price;
                     if (cash - buy_cost < reserve) {
-                        float affordable = fmaxf(cash - reserve, 0.0f);
-                        delta = fminf(delta, (price > 0.0f) ? floorf(affordable / price) : 0.0f);
+                        __nv_bfloat16 affordable = bf16_fmax(cash - reserve, bf16_zero());
+                        delta = bf16_fmin(delta, (price > bf16_zero()) ? bf16_floor(affordable / price) : bf16_zero());
                     }
                 }
 
-                if (delta > 0.0f) {
+                if (delta > bf16_zero()) {
                     entry_price = price;
-                } else if (target_position == 0.0f) {
-                    entry_price = 0.0f;
+                } else if (target_position == bf16_zero()) {
+                    entry_price = bf16_zero();
                 }
-                cash -= delta * price;
+                cash = cash - delta * price;
                 position = position + delta;
             }
         }
         last_price = price;
 
         /* ---- Step 6: Barrier tracking ---- */
-        float old_barrier_entry = barrier_st[0];
-        if (entry_price > 0.0f && old_barrier_entry <= 0.0f) {
-            barrier_init(barrier_st, barrier_config_f, entry_price, global_bar);
+        __nv_bfloat16 old_barrier_entry = barrier_st[0];
+        if (entry_price > bf16_zero() && old_barrier_entry <= bf16_zero()) {
+            barrier_init_bf16(barrier_st, barrier_config, entry_price, global_bar);
         }
 
-        int barrier_label = barrier_check(barrier_st, price, global_bar, position);
+        int barrier_label = barrier_check_bf16(barrier_st, price, global_bar, position);
         if (barrier_label != 0) {
-            barrier_reset(barrier_st);
+            barrier_reset_bf16(barrier_st);
         }
 
         /* ---- Step 7: Diversity entropy penalty ---- */
-        float div_penalty = diversity_entropy(div_window, div_meta, action_idx);
+        __nv_bfloat16 div_penalty = diversity_entropy_bf16(div_window, div_meta, action_idx);
 
         /* ---- Step 8: Mark-to-market + build next_state for curiosity ---- */
-        float next_price = (next_close_raw != 0.0f) ? next_close_raw : next_close;
-        if (next_price <= 0.0f) next_price = price;
-        float next_value = cash + position * next_price;
-        float next_norm  = next_value / initial_cap;
+        __nv_bfloat16 next_price = (next_close_raw != bf16_zero()) ? next_close_raw : next_close;
+        if (next_price <= bf16_zero()) next_price = price;
+        __nv_bfloat16 next_value = cash + position * next_price;
+        __nv_bfloat16 next_norm  = next_value / initial_cap;
 
         int next_bar = global_bar + 1;
         if (next_bar < total_bars) {
             int nmf_off = next_bar * MARKET_DIM;
             for (int i = 0; i < MARKET_DIM; i++)
-                next_state[i] = (float)market_features[nmf_off + i];
+                next_state[i] = market_features[nmf_off + i];
         } else {
             for (int i = 0; i < MARKET_DIM; i++)
                 next_state[i] = state[i];
         }
-        float next_max_pos_norm = (next_price > 0.0f) ? initial_cap / next_price : 1.0f;
+        __nv_bfloat16 next_max_pos_norm = (next_price > bf16_zero()) ? initial_cap / next_price : bf16_one();
         next_state[MARKET_DIM + 0] = next_norm;
         next_state[MARKET_DIM + 1] = position / next_max_pos_norm;
         next_state[MARKET_DIM + 2] = spread;
 
         /* ---- Step 9: Curiosity inference ---- */
-        float curiosity_reward = 0.0f;
+        __nv_bfloat16 curiosity_reward = bf16_zero();
         if (cur_w1 != 0) {
-            curiosity_reward = curiosity_inference(
+            curiosity_reward = curiosity_inference_bf16(
                 state, next_state, action_idx,
                 cur_w1, cur_b1, cur_w2, cur_b2,
-                cur_scratch, curiosity_max_reward
+                cur_scratch, bf_curiosity_max_rew
             );
         }
 
         /* ---- Step 10: Risk penalty (drawdown) ---- */
-        float pnl_reward = 0.0f;
-        if (current_norm > 0.0f) {
+        __nv_bfloat16 pnl_reward = bf16_zero();
+        if (current_norm > bf16_zero()) {
             pnl_reward = (next_norm - current_norm) / current_norm;
         }
 
-        float abs_pos = fabsf(pos_norm);
-        if (abs_pos > 0.8f) {
-            pnl_reward -= (abs_pos - 0.8f) * 5.0f * risk_weight;
+        __nv_bfloat16 abs_pos = bf16_fabs(pos_norm);
+        if (abs_pos > bf16(0.8f)) {
+            pnl_reward = pnl_reward - (abs_pos - bf16(0.8f)) * bf16(5.0f) * bf_risk_weight;
         }
 
         /* ---- Step 11: Reward combination ---- */
-        float barrier_mult = 1.0f;
+        __nv_bfloat16 barrier_mult = bf16_one();
         if (barrier_label != 0) {
-            barrier_mult = 1.0f + barrier_scale * (float)barrier_label;
+            barrier_mult = bf16_one() + bf_barrier_scale * bf16((float)barrier_label);
         }
 
-        float combined_reward = pnl_reward * barrier_mult
-                              + diversity_scale * div_penalty
-                              + curiosity_scale * curiosity_reward;
+        __nv_bfloat16 combined_reward = pnl_reward * barrier_mult
+                              + bf_diversity_scale * div_penalty
+                              + bf_curiosity_scale * curiosity_reward;
 
         /* ---- Step 12: Episode done check ---- */
         step_in_episode++;
@@ -579,17 +781,17 @@ extern "C" __global__ void ppo_full_experience_kernel(
 
         /* Store for GAE backward scan */
         gae_rewards[t] = combined_reward;
-        gae_dones[t]   = (float)done;
+        gae_dones[t]   = bf16((float)done);
 
         /* ---- Step 14: Episode reset on done ---- */
         if (done) {
             cash = initial_cap;
-            position = 0.0f;
-            entry_price = 0.0f;
-            cum_costs = 0.0f;
-            last_price = 0.0f;
+            position = bf16_zero();
+            entry_price = bf16_zero();
+            cum_costs = bf16_zero();
+            last_price = bf16_zero();
             step_in_episode = 0;
-            barrier_reset(barrier_st);
+            barrier_reset_bf16(barrier_st);
             /* Reset diversity window */
             for (int i = 0; i < DIVERSITY_WINDOW; i++)
                 div_window[i] = 0;
@@ -607,18 +809,18 @@ extern "C" __global__ void ppo_full_experience_kernel(
     int last_out_off = tid * actual_L + last_t;
     if (out_dones[last_out_off] == 1) {
         /* Last step was terminal — bootstrap value is 0 */
-        gae_values[actual_L] = 0.0f;
+        gae_values[actual_L] = bf16_zero();
     } else {
         /* Last step was not terminal — run critic on next_state for bootstrap */
         int last_global_bar = ep_start + last_t;
         int next_bar_boot = last_global_bar + 1;
 
         /* Build bootstrap next_state */
-        float boot_state[STATE_DIM];
+        __nv_bfloat16 boot_state[STATE_DIM];
         if (next_bar_boot < total_bars) {
             int nmf_off = next_bar_boot * MARKET_DIM;
             for (int i = 0; i < MARKET_DIM; i++)
-                boot_state[i] = (float)market_features[nmf_off + i];
+                boot_state[i] = market_features[nmf_off + i];
         } else {
             /* Reuse last state market features from output */
             for (int i = 0; i < MARKET_DIM; i++)
@@ -626,17 +828,17 @@ extern "C" __global__ void ppo_full_experience_kernel(
         }
 
         /* Approximate portfolio features from current thread state */
-        float boot_price_raw = 0.0f;
+        __nv_bfloat16 boot_price_raw = bf16_zero();
         if (next_bar_boot < total_bars) {
             int t_off_boot = next_bar_boot * 4;
-            boot_price_raw = (float)targets[t_off_boot + 2];
-            if (boot_price_raw <= 0.0f) boot_price_raw = (float)targets[t_off_boot + 0];
+            boot_price_raw = targets[t_off_boot + 2];
+            if (boot_price_raw <= bf16_zero()) boot_price_raw = targets[t_off_boot + 0];
         }
-        if (boot_price_raw <= 0.0f) boot_price_raw = 1.0f;
+        if (boot_price_raw <= bf16_zero()) boot_price_raw = bf16_one();
 
-        float boot_value = cash + position * boot_price_raw;
-        float boot_norm  = boot_value / initial_cap;
-        float boot_max_pos = (boot_price_raw > 0.0f) ? initial_cap / boot_price_raw : 1.0f;
+        __nv_bfloat16 boot_value = cash + position * boot_price_raw;
+        __nv_bfloat16 boot_norm  = boot_value / initial_cap;
+        __nv_bfloat16 boot_max_pos = (boot_price_raw > bf16_zero()) ? initial_cap / boot_price_raw : bf16_one();
 
         boot_state[MARKET_DIM + 0] = boot_norm;
         boot_state[MARKET_DIM + 1] = position / boot_max_pos;
@@ -654,7 +856,7 @@ extern "C" __global__ void ppo_full_experience_kernel(
     compute_gae_backward(
         gae_rewards, gae_values, gae_dones,
         gae_advantages, gae_returns,
-        actual_L, gamma, gae_lambda
+        actual_L, bf_gamma, bf_gae_lambda
     );
 
     /* Step 3: Write advantages and returns to output buffers */