[None][perf] DSV4 o_proj: fuse fp8/UE8M0 quantize into o_a proj and introduce splitk cutedsl o_b proj#15257
Conversation
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
from
a21cd4c to
45a2bc2
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #53573 [ run ] triggered by Bot. Commit: |
|
PR_Github #53573 [ run ] completed with state
|
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
3 times, most recently
from
26b34dc to
970315b
Compare
|
/bot run |
|
PR_Github #56564 [ run ] triggered by Bot. Commit: |
|
PR_Github #56564 [ run ] completed with state
|
|
/bot run |
|
PR_Github #56623 [ run ] triggered by Bot. Commit: |
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
from
970315b to
e00062c
Compare
|
/bot run |
|
PR_Github #56623 [ run ] completed with state
|
|
PR_Github #56628 [ run ] triggered by Bot. Commit: |
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
from
e00062c to
df9dded
Compare
|
/bot run |
|
PR_Github #56632 [ run ] triggered by Bot. Commit: |
|
PR_Github #56628 [ run ] completed with state |
|
PR_Github #56632 [ run ] completed with state
|
|
/bot run |
|
PR_Github #56781 [ run ] triggered by Bot. Commit: |
|
PR_Github #56781 [ run ] completed with state
|
|
/bot run |
|
PR_Github #56820 [ run ] triggered by Bot. Commit: |
|
PR_Github #56820 [ run ] completed with state
|
|
/bot run |
Fuse o_a FP8 and UE8M0 quantization into the CuTe DSL epilogue. Add a TRT-LLM-owned SM100 O_b split-K kernel using DeepGEMM MMA/TMA building blocks, and fold split-partial reduction into the mHC half-MMA and half-FMA paths. Keep the optimization opt-in and cover dispatch, producer correctness, and fused reduction. Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
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>
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
from
1751c6f to
7716c86
Compare
|
/bot kill |
|
/bot run --disable-fail-fast |
|
PR_Github #57970 [ run ] triggered by Bot. Commit: |
|
PR_Github #57971 [ kill ] triggered by Bot. Commit: |
|
PR_Github #57970 [ run ] completed with state |
|
PR_Github #57971 [ kill ] completed with state |
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
|
/bot run --disable-fail-fast |
|
PR_Github #57991 [ run ] triggered by Bot. Commit: |
|
PR_Github #57991 [ run ] completed with state
|
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
|
/bot run --disable-fail-fast |
|
PR_Github #58108 [ run ] triggered by Bot. Commit: |
|
@mingyangHao Should this PR target main instead? |
|
@longlee0622 Yes, I will change the target of this PR later. |
|
PR_Github #58108 [ run ] completed with state
|
Signed-off-by: Mingyang Hao <200044211+mingyangHao@users.noreply.github.com>
mingyangHao
force-pushed
the
user/mingyangh/dsv4-oproj-fp8-fuse
branch
from
af97041 to
c797edf
Compare
|
/bot run --disable-fail-fast |
|
/bot run |
|
PR_Github #58231 [ run ] triggered by Bot. Commit: |
|
PR_Github #58231 [ run ] completed with state |
|
We no longer use the ds v4 side branch. Closing this PR. Please prepare PR against main. |
3 participants
Fold the o_lora bf16->fp8 + 1x128 packed-UE8M0 quantization into the o_a blockwise GEMM epilogue (single-pass, tree-reduce amax), so o_a emits fp8+sf directly and o_b (DeepGEMM 1d1d) consumes it without the separate fp8_quantize_1x128_packed kernel and the bf16 o_lora HBM round-trip.
Validated byte-exact (G=1/G=16/tail-M/negative guards) and net-positive (+1.6us/o_a in decode, nsys hardware-timestamp measured). Default OFF -> zero functional change unless DSV4_FUSE_OPROJ=1.
@coderabbitai summary
Perf
O_a proj
O_b proj
Test Coverage
PR Checklist
Please review the following before submitting your PR:
PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
If PR introduces API changes, an appropriate PR label is added - either
api-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin the PR title.Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.
GitHub Bot Help
To see a list of available CI bot commands, please comment
/bot help.