// Custom atomic min/max for float (must be defined BEFORE the kernel).
__device__ float atomicMin_float(float* addr, float val) {
    int* addr_as_int = (int*)addr;
    int old = *addr_as_int, assumed;
    do {
        assumed = old;
        old = atomicCAS(addr_as_int, assumed,
            __float_as_int(fminf(val, __int_as_float(assumed))));
    } while (assumed != old);
    return __int_as_float(old);
}

__device__ float atomicMax_float(float* addr, float val) {
    int* addr_as_int = (int*)addr;
    int old = *addr_as_int, assumed;
    do {
        assumed = old;
        old = atomicCAS(addr_as_int, assumed,
            __float_as_int(fmaxf(val, __int_as_float(assumed))));
    } while (assumed != old);
    return __int_as_float(old);
}

// Reduce per-experience rewards and actions into a compact summary.
// One block, parallel reduction across N elements.
extern "C" __global__ void monitoring_reduce(
    const float* __restrict__ rewards,   // [N]
    const int*   __restrict__ actions,   // [N]
    float*       summary,                // [16]: mean, std, min, max, sharpe, counts[9], total, _pad
    int          N
) {
    __shared__ float s_sum;
    __shared__ float s_sq_sum;
    __shared__ float s_min;
    __shared__ float s_max;
    __shared__ int   s_counts[9];

    int tid = threadIdx.x;
    int stride = blockDim.x;

    // Init shared memory
    if (tid == 0) {
        s_sum = 0.0f; s_sq_sum = 0.0f;
        s_min = 1e30f; s_max = -1e30f;
        for (int i = 0; i < 9; i++) s_counts[i] = 0;
    }
    __syncthreads();

    // Thread-local accumulators
    float local_sum = 0.0f, local_sq = 0.0f;
    float local_min = 1e30f, local_max = -1e30f;
    int local_counts[9] = {0, 0, 0, 0, 0, 0, 0, 0, 0};

    for (int i = tid; i < N; i += stride) {
        float r = rewards[i];
        local_sum += r;
        local_sq += r * r;
        local_min = fminf(local_min, r);
        local_max = fmaxf(local_max, r);
        /* Decode exposure index from factored action for monitoring.
         * Factored action = exposure * (b1*b2) + order * b2 + urgency. */
        int a = actions[i];
        int exp_idx = a / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
        if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS) local_counts[exp_idx]++;
    }

    // Warp-level reduction before atomic to shared (reduces contention 32x)
    for (int mask = 16; mask >= 1; mask >>= 1) {
        local_sum += __shfl_xor_sync(0xFFFFFFFF, local_sum, mask);
        local_sq  += __shfl_xor_sync(0xFFFFFFFF, local_sq, mask);
        local_min  = fminf(local_min, __shfl_xor_sync(0xFFFFFFFF, local_min, mask));
        local_max  = fmaxf(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, mask));
        for (int i = 0; i < 9; i++)
            local_counts[i] += __shfl_xor_sync(0xFFFFFFFF, local_counts[i], mask);
    }
    // Only lane 0 of each warp atomicAdd (8 warp leaders instead of 256 threads)
    if ((tid & 31) == 0) {
        atomicAdd(&s_sum, local_sum);
        atomicAdd(&s_sq_sum, local_sq);
        atomicMin_float(&s_min, local_min);
        atomicMax_float(&s_max, local_max);
        for (int i = 0; i < 9; i++) atomicAdd(&s_counts[i], local_counts[i]);
    }
    __syncthreads();

    // Thread 0 writes summary
    if (tid == 0) {
        float mean = s_sum / (float)N;
        float var = s_sq_sum / (float)N - mean * mean;
        float std = sqrtf(fmaxf(var, 0.0f));
        summary[0] = mean;
        summary[1] = std;
        summary[2] = s_min;
        summary[3] = s_max;
        summary[4] = (std > 1e-8f) ? mean / std : 0.0f;  // Sharpe estimate
        for (int i = 0; i < 9; i++) summary[5 + i] = (float)s_counts[i];
        summary[14] = (float)N;
        summary[15] = 0.0f;  // padding
    }
}
