microsoft/DeepSpeed

Pull Requests Commits

refactor(interface): portable segKI backend interface — grouped_linear dispatcher + capabilities

delock committed 7 days ago

0a41fc28

cleanup: drop the dead non-graph decode_step; guard lm_head install against TP

delock committed 7 days ago

c8e3bf63

docs: full temperature=1 comparison table (RL/OPD scenario)

delock committed 7 days ago

52b560a9

docs: sampling (temp=1) is free at all batch sizes; sampler distribution verified

delock committed 7 days ago

fd9275a2

docs: full comparison table (eager/graph/segKI/vLLM, bs 1..32)

delock committed 7 days ago

35c50bbe

docs: sub-16 grouped GEMM wins (one tile per warp)

delock committed 7 days ago

fc6107fa

perf: lower the grouped-GEMM gate to all decode batches (2..32)

delock committed 7 days ago

c1a488b3

perf(kernels): one wmma tile per warp (8 warps, TPB 256)

delock committed 7 days ago

585a27a2

docs: bs=16 comparison (112% of vLLM at the grouped-GEMM boundary)

delock committed 7 days ago

9c1acf34

docs: vLLM small-batch baselines (94%/88% at bs=1/4)

delock committed 7 days ago

246f5d86

docs: record lm_head + in-place recurrent results (neutral; overlap lesson)

delock committed 7 days ago

f68cef26

perf: gate the in-place GDN recurrent step to batches >= 16

delock committed 7 days ago

4bdf5d9c

perf(kernels): in-place GDN recurrent step — half the state traffic

delock committed 7 days ago

4af76e6f

perf(kernels): route lm_head decode batches through the grouped kernel

delock committed 7 days ago

7347a622

docs: memcpy root cause — GDN recurrent-state copy is HF's cudagraph design

delock committed 7 days ago

43d5e90f

docs: record cp.async results (bs=32 = 136% of vLLM)

delock committed 7 days ago

25fbd655

perf(kernels): cp.async double-buffered staging with dynamic smem

delock committed 7 days ago

3c1830db

docs: record K-fusion and deeper-tile results (bs=32 = 117% of vLLM)

delock committed 8 days ago

ee15a64f

perf(kernels): deeper grouped-GEMM tiles (KT 128, N 64, 2 wmma tiles/warp)

delock committed 8 days ago

fecd7f1c

perf(kernels): fuse K-side norm + partial RoPE + cache write into decode_attn

delock committed 8 days ago

c2f6c68e

docs: record GEMV vectorization results (+1.2% at bs=1)

delock committed 8 days ago

ead1b813

perf(kernels): vectorized 16-byte loads in the b=1 GEMV kernels

delock committed 8 days ago

8ade5036

docs: record vectorized-staging results (bs=32 beats vLLM)

delock committed 8 days ago

df136a1b

perf(kernels): vectorized 16-byte staging in the grouped GEMM kernels

delock committed 8 days ago

2da647d2

fix: restore the qkv split lines dropped from the b=1 triple_gemv branch

delock committed 8 days ago

c461b841

docs(reference): explicit specs for the grouped-GEMM kernels

delock committed 9 days ago

cc3e8042

perf(kernels): WMMA grouped GEMM, full 3D dispatch, strided decode_attn

delock committed 10 days ago

f87077b0

perf: gate grouped GEMM to batches 16..32

delock committed 10 days ago

1924a0df

feat(kernels): grouped GEMM for decode batches 2..32 — one launch, no weight copies

delock committed 10 days ago

d9abdac2

fix: strided decode_attn query + transformers 5.19 GDN scan fallback

delock committed 10 days ago

c526daf6

Older