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>
ea31ebfe38