رمز الخطأ: 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 الخلفية: وحدة معالجة الموتّرات

نظرة عامة

‫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. تقليل أحجام البيانات المؤقتة (مستخدمو 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. إبلاغ عن خطأ: إذا لم تحلّ الخطوات السابقة المشكلة، من المحتمل أن يكون السبب خطأ في برنامج التجميع. يُرجى تقديم تقرير خطأ.