Harwinder Sekhon
02589cbb88
CuPy from pip had never compiled a kernel on the dev box (no CUDA_PATH for the splayed CUDA-12 wheel layout); the probe swallowed the error and every book to date was built on the NumPy path under a "GPU backtester" name. - backtest/gpu_kernels.py (new): lazily probed, never-silent BACKEND state (CUDA_PATH pointed at the nvidia-cuda-runtime-cu12 header wheel; the probe compiles a kernel and keeps the failure reason), --backend auto|gpu|cpu and FB2_BACKEND, and three CUDA kernels compiled with --fmad=false: ewm_batch (the _gpu_ewm_alpha recurrence, one thread per series), roll_mean_var (pandas 3.0.2 roll_mean/roll_var replica, Kahan + Welford + unstable recompute), resolve_trades (trailing and static exits, one thread per trade). - gpu_backtest.py: precompute_indicators and simulate_trades dispatch to the GPU twins or to the untouched NumPy paths; caches are host NumPy and device memory is released before the replay workers fork; the Backend: line names the device or the fallback reason. - test_all.py Module 36 (7 tests): a visible CUDA device must never fall back silently; kernels vs their CPU references on synthetic edge cases; backend selection; import guards (cupy banned from the live roots). Suite 536. - CLAUDE.md, instructions_for_backtest.txt, research note with the evidence. Verification: all five reference MD5s reproduced on the GPU and with --backend cpu (ten runs, GPU cache fallbacks 0); precompute GPU == CPU on all 42 cached windows (2.34 M bars); suite green with GPU, with CUDA_VISIBLE_DEVICES="" and with FB2_BACKEND=cpu; --workers 1 and --record pairs equal. The per-bar replay stays the live Python signal code on CPU workers by design; wall time is roughly unchanged (precompute ~3x faster).
700 lines
30 KiB
Python
700 lines
30 KiB
Python
"""
|
|
backtest/gpu_kernels.py — the CUDA side of the backtester.
|
|
|
|
Owns two things:
|
|
|
|
1. ``BACKEND`` — the one CuPy/CUDA state object the backtester consults.
|
|
It is probed lazily (never at import: test_all.py, analyze_signal_log.py
|
|
and replay_record import the backtester without wanting a CUDA context)
|
|
and it never fails silently: when CUDA is unusable the reason is kept in
|
|
``BACKEND.reason`` and printed in the ``Backend:`` log line. Before
|
|
2026-09-18 the probe swallowed "Failed to auto-detect CUDA root
|
|
directory" and every book was built on the NumPy path while the file
|
|
was called a GPU backtester.
|
|
|
|
CuPy from pip (cupy-cuda12x) needs the CUDA headers to JIT kernels. They
|
|
come from the ``nvidia-cuda-runtime-cu12`` wheel, whose splayed layout
|
|
has no single root, so CuPy cannot find them on its own: this module
|
|
points ``CUDA_PATH`` at that wheel when the variable is unset.
|
|
|
|
2. Three CUDA kernels, each a bit-for-bit replica of the CPU code it
|
|
replaces (proven on every cached window, see
|
|
backtest/research/gpu_backend_2026-09-18.md):
|
|
|
|
- ``ewm_batch`` — pandas ``ewm(alpha, min_periods, adjust=False)``
|
|
as implemented by ``gpu_backtest._gpu_ewm_alpha``;
|
|
one CUDA thread per series (the recurrence is
|
|
order-dependent, so a parallel scan would change
|
|
the last bits — that is not acceptable here).
|
|
- ``roll_mean_var`` — pandas 3.0.2 ``rolling(window).mean()`` and
|
|
``.var(ddof)`` (Kahan-compensated add/remove and
|
|
Welford with the numerically-unstable recompute),
|
|
one thread per series.
|
|
- ``resolve_trades``— ``gpu_backtest._simulate_trailing`` and the static
|
|
SL/TP1 branch of ``simulate_trades``, one thread per
|
|
trade.
|
|
|
|
Every kernel is compiled with ``--fmad=false`` so ``a*b + c`` rounds
|
|
twice exactly like the NumPy/Python code, and every kernel object lives
|
|
in a module global (a forked worker must never drop the last reference
|
|
and unload a CUDA module it did not create). The backtester's forked
|
|
replay workers never touch the device: precompute results are moved to
|
|
host NumPy arrays before the fork.
|
|
|
|
Import as ``from backtest import gpu_kernels`` (package form — test_all.py
|
|
puts backtest/ on sys.path and imports the backtester under two names; the
|
|
backend state must live in one module object).
|
|
"""
|
|
|
|
from __future__ import annotations
|
|
|
|
import contextlib
|
|
import importlib.util
|
|
import logging
|
|
import os
|
|
import sys
|
|
from typing import Dict, List, Optional, Sequence
|
|
|
|
import numpy as np
|
|
|
|
logger = logging.getLogger("gpu_backtest.kernels")
|
|
|
|
ENV_BACKEND = "FB2_BACKEND" # auto | gpu | cpu
|
|
KERNEL_OPTIONS = ("--fmad=false",) # no FMA contraction: two roundings, like the CPU
|
|
_MODES = ("auto", "gpu", "cpu")
|
|
|
|
# Exit-reason codes written by resolve_trades (see the host mapping in
|
|
# gpu_backtest._resolve_trades_gpu).
|
|
REASON_NO_DATA = 0
|
|
REASON_SL = 1
|
|
REASON_TRAIL = 2
|
|
REASON_TP1 = 3
|
|
REASON_TIMEOUT = 4
|
|
REASON_NO_DATA_AFTER_TIMEOUT = 5 # static path: bars_held/exit_time set, price None
|
|
REASON_NAMES = {
|
|
REASON_NO_DATA: "NO_DATA", REASON_SL: "SL", REASON_TRAIL: "TRAIL",
|
|
REASON_TP1: "TP1", REASON_TIMEOUT: "TIMEOUT",
|
|
REASON_NO_DATA_AFTER_TIMEOUT: "NO_DATA",
|
|
}
|
|
|
|
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
# CUDA root discovery for pip-installed CuPy
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
|
|
def _cuda_runtime_wheel_dir() -> Optional[str]:
|
|
"""Directory of the nvidia-cuda-runtime-cu12 wheel (holds include/)."""
|
|
try:
|
|
spec = importlib.util.find_spec("nvidia.cuda_runtime")
|
|
except (ImportError, ValueError):
|
|
spec = None
|
|
candidates: List[str] = []
|
|
if spec is not None and spec.submodule_search_locations:
|
|
candidates.extend(spec.submodule_search_locations)
|
|
try:
|
|
import site
|
|
for sp in list(site.getsitepackages()) + [site.getusersitepackages()]:
|
|
candidates.append(os.path.join(sp, "nvidia", "cuda_runtime"))
|
|
except Exception:
|
|
pass
|
|
for d in candidates:
|
|
if os.path.isfile(os.path.join(d, "include", "cuda_runtime.h")):
|
|
return d
|
|
return None
|
|
|
|
|
|
def ensure_cuda_path() -> str:
|
|
"""Point CUDA_PATH at the runtime-header wheel when it is unset.
|
|
|
|
Returns a short note for the log ("" when nothing was changed). Runs at
|
|
import (cheap: no CuPy import, no CUDA context) so that any later kernel
|
|
compile — including one triggered by a test that imports cupy itself —
|
|
finds the headers.
|
|
"""
|
|
current = os.environ.get("CUDA_PATH", "")
|
|
if current and os.path.isdir(current):
|
|
return ""
|
|
wheel = _cuda_runtime_wheel_dir()
|
|
if wheel is None:
|
|
return ""
|
|
os.environ["CUDA_PATH"] = wheel
|
|
# CuPy caches a failed lookup; if it was imported before this ran, drop
|
|
# that cache so the new CUDA_PATH is seen (private, hence guarded).
|
|
env_mod = sys.modules.get("cupy._environment")
|
|
if env_mod is not None and getattr(env_mod, "_cuda_path", "") is None:
|
|
env_mod._cuda_path = ""
|
|
return f"CUDA_PATH set to the nvidia-cuda-runtime wheel: {wheel}"
|
|
|
|
|
|
_CUDA_PATH_NOTE = ensure_cuda_path()
|
|
|
|
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
# Backend state
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
|
|
_PROBE_SRC = r'''
|
|
extern "C" __global__
|
|
void fb2_probe(const double* __restrict__ x, double* __restrict__ y, const int n) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) y[i] = x[i] * 2.0 + 1.0;
|
|
}
|
|
'''
|
|
|
|
|
|
class Backend:
|
|
"""CuPy availability, selection mode and the device description."""
|
|
|
|
def __init__(self) -> None:
|
|
mode = os.environ.get(ENV_BACKEND, "auto").strip().lower() or "auto"
|
|
if mode not in _MODES:
|
|
logger.warning("%s=%r is not one of %s; using auto", ENV_BACKEND, mode, _MODES)
|
|
mode = "auto"
|
|
self.mode = mode
|
|
self._probed = False
|
|
self._ok = False
|
|
self.reason = ""
|
|
self.device_name = ""
|
|
self.cupy_version = ""
|
|
self.cp = None # the cupy module once probed OK
|
|
|
|
# ── selection ────────────────────────────────────────────────────
|
|
def set_mode(self, mode: str) -> None:
|
|
mode = (mode or "auto").strip().lower()
|
|
if mode not in _MODES:
|
|
raise ValueError(f"backend mode must be one of {_MODES}, got {mode!r}")
|
|
self.mode = mode
|
|
if mode == "gpu" and not self.available:
|
|
raise RuntimeError(
|
|
f"--backend gpu requested but CUDA is unavailable: {self.reason}")
|
|
|
|
def force_cpu(self) -> None:
|
|
self.mode = "cpu"
|
|
|
|
@contextlib.contextmanager
|
|
def cpu_only(self):
|
|
"""Temporarily route everything to the NumPy twins (tests)."""
|
|
old = self.mode
|
|
self.mode = "cpu"
|
|
try:
|
|
yield
|
|
finally:
|
|
self.mode = old
|
|
|
|
@property
|
|
def available(self) -> bool:
|
|
if self.mode == "cpu":
|
|
return False
|
|
self._probe()
|
|
return self._ok
|
|
|
|
def describe(self) -> str:
|
|
"""The text after ``Backend:`` in the run log and the --record manifest.
|
|
|
|
Prefixes are load-bearing: "CuPy GPU" / "NumPy CPU fallback" are what
|
|
every report and the Module 33 manifest fixture use."""
|
|
if self.available:
|
|
return f"CuPy GPU ({self.device_name}, cupy {self.cupy_version})"
|
|
if self.mode == "cpu":
|
|
return f"NumPy CPU fallback (forced: --backend cpu / {ENV_BACKEND}=cpu)"
|
|
return f"NumPy CPU fallback ({self.reason})"
|
|
|
|
# ── probe ────────────────────────────────────────────────────────
|
|
def _probe(self) -> None:
|
|
if self._probed:
|
|
return
|
|
self._probed = True
|
|
try:
|
|
import cupy as cp # type: ignore
|
|
a = cp.asarray(np.zeros(4, dtype=np.float64))
|
|
_ = cp.cumsum(a).get()
|
|
# The compile probe is the step that used to fail: an importable
|
|
# CuPy still cannot JIT without the CUDA headers.
|
|
k = cp.RawKernel(_PROBE_SRC, "fb2_probe", options=KERNEL_OPTIONS)
|
|
out = cp.empty(4, dtype=np.float64)
|
|
k((1,), (4,), (a, out, np.int32(4)))
|
|
cp.cuda.Device().synchronize()
|
|
if not np.array_equal(out.get(), np.ones(4)):
|
|
raise RuntimeError("probe kernel returned wrong values")
|
|
props = cp.cuda.runtime.getDeviceProperties(cp.cuda.Device().id)
|
|
name = props.get("name", b"")
|
|
if isinstance(name, bytes):
|
|
name = name.decode("utf-8", "replace")
|
|
self.cp = cp
|
|
self._ok = True
|
|
self.device_name = str(name)
|
|
self.cupy_version = str(cp.__version__)
|
|
logger.info("CUDA backend ready: %s, cupy %s%s", self.device_name,
|
|
self.cupy_version,
|
|
f" ({_CUDA_PATH_NOTE})" if _CUDA_PATH_NOTE else "")
|
|
except Exception as e: # noqa: BLE001 — any failure means CPU, loudly
|
|
first = str(e).strip().splitlines()[0] if str(e).strip() else ""
|
|
self.reason = f"{type(e).__name__}: {first}"[:240]
|
|
self._ok = False
|
|
logger.warning("CUDA backend unavailable, NumPy CPU fallback: %s", self.reason)
|
|
|
|
def release_device_memory(self) -> None:
|
|
"""Synchronise and return pooled device/pinned memory before a fork."""
|
|
if not self._ok or self.cp is None:
|
|
return
|
|
cp = self.cp
|
|
cp.cuda.Device().synchronize()
|
|
cp.get_default_memory_pool().free_all_blocks()
|
|
cp.get_default_pinned_memory_pool().free_all_blocks()
|
|
|
|
|
|
BACKEND = Backend()
|
|
|
|
|
|
def _require_cp():
|
|
if not BACKEND.available:
|
|
raise RuntimeError(f"CUDA backend unavailable: {BACKEND.reason or BACKEND.mode}")
|
|
return BACKEND.cp
|
|
|
|
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
# Kernel sources
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
|
|
# Replica of gpu_backtest._gpu_ewm_alpha: y[i] = alpha*x[i] + (1-alpha)*y[i-1],
|
|
# seeded at the first non-NaN x; a NaN input carries y forward without
|
|
# advancing the non-NaN counter; output is NaN until min_periods non-NaN
|
|
# inputs were seen. Rounding intrinsics pin the two-multiply-one-add order.
|
|
_EWM_SRC = r'''
|
|
extern "C" __global__
|
|
void fb2_ewm_batch(const double* __restrict__ x, double* __restrict__ y,
|
|
const long long* __restrict__ offs, const long long* __restrict__ lens,
|
|
const double* __restrict__ alphas, const int* __restrict__ minp,
|
|
const int n_series) {
|
|
int s = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (s >= n_series) return;
|
|
const double alpha = alphas[s];
|
|
const double one_minus = __dsub_rn(1.0, alpha);
|
|
const int mp = minp[s];
|
|
const double* xs = x + offs[s];
|
|
double* ys = y + offs[s];
|
|
const long long n = lens[s];
|
|
double last_y = nan("");
|
|
int seen = 0;
|
|
for (long long i = 0; i < n; ++i) {
|
|
const double xi = xs[i];
|
|
if (isnan(xi)) {
|
|
ys[i] = (!isnan(last_y) && seen >= mp) ? last_y : nan("");
|
|
continue;
|
|
}
|
|
if (isnan(last_y)) last_y = xi;
|
|
else last_y = __dadd_rn(__dmul_rn(alpha, xi), __dmul_rn(one_minus, last_y));
|
|
seen += 1;
|
|
ys[i] = (seen >= mp) ? last_y : nan("");
|
|
}
|
|
}
|
|
'''
|
|
|
|
# Replica of pandas 3.0.2 pandas/_libs/window/aggregations.pyx roll_mean and
|
|
# roll_var for a fixed window, closed='right', min_periods=window (the
|
|
# defaults of Series.rolling(window)). Kahan-compensated running sum with
|
|
# the sign count and consecutive-same-value guard (GH#42064), Welford's
|
|
# online variance with Kahan compensation and the InvCondTol recompute.
|
|
_ROLL_SRC = r'''
|
|
#define FB2_EPS 2.220446049250313e-16
|
|
#define FB2_INVCONDTOL (FB2_EPS * 1e3)
|
|
|
|
__device__ __forceinline__ void fb2_add_mean(double val, long long* nobs, double* sum_x,
|
|
long long* neg_ct, double* comp, long long* ncsv, double* prev_value) {
|
|
if (val == val) {
|
|
*nobs += 1;
|
|
double y = val - *comp;
|
|
double t = *sum_x + y;
|
|
*comp = t - *sum_x - y;
|
|
*sum_x = t;
|
|
if (signbit(val)) *neg_ct += 1;
|
|
if (val == *prev_value) *ncsv += 1; else *ncsv = 1;
|
|
*prev_value = val;
|
|
}
|
|
}
|
|
|
|
__device__ __forceinline__ void fb2_remove_mean(double val, long long* nobs, double* sum_x,
|
|
long long* neg_ct, double* comp) {
|
|
if (val == val) {
|
|
*nobs -= 1;
|
|
double y = -val - *comp;
|
|
double t = *sum_x + y;
|
|
*comp = t - *sum_x - y;
|
|
*sum_x = t;
|
|
if (signbit(val)) *neg_ct -= 1;
|
|
}
|
|
}
|
|
|
|
__device__ __forceinline__ void fb2_add_var(double val, double* nobs, double* mean_x,
|
|
double* ssqdm_x, double* comp, bool* unstable) {
|
|
if (val != val) return;
|
|
double prev_m2 = *ssqdm_x;
|
|
*nobs += 1;
|
|
double prev_mean = *mean_x - *comp;
|
|
double y = val - *comp;
|
|
double t = y - *mean_x;
|
|
*comp = t + *mean_x - y;
|
|
double delta = t;
|
|
if (*nobs != 0) *mean_x = *mean_x + delta / *nobs; else *mean_x = 0;
|
|
*ssqdm_x = *ssqdm_x + (val - prev_mean) * (val - *mean_x);
|
|
if (prev_m2 * FB2_INVCONDTOL > *ssqdm_x) *unstable = true;
|
|
}
|
|
|
|
__device__ __forceinline__ void fb2_remove_var(double val, double* nobs, double* mean_x,
|
|
double* ssqdm_x, double* comp, bool* unstable) {
|
|
if (val == val) {
|
|
double prev_m2 = *ssqdm_x;
|
|
*nobs -= 1;
|
|
if (*nobs != 0) {
|
|
double prev_mean = *mean_x - *comp;
|
|
double y = val - *comp;
|
|
double t = y - *mean_x;
|
|
*comp = t + *mean_x - y;
|
|
double delta = t;
|
|
*mean_x = *mean_x - delta / *nobs;
|
|
*ssqdm_x = *ssqdm_x - (val - prev_mean) * (val - *mean_x);
|
|
if (prev_m2 * FB2_INVCONDTOL > *ssqdm_x) *unstable = true;
|
|
} else {
|
|
*mean_x = 0; *ssqdm_x = 0; *unstable = false;
|
|
}
|
|
}
|
|
}
|
|
|
|
extern "C" __global__
|
|
void fb2_roll_mean_var(const double* __restrict__ x, double* __restrict__ out_mean,
|
|
double* __restrict__ out_var,
|
|
const long long* __restrict__ offs, const long long* __restrict__ lens,
|
|
const int n_series, const int window, const int ddof) {
|
|
int s = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (s >= n_series) return;
|
|
const double* xs = x + offs[s];
|
|
double* om = out_mean + offs[s];
|
|
double* ov = out_var + offs[s];
|
|
const long long n = lens[s];
|
|
if (n <= 0) return;
|
|
const long long minp = window; // min_periods defaults to the window
|
|
// roll_mean state
|
|
double comp_add = 0, comp_rem = 0, sum_x = 0, prev_value = xs[0];
|
|
long long nobs = 0, neg_ct = 0, ncsv = 0;
|
|
// roll_var state (nobs is a double in pandas' roll_var)
|
|
double vnobs = 0, mean_x = 0, ssqdm_x = 0, vcomp_add = 0, vcomp_rem = 0;
|
|
bool unstable = false;
|
|
for (long long i = 0; i < n; ++i) {
|
|
long long st = i - window + 1; if (st < 0) st = 0; // start[i]
|
|
long long en = i + 1; // end[i]
|
|
long long pst = (i - 1) - window + 1; if (pst < 0) pst = 0; // start[i-1]
|
|
long long pen = i; // end[i-1]
|
|
// ---- mean
|
|
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);
|
|
}
|
|
double rm;
|
|
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 = nan("");
|
|
}
|
|
om[i] = rm;
|
|
// ---- var
|
|
bool recompute = (i == 0);
|
|
if (!recompute) {
|
|
for (long long j = pst; j < st; ++j)
|
|
fb2_remove_var(xs[j], &vnobs, &mean_x, &ssqdm_x, &vcomp_rem, &unstable);
|
|
for (long long j = pen; j < en; ++j)
|
|
fb2_add_var(xs[j], &vnobs, &mean_x, &ssqdm_x, &vcomp_add, &unstable);
|
|
}
|
|
if (recompute || unstable) {
|
|
mean_x = ssqdm_x = vnobs = vcomp_add = vcomp_rem = 0;
|
|
for (long long j = st; j < en; ++j)
|
|
fb2_add_var(xs[j], &vnobs, &mean_x, &ssqdm_x, &vcomp_add, &unstable);
|
|
unstable = false;
|
|
}
|
|
ov[i] = (vnobs >= (double)minp && vnobs > (double)ddof)
|
|
? ssqdm_x / (vnobs - (double)ddof) : nan("");
|
|
}
|
|
}
|
|
'''
|
|
|
|
# Replica of gpu_backtest._simulate_trailing (trailing != 0) and of the
|
|
# static SL/TP1 branch of gpu_backtest.simulate_trades (trailing == 0).
|
|
# Python max/min are written out as (b > a) ? b : a so NaN handling matches.
|
|
_TRADES_SRC = r'''
|
|
#define R_NO_DATA 0
|
|
#define R_SL 1
|
|
#define R_TRAIL 2
|
|
#define R_TP1 3
|
|
#define R_TIMEOUT 4
|
|
#define R_NO_DATA_AFTER_TIMEOUT 5
|
|
|
|
__device__ __forceinline__ void fb2_excursions(const double* h, const double* l, long long bars,
|
|
bool buy, double entry, double risk, double* mfe_out, double* mae_out) {
|
|
// np.nanmax(h[:bars]) / np.nanmin(l[:bars]); all-NaN gives NaN.
|
|
double hmax = nan(""), lmin = nan("");
|
|
for (long long j = 0; j < bars; ++j) {
|
|
double hv = h[j], lv = l[j];
|
|
if (!isnan(hv) && (isnan(hmax) || hv > hmax)) hmax = hv;
|
|
if (!isnan(lv) && (isnan(lmin) || lv < lmin)) lmin = lv;
|
|
}
|
|
double mfe, mae;
|
|
if (buy) { mfe = (hmax - entry) / risk; mae = (entry - lmin) / risk; }
|
|
else { mfe = (entry - lmin) / risk; mae = (hmax - entry) / risk; }
|
|
*mfe_out = (mfe > 0.0) ? mfe : 0.0; // Python max(0.0, x): NaN -> 0.0
|
|
*mae_out = (mae > 0.0) ? mae : 0.0;
|
|
}
|
|
|
|
extern "C" __global__
|
|
void fb2_resolve_trades(const double* __restrict__ opens, const double* __restrict__ highs,
|
|
const double* __restrict__ lows, const double* __restrict__ closes,
|
|
const long long* __restrict__ valid_lens,
|
|
const double* __restrict__ entries, const double* __restrict__ sls,
|
|
const double* __restrict__ tps, const signed char* __restrict__ is_buy,
|
|
const double* __restrict__ arms, const double* __restrict__ dists,
|
|
const double spread, const int trailing, const int max_bars,
|
|
const int n_trades,
|
|
long long* __restrict__ exit_idx_out, double* __restrict__ exit_price_out,
|
|
int* __restrict__ reason_out, long long* __restrict__ bars_held_out,
|
|
double* __restrict__ r_out, double* __restrict__ mfe_out,
|
|
double* __restrict__ mae_out) {
|
|
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i >= n_trades) return;
|
|
const long long base = (long long)i * max_bars;
|
|
const double* o = opens + base;
|
|
const double* h = highs + base;
|
|
const double* l = lows + base;
|
|
const double* c = closes + base;
|
|
const long long n = valid_lens[i];
|
|
const bool buy = is_buy[i] != 0;
|
|
const double entry = entries[i];
|
|
const double sl = sls[i];
|
|
const double tp = tps[i];
|
|
// defaults (Trade dataclass defaults / "not set")
|
|
exit_idx_out[i] = -1; exit_price_out[i] = nan(""); reason_out[i] = R_NO_DATA;
|
|
bars_held_out[i] = 0; r_out[i] = 0.0; mfe_out[i] = 0.0; mae_out[i] = 0.0;
|
|
if (n == 0) return;
|
|
|
|
if (trailing) {
|
|
const double risk = fabs(entry - sl);
|
|
if (risk <= 0) return; // NO_DATA
|
|
const double arm = arms[i], dist = dists[i];
|
|
double cur_sl = sl, peak = entry;
|
|
bool have_price = false; double exit_price = nan("");
|
|
long long exit_idx = n - 1; int reason = R_TIMEOUT;
|
|
for (long long j = 0; j < n; ++j) {
|
|
if (isnan(h[j])) {
|
|
if (j == 0) { exit_idx = 0; reason = R_NO_DATA; }
|
|
else { exit_idx = j - 1; exit_price = c[exit_idx]; have_price = true; }
|
|
break;
|
|
}
|
|
const bool hit = buy ? (l[j] <= cur_sl) : (h[j] >= cur_sl);
|
|
if (hit) {
|
|
exit_price = cur_sl; have_price = true;
|
|
if ((buy && o[j] < cur_sl) || (!buy && o[j] > cur_sl)) exit_price = o[j];
|
|
reason = (cur_sl == sl) ? R_SL : R_TRAIL;
|
|
exit_idx = j;
|
|
break;
|
|
}
|
|
const double fav = buy ? h[j] : l[j];
|
|
peak = buy ? ((fav > peak) ? fav : peak) : ((fav < peak) ? fav : peak);
|
|
const double exc = (buy ? (peak - entry) : (entry - peak)) / risk;
|
|
if (exc >= arm) {
|
|
const double lvl = buy ? (peak - dist * risk) : (peak + dist * risk);
|
|
cur_sl = buy ? ((lvl > cur_sl) ? lvl : cur_sl) : ((lvl < cur_sl) ? lvl : cur_sl);
|
|
}
|
|
}
|
|
if (reason == R_NO_DATA) return;
|
|
if (!have_price) exit_price = c[exit_idx];
|
|
if (!isfinite(exit_price)) return; // NO_DATA, nothing else set
|
|
const long long bars = exit_idx + 1;
|
|
reason_out[i] = reason; bars_held_out[i] = bars; exit_idx_out[i] = exit_idx;
|
|
exit_price_out[i] = exit_price;
|
|
const double pnl = buy ? (exit_price - entry) : (entry - exit_price);
|
|
r_out[i] = (pnl - spread) / risk;
|
|
fb2_excursions(h, l, bars, buy, entry, risk, &mfe_out[i], &mae_out[i]);
|
|
return;
|
|
}
|
|
|
|
// ---- static SL / TP1: first-hit scan over the whole window, NaN highs masked
|
|
long long sl_idx = max_bars, tp_idx = max_bars;
|
|
for (long long j = 0; j < max_bars; ++j) {
|
|
if (isnan(h[j])) continue;
|
|
const bool sl_hit = buy ? (l[j] <= sl) : (h[j] >= sl);
|
|
const bool tp_hit = buy ? (h[j] >= tp) : (l[j] <= tp);
|
|
if (sl_hit && j < sl_idx) sl_idx = j;
|
|
if (tp_hit && j < tp_idx) tp_idx = j;
|
|
}
|
|
long long exit_idx; double exit_price; int reason; long long bars;
|
|
if (sl_idx < n && sl_idx <= tp_idx) {
|
|
reason = R_SL; bars = sl_idx + 1; exit_price = sl; exit_idx = sl_idx;
|
|
const double bar_open = o[sl_idx];
|
|
if (!isnan(bar_open)) {
|
|
if ((buy && bar_open < sl) || (!buy && bar_open > sl)) exit_price = bar_open;
|
|
}
|
|
} else if (tp_idx < n) {
|
|
reason = R_TP1; bars = tp_idx + 1; exit_price = tp; exit_idx = tp_idx;
|
|
const double bar_open = o[tp_idx];
|
|
if (!isnan(bar_open)) {
|
|
if ((buy && bar_open > tp) || (!buy && bar_open < tp)) exit_price = bar_open;
|
|
}
|
|
} else {
|
|
reason = R_TIMEOUT; bars = n; exit_idx = n - 1; exit_price = c[exit_idx];
|
|
if (!isfinite(exit_price)) {
|
|
reason_out[i] = R_NO_DATA_AFTER_TIMEOUT; bars_held_out[i] = bars;
|
|
exit_idx_out[i] = exit_idx;
|
|
return; // exit_price None, R untouched
|
|
}
|
|
}
|
|
reason_out[i] = reason; bars_held_out[i] = bars; exit_idx_out[i] = exit_idx;
|
|
exit_price_out[i] = exit_price;
|
|
const double risk = fabs(entry - sl);
|
|
if (risk > 0) {
|
|
const double pnl = buy ? (exit_price - entry) : (entry - exit_price);
|
|
r_out[i] = (pnl - spread) / risk;
|
|
}
|
|
if (bars > 0 && risk > 0)
|
|
fb2_excursions(h, l, bars, buy, entry, risk, &mfe_out[i], &mae_out[i]);
|
|
}
|
|
'''
|
|
|
|
_KERNELS: Dict[str, object] = {} # module-level: never collected in a child
|
|
|
|
|
|
def _kernel(name: str):
|
|
cp = _require_cp()
|
|
k = _KERNELS.get(name)
|
|
if k is None:
|
|
src = {"ewm": _EWM_SRC, "roll": _ROLL_SRC, "trades": _TRADES_SRC}[name]
|
|
fn = {"ewm": "fb2_ewm_batch", "roll": "fb2_roll_mean_var",
|
|
"trades": "fb2_resolve_trades"}[name]
|
|
k = cp.RawKernel(src, fn, options=KERNEL_OPTIONS)
|
|
_KERNELS[name] = k
|
|
return k
|
|
|
|
|
|
def _ragged(cp, series: Sequence) -> tuple:
|
|
"""Concatenate 1-D float64 series on the device; return (flat, offs, lens)."""
|
|
arrs = [cp.asarray(s, dtype=cp.float64).ravel() for s in series]
|
|
lens = np.array([int(a.shape[0]) for a in arrs], dtype=np.int64)
|
|
offs = np.zeros(len(arrs), dtype=np.int64)
|
|
if len(arrs) > 1:
|
|
offs[1:] = np.cumsum(lens)[:-1]
|
|
flat = cp.concatenate(arrs) if arrs else cp.empty(0, dtype=cp.float64)
|
|
return flat, cp.asarray(offs), cp.asarray(lens), offs, lens
|
|
|
|
|
|
def _launch_1d(n: int, threads: int = 128) -> tuple:
|
|
return ((max(1, (n + threads - 1) // threads),), (threads,))
|
|
|
|
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
# Public wrappers (device in/out; callers .get() once)
|
|
# ═════════════════════════════════════════════════════════════════════
|
|
|
|
def ewm_batch_dev(series: Sequence, alphas: Sequence[float],
|
|
min_periods: Sequence[int]) -> List:
|
|
"""EWM of every series in one launch. Inputs may be NumPy or CuPy;
|
|
returns CuPy views (one per series)."""
|
|
cp = _require_cp()
|
|
if len(series) != len(alphas) or len(series) != len(min_periods):
|
|
raise ValueError("series, alphas and min_periods must have equal length")
|
|
if not series:
|
|
return []
|
|
flat, d_offs, d_lens, offs, lens = _ragged(cp, series)
|
|
out = cp.empty_like(flat)
|
|
k = _kernel("ewm")
|
|
grid, block = _launch_1d(len(series))
|
|
k(grid, block, (flat, out, d_offs, d_lens,
|
|
cp.asarray(np.asarray(alphas, dtype=np.float64)),
|
|
cp.asarray(np.asarray(min_periods, dtype=np.int32)),
|
|
np.int32(len(series))))
|
|
return [out[o:o + n] for o, n in zip(offs, lens)]
|
|
|
|
|
|
def ewm_batch(series: Sequence[np.ndarray], alphas: Sequence[float],
|
|
min_periods: Sequence[int]) -> List[np.ndarray]:
|
|
"""Host convenience: NumPy in, NumPy out."""
|
|
return [y.get() for y in ewm_batch_dev(series, alphas, min_periods)]
|
|
|
|
|
|
def roll_mean_var_dev(series: Sequence, window: int, ddof: int = 0) -> tuple:
|
|
"""pandas rolling(window).mean() and .var(ddof) of every series in one
|
|
launch; returns (means, vars) as lists of CuPy views."""
|
|
cp = _require_cp()
|
|
if not series:
|
|
return [], []
|
|
flat, d_offs, d_lens, offs, lens = _ragged(cp, series)
|
|
out_mean = cp.empty_like(flat)
|
|
out_var = cp.empty_like(flat)
|
|
k = _kernel("roll")
|
|
grid, block = _launch_1d(len(series))
|
|
k(grid, block, (flat, out_mean, out_var, d_offs, d_lens,
|
|
np.int32(len(series)), np.int32(window), np.int32(ddof)))
|
|
means = [out_mean[o:o + n] for o, n in zip(offs, lens)]
|
|
vars_ = [out_var[o:o + n] for o, n in zip(offs, lens)]
|
|
return means, vars_
|
|
|
|
|
|
def roll_std_from_var_dev(var):
|
|
"""pandas zsqrt: sqrt(var) with a negative variance clipped to 0."""
|
|
cp = _require_cp()
|
|
return cp.where(var < 0, 0.0, cp.sqrt(var))
|
|
|
|
|
|
def roll_mean_var(series: Sequence[np.ndarray], window: int, ddof: int = 0) -> tuple:
|
|
"""Host convenience: (means, vars, stds) as NumPy arrays."""
|
|
means, vars_ = roll_mean_var_dev(series, window, ddof)
|
|
stds = [roll_std_from_var_dev(v) for v in vars_]
|
|
return ([m.get() for m in means], [v.get() for v in vars_],
|
|
[s.get() for s in stds])
|
|
|
|
|
|
def resolve_trades(opens: np.ndarray, highs: np.ndarray, lows: np.ndarray,
|
|
closes: np.ndarray, valid_lens: np.ndarray,
|
|
entries: np.ndarray, sls: np.ndarray, tps: np.ndarray,
|
|
is_buy: np.ndarray, arms: np.ndarray, dists: np.ndarray,
|
|
spread: float, trailing: bool) -> Dict[str, np.ndarray]:
|
|
"""Resolve every trade's exit on the device; one thread per trade.
|
|
|
|
``opens/highs/lows/closes`` are (n_trades, max_bars) float64 windows
|
|
padded with NaN, ``valid_lens`` the real length of each. Returns host
|
|
arrays: exit_idx, exit_price, reason (REASON_* codes), bars_held, r,
|
|
mfe, mae.
|
|
"""
|
|
cp = _require_cp()
|
|
n_trades, max_bars = opens.shape
|
|
if n_trades == 0:
|
|
empty = np.zeros(0)
|
|
return {k: empty for k in ("exit_idx", "exit_price", "reason",
|
|
"bars_held", "r", "mfe", "mae")}
|
|
d = lambda a, dt: cp.asarray(np.ascontiguousarray(a, dtype=dt)) # noqa: E731
|
|
d_o, d_h, d_l, d_c = (d(opens, np.float64), d(highs, np.float64),
|
|
d(lows, np.float64), d(closes, np.float64))
|
|
o_exit_idx = cp.empty(n_trades, dtype=np.int64)
|
|
o_exit_price = cp.empty(n_trades, dtype=np.float64)
|
|
o_reason = cp.empty(n_trades, dtype=np.int32)
|
|
o_bars = cp.empty(n_trades, dtype=np.int64)
|
|
o_r = cp.empty(n_trades, dtype=np.float64)
|
|
o_mfe = cp.empty(n_trades, dtype=np.float64)
|
|
o_mae = cp.empty(n_trades, dtype=np.float64)
|
|
k = _kernel("trades")
|
|
grid, block = _launch_1d(n_trades)
|
|
k(grid, block, (d_o, d_h, d_l, d_c, d(valid_lens, np.int64),
|
|
d(entries, np.float64), d(sls, np.float64), d(tps, np.float64),
|
|
d(is_buy, np.int8), d(arms, np.float64), d(dists, np.float64),
|
|
np.float64(spread), np.int32(1 if trailing else 0),
|
|
np.int32(max_bars), np.int32(n_trades),
|
|
o_exit_idx, o_exit_price, o_reason, o_bars, o_r, o_mfe, o_mae))
|
|
cp.cuda.Device().synchronize()
|
|
return dict(exit_idx=o_exit_idx.get(), exit_price=o_exit_price.get(),
|
|
reason=o_reason.get(), bars_held=o_bars.get(), r=o_r.get(),
|
|
mfe=o_mfe.get(), mae=o_mae.get())
|