رمز الخطأ: E3000

الفئة: وقت التجميع: تعذُّر تخصيص SparseCore

يحدث هذا الخطأ عندما يتعذّر على برنامج التجميع XLA:SparseCore تخصيص كتلة متجاورة من الذاكرة في مساحة الذاكرة المحدّدة التي يتطلّبها برنامج SparseCore الحالي.

أمثلة على رسائل الخطأ:

INTERNAL:Failed to run pass pipeline. Hlo-Op: result.1:279:1: error: 'memref.alloca' op current allocation offset upper bound (140704 words) exceeds the legitimate user allocatable offset upper bound (131071 words) in memory space 201 when allocating 23440 words. result.1:279:1: note: see current operation: %232 = "memref.alloca"() <{operandSegmentSizes = array<i32: 0, 0>}> : () -> memref<23440xf32, 201>

البرامج الخلفية في XLA: وحدة معالجة الموتّرات (TPU)

نظرة عامة

‫SparseCore (SC) هو معالج بيانات متخصّص لأحمال العمل المتفرقة. وتعتمد على تسلسلات هرمية محددة للذاكرة من أجل إدارة نقل البيانات بكفاءة. يحاول برنامج التجميع XLA تحديد حجم المخازن المؤقتة وتخصيصها بشكل ثابت استنادًا إلى حدود الأجهزة والأشكال التي يحدّدها المستخدم. يشير هذا الخطأ إلى حدوث حالة نفاد الذاكرة (OOM) أثناء مرحلة التخصيص هذه.

تحدّد رسالة الخطأ عادةً معرّف مساحة تخزين. في ما يلي مساحات الذاكرة الشائعة ورموزها العددية:

معرّف مساحة الذاكرة الاسم الوصف
0 Smem ذاكرة الأعداد القياسية المحلية تُستخدَم لتسجيل القيم العددية والتحكّم في التدفق.
201 TileSpmem ذاكرة لوحة الخدش الخاصة بالنافذة ذاكرة وصول عشوائي ثابتة (SRAM) محلية وسريعة متاحة لقطعة نظام معيّن (SC).
202 Spmem Shared Scratchpad Memory. تُستخدَم هذه السمة لتنظيم البيانات (المدخلات والمخرجات والبيانات الوسيطة) بشكل انتهازي لإخفاء وقت استجابة الذاكرة ذات النطاق الترددي العالي.
203 HBM ذاكرة الوصول العشوائي العالية النطاق الترددي ذاكرة مشتركة كبيرة تُستخدم لتضمين الجداول والمكدّسات والمجموعات.
204 علامات المزامنة عناصر المزامنة الأساسية المستخدَمة للتنسيق

للحصول على معلومات تفصيلية حول SC والتسلسل الهرمي للذاكرة، يُرجى الرجوع إلى مستندات SparseCore.

تصحيح الأخطاء

يعتمد الحل على مساحة الذاكرة التي تعذّر تخصيصها.

السيناريو 1 أخطاء تخصيص HBM

معرّف مساحة الذاكرة: 203

يحدث هذا الخطأ إذا كان حجم عملية التخصيص المؤقتة الواحدة التي طلبها برنامج SparseCore أكبر من أن يتسع لذاكرة النطاق الترددي العالي المتاحة. في أحمال العمل المضمّنة العادية والمجموعات التي تم نقلها إلى SC، يمكن أن تؤدي الأقسام الكبيرة جدًا لكل نواة أو مواصفات التقسيم غير الصحيحة إلى إجبار المترجم على طلب مخازن مؤقتة ضخمة.

الإجراءات المقترَحة:

  • التحقّق من التقسيم: تأكَّد من تقسيم/تجزئة جداول التضمين وموترات الإدخال/الإخراج في SC بشكلٍ صحيح. إذا كان هناك مركز واحد مسؤول عن كمية كبيرة جدًا من البيانات، قد يتعذّر التخصيص.
  • تعديل الحدود: راجِع max_ids_per_partition وmax_unique_ids_per_partition. إذا تم ضبط هذه القيم على مستوى مرتفع بدون داعٍ، سيحجز المترجم البرمجي مقدارًا من الذاكرة أكبر من اللازم. راجِع مقالة كيفية تحويل الحدود إلى جداول.

السيناريو 2. أعطال الذاكرة الداخلية

معرّفات مساحة الذاكرة: 0 و201 و202 و204

تحدث أخطاء التخصيص في Smem أو TileSpmem أو Spmem أو Sync Flags عادةً بسبب أخطاء في برنامج التجميع أو قيود في استراتيجية التخصيص، حيث لا يراعي برنامج التجميع جميع متطلبات الذاكرة.

الإجراءات المقترَحة:

  1. عزل عملية XLA التي تعذّر تنفيذها: لتحديد عملية SC HLO أو نواة Mosaic المحدّدة التي تسبّبت في تعذّر التنفيذ، أنشئ تمثيلات وسيطة للمترجم:
    • Dump SparseCore MLIR: اضبط العلامة --xla_sc_dump_mlir_to=/path/to/dump. يؤدي ذلك إلى إنشاء رمز MLIR لبرنامج SparseCore، ما يتيح لك معرفة حجم التخصيص الذي يتطابق مع رسالة الخطأ.
    • Dump Mosaic LLO: بالنسبة إلى النواة المخصّصة، استخدِم --xla_mosaic_dump_to=/path/to/dump لفحص جميع برامج Low Level Optimizer (LLO) التي تم إنشاؤها بواسطة Mosaic.
  2. تقليل أحجام Scratch (مستخدمو Pallas): إذا حدث الخطأ في إحدى وحدات Mosaic الأساسية، راجِع إعدادات scratch_shapes. تأكَّد من أنّ طلباتك pltpu.SMEM تتوافق مع مواصفات الأجهزة الخاصة بجيل TPU الذي تستخدمه.
  3. إيقاف ميزة "التفريغ الجماعي": إذا كان الخطأ ناتجًا عن عمليات جماعية تم تفريغها من SC، جرِّب إيقاف ميزات التفريغ من SC:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. إبلاغ عن خطأ: إذا لم تحلّ الخطوات أعلاه المشكلة، من المحتمل أن يكون هناك خطأ في المحول البرمجي. يُرجى تقديم تقرير خطأ.