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

Posted 2023-03-16

【中文标题】如何在 OpenCL 中使用本地内存？

【英文标题】：How do I use local memory in OpenCL?

【发布时间】：2011-02-02 06:30:46

【问题描述】：

我最近一直在玩 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];

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

【参考方案1】：

查看 NVIDIA 或 AMD SDK 中的示例，它们应该会为您指明正确的方向。例如，矩阵转置会使用本地内存。

使用平方内核，您可以将数据暂存到中间缓冲区中。记得传入附加参数。

__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];
    

【讨论】：

我已经阅读了 NVIDIA 介绍材料，但我仍然觉得这些示例过于复杂。我正在寻找一个超级简单的一维示例，使用本地内存来弄湿我的脚。

感谢您在上次编辑中添加代码！不过，我似乎无法让您的内核正常工作.... 我将如何使用 clSetKernelArg() 作为临时对象？我需要对临时使用 clCreateBuffer() 吗？此外，您的内核中有一些拼写错误：“temp * temp”应该是“temp[ltid] * temp[ltid]”，并且应该在最后一行之前插入右大括号。

我更正了拼写错误 - 适合在 ipod 上打字。但是，您的 clSetKernelArg 没有分配足够的内存，每个线程需要一个 cl_float 空间（您只分配了一个浮点数）。尝试：clSetKernelArg(kernel, 2, sizeof(cl_float) * local_work_size[0], NULL); 其中local_work_size[0] 是维度 0 中的工作组大小。

请注意，您可以使用限定符 __local 将变量声明为本地变量。例如，您可以执行__local float values[GROUP_SIZE];，然后让每个线程写入values[get_local_id(0)] = ...。不需要通过传递给内核的指针来访问本地内存。

记住本地存储只对同一个工作组可见。因此，您的工作组大小需要大于 1 才能跨多个线程共享。

【参考方案2】：

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

__local float localBuffer[1024];

由于 clSetKernelArg 调用较少，这会删除代码。

【讨论】：

这是真的，但如果您不必在运行时知道数组的大小，它会更有用。当在对象类中封装 OpenCL 功能时，这是可取的。例如，参见上面 EdwardLuong 的评论；如果他的建议可行（似乎不适用于我的硬件），那就太好了。谢谢。

【参考方案3】：

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

如果您不小心使用本地内存，有时您的性能可能会比使用全局内存更差。