Using OpenCL Sub‑Buffers for Parallel Kernel Execution: A Practical Guide

Why Sub‑Buffers Matter

When a single device memory object is shared by several kernels, copying data between buffers is often the simplest solution. Sub‑buffers avoid that copy by carving the parent buffer into non‑overlapping regions that can be passed to kernels as independent cl_mem objects. This enables true concurrent execution on out‑of‑order queues and reduces memory traffic.

Creating a Sub‑Buffer

The API call is clCreateSubBuffer. It requires the parent buffer, a flag specifying the creation type, and a region descriptor.

cl_mem parent = clCreateBuffer(context, CL_MEM_READ_WRITE, 1024*1024, NULL, &err); // 1 MiB

cl_buffer_region region = { .origin = 0, .size = 512*1024 }; // first half
cl_mem sub1 = clCreateSubBuffer(parent, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);

region.origin = 512*1024; // second half
cl_mem sub2 = clCreateSubBuffer(parent, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);

Key points:

• CL_BUFFER_CREATE_TYPE_REGION tells the runtime that the sub‑buffer is a contiguous slice.
• The origin and size must be multiples of the device’s CL_DEVICE_MEM_BASE_ADDR_ALIGN value.
• Sub‑buffers inherit the parent’s memory flags but are otherwise independent objects.

Passing Sub‑Buffers to Kernels

Sub‑buffers are treated like any other buffer when setting kernel arguments:

clSetKernelArg(kernel1, 0, sizeof(cl_mem), &sub1);
clSetKernelArg(kernel2, 0, sizeof(cl_mem), &sub2);

Enqueue each kernel on a separate out‑of‑order queue:

cl_command_queue q1 = clCreateCommandQueue(context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
cl_command_queue q2 = clCreateCommandQueue(context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);

clEnqueueNDRangeKernel(q1, kernel1, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);
clEnqueueNDRangeKernel(q2, kernel2, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);

Because the sub‑buffers occupy disjoint memory ranges, the device can run both kernels simultaneously, provided it supports concurrent execution on out‑of‑order queues.

Reading Back the Results

When the kernels finish, read the entire parent buffer using a normal read operation and verify the two halves contain the expected values. You cannot use clEnqueueReadBuffer directly on a sub‑buffer; instead you read the parent and slice the data in host code.

clEnqueueReadBuffer(q1, parent, CL_TRUE, 0, 1024*1024, hostData, 0, NULL, NULL);

Common Mistakes and How to Avoid Them

• Misaligned offsets: If origin or size is not a multiple of CL_DEVICE_MEM_BASE_ADDR_ALIGN, clCreateSubBuffer returns CL_INVALID_VALUE. Query the alignment with clGetDeviceInfo before creating sub‑buffers.
• Overlapping regions: Accidentally overlapping sub‑buffers introduces race conditions. Keep a record of each region’s bounds and assert that they do not intersect.
• Reference counting errors: Sub‑buffers depend on the parent. Releasing the parent while a kernel still uses a sub‑buffer causes CL_INVALID_MEM_OBJECT. Release sub‑buffers first, then the parent.
• Using read/write buffer on sub‑buffer: clEnqueueReadBuffer/clEnqueueWriteBuffer do not accept sub‑buffers. Use clEnqueueReadBufferRect/clEnqueueWriteBufferRect or map/unmap with an offset.
• Assuming image sub‑buffers work on all devices: Only buffer objects are guaranteed to support sub‑buffers. Image sub‑buffers may be unsupported on pre‑OpenCL 1.2 GPUs.

Verifying the Feature Works

1. Compile the example code above and run it on a device that reports CL_DEVICE_PARTITION_HANDLE_CAPABILITIES or supports out‑of‑order execution.
2. After kernel execution, inspect hostData. The first half should contain the value written by kernel1 and the second half the value written by kernel2.
3. Call clGetMemObjectInfo with CL_MEM_ASSOCIATED_MEMOBJECT on each sub‑buffer; the returned handle should match the parent buffer’s ID.
4. Try creating a sub‑buffer with an offset that is not a multiple of the alignment. The call should return CL_INVALID_VALUE, confirming the alignment rule.
5. Measure performance by timing the sequential execution of both kernels on the parent buffer versus the concurrent execution using sub‑buffers. A noticeable speedup indicates that the device is truly exploiting parallelism.

When Sub‑Buffers Are Not the Right Tool

If the device does not support out‑of‑order queues or if the kernels require overlapping data, sub‑buffers may not provide a benefit. In such cases, traditional separate buffers or double‑buffering with explicit copies might be simpler and more portable.

Practical Checklist

• Check CL_DEVICE_MEM_BASE_ADDR_ALIGN and align origin and size.
• Ensure sub‑buffers do not overlap unless intentional.
• Release sub‑buffers before the parent.
• Use out‑of‑order queues for concurrent execution.
• Verify with clGetMemObjectInfo that sub‑buffers are correctly associated.

Conclusion

Sub‑buffers give you a lightweight, copy‑free way to partition a large buffer and run kernels in parallel on the same device memory space. By respecting alignment, lifetime, and queue requirements, you can achieve higher throughput on streaming workloads while keeping memory usage minimal.