// crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu
//
// SP5 Task A4: Pearl 4 — per-param-group Adam β1, β2, ε from gradient
// direction-stability + L2 norm.
// 8-thread single-block (one per param group: DqnTrunk/Value/Branches/IQN/
// IqlHigh/IqlLow/Attn/Curiosity).
//
// Producer signal: per-group cosine similarity of current vs previous-step
// gradient direction (written by grad_cosine_sim_kernel). High stability
// (cosine→1) → β2 → 0.9995 (long memory); low stability (cosine→0 or →-1)
// → β2 → 0.99 (short memory). β1 likewise scales with stability.
//
// Structural envelopes (Invariant 1 anchors per spec line 88):
//   β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6]
//
// THEORETICAL CAVEAT: adaptive β breaks Adam's constant-β convergence
// proof. Mitigations:
//   - Structural β envelopes (Invariant 1 anchors)
//   - ALPHA_META=5e-4 in apply_pearls_ad_kernel (half the SP4 default → slow change rate)
//   - Pearl A sentinel bootstrap on fold boundary (cosine_sim=0 → envelope midpoints)
// If Layer C destabilization observed, fall-back path: revert to ε-only adaptive.

#include <cuda_runtime.h>

extern "C" __global__ void pearl_4_adam_hparams_update(
    const float* __restrict__ grad_cosine_sim_per_group,  // [8] from grad_cosine_sim_kernel
    const float* __restrict__ grad_l2_norm_per_group,     // [8] from grad_cosine_sim_kernel
    float* __restrict__ scratch_out,
    int   beta1_idx_base,
    int   beta2_idx_base,
    int   eps_idx_base
) {
    int g = threadIdx.x;
    if (blockIdx.x != 0 || g >= 8) return;

    float stability = fmaxf(0.0f, fminf(1.0f, grad_cosine_sim_per_group[g]));
    float grad_norm = grad_l2_norm_per_group[g];

    // Structural envelopes (Invariant 1 anchors per spec line 88).
    // β1 scales linearly with stability: [0.85, 0.95]
    float beta1 = 0.85f + 0.10f * stability;
    // β2 scales linearly with stability: [0.99, 0.9995]
    float beta2 = 0.99f + 0.0095f * stability;
    // ε scales with grad norm to maintain numerical headroom; envelope clamp.
    // grad_norm * 1e-7 maps typical grad norms [1e-3, 10.0] to [1e-10, 1e-6].
    float eps = fmaxf(1e-10f, fminf(1e-6f, grad_norm * 1e-7f));

    scratch_out[beta1_idx_base + g] = beta1;
    scratch_out[beta2_idx_base + g] = beta2;
    scratch_out[eps_idx_base + g]   = eps;

    // Make writes globally visible before apply_pearls_ad_kernel consumes them.
    __threadfence_system();
}
