cxx/stdlib/work_item.txt

// Copyright 2017-2020 The Khronos Group. This work is licensed under a
// Creative Commons Attribution 4.0 International License; see
// http://creativecommons.org/licenses/by/4.0/

[[work-item-functions]]
=== Work-Item Functions

This section describes the library of 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.

[[header-opencl_work_item-synopsis]]
==== Header <opencl_work_item> Synopsis
[source]
----
namespace cl
{
uint get_work_dim();
size_t get_global_size(uint dimindx);
size_t get_global_id(uint dimindx);
size_t get_local_size(uint dimindx);
size_t get_enqueued_local_size(uint dimindx);
size_t get_local_id(uint dimindx);
size_t get_num_groups(uint dimindx);
size_t get_group_id(uint dimindx);
size_t get_global_offset(uint dimindx);
size_t get_global_linear_id();
size_t get_local_linear_id();
size_t get_sub_group_size();
size_t get_max_sub_group_size();
size_t get_num_sub_groups();
size_t get_enqueued_num_sub_groups();
size_t get_sub_group_id();
size_t get_sub_group_local_id();

}
----

[[work-item-operations]]
==== Work item operations

[[get_work_dim]]
===== get_work_dim
[source]
----
uint get_work_dim()
----

Returns the number of dimensions in use.
This is the value given to the `work_dim` argument specified in `clEnqueueNDRangeKernel`.

[[get_global_size]]
===== get_global_size
[source]
----
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`.

[[get_global_id]]
===== get_global_id
[source]
----
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`.

[[get_local_size]]
===== get_local_size
[source]
----
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 size [[ftnref19]] <<ftn19,[19]>>, 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`.

[[get_enqueued_local_size]]
===== get_enqueued_local_size
[source]
----
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`.

[[get_local_id]]
===== get_local_id
[source]
----
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`.

[[get_num_groups]]
===== get_num_groups
[source]
----
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`.

[[get_group_id]]
===== get_group_id
[source]
----
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`.

[[get_global_offset]]
===== get_global_offset
[source]
----
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`.

[[get_global_linear_id]]
===== get_global_linear_id
[source]
----
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) )`.

[[get_local_linear_id]]
===== get_local_linear_id
[source]
----
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)`

[[get_sub_group_size]]
===== get_sub_group_size
[source]
----
size_t get_sub_group_size()
----

Returns the number of work-items in the sub-group.
This value is no more than the maximum sub-group size and is implementation-defined based on a combination of the compiled kernel and the dispatch dimensions.
This will be a constant value for the lifetime of the sub-group.

[[get_max_sub_group_size]]
===== get_max_sub_group_size
[source]
----
size_t get_max_sub_group_size()
----

Returns the maximum size of a sub-group within the dispatch.
This value will be invariant for a given set of dispatch dimensions and a kernel object compiled for a given device.

[[get_num_sub_groups]]
===== get_num_sub_groups
[source]
----
size_t get_num_sub_groups()
----

Returns the number of sub-groups that the current work-group is divided into.

This number will be constant for the duration of a work-group's execution.
If the kernel is executed with a non-uniform work-group size <<ftn17,[17]>> values for any dimension, calls to this built-in from some work-groups may return different values than calls to this built-in from other work-groups.

[[get_enqueued_num_sub_groups]]
===== get_enqueued_num_sub_groups
[source]
----
size_t get_enqueued_num_sub_groups()
----

Returns the same value as that returned by `get_num_sub_groups()` 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 sub groups in each of the work groups that make up the uniform region of the global range.

[[get_sub_group_id]]
===== get_sub_group_id
[source]
----
size_t get_sub_group_id()
----

`get_sub_group_id()` returns the sub-group ID which is a number from `0 ... get_num_sub_groups()-1`.

For `clEnqueueTask`, this returns `0`.

[[get_sub_group_local_id]]
===== get_sub_group_local_id
[source]
----
size_t get_sub_group_local_id()
----

Returns the unique work-item ID within the current sub-group.
The mapping from `get_local_id(dimindx)` to `get_sub_group_local_id()` will be invariant for the lifetime of the work-group.