Choosing Work‑Group Size and Local Memory for OpenCL Kernels

Quick Answer

Select a local work‑group size that evenly divides the global size and stays within the device’s CL_DEVICE_MAX_WORK_GROUP_SIZE and CL_DEVICE_LOCAL_MEM_SIZE. Launch the kernel with clEnqueueNDRangeKernel and declare any on‑chip buffers using __local inside the kernel.

Worked Example

Host‑side setup

cl_int err;
cl_device_id dev;
cl_context ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, &err);
cl_command_queue queue = clCreateCommandQueue(ctx, dev, 0, &err);

// 1. Query device limits
size_t maxWGSize;
cl_ulong localMemSize;
err = clGetDeviceInfo(dev, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWGSize), &maxWGSize, NULL);
err |= clGetDeviceInfo(dev, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(localMemSize), &localMemSize, NULL);
if (err != CL_SUCCESS) { /* handle error */ }

// 2. Choose sizes – example values that satisfy the limits
const size_t globalSize = 1024;          // total work‑items
const size_t localSize  = 256;           // work‑items per work‑group
if (globalSize % localSize != 0) { /* error: not a multiple */ }
if (localSize > maxWGSize) { /* error: exceeds device limit */ }

// 3. Allocate a buffer that will be processed
float *hostData = (float*)malloc(globalSize * sizeof(float));
// …initialise hostData…
cl_mem buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                            globalSize * sizeof(float), hostData, &err);

// 4. Create program & kernel (omitted for brevity)
cl_kernel kernel = clCreateKernel(program, "reduce", &err);

// 5. Set kernel arguments – note the third argument is the local memory size
size_t localMemBytes = localSize * sizeof(float);
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf);
err |= clSetKernelArg(kernel, 1, localMemBytes, NULL); // __local buffer
if (err != CL_SUCCESS) { /* handle error */ }

// 6. Enqueue the kernel
size_t global[1] = { globalSize };
size_t local[1]  = { localSize };
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, local, 0, NULL, NULL);
clFinish(queue);

Kernel code

__kernel void reduce(__global float *data, __local float *scratch) {
    int gid = get_global_id(0);
    int lid = get_local_id(0);
    // Load data into fast local memory
    scratch[lid] = data[gid];
    barrier(CLK_LOCAL_MEM_FENCE);

    // Simple tree reduction in local memory
    for (int s = get_local_size(0) / 2; s > 0; s >>= 1) {
        if (lid < s) {
            scratch[lid] += scratch[lid + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    // Write back the per‑work‑group result
    if (lid == 0) {
        data[get_group_id(0)] = scratch[0];
    }
}

Device Limits to Check

• CL_DEVICE_MAX_WORK_GROUP_SIZE – maximum total work‑items per work‑group (often 1024 on modern GPUs).
• CL_DEVICE_MAX_WORK_ITEM_SIZES – per‑dimension caps; for a 1‑D launch only the first entry matters.
• CL_DEVICE_LOCAL_MEM_SIZE – bytes of on‑chip local memory per compute unit (e.g., 48 KB). The kernel’s __local allocation plus any driver overhead must fit.

Always query these values at runtime; hard‑coding them makes the program fragile across hardware generations.

Common Mistakes

1. Mismatched global/local sizes – globalSize % localSize != 0 leads to CL_INVALID_WORK_GROUP_SIZE or undefined behaviour.
2. Exceeding per‑dimension limits – even if the total fits, a single dimension larger than CL_DEVICE_MAX_WORK_ITEM_SIZES[0] fails.
3. Oversubscribing local memory – requesting more than CL_DEVICE_LOCAL_MEM_SIZE causes CL_OUT_OF_RESOURCES at launch.
4. Missing or stale kernel arguments – forgetting clSetKernelArg for the __local buffer (argument 1 in the example) yields CL_INVALID_ARGUMENT.
5. Reusing a buffer after it has been released – the buffer object must stay alive until the kernel finishes; otherwise the runtime may read freed memory.

Verification Checklist

• Run the program on the target device and confirm clEnqueueNDRangeKernel returns CL_SUCCESS.
• Compare the kernel output against a known‑good reference (e.g., a CPU reduction).
• Print the queried maxWGSize and localMemSize and verify they are ≥ the values used.
• Use a profiler (e.g., Intel VTune, NVIDIA Nsight) to ensure the kernel achieves expected occupancy; low occupancy often signals a work‑group size that is too large or too small for the hardware.