C Specification

To enqueue a command to execute a kernel on a device, call the function

// Provided by CL_VERSION_1_0
cl_int clEnqueueNDRangeKernel(
    cl_command_queue command_queue,
    cl_kernel kernel,
    cl_uint work_dim,
    const size_t* global_work_offset,
    const size_t* global_work_size,
    const size_t* local_work_size,
    cl_uint num_events_in_wait_list,
    const cl_event* event_wait_list,
    cl_event* event);

Parameters

  • command_queue is a valid host command-queue. The kernel will be queued for execution on the device associated with command_queue.

  • kernel is a valid kernel object. The OpenCL context associated with kernel and command-queue must be the same.

  • work_dim is the number of dimensions used to specify the global work-items and work-items in the work-group. work_dim must be greater than zero and less than or equal to CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.

  • global_work_offset can be used to specify an array of work_dim unsigned values that describe the offset used to calculate the global ID of a work-item. If global_work_offset is NULL, the global IDs start at offset (0, 0, 0). global_work_offset must be NULL before version 1.1.

  • global_work_size points to an array of work_dim unsigned values that describe the number of global work-items in work_dim dimensions that will execute the kernel function. The total number of global work-items is computed as global_work_size[0] × …​ × global_work_size[work_dim - 1]. If the device associated with command_queue is an OpenCL 2.1 or newer device, and global_work_size is NULL or the value in any passed dimension is zero, then the kernel command will trivially succeed after its event dependencies are satisfied and subsequently update its completion event. The behavior in this situation is similar to that of an enqueued marker, except that unlike a marker, an enqueued kernel with no events passed to event_wait_list may run at any time.

  • local_work_size points to an array of work_dim unsigned values that describe the number of work-items that make up a work-group (also referred to as the size of the work-group) that will execute the kernel specified by kernel. The total number of work-items in a work-group is computed as local_work_size[0] × …​ × local_work_size[work_dim - 1]. The total number of work-items in the work-group must be less than or equal to the CL_KERNEL_WORK_GROUP_SIZE value specified in the Kernel Object Device Queries table, and the number of work-items specified in local_work_size[0], …​, local_work_size[work_dim - 1] must be less than or equal to the corresponding values specified by CL_DEVICE_MAX_WORK_GROUP_SIZES[0], …​, CL_DEVICE_MAX_WORK_GROUP_SIZES[work_dim - 1]. The explicitly specified local_work_size will be used to determine how to break the global work-items specified by global_work_size into appropriate work-group instances.

  • event_wait_list and num_events_in_wait_list specify events that need to complete before this particular command can be executed. If event_wait_list is NULL, then this particular command does not wait on any event to complete. If event_wait_list is NULL, num_events_in_wait_list must be 0. If event_wait_list is not NULL, the list of events pointed to by event_wait_list must be valid and num_events_in_wait_list must be greater than 0. The events specified in event_wait_list act as synchronization points. The context associated with events in event_wait_list and command_queue must be the same. The memory associated with event_wait_list can be reused or freed after the function returns.

  • event returns an event object that identifies this command and can be used to query or wait for this command to complete. If event is NULL or the enqueue is unsuccessful, no event will be created and therefore it will not be possible to query the status of this command or to wait for this command to complete. If event_wait_list and event are not NULL, event must not refer to an element of the event_wait_list array.

Description

An ND-range kernel command may require uniform work-groups or may support non-uniform work-groups. To support non-uniform work-groups:

  1. The device associated with command_queue must support non-uniform work-groups.

  2. The program object associated with kernel must support non-uniform work-groups. Specifically, this means:

    1. If the program was created with clCreateProgramWithSource, then the program must be compiled for OpenCL C 2.0 or newer (e.g. by using the -cl-std=CL2.0 or newer build option), and without the -cl-uniform-work-group-size build option.

    2. If the program was created with clCreateProgramWithIL or clCreateProgramWithBinary, then the program must be compiled or built without the -cl-uniform-work-group-size build options.

    3. If the program was created using clLinkProgram, then all input programs must support non-uniform work-groups.

If non-uniform work-groups are supported, any single dimension for which the global size is not divisible by the local size will be partitioned into two regions. One region will have work-groups that have the same number of work-items as was specified by the local size parameter in that dimension. The other region will have work-groups with less than the number of work items specified by the local size parameter in that dimension. The global IDs and group IDs of the work-items in the first region will be numerically lower than those in the second, and the second region will be at most one work-group wide in that dimension. Work-group sizes could be non-uniform in multiple dimensions, potentially producing work-groups of up to 4 different sizes in a 2D range and 8 different sizes in a 3D range.

If non-uniform work-groups are supported and local_work_size is NULL, the OpenCL runtime may choose a uniform or non-uniform work-group size.

Otherwise, when non-uniform work-groups are not supported, the size of each work-group must be uniform. If local_work_size is specified, the values specified in global_work_size[0], …​, global_work_size[work_dim - 1] must be evenly divisible by the corresponding values specified in local_work_size[0], …​, local_work_size[work_dim - 1]. If local_work_size is NULL, the OpenCL runtime must choose a uniform work-group size.

The work-group size to be used for kernel can also be specified in the program source or intermediate language. In this case the size of work-group specified by local_work_size must match the value specified in the program source.

These work-group instances are executed in parallel across multiple compute units or concurrently on the same compute unit.

Each work-item is uniquely identified by a global identifier. The global ID, which can be read inside the kernel, is computed using the value given by global_work_size and global_work_offset. In addition, a work-item is also identified within a work-group by a unique local ID. The local ID, which can also be read by the kernel, is computed using the value given by local_work_size. The starting local ID is always (0, 0, …​, 0).

clEnqueueNDRangeKernel returns CL_SUCCESS if the kernel-instance was successfully queued. Otherwise, it returns one of the following errors:

  • CL_INVALID_COMMAND_QUEUE

    • if command_queue is not a valid host command-queue

  • CL_INVALID_KERNEL

    • if kernel is not a valid kernel

  • CL_INVALID_PROGRAM_EXECUTABLE

    • if there is no successfully built program executable available for the device associated with command_queue

  • CL_INVALID_CONTEXT

    • if the context associated with command_queue and kernel are not the same

    • if the context associated with command_queue and events in event_wait_list are not the same

  • CL_INVALID_KERNEL_ARGS

    • if any kernel arguments for kernel have not been set

  • CL_INVALID_WORK_DIMENSION

    • if work_dim is not valid for the device associated with command_queue (is greater than the value returned for CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS)

  • CL_INVALID_GLOBAL_OFFSET

    • if global_work_offset is not NULL. This error condition does not apply when the device associated with command_queue supports OpenCL 1.1 or newer.

    • if the value specified in global_work_size plus the corresponding value in global_work_offset for any dimensions is greater than the maximum value representable by size_t on the device associated with command_queue

  • CL_INVALID_GLOBAL_WORK_SIZE

    • if global_work_size is NULL. This error condition does not apply when the device associated with command_queue supports OpenCL 2.1 or newer.

    • if any of the values specified in global_work_size[0], …​ global_work_size[work_dim - 1] are zero. This error condition does not apply when the device associated with command_queue supports OpenCL 2.1 or newer.

    • if any of the values specified in global_work_size[0], …​ global_work_size[work_dim - 1] exceed the maximum value representable by size_t on the device associated with command_queue

  • CL_INVALID_WORK_GROUP_SIZE

    • if local_work_size is not NULL, if the work-group size must be uniform, and if the global_work_size is not evenly divisible by the local_work_size

    • if local_work_size is not NULL and if the total number of work-items in the work-group is zero

    • if local_work_size is not NULL and if the total number of work-items in the work-group is greater than the maximum work-group size supported for kernel on the device associated with command_queue (is greater than the value returned for CL_KERNEL_WORK_GROUP_SIZE)

    • if local_work_size is not NULL and if the local_work_size does not match the required work-group size for kernel

    • if local_work_size is not NULL and if the local_work_size is not consistent with the required number of sub-groups for kernel

  • CL_INVALID_WORK_ITEM_SIZE

    • if the number of work-items specified in any dimension of local_work_size is not valid for the device associated with command_queue (is greater than the corresponding value returned for CL_DEVICE_MAX_WORK_GROUP_SIZES)

  • CL_MISALIGNED_SUB_BUFFER_OFFSET

    • if a kernel argument for kernel is a sub-buffer object and the offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN for the device associated with command_queue. This error code is missing before version 1.1.

  • CL_INVALID_IMAGE_SIZE

    • if a kernel argument for kernel is an image and the dimensions of the image, such as the image width or image height, are not supported by the device associated with command_queue

  • CL_IMAGE_FORMAT_NOT_SUPPORTED

    • if a kernel argument for kernel is an image and the format of the image, such as the image channel order or image channel data type, are not supported by the device associated with command_queue

  • CL_MEM_OBJECT_ALLOCATION_FAILURE

    • if there is a failure to allocate memory for the data store associated with any buffer or image object kernel arguments for kernel

  • CL_INVALID_EVENT_WAIT_LIST

    • if event_wait_list is NULL and num_events_in_wait_list is greater than zero

    • if event_wait_list is not NULL and num_events_in_wait_list is zero

    • if event objects in event_wait_list are not valid events

  • CL_INVALID_OPERATION

    • if SVM pointers are set as arguments for kernel and the device associated with command_queue does not support SVM

    • if system pointers are set as arguments for kernel and the device associated with command_queue does not support fine-grain system SVM

  • CL_OUT_OF_RESOURCES

    • if there is a failure to allocate resources required by the OpenCL implementation on the device

  • CL_OUT_OF_HOST_MEMORY

    • if there is a failure to allocate resources required by the OpenCL implementation on the host

See Also

CL_VERSION_1_0, cl_command_queue, cl_event, cl_kernel, cl_uint

Document Notes

For more information, see the OpenCL Specification

This page is extracted from the OpenCL Specification. Fixes and changes should be made to the Specification, not directly.

Copyright

Copyright 2014-2026 The Khronos Group Inc.

SPDX-License-Identifier: CC-BY-4.0