รหัสข้อผิดพลาด: 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 หน่วยความจำ Scratchpad ที่ใช้ร่วมกัน ใช้เพื่อจัดเตรียมข้อมูล (อินพุต เอาต์พุต ตัวกลาง) ตามโอกาสเพื่อซ่อนเวลาในการตอบสนองของ HBM
203 HBM หน่วยความจำแบนด์วิดท์สูง หน่วยความจำที่แชร์ขนาดใหญ่ที่ใช้สำหรับตารางการฝัง ฮีป และสแต็ก
204 Sync Flags Primitive การซิงโครไนซ์ที่ใช้ในการประสานงาน

ดูข้อมูลเจาะลึกเกี่ยวกับ 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 ที่เฉพาะเจาะจงซึ่งทำให้เกิดความล้มเหลว ให้สร้างการแสดงผลของคอมไพเลอร์ระดับกลาง ดังนี้
    • ส่งออก MLIR ของ SparseCore: ตั้งค่าสถานะ --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. รายงานข้อบกพร่อง: หากขั้นตอนข้างต้นไม่ช่วยแก้ปัญหา แสดงว่าอาจเป็นข้อบกพร่องของคอมไพเลอร์ โปรดส่งรายงานข้อบกพร่อง