Código do erro: E3000

Categoria:tempo de compilação: falha na alocação do SparseCore

Esse erro ocorre quando o compilador XLA:SparseCore não consegue alocar um bloco contíguo de memória no espaço especificado exigido pelo programa SparseCore atual.

Exemplos de mensagens de erro:

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>

Backends do XLA:TPU

Visão geral

O SparseCore (SC) é um processador especializado para cargas de trabalho esparsas. Ele depende de hierarquias de memória específicas para gerenciar a movimentação de dados com eficiência. O compilador XLA tenta dimensionar e alocar buffers de forma estática com base em limites de hardware e formatos definidos pelo usuário. Esse erro indica uma condição de falta de memória (OOM) durante essa fase de alocação.

A mensagem de erro geralmente especifica um ID de espaço de memória. Confira abaixo os espaços de memória comuns e as codificações de números inteiros deles:

ID do espaço de memória Nome Descrição
0 Smem (em inglês) Memória escalar local. Usado para registros escalares e fluxo de controle.
201 TileSpmem Memória de bloco de notas específica para blocos. SRAM local e rápida disponível para um bloco específico do SC.
202 Spmem Memória de bloco de notas compartilhada. Usado para armazenar dados (entradas, saídas, intermediários) de maneira oportunista para ocultar a latência do HBM.
203 HBM Memória de largura de banda alta. Memória compartilhada grande usada para incorporar tabelas, heaps e stacks.
204 Flags de sincronização Primitivas de sincronização usadas para coordenação.

Para saber mais sobre o SC e a hierarquia de memória dele, consulte a documentação do SparseCore.

Depuração

A resolução depende de qual espaço de memória falhou na alocação.

Cenário 1. Falhas de alocação de HBM (ID do espaço de memória: 203)

Esse erro ocorre se uma única alocação temporária solicitada pelo programa SparseCore for muito grande para caber na HBM disponível. Em cargas de trabalho de incorporação padrão e coletivos descarregados do SC, partições por núcleo extremamente grandes ou especificações de fragmentação incorretas podem forçar o compilador a solicitar buffers enormes.

Ações recomendadas:

  • Verifique o sharding:confira se as tabelas de incorporação e os tensores de entrada/saída do SC estão particionados/fragmentados corretamente. Se um único núcleo for responsável por muitos dados, a alocação poderá falhar.
  • Ajustar limites:analise max_ids_per_partition e max_unique_ids_per_partition. Se esses valores forem definidos de maneira desnecessariamente alta, o compilador vai reservar mais memória do que o necessário. Consulte Como os limites se traduzem em tabelas.

Cenário 2. Falhas na memória interna (IDs de espaço de memória: 0, 201, 202, 204)

Falhas de alocação em Smem, TileSpmem, Spmem ou Flags de sincronização geralmente ocorrem devido a bugs do compilador ou limitações na estratégia de alocação, em que o compilador não considera todos os requisitos de memória.

Ações recomendadas:

  1. Isolar a operação XLA com falha:para identificar o SC HLO ou o kernel do Mosaic específico que está causando a falha, gere as representações intermediárias do compilador:
    • Despejar MLIR do SparseCore:defina a flag --xla_sc_dump_mlir_to=/path/to/dump. Isso gera o MLIR do programa SparseCore, permitindo que você veja qual tamanho de alocação corresponde à mensagem de erro.
    • Despejar LLO do Mosaic:para kernels personalizados, use --xla_mosaic_dump_to=/path/to/dump para inspecionar todos os programas do otimizador de baixo nível (LLO) emitidos pelo Mosaic.
  2. Reduza os tamanhos de scratch (usuários do Pallas): se a falha ocorrer em um kernel do Mosaic, revise sua configuração scratch_shapes. Verifique se as solicitações de pltpu.SMEM se encaixam nas especificações de hardware da geração de TPU específica.
  3. Desative o descarregamento coletivo:se o erro surgir de uma operação coletiva descarregada pelo SC, tente desativar os recursos de descarregamento do SC:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Registre um bug:se as etapas acima não resolverem o problema, provavelmente é um bug do compilador. Envie um relatório do bug.