// crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu
//
// SP15 Phase 1.5 — dd_pct foundational state input concat kernel.
//
// LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after this
// commit. Greenfield OK per spec Q1.
//
// Per spec §6.5: dd_pct (slot 406, written by Task 1.3's dd_state_kernel)
// gets concatenated to the trunk forward input as the last dim. Eval-time
// policy SEES drawdown context on every forward pass; Phase 3 teachings
// can condition on dd_pct directly via state, not just reward modulation.
//
// Kernel contract (Pearl: feedback_no_partial_refactor):
//   Input:  states_buf [B, state_dim_padded]  — existing trunk input
//           isv[DD_PCT_INDEX=406]             — dd_pct ∈ [0, 1] (broadcast)
//           B, state_dim_padded
//   Output: concat_buf [B, state_dim_padded + 1]  — pre-allocated
//           concat_buf[b, 0..state_dim_padded]   = states_buf[b, 0..state_dim_padded]
//           concat_buf[b, state_dim_padded]      = isv[DD_PCT_INDEX]
//
// The kernel produces the concatenated buffer; trunk consumer migration
// (re-pointing forward_online to consume concat_buf instead of states_buf
// and bumping s1_input_dim accordingly) lands as a follow-up atomic
// commit per the established Phase 1 precedent (Tasks 1.1-1.4 all defer
// consumer migration to follow-up commits per feedback_no_partial_refactor).
//
// Layout: row-major [B, state_dim_padded] for input; row-major
// [B, state_dim_padded + 1] for output. The +1 last column is the
// dd_pct broadcast slot. No atomicAdd (pure scatter-copy).

extern "C" __global__ void dd_pct_concat_kernel(
    const float* __restrict__ states_buf,
    const float* __restrict__ isv,
    float* __restrict__       concat_buf,
    int                       batch_size,
    int                       state_dim_padded
) {
    const int b      = blockIdx.x;
    const int tid    = threadIdx.x;
    const int n_tid  = blockDim.x;

    if (b >= batch_size) return;

    const float dd_pct = isv[406];  // DD_PCT_INDEX (compile-time constant)

    const int out_stride = state_dim_padded + 1;
    const int in_offset  = b * state_dim_padded;
    const int out_offset = b * out_stride;

    // 1. Copy the existing state vector — full padded width (including
    //    pre-existing pad zeros from the producer).
    for (int i = tid; i < state_dim_padded; i += n_tid) {
        concat_buf[out_offset + i] = states_buf[in_offset + i];
    }

    // 2. Last column = dd_pct broadcast.  Single thread writes the +1 slot.
    if (tid == 0) {
        concat_buf[out_offset + state_dim_padded] = dd_pct;
    }
}
