Do an inclusive prefixsum operation of all values in workitems in the workgroup
gentype work_group_scan_inclusive_<op>(gentype x)
Description
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.
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.
