假设我想在一个大型的固定对象上执行并行计算,例如一个固定的大型稀疏(有向)图,或任何类似类型的对象。
为了对这个图或对象进行任何合理的计算,比如在图中随机漫游,将图放在全局内存中大概是不可能的,因为速度原因。
这就留下了本地/私有内存。如果我对GPU架构的理解是正确的,那么(只读)访问本地内存或私有内存的速度几乎没有差别,对吗?我不愿意将图形复制到私有内存中,因为这意味着每个单独的工作单元都必须存储整个图形,这可能会非常快地消耗GPU的内存(对于非常大的图形,甚至会减少可使用的核心数量和/或使操作系统不稳定)。
那么,假设我在本地与私有的读取速度上是正确的,那么我如何在实践中做到这一点呢?例如,如果为了简化,我将图简化为int[] from和int[] to (存储每个有向边的开始和结束),我当然可以使内核看起来像这样
computeMe(__local const int *to, __local const int *from, __global int *result) {
//...
}但是我不知道如何从JOCL调用它,因为那里没有给出private/local/global修饰符。
本地变量是否会自动写入每个本地工作组的内存?或者这是如何工作的?我根本不清楚我应该如何正确地分配内存。
发布于 2012-01-09 11:26:28
您不能从主机传递本地内存参数的值。主机不能读/写本地内存。要使用本地内存,您仍然需要将数据作为全局传入,然后在使用它之前从全局复制到本地。只有当您多次读取数据时,这才是有益的。
常量内存怎么样?如果您的输入数据没有变化,并且不是太大,那么将您的输入数据放入常量内存中可能会给您带来相当大的加速。可用的常量内存通常在16K到64K之间。
computeMe(__constant int *to, __constant int *from, __global int *result) {
//...
}编辑(添加引用):
有关在OpenCL中使用__local内存的示例,请参见here。
有关NVidia硬件的更多性能详细信息,请访问NVidia OpenCL best practices guide (PDF)。在那里,有更多关于内存类型之间性能差异的信息。
发布于 2013-05-12 01:58:45
你写道:“出于速度的原因,将图表放入全局内存大概是不可能的。”-好吧,你没有太多的其他选择。我的意思是数据一般都在全局内存中。
(作为附注-在特定情况下,您可以将其重新转换为纹理(如果元素格式合适)。此外,nvidia上所谓的“常量”内存针对“广播”类型的操作进行了优化,这意味着所有线程都从相同的位置读取,但我猜情况并非如此。我建议在一开始就远离这些类型。)
好吧,作为建议,首先试着简单地使用“全局内存”。本地内存的“生命周期”仅在内核执行期间。只有当您多次重用相同的数据元素(将其视为显式预加载的高速缓存)时,它才是合理的。
另外,本地内存被限制在16-48K字节,所以它只能存储你的数据的一部分。试着把你的图分解成适合这些块的子图。
在您的表示中,可以将边(from[]、to[])划分为固定大小的组。
通用模式是
步骤1.从全局复制到本地
your_local_array get_local_id(0) = input_global_mem get_global_id(0)
步骤2.确保每个线程都做了操作屏障(本地内存栅栏)
现在,工作项(线程)可以在本地内存中加载的子图上工作。
请记住,本地mem将只包含整个图的有限部分。如果您需要从任何线程访问任意节点,则上面的模式将不可用。
我建议开始在不使用本地内存(直接从全局读取)的情况下对算法进行实验,并确保它正常工作(通常会有一些惊喜)。稍后,您可以确定哪些数据部分可以存储在本地内存中,以提高速度。
https://stackoverflow.com/questions/8767159
复制相似问题