AMD 设备上的物理内存：本地内存与私有内存

Question

我正在 OpenCL 中编写一个算法，其中我需要每个工作单元记住相当一部分数据，例如 long[70] 和 long[200]< 之间的数据/code> 每个内核左右。

最近的 AMD 设备拥有 32 KiB __local 内存，（对于每个内核给定的数据量）足以存储 20-58 个工作单元的信息。然而，根据我对架构的理解（尤其是这张图)，每个着色器核心还拥有专用的私有内存。然而我找不到它的大小。

谁能告诉我如何找出每个内核有多少私有内存？

我对 HD7970 特别好奇，因为我计划很快购买其中一些。

编辑：问题已解决，答案在此处附录D中。

Answer 1

答案是由用户talonmies在评论中给出的，所以我将在这里写一个新答案来结束问题。

这些值可以在 AMD APP OpenCL 编程指南 http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_pa​​rallel_processing_opencl_programming_guide.pdf（nVidia 也有类似的文档）。显然，AMD 设备的寄存器为 128 位 (4x32)，所有现代高端设备有 16384 个寄存器，因此每个计算单元有 256KB 的内存。

Answer 2

我认为你正在寻找 __local 内存。这就是 32KB 本地数据存储所指的。我认为您无法轮询设备来获取私有内存量。

您可以传入 NULL long* cl_mem 引用来分配内存。我认为最好为每个 WI 使用静态内存量。假设每个工作项都需要 long[200]，您将使用下面的代码。将工作划分为具有相同（或相似）内存要求的组也是一个好主意，以便充分利用 LDS 内存。

void __kernel(__local long* localMem, const int localMemPerItem
       //more args...
       )
{
  //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem
  //this work item has access to all of it, but can choose to restrict
  //itself to only the portion it needs.
  //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem)
  int startIndex=localMemPerItem*get_local_id(0);
  //use localMem[startIndex+ ... ]
}

Answer 3

要回答 79xx 系列卡中的寄存器文件有多大，因为它基于 GCN 架构，根据 anandtech 中的图像，它是 64KB： http://www.anandtech.com/print/5261

要回答您的问题如何找出每个内核使用了多少内存..您可以在您的内核上运行 AMD APP Profiler，它在内核占用部分告诉您内核使用了多少空间。