52 Commits

Include renames

Author	SHA1	Message	Date
Omar SobhandClaude Sonnet 4.6	960fd73c82	feat(batch18): contrastive losses, feature distillation, advanced data samplers ... - ContrastiveLoss: NT-Xent/SimCLR (arXiv:2002.05709), InfoNCE in-batch (arXiv:1807.03748), SupCon with multi-positive P(i) (arXiv:2004.11362); L2-normalize + log-sum-exp stable; 25 tests - FeatureDistillation: FitNets hint L2 (arXiv:1412.6550), Attention Transfer spatial map matching (arXiv:1612.03928), RKD distance+angle (arXiv:1904.05068) with Huber loss and LCG triplet subsampling; 24 tests - Data samplers: TemperatureSampler (log-space multinomial), ImportanceSampler (easy/hard weighting), StratifiedSampler (equal/proportional), HardNegativeMiner (O(n²) cosine), CurriculumSampler (percentile threshold ramp); 23 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 12:35:31 +00:00
Omar SobhandClaude Sonnet 4.6	bb5f5c519f	feat(batch17): RoPE scaling extensions, DPO loss, label smoothing + focal loss ... - RopeTable/RopeScaler: linear interpolation, dynamic NTK (base scaling), YaRN per-frequency blending with ramp fn + temperature correction (arXiv:2309.00071); apply_to_sequence multi-head; 23 tests - DpoLoss: log-sigmoid DPO (arXiv:2305.18290), IPO squared variant (arXiv:2310.12036), robust DPO label smoothing; implicit reward tracking; DpoAccumulator with preference accuracy; 23 tests - LossFunctions: label-smoothed CE (Szegedy 2016), focal loss (Lin 2017 arXiv:1708.02002), smoothed focal, binary CE (stable), binary focal; Reduction::Mean/Sum/None; 22 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 07:46:08 +00:00
Omar SobhandClaude Sonnet 4.6	54a9652041	feat(batch16): SOAP optimizer, lookahead decoding, SWA+SWAG ... - SoapOptimizer: Adam in Shampoo eigenbasis (arXiv:2409.11321); Jacobi eigendecomposition for L/R Kronecker factors; projection G_hat=Q_L^T@G@Q_R, bias-corrected Adam, unproject U=Q_L@U_hat@Q_R^T; 1D plain Adam fallback; 19 tests - LookaheadDecoder: NGramCache (FIFO eviction, count-sorted candidates); draft-then-verify loop; auto-cache update on accepted tokens; LookaheadStats with avg_tokens_per_step; 22 tests - SwaTrainer+SwagBuffer: cyclic cosine LR schedule; online incremental mean (SwaBuffer); E[θ²]-E[θ]² diagonal variance + low-rank deviation columns; Box-Muller SWAG sample; 29 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 06:43:23 +00:00
Omar SobhandClaude Sonnet 4.6	6daa20c94b	feat(batch15): Shampoo optimizer, beam search decoder, sliding window attention ... - ShampooOptimizer: Kronecker-factored L/R preconditioners; Schulz iteration for A^{-1/4} (two-pass: inv_sqrt then inv_sqrt of sqrt); spectral-norm normalization; large-dim SGD fallback; 16 tests - BeamSearchDecoder: length normalization (Wu et al. α); n-gram blocking; EOS suppression before min_length; DiverseBeamSearchDecoder with per-group diversity penalty; 20 tests - SlidingWindowAttention: causal/bidir window; global tokens attend to all; O(n·W) forward_single_head + multi-head forward; AttentionStats sparsity; WindowMask; 20 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 06:26:31 +00:00
Omar SobhandClaude Sonnet 4.6	e00a6da018	feat(batch14): Muon optimizer, logit processors, per-token activation quantization ... - MuonOptimizer: Nesterov momentum + quintic Newton-Schulz orthogonalization (arXiv:2409.20325); 5-iteration NS maps gradient to near-orthogonal matrix; 1D fallback skips NS; decoupled weight decay; 16 tests - LogitProcessorList: temperature, top-k, top-p nucleus, min-p, repetition/ presence/frequency penalty, eta-sampling; softmax/log_softmax/argmax/ sample_token helpers; 39 tests - ActivationQuantizer: per-token dynamic INT8/FP8E4M3 scaling for inference activations; per-tensor mode; dequantize; max_error diagnostic; 19 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 06:11:25 +00:00
Omar SobhandClaude Sonnet 4.6	9c3b9f82f0	feat(batch13): EMA model weights, cross-layer weight sharing, schedule-free optimizer ... - ModelEma: decay-weighted shadow weights with warmup ramp, bias correction, apply/restore swap for eval, and shadow_drift L2 diagnostic - SharedLayerStack: FullSharing/GroupedSharing/AlternatingPairs strategies (ALBERT-style); memory_reduction_ratio(); LCG-seeded SharedFfnWeight - ScheduleFreeOptimizer: Defazio 2024 z/x dual sequences, c_t cubic interpolation coefficient, Adam+SGD variants, weight decay 48 tests + 5 doctests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 06:05:58 +00:00
Omar SobhandClaude Sonnet 4.6	924c237096	feat(batch12): online quant calibration, draft distillation loss, gradient noise scale ... - Online quantization calibration (rtx-compress): OnlineCalibrator with MaxAbs, EmaMaxAbs{momentum}, Percentile{percentile,bins} methods; streaming observe(); quantize_int8/dequantize_int8; int8_maxabs/int8_ema/fp8_maxabs convenience ctors; ModelCalibrator tracks all tensors; 17 tests + 2 doctests - Draft distillation loss (rtx-transformers): KL(p_target‖p_draft) + CE hard-label with temperature scaling; log_softmax/softmax/kl_divergence/token_acceptance_estimate primitives; DistillAccumulator for epoch-level tracking; normalize_by_length; 13 tests + 6 doctests - Gradient noise scale (rtx-transformers): GradientNoiseScale with McCandlish 2018 two-point B_noise estimator + Welford single-pass mode; EMA smoothing; should_increase/ decrease_batch signals; GnsTracker with bounded history + trend detection; 16 tests + 2 doctests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 05:37:19 +00:00
Omar SobhandClaude Sonnet 4.6	4ae0c34537	feat(batch11): WSD LR scheduler, KV CPU offloading, GQA KV head expansion ... - WSD scheduler (Warmup-Stable-Decay / trapezoidal): linear warmup → constant plateau → cosine/linear/sqrt decay; extend_stable() adds steps mid-run without restart; phase_at()/decay_progress() introspection; 26 tests + 2 doctests - KV CPU offloading: KvCpuOffloadManager LRU-based GPU→CPU page spill with on-demand prefetch; insert() auto-offloads when at gpu_page_limit; stats() with hit rate and utilization; 14 tests - GQA KV head expansion: GqaConfig validates num_q_heads/num_kv_heads divisibility; expand_kv_heads() tiles KV [batch,kv_heads,seq,dim]→[batch,q_heads,seq,dim]; gqa_attention_cpu() with numerically stable softmax + causal mask; 15 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 05:28:27 +00:00
Omar SobhandClaude Sonnet 4.6	be3e3965b1	feat(batch10): attention sinks (StreamingLLM), chunked prefill, per-layer LR decay ... - Attention sinks (arXiv:2309.17453): AttentionSinkEviction always retains first sink_size KV positions + last window_size; evicts middle band in O(evict_count); select_evict_positions/should_retain consistent; 15 tests - Chunked prefill (vLLM arXiv:2309.06180): ChunkedPrefillScheduler splits long prompts into chunk_size=512 chunks interleaved with decode steps (max 128 decode tokens/step); PrefillChunkState tracks progress/remaining/completion; drain_completed(); 14 tests including 1500-token→3-chunk coverage - Per-layer LR decay (ULMFiT / discriminative fine-tuning): LayerLrDecayConfig with base_lr * decay_rate^(num_layers-1-depth); LayerLrDecayBuilder parses layer/layers/ h/blocks/bracket notation param names; LayerLrScheduler with outer multiplier for cosine/linear schedule composition; 14 tests + 1 doctest Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 04:53:53 +00:00
Omar SobhandClaude Sonnet 4.6	ef3cdb1e1a	feat(batch9): token merging (ToMe), grad accum per-step norm, speculative streaming ... - Token Merging (ToMe, arXiv:2210.09461): bipartite soft matching via priority-queue second-chance loop; apply_merge/apply_unmerge; TokenMergingLayer::forward(); 17 unit tests + 2 doctests; 75% merge at r=32/seq=64 - Gradient accumulation per-step normalization: NormalizationStrategy {EndOfAccumulation, PerStep, None}; with_per_step_normalization(); compute_gradient_norm() L2 norm; PerStep divides by fixed accumulation_steps before add (not end-of-batch); 11 tests including equivalence proof vs EndOfAccumulation - Speculative streaming: SpeculativeStreamer + mpsc::Receiver<StreamedToken>; notify_step() sends accepted draft tokens + optional continuation immediately; StreamStats with Welford online mean latency; collect_stream() test helper; 17 async tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 04:38:20 +00:00
Omar SobhandClaude Sonnet 4.6	7c8e9a8a35	feat(batch8): multi-token prediction heads, sparse attention, length bucketing ... Multi-Token Prediction heads (rtx-transformers/gpt): - MtpConfig { num_future_tokens=4, loss_weight=0.3 }; MultiTokenPredictionHead with k independent [hidden, vocab] weight matrices; forward() → Vec<Tensor> - compute_loss(): log-softmax NLL for each offset 1..k; weighted by loss_weight; MtpLossResult with per_head_losses + is_valid(); 12 tests Sparse attention (rtx-transformers/layers): - SparseAttentionMask: local window (radius w), global tokens (first g attend all + all attend them), random long-range (r symmetric positions per token) - LCG seeded for reproducibility; apply_to_scores() masks to -inf; to_additive_bias() - SparseAttentionLayer::forward_cpu() skips masked pairs early; numerically-stable softmax; sparsity 75% at n=512, 87% at n=1024, 97% at n=4096; 14 tests Sequence length bucketing (rtx-transformers/training): - LengthGroupedSampler: Fisher-Yates per-bucket shuffle, token-budget batching, overflow bucket for long sequences; padding_efficiency() vs baseline_efficiency() - pack_into_batch(): greedy first-fit packing; naive_padding_ratio() baseline metric - Measured 2.2× padding reduction on power-law data (26%→67% efficiency); 18 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 04:16:14 +00:00
Omar SobhandClaude Sonnet 4.6	b45a58792d	feat(batch7): interleaved 1F1B, attention-selective checkpointing, flash decoding ... Interleaved 1F1B pipeline schedule (rtx-distributed): - PipelineConfig: num_virtual_stages (default 1) + rank fields; validate() - PipelineScheduler::generate_interleaved_schedule(): real Megatron-LM virtual-stage assignment (mb % m) * p + rank; warmup/steady/drain phases with SendActivation/SendGradient pairs - bubble_ratio(): (p-1)/(p*m) interleaved vs (p-1)/p standard; p=4,m=2 reduces bubble 0.750 → 0.375; 4 new tests, 24 total pass Attention-selective activation checkpointing (rtx-distributed): - CheckpointPolicy::AttentionSelective { attention_patterns } — name-match on attn/attention/self_attn/cross_attn/mha; ~40% memory savings - CheckpointPolicy::Adaptive: replaced layer%2 stub with 3-tier heuristic (>4096MB→sqrt(n), >1024MB→every-other, ≤1024MB→all) - MemoryAwareCheckpointer: AtomicUsize pressure tracking, fallback-to-all when over target; re-exported from crate root; 14 new tests, 29 total pass Flash decoding (rtx-flash-attention): - flash_decode_cpu(): split-K attention with log-sum-exp chunk reduction; matches naive attention within 1e-4 for all tested configs - FlashDecodeKernel wrapper; num_splits_for_seq_len heuristic (256 tok/chunk) - flash_decode_forward.cu: 2-phase CUDA (per-chunk partial + reduce kernel) - SdpaBackend::FlashDecode: score 0.97 for seq_q=1 && kv>=1024; up to 50× speedup at 32K tokens; selected over other backends for long-context decode - 10 unit tests + 3 doctests + 1 backend selector test; all pass Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 04:06:55 +00:00
Omar SobhandClaude Sonnet 4.6	0e1d6a74b6	feat(batch6): windowed acceptance metrics, KV INT8 quant, col/row-parallel linear ... Windowed acceptance rate (rtx-inference/speculative): - WindowedAcceptanceTracker: O(1) VecDeque sliding window, p50/p95/min/max - AcceptanceTrend enum (Rising/Falling/Stable, ±0.05 threshold) - AcceptanceDashboard aggregator; wired into PerformanceMetrics::update() and dashboard(); 13 tests KV cache INT8 quantization (rtx-inference/cache): - KvCacheQuantMode { None, Int8 { scale_per_token }, Fp8E4M3 } enum - KvQuantizer::encode/decode: symmetric per-block INT8 (scale=max_abs/127) gives 4× compression vs f32; Fp8E4M3 CPU proxy, GPU path reserved - QuantizedKvBlock carries data+scale+mode; KvCacheConfig::quant_mode defaulting to None; 14 tests ColParallel + RowParallel linear (rtx-distributed): - ColParallelLinear: shards weight rows across TP ranks, forward_cpu() batch matmul + per-shard bias; no AllReduce (output shards concatenated) - RowParallelLinear: shards weight cols across TP ranks, forward_cpu() partial sum + bias on rank 0 only; async forward() calls ProcessGroup AllReduce for real NCCL path; CPU sim is no-op - TensorParallel::matmul() replaced zeros stub with ColParallelLinear(tp=1) - col→row roundtrip verified within 1e-3; 9 tests Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 03:54:07 +00:00
Omar SobhandClaude Sonnet 4.6	ef786c0ab1	feat(batch5): mid-batch injection, PagedAttn v2 defrag, fused RoPE kernel ... Continuous batching (rtx-serving-api): - ContinuousBatchingConfig: enable_mid_batch_injection (default true), injection_check_interval (default 1), max_injections_per_step (default 4) - ContinuousBatchingController: inject_into_active_batch() + try_inject_pending() allow new sequences to join a running decode batch after each step - BatchingError::BatchFull variant; 3 new tests PagedAttention v2 defrag (rtx-memory): - PageTable::fragmentation_ratio() — hole-counting (sandwiched free pages / total) - PageTable::defragment() — in-place left-compaction of physical page metadata, consistent lock order (free_pages -> physical_pages -> sequences); GPU KV copy stub comment; DefragStats return value; re-exported from lib.rs - 4 defrag tests; fixed 2 pre-existing compile errors in gpu_oom.rs + gpu_transfer.rs - 192 tests pass Fused RoPE kernel (rtx-transformers): - build_cos_sin_table() + rope_forward_cpu() CPU reference (norm-preserving) - RopeFusedKernel wrapper; rope_forward.cu CUDA kernel (1 block per (B,H,T), 1 thread per dim pair, NVRTC compiled) - Replaced apply_rope_rotation() mul_scalar(0.99) stub with real pairwise rotation - build.rs for NVRTC kernel tracking; layers/mod.rs wired; 8 tests pass Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 01:40:47 +00:00
Omar SobhandClaude Sonnet 4.6	d6769ef641	feat(perf): GPU perf batch 4 — SmoothQuant INT8 forward, varlen FA, inference graph capture ... SmoothQuant INT8 linear forward (rtx-compress) - `advanced.rs`: `SmoothQuantizedLayer::forward_raw()` — smooth activations ÷ scales, INT8-quantize both sides (range −127…127 symmetric), INT8 GEMM w/ i32 accumulation, dequantize: acc * act_scale * weight_scale; `out_features()` / `in_features()` helpers - `int8_matmul.rs`: `int8_matvec` + `int8_gemm` (i32 accumulation); 7 unit tests - 4 forward_raw tests: shape, identity layer, scale effect, manual verification - 11 new tests; total 112 pass Variable-length packed flash attention (rtx-flash-attention) [commit 80d7c7f] - `flash_attention_varlen.cu`: WGMMA-compatible CUDA kernel, BLOCK_Q/K=64, block=(128,1,1), grid=(ceil(max_seqlen/64), heads, 1); linear cu_seqlens scan for sequence-to-block mapping; 24KB shared memory (3 × 64 × 64 × 2 bytes); early-exit for past-end blocks - `flash_varlen_forward.rs`: `varlen_attention_cpu` O(n²) reference + `#[cfg(cuda)]` `FlashVarlenKernel` NVRTC wrapper; `SdpaBackend::VarLen` added to backend_selector - 8 CPU tests: single-seq matches regular attn, two seqs independent, causal mask, softmax sums to 1, empty sequence handled, output shape; total 50 pass Inference CUDA graph capture (rtx-inference) - `inference_graph.rs`: `InferenceGraphCapture` + `StepMode` {Warmup, Capture, Replay}; state machine: N warmup steps → capture once → replay forever; `check_static_shape()` invalidates on batch/step change; `record_capture(graph_id)` stores graph - `batch_processor.rs`: `graph_capture: Mutex<InferenceGraphCapture>` field (#[cfg(cuda)]); `advance()` wired at line 852 in `execute_batch_inference`; stream TODO matches training_loop.rs pattern; config gains `enable_decode_graphs`/`graph_warmup_steps` - 8 pure-logic tests; total 93 pass; 4 integration test literals fixed Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 01:11:08 +00:00
Omar SobhandClaude Sonnet 4.6	80d7c7fb6c	feat(flash-attention): add varlen packed-sequence support ... Implements variable-length (varlen) FlashAttention that processes mixed-length batches without padding waste: - New CUDA kernel flash_attention_varlen_forward with BLOCK_Q=64 / BLOCK_K=64 tiling; grid=(ceil(max_seqlen_q/64), num_heads, 1). Each block uses a linear scan over cu_seqlens_q to identify its owning sequence and exits early when past sequence end. - New Rust module flash_varlen_forward: always-compiled CPU simulation (varlen_attention_cpu) for testing + #[cfg(cuda)] FlashVarlenKernel. - SdpaBackend::VarLen variant added to backend_selector. - 8 new CPU-only tests; total test count: 50. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 01:09:39 +00:00
Omar SobhandClaude Sonnet 4.6	311eb23dbd	feat(perf): GPU perf batch 3 — W4A16 AWQ matmul, FSDP2 hooks, RMSNorm+SwiGLU fused kernel ... W4A16 AWQ on-the-fly dequantize+GEMM (rtx-compress) - `w4a16_matmul.rs`: `AWQQuantizedWeightExt` trait + `matmul_cpu()` — group-aligned inner loop, f64 accumulation, low-nibble-first INT4 unpacking matching mx_kernels.cu - `cuda_kernels/w4a16_gemm.cu`: `w4a16_dequant_gemm` kernel, one thread per (batch, out_col), 8-INT4-per-iteration inner loop with `__ldg()` cache hints, BF16 scale decode, f32 accumulate - `quantization/mod.rs`: exports `w4a16_matmul_cpu`, `AWQQuantizedWeightExt` - Fixed 2 pre-existing pruning compile errors - 12 tests: nibble unpack, identity weights, shape, vs-dequant (tol=1e-3), batch=1, zeros FSDP2 forward/backward hooks (rtx-distributed) - `fsdp2.rs`: `update_local_shard()` on `Fsdp2ShardedParam`; sync `all_gather()` + `reduce_scatter_gradient()` using `ProcessGroup::{all_gather,reduce_scatter}` - `pre_forward_hook()` — all-gathers every param (or copies shard in single-process) - `post_backward_hook()` — reduce-scatters gradients, zero_grad, re-shards cache - `step(optimizer_fn)` — applies optimizer closure to each local shard - `make_fsdp2_module()` top-level factory; `Fsdp2MemoryStats` gains 5 new fields incl. `memory_reduction_ratio ≈ world_size` - 6 new tests (end-to-end training step included); total 444 pass RMSNorm+SwiGLU fused CUDA kernel (rtx-fusion) - `cuda/rms_norm_swiglu_fused.cu`: `rms_norm_kernel` + `rms_norm_swiglu_fused`; shared-mem warp reduction (block_x floats), launch: grid=(batch,1,1), block=(min(hidden,1024),1,1) - `cuda_kernels/rms_norm_fused.rs`: CPU reference `rms_norm_cpu`/`swiglu_cpu`/ `rms_norm_swiglu_cpu`; `#[cfg(feature="cuda")] RmsNormFusedKernel` NVRTC wrapper - `codegen/cubecl.rs`: replaced RmsNorm comment stub with cfg-gated NVRTC dispatch - `backend.rs` + `tensor.rs`: added 15 missing `Backend` trait impls (sin/cos/relu/conv2d/…) that blocked test compilation - `Cargo.toml`: added rtx-fusion to workspace members - 8 new tests (PyTorch-formula verified: x=[1,2,3,4] → [0.365, 0.730, 1.095, 1.461]); total 103 pass Test results: 12 + 444 + 103 = 559 tests, 0 failures Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-27 00:18:20 +00:00
Omar SobhandClaude Sonnet 4.6	a7d9969702	feat(galore2): implement GaLore-2 optimizer with 12 pure-CPU tests ... Adds GaLoreAdamW to crates/training/rtx-transformers — a memory-efficient AdamW variant that reduces optimizer state by projecting gradients to a low-rank subspace and periodically refreshing it via randomised SVD. Key facts verified by tests: - Memory formula: for [rows×cols] param with rank r, GaLore stores: rows*r + 2*r*cols f32 elements AdamW stores: 2*rows*cols f32 elements For [256×256] r=64: ratio=0.375 (62.5% reduction) For [4096×4096] r=128: ratio<10% (>90% reduction) - Subspace refresh triggers when (step - last_refresh) >= update_proj_gap - Momentum inheritance: m_new = new_Q^T @ old_Q @ m_old preserves direction - Small params (< min_param_size=4096 elements) fall back to standard AdamW Files changed: - crates/training/rtx-transformers/src/optimizers/galore.rs (new) - crates/training/rtx-transformers/src/optimizers/mod.rs (mod + re-exports) Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 22:41:42 +00:00
Omar SobhandClaude Sonnet 4.6	082a50e3a0	feat(perf): GPU perf batch 1 — wire GPU execution paths for FP8, FA3, CUDA Graphs, SnapKV ... FP8 GPU FFI (rtx-tensor) - `fp8_cast.rs`: replaced `not_implemented` stubs with real cudarc 0.18.2 PTX launches; `cast_bf16_to_fp8_e4m3` and `cast_fp8_e4m3_to_bf16` now dispatch to NVCC-compiled `fp8_cast.ptx` via `LazyLock` module cache, matching the `inplace_ops` pattern - `build.rs`: `create_dummy_ptx` now emits `fp8_cast.ptx` alongside `element_wise.ptx` so `include_str!` resolves cleanly when NVCC is absent FlashAttention-3 typed kernel launch (rtx-flash-attention) - `flash_v3_forward.rs`: `forward()` now takes typed `CudaSlice<bf16>` Q/K/V/O + `CudaSlice<f32>` LSE buffer; dispatches via `stream.launch_builder` with block_dim=(128,1,1), grid_dim=(ceil(seq_len/64), batch*heads, 1), shared_mem_bytes=0 (PTX metadata-resolved) - `simple.rs`: added `has_flash_v3()` + `flash_attention_v3_forward_raw()` dispatch - `Cargo.toml`: `half` added as optional cuda-gated dependency CUDA Graphs stream threading (rtx-transformers) - `training_loop.rs`: added `cuda_stream: Option<CudaStreamHandle>` field; `set_cuda_backend()` now creates a non-default capture stream; capture step calls real `begin_capture(stream)` + `end_capture(stream)`; added `set_cuda_stream()` override; replay unchanged (no stream needed) SnapKV + prefix cache BatchScheduler wiring (rtx-inference) - `scheduler.rs`: added `prefix_hit_pages: Option<Vec<PageId>>` + `evicted_positions: Vec<usize>` to `SchedulerRequest`; `BatchScheduler` gains `kv_cache` + `snapkv_eviction` fields; `submit_request` does non-blocking `try_lock` prefix lookup; added `notify_prefill_complete` (registers prefix + runs `select_evict_positions`), `set_kv_cache`, `set_snapkv_eviction`, `get_evicted_positions`, `get_prefix_hit_pages` — +5 new integration tests Test results: 22 + 42 + 75 + 102 = 241 tests, 0 failures Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 20:39:58 +00:00
Omar SobhandClaude Sonnet 4.6	1eb89c5b2b	feat(perf): GPU perf batch 1 — FP8, FA3 Blackwell, CUDA Graphs, SnapKV, prefix cache ... Item 1 — CUDA Graphs wiring (rtx-transformers) - Added `enable_cuda_graphs: bool` (default false) + `cuda_graph_warmup_iters: usize` (default 3) to `TrainingConfig` - Wired 3-phase state machine into `training_loop.rs` (warmup → capture → replay) gated on `#[cfg(feature = "cuda")]`; stream plumbing stubbed with TODO pending `CudaStreamHandle` threading Item 2 — FP8 E4M3/E5M2 training infrastructure (rtx-tensor, rtx-transformers) - `fp8_cast.cu`: dual-path CUDA kernels — SM_89+ uses `<cuda_fp8.h>` native `__nv_cvt_*` intrinsics; older SM uses software bit-manipulation fallback - `fp8_cast.rs`: host-side CPU casting + `#[cfg(feature = "cuda")]` GPU stubs - `fp8_gemm.rs`: bit-accurate E4M3 decoder/encoder, BF16 round-trip utils, CPU reference matmul with cuBLASLt GPU path documented inline; 12 unit tests - `training_config.rs`: `fp8_training: bool`, `fp8_e4m3_forward: bool` - `linear.rs` (modular): `fp8_mode: bool` field + forward dispatch stub - build.rs: registers `fp8_cast.cu` alongside existing `element_wise.cu` - 22 FP8 unit tests — all pass Item 3 — FlashAttention-3 Blackwell (WGMMA + TMA + warp specialization) - `flash_attention_v3_forward.cu`: SM_90+ warp-specialised producer/consumer kernel (producer TMA-loads K/V tiles, consumers run WMMA as portable WGMMA proxy); SM_89+ FP8 header path; SM<90 standard FA2-style WMMA fallback - `flash_v3_forward.rs`: NVRTC wrapper (`compile_ptx` via `include_str!`), `FlashV3ForwardKernel::new/is_supported/forward`; 6 unit tests - `backend_selector.rs`: `SdpaBackend::FlashAttentionV3`, `for_compute_capability`, `supports_flash_v3` (major >= 9), FA3 scoring (0.98/0.90/0.70), 2× speedup estimate - `kernels/simple.rs`: `v3_kernel: Option<FlashV3ForwardKernel>` in `FlashCudaKernels` - Fixed pre-existing `Device::Cpu` cfg-gate bug in `tensor/creation.rs` - 8 new FA3 backend tests + 6 kernel unit tests; 50 total pass Item 4 — SnapKV attention-score eviction + prefix caching (rtx-inference, rtx-serving-api) - `prefix_index.rs`: `PrefixIndex` with 8MB Zobrist hash table (Knuth MMIX LCG seed), `compute_hash/lookup/insert/remove/remove_page`; 10 unit tests - `eviction.rs`: `AttentionScoreEviction` struct — `accumulate_scores` + `select_evict_positions` (retain top keep_ratio + last recent_window); 7 unit tests - `types.rs`: `EvictionPolicy::AttentionScore { keep_ratio, recent_window }` + `KvCacheConfig::enable_prefix_caching` - `paged_kv_cache.rs`: `prefix_index: Option<PrefixIndex>` + `lookup_prefix / register_prefix / unregister_prefix_page / prefix_caching_enabled` methods - `config.rs` (serving-api): `enable_prefix_sharing: true` (was false), `snapkv_keep_ratio: 0.6`, `snapkv_recent_window: 32` - Fixed 12 pre-existing test errors (spurious `.await` on sync constructors) - 17 SnapKV/prefix tests pass Total: 918 lib tests pass across rtx-tensor, rtx-flash-attention, rtx-transformers, rtx-inference. Zero new failures. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 19:56:39 +00:00
Omar SobhandClaude Sonnet 4.6	c82d26d6e7	style: rustfmt formatting pass on rtx-tensor and rtx-flash-attention ... Import reordering, long-line reformatting — no logic changes. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 18:11:10 +00:00
Omar SobhandClaude Sonnet 4.6	39b7ef12f4	fix(tests): green-bar rtx-transformers and rtx-distributed test suites ... - Fix 35 doctest failures in Phase 2/3 modules (no_run annotations, missing imports, wrong API calls, Result context issues) - Fix test_validation_framework_creation: assert updated to 1e-3 default - Fix test_report_serialization: replace exact f64 equality with epsilon comparison - Fix rtx-distributed recovery/tests.rs: add missing ProcessGroup import, use recovery_stats().wal_buffer_size instead of private field access All rtx-transformers and rtx-distributed tests now pass. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 17:56:39 +00:00
Omar SobhandClaude Sonnet 4.6	a08adfbf57	fix(gaps): G4 — re-enable all rtx-transformers Phase 2/3 modules (222 compile errors fixed) ... Uncommented all deferred modules in lib.rs and fixed API drift across ~60 files in 9 module groups: continual, curriculum, meta, modular, neural_ode, graph, kan, perceiver, distributed/pipeline_parallelism. Common patterns fixed across modules: - Tensor::randn/zeros/ones([a,b]) → (&[a,b], device)? (slice + Result) - Result<T, TensorError> → .map_err(Into::into)? in TransformerError contexts - Device by value → &device references - &Tensor where Tensor expected → .clone() - tensor.relu()/tanh()/sigmoid() as methods not ops functions - Tensor arithmetic returning Result: (a + b)? → (a.clone() + b)? - shape literals → shape.dims() for Shape type - sum(n) → sum(Some(n)), mean(None) → mean(&[], false) - i64 indices → usize where required - backward(x) → backward(x, None) - Borrow conflicts on self.field resolved by extracting to locals before mut borrow - BatchingStats private fields → pub(crate) - TransformerError::Serialization → ::SerializationError - Add scalar to tensor: (t + 0.1)? → t.add_scalar(0.1)? Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 16:08:02 +00:00
Omar SobhandClaude Sonnet 4.6	228137555f	fix(gaps): G0/G2/G5/G8 — eliminate unimplemented! panics, re-enable rtx-distributed, rtx-tts, fix multimodal forward ... G0 (Critical): Replace 45 unimplemented!() panics across three GPU backends - rtx-backend-cuda: sin/cos/tanh via PTX, relu/sigmoid/leaky_relu/elu via activation.rs, pow/clamp/gt_scalar via unary.rs, var/var_dim host-side, conv2d/max_pool2d/avg_pool2d CPU fallback in new ops/conv.rs; new PTX kernels in element_wise.cu - rtx-backend-rocm: all 15 ops via CPU round-trip (to_vec → compute → from_slice) - rtx-backend-sycl: all 15 ops via CPU round-trip (to_host → compute → from_data) G2 (High): Re-add rtx-distributed to workspace - Vendor 4 minimal RNCCL stub crates at crates/vendor/rnccl/* - Update rtx-distributed RNCCL path deps to point at stubs (../../../../RNCCL/* → ../../vendor/rnccl/*) - Remove rtx-distributed from workspace exclude list, add to members G5 (Medium): Re-enable rtx-tts (213 tests restored) - Fix 15 rtx-nn API drift issues: LayerNorm::new, Conv1d::from_config, Conv1dPadding::Zeros, Dropout::new(p, device), tensor methods (relu/tanh/sigmoid/cat/stack), squeeze(Some(n)), to_vec() turbofish removal, Tensor::randn with &[...] slices G8 (Low): Quantum stubs + multimodal forward bug - rtx-timeseries: remove dead quantum/neuromorphic TODO comment blocks (no module files exist) - rtx-multimodal/fusion/transformer.rs: wire TransformerBlock loop in forward() - rtx-multimodal/fusion/strategies.rs: wire bottleneck_layers loop in forward() - rtx-transformers/architectures/transformer_block.rs: add forward() method (pre-norm residuals; full attention+FFN pending when those sub-layers are wired) Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 13:40:23 +00:00
Omar SobhandClaude Sonnet 4.6	c2f4796871	fix(rtx-flash-attention): sm_120 for Blackwell, robust nvcc path resolution ... build.rs was hardcoded for sm_90 (incorrectly labelled Ada Lovelace/RTX 5090). Fix for RTX 5060 Ti (sm_120, Blackwell): - Auto-detect SM via CUDA_ARCH env var (default sm_120); compute_ prefix derived automatically so compute_120/sm_120 are no longer hardcoded. - nvcc resolution: try PATH first, then CUDA_PATH/bin/nvcc, CUDA_HOME, and common installation prefixes — no longer panics when nvcc is at /usr/local/cuda-13.1/bin but not in $PATH. - PTX version: sm_100+ → .version 8.0 (PTX ISA 8.0 for Blackwell). - No-GPU branch: remove the warning — CPU fallback is valid, there is no reason to warn every build when cuda/metal are intentionally off. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-26 02:33:37 +00:00
Omar SobhandClaude Sonnet 4.6	470fe07144	D350: GPU backbone training benchmark (RTX 5060 Ti) ... Adds d350_gpu_backbone_training — 200-step Adam loop on a 4-regime corpus that prefers Device::cuda(0) and falls back gracefully to CPU. Measured numbers: - CPU (DIM=16): 886 steps/s, MSE 0.2163 → 0.0024 - GPU (DIM=16): 803 steps/s, MSE 0.2163 → 0.0024 GPU is marginally slower at DIM=16 because the SSM scan and conv1d remain on CPU in both paths; cuBLAS only helps the four linear projections, which are tiny at dim=16. The GPU advantage emerges at larger dims (≥256) where the projections dominate. Correctness is identical on both devices. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-23 21:48:04 +00:00
Omar SobhandClaude Sonnet 4.6	f751414c38	D309: fix MambaRecurrence step — conv1d history buffer + new ergonomic API ... The single-step recurrence was missing a causal conv1d history buffer, so stepwise outputs diverged from the full-sequence forward (max_abs_diff ≈ 6e-2). Add `conv_buf: Vec<f32>` to `MambaState` (oldest-first per channel), thread it through `MambaRecurrence::step` so the kernel sees the correct `kc-1` prior x_in values, and shift the buffer after each step. Ergonomic additions: - `MambaRecurrence::init_state()` — zero-initialised state with correct dims - `MambaRecurrence::state_size()` / `d_model()` — accessor methods - `MambaState::hidden()` — slice accessor for the SSM h vector - `MambaState: PartialEq` — enables determinism assertions in tests Both D309 tests now pass: `step_matches_full_forward` and `fresh_state_is_zero_and_deterministic`. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-23 17:41:27 +00:00
Omar SobhandClaude Sonnet 4.6	24bd5cf9dc	feat(layers): add mamba_step, cloned/gated memory updaters, set_encoder_teacher ... Implements four new rtx-transformers layers required by omni-cortex's omni-think crate: MambaRecurrence/MambaState (single-token S6 recurrent step), ClonedMemoryUpdater (linear-tanh cell with SGD + rollout refinement), GatedMemoryUpdater (GRU-style cell with full backward pass), and SetEncoderTeacher (time-parallel set encoder with named_params persistence API). Fixes 10 compile errors in omni-think. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>	2026-06-22 11:38:42 +00:00
osobhandClaude Opus 4.8	0e3ff0a1b9	SMT D319 (rustytorch): learned-[MEM]-query (content-addressable) teacher pool ... Adds an opt-in learned-query attention pool to SetEncoderTeacher (SetEncoderConfig::with_learned_pool): pool = softmax(q_mem·hᵀ)·h instead of the fixed mean/decay pool — Isola's transformer-teacher [MEM] query. New q_mem param (threaded through graph/train_step/run/named_params; only updated in learned-pool mode). Default off → existing teachers byte-unchanged. Finding (test teacher_content_addresses_selective_retrieval): on a selective- retrieval task (signal in one marked token among distractors) BOTH the mean-pool and learned-query teachers recover the marked token to low MSE (~0.0006 / ~0.002) — because the self-attention layer already routes the marked token's signal to every position before the pool. So the mean-pool was NOT the recall bottleneck (correcting the D318 hypothesis): the teacher can content-address; the real recall bottleneck is the recurrent *cell* that imitates it. The learned-query pool is shipped as an equally-capable, faithful-to-the-talk alternative. All 5 teacher unit tests pass; clippy(-D)/fmt clean. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-17 09:31:18 -07:00
osobhandClaude Opus 4.8	ca3f12c6c8	SMT D316 (rustytorch): train_rollout — truncated K-step BPTT for the memory cell ... ClonedMemoryUpdater::train_rollout rolls the cell forward over K inputs on its OWN memory (not teacher-forced) and backprops the accumulated predict-the-future loss through the whole K-step graph — the tape's first multi-step training path, directly optimizing the free-rollout behavior the cell is evaluated on (vs the one-step BC/DAgger paths). To stay within the finite-diff-gated op set (matmul/gelu/add/mul/sub/sum — no tensor concat), W_in is split into its memory rows (applied to M) and input rows (applied to x); the two gradient halves are re-stacked for the Adam update. Non-finite guard + gradient clip as in the other training paths. Test: rollout training reduces the K-step loss (>2x). clippy(-D)/fmt clean. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-17 04:12:46 -07:00
osobhandClaude Opus 4.8	6b7b86fe34	SMT D315 (rustytorch): configurable recency-pool decay (sharp vs smooth oracle) ... SetEncoderConfig gains `recency_decay` (default 0.85) + `with_recency_decay`, threaded into pool_weights. A smaller decay concentrates the recency pool on the most recent tokens — a *sharper* oracle that reacts fast to regime switches (less denoising). Existing teachers default to 0.85 (unchanged). Lets omni-think tune the teacher's reaction speed for non-stationary streams. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-17 03:59:36 -07:00
osobhandClaude Opus 4.8	96809d48a1	SMT D314 (rustytorch): gradcheck tanh/sigmoid + GatedMemoryUpdater (GRU cell) ... F.1 — finite-difference gradchecks for tanh, sigmoid, and the exact gated composite (1-z)*m + z*c (z=sigmoid, c=tanh) added to tape_cpu_gradcheck.rs. All pass (rel-err < 2e-2) — the activation VJPs that already existed are now proven correct on the real CpuBackend (E0 discipline), so the gated cell can rely on them. F.2 — GatedMemoryUpdater: a GRU-style gated recurrent memory cell. z = sigmoid([M‖x]·W_z); c = tanh([M‖x]·W_c); M_t = (1-z)⊙M + z⊙c The convex update is a non-expansion (|M_t| ≤ max(|M_0|, 1)), so free rollout stays bounded with NO clamp — and the learned gate can both jump at a regime switch (z≈1) and hold+denoise in steady state (z≈0), which the residual+leaky ClonedMemoryUpdater cannot. Same method surface (new/step/predict/train_step/ train_step_memory) so the omni-think facade is cell-generic. Tests: gated cell trains (loss drops); 500-step free rollout stays in [-1,1] without a clamp. clippy(-D)/fmt clean. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 21:09:10 -07:00
osobhandClaude Opus 4.8	efd0dc9f9b	SMT D312: timestamp-embedded (recency) mode for the predictive-state teacher ... SetEncoderConfig gains an opt-in `recency` flag (default false = unchanged mean-pool stationary behavior). When set, the encoder adds sinusoidal timestamp embeddings to the token embeddings (so attention can reason about position) and pools with an exponential-decay (recent-weighted) reduction instead of a uniform mean — a recent-window sufficient statistic that tracks non-stationary signals. Both use only existing-VJP tape ops (add + matmul/softmax); no new params. Test: recency-mode teacher trains end-to-end (loss >5x drop). Existing stationary tests unchanged. clippy(-D)/fmt clean. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 18:34:20 -07:00
osobhandClaude Opus 4.8	1d306ba94b	SMT E4: stabilize the cloned updater for DAgger (guard/clip + memory-only step) ... Two robustness additions for on-policy distillation, which trains on the policy's own (clamped, sometimes extreme) visited states: - train_step now skips non-finite-loss steps and clips gradients, so an extreme rollout state can't poison every weight via Adam. - train_step_memory: a memory-only correction at a caller-chosen lr — trains just the recurrence (W_in, W_mem) to map (prev,x)->target_mem, leaving the BC-trained readout (W_read) intact. This is the key to DAgger working: it pulls the free-running memory back toward the oracle trajectory without the readout retraining that otherwise blows the predictions up. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 16:03:46 -07:00
osobhandClaude Opus 4.8	e888f7fe4c	SMT E3b: leaky + clamped memory updater for bounded free rollout ... Make the cloned updater's recurrence M_t = ρ·M_{t-1} + Δ a contraction (ρ=0.9) and clamp the free-running state (|M| ≤ 4). One-step BC has no signal against autoregressive drift (that is E4/DAgger's job); these keep a free rollout numerically bounded — no 1e25 blow-up — so E4 has a stable base to refine. Training is teacher-forced on the bounded oracle memory, so neither ever binds during training. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 14:51:55 -07:00
osobhandClaude Opus 4.8	0ff3fd4e1f	SMT E3b: ClonedMemoryUpdater — recurrent updater cloned from the oracle (no BPTT) ... A trainable recurrent cell g(M_{t-1}, x_t) -> M_t (gelu MLP, residual memory update) plus a predict-the-future readout, trained on Autodiff<CpuBackend> with the same host-Adam pattern as SetEncoderTeacher. train_step is one-step behavioral cloning, TEACHER-FORCED on the oracle memory: each example feeds the oracle's previous memory and regresses (new_mem -> target_mem) + (readout -> next latent) in a single tape graph. No rollout, no backprop-through-time. Off-tape step/predict use a gelu byte-identical to the CPU tape's (tanh approx, sqrt(2/pi), 0.044715) so rollout matches training. The residual update M_t = M_{t-1} + delta lets the recurrence rule stay length-invariant — the basis for extrapolating past the teacher's horizon (E3b generalization test, omni-cortex side). Tests: updater learns a one-step transition (>10x loss drop); off-tape step/predict shapes + finiteness. clippy(-D)/fmt clean. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 14:30:20 -07:00
osobhandClaude Opus 4.8	51a9185056	SMT E3a: expose Mamba's recurrent state as a single-step updater (equivalence-pinned) ... MambaRecurrence reconstructs MambaBlock's selective-scan as an explicit step(state, x_t) -> (state', y_t) recurrence over host f32 weights, with the recurrent state (h, conv-ring) carried between steps. This is the substrate the SMT memory updater is behaviorally cloned on (E3b) — no BPTT, one step at a time. Dimensions are derived from the persisted tensor shapes, so mamba.rs is left untouched (it is near the 1250-LOC cap). silu/softplus are byte-identical copies of the forward's. Equivalence pin (the plan's highest-risk item): stepping a window one token at a time from a zero state reproduces the full-sequence forward EXACTLY — max_abs_diff = 0.0 (bit-identical), on an active_block with wide Delta so the scan genuinely drives the output. Plus a fresh-state determinism test. clippy(-D warnings) + fmt clean on the new module and test. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 11:00:12 -07:00
osobhandClaude Opus 4.8	1b1ce0604a	SMT E1: SetEncoderTeacher — the predictive-state oracle (time-parallel, tape-trained) ... Promotes the proven attention set-encoder smoke test (PR #4, tape_train_smoke.rs) into a reusable rtx-transformers layer for Supervised Memory Training. The teacher maps a window of past tokens to a fixed-size memory M via embed → self-attention → residual → mean-pool → memory projection, with a decoder head supervised by predict-the-future MSE so that M becomes a sufficient statistic of the past. - Trains end-to-end on Autodiff<CpuBackend> (the gradient-correct real backend from PR #3), with a self-contained deterministic host-side Adam. - Time-parallel by construction (one window → one memory, no recurrence to unroll) — this is the oracle whose trajectory the recurrent Mamba updater is later behaviorally cloned against, so the recurrent net never needs BPTT. - Exposes named_params/set_named_params so the caller (omni-think's PredictiveStateTeacher facade) owns safetensors persistence + BLAKE3 sealing. - Adds rtx-backend + rtx-backend-cpu deps (the tape needs a concrete backend). Tests: teacher trains (loss >5x drop), encode is deterministic + fixed-size, params round-trip. fmt + clippy(-D warnings) clean on the new module. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-16 09:37:55 -07:00
osobh	5c2ee63d53	GPU Mamba forward + rtx-tensor device-pointer API + matrix exponential ... Persist accumulated WIP across rtx-tensor / rtx-transformers / rtx-runtime. Two related feature groups: GPU enablement (unblocks the Phase-3 spec §8 CUDA Mamba path): - rtx-tensor: `Tensor::cuda_device_ptr()` + storage GPU-buffer accessors (`raw_ptr.rs`) expose the raw CUdeviceptr that kernel launches need — the "rtx-tensor GPU memory access API" the Mamba CUDA kernels were blocked on. - rtx-transformers: `MambaBlock::forward_cuda` runs the four linear projections through cuBLAS on-device (in/out/x/dt_proj), keeping the selective scan + conv1d + activations on CPU; dispatched automatically from `forward` when on a CUDA device under the `cuda` feature. Updated `mamba_cuda_kernels.rs` accordingly. - supporting plumbing in rtx-runtime stream/bridge and rtx-tensor storage/conversion/concatenation/creation + rtx-flash-attention. Linear algebra (rtx-tensor): - `linalg/matrix_exp.rs`: real matrix exponential via scaling-and-squaring with a degree-13 Padé approximant (Higham 2005), f64 internally. - `complex/linalg.rs`: complex matmul/adjoint, Hermitian eigendecomposition (`ComplexEigenResult`), and the complex matrix exponential, nalgebra-backed. - tests for both. Builds verified on the CPU path (`cargo check -p rtx-tensor -p rtx-transformers -p rtx-runtime -p rtx-flash-attention` clean). The `cuda` feature and the rtx-backend-cuda NVCC build remain unbuildable on this host (CUDA/glibc header mismatch) — pre-existing and unrelated to these changes.	2026-06-15 18:51:30 -07:00
omar sobhandClaude Opus 4.8	bc82745d8f	Phase 4a: trained real SSM beats memoryless baseline on a temporal task ... End-to-end demonstration that training the real selective-scan Mamba buys a genuine temporal-modeling win — the payoff of M1–M3. Task: next-token prediction on a multi-regime sequence (x[t] = μ_regime + noise). Predicting x[t+1] inside a regime requires integrating recent history to average out the noise — a memoryless model can't. Result (same sequence, all three): persistence (memoryless) MSE = 0.0895 untrained Mamba MSE = 0.2167 trained Mamba (400 Adam) MSE = 0.0002 The trained backbone integrates history to de-noise the regime mean, beating the memoryless persistence baseline by ~450× and improving ~1000× over its untrained self. (Single-sequence fit: demonstrates the SSM's temporal-modeling capacity, not held-out generalization.) This replaces the old non-result ("random Mamba 6.5% vs linear 48%") with a real "trained SSM exploits temporal structure" demonstration. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-03 12:10:19 +00:00
omar sobhandClaude Opus 4.8	639937e9f5	Mamba M3.5: canonical numerically-stable initialization ... Plain randn(0,1) init made A=-exp(A_log) and Δ wildly large, overflowing the real exp(Δ·A) scan to NaN (the stub never cared — it discarded these weights). new()/new_seeded() now share a build() with canonical S6 init: - A_log = ln(1..=d_state) ⇒ A = -(1..=d_state), bounded - dt_bias so softplus(dt_bias) ≈ 0.01 (small, stable Δ; near-identity scan at init — intentional for gradient flow) - D = 1, zero conv bias, projections scaled by 1/√fan_in (capped 0.5) This fixes the NaN that broke omni-cortex's d231 action-conditioned predictor training (now green). Seeded determinism preserved. Tests: active_block helper (Δ overridden to ≈0.69) exercises the scan-active regime so the liveness check can observe each weight; the training test asserts a seed-varying backbone weight (conv1d_weight) moves. All 6 selective-scan tests green incl. the finite-diff grad check. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-03 11:30:33 +00:00
omar sobhandClaude Opus 4.8	33cf9bf731	Mamba M3: trainable — Adam loop reduces loss and moves weights ... Demonstrates end-to-end trainability of the real selective scan. A self-contained Adam loop (the rtx-transformers AdamOptimizer has no public gradient setter — the spec permits a bespoke loop) fits a teacher block's output on a fixed input: forward → MSE → analytic backward → Adam step → rebuild. Over 200 steps the loss drops >50% and the backbone weight A_log moves, confirming gradients actually train the model (not just the head). All 6 selective-scan tests green. The production AdamOptimizer can be wired once it exposes a gradient setter; the M2 backward already returns grads in its HashMap shape. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-03 11:15:45 +00:00
omar sobhandClaude Opus 4.8	199130fa7d	Mamba M2: analytic backward (gradient-checked) ... Hand-written VJP of the selective-scan forward (Phase-3 spec §5/§7): MambaBlock::backward(x, d_out) -> per-parameter gradients keyed by persistence name, summed over the batch. Differentiates the scan analytically (reverse-time recurrence over the cached h trajectory) rather than through the immature rtx-tensor autograd tape. Covers every parameter: in_proj, conv1d_weight, conv1d_bias, A_log (via A=-exp(A_log) ⇒ dA_log = dA·A), x_proj, dt_proj, dt_bias, D, out_proj. Adds stable sigmoid_f32 / silu_grad_f32 helpers. New test analytic_gradients_match_finite_differences: on a small well-conditioned instance, ≥30 sampled grad elements across all 9 params match central finite differences within (5e-3 + 5e-2·|fd|). All 5 selective-scan tests green. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-03 11:13:46 +00:00
omar sobhandClaude Opus 4.8	00cd527ed4	Mamba M1: real selective-scan forward (replaces the passthrough stub) ... The MambaBlock forward was a stub — SelectiveScan::forward passed input through, discretize returned zeros, conv1d was a no-op, and B/C were randn per call, leaving conv1d_weight/dt_proj/A_log as dead weights with zero temporal mixing. This implements a genuine S6 selective-scan: - New params: x_proj [d_inner, dt_rank+2*d_state] (data-dependent dt,B,C), dt_bias [d_inner], D [d_inner] (skip). Added to new()/new_seeded() and the persistence contract (persistence_tensors/from_persistence_tensors). - Real forward (CPU f32, looped — backbone is small): in_proj -> causal depthwise conv1d -> SiLU -> x_proj->(dt,B,C) -> delta=softplus(dt.dt_proj +dt_bias) -> A=-exp(A_log) -> sequential scan h=dA.h+dBu, y=sum C.h + D.u -> gate by SiLU(z) -> out_proj. Residual moved OUT (canonical). Numerically-stable silu_f32/softplus_f32 helpers. - The scan runs inline (not via the immature rtx-tensor autograd tape); the analytic backward lands in M2 per docs/phase3_real_ssm_spec.md. New tests/real_selective_scan.rs (4 cases, all green): liveness (each formerly-dead weight now moves the output), causality (no future leakage), seeded determinism, and finite/non-constant output. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>	2026-06-03 11:02:46 +00:00
Omar Sobh	ff81df62ce	chore: commit local changes before node reformat	2026-05-10 14:11:34 -07:00
redclawsystems	ae53983c03	style: cargo fmt --all (18 files) ... Auto-merged by ci-doctor.	2026-05-07 16:30:04 +00:00
osobhandClaude Opus 4.7	3e3e8819f6	rtx-tensor + rtx-transformers: deterministic init + Mamba weight persistence ... Two coordinated additions for the omni-cortex D249/D250 work: rtx-tensor: Tensor::randn_seeded(shape, device, seed) — like randn() but routes through StdRng::seed_from_u64(seed) so two calls with (shape, device, seed) produce bit-exact identical tensors. CPU is the canonical generator; GPU calls go via to_device transfer. Required for reproducible model init. rtx-transformers: MambaBlock gains: - new_seeded(config, device, seed) — every internal weight tensor initialised via randn_seeded() with per-tensor SplitMix64-derived seeds. Two calls with the same (config, device, seed) → bit- exact identical block. - persistence_tensors() -> Vec<(&'static str, &Tensor)> — read-only view of the six (or seven, with conv_bias) internal weight tensors with canonical names (in_proj, conv1d_weight, conv1d_bias?, A_log, dt_proj, out_proj). Stable across versions so safetensors round-trip works. - from_persistence_tensors(config, device, HashMap<String, Tensor>) — rebuild a MambaBlock from a name → tensor map. Validates each tensor's shape against the config and surfaces clean errors on mismatch (so wrong-DIM safetensors loads fail explicitly). These three primitives together give omni-cortex's D249 (operator- seeded determinism) and D250 (safetensors round-trip + BLAKE3 hash pin) clean library hooks without exposing MambaBlock's private fields. Co-Authored-By: Claude Opus 4.7 (1M context) <[email protected]>	2026-05-04 05:26:27 -07:00
Omar Sobh	623e6679d5	fix(rtx-distributed): ring_allreduce delegates to ProcessGroup for NCCL dispatch ... ring_allreduce() contained its own simulation that multiplied each gradient value by world_size (to fake an AllReduce sum), bypassing the ProcessGroup dispatch entirely. This meant the overlapped synchronization path never used NCCL or RNCCL, even when those features were compiled in. Replace the hand-rolled simulation with a call to self.process_group.allreduce(tensor, ReduceOp::Sum) so the overlapped path uses the same backend as synchronize_gradients_sequential. The communication latency sleep is kept for benchmarking purposes. Add test_ring_allreduce_matches_sequential_path to verify both paths produce identical gradient values under CPU simulation. Closes #10	2026-05-02 05:58:10 -07:00
Omar Sobh	16161bb9df	deps: align all 56 per-crate Cargo.toml files to thiserror v2 ... The workspace root was upgraded to thiserror = "2" in an earlier commit, but 56 per-crate Cargo.toml files still independently declared "1.0". These crates do not use workspace.dependencies inheritance for thiserror. All updated to thiserror = "2" for complete fleet alignment. Includes: rtx-backend, rtx-tensor, rtx-losses, rtx-backend-cuda/rocm/metal, all training crates (rtx-auto, rtx-rl, rtx-distributed, rtx-federated, etc.), specialized crates (rtx-science, rtx-platform, rtx-nmf, rtx-neuro-*), production crates (rtx-streaming, rtx-serving-api), and all demo crates. cargo check --workspace: PASSES.	2026-04-26 11:45:14 -07:00
Omar Sobh	a88d254518	rust-scan: edition 2024 clippy clean, workspace lint fixes 2026-04-25	2026-04-25 22:25:49 -07:00