8

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

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

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

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

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

4

3 回答 3

4

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

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

于 2012-03-01T12:16:47.640 回答
0

我认为您正在寻找 __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+ ... ]
}
于 2012-02-17T18:41:18.013 回答
0

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

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

于 2012-02-20T14:50:24.780 回答