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