diff --git a/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu b/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu index 759d07afe..bacffc508 100644 --- a/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu +++ b/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu @@ -1,21 +1,25 @@ -// bce_loss_multi_horizon.cu — Kendall σ-weighted multi-horizon BCE. +// bce_loss_multi_horizon.cu — base-weight-only multi-horizon BCE. // -// Single source of truth for the multi-horizon BCE. Implements -// Kendall homoscedastic-uncertainty weighting per [Kendall+Gal+Cipolla -// 2018]: each horizon has a learnable `log_sigma_h` scalar that -// balances its loss contribution. Replaces the prior plain-pooled BCE. +// Single source of truth for the multi-horizon BCE. Loss is the +// base-weight-scaled per-horizon mean of unweighted BCE — no Kendall +// σ damping (2026-05-18). Per-horizon prioritization is done outside +// this kernel by the ISV-driven λ controller scaling the trunk +// grad_h contribution in heads_grn_bwd (per `pearl_adam_normalizes_loss_weights` +// — affect the gradient, not the loss aggregate). // // Math: // // raw_bce_h = Σ_{i in horizon h, m_i = 1} L_i (unweighted sum) // count_h = #{i in horizon h : m_i = 1} // mean_bce_h = raw_bce_h / count_h -// w_h = base_weight_h / (2 * exp(2 * log_sigma_h)) -// total_loss = Σ_h [ w_h * mean_bce_h + log_sigma_h ] +// w_h = base_weight_h (typically 1) +// total_loss = Σ_h [ w_h * mean_bce_h ] // // Gradients: // d L / d p_i = m_i * (w_h / count_h) * (p − y) / (p (1 − p)) -// d L / d log_sigma_h = 1 − 2 * w_h * mean_bce_h +// d L / d log_sigma_h = 0 (σ is a sidecar — emitted by +// horizon_ema_and_lambda for telemetry, +// not used in the loss path) // // Per `feedback_nvidia_grade_perf_for_kernels`: // - Warp-shuffle reduction (`__shfl_xor_sync`), NOT block tree-reduce. @@ -138,14 +142,21 @@ extern "C" __global__ void bce_multi_horizon_forward_backward( mean_bce_h_shared[h] = mean_bce; mean_bce_per_h_out[h] = mean_bce; - const float ls = log_sigma_h[h]; + // σ sidecar (2026-05-18): Kendall σ damped every horizon + // by ~30% (σ ≈ √mean_bce ≈ 0.84 → w_h ≈ 0.71) and biased + // the encoder away from h6000 — empirically -2pp mean_auc + // vs Phase 1+2+3 across multiple architectures. The + // Kendall framework's assumption ("damp the noisier + // task") inverts what we need ("learn the harder + // horizon harder"). σ is now emitted by + // horizon_ema_and_lambda for telemetry only and does + // NOT participate in the BCE forward weighting or + // backward gradient. `log_sigma_h` arg is preserved + // for kernel-signature stability. const float bw = (base_weights != nullptr) ? base_weights[h] : 1.0f; - // w_h = bw / (2 * exp(2*ls)) = 0.5 * bw * exp(-2 * ls) - const float w_h = 0.5f * bw * expf(-2.0f * ls); - // Per-element grad prefactor = w_h / count_h. + const float w_h = bw; w_eff_shared[h] = (c_h > 0.0f) ? w_h / c_h : 0.0f; - // ∂L/∂log_sigma_h = 1 − 2 * w_h * mean_bce_h. - d_log_sigma_h[h] = 1.0f - 2.0f * w_h * mean_bce; + d_log_sigma_h[h] = 0.0f; } } int v_valid = (lane < BCS_N_WARPS) ? s_valid[lane] : 0; @@ -158,14 +169,15 @@ extern "C" __global__ void bce_multi_horizon_forward_backward( __syncthreads(); // Total loss — small serial reduce by thread 0 over 5 horizons. + // σ-Kendall regularizer (log σ_h term) intentionally absent — see + // sidecar note above. Loss is unweighted per-horizon BCE sum, + // scaled only by `base_weights[h]` (typically uniform [1,1,1,1,1]). if (tid == 0) { float total_loss = 0.0f; #pragma unroll for (int h = 0; h < BCS_N_HORIZONS; ++h) { - const float ls = log_sigma_h[h]; const float bw = (base_weights != nullptr) ? base_weights[h] : 1.0f; - const float w_h = 0.5f * bw * expf(-2.0f * ls); - total_loss += w_h * mean_bce_h_shared[h] + ls; + total_loss += bw * mean_bce_h_shared[h]; } loss_out[0] = total_loss; }