OpenCL Local Memory Tiling: Speeding Up Matrix Multiplication with Work-Group Shared Memory

If your OpenCL kernel re-reads the same global memory many times, the single biggest win available is usually not a better algorithm — it is moving the hot data into local memory, the fast on-chip scratchpad shared by all work-items in a work-group. Matrix multiplication is the canonical case: a naive kernel reads each element of A and B roughly N times from global memory, while a tiled kernel reads each element once per tile and serves the rest of the reads from local memory, which typically has latency orders of magnitude lower than global memory.

Why local memory helps

OpenCL executes a kernel as many work-items, organized into work-groups. Work-items in the same group share a region of memory declared with the __local qualifier. This memory lives on-chip (on GPUs it maps to shared memory per compute unit), so reads and writes to it avoid the long trip to device DRAM.

Tiling exploits data reuse. In C = A × B, every row of A is reused across all columns of C, and every column of B is reused across all rows. If a work-group of 16×16 work-items cooperatively loads a 16×16 tile of A and a 16×16 tile of B into local memory, each loaded element serves 16 multiply-accumulates instead of one. Global memory traffic drops by roughly a factor of the tile width.

A worked tiled matrix multiplication kernel

The sketch below assumes square matrices whose dimension is a multiple of the tile size (16), OpenCL 1.2 as the portability baseline, and row-major storage. Boundary handling for arbitrary sizes is covered in the mistakes section.

#define TILE 16

__kernel void matmul_tiled(__global const float *A,
                           __global const float *B,
                           __global float *C,
                           const int N)
{
    // Shared tiles, padded to reduce bank conflicts on some GPUs
    __local float As[TILE][TILE + 1];
    __local float Bs[TILE][TILE + 1];

    const int lx = get_local_id(0);
    const int ly = get_local_id(1);
    const int row = get_global_id(1);
    const int col = get_global_id(0);

    float sum = 0.0f;

    for (int t = 0; t < N / TILE; t++) {
        // Cooperatively load one tile of A and one of B.
        // Adjacent work-items read adjacent addresses (coalesced).
        As[ly][lx] = A[row * N + t * TILE + lx];
        Bs[ly][lx] = B[(t * TILE + ly) * N + col];

        // All work-items must finish writing before anyone reads.
        barrier(CLK_LOCAL_MEM_FENCE);

        for (int k = 0; k < TILE; k++)
            sum += As[ly][k] * Bs[k][lx];

        // Prevent the next iteration's loads from overwriting
        // tiles still being read.
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    C[row * N + col] = sum;
}

Three details do the real work:

• Cooperative loading. Each work-item loads exactly one element of each tile, so a 16×16 group fills both tiles with 256 coalesced global reads instead of 2 × 16 × 16 × 16 scattered ones.
• Two barriers per iteration. The first guarantees the tiles are fully written before any work-item reads them. The second guarantees no work-item starts overwriting the tiles for the next iteration while others are still accumulating. Dropping the second barrier is a classic source of intermittent wrong results.
• Row padding. As[TILE][TILE + 1] wastes a little local memory but shifts consecutive rows to different memory banks on architectures where same-bank column accesses serialize. Whether this matters is device-dependent; measure both ways.

Checking that your tile size actually fits

Local memory is a scarce, device-specific resource. Before committing to a tile size, query the device from host code (requires an initialized context and device; no special privileges):

cl_ulong localBytes;
clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE,
                sizeof(localBytes), &localBytes, NULL);

cl_ulong maxGroup;
clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE,
                sizeof(maxGroup), &maxGroup, NULL);

size_t kernelGroup;
clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE,
                         sizeof(kernelGroup), &kernelGroup, NULL);

Two 16×17 padded float tiles need about 2 × 16 × 17 × 4 ≈ 2.2 KB — trivially within the 32–64 KB typical of GPUs. But larger tiles grow quadratically: 32×32 padded tiles need ~8.7 KB, and local memory is shared between all concurrently resident work-groups on a compute unit, so big tiles can reduce occupancy (how many groups run at once) and hurt latency hiding. The largest work-group is not automatically the fastest; benchmark 8×8, 16×16, and 32×8 shapes. Note that CL_KERNEL_WORK_GROUP_SIZE can be smaller than the device maximum because it accounts for the kernel's register and local memory usage — a launch exceeding it fails with CL_INVALID_WORK_GROUP_SIZE.

Limits and common mistakes

• Barriers in divergent control flow. Every work-item in the group must reach the same barrier() call. A barrier inside an if that only some work-items take, or inside a loop with a work-item-dependent trip count, causes undefined behavior or deadlock. This is the most common tiling correctness bug.
• Tiling without reuse. For element-wise kernels (vector add, scaling), each global element is read exactly once; staging it through local memory only adds barrier and copy overhead. Skip tiling when there is no reuse pattern.
• Ignoring boundaries. The sketch assumes N is a multiple of TILE. For arbitrary sizes, guard the loads (As[ly][lx] = (row < N && t*TILE+lx < N) ? A[...] : 0.0f;) and the final store, or pad the matrices on the host.
• Assuming portability of tuning. Local memory capacity, bank behavior, and work-group limits vary widely across GPU, CPU, and FPGA devices. A tile size tuned on one device may fail to launch or run slowly on another. Keep the tile size a compile-time parameter and re-tune per target.
• Version assumptions. Everything here is OpenCL 1.2, the most portable baseline. Features like 2.x shared virtual memory are irrelevant to tiling and not universally supported; check device capabilities before relying on anything newer.

Verifying correctness and speedup

Validate the tiled kernel against a simple triple-loop CPU reference on small random matrices, comparing within a floating-point tolerance (e.g., relative error below 1e-4 for 1024×1024 float matrices). Then benchmark against a naive global-memory kernel using event profiling — create the queue with CL_QUEUE_PROFILING_ENABLE and read elapsed time via clGetEventProfilingInfo — across several matrix sizes and work-group shapes. Actual speedup is device- and size-dependent, so treat any fixed number with suspicion and measure on your target hardware.