gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat(bf16): Phase 1 kernel rewrite — experience, ppo, backtest native BF16

Browse Source 
4 critical kernel files rewritten (623 float→bf16):
- experience_kernels.cu: portfolio_states, out_rewards, out_dones,
  q_values all __nv_bfloat16*. Trade physics interop via scoped
  float temporaries at call boundary.
- ppo_experience_kernel.cu: actor/critic forward, softmax, GAE,
  curiosity, barrier, diversity all native bf16.
- backtest_env_kernel.cu: portfolio state, rewards, returns all bf16.
- backtest_gather_kernel.cu: states_out and all feature locals bf16.

Fix actions staging: i32 actions need size_of::<i32>() not bf16 size.

1722/1722 unit tests pass. Smoke test progresses past experience
collection (was ILLEGAL_ADDRESS, now reaches fused training step).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-28 14:46:36 +01:00
parent d83c1d4c48
commit 727308b891
5 changed files with 881 additions and 654 deletions
Download Patch File Download Diff File Expand all files Collapse all files

175
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
View File

@@ -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

160
crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
View File

@@ -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

580
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

File diff suppressed because it is too large Load Diff

4
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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;

616
crates/ml/src/cuda_pipeline/ppo_experience_kernel.cu
View File

@@ -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 */