-
Notifications
You must be signed in to change notification settings - Fork 681
[Cherry-Pick][Optimization]Decode attention support(#5767) #5833
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[Cherry-Pick][Optimization]Decode attention support(#5767) #5833
Conversation
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## release/online/20251131 #5833 +/- ##
==========================================================
Coverage ? 58.39%
==========================================================
Files ? 324
Lines ? 39333
Branches ? 5931
==========================================================
Hits ? 22969
Misses ? 14522
Partials ? 1842
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Sentry. 🚀 New features to boost your workflow:
|
|
Thanks for your contribution! |
| const int max_just_dec_len_this_time = set_max_lengths.data<int>()[4]; | ||
|
|
||
| if (max_just_dec_len_this_time > 0) { | ||
| if (speculate_decoder) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
这两个分支后续需要统一
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
下一步这个算子会拆出来 统一去掉speculate_decoder分支
| cudaDeviceGetAttribute(&sm_cout, cudaDevAttrMultiProcessorCount, device)); | ||
|
|
||
| dim3 grids( | ||
| sm_cout * |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
这里最好可以预留一个设置位,如果没有设置就按照sm_count来,TBO下可能会用到。
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
目前kerenl需要提前自适应搜索chunk 搜索依赖这个block数 没办法提前拿到 我理解可以以经验值来做支持
| bool IsDynamicC8 = false> | ||
| __global__ void decode_append_attention_c8_kernel( | ||
| const __grid_constant__ AttentionParams<T, CacheT> params | ||
| // const __grid_constant__ CUtensorMap key_tensor_map, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
TMA的代码不用删,入参保留,增加一个模板参数,kernel里通过constexpr判断是否走TMA加载,方便后续调试。
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
TMA代码涉及到sync操作的顺序 两种同时支持会引入比较多的开关 影响代码简洁 以注释形式保留了 调试可以打开注释做简单改动跑通
yongqiangma
left a comment
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
2e04b4e
into
PaddlePaddle:release/online/20251131
Motivation
Modifications
Usage or Command
Accuracy Tests
Checklist
[FDConfig],[APIServer],[Engine],[Scheduler],[PD Disaggregation],[Executor],[Graph Optimization],[Speculative Decoding],[RL],[Models],[Quantization],[Loader],[OP],[KVCache],[DataProcessor],[BugFix],[Docs],[CI],[Optimization],[Feature],[Benchmark],[Others],[XPU],[HPU],[GCU],[DCU],[Iluvatar],[Metax]]pre-commitbefore commit.releasebranch, make sure the PR has been submitted to thedevelopbranch, then cherry-pick it to thereleasebranch with the[Cherry-Pick]PR tag.