Update docs

Author: LeiWang1999

api/tilelang.html

@@ -823,6 +823,8 @@ <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.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.sync_thread_partial"><code class="docutils literal notranslate"><span class="pre">sync_thread_partial()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.sync_threads"><code class="docutils literal notranslate"><span class="pre">sync_threads()</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

@@ -685,6 +685,30 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.sync_thread_partial">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">sync_thread_partial</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.sync_thread_partial" title="Permalink to this definition">#</a></dt>
+<dd><p>Synchronize threads within a warp.</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 synchronize</p>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A handle to the synchronization 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.sync_threads">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">sync_threads</span></span><span class="sig-paren">(</span><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.sync_threads" title="Permalink to this definition">#</a></dt>
+<dd><p>Synchronize all threads in a warp.</p>
+</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>
@@ -827,6 +851,8 @@
 <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.sync_thread_partial"><code class="docutils literal notranslate"><span class="pre">sync_thread_partial()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.sync_threads"><code class="docutils literal notranslate"><span class="pre">sync_threads()</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.html

@@ -599,6 +599,8 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li class="toctree-l2"><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-l2"><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-l2"><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-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.sync_thread_partial"><code class="docutils literal notranslate"><span class="pre">sync_thread_partial()</span></code></a></li>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.sync_threads"><code class="docutils literal notranslate"><span class="pre">sync_threads()</span></code></a></li>
 <li class="toctree-l2"><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-l2"><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-l2"><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.logical.html

@@ -406,7 +406,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">0x7f8d07ce1900&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">0x7f7be01db520&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>
@@ -420,7 +420,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">0x7f8d07ce1900&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">0x7f7be01db520&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.proxy.html

@@ -454,7 +454,7 @@
 <p>Buffer proxy class for constructing tir buffer.</p>
 <dl class="py method">
 <dt class="sig sig-object py" id="tilelang.language.proxy.BufferProxy.from_ptr">
-<span class="sig-name descname"><span class="pre">from_ptr</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">pointer_var</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Var</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">shape</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">tuple</span><span class="p"><span class="pre">[</span></span><span class="pre">tvm.ir.expr.PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="p"><span class="pre">...</span></span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</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">'float32'</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">&lt;tilelang.language.proxy.BufferProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f8d07ce15a0&gt;</span></span></span><a class="headerlink" href="#tilelang.language.proxy.BufferProxy.from_ptr" title="Permalink to this definition">#</a></dt>
+<span class="sig-name descname"><span class="pre">from_ptr</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">pointer_var</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Var</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">shape</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">tuple</span><span class="p"><span class="pre">[</span></span><span class="pre">tvm.ir.expr.PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="p"><span class="pre">...</span></span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</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">'float32'</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">&lt;tilelang.language.proxy.BufferProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f7be01db550&gt;</span></span></span><a class="headerlink" href="#tilelang.language.proxy.BufferProxy.from_ptr" title="Permalink to this definition">#</a></dt>
 <dd><p>Create a buffer from a pointer, shape, and data type.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>

genindex.html

@@ -3033,6 +3033,10 @@ <h2>S</h2>
         <li><a href="api/tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.cudaDeviceProp.surfaceAlignment">surfaceAlignment (tilelang.carver.arch.driver.cuda_driver.cudaDeviceProp attribute)</a>
 </li>
         <li><a href="api/tilelang.language.html#tilelang.language.symbolic">symbolic() (in module tilelang.language)</a>
+</li>
+        <li><a href="api/tilelang.language.builtin.html#tilelang.language.builtin.sync_thread_partial">sync_thread_partial() (in module tilelang.language.builtin)</a>
+</li>
+        <li><a href="api/tilelang.language.builtin.html#tilelang.language.builtin.sync_threads">sync_threads() (in module tilelang.language.builtin)</a>
 </li>
     </ul></td>
   </tr></table>