Update docs

Author: LeiWang1999

api/tilelang.html

@@ -862,6 +862,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.customize.html">tilelang.language.customize module</a><ul>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_add"><code class="docutils literal notranslate"><span class="pre">atomic_add()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_addx2"><code class="docutils literal notranslate"><span class="pre">atomic_addx2()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_addx4"><code class="docutils literal notranslate"><span class="pre">atomic_addx4()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.clamp"><code class="docutils literal notranslate"><span class="pre">clamp()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.dp4a"><code class="docutils literal notranslate"><span class="pre">dp4a()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.reshape"><code class="docutils literal notranslate"><span class="pre">reshape()</span></code></a></li>

api/tilelang.language.customize.html

@@ -450,6 +450,26 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.customize.atomic_addx4">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.customize.</span></span><span class="sig-name descname"><span class="pre">atomic_addx4</span></span><span class="sig-paren">(</span><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">Buffer</span></span></em>, <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">PrimExpr</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.customize.atomic_addx4" title="Permalink to this definition">#</a></dt>
+<dd><p>Perform an atomic addition operation with double-width operands.</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>dst</strong> (<em>Buffer</em>) – Destination buffer where the atomic addition will be performed</p></li>
+<li><p><strong>value</strong> (<em>PrimExpr</em>) – Value to be atomically added (double-width)</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>Handle to the double-width atomic addition operation</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>PrimExpr</p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.customize.clamp">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.customize.</span></span><span class="sig-name descname"><span class="pre">clamp</span></span><span class="sig-paren">(</span><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">PrimExpr</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">min_val</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">PrimExpr</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">max_val</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">PrimExpr</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.customize.clamp" title="Permalink to this definition">#</a></dt>
@@ -590,6 +610,7 @@
 <li><a class="reference internal" href="#">tilelang.language.customize module</a><ul>
 <li><a class="reference internal" href="#tilelang.language.customize.atomic_add"><code class="docutils literal notranslate"><span class="pre">atomic_add()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.customize.atomic_addx2"><code class="docutils literal notranslate"><span class="pre">atomic_addx2()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.customize.atomic_addx4"><code class="docutils literal notranslate"><span class="pre">atomic_addx4()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.customize.clamp"><code class="docutils literal notranslate"><span class="pre">clamp()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.customize.dp4a"><code class="docutils literal notranslate"><span class="pre">dp4a()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.customize.reshape"><code class="docutils literal notranslate"><span class="pre">reshape()</span></code></a></li>

api/tilelang.language.html

@@ -627,6 +627,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li class="toctree-l1"><a class="reference internal" href="tilelang.language.customize.html">tilelang.language.customize module</a><ul>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_add"><code class="docutils literal notranslate"><span class="pre">atomic_add()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_addx2"><code class="docutils literal notranslate"><span class="pre">atomic_addx2()</span></code></a></li>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.atomic_addx4"><code class="docutils literal notranslate"><span class="pre">atomic_addx4()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.clamp"><code class="docutils literal notranslate"><span class="pre">clamp()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.dp4a"><code class="docutils literal notranslate"><span class="pre">dp4a()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.customize.html#tilelang.language.customize.reshape"><code class="docutils literal notranslate"><span class="pre">reshape()</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">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>
+<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">0x7f36c2ffee30&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">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>
+<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">0x7f36c2ffee30&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">0x7f7e21b9b7c0&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">0x7f36c2ffee00&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

@@ -487,10 +487,10 @@ <h2>A</h2>
 </li>
         <li><a href="api/tilelang.language.logical.html#tilelang.language.logical.any_of">any_of() (in module tilelang.language.logical)</a>
 </li>
-    </ul></td>
-    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="api/tilelang.language.tir.op.html#tilelang.language.tir.op.anylist_getitem">anylist_getitem() (in module tilelang.language.tir.op)</a>
 </li>
+    </ul></td>
+    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="api/tilelang.language.tir.op.html#tilelang.language.tir.op.anylist_resetitem">anylist_resetitem() (in module tilelang.language.tir.op)</a>
 </li>
         <li><a href="api/tilelang.language.tir.op.html#tilelang.language.tir.op.anylist_setitem_call_cpacked">anylist_setitem_call_cpacked() (in module tilelang.language.tir.op)</a>
@@ -546,6 +546,8 @@ <h2>A</h2>
         <li><a href="api/tilelang.language.customize.html#tilelang.language.customize.atomic_add">atomic_add() (in module tilelang.language.customize)</a>
 </li>
         <li><a href="api/tilelang.language.customize.html#tilelang.language.customize.atomic_addx2">atomic_addx2() (in module tilelang.language.customize)</a>
+</li>
+        <li><a href="api/tilelang.language.customize.html#tilelang.language.customize.atomic_addx4">atomic_addx4() (in module tilelang.language.customize)</a>
 </li>
         <li><a href="api/tilelang.utils.tensor.html#tilelang.utils.tensor.TensorSupplyType.Auto">Auto (tilelang.utils.tensor.TensorSupplyType attribute)</a>
 </li>