如何确定FPGA内核可使用的最大本地内存？解析内存资源差异

1. 本地内存查询值与厂商eSRAM资源差异的原因 #

你查到的16KB是单个工作组（work-group）能使用的本地内存上限，而厂商文档里的47.25Mb是整个FPGA的eSRAM总资源——这俩不是一回事。

FPGA的eSRAM是全局片上资源，会被编译器分配给多个工作组的本地内存、全局内存缓存、内核流水线寄存器、私有内存等多个部分，不会全部留给单个工作组。Intel FPGA的SYCL/OpenCL实现里，单工作组的本地内存默认配额就是16KB，这个限制是为了避免单个工作组占用过多片上资源，导致整个内核无法布局到FPGA上。

2. 片上内存、本地内存与SYCL内核的关系 #

先明确几个核心概念：

• 片上内存：FPGA内部集成的所有高速存储，包括eSRAM、MLAB（小型查找表内存）、寄存器堆，是区别于外部DDR的本地高速存储。
• 本地内存：SYCL/OpenCL层面的概念，属于片上内存的一部分，是同一个工作组内所有工作项共享的高速存储，生命周期和工作组绑定，由编译器从eSRAM/MLAB中分配，用来在工作组内共享数据、暂存中间结果，大幅降低对外部全局内存的访问延迟。
• 全局内存：一般指FPGA外接的DDR内存，容量大但延迟高，是主机和内核之间交换数据的主要载体。

三者的绑定关系：

• 当你在SYCL内核里用local修饰声明数组时，编译器会为每个工作组从eSRAM中划出对应大小的存储块，这部分就是本地内存。
• 片上内存除了供给本地内存，还要用来实现全局内存的缓存（比如你查到的32KB全局缓存）、内核里的流水线寄存器、私有内存（单个工作项独有的存储，通常用寄存器实现）等。

3. SYCL/OpenCL内核中数组最大尺寸的影响因素 #

不同类型的数组，限制因素不一样：

• 本地内存数组（local修饰）：
  • 首先受单工作组的本地内存配额（16KB）限制，单个工作组里的所有本地内存数组总大小不能超过这个值；
  • 其次受整个FPGA的片上内存总资源（47.25Mb）限制，如果所有工作组的本地内存总和超过总片上内存，编译器会报错无法完成布局。
• 全局内存数组：最大尺寸由FPGA外接的DDR容量决定（你查到的34359737344字节≈32GB，就是DDR的总容量）。
• 私有内存数组（无修饰，工作项独有）：最大尺寸受单个工作项可用的寄存器资源限制，因为私有内存通常用寄存器实现。如果私有数组太大，寄存器资源耗尽，编译器会把溢出的部分放到本地内存或全局内存，直接拉低性能。

原SYCL查询代码 #

std::cout << "Local memory size: "
          << device.get_info<sycl::info::device::local_mem_size>()
          << std::endl;

std::cout << "Global memory size: "
          << device.get_info<sycl::info::device::global_mem_size>()
          << std::endl;

std::cout << "Global memory cache size: "
          << device.get_info<sycl::info::device::global_mem_cache_size>()
          << std::endl;

查询输出 #

Local memory size: 16384
Global memory size: 34359737344
Global memory cache size: 32768

内容的提问来源于stack exchange，提问作者Ypro