/* ══════════════════════════════════════════════════════════════════════════
 * SP13 Layer B — Commit B1.1b producer kernel for aux next-bar sign label
 * (2026-05-05).
 *
 * Replaces the B0 `alloc_zeros::<i32>(total)` placeholder for
 * `aux_sign_labels` in `GpuExperienceCollector::collect_experiences_gpu`.
 * The B0 plumbing carried `0` everywhere; with B1.1a the K=2 softmax CE
 * head consumes -1/0/1 labels (-1 = mask/skip, 0 = down/flat, 1 = up).
 * Without this producer, the model trains on "all bars are class 0" — the
 * known degraded behavior between B1.1a and B1.1b. This kernel computes
 * the real label per experience from the price trajectory.
 *
 * Per-thread O(1) map: read `targets[bar*6+2]` (raw_close column — see
 * `dt_kernels.cu:803` and `scripted_policy_kernel.cu:56` for the canonical
 * column layout: 0=preproc_close, 1=preproc_next, 2=raw_close, 3=raw_next,
 * 4=raw_open, 5=mid_open) at `bar` and `bar+lookahead`, classify by
 * strict greater-than — flat/down both map to label 0 because the K=2
 * softmax can only distinguish "up" from "not-up". The skip sentinel -1
 * fires when the lookahead window would run past `total_bars`; the CE
 * loss reduce kernel masks it (`label < 0` → row excluded from
 * mean + B_valid count per `aux_heads_kernel.cu`).
 *
 * Pure map; no atomicAdd, no reduction, no shared memory — block
 * tree-reduce N/A here per `feedback_no_atomicadd.md`. GPU-only data
 * read/write per `feedback_cpu_is_read_only.md`. Inputs are mapped-pinned
 * (targets) or device-allocated (bar_indices, out_labels) per
 * `feedback_no_htod_htoh_only_mapped_pinned.md`.
 *
 * Launch config: grid=(ceil(total / 256)), block=(256). Mirrors
 * `gather_i32_scalar` (graph_utility_kernels.cu:96) and
 * `hindsight_relabel_kernel` (experience_kernels.cu) — same tiling
 * convention used everywhere a per-experience map kernel runs in this
 * pipeline.
 *
 * Args:
 *   targets       — `[total_bars, 6]` mapped-pinned f32 (raw_close at col 2).
 *   bar_indices   — `[total]` mapped-pinned i32, bar index for each
 *                   experience (constructed by host in
 *                   `collect_experiences_gpu`'s pre-existing
 *                   `bar_indices_pinned` fill loop).
 *   out_labels    — `[total]` device i32; receives -1/0/1 per experience.
 *   total         — number of experiences.
 *   total_bars    — total bars in the data series (for skip-sentinel
 *                   bound check: bar+lookahead >= total_bars ⇒ -1).
 *   lookahead     — number of bars to look ahead (structural design
 *                   constant, value passed in by the host launcher;
 *                   must match the trainer's expectation since the K=2
 *                   head's "up" label semantics encode the same
 *                   lookahead window).
 * ══════════════════════════════════════════════════════════════════════════ */
extern "C" __global__
void aux_sign_label_kernel(
    const float* __restrict__ targets,
    const int*   __restrict__ bar_indices,
    int*         __restrict__ out_labels,
    int total,
    int total_bars,
    int lookahead)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= total) return;

    int bar = bar_indices[tid];
    if (bar < 0 || bar + lookahead >= total_bars) {
        out_labels[tid] = -1;  // skip sentinel — no lookahead window available
        return;
    }

    float p_now = targets[bar * 6 + 2];
    float p_fut = targets[(bar + lookahead) * 6 + 2];

    // Strict greater-than tie-break: ties (p_fut == p_now, e.g. flat)
    // map to "down"=0 because the K=2 head can only distinguish "up"
    // from "not-up". Asymmetry is intentional and matches the brief.
    out_labels[tid] = (p_fut > p_now) ? 1 : 0;
}
