/**
 * Signal adapter CUDA kernels for converting model outputs to trading signals.
 *
 * Three fused kernels replacing Candle tensor operations:
 *
 * 1. ppo_to_exposure_scores_kernel:
 *    Aggregates PPO 45-action softmax probs into 5 exposure scores.
 *    Input: [batch * 45] float probs
 *    Output: [batch * 5] float scores (sum of 9 order/urgency combos per bucket)
 *
 * 2. signal_to_action_scores_kernel:
 *    Maps scalar predictions to [batch, 5] one-hot-ish action scores via thresholds.
 *    Input: [batch] float predictions
 *    Output: [batch * 5] float scores (one-hot based on threshold buckets)
 *
 * 3. tft_quantile_extract_kernel:
 *    Extracts median signal from TFT quantile predictions.
 *    Input: [batch * horizon * num_quantiles] float
 *    Output: [batch] float median values (quantile index 1, horizon index 0)
 *
 * None of these kernels require common_device_functions.cuh (standalone).
 * Launch config: grid=(ceil(batch/256), 1, 1), block=(256, 1, 1).
 * One thread per batch element.
 */

extern "C" __global__ void ppo_to_exposure_scores_kernel(
    const float* __restrict__ probs,   /* [batch * 45] softmax probs */
    float* __restrict__ out_scores,    /* [batch * 5]  exposure scores */
    int batch
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= batch) return;

    /* For each of 5 exposure buckets, sum over 9 order/urgency combos.
     * probs layout: [batch, 45] where 45 = 5 exposure * 9 (3 order * 3 urgency)
     * bucket e covers indices [e*9 .. e*9+8] within each batch row. */
    int base = b * 45;
    for (int e = 0; e < 5; ++e) {
        float sum = 0.0f;
        int offset = base + e * 9;
        for (int j = 0; j < 9; ++j) {
            sum += probs[offset + j];
        }
        out_scores[b * 5 + e] = sum;
    }
}

extern "C" __global__ void signal_to_action_scores_kernel(
    const float* __restrict__ predictions,  /* [batch] scalar signals */
    float* __restrict__ out_scores,         /* [batch * 5] one-hot scores */
    float high_threshold_bps,
    float low_threshold_bps,
    int batch
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= batch) return;

    float pred = predictions[b];
    int out_base = b * 5;

    /* Zero all 5 scores, then set the matching bucket to 1.0 */
    out_scores[out_base + 0] = 0.0f;
    out_scores[out_base + 1] = 0.0f;
    out_scores[out_base + 2] = 0.0f;
    out_scores[out_base + 3] = 0.0f;
    out_scores[out_base + 4] = 0.0f;

    if (pred < -high_threshold_bps) {
        out_scores[out_base + 0] = 1.0f;  /* Short100 */
    } else if (pred < -low_threshold_bps) {
        out_scores[out_base + 1] = 1.0f;  /* Short50 */
    } else if (pred > high_threshold_bps) {
        out_scores[out_base + 4] = 1.0f;  /* Long100 */
    } else if (pred > low_threshold_bps) {
        out_scores[out_base + 3] = 1.0f;  /* Long50 */
    } else {
        out_scores[out_base + 2] = 1.0f;  /* Flat */
    }
}

extern "C" __global__ void tft_quantile_extract_kernel(
    const float* __restrict__ quantiles,  /* [batch * horizon * num_q] */
    float* __restrict__ out_signal,       /* [batch] median values */
    int horizon,
    int num_quantiles,
    int batch
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= batch) return;

    /* Extract quantile index 1 (median) from horizon step 0.
     * Layout: quantiles[b * horizon * num_q + 0 * num_q + 1] */
    int idx = b * horizon * num_quantiles + 1;  /* horizon=0, quantile=1 */
    out_signal[b] = quantiles[idx];
}
