OpenCL Local Memory: When Caching Data in the Work-Group Actually Pays Off

Your OpenCL kernel is mathematically correct, the compiler isn't complaining, and yet it runs at a fraction of the device's peak throughput. In many cases the bottleneck isn't arithmetic at all — it's that every work-item is fetching the same data from global memory over and over. The fix is usually the same pattern: stage a tile of data into local memory once, let the whole work-group reuse it, and only go back to global memory for the next tile. This post walks through when that trade is worth making, using a simple 1D blur as the worked example.

The memory hierarchy in practical terms

OpenCL exposes a few distinct memory spaces, but two matter most for bandwidth-bound kernels:

• Global memory (__global) is the large device DRAM. Every work-item can see it, but a load can cost hundreds of cycles of latency, and the bandwidth is shared across the entire device.
• Local memory (__local) is a small, fast scratchpad — often tens of kilobytes — shared by the work-items within a single work-group. On GPUs it typically maps to on-chip SRAM, so accesses are dramatically cheaper than a round trip to DRAM.

The catch: local memory is per-work-group, not per-device, and it's small. Query the actual budget with clGetDeviceInfo using CL_DEVICE_LOCAL_MEM_SIZE before assuming anything — the value varies widely between vendors and even between generations from the same vendor.

The tiling pattern

The core idea is simple. If each work-item in a group reads neighboring elements — as in stencils, convolutions, and matrix multiplication — the group's combined read set overlaps heavily. Instead of letting every work-item hit global memory for its neighbors, you:

1. Cooperatively load the union of everything the group needs into a __local array (the "tile," plus any halo/edge elements).
2. Synchronize so the tile is fully populated.
3. Compute, reading only from local memory.

A global read that would have happened N times now happens once. Whether that wins depends on the reuse factor — a kernel where each element is read exactly once gains nothing from tiling and may even lose time to the extra copy.

Worked example: a 1D blur

Consider a 3-tap blur, out[i] = (in[i-1] + in[i] + in[i+1]) / 3. A naive kernel does three global loads per work-item, and adjacent work-items re-read two of the same elements. The tiled version loads one element per work-item plus two halo elements:

__kernel void blur_tiled(__global const float* in,
                         __global float* out,
                         int n,
                         __local float* tile)
{
    int lid = get_local_id(0);
    int gid = get_global_id(0);
    int lsz = get_local_size(0);

    // Main body: one global load per work-item
    tile[lid + 1] = (gid < n) ? in[gid] : 0.0f;

    // Halo: first and last work-item fetch the edges
    if (lid == 0)
        tile[0] = (gid > 0) ? in[gid - 1] : 0.0f;
    if (lid == lsz - 1)
        tile[lsz + 1] = (gid + 1 < n) ? in[gid + 1] : 0.0f;

    barrier(CLK_LOCAL_MEM_FENCE);

    if (gid < n)
        out[gid] = (tile[lid] + tile[lid + 1] + tile[lid + 2]) / 3.0f;
}

The host side allocates the local buffer dynamically, e.g. with clSetKernelArg(kernel, 3, (localSize + 2) * sizeof(float), NULL). Run this from your host application after creating the kernel from a program object; you need a context and command queue already set up. A meaningful check is to compare output against a CPU reference implementation for a small input, then benchmark both the naive and tiled kernels on realistic sizes with a vendor profiler (NVIDIA Nsight, Radeon GPU Profiler, or Intel's tools) to confirm global-load counts actually dropped.

Why the barrier matters — and how it bites

The barrier(CLK_LOCAL_MEM_FENCE) call guarantees that every work-item in the group has written its portion of the tile before anyone reads a neighbor's portion. Two rules follow:

• All work-items in the group must reach the barrier. A barrier inside divergent control flow (e.g., inside an if that only some work-items take) is undefined behavior and can hang the kernel. Note the halo loads above are guarded, but the barrier is not.
• Barriers only synchronize within a work-group. There is no portable global synchronization inside a single kernel launch; cross-group coordination requires separate kernel launches or atomics.

The trade-off: occupancy

Local memory is not free. Each work-group's allocation is carved out of the CU's (compute unit's) fixed scratchpad. If a CU has 64 KB of local memory and each work-group requests 32 KB, at most two groups can be resident on that CU at once — regardless of how many the register file or work-item limits would otherwise allow. Fewer resident groups means fewer warps/wavefronts available to hide latency, which can erase the bandwidth savings.

Practical guidance:

• Keep tiles as small as the algorithm's reuse pattern allows.
• Check occupancy effects in a profiler rather than reasoning from first principles; the interaction with register pressure is hardware-specific.
• Watch for bank conflicts on GPUs where local memory is banked — strided access patterns can serialize what should be parallel local reads.
• Don't tile data with no reuse. The copy cost is pure overhead in that case.

Closing: measure, don't assume

The decision procedure is short: identify whether your kernel re-reads global data within a work-group; if yes, stage that data in local memory with a single barrier; then verify with CL_DEVICE_LOCAL_MEM_SIZE that your tile fits and with a profiler that occupancy and bandwidth actually improved. Local memory is one of the highest-leverage optimizations in OpenCL, but only when the reuse is real — otherwise you've just added a copy and a synchronization point to a kernel that was fine as it was.