Kod błędu: E3000

Kategoria: Czas kompilacji: błąd przydzielania SparseCore

Ten błąd występuje, gdy kompilator XLA:SparseCore nie może przydzielić ciągłego bloku pamięci w określonej przestrzeni pamięci wymaganej przez bieżący program SparseCore.

Przykładowe komunikaty o błędach:

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>

Backendy XLA: TPU

Przegląd

SparseCore (SC) to specjalistyczny procesor do zadań rzadkich. Wykorzystuje ona określone hierarchie pamięci, aby efektywnie zarządzać przenoszeniem danych. Kompilator XLA próbuje statycznie określać rozmiar buforów i przydzielać je na podstawie ograniczeń sprzętowych i kształtów zdefiniowanych przez użytkownika. Ten błąd wskazuje na brak pamięci podczas tej fazy przydzielania.

W komunikacie o błędzie zwykle podany jest identyfikator przestrzeni pamięci. Poniżej znajdziesz typowe przestrzenie pamięci i ich kodowanie w postaci liczb całkowitych:

Identyfikator miejsca w pamięci Nazwa Opis
0 Smem Lokalna pamięć skalarna. Używane w przypadku rejestrów skalarnych i sterowania przepływem.
201 TileSpmem Pamięć podręczna specyficzna dla kafelka. Szybka, lokalna pamięć SRAM dostępna dla konkretnego bloku SC.
202 Spmem Współdzielona pamięć robocza. Służy do tymczasowego przechowywania danych (wejściowych, wyjściowych i pośrednich), aby ukryć opóźnienie HBM.
203 HBM Pamięć o dużej przepustowości. Duża pamięć współdzielona używana na potrzeby tabel z osadzaniem, stert i stosów.
204 Flagi synchronizacji Prymitywy synchronizacji używane do koordynacji.

Szczegółowe informacje o SC i jego hierarchii pamięci znajdziesz w dokumentacji SparseCore.

Debugowanie

Rozwiązanie zależy od tego, w którym obszarze pamięci nie udało się przydzielić miejsca.

Scenariusz 1. Błędy przydzielania HBM

Identyfikator miejsca w pamięci: 203

Ten błąd występuje, gdy pojedyncza tymczasowa alokacja żądana przez program SparseCore jest zbyt duża, aby zmieścić się w dostępnej pamięci HBM. W przypadku standardowych obciążeń związanych z osadzaniem i kolektywów przeniesionych do SC bardzo duże partycje na rdzeń lub nieprawidłowe specyfikacje dzielenia mogą wymusić na kompilatorze żądanie ogromnych buforów.

Zalecane działania:

  • Sprawdź dzielenie na fragmenty: upewnij się, że tabele z osadzaniem i tensory wejściowe/wyjściowe SC są prawidłowo podzielone na fragmenty. Jeśli jeden rdzeń odpowiada za zbyt dużą ilość danych, przydzielenie może się nie udać.
  • Dostosuj limity: sprawdź max_ids_per_partition i max_unique_ids_per_partition. Jeśli te wartości są niepotrzebnie wysokie, kompilator rezerwuje więcej pamięci niż jest to konieczne. Zobacz Jak limity przekładają się na tabele.

Sytuacja 2. Błędy pamięci wewnętrznej

Identyfikatory przestrzeni pamięci: 0, 201, 202, 204

Błędy alokacji w przypadku Smem, TileSpmem, Spmem lub Sync Flags zwykle występują z powodu błędów kompilatora lub ograniczeń strategii alokacji, gdy kompilator nie uwzględnia wszystkich wymagań dotyczących pamięci.

Zalecane działania:

  1. Wyodrębnij operację XLA, która powoduje błąd: aby zidentyfikować konkretny SC HLO lub jądro Mosaic, które powoduje błąd, wygeneruj pośrednie reprezentacje kompilatora:
    • Zrzut SparseCore MLIR: ustaw flagę --xla_sc_dump_mlir_to=/path/to/dump. Spowoduje to wygenerowanie kodu MLIR programu SparseCore, dzięki czemu zobaczysz, który rozmiar przydziału pasuje do komunikatu o błędzie.
    • Dump Mosaic LLO: w przypadku niestandardowych jąder użyj --xla_mosaic_dump_to=/path/to/dump, aby sprawdzić wszystkie programy optymalizatora niskiego poziomu (LLO) wygenerowane przez Mosaic.
  2. Zmniejsz rozmiary obszaru roboczego (użytkownicy Pallas): jeśli błąd występuje w jądrze Mosaic, sprawdź konfigurację scratch_shapes. Sprawdź, czy pltpu.SMEMżądania mieszczą się w specyfikacjach sprzętowych dla Twojej generacji TPU.
  3. Wyłączanie zbiorczego przenoszenia: jeśli błąd wynika z przeniesienia operacji zbiorczych do SC, spróbuj wyłączyć funkcje przenoszenia do SC:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Zgłoś błąd: jeśli powyższe czynności nie rozwiążą problemu, prawdopodobnie jest to błąd kompilatora. Zgłoś błąd.