Hata kodu: E3000

Kategori: Derleme Zamanı: SparseCore Allocation Failure

Bu hata, XLA:SparseCore derleyicisi geçerli SparseCore programının gerektirdiği belirtilen bellek alanında bitişik bir bellek bloğu ayıramadığında oluşur.

Örnek hata mesajları:

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 arka uçları: TPU

Genel Bakış

SparseCore (SC), seyrek iş yükleri için özel bir işlemcidir. Veri hareketini verimli bir şekilde yönetmek için belirli bellek hiyerarşilerine dayanır. XLA derleyicisi, donanım sınırlarına ve kullanıcı tanımlı şekillere göre arabellekleri statik olarak boyutlandırmaya ve ayırmaya çalışır. Bu hata, bu ayırma aşamasında bellek yetersizliği (OOM) durumu olduğunu gösterir.

Hata mesajında genellikle bir bellek alanı kimliği belirtilir. Yaygın bellek alanları ve bunların tam sayı kodlamaları aşağıda verilmiştir:

Bellek alanı kimliği Ad Açıklama
0 Smem Yerel skaler bellek. Skaler kayıtlar ve kontrol akışı için kullanılır.
201 TileSpmem Döşemeye özel not defteri belleği. Belirli bir SC kutusunda kullanılabilen hızlı ve yerel SRAM.
202 Spmem Paylaşılan Scratchpad Belleği. HBM gecikmesini gizlemek için verileri (girişler, çıkışlar, ara veriler) fırsatçı bir şekilde hazırlamak için kullanılır.
203 HBM Yüksek bant genişlikli bellek. Tabloları, yığınları ve yığınları yerleştirmek için kullanılan büyük, paylaşılan bellek.
204 Senkronizasyon İşaretleri Koordinasyon için kullanılan senkronizasyon temel öğeleri.

SC ve bellek hiyerarşisi hakkında ayrıntılı inceleme için SparseCore belgelerine bakın.

Hata ayıklama

Çözüm, hangi bellek alanının ayırmada başarısız olduğuna bağlıdır.

1. senaryo HBM ayırma hataları (Bellek Alanı Kimliği: 203)

Bu hata, SparseCore programı tarafından istenen tek bir geçici ayırma işlemi, kullanılabilir HBM'ye sığamayacak kadar büyükse oluşur. Standart yerleştirme iş yüklerinde ve SC'nin boşaltıldığı kolektiflerde, çekirdek başına son derece büyük bölümler veya yanlış parçalama özellikleri, derleyiciyi büyük arabellekler istemeye zorlayabilir.

Önerilen işlemler:

  • Parçalama işlemini kontrol edin: Yerleştirme tablolarınızın ve SC giriş/çıkış tensörlerinizin doğru şekilde bölümlendirildiğinden/parçalandığından emin olun. Tek bir çekirdek çok fazla veriden sorumluyorsa tahsis başarısız olabilir.
  • Sınırları ayarlama: max_ids_per_partition ve max_unique_ids_per_partition'ü inceleyin. Bu değerler gereksiz yere yüksek ayarlanırsa derleyici, gerekenden daha fazla bellek ayırır. Sınırlar tablolara nasıl yansıtılır? başlıklı makaleyi inceleyin.

2. Senaryo. Dahili bellek hataları (Bellek alanı kimlikleri: 0, 201, 202, 204)

Smem, TileSpmem, Spmem veya Sync Flags'deki tahsis hataları genellikle derleyici hatalarından veya tahsis stratejisindeki sınırlamalardan kaynaklanır. Bu durumda derleyici, tüm bellek gereksinimlerini hesaba katamaz.

Önerilen işlemler:

  1. Başarısız olan XLA işlemini yalıtın: Başarısızlığa neden olan belirli SC HLO veya Mosaic çekirdeğini belirlemek için ara derleyici temsillerini oluşturun:
    • Dump SparseCore MLIR: İşareti ayarlayın --xla_sc_dump_mlir_to=/path/to/dump. Bu, SparseCore programının MLIR'sini oluşturur. Böylece, hangi ayırma boyutunun hata mesajıyla eşleştiğini görebilirsiniz.
    • Dump Mosaic LLO: Özel çekirdekler için, Mosaic tarafından yayınlanan tüm Low Level Optimizer (LLO) programlarını incelemek üzere --xla_mosaic_dump_to=/path/to/dump kullanın.
  2. Çizik boyutlarını küçültün (Pallas kullanıcıları): Hata bir Mosaic çekirdeğinde meydana gelirse scratch_shapes yapılandırmanızı inceleyin. pltpu.SMEM İsteklerinizin, TPU'nuzun donanım spesifikasyonlarına uygun olduğundan emin olun.
  3. Toplu boşaltmayı devre dışı bırakın: Hata, boşaltılmış bir SC toplu işlemlerinden kaynaklanıyorsa SC boşaltma özelliklerini devre dışı bırakmayı deneyin:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Hata bildirin: Yukarıdaki adımlar sorunu çözmezse büyük olasılıkla derleyici hatası vardır. Lütfen hata raporu gönderin.