کد خطا: 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 اسم حافظه اسکالر محلی. برای رجیسترهای اسکالر و جریان کنترل استفاده می‌شود.
۲۰۱ کاشی‌اسپم حافظه‌ی مخصوص هر کاشی. SRAM محلی و سریعی که در اختیار یک کاشی SC خاص قرار دارد.
۲۰۲ اسپم حافظه اشتراکی Scratchpad. برای طبقه‌بندی فرصت‌طلبانه داده‌ها (ورودی‌ها، خروجی‌ها، واسطه‌ها) جهت پنهان کردن تأخیر HBM استفاده می‌شود.
۲۰۳ اچ بی ام حافظه با پهنای باند بالا. حافظه اشتراکی بزرگ که برای جاسازی جداول، هیپ‌ها و استک‌ها استفاده می‌شود.
۲۰۴ پرچم‌های همگام‌سازی عناصر اولیه همگام‌سازی که برای هماهنگی استفاده می‌شوند.

برای بررسی عمیق‌تر SC و سلسله مراتب حافظه آن، به مستندات SparseCore مراجعه کنید.

اشکال‌زدایی

وضوح بستگی به این دارد که کدام فضای حافظه تخصیص داده نشده است.

سناریوی ۱. شکست در تخصیص HBM

شناسه فضای حافظه: 203

این خطا زمانی رخ می‌دهد که یک تخصیص موقت واحد درخواست شده توسط برنامه SparseCore برای قرار گرفتن در HBM موجود، بسیار بزرگ باشد. در بارهای کاری تعبیه استاندارد و مجموعه‌های تخلیه شده SC، پارتیشن‌های بسیار بزرگ به ازای هر هسته یا مشخصات نادرست sharding می‌توانند کامپایلر را مجبور به درخواست بافرهای عظیم کنند.

اقدامات توصیه شده:

  • بررسی تقسیم‌بندی: مطمئن شوید که جداول جاسازی و تانسورهای ورودی/خروجی SC شما به درستی تقسیم‌بندی/تقسیم‌بندی شده‌اند. اگر یک هسته مسئول داده‌های زیادی باشد، تخصیص ممکن است با شکست مواجه شود.
  • تنظیم محدودیت‌ها: max_ids_per_partition و max_unique_ids_per_partition را بررسی کنید. اگر این موارد بی‌جهت بالا تنظیم شده باشند، کامپایلر حافظه بیشتری از حد مورد نیاز را رزرو می‌کند. به نحوه تبدیل محدودیت‌ها به جداول مراجعه کنید.

سناریو ۲. خرابی حافظه داخلی

شناسه‌های فضای حافظه: 0، 201، 202، 204

خرابی‌های تخصیص در Smem ، TileSmem ، 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 برای بررسی تمام برنامه‌های بهینه‌ساز سطح پایین (LLO) منتشر شده توسط Mosaic استفاده کنید.
  2. کاهش اندازه خراش (کاربران پالاس): اگر خطا در هسته 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. ثبت اشکال: اگر مراحل بالا مشکل را حل نکرد، احتمالاً اشکال از کامپایلر است. لطفاً یک گزارش اشکال ثبت کنید.