How do I use local memory in OpenCL?

Question

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.

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.

Tom · Accepted Answer · 2015-05-02 16:22:22Z

35

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

edited May 2, 2015 at 16:22

answered Apr 2, 2010 at 10:20

Tom

21.2k4 gold badges48 silver badges56 bronze badges

Sign up to request clarification or add additional context in comments.

11 Comments

splicer Over a year ago

I've read through the NVIDIA introductory material, and I still find the examples too complex. I'm looking for an über-simple 1-dimensional example of using local memory to get my feet wet.

splicer Over a year ago

Thanks for adding code in your last edit! I can't seem to get your kernel working though.... How would I use clSetKernelArg() for temp? Do I need to use clCreateBuffer() for temp? Also, there are a few typos in your kernel: "temp * temp" should be "temp[ltid] * temp[ltid]", and a closing brace should be inserted before the last line.

Tom Over a year ago

I corrected the typos - serves me right for typing on ipod. Your clSetKernelArg is not allocating enough memory though, you need space for one cl_float per thread (you have only allocated one float). Try: clSetKernelArg(kernel, 2, sizeof(cl_float) * local_work_size[0], NULL); where local_work_size[0] is the work group size in dimension 0.

user333 Over a year ago

Note that you can declare variables as local with the qualifier __local. For example, you could do __local float values[GROUP_SIZE]; then have each thread write values[get_local_id(0)] = .... Local memory doesn't need to be reached via a pointer passed into the kernel.

Tim Child Over a year ago

Remember local storage is only visible to the same work group. So you work group size needs to be > 1 to share it across multiple threads.

|

Rick-Rainer Ludwig · Accepted Answer · 2011-06-13 17:23:51Z

31

There is another possibility to do this, if the size of the local memory is constant. Without using a pointer in the kernels parameter list, the local buffer can be declared within the kernel just by declaring it __local:

__local float localBuffer[1024];

This removes code due to less clSetKernelArg calls.

answered Jun 13, 2011 at 17:23

Rick-Rainer Ludwig

2,3992 gold badges27 silver badges42 bronze badges

1 Comment

Bruce Dean Over a year ago

This is true but it would be way more useful if you didn't have to know the size of the array at run-time. This is desirable when encapsulating OpenCL functionality within and Object Class. E.g., see EdwardLuong's comment above; it would be great if his suggestion could work (does not seem to work for my hardware). Thanks.

Hunter Wang · Accepted Answer · 2013-07-03 20:38:55Z

5

In OpenCL local memory is meant to share data across all work items in a workgroup. And it usually requires to do a barrier call before the local memory data can be used (for example, one work item wants to read a local memory data that is written by the other work items). Barrier is costly in hardware. Keep in mind, local memory should be used for repeated data read/write. Bank conflict should be avoided as much as possible.

If you are not careful with local memory, you may end up with worse performance some time than using global memory.

answered Jul 3, 2013 at 20:38

Hunter Wang

1771 silver badge5 bronze badges

Collectives™ on Stack Overflow

How do I use local memory in OpenCL?

3 Answers 3

11 Comments

1 Comment

Comments

Your Answer

Linked

Hot Network Questions

Collectives™ on Stack Overflow

3 Answers 3

11 Comments

1 Comment

Comments

Your Answer

Sign up or log in

Post as a guest

Linked

Related