Using OpenCL Local Memory to Cut Global Memory Traffic in Tiled Kernels

Problem: Global Memory Bandwidth Bottleneck

In many data‑parallel kernels the dominant cost is not the arithmetic itself but the repeated reads and writes to global memory. Global memory has high latency and limited bandwidth compared to the on‑chip scratchpad that OpenCL exposes as local memory (also called LDS on AMD hardware). When each work‑item fetches the same data independently, the memory subsystem becomes saturated and the compute units spend cycles waiting for data.

Solution: Cache Data in OpenCL Local Memory

Local memory provides a high‑bandwidth, low‑latency region that is shared by all work‑items in a single work‑group. By loading a tile of data from global memory into this shared region once, and then having all work‑items read from local memory, the number of global memory transactions can be reduced dramatically. The pattern is:

1. Each work‑item loads a portion of the tile into __local storage.
2. A work‑group barrier (cl_barrier(CLK_LOCAL_MEM_FENCE)) guarantees that every item has finished its load before any item starts computing.
3. The kernel performs its computation using the data now resident in local memory.
4. Optionally, a second barrier can be used before writing results back to global memory.

Correct tiling requires that the tile size fits within the available local memory, which is device‑specific and typically ranges from a few kilobytes to 64 KB.

Worked Example: Tiled Matrix Multiplication

Consider a simple square matrix multiplication C = A × B where each element of C is the dot product of a row of A and a column of B. A naïve kernel would read each needed element of A and B from global memory for every work‑item, leading to O(N³) global memory accesses. By tiling the matrices into blocks that fit in local memory, each block of A and B is loaded once per work‑group and reused.

__kernel void matmul_tiled(
    const __global float* A,
    const __global float* B,
    __global float* C,
    const int N,
    const int TILE)   // TILE must divide N and fit in local memory
{
    int row = get_group_id(1) * TILE + get_local_id(1);
    int col = get_group_id(0) * TILE + get_local_id(0);

    __local float As[TILE][TILE];
    __local float Bs[TILE][TILE];

    float acc = 0.0f;
    for (int t = 0; t < N; t += TILE) {
        // Load tile of A and B into local memory
        int aRow = row;
        int aCol = t + get_local_id(0);
        int bRow = t + get_local_id(1);
        int bCol = col;
        if (aRow < N && aCol < N) As[get_local_id(1)][get_local_id(0)] = A[aRow * N + aCol];
        else                        As[get_local_id(1)][get_local_id(0)] = 0.0f;
        if (bRow < N && bCol < N) Bs[get_local_id(1)][get_local_id(0)] = B[bRow * N + bCol];
        else                        Bs[get_local_id(1)][get_local_id(0)] = 0.0f;

        cl_barrier(CLK_LOCAL_MEM_FENCE);

        // Compute partial dot product using local memory
        for (int k = 0; k < TILE; ++k) {
            acc += As[get_local_id(1)][k] * Bs[k][get_local_id(0)];
        }

        cl_barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (row < N && col < N) C[row * N + col] = acc;
}

Key points in the example:

• The __local qualifier declares the tile arrays As and Bs that reside in LDS.
• Each work‑item loads one element of A and one element of B into the shared tile.
• The first barrier ensures the tile is fully populated before any work‑item starts the inner product loop.
• The second barrier guarantees that all work‑items have finished using the tile before the next iteration overwrites it.
• If TILE is chosen too large, the local memory arrays may exceed the device limit, causing compilation errors or runtime failures.

Trade‑offs and Practical Limits

• Bank conflicts: Local memory is divided into banks; simultaneous accesses to the same bank serialize. Access patterns with stride‑1 or power‑of‑two strides can cause conflicts. Choosing tile dimensions that are not multiples of the bank count (often 16 or 32) helps mitigate this.
• Occupancy impact: Each work‑group reserves a portion of local memory. Over‑allocating reduces the number of work‑groups that can reside simultaneously on a compute unit, lowering occupancy and hiding latency less effectively.
• Synchronization cost: Barriers incur overhead; if the tile is too small, the barrier frequency may outweigh the benefit of reduced global traffic.
• Correctness: Missing or mismatched barriers lead to race conditions or deadlocks, especially when the work‑group size is not a multiple of the hardware wavefront/warp size.

Closing: How to Verify and Tune

1. Check the device’s local memory size with clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, ...) and ensure TILE*TILE*sizeof(float) stays comfortably below that limit.
2. Run the kernel with a profiler (e.g., Intel VTune, AMD Radeon GPU Profiler) and compare global memory bandwidth versus local memory hit rate. Look for a drop in global memory transactions proportional to the tile size.
3. Vary the local work‑size (which must be a multiple of the warp/wavefront, typically 32 or 64) and observe changes in occupancy and execution time. The sweet spot balances tile reuse, bank conflict avoidance, and sufficient active work‑groups.
4. If bank conflicts are suspected, experiment with padding the local arrays (e.g., __local float As[TILE][TILE+1]) to change stride patterns and re‑measure.

By deliberately moving reusable data into the fast scratchpad offered by OpenCL local memory and synchronizing with barriers, many memory‑bound kernels can achieve significant speed‑ups. The technique requires careful sizing, awareness of bank architecture, and validation with profiling tools, but the payoff—lower latency and higher compute utilization—is often well worth the effort.