""" 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"} # The grading rules K2 implements (fb2_aggregate: 0 = the shipped pair rule, # 1 = min conviction, 2 = product conviction). Pinned by name, never by the # position in signal_aggregator.GRADING_MODES; pack_params checks the two agree. GRADING_MODE_CODES = {"pair": 0, "min": 1, "product": 2} # ── 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", # grading_mode research thresholds (signal_aggregator.CONVICTION_THRESHOLDS; # all 0.0 under "pair"): aplus, a, b, c, strong, buy, boost_a, boost_b "conv_aplus", "conv_a", "conv_b", "conv_c", "conv_strong", "conv_buy", "conv_boost_a", "conv_boost_b", ) 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", "grading_mode", # 0 pair (shipped) / 1 min / 2 product ) # 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, } # grading rule: the live module is the single source of the mode list and # the thresholds (local import: this module must stay a leaf for the live # code, the reverse direction is the banned one) from signal_aggregator import CONVICTION_THRESHOLDS, GRADING_MODES, grading_mode_of assert set(GRADING_MODES) == set(GRADING_MODE_CODES), ( f"signal_aggregator.GRADING_MODES {GRADING_MODES} != the kernel's {sorted(GRADING_MODE_CODES)}") grading_mode = grading_mode_of(s) conv_t = (0.0,) * 8 if grading_mode == "pair" else CONVICTION_THRESHOLDS[grading_mode] for name, val in zip(("conv_aplus", "conv_a", "conv_b", "conv_c", "conv_strong", "conv_buy", "conv_boost_a", "conv_boost_b"), conv_t): f[name] = val 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), "grading_mode": GRADING_MODE_CODES[grading_mode], } 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; } // grading_mode 0 = the shipped pair rule (signal_aggregator._signal_pair / // _quality_pair, literals verbatim); 1 / 2 = the research conviction number // (signal_aggregator.conviction: PYMIN in the same argument order as // Python's min(abs, 0.6 * conf); product as (abs * conf) / 100, never // re-associated; thresholds ride PF from CONVICTION_THRESHOLDS). const int gmode = PI[PI_grading_mode]; const double abs_score0 = fabs(score); int final_signal; int quality; if (gmode == 0) { 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; 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; } else { const double conv = (gmode == 1) ? PYMIN(abs_score0, 0.6 * confidence) : abs_score0 * confidence / 100.0; if (score > 0.0 && conv >= PF[PF_conv_strong]) final_signal = 2; else if (score > 0.0 && conv >= PF[PF_conv_buy]) final_signal = 1; else if (score < 0.0 && conv >= PF[PF_conv_strong]) final_signal = -2; else if (score < 0.0 && conv >= PF[PF_conv_buy]) final_signal = -1; else final_signal = 0; if (conv >= PF[PF_conv_aplus] && mtf) quality = Q_APLUS; else if (conv >= PF[PF_conv_a] && mtf) quality = Q_A; else if (conv >= PF[PF_conv_b]) quality = Q_B; else if (conv >= PF[PF_conv_c]) 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); bool boost_a, boost_b; if (gmode == 0) { boost_a = abs_score >= 30.0 && confidence >= 50.0 && mtf; boost_b = abs_score >= 25.0 && confidence >= 45.0; } else { const double conv6c = (gmode == 1) ? PYMIN(abs_score, 0.6 * confidence) : abs_score * confidence / 100.0; boost_a = conv6c >= PF[PF_conv_boost_a] && mtf; boost_b = conv6c >= PF[PF_conv_boost_b]; } if (boost_a) quality = Q_A; else if (boost_b) 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()