blob: 099b782ca996537b6c2886bf397a2fc799b9ac01 [file] [log] [blame]
<?xml version="1.0" encoding="UTF-8"?>
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "xhtml1-transitional.dtd">
<!-- saved from url=(0013)about:internet -->
<?xml-stylesheet type="text/xsl" href="mathml.xsl"?><html xmlns="http://www.w3.org/1999/xhtml" xmlns:pref="http://www.w3.org/2002/Math/preference" xmlns:xlink="http://www.w3.org/1999/xlink" pref:renderer="mathplayer-dl">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8" />
<style xmlns="" type="text/css">
/* This style sets a margin around the entire page */
html, body {
margin: 10px;
}
p {
font: normal 16px verdana, sans-serif;
margin: 0;
padding-bottom:12px;
}
h1 {
font: bold 25px verdana, sans-serif;
margin-top: 0;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
h2 {
font: bold 19px verdana, sans-serif;
margin-top: 28px;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
h3 {
font: bold 19px verdana, sans-serif !important;
margin-top: 28px;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
li {
font: normal 16px verdana, sans-serif;
margin-top: 0;
margin-bottom: 18px;
padding-top: 0;
padding-bottom: 0;
}
.pdparam {
font: italic 16px verdana, sans-serif;
}
.term {
font: italic 16px verdana, sans-serif;
font-weight: normal;
}
.type {
font: normal 16px verdana, sans-serif !important;
}
.parameter {
font-style: italic;
}
a:link, a:visited {
color: blue;
text-decoration: none;
font: normal 16px;
}
a:hover {
background-color: #FFFF99;
font: normal 16px;
}
div.funcsynopsis {
text-align: left;
background-color: #e6e6e6;
font: normal 16px verdana, sans-serif;
padding-top: 10px;
padding-bottom: 10px;
}
div.funcsynopsis table {
border-collapse: separate;
font: normal 16px verdana, sans-serif;
}
div.funcsynopsis td {
background-color: #e6e6e6;
border: 0 solid #000;
padding: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 {
font-family: verdana, sans-serif;
font-size: 16px;
}
code.constant {
font: normal 16px courier new, monospace !important;
}
span.errorname {
font: normal 16px verdana, sans-serif !important;
}
code.function {
font: bold 16px verdana, sans-serif !important;
}
b.fsfunc {
font: bold 16px verdana, sans-serif !important;
}
code.varname {
font: italic 16px verdana, sans-serif;
}
code.replaceable {
font: italic 16px courier new, monospace;
}
code.funcdef {
font: normal 16px verdana, sans-serif !important;
}
.citerefentry {
font: normal 16px verdana, sans-serif !important;
}
.parameter {
font-style: italic;
}
code.fsfunc {
font: normal 16px verdana, sans-serif !important;
}
/* PARAMETER: This style controls spacing between the terms in Parameter section */
dt {
margin-top: 15px;
}
/* TABLES: These styles apply to all tables OTHER than the Synopsis and Example tables */
div.refsect1 table {
width: 100%;
margin-top: 10px;
background-color: #FFF;
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 th {
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: bold 16px verdana, sans-serif;
}
div.refsect1 td {
background-color: #FFF;
padding: 5px;
vertical-align: text-top;
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 p{
font: normal 16px verdana, sans-serif;
margin-top: 8px;
margin-bottom: 8px;
padding-top: 0;
padding-bottom: 0;
}
/* EXAMPLE: These styles apply only to the Example section */
div.refsect2 {
font: normal 16px courier new, monospace !important;
}
div.refsect2 table {
margin-top: 0;
background-color: #e6e6e6;
width: 100%;
border: 0 solid #000;
padding: 2px;
font: normal 16px courier new, monospace !important;
}
div.refsect2 td {
background-color: #e6e6e6;
font: normal 16px courier new, monospace !important;
white-space:pre;
}
/* COPYRIGHT: This style formats the text of the copyright statement at the bottom of the page */
div.refsect3 {
font: normal 11px verdana, sans-serif;
margin-top: 50px;
margin-bottom: 20px;
padding-top: 0;
padding-bottom: 0;
}
</style>
<title>work_group_reduce</title>
<meta name="generator" content="DocBook XSL Stylesheets V1.79.1" />
<meta name="keywords" content="work_group_reduce" />
</head>
<body>
<div class="refentry">
<a id="id-1"></a>
<div class="titlepage"></div>
<div xmlns="" class="refnamediv">
<a xmlns="http://www.w3.org/1999/xhtml" id="work_group_reduce"></a>
<h1>work_group_reduce_&lt;op&gt;</h1>
<p>
Return result of reduction operation for work-items in a work-group.
</p>
</div>
<div class="refsynopsisdiv">
<h2></h2>
<div class="funcsynopsis">
<table xmlns="" border="0" summary="Function synopsis" cellspacing="0" cellpadding="0">
<tr valign="bottom">
<td>
<code xmlns="http://www.w3.org/1999/xhtml" class="funcdef">
gentype
<strong class="fsfunc">
work_group_reduce_&lt;op&gt;
</strong>
(</code>
<td>gentype  <var xmlns="http://www.w3.org/1999/xhtml" class="pdparam">x</var><code>)</code></td>
</td>
</tr>
</table>
</div>
</div>
<div class="refsect1">
<a id="description"></a>
<h2>Description</h2>
<p>
Return result of reduction operation specified
by <span class="emphasis"><em>&lt;op&gt;</em></span> for all values of <code class="varname">x</code>
specified by work-items in a work-group.
</p>
<h4><a id="id-1.5.3"></a>General information about work-group functions</h4>
<p>
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 <span class="type">half</span> (if the
<a class="citerefentry" href="cl_khr_fp16.html"><span class="citerefentry"><span class="refentrytitle">cl_khr_fp16</span></span></a>
extension is supported),
<span class="type">int</span>, <span class="type">uint</span>, <span class="type">long</span>,
<span class="type">ulong</span>, <span class="type">float</span> or <span class="type">double</span>
(if double precision is supported) as the type
for the arguments.
</p>
<p>
The <span class="emphasis"><em>&lt;op&gt;</em></span> in
<code class="function">work_group_reduce_<span class="emphasis"><em>&lt;op&gt;</em></span></code>,
<code class="function">work_group_scan_exclusive_<span class="emphasis"><em>&lt;op&gt;</em></span></code> and
<code class="function">work_group_scan_inclusive_<span class="emphasis"><em>&lt;op&gt;</em></span></code>
defines the operator and can be <code class="code">add</code>, <code class="code">min</code> or <code class="code">max</code>.
</p>
<p>
The inclusive scan operation takes a binary operator <code class="code">op</code>
with an identity <code class="code">I</code> and <code class="code">n</code> (where
<code class="code">n</code> is the size of the work-group) elements
<code class="code">[a0, a1, ... an-1]</code> and returns
<code class="code">[a0, (a0 <span class="emphasis"><em>op</em></span> a1), ... (a0 <span class="emphasis"><em>op</em></span> a1
<span class="emphasis"><em>op</em></span> ... <span class="emphasis"><em>op</em></span> an-1)]</code>.
If &lt;<code class="code">op</code>&gt; = <code class="code">add</code>, the identity
<code class="code">I</code> is 0. If &lt;<code class="code">op</code>&gt; = <code class="code">min</code>,
the identity <code class="code">I</code> is <code class="constant">INT_MAX</code>,
<code class="constant">UINT_MAX</code>, <code class="constant">LONG_MAX</code>,
<code class="constant">ULONG_MAX</code>, for <span class="type">int</span>,
<span class="type">uint</span>, <span class="type">long</span>, <span class="type">ulong</span> types
and is <code class="code">+INF</code> for floating-point types. Similarly
if &lt;<code class="code">op</code>&gt; = max, the identity <code class="code">I</code> is
<code class="constant">INT_MIN</code>, 0, <code class="constant">LONG_MIN</code>,
0 and <code class="code">-INF</code>.
</p>
<p>
Consider the following example:
</p>
<p>
</p>
<div class="literallayout">
<p>
<code class="code"><br />
void foo(int *p)<br />
{<br />
    ... <br />
    int prefix_sum_val = work_group_scan_inclusive_add(<br />
              p[get_local_id(0)]);<br />
}</code>
</p>
</div>
<p>
</p>
<p>
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
<code class="function">work_group_scan_inclusive_add</code>
with 3 and returns 3. Work-item 1 calls
<code class="function">work_group_scan_inclusive_add</code> with 1 and
returns 4. The full set of values returned by
<code class="function">work_group_scan_inclusive_add</code>
for work-items 0 ... 7 are [3 4 11 11 14 16 22 25].
</p>
<p>
The exclusive scan operation takes a binary
associative operator <code class="varname">op</code> with an
identity <code class="code">I</code> and <code class="code">n</code>
(where <code class="code">n</code> is the size of the work-group)
elements <code class="code">[a0, a1, ... an-1]</code> and returns <code class="code">[I, a0,
(a0 <span class="emphasis"><em>op</em></span> a1), ... (a0 <span class="emphasis"><em>op</em></span> a1
<span class="emphasis"><em>op</em></span> ... <span class="emphasis"><em>op</em></span> an-2)]</code>. 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].
</p>
<p>
NOTE: The order of floating-point operations is not guaranteed for the
<code class="function">work_group_reduce_<span class="emphasis"><em>&lt;op&gt;</em></span></code>,
<code class="function">work_group_scan_inclusive_<span class="emphasis"><em>&lt;op&gt;</em></span></code> and
<code class="function">work_group_scan_exclusive_<span class="emphasis"><em>&lt;op&gt;</em></span></code>
built-in functions that operate on <span class="type">half</span>, <span class="type">float</span> and
<span class="type">double</span> data types.
The order of these floating-point operations is also non-deterministic
for a given workgroup.
</p>
</div>
<div class="refsect1">
<a id="specification"></a>
<h2>Specification</h2>
<p>
<img src="pdficon_small1.gif" />
<a href="https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#namedest=work_group_reduce" target="OpenCL Spec">OpenCL Specification</a>
</p>
</div>
<div class="refsect1">
<a id="seealso"></a>
<h2>Also see</h2>
<p>
<a class="citerefentry" href="work_group_scan_exclusive.html"><span class="citerefentry"><span class="refentrytitle">work_group_scan_exclusive</span></span></a>,
<a class="citerefentry" href="work_group_scan_inclusive.html"><span class="citerefentry"><span class="refentrytitle">work_group_scan_inclusive</span></span></a>
</p>
</div>
<div xmlns="" class="refsect3" lang="en" xml:lang="en"><a xmlns="http://www.w3.org/1999/xhtml" id="Copyright"></a><h4 xmlns="http://www.w3.org/1999/xhtml"></h4><img xmlns="http://www.w3.org/1999/xhtml" src="KhronosLogo.jpg" /><p xmlns="http://www.w3.org/1999/xhtml"></p>Copyright © 2007-2015 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.
</div>
</div>
</body>
</html>