
هندسة عكسية لقاموس تعليمات NVIDIA SASS وتدقيق النواة (kernel) والتعرف على الأنماط عبر معماريات GPU.
هندسة عكسية لـ NVIDIA SASS من النوى المتحكم بها إلى تدقيقات الإنتاج.
المقال 1 · المقال 2 · قاعدة المعرفة · مكتبة الأنماط · قاموس تعليمات SM120 · ملاحظات الترميز · ابدأ من هنا · هيكل المشروع · فصول نوى التنسور · المساهمة
SASS King هو مشروع منهجي للهندسة العكسية لـ NVIDIA SASS، وهي مجموعة التعليمات الأصلية لوحدة معالجة الرسوميات (GPU) المنبعثة داخل ملفات CUDA الثنائية المترجمة. يبدأ المشروع بأجهزة Blackwell الاستهلاكية SM120 / SM120a ويتوسع نحو مكتبة ISA شاملة عبر البنى ومكتبة أنماط بمرور الوقت.
الهدف عملي: مساعدة مهندس النوى على فتح تفريغ SASS، والتعرف على أنماط المترجم، وتحديد الهياكل ذات الصلة بالأداء، وربط الثنائي بقرارات التحسين على مستوى المصدر.
أكمل المشروع مكتبة الأنماط الأولية للمرحلة الثالثة: 29 توقيع SASS قابل لإعادة الاستخدام موثّقة الآن تحت patterns/، مع الاحتفاظ بـ knowledge/FINDINGS.md كمسار أدلة كامل. الخطوة الرئيسية التالية هي المرحلة الرابعة: تطبيق تلك الأنماط على نوى الإنتاج الحقيقية.
المستودع منظم كخط أنابيب للأدلة:
corpus/ نوى محكومة وأدلة SASS خام
knowledge/ نتائج على مستوى المشروع، ملاحظات التعليمات، وملاحظات الترميز
patterns/ توقيعات تدقيق قابلة لإعادة الاستخدام من المرحلة الثالثة
production/ تدقيقات نوى حقيقية من المرحلة الرابعة
آخر عمل واسع النطاق للهندسة العكسية لـ SASS العامة المشابه في الروح كان لـ Jia et al. على Volta و Turing في 2018. لقد غيرت Ampere و Hopper و Blackwell مزيج التعليمات بشكل جوهري: مسارات النسخ غير المتزامن، عائلات نوى التنسور، تعليمات تحميل/تخزين المصفوفات، أشكال MMA المتناثرة والموزونة، وتدفقات السجلات الموحدة الجديدة.
يسد SASS King هذه الفجوة من خلال الجمع بين النوى الدقيقة المتحكم بها، وقراءة SASS الخام، والمسابير الزمنية، وتدقيقات نوى الإنتاج.
مكتبة الأنماط الرسمية هي الناتج الرئيسي للمرحلة الثالثة. تحوّل الأدلة المحلية لكل فصل إلى توقيعات تدقيق قابلة لإعادة الاستخدام، بحيث يمكن للتدقيق أن يستشهد بنمط مسمى بدلاً من إعادة كتابة مسار البحث الكامل في كل مرة.
تعتبر المرحلة الثالثة مكتملة لأن:
knowledge/FINDINGS.md؛patterns/README.md؛تشمل كل صفحة نمط:
استخدم patterns/README.md كفهرس مخصص للتدقيق. استخدم knowledge/FINDINGS.md عندما تحتاج إلى سياق البحث الأطول وراء نمط ما.
لا تدّعي المرحلة الثالثة أن كل سلوك NVIDIA SASS قد تم فك شفرته. إنها تؤسس طبقة أنماط SM120 / SM120a قابلة لإعادة الاستخدام كافية لبدء تدقيقات الإنتاج اليدوية. فك ترميز التخطيط الزمني، وتحديد مواضع بتات رموز التحكم الكاملة، وإعداد تقارير cubin الآلية، وإعادة التشغيل عبر البنى تظل أعمالاً مستقبلية.
كتابات عامة:
التباين المتحكم به. نواتان تختلفان بمتغير واحد بالضبط: dtype، ترتيب المعاملات، عامل فك اللف، تخطيط الذاكرة، أو هدف الترجمة. يحدد فرق SASS قرار المترجم.
وسوم ادعاءات صارمة. كل ادعاء تقني يستخدم وسماً:
من الأعلى إلى الأسفل ومن الأسفل إلى الأعلى معاً. تقوم النوى الدقيقة بعزل التعليمات الفردية وقرارات المترجم. تظهر النوى الشبيهة بالإنتاج أي الأنماط مهمة في الكود الحقيقي.
تدقيقات تعتمد على الأنماط. يجب أن يستشهد تدقيق الإنتاج بصفحة PATTERN-NN رسمية فقط بعد مطابقة توقيع SASS المرئي ونقل حدود ثقته وأنماطه المضادة وفجواته المفتوحة.
يركز المسح الأول على خط أنابيب نوى التنسور والذاكرة لـ SM120:
HMMA، QMMA، OMMALDSM، STSMLDGSTS، LDGDEPBAR، DEPBARLDG، STG، LDS، STS، REDGBRA، EXIT، BSSY، ، لا يدّعي المشروع أن ISA مكتملة بعد. يتتبع القاموس العام ما تم ملاحظته وشرحه؛ الصفحات الأعمق تحت knowledge/encoding/ تتبع العائلات التي لديها أدلة كافية لتوثيق شبيه بالمطابق.
لا ينافس SASS King مُفكّكات SASS على مستوى البت. يستخدم المشروع التفريغات المحلية كدليل أساسي وقد يستخدم redplait/denvdis كفحص متبادل لحقول التعليمات، جداول الجدولة، المسندات، وتتبع السجلات. يمكن لـ denvdis التحقق من تفسيرات الترميز منخفضة المستوى؛ SASS King يملك أدلة التباين المتحكم به، طبقة الأنماط الدلالية، وتفسير تدقيق الإنتاج.
flowchart LR
P1["المرحلة 1<br/>نوى تعليمية<br/>01-12"] --> P2["المرحلة 2<br/>مجموعة نوى التنسور SM120<br/>13-25"]
P2 --> P25["المرحلة 2.5<br/>التحقق المتبادل مع denvdis<br/>الجزء الخلفي على مستوى البت"]
P25 --> P3["المرحلة 3<br/>مكتبة الأنماط<br/>توقيعات المترجم"]
P3 --> P4["المرحلة 4<br/>تدقيقات الإنتاج<br/>نوى حقيقية"]
P4 --> P5["المرحلة 5<br/>أداة التدقيق<br/>تقارير cubin"]
P5 --> P6["المرحلة 6<br/>إعادة التشغيل عبر البنى<br/>SM80/86/89/90a/100a/120"]
classDef done fill:#0b6d55,color:#fff,stroke:#0b6d55;
classDef active fill:#f4c95d,color:#111,stroke:#b89422;
classDef planned fill:#1f2937,color:#fff,stroke:#6b7280;
class P1,P2,P25,P3 done;
class P4 active;
class P5,P6 planned;
تؤسس النوى 01-12 مفاهيم SASS الأساسية: دمج FMA، سلوك لوحة النتائج، خفض الحلقات، الذاكرة المشتركة، الذاكرة العالمية، أساسيات الالتفاف، الرياضيات بطيئة المسار، وانسكابات الذاكرة المحلية.
تغطي النوى 13-25 مسار نوى التنسور الحالي لـ SM120:
التحقق من redplait/denvdis كخلفية فحص متبادل على مستوى البت لـ SM120 / SM120a قبل أن تعتمد تدقيقات الإنتاج على مكتبة الأنماط. يقوم المسح بتشغيل nvd -O، nvd -S، nvd -p، وحيث يكون مفيداً nvd -T على cubins أو تفريغات محلية ممثلة تغطي HMMA، QMMA، QMMA.SF، QMMA.SP، OMMA، LDSM، STSM b16/b8، LDGSTS، DEPBAR، وعلامات التفرع.
المخرجات هي knowledge/DENVDIS_INTEGRATION.md: جدول توافق واقعي من العائلة إلى حالة التعرف في denvdis، تغطية المُعدِّلات، حقول رموز التحكم المكشوفة، وإجراء SASS King. مخرجات denvdis هي أدلة داعمة، وليست بديلاً عن ملاحظات التفريغ المحلية.
صياغة الهياكل المتكررة في توقيعات قابلة لإعادة الاستخدام:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMA المتسلسلةSTSM -> BAR -> LDS -> STGتحتوي مكتبة المرحلة الثالثة الأولية على 29 صفحة نمط تحت patterns/. يبقى knowledge/FINDINGS.md سجل البحث ومصدر الحقيقة؛ patterns/ هي نقطة الدخول المخصصة للتدقيق.
المرحلة الثالثة مكتملة على مستوى المكتبة الأولية. العناصر المتبقية مثل فك ترميز التخطيط الزمني، ومواضع بتات رموز التحكم الكاملة، وإعادة التشغيل عبر البنى تُتتبع كفجوات أو مراحل مستقبلية، وليست عوائق لبدء تدقيقات الإنتاج في المرحلة الرابعة.
تطبيق مكتبة الأنماط على نوى حقيقية من مكتبات مثل FlashAttention و CUTLASS و xFormers و Transformer Engine و FlashInfer و llama.cpp / ggml و tinygrad والمشاريع ذات الصلة. الهدف هو تغطية تمثيلية حسب النمط الخوارزمي، وليس ملف Markdown واحد لكل نواة.
يجب أن يكون أول تسليم للمرحلة الرابعة تقرير تدقيق يدوي يقوم بما يلي:
PATTERN-NN المطابقة؛بناء خط أنابيب يأخذ cubin، ويكتشف الأنماط المعروفة، ويصدر تقريراً موجهاً نحو التحسين.
إعادة تطبيق المنهجية على أهداف إضافية:
.
├── corpus/ # نوى محكومة، تفريغات، وكتابات الفصول
│ ├── basics/ # النوى 01-08: أساسيات العددية/المتجهة والذاكرة
│ ├── warp_collectives/ # النوى 09-10: الخلط، التصويت، الاختزال
│ ├── math_and_spills/ # النوى 11-12: مسارات بطيئة وانسكابات
│ └── tensor_cores/ # النوى 13-25: دراسات نوى التنسور
├── knowledge/ # النتائج، القاموس، ملاحظات الترميز
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # مكتبة الأنماط الرسمية للمرحلة الثالثة
├── production/ # تدقيقات نوى الإنتاج للمرحلة الرابعة
├── docs/ # ملاحظات الإعداد والهيكل والإصدار
└── guide/ # الوحدة الفرعية لدليل قراءة SASS الخارجي
يحتوي كل مجلد فصل على نوى المصدر، وأعمال مترجمة عند الاقتضاء، وتفريغات SASS عندما تكون جزءاً من مجموعة الأدلة المُتحقق منها، وكتابة conclusion<N>.md.
للحصول على شرح أكثر اكتمالاً لما ينتمي إلى كل دليل، اقرأ هيكل المشروع.
cuobjdump --dump-sass للتفكيك الخام.gpuasm.com للوحات النتائج، والتوقفات، والضغط، وأسهم التبعية.%clock لقياسات زمن انتقال التعليمات.nvcc -Xptxas -v لبيانات السجلات والانسكابات.يعمل SASS King في طبقة الأنماط الخوارزمية: التعرف على كيفية تنظيم النوى المترجمة وربط تلك الهياكل بقرارات التحسين على مستوى المصدر.
المساهمات مرحب بها، خاصة:
انظر CONTRIBUTING.md للبيانات الوصفية المتوقعة ومعيار الكتابة.
Florian Mattana. florianmattana.com
| إذا كنت تريد... | ابدأ من هنا | ثم اقرأ |
|---|
| فهم المشروع في 10 دقائق | docs/README.md | docs/START_HERE.md، ثم docs/PROJECT_STRUCTURE.md |
| إعادة إنتاج الأدلة | corpus/README.md | فصل واحد conclusion*.md، ثم تفريغ .sass الخاص به |
| العثور على مصدر الحقيقة | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md، knowledge/encoding/README.md |
| التعرف على نمط في تفريغ جديد | patterns/README.md | صفحة patterns/NN-*.md المطابقة |
| بدء تدقيق إنتاج | production/README.md | صفحات PATTERN-NN المطابقة وأدلة المصدر |
| المساهمة بتصحيح أو تفريغ | CONTRIBUTING.md | docs/START_HERE.md |
| المجال | الحالة | المكان |
|---|
| نوى تعليم SM120 | مكتمل من خلال النوى 01-12 | corpus/basics/01_vector_add/ إلى corpus/math_and_spills/12_register_spill/ |
| دراسات نوى التنسور | مكتمل حتى النواة 25 | corpus/tensor_cores/ |
| النتائج العامة | مصدر حقيقة نشط | knowledge/FINDINGS.md |
| قاموس تعليمات SM120 | نشط، مدعوم بالأدلة | knowledge/SASS_INSTRUCTIONS_SM120.md |
| نماذج الترميز | بدأت مع LDSM، STSM، QMMA | knowledge/encoding/ |
| التحقق المتبادل مع denvdis | تم المسح الأولي؛ فجوات في رموز التحكم الأعمق باقية | knowledge/DENVDIS_INTEGRATION.md |
| مكتبة الأنماط | نموذج أولي للمرحلة الثالثة مكتمل | patterns/ |
| تدقيقات الإنتاج | المرحلة التالية | production/ |
| عائلة الأنماط | أمثلة | المكان |
|---|
| حسابات نوى التنسور | سلاسل تراكم HMMA، QMMA، OMMA؛ بيانات وصفية متناثرة؛ أجزاء ضيقة | patterns/02-* إلى patterns/04-*، patterns/10-*، patterns/21-* |
| ذاكرة المصفوفات والنهايات | LDSM، STSM، خطوط أنابيب النسخ غير المتزامن، نهايات اختزال REDG | patterns/05-*، patterns/06-*، patterns/07-*، patterns/28-* |
| تدفق التحكم | التفرع/إعادة التقارب، حواف الحلقة الخلفية، مخارج مشروطة، مصائد باردة، استدعاءات محلية | patterns/08-*، patterns/14-*، patterns/16-*، patterns/26-*، patterns/29-* |
| الذاكرة والسجلات | ذاكرة عامة متجهة، انسكابات، تخزين مؤقت للذاكرة المشتركة، واصفات، تدفق سجلات موحدة | patterns/09-*، patterns/11-*، patterns/17-*، patterns/19-*، patterns/20-* |
| الحساب والجدولة | دمج FFMA، ثوابت، مسارات بطيئة MUFU، لوحات النتائج، إعادة تدوير العمر | patterns/12-*، patterns/18-*، patterns/22-*، patterns/23-*، patterns/24-* |
| جماعات الالتفاف | اختزالات الالتفاف، أساسيات الخلط/التصويت/المطابقة/المزامنة | patterns/01-*، patterns/25-* |
| الوسم | المعنى |
|---|
[OBS] | ملاحظ مباشرة في تفريغ أو سجل أو مخرجات زمنية أو ملف تعريفي. |
[INF] | مستنتج من أدلة ملاحظة. |
[HYP] | معقول لكن غير مؤكد. |
[RES] | فرضية سابقة تم حلها بأدلة لاحقة. |
[GAP] | سؤال مفتوح موثق صراحة. |
BSYNCWARPSYNCSHFL، VOTE، REDUXS2UR، R2UR، UMOV، ULEA، LDCU| المرحلة | الحالة | المخرجات | لماذا يهم |
|---|
| 1. نوى تعليمية | تم | corpus/basics/، corpus/warp_collectives/، corpus/math_and_spills/ | يؤسس مفردات القراءة من تجارب CUDA-to-SASS المتحكم بها. |
| 2. مجموعة نوى التنسور SM120 | تم | corpus/tensor_cores/13_hmma_fp16/ إلى 25_stsm_epilogue/ | يلتقط أول مجموعة أدلة لنوى التنسور، ذاكرة المصفوفات، تدفق التحكم، والنهايات لـ SM120 / SM120a. |
| 2.5. التحقق المتبادل مع denvdis | تم المسح الأولي | knowledge/DENVDIS_INTEGRATION.md، knowledge/encoding/CONTROL_CODE.md | يستخدم denvdis كفحص متبادل على مستوى البت دون استبدال أدلة التفريغ المحلية. تظل مواقع بتات التوقف/العائد الكاملة مفتوحة. |
| 3. مكتبة الأنماط | النموذج الأولي مكتمل | patterns/ | يحوّل هياكل المترجم/SASS المتكررة إلى توقيعات قابلة لإعادة الاستخدام. |
| 4. تدقيقات الإنتاج | التالي | production/ | يختبر ما إذا كانت أنماط المجموعة تشرح النوى الحقيقية من مكتبات الإنتاج. |
| 5. أداة التدقيق | مخطط | خط أنابيب من cubin إلى تقرير | يجعل طبقة الأنماط قابلة للبرمجة والتكرار. |
| 6. إعادة التشغيل عبر البنى | مخطط | مقارنات SM80، SM86، SM89، SM90a، SM100a، SM120 | يفصل الحقائق الخاصة بالهندسة عن سلوك NVIDIA SASS العام. |
| النواة | الموضوع |
|---|
| 13 | خط الأساس HMMA، تخصيص السجل، سلاسل التراكم |
| 14 | خط الأساس QMMA FP8 / FP6 / FP4 |
| 15 | متغيرات MMA الضيقة |
| 16 | ذروة FP4 و OMMA/QMMA ذات المقياس الكتلي |
| 17 | سلوك LDSM وتحميل المصفوفات |
| 18 | أنابيب بلاطات MMA والتخزين المؤقت للنسخ غير المتزامن |
| 19 | بيانات تعريف MMA المتناثرة |
| 20 | تدفق التحكم وكشف الحواف الخلفية |
| 21 | التفرع وإعادة التقارب |
| 22 | سلوك تخزين المصفوفات STSM |
| 23 | مسابير تخطيط أجزاء FP4 / FP6 |
| 24 | تدقيق مصغر لـ GEMM الإنتاج |
| 25 | تخطيط نهاية STSM ودلالات إعادة التخزين |
| الهندسة | GPU تمثيلي | لماذا |
|---|
| SM80 | A100 | خط الأساس Ampere لمراكز البيانات |
| SM86 | RTX 3090 | مجموعة Ampere الاستهلاكية |
| SM89 | RTX 4090 | بطاقة استدلال استهلاكية شائعة |
| SM90a | H100 | TMA، WGMMA، تخصص الالتفاف، العناقيد |
| SM100a | B200 | tcgen05.mma، TMEM |
| SM120 | RTX 5070 Ti / 5090 | نقطة بداية Blackwell الاستهلاكية |