Update docs

Author: github-actions[bot]

_sources/autoapi/tilelang/language/reduce/index.rst.txt

@@ -26,6 +26,11 @@ Functions
    tilelang.language.reduce.cumsum_fragment
    tilelang.language.reduce.cumsum
    tilelang.language.reduce.finalize_reducer
+   tilelang.language.reduce.warp_reduce_sum
+   tilelang.language.reduce.warp_reduce_max
+   tilelang.language.reduce.warp_reduce_min
+   tilelang.language.reduce.warp_reduce_bitand
+   tilelang.language.reduce.warp_reduce_bitor
 
 
 Module Contents
@@ -236,3 +241,78 @@ Module Contents
    :rtype: tir.Call
 
 
+.. py:function:: warp_reduce_sum(value)
+
+   Perform warp reduction sum on a register value.
+
+   This function reduces a value across all threads in a warp using shuffle operations.
+   Each thread provides a  register `value`, and after the reduction, all threads
+   will have the sum of all values across the warp.
+
+   :param value: The input register value to reduce
+   :type value: tir.PrimExpr
+
+   :returns: The reduced sum value (same on all threads in the warp)
+   :rtype: tir.PrimExpr
+
+
+.. py:function:: warp_reduce_max(value)
+
+   Perform warp reduction max on a register value.
+
+   This function reduces a value across all threads in a warp using shuffle operations.
+   Each thread provides a  register `value`, and after the reduction, all threads
+   will have the max of all values across the warp.
+
+   :param value: The input register value to reduce
+   :type value: tir.PrimExpr
+
+   :returns: The reduced max value (same on all threads in the warp)
+   :rtype: tir.PrimExpr
+
+
+.. py:function:: warp_reduce_min(value)
+
+   Perform warp reduction min on a register value.
+
+   This function reduces a value across all threads in a warp using shuffle operations.
+   Each thread provides a  register `value`, and after the reduction, all threads
+   will have the min of all values across the warp.
+
+   :param value: The input register value to reduce
+   :type value: tir.PrimExpr
+
+   :returns: The reduced min value (same on all threads in the warp)
+   :rtype: tir.PrimExpr
+
+
+.. py:function:: warp_reduce_bitand(value)
+
+   Perform warp reduction bitwise-and on a register value.
+
+   This function reduces a value across all threads in a warp using shuffle operations.
+   Each thread provides a  register `value`, and after the reduction, all threads
+   will have the bitwise-and of all values across the warp.
+
+   :param value: The input register value to reduce
+   :type value: tir.PrimExpr
+
+   :returns: The reduced bitwise-and value (same on all threads in the warp)
+   :rtype: tir.PrimExpr
+
+
+.. py:function:: warp_reduce_bitor(value)
+
+   Perform warp reduction bitwise-or on a register value.
+
+   This function reduces a value across all threads in a warp using shuffle operations.
+   Each thread provides a  register `value`, and after the reduction, all threads
+   will have the bitwise-or of all values across the warp.
+
+   :param value: The input register value to reduce
+   :type value: tir.PrimExpr
+
+   :returns: The reduced bitwise-or value (same on all threads in the warp)
+   :rtype: tir.PrimExpr
+
+

autoapi/tilelang/language/reduce/index.html

@@ -541,6 +541,21 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 <tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.reduce.finalize_reducer" title="tilelang.language.reduce.finalize_reducer"><code class="xref py py-obj docutils literal notranslate"><span class="pre">finalize_reducer</span></code></a>(reducer)</p></td>
 <td><p>Finalize a reducer buffer by emitting the <cite>tl.finalize_reducer</cite> intrinsic.</p></td>
 </tr>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_sum" title="tilelang.language.reduce.warp_reduce_sum"><code class="xref py py-obj docutils literal notranslate"><span class="pre">warp_reduce_sum</span></code></a>(value)</p></td>
+<td><p>Perform warp reduction sum on a register value.</p></td>
+</tr>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_max" title="tilelang.language.reduce.warp_reduce_max"><code class="xref py py-obj docutils literal notranslate"><span class="pre">warp_reduce_max</span></code></a>(value)</p></td>
+<td><p>Perform warp reduction max on a register value.</p></td>
+</tr>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_min" title="tilelang.language.reduce.warp_reduce_min"><code class="xref py py-obj docutils literal notranslate"><span class="pre">warp_reduce_min</span></code></a>(value)</p></td>
+<td><p>Perform warp reduction min on a register value.</p></td>
+</tr>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_bitand" title="tilelang.language.reduce.warp_reduce_bitand"><code class="xref py py-obj docutils literal notranslate"><span class="pre">warp_reduce_bitand</span></code></a>(value)</p></td>
+<td><p>Perform warp reduction bitwise-and on a register value.</p></td>
+</tr>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_bitor" title="tilelang.language.reduce.warp_reduce_bitor"><code class="xref py py-obj docutils literal notranslate"><span class="pre">warp_reduce_bitor</span></code></a>(value)</p></td>
+<td><p>Perform warp reduction bitwise-or on a register value.</p></td>
+</tr>
 </tbody>
 </table>
 </div>
@@ -853,6 +868,106 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.reduce.warp_reduce_sum">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.reduce.</span></span><span class="sig-name descname"><span class="pre">warp_reduce_sum</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.reduce.warp_reduce_sum" title="Link to this definition">¶</a></dt>
+<dd><p>Perform warp reduction sum on a register value.</p>
+<p>This function reduces a value across all threads in a warp using shuffle operations.
+Each thread provides a  register <cite>value</cite>, and after the reduction, all threads
+will have the sum of all values across the warp.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> (<em>tir.PrimExpr</em>) – The input register value to reduce</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>The reduced sum value (same on all threads in the warp)</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.reduce.warp_reduce_max">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.reduce.</span></span><span class="sig-name descname"><span class="pre">warp_reduce_max</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.reduce.warp_reduce_max" title="Link to this definition">¶</a></dt>
+<dd><p>Perform warp reduction max on a register value.</p>
+<p>This function reduces a value across all threads in a warp using shuffle operations.
+Each thread provides a  register <cite>value</cite>, and after the reduction, all threads
+will have the max of all values across the warp.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> (<em>tir.PrimExpr</em>) – The input register value to reduce</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>The reduced max value (same on all threads in the warp)</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.reduce.warp_reduce_min">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.reduce.</span></span><span class="sig-name descname"><span class="pre">warp_reduce_min</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.reduce.warp_reduce_min" title="Link to this definition">¶</a></dt>
+<dd><p>Perform warp reduction min on a register value.</p>
+<p>This function reduces a value across all threads in a warp using shuffle operations.
+Each thread provides a  register <cite>value</cite>, and after the reduction, all threads
+will have the min of all values across the warp.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> (<em>tir.PrimExpr</em>) – The input register value to reduce</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>The reduced min value (same on all threads in the warp)</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.reduce.warp_reduce_bitand">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.reduce.</span></span><span class="sig-name descname"><span class="pre">warp_reduce_bitand</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.reduce.warp_reduce_bitand" title="Link to this definition">¶</a></dt>
+<dd><p>Perform warp reduction bitwise-and on a register value.</p>
+<p>This function reduces a value across all threads in a warp using shuffle operations.
+Each thread provides a  register <cite>value</cite>, and after the reduction, all threads
+will have the bitwise-and of all values across the warp.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> (<em>tir.PrimExpr</em>) – The input register value to reduce</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>The reduced bitwise-and value (same on all threads in the warp)</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.reduce.warp_reduce_bitor">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.reduce.</span></span><span class="sig-name descname"><span class="pre">warp_reduce_bitor</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.reduce.warp_reduce_bitor" title="Link to this definition">¶</a></dt>
+<dd><p>Perform warp reduction bitwise-or on a register value.</p>
+<p>This function reduces a value across all threads in a warp using shuffle operations.
+Each thread provides a  register <cite>value</cite>, and after the reduction, all threads
+will have the bitwise-or of all values across the warp.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> (<em>tir.PrimExpr</em>) – The input register value to reduce</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>The reduced bitwise-or value (same on all threads in the warp)</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
 </section>
 </section>
 
@@ -926,6 +1041,11 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <li><a class="reference internal" href="#tilelang.language.reduce.cumsum_fragment"><code class="docutils literal notranslate"><span class="pre">cumsum_fragment()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.reduce.cumsum"><code class="docutils literal notranslate"><span class="pre">cumsum()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.reduce.finalize_reducer"><code class="docutils literal notranslate"><span class="pre">finalize_reducer()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_sum"><code class="docutils literal notranslate"><span class="pre">warp_reduce_sum()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_max"><code class="docutils literal notranslate"><span class="pre">warp_reduce_max()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_min"><code class="docutils literal notranslate"><span class="pre">warp_reduce_min()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_bitand"><code class="docutils literal notranslate"><span class="pre">warp_reduce_bitand()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.reduce.warp_reduce_bitor"><code class="docutils literal notranslate"><span class="pre">warp_reduce_bitor()</span></code></a></li>
 </ul>
 </li>
 </ul>

genindex.html

@@ -7424,6 +7424,16 @@ <h2>W</h2>
 </li>
         </ul></li>
         <li><a href="autoapi/tilelang/intrinsics/mfma_macro_generator/index.html#tilelang.intrinsics.mfma_macro_generator.MatrixCoreIntrinEmitter.warp_cols">warp_cols (tilelang.intrinsics.mfma_macro_generator.MatrixCoreIntrinEmitter attribute)</a>
+</li>
+        <li><a href="autoapi/tilelang/language/reduce/index.html#tilelang.language.reduce.warp_reduce_bitand">warp_reduce_bitand() (in module tilelang.language.reduce)</a>
+</li>
+        <li><a href="autoapi/tilelang/language/reduce/index.html#tilelang.language.reduce.warp_reduce_bitor">warp_reduce_bitor() (in module tilelang.language.reduce)</a>
+</li>
+        <li><a href="autoapi/tilelang/language/reduce/index.html#tilelang.language.reduce.warp_reduce_max">warp_reduce_max() (in module tilelang.language.reduce)</a>
+</li>
+        <li><a href="autoapi/tilelang/language/reduce/index.html#tilelang.language.reduce.warp_reduce_min">warp_reduce_min() (in module tilelang.language.reduce)</a>
+</li>
+        <li><a href="autoapi/tilelang/language/reduce/index.html#tilelang.language.reduce.warp_reduce_sum">warp_reduce_sum() (in module tilelang.language.reduce)</a>
 </li>
         <li><a href="autoapi/tilelang/intrinsics/mfma_macro_generator/index.html#tilelang.intrinsics.mfma_macro_generator.MatrixCoreIntrinEmitter.warp_row_tiles">warp_row_tiles (tilelang.intrinsics.mfma_macro_generator.MatrixCoreIntrinEmitter attribute)</a>
 
@@ -7457,14 +7467,14 @@ <h2>W</h2>
 </li>
         <li><a href="autoapi/tilelang/language/builtin/index.html#tilelang.language.builtin.warpgroup_commit_batch">warpgroup_commit_batch() (in module tilelang.language.builtin)</a>
 </li>
+    </ul></td>
+    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="autoapi/tilelang/language/builtin/index.html#tilelang.language.builtin.warpgroup_fence_operand">warpgroup_fence_operand() (in module tilelang.language.builtin)</a>
 </li>
         <li><a href="autoapi/tilelang/language/builtin/index.html#tilelang.language.builtin.warpgroup_wait">warpgroup_wait() (in module tilelang.language.builtin)</a>
 </li>
         <li><a href="autoapi/tilelang/language/warpgroup/index.html#tilelang.language.warpgroup.WarpSpecialize">WarpSpecialize() (in module tilelang.language.warpgroup)</a>
 </li>
-    </ul></td>
-    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="autoapi/tilelang/transform/index.html#tilelang.transform.WarpSpecialized">WarpSpecialized() (in module tilelang.transform)</a>
 </li>
         <li><a href="autoapi/tilelang/transform/index.html#tilelang.transform.WarpSpecializedPipeline">WarpSpecializedPipeline() (in module tilelang.transform)</a>