Commit Graph
5 Commits
Author SHA1 Message Date
Omar SobhandClaude Fable 5 9969d8a661 feat(backend): add differentiable index_select / index_add row ops
Add two row-indexing ops along dim 0 to the `Backend` trait so gather /
scatter-add message passing (GNNs, segment softmax, bias tiling) can be
trained through `Autodiff<B>`:

- `index_select(tensor, indices)` — out[i, ..] = tensor[indices[i], ..]
- `index_add(tensor, indices, num_rows)` — out = zeros; out[idx[i], ..] += tensor[i, ..]

They are each other's adjoint, which is what the backward passes use.

Both trait methods have default bodies (host round-trip via to_data /
from_data) so every existing backend keeps compiling and is correct;
backends override with native kernels:

- rtx-backend-cpu: new ops/index.rs (rayon-parallel gather over output
  rows above a size threshold, sequential deterministic scatter-add),
  wired into CpuBackend and CpuBackendF64, with unit tests for D=1/2/3,
  duplicates, untouched rows, empty inputs, bounds panics and adjointness.
- rtx-autograd: Autodiff<B> overrides both ops and records
  IndexSelectBackward / IndexAddBackward (new ops/index.rs); finite-
  difference gradchecks on the real CpuBackend cover repeated-index
  accumulation, untouched-row zero grads, bias tiling via index_select
  of a [1,F] row, and a full per-segment softmax.
- rtx-fusion: forward both ops to the inner backend.

Co-Authored-By: Claude Fable 5 <[email protected]>
2026-08-20 11:05:42 -07:00
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 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
osobh d52d359f52 Fix 3 compile errors: rtx-metal (Linux cfg), rtx-onnx (ort API), rtx-fusion (edition)
- rtx-metal: fix MetalError import in sparse/conversion.rs non-macOS stub
- rtx-onnx: update session.rs and tensor_bridge.rs for ort 2.x API changes
- rtx-fusion: fix Cargo.toml package name
- rtx-hub: fix discovery.rs type mismatch
- Full workspace (80+ crates) now compiles clean on Linux
2026-03-15 17:33:54 -07:00
redclawsystems 4d88dc0584 Initial commit 2026-03-04 00:08:42 +00:00