如何在OpenCL中使用本地内存？

我最近一直在玩OpenCL，并且能够编写只使用全局内存的简单内核。 现在我想开始使用本地内存，但我似乎无法弄清楚如何使用get_local_size()和get_local_id()来一次计算一个“块”的输出。

例如，假设我想将Apple的OpenCL Hello World示例内核转换为使用本地内存的东西。 你会怎么做？ 这是原始的内核源码：

 __kernel square( __global float *input, __global float *output, const unsigned int count) { int i = get_global_id(0); if (i < count) output[i] = input[i] * input[i]; } 

如果这个例子不能很容易地转换成显示如何使用本地内存的东西，那么任何其他简单的例子都可以。

Answer 1

查看NVIDIA或AMD SDK中的示例，他们应该指出你正确的方向。 matrix转置将使用本地内存为例。

使用你的平方内核，你可以将数据放在中间缓冲区中。 记得传入附加参数。

 __kernel square( __global float *input, __global float *output, __local float *temp, const unsigned int count) { int gtid = get_global_id(0); int ltid = get_local_id(0); if (gtid < count) { temp[ltid] = input[gtid]; // if the threads were reading data from other threads, then we would // want a barrier here to ensure the write completes before the read output[gtid] = temp[ltid] * temp[ltid]; } } 

Answer 2

如果本地内存的大小不变，还有另一种可能性。 如果不在kernels参数列表中使用指针，只需声明__local即可在内核中声明本地缓冲区：

 __local float localBuffer[1024]; 

这会由于clSetKernelArg调用较less而删除代码。

Answer 3

在OpenCL本地内存中是为了在工作组中的所有工作项目之间共享数据。 通常需要在使用本地内存数据之前进行屏障调用（例如，一个工作项需要读取由其他工作项写入的本地内存数据）。 屏障在硬件上是昂贵的。 请记住，本地内存应该用于重复的数据读/写。 应尽可能避免银行冲突。

如果您对本地内存不够小心，则可能会比使用全局内存在性能上差一些。