ViperEkura/AstrAI
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Commit Graph

  • 432dfec3c2 refactor: collapse fp8 recipe hierarchy and state property layers main ViperEkura 2026-08-29 20:03:32 +08:00
  • e3c3e28a11 docs: fix stale developer documentation claims ViperEkura 2026-08-29 14:34:07 +08:00
  • a7d4cb25c5 docs: scope trainer environment variables per job ViperEkura 2026-08-29 14:01:40 +08:00
  • 0546331637 fix: skip gradient checkpointing log when no modules configured ViperEkura 2026-08-29 11:18:39 +08:00
  • 962c10c52b perf: fold the delayed-scaling ring update into the quantize kernel ViperEkura 2026-08-29 09:16:22 +08:00
  • 1cf7d6c76b perf: fill steady-state decode input ids via d2d copy ViperEkura 2026-08-28 22:55:35 +08:00
  • 36e39496d4 perf: vectorize tiled fp8 transpose quantize and arm amax via memset ViperEkura 2026-08-28 22:39:02 +08:00
  • a1a1a6bf0f perf: pack gqa q-heads per prefill block to reuse kv tiles ViperEkura 2026-08-31 13:38:48 +08:00
  • 7dd184a4e5 refactor: split fp8 gemm device code into layered headers ViperEkura 2026-08-28 17:29:13 +08:00
  • bf239d194c refactor: dedupe fp8 kernel helpers and trim comments ViperEkura 2026-08-28 16:45:21 +08:00
  • 8a353117ea perf: transpose-quantize backward operands to route all gemms nt ViperEkura 2026-08-28 16:11:22 +08:00
  • 04a8e2517a perf: split fp8 gemm cta plan by operand layout ViperEkura 2026-08-28 12:44:34 +08:00
  • c4f7f82725 refactor: drop dead fp8 gemm knobs and dedupe ring depth logic ViperEkura 2026-08-28 01:55:30 +08:00
  • fac9d07542 refactor: fp8 gemm policy layering with swap-NN and narrow-N ctas ViperEkura 2026-08-28 01:21:55 +08:00
  • bbb2d95256 fix: report true rank in logs via dist-aware helpers ViperEkura 2026-08-27 13:39:38 +08:00
  • 1c04a0b9fa refactor: move runtime parsers into scripts/docker ViperEkura 2026-08-27 12:33:32 +08:00
  • ba8beb81be fix: serve and train reuse the built image and expose GPUs correctly ViperEkura 2026-08-27 12:22:02 +08:00
  • 7cfcc6c86a Merge pull request #26 from Cytosine-code/fix/cuda-kernel-install ViperEkura 2026-08-27 05:07:44 +08:00
  • f4c44ebf1c docs: correct fp8 kernel descriptions in cuda_kernels guide ViperEkura 2026-08-27 05:07:41 +08:00
  • f86f605f5f fix: make CUDA kernel installation reliable Cytosine 2026-08-27 02:34:16 +08:00
  • 4c82d5d84b perf: dispatch non-128-divisible shapes to the 64x64 cta ViperEkura 2026-08-26 19:02:50 +08:00
  • 76aa4edc9f perf: skip grad-bias reduce for bias-free linears ViperEkura 2026-08-26 18:50:54 +08:00
  • 6354dbe8bc perf: interleave prefetch into the mma phase ViperEkura 2026-08-26 16:36:43 +08:00
  • f45230fb2c perf: pair b fragments into ldmatrix x4 loads ViperEkura 2026-08-26 15:43:09 +08:00
  • d4534be8ca perf: single-wave big cta takes a stage deeper pipeline ViperEkura 2026-08-26 15:00:02 +08:00
  • 1d57588d27 fix: free default ubuntu uid/gid before creating astrai user ViperEkura 2026-08-26 15:00:15 +08:00
  • a92bf79295 perf: fast interior loop on the big cta and fused epilogue bias ViperEkura 2026-08-26 14:52:06 +08:00
  • f7d96455a5 perf: base-pair fragment addressing and full-ring small cta ViperEkura 2026-08-26 14:10:14 +08:00
  • 8cfe7536ea perf: retarget fp8 gemm tile dispatch and prune dead configs ViperEkura 2026-08-26 08:24:51 +08:00
  • a8b63fa362 chore: untrack fp8 sweep tool ViperEkura 2026-08-26 07:22:07 +08:00
  • 4d6a244093 perf: fp8 batched gemm and measured dispatch table ViperEkura 2026-08-26 06:45:14 +08:00
  • 01eacbde51 perf: speed up fp8 gemm across small and large shapes ViperEkura 2026-08-25 22:24:51 +08:00
  • 057c0d33df refactor: split quantize into multi-type primitive ViperEkura 2026-08-25 20:07:17 +08:00
  • 3e57cc8069 feat: fp8 backward reuses pre-quantized operands ViperEkura 2026-08-25 17:24:20 +08:00
  • 2eeac02d70 perf: overlap direct crosswise loads with fp8 gemm mma ViperEkura 2026-08-25 14:22:51 +08:00
  • 5e76fbd1bf perf: fp8 rings, lean autocast, gemm staging ViperEkura 2026-08-25 14:22:51 +08:00
  • 4dc5e923e0 perf: finalize fp8 scale rings inside quantize kernels ViperEkura 2026-08-25 11:07:31 +08:00
  • 998b443aa3 refactor: dedupe fp8 meta state into per-operand rings ViperEkura 2026-08-24 21:19:45 +08:00
  • cebdd45d3a perf: batch crosswise stage loads in fp8 gemm ViperEkura 2026-08-24 19:49:10 +08:00
  • 7da1439c9e feat: static fp8 weights and bias with fused epilogue ViperEkura 2026-08-24 19:23:11 +08:00
  • 29e5f571af fix: own fp8 linear backward via autograd Function ViperEkura 2026-08-24 19:03:01 +08:00
  • 74e694921c perf: speed up fp8 gemm tiles and scheduling ViperEkura 2026-08-24 18:29:55 +08:00
  • d5067af064 refactor: harden param PODs and CUTLASS-style fp8 layout tags ViperEkura 2026-08-24 15:27:52 +08:00
  • f6db546578 fix: zero-init AttentionParams in pure C tests ViperEkura 2026-08-24 14:55:32 +08:00
  • 31ca357c61 refactor: namespace csrc kernels and extract common helpers ViperEkura 2026-08-24 14:49:55 +08:00
  • 34471252ab perf: pipeline fp8 gemm fragment loads and pack bf16 epilogue ViperEkura 2026-08-23 21:31:20 +08:00
  • aa08479285 perf: widen fp8 gemm tile and load fragments with ldmatrix ViperEkura 2026-08-23 21:12:56 +08:00
  • 4b10d3ca37 perf: vectorize fp8 quantize and swizzle gemm smem ViperEkura 2026-08-23 20:23:57 +08:00
  • 2bc4d2b8a8 refactor: unify kernel module loading and packaging ViperEkura 2026-08-23 15:57:04 +08:00
  • 4244df2785 perf: pure FP8 fwd/bwd and lean non-transposed GEMM - drop the fused kernel; forward/backward are quantize + a pre-quantized GEMM - rename module fp8_mm -> fp8_ops (mm.cu -> ops.cu) - kernels/launchers fp8_gemm_kernel / launch_fp8_gemm; drop PqTraits/gather_trans/pack_fp8x4_vector - remove the in-kernel transposed-operand branches (TransA/TransB) - backward: quantize g once (amax_g here), explicit fp8 transposes, fast non-transposed GEMMs (dX = g@w^T, dW = g^T@x^T) - each pass uses a single FP8 format (E4M3 fwd / E5M2 bwd) ViperEkura 2026-08-23 15:37:42 +08:00
  • a29bdfae46 refactor: rework attention backend resolution ViperEkura 2026-08-23 14:47:02 +08:00
  • 10fec8dca1 docs: fix stale docs and align with code ViperEkura 2026-08-23 14:23:28 +08:00
  • 75304d084d refactor: unify CUDA skip guards in tests ViperEkura 2026-08-22 21:08:06 +08:00
  • 16a55bb474 refactor: reorganize CUDA kernels into per-family directories ViperEkura 2026-08-22 20:39:10 +08:00
  • cb21af38ba feat: unify Docker serving configuration in YAML ViperEkura 2026-08-21 23:16:08 +08:00
  • dcc96de12a test: refactor tests and fix inference edge cases ViperEkura 2026-08-21 22:59:15 +08:00
  • 7d27f3e078 feat: load HuggingFace checkpoints via key/config conversion ViperEkura 2026-08-20 11:34:48 +08:00
  • 84753d3e08 refactor: deduplicate and restructure test suite ViperEkura 2026-08-20 01:52:49 +08:00
  • 53a7149577 feat: add distributed rank to logs ViperEkura 2026-08-20 01:37:00 +08:00
  • c79d34eee1 refactor: simplify training and inference interfaces ViperEkura 2026-08-19 20:14:53 +08:00
  • 398e8a3ea3 refactor: deduplicate low-risk code paths ViperEkura 2026-08-19 16:17:40 +08:00
  • f252af495c refactor: remove dead code and deduplicate scheduler setup ViperEkura 2026-08-19 14:58:37 +08:00
  • 00c2c80c8f feat: unify Docker training configuration in YAML ViperEkura 2026-08-19 14:04:35 +08:00
  • a6c6a54ace feat: read host training vars from YAML infra section ViperEkura 2026-08-19 01:26:50 +08:00
  • 3d3ea47d37 refactor: standardize packed 3d inference ViperEkura 2026-08-19 00:30:14 +08:00
  • f7f14d0e5f test: add online GRPO end-to-end training test ViperEkura 2026-08-19 00:07:16 +08:00
  • 7580d80d45 perf: accelerate FP8 backward with fused fast kernel ViperEkura 2026-08-18 23:43:06 +08:00
  • cb51a3587b perf: optimize fused FP8 GEMM kernel ViperEkura 2026-08-18 19:56:15 +08:00
  • 1bcd8f53ab perf: precompute ragged Q tile scheduling ViperEkura 2026-08-16 23:32:46 +08:00
  • 0d0dc64884 docs: explain extension layer boundaries ViperEkura 2026-08-16 21:30:03 +08:00
  • 6ac3b51496 refactor: separate extension ops and backends ViperEkura 2026-08-16 21:15:52 +08:00
  • 3406157431 refactor: standardize packed 3d inference ViperEkura 2026-08-16 13:15:23 +08:00
  • 0dd9a417b7 refactor: separate KV token address resolution ViperEkura 2026-08-15 22:59:35 +08:00
  • a01c1fd427 perf: bypass L1 for attention tile loads ViperEkura 2026-08-15 21:23:20 +08:00
  • f8d9ab344d refactor: remove unused streaming dataset ViperEkura 2026-08-15 20:55:08 +08:00
  • 3fb4b8ab13 perf: use int32 paged KV indices ViperEkura 2026-08-15 13:12:16 +08:00
  • b5afe3d7a4 perf: optimize small-head causal prefill ViperEkura 2026-08-14 23:25:49 +08:00
  • 69f35c46e0 fix: quantize amax from raw values, not scaled fp8 values - amax for delayed scale was the quantized max (always ~448), so scale collapsed to 1 - this made fp8 gradients diverge (cosine 0.05) and training stall - stop w/x transpose-quantize amax from polluting the grad scale ViperEkura 2026-08-14 12:38:07 +08:00
  • 0378e62e17 refactor: split fp8 into fp8_ops adapter and fp8 policy module - fp8_ops is the only module touching the pybind (kernel interface) - fp8.py keeps scaling state, delayed amax and aten::linear dispatch - remove circular imports between old fp8_ops/fp8_state/fp8_dispatch ViperEkura 2026-08-14 12:22:54 +08:00
  • 5244f1a8fc feat: add te-style scaled fp8 training via fp8_autocast - per-tensor scales applied inside cublasLt via A_SCALE/B_SCALE - delayed scaling: weight amax history ring, refresh every 16 steps - quantize kernels emit atomic amax, device-side scale updates - fp8_autocast context toggles aten::linear dispatch like torch.autocast - fallback to bf16 when M/N not 16-aligned (fp8 gemm constraint) - x/g scales delayed one step, reuse free atomic amax (no abs/max reduce) ViperEkura 2026-08-14 12:14:04 +08:00
  • 5104638447 perf: use fp8 tensor-core gemm in linear backward - dX/dW run as fp8 cublasLt gemms via fused transpose-cast - shared (m,k,n) algo cache for fwd/bwd, mutex-protected - bias add in-place on bf16 output, drop output copy ViperEkura 2026-08-14 10:43:03 +08:00
  • a711d9f478 perf: eliminate gemm output transpose via A/B swap ViperEkura 2026-08-14 01:41:45 +08:00
  • 15862d4b56 perf: fuse fp8 linear fwd and bwd into single kernel calls ViperEkura 2026-08-14 01:24:37 +08:00
  • f9efb705b8 perf: output fp8 gemm in bf16 instead of fp32 ViperEkura 2026-08-14 01:08:26 +08:00
  • c6a82a5029 refactor: align linear backward dtype with weight ViperEkura 2026-08-14 00:58:53 +08:00
  • a5b238dd86 feat: add fp8 training via cublasLt dispatch ViperEkura 2026-08-14 00:37:52 +08:00
  • da6d94492d fix: parse yaml floats with yaml 1.2 schema ViperEkura 2026-08-13 23:28:06 +08:00
  • 71b6e3aaaf feat: rework docker workflow for gpu-first training ViperEkura 2026-08-13 22:47:34 +08:00
  • f95722a277 feat: add containerized training workflow ViperEkura 2026-08-12 20:20:07 +08:00
  • 9f48cb8928 refactor: streamline Q block mapping ViperEkura 2026-08-10 08:39:43 +08:00
  • 9b58fef222 refactor: extract QTileMapper for prefill tile dispatch ViperEkura 2026-08-09 23:18:00 +08:00
  • c5fba9c238 perf: flatten paged prefill tile dispatch ViperEkura 2026-08-09 23:12:01 +08:00
  • cd31f1f62f refactor: tidy attention params and launcher interfaces ViperEkura 2026-08-09 20:52:06 +08:00
  • a5a3cc1fc2 refactor: unify attention param field names ViperEkura 2026-08-09 20:22:40 +08:00
  • d565d44c43 fix: harden attention kernel boundaries - fix scalar prefill head_dim=32 out-of-bounds via G=4 dispatch - fix MMA decode 4D mask head indexing and invalid-row mask access - add q_head/kv_head divisibility and head-dim contiguity checks - validate split-KV scratch and decode out_buf layout in bindings - set max dynamic shared memory for scalar decode D=256 - cover scalar prefill D=32 in pure C test ViperEkura 2026-08-09 14:52:27 +08:00
  • 596c35fd71 fix: report gradient snr in db ViperEkura 2026-08-09 13:40:27 +08:00
  • 47b3ed4e44 feat: propagate attention backend across scheduler threads ViperEkura 2026-08-09 13:29:30 +08:00
  • c1d05ae11d perf: benchmark decode via real inference engine ViperEkura 2026-08-09 11:47:14 +08:00
  • cf4f5ab9f6 feat: add persistent DataLoader workers ViperEkura 2026-08-09 11:38:50 +08:00
  • 3416f98c58 fix: wire benchmark cache selection ViperEkura 2026-08-09 10:56:05 +08:00