من الميزات المعمارية إلى بناء النظام البيئي، يقوم Muxi Dong Zhaohua بتحليل عميق لممارسات تطبيق TVM على وحدات معالجة الرسومات المحلية

في الخامس من يوليو، اختُتم بنجاحٍ الدورة السابعة من ملتقى تقنيات مُجمّعي الذكاء الاصطناعي، الذي استضافته شركة HyperAI. بدءًا من الابتكار الأساسي في بنية وحدة معالجة الرسومات (GPU) وصولًا إلى التصميم عالي المستوى لبيئة التجميع عبر الأجهزة؛ ومن تحسين مُشغّل الشريحة الواحدة إلى الإنجازات في التجميع الموزع متعدد العقد... اجتمع الممارسون والباحثون في مجال تجميع الذكاء الاصطناعي في هذه الولائم التكنولوجية الرائدة. كان موقع الفعالية غزيرًا بالحضور، وكان جو التواصل قويًا.

اتبع حساب WeChat العام "HyperAI Super Neuro" وأجب على الكلمة الرئيسية "0705 AI Compiler" للحصول على عرض تقديمي للمحاضر المعتمد PPT.

في هذا الحدث، قدم لنا المهندس المعماري Zhang Ning من AMD تحليلًا متعمقًا لأسرار تحسين الأداء لمترجم Triton على منصة AMD GPU، وكشف عن كيفية جعل كود Python يتحكم بسهولة في نواة GPU عالية الأداء؛ قدم المدير Dong Zhaohua من Muxi Integrated Circuit خبرة عملية في تطبيقات TVM على وحدات معالجة الرسومات المحلية، مما يدل على شرارات الاصطدام بين الرقائق المستقلة وأطر التجميع مفتوحة المصدر؛ كشف الباحث Zheng Size من ByteDance عن لغز Triton الموزع وشارك كيف يقوض Python سقف أداء الاتصالات الموزعة؛ أعاد TileLang الذي قدمه الدكتور Wang Lei من جامعة بكين تعريف حدود كفاءة تطوير المشغل.

في الخطاب الرئيسي "ممارسة تطبيق TVM على وحدة معالجة الرسوميات Muxi"،قدم دونج تشاو هوا، المدير الأول لشركة Muxi Integrated Circuit، الخصائص التقنية لمنتجات وحدة معالجة الرسوميات الخاصة بها، وحلول تكييف مُجمِّع TVM، وحالات التطبيق العملية، ورؤية البناء البيئي.وقد أظهرت الإنجازات التكنولوجية وإمكانات تطبيق وحدات معالجة الرسوميات المحلية في مجالات الحوسبة عالية الأداء والذكاء الاصطناعي.

قامت شركة HyperAI بجمع وتلخيص خطاب البروفيسور دونغ تشاو هوا دون المساس بالقصد الأصلي. فيما يلي نص الخطاب.

مقدمة عن وحدة معالجة الرسومات Muxi

تتضمن وحدة معالجة الرسومات Muxi حاليًا خطوط إنتاج متعددة، مثل سلسلة N وسلسلة C وسلسلة G، تغطي مجموعة واسعة من السيناريوهات، من تدريب الذكاء الاصطناعي والاستدلال إلى الحوسبة العلمية. ومن خلال بناء حزمة برمجية متعددة المستويات، تحقق تكاملاً سلسًا مع الأطر البرمجية السائدة. وبصفتها الوحدة الأساسية في حزمة البرامج، تتولى وحدة التجميع مسؤولية توفير واجهة برمجة سهلة الاستخدام، وتحسين التطبيقات عالية المستوى، وتوليد أكواد الآلة المقابلة وفقًا لهياكل الآلات المختلفة، وتسليمها إلى وحدة معالجة الرسومات للتنفيذ. وبعد تعديلات دقيقة أجراها المهندسون، وصل أداؤها إلى المستوى المتقدم الدولي، وأقامت علاقة تكيف مماثلة مع مكتبات الحوسبة السائدة في هذا المجال.

تتمتع وحدة معالجة الرسوميات Muxi بواجهة وظيفية غنية على مستوى التعليمات.واجهة MACA C التي طورناها بأنفسنا مبنية على امتداد لغة C، وتدمج عناصر نحوية لمجالات محددة، وتتمتع بتكافؤ وظيفي مع واجهات البرمجة الأساسية للشركات المصنعة الرئيسية، مما يُمكّن المطورين من إتمام عملية الترحيل والتكيف بسرعة. في الوقت نفسه، توفر الواجهة واجهات برمجة متنوعة مثل بايثون وتريتون وفورتران، وتدعم معايير البرمجة المتوازية مثل OpenACC وOpenCL، وتتميز بكفاءة ممتازة في توليد أكواد التوازي التلقائي.

أيضًا،يعتمد Muxi GPU على بنية GPGPU (وحدة معالجة الرسوميات للأغراض العامة).يدعم نظام التجميع القائم على LLVM تحسين العملية الكاملة من اللغات عالية المستوى إلى أكواد الآلة منخفضة المستوى، ويأخذ في الاعتبار كفاءة التطوير وأداء الأجهزة، ويوفر مجموعة برامج عالية الأداء.

تكييف TVM على وحدة معالجة الرسومات Muxi

بصفته مُجمِّعًا مفتوح المصدر للتعلم العميق، يُمكن لـ TVM تحويل نماذج التعلم العميق إلى شيفرة برمجية تعمل بكفاءة على مختلف الأجهزة. وقد طوّر فريق Muxi حلاً متكاملاً لتكييف TVM استنادًا إلى خصائص وحدة معالجة الرسومات الخاصة به لتحقيق تحسين كامل للعملية، بدءًا من تعريف النموذج وحتى تنفيذه على الأجهزة.

من منظور هندسة المترجم، تم تحقيق الدعم الكامل.ومن الناحية النظرية، يمكن ربطه بأربعة مستويات أساسية.

لتكييف واجهة C++، نريد تحويلها إلى لغة MACA لحل المشكلة. هذه العملية صعبة للغاية، وهناك بعض التحديات في تنفيذ التحويل التلقائي القائم على الأدوات.

إذا كان تجريد الكود عاليًا، فسيكون من الأسهل تحقيق التكيف عبر المستويات. بالإضافة إلى ذلك، عند الاتصال بـ LLVM، يجب الانتباه إلى مشاكل توافق الإصدارات - نظرًا لتعدد إصدارات LLVM، يعتمد تكييف إصدار معين على دعم الإصدار المقابل، وقد يؤدي عدم تطابق الإصدارات إلى عمليات تجميع غير طبيعية.

من حيث تكييف وحدة معالجة الرسوميات لـ Muxi Arch:

يضيف tvm.Target هدف MACA ويضيف الدعم لكل مرحلة: أولاً، أضف خط أنابيب هدف MACA في التحويل/الخفض لإعادة استخدام عملية وحدة معالجة الرسومات العامة؛ ثم أضف قواعد الجدولة لهدف MACA في مرحلة الضبط؛ وأخيرًا، أضف الدعم لـ CodeGenMACA وقم بتجميع كود MACAC.

بالإضافة إلى ذلك، تمت إضافة استخدام MACA Device وMACA Runtime API إلى tvm.Device، بما في ذلك عمليات الذاكرة على MACA Device وإطلاق kernel في وقت التشغيل.

الآن تدعم العديد من منتجاتنا الأجهزة الفرعية على مستوى TVM، ويقوم البرنامج الخلفي بإجراء عمليات تحسين لمنتجات مختلفة استنادًا إلى الأجهزة الفرعية.في المنتجات المختلفة، سيقوم المترجم تلقائيًا باختيار حل التكيف المقابل بناءً على الاختلاف في نوع الجهاز؛في الوقت نفسه، في سيناريوهات التجميع الدفعي، نحاول إنشاء تكوينات اختيار ثابتة للهندسة المعمارية المختلفة.أثناء مرحلة التجميع العامة، سيقوم المترجم بضبط قواعد التجميع المتعلقة بالوظيفة بشكل ديناميكي للهندسة المعمارية المختلفة استنادًا إلى التكوين المحدد.

من حيث تكيف المشغل، قمنا بإجراء تكوين على المستوى الأعلى، بما في ذلك بشكل أساسي:

* في الجدول السفلي، أضف الجدول إلى PrimFunc من خلال بدائية جدولة TIR

* في Split Mixed Module، أضف الهدف والمعلومات الأخرى، واحقن maca المضمنة، وأضف تعليمات المزامنة لـ MACA

* عند CodeGenMACA، قم بتضمين ملفات الرأس التي يعتمد عليها MACA، وقم بإنشاء استخدام واجهة برمجة تطبيقات maca wmma من التعليمات ذات الصلة بـ tir.mma، وأعلن عن متغيرات من أنواع مختلفة واستخدمها.

أثناء عملية التكيف، ومن أجل تحقيق أداء أفضل، قمنا بإجراء معالجة خاصة بناءً على خصائصها:

* غير قادر على تمكين tensorize أثناء الضبط: مجموعة معلمات مشغل conv2d ليست 1، ويتم استخدام تنفيذ مكتبة مشغل MACA بشكل مباشر في TOPI؛
* مشغلات مخصصة بعد استيراد نموذج onnx: مشغل Multi Head AttentionV1 مُحسَّن للغاية في بيئة تشغيل MACA onnx. يتم تضمين استدعاء المشغل في contrib، مما يسمح لـ TVM باستخدام تنفيذ المشغل عالي الأداء المُحسَّن يدويًا مباشرةً بعد استيراد نموذج onnx.

بالنسبة لنا،من المرجح أن يتم تنفيذ مشغلي التحسين المخصصين للبائعين في Python وMAC A:

* في واجهة Python، يتم تعريف Relax.Function من خلال الجمع بين المشغلات الأساسية؛ يستخدم tir.PrimFunc tir لتحديد تنفيذ المشغل وإضافة الجداول حسب الحاجة

* في واجهة MACA C، يقوم tir.call_packed بتغليف تنفيذ مكتبة المشغلات عالية الأداء ويستخدمها؛ يستخدم tir.call_kernel كود kernel المطبق في MACA C ويجمعه في استدعاء PackedFunc من خلال مكدس TVM.

بالإضافة إلى ذلك، من أجل إعطاء اللعب الكامل لخصائص الأجهزة الخاصة بوحدة معالجة الرسومات Muxi،لقد قام الفريق بتحسين خوارزمية جدولة TVM بشكل عميق:

* تمت إضافة دعم لنوع WMMA float32 لهدف MACA:

أولاً، دعم واجهة برمجة تطبيقات wmma من نوع float32 في MACA، وإضافة قاعدة tensorize التلقائية من نوع float32 في ScheduleRules في MACA، بحيث يمكن لـ TVM التعرف تلقائيًا على WMMA للأجهزة واستخدامها، وإضافة تحسين tensorize من نوع float32 المقابل في إطار عمل تحسين dlight لتحسين كفاءة تشغيل المصفوفة.

* تقييم تأثير النسخ غير المتزامن على خوارزمية الجدولة:

تحسين حسابات wmma المتعددة من تحميل مجموعة واحدة وحساب مجموعة واحدة إلى تحميل المجموعة التالية من البيانات بشكل غير متزامن وحساب المجموعة الحالية من البيانات بشكل متزامن، وتحسين كفاءة خط الأنابيب، وتمكين منطق تحسين خط الأنابيب البرمجي في MACA ScheduleRules، وإضافة حقن تعليمات النسخ غير المتزامنة وتوليد التعليمات البرمجية لهدف MACA

لقد قمنا أيضًا ببعض المحاولات لدعم أنواع البيانات الجديدة:تمكين التكيف مع هدف MACA في نظام DataType؛ ودعم منطق التوزيع التلقائي لنوع Float8 في MACA ScheduleRules، وتوسيع دعم TVM لأنواع البيانات المخصصة مثل Float8؛ وتنفيذ الدعم لتحويل نوع Float8 وتوليد كود التشغيل في CodeGenMACA، واستكمال تعريفات التشغيل ذات الصلة في maca_half_t.h.

تطبيق TVM على وحدة معالجة الرسومات Muxi

من حيث تصميم الإطار، نفذ الفريق طريقتين للوصول:الطريقة الأولى هي استيراد نموذج الشعلة مباشرة وتنفيذه في واجهة Relay الأمامية، والطريقة الثانية هي استخدام torch.compile لاستخدام TVM كواجهة خلفية.ويحقق اتصالاً فعالاً بين الإطار العلوي والأجهزة الأساسية.

في مرحلة تقييم الأداء، قمت باختيار ReseNet50 وBert كنماذج مرجعية، وقارنت أداء تجميع وتنفيذ Torch وTVM دون تحسين عميق.تظهر البيانات التجريبية أن TVM تتمتع بمزايا كبيرة في بعض الجوانب، وأن أداؤها يتفوق على الشعلة في بعض السيناريوهات.ويرجع ذلك إلى مرونة التمثيل الوسيط لـ TVM (IR) وتحسينه المستهدف لخصائص الأجهزة.

TileLang هي لغة خاصة بالمجال (DSL) في نظام TVM البيئي، تركز على التحسين المتطور للحوسبة الموترية.لقد قام فريقنا بإجراء تكيف وظيفي متعمق في الجوانب التالية:* دعم استخدام هدف MACA * إضافة CodeGenTileLangMACA لتوليد كود Kernel MACA C * استبدال mcTVM واستخدامه كتبعية * إضافة منطق معالجة هدف MACA في libgen والمحول والغلاف * إضافة تعريف هدف MACA باستخدام gemm في tl_template

من حيث التحسين، يدور العمل بشكل أساسي حول تنفيذ gemm في tl_template وتنفيذ الخوارزميات التي تتكيف مع خصائص Muxi GPU.

من حيث التفاعل بين البائعين والمجتمع، يجب أن يحقق تصميم وتطوير TileLang التوازن بين ثلاثة مبادئ رئيسية:

من حيث تصميم اللغة،أول شيء يجب حله هو تحقيق التوازن بين التجريد والأداء العالي.على وجه التحديد، عند التعامل مع وحدات حوسبة متعددة، ينبغي أن يوفر المُجمِّع آلية مرنة لاختيار الاستراتيجية، مما يسمح للمطورين ذوي الفهم العميق لخصائص الأجهزة الأساسية بتحديد مسارات محددة لتوليد الشيفرة البرمجية، مع دعم المطورين العاديين في الوقت نفسه لتحقيق برمجة فعّالة من خلال واجهات مُجرّدة دون الحاجة إلى الاهتمام بالتفاصيل الأساسية. هذا التوازن هو العنصر الأساسي الذي يجب على TileLang، كلغة تعريف محددة، التركيز عليه خلال مرحلة التصميم.

ثانيًا،يجب أن يؤخذ في الاعتبار التكوين المخصص للبائع وتوحيد معايير DSL.يجب أن يدعم المُجمِّع خيارات التحسين متعددة المستويات، على سبيل المثال، مما يسمح للمستخدمين المتقدمين بتعديل استراتيجيات التجميع الأساسية، مثل مستوى فك الحلقة والضغط الموروث من خلال المعلمات، وذلك للحصول على نتائج أفضل في توليد الشيفرة. في الوقت نفسه، بالنسبة لبنيات الأجهزة المختلفة، يجب على المُجمِّع توفير نصائح تحسين مُستهدفة وآليات شرح توضيحية لمساعدة المطورين على اختيار مسار التجميع الأمثل وفقًا لخصائص الأجهزة وتحسين كفاءة التطوير.

ثالث،يجب ضمان استمرارية الواجهة بين أجيال المنتجات.أثناء عملية تكرار المُجمِّعات وسلاسل أدوات اللغة، يجب ضمان التوافق مع الإصدارات السابقة لتصميم الواجهة - حيث يُمكن لمنطق التجميع والشفرة المُولَّدة للإصدار الحالي العمل بكفاءة في الجيل التالي من منتجات الأجهزة، وذلك لتجنب تكاليف إعادة بناء الشفرة الناتجة عن تكرار البنية. تُشكِّل هذه الاستمرارية أساس البناء البيئي، الذي يُمكنه تحقيق تراكم "إضافي" لوظائف سلسلة أدوات المُجمِّع بدلاً من الخسارة "الطرحية" الناتجة عن عدم التوافق. وفي الوقت نفسه، يُمكنه تقليل تكاليف تعلم المستخدم ونقله، وتجنب ارتباك الاستخدام.

التحديات والفرص

وأخيرًا، أود أن أتحدث عن التحديات والفرص التي تواجه التطور الحالي للصناعة.وتتمثل التحديات بشكل رئيسي في الجوانب التالية:

الأول هو أن الأطر والخوارزميات القائمة على التطبيقات تتغير بسرعة.مع التطور السريع لمجالات مثل التعلم العميق، تستمر دورة تحديث الأطر والخوارزميات عالية المستوى في التقصير، وقد أدى زيادة الوظائف والأداء إلى الضغط على تكيف المترجم - يحتاج مجتمع المترجم إلى إنشاء آلية تكيف فعالة للاستجابة بسرعة لاحتياجات الدعم للمشغلين الجدد ونماذج الحوسبة الجديدة.

ثانياً، يستمر تطور بنية الأجهزة.حاليًا، يدعم المُجمِّع ميزات بعض بنى وحدات معالجة الرسومات. في حال ظهور ميزات جديدة لبنى الأجهزة مستقبلًا، يجب أن يكون المُجمِّع قادرًا أيضًا على استيعاب ميزات الأجهزة غير المتجانسة.

ثالثًا، أن نماذج البرمجة مستمرة في التطور.من لغة C/C++ التقليدية إلى البرمجة الوظيفية الناشئة ونماذج البرمجة غير المتجانسة، فإن كيفية تحديد السلسلة البيئية المرتبطة بـ Python يمثل تحديًا كبيرًا.

وأخيرًا، يتعلق الأمر بالتوازن بين الدقة والأداء واستهلاك الطاقة.في التطبيقات العملية، لا يقتصر اهتمام المُجمِّعين على أداء الكود فحسب، بل يشمل أيضًا استهلاك طاقة الأجهزة، وهو أمر بالغ الأهمية. ترتبط هذه العوامل باختيار التعليمات اللاحقة وتصميم البنية.

مستقبل،نود أن نقوم ببعض البناء المشترك مع المجتمع:

فيما يتعلق باستراتيجية المصدر المفتوح، نخطط لإطلاق المكونات الأساسية لإطار العمل ومكتبة المشغل، بما في ذلك وحدات الحوسبة الرئيسية مثل FlashMLA. ومن خلال نموذج المصدر المفتوح، سنعزز التحسين التكراري لسلسلة أدوات المترجم، ونُنشئ تأثيرًا بيئيًا واسع النطاق.

ثانيًا، نأمل أن تتاح فرص تعاون أوثق في تطبيقات الصناعة وأطر عملها ومكتبات المشغلين والمُجمِّعات وهياكل الأجهزة. ومن خلال التبادلات التقنية المنتظمة (مثل منتديات الصناعة)، سنركز على القضايا الجوهرية مثل صعوبات تحسين التجميع واستراتيجيات جدولة المشغلين، ونعزز التعاون بين مختلف المجالات، ونستكشف الابتكارات التكنولوجية.

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