Mã lỗi: E3000

Danh mục: Thời gian biên dịch: Lỗi phân bổ SparseCore

Lỗi này xảy ra khi trình biên dịch XLA:SparseCore không thể phân bổ một khối bộ nhớ liền kề trong không gian bộ nhớ được chỉ định mà chương trình SparseCore hiện tại yêu cầu.

Thông báo lỗi mẫu:

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>

Phần phụ trợ XLA: TPU

Tổng quan

SparseCore (SC) là một bộ xử lý chuyên dụng cho các khối lượng công việc thưa thớt. Bộ xử lý này dựa vào các hệ thống phân cấp bộ nhớ cụ thể để quản lý việc di chuyển dữ liệu một cách hiệu quả. Trình biên dịch XLA cố gắng định cỡ và phân bổ bộ đệm một cách tĩnh dựa trên các giới hạn phần cứng và hình dạng do người dùng xác định. Lỗi này cho biết điều kiện Hết bộ nhớ (OOM) trong giai đoạn phân bổ này.

Thông báo lỗi thường chỉ định một mã không gian bộ nhớ. Dưới đây là các không gian bộ nhớ phổ biến và mã hoá số nguyên của chúng:

Mã không gian bộ nhớ Tên Mô tả
0 Smem Bộ nhớ vô hướng cục bộ. Dùng cho các thanh ghi vô hướng và luồng điều khiển.
201 TileSpmem Bộ nhớ tạm dành riêng cho ô xếp. SRAM cục bộ, nhanh chóng dành cho một ô xếp SC cụ thể.
202 Spmem Bộ nhớ tạm dùng chung. Dùng để dàn xếp dữ liệu (đầu vào, đầu ra, trung gian) một cách cơ hội nhằm ẩn độ trễ HBM.
203 HBM Bộ nhớ băng thông cao. Bộ nhớ dùng chung, lớn dùng cho các bảng nhúng, vùng nhớ heap và ngăn xếp.
204 Cờ đồng bộ hoá Các nguyên hàm đồng bộ hoá dùng để điều phối.

Để nghiên cứu chuyên sâu về SC và hệ thống phân cấp bộ nhớ của nó, hãy tham khảo tài liệu về SparseCore.

Gỡ lỗi

Giải pháp phụ thuộc vào không gian bộ nhớ không phân bổ được.

Trường hợp 1. Lỗi phân bổ HBM (Mã không gian bộ nhớ: 203)

Lỗi này xảy ra nếu một lần phân bổ tạm thời do chương trình SparseCore yêu cầu quá lớn để vừa với HBM có sẵn. Trong các khối lượng công việc nhúng tiêu chuẩn và các tập hợp được giảm tải SC, các phân vùng cực lớn trên mỗi lõi hoặc thông số kỹ thuật phân đoạn không chính xác có thể buộc trình biên dịch yêu cầu các vùng đệm lớn.

Hành động được đề xuất:

  • Kiểm tra phân đoạn: Đảm bảo các bảng nhúng và tensor đầu vào/đầu ra SC được phân vùng/phân đoạn đúng cách. Nếu một lõi chịu trách nhiệm cho quá nhiều dữ liệu, thì việc phân bổ có thể không thành công.
  • Điều chỉnh giới hạn: Xem lại max_ids_per_partitionmax_unique_ids_per_partition. Nếu các giá trị này được đặt quá cao một cách không cần thiết, thì trình biên dịch sẽ dành nhiều bộ nhớ hơn mức cần thiết. Tham khảo bài viết Cách giới hạn chuyển đổi thành bảng.

Trường hợp 2. Lỗi bộ nhớ trong (Mã không gian bộ nhớ: 0, 201, 202, 204)

Lỗi phân bổ trong Smem, TileSpmem, Spmem hoặc Cờ đồng bộ hoá thường xảy ra do lỗi trình biên dịch hoặc giới hạn trong chiến lược phân bổ, trong đó trình biên dịch không tính đến tất cả các yêu cầu về bộ nhớ.

Hành động được đề xuất:

  1. Cô lập thao tác XLA không thành công: Để xác định HLO SC cụ thể hoặc kernel Mosaic gây ra lỗi, hãy tạo các biểu diễn trình biên dịch trung gian:
    • Kết xuất MLIR SparseCore: Đặt cờ --xla_sc_dump_mlir_to=/path/to/dump. Thao tác này sẽ tạo MLIR của chương trình SparseCore, cho phép bạn xem kích thước phân bổ nào khớp với thông báo lỗi.
    • Kết xuất LLO Mosaic: Đối với các kernel tuỳ chỉnh, hãy sử dụng --xla_mosaic_dump_to=/path/to/dump để kiểm tra tất cả các chương trình Trình tối ưu hoá cấp thấp (LLO) do Mosaic phát ra.
  2. Giảm kích thước tạm thời (người dùng Pallas): Nếu lỗi xảy ra trong một kernel Mosaic, hãy xem lại cấu hình scratch_shapes. Đảm bảo rằng các yêu cầu pltpu.SMEM của bạn phù hợp với thông số kỹ thuật phần cứng cho thế hệ TPU cụ thể của bạn.
  3. Tắt tính năng giảm tải tập thể: Nếu lỗi phát sinh từ một thao tác tập thể được giảm tải SC, hãy thử tắt các tính năng giảm tải SC:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. Báo lỗi: Nếu các bước trên không giải quyết được vấn đề, thì có thể là lỗi trình biên dịch. Vui lòng gửi báo cáo lỗi.