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

Question

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

例如，假设我想要将苹果的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];
}

如果这个示例不能轻松转换为展示如何利用本地内存的示例，任何其他简单的示例也可以。

- splicer

Answer 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];
    }
}

- Tom

---

我已经阅读了英伟达的介绍材料，但我仍然觉得这些示例太过复杂。我正在寻找一个超级简单的一维示例，以使用本地内存来入门。 - splicer

谢谢您在上次编辑中添加代码！不过我好像无法让您的内核正常工作...我该如何使用clSetKernelArg()为temp设置参数？我需要为temp使用clCreateBuffer()吗？另外，您的内核中有一些拼写错误："temp * temp"应改为"temp[ltid] * temp[ltid]"，并且在最后一行之前应插入一个闭合括号。 - splicer

我纠正了错别字——因为使用iPod打字真是自作自受。但你的clSetKernelArg没有分配足够的内存，每个线程都需要一个cl_float的空间（你只分配了一个float）。尝试使用以下代码：clSetKernelArg(kernel, 2, sizeof(cl_float) * local_work_size[0], NULL); 其中local_work_size [0]是维度0的工作组大小。请注意不要改变原意。 - Tom

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

@mike：在上面的示例中，每个线程都使用本地内存的唯一部分作为自己的个人草稿板，因此不需要屏障。但是，如果线程将通过本地内存进行通信，即读取另一个线程写入的数据，则需要屏障。鉴于这是对初学者的回应，我应该提到这一点，因为这是未来明显的陷阱。 - Tom

Answer 2

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

__local float localBuffer[1024];

这将由于较少的 clSetKernelArg 调用而删除代码。

- Rick-Rainer Ludwig

---

这是正确的，但如果您不必在运行时知道数组的大小，那将更加有用。当封装Object Class中的OpenCL功能时，这是可取的。例如，请参见EdwardLuong上面的评论；如果他的建议可以实现就太好了（似乎对我的硬件不起作用）。谢谢。 - Bruce Dean

Answer 3

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

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

- Hunter Wang