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

feat(sp4): Tasks A14+A15 — Pearl C engagement counter in 5 Adam kernels

Browse Source 
Layer A scaffolding for post-clamp feedback-loop self-correction.

A14: All 5 Adam kernels (dqn/iqn/iql/attn/curiosity) extended with
register-then-tree-reduce engagement counter. Per-thread `local_engage`
set on clamp engagement; block tree-reduce (no atomicAdd, mirrors
dqn_grad_norm_kernel pattern); single block-leader writes per-block
count to clamp_engage_per_block_buf[engage_buf_offset + blockIdx.x].
`engage_buf_offset == -1 (SP4_ENGAGE_OFFSET_DISABLED)` skips writeback.

A15: Host-side rate-deficit check in pearl_c_post_adam_engagement_check.
Sums per-block counts via mapped-pinned zero-copy reads, computes
engagement_rate, rate_deficit = rate - 0.01. Applies Pearls A+D via
pearl_c_rate_deficit_ema (host-only [f32;8]) + pearl_c_rate_deficit_state_buf
(MappedF32Buffer[24] = 8 groups × 3 Wiener floats). Wired post-graph
in FusedTrainingCtx::run_full_step for groups 0/3/4/5/6 and
post-curiosity_train in training_loop for group 7.

3 design issues resolved per Layer A scope:
1. Curiosity sub-launches: SP4_ENGAGE_BUF_LEN extended 2048->2816
   (= 11 x 256). Curiosity sub-launches use offsets 1792/2048/2304/2560
   (W1/b1/W2/b2). Host-side check sums all 4. Allocation +
   reset-registry updated.
2. TLOB/Attn sharing attn_adam_kernel: TLOB passes
   SP4_ENGAGE_OFFSET_DISABLED=-1 (silently skips Pearl C). Attn writes
   to its slot (offset 6x256=1536). Layer B can refactor if needed.
3. DQN main Adam covers groups 0/1/2 in single launch: Layer A accounts
   only under group 0 (DqnTrunk). Groups 1/2 Pearl C tracking deferred
   to Layer B's per-group sub-launch decision.

Wiring path: GpuDqnTrainer exposes nan_flags_buf_ptr() +
clamp_engage_per_block_buf_dev_ptr(). FusedTrainingCtx adds
wire_aux_trainer_pearl_c_buffers() for IQN/IQL hi+lo/Attn.
GpuExperienceCollector adds set_curiosity_pearl_c_buffers() for the
curiosity trainer. Wiring fires once in init_gpu_experience_collector.

State reset registry (Task A12 follow-up):
- sp4_clamp_engage_counters description updated to reflect 2816 buffer
  length and 4 distinct curiosity sub-launch offsets
- New entries: sp4_pearl_c_rate_deficit_state (24 mapped-pinned floats,
  Pearls A+D Wiener state) and sp4_pearl_c_rate_deficit_ema (host-only
  [f32; 8] EMA surrogate). Both reset to zero at fold boundary so Pearl
  A's first-observation sentinel fires on the new fold's first
  engagement-rate-deficit observation. Dispatch arms wired in
  reset_named_state alongside existing sp4_* entries.

Layer A: Pearl C is observability scaffolding. Mech 9's clamp still
uses hardcoded 100xQ_ABS_REF. Engagement counters fire correctly,
rate_deficit EMA tracks. Force-bump branch logs via tracing::debug only;
Layer B's atomic flip activates the actual ISV mutation.

cargo check --lib --tests clean. state_reset_registry tests pass.
sp4_isv_slots::tests::pearl_c_engage_buf_layout pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 01:42:40 +02:00
parent 91535deb1a
commit ea31ebfe38
19 changed files with 1291 additions and 176 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long