Codice di errore: E3000

Categoria: Tempo di compilazione: errore di allocazione SparseCore

Questo errore si verifica quando il compilatore XLA:SparseCore non è in grado di allocare un blocco di memoria contiguo nello spazio di memoria specificato richiesto dal programma SparseCore corrente.

Esempi di messaggi di errore:

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>

Backend XLA: TPU

Panoramica

SparseCore (SC) è un processore specializzato per carichi di lavoro sparsi. Si basa su gerarchie di memoria specifiche per gestire in modo efficiente lo spostamento dei dati. Il compilatore XLA tenta di dimensionare e allocare staticamente i buffer in base ai limiti hardware e alle forme definite dall'utente. Questo errore indica una condizione di esaurimento della memoria (OOM) durante questa fase di allocazione.

Il messaggio di errore in genere specifica un ID spazio di memoria. Di seguito sono riportati gli spazi di memoria comuni e le relative codifiche numeriche:

ID spazio di memoria Nome Descrizione
0 Smem Memoria scalare locale. Utilizzato per i registri scalari e il flusso di controllo.
201 TileSpmem Memoria Scratchpad specifica per il riquadro. SRAM locale veloce disponibile per un riquadro SC specifico.
202 Spmem Memoria scratchpad condivisa. Utilizzato per organizzare in modo opportunistico i dati (input, output, intermedi) per nascondere la latenza HBM.
203 HBM Memoria ad alta larghezza di banda. Memoria condivisa di grandi dimensioni utilizzata per incorporare tabelle, heap e stack.
204 Sync Flags Primitive di sincronizzazione utilizzate per il coordinamento.

Per un approfondimento su SparseCore e sulla sua gerarchia di memoria, consulta la documentazione di SparseCore.

Debug

La risoluzione dipende dallo spazio di memoria per cui l'allocazione non è riuscita.

Scenario 1 Errori di allocazione HBM (ID spazio di memoria: 203)

Questo errore si verifica se una singola allocazione temporanea richiesta dal programma SparseCore è troppo grande per essere inserita nella HBM disponibile. Nei carichi di lavoro di incorporamento standard e nei collettivi con SC offload, partizioni per core estremamente grandi o specifiche di sharding errate possono costringere il compilatore a richiedere buffer enormi.

Azioni consigliate:

  • Controlla lo sharding:assicurati che le tabelle di incorporamento e i tensori di input/output SC siano partizionati/suddivisi correttamente. Se un singolo core è responsabile di troppi dati, l'allocazione potrebbe non riuscire.
  • Modifica i limiti: rivedi max_ids_per_partition e max_unique_ids_per_partition. Se questi valori sono impostati su un valore inutilmente elevato, il compilatore riserva più memoria del necessario. Consulta la sezione Come vengono tradotti i limiti nelle tabelle.

Scenario 2. Errori della memoria interna (ID spazio di memoria: 0, 201, 202, 204)

Gli errori di allocazione in Smem, TileSpmem, Spmem o Sync Flags si verificano in genere a causa di bug del compilatore o limitazioni nella strategia di allocazione, in cui il compilatore non tiene conto di tutti i requisiti di memoria.

Azioni consigliate:

  1. Isola l'operazione XLA non riuscita:per identificare l'HLO SC o il kernel Mosaic specifico che causa l'errore, genera le rappresentazioni intermedie del compilatore:
    • Dump SparseCore MLIR: imposta il flag --xla_sc_dump_mlir_to=/path/to/dump. In questo modo viene generato l'MLIR del programma SparseCore, che ti consente di vedere quale dimensione di allocazione corrisponde al messaggio di errore.
    • Dump Mosaic LLO: per i kernel personalizzati, utilizza --xla_mosaic_dump_to=/path/to/dump per esaminare tutti i programmi Low Level Optimizer (LLO) emessi da Mosaic.
  2. Ridurre le dimensioni dei scratch (utenti Pallas): se l'errore si verifica all'interno di un kernel Mosaic, rivedi la configurazione di scratch_shapes. Assicurati che le tue richieste pltpu.SMEM rientrino nelle specifiche hardware per la tua generazione di TPU specifica.
  3. Disattiva l'offload collettivo:se l'errore deriva da operazioni collettive con offload SC, prova a disattivare le funzionalità di offload SC:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Segnala un bug:se i passaggi precedenti non risolvono il problema, è probabile che si tratti di un bug del compilatore. Invia una segnalazione di bug.