How Does OpenCL Allocate and Access Global Memory?

OpenCL manages data across heterogeneous architectures by allocating memory objects called buffers on a compute device and mapping or copying data between the host and global device memory. This article explains how the host runtime creates global buffers using context APIs, how data transfers occur via command queues, how kernel functions declare and access this memory through address space qualifiers, and best practices for achieving optimal memory bandwidth.

The OpenCL Memory Model and Global Memory

OpenCL defines a tiered memory hierarchy comprising private, local, constant, and global memory spaces. Global memory represents the largest storage pool available on a compute device, such as a GPU or multi-core CPU. It persists across kernel executions within the same context and is accessible by all work-items across all work-groups. Because the host (CPU) and device often possess physically separate memory spaces, OpenCL uses an explicit memory management model where the host coordinates buffer creation, data migration, and kernel argument binding.

Allocating Global Memory Buffers on the Host

Buffer allocation takes place on the host using the OpenCL runtime API. The primary function responsible for creating a global memory buffer is clCreateBuffer.

cl_mem clCreateBuffer(cl_context context,
                      cl_mem_flags flags,
                      size_t size,
                      void *host_ptr,
                      cl_int *errcode_ret);

When invoking clCreateBuffer, developers specify allocation characteristics using memory flags:

• CL_MEM_READ_WRITE: The default access flag, allowing the compute device to both read from and write to the buffer.
• CL_MEM_READ_ONLY and CL_MEM_WRITE_ONLY: Restrict device access permissions, enabling driver-level optimizations.
• CL_MEM_COPY_HOST_PTR: Allocates device memory and immediately initializes it with data from a host pointer (host_ptr).
• CL_MEM_USE_HOST_PTR: Instructs the OpenCL implementation to use the host memory referenced by host_ptr directly, if supported by the architecture, reducing allocation overhead.
• CL_MEM_ALLOC_HOST_PTR: Allocates memory accessible by the host, useful for setting up pinned or zero-copy memory transfers.

Transferring Data to and from Global Buffers

Creating a buffer reserves memory on the target device context, but populating and retrieving data requires queueing explicit memory commands. The host program issues these operations to a cl_command_queue.

Explicit Enqueued Transfers

To write host data into an allocated buffer, the host enqueues a write operation:

clEnqueueWriteBuffer(queue, device_buffer, CL_TRUE, 0, buffer_size, host_array, 0, NULL, NULL);

The third argument specifies whether the operation is blocking (CL_TRUE) or non-blocking (CL_FALSE). A blocking write pauses host execution until the memory copy completes, ensuring the host memory buffer can be safely reused or modified immediately. After kernel execution finishes, clEnqueueReadBuffer transfers processed results back from the device global buffer to host memory.

Zero-Copy and Memory Mapping

For architectures where host and device share physical memory (such as integrated graphics or APUs), explicit copies introduce unnecessary latency. OpenCL provides memory mapping functions:

void *mapped_ptr = clEnqueueMapBuffer(queue, device_buffer, CL_TRUE, CL_MAP_WRITE, 0, buffer_size, 0, NULL, NULL, &err);
// Manipulate mapped_ptr directly on host
clEnqueueUnmapMemObject(queue, device_buffer, mapped_ptr, 0, NULL, NULL);

Mapping creates a pointer in the host address space directly bound to the device buffer, avoiding explicit duplication.

Passing Buffers to Compute Kernels

Before a kernel can operate on a global buffer, the host program must associate the buffer object with a specific kernel argument index using clSetKernelArg:

clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&device_buffer);

This binds the allocated memory object to the corresponding parameter in the OpenCL C kernel signature.

Accessing Global Memory Inside the Kernel

Inside OpenCL C kernel code, global memory is accessed by declaring pointer arguments with the __global (or global) address space qualifier.

__kernel void vector_add(__global const float *a,
                         __global const float *b,
                         __global float *result,
                         const unsigned int count) {
    int id = get_global_id(0);
    if (id < count) {
        result[id] = a[id] + b[id];
    }
}

Every work-item computes its unique index via built-in functions like get_global_id(). The pointer dereference reads from or writes to the underlying global memory address. Work-items can read across arbitrary offsets, but access patterns heavily dictate hardware throughput.

Performance Considerations for Global Memory Access

Global memory possesses high latency compared to local memory or registers. Maximizing performance requires optimizing memory transactions:

• Coalesced Memory Access: On GPU hardware, adjacent work-items within a warp or wavefront should access contiguous memory locations. When work-item i accesses address base + i, the hardware consolidates multiple requests into a single wide memory transaction.
• Alignment: Memory buffers should be aligned to hardware cache line boundaries (typically 64 or 128 bytes) to avoid split transactions.
• Vectorized Data Types: Utilizing built-in vector types such as float4 or int2 allows the memory controller to issue single, wider load and store instructions, improving throughput on devices with native vector registers.