1

假设我想对一个大型固定对象执行并行计算，例如一个固定的大型稀疏（有向）图，或任何类似类型的对象。

要对此图或对象进行任何合理的计算，例如图中的随机游走，出于速度原因，将图放在全局内存中可能是不可能的。

这留下了本地/私人内存。如果我对 GPU 架构的理解正确，那么（只读）访问本地或私有内存之间几乎没有速度差异，对吗？我不愿意将图形复制到私有内存，因为这意味着每个工作单元都必须存储整个图形，这可能会很快耗尽 GPU 的内存（对于非常大的图形，甚至会减少可以使用和/或使操作系统不稳定）。

所以，假设我在本地与私有的读取速度上是正确的，我该如何在实践中做到这一点？如果例如为了简化我将图简化为 anint[] from和 an int[] to（存储每个有向边的开始和结束），我当然可以使内核看起来像这样

computeMe(__local const int *to, __local const int *from, __global int *result) {
     //...
}

但我不知道我应该如何从 JOCL 调用它，因为那里没有给出私有/本地/全局修饰符。

局部变量会自动写入每个本地工作组的内存吗？或者这是如何工作的？我完全不清楚我应该如何正确地进行内存分配。

4

2 回答 2

3

您不能从主机传递本地内存参数的值。主机无法读取/写入本地内存。要使用本地内存，您仍然需要将数据作为全局传递，然后在使用之前从全局复制到本地。这仅在您多次读取数据时才有用。

持续记忆怎么样？如果您的输入数据没有改变并且它不是太大，将您的输入数据放入常量内存可能会给您带来相当大的加速。可用的常量内存通常在 16K 到 64K 左右。

computeMe(__constant int *to, __constant int *from, __global int *result) {
 //...
}

编辑（添加参考）：

有关在 OpenCL 中使用 __local 内存的示例，请参见此处。

对于 NVidia 硬件，更多性能细节是NVidia OpenCL 最佳实践指南 (PDF)。在那里，有更多关于内存类型之间性能差异的信息。

于 2012-01-09T03:26:28.203 回答
1

您写道“出于速度原因，将图形放入全局内存中可能是不可能的。” - 好吧，你没有太多其他选择。我的意思是数据通常在全局内存中。

（作为旁注 - 在特定情况下，您可能会将其重新转换为纹理（如果元素格式合适）。此外，nvidia 上所谓的“常量”内存针对“广播”类型的操作进行了优化，这意味着所有线程都从同一个读取位置，我猜不是这样。我建议一开始就远离这些类型。）

好的，作为建议，首先尝试简单地使用“全局内存”。本地内存的“生命周期”仅在内核执行期间。仅当您多次重复使用相同的数据元素时才合理（将其视为您明确预加载的高速缓存）。

本地内存也被限制在 16-48kbytes 左右，所以它只能存储你的一部分数据。尝试将您的图分解为适合这些块的子图。

在您的表示中，您可以将边（从 []，到 []）划分为固定大小的组。

通用模式是

步骤 1. 从全局复制到本地

your_local_array[get_local_id(0)] = input_global_mem[get_global_id(0)]

步骤 2. 确保每个线程都执行操作屏障（本地内存栅栏）

现在，工作项（线程）可以在本地内存中加载的子图上工作。

请记住，本地内存将仅包含整个图形的有限部分。如果您需要从任何线程访问任意节点，上述模式将不可用。

我建议在不使用本地内存（直接从全局读取）的情况下开始对算法进行实验，并确保它正常工作（通常在路上会有一些惊喜）。稍后您可以确定哪些数据部分可以存储在本地内存中以加快速度。

于 2013-05-11T17:58:44.647 回答