Update docs

Author: LeiWang1999

api/tilelang.html

@@ -841,6 +841,7 @@ <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_global"><code class="docutils literal notranslate"><span class="pre">sync_global()</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>

api/tilelang.language.builtin.html

@@ -691,6 +691,12 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.sync_global">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">sync_global</span></span><span class="sig-paren">(</span><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.sync_global" title="Permalink to this definition">#</a></dt>
+<dd><p>Synchronize all threads in a block.</p>
+</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>
@@ -857,6 +863,7 @@
 <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_global"><code class="docutils literal notranslate"><span class="pre">sync_global()</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>

api/tilelang.language.html

@@ -606,6 +606,7 @@ <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_global"><code class="docutils literal notranslate"><span class="pre">sync_global()</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>

api/tilelang.language.logical.html

@@ -412,7 +412,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">0x7f1488881ba0&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">0x7f7e21b9b730&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>
@@ -426,7 +426,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">0x7f1488881ba0&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">0x7f7e21b9b730&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

@@ -460,7 +460,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">0x7f1488881b70&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">0x7f7e21b9b7c0&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

@@ -3069,6 +3069,8 @@ <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_global">sync_global() (in module tilelang.language.builtin)</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>