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 回答