Skip to content

deepseek v4 fp8_einsum enable#960

Draft
ganyi1996ppo wants to merge 3 commits into
mainfrom
ganyi/deepseek_v4_einsum
Draft

deepseek v4 fp8_einsum enable#960
ganyi1996ppo wants to merge 3 commits into
mainfrom
ganyi/deepseek_v4_einsum

Conversation

@ganyi1996ppo
Copy link
Copy Markdown
Contributor

@ganyi1996ppo ganyi1996ppo commented May 28, 2026

Motivation

depends on ROCm/aiter#3369

Test Plan

Test Result

for small conc like 8, basically same:
PR

============ Serving Benchmark Result ============
Successful requests:                     32        
Failed requests:                         0         
Maximum request concurrency:             8         
Benchmark duration (s):                  63.20     
Total input tokens:                      32768     
Total generated tokens:                  32768     
Request throughput (req/s):              0.51      
Output token throughput (tok/s):         518.49    
Peak output token throughput (tok/s):    544.00    
Peak concurrent requests:                16.00     
Total token throughput (tok/s):          1036.99   
---------------Time to First Token----------------
Mean TTFT (ms):                          291.49    
Median TTFT (ms):                        301.98    
P99 TTFT (ms):                           308.53    
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          15.16     
Median TPOT (ms):                        15.15     
P99 TPOT (ms):                           15.30     
---------------Inter-token Latency----------------
Mean ITL (ms):                           15.14     
Median ITL (ms):                         15.16     
P99 ITL (ms):                            16.16     
==================================================

before

=========== Serving Benchmark Result ============
Successful requests:                     32        
Failed requests:                         0         
Maximum request concurrency:             8         
Benchmark duration (s):                  62.68     
Total input tokens:                      32768     
Total generated tokens:                  32768     
Request throughput (req/s):              0.51      
Output token throughput (tok/s):         522.79    
Peak output token throughput (tok/s):    552.00    
Peak concurrent requests:                16.00     
Total token throughput (tok/s):          1045.57   
---------------Time to First Token----------------
Mean TTFT (ms):                          288.79    
Median TTFT (ms):                        277.55    
P99 TTFT (ms):                           367.73    
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          15.03     
Median TPOT (ms):                        15.03     
P99 TPOT (ms):                           15.15     
---------------Inter-token Latency----------------
Mean ITL (ms):                           15.02     
Median ITL (ms):                         15.04     
P99 ITL (ms):                            15.84     
==================================================

for large conc like 64, small improvement:
this PR

============ Serving Benchmark Result ============
Successful requests:                     256       
Failed requests:                         0         
Maximum request concurrency:             64        
Benchmark duration (s):                  136.14    
Total input tokens:                      2097152   
Total generated tokens:                  262144    
Request throughput (req/s):              1.88      
Output token throughput (tok/s):         1925.52   
Peak output token throughput (tok/s):    3456.00   
Peak concurrent requests:                128.00    
Total token throughput (tok/s):          17329.64  
---------------Time to First Token----------------
Mean TTFT (ms):                          7411.25   
Median TTFT (ms):                        7397.39   
P99 TTFT (ms):                           14209.87  
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          26.02     
Median TPOT (ms):                        25.99     
P99 TPOT (ms):                           32.28     
---------------Inter-token Latency----------------
Mean ITL (ms):                           25.99     
Median ITL (ms):                         19.35     
P99 ITL (ms):                            21.45     
==================================================

before

============ Serving Benchmark Result ============
Successful requests:                     256       
Failed requests:                         0         
Maximum request concurrency:             64        
Benchmark duration (s):                  139.50    
Total input tokens:                      2097152   
Total generated tokens:                  262144    
Request throughput (req/s):              1.84      
Output token throughput (tok/s):         1879.18   
Peak output token throughput (tok/s):    4022.00   
Peak concurrent requests:                128.00    
Total token throughput (tok/s):          16912.66  
---------------Time to First Token----------------
Mean TTFT (ms):                          7486.73   
Median TTFT (ms):                        7466.13   
P99 TTFT (ms):                           14321.71  
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          26.76     
Median TPOT (ms):                        26.78     
P99 TPOT (ms):                           33.13     
---------------Inter-token Latency----------------
Mean ITL (ms):                           26.73     
Median ITL (ms):                         20.10     
P99 ITL (ms):                            21.43     
==================================================

64 decode:
bf16 einsum
image
fp8 einsum
image

prefill 16384:
bf16 einsum
image

fp8_einsum
image

Submission Checklist

cos_mask = s_mask[:, None]
cos = tl.load(cos_ptr + cos_offs, mask=cos_mask, other=0.0)
sin = tl.load(sin_ptr + cos_offs, mask=cos_mask, other=0.0)
even_mask = (rope_d_offs % 2 == 0)[None, :]
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

⚠️ [ruff] <F841> reported by reviewdog 🐶
Local variable even_mask is assigned to but never used

Suggested change
even_mask = (rope_d_offs % 2 == 0)[None, :]
(rope_d_offs % 2 == 0)[None, :]

Signed-off-by: ganyi <ygan@amd.com>
@ganyi1996ppo ganyi1996ppo force-pushed the ganyi/deepseek_v4_einsum branch from 17a2738 to 7503fdd Compare May 28, 2026 08:56
Signed-off-by: ganyi <ygan@amd.com>
Signed-off-by: ganyi <ygan@amd.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant