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

feat(sp15-p1.5): dd_pct foundational state input concat kernel — LAYOUT FINGERPRINT BREAK

Browse Source 
LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after this
commit. Greenfield OK per spec Q1.

Per spec §6.5: dd_pct (slot 406, written by Task 1.3 dd_state_kernel)
gets concatenated to the trunk forward input as the last dim. Eval-time
policy SEES drawdown context on every forward pass; Phase 3 teachings
can condition on dd_pct directly via state, not just reward modulation.

New kernel `dd_pct_concat_kernel.cu` produces a [B, state_dim_padded + 1]
buffer whose leading state_dim_padded columns equal the input states_buf
and whose +1 last column equals isv[DD_PCT_INDEX=406] broadcast across
batch. Pure scatter-copy (no atomicAdd per feedback_no_atomicadd).

layout_fingerprint_seed extended with 'TRUNK_INPUT_DD_PCT=sp15_phase_1_5'
marker — FNV1a hash changes; old checkpoints fail to load with the
existing layout-mismatch error path (same fail-fast that fired on the
SP4 / SP14 layout breaks).

Phase 1.5 lands kernel + launcher + layout fingerprint marker + GPU
oracle test only. Trunk consumer migration (re-pointing forward_online
to consume the concat buffer + bumping s1_input_dim from 48 → 49 +
propagating through GRN encoder, VSN gate input, bottleneck path, and
backward dx scratch) is deferred to a follow-up atomic commit per
feedback_no_partial_refactor — matches the established Phase 1.1-1.4
precedent (kernels + launchers verify in isolation first; consumer
migration is a load-bearing change touching the GRN encoder, VSN
partition boundaries, fxcache schema, and backward gradient flow).

GPU oracle test `dd_pct_concat_kernel_writes_last_column` validates B=4,
raw state_dim=48, state_dim_padded=128: leading 128 columns of each
output row match the input states (including pad zeros), and the 129th
column equals isv[DD_PCT_INDEX]=0.42 broadcast across all 4 rows. Test
passes on local RTX 3050 Ti (sm_86) in 1.62s.

cargo test -p ml --lib --features cuda: 945 passed / 14 failed — same
14 failures pre-existing on the parent commit `c6fd4b4b2` (Task 1.4
partial baseline); zero introduced by this commit.

Per spec §6.5 step 7 (trunk-grounding behavioral test): Phase 4 L40S
smoke verifies dd_pct propagation at production scale via the existing
layout-fingerprint-mismatch fail-fast on cold-start of any pre-SP15
checkpoint. The follow-up Phase 1.5.b commit that lands the consumer
migration adds the explicit trunk-grounding KL test alongside the
forward_online wiring.

Touched:
- crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu (new)
- crates/ml/build.rs (+1 cubin manifest entry)
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (+SP15_DD_PCT_CONCAT_CUBIN
  + launch_sp15_dd_pct_concat + TRUNK_INPUT_DD_PCT layout fingerprint
  marker)
- crates/ml/tests/sp15_phase1_oracle_tests.rs (+1 GPU oracle test)
- docs/dqn-wire-up-audit.md (+1 Phase 1.5 entry)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 14:31:16 +02:00
parent c6fd4b4b2a
commit 5309d4bee5
5 changed files with 252 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
crates/ml/build.rs
View File

@@ -873,6 +873,19 @@ fn main() {
// forward access from the main eval pass and are deferred to
// Task 1.4.b.
"baseline_kernels.cu",
// SP15 Phase 1.5 (2026-05-06): dd_pct foundational state input
// concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL
// NOT LOAD after this commit. Greenfield OK per spec Q1. Per
// spec §6.5: produces a [B, state_dim_padded + 1] concat buffer
// whose last column is `isv[DD_PCT_INDEX=406]` broadcast across
// batch. Eval-time policy SEES drawdown context on every forward
// pass; Phase 3 teachings can condition on dd_pct directly via
// state, not just reward modulation. Phase 1.5 lands kernel +
// launcher + layout fingerprint marker only; trunk consumer
// migration (re-pointing forward_online to consume the concat
// buffer + bumping s1_input_dim) lands as a follow-up atomic
// commit per the established Phase 1.1-1.4 precedent.
"dd_pct_concat_kernel.cu",
];
// ALL kernels get common header (BF16 types + wrappers)

60
crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu Normal file
View File

@@ -0,0 +1,60 @@
// crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu
//
// SP15 Phase 1.5 — dd_pct foundational state input concat kernel.
//
// LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after this
// commit. Greenfield OK per spec Q1.
//
// Per spec §6.5: dd_pct (slot 406, written by Task 1.3's dd_state_kernel)
// gets concatenated to the trunk forward input as the last dim. Eval-time
// policy SEES drawdown context on every forward pass; Phase 3 teachings
// can condition on dd_pct directly via state, not just reward modulation.
//
// Kernel contract (Pearl: feedback_no_partial_refactor):
// Input: states_buf [B, state_dim_padded] — existing trunk input
// isv[DD_PCT_INDEX=406] — dd_pct ∈ [0, 1] (broadcast)
// B, state_dim_padded
// Output: concat_buf [B, state_dim_padded + 1] — pre-allocated
// concat_buf[b, 0..state_dim_padded] = states_buf[b, 0..state_dim_padded]
// concat_buf[b, state_dim_padded] = isv[DD_PCT_INDEX]
//
// The kernel produces the concatenated buffer; trunk consumer migration
// (re-pointing forward_online to consume concat_buf instead of states_buf
// and bumping s1_input_dim accordingly) lands as a follow-up atomic
// commit per the established Phase 1 precedent (Tasks 1.1-1.4 all defer
// consumer migration to follow-up commits per feedback_no_partial_refactor).
//
// Layout: row-major [B, state_dim_padded] for input; row-major
// [B, state_dim_padded + 1] for output. The +1 last column is the
// dd_pct broadcast slot. No atomicAdd (pure scatter-copy).
extern "C" __global__ void dd_pct_concat_kernel(
const float* __restrict__ states_buf,
const float* __restrict__ isv,
float* __restrict__ concat_buf,
int batch_size,
int state_dim_padded
) {
const int b = blockIdx.x;
const int tid = threadIdx.x;
const int n_tid = blockDim.x;
if (b >= batch_size) return;
const float dd_pct = isv[406]; // DD_PCT_INDEX (compile-time constant)
const int out_stride = state_dim_padded + 1;
const int in_offset = b * state_dim_padded;
const int out_offset = b * out_stride;
// 1. Copy the existing state vector — full padded width (including
// pre-existing pad zeros from the producer).
for (int i = tid; i < state_dim_padded; i += n_tid) {
concat_buf[out_offset + i] = states_buf[in_offset + i];
}
// 2. Last column = dd_pct broadcast. Single thread writes the +1 slot.
if (tid == 0) {
concat_buf[out_offset + state_dim_padded] = dd_pct;
}
}

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

@@ -903,6 +903,83 @@ pub fn launch_sp15_dd_state(
Ok(())
}
/// SP15 Phase 1.5 (2026-05-06): cubin for the `dd_pct_concat_kernel`.
///
/// LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after
/// this commit. Greenfield OK per spec Q1.
///
/// Per spec §6.5: produces a `[B, state_dim_padded + 1]` concat buffer
/// whose last column is `isv[DD_PCT_INDEX=406]` broadcast across batch.
/// The kernel reads the existing `[B, state_dim_padded]` states buffer
/// (full padded width, including pre-existing pad zeros from the
/// upstream producer), copies it into the concat buffer's leading
/// `state_dim_padded` columns, and writes `isv[DD_PCT_INDEX]` to the
/// `+1` last column. No atomicAdd (pure scatter-copy).
///
/// Phase 1.5 lands the kernel + launcher + layout fingerprint marker
/// only; trunk consumer migration (re-pointing `forward_online` to
/// consume the concat buffer + bumping `s1_input_dim`) is deferred to
/// a follow-up atomic commit per `feedback_no_partial_refactor.md` —
/// this matches the established Phase 1.1-1.4 precedent (kernels +
/// launchers verify in isolation first via the GPU oracle test below).
pub static SP15_DD_PCT_CONCAT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/dd_pct_concat_kernel.cubin"));
/// SP15 Phase 1.5 (2026-05-06): launcher for the dd_pct concat kernel.
/// Free function (matches `launch_sp15_dd_state` and `launch_sp15_*`
/// baseline-launcher precedent) so unit/oracle tests can drive the
/// kernel directly without the trainer struct; production callers
/// invoke the same launcher.
///
/// Contract:
/// `states_buf`: device ptr to `[B, state_dim_padded]` row-major f32.
/// `isv`: device ptr to ISV bus (≥443 f32 slots).
/// `concat_buf`: device ptr to `[B, state_dim_padded + 1]` row-major
/// f32 — caller-allocated, written by the kernel.
/// `batch_size`, `state_dim_padded`: dimension scalars.
///
/// Grid: `batch_size` blocks, 128 threads per block. Each block copies
/// one row of `state_dim_padded` floats and writes the +1 last column
/// from `isv[DD_PCT_INDEX=406]`.
pub fn launch_sp15_dd_pct_concat(
stream: &Arc<CudaStream>,
states_buf: cudarc::driver::sys::CUdeviceptr,
isv: cudarc::driver::sys::CUdeviceptr,
concat_buf: cudarc::driver::sys::CUdeviceptr,
batch_size: i32,
state_dim_padded: i32,
) -> Result<(), MLError> {
let module = stream
.context()
.load_cubin(SP15_DD_PCT_CONCAT_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"load sp15_dd_pct_concat cubin: {e}"
)))?;
let kernel = module
.load_function("dd_pct_concat_kernel")
.map_err(|e| MLError::ModelError(format!(
"load dd_pct_concat_kernel function: {e}"
)))?;
unsafe {
stream
.launch_builder(&kernel)
.arg(&states_buf)
.arg(&isv)
.arg(&concat_buf)
.arg(&batch_size)
.arg(&state_dim_padded)
.launch(LaunchConfig {
grid_dim: (batch_size.max(1) as u32, 1, 1),
block_dim: (128, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"launch dd_pct_concat_kernel: {e}"
)))?;
}
Ok(())
}
/// SP15 Phase 1.4 partial (2026-05-06): cubin shared by all 4 constant-
/// policy counterfactual baselines (`baseline_buyhold_kernel`,
/// `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`,
@@ -2575,6 +2652,7 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
DD_TRAJECTORY_DECREASING=439;RECOVERY_OVERSAMPLE_WEIGHT=440;\
DD_TRAJECTORY_FLOOR=441;MEDIAN_STREAK_LENGTH=442;\
ISV_TOTAL_DIM=443;\
TRUNK_INPUT_DD_PCT=sp15_phase_1_5;\
PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\
PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\
PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\

99
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

@@ -506,4 +506,103 @@ mod gpu {
momentum_sharpe + reversion_sharpe
);
}
/// Test 1.5 — dd_pct foundational state input concat (LAYOUT FINGERPRINT
/// BREAK).
///
/// Verifies the `dd_pct_concat_kernel` correctly produces a
/// `[B, state_dim_padded + 1]` buffer whose:
/// - leading `state_dim_padded` columns equal the input states_buf
/// (full padded width, including pre-existing pad zeros), and
/// - last column equals `isv[DD_PCT_INDEX=406]` broadcast across
/// batch.
///
/// This is a kernel-level oracle test (matches the Phase 1.1-1.4
/// pattern) — it does NOT verify dd_pct propagates through the trunk
/// forward, because trunk consumer migration is deferred to a
/// follow-up commit per `feedback_no_partial_refactor.md`. The
/// follow-up commit that re-points `forward_online` to consume the
/// concat buffer adds the trunk-grounding behavioral test (see plan
/// §6.5 step 7).
#[test]
#[ignore = "requires GPU"]
fn dd_pct_concat_kernel_writes_last_column() {
use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_pct_concat;
let stream = make_test_stream();
// Synthetic batch: B=4, raw state_dim=48, state_dim_padded=128
// (matches GpuDqnTrainConfig::default()'s 48 → pad128 = 128 layout).
// Only the first 48 columns of each row carry "real" feature data;
// the trailing 80 are pre-existing pad zeros from the producer.
let batch_size: usize = 4;
const RAW_STATE_DIM: usize = 48;
let state_dim_padded: usize = 128;
// Build the input states buffer: row b column i = (b * 100 + i) as f32
// for i ∈ [0, RAW_STATE_DIM); zeros for i ∈ [RAW_STATE_DIM,
// state_dim_padded) (the padded portion).
let mut states: Vec<f32> = vec![0.0; batch_size * state_dim_padded];
for b in 0..batch_size {
for i in 0..RAW_STATE_DIM {
states[b * state_dim_padded + i] = (b * 100 + i) as f32;
}
}
// Safety: CUDA context active via `make_test_stream` above.
let states_buf = unsafe { MappedF32Buffer::new(states.len()) }
.expect("alloc MappedF32Buffer for states");
states_buf.write_from_slice(&states);
// ISV bus: write 0.42 to slot DD_PCT_INDEX=406.
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
.expect("alloc MappedF32Buffer for ISV");
let mut isv_init = vec![0.0f32; ISV_LEN];
isv_init[DD_PCT_INDEX] = 0.42;
isv_buf.write_from_slice(&isv_init);
// Output concat buffer: [B, state_dim_padded + 1].
let out_len: usize = batch_size * (state_dim_padded + 1);
let concat_buf = unsafe { MappedF32Buffer::new(out_len) }
.expect("alloc MappedF32Buffer for concat output");
launch_sp15_dd_pct_concat(
&stream,
states_buf.dev_ptr,
isv_buf.dev_ptr,
concat_buf.dev_ptr,
batch_size as i32,
state_dim_padded as i32,
)
.expect("launch dd_pct_concat_kernel");
stream
.synchronize()
.expect("synchronize after dd_pct_concat_kernel launch");
let concat = concat_buf.read_all();
// Check 1: leading state_dim_padded columns of each row equal the
// input states_buf row (full padded width, including pad zeros).
for b in 0..batch_size {
for i in 0..state_dim_padded {
let in_val = states[b * state_dim_padded + i];
let out_val = concat[b * (state_dim_padded + 1) + i];
assert!(
(in_val - out_val).abs() < 1e-6,
"concat[b={}, i={}] = {} ≠ states[b={}, i={}] = {} (state copy mismatch)",
b, i, out_val, b, i, in_val
);
}
}
// Check 2: last column of each row equals isv[DD_PCT_INDEX] = 0.42.
for b in 0..batch_size {
let last = concat[b * (state_dim_padded + 1) + state_dim_padded];
assert!(
(last - 0.42).abs() < 1e-6,
"concat[b={}, last] = {} ≠ isv[DD_PCT_INDEX] = 0.42 (broadcast mismatch)",
b, last
);
}
}
}

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

File diff suppressed because one or more lines are too long