diff --git a/CLAUDE.md b/CLAUDE.md index 26766ce..8bef4da 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -23,7 +23,7 @@ cp config.json.example config.json # then fill i venv/bin/pip install cupy-cuda12x "nvidia-cuda-runtime-cu12==12.9.*" # optional, dev box: backtester GPU backend # (CuPy from pip needs that header wheel to compile kernels) venv/bin/python test_all.py # QA suite, no network; prints its own count - # (536 tests / 35 modules as of 2026-09-18) + # (552 tests / 36 modules as of 2026-09-18) venv/bin/python main.py --pair XAUUSD --timeframe H1 --detail # one signal report in the terminal # --json full dump, --panel adds the advisory LLM panel, # --timeframe M5 | M15 | H1 | H4 @@ -35,6 +35,9 @@ venv/bin/python dashboard_publisher.py --once # one scan + venv/bin/python backtest/gpu_backtest.py --pair XAUUSD --days 365 venv/bin/python backtest/gpu_backtest.py --pair XAUUSD --start 2024-09-05 --end 2026-09-05 \ --primary-tf 15min --spread 0.65 --workers 12 --export-csv book.csv +venv/bin/python backtest/cudaBasedBacktesting.py --pair XAUUSD --start 2024-09-05 --end 2026-09-05 \ + --primary-tf 15min --spread 0.65 --export-csv book.csv # research twin: same flags, every stage on + # the GPU, seconds per book, same MD5 (section 9) ``` Backtest flags worth knowing: `--spread` (round-trip cost in dollars; the default is 0.30, the @@ -273,6 +276,8 @@ Arrows show who imports whom (generated from the code on 2026-09-15). |---|---| | `backtest/gpu_backtest.py` | The backtester. The indicator precompute over the full history and the resolution of every trade run on the GPU when CUDA works (`--backend`, `Backend:` log line names the device or the fallback reason); the per-bar replay through `analyze_from_dataframes` is the live Python signal code in forked CPU workers by design and is 80–95% of wall time | | `backtest/gpu_kernels.py` | CuPy backend state (lazy probe, never silent, sets `CUDA_PATH` to the pip header wheel) and the three CUDA kernels: EWM batch, pandas rolling mean/var replica, trade resolver. Each is a bit-for-bit twin of the NumPy/pandas code kept in `gpu_backtest.py` (`--fmad=false`, one thread per series or trade); the twins are pinned by Module 36 and by the reference MD5s | +| `backtest/cudaBasedBacktesting.py` | Research-only full-GPU twin of the backtester (2026-09-18): same CLI, same cache files, same 63-column book byte for byte, but the per-bar decision runs in CUDA kernels instead of the live Python code, so a book takes seconds (M5 two-year 2.4 s against 115 s). Reuses the loader, the precompute, the resolver, the metrics and the writers from `gpu_backtest.py`; the host only turns the kernels' numeric codes into the exact strings. No CPU path: `--backend cpu` is refused. Its numbers count as evidence only while the MD5 parity with `gpu_backtest.py` holds (section 7) | +| `backtest/cuda_signal_kernels.py` | The CUDA decision stage: K1 (`engine.py` plus the `GPUTechnicalEngine` overrides, one thread per bar), K2 (`analyze_from_dataframes` end to end with the gates, the setup generators, XAU_PB, XAU_CP, the regime filter and the replay's per-bar rules, one thread per primary bar), K3 (`position_sizing.build_features`), K4 (cooldown), K5 (prefix median), K6 (recorder bar rows); the config packer and the name lists the `#define`s are generated from. A leaf module that never imports the live code; pinned by Module 37 | | `backtest/replay_record.py` | Opt-in `--record` recorder: why every accepted trade was taken, plus calendar events. Book-neutral by contract | | `backtest/replay_scorecard.py` | Read-only reader of recorded runs (`DIR --check` verifies its invariants) | | `scripts/ff_calendar_scraper.py` | Builds `backtest/cache/calendar.parquet` (Forex Factory events) for the recorder | @@ -287,7 +292,7 @@ Arrows show who imports whom (generated from the code on 2026-09-15). | `app_config.py` | Loads `config.json` once per process; env var → config.json → refuse; validates keys | | `config.json.example` | The schema of `config.json` (which is gitignored) | | `deploy/README-dashboard.md`, `deploy/gold-dashboard.service` | Live-server checklist and systemd unit for the dashboard daemon | -| `test_all.py` | The QA suite (35 modules; Module 36 is the GPU backend parity module) | +| `test_all.py` | The QA suite (36 modules; Module 36 is the GPU backend parity module, Module 37 the CUDA decision-stage parity module) | --- @@ -365,9 +370,12 @@ day, 30 kept), `logs/dashboard.log`. Read the databases with `mode=ro`, never `i ## 6. Testing - `venv/bin/python test_all.py` runs everything with mocked network. Expect - `RESULTS: 536 passed, 0 failed, 536 total` (2026-09-18). On a box with a visible CUDA device + `RESULTS: 552 passed, 0 failed, 552 total` (2026-09-18). On a box with a visible CUDA device Module 36 fails if the GPU backend falls back (the reason is in the message); with - `FB2_BACKEND=cpu` or no CUDA the GPU-path tests print a notice and the NumPy twins are tested. The suite's hermetic app config points + `FB2_BACKEND=cpu` or no CUDA the GPU-path tests print a notice and the NumPy twins are tested. + Module 37 (the CUDA decision stage) adds about 85 s on the GPU because its references replay + two synthetic fixtures through the live signal code bar by bar; without CUDA each of its device + tests prints a notice and only the arithmetic rule, packing, flag and source-guard tests run. The suite's hermetic app config points both database writers at a suite-private temp dir; before 2026-09-17 the RSS fixture tests wrote fixture headlines into the real `logs/newsdata.db`, and a Module 32 test now pins the isolation. One Twelve Data interval-map test reaches the network and can time out on a slow connection; a rerun clears it. @@ -416,7 +424,13 @@ Process rules that go with them: every backtest report follows `backtest/instructions_for_backtest.txt`; every candidate is compared with a control run on the same code and cache, never with an old book directly; the `Backend:` line (since 2026-09-18 it names the CUDA device or the fallback reason) and the "GPU cache fallbacks: 0" line are part of -the evidence; keep `docs/` in sync when behaviour changes. +the evidence; keep `docs/` in sync when behaviour changes. Any change to `engine.py`, +`signal_aggregator.py`, `chart_patterns.py`, `position_sizing.py`, the recorder's row layout or +the `GPUTechnicalEngine` overrides is also a **parity event for the CUDA research twin** +(`backtest/cudaBasedBacktesting.py`): mirror it in `backtest/cuda_signal_kernels.py` in the same +commit, keep Module 37 green, and regenerate at least the M15 OOS book with both scripts and show +equal MD5s before a number from the twin is used as evidence. Until then its books are not +evidence. --- @@ -480,12 +494,20 @@ H1 12-month (2025-08-04 → 2026-08-04) 91 / +22.2R / PF 1.79. All five MD5s abo reproduced on 2026-09-15 after the cache prune and on 2026-09-18 on both backends (`backtest/research/gpu_backend_2026-09-18.md`). +The CUDA research twin (`backtest/cudaBasedBacktesting.py`, section 3) reproduced all five MD5s +above on 2026-09-18 in 0.8–2.4 s per book, plus the four 6-month acceptance books (2026-03-18 → +2026-09-18, M5 and M15 at $0.65 and $0.30: `86111c36…`, `e70d0252…`, `b8c98083…`, `e66b5fbd…`, +each equal to the reference script's), the M15 OOS `--export-features` CSV (`6239d197…`) and +`--record --record-bars` tables, and ten research-flag variants +(`backtest/research/cuda_backtester_2026-09-18.md`). + Kept cache windows (`backtest/cache/oanda_XAU_USD___.parquet`; the padded start is the window start minus 62 days for M5/M15/H1, 234 for H4 and 1,400 for D1, so the D1 and H4 files of one run have different start dates than its M5 file): 3-year `20230705_20260905`, 2-year `20240705_20260905`, OOS `20240331_20250330`, M5 12-month `20250613_20260814`, M15/H1 12-month `20250603_20260804`, 5-year `20210531_20260804`, 7.6-year `20181031_20260807`, the -2026-08-14 parity window `20250814_20260814`, their D1/H4 pads, and `calendar.parquet`. +2026-08-14 parity window `20250814_20260814`, the 6-month CUDA acceptance window +`20260115_20260918` (M5/M15/H1), their D1/H4 pads, and `calendar.parquet`. The research recorder (`--record DIR [--record-bars]`) writes Parquet files describing every accepted trade (indicator votes on all three timeframes, penalties, setup, outcome, calendar @@ -541,6 +563,20 @@ MD5s. Scorecards from it are in-book statistics, never decisions. - The replay workers are forked after the CUDA context exists and never touch the device (every cache array is host NumPy, asserted before the fork). Python 3.12 may warn about forking a multi-threaded process under `-X dev`; that is expected. +- `backtest/cudaBasedBacktesting.py` is a CUDA re-implementation of the decision stage, exempt + from "replays go through the real signal code" only while its book MD5 equals + `gpu_backtest.py`'s. A signal change that is not mirrored in `cuda_signal_kernels.py` makes the + twin drift silently: no test compares it with the live code except Module 37 and the paired + MD5 runs. It has no CPU fallback, refuses `--backend cpu`, and every description, warning and + reason string is produced on the host from the same f-string templates as the live code; the + kernels emit codes and numbers only. +- The kernels keep Python's evaluation order on purpose: `min`/`max` as argument-order ternaries + (never `fmin`/`fmax`), numpy's pairwise summation for `mean`/`nanmean`, pandas' online rolling + and EWM recurrences, Python 3.12's Neumaier-compensated `sum()` for the multi-timeframe mean, + `round(x, 2)` through an error-free product, `atr_percentile = count / 50 * 100` (not + `count * 2`: 15/50·100 = 30.000000000000004 flips "low" to "normal"). A cleaner formulation + flips a threshold somewhere in 500k bars. The as-of alignment compares int64 nanoseconds; a + mixed-unit index once returned slice length 0 for every confirmation frame. - The dashboard's due-timeframe check uses the raw wake time; flooring `wake − buffer` skipped every H1 and H4 scan once (2026-08-02). - The agent panel's price context reads the frames `analyze_pair` stashes on `report._frames`. @@ -561,6 +597,7 @@ One line per validated change or verdict, newest first. Full evidence in the nam | Date | Decision | Where the evidence is | |---|---|---| +| 2026-09-18 | Full-GPU research twin SHIPPED: `backtest/cudaBasedBacktesting.py` + `backtest/cuda_signal_kernels.py` (owner decisions: research only, `gpu_backtest.py` stays the base reference, every research flag supported, the one exemption from instructions item 6 recorded there). Proof: the five reference MD5s, four 6-month acceptance books (2026-03-18 → 2026-09-18, M5 and M15, $0.65 and $0.30), the M15 OOS `--export-features` CSV and `--record --record-bars` tables, and ten research-flag variants, all byte-identical to the reference script; Module 37 (16 tests) pins every stage against the live code. Books take 0.8–2.4 s instead of 13–115 s. Rule added: any signal change is a parity event for the twin | `backtest/research/cuda_backtester_2026-09-18.md` | | 2026-09-18 | GPU backend made real: CuPy had never compiled a kernel on the dev box (`CUDA_PATH` missing for the pip layout), so every book to date was NumPy. SHIPPED: `backtest/gpu_kernels.py` (lazy, loud backend probe; `nvidia-cuda-runtime-cu12` header wheel; CUDA kernels for the EWM batch, the pandas rolling mean/var replica and the trade resolver, each a bit-for-bit twin of the CPU code), `--backend auto/gpu/cpu` + `FB2_BACKEND`, Module 36. All five reference MD5s reproduced on the GPU and with `--backend cpu`; precompute GPU == CPU on all 42 cached windows (2.34 M bars). Wall time roughly unchanged: the replay is CPU by design | `backtest/research/gpu_backend_2026-09-18.md` | | 2026-09-17 | Local news scorer evaluated (the owner's scikit TF-IDF model in `~/Development/ML-NEWS`, trained on 606 DeepSeek gold labels): golden set 39/78 vs DeepSeek 78/78 vs keywords 24/78; out-of-sample agreement with DeepSeek 64% (DeepSeek agrees with itself 90%); cleaning / negation marking / label weighting move nothing; backfilling the 961 missing gold labels lifts macro-headline agreement 38% → 52% but only reshuffles golden-set errors. Verdict: a fallback, not a replacement. SHIPPED to grow a complete corpus: `xau_score` archived for every headline, no-currency headlines scored and archived (never shown), every new headline sent in its cycle (25 per call). Dashboard process made database-read-only (`set_db_writes_for_process`); the bot is the single writer. Books unaffected (news never enters signals) | `~/Development/ML-NEWS/eval/RESULTS_2026-09-17.md` (evaluation scripts live outside the repo by owner instruction) | | 2026-09-17 | Live review of two weeks of `signals.db`: live == backtest bar for bar (26/29 trades matched, R within 0.14); the raw advisory −7.15R is cluster re-entries the validated books never take (cooldown view −1.6R, backtester +1.1R on the window, an ordinary fortnight); composite score inversely related to the 24 h move on all timeframes (two-week bear reading, test in-book with `--record-bars`); big-candle gate re-tested and REJECTED again (removes +32R / +60R of net-positive trades). Agent panel v2 SHIPPED: probabilities anchored to book base rates, code-computed decision context, no news doctrine, mean aggregator (`verdict_source: "mean"`), prompt version logged, offline replay harness; owner decision the same day: only the macro persona runs live (`personas: ["macro"]`, 1 call per panel; technical / risk / contrarian carried no information in either evaluation, kept for research); backtester `strategy` tag fixed (label-only, books re-pinned). Delivery findings: dashboard daemon on stale code since Sep 10, bot outage Sep 9 | `backtest/research/signal_log_review_2026-09-17.md`, `backtest/research/panel_v2_2026-09-17.md` | diff --git a/backtest/cudaBasedBacktesting.py b/backtest/cudaBasedBacktesting.py new file mode 100644 index 0000000..32c790a --- /dev/null +++ b/backtest/cudaBasedBacktesting.py @@ -0,0 +1,1028 @@ +#!/usr/bin/env python3 +""" +cudaBasedBacktesting.py — the full-GPU twin of backtest/gpu_backtest.py. + +Same command line, same cache files, same exported book (byte for byte — +the MD5 of the CSV is the acceptance gate), but every stage of the +computation runs on the GPU: the indicator precompute (shared with +gpu_backtest.py), the per-bar decision (backtest/cuda_signal_kernels.py, a +CUDA twin of engine.py + signal_aggregator.py + chart_patterns.py), the +entry cooldown, the decision-time features (--export-features), the +recorder rows (--record, --record-bars) and the trade resolution (shared). + +Research tool only (owner decision 2026-09-18): live scans stay on the CPU +code and gpu_backtest.py stays in the tree as the base reference. Any change +to the signal logic is a parity event for this file — regenerate a book with +both scripts and show equal MD5s before using numbers from here as evidence +(backtest/instructions_for_backtest.txt). + +What stays on the host, on purpose: loading and cleaning the parquet cache, +the calendar/session flags (pandas time-zone arithmetic over the whole +index), turning the kernels' numeric codes into the exact strings the +reference script writes (indicator descriptions, warnings, reasons), and +the report / CSV / parquet writers, all of which are reused from the +reference script unchanged. + +There is no CPU fallback here by design: the twin IS the live code. +""" + +from __future__ import annotations + +import logging +import sys +import time +import types +from pathlib import Path +from typing import Dict, List, Optional, Sequence, Tuple + +import numpy as np +import pandas as pd + +PROJECT_ROOT = Path(__file__).resolve().parent.parent +if str(PROJECT_ROOT) not in sys.path: + sys.path.insert(0, str(PROJECT_ROOT)) + +from config import PAIR_INFO, PlatformConfig, Timeframe # noqa: E402 +from providers import make_data_provider # noqa: E402 +from providers.base import BaseDataProvider # noqa: E402 +from signal_aggregator import SignalAggregator # noqa: E402 (fingerprints + the two static cleanups only) +import position_sizing # noqa: E402 +from backtest import gpu_kernels as _gk # noqa: E402 +from backtest import cuda_signal_kernels as ck # noqa: E402 +from backtest import replay_record # noqa: E402 +from backtest.gpu_backtest import ( # noqa: E402 + BAR_SECONDS, _NY_TZ, BacktestConfig, Trade, _apply_config_overrides, _apply_fixed_rr_tp, + _cache_path, _DAY_NAMES, _history_pad_days, _parse_args, _reanchor_entries_to_raw, _resolve_pair, + _resolve_tf, compute_metrics, fetch_history, precompute_indicators, print_report, + print_trade_reasoning, simulate_trades, trades_to_dataframe, +) + +logger = logging.getLogger("cudaBasedBacktesting") +BACKEND = _gk.BACKEND + + +# ═════════════════════════════════════════════════════════════════════ +# Text: the exact templates of engine.py / signal_aggregator.py +# ═════════════════════════════════════════════════════════════════════ + +_RSI_SUFFIX = { + 0: "Overbought", 1: "Oversold", 2: "Bullish momentum zone", 3: "Bearish momentum zone", 4: "Neutral", + 5: "Pullback zone in uptrend (ideal entry)", 6: "Deep pullback in uptrend", 7: "Momentum ok in uptrend", + 8: "Extended in uptrend, wait for pullback", 9: "Oversold in uptrend, possible reversal", + 10: "Bounce zone in downtrend (ideal short)", 11: "Deep bounce in downtrend", 12: "Momentum ok in downtrend", + 13: "Extended in downtrend, wait for bounce", 14: "Overbought in downtrend, possible reversal", + 15: "Overbought while trend forming", 16: "Oversold while trend forming", 17: "No RSI edge while trend forms", + 18: "Overbought in range", 19: "Oversold in range", 20: "Upper range zone", 21: "Lower range zone", + 22: "Mid-range, no edge", +} +_MACD_DESC = { + 0: "MACD histogram accelerating ↑", 1: "MACD histogram crossed above zero", + 2: "MACD histogram accelerating ↓", 3: "MACD histogram crossed below zero", + 4: "MACD histogram positive but fading", 5: "MACD histogram negative but fading", + 6: "MACD histogram flat", +} +_BB_DESC = {0: "Price above upper BB", 1: "Price below lower BB", 2: "Price near upper BB", + 3: "Price near lower BB", 4: "Price mid-band"} +_KC_DESC = { + 0: "Keltner SQUEEZE RELEASED — bullish breakout", 1: "Keltner SQUEEZE RELEASED — bearish breakout", + 2: "Keltner squeeze released — direction unclear", 3: "Keltner SQUEEZE active — breakout imminent", + 4: "Price above Keltner upper — strong uptrend continuation", + 5: "Price below Keltner lower — strong downtrend continuation", + 6: "Price within Keltner Channels", +} +_EMA2_DESC = { + 0: "EMA200 data insufficient", 1: "Full bullish: price > EMA50 > EMA200", + 2: "Emerging uptrend: price > EMA50, EMA50 < EMA200", 3: "Full bearish: price < EMA50 < EMA200", + 4: "Emerging downtrend: price < EMA50, EMA50 > EMA200", 5: "EMA alignment mixed", +} +_DIV_ITEMS = ("Bullish RSI divergence (reversal warning)", "Bearish RSI divergence (reversal warning)", + "Bullish MACD divergence (reversal warning)", "Bearish MACD divergence (reversal warning)") +_CS_ITEMS = ("Doji (indecision)", "Hammer (bullish reversal)", "Hanging Man (bearish warning)", + "Shooting Star (bearish reversal)", "Inverted Hammer (bullish)", "Bullish Engulfing", + "Bearish Engulfing", "Bullish Marubozu (strong buyers)", "Bearish Marubozu (strong sellers)", + "(at support)", "(at resistance)") +_ADX_STRENGTH = ("Very strong", "Strong", "Weak/Ranging") +_ADX_DIR = ("bullish", "bearish") + + +class K1Row: + """One timeframe bar's K1 outputs as Python scalars (name -> value).""" + + __slots__ = ("i", "f") + + def __init__(self, ints: Dict[str, int], floats: Dict[str, float]): + self.i, self.f = ints, floats + + +def k1_rows(an: "ck.TFAnalysis", idx: np.ndarray) -> List[K1Row]: + """Gather K1 outputs for the given bar indexes (one host transfer).""" + idx = np.asarray(idx, dtype=np.int64) + if len(idx) == 0: + return [] + fi = an.OF[:, idx].get() # [K1F_N, m] + ii = an.OI[:, idx].get() # [K1I_N, m] + rows = [] + for m in range(len(idx)): + rows.append(K1Row({n: int(ii[k, m]) for k, n in enumerate(ck.K1_INT_FIELDS)}, + {n: float(fi[k, m]) for k, n in enumerate(ck.K1_FLOAT_FIELDS)})) + return rows + + +def describe_indicators(row: K1Row, ema_mode_2ema: bool) -> Dict[str, Dict]: + """The scored indicators of one analysis in engine insertion order: + name -> {value, signal (int), confidence, description, weight, + effective_weight, weighted_score}. Exactly what result.indicators + carries minus the weight-0 ATR entry.""" + i, f = row.i, row.f + out: Dict[str, Dict] = {} + if not i["valid"]: + return out + if i["ema_present"]: + if ema_mode_2ema: + desc = _EMA2_DESC[i["ema2_desc"]] + desc += {1: " | Golden cross (50/200)", 2: " | Death cross (50/200)"}.get(i["ema_cross"], "") + value = 0 if i["ema2_desc"] == 0 else f["ema_value"] + else: + cross = {1: " | Golden cross (fast/medium)", 2: " | Death cross (fast/medium)"}.get(i["ema_cross"], "") + ribbon = {1: " | Ribbon COMPRESSED (trend change imminent)", + 2: " | Ribbon EXPANDING (new trend starting)"}.get(i["ema_ribbon"], "") + desc = f"EMA alignment: {i['ema_align']}/{i['ema_max_checks']}{cross}{ribbon}" + value = f["ema_value"] + out["ema_trend"] = dict(value=value, signal=i["ema_signal"], confidence=f["ema_conf"], description=desc, + weight=1.0, effective_weight=f["ew_ema"], weighted_score=f["ws_ema"]) + if i["rsi_present"]: + v = f["rsi_value"] + out["rsi"] = dict(value=v, signal=i["rsi_signal"], confidence=f["rsi_conf"], + description=f"RSI {v:.1f} — {_RSI_SUFFIX[i['rsi_desc']]}", + weight=1.5, effective_weight=f["ew_rsi"], weighted_score=f["ws_rsi"]) + if i["bb_present"]: + v = f["bb_value"] + note = " | SQUEEZE detected (expect breakout)" if i["bb_squeeze"] else "" + out["bollinger"] = dict(value=v, signal=i["bb_signal"], confidence=f["bb_conf"], + description=f"{_BB_DESC[i['bb_desc']]} ({v:.2f}){note}", + weight=1.5, effective_weight=f["ew_bb"], weighted_score=f["ws_bb"]) + if i["adx_present"]: + slope = f["adx_slope"] + sd = {1: " | TREND EMERGING (ADX rising fast)", + 2: f" | Trend strengthening (slope +{slope:.1f})", + 3: f" | Trend weakening (slope {slope:.1f})"}.get(i["adx_slope_desc"], "") + desc = (f"ADX {f['adx_value']:.1f} ({_ADX_STRENGTH[i['adx_strength']]} {_ADX_DIR[i['adx_dir']]}) | " + f"+DI={f['plus_di']:.1f} -DI={f['minus_di']:.1f}{sd}") + out["adx"] = dict(value=f["adx_value"], signal=i["adx_signal"], confidence=f["adx_conf"], description=desc, + weight=0.5, effective_weight=f["ew_adx"], weighted_score=f["ws_adx"]) + if i["macd_present"]: + out["macd"] = dict(value=f["macd_value"], signal=i["macd_signal"], confidence=f["macd_conf"], + description=_MACD_DESC[i["macd_desc"]], + weight=1.5, effective_weight=f["ew_macd"], weighted_score=f["ws_macd"]) + if i["roc_present"]: + v = f["roc_value"] + d = i["roc_desc"] + desc = {0: f"ROC +{v:.2f}% — Strong momentum acceleration", 1: f"ROC +{v:.2f}% — Rising momentum", + 2: f"ROC {v:.2f}% — Strong bearish momentum", 3: f"ROC {v:.2f}% — Falling momentum", + 4: f"ROC {v:+.2f}% — No clear momentum"}[d] + out["roc"] = dict(value=v, signal=i["roc_signal"], confidence=f["roc_conf"], description=desc, + weight=1.5, effective_weight=f["ew_roc"], weighted_score=f["ws_roc"]) + if i["kc_present"]: + out["keltner"] = dict(value=f["kc_value"], signal=i["kc_signal"], confidence=f["kc_conf"], + description=_KC_DESC[i["kc_desc"]], + weight=1.5, effective_weight=f["ew_kc"], weighted_score=f["ws_kc"]) + if i["div_present"]: + items = [s for b, s in enumerate(_DIV_ITEMS) if i["div_mask"] & (1 << b)] + out["divergence"] = dict(value=len(items), signal=i["div_signal"], confidence=70, + description=" | ".join(items), + weight=1.5, effective_weight=f["ew_div"], weighted_score=f["ws_div"]) + if i["cs_present"]: + items = [s for b, s in enumerate(_CS_ITEMS) if i["cs_mask"] & (1 << b)] + out["candlestick"] = dict(value=f["cs_value"], signal=i["cs_signal"], confidence=f["cs_conf"], + description=" | ".join(items), + weight=1.0, effective_weight=f["ew_cs"], weighted_score=f["ws_cs"]) + return out + + +def pattern_signals(row: K1Row) -> List[str]: + if not row.i["valid"] or not row.i["cs_present"]: + return [] + return [s for b, s in enumerate(_CS_ITEMS) if row.i["cs_mask"] & (1 << b)] + + +def key_levels(row: K1Row, ema_mode_2ema: bool) -> Dict[str, float]: + """result.key_levels in insertion order.""" + i, f = row.i, row.f + kl: Dict[str, float] = {} + if not i["valid"]: + return kl + if i["atr_present"]: + kl["atr"] = f["atr"] + kl["atr_pct"] = f["atr_pct"] + if i["ema_present"]: + if i["ema_trend_present"]: + kl["ema_trend"] = f["ema_trend"] + if not ema_mode_2ema: + kl["ema_fast"] = f["ema_fast"] + kl["ema_medium"] = f["ema_medium"] + kl["ema_slow"] = f["ema_slow"] + if i["bb_present"]: + kl["bb_upper"] = f["bb_upper"]; kl["bb_lower"] = f["bb_lower"]; kl["bb_mid"] = f["bb_mid"] + if i["kc_present"]: + kl["kc_upper"] = f["kc_upper"]; kl["kc_lower"] = f["kc_lower"]; kl["kc_mid"] = f["kc_mid"] + if not i["sh_fallback"]: + for a in range(3): + if i["sh_mask"] & (1 << a): + kl[f"swing_high_{a + 1}"] = f[f"sh{a + 1}"] + if not i["sl_fallback"]: + for a in range(3): + if i["sl_mask"] & (1 << a): + kl[f"swing_low_{a + 1}"] = f[f"sl{a + 1}"] + if i["sh_fallback"]: + kl["swing_high_1"] = f["sh1"] + if i["sl_fallback"]: + kl["swing_low_1"] = f["sl1"] + return kl + + +def format_warning(code: int, arg: float, ctx: Dict) -> str: + """signal_aggregator's warning text for one K2 warning code.""" + n = ck.WARN_NAMES[code] + if n == "conflict": + return "Timeframes CONFLICT — mixed signals across timeframes" + if n == "counter_trend": + return "Counter-trend signal — reduced confidence" + if n == "extreme_vol": + return "EXTREME volatility — wider stops needed" + if n == "low_vol": + return "Low volatility — potential for breakout" + if n == "ema200_near": + return f"NEAR EMA200 ({arg:.2f}%) — chop risk" + if n == "ema200_far": + return f"FAR FROM EMA200 ({arg:.2f}%) — over-extended" + if n == "exhaustion_low": + return "TREND EXHAUSTION — price extended >3.5 ATR past swing low" + if n == "exhaustion_high": + return "TREND EXHAUSTION — price extended >3.5 ATR past swing high" + if n == "no_pullback": + return f"NO PULLBACK — price {arg:.1f} ATR from EMA21" + if n == "weak_pullback": + return f"Weak pullback — price {arg:.1f} ATR from EMA21" + if n == "rsi_ext_up": + return f"Extended — RSI {arg:.0f} in uptrend" + if n == "rsi_ext_down": + return f"Extended — RSI {arg:.0f} in downtrend" + if n == "exhaustion_block": + return f"EXHAUSTION BLOCK — {int(arg)} exhaustion signals fired, trade rejected" + if n == "session_filter": + return f"SESSION FILTER: {int(arg):02}:00 UTC blocked" + if n == "day_filter": + day_names = {0: "Mon", 1: "Tue", 2: "Wed", 3: "Thu", 4: "Fri", 5: "Sat", 6: "Sun"} + return f"DAY FILTER: {day_names.get(int(arg), '?')} not in allowed days" + if n == "exhausted_trend": + return "EXHAUSTED TREND FILTER: C-grade blocked in exhausted trend" + if n == "sell_quality": + return f"SELL FILTER: quality {ck.QUALITY_NAMES[int(arg)]} below minimum for SELL" + if n == "mtf_opposition": + return "MTF OPPOSITION: majority of timeframes point against signal — no trade" + if n == "sell_vol": + return f"SELL FILTER: volatility '{ck.VOL_NAMES[int(arg)]}' too quiet for shorts" + if n == "regime_blocked": + return f"REGIME FILTER: '{ck.REGIME_NAMES[int(arg)]}' blocked for XAU" + if n == "sell_regime": + return f"SELL FILTER: regime '{ck.REGIME_NAMES[int(arg)]}' blocked for shorts" + if n == "sell_hour": + return f"SELL FILTER: {int(arg):02}:00 UTC entry blocked for shorts" + if n == "big_candle": + return (f"BIG CANDLE FILTER: {arg:.1f}x ATR bar at {ctx['bc_ts']} within " + f"{ctx['big_candle_lookback_hours']:g}h — no entries") + if n == "m5_score": + return f"M5 FILTER: |score| {arg:.1f} < {ctx['m5_min_abs_score']:.0f} minimum" + if n == "m5_mtf": + return "M5 FILTER: no MTF agreement — M5 signals require confirmation" + if n == "m5_conviction": + return "M5 FILTER: no confirmation TF shows directional conviction" + if n == "m5_sell_quality": + return (f"M5 FILTER: SELL quality {ck.QUALITY_NAMES[int(arg)]} capped — " + f"high-conviction M5 shorts enter at extremes") + if n == "m5_vol": + return (f"M5 FILTER: '{ck.VOL_NAMES[int(arg)]}' volatility — " + f"chop-transition regime, stops get swept") + if n == "m5_sell_rsi": + return (f"M5 FILTER: SELL blocked — RSI {arg:.0f} < " + f"{ctx['m5_sell_min_rsi']:.0f}, dip already extended") + if n == "spread_gate_atr" or n == "spread_gate_generic": + return (f"SPREAD GATE: target ${arg:.2f} < ${ctx['min_tp']:.2f} minimum — " + f"costs would eat the edge") + if n == "low_atr": + return f"LOW ATR FILTER: ATR ${arg:.2f} < ${ctx['min_atr_threshold']:.2f} minimum" + if n == "rf_unavailable": + return f"Regime filter requested but {ctx['rf_tf']} data unavailable" + if n == "rf_no_adx": + return f"Regime filter: ADX could not be computed on {ctx['rf_tf']}" + if n == "rf_blocked": + return (f"REGIME FILTER: {ctx['rf_tf']} ADX={arg:.1f} < " + f"{ctx['regime_filter_min_adx']:.0f} — market choppy, skipping signal") + raise ValueError(f"unknown warning code {code}") + + +def collect_reasons(inds: Dict[str, Dict]) -> Tuple[List[str], List[str]]: + """signal_aggregator._collect_reasons (news never contributes here).""" + bullish: List[str] = [] + bearish: List[str] = [] + for name, ind in inds.items(): + sig = ind["signal"] + label = ck.INDICATOR_LABELS[name] + if sig > 0: + prefix = "STRONG: " if sig == 2 else "" + bullish.append(f"{prefix}{label}: {ind['description']}") + elif sig < 0: + prefix = "STRONG: " if sig == -2 else "" + bearish.append(f"{prefix}{label}: {ind['description']}") + return bullish, bearish + + +# ═════════════════════════════════════════════════════════════════════ +# Host flags: the calendar / session rules of _replay_span, vectorised +# ═════════════════════════════════════════════════════════════════════ + +def host_flags(primary_index: pd.DatetimeIndex, platform_cfg: PlatformConfig, + bt_cfg: BacktestConfig) -> Dict[str, np.ndarray]: + """Per primary bar: the replay skip reason (1 WEEKEND_FILTER, 2 DAY_FILTER, + 3 SESSION_WINDOW, 0 analysed), the bar-label hour/weekday (gate 6b) and + the bar-close hour/weekday/minute/day (gate 6g3, XAU_CP, features).""" + n = len(primary_index) + primary_dur = pd.Timedelta(seconds=BAR_SECONDS[bt_cfg.primary_tf]) + decision = primary_index + primary_dur + skip = np.zeros(n, dtype=np.int8) + is_xau = PAIR_INFO[bt_cfg.pair]["display"] == "XAU/USD" + xau = platform_cfg.xau_strategy + if is_xau: + if xau.weekend_filter: + ny = decision.tz_convert(_NY_TZ) + wd = np.asarray(ny.weekday) + hr = np.asarray(ny.hour) + weekend = ((wd == 4) & (hr >= 17)) | (wd == 5) | ((wd == 6) & (hr < 18)) + skip[weekend] = 1 + if xau.allowed_weekdays: + day = ~np.isin(np.asarray(decision.weekday), list(xau.allowed_weekdays)) + skip[day & (skip == 0)] = 2 + flt = platform_cfg.backtest + if flt.session_filter: + h = np.asarray(decision.hour) + blocked = xau.blocked_hours_utc if is_xau else flt.blocked_hours_utc + sess = ~((h >= flt.session_start_hour) & (h < flt.session_end_hour)) | np.isin(h, list(blocked)) + skip[sess & (skip == 0)] = 3 + return { + "skip_reason": skip, + "label_hour": np.asarray(primary_index.hour, dtype=np.int8), + "label_wd": np.asarray(primary_index.weekday, dtype=np.int8), + "close_hour": np.asarray(decision.hour, dtype=np.int8), + "close_wd": np.asarray(decision.weekday, dtype=np.int8), + "close_minute": np.asarray(decision.minute, dtype=np.int8), + "close_day": np.asarray(decision.day, dtype=np.int8), + } + + +def replay_bounds(primary_df: pd.DataFrame, bt_cfg: BacktestConfig) -> Tuple[int, int]: + bt_start_ts = pd.Timestamp(bt_cfg.start) + bt_end_ts = pd.Timestamp(bt_cfg.end) + first_eligible_pos = max(bt_cfg.warmup_bars, int(primary_df.index.searchsorted(bt_start_ts))) + last_pos = int(primary_df.index.searchsorted(bt_end_ts, side="right")) + last_pos = min(last_pos, len(primary_df)) + return first_eligible_pos, last_pos + + +# ═════════════════════════════════════════════════════════════════════ +# The pipeline +# ═════════════════════════════════════════════════════════════════════ + +class _Run: + """Everything one replay produces on the device plus the host handles + the builders need.""" + + +def refuse_unsupported(platform_cfg: PlatformConfig, bt_cfg: BacktestConfig) -> None: + t = platform_cfg.technical + x = platform_cfg.xau_strategy + problems = [] + if len(platform_cfg.confirmation_timeframes) > ck.MAX_CONFIRM: + problems.append(f"{len(platform_cfg.confirmation_timeframes)} confirmation timeframes " + f"(the CUDA aggregate gathers at most {ck.MAX_CONFIRM})") + if t.divergence_lookback >= ck.MAX_DIV_LOOKBACK: + problems.append(f"technical.divergence_lookback={t.divergence_lookback} " + f"(CUDA limit {ck.MAX_DIV_LOOKBACK - 1})") + if max(t.swing_sr_lookback, 120) >= ck.MAX_SR_LOOKBACK: + problems.append(f"technical.swing_sr_lookback={t.swing_sr_lookback} " + f"(CUDA limit {ck.MAX_SR_LOOKBACK - 1})") + if x.cp_lookback + 1 > ck.MAX_CP_LOOKBACK: + problems.append(f"xau_strategy.cp_lookback={x.cp_lookback} (CUDA limit {ck.MAX_CP_LOOKBACK - 1})") + if x.big_candle_atr_multiple > 0: + bar_secs = BaseDataProvider.CANDLE_SECONDS[platform_cfg.primary_timeframe] + n_bc = max(1, int(x.big_candle_lookback_hours * 3600 // bar_secs)) + if n_bc + 3 * t.atr_period > ck.FEAT_MAX_P: + problems.append(f"big-candle window {n_bc + 3 * t.atr_period} bars (CUDA limit {ck.FEAT_MAX_P})") + if t.ema_mode not in ("4ema", "2ema"): + problems.append(f"technical.ema_mode={t.ema_mode!r}") + if problems: + raise SystemExit("cudaBasedBacktesting.py cannot run this configuration: " + + "; ".join(problems) + ". Use backtest/gpu_backtest.py.") + + +def analyze_frames(bt_cfg: BacktestConfig, platform_cfg: PlatformConfig, + dfs_full: Dict[Timeframe, pd.DataFrame], timeframes: Sequence[Timeframe]) -> _Run: + """Precompute -> K1 per timeframe -> alignment -> host flags -> K2 -> + cooldown, on already loaded (cleaned) frames. Returns the run context + the builders read. Shared by run_backtest_cuda and the parity tests.""" + cp = _gk._require_cp() + flt = platform_cfg.backtest + t0 = time.perf_counter() + caches = {tf: precompute_indicators(dfs_full[tf], platform_cfg.technical, gpu_log=True) for tf in timeframes} + logger.info(f"Indicator pre-compute took {time.perf_counter() - t0:.2f}s (CuPy GPU)") + + pf, pi, rm = ck.pack_params(platform_cfg, bt_cfg) + t1 = time.perf_counter() + dev = {tf: ck.upload_cache(caches[tf], dfs_full[tf]) for tf in timeframes} + an = {tf: ck.analyze_tf(dev[tf], pf, pi, rm, tf.value == "5min") for tf in timeframes} + logger.info(f"K1: {sum(dev[tf]['n'] for tf in timeframes)} bars of " + f"{len(timeframes)} timeframes analysed on the device in {time.perf_counter() - t1:.3f}s") + + primary_df = dfs_full[bt_cfg.primary_tf] + primary_dev = dev[bt_cfg.primary_tf] + primary_secs = BAR_SECONDS[bt_cfg.primary_tf] + cuts = {tf: ck.align_cut(dev[tf]["index_ns"], primary_dev["index_ns"], primary_secs, BAR_SECONDS[tf]) + for tf in timeframes if tf != bt_cfg.primary_tf} + confirms = [(an[tf], cuts[tf]) for tf in bt_cfg.confirmation_tfs] + rf = None + if flt.regime_filter: + rtf = flt.regime_filter_timeframe + rf_cut = cuts[rtf] if rtf != bt_cfg.primary_tf else cp.arange(1, primary_dev["n"] + 1, dtype=cp.int64) + rf = (dev[rtf]["adx"], rf_cut) + flags = host_flags(primary_df.index, platform_cfg, bt_cfg) + first_eligible_pos, last_pos = replay_bounds(primary_df, bt_cfg) + logger.info(f"Replaying {bt_cfg.pair} {bt_cfg.primary_tf.value}: " + f"bars {first_eligible_pos}–{last_pos} of {len(primary_df)}") + n_bars = last_pos - first_eligible_pos + + t_loop = time.perf_counter() + k2 = ck.aggregate(primary_dev, an[bt_cfg.primary_tf], confirms, rf, flags, pf, pi, + first_eligible_pos, last_pos) + logger.info(f"K2: {n_bars} primary bars aggregated on the device in {time.perf_counter() - t_loop:.3f}s") + logger.info("GPU cache fallbacks: 0 (no per-bar recompute path in cudaBasedBacktesting)") + + status = k2.i("status").get()[:n_bars] if n_bars > 0 else np.zeros(0, dtype=np.int32) + cand_pos = (np.flatnonzero(status == 2) + first_eligible_pos).astype(np.int64) + accepted_mask = ck.cooldown(cand_pos, bt_cfg.cooldown_bars) + acc_pos = cand_pos[accepted_mask.astype(bool)] + + skipped_reasons: Dict[str, int] = {} + skip = flags["skip_reason"][first_eligible_pos:last_pos] + order: List[str] = [] + for code in skip: + if code: + name = ck.SKIP_NAMES[int(code)] + if name not in skipped_reasons: + order.append(name) + skipped_reasons[name] = skipped_reasons.get(name, 0) + 1 + skipped_reasons = {name: skipped_reasons[name] for name in order} + + ctx = _Run() + ctx.bt_cfg, ctx.platform_cfg, ctx.timeframes = bt_cfg, platform_cfg, list(timeframes) + ctx.dfs_full, ctx.caches, ctx.dev, ctx.an, ctx.cuts, ctx.k2, ctx.flags = dfs_full, caches, dev, an, cuts, k2, flags + ctx.pf, ctx.pi, ctx.rm = pf, pi, rm + ctx.first_pos, ctx.last_pos, ctx.n_bars = first_eligible_pos, last_pos, n_bars + ctx.status, ctx.cand_pos, ctx.acc_pos, ctx.skipped_reasons = status, cand_pos, acc_pos, skipped_reasons + ctx.t_loop = t_loop + return ctx + + +def run_backtest_cuda(bt_cfg: BacktestConfig, platform_cfg: PlatformConfig) -> Tuple[List[Trade], Dict[Timeframe, pd.DataFrame]]: + """gpu_backtest.run_backtest, with the replay on the device.""" + cp = _gk._require_cp() + if platform_cfg.primary_timeframe != bt_cfg.primary_tf: + raise ValueError( + f"primary_tf mismatch: BacktestConfig={bt_cfg.primary_tf.value} " + f"vs PlatformConfig={platform_cfg.primary_timeframe.value}") + refuse_unsupported(platform_cfg, bt_cfg) + provider = make_data_provider(platform_cfg) + + # ── Load all required timeframes (same files as the reference script) ── + timeframes = [bt_cfg.primary_tf] + list(bt_cfg.confirmation_tfs) + flt = platform_cfg.backtest + if flt.regime_filter and flt.regime_filter_timeframe not in timeframes: + timeframes.append(flt.regime_filter_timeframe) + dfs_full: Dict[Timeframe, pd.DataFrame] = {} + cache_names: Dict[str, str] = {} + tf_starts = {} + from datetime import timedelta + for tf in timeframes: + tf_start = bt_cfg.start - timedelta(days=_history_pad_days(tf, bt_cfg.warmup_bars)) + tf_starts[tf] = tf_start + cache_names[tf.value] = _cache_path(bt_cfg.pair, tf, tf_start, bt_cfg.end, provider.name).name + df = fetch_history(provider, bt_cfg.pair, tf, start=tf_start, end=bt_cfg.end, + use_cache=bt_cfg.use_cache) + if df.empty: + raise RuntimeError(f"No data for {bt_cfg.pair} {tf.value}") + ohlc = df[["Open", "High", "Low", "Close"]].to_numpy(dtype=np.float64) + if not np.isfinite(ohlc).all(): + raise RuntimeError(f"{tf.value}: the cleaned frame carries non-finite OHLC — " + f"the device replay assumes clean bars") + dfs_full[tf] = df + + ctx = analyze_frames(bt_cfg, platform_cfg, dfs_full, timeframes) + primary_df = dfs_full[bt_cfg.primary_tf] + first_eligible_pos, last_pos, n_bars = ctx.first_pos, ctx.last_pos, ctx.n_bars + acc_pos, cand_pos, skipped_reasons, t_loop = ctx.acc_pos, ctx.cand_pos, ctx.skipped_reasons, ctx.t_loop + trades = build_trades(ctx, acc_pos) + signal_ts_list = [primary_df.index[int(p)] for p in acc_pos] + + replay_seconds = time.perf_counter() - t_loop + logger.info( + f"Replay complete: {len(trades)} entries in {replay_seconds:.1f}s " + f"({n_bars / max(replay_seconds, 1e-6):.0f} bars/s, device)") + if skipped_reasons: + logger.info(f"Skipped bars by reason: {skipped_reasons}") + + # ── Resolve trade outcomes on the RAW price path (shared GPU resolver) ── + sim_df = fetch_history(provider, bt_cfg.pair, bt_cfg.primary_tf, start=tf_starts[bt_cfg.primary_tf], + end=bt_cfg.end, use_cache=bt_cfg.use_cache, clean=False) + _reanchor_entries_to_raw(trades, signal_ts_list, sim_df) + if bt_cfg.fixed_rr_tp > 0: + _apply_fixed_rr_tp(trades, bt_cfg.fixed_rr_tp) + logger.info( + f"RESEARCH override: fixed-RR TP = {bt_cfg.fixed_rr_tp} x risk; " + f"trailing exit DISABLED for all {len(trades)} trades (incl. XAU_PB)") + run_arm = 0.0 if bt_cfg.fixed_rr_tp > 0 else platform_cfg.signals.trail_arm_r + t_sim = time.perf_counter() + simulate_trades(trades, sim_df, max_bars=bt_cfg.max_hold_bars, spread=bt_cfg.spread, + trail_arm_r=run_arm, trail_dist_r=platform_cfg.signals.trail_dist_r) + logger.info(f"Trade simulation took {time.perf_counter() - t_sim:.2f}s (CuPy GPU)") + + # ── research flags ── + if bt_cfg.export_features and trades: + t_f = time.perf_counter() + attach_features(ctx, acc_pos, trades) + logger.info(f"K3: {len(trades)} feature vectors on the device in {time.perf_counter() - t_f:.3f}s") + if bt_cfg.record_dir: + t_r = time.perf_counter() + attach_snapshots(ctx, acc_pos, trades) + bar_chunks = build_bar_chunks(ctx) if bt_cfg.record_bars else None + logger.info(f"--record: snapshots{' + bar rows' if bt_cfg.record_bars else ''} built in " + f"{time.perf_counter() - t_r:.2f}s") + replay_record.write_run( + bt_cfg.record_dir, trades=trades, book_df=trades_to_dataframe(trades), + bt_cfg=bt_cfg, platform_cfg=platform_cfg, aggregator=SignalAggregator(platform_cfg), + run_info=dict(first_pos=first_eligible_pos, last_pos=last_pos, + n_primary_bars=len(primary_df), n_bars_in_window=n_bars, + skipped=skipped_reasons, n_candidates=int(len(cand_pos)), + workers=0, backend=BACKEND.describe(), + cache_files=cache_names, replay_seconds=replay_seconds, + cache_fallbacks=0), + sim_df=sim_df, bar_chunks=bar_chunks) + logger.info(f"--record: wrote run to {bt_cfg.record_dir}") + return trades, dfs_full + + +# ═════════════════════════════════════════════════════════════════════ +# Host builders: Trade objects, warnings, snapshots, features, bar rows +# ═════════════════════════════════════════════════════════════════════ + +def _k2_rows(k2: "ck.K2Out", positions: np.ndarray): + """K2 fields + warning codes/args for the given absolute bar positions.""" + idx = np.asarray(positions, dtype=np.int64) - k2.first_pos + if len(idx) == 0: + return [], [], np.zeros((0, ck.MAX_WARN), np.int16), np.zeros((0, ck.MAX_WARN), np.float64) + fi = k2.OF2[:, idx].get() + ii = k2.OI2[:, idx].get() + wc = k2.warn_code[idx].get() + wa = k2.warn_arg[idx].get() + ints = [{n: int(ii[k, m]) for k, n in enumerate(ck.K2_INT_FIELDS)} for m in range(len(idx))] + floats = [{n: float(fi[k, m]) for k, n in enumerate(ck.K2_FLOAT_FIELDS)} for m in range(len(idx))] + return ints, floats, wc, wa + + +def _warning_ctx(ctx, pos: int, r2i: Dict, r2f: Dict) -> Dict: + s = ctx.platform_cfg.signals + x = ctx.platform_cfg.xau_strategy + primary_df = ctx.dfs_full[ctx.bt_cfg.primary_tf] + return { + "min_tp": s.round_trip_cost * s.min_tp_cost_multiple, + "min_atr_threshold": x.min_atr_threshold, + "m5_min_abs_score": x.m5_min_abs_score, + "m5_sell_min_rsi": x.m5_sell_min_rsi, + "big_candle_lookback_hours": x.big_candle_lookback_hours, + "bc_ts": primary_df.index[r2i["bc_idx"]] if r2i["bc_idx"] >= 0 else None, + "rf_tf": ctx.platform_cfg.backtest.regime_filter_timeframe.value, + "regime_filter_min_adx": ctx.platform_cfg.backtest.regime_filter_min_adx, + } + + +_RF_CODES = {ck.WARN_CODES["rf_unavailable"], ck.WARN_CODES["rf_no_adx"], ck.WARN_CODES["rf_blocked"]} + + +def build_warnings(ctx, pos: int, r2i: Dict, r2f: Dict, codes, args) -> Tuple[List[str], List[str]]: + """(warnings, cleaned_warnings) of one analysed bar, after the XAU_CP + cleanup (when a pattern setup was accepted) and the stale spread-gate + step (when a setup existed at step 11), with the regime-filter lines + appended last like the replay does.""" + wctx = _warning_ctx(ctx, pos, r2i, r2f) + n = r2i["warn_n"] + main: List[str] = [] + rf: List[str] = [] + for q in range(n): + code = int(codes[q]) + text = format_warning(code, float(args[q]), wctx) + (rf if code in _RF_CODES else main).append(text) + ns = types.SimpleNamespace(warnings=main, cleaned_warnings=[], + trade_setup=object() if r2i["strategy_pre_rf"] != 0 else None) + if r2i["cp_accepted"]: + SignalAggregator._clean_accepted_cp_warnings(ns) + SignalAggregator._drop_stale_spread_gate_vetoes(ns) + return list(ns.warnings) + rf, list(ns.cleaned_warnings) + + +def build_reasons(ctx, r2i: Dict, r2f: Dict, inds: Dict[str, Dict]) -> Tuple[List[str], List[str]]: + bullish, bearish = collect_reasons(inds) + strat = r2i["strategy"] if r2i["strategy"] else r2i["strategy_pre_rf"] + if r2i["pb_accepted"]: + marker = (f"XAU PULLBACK STRATEGY: retrace {r2f['pb_retrace_atr']:.1f} " + f"ATR to EMA21/KC-mid, trend resumed") + (bullish if r2i["direction"] > 0 else bearish).insert(0, marker) + if r2i["cp_accepted"]: + atr = ctx_atr = None + pattern = ck.PATTERN_ORDER[r2i["cp_pattern"]] + direction = "BUY" if r2i["direction"] > 0 else "SELL" + marker = (f"XAU CHART PATTERN: {pattern} ({direction}) break of {r2f['cp_trigger']:.2f}, " + f"height {r2f['cp_height'] / r2f['_atr']:.1f} ATR") + (bullish if r2i["direction"] > 0 else bearish).insert(0, marker) + if r2i["rf_state"] == 4: + rtf = ctx.platform_cfg.backtest.regime_filter_timeframe.value + line = f"Regime: {rtf} ADX={r2f['rf_adx']:.1f} confirms trending market" + (bullish if r2f["final_score"] > 0 else bearish).append(line) + return bullish, bearish + + +def build_trades(ctx, acc_pos: np.ndarray) -> List[Trade]: + bt_cfg, platform_cfg = ctx.bt_cfg, ctx.platform_cfg + primary_df = ctx.dfs_full[bt_cfg.primary_tf] + pan = ctx.an[bt_cfg.primary_tf] + two = platform_cfg.technical.ema_mode == "2ema" + ints, floats, wc, wa = _k2_rows(ctx.k2, acc_pos) + prows = k1_rows(pan, acc_pos) + pip_value = PAIR_INFO[bt_cfg.pair]["pip_value"] + trades: List[Trade] = [] + for m, pos in enumerate(acc_pos): + pos = int(pos) + r2i, r2f, prow = ints[m], floats[m], prows[m] + r2f["_atr"] = prow.f["atr"] if prow.i["atr_present"] else 0.0 + direction = "BUY" if r2i["direction"] > 0 else "SELL" + tag = ck.STRATEGY_NAMES[r2i["strategy"]] or "GENERIC" + inds = describe_indicators(prow, two) + # gpu_backtest._build_indicator_breakdown + breakdown: Dict[str, Dict] = {} + agree = oppose = neutral = 0 + is_buy = direction == "BUY" + for name, ind in inds.items(): + sig = ind["signal"] + agrees = "neutral" + if sig > 0: + agrees = "agree" if is_buy else "oppose" + elif sig < 0: + agrees = "agree" if not is_buy else "oppose" + breakdown[name] = { + "value": float(ind["value"]) if ind["value"] is not None else None, + "signal": ck.SIGNAL_NAMES[sig], + "confidence": float(ind["confidence"]), + "weight": float(ind["weight"]), + "agrees": agrees, + "description": ind["description"], + } + if agrees == "agree": + agree += 1 + elif agrees == "oppose": + oppose += 1 + else: + neutral += 1 + total = len(breakdown) + bullish, bearish = build_reasons(ctx, r2i, r2f, inds) + reasons = bullish if direction == "BUY" else bearish + warnings, _cleaned = build_warnings(ctx, pos, r2i, r2f, wc[m], wa[m]) + # the setup's own pip value: 0.10 hard-coded by the XAU ATR generator, PAIR_INFO otherwise + setup_pip = 0.10 if tag == "XAU_ATR" else pip_value + # the setup's own trailing parameters (XAU_PB carries the pullback pair) + trail_arm = platform_cfg.signals.pb_trail_arm_r if tag == "XAU_PB" else 0.0 + trail_dist = platform_cfg.signals.pb_trail_dist_r if tag == "XAU_PB" else 0.0 + trades.append(Trade( + pair=bt_cfg.pair, + entry_time=primary_df.index[pos + 1], + entry_price=r2f["entry_price"], + sl=r2f["adj_sl"], tp1=r2f["adj_tp1"], tp2=r2f["adj_tp2"], + direction=direction, + pip_value=setup_pip, + risk_pips=r2f["sl_dist"] / setup_pip, + reward_pips=r2f["tp1_dist"] / setup_pip, + rr_planned=r2f["rr_rounded"], + quality=ck.QUALITY_NAMES[r2i["quality"]], + final_score=r2f["final_score"], + final_confidence=r2f["final_conf"], + market_regime=ck.REGIME_NAMES[prow.i["regime"]], + volatility_regime=ck.VOL_NAMES[prow.i["vol"]], + strategy_tag=tag, + entry_reasons=list(reasons), + entry_warnings=list(warnings), + indicator_breakdown=breakdown, + indicators_agreeing=agree, indicators_opposing=oppose, + indicators_neutral=neutral, indicators_total=total, + stake_multiplier=1.0, + trail_arm_r=trail_arm, trail_dist_r=trail_dist, + )) + return trades + + +def attach_features(ctx, acc_pos: np.ndarray, trades: List[Trade]) -> None: + bt_cfg = ctx.bt_cfg + confirms = [(ctx.an[tf], ctx.cuts[tf]) for tf in bt_cfg.confirmation_tfs] + # f_h1_ext_atr: the H1 frame wherever it sits; f_c_other_ext_atr: the first + # non-H1 confirmation with >= 21 bars (evaluated per bar on the device). + if bt_cfg.primary_tf == Timeframe.H1: + h1 = (1, None, None) + elif Timeframe.H1 in ctx.dfs_full and Timeframe.H1 in ctx.cuts: + h1 = (2, ctx.dev[Timeframe.H1], ctx.cuts[Timeframe.H1]) + else: + h1 = (0, None, None) + others = [(ctx.dev[tf], ctx.cuts[tf]) for tf in bt_cfg.confirmation_tfs if tf is not Timeframe.H1][:2] + F = ck.features(acc_pos, ctx.dev[bt_cfg.primary_tf], ctx.an[bt_cfg.primary_tf], ctx.k2, + len(bt_cfg.confirmation_tfs), confirms, h1, others, ctx.flags, ctx.pf, ctx.pi) + for t, row in zip(trades, F): + t.features = {name: float(row[k]) for k, name in enumerate(position_sizing.FEATURE_ORDER)} + + +def _tf_block(row: K1Row, tf_name: str, role: str, direction: str, two: bool) -> Dict: + inds = describe_indicators(row, two) + is_buy = direction == "BUY" + out_inds: Dict[str, Dict] = {} + for name, ind in inds.items(): + sig = ind["signal"] + base = float(ind["weight"]) + eff = replay_record._finite(ind["effective_weight"]) + out_inds[name] = { + "signal": sig, + "signal_name": ck.SIGNAL_NAMES[sig], + "value": replay_record._finite(ind["value"]), + "confidence": replay_record._finite(ind["confidence"]), + "base_weight": base, + "regime_mult": (eff / base) if (eff is not None and base) and base > 0 else None, + "eff_weight": eff, + "wscore": replay_record._finite(ind["weighted_score"]), + "agrees": replay_record._agrees(sig, direction), + "description": str(ind["description"] or ""), + } + kl = key_levels(row, two) + return { + "tf": tf_name, + "role": role, + "regime": ck.REGIME_NAMES[row.i["regime"]], + "vol": ck.VOL_NAMES[row.i["vol"]], + "composite_score": replay_record._finite(row.f["composite_score"]), + "composite_signal": int(row.i["composite_signal"]), + "confidence": replay_record._finite(row.f["composite_conf"]), + "atr": replay_record._finite(kl.get("atr")), + "atr_pct": replay_record._finite(kl.get("atr_pct")), + "pattern_signals": pattern_signals(row), + "indicators": out_inds, + } + + +def attach_snapshots(ctx, acc_pos: np.ndarray, trades: List[Trade]) -> None: + """replay_record.snapshot_report for every accepted trade, from the device outputs.""" + bt_cfg, platform_cfg = ctx.bt_cfg, ctx.platform_cfg + primary_df = ctx.dfs_full[bt_cfg.primary_tf] + primary_dur = pd.Timedelta(seconds=BAR_SECONDS[bt_cfg.primary_tf]) + two = platform_cfg.technical.ema_mode == "2ema" + pdev = ctx.dev[bt_cfg.primary_tf] + medians = ck.prefix_median(pdev["high"], pdev["low"], pdev["n"])[acc_pos].get() if len(acc_pos) else [] + ints, floats, wc, wa = _k2_rows(ctx.k2, acc_pos) + prows = k1_rows(ctx.an[bt_cfg.primary_tf], acc_pos) + conf_rows = {} + conf_cuts = {} + for tf in bt_cfg.confirmation_tfs: + cut = ctx.cuts[tf][acc_pos].get() if len(acc_pos) else np.zeros(0, np.int64) + conf_cuts[tf] = cut + present = cut > 0 + rows = k1_rows(ctx.an[tf], np.where(present, cut - 1, 0)) + conf_rows[tf] = rows + s = platform_cfg.signals + for m, (pos, t) in enumerate(zip(acc_pos, trades)): + pos = int(pos) + r2i, r2f, prow = ints[m], floats[m], prows[m] + r2f["_atr"] = prow.f["atr"] if prow.i["atr_present"] else 0.0 + direction = t.direction + inds = describe_indicators(prow, two) + bullish, bearish = build_reasons(ctx, r2i, r2f, inds) + warnings, cleaned = build_warnings(ctx, pos, r2i, r2f, wc[m], wa[m]) + tag = t.strategy_tag + pmask = r2i["penalty_mask"] + penalty_tags = {tagname: float(ck.PENALTY_AMOUNT[tagname]) for tagname in ck.PENALTY_SORTED + if pmask & (1 << ck.PENALTY_TAGS.index(tagname))} + mtf_signals = {bt_cfg.primary_tf.value: int(prow.i["composite_signal"])} + mtf_scores = {bt_cfg.primary_tf.value: replay_record._finite(prow.f["composite_score"])} + tfs = [_tf_block(prow, bt_cfg.primary_tf.value, "primary", direction, two)] + for k, tf in enumerate(bt_cfg.confirmation_tfs): + if conf_cuts[tf][m] == 0: + continue + crow = conf_rows[tf][m] + mtf_signals[tf.value] = int(crow.i["composite_signal"]) + mtf_scores[tf.value] = replay_record._finite(crow.f["composite_score"]) + if k + 1 < len(replay_record.ROLES): + tfs.append(_tf_block(crow, tf.value, replay_record.ROLES[k + 1], direction, two)) + trail_arm = s.pb_trail_arm_r if tag == "XAU_PB" else s.trail_arm_r + trail_dist = s.pb_trail_dist_r if tag == "XAU_PB" else s.trail_dist_r + t.snapshot = { + "bar_ts": str(primary_df.index[pos]), + "decision_time": pd.Timestamp(primary_df.index[pos] + primary_dur).isoformat(), + "final_signal": int(r2i["final_signal"]), + "final_score": replay_record._finite(r2f["final_score"]), + "final_confidence": replay_record._finite(r2f["final_conf"]), + "quality": ck.QUALITY_NAMES[r2i["quality"]], + "mtf_agreement": bool(r2i["mtf_agreement"]), + "mtf_signals": mtf_signals, + "mtf_scores": mtf_scores, + "penalty_tags": penalty_tags, + "penalty_total": replay_record._finite(r2f["penalty_total"]), + "momentum_fading": bool(r2i["momentum_fading"]), + "exhaustion_count": int(r2i["exhaustion_count"]), + "spread_estimate": replay_record._finite(float(medians[m])), + "current_price": replay_record._finite(float(primary_df["Close"].iloc[pos])), + "warnings": warnings, + "cleaned_warnings": cleaned, + "bullish_reasons": bullish, + "bearish_reasons": bearish, + "setup": { + "strategy": tag, + "setup_pattern": ck.PATTERN_ORDER[r2i["cp_pattern"]] if tag == "XAU_CP" else None, + "round_trip_cost": replay_record._finite(s.round_trip_cost), + "trail_arm_r": replay_record._finite(trail_arm), + "trail_dist_r": replay_record._finite(trail_dist), + "entry_price": replay_record._finite(r2f["setup_entry"]), + "stop_loss": replay_record._finite(r2f["setup_sl"]), + "take_profit_1": replay_record._finite(r2f["setup_tp1"]), + "take_profit_2": replay_record._finite(r2f["setup_tp2"]), + "risk_reward": replay_record._finite(r2f["rr_rounded"]), + }, + "tfs": tfs, + } + + +def build_bar_chunks(ctx): + """replay_record's (matrix float32 [n,160], ts_ns, pos, warning0) for + every analysed primary bar (the K6 rows plus the host warning text).""" + bt_cfg = ctx.bt_cfg + k2 = ctx.k2 + n_bars = k2.n_bars + if n_bars <= 0: + return None + status = k2.i("status").get()[:n_bars] + bar_pos = (np.flatnonzero(status >= 1) + k2.first_pos).astype(np.int64) + if len(bar_pos) == 0: + return None + confirms = [(ctx.an[tf], ctx.cuts[tf]) for tf in bt_cfg.confirmation_tfs] + mat = ck.bar_rows(bar_pos, ctx.dev[bt_cfg.primary_tf], ctx.an[bt_cfg.primary_tf], k2, + len(bt_cfg.confirmation_tfs), confirms) + ints, floats, wc, wa = _k2_rows(k2, bar_pos) + prows_atr = None + warning0: List[str] = [] + prim = ctx.an[bt_cfg.primary_tf] + atr_present = prim.i("atr_present")[bar_pos].get() + atr_vals = prim.f("atr")[bar_pos].get() + for m, pos in enumerate(bar_pos): + r2i, r2f = ints[m], floats[m] + if r2i["warn_n"] == 0: + warning0.append("") + continue + r2f["_atr"] = float(atr_vals[m]) if atr_present[m] else 0.0 + warnings, _ = build_warnings(ctx, int(pos), r2i, r2f, wc[m], wa[m]) + warning0.append(warnings[0] if warnings else "") + primary_index = ctx.dfs_full[bt_cfg.primary_tf].index + return [(mat, primary_index.as_unit("ns").asi8[bar_pos].astype(np.int64), bar_pos.astype(np.int64), warning0)] + + +# ═════════════════════════════════════════════════════════════════════ +# CLI (the reference script's parser, config block and output, verbatim) +# ═════════════════════════════════════════════════════════════════════ + +def require_gpu_or_exit(requested_backend: Optional[str]) -> None: + if requested_backend == "cpu" or BACKEND.mode == "cpu": + raise SystemExit( + "cudaBasedBacktesting.py runs the decision stage on CUDA and has no CPU path " + "(--backend cpu / FB2_BACKEND=cpu requested). Use backtest/gpu_backtest.py for the CPU twins.") + if requested_backend: + BACKEND.set_mode(requested_backend) + if not BACKEND.available: + raise SystemExit( + f"cudaBasedBacktesting.py needs a working CUDA backend: {BACKEND.describe()}. " + f"Use backtest/gpu_backtest.py.") + + +def main(argv: Optional[List[str]] = None) -> int: + args = _parse_args(argv) + logging.basicConfig( + level=args.log_level.upper(), + format="%(asctime)s [%(levelname)s] %(name)s: %(message)s", + ) + require_gpu_or_exit(args.backend) + + pair = _resolve_pair(args.pair) + primary_tf = _resolve_tf(args.primary_tf) + confirm_tfs = ( + [_resolve_tf(s.strip()) for s in args.confirm_tfs.split(",") if s.strip()] + if args.confirm_tfs is not None else None + ) + from pair_signals import PAIR_PRODUCERS # noqa: E402 + factory = PAIR_PRODUCERS.get(pair) + if factory is not None: + platform_cfg = factory(primary_tf, confirm_tfs) + else: + platform_cfg = PlatformConfig( + symbol=pair, + display_name=PAIR_INFO[pair]["display"], + primary_timeframe=primary_tf, + confirmation_timeframes=confirm_tfs or [Timeframe.H1, Timeframe.H4], + ) + from app_config import apply_pair_overrides # noqa: E402 + apply_pair_overrides(platform_cfg) + if args.fixed_rr_tp is not None: + if args.fixed_rr_tp <= 0: + raise SystemExit("--fixed-rr-tp must be > 0") + if (args.trail_arm is not None or args.trail_dist is not None or + any(s.split("=", 1)[0].strip().startswith(("signals.trail", "signals.pb_trail")) + for s in (args.set or []))): + raise SystemExit( + "--fixed-rr-tp conflicts with --trail-arm/--trail-dist/" + "--set signals.*trail* (the flag disables trailing entirely)") + if args.swing_sr_lookback > 0: + platform_cfg.technical.swing_sr_lookback = args.swing_sr_lookback + if args.trail_arm is not None: + platform_cfg.signals.trail_arm_r = args.trail_arm + if args.trail_dist is not None: + platform_cfg.signals.trail_dist_r = args.trail_dist + if args.cp: + platform_cfg.xau_strategy.cp_enabled = True + if args.cp_patterns: + platform_cfg.xau_strategy.cp_patterns = [ + s.strip() for s in args.cp_patterns.split(",") if s.strip() + ] + if args.big_candle is not None: + platform_cfg.xau_strategy.big_candle_atr_multiple = args.big_candle + if args.big_candle_hours is not None: + platform_cfg.xau_strategy.big_candle_lookback_hours = args.big_candle_hours + if args.set: + _apply_config_overrides(platform_cfg, args.set) + confirm_tfs = list(platform_cfg.confirmation_timeframes) + + from datetime import datetime, timedelta, timezone + end = ( + datetime.fromisoformat(args.end).replace(tzinfo=timezone.utc) + if args.end else datetime.now(timezone.utc) + ) + if args.start: + start = datetime.fromisoformat(args.start).replace(tzinfo=timezone.utc) + else: + start = end - timedelta(days=args.days) + spread = (args.spread if args.spread is not None + else PAIR_INFO[pair].get("default_spread", 0.30)) + bt_cfg = BacktestConfig( + pair=pair, start=start, end=end, primary_tf=primary_tf, confirmation_tfs=confirm_tfs, + warmup_bars=args.warmup_bars, cooldown_bars=args.cooldown_bars, max_hold_bars=args.max_hold_bars, + use_cache=not args.no_cache, + quality_filter=[q.strip() for q in args.quality.split(",") if q.strip()], + strategies_only=[("XAU_ATR" if s.strip() == "XAU" else s.strip()) + for s in args.strategies.split(",") if s.strip()], + spread=spread, workers=args.workers, fixed_rr_tp=args.fixed_rr_tp or 0.0, + export_features=args.export_features, record_dir=args.record, record_bars=args.record_bars, + ) + if args.workers: + logger.info(f"--workers {args.workers} is ignored: the replay runs on the device") + if args.record: + logger.info( + f"--record ON: decision snapshots → {args.record} " + f"(trades/indicators/events{'/bars' if args.record_bars else ''} parquet + manifest.json; " + f"the trade book is unchanged)") + if args.export_features: + logger.info( + f"--export-features ON: appending stake_multiplier, r_sized and " + f"{len(position_sizing.FEATURE_ORDER)} f_* feature columns to the " + f"CSV export (research flag; book itself is unchanged)") + logger.info( + f"Backend: {BACKEND.describe()} | " + f"Pair: {PAIR_INFO[pair]['display']} | Window: {start.date()} → {end.date()} | " + f"Confirm TFs: {','.join(tf.value for tf in confirm_tfs)} | Spread: {spread}" + ) + + trades, _ = run_backtest_cuda(bt_cfg, platform_cfg) + metrics = compute_metrics(trades) + print_report(pair, bt_cfg, metrics, trades=trades, detail=args.detail) + print_trade_reasoning(trades, limit=args.reasoning_samples) + if args.export_csv: + df = trades_to_dataframe(trades, include_features=args.export_features) + df.to_csv(args.export_csv, index=False) + logger.info(f"Exported {len(df)} trades → {args.export_csv}") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/backtest/cuda_signal_kernels.py b/backtest/cuda_signal_kernels.py new file mode 100644 index 0000000..a2f2f7c --- /dev/null +++ b/backtest/cuda_signal_kernels.py @@ -0,0 +1,2914 @@ +""" +backtest/cuda_signal_kernels.py — the CUDA decision stage of cudaBasedBacktesting.py. + +Every kernel here is a line-by-line twin of Python code that already exists in +the repo and stays the reference: + + K1 fb2_analyze_tf GPUTechnicalEngine's nine cached-indicator overrides + (backtest/gpu_backtest.py) plus the live engine code + those overrides do not cover (engine.py: divergence, + structure S/R, candlestick, regime, regime-aware RSI, + composite). One thread per bar of one timeframe; the + thread analyses the prefix [0..j]. + K2 fb2_aggregate signal_aggregator.analyze_from_dataframes end to end + (MTF merge, penalties, grade, gates 6b-6i, the three + setup generators, XAU_PB, XAU_CP, the stale-veto step) + plus the per-bar rules of gpu_backtest._replay_span. + One thread per primary bar. + K3 fb2_features position_sizing.build_features (65 features), one + thread per accepted trade (--export-features). + K4 fb2_cooldown gpu_backtest._apply_cooldown (single thread). + K5 fb2_prefix_median pandas Series.median of High-Low over every prefix + (the snapshot's spread_estimate, --record only). + K6 fb2_bar_rows replay_record.add_bar_row's 160-column float32 row + (--record-bars). + +Text is never built on the device: the kernels emit small variant codes plus +the numbers each template needs, and cudaBasedBacktesting.py applies the same +f-string templates as engine.py / signal_aggregator.py on the host. Because +the doubles are bit-identical the strings are identical. + +Rules that keep the twins bit-identical (see the plan and the tests in +test_all.py Module 37): compiled with --fmad=false; doubles everywhere; Python +min()/max() written as ternaries in argument order (never fmin/fmax); every +sum in the Python iteration order; numpy's pairwise summation replicated for +np.mean / np.nanmean; pandas' online rolling mean and ewm algorithms copied +verbatim (the rolling helpers are taken from gpu_kernels._ROLL_SRC at import +time); Python round(x, 2) reproduced exactly via an error-free product. + +This module imports no live writer and never imports engine, signal_aggregator, +gpu_backtest or replay_record (guard test): the name lists below are the only +contract between the CUDA source and the host, and the host generates the +#defines from them so a name cannot drift. +""" + +from __future__ import annotations + +import logging +import math +import os +from typing import Dict, List, Sequence, Tuple + +import numpy as np + +from backtest import gpu_kernels as _gk +from config import PAIR_INFO, Timeframe +from chart_patterns import PATTERN_DETECTORS + +logger = logging.getLogger("cuda_signal_kernels") + +BACKEND = _gk.BACKEND +KERNEL_OPTIONS = _gk.KERNEL_OPTIONS + +# ═════════════════════════════════════════════════════════════════════ +# Compile-time limits (thread-local array sizes). A config beyond them is +# refused by cudaBasedBacktesting.py before any kernel runs. +# ═════════════════════════════════════════════════════════════════════ + +MAX_CONFIRM = 4 # confirmation timeframes gathered by K2 +MAX_DIV_LOOKBACK = 64 # technical.divergence_lookback <= 63 +MAX_SR_LOOKBACK = 256 # swing lookback <= 255 (120 on M5 by rule) +MAX_CP_LOOKBACK = 128 # xau_strategy.cp_lookback + 1 <= 128 +MAX_WARN = 32 # warnings per bar +FEAT_MAX_P = 500 # position_sizing.MAX_LOOKBACK_BARS +FEAT_CONFIRM_TAIL = 60 # position_sizing.CONFIRM_LOOKBACK_BARS +FEAT_BURST_BARS = 200 # position_sizing.BURST_LOOKBACK_BARS +N_FEATURES = 65 +N_BAR_COLS = 160 + +# ═════════════════════════════════════════════════════════════════════ +# Codes shared with the host (and, where noted, with replay_record) +# ═════════════════════════════════════════════════════════════════════ + +# replay_record.REGIME_CODES / VOL_CODES / QUALITY_CODES / STRATEGY_CODES +REGIME_CODES = {"unknown": 0, "trending_up": 1, "trending_down": 2, + "trending": 3, "transitioning": 4, "ranging": 5} +REGIME_NAMES = {v: k for k, v in REGIME_CODES.items()} +VOL_CODES = {"low": 0, "normal": 1, "high": 2, "extreme": 3} +VOL_NAMES = {v: k for k, v in VOL_CODES.items()} +QUALITY_CODES = {"F": 0, "D": 1, "C": 2, "B": 3, "A": 4, "A+": 5} +QUALITY_NAMES = {v: k for k, v in QUALITY_CODES.items()} +STRATEGY_CODES = {"": 0, "GENERIC": 1, "XAU_ATR": 2, "XAU_PB": 3, "XAU_CP": 4} +STRATEGY_NAMES = {v: k for k, v in STRATEGY_CODES.items()} +SIGNAL_NAMES = {2: "STRONG_BUY", 1: "BUY", 0: "NEUTRAL", -1: "SELL", -2: "STRONG_SELL"} + +# Indicator columns of the regime table and of the eff/wscore outputs, in the +# engine's insertion order minus atr (weight 0). +INDICATOR_ORDER = ("ema_trend", "rsi", "bollinger", "adx", "macd", "roc", + "keltner", "divergence", "candlestick") +INDICATOR_LABELS = {"ema_trend": "EMA Trend", "rsi": "RSI", "bollinger": "Bollinger Bands", + "adx": "ADX", "macd": "MACD", "roc": "ROC", "keltner": "Keltner", + "divergence": "Divergence", "candlestick": "Candlestick"} +BASE_WEIGHTS = {"ema_trend": 1.0, "rsi": 1.5, "bollinger": 1.5, "adx": 0.5, "macd": 1.5, + "roc": 1.5, "keltner": 1.5, "divergence": 1.5, "candlestick": 1.0} +# engine.analyze's disabled_indicators names -> bit +DISABLED_BITS = {"atr": 0, "ema_trend": 1, "rsi": 2, "bollinger": 3, "adx": 4, "macd": 5, + "roc": 6, "keltner": 7, "divergence": 8, "structure": 9, "candlestick": 10} +PATTERN_ORDER = tuple(PATTERN_DETECTORS.keys()) # priority order of detect_chart_pattern + +# Penalty tags in APPEND order (signal_aggregator._compute_final_signal). +PENALTY_TAGS = ("mtf_disagreement", "counter_trend", "extreme_vol", "ema200_near", + "ema200_far", "exhaustion", "no_pullback", "weak_pullback", + "rsi_extended", "bb_overextended") +PENALTY_AMOUNT = {"mtf_disagreement": 15, "counter_trend": 15, "extreme_vol": 10, + "ema200_near": 5, "ema200_far": 5, "exhaustion": 8, "no_pullback": 8, + "weak_pullback": 5, "rsi_extended": 5, "bb_overextended": 8} +# The order after `penalties.sort(key=amount, reverse=True)` (stable): buckets +# 15, 10, 8, 5 each in append order. This is the decay order. +PENALTY_SORTED = tuple(sorted(PENALTY_TAGS, key=lambda t: -PENALTY_AMOUNT[t])) + +# Warning codes (K2 -> host). The host holds the templates. +WARN_NAMES = ( + "none", # 0 + "conflict", # 1 Timeframes CONFLICT + "counter_trend", # 2 + "extreme_vol", # 3 + "low_vol", # 4 + "ema200_near", # 5 arg ema_dist_pct + "ema200_far", # 6 arg ema_dist_pct + "exhaustion_low", # 7 + "exhaustion_high", # 8 + "no_pullback", # 9 arg pullback_dist + "weak_pullback", # 10 arg pullback_dist + "rsi_ext_up", # 11 arg rsi + "rsi_ext_down", # 12 arg rsi + "exhaustion_block", # 13 arg count + "session_filter", # 14 arg label hour + "day_filter", # 15 arg label weekday + "exhausted_trend", # 16 + "sell_quality", # 17 arg quality code + "mtf_opposition", # 18 + "sell_vol", # 19 arg vol code + "regime_blocked", # 20 arg regime code + "sell_regime", # 21 arg regime code + "sell_hour", # 22 arg close hour + "big_candle", # 23 arg ratio (bc_idx in K2 ints) + "m5_score", # 24 arg abs_score + "m5_mtf", # 25 + "m5_conviction", # 26 + "m5_sell_quality", # 27 arg quality code + "m5_vol", # 28 arg vol code + "m5_sell_rsi", # 29 arg rsi + "spread_gate_atr", # 30 arg tp_distance + "low_atr", # 31 arg atr + "spread_gate_generic", # 32 arg tp1_distance + "rf_unavailable", # 33 regime filter: TF missing + "rf_no_adx", # 34 regime filter: ADX None + "rf_blocked", # 35 arg adx +) +WARN_CODES = {n: i for i, n in enumerate(WARN_NAMES)} + +# ── K1 outputs (one column per bar of the timeframe) ───────────────── +K1_INT_FIELDS = ( + "valid", "atr_present", "regime", "vol", "composite_signal", + "ema_present", "ema_signal", "ema_align", "ema_max_checks", "ema_cross", "ema_ribbon", + "ema_trend_present", "ema2_desc", "kl_ema_fast_present", "kl_ema_present", + "rsi_present", "rsi_signal", "rsi_desc", + "macd_present", "macd_signal", "macd_desc", + "bb_present", "bb_signal", "bb_desc", "bb_squeeze", + "adx_present", "adx_signal", "adx_strength", "adx_dir", "adx_slope_desc", + "roc_present", "roc_signal", "roc_desc", + "kc_present", "kc_signal", "kc_desc", + "div_present", "div_signal", "div_mask", + "cs_present", "cs_signal", "cs_mask", "cs_bull", "cs_bear", + "sh_mask", "sl_mask", "sh_fallback", "sl_fallback", +) +K1_FLOAT_FIELDS = ( + "composite_score", "composite_conf", "atr", "atr_pct", "atr_percentile", + "ema_fast", "ema_medium", "ema_slow", "ema_trend", "ema_value", "ema_conf", + "rsi_value", "rsi_conf", + "macd_value", "macd_conf", + "bb_upper", "bb_lower", "bb_mid", "bb_value", "bb_conf", + "adx_value", "plus_di", "minus_di", "adx_slope", "adx_conf", + "roc_value", "roc_conf", + "kc_upper", "kc_lower", "kc_mid", "kc_value", "kc_conf", + "div_value", "cs_value", "cs_conf", + "sh1", "sh2", "sh3", "sl1", "sl2", "sl3", + "ew_ema", "ws_ema", "ew_rsi", "ws_rsi", "ew_bb", "ws_bb", "ew_adx", "ws_adx", + "ew_macd", "ws_macd", "ew_roc", "ws_roc", "ew_kc", "ws_kc", "ew_div", "ws_div", + "ew_cs", "ws_cs", +) +# ── K2 outputs (one column per primary bar in [first_pos, last_pos)) ── +K2_INT_FIELDS = ( + "status", "skip_reason", "mtf_agreement", "final_signal", "quality", + "exhaustion_count", "momentum_fading", "penalty_mask", "strategy", "direction", + "cp_pattern", "cp_accepted", "pb_accepted", "bc_idx", "rf_state", "warn_n", + "n_confirm_present", "strategy_pre_rf", +) +K2_FLOAT_FIELDS = ( + "primary_score", "primary_conf", "final_score", "final_conf", "penalty_total", + "setup_entry", "setup_sl", "setup_tp1", "setup_tp2", "rr_rounded", + "sl_dist", "tp1_dist", "tp2_dist", "entry_price", "adj_sl", "adj_tp1", "adj_tp2", + "cp_trigger", "cp_height", "pb_retrace_atr", "bc_ratio", "rf_adx", +) +SKIP_NAMES = {1: "WEEKEND_FILTER", 2: "DAY_FILTER", 3: "SESSION_WINDOW"} + +# ── Config packing ──────────────────────────────────────────────────── +PARAMS_F = ( + "rsi_overbought", "rsi_oversold", "rsi_neutral_upper", "rsi_neutral_lower", + "adx_strong_trend", "adx_very_strong", "adx_slope_emerging_threshold", "roc_threshold", + "min_rr_ratio", "tp_tightening_factor", "round_trip_cost", "min_tp_cost_multiple", + "min_sl_pips", "tp_floor_rr", + "tp_atr_multiplier", "sl_atr_multiplier", "min_atr_threshold", "v2_sl_atr_multiplier", + "m5_min_abs_score", "m5_min_confirmation_score", "m5_sell_min_rsi", + "m5_pullback_min_score", "m5_pullback_adx_min", + "m5_pullback_rsi_buy_lo", "m5_pullback_rsi_buy_hi", + "m5_pullback_rsi_sell_lo", "m5_pullback_rsi_sell_hi", + "m5_pullback_retrace_lo", "m5_pullback_retrace_hi", + "cp_tolerance_atr", "cp_min_height_atr", "cp_max_height_atr", "cp_hs_head_min_atr", + "cp_hs_shoulder_tol_atr", "cp_flag_pole_atr", "cp_flag_retrace_lo", "cp_flag_retrace_hi", + "big_candle_atr_multiple", "big_candle_lookback_hours", + "regime_filter_min_adx", "pip_value", "fb2_ema_lo", + "tp_scale_aplus", "tp_scale_a", "max_penalty", +) +PARAMS_I = ( + "ema_trend_period", "ema_mode_2ema", "ema_partial_neutral", "rsi_regime_aware", + "macd_as_pillar", "roc_period", "kc_period", "kc_atr_period", "divergence_lookback", + "swing_sr_lookback", "atr_period", "disabled_mask", + "mtf_agreement_threshold", "xau_enabled", "is_xau", "session_filter", + "xau_blocked_hours_mask", "bt_blocked_hours_mask", + "xau_allowed_weekdays_mask", "bt_allowed_weekdays_mask", + "sell_min_quality_mask", "sell_vol_mask", "blocked_regimes_mask", + "sell_blocked_regimes_mask", "sell_blocked_hours_mask", + "m5_require_mtf", "m5_sell_max_quality_mask", "m5_blocked_vol_mask", + "m5_pullback_enabled", "m5_pullback_touch_bars", + "cp_enabled", "cp_patterns_mask", "cp_lookback", "cp_swing_window", "cp_recency_bars", + "cp_min_sep_bars", "cp_flag_pole_bars", "cp_flag_min_bars", "cp_flag_max_bars", + "quality_filter_mask", "strategies_mask", "primary_is_m5", "primary_bar_secs", + "regime_filter", "n_confirm", +) + +# Fields of the four config dataclasses the kernels do NOT read, with the +# reason (checked by a Module 37 test so a new field cannot slip past unread). +CONFIG_FIELDS_HOST_ONLY = { + "TechnicalConfig": { + # precompute-only periods (the cached series carry them) + "ema_fast", "ema_medium", "ema_slow", "rsi_period", "macd_fast", "macd_slow", + "macd_signal", "bb_period", "bb_std", "adx_period", "kc_multiplier", + "adx_slope_lookback", + # packed by name below (listed here because the packer reads them + # under another name) + "disabled_indicators", "ema_mode", "rsi_mode", "regime_weights_override", + "ema_trend", + }, + "SignalConfig": { + # trade resolution / Trade stamps (host) + "trail_arm_r", "trail_dist_r", "pb_trail_arm_r", "pb_trail_dist_r", + }, + "XauStrategyConfig": { + # host calendar/session flags + "weekend_filter", "allowed_weekdays", "blocked_hours_utc", + # packed as masks / pairs under other names + "sell_min_quality", "sell_vol_regimes", "blocked_regimes", "sell_blocked_regimes", + "sell_blocked_hours_utc", "m5_sell_max_quality", "m5_blocked_vol_regimes", + "m5_pullback_rsi_buy", "m5_pullback_rsi_sell", "m5_pullback_retrace_atr", + "cp_patterns", "cp_flag_retrace", "enabled", + }, + "FilterConfig": { + # host calendar/session flags / packed under other names + "session_start_hour", "session_end_hour", "blocked_hours_utc", "allowed_weekdays", + "regime_filter_timeframe", + }, +} + + +def _mask(values, table) -> int: + m = 0 + for v in values or (): + code = table[v] + m |= (1 << int(code)) + return m + + +def _hours_mask(hours) -> int: + m = 0 + for h in hours or (): + h = int(h) + if 0 <= h < 24: + m |= (1 << h) + return m + + +def fb2_ema_lo_from_env() -> float: + """The FB2_EMA_LO research knob, read once (the CPU path reads it per bar).""" + raw = os.environ.get("FB2_EMA_LO", "0.5") + try: + return float(raw) + except (TypeError, ValueError): + logger.warning("FB2_EMA_LO=%r is not a number — using tuned default 0.5", raw) + return 0.5 + + +def regime_table(platform_cfg) -> np.ndarray: + """The 6 x 9 regime multiplier table (rows = REGIME_CODES, cols = + INDICATOR_ORDER; the 'unknown' row is all 1.0 like an empty dict).""" + from engine import regime_weight_table # local: keep the module import-light + table = regime_weight_table(platform_cfg.technical) + out = np.ones((6, len(INDICATOR_ORDER)), dtype=np.float64) + for regime, code in REGIME_CODES.items(): + weights = table.get(regime) + if not weights: + continue + for col, name in enumerate(INDICATOR_ORDER): + out[code, col] = float(weights.get(name, 1.0)) + return out + + +def pack_params(platform_cfg, bt_cfg) -> Tuple[np.ndarray, np.ndarray, np.ndarray]: + """(params_f float64[NF], params_i int32[NI], regime_mult float64[6*9]).""" + t = platform_cfg.technical + s = platform_cfg.signals + x = platform_cfg.xau_strategy + b = platform_cfg.backtest + pair_info = PAIR_INFO[bt_cfg.pair] + is_xau = pair_info["display"] == "XAU/USD" + + f = { + "rsi_overbought": t.rsi_overbought, "rsi_oversold": t.rsi_oversold, + "rsi_neutral_upper": t.rsi_neutral_upper, "rsi_neutral_lower": t.rsi_neutral_lower, + "adx_strong_trend": t.adx_strong_trend, "adx_very_strong": t.adx_very_strong, + "adx_slope_emerging_threshold": t.adx_slope_emerging_threshold, + "roc_threshold": t.roc_threshold, + "min_rr_ratio": s.min_rr_ratio, "tp_tightening_factor": s.tp_tightening_factor, + "round_trip_cost": s.round_trip_cost, "min_tp_cost_multiple": s.min_tp_cost_multiple, + "min_sl_pips": s.min_sl_pips, "tp_floor_rr": s.tp_floor_rr, + "tp_atr_multiplier": x.tp_atr_multiplier, "sl_atr_multiplier": x.sl_atr_multiplier, + "min_atr_threshold": x.min_atr_threshold, "v2_sl_atr_multiplier": x.v2_sl_atr_multiplier, + "m5_min_abs_score": x.m5_min_abs_score, + "m5_min_confirmation_score": x.m5_min_confirmation_score, + "m5_sell_min_rsi": x.m5_sell_min_rsi, + "m5_pullback_min_score": x.m5_pullback_min_score, + "m5_pullback_adx_min": x.m5_pullback_adx_min, + "m5_pullback_rsi_buy_lo": x.m5_pullback_rsi_buy[0], + "m5_pullback_rsi_buy_hi": x.m5_pullback_rsi_buy[1], + "m5_pullback_rsi_sell_lo": x.m5_pullback_rsi_sell[0], + "m5_pullback_rsi_sell_hi": x.m5_pullback_rsi_sell[1], + "m5_pullback_retrace_lo": x.m5_pullback_retrace_atr[0], + "m5_pullback_retrace_hi": x.m5_pullback_retrace_atr[1], + "cp_tolerance_atr": x.cp_tolerance_atr, "cp_min_height_atr": x.cp_min_height_atr, + "cp_max_height_atr": x.cp_max_height_atr, "cp_hs_head_min_atr": x.cp_hs_head_min_atr, + "cp_hs_shoulder_tol_atr": x.cp_hs_shoulder_tol_atr, "cp_flag_pole_atr": x.cp_flag_pole_atr, + "cp_flag_retrace_lo": x.cp_flag_retrace[0], "cp_flag_retrace_hi": x.cp_flag_retrace[1], + "big_candle_atr_multiple": x.big_candle_atr_multiple, + "big_candle_lookback_hours": x.big_candle_lookback_hours, + "regime_filter_min_adx": b.regime_filter_min_adx, + "pip_value": pair_info["pip_value"], + "fb2_ema_lo": fb2_ema_lo_from_env(), + "tp_scale_aplus": 2.0, "tp_scale_a": 1.5, "max_penalty": 30.0, + } + from providers.base import BaseDataProvider # CANDLE_SECONDS (the aggregator's clock) + bar_secs = BaseDataProvider.CANDLE_SECONDS[platform_cfg.primary_timeframe] + strategies = list(bt_cfg.strategies_only or []) + if strategies: + smask = 0 + for name in strategies: + if name in STRATEGY_CODES and name: + smask |= 1 << STRATEGY_CODES[name] + else: + smask = (1 << 5) - 1 + i = { + "ema_trend_period": t.ema_trend, + "ema_mode_2ema": 1 if t.ema_mode == "2ema" else 0, + "ema_partial_neutral": 1 if t.ema_partial_neutral else 0, + "rsi_regime_aware": 1 if t.rsi_mode == "regime_aware" else 0, + "macd_as_pillar": 1 if t.macd_as_pillar else 0, + "roc_period": t.roc_period, "kc_period": t.kc_period, "kc_atr_period": t.kc_atr_period, + "divergence_lookback": t.divergence_lookback, "swing_sr_lookback": t.swing_sr_lookback, + "atr_period": t.atr_period, + "disabled_mask": _mask([d for d in (t.disabled_indicators or []) if d in DISABLED_BITS], + DISABLED_BITS), + "mtf_agreement_threshold": s.mtf_agreement_threshold, + "xau_enabled": 1 if x.enabled else 0, + "is_xau": 1 if is_xau else 0, + "session_filter": 1 if b.session_filter else 0, + "xau_blocked_hours_mask": _hours_mask(x.blocked_hours_utc), + "bt_blocked_hours_mask": _hours_mask(b.blocked_hours_utc), + "xau_allowed_weekdays_mask": _hours_mask(x.allowed_weekdays), + "bt_allowed_weekdays_mask": _hours_mask(b.allowed_weekdays), + "sell_min_quality_mask": _mask(x.sell_min_quality, QUALITY_CODES), + "sell_vol_mask": _mask(x.sell_vol_regimes, VOL_CODES), + "blocked_regimes_mask": _mask(x.blocked_regimes, REGIME_CODES), + "sell_blocked_regimes_mask": _mask(x.sell_blocked_regimes, REGIME_CODES), + "sell_blocked_hours_mask": _hours_mask(x.sell_blocked_hours_utc), + "m5_require_mtf": 1 if x.m5_require_mtf else 0, + "m5_sell_max_quality_mask": _mask(x.m5_sell_max_quality, QUALITY_CODES), + "m5_blocked_vol_mask": _mask(x.m5_blocked_vol_regimes, VOL_CODES), + "m5_pullback_enabled": 1 if x.m5_pullback_enabled else 0, + "m5_pullback_touch_bars": x.m5_pullback_touch_bars, + "cp_enabled": 1 if x.cp_enabled else 0, + "cp_patterns_mask": _mask([p for p in (x.cp_patterns or []) if p in PATTERN_ORDER], + {p: k for k, p in enumerate(PATTERN_ORDER)}), + "cp_lookback": x.cp_lookback, "cp_swing_window": x.cp_swing_window, + "cp_recency_bars": x.cp_recency_bars, "cp_min_sep_bars": x.cp_min_sep_bars, + "cp_flag_pole_bars": x.cp_flag_pole_bars, "cp_flag_min_bars": x.cp_flag_min_bars, + "cp_flag_max_bars": x.cp_flag_max_bars, + "quality_filter_mask": _mask([q for q in bt_cfg.quality_filter if q in QUALITY_CODES], + QUALITY_CODES), + "strategies_mask": smask, + "primary_is_m5": 1 if platform_cfg.primary_timeframe == Timeframe.M5 else 0, + "primary_bar_secs": bar_secs, + "regime_filter": 1 if b.regime_filter else 0, + "n_confirm": len(platform_cfg.confirmation_timeframes), + } + pf = np.array([float(f[name]) for name in PARAMS_F], dtype=np.float64) + pi = np.array([int(i[name]) for name in PARAMS_I], dtype=np.int32) + return pf, pi, regime_table(platform_cfg).ravel() + + +def _defines() -> str: + lines = ["#define FB2_MAX_CONFIRM %d" % MAX_CONFIRM, + "#define FB2_MAX_DIV_LOOKBACK %d" % MAX_DIV_LOOKBACK, + "#define FB2_MAX_SR_LOOKBACK %d" % MAX_SR_LOOKBACK, + "#define FB2_MAX_CP_LOOKBACK %d" % MAX_CP_LOOKBACK, + "#define FB2_MAX_WARN %d" % MAX_WARN, + "#define FB2_FEAT_MAX_P %d" % FEAT_MAX_P, + "#define FB2_FEAT_CONFIRM_TAIL %d" % FEAT_CONFIRM_TAIL, + "#define FB2_FEAT_BURST_BARS %d" % FEAT_BURST_BARS, + "#define FB2_N_FEATURES %d" % N_FEATURES, + "#define FB2_N_BAR_COLS %d" % N_BAR_COLS] + for k, name in enumerate(PARAMS_F): + lines.append(f"#define PF_{name} {k}") + for k, name in enumerate(PARAMS_I): + lines.append(f"#define PI_{name} {k}") + for k, name in enumerate(K1_INT_FIELDS): + lines.append(f"#define K1I_{name} {k}") + lines.append(f"#define K1I_N {len(K1_INT_FIELDS)}") + for k, name in enumerate(K1_FLOAT_FIELDS): + lines.append(f"#define K1F_{name} {k}") + lines.append(f"#define K1F_N {len(K1_FLOAT_FIELDS)}") + for k, name in enumerate(K2_INT_FIELDS): + lines.append(f"#define K2I_{name} {k}") + lines.append(f"#define K2I_N {len(K2_INT_FIELDS)}") + for k, name in enumerate(K2_FLOAT_FIELDS): + lines.append(f"#define K2F_{name} {k}") + lines.append(f"#define K2F_N {len(K2_FLOAT_FIELDS)}") + for k, name in enumerate(WARN_NAMES): + lines.append(f"#define W_{name} {k}") + for name, bit in DISABLED_BITS.items(): + lines.append(f"#define DIS_{name} (1 << {bit})") + for k, name in enumerate(PENALTY_TAGS): + lines.append(f"#define PEN_{name} {k}") + for k, name in enumerate(PATTERN_ORDER): + lines.append(f"#define CP_{name} {k}") + for name, code in REGIME_CODES.items(): + lines.append(f"#define REG_{name} {code}") + for name, code in VOL_CODES.items(): + lines.append(f"#define VOL_{name} {code}") + for name, code in QUALITY_CODES.items(): + lines.append(f"#define Q_{name.replace('+', 'PLUS')} {code}") + for name, code in STRATEGY_CODES.items(): + lines.append(f"#define STRAT_{name or 'NONE'} {code}") + return "\n".join(lines) + "\n" + + +# ═════════════════════════════════════════════════════════════════════ +# CUDA sources +# ═════════════════════════════════════════════════════════════════════ + +# pandas' Kahan/sign-count rolling-mean helpers, taken verbatim from the +# proven kernel module so the two copies cannot diverge (a test pins it). +_ROLL_HELPERS = _gk._ROLL_SRC.split('extern "C" __global__')[0] + +_SRC_COMMON = r''' + +// Python min()/max() in argument order: keep the first argument unless the +// second is strictly smaller/greater (NaN never wins). +#define PYMIN(a, b) (((b) < (a)) ? (b) : (a)) +#define PYMAX(a, b) (((b) > (a)) ? (b) : (a)) +#define FB2_NAN nan("") + +__device__ __forceinline__ double fb2_finite(double x, double d) { + return isfinite(x) ? x : d; +} + +// CPython round(x, ndigits) for ndigits = 2 (scale 100) or 1 (scale 10): +// the exact binary value is rounded half-even to that many decimals and the +// nearest double of the result is returned. fma() gives the exact residual of +// the product, so a tie is recognised exactly; the correctly rounded division +// r/scale is the same double strtod() returns for the decimal string. +__device__ __forceinline__ double fb2_pyround(double x, double scale) { + if (!isfinite(x)) return x; + const double p = x * scale; + const double e = fma(x, scale, -p); + double r = rint(p); + const double d = p - r; + if (fabs(d) == 0.5 && e != 0.0 && ((e > 0.0) == (d > 0.0))) + r += (d > 0.0) ? 1.0 : -1.0; + return r / scale; +} + +// numpy pairwise summation (loops_utils.h.src, PW_BLOCKSIZE 128), stride 1. +// numpy's recursion (n2 = n/2 rounded down to a multiple of 8; left + right) +// is emulated with an explicit stack so no device function is recursive. +__device__ __forceinline__ double fb2_np_pairwise_leaf(const double* a, int n) { + if (n < 8) { + double res = 0.0; + for (int i = 0; i < n; ++i) res += a[i]; + return res; + } + double r0 = a[0], r1 = a[1], r2 = a[2], r3 = a[3], + r4 = a[4], r5 = a[5], r6 = a[6], r7 = a[7]; + int i; + for (i = 8; i < n - (n % 8); i += 8) { + r0 += a[i + 0]; r1 += a[i + 1]; r2 += a[i + 2]; r3 += a[i + 3]; + r4 += a[i + 4]; r5 += a[i + 5]; r6 += a[i + 6]; r7 += a[i + 7]; + } + double res = ((r0 + r1) + (r2 + r3)) + ((r4 + r5) + (r6 + r7)); + for (; i < n; ++i) res += a[i]; + return res; +} + +__device__ __forceinline__ double fb2_np_pairwise_sum(const double* a, int n) { + int st_start[40], st_len[40], st_phase[40]; + double st_left[40]; + int sp = 1; + st_start[0] = 0; st_len[0] = n; st_phase[0] = 0; st_left[0] = 0.0; + double ret = 0.0; + while (sp > 0) { + const int i = sp - 1; + if (st_phase[i] == 0) { + if (st_len[i] <= 128) { ret = fb2_np_pairwise_leaf(a + st_start[i], st_len[i]); sp -= 1; continue; } + int n2 = st_len[i] / 2; n2 -= n2 % 8; + st_phase[i] = 1; + st_start[sp] = st_start[i]; st_len[sp] = n2; st_phase[sp] = 0; st_left[sp] = 0.0; sp += 1; + } else if (st_phase[i] == 1) { + st_left[i] = ret; + int n2 = st_len[i] / 2; n2 -= n2 % 8; + st_phase[i] = 2; + st_start[sp] = st_start[i] + n2; st_len[sp] = st_len[i] - n2; st_phase[sp] = 0; st_left[sp] = 0.0; sp += 1; + } else { + ret = st_left[i] + ret; + sp -= 1; + } + } + return ret; +} + +// np.nanmean over a window: NaN -> 0 for the pairwise sum, divided by the +// count of finite values (NaN when there are none). +__device__ __forceinline__ double fb2_np_nanmean(const double* a, int n) { + double tmp[128]; + int cnt = 0; + for (int i = 0; i < n; ++i) { + const double v = a[i]; + if (v == v) { tmp[i] = v; cnt += 1; } else tmp[i] = 0.0; + } + if (cnt == 0) return FB2_NAN; + return fb2_np_pairwise_sum(tmp, n) / (double)cnt; +} + +// pandas ewm(alpha, min_periods, adjust=False, ignore_na=False).mean() — +// pandas/_libs/window/aggregations.pyx `ewm`, restarted at x[0]. +// Writes the whole series into out[0..n). +__device__ __forceinline__ void fb2_pd_ewm(const double* x, int n, double alpha, int minp, double* out) { + const double com = (1.0 - alpha) / alpha; + const double a = 1.0 / (1.0 + com); + const double old_wt_factor = 1.0 - a; + const double new_wt = a; + if (n <= 0) return; + double w = x[0]; + int nobs = (w == w) ? 1 : 0; + out[0] = (nobs >= minp) ? w : FB2_NAN; + double old_wt = 1.0; + for (int i = 1; i < n; ++i) { + const double cur = x[i]; + const bool is_obs = (cur == cur); + if (is_obs) nobs += 1; + if (w == w) { + old_wt *= old_wt_factor; // ignore_na=False: every step decays + if (is_obs) { + if (w != cur) { + w = old_wt * w + new_wt * cur; + w /= (old_wt + new_wt); + } + old_wt = 1.0; + } + } else if (is_obs) { + w = cur; + } + out[i] = (nobs >= minp) ? w : FB2_NAN; + } +} + +// position_sizing.wilder_atr on a tail that starts at bar `base`: +// tr[0] = high-low (the shifted close is NaN on the first row of the tail), +// then the pandas ewm above with alpha = 1/period, min_periods = period. +__device__ __forceinline__ void fb2_wilder_atr_tail(const double* high, const double* low, const double* close, + int base, int n, int period, double* out) { + double tr[FB2_FEAT_MAX_P]; + for (int i = 0; i < n; ++i) { + const int k = base + i; + const double tr1 = high[k] - low[k]; + if (i == 0) { tr[i] = tr1; continue; } + const double pc = close[k - 1]; + const double tr2 = fabs(high[k] - pc); + const double tr3 = fabs(low[k] - pc); + // pandas DataFrame.max(axis=1): NaN -> -inf then the plain maximum + double m = tr1; + if (tr2 > m) m = tr2; + if (tr3 > m) m = tr3; + tr[i] = m; + } + fb2_pd_ewm(tr, n, 1.0 / (double)period, period, out); +} + +// pandas Series.rolling(window).mean() restarted on a tail; returns the value +// at the tail's last bar (the online algorithm must run from the tail start). +__device__ __forceinline__ double fb2_roll_mean_last(const double* xs, int n, int window) { + const long long minp = window; + double comp_add = 0, comp_rem = 0, sum_x = 0, prev_value = xs[0]; + long long nobs = 0, neg_ct = 0, ncsv = 0; + double rm = FB2_NAN; + for (long long i = 0; i < n; ++i) { + long long st = i - window + 1; if (st < 0) st = 0; + long long en = i + 1; + long long pst = (i - 1) - window + 1; if (pst < 0) pst = 0; + long long pen = i; + if (i == 0) { + comp_add = comp_rem = sum_x = 0; nobs = neg_ct = 0; prev_value = xs[st]; ncsv = 0; + for (long long j = st; j < en; ++j) + fb2_add_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_add, &ncsv, &prev_value); + } else { + for (long long j = pst; j < st; ++j) + fb2_remove_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_rem); + for (long long j = pen; j < en; ++j) + fb2_add_mean(xs[j], &nobs, &sum_x, &neg_ct, &comp_add, &ncsv, &prev_value); + } + if (nobs >= minp && nobs > 0) { + rm = sum_x / (double)nobs; + if (ncsv >= nobs) rm = prev_value; + else if (neg_ct == 0 && rm < 0) rm = 0; + else if (neg_ct == nobs && rm > 0) rm = 0; + } else { + rm = FB2_NAN; + } + } + return rm; +} + +// engine._find_swing_points (NON-strict, NaN-skipping) on the window +// vals[0..L) — is index i a pivot? window=3. +__device__ __forceinline__ bool fb2_engine_pivot(const double* vals, int L, int i, int window, bool low) { + if (i < window || i >= L - window) return false; + const double c = vals[i]; + if (c != c) return false; + int nleft = 0, nright = 0; + bool ok = true; + for (int j = 1; j <= window; ++j) { + const double v = vals[i - j]; + if (v != v) continue; + nleft += 1; + if (low ? !(c <= v) : !(c >= v)) ok = false; + } + for (int j = 1; j <= window; ++j) { + if (i + j >= L) continue; + const double v = vals[i + j]; + if (v != v) continue; + nright += 1; + if (low ? !(c <= v) : !(c >= v)) ok = false; + } + if (nleft == 0 || nright == 0) return false; + return ok; +} + +// chart_patterns._swing_indices (STRICT, any NaN in the window rejects). +__device__ __forceinline__ bool fb2_strict_pivot(const double* vals, int n, int i, int w, bool find_high) { + if (i < w || i >= n - w) return false; + const double c = vals[i]; + if (c != c) return false; + double rest_max = -1.0/0.0, rest_min = 1.0/0.0; + for (int k = i - w; k <= i + w; ++k) { + const double v = vals[k]; + if (v != v) return false; + if (k == i) continue; + if (v > rest_max) rest_max = v; + if (v < rest_min) rest_min = v; + } + return find_high ? (c > rest_max) : (c < rest_min); +} +''' + +_SRC_K1 = r''' +#define K1F(name) OF[(size_t)(K1F_##name) * (size_t)n + (size_t)j] +#define K1I(name) OI[(size_t)(K1I_##name) * (size_t)n + (size_t)j] + +extern "C" __global__ +void fb2_analyze_tf(const double* __restrict__ open_, const double* __restrict__ high, + const double* __restrict__ low, const double* __restrict__ close, + const double* __restrict__ ema_fast, const double* __restrict__ ema_medium, + const double* __restrict__ ema_slow, const double* __restrict__ ema_trend, + const int has_ema_trend, + const double* __restrict__ rsi, const double* __restrict__ macd_hist, + const double* __restrict__ bb_mid, const double* __restrict__ bb_upper, + const double* __restrict__ bb_lower, const double* __restrict__ bb_band_width, + const double* __restrict__ bb_avg_width, const double* __restrict__ atr, + const double* __restrict__ adx, const double* __restrict__ plus_di, + const double* __restrict__ minus_di, const double* __restrict__ adx_slope, + const double* __restrict__ roc, const double* __restrict__ kc_mid, + const double* __restrict__ kc_upper, const double* __restrict__ kc_lower, + const double* __restrict__ kc_squeeze, + const double* __restrict__ PF, const int* __restrict__ PI, + const double* __restrict__ RM, const int tf_is_5min, const int n, + double* __restrict__ OF, int* __restrict__ OI) +{ + const int j = blockIdx.x * blockDim.x + threadIdx.x; + if (j >= n) return; + const int pos = j; + const int len = j + 1; + for (int f = 0; f < K1F_N; ++f) OF[(size_t)f * (size_t)n + (size_t)j] = FB2_NAN; + for (int f = 0; f < K1I_N; ++f) OI[(size_t)f * (size_t)n + (size_t)j] = 0; + K1I(regime) = REG_unknown; + K1I(vol) = VOL_normal; + K1F(composite_score) = 0.0; + K1F(composite_conf) = 0.0; + K1I(composite_signal) = 0; + if (len < 30) { K1I(valid) = 0; return; } + K1I(valid) = 1; + + const int dis = PI[PI_disabled_mask]; + const double latest = close[pos]; + + // ── ATR (weight 0; feeds ribbon, candlestick, volatility regime) ── + int atr_present = 0; + double atr_val = 0.0; + if (!(dis & DIS_atr)) { + atr_present = 1; + atr_val = atr[pos]; + if (atr_val != atr_val) atr_val = 0.0; + const double atr_pct = (latest != 0.0) ? (atr_val / latest) * 100.0 : 0.0; + double pct; + if (pos >= 49) { + int cnt = 0; + for (int k = pos - 49; k <= pos; ++k) if (atr[k] < atr_val) cnt += 1; + pct = ((double)cnt / 50.0) * 100.0; + } else { + pct = 50.0; + } + int vol; + if (pct > 80.0) vol = VOL_extreme; else if (pct > 60.0) vol = VOL_high; + else if (pct > 30.0) vol = VOL_normal; else vol = VOL_low; + K1I(vol) = vol; + K1F(atr) = atr_val; K1F(atr_pct) = atr_pct; K1F(atr_percentile) = pct; + } + K1I(atr_present) = atr_present; + const double atr_kl = atr_present ? atr_val : 0.0; // key_levels.get("atr", 0) + + // ── EMA trend ── + int ema_present = 0, ema_signal = 0; + if (!(dis & DIS_ema_trend)) { + ema_present = 1; + const double m = ema_medium[pos]; + const double s = ema_slow[pos]; + int t_ok = (has_ema_trend && !(ema_trend[pos] != ema_trend[pos])) ? 1 : 0; + double t = t_ok ? ema_trend[pos] : FB2_NAN; + if (t_ok && pos < PI[PI_ema_trend_period] - 1) { t_ok = 0; t = FB2_NAN; } + double conf, value; int sig, cross = 0; + if (PI[PI_ema_mode_2ema]) { + if (!t_ok) { + sig = 0; conf = 30.0; value = 0.0; + K1I(ema2_desc) = 0; + K1I(kl_ema_present) = 1; + K1F(ema_medium) = m; K1F(ema_slow) = s; + } else { + int d; + if (latest > s && s > t) { sig = 2; conf = 85.0; d = 1; } + else if (latest > s && s < t) { sig = 1; conf = 60.0; d = 2; } + else if (latest < s && s < t) { sig = -2; conf = 85.0; d = 3; } + else if (latest < s && s > t) { sig = -1; conf = 60.0; d = 4; } + else { sig = 0; conf = 40.0; d = 5; } + if (pos >= 9) { + const int lim = PYMIN(11, pos + 1); + for (int lb = 2; lb < lim; ++lb) { + const double ps = ema_slow[pos - lb + 1]; + const double pt = ema_trend[pos - lb + 1]; + if (ps <= pt && s > t) { cross = 1; conf = PYMIN(conf + 10.0, 95.0); break; } + if (ps >= pt && s < t) { cross = 2; conf = PYMIN(conf + 10.0, 95.0); break; } + } + } + value = (sig == 2) ? 1.0 : (sig == 1) ? 0.5 : (sig == 0) ? 0.0 : (sig == -1) ? -0.5 : -1.0; + K1I(ema2_desc) = d; + K1I(ema_trend_present) = 1; K1F(ema_trend) = t; + K1I(kl_ema_present) = 1; + K1F(ema_medium) = m; K1F(ema_slow) = s; + } + } else { + const double f = ema_fast[pos]; + int align = 0; + if (latest > f) align += 1; else if (latest < f) align -= 1; + if (f > m) align += 1; else if (f < m) align -= 1; + if (m > s) align += 1; else if (m < s) align -= 1; + if (t_ok) { if (s > t) align += 1; else if (s < t) align -= 1; } + const int maxc = t_ok ? 4 : 3; + const double normalized = (double)align / (double)maxc; + if (normalized > 0.6) { sig = 2; conf = PYMIN(90.0, 60.0 + normalized * 30.0); } + else if (normalized > 0.2) { sig = 1; conf = 55.0 + normalized * 20.0; } + else if (normalized < -0.6) { sig = -2; conf = PYMIN(90.0, 60.0 + fabs(normalized) * 30.0); } + else if (normalized < -0.2) { sig = -1; conf = 55.0 + fabs(normalized) * 20.0; } + else { sig = 0; conf = 40.0; } + if (PI[PI_ema_partial_neutral] && (sig == 1 || sig == -1)) { sig = 0; conf = 40.0; } + if (pos >= 2) { + const double pf = ema_fast[pos - 1], pm = ema_medium[pos - 1]; + if (pf <= pm && f > m) { cross = 1; conf = PYMIN(conf + 10.0, 95.0); } + else if (pf >= pm && f < m) { cross = 2; conf = PYMIN(conf + 10.0, 95.0); } + } + int ribbon = 0; + if (atr_kl > 0.0) { + const double mx = PYMAX(PYMAX(f, m), s); + const double mn = PYMIN(PYMIN(f, m), s); + const double rw = (mx - mn) / atr_kl; + if (rw < 1.0) ribbon = 1; + if (pos >= 5) { + const double pf = ema_fast[pos - 4], pm = ema_medium[pos - 4], ps = ema_slow[pos - 4]; + const double pmx = PYMAX(PYMAX(pf, pm), ps); + const double pmn = PYMIN(PYMIN(pf, pm), ps); + const double pw = (atr_kl > 0.0) ? (pmx - pmn) / atr_kl : 0.0; + if (pw < 1.0 && rw >= 1.0) { ribbon = 2; conf = PYMIN(conf + 10.0, 95.0); } + } + } + value = normalized; + K1I(ema_align) = align; K1I(ema_max_checks) = maxc; K1I(ema_ribbon) = ribbon; + if (t_ok) { K1I(ema_trend_present) = 1; K1F(ema_trend) = t; } + K1I(kl_ema_fast_present) = 1; K1I(kl_ema_present) = 1; + K1F(ema_fast) = f; K1F(ema_medium) = m; K1F(ema_slow) = s; + } + ema_signal = sig; + K1I(ema_present) = 1; K1I(ema_signal) = sig; K1I(ema_cross) = cross; + K1F(ema_value) = value; K1F(ema_conf) = conf; + } + + // ── RSI (classic pass) ── + int rsi_present = 0, rsi_signal = 0; double rsi_val = 50.0, rsi_conf = 30.0; int rsi_desc = 4; + if (!(dis & DIS_rsi)) { + rsi_present = 1; + rsi_val = rsi[pos]; + if (rsi_val != rsi_val) rsi_val = 50.0; + if (rsi_val >= PF[PF_rsi_overbought]) { rsi_signal = -1; rsi_conf = PYMIN(80.0, 50.0 + (rsi_val - PF[PF_rsi_overbought])); rsi_desc = 0; } + else if (rsi_val <= PF[PF_rsi_oversold]) { rsi_signal = 1; rsi_conf = PYMIN(80.0, 50.0 + (PF[PF_rsi_oversold] - rsi_val)); rsi_desc = 1; } + else if (rsi_val >= PF[PF_rsi_neutral_upper]) { rsi_signal = 1; rsi_conf = 45.0; rsi_desc = 2; } + else if (rsi_val <= PF[PF_rsi_neutral_lower]) { rsi_signal = -1; rsi_conf = 45.0; rsi_desc = 3; } + else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 4; } + } + + // ── Bollinger ── + int bb_present = 0, bb_signal = 0; double bb_value = FB2_NAN, bb_conf = 0.0; + if (!(dis & DIS_bollinger)) { + const double up = bb_upper[pos], lo = bb_lower[pos], sma = bb_mid[pos]; + if (!(up != up || lo != lo || sma != sma || sma == 0.0)) { + bb_present = 1; + const double range = up - lo; + const double bbpos = (range != 0.0) ? (latest - lo) / range : 0.5; + const double bw = bb_band_width[pos]; + int sq = 0; + if (pos >= 19) { + const double aw = bb_avg_width[pos]; + if (!(aw != aw) && aw > 0.0) sq = (bw < aw * 0.75) ? 1 : 0; + } + int d; + if (bbpos > 1.0) { bb_signal = -1; bb_conf = 65.0; d = 0; } + else if (bbpos < 0.0) { bb_signal = 1; bb_conf = 65.0; d = 1; } + else if (bbpos > 0.8) { bb_signal = -1; bb_conf = 45.0; d = 2; } + else if (bbpos < 0.2) { bb_signal = 1; bb_conf = 45.0; d = 3; } + else { bb_signal = 0; bb_conf = 30.0; d = 4; } + if (sq) bb_conf = PYMIN(bb_conf + 10.0, 85.0); + bb_value = bbpos; + K1I(bb_present) = 1; K1I(bb_signal) = bb_signal; K1I(bb_desc) = d; K1I(bb_squeeze) = sq; + K1F(bb_upper) = up; K1F(bb_lower) = lo; K1F(bb_mid) = sma; + K1F(bb_value) = bbpos; K1F(bb_conf) = bb_conf; + } + } + + // ── ADX ── + int adx_present = 0, adx_signal = 0; double adx_val = 0.0, adx_conf = 0.0, slope = 0.0; + if (!(dis & DIS_adx)) { + adx_present = 1; + adx_val = adx[pos]; if (adx_val != adx_val) adx_val = 0.0; + double pdi = plus_di[pos], mdi = minus_di[pos]; + if (pdi != pdi) pdi = 0.0; + if (mdi != mdi) mdi = 0.0; + int strength; + if (adx_val >= PF[PF_adx_very_strong]) { strength = 0; adx_conf = 80.0; } + else if (adx_val >= PF[PF_adx_strong_trend]) { strength = 1; adx_conf = 65.0; } + else { strength = 2; adx_conf = 40.0; } + int dir; + if (pdi > mdi) { adx_signal = (adx_val >= PF[PF_adx_strong_trend]) ? 1 : 0; dir = 0; } + else { adx_signal = (adx_val >= PF[PF_adx_strong_trend]) ? -1 : 0; dir = 1; } + slope = adx_slope[pos]; if (slope != slope) slope = 0.0; + int sd = 0; + const double thr = PF[PF_adx_slope_emerging_threshold]; + if (adx_val < PF[PF_adx_strong_trend] && slope > thr) { sd = 1; adx_conf = PYMIN(adx_conf + 15.0, 85.0); } + else if (adx_val >= PF[PF_adx_strong_trend] && slope > 0.0) { sd = 2; adx_conf = PYMIN(adx_conf + 5.0, 85.0); } + else if (adx_val >= PF[PF_adx_strong_trend] && slope < -thr) { sd = 3; adx_conf = PYMAX(adx_conf - 10.0, 30.0); } + K1I(adx_present) = 1; K1I(adx_signal) = adx_signal; K1I(adx_strength) = strength; + K1I(adx_dir) = dir; K1I(adx_slope_desc) = sd; + K1F(adx_value) = adx_val; K1F(plus_di) = pdi; K1F(minus_di) = mdi; + K1F(adx_slope) = slope; K1F(adx_conf) = adx_conf; + } + + // ── MACD (histogram always computed; pillar honours the disable) ── + int macd_present = 0, macd_signal = 0; double macd_conf = 0.0; + if (PI[PI_macd_as_pillar] && pos >= 2) { + const double hist_now = macd_hist[pos], hist_prev = macd_hist[pos - 1]; + if (!(hist_now != hist_now || hist_prev != hist_prev)) { + const bool crossed_up = (hist_prev <= 0.0 && 0.0 < hist_now); + const bool crossed_down = (hist_prev >= 0.0 && 0.0 > hist_now); + const bool accel_up = (hist_now > 0.0 && hist_now > hist_prev && hist_prev > 0.0); + const bool accel_down = (hist_now < 0.0 && hist_now < hist_prev && hist_prev < 0.0); + const bool fading_up = (hist_now > 0.0 && hist_now < hist_prev); + const bool fading_down = (hist_now < 0.0 && hist_now > hist_prev); + double hist_pct = 1.0; + if (pos >= 19) { + double win[20]; + for (int k = 0; k < 20; ++k) win[k] = fabs(macd_hist[pos - 19 + k]); + const double mean_abs = fb2_np_nanmean(win, 20); + hist_pct = (mean_abs > 0.0) ? fabs(hist_now) / mean_abs : 1.0; + } + int d; + if (crossed_up || accel_up) { macd_signal = accel_up ? 2 : 1; macd_conf = PYMIN(80.0, 55.0 + hist_pct * 10.0); d = accel_up ? 0 : 1; } + else if (crossed_down || accel_down) { macd_signal = accel_down ? -2 : -1; macd_conf = PYMIN(80.0, 55.0 + hist_pct * 10.0); d = accel_down ? 2 : 3; } + else if (fading_up) { macd_signal = 0; macd_conf = 35.0; d = 4; } + else if (fading_down) { macd_signal = 0; macd_conf = 35.0; d = 5; } + else { macd_signal = 0; macd_conf = 30.0; d = 6; } + if (!(dis & DIS_macd)) { + macd_present = 1; + K1I(macd_present) = 1; K1I(macd_signal) = macd_signal; K1I(macd_desc) = d; + K1F(macd_value) = hist_now; K1F(macd_conf) = macd_conf; + } + } + } + + // ── ROC ── + int roc_present = 0, roc_signal = 0; double roc_val = 0.0, roc_conf = 0.0; + if (!(dis & DIS_roc) && pos >= PI[PI_roc_period] + 4) { + roc_val = roc[pos]; + const double roc_prev = (pos > 0) ? roc[pos - 1] : FB2_NAN; + if (!(roc_val != roc_val || roc_prev != roc_prev)) { + roc_present = 1; + const bool rising = roc_val > roc_prev; + const double roc_abs = fabs(roc_val); + const double threshold = PF[PF_roc_threshold]; + int d; + if (roc_val > 0.0 && rising) { + if (roc_abs > threshold * 2.0) { roc_signal = 2; roc_conf = PYMIN(80.0, 60.0 + roc_abs * 5.0); d = 0; } + else { roc_signal = 1; roc_conf = PYMIN(70.0, 50.0 + roc_abs * 10.0); d = 1; } + } else if (roc_val < 0.0 && !rising) { + if (roc_abs > threshold * 2.0) { roc_signal = -2; roc_conf = PYMIN(80.0, 60.0 + roc_abs * 5.0); d = 2; } + else { roc_signal = -1; roc_conf = PYMIN(70.0, 50.0 + roc_abs * 10.0); d = 3; } + } else { roc_signal = 0; roc_conf = 30.0; d = 4; } + K1I(roc_present) = 1; K1I(roc_signal) = roc_signal; K1I(roc_desc) = d; + K1F(roc_value) = roc_val; K1F(roc_conf) = roc_conf; + } + } + + // ── Keltner ── + int kc_present = 0, kc_signal = 0; double kc_conf = 0.0, kc_value = 0.0; + if (!(dis & DIS_keltner)) { + const int guard = PYMAX(PI[PI_kc_period], PI[PI_kc_atr_period]) + 4; + if (pos >= guard) { + const double up = kc_upper[pos], lo = kc_lower[pos], mid = kc_mid[pos]; + if (!(up != up || lo != lo)) { + kc_present = 1; + const bool is_squeeze = (kc_squeeze[pos] != 0.0); + bool released = false; + if (pos >= 1) released = (kc_squeeze[pos - 1] != 0.0) && !is_squeeze; + int macd_direction = 0; + if (!(macd_hist[pos] != macd_hist[pos])) macd_direction = (macd_hist[pos] > 0.0) ? 1 : -1; + int d; + if (released) { + if (macd_direction > 0) { kc_signal = 2; d = 0; } + else if (macd_direction < 0) { kc_signal = -2; d = 1; } + else { kc_signal = 0; d = 2; } + kc_conf = 80.0; + } else if (is_squeeze) { kc_signal = 0; kc_conf = 60.0; d = 3; } + else if (latest > up) { kc_signal = 1; kc_conf = 65.0; d = 4; } + else if (latest < lo) { kc_signal = -1; kc_conf = 65.0; d = 5; } + else { kc_signal = 0; kc_conf = 30.0; d = 6; } + kc_value = latest - mid; + K1I(kc_present) = 1; K1I(kc_signal) = kc_signal; K1I(kc_desc) = d; + K1F(kc_upper) = up; K1F(kc_lower) = lo; K1F(kc_mid) = mid; + K1F(kc_value) = kc_value; K1F(kc_conf) = kc_conf; + } + } + } + + // ── Divergence (engine._detect_divergences, not overridden) ── + int div_present = 0, div_signal = 0, div_mask = 0; + if (!(dis & DIS_divergence)) { + const int lookback = PYMIN(PI[PI_divergence_lookback], len - 1); + if (lookback >= 5) { + const int base = pos + 1 - lookback; + for (int series = 0; series < 2; ++series) { + if (series == 0 && (dis & DIS_rsi)) continue; // _rsi_series is None + const double* ind = (series == 0) ? rsi : macd_hist; + for (int mode = 0; mode < 2; ++mode) { // 0 lows, 1 highs + const bool low = (mode == 0); + int ind_idx[FB2_MAX_DIV_LOOKBACK]; double ind_val[FB2_MAX_DIV_LOOKBACK]; + int n_ind = 0; + for (int i = 0; i < lookback; ++i) { + if (fb2_engine_pivot(ind + base, lookback, i, 3, low)) { + ind_idx[n_ind] = i; ind_val[n_ind] = ind[base + i]; n_ind += 1; + } + } + int n_pairs = 0; + double p1v = 0, i1v = 0, p2v = 0, i2v = 0; + for (int i = 0; i < lookback; ++i) { + if (!fb2_engine_pivot(close + base, lookback, i, 3, low)) continue; + if (n_ind == 0) continue; + int best = 0; int bestd = abs(ind_idx[0] - i); + for (int q = 1; q < n_ind; ++q) { + const int dd = abs(ind_idx[q] - i); + if (dd < bestd) { bestd = dd; best = q; } + } + if (bestd <= 2) { + p1v = p2v; i1v = i2v; + p2v = close[base + i]; i2v = ind_val[best]; + n_pairs += 1; + } + } + if (n_pairs >= 2) { + if (low) { if (p2v < p1v && i2v > i1v) div_mask |= (series == 0) ? 1 : 4; } + else { if (p2v > p1v && i2v < i1v) div_mask |= (series == 0) ? 2 : 8; } + } + } + } + if (div_mask) { + div_present = 1; + const bool is_bull = (div_mask & 5) != 0; + const bool is_bear = (div_mask & 10) != 0; + div_signal = (is_bull && !is_bear) ? 1 : (is_bear && !is_bull) ? -1 : 0; + const int cnt = ((div_mask & 1) ? 1 : 0) + ((div_mask & 2) ? 1 : 0) + + ((div_mask & 4) ? 1 : 0) + ((div_mask & 8) ? 1 : 0); + K1I(div_present) = 1; K1I(div_signal) = div_signal; K1I(div_mask) = div_mask; + K1F(div_value) = (double)cnt; + } + } + } + + // ── Structure S/R (engine._compute_structure_sr) ── + int sh_mask = 0, sl_mask = 0; + double sh[3] = {FB2_NAN, FB2_NAN, FB2_NAN}, sl[3] = {FB2_NAN, FB2_NAN, FB2_NAN}; + if (!(dis & DIS_structure)) { + int base_lb = PI[PI_swing_sr_lookback]; + if (tf_is_5min) base_lb = PYMAX(base_lb, 120); + const int lookback = PYMIN(base_lb, len - 1); + if (lookback >= 10) { + const double price = latest; + const int base = pos + 1 - lookback; + int n_above = 0; // highs + double above[3], below_max = FB2_NAN; bool has_below = false; + int n_lbelow = 0; // lows + double lbelow[3], labove_min = FB2_NAN; bool has_labove = false; + for (int i = 0; i < lookback; ++i) { + if (fb2_engine_pivot(high + base, lookback, i, 3, false)) { + const double v = high[base + i]; + if (v > price) { + // keep the 3 smallest (ascending) + if (n_above < 3) { above[n_above++] = v; } + else if (v < above[2]) { above[2] = v; } + for (int a = n_above - 1; a > 0; --a) if (above[a] < above[a-1]) { double t = above[a]; above[a] = above[a-1]; above[a-1] = t; } + } else { + if (!has_below || v > below_max) { below_max = v; has_below = true; } + } + } + if (fb2_engine_pivot(low + base, lookback, i, 3, true)) { + const double v = low[base + i]; + if (v < price) { + // keep the 3 largest (descending) + if (n_lbelow < 3) { lbelow[n_lbelow++] = v; } + else if (v > lbelow[2]) { lbelow[2] = v; } + for (int a = n_lbelow - 1; a > 0; --a) if (lbelow[a] > lbelow[a-1]) { double t = lbelow[a]; lbelow[a] = lbelow[a-1]; lbelow[a-1] = t; } + } else { + if (!has_labove || v < labove_min) { labove_min = v; has_labove = true; } + } + } + } + for (int a = 0; a < n_above; ++a) { sh[a] = above[a]; sh_mask |= (1 << a); } + for (int a = 0; a < n_lbelow; ++a) { sl[a] = lbelow[a]; sl_mask |= (1 << a); } + if (n_above == 0 && has_below) { sh[0] = below_max; sh_mask |= 1; K1I(sh_fallback) = 1; } + if (n_lbelow == 0 && has_labove) { sl[0] = labove_min; sl_mask |= 1; K1I(sl_fallback) = 1; } + } + } + K1I(sh_mask) = sh_mask; K1I(sl_mask) = sl_mask; + K1F(sh1) = sh[0]; K1F(sh2) = sh[1]; K1F(sh3) = sh[2]; + K1F(sl1) = sl[0]; K1F(sl2) = sl[1]; K1F(sl3) = sl[2]; + + // ── Candlestick (engine._detect_candlestick_patterns) ── + int cs_present = 0, cs_signal = 0; double cs_conf = 0.0, cs_value = 0.0; + if (!(dis & DIS_candlestick) && len >= 5) { + const double o = open_[pos], h = high[pos], l = low[pos], c = close[pos]; + const double prev_o = open_[pos - 1], prev_c = close[pos - 1]; + const double body = fabs(c - o); + const double upper_wick = h - PYMAX(o, c); + const double lower_wick = PYMIN(o, c) - l; + const double full_range = h - l; + if (full_range != 0.0) { + const double body_ratio = body / full_range; + int mask = 0, bullish = 0, bearish = 0; + if (body_ratio < 0.1) mask |= (1 << 0); + if (lower_wick > 2.0 * body && upper_wick < body * 0.5 && body_ratio > 0.1) { + if (c > o) { mask |= (1 << 1); bullish += 1; } else { mask |= (1 << 2); bearish += 1; } + } + if (upper_wick > 2.0 * body && lower_wick < body * 0.5 && body_ratio > 0.1) { + if (c < o) { mask |= (1 << 3); bearish += 1; } else { mask |= (1 << 4); bullish += 1; } + } + if (len >= 3) { + if (prev_c < prev_o && c > o && c > prev_o && o < prev_c) { mask |= (1 << 5); bullish += 1; } + else if (prev_c > prev_o && c < o && c < prev_o && o > prev_c) { mask |= (1 << 6); bearish += 1; } + } + if (body_ratio > 0.85) { + if (c > o) { mask |= (1 << 7); bullish += 1; } else { mask |= (1 << 8); bearish += 1; } + } + if (mask) { + cs_present = 1; + cs_signal = (bullish > bearish) ? 1 : (bearish > bullish) ? -1 : 0; + cs_conf = 50.0; + const double proximity = (atr_kl > 0.0) ? atr_kl * 0.5 : 0.0; + if (proximity > 0.0) { + bool near_support = false, near_resistance = false; + for (int a = 0; a < 3; ++a) { + if ((sl_mask & (1 << a)) && fabs(c - sl[a]) <= proximity) near_support = true; + if ((sh_mask & (1 << a)) && fabs(c - sh[a]) <= proximity) near_resistance = true; + } + if (bullish > bearish && near_support) { cs_conf += 20.0; mask |= (1 << 9); } + else if (bearish > bullish && near_resistance) { cs_conf += 20.0; mask |= (1 << 10); } + } + cs_value = (double)(bullish - bearish); + K1I(cs_present) = 1; K1I(cs_signal) = cs_signal; K1I(cs_mask) = mask; + K1I(cs_bull) = bullish; K1I(cs_bear) = bearish; + K1F(cs_value) = cs_value; K1F(cs_conf) = cs_conf; + } + } + } + + // ── Regime ── + int regime; + if (adx_present && adx_val >= PF[PF_adx_strong_trend]) { + if (ema_present && (ema_signal == 1 || ema_signal == 2)) regime = REG_trending_up; + else if (ema_present && (ema_signal == -1 || ema_signal == -2)) regime = REG_trending_down; + else regime = REG_trending; + } else if (adx_present && slope > PF[PF_adx_slope_emerging_threshold]) { + regime = REG_transitioning; + } else { + regime = REG_ranging; + } + K1I(regime) = regime; + + // ── Regime-aware RSI (replaces the classic vote) ── + if (PI[PI_rsi_regime_aware] && !(dis & DIS_rsi) && rsi_present) { + const double v = rsi_val; + if (regime == REG_trending_up) { + if (38.0 <= v && v <= 55.0) { rsi_signal = 2; rsi_conf = 75.0; rsi_desc = 5; } + else if (30.0 <= v && v < 38.0) { rsi_signal = 1; rsi_conf = 60.0; rsi_desc = 6; } + else if (55.0 < v && v <= 65.0) { rsi_signal = 1; rsi_conf = 50.0; rsi_desc = 7; } + else if (v > 65.0) { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 8; } + else if (v < 30.0) { rsi_signal = 0; rsi_conf = 25.0; rsi_desc = 9; } + } else if (regime == REG_trending_down) { + if (45.0 <= v && v <= 62.0) { rsi_signal = -2; rsi_conf = 75.0; rsi_desc = 10; } + else if (62.0 < v && v <= 70.0) { rsi_signal = -1; rsi_conf = 60.0; rsi_desc = 11; } + else if (35.0 <= v && v < 45.0) { rsi_signal = -1; rsi_conf = 50.0; rsi_desc = 12; } + else if (v < 35.0) { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 13; } + else if (v > 70.0) { rsi_signal = 0; rsi_conf = 25.0; rsi_desc = 14; } + } else if (regime == REG_transitioning || regime == REG_trending) { + if (v >= 70.0) { rsi_signal = -1; rsi_conf = 50.0; rsi_desc = 15; } + else if (v <= 30.0) { rsi_signal = 1; rsi_conf = 50.0; rsi_desc = 16; } + else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 17; } + } else { + if (v >= 70.0) { rsi_signal = -1; rsi_conf = 70.0; rsi_desc = 18; } + else if (v <= 30.0) { rsi_signal = 1; rsi_conf = 70.0; rsi_desc = 19; } + else if (v >= 60.0) { rsi_signal = -1; rsi_conf = 45.0; rsi_desc = 20; } + else if (v <= 40.0) { rsi_signal = 1; rsi_conf = 45.0; rsi_desc = 21; } + else { rsi_signal = 0; rsi_conf = 30.0; rsi_desc = 22; } + } + } + if (rsi_present) { + K1I(rsi_present) = 1; K1I(rsi_signal) = rsi_signal; K1I(rsi_desc) = rsi_desc; + K1F(rsi_value) = rsi_val; K1F(rsi_conf) = rsi_conf; + } + + // ── Composite (insertion order: ema, rsi, bollinger, adx, macd, roc, keltner, divergence, candlestick) ── + double total_score = 0.0, total_weight = 0.0, total_conf = 0.0; + const double* rm = RM + regime * 9; + #define FB2_VOTE(present, sig, conf, base, col, ewf, wsf) \ + if (present) { \ + const double w_ = (base) * rm[col]; \ + const double ws_ = ((double)(sig) * w_) * ((conf) / 100.0); \ + K1F(ewf) = w_; K1F(wsf) = ws_; \ + total_score += ws_; total_weight += w_; total_conf += (conf) * w_; \ + } + FB2_VOTE(ema_present, ema_signal, K1F(ema_conf), 1.0, 0, ew_ema, ws_ema) + FB2_VOTE(rsi_present, rsi_signal, rsi_conf, 1.5, 1, ew_rsi, ws_rsi) + FB2_VOTE(bb_present, bb_signal, bb_conf, 1.5, 2, ew_bb, ws_bb) + FB2_VOTE(adx_present, adx_signal, adx_conf, 0.5, 3, ew_adx, ws_adx) + FB2_VOTE(macd_present, macd_signal, macd_conf, 1.5, 4, ew_macd, ws_macd) + FB2_VOTE(roc_present, roc_signal, roc_conf, 1.5, 5, ew_roc, ws_roc) + FB2_VOTE(kc_present, kc_signal, kc_conf, 1.5, 6, ew_kc, ws_kc) + FB2_VOTE(div_present, div_signal, 70.0, 1.5, 7, ew_div, ws_div) + FB2_VOTE(cs_present, cs_signal, cs_conf, 1.0, 8, ew_cs, ws_cs) + #undef FB2_VOTE + double normalized, avg_conf; + if (total_weight > 0.0) { + normalized = (total_score / total_weight) * 50.0; + avg_conf = total_conf / total_weight; + } else { + normalized = 0.0; avg_conf = 0.0; + } + normalized = PYMAX(-100.0, PYMIN(100.0, normalized)); + int csig; + if (normalized >= 40.0) csig = 2; else if (normalized >= 15.0) csig = 1; + else if (normalized <= -40.0) csig = -2; else if (normalized <= -15.0) csig = -1; else csig = 0; + K1I(composite_signal) = csig; + K1F(composite_score) = normalized; + K1F(composite_conf) = avg_conf; +} +#undef K1F +#undef K1I +''' + +_SRC_K2 = r''' +#define KF(OFx, nx, jj, name) ((OFx)[(size_t)(K1F_##name) * (size_t)(nx) + (size_t)(jj)]) +#define KI(OIx, nx, jj, name) ((OIx)[(size_t)(K1I_##name) * (size_t)(nx) + (size_t)(jj)]) + +// CPython float floor division (float_floor_div / float_divmod). +__device__ __forceinline__ double fb2_py_floordiv(double vx, double wx) { + double mod = fmod(vx, wx); + double div = (vx - mod) / wx; + if (mod != 0.0) { + if ((wx < 0.0) != (mod < 0.0)) { mod += wx; div -= 1.0; } + } + double floordiv; + if (div != 0.0) { + floordiv = floor(div); + if (div - floordiv > 0.5) floordiv += 1.0; + } else { + floordiv = copysign(0.0, vx / wx); + } + return floordiv; +} + +// chart_patterns._first_close_beyond +__device__ __forceinline__ bool fb2_first_close_beyond(const double* c, int n, int start, double level, bool above) { + const double last = c[n - 1]; + const int m = (n - 1) - start; // len(prior) + if (above) { + if (!(last > level)) return false; + if (m <= 0) return true; + double mx = c[start]; + for (int k = start + 1; k < n - 1; ++k) if (c[k] > mx) mx = c[k]; + return mx <= level; + } + if (!(last < level)) return false; + if (m <= 0) return true; + double mn = c[start]; + for (int k = start + 1; k < n - 1; ++k) if (c[k] < mn) mn = c[k]; + return mn >= level; +} + +__device__ __forceinline__ double fb2_max_range(const double* a, int lo, int hi_incl) { + double m = a[lo]; + for (int k = lo + 1; k <= hi_incl; ++k) if (a[k] > m) m = a[k]; + return m; +} +__device__ __forceinline__ double fb2_min_range(const double* a, int lo, int hi_incl) { + double m = a[lo]; + for (int k = lo + 1; k <= hi_incl; ++k) if (a[k] < m) m = a[k]; + return m; +} + +// The last `want` strict pivots (chart_patterns._swing_indices), in order. +__device__ __forceinline__ int fb2_last_pivots(const double* v, int n, int w, bool find_high, int* out, int want) { + int cnt = 0; + for (int i = w; i < n - w; ++i) { + if (!fb2_strict_pivot(v, n, i, w, find_high)) continue; + if (cnt < want) { out[cnt++] = i; } + else { for (int a = 0; a < want - 1; ++a) out[a] = out[a + 1]; out[want - 1] = i; cnt += 1; } + } + return cnt; // total found; out holds the last min(cnt, want) +} + +struct fb2_cp_hit { int pattern; int dir; double anchor; double trigger; double height; }; + +__device__ __forceinline__ bool fb2_cp_double(const double* h, const double* l, const double* c, int n, double atr, + const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) { + const int w = PI[PI_cp_swing_window]; + const double* ext = bullish ? l : h; + int idx[2]; + const int cnt = fb2_last_pivots(ext, n, w, !bullish, idx, 2); + if (cnt < 2) return false; + const int i1 = idx[0], i2 = idx[1]; + if (i2 - i1 < PI[PI_cp_min_sep_bars]) return false; + if ((n - 1) - i2 > PI[PI_cp_recency_bars]) return false; + const double e1 = ext[i1], e2 = ext[i2]; + if (fabs(e1 - e2) > PF[PF_cp_tolerance_atr] * atr) return false; + double neckline, height, floor_; + if (bullish) { + neckline = fb2_max_range(h, i1, i2); + floor_ = PYMIN(e1, e2); + height = neckline - floor_; + } else { + neckline = fb2_min_range(l, i1, i2); + floor_ = PYMAX(e1, e2); + height = floor_ - neckline; + } + if (!(PF[PF_cp_min_height_atr] * atr <= height && height <= PF[PF_cp_max_height_atr] * atr)) return false; + if (bullish) { + if (i2 + 1 < n && fb2_min_range(l, i2 + 1, n - 1) < PYMIN(e1, e2) - 0.1 * atr) return false; + if (!fb2_first_close_beyond(c, n, i2, neckline, true)) return false; + out->pattern = CP_double_bottom; out->dir = 1; out->anchor = PYMIN(e1, e2); + out->trigger = neckline; out->height = height; + return true; + } + if (i2 + 1 < n && fb2_max_range(h, i2 + 1, n - 1) > PYMAX(e1, e2) + 0.1 * atr) return false; + if (!fb2_first_close_beyond(c, n, i2, neckline, false)) return false; + out->pattern = CP_double_top; out->dir = -1; out->anchor = PYMAX(e1, e2); + out->trigger = neckline; out->height = height; + return true; +} + +__device__ __forceinline__ bool fb2_cp_hs(const double* h, const double* l, const double* c, int n, double atr, + const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) { + const int w = PI[PI_cp_swing_window]; + const double* ext = bullish ? l : h; + int idx[3]; + const int cnt = fb2_last_pivots(ext, n, w, !bullish, idx, 3); + if (cnt < 3) return false; + const int i1 = idx[0], i2 = idx[1], i3 = idx[2]; + if ((n - 1) - i3 > PI[PI_cp_recency_bars]) return false; + const double s1 = ext[i1], head = ext[i2], s2 = ext[i3]; + const double hm = PF[PF_cp_hs_head_min_atr] * atr; + double neckline, height; + if (bullish) { + if (!(head < s1 - hm && head < s2 - hm)) return false; + if (fabs(s1 - s2) > PF[PF_cp_hs_shoulder_tol_atr] * atr) return false; + const double p1 = fb2_max_range(h, i1, i2); + const double p2 = fb2_max_range(h, i2, i3); + neckline = PYMIN(p1, p2); + height = neckline - head; + } else { + if (!(head > s1 + hm && head > s2 + hm)) return false; + if (fabs(s1 - s2) > PF[PF_cp_hs_shoulder_tol_atr] * atr) return false; + const double t1 = fb2_min_range(l, i1, i2); + const double t2 = fb2_min_range(l, i2, i3); + neckline = PYMAX(t1, t2); + height = head - neckline; + } + if (!(PF[PF_cp_min_height_atr] * atr <= height && height <= PF[PF_cp_max_height_atr] * atr)) return false; + if (bullish) { + if (i3 + 1 < n && fb2_min_range(l, i3 + 1, n - 1) < head - 0.1 * atr) return false; + if (!fb2_first_close_beyond(c, n, i3, neckline, true)) return false; + out->pattern = CP_inv_hs; out->dir = 1; out->anchor = s2; out->trigger = neckline; out->height = height; + return true; + } + if (i3 + 1 < n && fb2_max_range(h, i3 + 1, n - 1) > head + 0.1 * atr) return false; + if (!fb2_first_close_beyond(c, n, i3, neckline, false)) return false; + out->pattern = CP_hs_top; out->dir = -1; out->anchor = s2; out->trigger = neckline; out->height = height; + return true; +} + +__device__ __forceinline__ bool fb2_cp_flag(const double* h, const double* l, const double* c, int n, double atr, + const double* PF, const int* PI, bool bullish, fb2_cp_hit* out) { + const double lo_r = PF[PF_cp_flag_retrace_lo], hi_r = PF[PF_cp_flag_retrace_hi]; + if (bullish) { + int ip = 0; // np.argmax(h[:-1]) — first occurrence + for (int k = 1; k < n - 1; ++k) if (h[k] > h[ip]) ip = k; + const int flag_len = (n - 1) - ip; + if (!(PI[PI_cp_flag_min_bars] <= flag_len && flag_len <= PI[PI_cp_flag_max_bars])) return false; + const int pole_start = PYMAX(0, ip - PI[PI_cp_flag_pole_bars] + 1); + const double pole_low = fb2_min_range(l, pole_start, ip); + const double pole = h[ip] - pole_low; + if (pole < PF[PF_cp_flag_pole_atr] * atr) return false; + const double flag_low = fb2_min_range(l, ip + 1, n - 1); + const double retrace = (h[ip] - flag_low) / pole; + if (!(lo_r <= retrace && retrace <= hi_r)) return false; + if (!fb2_first_close_beyond(c, n, ip, h[ip], true)) return false; + out->pattern = CP_bull_flag; out->dir = 1; out->anchor = flag_low; out->trigger = h[ip]; out->height = pole; + return true; + } + int ip = 0; // np.argmin(l[:-1]) + for (int k = 1; k < n - 1; ++k) if (l[k] < l[ip]) ip = k; + const int flag_len = (n - 1) - ip; + if (!(PI[PI_cp_flag_min_bars] <= flag_len && flag_len <= PI[PI_cp_flag_max_bars])) return false; + const int pole_start = PYMAX(0, ip - PI[PI_cp_flag_pole_bars] + 1); + const double pole_high = fb2_max_range(h, pole_start, ip); + const double pole = pole_high - l[ip]; + if (pole < PF[PF_cp_flag_pole_atr] * atr) return false; + const double flag_high = fb2_max_range(h, ip + 1, n - 1); + const double retrace = (flag_high - l[ip]) / pole; + if (!(lo_r <= retrace && retrace <= hi_r)) return false; + if (!fb2_first_close_beyond(c, n, ip, l[ip], false)) return false; + out->pattern = CP_bear_flag; out->dir = -1; out->anchor = flag_high; out->trigger = l[ip]; out->height = pole; + return true; +} + +// chart_patterns.detect_chart_pattern on the prefix [0..pos] of a frame. +__device__ __forceinline__ bool fb2_detect_chart_pattern(const double* high, const double* low, const double* close, + int pos, double atr, const double* PF, const int* PI, + fb2_cp_hit* out) { + const int L = PI[PI_cp_lookback] + 1; + if (atr <= 0.0 || pos + 1 < L) return false; + const int base = pos + 1 - L; + const double* h = high + base; const double* l = low + base; const double* c = close + base; + const int mask = PI[PI_cp_patterns_mask]; + if ((mask & (1 << CP_inv_hs)) && fb2_cp_hs(h, l, c, L, atr, PF, PI, true, out)) return true; + if ((mask & (1 << CP_hs_top)) && fb2_cp_hs(h, l, c, L, atr, PF, PI, false, out)) return true; + if ((mask & (1 << CP_double_bottom)) && fb2_cp_double(h, l, c, L, atr, PF, PI, true, out)) return true; + if ((mask & (1 << CP_double_top)) && fb2_cp_double(h, l, c, L, atr, PF, PI, false, out)) return true; + if ((mask & (1 << CP_bull_flag)) && fb2_cp_flag(h, l, c, L, atr, PF, PI, true, out)) return true; + if ((mask & (1 << CP_bear_flag)) && fb2_cp_flag(h, l, c, L, atr, PF, PI, false, out)) return true; + return false; +} + +// signal_aggregator._detect_big_candle on the prefix [0..pos]; returns the +// absolute index of the bar (first occurrence of the max ratio) or -1. +__device__ __forceinline__ int fb2_big_candle(const double* high, const double* low, const double* close, int pos, + const double* PF, const int* PI, double* ratio_out) { + const double hours = PF[PF_big_candle_lookback_hours]; + const int bar_secs = PI[PI_primary_bar_secs]; + int nbc = (int)fb2_py_floordiv(hours * 3600.0, (double)bar_secs); + if (nbc < 1) nbc = 1; + const int period = PI[PI_atr_period]; + int tail_len = nbc + 3 * period; + if (tail_len > pos + 1) tail_len = pos + 1; + const int base = pos + 1 - tail_len; + // pandas ewm state (alpha = 1/period, min_periods = period) + const double alpha = 1.0 / (double)period; + const double com = (1.0 - alpha) / alpha; + const double a = 1.0 / (1.0 + com); + const double old_wt_factor = 1.0 - a, new_wt = a; + double w = FB2_NAN, old_wt = 1.0; int nobs = 0; + double prev_atr = FB2_NAN; // atr.shift(1) + double best = FB2_NAN; int best_idx = -1; + const int win_start = tail_len - nbc; // ratio.iloc[-n:] + for (int i = 0; i < tail_len; ++i) { + const int k = base + i; + const double tr1 = high[k] - low[k]; + double tr = tr1; + if (i > 0) { + const double pc = close[k - 1]; + const double tr2 = fabs(high[k] - pc), tr3 = fabs(low[k] - pc); + if (tr2 > tr) tr = tr2; + if (tr3 > tr) tr = tr3; + } + // ratio for this bar uses the PRIOR bar's atr output + if (i >= win_start) { + const double ratio = tr1 / prev_atr; + if (ratio == ratio) { + if (best_idx < 0 || ratio > best) { best = ratio; best_idx = k; } + } + } + // ewm update + double out; + if (i == 0) { + w = tr; nobs = (w == w) ? 1 : 0; old_wt = 1.0; + out = (nobs >= period) ? w : FB2_NAN; + } else { + const bool is_obs = (tr == tr); + if (is_obs) nobs += 1; + if (w == w) { + old_wt *= old_wt_factor; + if (is_obs) { + if (w != tr) { w = old_wt * w + new_wt * tr; w /= (old_wt + new_wt); } + old_wt = 1.0; + } + } else if (is_obs) { w = tr; } + out = (nobs >= period) ? w : FB2_NAN; + } + prev_atr = out; + } + if (best_idx < 0) return -1; + if (best < PF[PF_big_candle_atr_multiple]) return -1; + *ratio_out = best; + return best_idx; +} + +extern "C" __global__ +void fb2_aggregate(const double* __restrict__ open_p, const double* __restrict__ high_p, + const double* __restrict__ low_p, const double* __restrict__ close_p, const int n_p, + const double* __restrict__ OFp, const int* __restrict__ OIp, + const double* __restrict__ OFc0, const int* __restrict__ OIc0, const long long* __restrict__ cut0, const int nc0, + const double* __restrict__ OFc1, const int* __restrict__ OIc1, const long long* __restrict__ cut1, const int nc1, + const double* __restrict__ OFc2, const int* __restrict__ OIc2, const long long* __restrict__ cut2, const int nc2, + const double* __restrict__ OFc3, const int* __restrict__ OIc3, const long long* __restrict__ cut3, const int nc3, + const double* __restrict__ rf_adx, const long long* __restrict__ rf_cut, + const signed char* __restrict__ skip_reason, const signed char* __restrict__ label_hour, + const signed char* __restrict__ label_wd, const signed char* __restrict__ close_hour, + const double* __restrict__ PF, const int* __restrict__ PI, + const int first_pos, const int n_bars, + double* __restrict__ OF2, int* __restrict__ OI2, + short* __restrict__ warn_code, double* __restrict__ warn_arg) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n_bars) return; + const int pos = first_pos + i; +#define K2F(name) OF2[(size_t)(K2F_##name) * (size_t)n_bars + (size_t)i] +#define K2I(name) OI2[(size_t)(K2I_##name) * (size_t)n_bars + (size_t)i] +#define PRF(name) KF(OFp, n_p, pos, name) +#define PRI(name) KI(OIp, n_p, pos, name) + for (int f = 0; f < K2F_N; ++f) OF2[(size_t)f * (size_t)n_bars + (size_t)i] = FB2_NAN; + for (int f = 0; f < K2I_N; ++f) OI2[(size_t)f * (size_t)n_bars + (size_t)i] = 0; + int wn = 0; +#define WARN(code, arg) do { if (wn < FB2_MAX_WARN) { warn_code[(size_t)i * FB2_MAX_WARN + wn] = (short)(code); warn_arg[(size_t)i * FB2_MAX_WARN + wn] = (double)(arg); wn += 1; } } while (0) + + K2I(bc_idx) = -1; + K2I(skip_reason) = skip_reason[pos]; + if (skip_reason[pos] != 0) { K2I(status) = 0; return; } + + const int is_xau = PI[PI_is_xau]; + const int xau_enabled = PI[PI_xau_enabled]; + const double price = close_p[pos]; + + // ── primary analysis ── + const double primary_score = PRF(composite_score); + const double primary_conf = PRF(composite_conf); + const int primary_signal = PRI(composite_signal); + const int regime = PRI(regime); + const int vol = PRI(vol); + const int atr_present = PRI(atr_present); + const double atr_kl = atr_present ? PRF(atr) : 0.0; + const int rsi_present = PRI(rsi_present); + const double rsi_val = PRF(rsi_value); + const int bb_present = PRI(bb_present); + const double bb_pos = PRF(bb_value); + const bool trending_any = (regime == REG_trending_up || regime == REG_trending_down || regime == REG_trending); + K2F(primary_score) = primary_score; K2F(primary_conf) = primary_conf; + + // ── confirmation analyses (as-of aligned) ── + const double* OFc[FB2_MAX_CONFIRM] = {OFc0, OFc1, OFc2, OFc3}; + const int* OIc[FB2_MAX_CONFIRM] = {OIc0, OIc1, OIc2, OIc3}; + const long long* cuts[FB2_MAX_CONFIRM] = {cut0, cut1, cut2, cut3}; + const int ncs[FB2_MAX_CONFIRM] = {nc0, nc1, nc2, nc3}; + const int n_confirm = PI[PI_n_confirm]; + int sigs[1 + FB2_MAX_CONFIRM]; double scores[1 + FB2_MAX_CONFIRM]; + int n_sigs = 1; sigs[0] = primary_signal; scores[0] = primary_score; + for (int k = 0; k < n_confirm; ++k) { + const long long cut = cuts[k][pos]; + if (cut == 0) continue; + const long long idx = cut - 1; + sigs[n_sigs] = KI(OIc[k], ncs[k], idx, composite_signal); + scores[n_sigs] = KF(OFc[k], ncs[k], idx, composite_score); + n_sigs += 1; + } + K2I(n_confirm_present) = n_sigs - 1; + + // ── 3. MTF agreement ── + const int thr = PI[PI_mtf_agreement_threshold]; + bool mtf = false; + if (n_sigs >= 2) { + int bull = 0, bear = 0; + for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull += 1; else if (sigs[q] < 0) bear += 1; } + if (primary_score > 0.0) mtf = (bull >= thr); + else if (primary_score < 0.0) mtf = (bear >= thr); + else mtf = (bull >= thr) || (bear >= thr); + } + K2I(mtf_agreement) = mtf ? 1 : 0; + + // ── 6. _compute_final_signal ── + double score = primary_score; + double confidence = primary_conf; + if (mtf) { + confidence += 10.0; + // Python 3.12 builtin sum() over floats: int 0 + s0 exactly, then the + // Neumaier-compensated accumulation of the remaining items. + double f = scores[0], c = 0.0; + for (int q = 1; q < n_sigs; ++q) { + const double x = scores[q]; + const double t = f + x; + if (fabs(f) >= fabs(x)) c += (f - t) + x; + else c += (x - t) + f; + f = t; + } + if (c != 0.0 && isfinite(c)) f += c; + const double mtf_avg = f / (double)n_sigs; + score = score * 0.6 + mtf_avg * 0.4; + } + if (trending_any) { + if (rsi_present) { + if (regime == REG_trending_up && score > 0.0 && 38.0 <= rsi_val && rsi_val <= 55.0) confidence += 5.0; + else if (regime == REG_trending_down && score < 0.0 && 45.0 <= rsi_val && rsi_val <= 62.0) confidence += 5.0; + } + if (bb_present) { + if (regime == REG_trending_up && score > 0.0 && bb_pos < 0.4) confidence += 5.0; + else if (regime == REG_trending_down && score < 0.0 && bb_pos > 0.6) confidence += 5.0; + } + } + int pmask = 0; + if (!mtf) { + pmask |= (1 << PEN_mtf_disagreement); + if (n_sigs >= 2) { + bool any_pos = false, any_neg = false; + for (int q = 0; q < n_sigs; ++q) { if (scores[q] > 10.0) any_pos = true; if (scores[q] < -10.0) any_neg = true; } + if (any_pos && any_neg) WARN(W_conflict, 0.0); + } + } + if ((regime == REG_trending_up && score < 0.0) || (regime == REG_trending_down && score > 0.0)) { + WARN(W_counter_trend, 0.0); + pmask |= (1 << PEN_counter_trend); + } + if (vol == VOL_extreme) { WARN(W_extreme_vol, 0.0); pmask |= (1 << PEN_extreme_vol); } + else if (vol == VOL_low) { WARN(W_low_vol, 0.0); } + if (PRI(ema_trend_present)) { + const double ema200 = PRF(ema_trend); + if (ema200 != 0.0 && ema200 > 0.0 && price > 0.0) { + const double ema_dist_pct = fabs(price - ema200) / ema200 * 100.0; + if (ema_dist_pct < PF[PF_fb2_ema_lo]) { WARN(W_ema200_near, ema_dist_pct); pmask |= (1 << PEN_ema200_near); } + else if (ema_dist_pct > 2.5) { WARN(W_ema200_far, ema_dist_pct); pmask |= (1 << PEN_ema200_far); } + } + } + const int sh_mask = PRI(sh_mask), sl_mask = PRI(sl_mask); + const double sh[3] = {PRF(sh1), PRF(sh2), PRF(sh3)}; + const double sl[3] = {PRF(sl1), PRF(sl2), PRF(sl3)}; + if (atr_kl > 0.0 && score != 0.0) { + if (score > 0.0) { + bool any = false; double mn = 0.0; + for (int a = 0; a < 3; ++a) if (sl_mask & (1 << a)) { if (!any || sl[a] < mn) mn = sl[a]; any = true; } + if (any && price - mn > 3.5 * atr_kl) { WARN(W_exhaustion_low, 0.0); pmask |= (1 << PEN_exhaustion); } + } else { + bool any = false; double mx = 0.0; + for (int a = 0; a < 3; ++a) if (sh_mask & (1 << a)) { if (!any || sh[a] > mx) mx = sh[a]; any = true; } + if (any && mx - price > 3.5 * atr_kl) { WARN(W_exhaustion_high, 0.0); pmask |= (1 << PEN_exhaustion); } + } + } + if (PRI(kl_ema_present)) { + const double ema_medium = PRF(ema_medium); + if (atr_kl > 0.0 && ema_medium != 0.0 && trending_any) { + const double pullback_dist = fabs(price - ema_medium) / atr_kl; + if (pullback_dist > 2.5) { WARN(W_no_pullback, pullback_dist); pmask |= (1 << PEN_no_pullback); } + else if (pullback_dist > 2.0) { WARN(W_weak_pullback, pullback_dist); pmask |= (1 << PEN_weak_pullback); } + } + } + if (trending_any) { + const bool up_or_mixed = (regime == REG_trending_up || regime == REG_trending); + const bool down_or_mixed = (regime == REG_trending_down || regime == REG_trending); + if (rsi_present) { + if (up_or_mixed && score > 0.0 && rsi_val > 65.0) { WARN(W_rsi_ext_up, rsi_val); pmask |= (1 << PEN_rsi_extended); } + else if (down_or_mixed && score < 0.0 && rsi_val < 35.0) { WARN(W_rsi_ext_down, rsi_val); pmask |= (1 << PEN_rsi_extended); } + } + if (bb_present) { + if (up_or_mixed && score > 0.0 && bb_pos > 0.85) pmask |= (1 << PEN_bb_overextended); + else if (down_or_mixed && score < 0.0 && bb_pos < 0.15) pmask |= (1 << PEN_bb_overextended); + } + } + int exhaustion_count = 0; + if (pmask & (1 << PEN_exhaustion)) exhaustion_count += 1; + if (pmask & (1 << PEN_no_pullback)) exhaustion_count += 1; + if (pmask & (1 << PEN_rsi_extended)) exhaustion_count += 1; + if (pmask & (1 << PEN_bb_overextended)) exhaustion_count += 1; + bool momentum_fading = false; + if (exhaustion_count >= 2) { + bool roc_fading = false, macd_fading = false; + if (PRI(roc_present)) { + const double rv = PRF(roc_value); + if (score > 0.0 && rv < 0.0) roc_fading = true; + else if (score < 0.0 && rv > 0.0) roc_fading = true; + } + if (PRI(macd_present)) { + const int ms = PRI(macd_signal); + if (score > 0.0 && ms < 0) macd_fading = true; + else if (score < 0.0 && ms > 0) macd_fading = true; + } + momentum_fading = roc_fading || macd_fading; + } + const bool exhaustion_blocked = (exhaustion_count >= 2) && momentum_fading; + // decay pass in the stable-sorted order (15s, 10, 8s, 5s; append order inside a bucket) + const int order[10] = {PEN_mtf_disagreement, PEN_counter_trend, PEN_extreme_vol, + PEN_exhaustion, PEN_no_pullback, PEN_bb_overextended, + PEN_ema200_near, PEN_ema200_far, PEN_weak_pullback, PEN_rsi_extended}; + const double amounts[10] = {15.0, 15.0, 10.0, 8.0, 8.0, 8.0, 5.0, 5.0, 5.0, 5.0}; + const double MAXP = PF[PF_max_penalty]; + double total_penalty = 0.0, decay = 1.0; + for (int q = 0; q < 10; ++q) { + if (!(pmask & (1 << order[q]))) continue; + double applied = amounts[q] * decay; + if (total_penalty + applied > MAXP) applied = MAXP - total_penalty; + total_penalty += applied; + decay *= 0.5; + if (total_penalty >= MAXP) break; + } + confidence -= total_penalty; + confidence = PYMAX(0.0, PYMIN(95.0, confidence)); + if (exhaustion_blocked) { WARN(W_exhaustion_block, (double)exhaustion_count); confidence = 0.0; } + int final_signal; + if (score >= 35.0 && confidence >= 55.0) final_signal = 2; + else if (score >= 25.0 && confidence >= 50.0) final_signal = 1; + else if (score <= -35.0 && confidence >= 55.0) final_signal = -2; + else if (score <= -25.0 && confidence >= 50.0) final_signal = -1; + else final_signal = 0; + const double abs_score0 = fabs(score); + int quality; + if (abs_score0 >= 40.0 && confidence >= 65.0 && mtf) quality = Q_APLUS; + else if (abs_score0 >= 33.0 && confidence >= 55.0 && mtf) quality = Q_A; + else if (abs_score0 >= 28.0 && confidence >= 50.0) quality = Q_B; + else if (abs_score0 >= 25.0 && confidence >= 45.0) quality = Q_C; + else if (abs_score0 >= 15.0) quality = Q_D; + else quality = Q_F; + K2I(penalty_mask) = pmask; K2F(penalty_total) = total_penalty; + K2I(exhaustion_count) = exhaustion_count; K2I(momentum_fading) = momentum_fading ? 1 : 0; + + // ── 6b. session / day filter on the bar LABEL ── + if (PI[PI_session_filter]) { + const int bar_hour = label_hour[pos]; + const int bar_wd = label_wd[pos]; + const int hmask = is_xau ? PI[PI_xau_blocked_hours_mask] : PI[PI_bt_blocked_hours_mask]; + if (hmask & (1 << bar_hour)) { final_signal = 0; quality = Q_F; WARN(W_session_filter, (double)bar_hour); } + const int wmask = is_xau ? PI[PI_xau_allowed_weekdays_mask] : PI[PI_bt_allowed_weekdays_mask]; + if (wmask != 0 && !(wmask & (1 << bar_wd))) { final_signal = 0; quality = Q_F; WARN(W_day_filter, (double)bar_wd); } + } + // ── 6c. transitioning boost ── + if (regime == REG_transitioning && (quality == Q_B || quality == Q_C || quality == Q_D)) { + const double abs_score = fabs(score); + if (abs_score >= 30.0 && confidence >= 50.0 && mtf) quality = Q_A; + else if (abs_score >= 25.0 && confidence >= 45.0) quality = Q_B; + } + // ── 6d. exhausted trend cap ── + if (trending_any && exhaustion_count >= 2 && (quality == Q_C || quality == Q_D)) { + WARN(W_exhausted_trend, 0.0); quality = Q_D; + } + const bool xau_is_sell_dir = (final_signal < 0) || (final_signal == 0 && score < 0.0); +#define ACTIONABLE(q) ((q) >= Q_C) + // ── 6e. SELL quality gate ── + if (xau_enabled && is_xau && xau_is_sell_dir && ACTIONABLE(quality) + && !(PI[PI_sell_min_quality_mask] & (1 << quality))) { + WARN(W_sell_quality, (double)quality); quality = Q_D; + } + // ── 6f. MTF opposition ── + if (ACTIONABLE(quality)) { + int bull_n = 0, bear_n = 0; + for (int q = 0; q < n_sigs; ++q) { if (sigs[q] > 0) bull_n += 1; else if (sigs[q] < 0) bear_n += 1; } + const bool opposed = (bear_n >= thr && score > 0.0) || (bull_n >= thr && score < 0.0); + if (opposed) { WARN(W_mtf_opposition, 0.0); quality = Q_D; } + } + // ── 6g. SELL volatility gate ── + if (xau_enabled && is_xau && xau_is_sell_dir && ACTIONABLE(quality) + && !(PI[PI_sell_vol_mask] & (1 << vol))) { + WARN(W_sell_vol, (double)vol); quality = Q_D; + } + // ── 6g2 / 6g3. regime and SELL-hour gates ── + if (xau_enabled && is_xau && ACTIONABLE(quality)) { + if (PI[PI_blocked_regimes_mask] & (1 << regime)) { WARN(W_regime_blocked, (double)regime); quality = Q_D; } + else if (xau_is_sell_dir && (PI[PI_sell_blocked_regimes_mask] & (1 << regime))) { WARN(W_sell_regime, (double)regime); quality = Q_D; } + else if (xau_is_sell_dir && PI[PI_sell_blocked_hours_mask] != 0) { + const int ch = close_hour[pos]; + if (PI[PI_sell_blocked_hours_mask] & (1 << ch)) { WARN(W_sell_hour, (double)ch); quality = Q_D; } + } + } + // ── 6h. big-candle gate ── + bool big_candle_hit = false; + if (xau_enabled && is_xau && PF[PF_big_candle_atr_multiple] > 0.0) { + double ratio = 0.0; + const int idx = fb2_big_candle(high_p, low_p, close_p, pos, PF, PI, &ratio); + if (idx >= 0) { + big_candle_hit = true; + K2I(bc_idx) = idx; K2F(bc_ratio) = ratio; + WARN(W_big_candle, ratio); + if (ACTIONABLE(quality)) quality = Q_D; + } + } + // ── 6i. M5 noise gates ── + if (PI[PI_primary_is_m5] && is_xau && ACTIONABLE(quality)) { + const double abs_score = fabs(score); + if (abs_score < PF[PF_m5_min_abs_score]) { WARN(W_m5_score, abs_score); quality = Q_D; } + else if (PI[PI_m5_require_mtf] && !mtf) { WARN(W_m5_mtf, 0.0); quality = Q_D; } + else { + const double minc = PF[PF_m5_min_confirmation_score]; + const bool is_buy = score > 0.0; + bool confirmed = false; + for (int q = 1; q < n_sigs; ++q) { + if (is_buy && scores[q] >= minc) { confirmed = true; break; } + if (!is_buy && scores[q] <= -minc) { confirmed = true; break; } + } + if (!confirmed) { WARN(W_m5_conviction, 0.0); quality = Q_D; } + else { + const bool is_sell = score < 0.0; + if (is_sell && !(PI[PI_m5_sell_max_quality_mask] & (1 << quality))) { WARN(W_m5_sell_quality, (double)quality); quality = Q_D; } + else if (PI[PI_m5_blocked_vol_mask] & (1 << vol)) { WARN(W_m5_vol, (double)vol); quality = Q_D; } + else if (is_sell && rsi_present && rsi_val < PF[PF_m5_sell_min_rsi]) { WARN(W_m5_sell_rsi, rsi_val); quality = Q_D; } + } + } + } + + // ── 7. setups ── + int strategy = STRAT_NONE, direction = 0; + double setup_entry = FB2_NAN, setup_sl = FB2_NAN, setup_tp1 = FB2_NAN, setup_tp2 = FB2_NAN, rr_rounded = FB2_NAN; + const double pip_value = PF[PF_pip_value]; + const double min_rr = PF[PF_min_rr_ratio]; + const double tfac = PF[PF_tp_tightening_factor]; + const double min_tp = PF[PF_round_trip_cost] * PF[PF_min_tp_cost_multiple]; + if (quality >= Q_A && is_xau) { + const double atr = atr_kl; + if (atr > 0.0) { + const bool is_buy = (final_signal > 0) ? true : (final_signal < 0) ? false : (score >= 0.0); + double quality_scale = (quality == Q_APLUS) ? PF[PF_tp_scale_aplus] : PF[PF_tp_scale_a]; + if (quality == Q_A && exhaustion_count >= 2) quality_scale = 1.0; + const double tp_mult = PF[PF_tp_atr_multiplier] * quality_scale; + double sl_, tp1, tp2; + if (is_buy) { sl_ = price - PF[PF_sl_atr_multiplier] * atr; tp1 = price + tp_mult * atr; tp2 = price + tp_mult * 1.5 * atr; } + else { sl_ = price + PF[PF_sl_atr_multiplier] * atr; tp1 = price - tp_mult * atr; tp2 = price - tp_mult * 1.5 * atr; } + tp1 = price + (tp1 - price) * tfac; + tp2 = price + (tp2 - price) * tfac; + const double sl_distance = fabs(price - sl_); + const double tp_distance = fabs(tp1 - price); + const 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 { + 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); + const double tp1_distance = fabs(tp1 - entry); + const 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 { + 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; + } + const double rr = (sl_dist > 0.0) ? height / sl_dist : 0.0; + if (!(rr < min_rr - 1e-9) && !(height < min_tp)) { + 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() diff --git a/backtest/instructions_for_backtest.txt b/backtest/instructions_for_backtest.txt index 30e9146..e43d084 100644 --- a/backtest/instructions_for_backtest.txt +++ b/backtest/instructions_for_backtest.txt @@ -46,3 +46,22 @@ re-implementation. - Any change to a kernel or to its CPU twin is a parity event: re-run Module 36, regenerate the five reference books on BOTH backends and show the ten MD5s unchanged. Never replace a kernel with a differently ordered algorithm (parallel scan, two-pass variance) even if it is faster. + +--- 2026-09-18: CUDA research twin (backtest/cudaBasedBacktesting.py) --- +- Exemption from "replays go through the real signal code" above, granted by the owner for this + ONE research script: its per-bar decision is a CUDA re-implementation of engine.py, + signal_aggregator.py and chart_patterns.py (backtest/cuda_signal_kernels.py). gpu_backtest.py + stays the base reference and live scans never use the twin. +- Its numbers count as evidence only while parity holds. After ANY change to engine.py, + signal_aggregator.py, chart_patterns.py, position_sizing.py, the recorder's row layout or the + GPUTechnicalEngine overrides: mirror the change in cuda_signal_kernels.py in the same commit, + keep test_all.py Module 37 green, and regenerate at least the M15 OOS book (2024-06-01 -> + 2025-03-30, $0.65) with BOTH scripts and show equal MD5s before quoting a CUDA-script result. +- A report built from the twin quotes both scripts' "Backend:" lines (the twin's necessarily + starts with "CuPy GPU (") and the paired MD5 of at least one book from the reference script on + the same code and cache. Whole-window and half-window figures, costs and breakdowns are then + read from the twin's book exactly as from the reference's; the CSVs are identical. +- The twin has no CPU path (--backend cpu / FB2_BACKEND=cpu are refused), --workers is ignored, + and it refuses configurations beyond its compile-time limits (more than 4 confirmation + timeframes, divergence lookback >= 64, swing lookback >= 256, cp_lookback >= 128) with a + message that points back to gpu_backtest.py. diff --git a/backtest/research/cuda_backtester_2026-09-18.md b/backtest/research/cuda_backtester_2026-09-18.md new file mode 100644 index 0000000..c783344 --- /dev/null +++ b/backtest/research/cuda_backtester_2026-09-18.md @@ -0,0 +1,215 @@ +# Full-GPU research twin of the backtester — 2026-09-18 + +Behaviour-neutral by construction and by proof. `backtest/cudaBasedBacktesting.py` runs every stage +of a backtest on the GPU, including the per-bar decision that `gpu_backtest.py` still runs through +the live Python signal code. Nothing in the signal logic, the weights, the gates or the exit rules +moved; the proof is that both scripts export byte-identical books (same CSV MD5) on every window +and every research flag tried, listed in section 3. + +## 1. Owner decisions + +- Research and backtesting only. Live scans stay on the CPU signal code; the Telegram bot, the CLI + and the dashboard never import the new files (Module 37 sweeps the live roots). +- `gpu_backtest.py` stays in the tree as the **base reference**. The twin must reproduce its + exported book byte for byte for the same arguments; the MD5 of the CSV is the acceptance gate. +- `--record`, `--record-bars`, `--export-features` and `--fixed-rr-tp` all work on the GPU path. +- Acceptance: 6-month M5 and M15 books, window 2026-03-18 → 2026-09-18, both scripts, identical + CSV MD5 at $0.65 (headline) and $0.30 (comparability). That window was not cached; the first + control run fetched M5/M15/H1 from Oanda (padded 62 days, `oanda_XAU_USD__20260115_20260918`) + and both scripts read the same files afterwards. +- `backtest/instructions_for_backtest.txt` item 6 ("replays go through the real signal code, never + a re-implementation") is waived for this ONE research script; the dated addendum in that file + records the exemption and the new obligation: any signal change is a parity event, re-prove the + MD5 with both scripts before a number from the twin counts as evidence. + +## 2. What shipped + +| File | Role | +|---|---| +| `backtest/cudaBasedBacktesting.py` (1,028 lines) | Same CLI as the reference (it reuses `_parse_args`), same cache files, same report and CSV writers. Reuses by import everything that is not the replay: the loader and its cleaning, `precompute_indicators` (GPU), the trade resolver (GPU), `_reanchor_entries_to_raw`, `_apply_fixed_rr_tp`, `compute_metrics`, `print_report`, `trades_to_dataframe`, `replay_record.write_run`. Replaces `_replay_span` + the worker pool + `_apply_cooldown` with kernel launches and host builders that turn the kernels' numeric codes into the exact strings (indicator descriptions, warnings, reasons) with the same f-string templates as `engine.py` / `signal_aggregator.py`, then call the real `SignalAggregator._clean_accepted_cp_warnings` and `_drop_stale_spread_gate_vetoes`. No CPU path: `--backend cpu`, `FB2_BACKEND=cpu` or a failed probe exit with a message that points to `gpu_backtest.py`; configurations beyond the compile-time limits (more than 4 confirmation frames, divergence lookback ≥ 64, swing lookback ≥ 256, `cp_lookback` ≥ 128, a big-candle window over 500 bars) are refused the same way. | +| `backtest/cuda_signal_kernels.py` (2,914 lines, most of it CUDA C in strings) | K1 `fb2_analyze_tf` (one thread per bar of one timeframe; the nine `GPUTechnicalEngine` overrides plus the live engine code they do not cover: divergence, structure S/R, candlestick, regime, regime-aware RSI, composite, the `len < 30` empty analysis, `ema_mode="2ema"`). K2 `fb2_aggregate` (one thread per primary bar; `analyze_from_dataframes` end to end: the multi-timeframe merge, bonuses, penalties in the post-sort decay order with the 30 cap, exhaustion block, signal, grade, gates 6b–6i, the XAU_ATR / GENERIC / XAU_PB / XAU_CP generators, the NEUTRAL-at-grade-C direction, `_apply_regime_filter_live`, plus the replay's quality / strategy filters, next-open entry and distance-preserving re-anchor). K3 `fb2_features` (the 65 `position_sizing.build_features` columns per accepted trade). K4 cooldown (single thread, the only cross-bar rule). K5 prefix median (the snapshot's `spread_estimate`). K6 recorder bar rows (the 160-column float32 row). The config packer (`PARAMS_F` / `PARAMS_I` name lists → generated `#define`s, so a name cannot drift), the classification table of every config field the kernels do not read, and thin test launches of the device helpers. Compiled with `--fmad=false`; imports only numpy, `gpu_kernels`, `config`, `chart_patterns` (pattern priority order) and `position_sizing` constants. | +| `test_all.py` Module 37 (16 tests) | Section 4. | +| `backtest/instructions_for_backtest.txt` | The 2026-09-18 CUDA-twin addendum. | + +What stays on the host, on purpose: loading and cleaning the parquet cache; the calendar and +session flags (pandas time-zone arithmetic over the whole index: NY weekend, UTC weekday, blocked +close hours, and the three clocks the gates use: bar label for gate 6b, bar close for gate 6g3 and +XAU_CP, decision time for the replay skip); the text; the writers. + +Rules that keep the twins bit-identical (they cost most of the debugging time and are now pinned +by tests): doubles everywhere; Python `min()` / `max()` as argument-order ternaries, never +`fmin` / `fmax`; every sum in Python's iteration order; numpy's pairwise summation (8 accumulators, +128-blocks) for `np.mean` / `np.nanmean`; Python 3.12's Neumaier-compensated builtin `sum()` for +the multi-timeframe score mean; pandas' online rolling mean and EWM recurrences copied from +`gpu_kernels.py` for the restarted tails the features use; `round(x, 2)` reproduced exactly with an +error-free product (the host additionally re-checks Python's `round()` on the device doubles of +every accepted trade); `atr_percentile = count / 50 * 100` (not `count * 2`: 15/50·100 = +30.000000000000004 flips "low" to "normal"); the as-of alignment compares int64 nanoseconds. + +## 3. Evidence + +Environment: NVIDIA GeForce RTX 4070 SUPER, CuPy 14.0.1 (`cupy-cuda12x` with the +`nvidia-cuda-runtime-cu12` header wheel), NumPy 2.2.6, pandas 3.0.2, Python 3.12, 16 CPUs. Code: +this working tree on branch `sql-integration` on top of `bac11d8`. `Backend:` lines verbatim: +reference `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)`, twin +`CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)`; both log `GPU cache fallbacks: 0` (the +twin's line adds "no per-bar recompute path in cudaBasedBacktesting"). + +### 3.1 Acceptance: 6 months, both scripts, both costs + +Harness: `venv/bin/python backtest/gpu_backtest.py --pair XAUUSD --primary-tf --start +2026-03-18 --end 2026-09-18 --spread --workers 12 --export-csv ` for the control and +`venv/bin/python backtest/cudaBasedBacktesting.py` with the same arguments (no `--workers`) for the +twin; factory config, no overrides. + +| Primary | Cost | Control MD5 | Twin MD5 | Equal | Control wall | Twin wall | +|---|---|---|---|---|---|---| +| M5 | $0.65 | `86111c36c1a07e543a65f0cd09d56ebf` | `86111c36c1a07e543a65f0cd09d56ebf` | ✓ | 51.8 s (replay 43.8 s) | 1.27 s | +| M5 | $0.30 | `e70d025205844ecedc2b9e9c08f38f77` | `e70d025205844ecedc2b9e9c08f38f77` | ✓ | 47.1 s (replay 45.9 s) | 1.65 s | +| M15 | $0.65 | `b8c980833f1a32bd59930bd037c75b89` | `b8c980833f1a32bd59930bd037c75b89` | ✓ | 20.4 s (replay 18.7 s) | 1.24 s | +| M15 | $0.30 | `e66b5fbdaace1d3fda4a76a0e2d47d69` | `e66b5fbdaace1d3fda4a76a0e2d47d69` | ✓ | 14.8 s (replay 13.5 s) | 1.50 s | + +Because the CSVs are identical the figures below are the same from either script (`NO_DATA` rows +excluded; halves split at 2026-06-18). They are reported because the instructions require them, +not as a validation: nothing changed, and the window is 6 months (rule 1 needs 12). + +| Primary | Cost | Segment | Trades | Total R | Win rate | PF | Max DD (R) | +|---|---|---|---|---|---|---|---| +| M5 | $0.65 | whole | 283 | +41.67 | 56.5% | 1.42 | 10.1 | +| M5 | $0.65 | first half | 146 | +29.58 | 58.9% | 1.64 | 6.6 | +| M5 | $0.65 | second half | 137 | +12.09 | 54.0% | 1.23 | 8.3 | +| M5 | $0.30 | whole | 283 | +46.38 | 57.2% | 1.48 | 8.4 | +| M5 | $0.30 | first half | 146 | +32.07 | 60.3% | 1.71 | 6.4 | +| M5 | $0.30 | second half | 137 | +14.30 | 54.0% | 1.28 | 7.4 | +| M15 | $0.65 | whole | 238 | +16.35 | 60.1% | 1.19 | 14.9 | +| M15 | $0.65 | first half | 120 | +14.09 | 63.3% | 1.37 | 7.6 | +| M15 | $0.65 | second half | 118 | +2.26 | 56.8% | 1.05 | 14.4 | +| M15 | $0.30 | whole | 238 | +19.20 | 60.9% | 1.23 | 13.8 | +| M15 | $0.30 | first half | 120 | +15.37 | 65.0% | 1.41 | 7.1 | +| M15 | $0.30 | second half | 118 | +3.82 | 56.8% | 1.08 | 13.8 | + +By generator at $0.65: M5 GENERIC 163 trades +42.4R, XAU_CP 104 −2.7R, XAU_ATR 16 +2.0R (BUY 244 ++25.7R, SELL 39 +16.0R); M15 GENERIC 149 +26.3R, XAU_ATR 89 −10.0R (BUY 144 +9.7R, SELL 94 ++6.7R). The M15 second half is the same weak stretch the 2026-09-17 live review saw; no decision +follows from a 6-month window. + +### 3.2 The five reference books (kept cache, `--spread 0.65`, factory config) + +| Book | Window | Twin MD5 | Pinned (CLAUDE.md §9) | Twin wall | +|---|---|---|---|---| +| M5 W2 | 2024-09-05 → 2026-09-05 | `49f8e602fc6ae769713c94ef7b257810` | ✓ | 2.4 s (reference 114.8 s) | +| M15 W2 | 2024-09-05 → 2026-09-05 | `b4aa07e33e17e23774c486a6900d673b` | ✓ | 2.3 s (37.2 s) | +| M15 OOS | 2024-06-01 → 2025-03-30 | `9fd624fb387c36052122bd1830cf9c32` | ✓ | 1.5 s (15.6 s) | +| H1 W3 | 2023-09-05 → 2026-09-05 | `7b848c917dcb5c371ecbb9f8f7292278` | ✓ | 1.3 s (13.0 s) | +| H4 W3 | 2023-09-05 → 2026-09-05 | `186e677fbbd53cca735c94094c23ce05` | ✓ | 0.8 s (21.2 s) | + +Reference wall times are the GPU-backend rows of `gpu_backend_2026-09-18.md`. On the twin the +decision stage itself is milliseconds (M5 W2: K1 218,088 bars of three timeframes in 0.020 s, K2 +141,819 primary bars in 0.002 s); what remains is loading and cleaning the parquet files, the +indicator precompute (1.3 s on the two-year M5 window), the host formatting of the accepted trades +and the CSV write. + +### 3.3 Research flags and knobs (M15 OOS window unless stated; both scripts, `--spread 0.65`) + +| Variant | Rows | Twin MD5 (prefix) | Equal | +|---|---|---|---| +| `--export-features` (130 columns) | 355 | `6239d197` | ✓ | +| `--record DIR --record-bars` | 355 | book `9fd624fb`; `trades` (355 × 138), `indicators` (8,120 × 15), `events` (6,902 × 14), `bars` (12,315 × 172) parquet tables equal via `DataFrame.equals`; manifests differ only in `cli`, `created_at`, `replay_seconds`, `workers` | ✓ | +| `--fixed-rr-tp 2.0` | 355 | `95a7d2ab` | ✓ | +| `--quality A+,A` | 249 | `1883c346` | ✓ | +| `--strategies XAU` | 3 | `c07a4025` | ✓ | +| `--cp` (all six chart patterns on M15) | 396 | `7ff51617` | ✓ | +| `--big-candle 3.0` | 104 | `76f76f3b` | ✓ | +| `--set backtest.regime_filter=true` (H4 ADX gate, a fourth frame) | 179 | `0a65af6b` | ✓ | +| `--set technical.ema_mode=2ema` | 368 | `faaa7e86` | ✓ | +| M5, `--set xau_strategy.m5_pullback_enabled=true --quality A+,A,B,C,D` | 279 | `1811117a` | ✓ | +| `--trail-arm 0.5 --trail-dist 0.9 --cooldown-bars 3 --max-hold-bars 150` | 414 | `9665d753` | ✓ | +| H1, `--swing-sr-lookback 80 --warmup-bars 300` | 75 | `f8ef2b9f` | ✓ | + +One variant could not run on either script: H1 with `--confirm-tfs 15min,4h,1d` on the OOS +window needs a D1 file padded 1,400 days (2020-08-01) that is not cached, and Oanda returned an +empty page for that start; both scripts fail identically in `fetch_range`. Not a parity issue. + +### 3.4 Test suite, three ways (final tree) + +| Run | Result | Wall | Module 37 note | +|---|---|---|---| +| default (GPU) | `RESULTS: 552 passed, 0 failed, 552 total` | 116 s | every device test exercised | +| `CUDA_VISIBLE_DEVICES=""` | `RESULTS: 552 passed, 0 failed, 552 total` | 30 s | `CUDA unavailable here: NumPy CPU fallback (CUDARuntimeError: cudaErrorNoDevice: no CUDA-capable device is detected)` printed; device tests print a notice | +| `FB2_BACKEND=cpu` | `RESULTS: 552 passed, 0 failed, 552 total` | 31 s | `NumPy CPU fallback (forced: --backend cpu / FB2_BACKEND=cpu)` printed; the twin's refusal of a cpu-forced backend is exercised for real | + +536 tests before this change, 16 added by Module 37. + +## 4. Module 37 (what each test pins) + +1. `fb2_round2` / `round1` against Python `round(x, 2)` / `round(x, 1)` on 400k doubles including + every `.xx5` half in ±300 and signed zero; without CUDA the rule is checked in exact arithmetic. +2. Parameter packing round-trips every decision leaf of the four factories (M5 blocked hours with + 9 and without 21, M15 SELL hours 15/17 and the RSI × 0.5 table, H1 empty SELL quality, M15 + `tp_floor_rr` 1.0, the six-regime × nine-indicator table, `FB2_EMA_LO` read once with the same + fallback); every field of the four config dataclasses is either packed or classified host-only, + so a new field fails the suite until it is placed; the `#define`s are generated from the lists. +3. K1 against `GPUTechnicalEngine.analyze` on every bar of a 600-bar fixture (doji, hammer, + shooting star, engulfing, marubozu, a flat stretch, plateaus) for 5min/15min/1h/4h under their + factories, partial-EMA off, `2ema`, classic RSI, disabled indicators and a 150-bar frame: valid + flag, composite, regime, volatility, every indicator's value / signal / confidence / description + / effective weight / weighted score, key levels, pattern signals. +4. The arithmetic twins: numpy pairwise sums and `Series.mean` for n = 1…600, `nanmean` with NaNs, + EWM and rolling means restarted on tails, Wilder ATR on tails, `_ext_vs_sma`, the prefix median + with ties, CPython float floor division; and the rolling helper text is the prefix of + `gpu_kernels._ROLL_SRC`. +5. `align_cut` against the replay's `searchsorted` rule on gapped, late-starting M15/H1/H4/D1 + frames for M5 and H1 primaries (`cut == 0` included). +6. `host_flags` against `_is_xau_blocked_by_calendar` plus the session block bar by bar across + both 2026 DST shifts, an H4 primary (label hour ≠ close hour), the day filter, and each filter + switched off. +7. K2 against `analyze_from_dataframes` on every analysed bar: the M5 factory over a 21-day + synthetic fixture (4,046 bars: all six grades, GENERIC and XAU_CP setups, both directions, the + mixed `trending` regime with both score signs, 23 distinct warnings); the M5 research knobs + (all six chart patterns, big-candle 3.0, the H4 regime filter, `FB2_EMA_LO=0.9`, grades to D; + every second bar); the M15 factory (1,182 bars, 62 candidates, XAU_ATR and GENERIC, the SELL, + exhaustion-block and spread gates); the H1 factory (shorts blocked) and the H4 factory with a + D1 confirmation shorter than 30 bars on a 30-day fixture. Compared per bar: score, confidence, + signal, grade, agreement, exhaustion count, momentum fading, penalty total and tags, the full + warning and cleaned-warning lists, both reason lists, price, spread estimate, regime, + volatility, the setup (strategy, direction, entry, stop, targets, rounded RR, pattern) and the + replay's candidate fields. +8. The device chart-pattern detector against `detect_chart_pattern` with all six patterns on the + fixture and on its mirror image (bear patterns), zero ATR included; the device big-candle + detector against `_detect_big_candle` (index and ratio) with injected abnormal bars. +9. K4 against `_apply_cooldown` (cooldown 0 / 1 / 5 / 17, empty, a dense run). +10. `Trade` rows, the 65 features, the recorder snapshots and the per-bar float32 matrix against a + real `_replay_span` over the whole M5 fixture with the pullback path re-enabled (XAU_PB, XAU_CP + and GENERIC trades; `dataclasses.asdict` field by field, exact floats, tz-aware entry times, the + CSV frame, dtypes and bytes with the research columns). +11. End to end: `cudaBasedBacktesting.main` against `gpu_backtest.main` on the fixture written to a + temporary cache directory (M15, two weeks): plain book MD5 equal; `--fixed-rr-tp 2.0 + --export-features --record DIR --record-bars` MD5 equal and the four parquet tables equal. +12. Refusals: `--backend cpu`, a cpu-forced backend and a failed probe exit before the cache is + touched; each compile-time limit is refused with a message naming the knob and pointing to + `gpu_backtest.py`, the largest supported values pass; the parser is the reference script's; + `require_gpu_or_exit("gpu")` and the `Backend:` text on a working box. +13. Source guards: the script never calls `analyze_from_dataframes` or `_replay_span`, never + imports cupy or a live writer; the kernel module stays a leaf (no `engine`, + `signal_aggregator`, `gpu_backtest`, `replay_record`, live writer or ML import at module level), + pins `options=KERNEL_OPTIONS` (`--fmad=false`) and contains no `fmin(` / `fmax(`; the shared + code tables equal the recorder's; the live roots import neither new module. + +On the GPU the module takes about 85 s, almost all of it the CPU references (about 5 ms per bar +through the live signal code). Without CUDA the device tests print a notice and the module takes +about a second. + +## 5. Honest notes + +- The twin is a second implementation of the decision stage. The MD5 gate proves equality on the + windows and flags tried, not in general; a future signal change that is not mirrored in the + kernels will make the twin drift silently. That is why the instructions addendum and CLAUDE.md + section 7 make every signal change a parity event, and why the reference script stays. +- Speed comes from the batch, not from the per-bar work: the decision kernels run one thread per + bar over the whole window at once. The remaining wall time (about a second per book) is parquet + loading, cleaning, the precompute and the host formatting of the accepted trades. +- `--workers` is accepted and ignored (logged); the twin never forks. +- Nothing in `engine.py`, `signal_aggregator.py`, `chart_patterns.py`, `position_sizing.py`, + `replay_record.py`, `gpu_backtest.py` or `gpu_kernels.py` changed. `git diff --stat` touches + `test_all.py`, `CLAUDE.md` and `backtest/instructions_for_backtest.txt` and adds the two new + modules and this note. diff --git a/backtest/research/m15_5yr_cuda_2026-09-18.md b/backtest/research/m15_5yr_cuda_2026-09-18.md new file mode 100644 index 0000000..93b354e --- /dev/null +++ b/backtest/research/m15_5yr_cuda_2026-09-18.md @@ -0,0 +1,149 @@ +# M15 five-year book on the CUDA research twin — 2026-09-18 + +Owner request: a five-year M15 backtest with `backtest/cudaBasedBacktesting.py`. This window goes +past the three-year cap the owner set on 2026-09-05 (CLAUDE.md rule 1); it was run as an explicit +owner-requested exception. It is a read-out of the shipped configuration on a longer window, not +a candidate change: nothing was tuned, nothing ships from it. + +## 1. Runs + +Window 2021-08-01 → 2026-08-04 (the kept 5-year cache, padded files +`oanda_XAU_USD_{5min,15min,1h}_20210531_20260804.parquet`; 0 cache fallbacks; 122,489 M15 bars, +118,389 replayed; skipped 43,245 `SESSION_WINDOW`, 247 `WEEKEND_FILTER`). Factory config, no +overrides. + +``` +venv/bin/python backtest/cudaBasedBacktesting.py --pair XAUUSD --primary-tf 15min --start 2021-08-01 --end 2026-08-04 --spread 0.65 --export-csv cuda_m15_5yr_0.65.csv +venv/bin/python backtest/cudaBasedBacktesting.py --pair XAUUSD --primary-tf 15min --start 2021-08-01 --end 2026-08-04 --spread 0.30 --export-csv cuda_m15_5yr_0.30.csv +venv/bin/python backtest/gpu_backtest.py --pair XAUUSD --primary-tf 15min --start 2021-08-01 --end 2026-08-04 --spread 0.65 --workers 12 --export-csv ctrl_m15_5yr_0.65.csv # paired reference +``` + +| Run | Backend line (verbatim) | Wall | Book MD5 | +|---|---|---|---| +| twin $0.65 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)` | 6.6 s (K1 520,523 bars 0.046 s, K2 118,389 bars 0.001 s, trade sim 0.80 s) | `d5e432bdb3dc9923e00594125b4867dd` | +| twin $0.30 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)` | 5.3 s | `4ac5c0307123557c0740e57f5b91ff72` | +| reference $0.65 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)`, `GPU cache fallbacks: 0` | 93.7 s (replay 89.2 s, 12 workers) | `d5e432bdb3dc9923e00594125b4867dd` — identical to the twin's (`cmp` equal) | + +The paired reference run is the parity evidence the 2026-09-18 addendum in +`backtest/instructions_for_backtest.txt` requires before a number from the twin is quoted. + +## 2. Results + +1,822 trades at both costs (1,087 `TRAIL`, 666 `SL`, 69 `TIMEOUT`, no `NO_DATA`). Halves split at +2024-02-02. + +| Cost | Segment | Trades | Total R | Win rate | PF | Max DD (R) | +|---|---|---|---|---|---|---| +| $0.65 | whole 2021-08-01 → 2026-08-04 | 1822 | +157.43 | 56.8% | 1.22 | 43.6 | +| $0.65 | first half 2021-08-01 → 2024-02-01 | 704 | −9.61 | 52.4% | 0.97 | 43.6 | +| $0.65 | second half 2024-02-02 → 2026-08-04 | 1118 | +167.04 | 59.5% | 1.40 | 17.1 | +| $0.30 | whole | 1822 | +196.77 | 57.7% | 1.28 | 36.3 | +| $0.30 | first half | 704 | +9.38 | 54.0% | 1.03 | 36.3 | +| $0.30 | second half | 1118 | +187.39 | 60.1% | 1.46 | 14.9 | + +Read against rule 1 for the record: the five-year window does **not** pass it (the first half is +negative at $0.65 and the whole-window PF is 1.22). The shipped M15 configuration was validated on +2024-09-05 → 2026-09-05 (+146.1R, PF 1.43) and the 2024-06 → 2025-03 out-of-sample window +(+60.7R, PF 1.49); the extra 2021–2023 years are where the R is missing, which is the same picture +the Oanda re-pin of 2026-08-04 gave ("all timeframes earn only 2024+") and the reason for the +three-year cap. + +### Per year + +| Year | Trades | $0.65 R | WR | PF | Max DD | $0.30 R | PF | +|---|---|---|---|---|---|---|---| +| 2021 (Aug–Dec) | 97 | +1.17 | 47.4% | 1.03 | 15.5 | +3.85 | 1.10 | +| 2022 | 318 | −15.37 | 51.9% | 0.89 | 39.1 | −6.73 | 0.95 | +| 2023 | 264 | +10.91 | 55.7% | 1.10 | 15.7 | +17.86 | 1.17 | +| 2024 | 397 | +52.86 | 58.7% | 1.36 | 17.1 | +62.53 | 1.44 | +| 2025 | 480 | +77.35 | 60.0% | 1.43 | 14.8 | +85.69 | 1.48 | +| 2026 (Jan–Aug 4) | 266 | +30.51 | 58.3% | 1.30 | 13.2 | +33.57 | 1.34 | + +The 43.6R drawdown is entirely 2021-Q4 → 2022-Q3 (quarterly R: −13.5, +13.5, −7.6, −22.6); from +2024 on the worst quarter is −3.4R (2024-Q2) and the worst since is −8.3R for the five weeks of +2026-Q3 in the window. + +### Breakdowns ($0.65) + +| Grade | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| A+ | 36 | +5.48 | 63.9% | 1.41 | 3.3 | +| A | 970 | +92.30 | 56.5% | 1.24 | 21.3 | +| B | 728 | +41.16 | 56.5% | 1.14 | 31.3 | +| C | 88 | +18.49 | 59.1% | 1.55 | 4.8 | + +| Generator | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| GENERIC | 1658 | +153.28 | 57.0% | 1.24 | 43.6 | +| XAU_ATR | 164 | +4.15 | 54.3% | 1.06 | 17.0 | + +XAU_ATR produces no trade before 2024 (5 in 2024, 48 in 2025, 111 in 2026). The gold ATR setup +is gated by dollar-denominated floors (`min_atr_threshold`, the $8 stop floor, the 30 × $0.30 +target floor); at 2021–2023 gold prices and volatility they blocked the generator and the GENERIC +swing setup took every A-grade signal. The pre-2024 book therefore runs under an effectively +different setup mix. Observation only; no proposal. + +| Direction | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| BUY | 1212 | +147.91 | 58.3% | 1.32 | 48.3 | +| SELL | 610 | +9.52 | 53.8% | 1.04 | 36.2 | + +Direction by year (trades / R): BUY 2021 58 / −9.8, 2022 215 / −28.9, 2023 166 / +10.9, 2024 +273 / +56.8, 2025 344 / +95.2, 2026 156 / +23.7; SELL 2021 39 / +10.9, 2022 103 / +13.6, 2023 +98 / 0.0, 2024 124 / −3.9, 2025 136 / −17.9, 2026 110 / +6.8. Shorts carried the book in the +2021–2022 bear phase and lost in the 2024–2025 rally; the parked cycle-2 candidate (block M15 +grade-A SELLs, `cycle_2026-09-10.md`) is about the recent years only and this table does not +change its status. + +| Regime | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| ranging | 168 | +45.30 | 61.3% | 1.73 | 7.7 | +| transitioning | 400 | +51.25 | 58.2% | 1.33 | 7.0 | +| trending_up | 788 | +55.43 | 56.9% | 1.18 | 37.3 | +| trending_down | 466 | +5.44 | 53.6% | 1.03 | 26.0 | + +(The mixed `trending` regime is blocked on M15, so it has no row.) + +| Volatility | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| low | 193 | +34.75 | 60.1% | 1.48 | 11.2 | +| normal | 251 | +41.65 | 61.0% | 1.50 | 20.6 | +| high | 393 | +2.77 | 53.7% | 1.02 | 19.2 | +| extreme | 985 | +78.27 | 56.2% | 1.20 | 32.0 | + +| Entry hour UTC | Trades | Total R | WR | PF | | Entry hour UTC | Trades | Total R | WR | PF | +|---|---|---|---|---|---|---|---|---|---|---| +| 03 | 98 | +8.18 | 60.2% | 1.24 | | 14 | 381 | +24.60 | 55.4% | 1.16 | +| 04 | 38 | +4.62 | 55.3% | 1.29 | | 15 | 146 | +17.60 | 57.5% | 1.30 | +| 05 | 77 | +16.50 | 61.0% | 1.58 | | 17 | 79 | +27.77 | 68.4% | 2.20 | +| 06 | 117 | +27.74 | 66.7% | 1.85 | | 18 | 112 | −6.65 | 50.9% | 0.87 | +| 07 | 116 | −9.29 | 50.0% | 0.83 | | 19 | 80 | +10.95 | 62.5% | 1.40 | +| 08 | 158 | +1.62 | 53.2% | 1.03 | | 20 | 55 | +11.68 | 61.8% | 1.66 | +| 09 | 145 | −7.31 | 49.7% | 0.89 | | 22 | 40 | +12.48 | 67.5% | 2.01 | +| 12 | 180 | +16.94 | 54.4% | 1.22 | | | | | | | + +(Entry hour is the open of the bar after the signal bar; the blocked close hours 0, 1, 2, 10, 11, +13, 16, 21 leave the gaps.) + +| Weekday | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| Mon | 309 | +25.97 | 57.3% | 1.21 | 15.8 | +| Tue | 394 | +40.58 | 60.2% | 1.29 | 17.4 | +| Wed | 386 | −18.07 | 50.5% | 0.90 | 25.4 | +| Thu | 356 | +50.99 | 56.5% | 1.37 | 22.3 | +| Fri | 358 | +47.13 | 58.4% | 1.35 | 19.3 | +| Sun | 19 | +10.82 | 78.9% | 3.64 | 2.1 | + +Tail: the top 5% of trades (91) carry +249.8R of the +157.4R total; median R +0.113; largest win ++9.07R, largest loss −2.04R; mean hold 70.5 bars. Same fat-right-tail shape as the 3-year cycle-1 +reading (94–96% of R in the top 5% there). + +## 3. What to do with this + +Nothing. The numbers describe the shipped configuration on a window the owner does not use for +decisions; the per-hour, per-weekday and per-direction tables are in-book statistics on a +five-year mix of setup gates and must not be turned into filters (rule 1 needs a control-vs- +candidate run on the validated windows, and the 2026-08-04 age/stretch and 2026-09-06 sizing +verdicts already refute the obvious ones). If the owner wants the pre-2024 weakness explained, +the honest next step is a `--record --record-bars` run on this window with the twin (seconds) and +a scorecard on the dollar-floor gates by year, not a parameter change. diff --git a/backtest/research/m5_5yr_cuda_2026-09-18.md b/backtest/research/m5_5yr_cuda_2026-09-18.md new file mode 100644 index 0000000..676fb22 --- /dev/null +++ b/backtest/research/m5_5yr_cuda_2026-09-18.md @@ -0,0 +1,156 @@ +# M5 five-year book on the CUDA research twin — 2026-09-18 + +Owner request, the companion of `m15_5yr_cuda_2026-09-18.md`: a five-year M5 backtest with +`backtest/cudaBasedBacktesting.py`. Same owner-requested exception to the three-year cap +(CLAUDE.md rule 1), same status: a read-out of the shipped configuration on a longer window, +nothing tuned, nothing ships. + +## 1. Runs + +Window 2021-08-01 → 2026-08-04 (kept 5-year cache, padded files +`oanda_XAU_USD_{5min,15min,1h}_20210531_20260804.parquet`; 0 cache fallbacks; 367,394 M5 bars, +355,096 replayed; skipped 139,064 `SESSION_WINDOW`, 247 `WEEKEND_FILTER`). Factory config (M5: +hour 9 blocked and 21 open, chart-pattern entries on, pullbacks off), no overrides. + +``` +venv/bin/python backtest/cudaBasedBacktesting.py --pair XAUUSD --primary-tf 5min --start 2021-08-01 --end 2026-08-04 --spread 0.65 --export-csv cuda_m5_5yr_0.65.csv +venv/bin/python backtest/cudaBasedBacktesting.py --pair XAUUSD --primary-tf 5min --start 2021-08-01 --end 2026-08-04 --spread 0.30 --export-csv cuda_m5_5yr_0.30.csv +venv/bin/python backtest/gpu_backtest.py --pair XAUUSD --primary-tf 5min --start 2021-08-01 --end 2026-08-04 --spread 0.65 --workers 12 --export-csv ctrl_m5_5yr_0.65.csv # paired reference +``` + +| Run | Backend line (verbatim) | Wall | Book MD5 | +|---|---|---|---| +| twin $0.65 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)` | 5.2 s (K1 520,523 bars 0.045 s, K2 355,096 bars 0.004 s, trade sim 0.38 s) | `0429885713ca34653a39754c013db501` | +| twin $0.30 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)` | 5.6 s | `87bfdba14090459f715e1ee1f751a6cb` | +| reference $0.65 | `CuPy GPU (NVIDIA GeForce RTX 4070 SUPER, cupy 14.0.1)`, `GPU cache fallbacks: 0` | 353.2 s (replay 348.7 s, 12 workers, 1,018 bars/s) | `0429885713ca34653a39754c013db501` — identical to the twin's (`cmp` equal) | + +## 2. Results + +905 trades at both costs (561 `TRAIL`, 290 `SL`, 54 `TIMEOUT`, no `NO_DATA`). Halves split at +2024-02-02. + +| Cost | Segment | Trades | Total R | Win rate | PF | Max DD (R) | +|---|---|---|---|---|---|---| +| $0.65 | whole 2021-08-01 → 2026-08-04 | 905 | +127.08 | 51.4% | 1.39 | 13.5 | +| $0.65 | first half 2021-08-01 → 2024-02-01 | 129 | −8.98 | 43.4% | 0.82 | 12.4 | +| $0.65 | second half 2024-02-02 → 2026-08-04 | 776 | +136.06 | 52.7% | 1.49 | 9.6 | +| $0.30 | whole | 905 | +147.83 | 52.5% | 1.46 | 10.6 | +| $0.30 | first half | 129 | −5.20 | 45.0% | 0.89 | 10.6 | +| $0.30 | second half | 776 | +153.03 | 53.7% | 1.57 | 8.9 | + +Against rule 1 for the record: the five-year window fails it (first half negative), on a first +half of only 129 trades. The shipped M5 configuration was validated on 2024-09-05 → 2026-09-05 +(+146.1R, PF 1.52, 781 trades) and the 12-month window (+108.8R, PF 1.54). + +### Per year + +| Year | Trades | $0.65 R | WR | PF | Max DD | $0.30 R | PF | +|---|---|---|---|---|---|---|---| +| 2021 (Aug–Dec) | 12 | −1.99 | 41.7% | 0.59 | 2.5 | −1.60 | 0.65 | +| 2022 | 50 | −2.18 | 44.0% | 0.90 | 11.1 | −0.71 | 0.97 | +| 2023 | 61 | −2.75 | 42.6% | 0.87 | 10.3 | −1.00 | 0.95 | +| 2024 | 102 | +0.97 | 49.0% | 1.02 | 8.2 | +4.30 | 1.11 | +| 2025 | 332 | +40.54 | 50.6% | 1.34 | 9.6 | +48.63 | 1.42 | +| 2026 (Jan–Aug 4) | 348 | +92.49 | 55.7% | 1.77 | 9.2 | +98.21 | 1.83 | + +Two things the M15 window did not show as sharply: + +1. **M5 barely trades before 2025.** 225 trades in the first 3.4 years against 680 in the last + 1.6. Trade generation on M5 is governed by dollar-denominated floors: target 1 must clear + 30 × $0.30 = $9 (`SPREAD GATE`), the generic and chart-pattern stops are floored at $8, and a + chart pattern needs a $9 height. At 2021–2023 prices and volatility an M5 ATR of $1.5–3 puts a + 1.5 × ATR target far below the floor, so the spread gate vetoed almost every setup and only + high-volatility bursts survived. The pre-2025 M5 book is therefore a burst-only sample, not the + shipped strategy under normal conditions. +2. **One quarter carries the window.** 2026-Q1 made +90.2R of the +127.1R total (71%); 2025 Q1–Q2 + another +32.3R. Quarterly R outside those three quarters ranges −6.5 … +5.8. This is the + fat-right-tail shape of every M5 reading (cycle 1: the top 5% of trades = 94–96% of R; here the + top 5%, 45 trades, carry +145.8R of +127.1R) compounded by a single very favourable quarter. + +### Breakdowns ($0.65) + +| Grade | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| A+ | 1 | −0.03 | 0.0% | 0.00 | 0.0 | +| A | 247 | +53.72 | 55.1% | 1.69 | 7.2 | +| B | 637 | +73.66 | 50.2% | 1.30 | 16.4 | +| C | 20 | −0.26 | 45.0% | 0.97 | 3.2 | + +| Generator | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| GENERIC | 447 | +70.36 | 52.8% | 1.46 | 8.1 | +| XAU_CP | 427 | +48.52 | 49.9% | 1.29 | 14.4 | +| XAU_ATR | 31 | +8.20 | 51.6% | 1.89 | 3.8 | + +Generator by year (trades / R): GENERIC 2021 4 / +0.4, 2022 23 / +4.4, 2023 32 / −2.9, 2024 +37 / +6.6, 2025 170 / +18.1, 2026 181 / +43.8; XAU_CP 2021 8 / −2.4, 2022 27 / −6.5, 2023 +29 / +0.2, 2024 65 / −5.7, 2025 159 / +21.6, 2026 139 / +41.3; XAU_ATR appears only in 2025 (3) +and 2026 (28), the same dollar-floor effect as on M15. The chart-pattern path is negative through +2024 (−14.4R over 129 trades) and positive since (+62.9R over 298); its validation +(2026-08-04, +38.9R of the 12-month book) covered the recent regime only. + +| Direction | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| BUY | 808 | +119.68 | 51.4% | 1.41 | 15.2 | +| SELL | 97 | +7.40 | 51.5% | 1.21 | 11.4 | + +Shorts are nearly absent before 2025 (9 trades in 2021–2024) under the M5 SELL gates; 2025 32 / +−2.8R, 2026 56 / +12.0R. + +| Regime | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| trending_up | 481 | +75.44 | 53.4% | 1.44 | 11.9 | +| ranging | 174 | +24.78 | 50.0% | 1.36 | 8.3 | +| trending_down | 94 | +6.61 | 51.1% | 1.19 | 12.4 | +| transitioning | 84 | +19.59 | 46.4% | 1.69 | 6.9 | +| trending (mixed) | 72 | +0.67 | 47.2% | 1.02 | 10.5 | + +| Volatility | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| low | 250 | +62.92 | 51.6% | 1.70 | 12.4 | +| normal | 169 | +0.69 | 47.9% | 1.01 | 12.5 | +| extreme | 486 | +63.47 | 52.5% | 1.38 | 13.3 | + +(`high` has no row: the M5 chop-transition gate blocks it.) + +| Entry hour UTC | Trades | Total R | WR | PF | | Entry hour UTC | Trades | Total R | WR | PF | +|---|---|---|---|---|---|---|---|---|---|---| +| 03 | 52 | +8.49 | 50.0% | 1.41 | | 15 | 121 | +12.81 | 53.7% | 1.30 | +| 04 | 31 | +3.22 | 61.3% | 1.27 | | 17 | 47 | −2.61 | 48.9% | 0.86 | +| 05 | 50 | +7.15 | 52.0% | 1.32 | | 18 | 51 | +13.17 | 58.8% | 1.76 | +| 06 | 56 | +24.31 | 58.9% | 2.47 | | 19 | 51 | +2.96 | 49.0% | 1.18 | +| 07 | 52 | +2.50 | 53.8% | 1.14 | | 20 | 43 | +31.45 | 46.5% | 2.80 | +| 08 | 63 | −4.88 | 41.3% | 0.81 | | 21 | 12 | +16.02 | 75.0% | 6.16 | +| 12 | 69 | +7.01 | 53.6% | 1.28 | | 22 | 37 | +4.14 | 45.9% | 1.32 | +| 14 | 170 | +1.37 | 47.6% | 1.02 | | | | | | | + +(Entry hour is the open of the bar after the signal bar; M5 blocks close hours 0, 1, 2, 9, 10, +11, 13, 16 and leaves 21 open.) + +| Weekday | Trades | Total R | WR | PF | Max DD | +|---|---|---|---|---|---| +| Mon | 172 | +20.80 | 55.8% | 1.36 | 12.3 | +| Tue | 180 | +24.22 | 52.8% | 1.36 | 13.2 | +| Wed | 200 | +30.84 | 48.5% | 1.43 | 11.4 | +| Thu | 167 | +5.27 | 48.5% | 1.08 | 13.3 | +| Fri | 177 | +41.41 | 50.3% | 1.62 | 9.2 | +| Sun | 9 | +4.55 | 77.8% | 3.22 | 1.0 | + +Tail: median R +0.031; largest win +5.53R, largest loss −1.84R; mean hold 76.6 bars. + +## 3. What to do with this + +Nothing ships. Two honest readings for the owner: + +- The five-year M5 figure (+127R) is not "five years of edge"; it is the validated 2024-09 → + 2026-09 book plus three and a half years in which the dollar floors kept M5 almost out of the + market and the few trades that got through lost a little. The three-year cap already keeps + decisions inside the period where the gates behave as designed. +- The concentration in 2026-Q1 (71% of the window's R) is a fact about the shipped M5 edge, not + about the window length. It is consistent with every earlier M5 reading and is the reason the + advisory outcomes over a quiet fortnight can look nothing like the book. + +If the owner wants the floor effect measured rather than inferred, a `--record --record-bars` +run of this window on the twin (seconds) with a by-year count of `SPREAD GATE` and `LOW ATR +FILTER` warnings per analysed bar would settle it. Changing the floors is a rule-1 event in its +own right (`round_trip_cost` drives the live target floor, CLAUDE.md rule 1) and is not proposed. diff --git a/test_all.py b/test_all.py index f99a1d5..f326989 100644 --- a/test_all.py +++ b/test_all.py @@ -13789,6 +13789,1084 @@ def _(): assert "cupy" not in open("requirements.txt").read() +# ══════════════════════════════════════════════════════════════════════ +# MODULE 37: CUDA decision stage parity (backtest/cuda_signal_kernels.py +# + backtest/cudaBasedBacktesting.py, 2026-09-18) +# The research-only full-GPU twin must reproduce gpu_backtest.py's book +# byte for byte. Every stage is pinned here against the live Python code +# it replicates: the kernels' arithmetic twins, K1 against the engine, +# K2 against analyze_from_dataframes, the host builders against +# _replay_span, and the two scripts end to end. GPU-path tests print a +# notice and return without CUDA, like Module 36. +# ══════════════════════════════════════════════════════════════════════ + +print("\n═══ Module 37: CUDA decision stage parity ═══") + +import contextlib as _ctx37 +import dataclasses as _dc37 +import hashlib as _hashlib37 +import io as _io37 +from fractions import Fraction as _Fr37 +from zoneinfo import ZoneInfo as _ZoneInfo37 + +import chart_patterns as _cp37 +import position_sizing as _ps37 +from providers import make_data_provider as _mdp37 +from signal_aggregator import SignalAggregator as _SA37 +from pair_signals.xauusd import make_config as _mk37 +from backtest import gpu_kernels as _gk37 +from backtest import gpu_backtest as _gb37 +from backtest import cuda_signal_kernels as _ck37 +from backtest import cudaBasedBacktesting as _cb37 +from backtest import replay_record as _rr37 + +_GPU37 = _gk37.BACKEND.available +if not _GPU37: + print(f" (CUDA unavailable here: {_gk37.BACKEND.describe()} — the device " + f"parity tests below print a notice instead of running)") + +_NY37 = _ZoneInfo37("America/New_York") +_M37 = {} # fixtures shared by the tests below (built once) + + +def _same37(a, b) -> bool: + """Ordered, NaN-aware equality for the plain-data trees the builders emit.""" + if isinstance(a, float) and isinstance(b, float): + return a == b or (np.isnan(a) and np.isnan(b)) + if isinstance(a, dict) and isinstance(b, dict): + return list(a) == list(b) and all(_same37(a[k], b[k]) for k in a) + if isinstance(a, (list, tuple)) and isinstance(b, (list, tuple)): + return len(a) == len(b) and all(_same37(x, y) for x, y in zip(a, b)) + return a == b + + +def _m37_m5_fixture(days=21, seed=11, start="2026-03-02"): + """Seeded M5 bars over `days` calendar days with the New-York weekend + removed and four alternating regimes (trend up, chop, trend down, + high-volatility chop), so the aggregate comparison sees every grade, + both directions, the mixed `trending` regime and every setup generator.""" + rng = np.random.default_rng(seed) + idx = pd.date_range(start, periods=days * 288, freq="5min", tz="UTC") + ny = idx.tz_convert(_NY37) + wd, hr = np.asarray(ny.weekday), np.asarray(ny.hour) + idx = idx[~(((wd == 4) & (hr >= 17)) | (wd == 5) | ((wd == 6) & (hr < 18)))] + n = len(idx) + drift = np.zeros(n) + vol = np.full(n, 1.2) + for seg, a in enumerate(range(0, n, 700)): + kind = seg % 4 + drift[a:a + 700] = {0: 0.25, 1: 0.0, 2: -0.25, 3: 0.05}[kind] + vol[a:a + 700] = {0: 1.5, 1: 0.8, 2: 1.8, 3: 3.0}[kind] + close = 3300 + np.cumsum(drift + rng.normal(0, 1, n) * vol) + op = np.concatenate([[close[0]], close[:-1]]) + rng.normal(0, 0.2, n) + hi = np.maximum(op, close) + np.abs(rng.normal(0, 1, n)) * vol + lo = np.minimum(op, close) - np.abs(rng.normal(0, 1, n)) * vol + df = pd.DataFrame({"Open": op, "High": hi, "Low": lo, "Close": close, + "Volume": 100.0}, index=idx) + df.index.name = "Datetime" + return df + + +def _m37_resample(df, rule): + agg = df.resample(rule, label="left", closed="left").agg( + {"Open": "first", "High": "max", "Low": "min", "Close": "last", "Volume": "sum"}) + return agg.dropna(subset=["Open"]) + + +def _m37_frames(key="main"): + """M5/M15/H1/H4/D1 frames. 'main' is the 21-day fixture (M5 and M15 + primaries); 'slow' a 30-day fixture whose D1 frame has fewer than 30 + bars (H1 and H4 primaries; the empty-analysis path on a confirmation).""" + if key not in _M37: + m5 = (_m37_m5_fixture() if key == "main" + else _m37_m5_fixture(days=30, seed=12, start="2026-04-06")) + _M37[key] = { + Timeframe.M5: m5, Timeframe.M15: _m37_resample(m5, "15min"), + Timeframe.H1: _m37_resample(m5, "1h"), Timeframe.H4: _m37_resample(m5, "4h"), + Timeframe.D1: _m37_resample(m5, "1D"), + } + return _M37[key] + + +def _m37_bt(cfg, dfs, **kw): + m5 = dfs[Timeframe.M5] + kw.setdefault("warmup_bars", 250) + kw.setdefault("quality_filter", ["A+", "A", "B", "C"]) + return _gb37.BacktestConfig( + pair=XAUUSD, start=m5.index[0].to_pydatetime(), + end=(m5.index[-1] + pd.Timedelta(minutes=5)).to_pydatetime(), + primary_tf=cfg.primary_timeframe, + confirmation_tfs=list(cfg.confirmation_timeframes), **kw) + + +def _m37_timeframes(cfg, bt): + tfs = [bt.primary_tf] + list(bt.confirmation_tfs) + flt = cfg.backtest + if flt.regime_filter and flt.regime_filter_timeframe not in tfs: + tfs.append(flt.regime_filter_timeframe) + return tfs + + +def _m37_aggregator(cfg, dfs, tfs): + """The reference: a real SignalAggregator over the cached-indicator engine.""" + caches = {tf.value: _gb37.precompute_indicators(dfs[tf], cfg.technical) for tf in tfs} + agg = _SA37(cfg) + agg.tech_engine = _gb37.GPUTechnicalEngine(cfg.technical, caches) + return agg + + +def _m37_reference(cfg, bt, dfs, positions): + """CPU truth per primary bar: ('skip', reason) from _replay_span's + calendar/session rules, or ('rep', {...}) from analyze_from_dataframes + on the real slices, plus the candidate fields the replay derives.""" + tfs = _m37_timeframes(cfg, bt) + agg = _m37_aggregator(cfg, dfs, tfs) + primary_df = dfs[bt.primary_tf] + primary_dur = pd.Timedelta(seconds=_gb37.BAR_SECONDS[bt.primary_tf]) + flt = cfg.backtest + out = {} + for pos in positions: + pos = int(pos) + decision_time = primary_df.index[pos] + primary_dur + block = _gb37._is_xau_blocked_by_calendar(decision_time, cfg) + if block: + out[pos] = ("skip", block) + continue + if flt.session_filter: + h = decision_time.hour + if (not (flt.session_start_hour <= h < flt.session_end_hour) + or h in cfg.xau_strategy.blocked_hours_utc): + out[pos] = ("skip", "SESSION_WINDOW") + continue + slice_dfs = {} + for tf in tfs: + full = dfs[tf] + cutoff = decision_time - pd.Timedelta(seconds=_gb37.BAR_SECONDS[tf]) + cut = int(full.index.searchsorted(cutoff, side="right")) + if cut: + slice_dfs[tf] = full.iloc[:cut] + report = agg.analyze_from_dataframes(symbol=bt.pair, dfs=slice_dfs, + context=None, use_news=False) + if flt.regime_filter: + agg._apply_regime_filter_live(report, slice_dfs) + s = report.trade_setup + cand = None + if s is not None and report.signal_quality in bt.quality_filter: + tag = str(getattr(s, "strategy", "") or "GENERIC") + if ((not bt.strategies_only or tag in bt.strategies_only) + and pos + 1 < len(primary_df)): + ep = float(primary_df["Open"].iloc[pos + 1]) + sld = abs(s.entry_price - s.stop_loss) + t1 = abs(s.take_profit_1 - s.entry_price) + t2 = abs(s.take_profit_2 - s.entry_price) + cand = ((ep, ep - sld, ep + t1, ep + t2, sld, t1, t2) if s.direction == "BUY" + else (ep, ep + sld, ep - t1, ep - t2, sld, t1, t2)) + out[pos] = ("rep", dict( + final_score=float(report.final_score), final_conf=float(report.final_confidence), + final_signal=report.final_signal.value, quality=report.signal_quality, + mtf=bool(report.mtf_agreement), exh=int(report.exhaustion_count), + fading=bool(report.momentum_fading), ptotal=float(report.penalty_total), + ptags=dict(report.penalty_tags), warnings=list(report.warnings), + cleaned=list(report.cleaned_warnings), bull=list(report.bullish_reasons), + bear=list(report.bearish_reasons), price=float(report.current_price), + spread=float(report.spread_estimate), regime=report.market_regime, + vol=report.volatility_regime, + setup=None if s is None else ( + s.strategy, s.direction, float(s.entry_price), float(s.stop_loss), + float(s.take_profit_1), float(s.take_profit_2), float(s.risk_reward), + s.setup_pattern), + cand=cand, + # non-vacuity bookkeeping only (never compared) + _short_confirm=any(len(slice_dfs[tf]) < 30 for tf in bt.confirmation_tfs + if tf in slice_dfs))) + return out + + +def _m37_device(cfg, bt, dfs): + """The same per-bar view built from the device outputs through the + script's own host formatters. Returns (rows, run context).""" + tfs = _m37_timeframes(cfg, bt) + ctx = _cb37.analyze_frames(bt, cfg, dfs, tfs) + positions = np.arange(ctx.first_pos, ctx.last_pos, dtype=np.int64) + ints, floats, wc, wa = _cb37._k2_rows(ctx.k2, positions) + prows = _cb37.k1_rows(ctx.an[bt.primary_tf], positions) + pdev = ctx.dev[bt.primary_tf] + med = _ck37.prefix_median(pdev["high"], pdev["low"], pdev["n"]).get() + two = cfg.technical.ema_mode == "2ema" + primary_df = dfs[bt.primary_tf] + out = {} + for m, pos in enumerate(positions): + pos = int(pos) + r2i, r2f, prow = ints[m], floats[m], prows[m] + if r2i["status"] == 0: + out[pos] = ("skip", _ck37.SKIP_NAMES[r2i["skip_reason"]]) + continue + r2f["_atr"] = prow.f["atr"] if prow.i["atr_present"] else 0.0 + inds = _cb37.describe_indicators(prow, two) + bull, bear = _cb37.build_reasons(ctx, r2i, r2f, inds) + warnings, cleaned = _cb37.build_warnings(ctx, pos, r2i, r2f, wc[m], wa[m]) + setup = None + if r2i["strategy"]: + tag = _ck37.STRATEGY_NAMES[r2i["strategy"]] + setup = (tag, "BUY" if r2i["direction"] > 0 else "SELL", + r2f["setup_entry"], r2f["setup_sl"], r2f["setup_tp1"], r2f["setup_tp2"], + r2f["rr_rounded"], + _ck37.PATTERN_ORDER[r2i["cp_pattern"]] if tag == "XAU_CP" else None) + cand = None + if r2i["status"] == 2: + cand = (r2f["entry_price"], r2f["adj_sl"], r2f["adj_tp1"], r2f["adj_tp2"], + r2f["sl_dist"], r2f["tp1_dist"], r2f["tp2_dist"]) + pmask = r2i["penalty_mask"] + ptags = {t: float(_ck37.PENALTY_AMOUNT[t]) for t in _ck37.PENALTY_SORTED + if pmask & (1 << _ck37.PENALTY_TAGS.index(t))} + out[pos] = ("rep", dict( + final_score=r2f["final_score"], final_conf=r2f["final_conf"], + final_signal=r2i["final_signal"], quality=_ck37.QUALITY_NAMES[r2i["quality"]], + mtf=bool(r2i["mtf_agreement"]), exh=r2i["exhaustion_count"], + fading=bool(r2i["momentum_fading"]), ptotal=r2f["penalty_total"], ptags=ptags, + warnings=warnings, cleaned=cleaned, bull=bull, bear=bear, + price=float(primary_df["Close"].iloc[pos]), spread=float(med[pos]), + regime=_ck37.REGIME_NAMES[prow.i["regime"]], vol=_ck37.VOL_NAMES[prow.i["vol"]], + setup=setup, cand=cand, + _warn_names=[_ck37.WARN_NAMES[int(c)] for c in wc[m][:r2i["warn_n"]]])) + return out, ctx + + +def _m37_compare(label, ref, got): + """Zero tolerance: every compared field of every bar. Returns the + non-vacuity tallies so the caller can assert the fixture earned its keep.""" + mismatches, first = 0, None + tally = {"status": set(), "quality": set(), "strategy": set(), "direction": set(), + "warn": set(), "regime_signs": set(), "cands": 0, "short_confirm": 0} + for pos, r in ref.items(): + g = got.get(pos) + if g is None or r[0] != g[0] or (r[0] == "skip" and r[1] != g[1]): + mismatches += 1 + first = first or (pos, r if r[0] == "skip" else r[0], g if g is None or g[0] == "skip" else g[0]) + continue + if r[0] == "skip": + tally["status"].add(r[1]) + continue + rd, gd = r[1], g[1] + diffs = [f"{k}: ref={rd[k]!r} dev={gd[k]!r}" for k in rd + if not k.startswith("_") and rd[k] != gd[k]] + if diffs: + mismatches += 1 + first = first or (pos, diffs[:3]) + tally["status"].add("candidate" if rd["cand"] else "analysed") + tally["quality"].add(rd["quality"]) + if rd["setup"]: + tally["strategy"].add(rd["setup"][0]) + tally["direction"].add(rd["setup"][1]) + tally["warn"].update(gd["_warn_names"]) + if rd["regime"] == "trending" and rd["final_score"] != 0: + tally["regime_signs"].add("+" if rd["final_score"] > 0 else "-") + tally["cands"] += bool(rd["cand"]) + tally["short_confirm"] += bool(rd["_short_confirm"]) + assert mismatches == 0, f"{label}: {mismatches} of {len(ref)} bars differ; first: {first}" + return tally + + +@test("CUDA stage: fb2_round2/round1 == Python round(x, 2/1) on 400k doubles incl. exact halves (rule checked exactly without CUDA)") +def _(): + rng = np.random.default_rng(371) + xs = np.concatenate([ + rng.uniform(1000, 5000, 120_000), rng.normal(0, 3, 100_000), + np.arange(-60_000, 60_000) / 200.0, # k/200: the .xx5 halves + np.arange(-10_000, 10_000) / 20.0, # k/20: the .x5 halves for round(x, 1) + rng.uniform(-1e9, 1e9, 20_000), rng.uniform(-1e-3, 1e-3, 20_000), + np.array([0.0, -0.0, 0.125, 0.135, 2.675, 1.005, 3299.995, 3300.005, -2.5, 2.5]), + ]) + + def rule(x: float, scale: float) -> float: + # the kernel's rule in exact arithmetic: rint(x*scale), then fix a + # tie by the sign of the exact residual of the product + p = x * scale + e = _Fr37(x) * int(scale) - _Fr37(p) + r = float(round(p)) + d = p - r + if abs(d) == 0.5 and e != 0 and ((e > 0) == (d > 0)): + r += 1.0 if d > 0 else -1.0 + return r / scale + + ref2 = np.array([round(float(x), 2) for x in xs]) + ref1 = np.array([round(float(x), 1) for x in xs]) + sample = xs[::7] + assert all(rule(float(x), 100.0) == round(float(x), 2) for x in sample) + assert all(rule(float(x), 10.0) == round(float(x), 1) for x in sample) + assert rule(2.675, 100.0) == 2.67 and rule(0.125, 100.0) == 0.12 and rule(0.135, 100.0) == 0.14 + if not _GPU37: + print(" (no CUDA: rule checked in exact arithmetic only)") + return + got2 = _ck37.t_round(xs, 100.0) + got1 = _ck37.t_round(xs, 10.0) + bad2 = np.flatnonzero(got2 != ref2) + bad1 = np.flatnonzero(got1 != ref1) + assert len(bad2) == 0, f"round2: {len(bad2)} differ, e.g. x={xs[bad2[0]]!r} dev={got2[bad2[0]]!r} py={ref2[bad2[0]]!r}" + assert len(bad1) == 0, f"round1: {len(bad1)} differ, e.g. x={xs[bad1[0]]!r} dev={got1[bad1[0]]!r} py={ref1[bad1[0]]!r}" + # signed zero survives like Python's + neg_zero = np.flatnonzero((xs == 0.0) & np.signbit(xs)) + assert len(neg_zero) == 1 and np.signbit(got2[neg_zero[0]]) and np.signbit(round(-0.0, 2)) + + +@test("CUDA stage: param packing round-trips the four factories' decision leaves; every config field is classified") +def _(): + import dataclasses as _dc + packed = set(_ck37.PARAMS_F) | set(_ck37.PARAMS_I) + assert len(packed) == len(_ck37.PARAMS_F) + len(_ck37.PARAMS_I), "duplicate parameter name" + for cls in (TechnicalConfig, SignalConfig, XauStrategyConfig, FilterConfig): + names = {f.name for f in _dc.fields(cls)} + host = _ck37.CONFIG_FIELDS_HOST_ONLY[cls.__name__] + assert not (names - packed - host), ( + f"{cls.__name__}: new field(s) {sorted(names - packed - host)} are read by neither the kernels nor " + f"the host — pack them or list them in CONFIG_FIELDS_HOST_ONLY with a reason") + assert not (host - names), f"{cls.__name__}: stale host-only entries {sorted(host - names)}" + assert not (host & packed & names), f"{cls.__name__}: {sorted(host & packed & names)} both packed and host-only" + dfs = _m37_frames("main") + + def bits(mask, n=24): + return sorted(h for h in range(n) if (mask >> h) & 1) + + for tf in (Timeframe.M5, Timeframe.M15, Timeframe.H1, Timeframe.H4): + cfg = _mk37(tf) + bt = _m37_bt(cfg, dfs, quality_filter=["A+", "A", "B"], strategies_only=["XAU_CP", "GENERIC"]) + pf, pi, rm = _ck37.pack_params(cfg, bt) + assert pf.dtype == np.float64 and pi.dtype == np.int32 and rm.dtype == np.float64 + assert len(pf) == len(_ck37.PARAMS_F) and len(pi) == len(_ck37.PARAMS_I) and rm.shape == (54,) + f = dict(zip(_ck37.PARAMS_F, pf.tolist())) + i = dict(zip(_ck37.PARAMS_I, pi.tolist())) + x, s, t = cfg.xau_strategy, cfg.signals, cfg.technical + assert bits(i["xau_blocked_hours_mask"]) == sorted(x.blocked_hours_utc) + assert bits(i["sell_blocked_hours_mask"]) == sorted(x.sell_blocked_hours_utc or []) + assert bits(i["sell_min_quality_mask"], 6) == sorted(_ck37.QUALITY_CODES[q] for q in x.sell_min_quality) + assert bits(i["quality_filter_mask"], 6) == [3, 4, 5] # B, A, A+ + assert bits(i["strategies_mask"], 5) == [1, 4] # GENERIC, XAU_CP + assert i["primary_is_m5"] == (tf == Timeframe.M5) and i["n_confirm"] == 2 + assert i["primary_bar_secs"] == _gb37.BAR_SECONDS[tf] + assert i["ema_partial_neutral"] == int(t.ema_partial_neutral) + assert i["cp_enabled"] == int(x.cp_enabled) and i["m5_pullback_enabled"] == int(x.m5_pullback_enabled) + assert bits(i["blocked_regimes_mask"], 6) == sorted(_ck37.REGIME_CODES[r] for r in x.blocked_regimes) + assert bits(i["sell_blocked_regimes_mask"], 6) == sorted(_ck37.REGIME_CODES[r] for r in x.sell_blocked_regimes) + assert bits(i["cp_patterns_mask"], 6) == sorted(_ck37.PATTERN_ORDER.index(p) for p in x.cp_patterns) + assert f["round_trip_cost"] == s.round_trip_cost and f["min_tp_cost_multiple"] == s.min_tp_cost_multiple + assert f["tp_floor_rr"] == s.tp_floor_rr and f["pip_value"] == PAIR_INFO[XAUUSD]["pip_value"] + assert f["fb2_ema_lo"] == 0.5 and f["max_penalty"] == 30.0 + assert i["divergence_lookback"] == t.divergence_lookback and i["swing_sr_lookback"] == t.swing_sr_lookback + # the regime table: rows in REGIME_CODES order, columns in INDICATOR_ORDER + from engine import regime_weight_table as _rwt + table = _rwt(t) + grid = rm.reshape(6, 9) + for regime, code in _ck37.REGIME_CODES.items(): + for col, name in enumerate(_ck37.INDICATOR_ORDER): + want = table.get(regime, {}).get(name, 1.0) + assert grid[code, col] == want, (tf, regime, name, grid[code, col], want) + # the documented per-timeframe overrides really arrive in the arrays + m5 = dict(zip(_ck37.PARAMS_I, _ck37.pack_params(_mk37(Timeframe.M5), _m37_bt(_mk37(Timeframe.M5), dfs))[1].tolist())) + m15 = _ck37.pack_params(_mk37(Timeframe.M15), _m37_bt(_mk37(Timeframe.M15), dfs)) + h1 = dict(zip(_ck37.PARAMS_I, _ck37.pack_params(_mk37(Timeframe.H1), _m37_bt(_mk37(Timeframe.H1), dfs))[1].tolist())) + assert bits(m5["xau_blocked_hours_mask"]) == [0, 1, 2, 9, 10, 11, 13, 16] + assert bits(m5["sell_blocked_regimes_mask"], 6) == [_ck37.REGIME_CODES["transitioning"]] + assert m5["ema_partial_neutral"] == 1 and m5["cp_enabled"] == 1 and m5["m5_pullback_enabled"] == 0 + m15i = dict(zip(_ck37.PARAMS_I, m15[1].tolist())) + assert bits(m15i["sell_blocked_hours_mask"]) == [15, 17] + assert bits(m15i["blocked_regimes_mask"], 6) == [_ck37.REGIME_CODES["trending"]] + grid15 = m15[2].reshape(6, 9) + grid5 = _ck37.pack_params(_mk37(Timeframe.M5), _m37_bt(_mk37(Timeframe.M5), dfs))[2].reshape(6, 9) + rsi = _ck37.INDICATOR_ORDER.index("rsi") + for code in range(1, 6): # every real regime: RSI halved, others equal + assert grid15[code, rsi] == grid5[code, rsi] * 0.5, code + other = [c for c in range(9) if c != rsi] + assert np.array_equal(grid15[code, other], grid5[code, other]), code + assert dict(zip(_ck37.PARAMS_F, m15[0].tolist()))["tp_floor_rr"] == 1.0 + assert h1["sell_min_quality_mask"] == 0, "H1 blocks every short" + # FB2_EMA_LO: read once on the host, bad values fall back like the CPU path + with patch.dict(os.environ, {"FB2_EMA_LO": "0.9"}): + assert _ck37.fb2_ema_lo_from_env() == 0.9 + assert dict(zip(_ck37.PARAMS_F, _ck37.pack_params(_mk37(Timeframe.M5), _m37_bt(_mk37(Timeframe.M5), dfs))[0].tolist()))["fb2_ema_lo"] == 0.9 + with patch.dict(os.environ, {"FB2_EMA_LO": "wide"}): + assert _ck37.fb2_ema_lo_from_env() == 0.5 + # the kernel #defines are generated from the same name lists + src = _ck37._defines() + for k, name in enumerate(_ck37.PARAMS_F): + assert f"#define PF_{name} {k}\n" in src, name + for k, name in enumerate(_ck37.PARAMS_I): + assert f"#define PI_{name} {k}\n" in src, name + assert f"#define FB2_MAX_CONFIRM {_ck37.MAX_CONFIRM}" in src + + +def _k1_fixture37(n=600, seed=36): + rng = np.random.default_rng(seed) + idx = pd.date_range("2025-01-06", periods=n, freq="5min", tz="UTC") + close = 3300 + np.cumsum(rng.normal(0, 2.0, n)) + op = close + rng.normal(0, 0.5, n) + hi = np.maximum(op, close) + np.abs(rng.normal(0, 1.5, n)) + lo = np.minimum(op, close) - np.abs(rng.normal(0, 1.5, n)) + # candlestick specials, a flat stretch (zero range, plateau pivots), equal highs / lows + op[100] = close[100]; hi[100] = close[100] + 3; lo[100] = close[100] - 3 # doji + op[110] = close[110] - 0.5; lo[110] = op[110] - 6; hi[110] = close[110] + 0.1 # hammer + close[120] = op[120] - 0.5; hi[120] = op[120] + 6; lo[120] = close[120] - 0.1 # shooting star + op[130] = close[129] - 1; close[130] = op[129] + 1; hi[130] = close[130] + .2; lo[130] = op[130] - .2 + op[140] = lo[140] + 0.01; close[140] = hi[140] - 0.01 # marubozu + if n > 262: + hi[250:262] = lo[250:262] = op[250:262] = close[250:262] = close[249] + if n > 326: + hi[300:306] = hi[300]; lo[320:326] = lo[320] + return pd.DataFrame({"Open": op, "High": hi, "Low": lo, "Close": close, "Volume": 1.0}, index=idx) + + +@test("CUDA stage: K1 == GPUTechnicalEngine.analyze on every bar (4 factories, partial-EMA off, 2ema, classic RSI, disabled indicators, short frame)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + df = _k1_fixture37() + cases = [] + for tf_name, primary in (("5min", Timeframe.M5), ("15min", Timeframe.M15), + ("1h", Timeframe.H1), ("4h", Timeframe.H4)): + cases.append((f"{tf_name} under the {primary.value} factory", tf_name, _mk37(primary), df)) + c = _mk37(Timeframe.M5); c.technical.ema_partial_neutral = False + cases.append(("5min partial_neutral=False", "5min", c, df)) + c = _mk37(Timeframe.H1); c.technical.ema_mode = "2ema" + cases.append(("1h ema_mode=2ema", "1h", c, df)) + c = _mk37(Timeframe.H1); c.technical.rsi_mode = "classic" + cases.append(("1h rsi classic", "1h", c, df)) + c = _mk37(Timeframe.H1); c.technical.disabled_indicators = ["macd", "divergence"] + cases.append(("1h macd+divergence disabled", "1h", c, df)) + cases.append(("15min short frame (no EMA200)", "15min", _mk37(Timeframe.M15), _k1_fixture37(150, seed=5))) + seen_patterns, seen_regimes, seen_div = set(), set(), 0 + for label, tf_name, cfg, frame in cases: + tech = cfg.technical + cache = _gb37.precompute_indicators(frame, tech) + eng = _gb37.GPUTechnicalEngine(tech, {tf_name: cache}) + pf, pi, rm = _ck37.pack_params(cfg, _m37_bt(cfg, _m37_frames("main"))) + an = _ck37.analyze_tf(_ck37.upload_cache(cache, frame), pf, pi, rm, tf_name == "5min") + rows = _cb37.k1_rows(an, np.arange(len(frame))) + two = tech.ema_mode == "2ema" + for pos in range(len(frame)): + ref = eng.analyze(frame.iloc[:pos + 1], XAUUSD, tf_name) + row = rows[pos] + diffs = [] + got = {"valid": bool(row.i["valid"]), "regime": _ck37.REGIME_NAMES[row.i["regime"]], + "vol": _ck37.VOL_NAMES[row.i["vol"]], "score": row.f["composite_score"], + "conf": row.f["composite_conf"], "sig": row.i["composite_signal"]} + exp = {"valid": pos + 1 >= 30, "regime": ref.market_regime, "vol": ref.volatility_regime, + "score": float(ref.composite_score), "conf": float(ref.confidence), + "sig": ref.composite_signal.value} + diffs += [f"{k}: dev={got[k]!r} ref={exp[k]!r}" for k in exp if got[k] != exp[k]] + inds = _cb37.describe_indicators(row, two) + ref_inds = {k: v for k, v in ref.indicators.items() if v.weight != 0 and not k.startswith("_")} + if list(inds) != list(ref_inds): + diffs.append(f"indicator set {list(inds)} vs {list(ref_inds)}") + else: + for k in ref_inds: + r, g = ref_inds[k], inds[k] + for fld, gv, rv in (("value", float(g["value"]), float(r.value)), + ("signal", g["signal"], r.signal.value), + ("confidence", float(g["confidence"]), float(r.confidence)), + ("description", g["description"], r.description), + ("effective_weight", g["effective_weight"], float(r.effective_weight)), + ("weighted_score", g["weighted_score"], float(r.weighted_score))): + if gv != rv: + diffs.append(f"{k}.{fld}: dev={gv!r} ref={rv!r}") + kl, rk = _cb37.key_levels(row, two), dict(ref.key_levels) + if list(kl) != list(rk) or any(kl[k] != rk[k] for k in rk): + diffs.append(f"key_levels dev={kl} ref={rk}") + if _cb37.pattern_signals(row) != list(ref.pattern_signals): + diffs.append("pattern_signals") + assert not diffs, f"{label}: bar {pos}: {diffs[:4]}" + seen_patterns.update(ref.pattern_signals) + seen_regimes.add(ref.market_regime) + seen_div += "divergence" in ref_inds + assert len(seen_patterns) >= 5, seen_patterns + assert {"ranging", "trending_up", "trending_down"} <= seen_regimes, seen_regimes + assert seen_div > 0, "no divergence vote in the fixture" + + +@test("CUDA stage: summation / rolling / EWM / Wilder / median / floordiv twins are bit-identical to numpy & pandas") +def _(): + # the rolling helpers are the SAME text the Module-36 kernels compile + assert _gk37._ROLL_SRC.startswith(_ck37._ROLL_HELPERS) and len(_ck37._ROLL_HELPERS) > 500 + assert 'extern "C" __global__' not in _ck37._ROLL_HELPERS + assert "options=KERNEL_OPTIONS" in Path("backtest/cuda_signal_kernels.py").read_text(encoding="utf-8") + if not _GPU37: + print(" (no CUDA: kernels not exercised)") + return + rng = np.random.default_rng(7) + # numpy pairwise sum (np.add.reduce) and pandas Series.mean for n = 1..600 + series = [rng.normal(3300, 5, n) for n in range(1, 601)] + [rng.uniform(0, 50, n) for n in range(1, 601)] + dev = _ck37.t_pairwise(series) + assert np.array_equal(dev, np.array([np.add.reduce(s) for s in series])) + assert np.array_equal(np.array([d / len(s) for d, s in zip(dev, series)]), + np.array([pd.Series(s).mean() for s in series])) + # np.nanmean on 20-windows with NaNs (MACD hist_pct) + wins = [] + for t in range(2000): + w = np.abs(rng.normal(0, 3, 20)) + if t % 4 == 1: + w[rng.integers(0, 20, 3)] = np.nan + if t % 7 == 3: + w[:] = np.nan + wins.append(w) + dev = _ck37.t_nanmean(wins) + import warnings as _w + with _w.catch_warnings(): + _w.simplefilter("ignore") + ref = np.array([np.nanmean(w) for w in wins]) + assert np.array_equal(dev, ref, equal_nan=True) + # EWM restarted on tails (alpha 1/14, min_periods 14) incl. flat runs + tails = [] + for t in range(400): + n = int(rng.integers(15, 500)); x = 3300 + np.cumsum(rng.normal(0, 2, n)) + if t % 3 == 0 and n > 80: + x[40:80] = x[39] + tails.append(x) + for d, s in zip(_ck37.t_ewm(tails, 1 / 14, 14), tails): + assert np.array_equal(d, pd.Series(s).ewm(alpha=1 / 14, min_periods=14, adjust=False).mean().to_numpy(), + equal_nan=True) + # rolling(20).mean() on 60-bar tails (features' _ext_vs_sma) incl. flat windows + tails = [3300 + np.cumsum(rng.normal(0, 2, int(rng.integers(21, 61)))) for _ in range(1500)] + for t in range(0, 1500, 5): + tails[t][5:25] = tails[t][4] + assert np.array_equal(_ck37.t_rollmean_last(tails, 20), + np.array([pd.Series(x).rolling(20).mean().iloc[-1] for x in tails])) + # Wilder ATR restarted on tails and position_sizing._ext_vs_sma on prefixes + n = 3000 + close = 3300 + np.cumsum(rng.normal(0, 2, n)); op = close + rng.normal(0, .5, n) + hi = np.maximum(op, close) + np.abs(rng.normal(0, 1.5, n)); lo = np.minimum(op, close) - np.abs(rng.normal(0, 1.5, n)) + hi[500:520] = lo[500:520] = close[500:520] = close[499] + df = pd.DataFrame({"Open": op, "High": hi, "Low": lo, "Close": close}) + bases, lens = [], [] + for t in range(300): + L = int(rng.integers(21, 500)); bases.append(int(rng.integers(0, n - L))); lens.append(L) + for d, b, L in zip(_ck37.t_wilder(hi, lo, close, bases, lens, 14), bases, lens): + assert np.array_equal(d, _ps37.wilder_atr(df.iloc[b:b + L], 14).to_numpy(), equal_nan=True) + lens2 = [int(x) for x in rng.integers(1, n, 300)] + [20, 21, 22, 59, 60, 61] + for d_ in (1.0, -1.0): + assert np.array_equal(_ck37.t_ext_vs_sma(hi, lo, close, lens2, d_), + np.array([_ps37._ext_vs_sma(df.iloc[:L], d_) for L in lens2])) + # K5: pandas Series.median of High-Low over every prefix (ties included) + m5 = _m37_frames("main")[Timeframe.M5].iloc[:1200].copy() + m5.iloc[100:140, m5.columns.get_loc("High")] = m5["Low"].iloc[100:140].to_numpy() + 2.0 # 40 equal ranges + med = _ck37.prefix_median(_ck37._dev(m5["High"].to_numpy(), np.float64), + _ck37._dev(m5["Low"].to_numpy(), np.float64), len(m5)).get() + rng_ = (m5["High"] - m5["Low"]) + assert np.array_equal(med, np.array([rng_.iloc[:k + 1].median() for k in range(len(m5))])) + # CPython float floor division (sign fix-ups, exact halves, negative zero) + a = np.concatenate([rng.uniform(-1e6, 1e6, 20000), rng.integers(-500, 500, 5000).astype(float), [0.0, -0.0, 7.0, -7.0]]) + b = np.concatenate([rng.uniform(-1e3, 1e3, 20000), rng.integers(1, 9, 5000).astype(float) * np.where(rng.random(5000) > .5, 1, -1), [3.0, 3.0, 2.0, 2.0]]) + b[b == 0] = 1.0 + got = _ck37.t_floordiv(a, b) + ref = np.array([float(x) // float(y) for x, y in zip(a, b)]) + assert np.array_equal(got, ref), int(np.sum(got != ref)) + assert np.array_equal(np.signbit(got), np.signbit(ref)) + + +@test("CUDA stage: align_cut == _replay_span's slice length on gapped M15/H1/H4/D1 frames (late start, cut == 0, weekend)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + dfs = _m37_frames("main") + for primary_tf in (Timeframe.M5, Timeframe.H1): + primary = dfs[primary_tf].index + p_secs = _gb37.BAR_SECONDS[primary_tf] + p_dev = _ck37._dev(primary.as_unit("ns").asi8, np.int64) + for tf in (Timeframe.M15, Timeframe.H1, Timeframe.H4, Timeframe.D1): + idx = dfs[tf].index[5:] # starts after the primary: cut == 0 at first + idx = idx.delete(list(range(40, min(47, len(idx) - 1)))) # an interior gap + tf_secs = _gb37.BAR_SECONDS[tf] + # the replay's per-bar rule, vectorised + cutoff = primary + pd.Timedelta(seconds=p_secs) - pd.Timedelta(seconds=tf_secs) + ref = np.asarray(idx.searchsorted(cutoff, side="right"), dtype=np.int64) + got = _ck37.align_cut(_ck37._dev(idx.as_unit("ns").asi8, np.int64), p_dev, p_secs, tf_secs).get() + assert np.array_equal(got, ref), (primary_tf, tf, int(np.sum(got != ref))) + # and the scalar call _replay_span really makes, on a sample + for pos in range(0, len(primary), max(1, len(primary) // 60)): + decision_time = primary[pos] + pd.Timedelta(seconds=p_secs) + cut = int(idx.searchsorted(decision_time - pd.Timedelta(seconds=tf_secs), side="right")) + assert got[pos] == cut, (primary_tf, tf, pos, got[pos], cut) + if tf_secs > p_secs: + assert got[0] == 0 and got[-1] == len(idx), (tf, got[0], got[-1]) + assert (np.diff(got) >= 0).all() + + +@test("CUDA stage: host_flags == the replay's calendar/session filters bar by bar (both 2026 DST shifts, day filter, filters off)") +def _(): + idx = pd.DatetimeIndex([]).tz_localize("UTC") if False else None + parts = [pd.date_range("2026-03-06", "2026-03-10", freq="5min", tz="UTC"), # spring forward (Mar 8) + pd.date_range("2026-10-30", "2026-11-03", freq="5min", tz="UTC"), # fall back (Nov 1) + pd.date_range("2026-06-01", "2026-06-08", freq="5min", tz="UTC")] # a plain week + idx = parts[0].append(parts[1]).append(parts[2]) + + def scalar(ts, cfg, bt): + decision = ts + pd.Timedelta(seconds=_gb37.BAR_SECONDS[bt.primary_tf]) + block = _gb37._is_xau_blocked_by_calendar(decision, cfg) + if block: + return block + flt = cfg.backtest + if flt.session_filter: + h = decision.hour + if not (flt.session_start_hour <= h < flt.session_end_hour): + return "SESSION_WINDOW" + if h in cfg.xau_strategy.blocked_hours_utc: + return "SESSION_WINDOW" + return None + + def check(cfg, primary_tf, index): + bt = _gb37.BacktestConfig(pair=XAUUSD, start=index[0].to_pydatetime(), end=index[-1].to_pydatetime(), + primary_tf=primary_tf, confirmation_tfs=list(cfg.confirmation_timeframes)) + flags = _cb37.host_flags(index, cfg, bt) + got = [_ck37.SKIP_NAMES.get(int(c)) for c in flags["skip_reason"]] + ref = [scalar(ts, cfg, bt) for ts in index] + bad = [k for k in range(len(index)) if got[k] != ref[k]] + assert not bad, f"{primary_tf.value}: {len(bad)} bars differ, first {index[bad[0]]}: dev={got[bad[0]]} ref={ref[bad[0]]}" + dur = pd.Timedelta(seconds=_gb37.BAR_SECONDS[primary_tf]) + close = index + dur + assert np.array_equal(flags["label_hour"], np.asarray(index.hour)) and np.array_equal(flags["label_wd"], np.asarray(index.weekday)) + assert np.array_equal(flags["close_hour"], np.asarray(close.hour)) and np.array_equal(flags["close_wd"], np.asarray(close.weekday)) + assert np.array_equal(flags["close_minute"], np.asarray(close.minute)) and np.array_equal(flags["close_day"], np.asarray(close.day)) + assert all(f.dtype == np.int8 for f in flags.values()) + return set(ref) + + seen = set() + seen |= check(_mk37(Timeframe.M5), Timeframe.M5, idx) + seen |= check(_mk37(Timeframe.M15), Timeframe.M15, idx[::3]) + h4_idx = pd.date_range("2026-03-01", "2026-03-31", freq="4h", tz="UTC") + seen |= check(_mk37(Timeframe.H4), Timeframe.H4, h4_idx) # label hour != close hour + assert seen == {None, "WEEKEND_FILTER", "SESSION_WINDOW"}, seen + c = _mk37(Timeframe.M5); c.xau_strategy.allowed_weekdays = [0, 1, 2, 3, 4] # no Sunday evening + assert "DAY_FILTER" in check(c, Timeframe.M5, idx) + c = _mk37(Timeframe.M5); c.xau_strategy.weekend_filter = False + assert "WEEKEND_FILTER" not in check(c, Timeframe.M5, idx) + c = _mk37(Timeframe.M5); c.backtest.session_filter = False + assert check(c, Timeframe.M5, idx) == {None, "WEEKEND_FILTER"} + # the flag builder's weekend rule is the same rule Module 21 sweeps + ny_close = pd.Timestamp("2026-03-06 21:59", tz="UTC") # Fri 16:59 NY (EST): open + ny_after = pd.Timestamp("2026-03-06 22:00", tz="UTC") # Fri 17:00 NY: closed + bt = _gb37.BacktestConfig(pair=XAUUSD, start=idx[0].to_pydatetime(), end=idx[-1].to_pydatetime(), primary_tf=Timeframe.M5) + f = _cb37.host_flags(pd.DatetimeIndex([ny_close - pd.Timedelta(minutes=5), ny_after - pd.Timedelta(minutes=5)]), _mk37(Timeframe.M5), bt) + assert list(f["skip_reason"]) == [0, 1] + + +def _k2_case37(label, primary, tweak, bt_kw, frames, stride, ema_lo, min_bars=100): + dfs = _m37_frames(frames) + with patch.dict(os.environ, {"FB2_EMA_LO": ema_lo} if ema_lo else {}, clear=False): + if not ema_lo: + os.environ.pop("FB2_EMA_LO", None) + cfg = _mk37(primary) + if tweak: + tweak(cfg) + bt = _m37_bt(cfg, dfs, **bt_kw) + got, ctx = _m37_device(cfg, bt, dfs) + positions = np.arange(ctx.first_pos, ctx.last_pos)[::stride] + ref = _m37_reference(cfg, bt, dfs, positions) + assert len(ref) >= min_bars, (label, len(ref)) + tally = _m37_compare(label, ref, got) + tally["n"] = len(ref) + tally["n_analysed"] = int((ctx.status >= 1).sum()) + tally["skipped"] = dict(ctx.skipped_reasons) + return tally + + +@test("CUDA stage: K2 == analyze_from_dataframes on every bar — M5 factory (all grades, GENERIC + XAU_CP, both directions)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + t = _k2_case37("M5 factory", Timeframe.M5, None, {}, "main", 1, None) + assert {"A+", "A", "B", "C", "D", "F"} <= t["quality"], t["quality"] + assert {"GENERIC", "XAU_CP"} <= t["strategy"], t["strategy"] + assert t["status"] == {"WEEKEND_FILTER", "SESSION_WINDOW", "analysed", "candidate"}, t["status"] + assert t["cands"] >= 20, t["cands"] + assert t["regime_signs"] == {"+", "-"}, t["regime_signs"] + # ("m5_score" is unreachable by construction since the 2026-07-30 relaxation: + # grade C already needs |score| >= 25 = m5_min_abs_score.) + must = {"session_filter", "conflict", "m5_mtf", "counter_trend", "exhaustion_low", + "exhaustion_high", "no_pullback", "weak_pullback", "rsi_ext_up", "rsi_ext_down", + "ema200_near", "low_vol", "extreme_vol", "spread_gate_generic"} + assert must <= t["warn"], (sorted(must - t["warn"]), sorted(t["warn"])) + print(f" ({t['n']} bars: {len(t['warn'])} distinct warnings, setups {sorted(t['strategy'])}, " + f"directions {sorted(t['direction'])}, {t['cands']} candidates)") + + +def _tweak_research37(cfg): + cfg.xau_strategy.cp_patterns = list(_ck37.PATTERN_ORDER) # all six detectors + cfg.xau_strategy.big_candle_atr_multiple = 3.0 # the shipped-off gate, on + cfg.backtest.regime_filter = True # H4 ADX filter (a 4th frame) + + +@test("CUDA stage: K2 == analyze_from_dataframes — M5 research knobs (six CP patterns, big-candle gate, H4 regime filter, FB2_EMA_LO=0.9, grades to D)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + t = _k2_case37("M5 research", Timeframe.M5, _tweak_research37, + {"quality_filter": ["A+", "A", "B", "C", "D"]}, "main", 2, "0.9") + assert {"big_candle", "rf_blocked"} & t["warn"], sorted(t["warn"]) + assert "XAU_CP" in t["strategy"], t["strategy"] + print(f" ({t['n']} bars: {len(t['warn'])} distinct warnings incl. " + f"{sorted(t['warn'] & {'big_candle', 'rf_blocked', 'rf_unavailable', 'rf_no_adx'})})") + + +@test("CUDA stage: K2 == analyze_from_dataframes — M15 factory (XAU_ATR + GENERIC, SELL gates, exhaustion block, spread gates)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + t = _k2_case37("M15 factory", Timeframe.M15, None, {}, "main", 1, None) + assert {"XAU_ATR", "GENERIC"} <= t["strategy"], t["strategy"] + assert {"A+", "A", "B", "C", "D", "F"} <= t["quality"], t["quality"] + must = {"sell_quality", "sell_vol", "spread_gate_atr", "spread_gate_generic", + "exhausted_trend", "exhaustion_block"} + assert must <= t["warn"], (sorted(must - t["warn"]), sorted(t["warn"])) + assert t["cands"] >= 40, t["cands"] + print(f" ({t['n']} bars, {t['cands']} candidates, directions {sorted(t['direction'])})") + + +@test("CUDA stage: K2 == analyze_from_dataframes — H1 and H4 factories (H1 shorts blocked; D1 confirmation shorter than 30 bars)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + t1 = _k2_case37("H1 factory", Timeframe.H1, None, {}, "slow", 1, None) + assert t1["cands"] > 0 and t1["direction"] == {"BUY"}, (t1["cands"], t1["direction"]) + assert "sell_quality" in t1["warn"], sorted(t1["warn"]) + t4 = _k2_case37("H4 factory", Timeframe.H4, None, {"warmup_bars": 60}, "slow", 1, None, min_bars=50) + assert t4["short_confirm"] > 0, "the D1 frame must be shorter than 30 bars somewhere" + assert t4["n_analysed"] > 20, t4["n_analysed"] + print(f" (H1 {t1['n']} bars / {t1['cands']} candidates; H4 {t4['n']} bars, " + f"{t4['short_confirm']} with an empty D1 analysis)") + + +@test("CUDA stage: device chart-pattern detector == detect_chart_pattern (six patterns, mirrored series, zero ATR); big candle == _detect_big_candle") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + dfs = _m37_frames("main") + m5 = dfs[Timeframe.M5] + cfg = _mk37(Timeframe.M5) + x = cfg.xau_strategy + x.cp_patterns = list(_ck37.PATTERN_ORDER) + x.big_candle_atr_multiple = 3.0 + pf, pi, _rm = _ck37.pack_params(cfg, _m37_bt(cfg, dfs)) + atr_full = _gb37.precompute_indicators(m5, cfg.technical).atr + mirrored = pd.DataFrame({"Open": 6600.0 - m5["Open"], "High": 6600.0 - m5["Low"], + "Low": 6600.0 - m5["High"], "Close": 6600.0 - m5["Close"], + "Volume": m5["Volume"]}, index=m5.index) + seen = {} + for label, df in (("fixture", m5), ("mirrored", mirrored)): + h, l, c = (df[k].to_numpy(dtype=np.float64) for k in ("High", "Low", "Close")) + pos_list = np.arange(x.cp_lookback + 1, len(df)) + pos_list = pos_list[np.isfinite(atr_full[pos_list]) & (atr_full[pos_list] > 0)] + atrs = atr_full[pos_list].copy() + atrs[:5] = 0.0 # atr <= 0 -> no pattern + pattern, vals = _ck37.t_chart_pattern(h, l, c, pos_list, atrs, pf, pi) + for k, pos in enumerate(pos_list): + ref = _cp37.detect_chart_pattern(df.iloc[:pos + 1], float(atrs[k]), x) + if ref is None: + assert pattern[k] == -1, (label, int(pos), _ck37.PATTERN_ORDER[pattern[k]]) + continue + assert pattern[k] >= 0, (label, int(pos), ref["pattern"]) + name = _ck37.PATTERN_ORDER[pattern[k]] + got = (name, "BUY" if vals[k, 0] > 0 else "SELL", vals[k, 1], vals[k, 2], vals[k, 3]) + want = (ref["pattern"], ref["direction"], ref["anchor"], ref["trigger"], ref["height"]) + assert got == want, (label, int(pos), got, want) + seen[name] = seen.get(name, 0) + 1 + assert {"inv_hs", "double_bottom", "bull_flag", "hs_top", "double_top", "bear_flag"} <= set(seen), seen + # big candle: inject two abnormal bars; the device returns the bar index and ratio + df2 = m5.copy() + col_h = df2.columns.get_loc("High") + df2.iloc[1000, col_h] += 40.0 + df2.iloc[2500, col_h] += 60.0 + agg = _SA37(cfg) + h, l, c = (df2[k].to_numpy(dtype=np.float64) for k in ("High", "Low", "Close")) + pos_list = np.concatenate([np.arange(0, 60), np.arange(990, 1300, 3), np.arange(2490, 2700, 5), [len(df2) - 1]]) + idx, ratio = _ck37.t_big_candle(h, l, c, pos_list, pf, pi) + hits = misses = 0 + for k, pos in enumerate(pos_list): + ref = agg._detect_big_candle(df2.iloc[:pos + 1]) + if ref is None: + assert idx[k] == -1, (int(pos), int(idx[k]), ratio[k]) + misses += 1 + else: + assert idx[k] >= 0 and df2.index[idx[k]] == ref[0] and ratio[k] == ref[1], (int(pos), idx[k], ratio[k], ref) + hits += 1 + assert hits > 5 and misses > 5, (hits, misses) + + +@test("CUDA stage: K4 cooldown == _apply_cooldown (cooldown 0/1/5/17, empty)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + rng = np.random.default_rng(374) + for cd in (0, 1, 5, 17): + for trial in range(4): + positions = np.unique(rng.integers(0, 3000, int(rng.integers(1, 400)))).astype(np.int64) + ref = [p for p, _ in _gb37._apply_cooldown([(int(p), None) for p in positions], cd)] + mask = _ck37.cooldown(positions, cd) + assert mask.dtype == np.int8 and len(mask) == len(positions) + assert positions[mask.astype(bool)].tolist() == ref, (cd, trial) + assert len(_ck37.cooldown(np.zeros(0, dtype=np.int64), 5)) == 0 + # dense run: exactly every 5th survives + dense = np.arange(100, 160, dtype=np.int64) + assert dense[_ck37.cooldown(dense, 5).astype(bool)].tolist() == list(range(100, 160, 5)) + + +@test("CUDA stage: Trade rows, decision features, snapshots and bar rows == _replay_span (M5, pullback on, --export-features --record-bars)") +def _(): + if not _GPU37: + print(" (no CUDA: kernel not exercised)") + return + dfs = _m37_frames("main") + cfg = _mk37(Timeframe.M5) + cfg.xau_strategy.m5_pullback_enabled = True + bt = _m37_bt(cfg, dfs, quality_filter=["A+", "A", "B", "C", "D"], record_dir="unused", + record_bars=True, export_features=True) + tfs = _m37_timeframes(cfg, bt) + ctx = _cb37.analyze_frames(bt, cfg, dfs, tfs) + agg = _m37_aggregator(cfg, dfs, tfs) + try: + _gb37._REPLAY_CTX.update(primary_df=dfs[Timeframe.M5], dfs_full=dfs, timeframes=tfs, + aggregator=agg, platform_cfg=cfg, bt_cfg=bt, + primary_dur=pd.Timedelta(minutes=5)) + skipped, cands, bars, fallbacks = _gb37._replay_span((ctx.first_pos, ctx.last_pos)) + finally: + _gb37._REPLAY_CTX.clear() + assert fallbacks == 0 + assert skipped == ctx.skipped_reasons, (skipped, ctx.skipped_reasons) + assert [p for p, _ in cands] == ctx.cand_pos.tolist(), "the candidate set (before cooldown) differs" + assert len(cands) >= 30, len(cands) + trades = _cb37.build_trades(ctx, ctx.cand_pos) + _cb37.attach_features(ctx, ctx.cand_pos, trades) + _cb37.attach_snapshots(ctx, ctx.cand_pos, trades) + for (pos, t_ref), t_dev in zip(cands, trades): + a, b = _dc37.asdict(t_ref), _dc37.asdict(t_dev) + fa, fb = a.pop("features"), b.pop("features") + sa, sb = a.pop("snapshot"), b.pop("snapshot") + diff = [k for k in a if not _same37(a[k], b[k])] + assert not diff, (pos, [(k, a[k], b[k]) for k in diff][:3]) + assert list(fa) == list(fb) == list(_ps37.FEATURE_ORDER) + fd = [(k, fa[k], fb[k]) for k in fa if fa[k] != fb[k]] + assert not fd, (pos, fd[:3]) + assert all(isinstance(v, float) and np.isfinite(v) for v in fb.values()) + assert _same37(sa, sb), (pos, [k for k in sa if not _same37(sa.get(k), sb.get(k))]) + assert isinstance(t_dev.entry_time, pd.Timestamp) and t_dev.entry_time.tz is not None + tags = {t.strategy_tag for t in trades} + assert {"XAU_PB", "XAU_CP", "GENERIC"} <= tags, tags + assert any(t.trail_arm_r > 0 for t in trades if t.strategy_tag == "XAU_PB") + assert any(len(t.snapshot["tfs"]) == 3 for t in trades) + # the per-analysed-bar float32 matrix (--record-bars) + mat_r, ts_r, pos_r, w0_r = bars + (mat_d, ts_d, pos_d, w0_d), = _cb37.build_bar_chunks(ctx) + assert mat_d.dtype == mat_r.dtype == np.float32 and mat_d.shape == mat_r.shape == (len(pos_r), len(_rr37.BAR_COLS)) + bad_cols = [c for k, c in enumerate(_rr37.BAR_COLS) if not np.array_equal(mat_d[:, k], mat_r[:, k], equal_nan=True)] + assert not bad_cols, bad_cols[:5] + assert np.array_equal(ts_d, ts_r) and np.array_equal(pos_d, pos_r) and list(w0_d) == list(w0_r) + assert sum(1 for w in w0_d if w) > 100 and sum(1 for w in w0_d if not w) > 0 + # the CSV builder sees identical rows: same frame, same dtypes, same bytes (63 + research columns) + df_r = _gb37.trades_to_dataframe([t for _, t in cands], include_features=True) + df_d = _gb37.trades_to_dataframe(trades, include_features=True) + pd.testing.assert_frame_equal(df_r, df_d, check_exact=True) + assert df_r.to_csv(index=False) == df_d.to_csv(index=False) + assert list(df_d.dtypes.astype(str)) == list(df_r.dtypes.astype(str)) + assert len(_gb37.trades_to_dataframe(trades).columns) == 63 + # no candidates: the builders return empty, never raise + assert _cb37.build_trades(ctx, np.zeros(0, dtype=np.int64)) == [] + _cb37.attach_features(ctx, np.zeros(0, dtype=np.int64), []) + + +@test("CUDA stage: end-to-end book == gpu_backtest.main on a temp cache (M15 plain; --fixed-rr-tp --export-features --record --record-bars)") +def _(): + if not _GPU37: + print(" (no CUDA: script not exercised)") + return + dfs = _m37_frames("main") + cfg = _mk37(Timeframe.M15) + start = datetime(2026, 3, 9, tzinfo=timezone.utc) + end = datetime(2026, 3, 23, tzinfo=timezone.utc) + provider_name = _mdp37(cfg).name + with tempfile.TemporaryDirectory() as d: + cache = Path(d) / "cache" + cache.mkdir() + with patch.object(_gb37, "CACHE_DIR", cache): + for tf in [Timeframe.M15] + list(cfg.confirmation_timeframes): + tf_start = start - timedelta(days=_gb37._history_pad_days(tf, 250)) + dfs[tf].to_parquet(_gb37._cache_path(XAUUSD, tf, tf_start, end, provider_name)) + cal = Path(d) / "cal.parquet" + pd.DataFrame([_m33_cal_row(1, pd.Timestamp("2026-03-06 13:30", tz="UTC"), + "Non-Farm Employment Change")]).to_parquet(cal, index=False) + base = ["--pair", "XAUUSD", "--primary-tf", "15min", "--start", "2026-03-09", + "--end", "2026-03-23", "--spread", "0.65", "--workers", "1"] + + def run(mod, out, extra): + with patch.object(_gb37, "CACHE_DIR", cache), _ctx37.redirect_stdout(_io37.StringIO()): + rc = mod.main(base + extra + ["--export-csv", str(out)]) + assert rc == 0 + return _hashlib37.md5(out.read_bytes()).hexdigest() + + md5_ref = run(_gb37, Path(d) / "ctrl.csv", []) + md5_dev = run(_cb37, Path(d) / "cuda.csv", []) + book = pd.read_csv(Path(d) / "cuda.csv") + assert len(book) >= 20 and len(book.columns) == 63, book.shape + assert md5_ref == md5_dev, f"plain book differs: {md5_ref} vs {md5_dev}" + saved = _rr37.CALENDAR_PARQUET + try: + _rr37.CALENDAR_PARQUET = cal + flags = ["--fixed-rr-tp", "2.0", "--export-features", "--record-bars"] + md5_ref_f = run(_gb37, Path(d) / "ctrl_f.csv", flags + ["--record", str(Path(d) / "rec_ctrl")]) + md5_dev_f = run(_cb37, Path(d) / "cuda_f.csv", flags + ["--record", str(Path(d) / "rec_cuda")]) + finally: + _rr37.CALENDAR_PARQUET = saved + assert md5_ref_f == md5_dev_f, f"research-flag book differs: {md5_ref_f} vs {md5_dev_f}" + assert md5_ref_f != md5_ref, "the feature columns / fixed target must change the file" + feat = pd.read_csv(Path(d) / "cuda_f.csv") + assert len(feat.columns) == 63 + 2 + len(_ps37.FEATURE_ORDER), len(feat.columns) + for name in ("trades", "indicators", "events", "bars"): + a = pd.read_parquet(Path(d) / "rec_cuda" / f"{name}.parquet") + b = pd.read_parquet(Path(d) / "rec_ctrl" / f"{name}.parquet") + assert a.shape == b.shape and a.equals(b), (name, a.shape, b.shape) + ma = json.load(open(Path(d) / "rec_cuda" / "manifest.json")) + mb = json.load(open(Path(d) / "rec_ctrl" / "manifest.json")) + assert ma["book_csv_md5"] == mb["book_csv_md5"] + assert ma["backend"].startswith("CuPy GPU (") and ma["workers"] == 0 and mb["workers"] == 1 + assert ma["cache_fallbacks"] == mb["cache_fallbacks"] == 0 + differing = sorted(k for k in set(ma) | set(mb) if ma.get(k) != mb.get(k)) + assert set(differing) <= {"cli", "created_at", "replay_seconds", "workers"}, differing + + +@test("CUDA stage: refusals — --backend cpu / a cpu-forced backend / unsupported lookbacks exit before the cache is touched; the Backend line names the device") +def _(): + with patch.object(_cb37, "fetch_history", side_effect=AssertionError("the cache was touched")): + try: + with _ctx37.redirect_stdout(_io37.StringIO()): + _cb37.main(["--pair", "XAUUSD", "--backend", "cpu", "--primary-tf", "15min"]) + raise AssertionError("--backend cpu must be refused") + except SystemExit as e: + assert "no CPU path" in str(e) and "gpu_backtest.py" in str(e), e + with _gk37.BACKEND.cpu_only(): + try: + _cb37.require_gpu_or_exit(None) + raise AssertionError("a cpu-forced backend must be refused") + except SystemExit as e: + assert "no CPU path" in str(e), e + dead = _gk37.Backend() + dead.mode = "auto" # not the suite's FB2_BACKEND: the probe-failure branch + dead._probed, dead._ok, dead.reason = True, False, "synthetic probe failure" + with patch.object(_cb37, "BACKEND", dead): + try: + _cb37.require_gpu_or_exit(None) + raise AssertionError("an unavailable backend must be refused") + except SystemExit as e: + assert "synthetic probe failure" in str(e) and "gpu_backtest.py" in str(e), e + dfs = _m37_frames("main") + bt = _m37_bt(_mk37(Timeframe.M15), dfs) + _cb37.refuse_unsupported(_mk37(Timeframe.M15), bt) # every factory config is supported + for tf in (Timeframe.M5, Timeframe.H1, Timeframe.H4): + _cb37.refuse_unsupported(_mk37(tf), bt) + cases = ( + (lambda c: setattr(c.technical, "divergence_lookback", _ck37.MAX_DIV_LOOKBACK), "divergence_lookback"), + (lambda c: setattr(c.technical, "swing_sr_lookback", _ck37.MAX_SR_LOOKBACK), "swing_sr_lookback"), + (lambda c: setattr(c.xau_strategy, "cp_lookback", _ck37.MAX_CP_LOOKBACK), "cp_lookback"), + (lambda c: setattr(c.technical, "ema_mode", "3ema"), "ema_mode"), + (lambda c: setattr(c, "confirmation_timeframes", + [Timeframe.M5, Timeframe.H1, Timeframe.H4, Timeframe.D1, Timeframe.M30]), "confirmation timeframes"), + (lambda c: (setattr(c.xau_strategy, "big_candle_atr_multiple", 3.0), + setattr(c.xau_strategy, "big_candle_lookback_hours", 48.0)), "big-candle window"), + ) + for mutate, needle in cases: + c = _mk37(Timeframe.M5) + mutate(c) + try: + _cb37.refuse_unsupported(c, bt) + raise AssertionError(f"{needle} beyond the compile-time limit must be refused") + except SystemExit as e: + assert needle in str(e) and "gpu_backtest.py" in str(e), (needle, e) + # the largest supported values pass + c = _mk37(Timeframe.M5) + c.technical.divergence_lookback = _ck37.MAX_DIV_LOOKBACK - 1 + c.technical.swing_sr_lookback = _ck37.MAX_SR_LOOKBACK - 1 + c.xau_strategy.cp_lookback = _ck37.MAX_CP_LOOKBACK - 1 + _cb37.refuse_unsupported(c, bt) + # the parser is the reference script's own + assert _cb37._parse_args is _gb37._parse_args + args = _cb37._parse_args(["--pair", "XAUUSD", "--backend", "gpu", "--record", "/tmp/x", "--record-bars"]) + assert args.backend == "gpu" and args.record_bars is True + if _GPU37: + _cb37.require_gpu_or_exit(None) + try: + _cb37.require_gpu_or_exit("gpu") + assert _gk37.BACKEND.mode == "gpu" + finally: + _gk37.BACKEND.set_mode("auto") + assert _gk37.BACKEND.describe().startswith("CuPy GPU ("), _gk37.BACKEND.describe() + assert _cb37.BACKEND is _gk37.BACKEND + + +@test("CUDA stage: source guards — no replay re-entry in the script, kernel module stays a leaf, --fmad=false, no fmin/fmax, live roots untouched") +def _(): + script = Path("backtest/cudaBasedBacktesting.py").read_text(encoding="utf-8") + for banned in ("analyze_from_dataframes(", "_replay_span(", "import cupy", "_REPLAY_CTX", + "import signal_log", "NewsEngine(", "news_archive", "import violent_watch"): + assert banned not in script, f"cudaBasedBacktesting.py must not contain {banned!r}" + for required in ("from backtest import gpu_kernels", "from backtest import cuda_signal_kernels", + "_clean_accepted_cp_warnings", "_drop_stale_spread_gate_vetoes", + "GPU cache fallbacks: 0", "def require_gpu_or_exit", "def refuse_unsupported"): + assert required in script, f"missing: {required}" + kpath = Path("backtest/cuda_signal_kernels.py") + ksrc = kpath.read_text(encoding="utf-8") + banned = {"signal_log", "news_engine", "providers", "telegram_bot", "news_archive", "agent_panel", + "violent_watch", "gpu_backtest", "engine", "signal_aggregator", "replay_record", + "cudaBasedBacktesting", "sklearn", "scipy", "lightgbm", "torch", "cupy"} + seen = set() + for node in ast.walk(ast.parse(ksrc, str(kpath))): + if isinstance(node, ast.Import): + seen.update(a.name.split(".")[0] for a in node.names) + elif isinstance(node, ast.ImportFrom): + # local imports inside pack_params / regime_table are the only exceptions + # (engine.regime_weight_table, providers.base.CANDLE_SECONDS): module level only + if node.col_offset == 0: + seen.add((node.module or "").split(".")[0]) + assert not (seen & banned), seen & banned + assert {"numpy", "config", "chart_patterns", "position_sizing", "backtest"} <= seen, seen + assert "options=KERNEL_OPTIONS" in ksrc and "KERNEL_OPTIONS = _gk.KERNEL_OPTIONS" in ksrc + assert _ck37.KERNEL_OPTIONS == ("--fmad=false",) + import re as _re + assert not _re.search(r"\bfmin\s*\(|\bfmax\s*\(", ksrc), "Python min/max must be the ternary twins, never fmin/fmax" + assert "#define PYMIN(a, b) (((b) < (a)) ? (b) : (a))" in ksrc + assert "#define PYMAX(a, b) (((b) > (a)) ? (b) : (a))" in ksrc + # every K1 / K2 field and warning name used by the host exists exactly once in the lists + for lst in (_ck37.K1_INT_FIELDS, _ck37.K1_FLOAT_FIELDS, _ck37.K2_INT_FIELDS, _ck37.K2_FLOAT_FIELDS, _ck37.WARN_NAMES): + assert len(set(lst)) == len(lst) + for name in _ck37.WARN_NAMES[1:]: + assert f'"{name}"' in script or f"'{name}'" in script or name in ("big_candle",), name + # the codes the recorder shares (replay_record adds an unused '' quality and the 'XAU' alias) + assert _ck37.REGIME_CODES == _rr37.REGIME_CODES and _ck37.VOL_CODES == _rr37.VOL_CODES + for ours, theirs in ((_ck37.QUALITY_CODES, _rr37.QUALITY_CODES), (_ck37.STRATEGY_CODES, _rr37.STRATEGY_CODES)): + assert all(theirs[k] == v for k, v in ours.items()), (ours, theirs) + assert set(theirs) - set(ours) <= {"", "XAU"}, set(theirs) - set(ours) + assert _ck37.N_BAR_COLS == len(_rr37.BAR_COLS) and _ck37.N_FEATURES == len(_ps37.FEATURE_ORDER) + assert _ck37.PATTERN_ORDER == tuple(_cp37.PATTERN_DETECTORS) + # live roots stay free of the research modules and of cupy (Modules 32 / 36 sweep cupy too) + offenders = [] + for root in (Path("."), Path("providers"), Path("pair_signals")): + for path in sorted(root.glob("*.py")): + if "venv" in path.parts: + continue + for node in ast.walk(ast.parse(path.read_text(encoding="utf-8"), str(path))): + names = ([a.name for a in node.names] if isinstance(node, ast.Import) + else [node.module or ""] if isinstance(node, ast.ImportFrom) else []) + offenders += [f"{path}:{node.lineno}" for nm in names + if nm.split(".")[-1] in ("cuda_signal_kernels", "cudaBasedBacktesting") or nm.split(".")[0] == "cupy"] + assert not offenders, offenders + assert "cupy" not in open("requirements.txt").read() + + +if _GPU37: + _gk37.BACKEND.release_device_memory() + + # ══════════════════════════════════════════════════════════════════════ # SUMMARY # ══════════════════════════════════════════════════════════════════════