## Description

The following table describes the list of built-in work-item functions that can be used to query the number of dimensions, the global and local work size specified to clEnqueueNDRangeKernel, and the global and local identifier of each work-item when this kernel is being executed on a device.

 Function Description uint get_work_dim() Returns the number of dimensions in use. This is the value given to the work_dim argument specified in clEnqueueNDRangeKernel. size_t get_global_size(uint dimindx) Returns the number of global work-items specified for dimension identified by dimindx. This value is given by the global_work_size argument to clEnqueueNDRangeKernel. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_global_size() returns 1. size_t get_global_id(uint dimindx) Returns the unique global work-item ID value for dimension identified by dimindx. The global work-item ID specifies the work-item ID based on the number of global work-items specified to execute the kernel. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_global_id() returns 0. size_t get_local_size(uint dimindx) Returns the number of local work-items specified in dimension identified by dimindx. This value is at most the value given by the local_work_size argument to clEnqueueNDRangeKernel if local_work_size is not NULL; otherwise the OpenCL implementation chooses an appropriate local_work_size value which is returned by this function. If the kernel is executed with a non-uniform work-group size28, calls to this built-in from some work-groups may return different values than calls to this built-in from other work-groups. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_local_size() returns 1. size_t get_enqueued_local_size( uint dimindx) Returns the same value as that returned by get_local_size(dimindx) if the kernel is executed with a uniform work-group size. If the kernel is executed with a non-uniform work-group size, returns the number of local work-items in each of the work-groups that make up the uniform region of the global range in the dimension identified by dimindx. If the local_work_size argument to clEnqueueNDRangeKernel is not NULL, this value will match the value specified in local_work_size[dimindx]. If local_work_size is NULL, this value will match the local size that the implementation determined would be most efficient at implementing the uniform region of the global range. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_enqueued_local_size() returns 1. size_t get_local_id(uint dimindx) Returns the unique local work-item ID, i.e. a work-item within a specific work-group for dimension identified by dimindx. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_local_id() returns 0. size_t get_num_groups(uint dimindx) Returns the number of work-groups that will execute a kernel for dimension identified by dimindx. Valid values of dimindx are 0 to get_work_dim() - 1. For other values of dimindx, get_num_groups() returns 1. size_t get_group_id(uint dimindx) get_group_id returns the work-group ID which is a number from 0 .. get_num_groups(dimindx) - 1. Valid values of dimindx are 0 to get_work_dim() - 1. For other values, get_group_id() returns 0. size_t get_global_offset(uint dimindx) get_global_offset returns the offset values specified in global_work_offset argument to clEnqueueNDRangeKernel. Valid values of dimindx are 0 to get_work_dim() - 1. For other values, get_global_offset() returns 0. size_t get_global_linear_id() Returns the work-items 1-dimensional global ID. For 1D work-groups, it is computed as get_global_id(0) - get_global_offset(0). For 2D work-groups, it is computed as (get_global_id(1) - get_global_offset(1)) * get_global_size(0) + (get_global_id(0) - get_global_offset(0)). For 3D work-groups, it is computed as get_global_id(2) - get_global_offset(2 * get_global_size(1) * get_global_size(0)) + get_global_id(1) - get_global_offset(1 * get_global_size (0)) + (get_global_id(0) - get_global_offset(0)). size_t get_local_linear_id() Returns the work-items 1-dimensional local ID. For 1D work-groups, it is the same value as get_local_id(0). For 2D work-groups, it is computed as get_local_id(1) * get_local_size(0) + get_local_id(0). For 3D work-groups, it is computed as (get_local_id(2) * get_local_size(1) * get_local_size(0)) + (get_local_id(1) * get_local_size(0)) + get_local_id(0).

[28] I.e. the global_work_size values specified to clEnqueueNDRangeKernel are not evenly divisible by the local_work_size values for each dimension.