Commit Graph
7 Commits
Author SHA1 Message Date
osobhandClaude Sonnet 5 4aaa36a57a style: cargo fmt --workspace (whitespace/wrapping only, no semantic change)
Whole-workspace rustfmt pass picked up while iterating on Mamba GPU
backward work. Verified formatting-only via diff sampling; no logic
changed.

Co-Authored-By: Claude Sonnet 5 <[email protected]>
2026-08-10 07:09:36 -07:00
Omar SobhandClaude Sonnet 4.6 e00a6da018 feat(batch14): Muon optimizer, logit processors, per-token activation quantization
CI / Format Check (push) Failing after 11s
Documentation / Build API Documentation (push) Failing after 20s
CI / Clippy Check (push) Failing after 19s
Documentation / Build User Guide (push) Successful in 12s
GPU Tests / Check GPU Availability (push) Successful in 0s
CI / Build (ubuntu-latest) (push) Failing after 54s
CI / Build CPU-Only (Explicit) (push) Failing after 1m8s
GPU Tests / CUDA Tests (12.1) (push) Has been skipped
GPU Tests / CUDA Tests (11.8) (push) Has been skipped
Performance Benchmarks / Run Benchmarks (push) Successful in 7m59s
CI / Build (macos-latest) (push) Failing after 56s
CI / Test (macos-latest) (push) Has been skipped
CI / Test (ubuntu-latest) (push) Has been skipped
CI / Python Bindings (maturin) (macos-latest) (push) Has been skipped
CI / Python Bindings (maturin) (ubuntu-latest) (push) Has been skipped
CI / WASM Build + Size Check (push) Has been skipped
CI / Distributed Training Tests (push) Has been skipped
CI / CI Success (push) Failing after 0s
GPU Tests / Metal Tests (push) Has been skipped
- 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 924c237096 feat(batch12): online quant calibration, draft distillation loss, gradient noise scale
Documentation / Build User Guide (push) Successful in 12s
CI / Build (ubuntu-latest) (push) Failing after 1m3s
CI / Format Check (push) Failing after 19s
CI / Clippy Check (push) Failing after 18s
Documentation / Build API Documentation (push) Failing after 35s
CI / Build CPU-Only (Explicit) (push) Failing after 1m17s
Performance Benchmarks / Run Benchmarks (push) Successful in 2m10s
CI / Build (macos-latest) (push) Failing after 57s
CI / Test (macos-latest) (push) Has been skipped
CI / Test (ubuntu-latest) (push) Has been skipped
CI / Python Bindings (maturin) (macos-latest) (push) Has been skipped
CI / Python Bindings (maturin) (ubuntu-latest) (push) Has been skipped
CI / WASM Build + Size Check (push) Has been skipped
CI / Distributed Training Tests (push) Has been skipped
CI / CI Success (push) Failing after 1s
- 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 d6769ef641 feat(perf): GPU perf batch 4 — SmoothQuant INT8 forward, varlen FA, inference graph capture
GPU Tests / Check GPU Availability (push) Successful in 1s
GPU Tests / CUDA Tests (12.1) (push) Has been skipped
GPU Tests / CUDA Tests (11.8) (push) Has been skipped
Documentation / Build API Documentation (push) Failing after 8s
CI / Build (ubuntu-latest) (push) Failing after 8s
CI / Build CPU-Only (Explicit) (push) Failing after 9s
Documentation / Build User Guide (push) Successful in 10s
Performance Benchmarks / Run Benchmarks (push) Successful in 1m10s
CI / Format Check (push) Failing after 10s
CI / Clippy Check (push) Failing after 16s
CI / Build (macos-latest) (push) Failing after 49s
CI / Test (macos-latest) (push) Has been skipped
CI / Test (ubuntu-latest) (push) Has been skipped
CI / Python Bindings (maturin) (macos-latest) (push) Has been skipped
CI / Python Bindings (maturin) (ubuntu-latest) (push) Has been skipped
CI / WASM Build + Size Check (push) Has been skipped
CI / Distributed Training Tests (push) Has been skipped
CI / CI Success (push) Failing after 0s
GPU Tests / Metal Tests (push) Has been skipped
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 311eb23dbd feat(perf): GPU perf batch 3 — W4A16 AWQ matmul, FSDP2 hooks, RMSNorm+SwiGLU fused kernel
CI / Format Check (push) Failing after 5s
CI / Clippy Check (push) Failing after 7s
CI / Build CPU-Only (Explicit) (push) Failing after 7s
Documentation / Build User Guide (push) Successful in 5s
CI / Build (ubuntu-latest) (push) Failing after 7m36s
Documentation / Build API Documentation (push) Failing after 8s
Performance Benchmarks / Run Benchmarks (push) Successful in 1m51s
CI / Python Bindings (maturin) (macos-latest) (push) Has been skipped
CI / Python Bindings (maturin) (ubuntu-latest) (push) Has been skipped
CI / WASM Build + Size Check (push) Has been skipped
CI / Distributed Training Tests (push) Has been skipped
CI / Build (macos-latest) (push) Failing after 49s
CI / Test (macos-latest) (push) Has been skipped
CI / Test (ubuntu-latest) (push) Has been skipped
CI / CI Success (push) Failing after 1s
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
osobhandClaude Opus 4.6 02d382d5f6 style: apply rustfmt across all crates and demos
Consistent formatting pass: line wrapping, import sorting, trailing
whitespace removal, let-chain indentation, merged derive attributes,
and unsafe block reformatting.

Co-Authored-By: Claude Opus 4.6 (1M context) <[email protected]>
2026-04-12 07:01:58 -07:00
redclawsystems 4d88dc0584 Initial commit 2026-03-04 00:08:42 +00:00