// crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu
//
// SP15 Wave 3a (2026-05-06) — post-loop derivation of per-bar position
// state from the recorded factored-action history.
//
// Reads `actions_history_buf[w * max_len + t]` (factored ints recorded
// by `experience_env_step` / `backtest_env_step`) and emits three f32
// streams per (window, bar) consumed by Wave 3b's cost-net sharpe path:
//
//   position_history_out[w * max_len + t]  ∈ {-1.0, 0.0, +1.0}
//   side_ind_out[w * max_len + t]           1.0 when position changed
//                                           from prev step, else 0.0
//   rt_ind_out[w * max_len + t]             1.0 when transitioning to
//                                           flat from a non-flat
//                                           position (round-trip
//                                           close), else 0.0
//
// Single-block-per-window launch with one thread per window. Each thread
// walks bars sequentially because position(t) depends on position(t-1)
// — there's no parallelism within a window for this state machine. The
// per-window walks ARE parallel across windows (n_windows blocks).
// No atomicAdd per `feedback_no_atomicadd` (the kernel doesn't reduce;
// it's a pure scan).
//
// Sentinel handling: `actions_history_buf` is initialised to -1 by
// `experience_kernels.cu` (per the actions_history measurement-bug fix
// from a prior commit). Bars where `action == -1` (never written by
// env_step — typically because the window terminated early) are treated
// as "no position change" — we keep the prior position. For windows
// shorter than `max_len`, the trailing -1 sentinels emit position =
// last-known-position with side_ind=0, rt_ind=0; downstream consumers
// (Wave 3b) gate on `window_lens[w]` to suppress these tail bars.

#include "action_decoding_helpers.cuh"

extern "C" __global__ void sp15_position_history_derivation_kernel(
    const int*   __restrict__ actions_history_buf,  /* [n_windows * max_len] */
    int                       n_windows,
    int                       max_len,
    int                       b1_size,
    int                       b2_size,
    int                       b3_size,
    float* __restrict__       position_history_out, /* [n_windows * max_len] */
    float* __restrict__       side_ind_out,         /* [n_windows * max_len] */
    float* __restrict__       rt_ind_out            /* [n_windows * max_len] */
) {
    const int w = (int)blockIdx.x;
    if (w >= n_windows) return;
    if (threadIdx.x != 0) return;  /* one thread per block per window */

    int prev_position = 0;  /* every window starts flat by convention */
    for (int t = 0; t < max_len; ++t) {
        const int idx = w * max_len + t;
        const int action = actions_history_buf[idx];

        int new_position;
        if (action < 0) {
            /* sentinel — keep prior position, emit no-change indicators */
            new_position = prev_position;
        } else {
            new_position = factored_action_to_position(
                action, b1_size, b2_size, b3_size, prev_position
            );
        }

        const int changed   = (new_position != prev_position) ? 1 : 0;
        const int round_trip = (changed && new_position == 0 && prev_position != 0) ? 1 : 0;

        position_history_out[idx] = (float)new_position;
        side_ind_out[idx]         = (float)changed;
        rt_ind_out[idx]           = (float)round_trip;

        prev_position = new_position;
    }
    __threadfence_system();
}
