Kod błędu: E3000

Kategoria: czas kompilacji: nieudana alokacja SparseCore

Ten błąd występuje, gdy kompilator XLA:SparseCore nie może przydzielić ciągłego bloku pamięci w określonym obszarze pamięci wymaganym 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 wyspecjalizowany procesor do zadań rzadkich. Wykorzystuje on określone hierarchie pamięci do efektywnego zarządzania przepływem danych. Kompilator XLA próbuje statycznie określić rozmiar buforów i przydzielić je na podstawie limitów sprzętowych i kształtów zdefiniowanych przez użytkownika. Ten błąd wskazuje na stan braku pamięci podczas tej fazy alokacji.

Komunikat o błędzie zwykle zawiera identyfikator obszaru pamięci. Poniżej znajdziesz typowe obszary pamięci i ich kodowanie liczbami całkowitymi:

Identyfikator obszaru pamięci Nazwa Opis
0 Smem Lokalna pamięć skalarna. Używana do rejestrów skalarnych i przepływu sterowania.
201 TileSpmem Pamięć podręczna specyficzna dla kafelka. Szybka, lokalna pamięć SRAM dostępna dla konkretnego kafelka SC.
202 Spmem Współdzielona pamięć podręczna. Używana do tymczasowego przechowywania danych (wejściowych, wyjściowych, pośrednich) w celu ukrycia opóźnienia HBM.
203 HBM Pamięć o dużej przepustowości. Duża, współdzielona pamięć używana do tabel osadzania, 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ć pamięci.

Sytuacja 1. Błędy alokacji HBM (identyfikator obszaru pamięci: 203)

Ten błąd występuje, gdy pojedyncza tymczasowa alokacja wymagana przez program SparseCore jest zbyt duża, aby zmieścić się w dostępnej pamięci HBM. W standardowych zadaniach osadzania i zbiorach SC przenoszonych do pamięci podręcznej bardzo duże partycje na rdzeń lub nieprawidłowe specyfikacje fragmentowania mogą wymusić na kompilatorze żądanie dużych buforów.

Zalecane działania:

  • Sprawdź fragmentowanie: upewnij się, że tabele osadzania i tensory wejściowe/wyjściowe SC są prawidłowo podzielone na partycje lub fragmenty. Jeśli jeden rdzeń odpowiada za zbyt dużą ilość danych, alokacja może się nie powieść.
  • Dostosuj limity: sprawdź max_ids_per_partition i max_unique_ids_per_partition. Jeśli te wartości są ustawione niepotrzebnie wysoko, kompilator rezerwuje więcej pamięci niż jest to potrzebne. Więcej informacji znajdziesz w artykule Jak limity przekładają się na tabele.

Sytuacja 2. Błędy pamięci wewnętrznej (identyfikatory obszarów pamięci: 0, 201, 202, 204)

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

Zalecane działania:

  1. Izoluj operację XLA, która powoduje błąd: aby zidentyfikować konkretny jądro SC HLO lub 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 MLIR programu SparseCore, co pozwoli Ci zobaczyć, który rozmiar alokacji odpowiada komunikatowi o błędzie.
    • Zrzut Mosaic LLO: w przypadku niestandardowych jąder użyj --xla_mosaic_dump_to=/path/to/dump, aby sprawdzić wszystkie programy Low Level Optimizer (LLO) emitowane przez Mosaic.
  2. Zmniejsz rozmiary pamięci podręcznej (użytkownicy Pallas): jeśli błąd występuje w jądrze Mosaic, sprawdź konfigurację scratch_shapes. Upewnij się, że żądania pltpu.SMEM mieszczą się w specyfikacjach sprzętowych dla konkretnej generacji TPU.
  3. Wyłącz odciążanie zbiorowe: Jeśli błąd wynika z operacji zbiorowych odciążonych przez SC, spróbuj wyłączyć funkcje odciążania 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 kroki nie rozwiążą problemu, prawdopodobnie jest to błąd kompilatora. Wyślij raport o błędzie.