Commit 1751c6f
committed
[None][perf] Replace DSV4 O_b split-K with CuTe DSL
Replace the dedicated C++/THOP O_b split-K GEMM with a CuTe DSL implementation built on the persistent SM100 block-scaled GEMM. Consume packed K128 UE8M0 scales directly, issue runtime-N tcgen05 MMA, and preserve BF16 split-major partials for the existing mHC fused reduction.
Tune the packed-scale, dynamic-N, scheduler, and epilogue paths against current DeepGEMM. Dispatch SK2 through M=128, use CuTe SK1 for supported M>=256 shapes, and retain DeepGEMM as the unsupported-shape fallback. Add tactic, dispatch, packed-scale, and split-partial correctness coverage while removing the obsolete C++ kernel and registration.
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>1 parent f2f16e7 commit 1751c6f
13 files changed
Lines changed: 1916 additions & 1388 deletions
File tree
- cpp/tensorrt_llm
- kernels/mhcKernels
- thop
- tensorrt_llm/_torch
- custom_ops
- cute_dsl_kernels/blackwell
- blockwise_gemm
- modules
- tests/unittest/_torch/attention/sparse/deepseek_v4
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
19 | 19 | | |
20 | 20 | | |
21 | 21 | | |
22 | | - | |
| 22 | + | |
23 | 23 | | |
24 | 24 | | |
25 | 25 | | |
| |||
28 | 28 | | |
29 | 29 | | |
30 | 30 | | |
31 | | - | |
| 31 | + | |
32 | 32 | | |
33 | 33 | | |
34 | 34 | | |
Lines changed: 0 additions & 208 deletions
This file was deleted.
0 commit comments