错误代码:E1001

类别: 编译时:作用域 Vmem OOM

此错误表示程序需要的作用域向量内存 (Vmem) 超出了分配量。

示例出错提示

RESOURCE_EXHAUSTED: Ran out of memory in memory space vmem while allocating on stack for %my-custom-kernel = bf16[2048,4096]{1,0:T(8,128)(2,1)} custom-call(...) ...

XLA 后端: TPU

概览

TPU 具有向量内存 (VMEM),这是一种仅供 TensorCore (TC) 使用的本地暂存内存。编译器会针对不同类型的分配管理 Vmem:

  • 指令作用域分配: 在执行单个 HLO 指令时,Vmem 中的临时存储空间。这包括操作数跨度缓冲区(例如,用于双缓冲)和寄存器溢出。
  • 程序作用域分配: 超出单个 HLO 指令范围的分配。这些通常是 HLO 临时变量和中间结果,它们是 HLO 指令的输入和/或输出。

当指令作用域分配超出该指令的分配限制时,就会发生编译时作用域 Vmem OOM。此限制由以下因素控制

这些错误通常是由内部编译器 bug 或自定义内核超出其分配限制引起的。

调试

仔细分析错误消息,以确定错误是源于自定义内核还是标准 HLO。由自定义内核引起的错误应具有以下签名:

Ran out of memory in memory space vmem while allocating on stack for %my-custom-call = <output-shape> custom-call(<params>), custom_call_target="tpu_custom_call" ...
  • 自定义内核作用域 Vmem OOM:如果错误指向自定义内核 → 跳转到 重新调整内核
  • 非内核 Vmem 问题:如果 Vmem Omem 是由 非自定义内核操作引起的,则可能是内部编译器 bug。 请针对 XLA 提交 bug 报告,并附上 HLO 转储。

重新调整内核

如果错误源于自定义内核,请使用以下技术来降低内核的内存要求: