如何在OpenCL中使用本地内存? [英] How do I use local memory in OpenCL?

问题描述

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

I've been playing with OpenCL recently, and I'm able to write simple kernels that use only global memory. Now I'd like to start using local memory, but I can't seem to figure out how to use get_local_size() and get_local_id() to compute one "chunk" of output at a time.

例如，假设我要将Apple的OpenCL Hello World示例内核转换为使用本地内存的内核.你会怎么做?这是原始的内核源代码:

For example, let's say I wanted to convert Apple's OpenCL Hello World example kernel to something the uses local memory. How would you do it? Here's the original kernel source:

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

如果无法轻松地将该示例转换为可以显示如何使用本地内存的内容，则可以使用任何其他简单示例.

If this example can't easily be converted into something that shows how to make use of local memory, any other simple example will do.

推荐答案

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

Check out the samples in the NVIDIA or AMD SDKs, they should point you in the right direction. Matrix transpose would use local memory for example.

使用平方内核，可以将数据暂存到中间缓冲区中.记住要传递附加参数.

Using your squaring kernel, you could stage the data in an intermediate buffer. Remember to pass in the additional parameter.

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

这篇关于如何在OpenCL中使用本地内存?的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！