harry/ForexBot2
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
1454ea1838b0c4cfdfd6d171854fa29aa02feaad
ForexBot2/backtest/cuda_signal_kernels.py
T
Harwinder Sekhon 1454ea1838 Quant review 2026-09-26: alert cooldown, plain-language alerts, M15 grade-A SELL block, 1R printed targets, panel/breach fixes 
Live signals.db review (2026-09-04 → 09-25, 100 unique setups replayed on the
scans' stored candles): the stop-hunt claim is refuted (71% of stop-outs ran
≥ 1R past the stop, the same as the books' 70–72%); −25.1R of the −37.8R raw
live book were consecutive-bar re-entries the backtester's 5-bar cooldown never
takes; the cooldown view sits at the 3rd–5th percentile of the books' three-week
windows.

Delivery (book-neutral):
- telegram_bot: 5-bar alert cooldown per timeframe (_ALERT_COOLDOWN_BARS =
  backtest.cooldown_bars); setup block prints the trail first as the tested
  exit, TP1/TP2 as guides with their R, "Planned TP1 RR", and one validated-book
  line per setup class (agent_panel_base_rates.json); the informational 100-pip
  TP line is removed (owner request).
- AI block rewritten for a new trader: "AI check: the odds look better than /
  worse than / about usual", the chance "in 100" against the usual odds, one
  reason for and one against, no verdict word or confidence figure; rendered
  whenever the verdict carries a probability. agent_panel: persona prompts ban
  trading jargon (PANEL_PROMPT_VERSION v2.1-2026-09-26); the macro-source
  verdict now carries its persona's p_win (every live v2 row had
  panel_probability NULL). Dashboard and guides mirror the wording.
- telegram_bot: level-breach rescans skip the settlement-pause window (the
  20:45 M15 bar was re-logged at 22:01/22:06/22:11 every evening).

Signal (pre-registered rule-1 campaign, CUDA twin paired with gpu_backtest.py):
- M15 grade-A SELLs blocked (sell_min_quality ["A+","B"], M15 override 6):
  W2 813/+154.15R PF 1.53 vs 931/+146.14R PF 1.43, OOS +68.5 vs +60.67,
  W3 max DD 20.9 vs 19.3R. PASSED a–f.
- Printed target 1 floored at 1R in every generator (signals.tp_print_floor_rr),
  applied AFTER the min-RR and spread gates so no trade changes; mirrored in
  cuda_signal_kernels.py K2. Label-only re-pin of all four timeframes (only
  tp1/rr_planned columns move). A floor applied before the gate cost −24.6R and
  was rejected.
- FAILED: M15 min_rr_ratio 1.0 (trade-identical, −2.0R). Owner's call: M5
  tp_floor_rr 1.0 (+48.9R W2, max DD 10.1 → 15.9R), not shipped.

New pins: M5 W2 e599c8a8…, M15 W2 5b65ecf6…, M15 OOS 09b8a8e7…, H1 W3
a44fd19b…, H4 W3 f3587bae…; twin and reference MD5s equal on every window.
Base-rate table rebuilt from the new pins. Suite 558 passed.

Evidence: backtest/research/rr_and_sell_campaign_2026-09-26.md (plus the
2026-09-18 loss-anatomy and wrong-direction reports it builds on).

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-09-26 21:06:05 -04:00

2944 lines
149 KiB
Python
Raw Blame History

"""
backtest/cuda_signal_kernels.py — the CUDA decision stage of cudaBasedBacktesting.py.
Every kernel here is a line-by-line twin of Python code that already exists in
the repo and stays the reference:
K1 fb2_analyze_tf GPUTechnicalEngine's nine cached-indicator overrides
(backtest/gpu_backtest.py) plus the live engine code
those overrides do not cover (engine.py: divergence,
structure S/R, candlestick, regime, regime-aware RSI,
composite). One thread per bar of one timeframe; the
thread analyses the prefix [0..j].
K2 fb2_aggregate signal_aggregator.analyze_from_dataframes end to end
(MTF merge, penalties, grade, gates 6b-6i, the three
setup generators, XAU_PB, XAU_CP, the stale-veto step)
plus the per-bar rules of gpu_backtest._replay_span.
One thread per primary bar.
K3 fb2_features position_sizing.build_features (65 features), one
thread per accepted trade (--export-features).
K4 fb2_cooldown gpu_backtest._apply_cooldown (single thread).
K5 fb2_prefix_median pandas Series.median of High-Low over every prefix
(the snapshot's spread_estimate, --record only).
K6 fb2_bar_rows replay_record.add_bar_row's 160-column float32 row
(--record-bars).
Text is never built on the device: the kernels emit small variant codes plus
the numbers each template needs, and cudaBasedBacktesting.py applies the same
f-string templates as engine.py / signal_aggregator.py on the host. Because
the doubles are bit-identical the strings are identical.
Rules that keep the twins bit-identical (see the plan and the tests in
test_all.py Module 37): compiled with --fmad=false; doubles everywhere; Python
min()/max() written as ternaries in argument order (never fmin/fmax); every
sum in the Python iteration order; numpy's pairwise summation replicated for
np.mean / np.nanmean; pandas' online rolling mean and ewm algorithms copied
verbatim (the rolling helpers are taken from gpu_kernels._ROLL_SRC at import
time); Python round(x, 2) reproduced exactly via an error-free product.
This module imports no live writer and never imports engine, signal_aggregator,
gpu_backtest or replay_record (guard test): the name lists below are the only
contract between the CUDA source and the host, and the host generates the
#defines from them so a name cannot drift.
"""
from __future__ import annotations
import logging
import math
import os
from typing import Dict, List, Sequence, Tuple
import numpy as np
from backtest import gpu_kernels as _gk
from config import PAIR_INFO, Timeframe
from chart_patterns import PATTERN_DETECTORS
logger = logging.getLogger("cuda_signal_kernels")
BACKEND = _gk.BACKEND
KERNEL_OPTIONS = _gk.KERNEL_OPTIONS
# ═════════════════════════════════════════════════════════════════════
# Compile-time limits (thread-local array sizes). A config beyond them is
# refused by cudaBasedBacktesting.py before any kernel runs.
# ═════════════════════════════════════════════════════════════════════
MAX_CONFIRM = 4 # confirmation timeframes gathered by K2
MAX_DIV_LOOKBACK = 64 # technical.divergence_lookback <= 63
MAX_SR_LOOKBACK = 256 # swing lookback <= 255 (120 on M5 by rule)
MAX_CP_LOOKBACK = 128 # xau_strategy.cp_lookback + 1 <= 128
MAX_WARN = 32 # warnings per bar
FEAT_MAX_P = 500 # position_sizing.MAX_LOOKBACK_BARS
FEAT_CONFIRM_TAIL = 60 # position_sizing.CONFIRM_LOOKBACK_BARS
FEAT_BURST_BARS = 200 # position_sizing.BURST_LOOKBACK_BARS
N_FEATURES = 65
N_BAR_COLS = 160
# ═════════════════════════════════════════════════════════════════════
# Codes shared with the host (and, where noted, with replay_record)
# ═════════════════════════════════════════════════════════════════════
# replay_record.REGIME_CODES / VOL_CODES / QUALITY_CODES / STRATEGY_CODES
REGIME_CODES = {"unknown": 0, "trending_up": 1, "trending_down": 2,
"trending": 3, "transitioning": 4, "ranging": 5}
REGIME_NAMES = {v: k for k, v in REGIME_CODES.items()}
VOL_CODES = {"low": 0, "normal": 1, "high": 2, "extreme": 3}
VOL_NAMES = {v: k for k, v in VOL_CODES.items()}
QUALITY_CODES = {"F": 0, "D": 1, "C": 2, "B": 3, "A": 4, "A+": 5}
QUALITY_NAMES = {v: k for k, v in QUALITY_CODES.items()}
STRATEGY_CODES = {"": 0, "GENERIC": 1, "XAU_ATR": 2, "XAU_PB": 3, "XAU_CP": 4}
STRATEGY_NAMES = {v: k for k, v in STRATEGY_CODES.items()}
SIGNAL_NAMES = {2: "STRONG_BUY", 1: "BUY", 0: "NEUTRAL", -1: "SELL", -2: "STRONG_SELL"}
# Indicator columns of the regime table and of the eff/wscore outputs, in the
# engine's insertion order minus atr (weight 0).
INDICATOR_ORDER = ("ema_trend", "rsi", "bollinger", "adx", "macd", "roc",
"keltner", "divergence", "candlestick")
INDICATOR_LABELS = {"ema_trend": "EMA Trend", "rsi": "RSI", "bollinger": "Bollinger Bands",
"adx": "ADX", "macd": "MACD", "roc": "ROC", "keltner": "Keltner",
"divergence": "Divergence", "candlestick": "Candlestick"}
BASE_WEIGHTS = {"ema_trend": 1.0, "rsi": 1.5, "bollinger": 1.5, "adx": 0.5, "macd": 1.5,
"roc": 1.5, "keltner": 1.5, "divergence": 1.5, "candlestick": 1.0}
# engine.analyze's disabled_indicators names -> bit
DISABLED_BITS = {"atr": 0, "ema_trend": 1, "rsi": 2, "bollinger": 3, "adx": 4, "macd": 5,
"roc": 6, "keltner": 7, "divergence": 8, "structure": 9, "candlestick": 10}
PATTERN_ORDER = tuple(PATTERN_DETECTORS.keys()) # priority order of detect_chart_pattern
# Penalty tags in APPEND order (signal_aggregator._compute_final_signal).
PENALTY_TAGS = ("mtf_disagreement", "counter_trend", "extreme_vol", "ema200_near",
"ema200_far", "exhaustion", "no_pullback", "weak_pullback",
"rsi_extended", "bb_overextended")
PENALTY_AMOUNT = {"mtf_disagreement": 15, "counter_trend": 15, "extreme_vol": 10,
"ema200_near": 5, "ema200_far": 5, "exhaustion": 8, "no_pullback": 8,
"weak_pullback": 5, "rsi_extended": 5, "bb_overextended": 8}
# The order after `penalties.sort(key=amount, reverse=True)` (stable): buckets
# 15, 10, 8, 5 each in append order. This is the decay order.
PENALTY_SORTED = tuple(sorted(PENALTY_TAGS, key=lambda t: -PENALTY_AMOUNT[t]))
# Warning codes (K2 -> host). The host holds the templates.
WARN_NAMES = (
"none", # 0
"conflict", # 1 Timeframes CONFLICT
"counter_trend", # 2
"extreme_vol", # 3
"low_vol", # 4
"ema200_near", # 5 arg ema_dist_pct
"ema200_far", # 6 arg ema_dist_pct
"exhaustion_low", # 7
"exhaustion_high", # 8
"no_pullback", # 9 arg pullback_dist
"weak_pullback", # 10 arg pullback_dist
"rsi_ext_up", # 11 arg rsi
"rsi_ext_down", # 12 arg rsi
"exhaustion_block", # 13 arg count
"session_filter", # 14 arg label hour
"day_filter", # 15 arg label weekday
"exhausted_trend", # 16
"sell_quality", # 17 arg quality code
"mtf_opposition", # 18
"sell_vol", # 19 arg vol code
"regime_blocked", # 20 arg regime code
"sell_regime", # 21 arg regime code
"sell_hour", # 22 arg close hour
"big_candle", # 23 arg ratio (bc_idx in K2 ints)
"m5_score", # 24 arg abs_score
"m5_mtf", # 25
"m5_conviction", # 26
"m5_sell_quality", # 27 arg quality code
"m5_vol", # 28 arg vol code
"m5_sell_rsi", # 29 arg rsi
"spread_gate_atr", # 30 arg tp_distance
"low_atr", # 31 arg atr
"spread_gate_generic", # 32 arg tp1_distance
"rf_unavailable", # 33 regime filter: TF missing
"rf_no_adx", # 34 regime filter: ADX None
"rf_blocked", # 35 arg adx
)
WARN_CODES = {n: i for i, n in enumerate(WARN_NAMES)}
# ── K1 outputs (one column per bar of the timeframe) ─────────────────
K1_INT_FIELDS = (
"valid", "atr_present", "regime", "vol", "composite_signal",
"ema_present", "ema_signal", "ema_align", "ema_max_checks", "ema_cross", "ema_ribbon",
"ema_trend_present", "ema2_desc", "kl_ema_fast_present", "kl_ema_present",
"rsi_present", "rsi_signal", "rsi_desc",
"macd_present", "macd_signal", "macd_desc",
"bb_present", "bb_signal", "bb_desc", "bb_squeeze",
"adx_present", "adx_signal", "adx_strength", "adx_dir", "adx_slope_desc",
"roc_present", "roc_signal", "roc_desc",
"kc_present", "kc_signal", "kc_desc",
"div_present", "div_signal", "div_mask",
"cs_present", "cs_signal", "cs_mask", "cs_bull", "cs_bear",
"sh_mask", "sl_mask", "sh_fallback", "sl_fallback",
)
K1_FLOAT_FIELDS = (
"composite_score", "composite_conf", "atr", "atr_pct", "atr_percentile",
"ema_fast", "ema_medium", "ema_slow", "ema_trend", "ema_value", "ema_conf",
"rsi_value", "rsi_conf",
"macd_value", "macd_conf",
"bb_upper", "bb_lower", "bb_mid", "bb_value", "bb_conf",
"adx_value", "plus_di", "minus_di", "adx_slope", "adx_conf",
"roc_value", "roc_conf",
"kc_upper", "kc_lower", "kc_mid", "kc_value", "kc_conf",
"div_value", "cs_value", "cs_conf",
"sh1", "sh2", "sh3", "sl1", "sl2", "sl3",
"ew_ema", "ws_ema", "ew_rsi", "ws_rsi", "ew_bb", "ws_bb", "ew_adx", "ws_adx",
"ew_macd", "ws_macd", "ew_roc", "ws_roc", "ew_kc", "ws_kc", "ew_div", "ws_div",
"ew_cs", "ws_cs",
)
# ── K2 outputs (one column per primary bar in [first_pos, last_pos)) ──
K2_INT_FIELDS = (
"status", "skip_reason", "mtf_agreement", "final_signal", "quality",
"exhaustion_count", "momentum_fading", "penalty_mask", "strategy", "direction",
"cp_pattern", "cp_accepted", "pb_accepted", "bc_idx", "rf_state", "warn_n",
"n_confirm_present", "strategy_pre_rf",
)
K2_FLOAT_FIELDS = (
"primary_score", "primary_conf", "final_score", "final_conf", "penalty_total",
"setup_entry", "setup_sl", "setup_tp1", "setup_tp2", "rr_rounded",
"sl_dist", "tp1_dist", "tp2_dist", "entry_price", "adj_sl", "adj_tp1", "adj_tp2",
"cp_trigger", "cp_height", "pb_retrace_atr", "bc_ratio", "rf_adx",
)
SKIP_NAMES = {1: "WEEKEND_FILTER", 2: "DAY_FILTER", 3: "SESSION_WINDOW"}
# ── Config packing ────────────────────────────────────────────────────
PARAMS_F = (
"rsi_overbought", "rsi_oversold", "rsi_neutral_upper", "rsi_neutral_lower",
"adx_strong_trend", "adx_very_strong", "adx_slope_emerging_threshold", "roc_threshold",
"min_rr_ratio", "tp_tightening_factor", "round_trip_cost", "min_tp_cost_multiple",
"min_sl_pips", "tp_floor_rr", "tp_print_floor_rr",
"tp_atr_multiplier", "sl_atr_multiplier", "min_atr_threshold", "v2_sl_atr_multiplier",
"m5_min_abs_score", "m5_min_confirmation_score", "m5_sell_min_rsi",
"m5_pullback_min_score", "m5_pullback_adx_min",
"m5_pullback_rsi_buy_lo", "m5_pullback_rsi_buy_hi",
"m5_pullback_rsi_sell_lo", "m5_pullback_rsi_sell_hi",
"m5_pullback_retrace_lo", "m5_pullback_retrace_hi",
"cp_tolerance_atr", "cp_min_height_atr", "cp_max_height_atr", "cp_hs_head_min_atr",
"cp_hs_shoulder_tol_atr", "cp_flag_pole_atr", "cp_flag_retrace_lo", "cp_flag_retrace_hi",
"big_candle_atr_multiple", "big_candle_lookback_hours",
"regime_filter_min_adx", "pip_value", "fb2_ema_lo",
"tp_scale_aplus", "tp_scale_a", "max_penalty",
)
PARAMS_I = (
"ema_trend_period", "ema_mode_2ema", "ema_partial_neutral", "rsi_regime_aware",
"macd_as_pillar", "roc_period", "kc_period", "kc_atr_period", "divergence_lookback",
"swing_sr_lookback", "atr_period", "disabled_mask",
"mtf_agreement_threshold", "xau_enabled", "is_xau", "session_filter",
"xau_blocked_hours_mask", "bt_blocked_hours_mask",
"xau_allowed_weekdays_mask", "bt_allowed_weekdays_mask",
"sell_min_quality_mask", "sell_vol_mask", "blocked_regimes_mask",
"sell_blocked_regimes_mask", "sell_blocked_hours_mask",
"m5_require_mtf", "m5_sell_max_quality_mask", "m5_blocked_vol_mask",
"m5_pullback_enabled", "m5_pullback_touch_bars",
"cp_enabled", "cp_patterns_mask", "cp_lookback", "cp_swing_window", "cp_recency_bars",
"cp_min_sep_bars", "cp_flag_pole_bars", "cp_flag_min_bars", "cp_flag_max_bars",
"quality_filter_mask", "strategies_mask", "primary_is_m5", "primary_bar_secs",
"regime_filter", "n_confirm",
)
# Fields of the four config dataclasses the kernels do NOT read, with the
# reason (checked by a Module 37 test so a new field cannot slip past unread).
CONFIG_FIELDS_HOST_ONLY = {
"TechnicalConfig": {
# precompute-only periods (the cached series carry them)
"ema_fast", "ema_medium", "ema_slow", "rsi_period", "macd_fast", "macd_slow",
"macd_signal", "bb_period", "bb_std", "adx_period", "kc_multiplier",
"adx_slope_lookback",
# packed by name below (listed here because the packer reads them
# under another name)
"disabled_indicators", "ema_mode", "rsi_mode", "regime_weights_override",
"ema_trend",
},
"SignalConfig": {
# trade resolution / Trade stamps (host)
"trail_arm_r", "trail_dist_r", "pb_trail_arm_r", "pb_trail_dist_r",
},
"XauStrategyConfig": {
# host calendar/session flags
"weekend_filter", "allowed_weekdays", "blocked_hours_utc",
# packed as masks / pairs under other names
"sell_min_quality", "sell_vol_regimes", "blocked_regimes", "sell_blocked_regimes",
"sell_blocked_hours_utc", "m5_sell_max_quality", "m5_blocked_vol_regimes",
"m5_pullback_rsi_buy", "m5_pullback_rsi_sell", "m5_pullback_retrace_atr",
"cp_patterns", "cp_flag_retrace", "enabled",
},
"FilterConfig": {
# host calendar/session flags / packed under other names
"session_start_hour", "session_end_hour", "blocked_hours_utc", "allowed_weekdays",
"regime_filter_timeframe",
},
}
def _mask(values, table) -> int:
m = 0
for v in values or ():
code = table[v]
m |= (1 << int(code))
return m
def _hours_mask(hours) -> int:
m = 0
for h in hours or ():
h = int(h)
if 0 <= h < 24:
m |= (1 << h)
return m
def fb2_ema_lo_from_env() -> float:
"""The FB2_EMA_LO research knob, read once (the CPU path reads it per bar)."""
raw = os.environ.get("FB2_EMA_LO", "0.5")
try:
return float(raw)
except (TypeError, ValueError):
logger.warning("FB2_EMA_LO=%r is not a number — using tuned default 0.5", raw)
return 0.5
def regime_table(platform_cfg) -> np.ndarray:
"""The 6 x 9 regime multiplier table (rows = REGIME_CODES, cols =
INDICATOR_ORDER; the 'unknown' row is all 1.0 like an empty dict)."""
from engine import regime_weight_table # local: keep the module import-light
table = regime_weight_table(platform_cfg.technical)
out = np.ones((6, len(INDICATOR_ORDER)), dtype=np.float64)
for regime, code in REGIME_CODES.items():
weights = table.get(regime)
if not weights:
continue
for col, name in enumerate(INDICATOR_ORDER):
out[code, col] = float(weights.get(name, 1.0))
return out
def pack_params(platform_cfg, bt_cfg) -> Tuple[np.ndarray, np.ndarray, np.ndarray]:
"""(params_f float64[NF], params_i int32[NI], regime_mult float64[6*9])."""
t = platform_cfg.technical
s = platform_cfg.signals
x = platform_cfg.xau_strategy
b = platform_cfg.backtest
pair_info = PAIR_INFO[bt_cfg.pair]
is_xau = pair_info["display"] == "XAU/USD"
f = {
"rsi_overbought": t.rsi_overbought, "rsi_oversold": t.rsi_oversold,
"rsi_neutral_upper": t.rsi_neutral_upper, "rsi_neutral_lower": t.rsi_neutral_lower,
"adx_strong_trend": t.adx_strong_trend, "adx_very_strong": t.adx_very_strong,
"adx_slope_emerging_threshold": t.adx_slope_emerging_threshold,
"roc_threshold": t.roc_threshold,
"min_rr_ratio": s.min_rr_ratio, "tp_tightening_factor": s.tp_tightening_factor,
"round_trip_cost": s.round_trip_cost, "min_tp_cost_multiple": s.min_tp_cost_multiple,
"min_sl_pips": s.min_sl_pips, "tp_floor_rr": s.tp_floor_rr,
"tp_print_floor_rr": max(s.tp_floor_rr, getattr(s, "tp_print_floor_rr", 0.0)),
"tp_atr_multiplier": x.tp_atr_multiplier, "sl_atr_multiplier": x.sl_atr_multiplier,
"min_atr_threshold": x.min_atr_threshold, "v2_sl_atr_multiplier": x.v2_sl_atr_multiplier,
"m5_min_abs_score": x.m5_min_abs_score,
"m5_min_confirmation_score": x.m5_min_confirmation_score,
"m5_sell_min_rsi": x.m5_sell_min_rsi,
"m5_pullback_min_score": x.m5_pullback_min_score,
"m5_pullback_adx_min": x.m5_pullback_adx_min,
"m5_pullback_rsi_buy_lo": x.m5_pullback_rsi_buy[0],
"m5_pullback_rsi_buy_hi": x.m5_pullback_rsi_buy[1],
"m5_pullback_rsi_sell_lo": x.m5_pullback_rsi_sell[0],
"m5_pullback_rsi_sell_hi": x.m5_pullback_rsi_sell[1],
"m5_pullback_retrace_lo": x.m5_pullback_retrace_atr[0],
"m5_pullback_retrace_hi": x.m5_pullback_retrace_atr[1],
"cp_tolerance_atr": x.cp_tolerance_atr, "cp_min_height_atr": x.cp_min_height_atr,
"cp_max_height_atr": x.cp_max_height_atr, "cp_hs_head_min_atr": x.cp_hs_head_min_atr,
"cp_hs_shoulder_tol_atr": x.cp_hs_shoulder_tol_atr, "cp_flag_pole_atr": x.cp_flag_pole_atr,
"cp_flag_retrace_lo": x.cp_flag_retrace[0], "cp_flag_retrace_hi": x.cp_flag_retrace[1],
"big_candle_atr_multiple": x.big_candle_atr_multiple,
"big_candle_lookback_hours": x.big_candle_lookback_hours,
"regime_filter_min_adx": b.regime_filter_min_adx,
"pip_value": pair_info["pip_value"],
"fb2_ema_lo": fb2_ema_lo_from_env(),
"tp_scale_aplus": 2.0, "tp_scale_a": 1.5, "max_penalty": 30.0,
}
from providers.base import BaseDataProvider # CANDLE_SECONDS (the aggregator's clock)
bar_secs = BaseDataProvider.CANDLE_SECONDS[platform_cfg.primary_timeframe]
strategies = list(bt_cfg.strategies_only or [])
if strategies:
smask = 0
for name in strategies:
if name in STRATEGY_CODES and name:
smask |= 1 << STRATEGY_CODES[name]
else:
smask = (1 << 5) - 1
i = {
"ema_trend_period": t.ema_trend,
"ema_mode_2ema": 1 if t.ema_mode == "2ema" else 0,
"ema_partial_neutral": 1 if t.ema_partial_neutral else 0,
"rsi_regime_aware": 1 if t.rsi_mode == "regime_aware" else 0,
"macd_as_pillar": 1 if t.macd_as_pillar else 0,
"roc_period": t.roc_period, "kc_period": t.kc_period, "kc_atr_period": t.kc_atr_period,
"divergence_lookback": t.divergence_lookback, "swing_sr_lookback": t.swing_sr_lookback,
"atr_period": t.atr_period,
"disabled_mask": _mask([d for d in (t.disabled_indicators or []) if d in DISABLED_BITS],
DISABLED_BITS),
"mtf_agreement_threshold": s.mtf_agreement_threshold,
"xau_enabled": 1 if x.enabled else 0,
"is_xau": 1 if is_xau else 0,
"session_filter": 1 if b.session_filter else 0,
"xau_blocked_hours_mask": _hours_mask(x.blocked_hours_utc),
"bt_blocked_hours_mask": _hours_mask(b.blocked_hours_utc),
"xau_allowed_weekdays_mask": _hours_mask(x.allowed_weekdays),
"bt_allowed_weekdays_mask": _hours_mask(b.allowed_weekdays),
"sell_min_quality_mask": _mask(x.sell_min_quality, QUALITY_CODES),
"sell_vol_mask": _mask(x.sell_vol_regimes, VOL_CODES),
"blocked_regimes_mask": _mask(x.blocked_regimes, REGIME_CODES),
"sell_blocked_regimes_mask": _mask(x.sell_blocked_regimes, REGIME_CODES),
"sell_blocked_hours_mask": _hours_mask(x.sell_blocked_hours_utc),
"m5_require_mtf": 1 if x.m5_require_mtf else 0,
"m5_sell_max_quality_mask": _mask(x.m5_sell_max_quality, QUALITY_CODES),
"m5_blocked_vol_mask": _mask(x.m5_blocked_vol_regimes, VOL_CODES),
"m5_pullback_enabled": 1 if x.m5_pullback_enabled else 0,
"m5_pullback_touch_bars": x.m5_pullback_touch_bars,
"cp_enabled": 1 if x.cp_enabled else 0,
"cp_patterns_mask": _mask([p for p in (x.cp_patterns or []) if p in PATTERN_ORDER],
{p: k for k, p in enumerate(PATTERN_ORDER)}),
"cp_lookback": x.cp_lookback, "cp_swing_window": x.cp_swing_window,
"cp_recency_bars": x.cp_recency_bars, "cp_min_sep_bars": x.cp_min_sep_bars,
"cp_flag_pole_bars": x.cp_flag_pole_bars, "cp_flag_min_bars": x.cp_flag_min_bars,
"cp_flag_max_bars": x.cp_flag_max_bars,
"quality_filter_mask": _mask([q for q in bt_cfg.quality_filter if q in QUALITY_CODES],
QUALITY_CODES),
"strategies_mask": smask,
"primary_is_m5": 1 if platform_cfg.primary_timeframe == Timeframe.M5 else 0,
"primary_bar_secs": bar_secs,
"regime_filter": 1 if b.regime_filter else 0,
"n_confirm": len(platform_cfg.confirmation_timeframes),
}
pf = np.array([float(f[name]) for name in PARAMS_F], dtype=np.float64)
pi = np.array([int(i[name]) for name in PARAMS_I], dtype=np.int32)
return pf, pi, regime_table(platform_cfg).ravel()
def _defines() -> str:
lines = ["#define FB2_MAX_CONFIRM %d" % MAX_CONFIRM,
"#define FB2_MAX_DIV_LOOKBACK %d" % MAX_DIV_LOOKBACK,
"#define FB2_MAX_SR_LOOKBACK %d" % MAX_SR_LOOKBACK,
"#define FB2_MAX_CP_LOOKBACK %d" % MAX_CP_LOOKBACK,
"#define FB2_MAX_WARN %d" % MAX_WARN,
"#define FB2_FEAT_MAX_P %d" % FEAT_MAX_P,
"#define FB2_FEAT_CONFIRM_TAIL %d" % FEAT_CONFIRM_TAIL,
"#define FB2_FEAT_BURST_BARS %d" % FEAT_BURST_BARS,
"#define FB2_N_FEATURES %d" % N_FEATURES,
"#define FB2_N_BAR_COLS %d" % N_BAR_COLS]
for k, name in enumerate(PARAMS_F):
lines.append(f"#define PF_{name} {k}")
for k, name in enumerate(PARAMS_I):
lines.append(f"#define PI_{name} {k}")
for k, name in enumerate(K1_INT_FIELDS):
lines.append(f"#define K1I_{name} {k}")
lines.append(f"#define K1I_N {len(K1_INT_FIELDS)}")
for k, name in enumerate(K1_FLOAT_FIELDS):
lines.append(f"#define K1F_{name} {k}")
lines.append(f"#define K1F_N {len(K1_FLOAT_FIELDS)}")
for k, name in enumerate(K2_INT_FIELDS):
lines.append(f"#define K2I_{name} {k}")
lines.append(f"#define K2I_N {len(K2_INT_FIELDS)}")
for k, name in enumerate(K2_FLOAT_FIELDS):
lines.append(f"#define K2F_{name} {k}")
lines.append(f"#define K2F_N {len(K2_FLOAT_FIELDS)}")
for k, name in enumerate(WARN_NAMES):
lines.append(f"#define W_{name} {k}")
for name, bit in DISABLED_BITS.items():
lines.append(f"#define DIS_{name} (1 << {bit})")
for k, name in enumerate(PENALTY_TAGS):
lines.append(f"#define PEN_{name} {k}")
for k, name in enumerate(PATTERN_ORDER):
lines.append(f"#define CP_{name} {k}")
for name, code in REGIME_CODES.items():
lines.append(f"#define REG_{name} {code}")
for name, code in VOL_CODES.items():
lines.append(f"#define VOL_{name} {code}")
for name, code in QUALITY_CODES.items():
lines.append(f"#define Q_{name.replace('+', 'PLUS')} {code}")
for name, code in STRATEGY_CODES.items():
lines.append(f"#define STRAT_{name or 'NONE'} {code}")
return "\n".join(lines) + "\n"
# ═════════════════════════════════════════════════════════════════════
# CUDA sources
# ═════════════════════════════════════════════════════════════════════
# pandas' Kahan/sign-count rolling-mean helpers, taken verbatim from the
# proven kernel module so the two copies cannot diverge (a test pins it).
_ROLL_HELPERS = _gk._ROLL_SRC.split('extern "C" __global__')[0]
_SRC_COMMON = r'''
// Python min()/max() in argument order: keep the first argument unless the
// second is strictly smaller/greater (NaN never wins).
#define PYMIN(a, b) (((b) < (a)) ? (b) : (a))
#define PYMAX(a, b) (((b) > (a)) ? (b) : (a))
#define FB2_NAN nan("")
__device__ __forceinline__ double fb2_finite(double x, double d) {
return isfinite(x) ? x : d;
}
// CPython round(x, ndigits) for ndigits = 2 (scale 100) or 1 (scale 10):
// the exact binary value is rounded half-even to that many decimals and the
// nearest double of the result is returned. fma() gives the exact residual of
// the product, so a tie is recognised exactly; the correctly rounded division
// r/scale is the same double strtod() returns for the decimal string.
__device__ __forceinline__ double fb2_pyround(double x, double scale) {
if (!isfinite(x)) return x;
const double p = x * scale;
const double e = fma(x, scale, -p);
double r = rint(p);
const double d = p - r;
if (fabs(d) == 0.5 && e != 0.0 && ((e > 0.0) == (d > 0.0)))
r += (d > 0.0) ? 1.0 : -1.0;
return r / scale;
}
// numpy pairwise summation (loops_utils.h.src, PW_BLOCKSIZE 128), stride 1.
// numpy's recursion (n2 = n/2 rounded down to a multiple of 8; left + right)
// is emulated with an explicit stack so no device function is recursive.
__device__ __forceinline__ double fb2_np_pairwise_leaf(const double* a, int n) {
if (n < 8) {
double res = 0.0;
for (int i = 0; i < n; ++i) res += a[i];
return res;
}
double r0 = a[0], r1 = a[1], r2 = a[2], r3 = a[3],
r4 = a[4], r5 = a[5], r6 = a[6], r7 = a[7];
int i;
for (i = 8; i < n - (n % 8); i += 8) {
r0 += a[i + 0]; r1 += a[i + 1]; r2 += a[i + 2]; r3 += a[i + 3];
r4 += a[i + 4]; r5 += a[i + 5]; r6 += a[i + 6]; r7 += a[i + 7];
}
double res = ((r0 + r1) + (r2 + r3)) + ((r4 + r5) + (r6 + r7));
for (; i < n; ++i) res += a[i];
return res;
}
__device__ __forceinline__ double fb2_np_pairwise_sum(const double* a, int n) {
int st_start[40], st_len[40], st_phase[40];
double st_left[40];
int sp = 1;
st_start[0] = 0; st_len[0] = n; st_phase[0] = 0; st_left[0] = 0.0;
double ret = 0.0;
while (sp > 0) {
const int i = sp - 1;
if (st_phase[i] == 0) {
if (st_len[i] <= 128) { ret = fb2_np_pairwise_leaf(a + st_start[i], st_len[i]); sp -= 1; continue; }
int n2 = st_len[i] / 2; n2 -= n2 % 8;
st_phase[i] = 1;
st_start[sp] = st_start[i]; st_len[sp] = n2; st_phase[sp] = 0; st_left[sp] = 0.0; sp += 1;
} else if (st_phase[i] == 1) {
st_left[i] = ret;
int n2 = st_len[i] / 2; n2 -= n2 % 8;
st_phase[i] = 2;
st_start[sp] = st_start[i] + n2; st_len[sp] = st_len[i] - n2; st_phase[sp] = 0; st_left[sp] = 0.0; sp += 1;
} else {
ret = st_left[i] + ret;
sp -= 1;
}
}
return ret;
}
// np.nanmean over a window: NaN -> 0 for the pairwise sum, divided by the
// count of finite values (NaN when there are none).
__device__ __forceinline__ double fb2_np_nanmean(const double* a, int n) {
double tmp[128];
int cnt = 0;
for (int i = 0; i < n; ++i) {
const double v = a[i];
if (v == v) { tmp[i] = v; cnt += 1; } else tmp[i] = 0.0;
}
if (cnt == 0) return FB2_NAN;
return fb2_np_pairwise_sum(tmp, n) / (double)cnt;
}
// pandas ewm(alpha, min_periods, adjust=False, ignore_na=False).mean() —
// pandas/_libs/window/aggregations.pyx `ewm`, restarted at x[0].
// Writes the whole series into out[0..n).
__device__ __forceinline__ void fb2_pd_ewm(const double* x, int n, double alpha, int minp, double* out) {
const double com = (1.0 - alpha) / alpha;
const double a = 1.0 / (1.0 + com);
const double old_wt_factor = 1.0 - a;
const double new_wt = a;
if (n <= 0) return;
double w = x[0];
int nobs = (w == w) ? 1 : 0;
out[0] = (nobs >= minp) ? w : FB2_NAN;
double old_wt = 1.0;
for (int i = 1; i < n; ++i) {
const double cur = x[i];
const bool is_obs = (cur == cur);
if (is_obs) nobs += 1;
if (w == w) {
old_wt *= old_wt_factor; // ignore_na=False: every step decays
if (is_obs) {
if (w != cur) {
w = old_wt * w + new_wt * cur;
w /= (old_wt + new_wt);
}
old_wt = 1.0;
}
} else if (is_obs) {
w = cur;
}
out[i] = (nobs >= minp) ? w : FB2_NAN;
}
}
// position_sizing.wilder_atr on a tail that starts at bar `base`:
// tr[0] = high-low (the shifted close is NaN on the first row of the tail),
// then the pandas ewm above with alpha = 1/period, min_periods = period.
__device__ __forceinline__ void fb2_wilder_atr_tail(const double* high, const double* low, const double* close,
int base, int n, int period, double* out) {
double tr[FB2_FEAT_MAX_P];
for (int i = 0; i < n; ++i) {
const int k = base + i;
const double tr1 = high[k] - low[k];
if (i == 0) { tr[i] = tr1; continue; }
const double pc = close[k - 1];
const double tr2 = fabs(high[k] - pc);
const double tr3 = fabs(low[k] - pc);
// pandas DataFrame.max(axis=1): NaN -> -inf then the plain maximum
double m = tr1;
if (tr2 > m) m = tr2;
if (tr3 > m) m = tr3;
tr[i] = m;
}
fb2_pd_ewm(tr, n, 1.0 / (double)period, period, out);
}
// pandas Series.rolling(window).mean() restarted on a tail; returns the value
// at the tail's last bar (the online algorithm must run from the tail start).
__device__ __forceinline__ double fb2_roll_mean_last(const double* xs, int n, int window) {
const long long minp = window;
double comp_add = 0, comp_rem = 0, sum_x = 0, prev_value = xs[0];
long long nobs = 0, neg_ct = 0, ncsv = 0;
double rm = FB2_NAN;
for (long long i = 0; i < n; ++i) {
long long st = i - window + 1; if (st < 0) st = 0;
long long en = i + 1;
long long pst = (i - 1) - window + 1; if (pst < 0) pst = 0;
long long pen = i;
if (i == 0) {
comp_add = comp_rem = sum_x = 0; nobs = neg_ct = 0; prev_value = xs[st]; ncsv = 0;
for (long long j = st; j < en; ++j)
fb2_add_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_add, &ncsv, &prev_value);
} else {
for (long long j = pst; j < st; ++j)
fb2_remove_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_rem);
for (long long j = pen; j < en; ++j)
fb2_add_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_add, &ncsv, &prev_value);
}
if (nobs >= minp && nobs > 0) {
rm = sum_x / (double)nobs;
if (ncsv >= nobs) rm = prev_value;
else if (neg_ct == 0 && rm < 0) rm = 0;
else if (neg_ct == nobs && rm > 0) rm = 0;
} else {
rm = FB2_NAN;
}
}
return rm;
}
// engine._find_swing_points (NON-strict, NaN-skipping) on the window
// vals[0..L) — is index i a pivot? window=3.
__device__ __forceinline__ bool fb2_engine_pivot(const double* vals, int L, int i, int window, bool low) {
if (i < window || i >= L - window) return false;
const double c = vals[i];
if (c != c) return false;
int nleft = 0, nright = 0;
bool ok = true;
for (int j = 1; j <= window; ++j) {
const double v = vals[i - j];
if (v != v) continue;
nleft += 1;
if (low ? !(c <= v) : !(c >= v)) ok = false;
}
for (int j = 1; j <= window; ++j) {
if (i + j >= L) continue;
const double v = vals[i + j];
if (v != v) continue;
nright += 1;
if (low ? !(c <= v) : !(c >= v)) ok = false;
}
if (nleft == 0 || nright == 0) return false;
return ok;
}
// chart_patterns._swing_indices (STRICT, any NaN in the window rejects).
__device__ __forceinline__ bool fb2_strict_pivot(const double* vals, int n, int i, int w, bool find_high) {
if (i < w || i >= n - w) return false;
const double c = vals[i];
if (c != c) return false;
double rest_max = -1.0/0.0, rest_min = 1.0/0.0;
for (int k = i - w; k <= i + w; ++k) {
const double v = vals[k];
if (v != v) return false;
if (k == i) continue;
if (v > rest_max) rest_max = v;
if (v < rest_min) rest_min = v;
}
return find_high ? (c > rest_max) : (c < rest_min);
}
'''
_SRC_K1 = r'''
#define K1F(name) OF[(size_t)(K1F_##name) * (size_t)n + (size_t)j]
#define K1I(name) OI[(size_t)(K1I_##name) * (size_t)n + (size_t)j]
extern "C" __global__
void fb2_analyze_tf(const double* __restrict__ open_, const double* __restrict__ high,
const double* __restrict__ low, const double* __restrict__ close,
const double* __restrict__ ema_fast, const double* __restrict__ ema_medium,
const double* __restrict__ ema_slow, const double* __restrict__ ema_trend,
const int has_ema_trend,
const double* __restrict__ rsi, const double* __restrict__ macd_hist,
const double* __restrict__ bb_mid, const double* __restrict__ bb_upper,
const double* __restrict__ bb_lower, const double* __restrict__ bb_band_width,
const double* __restrict__ bb_avg_width, const double* __restrict__ atr,
const double* __restrict__ adx, const double* __restrict__ plus_di,
const double* __restrict__ minus_di, const double* __restrict__ adx_slope,
const double* __restrict__ roc, const double* __restrict__ kc_mid,
const double* __restrict__ kc_upper, const double* __restrict__ kc_lower,
const double* __restrict__ kc_squeeze,
const double* __restrict__ PF, const int* __restrict__ PI,
const double* __restrict__ RM, const int tf_is_5min, const int n,
double* __restrict__ OF, int* __restrict__ OI)
{
const int j = blockIdx.x * blockDim.x + threadIdx.x;
if (j >= n) return;
const int pos = j;
const int len = j + 1;
for (int f = 0; f < K1F_N; ++f) OF[(size_t)f * (size_t)n + (size_t)j] = FB2_NAN;
for (int f = 0; f < K1I_N; ++f) OI[(size_t)f * (size_t)n + (size_t)j] = 0;
K1I(regime) = REG_unknown;
K1I(vol) = VOL_normal;
K1F(composite_score) = 0.0;
K1F(composite_conf) = 0.0;
K1I(composite_signal) = 0;
if (len < 30) { K1I(valid) = 0; return; }
K1I(valid) = 1;
const int dis = PI[PI_disabled_mask];
const double latest = close[pos];
// ── ATR (weight 0; feeds ribbon, candlestick, volatility regime) ──
int atr_present = 0;
double atr_val = 0.0;
if (!(dis & DIS_atr)) {
atr_present = 1;
atr_val = atr[pos];
if (atr_val != atr_val) atr_val = 0.0;
const double atr_pct = (latest != 0.0) ? (atr_val / latest) * 100.0 : 0.0;
double pct;
if (pos >= 49) {
int cnt = 0;
for (int k = pos - 49; k <= pos; ++k) if (atr[k] < atr_val) cnt += 1;
pct = ((double)cnt / 50.0) * 100.0;
} else {
pct = 50.0;
}
int vol;
if (pct > 80.0) vol = VOL_extreme; else if (pct > 60.0) vol = VOL_high;
else if (pct > 30.0) vol = VOL_normal; else vol = VOL_low;
K1I(vol) = vol;
K1F(atr) = atr_val; K1F(atr_pct) = atr_pct; K1F(atr_percentile) = pct;
}
K1I(atr_present) = atr_present;
const double atr_kl = atr_present ? atr_val : 0.0; // key_levels.get("atr", 0)
// ── EMA trend ──
int ema_present = 0, ema_signal = 0;
if (!(dis & DIS_ema_trend)) {
ema_present = 1;
const double m = ema_medium[pos];
const double s = ema_slow[pos];
int t_ok = (has_ema_trend && !(ema_trend[pos] != ema_trend[pos])) ? 1 : 0;
double t = t_ok ? ema_trend[pos] : FB2_NAN;
if (t_ok && pos < PI[PI_ema_trend_period] - 1) { t_ok = 0; t = FB2_NAN; }
double conf, value; int sig, cross = 0;
if (PI[PI_ema_mode_2ema]) {
if (!t_ok) {
sig = 0; conf = 30.0; value = 0.0;
K1I(ema2_desc) = 0;
K1I(kl_ema_present) = 1;
K1F(ema_medium) = m; K1F(ema_slow) = s;
} else {
int d;
if (latest > s && s > t) { sig = 2; conf = 85.0; d = 1; }
else if (latest > s && s < t) { sig = 1; conf = 60.0; d = 2; }
else if (latest < s && s < t) { sig = -2; conf = 85.0; d = 3; }
else if (latest < s && s > t) { sig = -1; conf = 60.0; d = 4; }
else { sig = 0; conf = 40.0; d = 5; }
if (pos >= 9) {
const int lim = PYMIN(11, pos + 1);
for (int lb = 2; lb < lim; ++lb) {
const double ps = ema_slow[pos - lb + 1];
const double pt = ema_trend[pos - lb + 1];
if (ps <= pt && s > t) { cross = 1; conf = PYMIN(conf + 10.0, 95.0); break; }
if (ps >= pt && s < t) { cross = 2; conf = PYMIN(conf + 10.0, 95.0); break; }
}
}
value = (sig == 2) ? 1.0 : (sig == 1) ? 0.5 : (sig == 0) ? 0.0 : (sig == -1) ? -0.5 : -1.0;
K1I(ema2_desc) = d;
K1I(ema_trend_present) = 1; K1F(ema_trend) = t;
K1I(kl_ema_present) = 1;
K1F(ema_medium) = m; K1F(ema_slow) = s;
}
} else {
const double f = ema_fast[pos];
int align = 0;
if (latest > f) align += 1; else if (latest < f) align -= 1;
if (f > m) align += 1; else if (f < m) align -= 1;
if (m > s) align += 1; else if (m < s) align -= 1;
if (t_ok) { if (s > t) align += 1; else if (s < t) align -= 1; }
const int maxc = t_ok ? 4 : 3;
const double normalized = (double)align / (double)maxc;
if (normalized > 0.6) { sig = 2; conf = PYMIN(90.0, 60.0 + normalized * 30.0); }
else if (normalized > 0.2) { sig = 1; conf = 55.0 + normalized * 20.0; }
else if (normalized < -0.6) { sig = -2; conf = PYMIN(90.0, 60.0 + fabs(normalized) * 30.0); }
else if (normalized < -0.2) { sig = -1; conf = 55.0 + fabs(normalized) * 20.0; }
else { sig = 0; conf = 40.0; }
if (PI[PI_ema_partial_neutral] && (sig == 1 || sig == -1)) { sig = 0; conf = 40.0; }
if (pos >= 2) {
const double pf = ema_fast[pos - 1], pm = ema_medium[pos - 1];
if (pf <= pm && f > m) { cross = 1; conf = PYMIN(conf + 10.0, 95.0); }
else if (pf >= pm && f < m) { cross = 2; conf = PYMIN(conf + 10.0, 95.0); }
}
int ribbon = 0;
if (atr_kl > 0.0) {
const double mx = PYMAX(PYMAX(f, m), s);
const double mn = PYMIN(PYMIN(f, m), s);
const double rw = (mx - mn) / atr_kl;
if (rw < 1.0) ribbon = 1;
if (pos >= 5) {
const double pf = ema_fast[pos - 4], pm = ema_medium[pos - 4], ps = ema_slow[pos - 4];
const double pmx = PYMAX(PYMAX(pf, pm), ps);
const double pmn = PYMIN(PYMIN(pf, pm), ps);
const double pw = (atr_kl > 0.0) ? (pmx - pmn) / atr_kl : 0.0;
if (pw < 1.0 && rw >= 1.0) { ribbon = 2; conf = PYMIN(conf + 10.0, 95.0); }
}
}
value = normalized;
K1I(ema_align) = align; K1I(ema_max_checks) = maxc; K1I(ema_ribbon) = ribbon;
if (t_ok) { K1I(ema_trend_present) = 1; K1F(ema_trend) = t; }
K1I(kl_ema_fast_present) = 1; K1I(kl_ema_present) = 1;
K1F(ema_fast) = f; K1F(ema_medium) = m; K1F(ema_slow) = s;
}
ema_signal = sig;
K1I(ema_present) = 1; K1I(ema_signal) = sig; K1I(ema_cross) = cross;
K1F(ema_value) = value; K1F(ema_conf) = conf;
}
// ── RSI (classic pass) ──
int rsi_present = 0, rsi_signal = 0; double rsi_val = 50.0, rsi_conf = 30.0; int rsi_desc = 4;
if (!(dis & DIS_rsi)) {
rsi_present = 1;
rsi_val = rsi[pos];
if (rsi_val != rsi_val) rsi_val = 50.0;
if (rsi_val >= PF[PF_rsi_overbought]) { rsi_signal = -1; rsi_conf = PYMIN(80.0, 50.0 + (rsi_val - PF[PF_rsi_overbought])); rsi_desc = 0; }
else if (rsi_val <= PF[PF_rsi_oversold]) { rsi_signal = 1; rsi_conf = PYMIN(80.0, 50.0 + (PF[PF_rsi_oversold] - rsi_val)); rsi_desc = 1; }
else if (rsi_val >= PF[PF_rsi_neutral_upper]) { rsi_signal = 1; rsi_conf = 45.0; rsi_desc = 2; }
else if (rsi_val <= PF[PF_rsi_neutral_lower]) { rsi_signal = -1; rsi_conf = 45.0; rsi_desc = 3; }
else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 4; }
}
// ── Bollinger ──
int bb_present = 0, bb_signal = 0; double bb_value = FB2_NAN, bb_conf = 0.0;
if (!(dis & DIS_bollinger)) {
const double up = bb_upper[pos], lo = bb_lower[pos], sma = bb_mid[pos];
if (!(up != up || lo != lo || sma != sma || sma == 0.0)) {
bb_present = 1;
const double range = up - lo;
const double bbpos = (range != 0.0) ? (latest - lo) / range : 0.5;
const double bw = bb_band_width[pos];
int sq = 0;
if (pos >= 19) {
const double aw = bb_avg_width[pos];
if (!(aw != aw) && aw > 0.0) sq = (bw < aw * 0.75) ? 1 : 0;
}
int d;
if (bbpos > 1.0) { bb_signal = -1; bb_conf = 65.0; d = 0; }
else if (bbpos < 0.0) { bb_signal = 1; bb_conf = 65.0; d = 1; }
else if (bbpos > 0.8) { bb_signal = -1; bb_conf = 45.0; d = 2; }
else if (bbpos < 0.2) { bb_signal = 1; bb_conf = 45.0; d = 3; }
else { bb_signal = 0; bb_conf = 30.0; d = 4; }
if (sq) bb_conf = PYMIN(bb_conf + 10.0, 85.0);
bb_value = bbpos;
K1I(bb_present) = 1; K1I(bb_signal) = bb_signal; K1I(bb_desc) = d; K1I(bb_squeeze) = sq;
K1F(bb_upper) = up; K1F(bb_lower) = lo; K1F(bb_mid) = sma;
K1F(bb_value) = bbpos; K1F(bb_conf) = bb_conf;
}
}
// ── ADX ──
int adx_present = 0, adx_signal = 0; double adx_val = 0.0, adx_conf = 0.0, slope = 0.0;
if (!(dis & DIS_adx)) {
adx_present = 1;
adx_val = adx[pos]; if (adx_val != adx_val) adx_val = 0.0;
double pdi = plus_di[pos], mdi = minus_di[pos];
if (pdi != pdi) pdi = 0.0;
if (mdi != mdi) mdi = 0.0;
int strength;
if (adx_val >= PF[PF_adx_very_strong]) { strength = 0; adx_conf = 80.0; }
else if (adx_val >= PF[PF_adx_strong_trend]) { strength = 1; adx_conf = 65.0; }
else { strength = 2; adx_conf = 40.0; }
int dir;
if (pdi > mdi) { adx_signal = (adx_val >= PF[PF_adx_strong_trend]) ? 1 : 0; dir = 0; }
else { adx_signal = (adx_val >= PF[PF_adx_strong_trend]) ? -1 : 0; dir = 1; }
slope = adx_slope[pos]; if (slope != slope) slope = 0.0;
int sd = 0;
const double thr = PF[PF_adx_slope_emerging_threshold];
if (adx_val < PF[PF_adx_strong_trend] && slope > thr) { sd = 1; adx_conf = PYMIN(adx_conf + 15.0, 85.0); }
else if (adx_val >= PF[PF_adx_strong_trend] && slope > 0.0) { sd = 2; adx_conf = PYMIN(adx_conf + 5.0, 85.0); }
else if (adx_val >= PF[PF_adx_strong_trend] && slope < -thr) { sd = 3; adx_conf = PYMAX(adx_conf - 10.0, 30.0); }
K1I(adx_present) = 1; K1I(adx_signal) = adx_signal; K1I(adx_strength) = strength;
K1I(adx_dir) = dir; K1I(adx_slope_desc) = sd;
K1F(adx_value) = adx_val; K1F(plus_di) = pdi; K1F(minus_di) = mdi;
K1F(adx_slope) = slope; K1F(adx_conf) = adx_conf;
}
// ── MACD (histogram always computed; pillar honours the disable) ──
int macd_present = 0, macd_signal = 0; double macd_conf = 0.0;
if (PI[PI_macd_as_pillar] && pos >= 2) {
const double hist_now = macd_hist[pos], hist_prev = macd_hist[pos - 1];
if (!(hist_now != hist_now || hist_prev != hist_prev)) {
const bool crossed_up = (hist_prev <= 0.0 && 0.0 < hist_now);
const bool crossed_down = (hist_prev >= 0.0 && 0.0 > hist_now);
const bool accel_up = (hist_now > 0.0 && hist_now > hist_prev && hist_prev > 0.0);
const bool accel_down = (hist_now < 0.0 && hist_now < hist_prev && hist_prev < 0.0);
const bool fading_up = (hist_now > 0.0 && hist_now < hist_prev);
const bool fading_down = (hist_now < 0.0 && hist_now > hist_prev);
double hist_pct = 1.0;
if (pos >= 19) {
double win[20];
for (int k = 0; k < 20; ++k) win[k] = fabs(macd_hist[pos - 19 + k]);
const double mean_abs = fb2_np_nanmean(win, 20);
hist_pct = (mean_abs > 0.0) ? fabs(hist_now) / mean_abs : 1.0;
}
int d;
if (crossed_up || accel_up) { macd_signal = accel_up ? 2 : 1; macd_conf = PYMIN(80.0, 55.0 + hist_pct * 10.0); d = accel_up ? 0 : 1; }
else if (crossed_down || accel_down) { macd_signal = accel_down ? -2 : -1; macd_conf = PYMIN(80.0, 55.0 + hist_pct * 10.0); d = accel_down ? 2 : 3; }
else if (fading_up) { macd_signal = 0; macd_conf = 35.0; d = 4; }
else if (fading_down) { macd_signal = 0; macd_conf = 35.0; d = 5; }
else { macd_signal = 0; macd_conf = 30.0; d = 6; }
if (!(dis & DIS_macd)) {
macd_present = 1;
K1I(macd_present) = 1; K1I(macd_signal) = macd_signal; K1I(macd_desc) = d;
K1F(macd_value) = hist_now; K1F(macd_conf) = macd_conf;
}
}
}
// ── ROC ──
int roc_present = 0, roc_signal = 0; double roc_val = 0.0, roc_conf = 0.0;
if (!(dis & DIS_roc) && pos >= PI[PI_roc_period] + 4) {
roc_val = roc[pos];
const double roc_prev = (pos > 0) ? roc[pos - 1] : FB2_NAN;
if (!(roc_val != roc_val || roc_prev != roc_prev)) {
roc_present = 1;
const bool rising = roc_val > roc_prev;
const double roc_abs = fabs(roc_val);
const double threshold = PF[PF_roc_threshold];
int d;
if (roc_val > 0.0 && rising) {
if (roc_abs > threshold * 2.0) { roc_signal = 2; roc_conf = PYMIN(80.0, 60.0 + roc_abs * 5.0); d = 0; }
else { roc_signal = 1; roc_conf = PYMIN(70.0, 50.0 + roc_abs * 10.0); d = 1; }
} else if (roc_val < 0.0 && !rising) {
if (roc_abs > threshold * 2.0) { roc_signal = -2; roc_conf = PYMIN(80.0, 60.0 + roc_abs * 5.0); d = 2; }
else { roc_signal = -1; roc_conf = PYMIN(70.0, 50.0 + roc_abs * 10.0); d = 3; }
} else { roc_signal = 0; roc_conf = 30.0; d = 4; }
K1I(roc_present) = 1; K1I(roc_signal) = roc_signal; K1I(roc_desc) = d;
K1F(roc_value) = roc_val; K1F(roc_conf) = roc_conf;
}
}
// ── Keltner ──
int kc_present = 0, kc_signal = 0; double kc_conf = 0.0, kc_value = 0.0;
if (!(dis & DIS_keltner)) {
const int guard = PYMAX(PI[PI_kc_period], PI[PI_kc_atr_period]) + 4;
if (pos >= guard) {
const double up = kc_upper[pos], lo = kc_lower[pos], mid = kc_mid[pos];
if (!(up != up || lo != lo)) {
kc_present = 1;
const bool is_squeeze = (kc_squeeze[pos] != 0.0);
bool released = false;
if (pos >= 1) released = (kc_squeeze[pos - 1] != 0.0) && !is_squeeze;
int macd_direction = 0;
if (!(macd_hist[pos] != macd_hist[pos])) macd_direction = (macd_hist[pos] > 0.0) ? 1 : -1;
int d;
if (released) {
if (macd_direction > 0) { kc_signal = 2; d = 0; }
else if (macd_direction < 0) { kc_signal = -2; d = 1; }
else { kc_signal = 0; d = 2; }
kc_conf = 80.0;
} else if (is_squeeze) { kc_signal = 0; kc_conf = 60.0; d = 3; }
else if (latest > up) { kc_signal = 1; kc_conf = 65.0; d = 4; }
else if (latest < lo) { kc_signal = -1; kc_conf = 65.0; d = 5; }
else { kc_signal = 0; kc_conf = 30.0; d = 6; }
kc_value = latest - mid;
K1I(kc_present) = 1; K1I(kc_signal) = kc_signal; K1I(kc_desc) = d;
K1F(kc_upper) = up; K1F(kc_lower) = lo; K1F(kc_mid) = mid;
K1F(kc_value) = kc_value; K1F(kc_conf) = kc_conf;
}
}
}
// ── Divergence (engine._detect_divergences, not overridden) ──
int div_present = 0, div_signal = 0, div_mask = 0;
if (!(dis & DIS_divergence)) {
const int lookback = PYMIN(PI[PI_divergence_lookback], len - 1);
if (lookback >= 5) {
const int base = pos + 1 - lookback;
for (int series = 0; series < 2; ++series) {
if (series == 0 && (dis & DIS_rsi)) continue; // _rsi_series is None
const double* ind = (series == 0) ? rsi : macd_hist;
for (int mode = 0; mode < 2; ++mode) { // 0 lows, 1 highs
const bool low = (mode == 0);
int ind_idx[FB2_MAX_DIV_LOOKBACK]; double ind_val[FB2_MAX_DIV_LOOKBACK];
int n_ind = 0;
for (int i = 0; i < lookback; ++i) {
if (fb2_engine_pivot(ind + base, lookback, i, 3, low)) {
ind_idx[n_ind] = i; ind_val[n_ind] = ind[base + i]; n_ind += 1;
}
}
int n_pairs = 0;
double p1v = 0, i1v = 0, p2v = 0, i2v = 0;
for (int i = 0; i < lookback; ++i) {
if (!fb2_engine_pivot(close + base, lookback, i, 3, low)) continue;
if (n_ind == 0) continue;
int best = 0; int bestd = abs(ind_idx[0] - i);
for (int q = 1; q < n_ind; ++q) {
const int dd = abs(ind_idx[q] - i);
if (dd < bestd) { bestd = dd; best = q; }
}
if (bestd <= 2) {
p1v = p2v; i1v = i2v;
p2v = close[base + i]; i2v = ind_val[best];
n_pairs += 1;
}
}
if (n_pairs >= 2) {
if (low) { if (p2v < p1v && i2v > i1v) div_mask |= (series == 0) ? 1 : 4; }
else { if (p2v > p1v && i2v < i1v) div_mask |= (series == 0) ? 2 : 8; }
}
}
}
if (div_mask) {
div_present = 1;
const bool is_bull = (div_mask & 5) != 0;
const bool is_bear = (div_mask & 10) != 0;
div_signal = (is_bull && !is_bear) ? 1 : (is_bear && !is_bull) ? -1 : 0;
const int cnt = ((div_mask & 1) ? 1 : 0) + ((div_mask & 2) ? 1 : 0)
+ ((div_mask & 4) ? 1 : 0) + ((div_mask & 8) ? 1 : 0);
K1I(div_present) = 1; K1I(div_signal) = div_signal; K1I(div_mask) = div_mask;
K1F(div_value) = (double)cnt;
}
}
}
// ── Structure S/R (engine._compute_structure_sr) ──
int sh_mask = 0, sl_mask = 0;
double sh[3] = {FB2_NAN, FB2_NAN, FB2_NAN}, sl[3] = {FB2_NAN, FB2_NAN, FB2_NAN};
if (!(dis & DIS_structure)) {
int base_lb = PI[PI_swing_sr_lookback];
if (tf_is_5min) base_lb = PYMAX(base_lb, 120);
const int lookback = PYMIN(base_lb, len - 1);
if (lookback >= 10) {
const double price = latest;
const int base = pos + 1 - lookback;
int n_above = 0; // highs
double above[3], below_max = FB2_NAN; bool has_below = false;
int n_lbelow = 0; // lows
double lbelow[3], labove_min = FB2_NAN; bool has_labove = false;
for (int i = 0; i < lookback; ++i) {
if (fb2_engine_pivot(high + base, lookback, i, 3, false)) {
const double v = high[base + i];
if (v > price) {
// keep the 3 smallest (ascending)
if (n_above < 3) { above[n_above++] = v; }
else if (v < above[2]) { above[2] = v; }
for (int a = n_above - 1; a > 0; --a) if (above[a] < above[a-1]) { double t = above[a]; above[a] = above[a-1]; above[a-1] = t; }
} else {
if (!has_below || v > below_max) { below_max = v; has_below = true; }
}
}
if (fb2_engine_pivot(low + base, lookback, i, 3, true)) {
const double v = low[base + i];
if (v < price) {
// keep the 3 largest (descending)
if (n_lbelow < 3) { lbelow[n_lbelow++] = v; }
else if (v > lbelow[2]) { lbelow[2] = v; }
for (int a = n_lbelow - 1; a > 0; --a) if (lbelow[a] > lbelow[a-1]) { double t = lbelow[a]; lbelow[a] = lbelow[a-1]; lbelow[a-1] = t; }
} else {
if (!has_labove || v < labove_min) { labove_min = v; has_labove = true; }
}
}
}
for (int a = 0; a < n_above; ++a) { sh[a] = above[a]; sh_mask |= (1 << a); }
for (int a = 0; a < n_lbelow; ++a) { sl[a] = lbelow[a]; sl_mask |= (1 << a); }
if (n_above == 0 && has_below) { sh[0] = below_max; sh_mask |= 1; K1I(sh_fallback) = 1; }
if (n_lbelow == 0 && has_labove) { sl[0] = labove_min; sl_mask |= 1; K1I(sl_fallback) = 1; }
}
}
K1I(sh_mask) = sh_mask; K1I(sl_mask) = sl_mask;
K1F(sh1) = sh[0]; K1F(sh2) = sh[1]; K1F(sh3) = sh[2];
K1F(sl1) = sl[0]; K1F(sl2) = sl[1]; K1F(sl3) = sl[2];
// ── Candlestick (engine._detect_candlestick_patterns) ──
int cs_present = 0, cs_signal = 0; double cs_conf = 0.0, cs_value = 0.0;
if (!(dis & DIS_candlestick) && len >= 5) {
const double o = open_[pos], h = high[pos], l = low[pos], c = close[pos];
const double prev_o = open_[pos - 1], prev_c = close[pos - 1];
const double body = fabs(c - o);
const double upper_wick = h - PYMAX(o, c);
const double lower_wick = PYMIN(o, c) - l;
const double full_range = h - l;
if (full_range != 0.0) {
const double body_ratio = body / full_range;
int mask = 0, bullish = 0, bearish = 0;
if (body_ratio < 0.1) mask |= (1 << 0);
if (lower_wick > 2.0 * body && upper_wick < body * 0.5 && body_ratio > 0.1) {
if (c > o) { mask |= (1 << 1); bullish += 1; } else { mask |= (1 << 2); bearish += 1; }
}
if (upper_wick > 2.0 * body && lower_wick < body * 0.5 && body_ratio > 0.1) {
if (c < o) { mask |= (1 << 3); bearish += 1; } else { mask |= (1 << 4); bullish += 1; }
}
if (len >= 3) {
if (prev_c < prev_o && c > o && c > prev_o && o < prev_c) { mask |= (1 << 5); bullish += 1; }
else if (prev_c > prev_o && c < o && c < prev_o && o > prev_c) { mask |= (1 << 6); bearish += 1; }
}
if (body_ratio > 0.85) {
if (c > o) { mask |= (1 << 7); bullish += 1; } else { mask |= (1 << 8); bearish += 1; }
}
if (mask) {
cs_present = 1;
cs_signal = (bullish > bearish) ? 1 : (bearish > bullish) ? -1 : 0;
cs_conf = 50.0;
const double proximity = (atr_kl > 0.0) ? atr_kl * 0.5 : 0.0;
if (proximity > 0.0) {
bool near_support = false, near_resistance = false;
for (int a = 0; a < 3; ++a) {
if ((sl_mask & (1 << a)) && fabs(c - sl[a]) <= proximity) near_support = true;
if ((sh_mask & (1 << a)) && fabs(c - sh[a]) <= proximity) near_resistance = true;
}
if (bullish > bearish && near_support) { cs_conf += 20.0; mask |= (1 << 9); }
else if (bearish > bullish && near_resistance) { cs_conf += 20.0; mask |= (1 << 10); }
}
cs_value = (double)(bullish - bearish);
K1I(cs_present) = 1; K1I(cs_signal) = cs_signal; K1I(cs_mask) = mask;
K1I(cs_bull) = bullish; K1I(cs_bear) = bearish;
K1F(cs_value) = cs_value; K1F(cs_conf) = cs_conf;
}
}
}
// ── Regime ──
int regime;
if (adx_present && adx_val >= PF[PF_adx_strong_trend]) {
if (ema_present && (ema_signal == 1 || ema_signal == 2)) regime = REG_trending_up;
else if (ema_present && (ema_signal == -1 || ema_signal == -2)) regime = REG_trending_down;
else regime = REG_trending;
} else if (adx_present && slope > PF[PF_adx_slope_emerging_threshold]) {
regime = REG_transitioning;
} else {
regime = REG_ranging;
}
K1I(regime) = regime;
// ── Regime-aware RSI (replaces the classic vote) ──
if (PI[PI_rsi_regime_aware] && !(dis & DIS_rsi) && rsi_present) {
const double v = rsi_val;
if (regime == REG_trending_up) {
if (38.0 <= v && v <= 55.0) { rsi_signal = 2; rsi_conf = 75.0; rsi_desc = 5; }
else if (30.0 <= v && v < 38.0) { rsi_signal = 1; rsi_conf = 60.0; rsi_desc = 6; }
else if (55.0 < v && v <= 65.0) { rsi_signal = 1; rsi_conf = 50.0; rsi_desc = 7; }
else if (v > 65.0) { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 8; }
else if (v < 30.0) { rsi_signal = 0; rsi_conf = 25.0; rsi_desc = 9; }
} else if (regime == REG_trending_down) {
if (45.0 <= v && v <= 62.0) { rsi_signal = -2; rsi_conf = 75.0; rsi_desc = 10; }
else if (62.0 < v && v <= 70.0) { rsi_signal = -1; rsi_conf = 60.0; rsi_desc = 11; }
else if (35.0 <= v && v < 45.0) { rsi_signal = -1; rsi_conf = 50.0; rsi_desc = 12; }
else if (v < 35.0) { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 13; }
else if (v > 70.0) { rsi_signal = 0; rsi_conf = 25.0; rsi_desc = 14; }
} else if (regime == REG_transitioning || regime == REG_trending) {
if (v >= 70.0) { rsi_signal = -1; rsi_conf = 50.0; rsi_desc = 15; }
else if (v <= 30.0) { rsi_signal = 1; rsi_conf = 50.0; rsi_desc = 16; }
else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 17; }
} else {
if (v >= 70.0) { rsi_signal = -1; rsi_conf = 70.0; rsi_desc = 18; }
else if (v <= 30.0) { rsi_signal = 1; rsi_conf = 70.0; rsi_desc = 19; }
else if (v >= 60.0) { rsi_signal = -1; rsi_conf = 45.0; rsi_desc = 20; }
else if (v <= 40.0) { rsi_signal = 1; rsi_conf = 45.0; rsi_desc = 21; }
else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 22; }
}
}
if (rsi_present) {
K1I(rsi_present) = 1; K1I(rsi_signal) = rsi_signal; K1I(rsi_desc) = rsi_desc;
K1F(rsi_value) = rsi_val; K1F(rsi_conf) = rsi_conf;
}
// ── Composite (insertion order: ema, rsi, bollinger, adx, macd, roc, keltner, divergence, candlestick) ──
double total_score = 0.0, total_weight = 0.0, total_conf = 0.0;
const double* rm = RM + regime * 9;
#define FB2_VOTE(present, sig, conf, base, col, ewf, wsf) \
if (present) { \
const double w_ = (base) * rm[col]; \
const double ws_ = ((double)(sig) * w_) * ((conf) / 100.0); \
K1F(ewf) = w_; K1F(wsf) = ws_; \
total_score += ws_; total_weight += w_; total_conf += (conf) * w_; \
}
FB2_VOTE(ema_present, ema_signal, K1F(ema_conf), 1.0, 0, ew_ema, ws_ema)
FB2_VOTE(rsi_present, rsi_signal, rsi_conf, 1.5, 1, ew_rsi, ws_rsi)
FB2_VOTE(bb_present, bb_signal, bb_conf, 1.5, 2, ew_bb, ws_bb)
FB2_VOTE(adx_present, adx_signal, adx_conf, 0.5, 3, ew_adx, ws_adx)
FB2_VOTE(macd_present, macd_signal, macd_conf, 1.5, 4, ew_macd, ws_macd)
FB2_VOTE(roc_present, roc_signal, roc_conf, 1.5, 5, ew_roc, ws_roc)
FB2_VOTE(kc_present, kc_signal, kc_conf, 1.5, 6, ew_kc, ws_kc)
FB2_VOTE(div_present, div_signal, 70.0, 1.5, 7, ew_div, ws_div)
FB2_VOTE(cs_present, cs_signal, cs_conf, 1.0, 8, ew_cs, ws_cs)
#undef FB2_VOTE
double normalized, avg_conf;
if (total_weight > 0.0) {
normalized = (total_score / total_weight) * 50.0;
avg_conf = total_conf / total_weight;
} else {
normalized = 0.0; avg_conf = 0.0;
}
normalized = PYMAX(-100.0, PYMIN(100.0, normalized));
int csig;
if (normalized >= 40.0) csig = 2; else if (normalized >= 15.0) csig = 1;
else if (normalized <= -40.0) csig = -2; else if (normalized <= -15.0) csig = -1; else csig = 0;
K1I(composite_signal) = csig;
K1F(composite_score) = normalized;
K1F(composite_conf) = avg_conf;
}
#undef K1F
#undef K1I
'''
_SRC_K2 = r'''
#define KF(OFx, nx, jj, name) ((OFx)[(size_t)(K1F_##name) * (size_t)(nx) + (size_t)(jj)])
#define KI(OIx, nx, jj, name) ((OIx)[(size_t)(K1I_##name) * (size_t)(nx) + (size_t)(jj)])
// CPython float floor division (float_floor_div / float_divmod).
__device__ __forceinline__ double fb2_py_floordiv(double vx, double wx) {
double mod = fmod(vx, wx);
double div = (vx - mod) / wx;
if (mod != 0.0) {
if ((wx < 0.0) != (mod < 0.0)) { mod += wx; div -= 1.0; }
}
double floordiv;
if (div != 0.0) {
floordiv = floor(div);
if (div - floordiv > 0.5) floordiv += 1.0;
} else {
floordiv = copysign(0.0, vx / wx);
}
return floordiv;
}
// chart_patterns._first_close_beyond
__device__ __forceinline__ bool fb2_first_close_beyond(const double* c, int n, int start, double level, bool above) {
const double last = c[n - 1];
const int m = (n - 1) - start; // len(prior)
if (above) {
if (!(last > level)) return false;
if (m <= 0) return true;
double mx = c[start];
for (int k = start + 1; k < n - 1; ++k) if (c[k] > mx) mx = c[k];
return mx <= level;
}
if (!(last < level)) return false;
if (m <= 0) return true;
double mn = c[start];
for (int k = start + 1; k < n - 1; ++k) if (c[k] < mn) mn = c[k];
return mn >= level;
}
__device__ __forceinline__ double fb2_max_range(const double* a, int lo, int hi_incl) {
double m = a[lo];
for (int k = lo + 1; k <= hi_incl; ++k) if (a[k] > m) m = a[k];
return m;
}
__device__ __forceinline__ double fb2_min_range(const double* a, int lo, int hi_incl) {
double m = a[lo];
for (int k = lo + 1; k <= hi_incl; ++k) if (a[k] < m) m = a[k];
return m;
}
// The last `want` strict pivots (chart_patterns._swing_indices), in order.
__device__ __forceinline__ int fb2_last_pivots(const double* v, int n, int w, bool find_high, int* out, int want) {
int cnt = 0;
for (int i = w; i < n - w; ++i) {
if (!fb2_strict_pivot(v, n, i, w, find_high)) continue;
if (cnt < want) { out[cnt++] = i; }
else { for (int a = 0; a < want - 1; ++a) out[a] = out[a + 1]; out[want - 1] = i; cnt += 1; }
}
return cnt; // total found; out holds the last min(cnt, want)
}
struct fb2_cp_hit { int pattern; int dir; double anchor; double trigger; double height; };
__device__ __forceinline__ bool fb2_cp_double(const double* h, const double* l, const double* c, int n, double atr,
const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) {
const int w = PI[PI_cp_swing_window];
const double* ext = bullish ? l : h;
int idx[2];
const int cnt = fb2_last_pivots(ext, n, w, !bullish, idx, 2);
if (cnt < 2) return false;
const int i1 = idx[0], i2 = idx[1];
if (i2 - i1 < PI[PI_cp_min_sep_bars]) return false;
if ((n - 1) - i2 > PI[PI_cp_recency_bars]) return false;
const double e1 = ext[i1], e2 = ext[i2];
if (fabs(e1 - e2) > PF[PF_cp_tolerance_atr] * atr) return false;
double neckline, height, floor_;
if (bullish) {
neckline = fb2_max_range(h, i1, i2);
floor_ = PYMIN(e1, e2);
height = neckline - floor_;
} else {
neckline = fb2_min_range(l, i1, i2);
floor_ = PYMAX(e1, e2);
height = floor_ - neckline;
}
if (!(PF[PF_cp_min_height_atr] * atr <= height && height <= PF[PF_cp_max_height_atr] * atr)) return false;
if (bullish) {
if (i2 + 1 < n && fb2_min_range(l, i2 + 1, n - 1) < PYMIN(e1, e2) - 0.1 * atr) return false;
if (!fb2_first_close_beyond(c, n, i2, neckline, true)) return false;
out->pattern = CP_double_bottom; out->dir = 1; out->anchor = PYMIN(e1, e2);
out->trigger = neckline; out->height = height;
return true;
}
if (i2 + 1 < n && fb2_max_range(h, i2 + 1, n - 1) > PYMAX(e1, e2) + 0.1 * atr) return false;
if (!fb2_first_close_beyond(c, n, i2, neckline, false)) return false;
out->pattern = CP_double_top; out->dir = -1; out->anchor = PYMAX(e1, e2);
out->trigger = neckline; out->height = height;
return true;
}
__device__ __forceinline__ bool fb2_cp_hs(const double* h, const double* l, const double* c, int n, double atr,
const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) {
const int w = PI[PI_cp_swing_window];
const double* ext = bullish ? l : h;
int idx[3];
const int cnt = fb2_last_pivots(ext, n, w, !bullish, idx, 3);
if (cnt < 3) return false;
const int i1 = idx[0], i2 = idx[1], i3 = idx[2];
if ((n - 1) - i3 > PI[PI_cp_recency_bars]) return false;
const double s1 = ext[i1], head = ext[i2], s2 = ext[i3];
const double hm = PF[PF_cp_hs_head_min_atr] * atr;
double neckline, height;
if (bullish) {
if (!(head < s1 - hm && head < s2 - hm)) return false;
if (fabs(s1 - s2) > PF[PF_cp_hs_shoulder_tol_atr] * atr) return false;
const double p1 = fb2_max_range(h, i1, i2);
const double p2 = fb2_max_range(h, i2, i3);
neckline = PYMIN(p1, p2);
height = neckline - head;
} else {
if (!(head > s1 + hm && head > s2 + hm)) return false;
if (fabs(s1 - s2) > PF[PF_cp_hs_shoulder_tol_atr] * atr) return false;
const double t1 = fb2_min_range(l, i1, i2);
const double t2 = fb2_min_range(l, i2, i3);
neckline = PYMAX(t1, t2);
height = head - neckline;
}
if (!(PF[PF_cp_min_height_atr] * atr <= height && height <= PF[PF_cp_max_height_atr] * atr)) return false;
if (bullish) {
if (i3 + 1 < n && fb2_min_range(l, i3 + 1, n - 1) < head - 0.1 * atr) return false;
if (!fb2_first_close_beyond(c, n, i3, neckline, true)) return false;
out->pattern = CP_inv_hs; out->dir = 1; out->anchor = s2; out->trigger = neckline; out->height = height;
return true;
}
if (i3 + 1 < n && fb2_max_range(h, i3 + 1, n - 1) > head + 0.1 * atr) return false;
if (!fb2_first_close_beyond(c, n, i3, neckline, false)) return false;
out->pattern = CP_hs_top; out->dir = -1; out->anchor = s2; out->trigger = neckline; out->height = height;
return true;
}
__device__ __forceinline__ bool fb2_cp_flag(const double* h, const double* l, const double* c, int n, double atr,
const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) {
const double lo_r = PF[PF_cp_flag_retrace_lo], hi_r = PF[PF_cp_flag_retrace_hi];
if (bullish) {
int ip = 0; // np.argmax(h[:-1]) — first occurrence
for (int k = 1; k < n - 1; ++k) if (h[k] > h[ip]) ip = k;
const int flag_len = (n - 1) - ip;
if (!(PI[PI_cp_flag_min_bars] <= flag_len && flag_len <= PI[PI_cp_flag_max_bars])) return false;
const int pole_start = PYMAX(0, ip - PI[PI_cp_flag_pole_bars] + 1);
const double pole_low = fb2_min_range(l, pole_start, ip);
const double pole = h[ip] - pole_low;
if (pole < PF[PF_cp_flag_pole_atr] * atr) return false;
const double flag_low = fb2_min_range(l, ip + 1, n - 1);
const double retrace = (h[ip] - flag_low) / pole;
if (!(lo_r <= retrace && retrace <= hi_r)) return false;
if (!fb2_first_close_beyond(c, n, ip, h[ip], true)) return false;
out->pattern = CP_bull_flag; out->dir = 1; out->anchor = flag_low; out->trigger = h[ip]; out->height = pole;
return true;
}
int ip = 0; // np.argmin(l[:-1])
for (int k = 1; k < n - 1; ++k) if (l[k] < l[ip]) ip = k;
const int flag_len = (n - 1) - ip;
if (!(PI[PI_cp_flag_min_bars] <= flag_len && flag_len <= PI[PI_cp_flag_max_bars])) return false;
const int pole_start = PYMAX(0, ip - PI[PI_cp_flag_pole_bars] + 1);
const double pole_high = fb2_max_range(h, pole_start, ip);
const double pole = pole_high - l[ip];
if (pole < PF[PF_cp_flag_pole_atr] * atr) return false;
const double flag_high = fb2_max_range(h, ip + 1, n - 1);
const double retrace = (flag_high - l[ip]) / pole;
if (!(lo_r <= retrace && retrace <= hi_r)) return false;
if (!fb2_first_close_beyond(c, n, ip, l[ip], false)) return false;
out->pattern = CP_bear_flag; out->dir = -1; out->anchor = flag_high; out->trigger = l[ip]; out->height = pole;
return true;
}
// chart_patterns.detect_chart_pattern on the prefix [0..pos] of a frame.
__device__ __forceinline__ bool fb2_detect_chart_pattern(const double* high, const double* low, const double* close,
int pos, double atr, const double* PF, const int* PI,
fb2_cp_hit* out) {
const int L = PI[PI_cp_lookback] + 1;
if (atr <= 0.0 || pos + 1 < L) return false;
const int base = pos + 1 - L;
const double* h = high + base; const double* l = low + base; const double* c = close + base;
const int mask = PI[PI_cp_patterns_mask];
if ((mask & (1 << CP_inv_hs)) && fb2_cp_hs(h, l, c, L, atr, PF, PI, true, out)) return true;
if ((mask & (1 << CP_hs_top)) && fb2_cp_hs(h, l, c, L, atr, PF, PI, false, out)) return true;
if ((mask & (1 << CP_double_bottom)) && fb2_cp_double(h, l, c, L, atr, PF, PI, true, out)) return true;
if ((mask & (1 << CP_double_top)) && fb2_cp_double(h, l, c, L, atr, PF, PI, false, out)) return true;
if ((mask & (1 << CP_bull_flag)) && fb2_cp_flag(h, l, c, L, atr, PF, PI, true, out)) return true;
if ((mask & (1 << CP_bear_flag)) && fb2_cp_flag(h, l, c, L, atr, PF, PI, false, out)) return true;
return false;
}
// signal_aggregator._detect_big_candle on the prefix [0..pos]; returns the
// absolute index of the bar (first occurrence of the max ratio) or -1.
__device__ __forceinline__ int fb2_big_candle(const double* high, const double* low, const double* close, int pos,
const double* PF, const int* PI, double* ratio_out) {
const double hours = PF[PF_big_candle_lookback_hours];
const int bar_secs = PI[PI_primary_bar_secs];
int nbc = (int)fb2_py_floordiv(hours * 3600.0, (double)bar_secs);
if (nbc < 1) nbc = 1;
const int period = PI[PI_atr_period];
int tail_len = nbc + 3 * period;
if (tail_len > pos + 1) tail_len = pos + 1;
const int base = pos + 1 - tail_len;
// pandas ewm state (alpha = 1/period, min_periods = period)
const double alpha = 1.0 / (double)period;
const double com = (1.0 - alpha) / alpha;
const double a = 1.0 / (1.0 + com);
const double old_wt_factor = 1.0 - a, new_wt = a;
double w = FB2_NAN, old_wt = 1.0; int nobs = 0;
double prev_atr = FB2_NAN; // atr.shift(1)
double best = FB2_NAN; int best_idx = -1;
const int win_start = tail_len - nbc; // ratio.iloc[-n:]
for (int i = 0; i < tail_len; ++i) {
const int k = base + i;
const double tr1 = high[k] - low[k];
double tr = tr1;
if (i > 0) {
const double pc = close[k - 1];
const double tr2 = fabs(high[k] - pc), tr3 = fabs(low[k] - pc);
if (tr2 > tr) tr = tr2;
if (tr3 > tr) tr = tr3;
}
// ratio for this bar uses the PRIOR bar's atr output
if (i >= win_start) {
const double ratio = tr1 / prev_atr;
if (ratio == ratio) {
if (best_idx < 0 || ratio > best) { best = ratio; best_idx = k; }
}
}
// ewm update
double out;
if (i == 0) {
w = tr; nobs = (w == w) ? 1 : 0; old_wt = 1.0;
out = (nobs >= period) ? w : FB2_NAN;
} else {
const bool is_obs = (tr == tr);
if (is_obs) nobs += 1;
if (w == w) {
old_wt *= old_wt_factor;
if (is_obs) {
if (w != tr) { w = old_wt * w + new_wt * tr; w /= (old_wt + new_wt); }
old_wt = 1.0;
}
} else if (is_obs) { w = tr; }
out = (nobs >= period) ? w : FB2_NAN;
}
prev_atr = out;
}
if (best_idx < 0) return -1;
if (best < PF[PF_big_candle_atr_multiple]) return -1;
*ratio_out = best;
return best_idx;
}
extern "C" __global__
void fb2_aggregate(const double* __restrict__ open_p, const double* __restrict__ high_p,
const double* __restrict__ low_p, const double* __restrict__ close_p, const int n_p,
const double* __restrict__ OFp, const int* __restrict__ OIp,
const double* __restrict__ OFc0, const int* __restrict__ OIc0, const long long* __restrict__ cut0, const int nc0,
const double* __restrict__ OFc1, const int* __restrict__ OIc1, const long long* __restrict__ cut1, const int nc1,
const double* __restrict__ OFc2, const int* __restrict__ OIc2, const long long* __restrict__ cut2, const int nc2,
const double* __restrict__ OFc3, const int* __restrict__ OIc3, const long long* __restrict__ cut3, const int nc3,
const double* __restrict__ rf_adx, const long long* __restrict__ rf_cut,
const signed char* __restrict__ skip_reason, const signed char* __restrict__ label_hour,
const signed char* __restrict__ label_wd, const signed char* __restrict__ close_hour,
const double* __restrict__ PF, const int* __restrict__ PI,
const int first_pos, const int n_bars,
double* __restrict__ OF2, int* __restrict__ OI2,
short* __restrict__ warn_code, double* __restrict__ warn_arg)
{
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= n_bars) return;
const int pos = first_pos + i;
#define K2F(name) OF2[(size_t)(K2F_##name) * (size_t)n_bars + (size_t)i]
#define K2I(name) OI2[(size_t)(K2I_##name) * (size_t)n_bars + (size_t)i]
#define PRF(name) KF(OFp, n_p, pos, name)
#define PRI(name) KI(OIp, n_p, pos, name)
for (int f = 0; f < K2F_N; ++f) OF2[(size_t)f * (size_t)n_bars + (size_t)i] = FB2_NAN;
for (int f = 0; f < K2I_N; ++f) OI2[(size_t)f * (size_t)n_bars + (size_t)i] = 0;
int wn = 0;
#define WARN(code, arg) do { if (wn < FB2_MAX_WARN) { warn_code[(size_t)i * FB2_MAX_WARN + wn] = (short)(code); warn_arg[(size_t)i * FB2_MAX_WARN + wn] = (double)(arg); wn += 1; } } while (0)
K2I(bc_idx) = -1;
K2I(skip_reason) = skip_reason[pos];
if (skip_reason[pos] != 0) { K2I(status) = 0; return; }
const int is_xau = PI[PI_is_xau];
const int xau_enabled = PI[PI_xau_enabled];
const double price = close_p[pos];
// ── primary analysis ──
const double primary_score = PRF(composite_score);
const double primary_conf = PRF(composite_conf);
const int primary_signal = PRI(composite_signal);
const int regime = PRI(regime);
const int vol = PRI(vol);
const int atr_present = PRI(atr_present);
const double atr_kl = atr_present ? PRF(atr) : 0.0;
const int rsi_present = PRI(rsi_present);
const double rsi_val = PRF(rsi_value);
const int bb_present = PRI(bb_present);
const double bb_pos = PRF(bb_value);
const bool trending_any = (regime == REG_trending_up || regime == REG_trending_down || regime == REG_trending);
K2F(primary_score) = primary_score; K2F(primary_conf) = primary_conf;
// ── confirmation analyses (as-of aligned) ──
const double* OFc[FB2_MAX_CONFIRM] = {OFc0, OFc1, OFc2, OFc3};
const int* OIc[FB2_MAX_CONFIRM] = {OIc0, OIc1, OIc2, OIc3};
const long long* cuts[FB2_MAX_CONFIRM] = {cut0, cut1, cut2, cut3};
const int ncs[FB2_MAX_CONFIRM] = {nc0, nc1, nc2, nc3};
const int n_confirm = PI[PI_n_confirm];
int sigs[1 + FB2_MAX_CONFIRM]; double scores[1 + FB2_MAX_CONFIRM];
int n_sigs = 1; sigs[0] = primary_signal; scores[0] = primary_score;
for (int k = 0; k < n_confirm; ++k) {
const long long cut = cuts[k][pos];
if (cut == 0) continue;
const long long idx = cut - 1;
sigs[n_sigs] = KI(OIc[k], ncs[k], idx, composite_signal);
scores[n_sigs] = KF(OFc[k], ncs[k], idx, composite_score);
n_sigs += 1;
}
K2I(n_confirm_present) = n_sigs - 1;
// ── 3. MTF agreement ──
const int thr = PI[PI_mtf_agreement_threshold];
bool mtf = false;
if (n_sigs >= 2) {
int bull = 0, bear = 0;
for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull += 1; else if (sigs[q] < 0) bear += 1; }
if (primary_score > 0.0) mtf = (bull >= thr);
else if (primary_score < 0.0) mtf = (bear >= thr);
else mtf = (bull >= thr) || (bear >= thr);
}
K2I(mtf_agreement) = mtf ? 1 : 0;
// ── 6. _compute_final_signal ──
double score = primary_score;
double confidence = primary_conf;
if (mtf) {
confidence += 10.0;
// Python 3.12 builtin sum() over floats: int 0 + s0 exactly, then the
// Neumaier-compensated accumulation of the remaining items.
double f = scores[0], c = 0.0;
for (int q = 1; q < n_sigs; ++q) {
const double x = scores[q];
const double t = f + x;
if (fabs(f) >= fabs(x)) c += (f - t) + x;
else c += (x - t) + f;
f = t;
}
if (c != 0.0 && isfinite(c)) f += c;
const double mtf_avg = f / (double)n_sigs;
score = score * 0.6 + mtf_avg * 0.4;
}
if (trending_any) {
if (rsi_present) {
if (regime == REG_trending_up && score > 0.0 && 38.0 <= rsi_val && rsi_val <= 55.0) confidence += 5.0;
else if (regime == REG_trending_down && score < 0.0 && 45.0 <= rsi_val && rsi_val <= 62.0) confidence += 5.0;
}
if (bb_present) {
if (regime == REG_trending_up && score > 0.0 && bb_pos < 0.4) confidence += 5.0;
else if (regime == REG_trending_down && score < 0.0 && bb_pos > 0.6) confidence += 5.0;
}
}
int pmask = 0;
if (!mtf) {
pmask |= (1 << PEN_mtf_disagreement);
if (n_sigs >= 2) {
bool any_pos = false, any_neg = false;
for (int q = 0; q < n_sigs; ++q) { if (scores[q] > 10.0) any_pos = true; if (scores[q] < -10.0) any_neg = true; }
if (any_pos && any_neg) WARN(W_conflict, 0.0);
}
}
if ((regime == REG_trending_up && score < 0.0) || (regime == REG_trending_down && score > 0.0)) {
WARN(W_counter_trend, 0.0);
pmask |= (1 << PEN_counter_trend);
}
if (vol == VOL_extreme) { WARN(W_extreme_vol, 0.0); pmask |= (1 << PEN_extreme_vol); }
else if (vol == VOL_low) { WARN(W_low_vol, 0.0); }
if (PRI(ema_trend_present)) {
const double ema200 = PRF(ema_trend);
if (ema200 != 0.0 && ema200 > 0.0 && price > 0.0) {
const double ema_dist_pct = fabs(price - ema200) / ema200 * 100.0;
if (ema_dist_pct < PF[PF_fb2_ema_lo]) { WARN(W_ema200_near, ema_dist_pct); pmask |= (1 << PEN_ema200_near); }
else if (ema_dist_pct > 2.5) { WARN(W_ema200_far, ema_dist_pct); pmask |= (1 << PEN_ema200_far); }
}
}
const int sh_mask = PRI(sh_mask), sl_mask = PRI(sl_mask);
const double sh[3] = {PRF(sh1), PRF(sh2), PRF(sh3)};
const double sl[3] = {PRF(sl1), PRF(sl2), PRF(sl3)};
if (atr_kl > 0.0 && score != 0.0) {
if (score > 0.0) {
bool any = false; double mn = 0.0;
for (int a = 0; a < 3; ++a) if (sl_mask & (1 << a)) { if (!any || sl[a] < mn) mn = sl[a]; any = true; }
if (any && price - mn > 3.5 * atr_kl) { WARN(W_exhaustion_low, 0.0); pmask |= (1 << PEN_exhaustion); }
} else {
bool any = false; double mx = 0.0;
for (int a = 0; a < 3; ++a) if (sh_mask & (1 << a)) { if (!any || sh[a] > mx) mx = sh[a]; any = true; }
if (any && mx - price > 3.5 * atr_kl) { WARN(W_exhaustion_high, 0.0); pmask |= (1 << PEN_exhaustion); }
}
}
if (PRI(kl_ema_present)) {
const double ema_medium = PRF(ema_medium);
if (atr_kl > 0.0 && ema_medium != 0.0 && trending_any) {
const double pullback_dist = fabs(price - ema_medium) / atr_kl;
if (pullback_dist > 2.5) { WARN(W_no_pullback, pullback_dist); pmask |= (1 << PEN_no_pullback); }
else if (pullback_dist > 2.0) { WARN(W_weak_pullback, pullback_dist); pmask |= (1 << PEN_weak_pullback); }
}
}
if (trending_any) {
const bool up_or_mixed = (regime == REG_trending_up || regime == REG_trending);
const bool down_or_mixed = (regime == REG_trending_down || regime == REG_trending);
if (rsi_present) {
if (up_or_mixed && score > 0.0 && rsi_val > 65.0) { WARN(W_rsi_ext_up, rsi_val); pmask |= (1 << PEN_rsi_extended); }
else if (down_or_mixed && score < 0.0 && rsi_val < 35.0) { WARN(W_rsi_ext_down, rsi_val); pmask |= (1 << PEN_rsi_extended); }
}
if (bb_present) {
if (up_or_mixed && score > 0.0 && bb_pos > 0.85) pmask |= (1 << PEN_bb_overextended);
else if (down_or_mixed && score < 0.0 && bb_pos < 0.15) pmask |= (1 << PEN_bb_overextended);
}
}
int exhaustion_count = 0;
if (pmask & (1 << PEN_exhaustion)) exhaustion_count += 1;
if (pmask & (1 << PEN_no_pullback)) exhaustion_count += 1;
if (pmask & (1 << PEN_rsi_extended)) exhaustion_count += 1;
if (pmask & (1 << PEN_bb_overextended)) exhaustion_count += 1;
bool momentum_fading = false;
if (exhaustion_count >= 2) {
bool roc_fading = false, macd_fading = false;
if (PRI(roc_present)) {
const double rv = PRF(roc_value);
if (score > 0.0 && rv < 0.0) roc_fading = true;
else if (score < 0.0 && rv > 0.0) roc_fading = true;
}
if (PRI(macd_present)) {
const int ms = PRI(macd_signal);
if (score > 0.0 && ms < 0) macd_fading = true;
else if (score < 0.0 && ms > 0) macd_fading = true;
}
momentum_fading = roc_fading || macd_fading;
}
const bool exhaustion_blocked = (exhaustion_count >= 2) && momentum_fading;
// decay pass in the stable-sorted order (15s, 10, 8s, 5s; append order inside a bucket)
const int order[10] = {PEN_mtf_disagreement, PEN_counter_trend, PEN_extreme_vol,
PEN_exhaustion, PEN_no_pullback, PEN_bb_overextended,
PEN_ema200_near, PEN_ema200_far, PEN_weak_pullback, PEN_rsi_extended};
const double amounts[10] = {15.0, 15.0, 10.0, 8.0, 8.0, 8.0, 5.0, 5.0, 5.0, 5.0};
const double MAXP = PF[PF_max_penalty];
double total_penalty = 0.0, decay = 1.0;
for (int q = 0; q < 10; ++q) {
if (!(pmask & (1 << order[q]))) continue;
double applied = amounts[q] * decay;
if (total_penalty + applied > MAXP) applied = MAXP - total_penalty;
total_penalty += applied;
decay *= 0.5;
if (total_penalty >= MAXP) break;
}
confidence -= total_penalty;
confidence = PYMAX(0.0, PYMIN(95.0, confidence));
if (exhaustion_blocked) { WARN(W_exhaustion_block, (double)exhaustion_count); confidence = 0.0; }
int final_signal;
if (score >= 35.0 && confidence >= 55.0) final_signal = 2;
else if (score >= 25.0 && confidence >= 50.0) final_signal = 1;
else if (score <= -35.0 && confidence >= 55.0) final_signal = -2;
else if (score <= -25.0 && confidence >= 50.0) final_signal = -1;
else final_signal = 0;
const double abs_score0 = fabs(score);
int quality;
if (abs_score0 >= 40.0 && confidence >= 65.0 && mtf) quality = Q_APLUS;
else if (abs_score0 >= 33.0 && confidence >= 55.0 && mtf) quality = Q_A;
else if (abs_score0 >= 28.0 && confidence >= 50.0) quality = Q_B;
else if (abs_score0 >= 25.0 && confidence >= 45.0) quality = Q_C;
else if (abs_score0 >= 15.0) quality = Q_D;
else quality = Q_F;
K2I(penalty_mask) = pmask; K2F(penalty_total) = total_penalty;
K2I(exhaustion_count) = exhaustion_count; K2I(momentum_fading) = momentum_fading ? 1 : 0;
// ── 6b. session / day filter on the bar LABEL ──
if (PI[PI_session_filter]) {
const int bar_hour = label_hour[pos];
const int bar_wd = label_wd[pos];
const int hmask = is_xau ? PI[PI_xau_blocked_hours_mask] : PI[PI_bt_blocked_hours_mask];
if (hmask & (1 << bar_hour)) { final_signal = 0; quality = Q_F; WARN(W_session_filter, (double)bar_hour); }
const int wmask = is_xau ? PI[PI_xau_allowed_weekdays_mask] : PI[PI_bt_allowed_weekdays_mask];
if (wmask != 0 && !(wmask & (1 << bar_wd))) { final_signal = 0; quality = Q_F; WARN(W_day_filter, (double)bar_wd); }
}
// ── 6c. transitioning boost ──
if (regime == REG_transitioning && (quality == Q_B || quality == Q_C || quality == Q_D)) {
const double abs_score = fabs(score);
if (abs_score >= 30.0 && confidence >= 50.0 && mtf) quality = Q_A;
else if (abs_score >= 25.0 && confidence >= 45.0) quality = Q_B;
}
// ── 6d. exhausted trend cap ──
if (trending_any && exhaustion_count >= 2 && (quality == Q_C || quality == Q_D)) {
WARN(W_exhausted_trend, 0.0); quality = Q_D;
}
const bool xau_is_sell_dir = (final_signal < 0) || (final_signal == 0 && score < 0.0);
#define ACTIONABLE(q) ((q) >= Q_C)
// ── 6e. SELL quality gate ──
if (xau_enabled && is_xau && xau_is_sell_dir && ACTIONABLE(quality)
&& !(PI[PI_sell_min_quality_mask] & (1 << quality))) {
WARN(W_sell_quality, (double)quality); quality = Q_D;
}
// ── 6f. MTF opposition ──
if (ACTIONABLE(quality)) {
int bull_n = 0, bear_n = 0;
for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull_n += 1; else if (sigs[q] < 0) bear_n += 1; }
const bool opposed = (bear_n >= thr && score > 0.0) || (bull_n >= thr && score < 0.0);
if (opposed) { WARN(W_mtf_opposition, 0.0); quality = Q_D; }
}
// ── 6g. SELL volatility gate ──
if (xau_enabled && is_xau && xau_is_sell_dir && ACTIONABLE(quality)
&& !(PI[PI_sell_vol_mask] & (1 << vol))) {
WARN(W_sell_vol, (double)vol); quality = Q_D;
}
// ── 6g2 / 6g3. regime and SELL-hour gates ──
if (xau_enabled && is_xau && ACTIONABLE(quality)) {
if (PI[PI_blocked_regimes_mask] & (1 << regime)) { WARN(W_regime_blocked, (double)regime); quality = Q_D; }
else if (xau_is_sell_dir && (PI[PI_sell_blocked_regimes_mask] & (1 << regime))) { WARN(W_sell_regime, (double)regime); quality = Q_D; }
else if (xau_is_sell_dir && PI[PI_sell_blocked_hours_mask] != 0) {
const int ch = close_hour[pos];
if (PI[PI_sell_blocked_hours_mask] & (1 << ch)) { WARN(W_sell_hour, (double)ch); quality = Q_D; }
}
}
// ── 6h. big-candle gate ──
bool big_candle_hit = false;
if (xau_enabled && is_xau && PF[PF_big_candle_atr_multiple] > 0.0) {
double ratio = 0.0;
const int idx = fb2_big_candle(high_p, low_p, close_p, pos, PF, PI, &ratio);
if (idx >= 0) {
big_candle_hit = true;
K2I(bc_idx) = idx; K2F(bc_ratio) = ratio;
WARN(W_big_candle, ratio);
if (ACTIONABLE(quality)) quality = Q_D;
}
}
// ── 6i. M5 noise gates ──
if (PI[PI_primary_is_m5] && is_xau && ACTIONABLE(quality)) {
const double abs_score = fabs(score);
if (abs_score < PF[PF_m5_min_abs_score]) { WARN(W_m5_score, abs_score); quality = Q_D; }
else if (PI[PI_m5_require_mtf] && !mtf) { WARN(W_m5_mtf, 0.0); quality = Q_D; }
else {
const double minc = PF[PF_m5_min_confirmation_score];
const bool is_buy = score > 0.0;
bool confirmed = false;
for (int q = 1; q < n_sigs; ++q) {
if (is_buy && scores[q] >= minc) { confirmed = true; break; }
if (!is_buy && scores[q] <= -minc) { confirmed = true; break; }
}
if (!confirmed) { WARN(W_m5_conviction, 0.0); quality = Q_D; }
else {
const bool is_sell = score < 0.0;
if (is_sell && !(PI[PI_m5_sell_max_quality_mask] & (1 << quality))) { WARN(W_m5_sell_quality, (double)quality); quality = Q_D; }
else if (PI[PI_m5_blocked_vol_mask] & (1 << vol)) { WARN(W_m5_vol, (double)vol); quality = Q_D; }
else if (is_sell && rsi_present && rsi_val < PF[PF_m5_sell_min_rsi]) { WARN(W_m5_sell_rsi, rsi_val); quality = Q_D; }
}
}
}
// ── 7. setups ──
int strategy = STRAT_NONE, direction = 0;
double setup_entry = FB2_NAN, setup_sl = FB2_NAN, setup_tp1 = FB2_NAN, setup_tp2 = FB2_NAN, rr_rounded = FB2_NAN;
const double pip_value = PF[PF_pip_value];
const double min_rr = PF[PF_min_rr_ratio];
const double tfac = PF[PF_tp_tightening_factor];
const double min_tp = PF[PF_round_trip_cost] * PF[PF_min_tp_cost_multiple];
if (quality >= Q_A && is_xau) {
const double atr = atr_kl;
if (atr > 0.0) {
const bool is_buy = (final_signal > 0) ? true : (final_signal < 0) ? false : (score >= 0.0);
double quality_scale = (quality == Q_APLUS) ? PF[PF_tp_scale_aplus] : PF[PF_tp_scale_a];
if (quality == Q_A && exhaustion_count >= 2) quality_scale = 1.0;
const double tp_mult = PF[PF_tp_atr_multiplier] * quality_scale;
double sl_, tp1, tp2;
if (is_buy) { sl_ = price - PF[PF_sl_atr_multiplier] * atr; tp1 = price + tp_mult * atr; tp2 = price + tp_mult * 1.5 * atr; }
else { sl_ = price + PF[PF_sl_atr_multiplier] * atr; tp1 = price - tp_mult * atr; tp2 = price - tp_mult * 1.5 * atr; }
tp1 = price + (tp1 - price) * tfac;
tp2 = price + (tp2 - price) * tfac;
const double sl_distance = fabs(price - sl_);
double tp_distance = fabs(tp1 - price);
double rr = (sl_distance > 0.0) ? tp_distance / sl_distance : 0.0;
if (!(rr < min_rr - 1e-9)) {
if (tp_distance < min_tp) { WARN(W_spread_gate_atr, tp_distance); }
else {
{ // printed-target floor AFTER the gates (signal_aggregator 2026-09-26)
const double floor_dist = sl_distance * PF[PF_tp_print_floor_rr];
if (tp_distance < floor_dist) {
tp_distance = floor_dist;
if (is_buy) { tp1 = price + floor_dist; tp2 = price + floor_dist * 1.5; }
else { tp1 = price - floor_dist; tp2 = price - floor_dist * 1.5; }
rr = tp_distance / sl_distance;
}
}
strategy = STRAT_XAU_ATR; direction = is_buy ? 1 : -1;
setup_entry = price;
setup_sl = fb2_pyround(sl_, 100.0);
setup_tp1 = fb2_pyround(tp1, 100.0);
setup_tp2 = fb2_pyround(tp2, 100.0);
rr_rounded = fb2_pyround(rr, 100.0);
}
}
}
}
if (strategy == STRAT_NONE && ACTIONABLE(quality)) {
const double atr = atr_kl;
bool ok = !(atr == 0.0) && !(pip_value <= 0.0);
if (ok && is_xau && atr < PF[PF_min_atr_threshold]) { WARN(W_low_atr, atr); ok = false; }
if (ok) {
const bool is_buy = (final_signal > 0) ? true : (final_signal < 0) ? false : (score >= 0.0);
const double min_sl_price = PF[PF_min_sl_pips] * pip_value;
const double tp_floor = PYMAX(min_rr, PF[PF_tp_floor_rr]);
const double pre = (tfac > 0.0) ? tp_floor / tfac : tp_floor;
const double tp_mult = 1.5, tp_swing_cap = 2.5;
const double sl_atr_floor = is_xau ? PF[PF_v2_sl_atr_multiplier] : 1.5;
const double sl_swing_buffer = 1.0;
double sl_price, sl_distance, tp1, tp2;
const double entry = price;
if (is_buy) {
bool any = false; double nearest = 0.0;
for (int a = 0; a < 3; ++a) if ((sl_mask & (1 << a)) && sl[a] < price) { if (!any || sl[a] > nearest) nearest = sl[a]; any = true; }
sl_price = any ? nearest - atr * sl_swing_buffer : price - atr * (1.0 + sl_swing_buffer);
const double min_sl = PYMAX(atr * sl_atr_floor, min_sl_price);
if (price - sl_price < min_sl) sl_price = price - min_sl;
sl_distance = price - sl_price;
double cand[3]; int nc = 0; const double lvl = price + atr * 0.3;
for (int a = 0; a < 3; ++a) if ((sh_mask & (1 << a)) && sh[a] > lvl) cand[nc++] = sh[a];
for (int a = 0; a < nc; ++a) for (int b = a + 1; b < nc; ++b) if (cand[b] < cand[a]) { double t = cand[a]; cand[a] = cand[b]; cand[b] = t; }
if (nc > 0 && cand[0] - price <= atr * tp_swing_cap && cand[0] - price >= sl_distance * pre) tp1 = cand[0];
else tp1 = price + PYMAX(atr * tp_mult, sl_distance * pre);
if (nc >= 2 && cand[1] - price <= atr * 5.0) tp2 = cand[1];
else tp2 = price + PYMAX(atr * (tp_mult + 1.5), sl_distance * pre * 1.5);
} else {
bool any = false; double nearest = 0.0;
for (int a = 0; a < 3; ++a) if ((sh_mask & (1 << a)) && sh[a] > price) { if (!any || sh[a] < nearest) nearest = sh[a]; any = true; }
sl_price = any ? nearest + atr * sl_swing_buffer : price + atr * (1.0 + sl_swing_buffer);
const double min_sl = PYMAX(atr * sl_atr_floor, min_sl_price);
if (sl_price - price < min_sl) sl_price = price + min_sl;
sl_distance = sl_price - price;
double cand[3]; int nc = 0; const double lvl = price - atr * 0.3;
for (int a = 0; a < 3; ++a) if ((sl_mask & (1 << a)) && sl[a] < lvl) cand[nc++] = sl[a];
for (int a = 0; a < nc; ++a) for (int b = a + 1; b < nc; ++b) if (cand[b] > cand[a]) { double t = cand[a]; cand[a] = cand[b]; cand[b] = t; }
if (nc > 0 && price - cand[0] <= atr * tp_swing_cap && price - cand[0] >= sl_distance * pre) tp1 = cand[0];
else tp1 = price - PYMAX(atr * tp_mult, sl_distance * pre);
if (nc >= 2 && price - cand[1] <= atr * 5.0) tp2 = cand[1];
else tp2 = price - PYMAX(atr * (tp_mult + 1.5), sl_distance * pre * 1.5);
}
tp1 = entry + (tp1 - entry) * tfac;
tp2 = entry + (tp2 - entry) * tfac;
if (is_buy) tp2 = PYMAX(tp2, tp1 + atr); else tp2 = PYMIN(tp2, tp1 - atr);
double tp1_distance = fabs(tp1 - entry);
double rr = (sl_distance > 0.0) ? tp1_distance / sl_distance : 0.0;
if (!(rr < min_rr - 1e-9)) {
if (tp1_distance < min_tp) { WARN(W_spread_gate_generic, tp1_distance); }
else {
{ // printed-target floor AFTER the gates (signal_aggregator 2026-09-26)
const double floor_dist = sl_distance * PF[PF_tp_print_floor_rr];
if (tp1_distance < floor_dist) {
tp1_distance = floor_dist;
if (is_buy) { tp1 = entry + floor_dist; tp2 = PYMAX(tp2, tp1 + atr); }
else { tp1 = entry - floor_dist; tp2 = PYMIN(tp2, tp1 - atr); }
rr = tp1_distance / sl_distance;
}
}
strategy = STRAT_GENERIC; direction = is_buy ? 1 : -1;
setup_entry = entry; setup_sl = sl_price; setup_tp1 = tp1; setup_tp2 = tp2;
rr_rounded = fb2_pyround(rr, 100.0);
}
}
}
}
if (strategy != STRAT_NONE && final_signal == 0) final_signal = direction;
// ── 9. M5 pullback (XAU_PB) ──
if (PI[PI_primary_is_m5] && is_xau && PI[PI_m5_pullback_enabled] && strategy == STRAT_NONE
&& quality != Q_F && !big_candle_hit) {
// _detect_m5_pullback
int pb_dir = 0; double pb_anchor = 0.0, pb_depth = 0.0;
const double atr = atr_kl;
if (pos + 1 >= 60 && atr != 0.0 && PRI(kl_ema_fast_present) && PRI(kl_ema_present)
&& PRI(ema_trend_present) && PRI(adx_present) && rsi_present) {
const double ema_f = PRF(ema_fast), ema_m = PRF(ema_medium), ema_s = PRF(ema_slow), ema_t = PRF(ema_trend);
if (!(PRF(adx_value) < PF[PF_m5_pullback_adx_min])) {
const double close_ = close_p[pos], open_ = open_p[pos];
const int n_touch = PI[PI_m5_pullback_touch_bars];
const double recent_low = fb2_min_range(low_p, pos - n_touch + 1, pos);
const double recent_high = fb2_max_range(high_p, pos - n_touch + 1, pos);
const double pb_low = fb2_min_range(low_p, pos - 7, pos);
const double pb_high = fb2_max_range(high_p, pos - 7, pos);
const double kc_mid = PRI(kc_present) ? PRF(kc_mid) : ema_m;
const double lo_atr = PF[PF_m5_pullback_retrace_lo], hi_atr = PF[PF_m5_pullback_retrace_hi];
const double rsi = rsi_val;
const int up_align = (ema_f > ema_m ? 1 : 0) + (ema_m > ema_s ? 1 : 0) + (ema_s > ema_t ? 1 : 0) + (close_ > ema_t ? 1 : 0);
if (up_align >= 3) {
const bool touched = recent_low <= PYMAX(ema_m, kc_mid);
const double swing_high = fb2_max_range(high_p, pos - 39, pos);
const double depth = (swing_high - pb_low) / atr;
const bool resumed = (close_ > open_) && (close_ > ema_f);
if (touched && resumed && lo_atr <= depth && depth <= hi_atr
&& PF[PF_m5_pullback_rsi_buy_lo] <= rsi && rsi <= PF[PF_m5_pullback_rsi_buy_hi]) {
pb_dir = 1; pb_anchor = pb_low; pb_depth = depth;
}
}
if (pb_dir == 0) {
const int dn_align = (ema_f < ema_m ? 1 : 0) + (ema_m < ema_s ? 1 : 0) + (ema_s < ema_t ? 1 : 0) + (close_ < ema_t ? 1 : 0);
if (dn_align >= 3) {
const bool touched = recent_high >= PYMIN(ema_m, kc_mid);
const double swing_low = fb2_min_range(low_p, pos - 39, pos);
const double depth = (pb_high - swing_low) / atr;
const bool resumed = (close_ < open_) && (close_ < ema_f);
if (touched && resumed && lo_atr <= depth && depth <= hi_atr
&& PF[PF_m5_pullback_rsi_sell_lo] <= rsi && rsi <= PF[PF_m5_pullback_rsi_sell_hi]) {
pb_dir = -1; pb_anchor = pb_high; pb_depth = depth;
}
}
}
}
}
if (pb_dir != 0) {
// _try_m5_pullback gates
const bool is_buy = pb_dir > 0;
bool ok = true;
const double signed_ = is_buy ? score : -score;
if (signed_ < PF[PF_m5_pullback_min_score]) ok = false;
if (ok) {
const double minc = PF[PF_m5_min_confirmation_score];
bool confirmed = false;
for (int q = 1; q < n_sigs; ++q) {
if ((is_buy && scores[q] >= minc) || (!is_buy && scores[q] <= -minc)) { confirmed = true; break; }
}
if (!confirmed) ok = false;
}
if (ok) {
int bull_n = 0, bear_n = 0;
for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull_n += 1; else if (sigs[q] < 0) bear_n += 1; }
if ((is_buy && bear_n >= thr) || (!is_buy && bull_n >= thr)) ok = false;
}
if (ok && (PI[PI_m5_blocked_vol_mask] & (1 << vol))) ok = false;
if (ok && !is_buy && !(PI[PI_sell_vol_mask] & (1 << vol))) ok = false;
if (ok) {
const double atr = atr_kl;
const double entry = price;
const double min_sl_price = PF[PF_min_sl_pips] * pip_value;
double sl_dist, sl_, tp1, tp2;
if (is_buy) {
sl_dist = PYMAX(entry - (pb_anchor - 0.5 * atr), min_sl_price);
sl_ = entry - sl_dist; tp1 = entry + 1.5 * sl_dist; tp2 = entry + 2.5 * sl_dist;
} else {
sl_dist = PYMAX((pb_anchor + 0.5 * atr) - entry, min_sl_price);
sl_ = entry + sl_dist; tp1 = entry - 1.5 * sl_dist; tp2 = entry - 2.5 * sl_dist;
}
if (!(1.5 * sl_dist < min_tp)) {
strategy = STRAT_XAU_PB; direction = pb_dir;
setup_entry = entry; setup_sl = sl_; setup_tp1 = tp1; setup_tp2 = tp2; rr_rounded = 1.5;
K2F(pb_retrace_atr) = pb_depth; K2I(pb_accepted) = 1;
if (quality < Q_B) quality = Q_B;
final_signal = pb_dir;
}
}
}
}
// ── 10. chart pattern (XAU_CP) ──
bool session_or_day_warned = false;
for (int q = 0; q < wn; ++q) {
const short c = warn_code[(size_t)i * FB2_MAX_WARN + q];
if (c == W_session_filter || c == W_day_filter || c == W_big_candle) session_or_day_warned = true;
}
if (is_xau && PI[PI_cp_enabled] && strategy == STRAT_NONE && !session_or_day_warned) {
const double atr = atr_kl;
if (!(atr < PF[PF_min_atr_threshold])) {
fb2_cp_hit hit;
if (fb2_detect_chart_pattern(high_p, low_p, close_p, pos, atr, PF, PI, &hit)) {
const bool is_buy = hit.dir > 0;
bool ok = true;
int bull_n = 0, bear_n = 0;
for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull_n += 1; else if (sigs[q] < 0) bear_n += 1; }
if ((is_buy && bear_n >= thr) || (!is_buy && bull_n >= thr)) ok = false;
if (ok && !is_buy) {
if (!(PI[PI_sell_vol_mask] & (1 << vol))) ok = false;
if (ok && PI[PI_sell_blocked_hours_mask] != 0) {
const int ch = close_hour[pos];
if (PI[PI_sell_blocked_hours_mask] & (1 << ch)) ok = false;
}
if (ok && PI[PI_primary_is_m5] && rsi_present && rsi_val < PF[PF_m5_sell_min_rsi]) ok = false;
}
if (ok && PI[PI_primary_is_m5] && (PI[PI_m5_blocked_vol_mask] & (1 << vol))) ok = false;
if (ok) {
const double entry = price;
const double min_sl_price = PF[PF_min_sl_pips] * pip_value;
const double height = hit.height;
double sl_dist, sl_, tp1, tp2;
if (is_buy) {
sl_dist = PYMAX(entry - (hit.anchor - 0.5 * atr), min_sl_price);
sl_ = entry - sl_dist; tp1 = entry + height; tp2 = entry + 1.5 * height;
} else {
sl_dist = PYMAX((hit.anchor + 0.5 * atr) - entry, min_sl_price);
sl_ = entry + sl_dist; tp1 = entry - height; tp2 = entry - 1.5 * height;
}
double rr = (sl_dist > 0.0) ? height / sl_dist : 0.0;
if (!(rr < min_rr - 1e-9) && !(height < min_tp)) {
{ // printed-target floor AFTER the gates (signal_aggregator 2026-09-26)
double tp1_dist = height;
const double floor_dist = sl_dist * PF[PF_tp_print_floor_rr];
if (tp1_dist < floor_dist) {
tp1_dist = floor_dist;
if (is_buy) { tp1 = entry + floor_dist; tp2 = entry + 1.5 * floor_dist; }
else { tp1 = entry - floor_dist; tp2 = entry - 1.5 * floor_dist; }
rr = tp1_dist / sl_dist;
}
}
strategy = STRAT_XAU_CP; direction = hit.dir;
setup_entry = entry; setup_sl = sl_; setup_tp1 = tp1; setup_tp2 = tp2;
rr_rounded = fb2_pyround(rr, 100.0);
K2I(cp_pattern) = hit.pattern; K2F(cp_trigger) = hit.trigger; K2F(cp_height) = height;
K2I(cp_accepted) = 1;
if (quality < Q_B) quality = Q_B;
final_signal = hit.dir;
}
}
}
}
}
// ── replay: the optional higher-TF regime filter (after analyze_from_dataframes) ──
K2I(strategy_pre_rf) = strategy;
if (PI[PI_regime_filter]) {
const long long cut = rf_cut[pos];
if (cut == 0) { K2I(rf_state) = 1; WARN(W_rf_unavailable, 0.0); }
else {
const int period = PI[PI_atr_period];
const double adxv = rf_adx[cut - 1];
if (cut < (long long)(period * 3) || adxv != adxv) { K2I(rf_state) = 2; WARN(W_rf_no_adx, 0.0); }
else {
K2F(rf_adx) = adxv;
if (adxv < PF[PF_regime_filter_min_adx]) {
K2I(rf_state) = 3; WARN(W_rf_blocked, adxv);
final_signal = 0; quality = Q_F; strategy = STRAT_NONE;
} else {
K2I(rf_state) = 4;
}
}
}
}
// ── outputs of the analysis ──
K2I(final_signal) = final_signal; K2I(quality) = quality;
K2F(final_score) = score; K2F(final_conf) = confidence;
K2I(strategy) = strategy; K2I(direction) = direction;
K2F(setup_entry) = setup_entry; K2F(setup_sl) = setup_sl; K2F(setup_tp1) = setup_tp1;
K2F(setup_tp2) = setup_tp2; K2F(rr_rounded) = rr_rounded;
K2I(warn_n) = wn;
// ── replay candidate filters (gpu_backtest._replay_span) ──
int status = 1;
if (strategy != STRAT_NONE
&& (PI[PI_quality_filter_mask] & (1 << quality))
&& (PI[PI_strategies_mask] & (1 << strategy))
&& pos + 1 < n_p) {
status = 2;
const double entry_price = open_p[pos + 1];
const double sl_distance = fabs(setup_entry - setup_sl);
const double tp1_distance = fabs(setup_tp1 - setup_entry);
const double tp2_distance = fabs(setup_tp2 - setup_entry);
K2F(entry_price) = entry_price;
K2F(sl_dist) = sl_distance; K2F(tp1_dist) = tp1_distance; K2F(tp2_dist) = tp2_distance;
if (direction > 0) {
K2F(adj_sl) = entry_price - sl_distance; K2F(adj_tp1) = entry_price + tp1_distance; K2F(adj_tp2) = entry_price + tp2_distance;
} else {
K2F(adj_sl) = entry_price + sl_distance; K2F(adj_tp1) = entry_price - tp1_distance; K2F(adj_tp2) = entry_price - tp2_distance;
}
}
K2I(status) = status;
#undef WARN
#undef ACTIONABLE
#undef K2F
#undef K2I
#undef PRF
#undef PRI
}
'''
_SRC_K3 = r'''
// position_sizing._ext_vs_sma on the frame prefix [0..len) of one timeframe.
__device__ __forceinline__ double fb2_ext_vs_sma(const double* high, const double* low, const double* close,
int len, double d) {
if (len < 21) return 0.0;
const int nt = (len < FB2_FEAT_CONFIRM_TAIL) ? len : FB2_FEAT_CONFIRM_TAIL;
const int base = len - nt;
const double sma = fb2_roll_mean_last(close + base, nt, 20);
double atr_tail[FB2_FEAT_MAX_P];
fb2_wilder_atr_tail(high, low, close, base, nt, 14, atr_tail);
const double atr_f = fb2_finite(atr_tail[nt - 1], 0.0);
if (atr_f <= 0.0) return 0.0;
return d * (fb2_finite(close[len - 1], 0.0) - fb2_finite(sma, 0.0)) / atr_f;
}
extern "C" __global__
void fb2_features(const long long* __restrict__ tpos, const int n_trades,
const double* __restrict__ open_p, const double* __restrict__ high_p,
const double* __restrict__ low_p, const double* __restrict__ close_p, const int n_p,
const double* __restrict__ OFp, const int* __restrict__ OIp,
const double* __restrict__ OF2, const int* __restrict__ OI2,
const int first_pos, const int n_bars,
const int n_confirm,
const double* __restrict__ OFc0, const long long* __restrict__ cut0, const int nc0,
const double* __restrict__ OFc1, const long long* __restrict__ cut1, const int nc1,
const int h1_mode, const double* __restrict__ h1_high, const double* __restrict__ h1_low,
const double* __restrict__ h1_close, const long long* __restrict__ h1_cut,
const int o0_mode, const double* __restrict__ o0_high, const double* __restrict__ o0_low,
const double* __restrict__ o0_close, const long long* __restrict__ o0_cut,
const int o1_mode, const double* __restrict__ o1_high, const double* __restrict__ o1_low,
const double* __restrict__ o1_close, const long long* __restrict__ o1_cut,
const signed char* __restrict__ close_hour, const signed char* __restrict__ close_wd,
const signed char* __restrict__ close_minute, const signed char* __restrict__ close_day,
const double* __restrict__ sin_tab, const double* __restrict__ cos_tab,
const double* __restrict__ PF, const int* __restrict__ PI,
double* __restrict__ F)
{
const int t = blockIdx.x * blockDim.x + threadIdx.x;
if (t >= n_trades) return;
const int pos = (int)tpos[t];
const int i2 = pos - first_pos;
#define PRF(name) KF(OFp, n_p, pos, name)
#define PRI(name) KI(OIp, n_p, pos, name)
#define R2F(name) OF2[(size_t)(K2F_##name) * (size_t)n_bars + (size_t)i2]
#define R2I(name) OI2[(size_t)(K2I_##name) * (size_t)n_bars + (size_t)i2]
double out[FB2_N_FEATURES];
for (int k = 0; k < FB2_N_FEATURES; ++k) out[k] = 0.0;
out[FEAT_f_bars_since_burst] = (double)FB2_FEAT_BURST_BARS;
int p0 = pos + 1 - FB2_FEAT_MAX_P; if (p0 < 0) p0 = 0;
const int nP = pos + 1 - p0;
const double d = (R2I(direction) > 0) ? 1.0 : -1.0;
const int atr_present = PRI(atr_present);
const double atr = atr_present ? fb2_finite(PRF(atr), 0.0) : 0.0;
const bool atr_ok = atr > 0.0;
const double atr_div = atr_ok ? atr : 1e-9;
#define NORM(v) (atr_ok ? (fb2_finite((v), 0.0) / atr_div) : 0.0)
const double px = fb2_finite(close_p[pos], 0.0);
out[FEAT_f_atr_pct] = atr_present ? fb2_finite(PRF(atr_pct), 0.0) : 0.0;
// _fill_time (decision time = bar close)
{
const int hour = close_hour[pos];
const int weekday = close_wd[pos];
const int minute = close_minute[pos];
const int day = close_day[pos];
out[FEAT_f_hour_utc] = (double)hour;
out[FEAT_f_hour_sin] = sin_tab[hour];
out[FEAT_f_hour_cos] = cos_tab[hour];
out[FEAT_f_weekday] = (double)weekday;
out[FEAT_f_sess_asian] = (3 <= hour && hour <= 6) ? 1.0 : 0.0;
out[FEAT_f_sess_london] = (7 <= hour && hour <= 12) ? 1.0 : 0.0;
out[FEAT_f_sess_overlap] = (13 <= hour && hour <= 16) ? 1.0 : 0.0;
out[FEAT_f_sess_ny] = (17 <= hour && hour <= 22) ? 1.0 : 0.0;
out[FEAT_f_minutes_to_hour] = (double)((60 - minute) % 60);
out[FEAT_f_day_of_month] = (double)day;
out[FEAT_f_first_friday] = (weekday == 4 && day <= 7) ? 1.0 : 0.0;
}
// _fill_geometry
double ranges[FB2_FEAT_MAX_P];
{
out[FEAT_f_atr] = PYMAX(0.0, fb2_finite(atr, 0.0));
for (int k = 0; k < nP; ++k) ranges[k] = high_p[p0 + k] - low_p[p0 + k];
const double last_range = fb2_finite(ranges[nP - 1], 0.0);
const double mean_range = fb2_finite(fb2_np_pairwise_sum(ranges, nP) / (double)nP, 0.0);
out[FEAT_f_atr_rel_500] = (mean_range > 0.0) ? (fb2_finite(atr, 0.0) / mean_range) : 0.0;
int cnt = 0;
for (int k = 0; k < nP; ++k) if (ranges[k] < last_range) cnt += 1;
out[FEAT_f_range_pctile_500] = (nP > 0) ? ((double)cnt / (double)nP) : 0.0;
out[FEAT_f_bar_range_atr] = NORM(last_range);
const double high_ = fb2_finite(high_p[pos], 0.0), low_ = fb2_finite(low_p[pos], 0.0), close_ = fb2_finite(close_p[pos], 0.0);
const double span = high_ - low_;
if (span > 0.0) out[FEAT_f_close_pos_in_bar] = (d > 0.0) ? (close_ - low_) / span : (high_ - close_) / span;
else out[FEAT_f_close_pos_in_bar] = 0.5;
const double entry = fb2_finite(R2F(setup_entry), 0.0);
const double stop = fb2_finite(R2F(setup_sl), 0.0);
const double tp1 = fb2_finite(R2F(setup_tp1), 0.0);
const double risk = fabs(entry - stop);
out[FEAT_f_risk_price] = risk;
out[FEAT_f_sl_atr] = NORM(risk);
out[FEAT_f_tp1_atr] = NORM(fabs(tp1 - entry));
out[FEAT_f_rr_planned] = fb2_finite(R2F(rr_rounded), 0.0);
}
// _fill_exhaustion
{
const int sh_mask = PRI(sh_mask), sl_mask = PRI(sl_mask);
const double sh[3] = {PRF(sh1), PRF(sh2), PRF(sh3)};
const double sl[3] = {PRF(sl1), PRF(sl2), PRF(sl3)};
if (d > 0.0) {
bool any = false; double mn = 0.0;
for (int a = 0; a < 3; ++a) if (sl_mask & (1 << a)) { if (!any) mn = sl[a]; else mn = PYMIN(mn, sl[a]); any = true; }
out[FEAT_f_swing_dist_atr] = any ? NORM(px - mn) : 0.0;
} else {
bool any = false; double mx = 0.0;
for (int a = 0; a < 3; ++a) if (sh_mask & (1 << a)) { if (!any) mx = sh[a]; else mx = PYMAX(mx, sh[a]); any = true; }
out[FEAT_f_swing_dist_atr] = any ? NORM(mx - px) : 0.0;
}
if (PRI(kl_ema_present)) out[FEAT_f_ema21_dist_atr] = NORM(d * (px - fb2_finite(PRF(ema_medium), 0.0)));
else out[FEAT_f_ema21_dist_atr] = 0.0;
if (PRI(ema_trend_present)) {
const double ema200_f = fb2_finite(PRF(ema_trend), 0.0);
if (ema200_f != 0.0) {
out[FEAT_f_ema200_dist_atr] = NORM(d * (px - ema200_f));
out[FEAT_f_ema200_dist_pct] = d * (px - ema200_f) / ema200_f * 100.0;
}
}
out[FEAT_f_exhaustion_count] = (double)R2I(exhaustion_count);
const int pm = R2I(penalty_mask);
out[FEAT_f_tag_exhaustion] = (pm & (1 << PEN_exhaustion)) ? 1.0 : 0.0;
out[FEAT_f_tag_no_pullback] = (pm & (1 << PEN_no_pullback)) ? 1.0 : 0.0;
out[FEAT_f_tag_weak_pullback] = (pm & (1 << PEN_weak_pullback)) ? 1.0 : 0.0;
out[FEAT_f_tag_rsi_extended] = (pm & (1 << PEN_rsi_extended)) ? 1.0 : 0.0;
out[FEAT_f_tag_bb_overextended] = (pm & (1 << PEN_bb_overextended)) ? 1.0 : 0.0;
out[FEAT_f_tag_ema200_near] = (pm & (1 << PEN_ema200_near)) ? 1.0 : 0.0;
out[FEAT_f_tag_ema200_far] = (pm & (1 << PEN_ema200_far)) ? 1.0 : 0.0;
out[FEAT_f_penalty_total] = fb2_finite(R2F(penalty_total), 0.0);
out[FEAT_f_momentum_fading] = R2I(momentum_fading) ? 1.0 : 0.0;
}
// _fill_quality
{
out[FEAT_f_score_dir] = d * fb2_finite(R2F(final_score), 0.0);
out[FEAT_f_confidence] = fb2_finite(R2F(final_conf), 0.0);
const int q = R2I(quality);
out[FEAT_f_quality_rank] = (q == Q_APLUS) ? 4.0 : (q == Q_A) ? 3.0 : (q == Q_B) ? 2.0 : (q == Q_C) ? 1.0 : 0.0;
int agree = 0, oppose = 0, total = 0;
const int pres[9] = {PRI(ema_present), PRI(rsi_present), PRI(bb_present), PRI(adx_present), PRI(macd_present),
PRI(roc_present), PRI(kc_present), PRI(div_present), PRI(cs_present)};
const int sg[9] = {PRI(ema_signal), PRI(rsi_signal), PRI(bb_signal), PRI(adx_signal), PRI(macd_signal),
PRI(roc_signal), PRI(kc_signal), PRI(div_signal), PRI(cs_signal)};
for (int k = 0; k < 9; ++k) {
if (!pres[k]) continue;
total += 1;
if (d > 0.0) { if (sg[k] > 0) agree += 1; else if (sg[k] < 0) oppose += 1; }
else { if (sg[k] < 0) agree += 1; else if (sg[k] > 0) oppose += 1; }
}
out[FEAT_f_ind_agree_frac] = total ? ((double)agree / (double)total) : 0.0;
out[FEAT_f_ind_oppose_frac] = total ? ((double)oppose / (double)total) : 0.0;
out[FEAT_f_is_buy] = (d > 0.0) ? 1.0 : 0.0;
const int st = R2I(strategy);
out[FEAT_f_strat_atr] = (st == STRAT_XAU_ATR) ? 1.0 : 0.0;
out[FEAT_f_strat_pb] = (st == STRAT_XAU_PB) ? 1.0 : 0.0;
out[FEAT_f_strat_cp] = (st == STRAT_XAU_CP) ? 1.0 : 0.0;
out[FEAT_f_strat_generic] = (st == STRAT_GENERIC) ? 1.0 : 0.0;
}
// _fill_mtf
{
out[FEAT_f_mtf_primary] = d * fb2_finite(PRF(composite_score), 0.0);
double c1 = 0.0, c2 = 0.0;
if (n_confirm >= 1) { const long long cut = cut0[pos]; if (cut > 0) c1 = KF(OFc0, nc0, cut - 1, composite_score); }
if (n_confirm >= 2) { const long long cut = cut1[pos]; if (cut > 0) c2 = KF(OFc1, nc1, cut - 1, composite_score); }
out[FEAT_f_mtf_c1] = (n_confirm >= 1) ? d * fb2_finite(c1, 0.0) : 0.0;
out[FEAT_f_mtf_c2] = (n_confirm >= 2) ? d * fb2_finite(c2, 0.0) : 0.0;
out[FEAT_f_mtf_agreement] = R2I(mtf_agreement) ? 1.0 : 0.0;
}
// _fill_indicators
{
const double rsi = PRI(rsi_present) ? fb2_finite(PRF(rsi_value), 50.0) : 50.0;
out[FEAT_f_rsi] = rsi;
out[FEAT_f_rsi_dir] = d * (rsi - 50.0);
const double bb = PRI(bb_present) ? fb2_finite(PRF(bb_value), 0.5) : 0.5;
out[FEAT_f_bb_pos] = bb;
out[FEAT_f_bb_pos_dir] = d * (bb - 0.5);
out[FEAT_f_adx] = PRI(adx_present) ? fb2_finite(PRF(adx_value), 0.0) : 0.0;
out[FEAT_f_kc_ext_atr] = NORM(d * (PRI(kc_present) ? fb2_finite(PRF(kc_value), 0.0) : 0.0));
out[FEAT_f_roc_dir] = d * (PRI(roc_present) ? fb2_finite(PRF(roc_value), 0.0) : 0.0);
out[FEAT_f_macd_hist_dir] = d * (PRI(macd_present) ? fb2_finite(PRF(macd_value), 0.0) : 0.0);
}
// _fill_path
{
const int lags[4] = {6, 12, 24, 48};
const int keys[4] = {FEAT_f_ret_6_atr, FEAT_f_ret_12_atr, FEAT_f_ret_24_atr, FEAT_f_ret_48_atr};
for (int k = 0; k < 4; ++k) {
const int lag = lags[k];
out[keys[k]] = (nP >= lag + 1) ? NORM(d * (close_p[pos] - close_p[pos - lag])) : 0.0;
}
double run = 0.0;
if (nP >= 2) {
const double last = close_p[pos] - close_p[pos - 1];
double sign = (last > 0.0) ? 1.0 : (last < 0.0) ? -1.0 : 0.0;
if (sign != 0.0) {
int count = 0;
for (int q = nP - 2; q >= 0; --q) {
const double value = close_p[p0 + q + 1] - close_p[p0 + q];
if ((value > 0.0 && sign > 0.0) || (value < 0.0 && sign < 0.0)) count += 1;
else break;
}
run = sign * (double)count;
}
}
out[FEAT_f_consec_closes_dir] = PYMAX(-20.0, PYMIN(20.0, d * run));
double atr_tail[FB2_FEAT_MAX_P];
fb2_wilder_atr_tail(high_p, low_p, close_p, p0, nP, 14, atr_tail);
int start = nP - FB2_FEAT_BURST_BARS; if (start < 0) start = 0;
for (int q = nP - 1; q >= start; --q) {
const double threshold = (q >= 1) ? atr_tail[q - 1] : FB2_NAN;
if (!isfinite(threshold) || threshold <= 0.0) continue;
if (ranges[q] >= 3.0 * threshold) { out[FEAT_f_bars_since_burst] = (double)((nP - 1) - q); break; }
}
}
// _fill_cross_tf
{
double h1v = 0.0;
if (h1_mode == 1) h1v = fb2_ext_vs_sma(high_p, low_p, close_p, pos + 1, d);
else if (h1_mode == 2) { const long long len = h1_cut[pos]; if (len > 0) h1v = fb2_ext_vs_sma(h1_high, h1_low, h1_close, (int)len, d); }
out[FEAT_f_h1_ext_atr] = h1v;
double ov = 0.0; bool done = false;
if (o0_mode == 1) { const long long len = o0_cut[pos]; if (len > 0 && len >= 21) { ov = fb2_ext_vs_sma(o0_high, o0_low, o0_close, (int)len, d); done = true; } }
if (!done && o1_mode == 1) { const long long len = o1_cut[pos]; if (len > 0 && len >= 21) { ov = fb2_ext_vs_sma(o1_high, o1_low, o1_close, (int)len, d); done = true; } }
out[FEAT_f_c_other_ext_atr] = ov;
}
for (int k = 0; k < FB2_N_FEATURES; ++k) {
const double dflt = (k == FEAT_f_bars_since_burst) ? (double)FB2_FEAT_BURST_BARS : 0.0;
F[(size_t)t * FB2_N_FEATURES + k] = fb2_finite(out[k], dflt);
}
#undef NORM
#undef PRF
#undef PRI
#undef R2F
#undef R2I
}
// gpu_backtest._apply_cooldown over sorted candidate positions (single thread).
extern "C" __global__
void fb2_cooldown(const long long* __restrict__ positions, const int n, const int cooldown_bars,
signed char* __restrict__ accepted)
{
if (blockIdx.x != 0 || threadIdx.x != 0) return;
long long last = -10000;
for (int k = 0; k < n; ++k) {
const long long pos = positions[k];
if (pos - last < cooldown_bars) { accepted[k] = 0; continue; }
accepted[k] = 1;
last = pos;
}
}
// pandas Series.median() of (High - Low) over every prefix, two heaps
// (max-heap of the lower half, min-heap of the upper half), single thread.
__device__ __forceinline__ void fb2_heap_push(double* h, int* n, double v, bool is_max) {
int i = (*n)++;
h[i] = v;
while (i > 0) {
const int p = (i - 1) / 2;
const bool up = is_max ? (h[i] > h[p]) : (h[i] < h[p]);
if (!up) break;
const double t = h[i]; h[i] = h[p]; h[p] = t; i = p;
}
}
__device__ __forceinline__ double fb2_heap_pop(double* h, int* n, bool is_max) {
const double top = h[0];
(*n)--;
h[0] = h[*n];
int i = 0;
for (;;) {
const int l = 2 * i + 1, r = 2 * i + 2;
int best = i;
if (l < *n && (is_max ? (h[l] > h[best]) : (h[l] < h[best]))) best = l;
if (r < *n && (is_max ? (h[r] > h[best]) : (h[r] < h[best]))) best = r;
if (best == i) break;
const double t = h[i]; h[i] = h[best]; h[best] = t; i = best;
}
return top;
}
extern "C" __global__
void fb2_prefix_median(const double* __restrict__ high, const double* __restrict__ low, const int n,
double* __restrict__ lo_heap, double* __restrict__ hi_heap,
double* __restrict__ median)
{
if (blockIdx.x != 0 || threadIdx.x != 0) return;
int nlo = 0, nhi = 0;
for (int k = 0; k < n; ++k) {
const double v = high[k] - low[k];
if (nlo == 0 || v <= lo_heap[0]) fb2_heap_push(lo_heap, &nlo, v, true);
else fb2_heap_push(hi_heap, &nhi, v, false);
if (nlo > nhi + 1) fb2_heap_push(hi_heap, &nhi, fb2_heap_pop(lo_heap, &nlo, true), false);
else if (nhi > nlo) fb2_heap_push(lo_heap, &nlo, fb2_heap_pop(hi_heap, &nhi, false), true);
median[k] = (nlo > nhi) ? lo_heap[0] : (lo_heap[0] + hi_heap[0]) / 2.0;
}
}
// replay_record.add_bar_row: the 160-column float32 row of one analysed bar.
extern "C" __global__
void fb2_bar_rows(const long long* __restrict__ bar_pos, const int n_rows,
const double* __restrict__ close_p, const int n_p,
const double* __restrict__ OFp, const int* __restrict__ OIp,
const double* __restrict__ OF2, const int* __restrict__ OI2,
const int first_pos, const int n_bars, const int n_confirm,
const double* __restrict__ OFc0, const int* __restrict__ OIc0, const long long* __restrict__ cut0, const int nc0,
const double* __restrict__ OFc1, const int* __restrict__ OIc1, const long long* __restrict__ cut1, const int nc1,
float* __restrict__ rows)
{
const int r = blockIdx.x * blockDim.x + threadIdx.x;
if (r >= n_rows) return;
const int pos = (int)bar_pos[r];
const int i2 = pos - first_pos;
float* row = rows + (size_t)r * FB2_N_BAR_COLS;
for (int k = 0; k < FB2_N_BAR_COLS; ++k) row[k] = nanf("");
#define R2F(name) OF2[(size_t)(K2F_##name) * (size_t)n_bars + (size_t)i2]
#define R2I(name) OI2[(size_t)(K2I_##name) * (size_t)n_bars + (size_t)i2]
row[0] = (float)close_p[pos];
row[1] = (float)R2F(final_score);
row[2] = (float)R2F(final_conf);
row[3] = (float)R2I(quality);
row[4] = (float)R2I(strategy);
row[5] = (R2I(strategy) != STRAT_NONE) ? 1.0f : 0.0f;
row[6] = (float)R2I(final_signal);
row[7] = R2I(mtf_agreement) ? 1.0f : 0.0f;
row[8] = (float)R2F(penalty_total);
row[9] = (float)R2I(exhaustion_count);
for (int role = 0; role < 3; ++role) {
const double* OFx; const int* OIx; long long jj; int nx;
if (role == 0) { OFx = OFp; OIx = OIp; jj = pos; nx = n_p; }
else {
const int k = role - 1;
if (k >= n_confirm) continue;
const long long cut = (k == 0) ? cut0[pos] : cut1[pos];
if (cut == 0) continue;
OFx = (k == 0) ? OFc0 : OFc1; OIx = (k == 0) ? OIc0 : OIc1; jj = cut - 1; nx = (k == 0) ? nc0 : nc1;
}
const int b = 10 + role * 5;
row[b + 0] = (float)KI(OIx, nx, jj, regime);
row[b + 1] = (float)KI(OIx, nx, jj, vol);
row[b + 2] = (float)KF(OFx, nx, jj, composite_score);
row[b + 3] = (float)KF(OFx, nx, jj, composite_conf);
row[b + 4] = KI(OIx, nx, jj, atr_present) ? (float)KF(OFx, nx, jj, atr) : nanf("");
// replay_record.INDICATORS order: ema_trend, rsi, macd, bollinger, adx, roc, keltner, divergence, candlestick
const int pres[9] = {KI(OIx, nx, jj, ema_present), KI(OIx, nx, jj, rsi_present), KI(OIx, nx, jj, macd_present),
KI(OIx, nx, jj, bb_present), KI(OIx, nx, jj, adx_present), KI(OIx, nx, jj, roc_present),
KI(OIx, nx, jj, kc_present), KI(OIx, nx, jj, div_present), KI(OIx, nx, jj, cs_present)};
const int sg[9] = {KI(OIx, nx, jj, ema_signal), KI(OIx, nx, jj, rsi_signal), KI(OIx, nx, jj, macd_signal),
KI(OIx, nx, jj, bb_signal), KI(OIx, nx, jj, adx_signal), KI(OIx, nx, jj, roc_signal),
KI(OIx, nx, jj, kc_signal), KI(OIx, nx, jj, div_signal), KI(OIx, nx, jj, cs_signal)};
const double val[9] = {KF(OFx, nx, jj, ema_value), KF(OFx, nx, jj, rsi_value), KF(OFx, nx, jj, macd_value),
KF(OFx, nx, jj, bb_value), KF(OFx, nx, jj, adx_value), KF(OFx, nx, jj, roc_value),
KF(OFx, nx, jj, kc_value), KF(OFx, nx, jj, div_value), KF(OFx, nx, jj, cs_value)};
const double cf[9] = {KF(OFx, nx, jj, ema_conf), KF(OFx, nx, jj, rsi_conf), KF(OFx, nx, jj, macd_conf),
KF(OFx, nx, jj, bb_conf), KF(OFx, nx, jj, adx_conf), KF(OFx, nx, jj, roc_conf),
KF(OFx, nx, jj, kc_conf), 70.0, KF(OFx, nx, jj, cs_conf)};
const double ew[9] = {KF(OFx, nx, jj, ew_ema), KF(OFx, nx, jj, ew_rsi), KF(OFx, nx, jj, ew_macd),
KF(OFx, nx, jj, ew_bb), KF(OFx, nx, jj, ew_adx), KF(OFx, nx, jj, ew_roc),
KF(OFx, nx, jj, ew_kc), KF(OFx, nx, jj, ew_div), KF(OFx, nx, jj, ew_cs)};
const double ws[9] = {KF(OFx, nx, jj, ws_ema), KF(OFx, nx, jj, ws_rsi), KF(OFx, nx, jj, ws_macd),
KF(OFx, nx, jj, ws_bb), KF(OFx, nx, jj, ws_adx), KF(OFx, nx, jj, ws_roc),
KF(OFx, nx, jj, ws_kc), KF(OFx, nx, jj, ws_div), KF(OFx, nx, jj, ws_cs)};
for (int k = 0; k < 9; ++k) {
if (!pres[k]) continue;
const int c = 25 + role * 45 + k * 5;
row[c + 0] = (float)sg[k];
row[c + 1] = (float)val[k];
row[c + 2] = (float)cf[k];
row[c + 3] = (float)ew[k];
row[c + 4] = (float)ws[k];
}
}
#undef R2F
#undef R2I
}
'''
_SRC_TESTS = r'''
extern "C" __global__
void fb2_t_round(const double* __restrict__ x, const int n, const double scale, double* __restrict__ out) {
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) out[i] = fb2_pyround(x[i], scale);
}
extern "C" __global__
void fb2_t_pairwise(const double* __restrict__ x, const long long* __restrict__ offs,
const long long* __restrict__ lens, const int n_series, double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) out[s] = fb2_np_pairwise_sum(x + offs[s], (int)lens[s]);
}
extern "C" __global__
void fb2_t_nanmean(const double* __restrict__ x, const long long* __restrict__ offs,
const long long* __restrict__ lens, const int n_series, double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) out[s] = fb2_np_nanmean(x + offs[s], (int)lens[s]);
}
extern "C" __global__
void fb2_t_ewm(const double* __restrict__ x, const long long* __restrict__ offs,
const long long* __restrict__ lens, const int n_series, const double alpha, const int minp,
double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) fb2_pd_ewm(x + offs[s], (int)lens[s], alpha, minp, out + offs[s]);
}
extern "C" __global__
void fb2_t_rollmean_last(const double* __restrict__ x, const long long* __restrict__ offs,
const long long* __restrict__ lens, const int n_series, const int window,
double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) out[s] = fb2_roll_mean_last(x + offs[s], (int)lens[s], window);
}
extern "C" __global__
void fb2_t_wilder(const double* __restrict__ high, const double* __restrict__ low, const double* __restrict__ close,
const long long* __restrict__ bases, const long long* __restrict__ lens, const int n_series,
const int period, const long long* __restrict__ out_offs, double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) fb2_wilder_atr_tail(high, low, close, (int)bases[s], (int)lens[s], period, out + out_offs[s]);
}
extern "C" __global__
void fb2_t_ext_vs_sma(const double* __restrict__ high, const double* __restrict__ low, const double* __restrict__ close,
const long long* __restrict__ lens, const int n_series, const double d, double* __restrict__ out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s < n_series) out[s] = fb2_ext_vs_sma(high, low, close, (int)lens[s], d);
}
extern "C" __global__
void fb2_t_floordiv(const double* __restrict__ a, const double* __restrict__ b, const int n, double* __restrict__ out) {
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) out[i] = fb2_py_floordiv(a[i], b[i]);
}
extern "C" __global__
void fb2_t_chart_pattern(const double* __restrict__ high, const double* __restrict__ low, const double* __restrict__ close,
const long long* __restrict__ pos_list, const int n_cases, const double* __restrict__ atrs,
const double* __restrict__ PF, const int* __restrict__ PI,
int* __restrict__ pattern, double* __restrict__ vals) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s >= n_cases) return;
fb2_cp_hit hit;
if (fb2_detect_chart_pattern(high, low, close, (int)pos_list[s], atrs[s], PF, PI, &hit)) {
pattern[s] = hit.pattern; vals[s * 4 + 0] = (double)hit.dir; vals[s * 4 + 1] = hit.anchor;
vals[s * 4 + 2] = hit.trigger; vals[s * 4 + 3] = hit.height;
} else {
pattern[s] = -1;
}
}
extern "C" __global__
void fb2_t_big_candle(const double* __restrict__ high, const double* __restrict__ low, const double* __restrict__ close,
const long long* __restrict__ pos_list, const int n_cases,
const double* __restrict__ PF, const int* __restrict__ PI,
int* __restrict__ idx_out, double* __restrict__ ratio_out) {
const int s = blockIdx.x * blockDim.x + threadIdx.x;
if (s >= n_cases) return;
double ratio = FB2_NAN;
idx_out[s] = fb2_big_candle(high, low, close, (int)pos_list[s], PF, PI, &ratio);
ratio_out[s] = ratio;
}
'''
# ═════════════════════════════════════════════════════════════════════
# Host side: source assembly, compilation, launch wrappers
# ═════════════════════════════════════════════════════════════════════
from position_sizing import (FEATURE_ORDER, MAX_LOOKBACK_BARS, CONFIRM_LOOKBACK_BARS,
BURST_LOOKBACK_BARS, BURST_ATR_MULTIPLE) # noqa: E402
assert len(FEATURE_ORDER) == N_FEATURES
assert MAX_LOOKBACK_BARS == FEAT_MAX_P and CONFIRM_LOOKBACK_BARS == FEAT_CONFIRM_TAIL
assert BURST_LOOKBACK_BARS == FEAT_BURST_BARS and BURST_ATR_MULTIPLE == 3.0
def _feature_defines() -> str:
return "".join(f"#define FEAT_{name} {k}\n" for k, name in enumerate(FEATURE_ORDER))
def full_source() -> str:
"""The complete CUDA source (defines + helpers + every kernel)."""
return (_defines() + _feature_defines() + _ROLL_HELPERS + "\n" + _SRC_COMMON
+ _SRC_K1 + _SRC_K2 + _SRC_K3 + _SRC_TESTS)
_MODULE = None # module-level: never collected in a child, compiled once
def _module():
global _MODULE
cp = _gk._require_cp()
if _MODULE is None:
_MODULE = cp.RawModule(code=full_source(), options=KERNEL_OPTIONS)
return _MODULE
def _fn(name: str):
return _module().get_function(name)
def _launch(n: int, threads: int = 128) -> tuple:
return ((max(1, (n + threads - 1) // threads),), (threads,))
def _dev(a, dtype):
cp = _gk._require_cp()
return cp.asarray(np.ascontiguousarray(a, dtype=dtype))
class TFAnalysis:
"""K1 output for one timeframe: OF float64[K1F_N, n], OI int32[K1I_N, n]."""
def __init__(self, OF, OI, n: int):
self.OF, self.OI, self.n = OF, OI, n
def f(self, name: str):
return self.OF[K1_FLOAT_FIELDS.index(name)]
def i(self, name: str):
return self.OI[K1_INT_FIELDS.index(name)]
def upload_cache(cache, df) -> Dict[str, object]:
"""Device copies of the precomputed IndicatorCache arrays and the OHLC."""
cp = _gk._require_cp()
out = {
"open": _dev(df["Open"].to_numpy(dtype=np.float64), np.float64),
"high": _dev(df["High"].to_numpy(dtype=np.float64), np.float64),
"low": _dev(df["Low"].to_numpy(dtype=np.float64), np.float64),
"close": _dev(df["Close"].to_numpy(dtype=np.float64), np.float64),
}
for name in ("ema_fast", "ema_medium", "ema_slow", "rsi", "macd_hist", "bb_mid", "bb_upper",
"bb_lower", "bb_band_width", "bb_avg_width", "atr", "adx", "plus_di", "minus_di",
"adx_slope", "roc", "kc_mid", "kc_upper", "kc_lower", "kc_squeeze"):
out[name] = _dev(getattr(cache, name), np.float64)
if cache.ema_trend is not None:
out["ema_trend"] = _dev(cache.ema_trend, np.float64)
out["has_ema_trend"] = 1
else:
out["ema_trend"] = out["ema_fast"] # never read (guarded by the flag)
out["has_ema_trend"] = 0
out["n"] = int(len(df))
out["index_ns"] = _dev(df.index.as_unit("ns").asi8, np.int64)
return out
def analyze_tf(dev: Dict, pf, pi, rm, tf_is_5min: bool) -> TFAnalysis:
cp = _gk._require_cp()
n = dev["n"]
OF = cp.empty((len(K1_FLOAT_FIELDS), n), dtype=cp.float64)
OI = cp.empty((len(K1_INT_FIELDS), n), dtype=cp.int32)
k = _fn("fb2_analyze_tf")
grid, block = _launch(n)
k(grid, block, (dev["open"], dev["high"], dev["low"], dev["close"],
dev["ema_fast"], dev["ema_medium"], dev["ema_slow"], dev["ema_trend"],
np.int32(dev["has_ema_trend"]),
dev["rsi"], dev["macd_hist"], dev["bb_mid"], dev["bb_upper"], dev["bb_lower"],
dev["bb_band_width"], dev["bb_avg_width"], dev["atr"], dev["adx"],
dev["plus_di"], dev["minus_di"], dev["adx_slope"], dev["roc"],
dev["kc_mid"], dev["kc_upper"], dev["kc_lower"], dev["kc_squeeze"],
_dev(pf, np.float64), _dev(pi, np.int32), _dev(rm, np.float64),
np.int32(1 if tf_is_5min else 0), np.int32(n), OF, OI))
cp.cuda.Device().synchronize()
return TFAnalysis(OF, OI, n)
def align_cut(tf_index_ns, primary_index_ns, primary_bar_secs: int, tf_bar_secs: int):
"""gpu_backtest._replay_span's per-bar slice length for one timeframe:
cut = searchsorted(tf_index, decision_time - tf_bar_dur, side='right') with
decision_time = primary label + primary_bar_dur. int64 nanoseconds."""
cp = _gk._require_cp()
decision = primary_index_ns + np.int64(primary_bar_secs) * np.int64(1_000_000_000)
cutoff = decision - np.int64(tf_bar_secs) * np.int64(1_000_000_000)
return cp.searchsorted(tf_index_ns, cutoff, side="right").astype(cp.int64)
class K2Out:
def __init__(self, OF2, OI2, warn_code, warn_arg, first_pos: int, n_bars: int):
self.OF2, self.OI2, self.warn_code, self.warn_arg = OF2, OI2, warn_code, warn_arg
self.first_pos, self.n_bars = first_pos, n_bars
def f(self, name: str):
return self.OF2[K2_FLOAT_FIELDS.index(name)]
def i(self, name: str):
return self.OI2[K2_INT_FIELDS.index(name)]
def aggregate(primary_dev: Dict, primary_an: TFAnalysis,
confirms: Sequence[Tuple[TFAnalysis, object]],
rf: object, flags: Dict[str, np.ndarray],
pf, pi, first_pos: int, last_pos: int) -> K2Out:
"""K2 over primary bars [first_pos, last_pos). `confirms` is a list of
(TFAnalysis, cut) in configuration order; `rf` is (adx_dev, cut) or None."""
cp = _gk._require_cp()
n_bars = int(last_pos - first_pos)
n_p = primary_dev["n"]
dummy_f = cp.zeros(1, dtype=cp.float64)
dummy_i = cp.zeros(1, dtype=cp.int32)
dummy_l = cp.zeros(1, dtype=cp.int64)
if len(confirms) > MAX_CONFIRM:
raise ValueError(f"at most {MAX_CONFIRM} confirmation timeframes are supported")
slots = []
for k in range(MAX_CONFIRM):
if k < len(confirms):
an, cut = confirms[k]
slots += [an.OF, an.OI, cut, np.int32(an.n)]
else:
slots += [dummy_f, dummy_i, dummy_l, np.int32(1)]
if rf is not None:
rf_adx, rf_cut = rf
else:
rf_adx, rf_cut = dummy_f, dummy_l
OF2 = cp.empty((len(K2_FLOAT_FIELDS), max(n_bars, 1)), dtype=cp.float64)
OI2 = cp.empty((len(K2_INT_FIELDS), max(n_bars, 1)), dtype=cp.int32)
warn_code = cp.zeros((max(n_bars, 1), MAX_WARN), dtype=cp.int16)
warn_arg = cp.zeros((max(n_bars, 1), MAX_WARN), dtype=cp.float64)
if n_bars > 0:
k = _fn("fb2_aggregate")
grid, block = _launch(n_bars)
k(grid, block, (primary_dev["open"], primary_dev["high"], primary_dev["low"], primary_dev["close"],
np.int32(n_p), primary_an.OF, primary_an.OI, *slots,
rf_adx, rf_cut,
_dev(flags["skip_reason"], np.int8), _dev(flags["label_hour"], np.int8),
_dev(flags["label_wd"], np.int8), _dev(flags["close_hour"], np.int8),
_dev(pf, np.float64), _dev(pi, np.int32),
np.int32(first_pos), np.int32(n_bars), OF2, OI2, warn_code, warn_arg))
cp.cuda.Device().synchronize()
return K2Out(OF2, OI2, warn_code, warn_arg, first_pos, n_bars)
def features(tpos, primary_dev: Dict, primary_an: TFAnalysis, k2: K2Out, n_confirm: int,
confirms: Sequence[Tuple[TFAnalysis, object]],
h1: Tuple[int, Dict, object], others: Sequence[Tuple[Dict, object]],
flags: Dict[str, np.ndarray], pf, pi) -> np.ndarray:
"""K3: the 65 features of every accepted trade (host float64[n, 65]).
`h1` = (mode, dev, cut): mode 0 none, 1 the primary itself, 2 a
confirmation frame; `others` = up to two (dev, cut) candidates for
f_c_other_ext_atr in configuration order."""
cp = _gk._require_cp()
n_trades = int(len(tpos))
if n_trades == 0:
return np.zeros((0, N_FEATURES), dtype=np.float64)
dummy_f = cp.zeros(1, dtype=cp.float64)
dummy_l = cp.zeros(1, dtype=cp.int64)
c = []
for k in range(2):
if k < len(confirms):
an, cut = confirms[k]
c += [an.OF, cut, np.int32(an.n)]
else:
c += [dummy_f, dummy_l, np.int32(1)]
h1_mode, h1_dev, h1_cut = h1
if h1_mode == 2:
h1_args = [np.int32(2), h1_dev["high"], h1_dev["low"], h1_dev["close"], h1_cut]
else:
h1_args = [np.int32(h1_mode), dummy_f, dummy_f, dummy_f, dummy_l]
o_args = []
for k in range(2):
if k < len(others):
dev, cut = others[k]
o_args += [np.int32(1), dev["high"], dev["low"], dev["close"], cut]
else:
o_args += [np.int32(0), dummy_f, dummy_f, dummy_f, dummy_l]
hours = np.arange(24)
sin_tab = np.array([math.sin(2.0 * math.pi * int(h) / 24.0) for h in hours], dtype=np.float64)
cos_tab = np.array([math.cos(2.0 * math.pi * int(h) / 24.0) for h in hours], dtype=np.float64)
F = cp.empty((n_trades, N_FEATURES), dtype=cp.float64)
k = _fn("fb2_features")
grid, block = _launch(n_trades)
k(grid, block, (_dev(tpos, np.int64), np.int32(n_trades),
primary_dev["open"], primary_dev["high"], primary_dev["low"], primary_dev["close"],
np.int32(primary_dev["n"]), primary_an.OF, primary_an.OI, k2.OF2, k2.OI2,
np.int32(k2.first_pos), np.int32(k2.n_bars), np.int32(n_confirm), *c, *h1_args, *o_args,
_dev(flags["close_hour"], np.int8), _dev(flags["close_wd"], np.int8),
_dev(flags["close_minute"], np.int8), _dev(flags["close_day"], np.int8),
_dev(sin_tab, np.float64), _dev(cos_tab, np.float64),
_dev(pf, np.float64), _dev(pi, np.int32), F))
cp.cuda.Device().synchronize()
return F.get()
def cooldown(positions: np.ndarray, cooldown_bars: int) -> np.ndarray:
cp = _gk._require_cp()
n = int(len(positions))
if n == 0:
return np.zeros(0, dtype=np.int8)
acc = cp.zeros(n, dtype=cp.int8)
k = _fn("fb2_cooldown")
k((1,), (1,), (_dev(positions, np.int64), np.int32(n), np.int32(cooldown_bars), acc))
cp.cuda.Device().synchronize()
return acc.get()
def prefix_median(high_dev, low_dev, n: int):
cp = _gk._require_cp()
lo = cp.empty(max(n, 1), dtype=cp.float64)
hi = cp.empty(max(n, 1), dtype=cp.float64)
med = cp.empty(max(n, 1), dtype=cp.float64)
k = _fn("fb2_prefix_median")
k((1,), (1,), (high_dev, low_dev, np.int32(n), lo, hi, med))
cp.cuda.Device().synchronize()
return med
def bar_rows(bar_pos, primary_dev: Dict, primary_an: TFAnalysis, k2: K2Out, n_confirm: int,
confirms: Sequence[Tuple[TFAnalysis, object]]) -> np.ndarray:
cp = _gk._require_cp()
n_rows = int(len(bar_pos))
if n_rows == 0:
return np.zeros((0, N_BAR_COLS), dtype=np.float32)
dummy_f = cp.zeros(1, dtype=cp.float64)
dummy_i = cp.zeros(1, dtype=cp.int32)
dummy_l = cp.zeros(1, dtype=cp.int64)
c = []
for k in range(2):
if k < len(confirms):
an, cut = confirms[k]
c += [an.OF, an.OI, cut, np.int32(an.n)]
else:
c += [dummy_f, dummy_i, dummy_l, np.int32(1)]
rows = cp.empty((n_rows, N_BAR_COLS), dtype=cp.float32)
k = _fn("fb2_bar_rows")
grid, block = _launch(n_rows)
k(grid, block, (_dev(bar_pos, np.int64), np.int32(n_rows), primary_dev["close"], np.int32(primary_dev["n"]),
primary_an.OF, primary_an.OI, k2.OF2, k2.OI2, np.int32(k2.first_pos), np.int32(k2.n_bars),
np.int32(n_confirm), *c, rows))
cp.cuda.Device().synchronize()
return rows.get()
# ── test helpers (thin launches of the device functions) ──────────────
def _ragged_host(series):
lens = np.array([int(len(s)) for s in series], dtype=np.int64)
offs = np.zeros(len(series), dtype=np.int64)
if len(series) > 1:
offs[1:] = np.cumsum(lens)[:-1]
flat = np.concatenate([np.asarray(s, dtype=np.float64) for s in series]) if series else np.zeros(0)
return flat, offs, lens
def t_round(x, scale: float) -> np.ndarray:
cp = _gk._require_cp()
x = np.asarray(x, dtype=np.float64)
out = cp.empty(len(x), dtype=cp.float64)
_fn("fb2_t_round")(*_launch(len(x)), (_dev(x, np.float64), np.int32(len(x)), np.float64(scale), out))
return out.get()
def t_pairwise(series) -> np.ndarray:
cp = _gk._require_cp()
flat, offs, lens = _ragged_host(series)
out = cp.empty(len(series), dtype=cp.float64)
_fn("fb2_t_pairwise")(*_launch(len(series)), (_dev(flat, np.float64), _dev(offs, np.int64), _dev(lens, np.int64),
np.int32(len(series)), out))
return out.get()
def t_nanmean(series) -> np.ndarray:
cp = _gk._require_cp()
flat, offs, lens = _ragged_host(series)
out = cp.empty(len(series), dtype=cp.float64)
_fn("fb2_t_nanmean")(*_launch(len(series)), (_dev(flat, np.float64), _dev(offs, np.int64), _dev(lens, np.int64),
np.int32(len(series)), out))
return out.get()
def t_ewm(series, alpha: float, minp: int) -> List[np.ndarray]:
cp = _gk._require_cp()
flat, offs, lens = _ragged_host(series)
out = cp.empty(len(flat), dtype=cp.float64)
_fn("fb2_t_ewm")(*_launch(len(series)), (_dev(flat, np.float64), _dev(offs, np.int64), _dev(lens, np.int64),
np.int32(len(series)), np.float64(alpha), np.int32(minp), out))
o = out.get()
return [o[a:a + n] for a, n in zip(offs, lens)]
def t_rollmean_last(series, window: int) -> np.ndarray:
cp = _gk._require_cp()
flat, offs, lens = _ragged_host(series)
out = cp.empty(len(series), dtype=cp.float64)
_fn("fb2_t_rollmean_last")(*_launch(len(series)), (_dev(flat, np.float64), _dev(offs, np.int64),
_dev(lens, np.int64), np.int32(len(series)),
np.int32(window), out))
return out.get()
def t_wilder(high, low, close, bases, lens, period: int) -> List[np.ndarray]:
cp = _gk._require_cp()
lens = np.asarray(lens, dtype=np.int64)
out_offs = np.zeros(len(lens), dtype=np.int64)
if len(lens) > 1:
out_offs[1:] = np.cumsum(lens)[:-1]
out = cp.empty(int(lens.sum()), dtype=cp.float64)
_fn("fb2_t_wilder")(*_launch(len(lens)), (_dev(high, np.float64), _dev(low, np.float64), _dev(close, np.float64),
_dev(bases, np.int64), _dev(lens, np.int64), np.int32(len(lens)),
np.int32(period), _dev(out_offs, np.int64), out))
o = out.get()
return [o[a:a + n] for a, n in zip(out_offs, lens)]
def t_ext_vs_sma(high, low, close, lens, d: float) -> np.ndarray:
cp = _gk._require_cp()
lens = np.asarray(lens, dtype=np.int64)
out = cp.empty(len(lens), dtype=cp.float64)
_fn("fb2_t_ext_vs_sma")(*_launch(len(lens)), (_dev(high, np.float64), _dev(low, np.float64), _dev(close, np.float64),
_dev(lens, np.int64), np.int32(len(lens)), np.float64(d), out))
return out.get()
def t_floordiv(a, b) -> np.ndarray:
cp = _gk._require_cp()
a = np.asarray(a, dtype=np.float64); b = np.asarray(b, dtype=np.float64)
out = cp.empty(len(a), dtype=cp.float64)
_fn("fb2_t_floordiv")(*_launch(len(a)), (_dev(a, np.float64), _dev(b, np.float64), np.int32(len(a)), out))
return out.get()
def t_chart_pattern(high, low, close, pos_list, atrs, pf, pi):
cp = _gk._require_cp()
n = len(pos_list)
pattern = cp.empty(n, dtype=cp.int32)
vals = cp.empty(n * 4, dtype=cp.float64)
_fn("fb2_t_chart_pattern")(*_launch(n), (_dev(high, np.float64), _dev(low, np.float64), _dev(close, np.float64),
_dev(pos_list, np.int64), np.int32(n), _dev(atrs, np.float64),
_dev(pf, np.float64), _dev(pi, np.int32), pattern, vals))
return pattern.get(), vals.get().reshape(n, 4)
def t_big_candle(high, low, close, pos_list, pf, pi):
cp = _gk._require_cp()
n = len(pos_list)
idx = cp.empty(n, dtype=cp.int32)
ratio = cp.empty(n, dtype=cp.float64)
_fn("fb2_t_big_candle")(*_launch(n), (_dev(high, np.float64), _dev(low, np.float64), _dev(close, np.float64),
_dev(pos_list, np.int64), np.int32(n),
_dev(pf, np.float64), _dev(pi, np.int32), idx, ratio))
return idx.get(), ratio.get()
Reference in New Issue View Git Blame Copy Permalink