microsoft/onnxruntime

Pull Requests Commits

Fix Windows CUDA build errors and warnings

tianleiwu committed 26 days ago

5cd90274

QMoE DSV4 DeepGEMM: make the expert count a template argument (32 and 64)

tianleiwu committed 54 days ago

87d2b8f0

LightningIndexer: a tensor-core scorer bounded from device memory

tianleiwu committed 55 days ago

ab422fe8

LightningIndexer: bound the fused scorer by rows, and default to cuBLAS

tianleiwu committed 55 days ago

8cfbab39

DSV4SwiGLU: make `up` optional and slice one fused projection internally

tianleiwu committed 56 days ago

ad950de2

Apply lintrunner, and stop clang-format breaking the FlashMLA launch

tianleiwu committed 56 days ago

11d86b95

PagedAttention: make the latent split's occupancy target measurable

tianleiwu committed 56 days ago

f8304035

LightningIndexer: bound the cache and score passes by the live rows

tianleiwu committed 56 days ago

9e9dfb2e

PagedAttention: bound the latent kernel by the live selection extent

tianleiwu committed 56 days ago

d3b6faeb

LightningIndexer: clamp startup-captured graphs

tianleiwu committed 57 days ago

63ba33da

LightningIndexer: score only the cache rows the step can reach

tianleiwu committed 57 days ago

e090fbf3

[CUDA] Let the custom all-reduce be restricted by message size, and trace its routing

tianleiwu committed 58 days ago

56544eb8

[CUDA] Widen the ArgMax block when the row count cannot fill the GPU

tianleiwu committed 58 days ago

6f9bea2f

Correct the measured verdict on the fp4 MoE GEMV register commit

tianleiwu committed 59 days ago

45e5cae9

[CUDA] Use FP8 e4m3 weights for the DSV4 DeepGEMM MoE path

tianleiwu committed 59 days ago

97aba237

Wire FlashMLA into PagedAttention's LATENT path behind ORT_USE_FLASH_MLA

tianleiwu committed 59 days ago

c69afdbb

Add the ORT-side FlashMLA launcher (still no call site)

tianleiwu committed 59 days ago

c933d07a

Vendor FlashMLA (bf16, headdim 576) into the CUDA EP build

tianleiwu committed 59 days ago

b6426177

Revert "Experiment: CtaM>1 rows per block in the MoE GEMV (LOSS, not wired up)"

tianleiwu committed 59 days ago

3132a5c3

Experiment: CtaM>1 rows per block in the MoE GEMV (LOSS, not wired up)

tianleiwu committed 59 days ago

e8c0412c

Decode fp4 weights one word at a time in the MoE GEMV inner loop

tianleiwu committed 59 days ago

eaba2610

[CUDA] Skip padded rows in the DeepGEMM MoE pack and SwiGLU kernels

tianleiwu committed 59 days ago

395c8b23

style(cuda): apply clang-format to DSV4 DeepGEMM and MoE sources

tianleiwu committed 59 days ago

527b7c34

Parallelize CUDA ArgMax/ArgMin over the reduction axis

tianleiwu committed 59 days ago

83744222

Revert "[experiment, LOSS] MoE decode GEMV: cap registers via __launch_bounds__"

tianleiwu committed 59 days ago

c2456927

[experiment, LOSS] MoE decode GEMV: cap registers via __launch_bounds__ to buy occupancy

tianleiwu committed 59 days ago

37768cef

refine deepgeem and update doc

tianleiwu committed 59 days ago

84643431

[CUDA] Restore fast BF16 MXFP4 MoE GEMV decode

tianleiwu committed 59 days ago

6df55927

[CUDA] Speed up the NVFP4 QMoE decode GEMV and enable it for MTP verify (#31159)

tianleiwu committed 59 days ago

1d9837ae

tianleiwu committed 59 days ago

407b0d3a

Older