DeepGEMM - مكتبة كفؤة لضرب المصفوفات بدقة FP8

مكتبة واضحة وفعالة لضرب المصفوفات بدقة FP8 مع تدرج دقيق، مُحسّنة لبنية NVIDIA Hopper

DeepGEMM Matrix Multiplication Visualization

تصميم بنية DeepGEMM

يتميز DeepGEMM ببنية مصممة بعناية ومُحسّنة لنواة التنسور من NVIDIA Hopper، مما يتيح ضرب المصفوفات بدقة FP8 بكفاءة.

DeepGEMM Architecture Diagram

تصميم JIT كامل

يستخدم DeepGEMM تصميم Just-In-Time (JIT) كامل بدون تجميع وقت التثبيت. يتم تجميع جميع النوى في وقت التشغيل باستخدام وحدة JIT خفيفة، مع معالجة أشكال GEMM وأحجام الكتل ومراحل خط الأنابيب كثوابت وقت التجميع، مما يوفر السجلات ويسمح بمزيد من تحسينات المُجمّع.

ميزات Hopper TMA

يستفيد بشكل كامل من مسرّع ذاكرة التنسور (TMA) في بنية Hopper لنقل البيانات بشكل أسرع وغير متزامن، بما في ذلك تحميلات TMA والتخزين والبث وجلب الواصفات المسبق، مما يعزز الأداء بشكل كبير.

تقنيات التحسين الأساسية

يستخدم DeepGEMM العديد من تقنيات التحسين المتقدمة التي تتجاوز مكتبات GEMM التقليدية لتحقيق أداء استثنائي.

DeepGEMM Optimization Techniques

تخصيص المؤشرات الدائمة

باتباع تصميم CUTLASS، يتداخل نقل البيانات وتعليمات MMA للنواة التنسورية وترقية نواة CUDA لتحسين كفاءة الحساب.

أحجام كتل غير متوافقة

يدعم أحجام كتل غير متوافقة (مثل 112) لاستخدام SM بشكل أفضل. على سبيل المثال، مع M=256، N=7168، استخدام BLOCK_M=128، BLOCK_N=112 يسمح بعمل 128 SM بدلاً من 112 مع أحجام كتل متوافقة.

تداخل FFMA SASS

يعدل تعليمات FFMA في الملف الثنائي المجمع عن طريق قلب بتات yield وreuse، مما يخلق المزيد من الفرص لتداخل تعليمات MMA مع تعليمات ترقية FFMA، مما يؤدي إلى تحسين الأداء بنسبة +10% في بعض الحالات.

مجدول كتل موحد محسّن

مجدول واحد لجميع النوى غير المجمعة والمجمعة، يستخدم التنقيط لتعزيز إعادة استخدام ذاكرة التخزين المؤقت L2، مما يحسن الأداء العام.

شرح الكود

تتكون وظيفة النواة الأساسية لـ DeepGEMM من حوالي 300 سطر من الكود فقط، بتصميم نظيف يسهل تعلم ضرب المصفوفات بدقة FP8 على Hopper وتقنيات التحسين.

// Simplified version of the core GEMM kernel
template <
    int BLOCK_M,
    int BLOCK_N,
    int BLOCK_K,
    int CLUSTER_M,
    int CLUSTER_N,
    int STAGES,
    int WARPGROUP_SIZE,
    bool ENABLE_INTERLEAVE
>
__global__ void gemm_fp8_fp8_bf16_nt_kernel(
    const void* __restrict__ lhs,
    const void* __restrict__ lhs_scale,
    const void* __restrict__ rhs,
    void* __restrict__ output,
    const void* __restrict__ rhs_scale,
    int m, int n, int k,
    int lhs_stride, int rhs_stride, int output_stride) {
    
    // Using TMA to load LHS, LHS scale factors, and RHS
    // Using CUDA cores for two-level accumulation (promotion)
    // Using TMA to store output
    
    // Persistent thread specialization implementation
    // Unaligned block size support
    // FFMA SASS interleaving optimization
}

ميزات بنية الكود

  • تصميم قالبي يدعم أحجام كتل مختلفة وتكوينات تحسين
  • استخدام كامل لميزات TMA في بنية Hopper لنقل البيانات بكفاءة
  • تراكم ثنائي المستوى لحل مشكلة التراكم غير الدقيق لنوى التنسور FP8
  • تنفيذ تخصيص المؤشرات الدائمة للحساب المحسّن وتداخل نقل البيانات
DeepGEMM Code Flow Diagram

تصميم الواجهة

  • GEMM كثيف عادي: باستخدام الدالة deep_gemm.gemm_fp8_fp8_bf16_nt
  • GEMM مجمع (تخطيط متصل): باستخدام الدالة m_grouped_gemm_fp8_fp8_bf16_nt_contiguous، مناسب لنماذج MoE حيث يشترك الخبراء في نفس الشكل
  • GEMM مجمع (تخطيط مقنع): باستخدام الدالة m_grouped_gemm_fp8_fp8_bf16_nt_masked، تستخدم في مرحلة فك ترميز الاستدلال مع تمكين رسومات CUDA حيث لا تعرف وحدة المعالجة المركزية عدد الرموز التي يتلقاها كل خبير

الأداء

أداء DeepGEMM قابل للمقارنة أو أفضل من المكتبات المضبوطة من قبل الخبراء لمختلف أشكال المصفوفات، مع تسريع يصل إلى 2.7x في بعض الحالات.

DeepGEMM Performance Comparison

أبرز نقاط الأداء

  • GEMM العادي: تسريع يصل إلى 2.7x في أشكال معينة
  • GEMM المجمع بتخطيط متصل: تسريع يصل إلى 1.2x
  • GEMM المجمع بتخطيط مقنع: أيضًا تسريع يصل إلى 1.2x
  • أداء حسابي يصل إلى 1358 TFLOPS تم اختباره على H800 SXM5
  • استخدام عالٍ لعرض النطاق الترددي للذاكرة، يصل إلى 2668 جيجابايت/ثانية

الميزات الرئيسية لـ DeepGEMM

أداء استثنائي

يوفر DeepGEMM تسريعًا يصل إلى 2.7x مقارنة بالمكتبات المضبوطة من قبل الخبراء على بنية NVIDIA Hopper، محققًا أكثر من 1350 TFLOPS

تحسينات متقدمة

يستخدم DeepGEMM تخصيص المؤشرات الدائمة وأحجام الكتل غير المتوافقة وتداخل FFMA SASS لتحقيق أقصى كفاءة حسابية

تكامل مرن

يتكامل DeepGEMM بسلاسة مع أطر التعلم العميق ومكتبات الحوسبة العلمية من خلال واجهات واضحة وموثقة جيدًا

تدرج دقيق

يتدرج DeepGEMM بكفاءة من المصفوفات الصغيرة إلى الحسابات واسعة النطاق باستخدام جدولة الكتل المحسنة وتقنيات التنقيط

دعم FP8 الأصلي

يوفر DeepGEMM دعمًا من الدرجة الأولى لدقة FP8، محسّن خصيصًا لنوى التنسور من NVIDIA Hopper

عرض نطاق ترددي عالٍ للذاكرة

يحقق DeepGEMM استخدامًا استثنائيًا لعرض النطاق الترددي للذاكرة يصل إلى 2668 جيجابايت/ثانية، مما يعظم قدرات الأجهزة

تصميم كود نظيف

يقدم DeepGEMM نواة أساسية موجزة من 300 سطر مع أنماط تصميم واضحة، مما يسهل فهمها وتوسيعها

GEMM مجمع محسّن

يتفوق DeepGEMM في عمليات GEMM المجمعة بتخطيط متصل ومقنع، مثالي لنماذج مزيج الخبراء (MoE)

الأسئلة الشائعة

ما هو DeepGEMM وما الذي يجعله مميزًا؟
DeepGEMM هو مكتبة فعالة لضرب المصفوفات بدقة FP8 محسّنة لبنية NVIDIA Hopper. ما يجعله مميزًا هو تصميمه النظيف (حوالي 300 سطر فقط من الكود الأساسي)، والأداء الاستثنائي (تسريع يصل إلى 2.7x مقارنة بالمكتبات المضبوطة من قبل الخبراء)، وتقنيات التحسين المتقدمة مثل تخصيص المؤشرات الدائمة وتداخل FFMA SASS.
لأي أجهزة تم تحسين DeepGEMM؟
تم تحسين DeepGEMM خصيصًا لبنية NVIDIA Hopper، مستفيدًا بشكل كامل من ميزات مسرّع ذاكرة التنسور (TMA) من Hopper ونوى التنسور لضرب المصفوفات بدقة FP8. تم اختباره على وحدات معالجة الرسومات H800 SXM5، محققًا أداءً حسابيًا يصل إلى 1358 TFLOPS.
كيف يحقق DeepGEMM مكاسب الأداء الخاصة به؟
يحقق DeepGEMM أداءه من خلال عدة تقنيات متقدمة: تصميم JIT كامل مع ثوابت وقت التجميع، وأحجام كتل غير متوافقة لاستخدام SM بشكل أفضل، وتداخل FFMA SASS، وتراكم ثنائي المستوى لحل مشاكل دقة FP8، ومجدول كتل موحد محسّن مع تنقيط لتعزيز إعادة استخدام ذاكرة التخزين المؤقت L2.
ما أنواع عمليات GEMM التي يدعمها DeepGEMM؟
يدعم DeepGEMM ثلاثة أنواع رئيسية من عمليات GEMM: GEMM كثيف عادي (باستخدام الدالة deep_gemm.gemm_fp8_fp8_bf16_nt)، وGEMM مجمع بتخطيط متصل، وGEMM مجمع بتخطيط مقنع. النوعان الأخيران مفيدان بشكل خاص لنماذج مزيج الخبراء (MoE).
هل DeepGEMM مناسب لنماذج مزيج الخبراء (MoE)؟
نعم، DeepGEMM مناسب للغاية لنماذج MoE. فهو يوفر عمليات GEMM مجمعة متخصصة بتخطيطات متصلة ومقنعة، والتي تم تحسينها لأنماط الحساب المحددة في بنيات MoE. يمكن لهذه العمليات أن توفر تسريعًا يصل إلى 1.2x مقارنة بالمكتبات الأخرى.
ما الدقة التي يدعمها DeepGEMM؟
يركز DeepGEMM بشكل أساسي على دقة FP8، والتي تزداد أهميتها للاستدلال الفعال لنماذج الذكاء الاصطناعي. ينفذ نهج تراكم ثنائي المستوى لحل مشاكل التراكم غير الدقيق التي يمكن أن تحدث مع نوى التنسور FP8، مما يضمن كلاً من الأداء والدقة.
كيف تقارن استخدام عرض النطاق الترددي للذاكرة في DeepGEMM مع المكتبات الأخرى؟
يحقق DeepGEMM استخدامًا استثنائيًا لعرض النطاق الترددي للذاكرة، يصل إلى 2668 جيجابايت/ثانية. هذا ممكن من خلال استخدامه الفعال لميزات TMA من Hopper لنقل البيانات بشكل أسرع وغير متزامن، بما في ذلك تحميلات TMA والتخزين والبث وجلب الواصفات المسبق.
هل يمكنني التعلم من تنفيذ DeepGEMM لتحسين كود CUDA الخاص بي؟
بالتأكيد! النواة الأساسية لـ DeepGEMM تتكون من حوالي 300 سطر فقط من الكود بتصميم نظيف، مما يجعلها مصدرًا تعليميًا ممتازًا. فهي توضح تقنيات تحسين CUDA المتقدمة، والاستخدام الفعال لميزات بنية Hopper، والنهج الذكية لضرب المصفوفات التي يمكن تطبيقها على مهام الحوسبة عالية الأداء الأخرى.