Update docs

Author: LeiWang1999

api/tilelang.html

@@ -781,6 +781,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 </li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.html#submodules">Submodules</a><ul>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.allocate.html">tilelang.language.allocate module</a><ul>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_barrier"><code class="docutils literal notranslate"><span class="pre">alloc_barrier()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_fragment"><code class="docutils literal notranslate"><span class="pre">alloc_fragment()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_local"><code class="docutils literal notranslate"><span class="pre">alloc_local()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_shared"><code class="docutils literal notranslate"><span class="pre">alloc_shared()</span></code></a></li>
@@ -1074,6 +1075,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerDeviceStorageAccessInfo"><code class="docutils literal notranslate"><span class="pre">LowerDeviceStorageAccessInfo()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerHopperIntrin"><code class="docutils literal notranslate"><span class="pre">LowerHopperIntrin()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerL2Persistent"><code class="docutils literal notranslate"><span class="pre">LowerL2Persistent()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerSharedBarrier"><code class="docutils literal notranslate"><span class="pre">LowerSharedBarrier()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerTileOp"><code class="docutils literal notranslate"><span class="pre">LowerTileOp()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.MakePackedAPI"><code class="docutils literal notranslate"><span class="pre">MakePackedAPI()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.MergeIfStmt"><code class="docutils literal notranslate"><span class="pre">MergeIfStmt()</span></code></a></li>

api/tilelang.language.allocate.html

@@ -427,6 +427,23 @@
 </dl>
 <p>Each function takes shape and dtype parameters and returns a TVM buffer object
 with the appropriate memory scope.</p>
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.allocate.alloc_barrier">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_barrier</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">arrive_count</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_barrier" title="Permalink to this definition">#</a></dt>
+<dd><p>Allocate a barrier buffer.</p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>arrive_count</strong> (<em>int</em>) – The number of threads that need to arrive at the barrier</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A TVM buffer object allocated as a barrier</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>T.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.allocate.alloc_fragment">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_fragment</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">scope</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">'local.fragment'</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_fragment" title="Permalink to this definition">#</a></dt>
@@ -568,6 +585,7 @@
           <div class="toc-tree">
             <ul>
 <li><a class="reference internal" href="#">tilelang.language.allocate module</a><ul>
+<li><a class="reference internal" href="#tilelang.language.allocate.alloc_barrier"><code class="docutils literal notranslate"><span class="pre">alloc_barrier()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.allocate.alloc_fragment"><code class="docutils literal notranslate"><span class="pre">alloc_fragment()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.allocate.alloc_local"><code class="docutils literal notranslate"><span class="pre">alloc_local()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.allocate.alloc_shared"><code class="docutils literal notranslate"><span class="pre">alloc_shared()</span></code></a></li>

api/tilelang.language.builtin.html

@@ -783,11 +783,12 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.wait_wgmma">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">wait_wgmma</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.wait_wgmma" title="Permalink to this definition">#</a></dt>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">wait_wgmma</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.wait_wgmma" title="Permalink to this definition">#</a></dt>
 <dd><p>Wait for WGMMA (Warp Group Matrix Multiply-Accumulate) operations to complete.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
-<dd class="field-odd"><p><strong>*args</strong> – Variable arguments specifying which operations to wait for</p>
+<dd class="field-odd"><p><strong>id</strong> – int
+The id of the WGMMA operation to wait for</p>
 </dd>
 <dt class="field-even">Returns<span class="colon">:</span></dt>
 <dd class="field-even"><p>A handle to the WGMMA wait operation</p>

api/tilelang.language.copy.html

@@ -502,7 +502,7 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.copy.copy">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.copy.</span></span><span class="sig-name descname"><span class="pre">copy</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">src</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">Buffer</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferLoad</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferRegion</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dst</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">Buffer</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferLoad</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">coalesced_width</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">int</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.copy.copy" 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">copy</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">src</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">Buffer</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferLoad</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferRegion</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dst</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">Buffer</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">BufferLoad</span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">coalesced_width</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">int</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>, <em class="sig-param"><span class="n"><span class="pre">disable_tma</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">bool</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">False</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.copy.copy" title="Permalink to this definition">#</a></dt>
 <dd><p>Copy data between memory regions.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>

api/tilelang.language.html

@@ -597,6 +597,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <div class="toctree-wrapper compound">
 <ul>
 <li class="toctree-l1"><a class="reference internal" href="tilelang.language.allocate.html">tilelang.language.allocate module</a><ul>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_barrier"><code class="docutils literal notranslate"><span class="pre">alloc_barrier()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_fragment"><code class="docutils literal notranslate"><span class="pre">alloc_fragment()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_local"><code class="docutils literal notranslate"><span class="pre">alloc_local()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.allocate.html#tilelang.language.allocate.alloc_shared"><code class="docutils literal notranslate"><span class="pre">alloc_shared()</span></code></a></li>

api/tilelang.language.logical.html

@@ -415,7 +415,7 @@
 <p>The language interface for tl programs.</p>
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.logical.all_of">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">all_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7fea3e8cad70&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.all_of" title="Permalink to this definition">#</a></dt>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">all_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f33837c2020&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.all_of" title="Permalink to this definition">#</a></dt>
 <dd><p>Check if all elements in the buffer are true.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
@@ -429,7 +429,7 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.logical.any_of">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">any_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7fea3e8cad70&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.any_of" title="Permalink to this definition">#</a></dt>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">any_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f33837c2020&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.any_of" title="Permalink to this definition">#</a></dt>
 <dd><p>Check if any element in the buffer is true.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>

api/tilelang.language.print.html

@@ -416,7 +416,7 @@
 It includes functionality to print variables, print values in buffers, and conditionally execute debug prints.</p>
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.print.print">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.print.</span></span><span class="sig-name descname"><span class="pre">print</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">obj</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Any</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">msg</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</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">''</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">PrimExpr</span></span></span><a class="headerlink" href="#tilelang.language.print.print" title="Permalink to this definition">#</a></dt>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.print.</span></span><span class="sig-name descname"><span class="pre">print</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">obj</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Any</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">msg</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</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">''</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">warp_group_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</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">0</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">warp_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</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">0</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">PrimExpr</span></span></span><a class="headerlink" href="#tilelang.language.print.print" title="Permalink to this definition">#</a></dt>
 <dd><p>A generic print function that handles both TIR buffers and primitive expressions.</p>
 <ul class="simple">
 <li><p>If the input is a TIR buffer, it prints its values, but only on the first thread (tx=0, ty=0, tz=0).</p></li>
@@ -427,6 +427,9 @@
 <dd class="field-odd"><ul class="simple">
 <li><p><strong>obj</strong> (<em>Any</em>) – The object to print. It can be either a tir.Buffer or tir.PrimExpr.</p></li>
 <li><p><strong>msg</strong> (<em>str</em>) – An optional message to include in the print statement.</p></li>
+<li><p><strong>warp_group_id</strong> (<em>int</em>) – The warp group id to print.</p></li>
+<li><p><strong>warp_id</strong> (<em>int</em>) – The warp id to print.</p></li>
+<li><p><strong>warp_id.</strong> (<em>print thread will be warp_group_id * warp_group_size +</em>) – </p></li>
 </ul>
 </dd>
 <dt class="field-even">Returns<span class="colon">:</span></dt>