Commit History
-
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
-
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
-
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