Update docs

Author: LeiWang1999

api/tilelang.html

@@ -807,6 +807,8 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 </ul>
 </li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.builtin.html">tilelang.language.builtin module</a><ul>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.barrier_arrive"><code class="docutils literal notranslate"><span class="pre">barrier_arrive()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.barrier_wait"><code class="docutils literal notranslate"><span class="pre">barrier_wait()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_list_of_mbarrier"><code class="docutils literal notranslate"><span class="pre">create_list_of_mbarrier()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_tma_descriptor"><code class="docutils literal notranslate"><span class="pre">create_tma_descriptor()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.dec_max_nreg"><code class="docutils literal notranslate"><span class="pre">dec_max_nreg()</span></code></a></li>
@@ -818,6 +820,9 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_wait_parity"><code class="docutils literal notranslate"><span class="pre">mbarrier_wait_parity()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.no_set_max_nreg"><code class="docutils literal notranslate"><span class="pre">no_set_max_nreg()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.set_max_nreg"><code class="docutils literal notranslate"><span class="pre">set_max_nreg()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.shfl_down"><code class="docutils literal notranslate"><span class="pre">shfl_down()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.shfl_up"><code class="docutils literal notranslate"><span class="pre">shfl_up()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.shfl_xor"><code class="docutils literal notranslate"><span class="pre">shfl_xor()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.tma_load"><code class="docutils literal notranslate"><span class="pre">tma_load()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.tma_store_arrive"><code class="docutils literal notranslate"><span class="pre">tma_store_arrive()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.tma_store_wait"><code class="docutils literal notranslate"><span class="pre">tma_store_wait()</span></code></a></li>

api/tilelang.language.builtin.html

@@ -404,6 +404,41 @@
           <section id="module-tilelang.language.builtin">
 <span id="tilelang-language-builtin-module"></span><h1>tilelang.language.builtin module<a class="headerlink" href="#module-tilelang.language.builtin" title="Permalink to this heading">#</a></h1>
 <p>The language interface for tl programs.</p>
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.barrier_arrive">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">barrier_arrive</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">barrier_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.barrier_arrive" title="Permalink to this definition">#</a></dt>
+<dd><p>Arrive at a memory barrier.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>barrier_id</strong> – Optional[int, PrimExpr]
+The memory barrier to arrive at</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.barrier_wait">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">barrier_wait</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">barrier_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">parity</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Optional</span><span class="p"><span class="pre">[</span></span><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Var</span><span class="p"><span class="pre">]</span></span><span class="p"><span class="pre">]</span></span></span><span class="w"> </span><span class="o"><span class="pre">=</span></span><span class="w"> </span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.barrier_wait" title="Permalink to this definition">#</a></dt>
+<dd><p>Wait for a memory barrier to complete.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><ul class="simple">
+<li><p><strong>barrier_id</strong> – Optional[int, PrimExpr]
+The memory barrier to wait on</p></li>
+<li><p><strong>parity</strong> – Optional[int, Var]
+The parity value to wait for</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A handle to the barrier wait operation</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.Call</p>
+</dd>
+</dl>
+<p>Current implementation is a sugar syntax for mbarrier_wait_parity, as we only support parity 0 and 1.</p>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.create_list_of_mbarrier">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">create_list_of_mbarrier</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="o"><span class="pre">*</span></span><span class="n"><span class="pre">args</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Any</span></span></em><span class="sig-paren">)</span> <span class="sig-return"><span class="sig-return-icon">&#x2192;</span> <span class="sig-return-typehint"><span class="pre">Call</span></span></span><a class="headerlink" href="#tilelang.language.builtin.create_list_of_mbarrier" title="Permalink to this definition">#</a></dt>
@@ -604,6 +639,52 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.shfl_down">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">shfl_down</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">offset</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.shfl_down" title="Permalink to this definition">#</a></dt>
+<dd><p>Perform a shuffle operation with down offset.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> – Optional[int, PrimExpr]
+The value to shuffle</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.shfl_up">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">shfl_up</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">offset</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.shfl_up" title="Permalink to this definition">#</a></dt>
+<dd><p>Perform a shuffle operation with up offset.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>value</strong> – Optional[int, PrimExpr]
+The value to shuffle</p>
+</dd>
+</dl>
+</dd></dl>
+
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.shfl_xor">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">shfl_xor</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">value</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">offset</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Union</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">Call</span><span class="p"><span class="pre">]</span></span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.shfl_xor" title="Permalink to this definition">#</a></dt>
+<dd><p>Perform a shuffle operation with XOR offset.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><ul class="simple">
+<li><p><strong>value</strong> – Optional[int, PrimExpr]
+The value to shuffle</p></li>
+<li><p><strong>offset</strong> – Optional[int, PrimExpr]
+The offset for the shuffle operation</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A handle to the shuffle operation</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tir.Call</p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.tma_load">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">tma_load</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="o"><span class="pre">*</span></span><span class="n"><span class="pre">args</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.tma_load" title="Permalink to this definition">#</a></dt>
@@ -730,6 +811,8 @@
           <div class="toc-tree">
             <ul>
 <li><a class="reference internal" href="#">tilelang.language.builtin module</a><ul>
+<li><a class="reference internal" href="#tilelang.language.builtin.barrier_arrive"><code class="docutils literal notranslate"><span class="pre">barrier_arrive()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.barrier_wait"><code class="docutils literal notranslate"><span class="pre">barrier_wait()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.create_list_of_mbarrier"><code class="docutils literal notranslate"><span class="pre">create_list_of_mbarrier()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.create_tma_descriptor"><code class="docutils literal notranslate"><span class="pre">create_tma_descriptor()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.dec_max_nreg"><code class="docutils literal notranslate"><span class="pre">dec_max_nreg()</span></code></a></li>
@@ -741,6 +824,9 @@
 <li><a class="reference internal" href="#tilelang.language.builtin.mbarrier_wait_parity"><code class="docutils literal notranslate"><span class="pre">mbarrier_wait_parity()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.no_set_max_nreg"><code class="docutils literal notranslate"><span class="pre">no_set_max_nreg()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.set_max_nreg"><code class="docutils literal notranslate"><span class="pre">set_max_nreg()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.shfl_down"><code class="docutils literal notranslate"><span class="pre">shfl_down()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.shfl_up"><code class="docutils literal notranslate"><span class="pre">shfl_up()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.shfl_xor"><code class="docutils literal notranslate"><span class="pre">shfl_xor()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.tma_load"><code class="docutils literal notranslate"><span class="pre">tma_load()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.tma_store_arrive"><code class="docutils literal notranslate"><span class="pre">tma_store_arrive()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.tma_store_wait"><code class="docutils literal notranslate"><span class="pre">tma_store_wait()</span></code></a></li>

api/tilelang.language.copy.html

@@ -427,7 +427,7 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.copy.buffer_region_to_tile_region">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.copy.</span></span><span class="sig-name descname"><span class="pre">buffer_region_to_tile_region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">buffer_region</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">BufferRegion</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.copy.buffer_region_to_tile_region" title="Permalink to this definition">#</a></dt>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.copy.</span></span><span class="sig-name descname"><span class="pre">buffer_region_to_tile_region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">buffer_region</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">BufferRegion</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">extents</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">List</span><span class="p"><span class="pre">[</span></span><span class="pre">PrimExpr</span><span class="p"><span class="pre">]</span></span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.copy.buffer_region_to_tile_region" title="Permalink to this definition">#</a></dt>
 <dd><p>Convert a buffer region to a tile region descriptor.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>