Update docs

Author: LeiWang1999

api/tilelang.html

@@ -807,8 +807,10 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.builtin.html">tilelang.language.builtin module</a><ul>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_list_of_mbarrier"><code class="docutils literal notranslate"><span class="pre">create_list_of_mbarrier()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_tma_descriptor"><code class="docutils literal notranslate"><span class="pre">create_tma_descriptor()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.dec_max_nreg"><code class="docutils literal notranslate"><span class="pre">dec_max_nreg()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.fence_proxy_async"><code class="docutils literal notranslate"><span class="pre">fence_proxy_async()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.get_mbarrier"><code class="docutils literal notranslate"><span class="pre">get_mbarrier()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.inc_max_nreg"><code class="docutils literal notranslate"><span class="pre">inc_max_nreg()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_arrive"><code class="docutils literal notranslate"><span class="pre">mbarrier_arrive()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_expect_tx"><code class="docutils literal notranslate"><span class="pre">mbarrier_expect_tx()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_wait_parity"><code class="docutils literal notranslate"><span class="pre">mbarrier_wait_parity()</span></code></a></li>
@@ -998,6 +1000,10 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l1"><a class="reference internal" href="tilelang.testing.html">tilelang.testing package</a><ul>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.testing.html#module-tilelang.testing">Module contents</a><ul>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.main"><code class="docutils literal notranslate"><span class="pre">main()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.requires_cuda_compute_version"><code class="docutils literal notranslate"><span class="pre">requires_cuda_compute_version()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.requires_cuda_compute_version_eq"><code class="docutils literal notranslate"><span class="pre">requires_cuda_compute_version_eq()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.requires_cuda_compute_version_ge"><code class="docutils literal notranslate"><span class="pre">requires_cuda_compute_version_ge()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.requires_cuda_compute_version_gt"><code class="docutils literal notranslate"><span class="pre">requires_cuda_compute_version_gt()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.testing.html#tilelang.testing.set_random_seed"><code class="docutils literal notranslate"><span class="pre">set_random_seed()</span></code></a></li>
 </ul>
 </li>

api/tilelang.jit.html

@@ -520,7 +520,8 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 “tl.disable_warp_specialized”: bool, default: False
 “tl.config_index_bitwidth”: int, default: None
 “tl.disable_dynamic_tail_split”: bool, default: False
-“tl.dynamic_vectorize_size_bits”: int, default: 128</p>
+“tl.dynamic_vectorize_size_bits”: int, default: 128
+“tl.disable_safe_memory_legalize”: bool, default: False</p>
 </dd>
 </dl>
 </div></blockquote>

api/tilelang.language.builtin.html

@@ -446,6 +446,12 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.dec_max_nreg">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">dec_max_nreg</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">reg_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.builtin.dec_max_nreg" title="Permalink to this definition">#</a></dt>
+<dd><p>Decrement the maximum number of registers to use.</p>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.fence_proxy_async">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">fence_proxy_async</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.fence_proxy_async" title="Permalink to this definition">#</a></dt>
@@ -480,6 +486,12 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.builtin.inc_max_nreg">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">inc_max_nreg</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">reg_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.builtin.inc_max_nreg" title="Permalink to this definition">#</a></dt>
+<dd><p>Increment the maximum number of registers to use.</p>
+</dd></dl>
+
 <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 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>
@@ -563,28 +575,25 @@
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.builtin.no_set_max_nreg">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">no_set_max_nreg</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.no_set_max_nreg" 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">no_set_max_nreg</span></span><span class="sig-paren">(</span><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.builtin.no_set_max_nreg" title="Permalink to this definition">#</a></dt>
 <dd><p>Disable the maximum register limit setting.</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 for the operation</p>
-</dd>
-<dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A handle to the register limit disable 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.set_max_nreg">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">set_max_nreg</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.set_max_nreg" title="Permalink to this definition">#</a></dt>
-<dd><p>Set the maximum number of registers to use.</p>
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.builtin.</span></span><span class="sig-name descname"><span class="pre">set_max_nreg</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">reg_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>, <em class="sig-param"><span class="n"><span class="pre">is_inc</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.set_max_nreg" title="Permalink to this definition">#</a></dt>
+<dd><p>Set the maximum number of registers to use.
+Detailed Documentation:
+<a class="reference external" href="https://docs.nvidia.com/cuda/parallel-thread-execution/#miscellaneous-instructions-setmaxnreg">https://docs.nvidia.com/cuda/parallel-thread-execution/#miscellaneous-instructions-setmaxnreg</a></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 register allocation limits</p>
+<dd class="field-odd"><ul class="simple">
+<li><p><strong>reg_count</strong> – int
+The number of registers to allocate</p></li>
+<li><p><strong>is_inc</strong> – int
+Whether to increment or decrement the register count
+0 if decrement, 1 if increment</p></li>
+</ul>
 </dd>
 <dt class="field-even">Returns<span class="colon">:</span></dt>
 <dd class="field-even"><p>A handle to the register setting operation</p>
@@ -723,8 +732,10 @@
 <li><a class="reference internal" href="#">tilelang.language.builtin module</a><ul>
 <li><a class="reference internal" href="#tilelang.language.builtin.create_list_of_mbarrier"><code class="docutils literal notranslate"><span class="pre">create_list_of_mbarrier()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.create_tma_descriptor"><code class="docutils literal notranslate"><span class="pre">create_tma_descriptor()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.dec_max_nreg"><code class="docutils literal notranslate"><span class="pre">dec_max_nreg()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.fence_proxy_async"><code class="docutils literal notranslate"><span class="pre">fence_proxy_async()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.get_mbarrier"><code class="docutils literal notranslate"><span class="pre">get_mbarrier()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.language.builtin.inc_max_nreg"><code class="docutils literal notranslate"><span class="pre">inc_max_nreg()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.mbarrier_arrive"><code class="docutils literal notranslate"><span class="pre">mbarrier_arrive()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.mbarrier_expect_tx"><code class="docutils literal notranslate"><span class="pre">mbarrier_expect_tx()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.builtin.mbarrier_wait_parity"><code class="docutils literal notranslate"><span class="pre">mbarrier_wait_parity()</span></code></a></li>

api/tilelang.language.html

@@ -585,8 +585,10 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li class="toctree-l1"><a class="reference internal" href="tilelang.language.builtin.html">tilelang.language.builtin module</a><ul>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_list_of_mbarrier"><code class="docutils literal notranslate"><span class="pre">create_list_of_mbarrier()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.create_tma_descriptor"><code class="docutils literal notranslate"><span class="pre">create_tma_descriptor()</span></code></a></li>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.dec_max_nreg"><code class="docutils literal notranslate"><span class="pre">dec_max_nreg()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.fence_proxy_async"><code class="docutils literal notranslate"><span class="pre">fence_proxy_async()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.get_mbarrier"><code class="docutils literal notranslate"><span class="pre">get_mbarrier()</span></code></a></li>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.inc_max_nreg"><code class="docutils literal notranslate"><span class="pre">inc_max_nreg()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_arrive"><code class="docutils literal notranslate"><span class="pre">mbarrier_arrive()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_expect_tx"><code class="docutils literal notranslate"><span class="pre">mbarrier_expect_tx()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.builtin.html#tilelang.language.builtin.mbarrier_wait_parity"><code class="docutils literal notranslate"><span class="pre">mbarrier_wait_parity()</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">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>
+<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">0x7f2167a05fc0&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">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>
+<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">0x7f2167a05fc0&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>