Description
The functionality described in this section requires
support for OpenCL C 2.0, or OpenCL C 3.0 or newer and the
__opencl_c_ feature.

This section describes builtin functions that perform collective options
across a workgroup.
These builtin functions must be encountered by all workitems in a
workgroup executing the kernel.
We use the generic type name gentype
to indicate the builtin data types
half
^{[1]}, int
, uint
, long
^{[2]}, ulong
, float
or double
^{[3]} as the type for the arguments.
Function 
Description 
int work_group_all(int predicate) 
Evaluates predicate for all workitems in the workgroup and returns a nonzero value if predicate evaluates to nonzero for all workitems in the workgroup. 
int work_group_any(int predicate) 
Evaluates predicate for all workitems in the workgroup and returns a nonzero value if predicate evaluates to nonzero for any workitems in the workgroup. 
gentype work_group_broadcast(gentype a, size_t local_id) 
Broadcast the value of a for workitem identified by local_id to all workitems in the workgroup. Behavior is undefined when the value of local_id is not equivalent for all workitems in the workgroup. Behavior is undefined when local_id is greater or equal to the workgroup size in the corresponding dimension. 
gentype work_group_reduce_<op>(gentype x) 
Return result of reduction operation specified by <op> for all values of x specified by workitems in a workgroup. 
gentype work_group_scan_exclusive_<op>(gentype x) 
Do an exclusive scan operation specified by <op> of all values specified by workitems in the workgroup. The scan results are returned for each workitem. The scan order is defined by increasing 1D linear global ID within the workgroup. 
gentype work_group_scan_inclusive_<op>(gentype x) 
Do an inclusive scan operation specified by <op> of all values specified by workitems in the workgroup. The scan results are returned for each workitem. The scan order is defined by increasing 1D linear global ID within the workgroup. 
The <op> in work_group_reduce_<op>, work_group_scan_exclusive_<op> and work_group_scan_inclusive_<op> defines the operator and can be add, min or max.
The inclusive scan operation takes a binary operator op with n (where n is the size of the workgroup) elements [a_{0}, a_{1}, … a_{n1}] and returns [a_{0}, (a_{0} op a_{1}), … (a_{0} op a_{1} op … op a_{n1})].
Consider the following example:
void foo(int *p)
{
...
int prefix_sum_val = work_group_scan_inclusive_add(
p[get_local_id(0)]);
}
For the example above, let’s assume that the workgroup size is 8 and p points to the following elements [3 1 7 0 4 1 6 3]. Workitem 0 calls work_group_scan_inclusive_add with 3 and returns 3. Workitem 1 calls work_group_scan_inclusive_add with 1 and returns 4. The full set of values returned by work_group_scan_inclusive_add for workitems 0 … 7 are [3 4 11 11 15 16 22 25].
The exclusive scan operation takes a binary associative operator op with
an identity I and n (where n is the size of the workgroup) elements [a_{0},
a_{1}, … a_{n1}] and returns [I, a_{0}, (a_{0} op a_{1}), … (a_{0} op
a_{1} op … op a_{n2})].
If op = add, the identity I is 0.
If op = min, the identity I is INT_MAX
, UINT_MAX
, LONG_MAX
,
ULONG_MAX
, for int
, uint
, long
, ulong
types and is +INF
for
floatingpoint types.
Similarly if op = max, the identity I is INT_MIN
, 0, LONG_MIN
, 0 and
INF
.
For the example above, the exclusive scan add operation on the ordered set
[3 1 7 0 4 1 6 3] would return [0 3 4 11 11 15 16 22].
The order of floatingpoint operations is not guaranteed for the
work_group_reduce_<op>, work_group_scan_inclusive_<op> and
work_group_scan_exclusive_<op> builtin functions that operate on 
Document Notes
For more information, see the OpenCL C Specification
This page is extracted from the OpenCL C Specification. Fixes and changes should be made to the Specification, not directly.