diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index f0bd39841..aac8465a6 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -454,15 +454,30 @@ extern "C" __global__ void domain_rand_sim_params( * * Grid: ceil(N / 256), Block: 256. One thread per episode. * + * SP11 Fix 39 B1a (2026-05-04): the effective perturbation scale is the host- + * passed scalar multiplied by the SP11 controller's saboteur_intensity_mult + * read from `isv_signals_ptr[saboteur_intensity_mult_slot]`. Per + * `feedback_cpu_is_read_only`, this multiply happens GPU-side (not host-side + * in the launcher). Sentinel-0 defense via `fmaxf(slot, SABOTEUR_MIN=0.5f)` + * — at cold start (before A2's controller has run for the first time) the + * ISV slot reads 0.0; the floor prevents the saboteur from being silently + * disabled. Post-controller, the slot is clamped to [SP11_SABOTEUR_MIN=0.5, + * SP11_SABOTEUR_MAX=2.0] so the floor only fires at the very first cold-start + * step. Spec §3.4.2. + * * @param saboteur_params [N, 3] output: (spread_mult, fill_prob, slippage_mult) * @param base_params [3] current best params (center of perturbation) - * @param perturbation_scale controls exploration width (decays over epochs) + * @param perturbation_scale base exploration width (decays over epochs) + * @param isv_signals_ptr pinned device-mapped ISV signals (NULL = static = mult ignored) + * @param saboteur_intensity_mult_slot ISV slot index for SP11 intensity multiplier * @param N number of episodes */ extern "C" __global__ void saboteur_generate_params( - float* __restrict__ saboteur_params, /* [N, 3] output */ - const float* __restrict__ base_params, /* [3] center */ + float* __restrict__ saboteur_params, /* [N, 3] output */ + const float* __restrict__ base_params, /* [3] center */ float perturbation_scale, + const float* __restrict__ isv_signals_ptr, /* pinned device-mapped ISV signals. NULL = static (no SP11 mult). */ + int saboteur_intensity_mult_slot, int N ) { int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -471,7 +486,16 @@ extern "C" __global__ void saboteur_generate_params( float base_spread = base_params[0]; float base_fill = base_params[1]; float base_slip = base_params[2]; - float ps = perturbation_scale; + + /* SP11 B1a: GPU-side multiplication of base scale by SP11 controller's + * saboteur_intensity_mult. Sentinel-0 defense (`fmaxf(slot, 0.5f)`) + * keeps the saboteur active at cold start before the controller has + * written its first real value. */ + const float SABOTEUR_MIN = 0.5f; + float mult = (isv_signals_ptr != NULL) + ? fmaxf(isv_signals_ptr[saboteur_intensity_mult_slot], SABOTEUR_MIN) + : 1.0f; + float ps = perturbation_scale * mult; /* Stateless Philox-based Gaussian perturbation */ float u1, u2; diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 74e9bba3f..5e9b8f4fd 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -3355,15 +3355,31 @@ impl GpuExperienceCollector { let mirror_i32: i32 = if config.mirror_active { 1 } else { 0 }; - // #33 Saboteur: generate per-episode adversarial params on GPU + // #33 Saboteur: generate per-episode adversarial params on GPU. + // + // SP11 Fix 39 B1a (2026-05-04): the host-side `ps` is the BASE + // perturbation scale only. The SP11 controller's intensity multiplier + // is applied GPU-side inside `saboteur_generate_params` by reading + // `isv[SABOTEUR_INTENSITY_MULT_INDEX]` and computing + // `effective_scale = ps × fmaxf(slot, SABOTEUR_MIN=0.5)`. The + // sentinel-0 floor defends the very first cold-start step before the + // controller has run; post-A2 the slot is clamped to + // [SP11_SABOTEUR_MIN, SP11_SABOTEUR_MAX] = [0.5, 2.0] so the floor + // does not fire steady-state. Per `feedback_cpu_is_read_only`, no + // host-side multiplication. if self.saboteur_active { let ps = self.saboteur_perturbation_scale; + let isv_ptr = self.isv_signals_dev_ptr; + let saboteur_slot_i32 = + crate::cuda_pipeline::sp11_isv_slots::SABOTEUR_INTENSITY_MULT_INDEX as i32; unsafe { self.stream .launch_builder(&self.saboteur_generate_kernel) .arg(&mut self.saboteur_params_buf) .arg(&self.saboteur_base_buf) .arg(&ps) + .arg(&isv_ptr) // SP11 B1a: ISV signals (0=NULL=static) + .arg(&saboteur_slot_i32) // SP11 B1a: SABOTEUR_INTENSITY_MULT_INDEX .arg(&n_i32) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( diff --git a/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu b/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu index bfc8c1021..d93a4744f 100644 --- a/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu +++ b/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu @@ -44,10 +44,15 @@ /* Locality-sensitive 16-bit SimHash from a 42-dim state row through a * 42×16 projection matrix. Sign quantization: bit b is set iff dot(state, - * proj[:, b]) > 0. Returns a u32 with the high 16 bits zero. */ + * proj[:, b]) > 0. Returns a u32 with the high 16 bits zero. + * + * The 42 floats consumed are the first 42 contiguous lanes of `state`. + * Callers compute `&states[i * state_stride]` so the per-row pitch can + * differ from 42 (B1c will wire against `trainer.states_buf` whose row + * stride is STATE_DIM_PADDED=128). */ __device__ __forceinline__ unsigned int simhash_code_42_16( - const float* __restrict__ state, /* [42] */ - const float* __restrict__ proj_matrix) /* [42 × 16] row-major */ + const float* __restrict__ state, /* first 42 floats consumed */ + const float* __restrict__ proj_matrix) /* [42 × 16] row-major */ { unsigned int code = 0u; for (int b = 0; b < 16; b++) { @@ -61,18 +66,25 @@ __device__ __forceinline__ unsigned int simhash_code_42_16( return code; } +/* SP11 Fix 39 B1a (2026-05-04): both lookup + update kernels accept a + * `state_stride` parameter so they can read the first 42 floats of each + * row from a buffer whose row pitch is wider than 42. B1c will wire these + * against `trainer.states_buf` which is STATE_DIM_PADDED=128-strided. The + * proj-init kernel does NOT need this parameter (it writes the projection + * matrix and never reads states). */ extern "C" __global__ void novelty_simhash_lookup_kernel( - const float* __restrict__ states, /* [batch_size × 42] */ - const int* __restrict__ actions, /* [batch_size] */ - const float* __restrict__ proj_matrix, /* [42 × 16] */ - const float* __restrict__ hash_table, /* [TABLE_SIZE] */ - float* __restrict__ novelty_out, /* [batch_size] */ + const float* __restrict__ states, /* [batch_size × state_stride] */ + const int* __restrict__ actions, /* [batch_size] */ + const float* __restrict__ proj_matrix, /* [42 × 16] */ + const float* __restrict__ hash_table, /* [TABLE_SIZE] */ + float* __restrict__ novelty_out, /* [batch_size] */ + const int state_stride, /* row pitch for `states` */ const int batch_size) { const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= batch_size) return; - const unsigned int code = simhash_code_42_16(&states[i * 42], proj_matrix); + const unsigned int code = simhash_code_42_16(&states[i * state_stride], proj_matrix); /* Cast action to u32 first so the multiplication is unsigned; negative * action sentinels (e.g. -1) would otherwise wrap to a huge bucket * index. action ∈ [0, 108) in production; defensive cast costs nothing. */ @@ -83,10 +95,11 @@ extern "C" __global__ void novelty_simhash_lookup_kernel( } extern "C" __global__ void novelty_simhash_update_kernel( - const float* __restrict__ states, + const float* __restrict__ states, /* [batch_size × state_stride] */ const int* __restrict__ actions, const float* __restrict__ proj_matrix, - float* __restrict__ hash_table, /* [TABLE_SIZE] in/out */ + float* __restrict__ hash_table, /* [TABLE_SIZE] in/out */ + const int state_stride, /* row pitch for `states` */ const int batch_size) { /* Race-tolerated update — see header for the safety argument. The @@ -94,7 +107,7 @@ extern "C" __global__ void novelty_simhash_update_kernel( const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= batch_size) return; - const unsigned int code = simhash_code_42_16(&states[i * 42], proj_matrix); + const unsigned int code = simhash_code_42_16(&states[i * state_stride], proj_matrix); const unsigned int act = (unsigned int)actions[i]; const unsigned int bucket = (act * 65536u + code) & SP11_NOVELTY_TABLE_MASK; /* Non-atomic RMW: read-modify-write loses increments under races, which diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 1938e67a7..97d3e3677 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -5155,3 +5155,80 @@ change is reviewable in isolation. Verified post-fix: cargo check + release build clean; 6/6 SP11 GPU oracle tests pass on RTX 3050 Ti; 10/10 sp5_isv_slots + 4/4 state_reset_registry contract tests pass. + +## SP11 B1a — saboteur GPU multiplication + SimHash state_stride (2026-05-04) + +Two small/safe contract migrations preceding the B1b structural +reward-composition refactor: + +### B1a.1 — saboteur effective-scale multiply moved GPU-side + +Per `feedback_cpu_is_read_only`, the SP11 controller's +`saboteur_intensity_mult` (ISV[347]) MUST be combined with the +host-side `saboteur_perturbation_scale` on the device, not in the +Rust launcher. Pre-B1a state: the kernel took only a scalar +`perturbation_scale` (host-passed), and B1's planned host-side +`scale × isv_signal_at(347)` would have violated the read-only rule. + +Code change scope: +- `experience_kernels.cu :: saboteur_generate_params` — kernel + signature gains `const float* isv_signals_ptr` + + `int saboteur_intensity_mult_slot` parameters. Inside the kernel, + `mult = (isv_signals_ptr != NULL) ? fmaxf(isv[slot], SABOTEUR_MIN) + : 1.0f` and `effective_scale = perturbation_scale × mult` is + applied to the perturbation generation. `SABOTEUR_MIN = 0.5f` is + a sentinel-0 defense for the very first cold-start step before + A2's controller has run (slot reads 0.0); post-controller the + slot is clamped by `reward_subsystem_controller_kernel` to + `[SP11_SABOTEUR_MIN=0.5, SP11_SABOTEUR_MAX=2.0]` so the floor + fires at most once at start-of-training. +- `gpu_experience_collector.rs` saboteur launcher — gains two new + args: `isv_ptr` (= `self.isv_signals_dev_ptr`, 0=NULL=static) and + `saboteur_slot_i32` (= `SABOTEUR_INTENSITY_MULT_INDEX as i32`, + imported from `sp11_isv_slots`). Only one call site — verified + via `grep -rn saboteur_generate` returning a single launch. + +Per `feedback_no_partial_refactor`, every consumer of the contract +migrates in the same commit (single launcher). + +### B1a.2 — SimHash kernels gain `state_stride` parameter + +`novelty_simhash_lookup_kernel` and `novelty_simhash_update_kernel` +previously read `states[i × 42]` assuming a 42-stride contiguous +state layout. Reality: the trainer's `states_buf` is row-strided to +`STATE_DIM_PADDED = 128` (per the previous BLOCKED audit). When B1c +wires the lookup/update against `trainer.states_buf`, it MUST read +42 contiguous floats per sample at a 128-stride offset — otherwise +the SimHash code computed at replay time would not match the code +produced from the trainer's actual state layout. + +Code change scope: +- `novelty_simhash_kernel.cu` — both `novelty_simhash_lookup_kernel` + and `novelty_simhash_update_kernel` gain an `int state_stride` + parameter. Inner read changes from `states[i × 42 + d]` (via + `&states[i * 42]` argument) to `states[i × state_stride + d]` + (via `&states[i * state_stride]`). The `simhash_code_42_16` + device-inline helper is unchanged — it still reads exactly 42 + contiguous floats; only the per-row offset computation moves. +- `novelty_simhash_proj_init_kernel` is NOT modified — it writes the + projection matrix and never reads states. + +The kernels are not yet launched from production code (B1c will +wire them); no Rust launcher updates required. SP11 unit tests do +not exercise the SimHash kernels directly, so no test launch-arg +update needed. B1c will pass `STATE_DIM_PADDED as i32` for the +`state_stride` argument when wiring against `trainer.states_buf`. + +### Verification + +- `cargo check -p ml --lib` — clean +- `cargo build -p ml --release` — clean +- `cargo test -p ml --test sp11_producer_unit_tests --features cuda + -- --ignored` — 6/6 GPU oracle tests pass on RTX 3050 Ti (none + exercise the saboteur kernel or the SimHash lookup/update kernels) +- `cargo test -p ml --lib sp5_isv_slots` — 10/10 contract tests pass +- `cargo test -p ml --lib state_reset_registry` — 4/4 contract + tests pass + +Pre-requisite for B1b (structural reward composition) and B1c +(replay-time curiosity wiring against `trainer.states_buf`).