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

【问题标题】：physical memory on AMD devices: local vs privateAMD 设备上的物理内存：本地与私有
【发布时间】：2012-03-09 01:20:57
【问题描述】：

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

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

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

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

编辑：问题已解决，答案是附录 D 中的here。

【问题讨论】：

我不相信私有内存是每个内核专用的 - 它映射到寄存器文件，它是每个计算单元资源的。每个工作项都从计算单元寄存器文件中分配寄存器，需要多少决定了在任何给定时刻飞行的波前数量。
从著名的随处可见的绘图codeproject.com/KB/showcase/Memory-Spaces/image001.jpg 得出结论，私有内存在物理上与__local 内存不同，不是吗？
是的，它们在物理上是不同的。在大多数现代 AMD 设备中，私有内存映射到计算单元寄存器文件，本地内存映射到计算单元级共享内存。一些早期的 OpenCL 兼容 GPU 没有共享内存，而本地内存只是 SDRAM。也不是每个核心，每个工作项对私有和每个工作组使用多少本地效果，每个计算单元运行的并发波前数量。
你还没有理解，我认为 - 私有内存（就像名字所说的那样）是每个工作项的私有内存。但它是从计算单元寄存器文件中分配给每个工作项的，该文件充当在给定计算单元上运行的所有工作项的公共资源池。而且我很确定 AMD 的编译器对每个工作单元设置了 256 个寄存器的硬性限制，而与 GPU 上寄存器文件的大小无关。
我认为每个寄存器都是一个 32 位的字。但请记住，代码中的所有其他变量也使用寄存器。我想我记得典型的 AMD GPU 每个计算单元有一个 64kb 的寄存器文件，需要由至少 4 或 8 个波前共享，每个波前有 64 个工作项。但我不太使用他们的硬件，所以这可能不正确。查看其 OpenCL SDK 中的当前发行说明。

标签： architecture opencl gpu gpgpu amd-processor

【解决方案1】：

我认为您正在寻找 __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+ ... ]
}

【讨论】：

你不能轮询它，但它存在吗？从著名的随处可见的绘图codeproject.com/KB/showcase/Memory-Spaces/image001.jpg 我假设每个工作单元上都有一组物理上独立的私有寄存器。不？我希望以某种方式比 CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem) 限制做得更好，因为它大约有一半的内核未使用。访问全局内存可能会太慢，即使它只是增加一个计数器。
我在这里找到了有关 cypress、cayman 和 fermi 寄存器大小的更多信息：realworldtech.com/page.cfm?ArticleID=RWT121410213827&p=11 您应该能够将一些大小合适的私有变量调整到该大小。我认为 LDS 仍然是你最好的选择。

【解决方案2】：

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

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

【讨论】：

哦，真的吗？这很奇怪。我以为找到了答案，但它是一个不同的答案。在附录 D 的 AMD OpenCL 编程指南developer.amd.com/sdks/amdappsdk/assets/… 中，有总的寄存器文件大小，所有现代设备都列为 256 KB。现在哪个是正确的？ :S
我相信两者都是正确的。据我了解，在 GCN 架构中，一个 SIMD 单元有一个 64kb 的寄存器文件，每个计算单元有 4 个 SIMD 单元，即。每个计算单元 4 * 64kb = 256kb 的总寄存器文件。

【解决方案3】：

答案由 cmets 中的用户 talonmies 给出，因此我将在此处将其写在新答案中以结束问题。

这些值可以在 AMD APP OpenCL 编程指南http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf 的附录 D 中找到（nVidia 也有类似的文档）。显然，对于 AMD 设备来说，一个寄存器是 128 位 (4x32)，而对于所有现代高端设备来说，有 16384 个寄存器，所以每个计算单元有 256KB，这是非常了不起的。

【讨论】：