// crates/ml/src/cuda_pipeline/dd_state_kernel.cu
//
// SP15 Phase 1.3 — per-step drawdown state tracking.
//
// Reads existing PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) from
// the position state buffer (canonical equity tracking — no new ISV slot
// needed; invariant verified during plan v2 critical review). Writes 6
// ISV slots per spec §6.3:
//
//   ISV[DD_CURRENT_INDEX=401]        : current_dd = max(0, (peak − equity) / peak)
//   ISV[DD_MAX_INDEX=402]            : running max of current_dd
//   ISV[DD_RECOVERY_BARS_INDEX=403]  : bars since last new high-water mark
//   ISV[DD_PERSISTENCE_INDEX=404]    : bars since last new high (twin of recovery for now;
//                                       split lifetime semantics may differ in later phases)
//   ISV[CALMAR_INDEX=405]            : floored max_dd (= max(dd_max, 1e-4)) — denominator
//                                       for the host composer's calmar = mean_pnl / value-here.
//                                       Kernel writes ONLY the denominator floor; the host
//                                       composes the calmar ratio so that mean_pnl source can
//                                       evolve without kernel changes. The 1e-4 floor
//                                       eliminates the saturation-at-100 artifact seen in the
//                                       train-dd4xl HEALTH_DIAG.
//   ISV[DD_PCT_INDEX=406]            : current_dd / max(dd_budget, 1e-4), clipped to 1.0
//
// Per-step state-machine kernel (single-thread, single-block). Mirrors the
// host-side draw-down accounting that previously lived in the CPU emit
// path; lifts the truth source onto the GPU per `feedback_cpu_is_read_only`.
//
// Layout fingerprint contract: PS_PEAK_EQUITY and PS_PREV_EQUITY are
// `#define`d in state_layout.cuh — use the macros so the kernel tracks
// any future portfolio-state reshuffle automatically.
//
// Allowed-write rule: the kernel writes ISV slots that the registry has
// FoldReset entries for (sentinels documented in state_reset_registry.rs)
// — pos_state[PS_*] writes go through the same slots that the env-step
// kernel already owns, so this kernel does NOT introduce new pos_state
// ownership; it only updates the existing high-water mark and per-step
// equity slots which the env-step kernel updates from gross PnL normally.
// Phase 1.3 lands the kernel + launcher + registry only (per
// `feedback_no_partial_refactor.md`); production wiring lands in a
// follow-up atomic commit alongside the HEALTH_DIAG composer.

#include "state_layout.cuh"

extern "C" __global__ void dd_state_kernel(
    float                pnl_step,
    float                dd_budget,
    float* __restrict__  isv,
    float* __restrict__  pos_state
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    // Update equity (PS_PREV_EQUITY = slot 9): new_equity = prev_equity + pnl_step.
    const float prev_equity = pos_state[PS_PREV_EQUITY];
    const float new_equity  = prev_equity + pnl_step;
    pos_state[PS_PREV_EQUITY] = new_equity;

    // Update peak (PS_PEAK_EQUITY = slot 7) and recovery counters.
    float peak = pos_state[PS_PEAK_EQUITY];
    if (new_equity > peak) {
        peak = new_equity;
        pos_state[PS_PEAK_EQUITY] = peak;
        // New high-water mark: zero recovery + persistence counters.
        isv[403] = 0.0f;  // DD_RECOVERY_BARS_INDEX
        isv[404] = 0.0f;  // DD_PERSISTENCE_INDEX
    } else {
        // Below peak: increment both counters by 1.
        isv[403] = isv[403] + 1.0f;  // DD_RECOVERY_BARS_INDEX
        isv[404] = isv[404] + 1.0f;  // DD_PERSISTENCE_INDEX
    }

    // current_dd = max(0, (peak − equity) / peak), guarded against peak ≤ 0.
    const float current_dd = (peak > 0.0f)
        ? fmaxf(0.0f, (peak - new_equity) / peak)
        : 0.0f;
    isv[401] = current_dd;  // DD_CURRENT_INDEX

    // Running max of current_dd.
    const float prior_max = isv[402];  // DD_MAX_INDEX
    const float new_max   = (current_dd > prior_max) ? current_dd : prior_max;
    isv[402] = new_max;

    // dd_pct = clip(current_dd / max(dd_budget, 1e-4), 0, 1).
    const float dd_pct = fminf(1.0f, current_dd / fmaxf(dd_budget, 1e-4f));
    isv[406] = dd_pct;  // DD_PCT_INDEX

    // CALMAR slot stores the floored max_dd denominator; host composes
    // calmar = mean_pnl / value-here. Floor eliminates the saturation-at-100
    // artifact that the legacy host-side composer exhibited when max_dd ≈ 0.
    isv[405] = fmaxf(new_max, 1e-4f);  // CALMAR_INDEX

    __threadfence_system();
}
