cl_ext_buffer_device_address(3) Manual Page

Shared Virtual Memory (SVM) introduced in OpenCL 2.0 is the first feature that enables raw pointers in the OpenCL standard. Its coarse-grain variant is relatively simple to implement on various platforms in terms of coherency requirements, but it requires mapping the buffer’s address range to the host virtual address space. However, various higher-level heterogeneous APIs present a memory allocation routine which can allocate device-only memory and provide raw addresses to it without guarentees of system-wide uniqueness. For example, minimal implementations of OpenMP’s omp_target_alloc() and CUDA/HIP’s cudaMalloc()/hipMalloc() do not require a shared address space between the host and the device.

Host-device unified addressing might not be a major implementation issue in systems which can provide virtual memory across the platform, but might bring challenges in cases where the device presents a global memory with a disjoint address space (that can also be a physical memory address space) or, for example, when a barebone embedded system lacks virtual memory support altogether. This extension is targeted to complement the OpenCL SVM extension by providing an additional lower-end step in the spectrum of type of pointers/buffers OpenCL can allocate.