/**
 * Standalone Polyak EMA kernel for target network soft-update.
 *
 * target[i] = (1 - tau) * target[i] + tau * online[i]
 *
 * Launch config: grid=(ceil(n/256), 1, 1), block=(256, 1, 1).
 */

extern "C" __global__
void dqn_ema_kernel(float* __restrict__ target,
                    const float* __restrict__ online,
                    const float* __restrict__ tau_buf,
                    int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float tau = tau_buf[0];
        float t = target[i];
        float o = online[i];
        target[i] = (1.0f - tau) * t + tau * o;
    }
}
