OpenCL: The Formation of a Heterogeneous Parallel Computing Standard and Its Early Implementations

OpenCL, short for Open Computing Language, is a standard that brings multicore central processing units, graphics processing units and other compute devices within a common programming interface. Apple proposed the initial draft, which was then developed collectively within Khronos. The original December 2008 announcement listed Apple, AMD, Intel, NVIDIA and IBM among the participating companies. Neil Trevett chaired the working group.

On December 9, 2008, Khronos announced in Singapore that version 1.0 had been approved and publicly released. Apple had submitted its draft approximately six months earlier. OpenCL entered software products for users in 2009, while maintenance of the specification continued. Document Revision 48 of the surviving 1.0 specification is dated October 6, 2009.

References: [1] [2]

As power consumption and other constraints limited further increases in clock frequency, adding processor cores and using GPU parallelism became important ways to increase computational throughput. CPUs and GPUs differ substantially in instructions, memory hierarchies and suitable workloads. Earlier programming tools often targeted a particular vendor or type of hardware, making it difficult to coordinate several devices within a common source-code framework.

OpenCL combines APIs for coordinating computation with a language for describing it. OpenCL C is based on a subset of C99, with additional qualifiers and built-in functions for parallel execution. A host program selects devices, establishes the execution environment and schedules operations; kernel functions running on the devices process the data assigned to them.

References: [1] [4]

The execution model in version 1.0 separates the host program from kernels. The host creates a context containing devices, programs and memory objects, then uses command queues to schedule data transfers, kernel execution and synchronization. Queues may execute commands in order, or, in implementation-supported modes, allow the program to specify dependencies explicitly.

A kernel launch defines a one-, two- or three-dimensional index space. Each index corresponds to a work-item. Work-items execute the same kernel code but access data using their own identifiers, and are organized into work-groups. Memory is divided into global and constant regions, local regions shared within a group, and private regions for individual work-items, allowing the program to express the scope of data sharing.

References: [1]

Consider adding two arrays of length n. The host places the inputs in memory objects and launches a one-dimensional set of work-items covering their elements. The kernel obtains index i with get_global_id(0) and executes c[i]=a[i]+b[i] when i is less than n. Each work-item handles a single output position, allowing the device to process many elements in parallel rather than assigning the entire sequential loop to one core.

This example follows the specification’s indexing and memory models. An application must also create buffers, set kernel arguments, submit commands and wait for results to become available before returning the output to the host. If more work-items are allocated than there are array elements, a bounds check prevents out-of-range access. Dependencies between operations are expressed through queue ordering or synchronization events.

References: [1]

On August 24, 2009, Apple announced that Mac OS X Snow Leopard would go on sale on August 28, listing OpenCL among its technologies. Developers could use it to apply GPU computing power to tasks beyond graphics rendering. The announcement documents an early path from a public specification to a software environment available to users.

By 2010, a paper by Stone, Gohara and Shi was comparing OpenCL models and implementations on multicore CPUs, GPUs and Cell processors. Its discussion of runtime compilation describes how applications supply kernel source to the target device’s compilation tools, which generate executable code for that implementation, rather than requiring every device to execute the same machine instructions.

References: [3] [4]

A common interface reduces the need to reorganize an entire application for each device, but achieving the best speed with the same kernel on different hardware still requires specific analysis. The 2010 paper explains that a memory-access pattern suited to one GPU may not suit a Cell processor or multicore CPU. A program can query device capabilities and select appropriate data layouts, work-groups and kernel versions.

OpenCL’s early contribution was to establish a low-level agreement for parallel computing across vendors. It unified the basic expression of device discovery, program compilation, memory management and task submission, while leaving implementers room for hardware-specific optimization. The standard’s release in 2008, product adoption in 2009 and subsequent research on actual architectures together shaped its early development.

References: [4]