manifoldbt/docs/gpu-derived-series-plan.md

# Plan — Séries dérivées hoistées sur GPU (« option C ») (2026-07-17)

Objectif : permettre aux ops fenêtrées du transpileur (`ZScore`, `RollingMean`,
`Rsi`, `Roc`, futurs `RollingStd/Sum/...`) de prendre une **expression** en
entrée, pas seulement une colonne brute — ce qui débloque l'exemple 16
(`zscore(price_ratio - hr_ratio, zwin)`) et toute la famille « indicateur d'un
spread/ratio » sur GPU.

## Le problème (analysé le 2026-07-17, voir la discussion du rapport)

Le kernel évalue en **streaming** : les intermédiaires vivent une barre dans des
registres. Une op fenêtrée doit évincer la valeur qui sort de la fenêtre
(`sum -= x[jj-w]`) : pour une colonne (`CCOL`) ou une série hoistée (`HIND`)
c'est un accès mémoire ; pour une expression dérivée, la valeur d'il y a `w`
barres a été jetée. Le CPU n'a pas ce problème : son évaluateur matérialise
chaque nœud du DAG en array.

Seul `EwmMean` accepte une expression aujourd'hui — une EMA ne regarde jamais
en arrière. `RollingMean`, `Rsi`, `ZScore` exigent tous une colonne brute ;
cette limite est **préexistante et structurelle**, pas un choix de mon ZScore.

## Le design

**Matérialiser l'expression d'entrée comme lignes `hind` supplémentaires**,
dédupliquées par résolution de paramètres, remplies par un kernel généré ; l'op
fenêtrée reste un **fold in-thread** (par combo) qui lit la ligne dérivée en
accès aléatoire (`HINDAT(j, jj)` / `HINDAT(j, jj-w)`).

```
stage 1  hoist_fill    : indicateurs feuilles (SMA/EMA/RSI/ROC/ZScore sur colonnes)
stage 2  derived_fill  : expressions POINTWISE sur (cols ∪ lignes stage 1)   ← NOUVEAU
sim      sweep_t       : folds fenêtrés in-thread lisant les lignes dérivées ← ÉTENDU
```

Décisions clés :

- **Éligibilité** : l'entrée doit être pointwise sur (colonnes brutes ∪ feuilles
  hoistables ∪ littéraux ∪ params). Sinon → CPU, raison nommée (Phase 0 style) :
  « windowed op over a non-pointwise expression ».
- **Dédup** : clé = empreinte structurelle de l'expression + bits résolus de
  CHAQUE param référencé (vecteur, pas un seul u64 — une expression peut lire
  plusieurs params). Même mécanique HashMap que `build_hoist_plan`. Nombre de
  lignes = résolutions distinctes de l'ENTRÉE (ex. 16 : #smooth distincts, PAS
  #(smooth × zwin) — la fenêtre reste par combo, in-thread). C'est ce qui rend
  la VRAM raisonnable à 1M combos.
- **`derived_fill` est un kernel GÉNÉRÉ** (comme `sweep_t`), pas un kernel fixe :
  le corps pointwise vient du MÊME `PipeCtx::emit` que le sim (mêmes sémantiques
  NaN, division-par-zéro → NaN, etc. — zéro seconde implémentation qui dérive).
  Pointwise ⇒ barres indépendantes ⇒ un thread par (ligne, barre), coalescé,
  ~1 ms pour 1000 lignes × 44k barres. Lancé entre hoist_fill et sweep_t.
- **Chaînage limité en V1** : dérivé = pointwise uniquement. Un fenêtré DANS un
  fenêtré sur expression (`zscore(sma(expr))`) → CPU, raison nommée. (Extensible
  plus tard en étages topologiques ; hors périmètre V1.)
- **Folds in-thread** : sémantique `*_no_nulls` de bt-expr AU COMPLET, y compris
  `nan_count` — voir P0 ci-dessous. État ZScore in-thread ≈ 4-6 registres.
- **fp32 scan** : lignes dérivées calculées en f64, stockées f32 en mode scan
  (même argument que le float-hind de la Part 5 : identique au cast côté
  lecteur).
- **Garde VRAM** : bytes(hind + derived) > free/2 → CPU, raison nommée.
- **Multi-asset** : hors V1 (comme l'exo), raison nommée.

## Bit-identité (la chaîne de preuve)

CPU : matérialise l'entrée (array) puis `zscore_no_nulls(vals)`. GPU : la ligne
dérivée doit être **bit-identique à l'array CPU** (mêmes ops pointwise dans le
même ordre via le même emit(), fmad off, mêmes sources : cols identiques,
lignes stage-1 déjà pinnées à bt-expr) ; puis le fold in-thread refait
exactement `zscore_no_nulls` sur les mêmes valeurs. Chaque maillon testé.

## Étapes

### P0 — Bug latent découvert par l'analyse (à faire AVANT tout)
Le fold SMA **in-thread** (`ss += CCOL(...)`, chemin non-hoisté = multi-asset)
n'a AUCUNE gestion NaN, alors que `rolling_mean_pair` (bt-expr) exclut les NaN
de ses sommes (`bad_a`/`bad_b`). Divergence silencieuse sur données à trous en
multi-asset — même famille que `804de5e`, invisible parce que les tests NaN
tournent en mono (hoisting actif) et que le e2e multi n'a pas de NaN.
- Auditer les 3 folds in-thread (SMA / RSI / ROC) contre bt-expr sur NaN.
- Test : e2e multi avec NaN closes (étendre `gpu_sweep_multi_e2e`), vérifier
  qu'il MORD (rituel du bug délibéré), corriger les folds.
- Estimation : ½ journée. Indépendant de C mais bloquant : les folds dérivés
  réutiliseront ce code.

### P1 — Enregistrement des lignes dérivées (host)
- `PipeCtx` : détection d'éligibilité (fn `is_pointwise_over_leaves`),
  empreinte, `derived_slot(expr) -> dj`, table `derived: Vec<DerivedSpec>`.
- `build_hoist_plan` étendu : résolution par combo des params de chaque
  expression dérivée, dédup, `hrow` étendu aux lignes dérivées, garde VRAM.

### P2 — Kernel `derived_fill` généré + wiring
- Génération du source (corps par ligne via emit(), switch sur le kind de
  ligne), compile NVRTC (cache par hash), launch stage 2, slot dans
  `MBT_GPU_PROBE`.
- `ptxas -v` sur le kernel généré (attendu trivial : pointwise, pas d'état).

### P3 — Folds in-thread sur lignes dérivées
- Macro `HINDAT(j, idx)` (HIND à index arbitraire).
- `ZScore` d'abord (débloque ex. 16), puis `RollingMean`, `Rsi`, `Roc` :
  entrée Column → chemins existants ; entrée dérivée-éligible → fold in-thread
  avec sémantique bt-expr complète.
- **GATE ptxas** : le sim est à 48/48 registres. Si le fold in-thread fait
  spiller, relâcher `__launch_bounds__` UNIQUEMENT pour les sources générées qui
  contiennent des folds in-thread (le bound est par source générée, donc par
  sweep) — mesurer, pas supposer.

### P4 — Tests (chaque test vérifié MORDANT via bug délibéré)
- `derived_fill_matches_bt_expr` : ligne dérivée vs array intermédiaire bt-expr,
  barre par barre, NaN gaps inclus (harnais `hoist_fill_matches_bt_expr` étendu).
- Fold in-thread vs bt-expr : zscore(expr) complet, barre par barre.
- e2e forme exemple 16 (exo + spread + zscore, périodes en params) :
  bit-identité CPU==GPU + non-vacuité + appel direct `run_sweep_lite_gpu`
  (une retombée = échec, pas une comparaison CPU-vs-CPU silencieuse).
- Retombées nommées : non-pointwise, fenêtré imbriqué, VRAM, multi-asset.

### P5 — Bench + rapport
- Exemple 16 réel (hashrate), grille ~10k : cible GPU >> CPU (aujourd'hui :
  2 589 c/s CPU, GPU retombe). Section dans `gpu_funding_bench.md`.
- Mettre à jour `docs/gpu-sweep-coverage-plan.md` (la Phase 2b absorbe C).

## Estimation & risques

| | |
|---|---|
| P0 | ½ j (correction comprise si le bug se confirme) |
| P1+P2 | ~1 j |
| P3 | ½ j + gate registres |
| P4+P5 | ½ j |
| **Total** | **~2-2,5 jours** |

Risques : (1) pression registres des folds in-thread — gate ptxas, bounds par
source ; (2) explosion du nombre de lignes dérivées si une expression référence
plusieurs params à forte cardinalité — garde VRAM + retombée nommée ; (3) toute
divergence NaN — couverte par les tests barre-par-barre ancrés bt-expr, jamais
par des miroirs.

---

## STATUT : LIVRÉ (2026-07-17)

- P0 : fold SMA in-thread corrigé (bad-count miroir de `rolling`), RSI/ROC
  audités conformes ; e2e `gpu_multi_sma_nan_volume_{2,5}` vérifié mordant.
- P1-P2 : `DerivedReg`/`derived_slot`, plan étendu (dédup multi-params, hrow
  absolu, garde VRAM sur le total), `derived_fill` généré par le même emit
  (macros remappées), `hoist_fill` a un stride explicite.
- P3 : fold ZScore in-thread sur ligne dérivée via `HINDAT` ; GATE ptxas :
  bounds relâchés à 7 (f64) / 8 (fp32) blocs pour les sources à folds, zéro
  spill ; fold-free inchangé (48 regs @10).
- P4 : `derived_fill_matches_bt_expr` (mord sur +1e-12 que l'e2e ne voit pas),
  e2e bit-identité + retombée nommée fenêtré-imbriqué.
- P5 : exemple 16 réel, 102x à 30k combos, Part 11 du rapport.
- Reste (hors V1, repris en Phase 2b) : `RollingMean`/`Rsi`/`Roc` sur entrée
  dérivée (mécanique identique au ZScore), multi-asset, chaînage fenêtré.