Muse Spark Open: 1,109 טוקנים בשנייה הם שינוי בקרנל, לא שינוי במודל

M​use Spark Open was released on October 15, 2026, and nothing about its weights changed this week. What changed is the software underneath them. On November 11, 2026 the SGLang team at LMSYS published an engineering post titled "M​use Spark Open Ultrafast: Pushing beyond 1,000 Tokens per second," and the headline number in it is decoding throughput on eight GB300s going from 517 to 1,109 tokens per second for the same model on the same hardware. The mechanism is a megakernel: one resident kernel per layer instead of roughly 1,450 separate kernel launches per decode step. So this is a serving story, not a model story, and the interesting question is not whether 1,109 is impressive — it is — but which of those numbers you are allowed to carry into a capacity plan. The short answer is fewer than you would like: the benchmark is one request, one dataset, batch size one, and no one outside the authors has reproduced it yet.

קו הבסיס שאיש לא מודד: 1,450 שיגורים לכל טוקן

A normal GPU kernel does one thing. It runs one matrix multiplication or one attention pass as a grid of thread blocks spread across the streaming multiprocessors, then exits. A serving engine like SGLang strings those kernels together into a chain, and for M​use Spark Open that chain is long: the post puts the count at about 1,450 kernel launches for each decoding step. Between two consecutive kernels the GPU drains and refills. SMs that have finished their tile sit idle waiting for the next launch, and inside an SM individual warps idle too, because a warp that has issued its loads has nothing to do until the data lands.

זו העלות שהמגהקרינל תוקף, והמחברים נעזרים באנלוגיה משנות ה-60 כדי להסביר אותה: מכונות מוקדמות הריצו עבודות אצווה, כשעבודה אחת החזיקה במכונה כולה, והמעבד עמד בטל בזמן שסרט הסתובב. ריבוי תוכניות תיקן זאת על ידי השארת כמה עבודות בזיכרון ומעבר לעבודה אחרת בכל פעם שהעבודה הנוכחית נחסמה. מגהקרינל הוא הרעיון הזה שהועבר ל-GPU — קרינל אחד נשאר פעיל בכל SM ומריץ עבודות רבות בבת אחת, כך שוורפים יכולים להתחיל בעבודה של הפעולה הבאה לפני שהפעולה הקודמת הסתיימה.

ההשפעה הנמדדת, מהתרשים הנלווה לפרק הביצועים, היא שלב פענוח שעובר מ-12.9 מ״ש ל-5.9 מ״ש. התרשים מפרק את השלב לשני חצאיו, טיוטה ואימות: 1.16 מ״ש ו-10.9 מ״ש עם 1,450 קרנלים בבסיס, ו-0.44 מ״ש ו-5.3 מ״ש עם 97 קרנלים במגה-קרנל. שנים-עשר נקודה תשע עד חמש נקודה תשע הם בערך פי 2.2, ו-1,109 חלקי 517 הם בערך פי 2.1, כך ששני יחסי הכותרת תואמים זה לזה. קיימת בדיקה שלישית לכל מי שמוכן לעשות את החלוקה: השלב מאמת 16 אסימונים ומקבל 6.6, כך ש-6.6 אסימונים ב-5.9 מ״ש יוצאים בערך 1,120 אסימונים לשנייה. החשבון עקבי פנימית, וזה הדבר הראשון שצריך לבסס לגבי מספר כזה, ולא האחרון.

דבר אחד שהמספרים לא אומרים הוא לאיזה ערך מגיעה המהירות לכל GPU. הפוסט מדווח על תפוקה מצטברת עבור כל פריסת שמונת ה-GPU. מחולק בשמונה, זה כ-139 טוקנים לשנייה לכל GPU — החישוב שלי על הנתון שלהם, לא מספר שהפוסט מפרסם, ואינו בר-השוואה למספר לכל GPU שאולי תראה מצוטט עבור מודל אחר במכונה אחרת.

היכן נמצא גבול הקרנל, ומדוע הוא עוצר שם

המגה-קרנל אינו בולע את כל המודל. כל הרצה מכסה את הקשב של שכבת דקודר אחת דרך השלכת ה-qkv וכתיבת ה-KV של השכבה הבאה, וזהו גבול מכוון ולא מגבלה של nerve. הסיבה שזה עובד היא שרוב מה שהקשב של שכבה קורא — היסטוריית ה-KV — נכתב בשלבים קודמים ואינו תלוי במשהו שנמצא בתהליך, כך שהרצה חדשה יכולה להתחיל לטעון אותו ברגע שהוא זמין, לפני שפלט השכבה הקודמת קיים. לשכבות הראשונה והאחרונה יש פעולות מעט שונות, והפוסט מציין גם זאת.

בתוך שיגור אחד הצורה קבועה: בלוק תהליכים אחד של שנים-עשר warp-ים בכל אחד מ-144 SM, שנשלחים באשכולות של ארבעה, תוך שימוש בזיכרון משותף מבוזר היכן שזמין reduction בעל השהיה נמוכה יותר. שנים-עשר ה-warp-ים אינם ניתנים להחלפה. warp 0 טוען KV ומשקלים ומנצל את המרווח שלו לטעינה מוקדמת של המשקלים של הפעולה הבאה. warp 1 ו-2 מריצים אקטיבציות וליבות tensor. warp 3 עד 11 מחלקים ביניהם את ה-attention, את ה-GEMM הצפופים, את ה-GEMM של המומחים ואת צד הקבלה של all-reduce ו-all-gather. בשכבות של חלון מחליק, warp 4 עד 11 טוענים KV משלהם במקום להמתין ל-warp 0.

Getting there required a change in how the model is split. M​use Spark Open is a mixture-of-experts model, routing every token to 8 of 256 experts, with experts stored in 4-bit NVFP4 and some layers left in bf16. It has 62 decoder layers, of which 46 use sliding-window attention over the last 2,048 tokens and 16 use global attention, plus a pooler projection that shrinks each token vector before the experts and an unpooler that restores it afterwards. The megakernel shards both MoE GEMMs on the N dimension, which forces an all-gather after each GEMM — an fp4 all-gather after the first, a bf16 one after the second. That is one more communication step than vanilla tensor parallelism, which shards on the intermediate dimension and needs a single all-reduce, but the authors argue it moves less data overall, scales better to the batch sizes they want next, and can be made bitwise identical to expert parallelism so it matches training numerics.

הפרט הזה חשוב יותר ממה שהוא נראה. צוותי אינפרנס בדרך כלל משלימים עם כך שקרנל מהיר יותר מפיק מספרים שונים במעט, ומכווננים בשקט את ספי ההערכה שלהם. קרנל ששווה ברמת הביטים למקבול בצד האימון מסיר את הוויכוח הזה לחלוטין, וזה בדיוק מסוג התכונות שגורמות לאופטימיזציה להיות מאומצת ולא רק להערץ.

גרעין בגודל כזה צריך לשלם שלושה חשבונות.

המגה-קרנל אינו ללא עלות. הפוסט גלוי לב באופן יוצא דופן לגבי התקורות, ואלה הם החלקים שכדאי לקרוא פעמיים, כי הם גם החלקים שמספרים לך כמה מקום עוד נותר.

החשבון הראשון הוא מטמון ההוראות. SM שומר במטמון כ-128 KB של הוראות, והקרנלים האלה הם באורך 215 עד 314 KB של קוד מכונה כל אחד, מפני שכל אחד מהם מכיל כל שלב של שכבה. השקה מקרבת רק את 32 עד 64 ה-KB הראשונים אל ה-SM-ים; כל מה שמעבר לחלון הזה נשלף שורת מטמון אחר שורת מטמון בשימוש הראשון. המחברים ביצעו על כך ניסוי נקי ודיווחו עליו: שני בלוקים של 100 KB, שכל אחד מהם לקח 12.4 µs לבדו, לקחו 19.5 µs כשחוברו גב אל גב, והקטנת בלוק אחד ל-20 KB החזירה את השני למהירותו לבדו. הוראות עוברות דרך אותו מטמון L2 כמו נתונים, כך שקרנל שגדל מעבר למטמון ההוראות מוצא את עצמו מתחרה בטעינות המשקולות של עצמו.

הסעיף השני הוא אחזור מוקדם של משקלים, והפוסט מתייחס אליו כאל בעיית מיקום עם שלושה פתרונות. הזיכרון המשותף מקבל טבעת של 5 × 32 KB לכל SM, כ-23 MB ברחבי ה-GPU. השאר הולך ל-L2, וה-L2 של GB300 מחזיק כ-135 MB, אם כי כל דבר ב-L2 צריך מעבר שני וקצר יותר אל הזיכרון המשותף, וה-L2 משותף עם כל דבר אחר שנמצא בפעילות. השניים מצטברים: לקראת סוף שכבה הטבעת כבר מחזיקה את משקלי ה-qkv הראשונים של השכבה הבאה. בנוסף, warp 3 קורא לאילו מומחים השלב הקודם ניתב, ומושך אל L2 את הפרוסות של ה-GPU הזה מתוך 14 המומחים הנפוצים ביותר, כ-50 MB לכל GPU, בהימור ששלבים עוקבים מנתבים באופן דומה. ההימור הזה שווה כ-0.2 ms מתוך שלב של 5.3 ms, ועלול לעלות ברוחב פס כשהוא שגוי.

החשבון השלישי הוא סנכרון, וזה החשבון שאין לו תשובה נקייה. כאשר warp ממתין ל-GPU אחר דרך NVLink, אין פסיקה ואין המתנה לכתובת בחומרה; ה-warp סקר דגל עד שהוא משתנה. המיקרו-בנצ'מרקים של הפוסט מעמידים סקר acquire בהיקף GPU על 141 עד 145 ns, דגל חד-כיווני בין שני GPUs באותו צומת על כ-1.9 µs, וסקר שממתין בתור אחרי העתקות bulk על 1.0 עד 2.6 µs. גדרות עולות 0.4 µs כשהן במצב סרק ו-3.6 עד 4.0 µs כשטעינות TMA פעילות. אפילו עם עבודה מחולקת באופן מושלם, GPUs יכולים לפגר זה אחרי זה בעד 4 µs, וכל עצירה בשכבה ממתינה לאיטי ביותר. המגה-קרנל מחליף גדרות בערכי זקיף הנכתבים לחריצי store רגילים — דגל מלא בנתונים, +0/−0 עבור ה-all-reduce, NaN שכולו אחדים עבור איסופי ה-bf16 — ומאפס את העותק הרדום של כל חריץ בזמן שהאחר נמצא בשימוש.

אלה המספרים שגורמים לתוצאה להרגיש כמו הנדסה ולא כמו קסם: {{1}}4 µs{{/1}} של סטייה בין GPUs לעומת צעד של {{2}}5.9 ms{{/2}} היא קטנה, אבל ארבע עצירות סנכרון בכל שכבה, כפול 62 שכבות, הן בדיוק סוג התקורה שמכלה רווח של {{3}}2×{{/3}} אם טועים בזה.

על מה עדיין מתבזבזות 5.9 ms

הפוסט כולל ביקורת עצמית, והשורה השימושית ביותר בו עוסקת בפער שאחרי ה-GEMM השני. בלוקים מסיימים את ה-GEMM הזה בכל נקודה בין 55 ל-64 מיקרושניות מתחילת השכבה, מהאחוזון ה-10 עד ה-90, אבל ה-combine ממתין לפלט של כל GPU, ולכן הוא לא יכול להתחיל לפני 65 מיקרושניות. מסוף GEMM 2 ועד סוף השכבה בלוק יושב בלי עבודה במשך כ-12.5 מיקרושניות בממוצע — כמעט מחצית מסך זמן הבטלה שלו. השרשרת הזו, ולא תקורת השקת הקרנל שהמגה-קרנל כבר ביטל, היא מה שהמחברים אומרים שהם בוחנים כעת.

הם גם מסבירים מדוע הקשב של השכבה הבאה אינו יכול פשוט להתחיל מוקדם יותר. PDL משגר אותו ב-79 µs, כ-4 µs לפני שהבלוק הראשון יוצא, אך הבלוקים שלו נוחתים רק כאשר הקודמים עוזבים את ה-SMs שלהם ב-84 µs, כך ששיגור מוקדם יותר לא יתחיל שום דבר מוקדם יותר. הפתרון הוא לבצע prefetch לשורות ה-KV הישנות שלו אל L2 ב-65 µs, שם הן ממתינות. זהו אותו דפוס כמו ה-prefetch של המשקולות: כאשר אינך יכול להסיר המתנה, אתה מעביר לתוכה עבודה.

פרופיילינג הוא הבעיה המבנית האחרת, והיא בעיה טובה: ה-kernel נראה כמו פס ארוך אחד לכל שכבה בעיני כל profiler חיצוני, ולכן הוא מודד את עצמו. Lane 0 של כל block כותב timestamp למערך גלובלי דרך אוגר ה-globaltimer, והמחברים מודדים את עלות האינסטרומנטציה הזו בכ-2.5 µs לכל השקת שכבה, 2.8% מצעד, כאשר הפלטים נשארים זהים ביט-ביט וקיים דגל build כדי לקמפל אותה החוצה. kernel שצריך לשאת profiler משל עצמו הוא kernel שהפך למערכת הפעלה, וזו בדיוק המסגור שבו הפוסט משתמש: תהליך הוא kernel אחד המחזיק בכל משאבי ה-GPU, threads הם warps ללא preemption, המתנה ל-I/O היא busy-waiting על דגל, DMA הוא TMA שנעקב על ידי mbarrier, הודעות הן stores רגילים על פני NVLink.

גם מודל הטיוטה נעשה מהיר יותר, וזה חצי מהניצחון

M​use Spark Open's decode path is speculative: a draft model proposes tokens and the big model verifies them in one pass. The draft here is DFlash V2, which guesses the next 15 tokens in a single pass, and the big model checks all of them in a pass over 16 tokens. The verification uses a tree of candidate tokens generated from the draft rather than a linear chain, which the post says is slightly better.

הצוות יישם את אותו טיפול מגה־קרנל גם על הטיוטה: חמש שכבות טיוטה רצות כל אחת כהרצה אחת, ועוד שלוש נוספות, מה שמעביר את הטיוטה מ־175 קרנלים ל־8 ומ־1.16 ms לכל צעד ל־440 µs, שיפור של פי 2.6 באותו חצי של הצעד. שווה להציב את נתוני הקבלה לצד ההשוואה הציבורית הקרובה ביותר שהפוסט מציע.

• Acceptance length — M​use Spark Open verifies 16 tokens per step and accepts 6.6 on average; Xiaomi's MiMo-V2.5-Pro, which crossed 1,000 tokens per second earlier, verifies 8 and accepts 6.3

• למה 16 — אימות של שישה־עשר טוקנים עולה יותר משום שהם מפעילים יותר מומחים, אבל לצוות היה קל יותר להגיע ל־6.6 עם עץ רחב יותר מאשר להגיע לאותו אורך עם שרשרת צרה יותר

• ערימת טיוטה — 8000 שורות של CUDA ו-PTX עבור המודל הגדול, כאשר הקוד פורסם ב-facebookresearch/sgl-muse-mega

קרא את ההשוואה בעיון, מכיוון שהיא אינה השוואת ביצועים בין שני מודלים. זוהי השוואה בין שתי תצורות פענוח ספקולטיבי, אחת שפורסמה על ידי Xiaomi ואחת על ידי צוות SGLang, על חומרה שונה ועומסי עבודה שונים. הטענה היחידה שהיא תומכת בה היא שעץ אימות רחב יותר הוא נתיב ישים לאותו אורך קבלה, וזו טענת תכנון, לא דירוג.

מהו המספר הזה, ומה הוא לא

החלק שתוכנית קיבולת חייבת להבין נכון, וזה החלק שהפוסט עצמו גלוי לגביו. הנתון 1,109 הוא החציון של 20 צעדי פענוח בבקשה בודדת של GSM8K, בדרגת מקבילות טנסור 8, על שמונה GB300, עם גודל אצווה אחד ופענוח ספקולטיבי פעיל. הפוסט מציין בבירור שהמגהקרנל תומך היום רק בגודל אצווה אחד, ושהרחבתו לאצוות גדולות יותר היא חלק העבודה הבא.

כל אחד מהאילוצים האלה מסיר טענה אחרת. גודל אצווה של אחד פירושו שזו אינה תוצאת תפוקה במובן שבו צוות serving משתמש במונח: אין בו רווח מקיבוץ לאצוות, אין תור, ואין ערבוב של אורכי פרומפטים. מערך נתונים אחד פירושו שהמספר קשור להתפלגות הפרומפט והפלט של GSM8K, ולתבנית ניתוב נמדדת של 49 מומחים בכל שכבה, שאותה השתמשו המחברים לאחר מכן כניתוב סינתטי בפרופיילינג שלהם עצמם. עשרים צעדים של בקשה אחת פירושם שהשונות חסומה במספר הצעדים, ולא במספר הבקשות.

ובקשה בודדת היא המקרה הגרוע ביותר לחפיפה, כפי שהמחברים עצמם אומרים: "בבקשה אחת כל המתנה נחשפת; עם יותר, כמעט תמיד יש עבודה מוכנה." לכן ה-2.1× של המגה-קרנל הוא מדידה של המשטר שבו הסתרת השהיה היא הקשה ביותר, והפוסט מפורש בכך שאצוות גדולות יותר אמורות להיות קלות יותר ולא קשות יותר. זהו טיעון לטובת התוצאה, והוא עדיין טיעון ולא מדידה.

מה שניתן לדעת היום, אם כן, הוא צר וממשי: הקוד קיים במאגר ציבורי, הפוסט מתאר את המנגנון בפירוט מספיק כדי לבדוק אותו, האריתמטיקה הפנימית מתיישבת לאורך שלושה יחסים בלתי תלויים, והמחברים תיעדו את שלושת התקורות שמגבילות רווחים נוספים. מה שאינו ניתן לדעת הוא אם צוות עצמאי משיג 1,109 על שמונת ה-GB300s שלו, אם הגישה שורדת באצווה אמיתית, ואם היא נכנסת לגרסת SGLang שפורסמה — הפוסט אינו טוען שכך קרה, והמאגר הוא תוצר מחקרי ולא חבילה שאפשר להתקין. ההשוואה שתוכל להכריע את הרב ביותר היא מול baseline מכויל באותו חומרה, שכן "שרשרת ברירת המחדל של SGLang" היא מטרה נעה, והצוות שכתב את המגה-קרנל גם מכייל את ה-baseline שהוא מנצח.

איך זה קשור ל־169 הטוקנים לשנייה שאפשר לשכור היום

Artificial Analysis records a very different number for M​use Spark Open: 169 tokens per second of output speed, measured end to end including time to first token, against a hosted endpoint, with a first-answer-token time of 13.81 seconds. That is not a contradiction and it is not a rebuke. It is a different measurement of a different deployment: one stream through somebody's managed API, versus one request through your own eight-GPU node with a research kernel.

הפער הוא הדבר השימושי ביותר בסיפור הזה עבור מי שאינו כבר בעל החומרה. אם נקודת קצה מגיבה בקצב של 169 אסימונים לשנייה ומערך של שמונה GPUs שנבנה ביד מגיב בקצב של 1,109, ההבדל אינו סיבה לבנות את המערך. שמונה GB300s הם רכישת הון הנמדדת במיליונים, הקרנל הוא קוד מחקרי, ואילוץ batch-size-one פירושו שהמערך כרגע טוב בדיוק לעומס העבודה שנקודת קצה מתארחת כבר מטפלת בו היטב: בקשה אחת בכל פעם. המקום שבו ה-megakernel באמת משנה את החישוב הוא עבור צוות שכבר קנה את ה-GPUs לצורך אימון או עבור עומס עבודה עם מגבלות שהיית נתונים, ועכשיו שואל איזו ניצולת הוא יכול להפיק מהם ב-batch נמוך. עבור אותו צוות, 2.1× על אותה סיליקון עם אותה נומריקה הוא מספר אמיתי עם חשבונית אמיתית מצורפת, והחשבונית כבר שולמה.

לכל השאר, הקריאה הכנה היא שהפוסט הזה הוא איתות לאן תשתיות ההסקה מתקדמות, לא מוצר שאפשר לאמץ השבוע. הכיוון תואם את מה שאותה מעבדה מפרסמת כל השנה: פחות קרנלים, גדולים יותר, יותר מהחלטת התזמון מועברת לתוך הקרנל, וה-GPU נתפס כמכונה שיש לתכנת ולא כמערכת של קריאות ספרייה שצריך לסדר ברצף.

That is also the right frame for our own position, which is worth stating plainly. M​use Spark Open is not on OrcaRouter; it is a download, a gated repository and a licence, not an endpoint we serve, and we are not going to describe it as one. Muse Spark 1.2 and Muse Spark 1.1 are the two Muse models OrcaRouter actually routes, and they are there at Meta's own list price because the provider's list price passes through rather than gets marked up. If your question is which model to put behind a product this quarter rather than which kernel to compile, that is the shorter path: one OpenAI-compatible endpoint across 200+ models, with the provider's list price passed through at 0% markup, plus automatic failover and a routing DSL that moves traffic between models by configuration rather than by redeploying code.

מה היה משנה את הקטע הזה

ארבעה דברים, לפי סדר יורד של עד כמה הם היו משנים את התמונה.

• אצווה גדולה מאחת. המחברים אומרים שאצוות גדולות יותר הן הצעד הבא וטוענים שהן קלות יותר, מפני שכאשר יותר טוקנים נמצאים בתהליך, חלק מהפלטים של הפעולה הקודמת כבר מוכנים בעוד שהשאר עדיין מגיעים. ברגע שקיים מספר באצווה, ההשוואה מפסיקה להיות "זרם אחד מול זרם אחד" והופכת לטענת תפוקה שצוות סרווינג יכול להשתמש בה.

• שחזור עצמאי. כל נתון כאן מגיע ממערכת הבדיקות של המחברים עצמם, על החומרה שלהם, עם הניתוב הסינתטי שלהם למעבר הפרופיילינג. זה נורמלי בעבודה על קרנל וזה עדיין מקור יחיד. צוות שני שמשיג 1,109, או שנכשל בכך, הוא העובדה שתכריע את העניין.

• מהדורת SGLang ששוחררה ושכוללת אותו. הפוסט מפרסם קוד ומתאר את העיצוב; הוא אינו אומר שה-megakernel החליף את נתיב הפענוח המוגדר כברירת מחדל עבור המודל הזה בגרסה ששוחררה. עד שזה יקרה, אימוצו משמעותו אימוץ של מאגר מחקרי.

• סגירת הפער בין GEMM-2 ל-combine. המחברים כבר הצביעו על 12.5 µs של זמן בטלה לכל בלוק שם כעל כמעט מחצית ממה שנותר. אם השרשרת הזו תתוקן, אותה חומרה תפיק זמן צעד טוב יותר באופן משמעותי, וה-2.1× יהפוך לרצפה של מספר גדול יותר ולא לכותרת הראשית.

Nothing in that list is a reason to discount the result. A team that writes 8,000 lines of CUDA and PTX to move a step from 12.9 ms to 5.9 ms, publishes the code, and then publishes the exact places where the remaining time goes is doing the thing the field claims to value and rarely does. The correct summary is not that M​use Spark Open got faster. It is that a specific serving stack, on specific hardware, at a specific batch size, got 2.1× faster, and the recipe for the next 2× is now public.