gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

diag+fix(dqn): Task 2.X ISV-adaptive magnitude mechanism — reveals direction-branch is the real blocker

Browse Source 
Per feedback_adaptive_not_tuned.md: adaptive signal-driven mechanism, zero
static tuning knobs. Extends the existing ISV bus with per-magnitude Q-mean
EMAs and an absolute-scale reference; C51 loss + gradient kernels now read
ISV at zero hot-path cost to modulate per-bin weight in response to observed
collapse severity. Weight is 1.0 when Q is healthy; scales up per-bin when
collapse signal fires; self-disables as training stabilises.

Additions:
  * ISV_DIM 13 → 17. New slots:
      [13] Q_MAG_MEAN_QUARTER: ema(mean Q(Quarter), tau=0.05)
      [14] Q_MAG_MEAN_HALF:    ema(mean Q(Half),    tau=0.05)
      [15] Q_MAG_MEAN_FULL:    ema(mean Q(Full),    tau=0.05)
      [16] Q_ABS_REF:          ema(max(|Q_mean[k]|), tau=0.05) — scale-invariant reference
  * q_mag_bin_means_reduce kernel (q_stats_kernel.cu) — one-block reduce
    computing per-mag Q-means from q_out_buf; output written to pinned
    scratch slots; drives the EMAs in isv_signal_update.
  * c51_loss_kernel::get_magnitude_bin_weight helper + matching inlined
    logic in c51_grad_kernel: composite collapse signal = min(1,
    frac_bin + (1 - learning_health)); bin_weight = 1.0 + collapse *
    mag_bias_signal[k] (bounded in [1, 2]); mag_bias_signal[k] = (k+1)/b1_size.
  * isv_signal_update extended with q_mag_means_ptr + q_abs_ref_ptr +
    mag_size kernel args.

Diagnostics (keystone finding below):
  * gpu_backtest_evaluator::read_eval_action_distribution_per_direction —
    4-bin per-direction count at eval (Short/Hold/Long/Flat fractions).
    This diagnostic flipped the task diagnosis.
  * DQNTrainer::last_eval_direction_dist accessor.
  * last_isv_magnitude_bin_q_means accessor.
  * EVAL_DIR_DIST + ISV_BIN_MEANS debug prints in magnitude_distribution smoke.
  * ef >= 0.05 smoke gate added (currently unreachable behind pre-existing
    eh+ef >= 0.30 gate; kept for future use).

Training-time outcome:
  Pre-fix  MAG_DIST: Quarter=0.60 Half=0.10 Full=0.23
  Post-fix MAG_DIST: Quarter=0.46 Half=0.24 Full=0.28  (2.4× Half lift,
                                                       Full unchanged)
  Pre-fix  EVAL_DIST: eq=1.000 eh=0.000 ef=0.000
  Post-fix EVAL_DIST: eq=0.981 eh=0.019 ef=0.000

Root cause revealed (why the adaptive fix couldn't lift ef off 0):

  EVAL_DIR_DIST: Short=0.045 Hold=0.115 Long=0.070 Flat=0.771

  ~88% of eval states have direction ∈ {Hold, Flat}. Kernel at
  experience_kernels.cu:~896 FORCES mag_idx=0 (Quarter) in those cases
  as a structural ABI invariant. Only ~11.5% of eval samples have a
  free magnitude choice. Upper bound on ef regardless of magnitude
  mechanism: ~0.11.

  The magnitude branch mechanism works as designed — it correctly
  rebalances per-bin Q-means and lifts the training-time Half share
  2.4×. But direction-branch collapse to Flat masks everything
  downstream. Task 2.X's magnitude-only scope cannot unblock eval ef.

  The real fix target is direction-branch eval collapse. Follow-up
  task "Task 2.Y make direction-branch trade" extends the same
  ISV-driven composite-signal mechanism to branch 0 (Short/Long vs
  Hold/Flat). Scoping doc to be written.

Smoke validation:
  magnitude_distribution  FAIL (pre-existing eh+ef >= 0.30 gate; same
                          fail mode as HEAD before this commit)
  reward_component_audit  PASS
  controller_activity     PASS
  exploration_coverage    PASS
  multi_fold_convergence  PASS (avg best_val_metric=0.039, within ±15%)

No config fields. No static tuning knobs. No feature flags. All
modulation flows through the ISV bus. Shape constants documented:
eps=1e-6 (numerical guard), alpha=0.05 (ema tau matching existing ISV
pattern), MAX_MAG=4 (branch-size ceiling, already established),
mag_bias_signal[k]=(k+1)/b1_size (architectural monotonicity w.r.t.
bin index as stake size).
This commit is contained in:
jgrusewski
2026-04-22 20:12:20 +02:00
parent a9a51e8fa0
commit fa8d546614
9 changed files with 652 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files

49
crates/ml/src/cuda_pipeline/c51_grad_kernel.cu
View File

@@ -27,7 +27,12 @@ extern "C" __global__ void c51_grad_kernel(
const float* __restrict__ per_sample_support, /* [B, 3] per-sample [v_min, v_max, delta_z] */
const float* __restrict__ liquid_mod, /* [4] pinned device-mapped per-branch modulators */
const float* __restrict__ atom_positions, /* [4, num_atoms] adaptive positions. NULL = linear. */
const float* __restrict__ q_mean_ema_ptr) /* [1] pinned device-mapped Q-mean EMA */
const float* __restrict__ q_mean_ema_ptr, /* [1] pinned device-mapped Q-mean EMA */
/* Task 2.X adaptive magnitude fix — ISV signal bus [ISV_DIM=15]. Slots
* [13]=q_mag_spread_ema, [14]=q_dir_spread_ema drive an adaptive bin
* weight applied to the magnitude branch (d==1). NULL → identity
* (safe: reduces to pre-fix behaviour). */
const float* __restrict__ isv_signals)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int total_elems = batch_size * num_atoms;
@@ -67,6 +72,48 @@ extern "C" __global__ void c51_grad_kernel(
d_combined += inv_batch * entropy_coeff * (1.0f + lp_clamped);
}
/* Task 2.X "make Full useful": scale magnitude-branch gradient by
* the same signal-driven bin weight applied to branch_ce in
* c51_loss_batched. Consistency is required — the forward CE and
* the backward gradient must share the same scalar multiplier so
* the update step follows the true gradient of the (scaled) loss.
*
* Composite signal — see c51_loss_kernel.cu::get_magnitude_bin_weight
* for the full rationale. Component A reads ISV slots [13..16]
* (per-mag Q-mean EMAs + |Q| ref) and fires when the taken bin's
* Q-mean lags max. Component B reads ISV slot [12] (learning_health)
* and fires when training has not stabilised. collapse_frac =
* min(1, frac_bin + frac_health). Bounded bin_weight ∈ [1, 2]. */
if (d == 1 && isv_signals != NULL) {
int a1 = branch_actions[1];
if (b1_size > 0 && a1 >= 0 && a1 < b1_size && a1 < 3) {
/* Component A: bin-specific bias-vs-argmax */
float q_q = isv_signals[13];
float q_h = isv_signals[14];
float q_f = isv_signals[15];
float q_abs_ref = isv_signals[16];
float means[3] = { q_q, q_h, q_f };
float max_mean = means[0];
for (int k = 1; k < 3; k++)
max_mean = fmaxf(max_mean, means[k]);
float frac_bin = 0.0f;
if (q_abs_ref > 1e-6f) {
float bias_gap = fmaxf(0.0f, max_mean - means[a1]);
frac_bin = fminf(1.0f, bias_gap / fmaxf(q_abs_ref, 1e-6f));
}
/* Component B: meta-health stress signal */
float health = fminf(1.0f, fmaxf(0.0f, isv_signals[12]));
float frac_health = 1.0f - health;
/* Composite */
float collapse_frac = fminf(1.0f, frac_bin + frac_health);
if (collapse_frac > 0.0f) {
float mag_bias_signal = (float)(a1 + 1) / (float)b1_size;
float bin_weight = 1.0f + collapse_frac * mag_bias_signal;
d_combined *= bin_weight;
}
}
}
d_val_sum += d_combined;
float branch_scale = branch_scales[b * 4 + d];

113
crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
View File

@@ -199,6 +199,107 @@ __device__ __forceinline__ float get_learning_health(const float* __restrict__ i
return (isv_signals_ptr != NULL) ? isv_signals_ptr[12] : 0.5f;
}
/* Task 2.X "make Full useful" — signal-driven magnitude bin weight.
*
* Composite collapse signal — combines two ISV-driven components so the
* mechanism has actual leverage in ALL smoke regimes observed during
* Task-2.X iteration (earlier formulations dormant because the chosen
* reference always matched or exceeded the target):
*
* ── Component A: bin-specific bias-vs-argmax (from per-bin Q-means) ──
* ISV slots [13..15] = per-magnitude-bin Q-value mean EMAs
* [13] Q_mean_ema(Quarter), [14] Q_mean_ema(Half), [15] Q_mean_ema(Full)
* ISV slot [16] = absolute-Q-scale reference EMA (max |Q_mean| across bins).
*
* max_mean = max(Q_mean[k])
* bias_gap_bin = max(0, max_mean − Q_mean[a1])
* frac_bin = min(1, bias_gap_bin / max(Q_abs_ref, eps))
*
* `frac_bin` is positive iff the TAKEN bin's Q-mean is below the best
* bin's Q-mean (self-disables when the taken bin IS the argmax).
*
* ── Component B: meta-health stress signal (from ISV learning health) ──
* ISV slot [12] = learning_health ∈ [0, 1]. 1 = healthy, 0 = collapsed.
*
* frac_health = (1 − learning_health) ∈ [0, 1]
*
* This engages UNIFORMLY across all magnitude samples (not only the
* ones where Q(a1) lags max), providing leverage even when the per-bin
* component A is zero for that specific sample. Meta-signal-driven —
* fades to zero as the model's training dynamics stabilise. Same slot
* that already drives label-smoothing and Boltzmann temperature scaling
* in this kernel (see LABEL_SMOOTHING_HEALTH_INDEX); reusing it keeps
* the ISV surface coherent.
*
* ── Composite ──
* collapse_frac = min(1, frac_bin + frac_health)
* bin_weight = 1 + collapse_frac * mag_bias_signal(a1) ∈ [1, 2]
*
* mag_bias_signal[a1] = (a1 + 1) / b1_size ∈ {0.33, 0.67, 1.00} for
* Quarter/Half/Full respectively. Higher-index bins correspond to
* higher intrinsic PnL variance — exactly the dimension C51's
* expected-Q under-prices. Bin-proportional weighting concentrates
* the correction on the under-represented tail.
*
* Properties:
* * Bounded ∈ [1, 2] — keeps magnitude-branch gradient in the same
* ballpark as other loss terms' budgets (direction, order, urgency,
* CQL, IQN, ensemble). The fix is a correction, not a replacement.
* * Dual self-disabling:
* - Component A: when Full becomes the argmax bin, frac_bin(Full)
* drops to 0; the mechanism fades as the Q-landscape rebalances.
* - Component B: when learning_health → 1 (training stabilises),
* frac_health → 0; the mechanism fades as training converges.
* * No hard-coded tuning target — every term is an ISV-driven signal
* or an architectural property (bin count, bin index).
*
* Earlier Task-2.X iterations (dir-branch Q-spread reference; per-bin
* Q-mean bias alone) had smoke-observed collapse_frac=0 throughout
* training, so the mechanism never fired. The composite signal above
* uses an AND/OR ensemble of two independent collapse indicators —
* either fires when its specific collapse mode is active, so the
* mechanism reliably engages in practice.
*
* Returns 1.0 if ISV is unavailable (safe identity). `eps=1e-6f` is the
* standard numerical-guard epsilon matching existing patterns in this
* file (see block_bellman_project_f, barrier_gradient_direction, etc.).
*/
__device__ __forceinline__ float get_magnitude_bin_weight(
const float* __restrict__ isv_signals_ptr,
int a1,
int b1_size
) {
if (isv_signals_ptr == NULL) return 1.0f;
if (b1_size <= 0 || a1 < 0 || a1 >= b1_size || a1 >= 3) return 1.0f;
/* Component A: bin-specific bias-vs-argmax */
float q_q = isv_signals_ptr[13];
float q_h = isv_signals_ptr[14];
float q_f = isv_signals_ptr[15];
float q_abs_ref = isv_signals_ptr[16];
float mean_by_bin[3] = { q_q, q_h, q_f };
float max_mean = mean_by_bin[0];
for (int k = 1; k < 3; k++)
max_mean = fmaxf(max_mean, mean_by_bin[k]);
float frac_bin = 0.0f;
if (q_abs_ref > 1e-6f) {
float bias_gap = fmaxf(0.0f, max_mean - mean_by_bin[a1]);
frac_bin = fminf(1.0f, bias_gap / fmaxf(q_abs_ref, 1e-6f));
}
/* Component B: meta-health stress signal */
float health = fminf(1.0f, fmaxf(0.0f, isv_signals_ptr[12]));
float frac_health = 1.0f - health;
/* Composite — bounded ∈ [0, 1]. Reaches 1 when either component
* saturates OR both components are partially active and sum ≥ 1. */
float collapse_frac = fminf(1.0f, frac_bin + frac_health);
if (collapse_frac <= 0.0f) return 1.0f;
float mag_bias_signal = (float)(a1 + 1) / (float)b1_size;
return 1.0f + collapse_frac * mag_bias_signal;
}
/* ══════════════════════════════════════════════════════════════════════
* MAIN KERNEL: c51_loss_batched (float arithmetic, BF16 I/O)
* ══════════════════════════════════════════════════════════════════════ */
@@ -713,6 +814,18 @@ extern "C" __global__ void c51_loss_batched(
for (int j = tid; j < num_atoms; j += BLOCK_THREADS)
local_ce -= shmem_lp[j] * shmem_current_lp[j];
float branch_ce = block_reduce_sum_f(local_ce, shmem_reduce, tid);
/* Task 2.X adaptive magnitude fix: scale magnitude-branch CE by a
* signal-driven bin weight derived from the ISV bus. Identity
* (bin_weight=1.0) when ISV reports magnitude Q-spread ≥ direction
* Q-spread — the mechanism self-disables once the magnitude head's
* Q-values differentiate. See get_magnitude_bin_weight. */
if (d == 1) {
float mag_bin_weight = get_magnitude_bin_weight(
isv_signals, branch_action[1], branch_sizes[1]);
branch_ce *= mag_bin_weight;
}
total_ce += branch_ce;
/* ═══ STEP f: Online-target Q-divergence for adaptive tau ══

50
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -5536,8 +5536,8 @@ extern "C" __global__ void risk_budget_backward(
/* ================================================================== */
extern "C" __global__ void isv_signal_update(
float* __restrict__ isv_signals, /* [12] pinned device-mapped */
float* __restrict__ isv_history, /* [K*12] pinned device-mapped */
float* __restrict__ isv_signals, /* [ISV_DIM=17] pinned device-mapped */
float* __restrict__ isv_history, /* [K*12] pinned — history only rotates slots [0..11] */
float* __restrict__ lagged_td_error, /* [1] pinned — for recursive confidence target */
const float* __restrict__ q_mean_ptr, /* [1] pinned — batch Q-mean */
const float* __restrict__ q_mean_ema_ptr, /* [1] pinned — Q-mean EMA */
@@ -5549,7 +5549,13 @@ extern "C" __global__ void isv_signal_update(
const float* __restrict__ loss_ptr, /* [1] pinned — total loss */
int K, /* temporal history length (4) */
const float* __restrict__ states_ptr, /* [B, state_dim] — for ADX/CUSUM regime signals. NULL = skip regime. */
int state_dim /* to index ADX at [40], CUSUM at [41] */
int state_dim, /* to index ADX at [40], CUSUM at [41] */
/* Task 2.X "make Full useful" — per-magnitude-bin Q-mean array +
* absolute-scale reference. Produced by q_mag_bin_means_reduce on
* the same stream before isv_signal_update. NULL = freeze slots. */
const float* __restrict__ q_mag_means_ptr, /* [mag_size] per-bin Q mean */
const float* __restrict__ q_abs_ref_ptr, /* [1] max(|Q_mean|) across mag bins */
int mag_size /* typically 3 (Quarter/Half/Full) */
) {
if (threadIdx.x != 0) return;
@@ -5624,6 +5630,44 @@ extern "C" __global__ void isv_signal_update(
isv_signals[10] = 0.0f;
isv_signals[11] = 0.5f; /* moderate stability */
}
/* ── Task 2.X "make Full useful": per-magnitude-bin Q-mean EMAs [13..15]
* and absolute-Q-scale reference EMA [16] ───────────────────────
*
* Signal-driven adaptive inputs for the C51 magnitude bin weight.
* [13] Q_mean_ema(Quarter) — persistent preference per bin.
* [14] Q_mean_ema(Half)
* [15] Q_mean_ema(Full)
* [16] |Q|-scale reference EMA (max over bins of |Q_mean|).
*
* Consumed by c51_loss_batched / c51_grad_kernel to derive
* max_mean = max(Q_mean(Quarter), Q_mean(Half), Q_mean(Full))
* bias_gap = max(0, max_mean − Q_mean(a1))
* collapse_frac = min(1, bias_gap / max(|Q|-scale, eps))
* bin_weight = 1 + collapse_frac * mag_bias_signal(a1)
*
* Directly targets C51's "high-variance bins under-priced" signature
* (experience_kernels.cu:~904). Self-disables when Full is the argmax
* — bias_gap becomes 0 for Full samples, bin_weight → 1.
*
* EMA tau=0.05 matches the existing alpha=0.05 pattern (20-step window).
* NULL-guarded — missing inputs leave the slots unchanged (no regression
* against callers that do not supply the pointers).
*
* NOT rotated through isv_history — already EMA-smoothed. isv_forward
* reads only the first 12 slots (the encoder/gate path is unchanged). */
if (q_mag_means_ptr != 0) {
int ub = (mag_size < 3) ? mag_size : 3; /* clamp to the 3 reserved slots */
for (int k = 0; k < ub; k++) {
int slot = 13 + k;
isv_signals[slot] = (1.0f - alpha) * isv_signals[slot]
+ alpha * q_mag_means_ptr[k];
}
}
if (q_abs_ref_ptr != 0) {
isv_signals[16] = (1.0f - alpha) * isv_signals[16]
+ alpha * fmaxf(q_abs_ref_ptr[0], 0.0f);
}
}
/* ================================================================== */

42
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -894,6 +894,48 @@ impl GpuBacktestEvaluator {
])
}
/// Task 2.X "make Full useful" diagnostic — read back the per-direction
/// action distribution from the most recent evaluation. Action encoding
/// is `dir*27 + mag*9 + ord*3 + urg`, so direction = `action / 27`
/// (0=Short, 1=Hold, 2=Long, 3=Flat).
///
/// Returns `[short, hold, long, flat]` summing to 1.0 over non-negative
/// action entries. When eval is strict-argmax and direction Q-values are
/// near-uniform, this distribution helps diagnose whether the eval
/// collapse is direction-side (e.g., `hold + flat > 0.95`) or whether
/// direction is healthy but magnitude itself is collapsing. Paired with
/// `read_eval_action_distribution_per_magnitude` to localise the
/// magnitude-only Task 2.X fix's leverage surface.
pub fn read_eval_action_distribution_per_direction(&self) -> Result<[f32; 4], MLError> {
let len = self.n_windows * self.max_len;
let mut host = vec![0_i32; len];
self.stream
.memcpy_dtoh(&self.actions_history_buf, &mut host)
.map_err(|e| MLError::ModelError(format!("actions_history dtoh: {e}")))?;
let mut counts = [0_u64; 4];
let mut total = 0_u64;
for &a in &host {
if a < 0 {
continue;
}
let dir = (a / 27) as usize;
if dir < 4 {
counts[dir] += 1;
total += 1;
}
}
if total == 0 {
return Ok([0.0; 4]);
}
let t = total as f32;
Ok([
counts[0] as f32 / t,
counts[1] as f32 / t,
counts[2] as f32 / t,
counts[3] as f32 / t,
])
}
/// Reset mutable evaluation state: portfolio, done flags, step returns, actions history.
/// Must be called before each `evaluate_dqn_graphed` when reusing the evaluator
/// across epochs — otherwise done_flags=1 from the previous evaluation causes

226
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -84,8 +84,44 @@ const OFI_EMBED_DIM: usize = 10; // OFI embedding width appended to history
/// Introspective State Vector (ISV) configuration.
const ISV_K: usize = 4; // Temporal ISV history length
const ISV_DIM: usize = 13; // ISV signal count (8 core + 4 regime awareness + 1 learning_health)
/// ISV signal count.
///
/// Layout:
/// [0..7] core training dynamics: q_drift, grad_norm_ema, td_err_ema,
/// ens_var_ema, ens_var_vel, reward_ema, atom_util_ema, loss_ema
/// [8..11] regime awareness (ADX/CUSUM): adx_ema, regime_disagree,
/// regime_vel_ema, regime_stability
/// [12] learning_health (written by separate kernel, consumed by c51/etc.)
/// [13..15] Task 2.X "make Full useful" — per-magnitude-bin Q-value EMAs.
/// [13] = Q_mean_ema(Quarter), [14] = Q_mean_ema(Half),
/// [15] = Q_mean_ema(Full). Populated by `q_mag_bin_means_reduce`
/// on the stats cadence, smoothed with tau=0.05. The C51 loss and
/// gradient kernels derive an adaptive bin weight from these
/// EMAs by computing the bias-vs-Full signature
/// collapse = max(0, max(Q_mean) − Q_mean(a1)) / max(|Q_ref|, eps)
/// applied ONLY to magnitude-branch samples (d==1). Self-disables
/// when Full becomes the argmax bin.
/// [16] q_abs_ref_ema — max(|Q_mean_ema[k]|) across the 3 magnitude
/// bins. Scale reference for the collapse-fraction normaliser
/// (keeps the bin-weight formula scale-invariant as Q-values
/// grow over training).
///
/// Slots [0..11] populated by `isv_signal_update` kernel and rotated through
/// `isv_history` for temporal decay in `isv_forward`. Slot [12] written
/// outside `isv_signal_update` and NOT rotated into history. Slots [13..16]
/// written by `isv_signal_update` (added with the Task 2.X adaptive magnitude
/// fix) but NOT rotated into history — they are already EMA-smoothed.
const ISV_DIM: usize = 17;
pub const LEARNING_HEALTH_INDEX: usize = 12;
/// Task 2.X "make Full useful": ISV slots for per-magnitude-bin Q-mean EMAs.
/// Read by `c51_loss_batched` and `c51_grad_kernel` to compute the
/// bias-vs-Full signal that drives adaptive bin weighting.
pub const Q_MAG_MEAN_QUARTER_INDEX: usize = 13;
pub const Q_MAG_MEAN_HALF_INDEX: usize = 14;
pub const Q_MAG_MEAN_FULL_INDEX: usize = 15;
/// Task 2.X "make Full useful": absolute-Q-scale reference used to
/// normalise the bias signal into a scale-invariant collapse fraction.
pub const Q_ABS_REF_INDEX: usize = 16;
const ISV_EMB_DIM: usize = 8; // ISV embedding output dimension (FC2 output, gate/gamma input)
/// First ISV tensor index in the flat param buffer.
/// ISV weights (68-79) are online-only — NOT synced to the target network.
@@ -1652,12 +1688,23 @@ pub struct GpuDqnTrainer {
atom_util_scratch_pinned: *mut f32,
atom_util_scratch_dev_ptr: u64,
// Task 2.X "make Full useful" — per-magnitude-bin Q-mean scratch arrays
// written by `q_mag_bin_means_reduce` (runs with q_stats on the stats
// cadence), read by isv_signal_update to drive the EMA slots [13..16].
// The means scratch is sized for magnitude-branch capacity (MAX_MAG=4
// in the kernel); the abs_ref scratch is a single scalar.
q_mag_means_scratch_pinned: *mut f32,
q_mag_means_scratch_dev_ptr: u64,
q_abs_ref_scratch_pinned: *mut f32,
q_abs_ref_scratch_dev_ptr: u64,
q_mag_bin_means_reduce_kernel: CudaFunction,
// ── ISV core buffers (pinned device-mapped, GPU read/write) ──
isv_signals_pinned: *mut f32, // [ISV_DIM = 13]
isv_signals_pinned: *mut f32, // [ISV_DIM = 15]
isv_signals_dev_ptr: u64,
isv_history_pinned: *mut f32, // [ISV_K * ISV_DIM = 52]
isv_history_pinned: *mut f32, // [ISV_K * 12 = 48] — history rotates slots [0..11] only
isv_history_dev_ptr: u64,
isv_decay_pinned: *mut f32, // [ISV_DIM = 13] learned decay weights
isv_decay_pinned: *mut f32, // [ISV_DIM = 15] learned decay weights
isv_decay_dev_ptr: u64,
lagged_td_error_pinned: *mut f32, // [1] recursive confidence target
lagged_td_error_dev_ptr: u64,
@@ -2142,6 +2189,12 @@ impl Drop for GpuDqnTrainer {
if !self.atom_util_scratch_pinned.is_null() {
let _ = unsafe { cudarc::driver::result::free_host(self.atom_util_scratch_pinned.cast()) };
}
if !self.q_mag_means_scratch_pinned.is_null() {
let _ = unsafe { cudarc::driver::result::free_host(self.q_mag_means_scratch_pinned.cast()) };
}
if !self.q_abs_ref_scratch_pinned.is_null() {
let _ = unsafe { cudarc::driver::result::free_host(self.q_abs_ref_scratch_pinned.cast()) };
}
if !self.isv_signals_pinned.is_null() {
let _ = unsafe { cudarc::driver::result::free_host(self.isv_signals_pinned.cast()) };
}
@@ -3014,12 +3067,64 @@ impl GpuDqnTrainer {
unsafe { *self.q_mean_ema_pinned = new_ema; }
}
/// Pure GPU ISV signal update — reads all pinned pointers, updates ISV [12] + history [K,12].
/// Includes 4 regime awareness signals computed from ADX/CUSUM in states buffer.
/// Task 2.X "make Full useful" — launch the per-magnitude-bin Q-mean
/// reducer over the current `q_out_buf`. Writes:
/// * `q_mag_means_scratch[0..mag_size]` = mean Q over batch for each
/// magnitude action (Quarter=0, Half=1, Full=2).
/// * `q_abs_ref_scratch[0]` = max(|mean Q|) across magnitude bins —
/// scale reference for the collapse-fraction normaliser.
///
/// These feed ISV slots [13..16] (EMA) via `update_isv_signals`. The
/// C51 loss / grad kernels then compute a bias-vs-Full signature
/// `max(0, max(Q_mean) − Q_mean(a1)) / Q_abs_ref` — directly targets
/// the "higher-variance bins systematically under-priced" bias cited
/// at experience_kernels.cu:~904 without requiring an inter-branch
/// reference (which failed during Task 2.X iteration 1 because the
/// direction branch collapses at the same cadence as magnitude).
///
/// Runs outside CUDA Graph capture on the stats cadence. Single-thread
/// single-block, ~N * mag_size f32 loads — microseconds at typical B.
pub(crate) fn launch_q_mag_bin_means_reduce(
&self,
batch_size: usize,
) -> Result<(), MLError> {
let total_actions = self.total_actions() as i32;
let n = batch_size as i32;
let mag_off = self.config.branch_0_size as i32;
let mag_size = self.config.branch_1_size as i32;
let q_out_buf_ptr = self.q_out_buf.raw_ptr();
unsafe {
self.stream
.launch_builder(&self.q_mag_bin_means_reduce_kernel)
.arg(&q_out_buf_ptr)
.arg(&self.q_mag_means_scratch_dev_ptr)
.arg(&self.q_abs_ref_scratch_dev_ptr)
.arg(&n)
.arg(&total_actions)
.arg(&mag_off)
.arg(&mag_size)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("q_mag_bin_means_reduce: {e}")))?;
}
Ok(())
}
/// Pure GPU ISV signal update — reads all pinned pointers, updates ISV
/// [ISV_DIM=17] + history [K,12]. Includes 4 regime awareness signals
/// computed from ADX/CUSUM in the states buffer. Slots [13..16] are the
/// Task 2.X "make Full useful" fix — per-magnitude-bin Q-mean EMAs and
/// absolute-scale reference, fed by `q_mag_bin_means_reduce` (launched
/// inside `reduce_current_q_stats` just before this kernel).
///
/// Call after graph_adam replay and stream sync, before next graph_forward.
pub(crate) fn update_isv_signals(&self) -> Result<(), MLError> {
let k = ISV_K as i32;
let state_dim = ml_core::state_layout::STATE_DIM as i32;
let mag_size_i32 = self.config.branch_1_size as i32;
unsafe {
self.stream.launch_builder(&self.isv_signal_update_kernel)
.arg(&self.isv_signals_dev_ptr)
@@ -3036,6 +3141,12 @@ impl GpuDqnTrainer {
.arg(&k)
.arg(&self.ptrs.states_buf) // states for ADX/CUSUM regime signals
.arg(&state_dim)
// Task 2.X "make Full useful" — per-magnitude-bin Q-mean
// scratch array + absolute-scale reference scalar feed ISV
// slots [13..15] (Q_mean per bin) and [16] (|Q| reference).
.arg(&self.q_mag_means_scratch_dev_ptr)
.arg(&self.q_abs_ref_scratch_dev_ptr)
.arg(&mag_size_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
@@ -6389,6 +6500,9 @@ impl GpuDqnTrainer {
// ── Compile expected Q-value + stats kernels (validation, not in CUDA Graph) ─
let expected_q_kernel = compile_expected_q_kernel(&stream)?;
let q_stats_kernel = compile_q_stats_kernel(&stream)?;
// Task 2.X "make Full useful" — per-magnitude-bin Q-mean reducer that
// feeds ISV slots [13..16] on the stats cadence.
let q_mag_bin_means_reduce_kernel = compile_q_mag_bin_means_kernel(&stream)?;
let q_stats_buf = stream.alloc_zeros::<f32>(7)
.map_err(|e| MLError::ModelError(format!("alloc q_stats_f32: {e}")))?;
let atom_stats_buf = stream.alloc_zeros::<f32>(2)
@@ -7273,6 +7387,53 @@ impl GpuDqnTrainer {
(host_ptr as *mut f32, dev_ptr_out)
};
// Task 2.X "make Full useful" — per-magnitude-bin Q-mean pinned
// array (sized for up to 4 magnitude bins to match
// MAX_MAGNITUDE_ACTIONS in experience_kernels.cu) plus the
// absolute-Q-scale reference scalar. Written by
// `q_mag_bin_means_reduce` (launched with q_stats); read by
// isv_signal_update to populate ISV slots [13..16].
let q_mag_means_slots: usize = 4;
let (q_mag_means_scratch_pinned, q_mag_means_scratch_dev_ptr) = {
let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut();
let mut dev_ptr_out: u64 = 0;
unsafe {
let rc = cudarc::driver::sys::cuMemAllocHost_v2(
&mut host_ptr,
q_mag_means_slots * std::mem::size_of::<f32>(),
);
assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for q_mag_means_scratch");
let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(
&mut dev_ptr_out,
host_ptr,
0,
);
assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for q_mag_means_scratch");
std::ptr::write_bytes(host_ptr as *mut u8, 0, q_mag_means_slots * std::mem::size_of::<f32>());
}
(host_ptr as *mut f32, dev_ptr_out)
};
let (q_abs_ref_scratch_pinned, q_abs_ref_scratch_dev_ptr) = {
let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut();
let mut dev_ptr_out: u64 = 0;
unsafe {
let rc = cudarc::driver::sys::cuMemAllocHost_v2(
&mut host_ptr,
std::mem::size_of::<f32>(),
);
assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for q_abs_ref_scratch");
let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(
&mut dev_ptr_out,
host_ptr,
0,
);
assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for q_abs_ref_scratch");
*(host_ptr as *mut f32) = 0.0;
}
(host_ptr as *mut f32, dev_ptr_out)
};
// ── ISV core buffers (pinned device-mapped) ──────────────────────
let (isv_signals_pinned, isv_signals_dev_ptr) = {
let num_bytes = ISV_DIM * std::mem::size_of::<f32>();
@@ -8176,6 +8337,11 @@ impl GpuDqnTrainer {
bw_d_glu_gate,
expected_q_kernel,
q_stats_kernel,
q_mag_bin_means_reduce_kernel,
q_mag_means_scratch_pinned,
q_mag_means_scratch_dev_ptr,
q_abs_ref_scratch_pinned,
q_abs_ref_scratch_dev_ptr,
q_stats_buf,
atom_stats_buf,
atom_stats_block_sums_buf,
@@ -8630,6 +8796,27 @@ impl GpuDqnTrainer {
}
}
/// Task 2.X "make Full useful" diagnostic — read the per-magnitude-bin
/// Q-mean EMAs and absolute-scale reference from ISV (pinned, zero-copy).
///
/// Returns `(q_mean_quarter, q_mean_half, q_mean_full, q_abs_ref)`. The
/// C51 loss / gradient kernels use these to compute the bin weight:
/// max_mean = max(q_mean_*)
/// bias_gap = max(0, max_mean − q_mean[a1])
/// collapse_frac = min(1, bias_gap / max(q_abs_ref, 1e-6))
/// bin_weight = 1 + collapse_frac * (a1 + 1) / b1_size
/// Returns all zeros if ISV is not allocated.
pub fn read_isv_magnitude_bin_q_means(&self) -> (f32, f32, f32, f32) {
if self.isv_signals_pinned.is_null() { return (0.0, 0.0, 0.0, 0.0); }
unsafe {
let q = *self.isv_signals_pinned.add(Q_MAG_MEAN_QUARTER_INDEX);
let h = *self.isv_signals_pinned.add(Q_MAG_MEAN_HALF_INDEX);
let f = *self.isv_signals_pinned.add(Q_MAG_MEAN_FULL_INDEX);
let r = *self.isv_signals_pinned.add(Q_ABS_REF_INDEX);
(q, h, f, r.max(0.0))
}
}
/// B4/G5: Read ISV health index and regime stability from pinned host memory.
/// Returns (health [0,1], regime_stability [0,1]). Falls back to (0.5, 0.5) if
/// pinned pointer is null.
@@ -10368,6 +10555,13 @@ impl GpuDqnTrainer {
// ISV: write Q-mean as reward proxy to pinned scratch
unsafe { *self.reward_scratch_pinned = host[3]; }
// Task 2.X "make Full useful" — per-magnitude-bin Q-mean reducer
// writes the 3 per-bin Q-means + absolute-scale reference to
// pinned scratch. Runs on the same stream; isv_signal_update
// reads the values directly via pinned device pointers (no sync).
let mag_means_batch_size = self.config.batch_size;
self.launch_q_mag_bin_means_reduce(mag_means_batch_size)?;
// ISV signal update — pure GPU, reads all pinned scalars, updates ISV vector
self.update_isv_signals()?;
@@ -12262,6 +12456,10 @@ impl GpuDqnTrainer {
.arg(&liquid_mod_buf_ptr)
.arg(&atom_positions_buf_ptr)
.arg(&self.q_mean_ema_dev_ptr) // unused but preserves 19-param graph node structure on Hopper
// Task 2.X adaptive magnitude fix — ISV signal bus; kernel reads
// slots [13] (q_mag_spread_ema) and [14] (q_dir_spread_ema) to
// derive the magnitude-branch adaptive bin weight.
.arg(&self.isv_signals_dev_ptr)
.launch(LaunchConfig {
grid_dim: (blocks, 1, 1),
block_dim: (256, 1, 1),
@@ -14039,6 +14237,22 @@ fn compile_q_stats_kernel(
.map_err(|e| MLError::ModelError(format!("q_stats_reduce load: {e}")))
}
/// Load the per-magnitude-bin Q-mean reducer from the q_stats cubin.
///
/// Task 2.X "make Full useful": this kernel computes batch-mean Q-values
/// per magnitude-bin (Quarter/Half/Full) plus an absolute-scale reference,
/// feeding the four ISV slots [13..16] that drive the adaptive C51
/// magnitude bin weight. See q_stats_kernel.cu for the kernel body.
fn compile_q_mag_bin_means_kernel(
stream: &Arc<CudaStream>,
) -> Result<CudaFunction, MLError> {
let context = stream.context();
let module = context.load_cubin(Q_STATS_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("q_mag_bin_means cubin load: {e}")))?;
module.load_function("q_mag_bin_means_reduce")
.map_err(|e| MLError::ModelError(format!("q_mag_bin_means_reduce load: {e}")))
}
// ── Shared memory sizing ────────────────────────────────────────────────────
/// Query the hardware's max shared memory per block via cuDeviceGetAttribute.

94
crates/ml/src/cuda_pipeline/q_stats_kernel.cu
View File

@@ -60,3 +60,97 @@ extern "C" __global__ void q_stats_reduce(
out[5] = (atom_stats != NULL) ? atom_stats[0] * inv_n / max_entropy : 0.0f; /* normalized entropy [0,1] */
out[6] = (atom_stats != NULL) ? atom_stats[1] * inv_n / (float)num_atoms : 0.0f; /* utilization fraction [0,1] */
}
/**
* Per-branch magnitude-Q mean reducer: computes batch-mean Q-values for each
* of the 3 magnitude bins separately (Quarter=0, Half=1, Full=2), plus the
* absolute Q-scale reference, from `q_values` of shape [N, total_actions]
* sliced at `mag_off..mag_off+mag_size`.
*
* Task 2.X "make Full useful" adaptive fix: feeds the ISV signal bus with
* the four scalars needed to compute a DIRECT anti-bias signal for C51's
* magnitude-branch gradient:
*
* bias_vs_full = max(0, max_mag(q_mean[k]) − q_mean[Full])
*
* When a higher-variance bin (Full=2) is priced below the max magnitude
* bin, `bias_vs_full > 0` — the exact signature of C51's structural
* variance bias (high-variance bins get systematically lower E[Q] because
* their support centering favours narrower distributions). When Full is
* already the argmax, `bias_vs_full == 0` and the mechanism self-disables.
*
* The reference scale `q_abs_ref` is the max of per-bin |Q_mean|, used to
* normalise `bias_vs_full` into a scale-invariant [0, 1] collapse fraction.
*
* Runs OUTSIDE the CUDA Graph (alongside q_stats_reduce). isv_signal_update
* folds the scalars into EMA slots [13] (Q(Quarter)), [14] (Q(Half)). No
* separate ISV slot for Q(Full) or the absolute scale — those are recovered
* inside the C51 kernels by reading per-sample q_values directly (the loss
* kernel already has access to the full magnitude-branch slice through its
* advantage inputs). The two EMAs we expose are the ones the bias detector
* NEEDS against the current batch's Q(Full): slow-moving references that
* reflect the policy's persistent preference rather than per-step noise.
*
* (Naming history: this kernel was originally scoped as
* `q_branch_spread_reduce` producing per-branch max-min spreads, but using
* direction-branch spread as the reference for magnitude collapse failed —
* observed `dir_spread=0.002 ≪ mag_spread=0.08` because the direction head
* itself collapses at the same cadence the magnitude head does, so the
* ratio-based formulation never triggered. The per-bin Q-mean
* formulation below targets the precise bias signature cited in the Task
* 2.X scoping doc and at experience_kernels.cu:~904 — high-variance bins
* under-priced by C51's expected-Q — without requiring an inter-branch
* reference.)
*
* Pure reduction — no training-path feedback and no atomicAdd. Single-block,
* single-thread to match q_stats_reduce's determinism.
*
* Output: q_mag_means_out[0] = Q_mean(Quarter) over N samples,
* q_mag_means_out[1] = Q_mean(Half),
* q_mag_means_out[2] = Q_mean(Full).
* q_abs_ref_out[0] = max(|Q_mean[k]|) across magnitude bins
* (absolute scale reference; floored at 1e-6
* downstream).
*
* Launch: grid=(1,1,1), block=(1,1,1). N=batch_size usually 256..16384 —
* stays well under any block timeout. Cost: ~N * mag_size FP32 loads,
* single-digit microseconds at typical B.
*/
extern "C" __global__ void q_mag_bin_means_reduce(
const float* __restrict__ q_values, /* [N, total_actions] */
float* __restrict__ q_mag_means_out, /* [mag_size] — typically [3] */
float* __restrict__ q_abs_ref_out, /* [1] absolute-value max of bin means */
int N,
int total_actions,
int mag_off, int mag_size)
{
if (threadIdx.x != 0) return;
/* Hard ceiling — covers any future magnitude-branch size extension
* without an unbounded stack array. Matches MAX_MAGNITUDE_ACTIONS=4
* used in experience_kernels.cu at the Boltzmann sampling site. */
const int MAX_MAG = 4;
float sums[MAX_MAG];
for (int k = 0; k < MAX_MAG; k++) sums[k] = 0.0f;
if (N <= 0 || total_actions <= 0 || mag_size <= 0 || mag_size > MAX_MAG
|| mag_off < 0 || mag_off + mag_size > total_actions) {
for (int k = 0; k < mag_size && k < MAX_MAG; k++) q_mag_means_out[k] = 0.0f;
q_abs_ref_out[0] = 0.0f;
return;
}
for (int i = 0; i < N; i++) {
const float* row = q_values + (long long)i * total_actions;
for (int k = 0; k < mag_size; k++) sums[k] += row[mag_off + k];
}
float inv_n = 1.0f / (float)N;
float abs_ref = 0.0f;
for (int k = 0; k < mag_size; k++) {
float mean_k = sums[k] * inv_n;
q_mag_means_out[k] = mean_k;
abs_ref = fmaxf(abs_ref, fabsf(mean_k));
}
q_abs_ref_out[0] = abs_ref;
}

10
crates/ml/src/trainers/dqn/fused_training.rs
View File

@@ -2607,6 +2607,16 @@ impl FusedTrainingCtx {
.map_err(|e| anyhow::anyhow!("launch_pad_states: {e}"))
}
/// Task 2.X diagnostic — read-only borrow of the inner `GpuDqnTrainer`.
/// Used by `DQNTrainer::last_isv_branch_q_spreads` to surface ISV
/// slots [13]/[14] to smoke tests without widening the field's
/// visibility. Do NOT add write paths through this accessor — the
/// trainer owns its CUDA context and mutation must go through
/// properly ordered submission paths.
pub(crate) fn trainer(&self) -> &GpuDqnTrainer {
&self.trainer
}
/// Raw pointer to trainer's states buffer.
pub(crate) fn trainer_states_buf_ptr(&self) -> u64 {
self.trainer.states_buf_ptr()

56
crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs
View File

@@ -65,6 +65,37 @@ fn test_magnitude_distribution() -> Result<()> {
// feedback_no_functionality_removal.md, the branch is never deleted as a fallback.
let [eq, eh, ef] = trainer.last_eval_magnitude_dist();
println!("[EVAL_DIST] Quarter={:.3} Half={:.3} Full={:.3}", eq, eh, ef);
// Task 2.X diagnostic — eval direction distribution. When eval collapses
// to Quarter==1 AND Hold+Flat dominate, the root cause is direction-side
// (magnitude gets forced to 0 by kernel for Hold/Flat), so a magnitude-
// only bin weight cannot shift eval_dist. Surfacing this lets the
// operator attribute the failure to direction vs magnitude.
if let Some(dir_dist) = trainer.last_eval_direction_dist() {
println!(
"[EVAL_DIR_DIST] Short={:.3} Hold={:.3} Long={:.3} Flat={:.3}",
dir_dist[0], dir_dist[1], dir_dist[2], dir_dist[3]
);
}
// Task 2.X diagnostic — ISV per-magnitude-bin Q-mean EMAs + |Q|-scale
// reference. The bin weight fires when Full is NOT the argmax bin of
// (Quarter / Half / Full), so the training-policy's Q-means directly
// drive the adaptive correction.
if let Some((qq, qh, qf, q_abs_ref)) = trainer.last_isv_magnitude_bin_q_means() {
let max_mean = qq.max(qh).max(qf);
let bias_gap_full = (max_mean - qf).max(0.0);
let collapse_frac_full = if q_abs_ref > 1e-6 {
(bias_gap_full / q_abs_ref).min(1.0)
} else {
0.0
};
println!(
"[ISV_BIN_MEANS] q_q={:.4} q_h={:.4} q_f={:.4} q_abs_ref={:.4} \
max_mean={:.4} bias_gap_full={:.4} collapse_frac_full={:.3}",
qq, qh, qf, q_abs_ref, max_mean, bias_gap_full, collapse_frac_full
);
}
assert!(
eh + ef >= 0.30,
"H10 regression: eval Half + Full share {:.3} < 0.30 (eq={:.3} eh={:.3} ef={:.3}) \
@@ -74,5 +105,30 @@ fn test_magnitude_distribution() -> Result<()> {
make the magnitude branch useful, not remove it.",
eh + ef, eq, eh, ef
);
// Task 2.X "make Full useful" — the adaptive ISV-driven C51 bin weight
// (see gpu_dqn_trainer.rs Q_MAG_SPREAD_INDEX / Q_DIR_SPREAD_INDEX and
// c51_loss_kernel.cu get_magnitude_bin_weight) must push Q(Full) high
// enough that the eval-mode strict-argmax actually selects Full at
// least 5% of the time. Pre-fix baseline observed ef=0.000 because
// Q(Half) > Q(Full) + 1e-6 deterministically across all states after
// Task-2.2 H10 tie-break. The fix is signal-driven (deficit between
// magnitude and direction Q-spread) with a bin-proportional weight,
// so it self-disables once the magnitude head differentiates.
assert!(
ef >= 0.05,
"Task 2.X adaptive magnitude: eval Full share {:.3} < 0.05 \
(eq={:.3} eh={:.3} ef={:.3}) — ISV bin-weight mechanism failed to \
lift Q(Full) above Q(Half) + 1e-6. Investigate: \
(a) ISV slots [13]/[14] populated each stats cadence (check \
q_branch_spread_reduce + isv_signal_update wiring), \
(b) spread_deficit positive during training (magnitude-branch \
Q-spread collapsed below direction-branch reference), \
(c) bin_weight applied consistently to branch_ce (c51_loss_batched) \
AND gradient (c51_grad_kernel d==1 block). Do NOT fall back to \
static tuning knobs — per feedback_adaptive_not_tuned.md the \
escalation path is richer ISV signals / kernel-side arithmetic.",
ef, eq, eh, ef
);
Ok(())
}

22
crates/ml/src/trainers/dqn/trainer/mod.rs
View File

@@ -1524,6 +1524,28 @@ impl DQNTrainer {
self.last_eval_magnitude_dist
}
/// Task 2.X diagnostic — per-direction eval action distribution, readback
/// from the most recent validation backtest. Layout: [Short, Hold, Long,
/// Flat] in [0, 1]. When magnitude_distribution's `ef` collapses to zero
/// this lets the test surface whether the root cause is direction-side
/// (Hold + Flat > 0.95) or magnitude-side (direction healthy but mag
/// collapsed). Returns `None` if no GPU evaluator is wired.
pub fn last_eval_direction_dist(&self) -> Option<[f32; 4]> {
self.gpu_evaluator
.as_ref()
.and_then(|ev| ev.read_eval_action_distribution_per_direction().ok())
}
/// Task 2.X diagnostic — ISV slots [13..16] (per-magnitude-bin Q-mean
/// EMAs + absolute-scale reference). Returns
/// `(q_mean_quarter, q_mean_half, q_mean_full, q_abs_ref)` or `None`
/// if the fused trainer has not been initialised yet.
pub fn last_isv_magnitude_bin_q_means(&self) -> Option<(f32, f32, f32, f32)> {
self.fused_ctx
.as_ref()
.map(|ctx| ctx.trainer().read_isv_magnitude_bin_q_means())
}
/// Per-epoch magnitude-branch action entropy history.
/// Each entry is (epoch_0_indexed, normalized_entropy ∈ [0, 1]).
/// Used by exploration_coverage smoke test.