Commit History
-
Merge pull request #22 from MotifTechnologies/jangwoong/mla-rope-fa4-port
Jangwoong KimUnverified signatureoptimize: MLA RoPE + FA4 wrapper Triton kernels (port from llm-training)5adea7d -
bench: MLA RoPE fused vs vanilla sweep
Jangwoong Kim +1: Claude Opus 4.6 (1M context)Adds cases/mla_rope.py (DiffCase) and registers it in run_cases.py with mla_rope-specific correctness configs and dummy `dim=[0]` (MLA head dims are fixed by motif3: H_q=80, H_kv=16, D_nope=128, D_rope=64, D_v=128). Also adds bench_mla_rope.yaml (single-GPU TrainJob) and a small mkdir in run_cases.py so triton's perf_report can savefig into a fresh path. B200, motif3 config (bs=8, sl=4096): fwd 6.90x, bwd 5.53x. Run: `python run_cases.py --case mla_rope --dtype bf16 --save-path ...` Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>536f0b2
-
style: fix yapf/isort/clang-format for CI --all-files
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>9dcee96 -
feat: add RMSNorm benchmark scripts and K8s job
- run_rms_bench.py: custom RMS benchmark for dims 512/1024/4096/16384 - run_bench.sh / run_and_wait.sh: K8s job apply + log streaming helpers - benchmark_rms_optim.yaml: TrainJob manifest for B200 benchmarks - bench_framework.py: minor fixes for benchmark runner Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>a5e85e1 -
style: apply yapf + isort formatting
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>60615a0
-
feat: replace triton do_bench with torch.profiler for kernel timing
Switch from triton.testing.do_bench to torch.profiler-based CUDA kernel time measurement. do_bench includes PyTorch CPU overhead in timing, leading to inaccurate results. torch.profiler sums only actual CUDA kernel durations for precise measurement. - Add profile_bench() with per-kernel breakdown and bandwidth output - Add GB/s and SpeedUp(ratio) columns to CSV reports - Fix grouped_mul_poly import path (activation.grouped_poly_norm) - Add docs/benchmark.md for benchmark system documentation - Update CLAUDE.md build system description Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>7d51e61
-
style: apply yapf, isort, and clang-format
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>6436ad6 -
fix: rename stale references and clean up Triton remnants
- Fix broken imports in benchmarks (grouped_fused_mul_poly_norm → fused_mul_grouped_poly_norm) - Rename GroupedTritonModule → GroupedCUDAModule in benchmarks - Rename _run_triton → _run_cuda in tests, update docstrings - Remove stale Triton comments in tests - Fix return type annotations in __init__.py (None → torch.Tensor) - Update README with new function name - Remove TRITON_PRINT_AUTOTUNING env var from profile_bwd.py Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>5a9d09d
-
feat: add GroupedFusedMulPolyNorm Triton kernel for MoE models (#16)
TaehyunKim +2: Claude Opus 4.6, github-actions[bot]Unverified signature* feat: add GroupedFusedMulPolyNorm Triton kernel for MoE models Fuses the full PolyNorm computation into two Triton kernels (fwd + bwd) with per-expert weights/bias and in-kernel binary search for expert mapping. Includes benchmarks, tests, and README documentation with B200 results. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add built binary [skip-build] --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>e195bbb
-
Revert "fix typo in readme (#7)" (#8)
TaehyunKimUnverified signatureThis reverts commit 2d926c3d4773b82c5a5f03904780a40d25a1afee.ddd119c -
fix typo in readme (#7)
TaehyunKim +1: github-actions[bot]Unverified signature* fix typo in readme Signed-off-by: ca1207 <ca1207zzz@gmail.com> * Add built binary [skip-build] --------- Signed-off-by: ca1207 <ca1207zzz@gmail.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>2d926c3
-
Fix fused add rms norm (#4)
TaehyunKim +1: TaehyunKimMotif Unverified signature* make fused add rms norm return 2 outputs * add builds * add plot --------- Co-authored-by: taehyun <taehyun.kim@motiftech.io>a1e5ca8
-
Add fusion (#3)
TaehyunKim +1: TaehyunKimMotif Unverified signature* vectorize and optimized block reduce * add benchmark test (w/o readme update) * implemented fused_mul_poly_norm Signed-off-by: taehyun <taehyun.kim@motiftech.io> * add_rms_norm added * deleted backward pass on fused add rms norm, split test and benchmarks Signed-off-by: taehyun <taehyun.kim@motiftech.io> * refactored benchmarks * add readme * fix readme * add build * fix readme * fix readme2 * add mi250 results * highlight used our kernel for baseline in fused performance * applied yapf --------- Signed-off-by: taehyun <taehyun.kim@motiftech.io> Co-authored-by: taehyun <taehyun.kim@motiftech.io>e5e2eeb