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

fix(sp11): B1a — saboteur GPU multiplication + SimHash state_stride

Browse Source 
Per feedback_cpu_is_read_only, saboteur effective scale now computed
on-device:
- saboteur_generate_params kernel signature gains isv ptr +
  saboteur_intensity_mult_slot parameter
- Kernel reads `mult = fmaxf(isv[saboteur_intensity_mult_slot],
  SABOTEUR_MIN)` (sentinel-0 defense for cold-start before A2's
  controller first runs) and applies `effective_scale = base × mult`
  to perturbation generation
- gpu_experience_collector.rs launcher updated; only one call site

SimHash state_stride parameter added to lookup + update kernels —
prepares for B1c replay-time curiosity wiring against trainer.
states_buf which is STATE_DIM_PADDED=128-strided. Kernel inner loop
reads `state[i × state_stride + d]` (was `i × 42 + d`). Proj-init
kernel unchanged (writes projection, doesn't read states).

Pre-requisite for B1c (curiosity wiring) and a small atomic step
toward full SP11 production behavior.

cargo check + build clean; 6/6 SP11 GPU tests + 14/14 contract tests
still pass.
This commit is contained in:
jgrusewski
2026-05-04 08:51:23 +02:00
parent 302992f63a
commit d5e1214f25
4 changed files with 147 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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;

18
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -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!(

37
crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu
View File

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

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

@@ -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`).