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.

هذه هي التكلفة التي يهاجمها الميغاكيرنل، ويستعين المؤلفون بتشبيه من ستينيات القرن العشرين لشرحها: كانت الآلات المبكرة تشغّل مهام الدفعات، إذ تمتلك مهمة واحدة الآلة بأكملها، ويظل المعالج خاملًا بينما يدور شريط. عالجت البرمجة المتعددة ذلك عبر إبقاء عدة مهام مقيمة والتبديل إلى أخرى كلما تعطّلت المهمة الحالية. والميغاكيرنل هو تلك الفكرة منقولة إلى GPU — إذ تبقى نواة واحدة مقيمة على كل SM وتشغّل مهام كثيرة في آن واحد، بحيث تستطيع الـ warps بدء عمل العملية التالية قبل انتهاء العملية السابقة.

الأثر المقاس، من الشكل المرافق لقسم الأداء، هو خطوة فك تشفير تنتقل من 12.9 مللي ثانية إلى 5.9 مللي ثانية. ويقسم ذلك الشكل الخطوة إلى نصفَي المسودة والتحقق: 1.16 مللي ثانية و10.9 مللي ثانية مع 1,450 نواة في خط الأساس، و0.44 مللي ثانية و5.3 مللي ثانية مع 97 نواة في النواة الضخمة. إن الانتقال من 12.9 إلى 5.9 يساوي نحو 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، تُطلق في مجموعات من أربع، باستخدام ذاكرة مشتركة موزعة حيث يتوفر اختزال أقل زمن استجابة. لا تكون الـ warps الاثنتا عشرة قابلة للتبادل. يحمّل warp 0 الـ KV والأوزان، ويقضي وقت فراغه في الجلب المسبق لأوزان العملية التالية. تشغّل warps 1 و2 التنشيطات وأنوية tensor. تتقاسم warps 3 حتى 11 الانتباه، وGEMMs الكثيفة، وGEMM الخبراء، واستقبالات all-reduce وall-gather. في طبقات النافذة المنزلقة، تحمّل warps 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 كيلوبايت من التعليمات، وهذه النوى البرمجية يبلغ حجم كل منها من 215 إلى 314 كيلوبايت من الشيفرة الآلية، لأن كل واحدة منها تحتوي على جميع أطوار الطبقة. عند الإطلاق، لا يُقرَّب من وحدات SM سوى أول 32 إلى 64 كيلوبايت؛ وكل ما يتجاوز هذه النافذة يُجلب سطرًا بسطر عند أول استخدام. أجرى المؤلفون تجربة نظيفة على هذا وأبلغوا عنها: كتلتان بحجم 100 كيلوبايت استغرقت كل منهما 12.4 ميكروثانية منفردة، واستغرقتا 19.5 ميكروثانية عند وصلهما الواحدة تلو الأخرى، وتقليص إحدى الكتلتين إلى 20 كيلوبايت أعاد الأخرى إلى سرعتها المنفردة. تمر التعليمات عبر ذاكرة L2 نفسها التي تمر عبرها البيانات، لذا فإن نواة برمجية يتجاوز حجمها ذاكرة التخزين المؤقت للتعليمات ينتهي بها الأمر إلى التنافس مع أحمال الأوزان الخاصة بها.

الفاتورة الثانية هي الجلب المسبق للأوزان، ويعالجها المقال باعتبارها مشكلة توزيع لها ثلاثة حلول. تحصل الذاكرة المشتركة على حلقة بسعة 5 × 32 كيلوبايت لكل SM، أي نحو 23 ميغابايت عبر الـ GPU. أما الباقي فيذهب إلى L2، وتحتوي ذاكرة L2 لدى GB300 على نحو 135 ميغابايت، لكن أي شيء في L2 يحتاج إلى رحلة ثانية أقصر إلى الذاكرة المشتركة، كما أن L2 مشتركة مع كل شيء آخر قيد التنفيذ. ويتكدّس الاثنان: قرب نهاية الطبقة، تحتوي الحلقة بالفعل على أوزان qkv الأولى للطبقة التالية. وفوق ذلك، يقرأ warp 3 الخبراء الذين وجّهت إليهم الخطوة السابقة ويسحب شرائح هذا الـ GPU من الخبراء الـ14 الأكثر استخدامًا إلى L2، بنحو 50 ميغابايت لكل GPU، راهنًا على أن الخطوات المتتالية توجّه بالطريقة نفسها. هذا الرهان يساوي نحو 0.2 ms من أصل خطوة مدتها 5.3 ms، وقد يكلّف عرض النطاق عندما يكون خاطئًا.

الفاتورة الثالثة هي المزامنة، وهي المسألة التي لا يوجد لها جواب نظيف. عندما ينتظر warp ما على GPU أخرى عبر NVLink، لا توجد مقاطعة ولا انتظار على عنوان في العتاد؛ بل يستقصي الـ warp راية حتى تتغير. تُقدّر القياسات المصغرة في المقال استقصاء acquire على نطاق GPU بـ 141 إلى 145 نانوثانية، وراية أحادية الاتجاه بين اثنين من الـ GPU على العقدة نفسها بنحو 1.9 ميكروثانية، واستقصاءً مصطفًا خلف نسخ جماعية بـ 1.0 إلى 2.6 ميكروثانية. تكلف الحواجز (fences) 0.4 ميكروثانية عند الخمول و3.6 إلى 4.0 ميكروثانية مع أحمال TMA قيد التنفيذ. وحتى مع توزيع مثالي للعمل، يمكن أن يتخلف بعض الـ GPU عن بعضها بما يصل إلى 4 ميكروثانية، وكل نقطة توقف في الطبقة تنتظر أبطأها. تستبدل النواة الضخمة (megakernel) الحواجز بقيم حارسة تُكتب في خانات تخزين عادية — راية مملوءة بالبيانات، و+0/−0 من أجل all-reduce، وNaN من كل الآحاد من أجل تجميعات bf16 — وتعيد ضبط النسخة الخاملة من كل خانة بينما تكون الأخرى قيد الاستخدام.

هذه هي الأرقام التي تجعل النتيجة تبدو هندسة لا سحرًا: انحراف GPU بمقدار 4 µs مقابل خطوة مدتها 5.9 ms صغير، لكن أربع نقاط توقف للمزامنة لكل طبقة مضروبة في 62 طبقة هو بالضبط نوع التكلفة الإضافية التي تلتهم فوزًا بمعامل 2× إذا أخطأت في ذلك.

على ماذا لا يزال يُنفَق 5.9 مللي ثانية

يتضمن المنشور نقدًا لنفسه، وأكثر عبارة مفيدة فيه تتعلق بالفجوة بعد GEMM الثاني. تنهي الكتل ذلك GEMM في أي وقت بين 55 و64 µs من بداية الطبقة، عند المئين العاشر إلى التسعين، لكن خطوة الدمج تنتظر مخرجات كل GPU، لذا لا يمكنها البدء حتى 65 µs. ومن نهاية GEMM 2 إلى نهاية الطبقة تظل الكتلة بلا عمل لنحو 12.5 µs في المتوسط — أي ما يقارب نصف إجمالي وقت خمولها. تلك السلسلة، وليس عبء إطلاق النواة الذي أزالته النواة الضخمة بالفعل، هي ما يقول المؤلفون إنهم ينظرون إليه لاحقًا.

كما يشرحون لماذا لا يمكن لانتباه الطبقة التالية أن يبدأ مبكرًا ببساطة. يُطلق PDL هذا الانتباه عند 79 µs، أي قبل خروج أول كتلة بنحو 4 µs، لكن كتله لا تحطّ إلا بعد أن تغادر الكتل السابقة وحدات SMs الخاصة بها عند 84 µs، لذا فإن الإطلاق في وقت أبكر لن يبدأ أي شيء أبكر. يتمثل الحل البديل في الجلب المسبق لصفوف KV القديمة الخاصة به إلى L2 عند 65 µs، حيث تنتظر. وهذا هو النمط نفسه المتبع في الجلب المسبق للأوزان: عندما لا تستطيع إزالة انتظار، تنقل العمل إلى داخله.

التنميط هو المشكلة البنيوية الأخرى، وهو مشكلة موفقة: إذ تبدو النواة لأي أداة تنميط خارجية كشريط طويل واحد لكل طبقة، لذا فهي تقيس زمنها بنفسها. يكتب المسار 0 في كل كتلة طابعًا زمنيًا إلى مصفوفة عامة عبر سجل globaltimer، ويقيس المؤلفون كلفة أدوات القياس هذه بنحو 2.5 µs لكل إطلاق طبقة، أي 2.8% من خطوة، مع الإبقاء على المخرجات متطابقة على مستوى البت ووجود علم بناء لاستبعادها عند الترجمة. النواة التي يتعين عليها حمل أداة تنميطها الخاصة هي نواة صارت نظام تشغيل، وهو بالضبط الإطار الذي يستخدمه المنشور: العملية هي نواة واحدة تمتلك جميع موارد GPU، والخيوط هي warps بلا استباق، وانتظار الإدخال/الإخراج هو انتظار مشغول على علم، وDMA هي TMA يتتبعها mbarrier، والرسائل مجرد عمليات تخزين عادية عبر 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.

طبّق الفريق المعالجة نفسها بالنواة الضخمة (megakernel) على المسودة: تعمل كل طبقة من طبقات المسودة الخمس بعملية إطلاق واحدة، إضافة إلى ثلاث أخرى، ما ينقل المسودة من 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، على ثماني وحدات GB300s، بحجم دفعة واحد وبتشغيل فك الترميز التخميني. يذكر المنشور بصراحة أن الـmegakernel يدعم اليوم حجم دفعة واحد فقط، وأن توسيعه ليشمل دفعات أكبر هو العمل التالي.

تزيل كل من هذه القيود ادعاءً مختلفًا. حجم الدفعة واحد يعني أن هذه ليست نتيجة إنتاجية بالمعنى الذي يستخدمه فريق تقديم الخدمة للكلمة: فلا يوجد فيها مكسب من التجميع في دفعات، ولا طابور، ولا خليط من أطوال المطالبات. ومجموعة بيانات واحدة تعني أن الرقم مرتبط بتوزيع المطالبات والمخرجات في GSM8K، وبنمط توجيه مُقاس لـ49 خبيرًا لكل طبقة استخدمه المؤلفون بعد ذلك كتوجيه اصطناعي في التوصيف الخاص بهم. أما عشرون خطوة لطلب واحد فتعني أن التباين محدود بعدد الخطوات، لا بعدد الطلبات.

وطلب واحد هو أسوأ حالة للتداخل، وهو ما يقول المؤلفون أنفسهم: «عند طلب واحد يكون كل انتظار مكشوفًا؛ ومع المزيد، يكاد يكون هناك دائمًا عمل جاهز.» لذا فإن نسبة 2.1× للـ megakernel هي قياس للنظام الذي يكون فيه إخفاء الكمون أصعب، والمنشور صريح في أن الدفعات الأكبر ينبغي أن تكون أسهل لا أصعب. وهذا حجة لصالح النتيجة، وهو لا يزال حجة وليس قياسًا.

ما يمكن معرفته اليوم، إذن، ضيّق وواقعي: الشيفرة موجودة في مستودع عام، والمنشور يصف الآلية بتفصيل كافٍ للتحقق منها، والحساب الداخلي متسق عبر ثلاث نسب مستقلة، وقد وثّق المؤلفون التكاليف الإضافية الثلاثة التي تحدّ من مزيد من المكاسب. وما لا يمكن معرفته هو ما إذا كان فريق مستقل يحصل على 1,109 على ثماني وحدات GB300 لديه، وما إذا كان النهج ينجو من دفعة حقيقية، وما إذا كان يستقر في إصدار SGLang مُصدَر — فالمنشور لا يزعم أنه فعل ذلك، والمستودع أثر بحثي وليس حزمة يمكن تثبيتها. والمقارنة التي من شأنها أن تحسم أكبر قدر من الغموض هي مقابل خط أساس مُعاير على العتاد نفسه، لأن «سلسلة SGLang الافتراضية» هدف متحرك، والفريق الذي كتب الـmegakernel هو نفسه من يعاير خط الأساس الذي يتفوق عليه.

كيف يرتبط ذلك بـ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 رمزًا في الثانية، واستجاب مكدس مكوّن يدويًا من ثماني وحدات GPU بمعدل 1,109، فإن الفرق ليس سببًا لبناء المكدس. ثماني وحدات GB300s عملية شراء رأسمالية تُقاس بالملايين، والنواة هي شيفرة بحثية، وقيود حجم الدفعة واحد تعني أن المكدس جيد حاليًا في بالضبط عبء العمل الذي تتعامل معه نقطة نهاية مستضافة جيدًا بالفعل: طلب واحد في كل مرة. حيث يغيّر الـ megakernel الحساب حقًا هو لفريق اشترى بالفعل وحدات GPU للتدريب أو لعبء عمل له قيود على إقامة البيانات، ويسأل الآن عن مقدار الاستفادة التي يمكنه الحصول عليها منها عند دفعة منخفضة. لذلك الفريق، 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-to-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.