[BugFix] Accuracy fix for llama4 int4 - improperly casted scales #16801
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
👋 Hi! Thank you for contributing to the vLLM project. 💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in #pr-reviews, coordinate on features in #feat- channels, or join special interest groups in #sig- channels. Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging. To run CI, PR reviewers can either: Add 🚀 |
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
csrc/moe/moe_wna16.cu
Outdated
| @@ -214,7 +211,7 @@ __global__ void moe_wna16_gemm_kernel( | |||
| const int32_t token_index = | |||
| sorted_token_ids[blockIdx.x * BLOCK_SIZE_M + m]; | |||
| if (mul_topk_weight) { | |||
| res[m] *= topk_weights[token_index]; | |||
| res[m] *= Dtype::num2float(topk_weights[token_index]); | |||
There was a problem hiding this comment.
If top_k = 2, then for each token, we would have 2 top_k_weights to be applied. Looking at the line here, it assumes topk_weights[token_index] is always a scalar value?
There was a problem hiding this comment.
its a bit weird but the way it works is that token_index ranges between 0 and topk * num_tokens, then topk_weights is of size topk * num_tokens, and basically we index the input via input[token_index / topk] but the topk_weights[token_index]
e.g. if we had the router output as topk 2 with 2 tokens:
weights = [
[0.1, 0.2],
[0.3, 0.4]
]
topk_ids = [
[1, 0],
[1, 2]
]
and if we had block size 1 for MoE align block size then we'd get:
sorted_token_ids = [1, 0, 2, 3]
expert_ids = [0, 1, 1, 2]
so token_index (a poor name) is actually the index from sorted_token_ids which is of size topk * num_tokens, meaning topk_weights[token_index] is still a scalar
There was a problem hiding this comment.
I see so basically for topk=2, the outer loop of for (int m = 0; m < num_valid_tokens; ++m) would iterate each token in the request twice, the first time it get token x first expert weight, and the second time, it get token x second expert weight. Is this the right understanding?
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
LucasWilkinson
force-pushed
the
lwilkinson/fix-llama4-int4
branch
from
2a1e54c to
ca7f666
Compare
|
Confirmed the fix recover llama4 int4 checkpoint accuracy back to normal. |
mgoin
approved these changes
Apr 17, 2025
There was a problem hiding this comment.
Thanks so much for this bundle of fixes, LGTM! cc @jinzhen-lin for the moe_wna16 cuda update
mgoin
added
bug
INT4 ScoutTP=8 after the changeTP=1 after the changeBefore the changeBF16 Scout TP=8After the changeBefore the change |
yangw-dev
pushed a commit
to yangw-dev/vllm
that referenced
this pull request
Apr 21, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Signed-off-by: Yang Wang <elainewy@meta.com>
heyselbi
pushed a commit
to red-hat-data-services/vllm
that referenced
this pull request
Apr 23, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
heyselbi
pushed a commit
to red-hat-data-services/vllm
that referenced
this pull request
Apr 24, 2025
SUMMARY: sync to upstream `v0.8.4` and cherry-pick of `7eb42556281d30436a3a988f2c9184ec63c59338`. the cherry-pick is @LucasWilkinson 's llama4 patch. GIT LOG: ```bash commit b197179 (HEAD -> sync-upstream-v0.8.4, origin/sync-upstream-v0.8.4) Author: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Date: Fri Apr 18 01:13:29 2025 -0400 [BugFix] Accuracy fix for llama4 int4 - improperly casted scales (vllm-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> commit 60267cc Author: andy-neuma <andy@neuralmagic.com> Date: Mon Apr 21 14:57:52 2025 -0400 remove duplicate entries commit 9d18b50 Merge: db0e117 dc1b4a6 Author: andy-neuma <andy@neuralmagic.com> Date: Mon Apr 21 14:50:01 2025 -0400 Merge remote-tracking branch 'upstream/v0.8.4' into sync-upstream-v0.8.4 commit db0e117 Author: andy-neuma <andy@neuralmagic.com> Date: Mon Apr 21 14:35:23 2025 -0400 Revert "Revert "[V1] DP scale-out (1/N): Use zmq ROUTER/DEALER sockets for input queue (vllm-project#15906)"" This reverts commit 296c657. commit dc1b4a6 (tag: v0.8.4, upstream/v0.8.4) Author: Russell Bryant <rbryant@redhat.com> Date: Sun Apr 13 22:13:38 2025 -0400 [Core][V0] Enable regex support with xgrammar (vllm-project#13228) Signed-off-by: Russell Bryant <rbryant@redhat.com> ``` COMMANDS: ```bash git fetch upstream git checkout -b sync-upstream-v0.8.4 git revert 296c657 git merge upstream/v0.8.4 git cherry-pick 7eb4255 ``` TEST PLAN: accept sync ... https://github.com/neuralmagic/nm-cicd/actions/runs/14581880024 release ... https://github.com/neuralmagic/nm-cicd/actions/runs/14596026989 --------- Signed-off-by: Tristan Leclercq <tristanleclercq@gmail.com> Signed-off-by: yihong0618 <zouzou0208@gmail.com> Signed-off-by: reidliu41 <reid201711@gmail.com> Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com> Signed-off-by: Jinzhen Lin <linjinzhen@hotmail.com> Signed-off-by: Jonghyun Choe <andy.choe729@gmail.com> Signed-off-by: Lu Fang <fanglu@fb.com> Signed-off-by: Hyesoo Yang <hyeygit@gmail.com> Signed-off-by: Ben Jackson <ben@ben.com> Signed-off-by: Roger Wang <ywang@roblox.com> Signed-off-by: Isotr0py <2037008807@qq.com> Signed-off-by: rongfu.leng <rongfu.leng@daocloud.io> Signed-off-by: Varun Sundar Rabindranath <varun@neuralmagic.com> Signed-off-by: paolovic <paul-philipp.luley@uzh.ch> Signed-off-by: Chengji Yao <chengjiyao@google.com> Signed-off-by: Kay Yan <kay.yan@daocloud.io> Signed-off-by: Woosuk Kwon <woosuk.kwon@berkeley.edu> Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk> Signed-off-by: shen-shanshan <467638484@qq.com> Signed-off-by: YamPengLi <yampayne.lyp@alibaba-inc.com> Signed-off-by: WangErXiao <863579016@qq.com> Signed-off-by: Aston Zhang <22279212+astonzhang@users.noreply.github.com> Signed-off-by: Chris Thi <chris.c.thi@gmail.com> Signed-off-by: drisspg <drisspguessous@gmail.com> Signed-off-by: Jon Swenson <jmswen@gmail.com> Signed-off-by: Keyun Tong <tongkeyun@gmail.com> Signed-off-by: Lu Fang <fanglu@meta.com> Signed-off-by: Xiaodong Wang <xdwang@meta.com> Signed-off-by: Yang Chen <yangche@fb.com> Signed-off-by: Ye (Charlotte) Qi <yeq@meta.com> Signed-off-by: Yong Hoon Shin <yhshin@meta.com> Signed-off-by: Zijing Liu <liuzijing2014@gmail.com> Signed-off-by: Lu Fang <lufang@fb.com> Signed-off-by: Lucia Fang <fanglu@fb.com> Signed-off-by: Gregory Shtrasberg <Gregory.Shtrasberg@amd.com> Signed-off-by: NickLucche <nlucches@redhat.com> Signed-off-by: Benjamin Chislett <benjamin.chislett@centml.ai> Signed-off-by: Nick Hill <nhill@redhat.com> Signed-off-by: Leon Seidel <leon.seidel@fau.de> Signed-off-by: mgoin <michael@neuralmagic.com> Signed-off-by: youkaichao <youkaichao@gmail.com> Signed-off-by: Miles Williams <42222518+mlsw@users.noreply.github.com> Signed-off-by: mgoin <mgoin64@gmail.com> Signed-off-by: Siyuan Liu <lsiyuan@google.com> Signed-off-by: Kebe <mail@kebe7jun.com> Signed-off-by: simon-mo <simon.mo@hey.com> Signed-off-by: Alex-Brooks <Alex.Brooks@ibm.com> Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com> Signed-off-by: imkero <kerorek@outlook.com> Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Signed-off-by: Russell Bryant <rbryant@redhat.com> Signed-off-by: Jee Jee Li <pandaleefree@gmail.com> Signed-off-by: Yue <yueshen@nvidia.com> Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com> Signed-off-by: kliuae <kuanfu.liu@embeddedllm.com> Signed-off-by: luka <luka@neuralmagic.com> Signed-off-by: lvfei.lv <lvfei.lv@alibaba-inc.com> Signed-off-by: Ajay Vohra <ajayvohr@amazon.com> Signed-off-by: Guillaume Calmettes <gcalmettes@scaleway.com> Signed-off-by: zh Wang <rekind133@outlook.com> Signed-off-by: Chendi Xue <chendi.xue@intel.com> Signed-off-by: Joe Runde <Joseph.Runde@ibm.com> Signed-off-by: zRzRzRzRzRzRzR <2448370773@qq.com> Signed-off-by: Aaron Ang <aaron.angyd@gmail.com> Signed-off-by: Benjamin Kitor <bkitor@gigaio.com> Signed-off-by: Michael Goin <mgoin64@gmail.com> Signed-off-by: Chenyaaang <chenyangli@google.com> Signed-off-by: cyy <cyyever@outlook.com> Signed-off-by: wineandchord <guoqizhou19@gmail.com> Signed-off-by: LiuXiaoxuanPKU <lilyliupku@gmail.com> Signed-off-by: Chih-Chieh-Yang <7364402+cyang49@users.noreply.github.com> Signed-off-by: look <eeslook@163.com> Signed-off-by: jadewang21 <jadewangcn@outlook.com> Signed-off-by: alexey-belyakov <alexey.belyakov@intel.com> Signed-off-by: jiang.li <jiang1.li@intel.com> Signed-off-by: DefTruth <qiustudent_r@163.com> Signed-off-by: chaow <chaow@amd.com> Signed-off-by: Tomasz Zielinski <tomasz.zielinski@intel.com> Signed-off-by: rzou <zou3519@gmail.com> Signed-off-by: Travis Johnson <tsjohnso@us.ibm.com> Signed-off-by: Christian Sears <csears@redhat.com> Signed-off-by: Gogs <gogs@fake.local> Signed-off-by: Yuan Tang <terrytangyuan@gmail.com> Signed-off-by: Tianer Zhou <ezhoureal@gmail.com> Signed-off-by: hzji210@gmail.com <hzji210@gmail.com> Signed-off-by: Jie Fu <jiefu@tencent.com> Signed-off-by: snowcharm <snowcharmqq@gmail.com> Signed-off-by: Ryan McConville <ryan@ryanmcconville.com> Co-authored-by: Tristan Leclercq <49700633+tristanleclercq@users.noreply.github.com> Co-authored-by: Kevin H. Luu <kevin@anyscale.com> Co-authored-by: yihong <zouzou0208@gmail.com> Co-authored-by: Reid <61492567+reidliu41@users.noreply.github.com> Co-authored-by: reidliu41 <reid201711@gmail.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com> Co-authored-by: Jinzhen Lin <linjinzhen@hotmail.com> Co-authored-by: Jonghyun Choe <andy.choe729@gmail.com> Co-authored-by: Lucia Fang <116399278+luccafong@users.noreply.github.com> Co-authored-by: Hyesoo Yang <45211235+hyeygit@users.noreply.github.com> Co-authored-by: Ben Jackson <ben@ben.com> Co-authored-by: Roger Wang <136131678+ywang96@users.noreply.github.com> Co-authored-by: Paul Schweigert <paul@paulschweigert.com> Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: rongfu.leng <rongfu.leng@daocloud.io> Co-authored-by: Varun Sundar Rabindranath <varunsundar08@gmail.com> Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com> Co-authored-by: paolovic <91155454+paolovic@users.noreply.github.com> Co-authored-by: paolovic <paul-philipp.luley@uzh.ch> Co-authored-by: Chengji Yao <chengjiyao@google.com> Co-authored-by: Woosuk Kwon <woosuk.kwon@berkeley.edu> Co-authored-by: Martin Hoyer <mhoyer@redhat.com> Co-authored-by: Kay Yan <kay.yan@daocloud.io> Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk> Co-authored-by: Shanshan Shen <467638484@qq.com> Co-authored-by: YamPengLi <yampayne.lyp@alibaba-inc.com> Co-authored-by: Cyrus Leung <cyrus.tl.leung@gmail.com> Co-authored-by: Robin <863579016@qq.com> Co-authored-by: Lu Fang <30275821+houseroad@users.noreply.github.com> Co-authored-by: Lu Fang <fanglu@fb.com> Co-authored-by: Roger Wang <ywang@roblox.com> Co-authored-by: Gregory Shtrasberg <156009573+gshtras@users.noreply.github.com> Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com> Co-authored-by: Benjamin Chislett <benjamin.chislett@centml.ai> Co-authored-by: Nick Hill <nhill@redhat.com> Co-authored-by: leon-seidel <83984854+leon-seidel@users.noreply.github.com> Co-authored-by: Driss Guessous <32754868+drisspg@users.noreply.github.com> Co-authored-by: Michael Goin <michael@neuralmagic.com> Co-authored-by: youkaichao <youkaichao@gmail.com> Co-authored-by: Miles Williams <42222518+mlsw@users.noreply.github.com> Co-authored-by: Satyajith Chilappagari <satchill@amazon.com> Co-authored-by: mgoin <mgoin64@gmail.com> Co-authored-by: Jennifer Zhao <ai.jenniferzhao@gmail.com> Co-authored-by: zxfan-cpu <zxfanzhang@tencent.com> Co-authored-by: Yong Hoon Shin <48474650+sarckk@users.noreply.github.com> Co-authored-by: Siyuan Liu <lsiyuan@google.com> Co-authored-by: Kebe <mail@kebe7jun.com> Co-authored-by: Simon Mo <simon.mo@hey.com> Co-authored-by: Alex Brooks <alex.brooks@ibm.com> Co-authored-by: TY-AMD <tianyuan.wu@amd.com> Co-authored-by: wang.yuqi <noooop@126.com> Co-authored-by: Kero Liang <kerorek@outlook.com> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Co-authored-by: Russell Bryant <rbryant@redhat.com> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com> Co-authored-by: yueshen2016 <39203804+yueshen2016@users.noreply.github.com> Co-authored-by: TJian <tunjian.tan@embeddedllm.com> Co-authored-by: Hongxia Yang <hongxia.yang@amd.com> Co-authored-by: kliuae <kuanfu.liu@embeddedllm.com> Co-authored-by: Luka Govedič <ProExpertProg@users.noreply.github.com> Co-authored-by: Accelerator1996 <lvfei.lv@alibaba-inc.com> Co-authored-by: ajayvohra2005 <ajayvohr@amazon.com> Co-authored-by: Guillaume Calmettes <gcalmettes@scaleway.com> Co-authored-by: zh Wang <rekind133@outlook.com> Co-authored-by: Chendi.Xue <chendi.xue@intel.com> Co-authored-by: Joe Runde <Joseph.Runde@ibm.com> Co-authored-by: Yuxuan Zhang <2448370773@qq.com> Co-authored-by: Aaron Ang <67321817+aaron-ang@users.noreply.github.com> Co-authored-by: Jintao <huangjintao@mail.ustc.edu.cn> Co-authored-by: Benjamin Kitor <bkitor@gigaio.com> Co-authored-by: Chenyaaang <42742451+Chenyaaang@users.noreply.github.com> Co-authored-by: cyyever <cyyever@outlook.com> Co-authored-by: Ye (Charlotte) Qi <yeq@meta.com> Co-authored-by: wineandchord <guoqizhou123123@qq.com> Co-authored-by: Nicolò Lucchesi <nicolo.lucchesi@gmail.com> Co-authored-by: Lily Liu <lilyliupku@gmail.com> Co-authored-by: Chih-Chieh Yang <7364402+cyang49@users.noreply.github.com> Co-authored-by: Yu Chin Fabian Lim <flim@sg.ibm.com> Co-authored-by: look <eeslook@163.com> Co-authored-by: WWW <jadewangcn@outlook.com> Co-authored-by: Alexey Belyakov <alexey.belyakov@intel.com> Co-authored-by: Li, Jiang <jiang1.li@intel.com> Co-authored-by: DefTruth <31974251+DefTruth@users.noreply.github.com> Co-authored-by: chaow-amd <chaow@amd.com> Co-authored-by: Tomasz Zielinski <85164140+tzielinski-habana@users.noreply.github.com> Co-authored-by: Richard Zou <zou3519@users.noreply.github.com> Co-authored-by: Travis Johnson <tsjohnso@us.ibm.com> Co-authored-by: Kai Wu <kaiwu@meta.com> Co-authored-by: Christian Sears <117944059+Chr1st1anSears@users.noreply.github.com> Co-authored-by: Gogs <gogs@fake.local> Co-authored-by: Yuan Tang <terrytangyuan@gmail.com> Co-authored-by: Tianer Zhou <ezhoureal@gmail.com> Co-authored-by: Huazhong Ji <hzji210@gmail.com> Co-authored-by: Jie Fu (傅杰) <jiefu@tencent.com> Co-authored-by: SnowCharm <qiuyilun@u.nus.edu> Co-authored-by: Ryan McConville <ryan@ryanmcconville.com> Co-authored-by: andy-neuma <andy@neuralmagic.com>
jikunshang
pushed a commit
to jikunshang/vllm
that referenced
this pull request
Apr 29, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
lk-chen
pushed a commit
to lk-chen/vllm
that referenced
this pull request
Apr 29, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
adobrzyn
pushed a commit
to HabanaAI/vllm-fork
that referenced
this pull request
Apr 30, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Signed-off-by: Agata Dobrzyniewicz <adobrzyniewicz@habana.ai>
dtrifiro
pushed a commit
to red-hat-data-services/vllm
that referenced
this pull request
Apr 30, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
RichardoMrMu
pushed a commit
to RichardoMrMu/vllm
that referenced
this pull request
May 12, 2025
…m-project#16801) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Signed-off-by: Mu Huai <tianbowen.tbw@antgroup.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
A few MoE fixes:
.data_ptr<float>()AttributeError: 'FusedMoE' object has no attribute 'params_dtype'moe_wna16_gemm)