Do an exclusive prefix-sum operation of all values in work-items in the work-group

```
gentype
```
| gentype x`)` |

Do an exclusive scan operation specified
by *<op>* of all values specified by work-items
in the work-group. The scan results are
returned for each work-item.

The scan order is defined by increasing 1D linear global ID within the work-group.

This built-in function must be encountered by all work-items in a work-group executing the kernel. We use the generic type name gentype to indicate the built-in 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_`

and
*<op>*`work_group_scan_inclusive_`

defines the operator and can be *<op>*`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 work-group) elements
`[a0, a1, ... an-1]`

and returns
`[a0, (a0 `

.
If <*op* a1), ... (a0 *op* a1
*op* ... *op* an-1)]`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 floating-point 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 work-group
size is 8 and p points to the following
elements [3 1 7 0 4 1 6 3]. Work-item 0 calls
`work_group_scan_inclusive_add`

with 3 and returns 3. Work-item 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 work-items 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 work-group)
elements `[a0, a1, ... an-1]`

and returns ```
[I, a0,
(a0
```

. 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].
*op* a1), ... (a0 *op* a1
*op* ... *op* an-2)]

NOTE: The order of floating-point operations is not guaranteed for the
`work_group_reduce_`

,
*<op>*`work_group_scan_inclusive_`

and
*<op>*`work_group_scan_exclusive_`

built-in functions that operate on half, float and
double data types.
The order of these floating-point operations is also non-deterministic
for a given workgroup.
*<op>*

Copyright © 2007-2013 The Khronos Group Inc.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and/or associated documentation files (the
"Materials"), to deal in the Materials without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Materials, and to
permit persons to whom the Materials are furnished to do so, subject to
the condition that this copyright notice and permission notice shall be included
in all copies or substantial portions of the Materials.