類別：編譯時間：SparseCore 分配失敗
如果 XLA:SparseCore 編譯器無法在目前 SparseCore 程式所需的指定記憶體空間中，分配連續的記憶體區塊，就會發生這項錯誤。
錯誤訊息範例：
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 後端：TPU
總覽
SparseCore (SC) 是專為稀疏工作負載設計的處理器，這項技術依賴特定記憶體階層，有效管理資料移動。XLA 編譯器會嘗試根據硬體限制和使用者定義的形狀，靜態調整緩衝區大小並分配緩衝區。這個錯誤表示在分配階段發生記憶體不足 (OOM) 狀況。
錯誤訊息通常會指定記憶體空間 ID。以下是常見的記憶體空間及其整數編碼：
| 記憶體空間 ID | 名稱 | 說明 |
|---|---|---|
| 0 | Smem | 本機純量記憶體。用於純量暫存器和控制流程。 |
| 201 | TileSpmem | 專屬的暫存記憶體。特定 SC 磁磚可用的快速本機 SRAM。 |
| 202 | Spmem | 共用暫存記憶體。用於暫存資料 (輸入、輸出、中介) 以隱藏 HBM 延遲。 |
| 203 | HBM | 高頻寬記憶體。用於嵌入資料表、堆積和堆疊的大型共用記憶體。 |
| 204 | 同步處理標記 | 用於協調的同步基本類型。 |
如要深入探索 SC 及其記憶體階層，請參閱 SparseCore 說明文件。
偵錯
解決方式取決於哪個記憶體空間的分配作業失敗。
情境 1：HBM 分配失敗 (記憶體空間 ID：203)
如果 SparseCore 程式要求的單一暫時分配空間過大，無法容納在可用的 HBM 中，就會發生這個錯誤。在標準嵌入工作負載和 SC 卸載集合中，極大的單一核心分割區或不正確的分片規格，可能會迫使編譯器要求大量緩衝區。
建議採取的行動：
- 檢查分片：確認嵌入資料表和 SC 輸入/輸出張量已正確分割/分片。如果單一核心負責的資料量過多，分配作業可能會失敗。
- 調整限制：查看
max_ids_per_partition和max_unique_ids_per_partition。如果這些值設定得過高，編譯器會保留超出需求的記憶體。請參閱「限制如何轉換為表格」。
情境 2：內部記憶體故障 (記憶體空間 ID：0、201、202、204)
Smem、TileSpmem、Spmem 或 Sync Flags 中的分配失敗通常是由於編譯器錯誤或分配策略的限制，編譯器無法考量所有記憶體需求。
建議採取的行動：
- 找出導致失敗的 XLA 作業：如要找出導致失敗的特定 SC HLO 或 Mosaic 核心，請產生中繼編譯器表示法：
- 傾印 SparseCore MLIR：設定
--xla_sc_dump_mlir_to=/path/to/dump旗標。這會產生 SparseCore 程式的 MLIR，方便您查看哪個分配大小與錯誤訊息相符。 - 傾印 Mosaic LLO：如果是自訂核心，請使用
--xla_mosaic_dump_to=/path/to/dump檢查 Mosaic 發出的所有低階最佳化工具 (LLO) 程式。
- 傾印 SparseCore MLIR：設定
- 縮減暫存大小 (Pallas 使用者)：如果失敗發生在 Mosaic 核心內，請檢查
scratch_shapes設定。請確保pltpu.SMEM要求符合特定 TPU 世代的硬體規格。 - 停用集體卸載：如果錯誤是由 SC 卸載的集體作業所引起，請嘗試停用 SC 卸載功能：
--xla_tpu_enable_sparse_core_collective_offload_all_gather=false--xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
- 回報錯誤：如果上述步驟無法解決問題，可能是編譯器錯誤。請提交錯誤報告。