Update docs

Author: LeiWang1999

api/tilelang.html

@@ -1044,6 +1044,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.InjectFenceProxy"><code class="docutils literal notranslate"><span class="pre">InjectFenceProxy()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.InjectPTXAsyncCopy"><code class="docutils literal notranslate"><span class="pre">InjectPTXAsyncCopy()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.InjectSoftwarePipeline"><code class="docutils literal notranslate"><span class="pre">InjectSoftwarePipeline()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.InjectTmaBarrier"><code class="docutils literal notranslate"><span class="pre">InjectTmaBarrier()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LayoutInference"><code class="docutils literal notranslate"><span class="pre">LayoutInference()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LegalizeSafeMemoryAccess"><code class="docutils literal notranslate"><span class="pre">LegalizeSafeMemoryAccess()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LegalizeVectorizedLoop"><code class="docutils literal notranslate"><span class="pre">LegalizeVectorizedLoop()</span></code></a></li>

api/tilelang.language.builtin.html

@@ -406,19 +406,27 @@
 <p>The language interface for tl programs.</p>
 <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></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.create_list_of_mbarrier" title="Permalink to this definition">#</a></dt>
-<dd><p>Create a list of memory barrier operations.</p>
+<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>
+<dd><p>Create a list of memory barrier handles.</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 passed to the memory barrier creation operation</p>
+<dd class="field-odd"><p><strong>*args</strong> (<em>list</em><em> or </em><em>Any</em>) – Either a single list of arguments, or multiple arguments directly.</p>
 </dd>
 <dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A handle to the created list of memory barriers</p>
+<dd class="field-even"><p>Handle to the created list of memory barriers.</p>
 </dd>
 <dt class="field-odd">Return type<span class="colon">:</span></dt>
-<dd class="field-odd"><p>tir.Call</p>
+<dd class="field-odd"><p>tvm.tir.Call</p>
+</dd>
+<dt class="field-even">Raises<span class="colon">:</span></dt>
+<dd class="field-even"><p><strong>TypeError</strong> – If the input is not a list or variadic arguments.</p>
 </dd>
 </dl>
+<p class="rubric">Examples</p>
+<div class="doctest highlight-default notranslate"><div class="highlight"><pre><span></span><span class="gp">&gt;&gt;&gt; </span><span class="n">create_list_of_mbarrier</span><span class="p">([</span><span class="mi">128</span><span class="p">,</span> <span class="mi">128</span><span class="p">])</span>
+<span class="gp">&gt;&gt;&gt; </span><span class="n">create_list_of_mbarrier</span><span class="p">(</span><span class="mi">128</span><span class="p">,</span> <span class="mi">128</span><span class="p">)</span>
+</pre></div>
+</div>
 </dd></dl>
 
 <dl class="py function">
@@ -474,7 +482,7 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.mbarrier_arrive">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">mbarrier_arrive</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">mbarrier</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></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.mbarrier_arrive" 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">mbarrier_arrive</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">mbarrier</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.mbarrier_arrive" title="Permalink to this definition">#</a></dt>
 <dd><p>Arrive at memory barrier.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
@@ -503,7 +511,7 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.mbarrier_wait_parity">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">mbarrier_wait_parity</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">mbarrier</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></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">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></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.mbarrier_wait_parity" 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">mbarrier_wait_parity</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">mbarrier</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">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></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.mbarrier_wait_parity" title="Permalink to this definition">#</a></dt>
 <dd><p>Wait for memory barrier parity condition.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>

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">0x7f3f5cc1a170&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">0x7f42caca3310&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">0x7f3f5cc1a170&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">0x7f42caca3310&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">0x7f3f5cc1a1a0&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">0x7f42caca32e0&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>

api/tilelang.transform.html

@@ -567,6 +567,19 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.transform.InjectTmaBarrier">
+<span class="sig-prename descclassname"><span class="pre">tilelang.transform.</span></span><span class="sig-name descname"><span class="pre">InjectTmaBarrier</span></span><span class="sig-paren">(</span><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.transform.InjectTmaBarrier" title="Permalink to this definition">#</a></dt>
+<dd><dl class="field-list simple">
+<dt class="field-odd">Returns<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>fpass</strong> – The result pass</p>
+</dd>
+<dt class="field-even">Return type<span class="colon">:</span></dt>
+<dd class="field-even"><p>tvm.transform.Pass</p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.transform.LayoutInference">
 <span class="sig-prename descclassname"><span class="pre">tilelang.transform.</span></span><span class="sig-name descname"><span class="pre">LayoutInference</span></span><span class="sig-paren">(</span><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.transform.LayoutInference" title="Permalink to this definition">#</a></dt>
@@ -885,6 +898,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li><a class="reference internal" href="#tilelang.transform.InjectFenceProxy"><code class="docutils literal notranslate"><span class="pre">InjectFenceProxy()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.transform.InjectPTXAsyncCopy"><code class="docutils literal notranslate"><span class="pre">InjectPTXAsyncCopy()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.transform.InjectSoftwarePipeline"><code class="docutils literal notranslate"><span class="pre">InjectSoftwarePipeline()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.transform.InjectTmaBarrier"><code class="docutils literal notranslate"><span class="pre">InjectTmaBarrier()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.transform.LayoutInference"><code class="docutils literal notranslate"><span class="pre">LayoutInference()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.transform.LegalizeSafeMemoryAccess"><code class="docutils literal notranslate"><span class="pre">LegalizeSafeMemoryAccess()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.transform.LegalizeVectorizedLoop"><code class="docutils literal notranslate"><span class="pre">LegalizeVectorizedLoop()</span></code></a></li>