קוד שגיאה: 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 Backends: TPU

סקירה כללית

‫SparseCore ‏ (SC) הוא מעבד ייעודי לעומסי עבודה דלילים. היא מסתמכת על היררכיות זיכרון ספציפיות כדי לנהל את תנועת הנתונים בצורה יעילה. הקומפיילר XLA מנסה להקצות באופן סטטי את הגודל של מאגרי נתונים זמניים על סמך מגבלות החומרה וצורות שהוגדרו על ידי המשתמש. השגיאה הזו מציינת מצב של חוסר זיכרון (OOM) במהלך שלב ההקצאה הזה.

בדרך כלל הודעת השגיאה מציינת מזהה של מרחב זיכרון. בהמשך מפורטים מרחבי הזיכרון הנפוצים וקידוד המספרים השלמים שלהם:

מזהה מרחב הזיכרון שם תיאור
0 Smem זיכרון סקלרי מקומי. משמש לרישום סקלרי ולזרימת בקרה.
201 TileSpmem זיכרון של טיוטה שספציפי למשבצת. זיכרון SRAM מהיר ומקומי שזמין למשבצת SC ספציפית.
202 Spmem זיכרון משותף של יומן הרשימות. התכונה הזו משמשת להכנת נתונים (קלט, פלט, נתונים זמניים) כדי להסתיר את זמן האחזור של HBM.
203 HBM זיכרון עם רוחב פס גבוה. זיכרון גדול ומשותף שמשמש להטמעת טבלאות, ערימות ומחסניות.
204 דגלים של סנכרון פרימיטיבים של סנכרון שמשמשים לתיאום.

למידע נוסף על SC והיררכיית הזיכרון שלו, אפשר לעיין במסמכי התיעוד של SparseCore.

ניפוי באגים

הפתרון תלוי בזיכרון שבו נכשלה ההקצאה.

תרחיש 1. כשלים בהקצאת HBM

מזהה מרחב הזיכרון: 203

השגיאה הזו מתרחשת אם הקצאה זמנית יחידה שנדרשת על ידי התוכנית SparseCore גדולה מדי מכדי להיכנס ל-HBM הזמין. במשימות סטנדרטיות של הטמעה ובמשימות של קולקטיבים שהועברו ל-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. דיווח על באג: אם השלבים שלמעלה לא פותרים את הבעיה, סביר להניח שמדובר בבאג בקומפיילר. מומלץ לשלוח דוח על באג.