Fehlercode: E3000

Kategorie:Kompilierungszeit: Fehler bei der Zuweisung von SparseCore

Dieser Fehler tritt auf, wenn der XLA:SparseCore-Compiler keinen zusammenhängenden Speicherblock im angegebenen Speicherbereich zuweisen kann, der für das aktuelle SparseCore-Programm erforderlich ist.

Beispiele für Fehlermeldungen:

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-Back-Ends:TPU

Übersicht

Der SparseCore (SC) ist ein spezieller Prozessor für spärliche Arbeitslasten. Dabei werden bestimmte Speicherhierarchien verwendet, um die Datenbewegung effizient zu verwalten. Der XLA-Compiler versucht, Puffer statisch zu dimensionieren und zuzuweisen, basierend auf Hardwarebeschränkungen und benutzerdefinierten Formen. Dieser Fehler weist auf einen Out of Memory (OOM)-Zustand während dieser Zuweisungsphase hin.

In der Fehlermeldung wird in der Regel eine Speicherplatz-ID angegeben. Im Folgenden finden Sie die gängigen Speicherbereiche und ihre Ganzzahlcodierungen:

Speicherbereichs-ID Name Beschreibung
0 Smem Lokaler Skalarspeicher. Wird für Skalarregister und die Ablaufsteuerung verwendet.
201 TileSpmem Kachelspezifischer Scratchpad-Speicher. Schneller, lokaler SRAM, der für eine bestimmte SC-Kachel verfügbar ist.
202 Spmem Gemeinsam genutzter Scratchpad-Speicher. Wird verwendet, um Daten (Eingaben, Ausgaben, Zwischenergebnisse) opportunistisch zu stagen, um die HBM-Latenz zu verbergen.
203 HBM Speicher mit hoher Bandbreite (High Bandwidth Memory, HBM) Großer, gemeinsam genutzter Speicher für Einbettungstabellen, Heaps und Stacks.
204 Synchronisierungs-Flags Synchronisationsprimitive, die für die Koordination verwendet werden.

Eine ausführliche Beschreibung von SC und der zugehörigen Speicherhierarchie finden Sie in der SparseCore-Dokumentation.

Debugging

Die Lösung hängt davon ab, in welchem Speicherbereich die Zuweisung fehlgeschlagen ist.

Szenario 1 Fehler bei der HBM-Zuweisung (Memory Space ID: 203)

Dieser Fehler tritt auf, wenn eine einzelne temporäre Zuweisung, die vom SparseCore-Programm angefordert wird, zu groß ist, um in den verfügbaren HBM zu passen. Bei Standard-Embedding-Arbeitslasten und SC-offload-Collectives können extrem große Partitionen pro Kern oder falsche Sharding-Spezifikationen den Compiler dazu zwingen, massive Puffer anzufordern.

Empfohlene Maßnahmen:

  • Sharding prüfen:Achten Sie darauf, dass Ihre Einbettungstabellen und SC-Ein-/Ausgabe-Tensoren richtig partitioniert/geshardet sind. Wenn ein einzelner Core für zu viele Daten verantwortlich ist, kann die Zuweisung fehlschlagen.
  • Limits anpassen:Prüfen Sie max_ids_per_partition und max_unique_ids_per_partition. Wenn diese unnötig hoch eingestellt sind, reserviert der Compiler mehr Speicher als erforderlich. Weitere Informationen finden Sie unter So werden Limits in Tabellen umgesetzt.

Beispielszenario 2: Fehler im internen Speicher (Speicherplatz-IDs: 0, 201, 202, 204)

Zuweisungsfehler in Smem, TileSpmem, Spmem oder Sync Flags treten in der Regel aufgrund von Compilerfehlern oder Einschränkungen in der Zuweisungsstrategie auf, bei denen der Compiler nicht alle Speicheranforderungen berücksichtigt.

Empfohlene Maßnahmen:

  1. Fehlerhaften XLA-Vorgang isolieren:Um den spezifischen SC-HLO- oder Mosaic-Kernel zu identifizieren, der den Fehler verursacht, generieren Sie die Zwischencompilerdarstellungen:
    • SparseCore-MLIR-Dump: Legen Sie das Flag --xla_sc_dump_mlir_to=/path/to/dump fest. Dadurch wird das MLIR des SparseCore-Programms generiert. So können Sie sehen, welche Zuweisungsgröße der Fehlermeldung entspricht.
    • Mosaic-LLO-Dump:Bei benutzerdefinierten Kernels können Sie mit --xla_mosaic_dump_to=/path/to/dump alle vom Low Level Optimizer (LLO) von Mosaic ausgegebenen Programme prüfen.
  2. Scratch-Größen reduzieren (Pallas-Nutzer): Wenn der Fehler in einem Mosaic-Kernel auftritt, überprüfen Sie die scratch_shapes-Konfiguration. Achten Sie darauf, dass Ihre pltpu.SMEM-Anfragen den Hardwarespezifikationen für Ihre spezifische TPU-Generation entsprechen.
  3. Kollektives Offload deaktivieren:Wenn der Fehler durch kollektive Vorgänge entsteht, die auf den Smart Chip ausgelagert wurden, versuchen Sie, die SC-Offload-Funktionen zu deaktivieren:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Fehler melden:Wenn das Problem durch die oben genannten Schritte nicht behoben wird, handelt es sich wahrscheinlich um einen Compilerfehler. Bitte reichen Sie einen Fehlerbericht ein.