OpenCL Memory Trade-offs: When Local Memory Beats Global for Data-Parallel Kernels

The Problem: Bandwidth Starvation in Data-Parallel Workloads

You've written an OpenCL kernel that looks correct. It reads from global memory, computes, writes back. But the profiler shows the compute units sitting idle while memory transactions pile up. The issue isn't arithmetic intensity—it's that every work-item fetches the same global data repeatedly, saturating the bus. This is the classic global memory bottleneck, and the fix isn't always obvious.

OpenCL's hierarchical memory model gives you a lever: local memory (on-chip SRAM shared within a work-group). Used correctly, it turns redundant global loads into a single cooperative fetch. Used poorly, it caps occupancy or introduces bank conflicts. This post walks through the trade-off with a concrete matrix multiplication example.

Memory Hierarchy in One Paragraph

Global memory is the large, high-latency DRAM visible to all work-items across all work-groups. Local memory is a small (typically 16–64 KB per compute unit), low-latency scratchpad shared only among work-items in the same work-group. Accesses to local memory don't go off-chip. The catch: you must explicitly load data from global to local, synchronize with a barrier, then compute. If the work-group's local memory usage exceeds CL_DEVICE_LOCAL_MEM_SIZE, the kernel simply won't launch.

Worked Example: Tiled Matrix Multiplication

Consider multiplying two N×N matrices A and B to produce C. The naive kernel has each work-item compute one element of C by iterating k from 0 to N-1, loading A[row][k] and B[k][col] from global memory every iteration. That's 2N global loads per output element.

Naive Global-Memory Kernel (Simplified)

__kernel void matmul_naive(__global const float* A,
                           __global const float* B,
                           __global float* C,
                           int N) {
    int row = get_global_id(0);
    int col = get_global_id(1);
    float sum = 0.0f;
    for (int k = 0; k < N; ++k) {
        sum += A[row * N + k] * B[k * N + col];
    }
    C[row * N + col] = sum;
}

Tiled Version Using Local Memory

We split the k loop into tiles of size TILE (e.g., 16 or 32). Each work-group loads a TILE×TILE block of A and a TILE×TILE block of B into local memory, synchronizes, computes partial dot products, then moves to the next tile.

#define TILE 16
__kernel void matmul_tiled(__global const float* A,
                           __global const float* B,
                           __global float* C,
                           int N) {
    __local float As[TILE][TILE];
    __local float Bs[TILE][TILE];

    int bx = get_group_id(0), by = get_group_id(1);
    int tx = get_local_id(0), ty = get_local_id(1);
    int row = by * TILE + ty;
    int col = bx * TILE + tx;

    float sum = 0.0f;
    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        // Cooperative load
        int tiledIndex = t * TILE;
        int aCol = tiledIndex + tx;
        int bRow = tiledIndex + ty;

        if (row < N && aCol < N)
            As[ty][tx] = A[row * N + aCol];
        else
            As[ty][tx] = 0.0f;

        if (bRow < N && col < N)
            Bs[ty][tx] = B[bRow * N + col];
        else
            Bs[ty][tx] = 0.0f;

        barrier(CLK_LOCAL_MEM_FENCE); // All writes visible

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

        barrier(CLK_LOCAL_MEM_FENCE); // Avoid overwrite next tile
    }

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

Each tile loads 2 × TILE² elements from global memory but performs TILE³ multiply-adds locally. For TILE=16, that's 512 global loads yielding 4,096 FLOPs—a 8× reduction in global traffic per FLOP.

Where the Trade-offs Bite

1. Local Memory Capacity Limits Work-Group Size

The tiled kernel above uses 2 × TILE² × sizeof(float) bytes of local memory. With TILE=32, that's 8 KB per work-group. If your GPU has 64 KB local memory per compute unit, you can fit at most 8 work-groups concurrently. If each work-group has 1024 work-items (32×32), that's 8,192 concurrent work-items per CU. But if you push TILE=64, local memory jumps to 32 KB—only 2 work-groups fit, cutting potential occupancy. Always query CL_DEVICE_LOCAL_MEM_SIZE and CL_DEVICE_MAX_WORK_GROUP_SIZE at runtime to pick a tile size that balances reuse and occupancy.

2. Bank Conflicts Can Erase Gains

Local memory is organized in banks (typically 32). If multiple work-items in a warp/wavefront access different addresses in the same bank simultaneously, the accesses serialize. In the tiled kernel, As[ty][k] and Bs[k][tx] are conflict-free for power-of-two TILE because consecutive threads access consecutive banks. But transposed access patterns or non-power-of-two tiles can cause 2-way or worse conflicts. Use clGetDeviceInfo with CL_DEVICE_LOCAL_MEM_TYPE (should be CL_LOCAL for dedicated SRAM) and profile with hardware counters to detect this.

3. No Cross-Work-Group Consistency

Local memory is private to a work-group. If your algorithm needs data sharing across work-groups, you must go through global memory (or use atomics/device-side enqueue). Don't assume a barrier synchronizes across work-groups—it doesn't.

How to Verify on Your Hardware

1. Query limits: Call clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, ...) and CL_DEVICE_MAX_WORK_GROUP_SIZE before choosing TILE.
2. Enable profiling: Create the command queue with CL_QUEUE_PROFILING_ENABLE. Record start/end events around clEnqueueNDRangeKernel and read CL_PROFILING_COMMAND_END - CL_PROFILING_COMMAND_START.
3. Compare kernels: Run both naive and tiled versions on the same input sizes (e.g., 512×512, 1024×1024). Expect 2–5× speedup on modern GPUs for compute-bound sizes; memory-bound small sizes may show less.
4. Check occupancy: Use vendor profiling tools (NVIDIA Nsight, AMD ROCm Profiler, Intel VTune) to see achieved vs. theoretical occupancy. If local memory pressure drops occupancy below ~50%, reduce TILE or work-group size.

Closing Checklist

• Is your kernel bandwidth-bound? (Compute units idle, memory throughput near peak)
• Can you identify a reusable data block per work-group?
• Does 2 × TILE² × element_size ≤ CL_DEVICE_LOCAL_MEM_SIZE / target_workgroups_per_CU?
• Are access patterns bank-conflict-free for your tile size?
• Did you profile both versions on the target device?

Local memory isn't free speed—it's a budget you spend for bandwidth reduction. Spend it where the reuse rate justifies the occupancy cost, and always measure.