Return result of reduction operation for workitems in a workgroup.
gentype work_group_reduce_<op>(gentype x)
Description
Return result of reduction operation specified by <op> for all values of x
specified by workitems in a workgroup.
This builtin function 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
(if the cl_khr_fp16
extension is supported), int
, uint
, long
, ulong
, float
or double
(if double precision is supported) as the type for the arguments.
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 an identity I
and n
(where n
is the size of the workgroup) elements [a0, a1, … an1]
and returns [a0, (a0 op a1), … (a0 op a1 op … op an1)]
.
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
.
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 14 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 [a0, a1, … an1]
and returns [I, a0, (a0 op a1), … (a0 op a1 op … op an2)]
.
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 14 16 22].
Note

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 half , float and double data types.
The order of these floatingpoint operations is also nondeterministic for a given workgroup.
