Update docs

Author: LeiWang1999

api/tilelang.autotuner.html

@@ -412,24 +412,39 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 </ul>
 </li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult"><code class="docutils literal notranslate"><span class="pre">AutotuneResult</span></code></a><ul>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.latency"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.latency</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.config"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.config</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.ref_latency"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.ref_latency</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.libcode"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.libcode</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.func"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.func</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.kernel"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.kernel</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.latency"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.latency</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.libcode"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.libcode</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.ref_latency"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.ref_latency</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id0"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.config</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id1"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.func</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id2"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.kernel</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id3"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.latency</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id4"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.libcode</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id5"><code class="docutils literal notranslate"><span class="pre">AutotuneResult.ref_latency</span></code></a></li>
 </ul>
 </li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext"><code class="docutils literal notranslate"><span class="pre">JITContext</span></code></a><ul>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.atol"><code class="docutils literal notranslate"><span class="pre">JITContext.atol</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.max_mismatched_ratio"><code class="docutils literal notranslate"><span class="pre">JITContext.max_mismatched_ratio</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.out_idx"><code class="docutils literal notranslate"><span class="pre">JITContext.out_idx</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.profiler"><code class="docutils literal notranslate"><span class="pre">JITContext.profiler</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.supply_type"><code class="docutils literal notranslate"><span class="pre">JITContext.supply_type</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.ref_prog"><code class="docutils literal notranslate"><span class="pre">JITContext.ref_prog</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.rtol"><code class="docutils literal notranslate"><span class="pre">JITContext.rtol</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.atol"><code class="docutils literal notranslate"><span class="pre">JITContext.atol</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.max_mismatched_ratio"><code class="docutils literal notranslate"><span class="pre">JITContext.max_mismatched_ratio</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.skip_check"><code class="docutils literal notranslate"><span class="pre">JITContext.skip_check</span></code></a></li>
-<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.supply_type"><code class="docutils literal notranslate"><span class="pre">JITContext.supply_type</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.profiler"><code class="docutils literal notranslate"><span class="pre">JITContext.profiler</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.JITContext.target"><code class="docutils literal notranslate"><span class="pre">JITContext.target</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id6"><code class="docutils literal notranslate"><span class="pre">JITContext.atol</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id7"><code class="docutils literal notranslate"><span class="pre">JITContext.max_mismatched_ratio</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id8"><code class="docutils literal notranslate"><span class="pre">JITContext.out_idx</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id9"><code class="docutils literal notranslate"><span class="pre">JITContext.profiler</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id10"><code class="docutils literal notranslate"><span class="pre">JITContext.ref_prog</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id11"><code class="docutils literal notranslate"><span class="pre">JITContext.rtol</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id12"><code class="docutils literal notranslate"><span class="pre">JITContext.skip_check</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id13"><code class="docutils literal notranslate"><span class="pre">JITContext.supply_type</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.autotuner.html#id14"><code class="docutils literal notranslate"><span class="pre">JITContext.target</span></code></a></li>
 </ul>
 </li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.autotuner.html#tilelang.autotuner.autotune"><code class="docutils literal notranslate"><span class="pre">autotune()</span></code></a></li>

api/tilelang.html

@@ -399,26 +399,103 @@
         <article role="main">
           <section id="module-tilelang.language.allocate">
 <span id="tilelang-language-allocate-module"></span><h1>tilelang.language.allocate module<a class="headerlink" href="#module-tilelang.language.allocate" title="Permalink to this heading">#</a></h1>
-<p>The language interface for tl programs.</p>
+<p>Memory allocation utilities for Tile-AI programs.</p>
+<p>This module provides a set of functions for allocating different types of memory buffers
+in Tile-AI programs. It wraps TVM’s buffer allocation functionality with convenient
+interfaces for different memory scopes.</p>
+<dl class="simple">
+<dt>Available allocation functions:</dt><dd><ul class="simple">
+<li><p>alloc_shared: Allocates shared memory buffers for inter-thread communication</p></li>
+<li><p>alloc_local: Allocates local memory buffers for thread-private storage</p></li>
+<li><p>alloc_fragment: Allocates fragment memory buffers for specialized operations</p></li>
+<li><p>alloc_var: Allocates single-element variable buffers</p></li>
+</ul>
+</dd>
+</dl>
+<p>Each function takes shape and dtype parameters and returns a TVM buffer object
+with the appropriate memory scope.</p>
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.allocate.alloc_fragment">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_fragment</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">scope</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">'local.fragment'</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_fragment" title="Permalink to this definition">#</a></dt>
-<dd></dd></dl>
+<dd><p>Allocate a fragment memory buffer for specialized operations.</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>shape</strong> (<em>tuple</em>) – The shape of the buffer to allocate</p></li>
+<li><p><strong>dtype</strong> (<em>str</em>) – The data type of the buffer (e.g., ‘float32’, ‘int32’)</p></li>
+<li><p><strong>scope</strong> (<em>str</em><em>, </em><em>optional</em>) – The memory scope. Defaults to “local.fragment”</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A TVM buffer object allocated in fragment memory</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>T.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.allocate.alloc_local">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_local</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">scope</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">'local'</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_local" title="Permalink to this definition">#</a></dt>
-<dd></dd></dl>
+<dd><p>Allocate a local memory buffer for thread-private storage.</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>shape</strong> (<em>tuple</em>) – The shape of the buffer to allocate</p></li>
+<li><p><strong>dtype</strong> (<em>str</em>) – The data type of the buffer (e.g., ‘float32’, ‘int32’)</p></li>
+<li><p><strong>scope</strong> (<em>str</em><em>, </em><em>optional</em>) – The memory scope. Defaults to “local”</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A TVM buffer object allocated in local memory</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>T.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.allocate.alloc_shared">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_shared</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">scope</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">'shared.dyn'</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_shared" title="Permalink to this definition">#</a></dt>
-<dd></dd></dl>
+<dd><p>Allocate a shared memory buffer for inter-thread communication.</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>shape</strong> (<em>tuple</em>) – The shape of the buffer to allocate</p></li>
+<li><p><strong>dtype</strong> (<em>str</em>) – The data type of the buffer (e.g., ‘float32’, ‘int32’)</p></li>
+<li><p><strong>scope</strong> (<em>str</em><em>, </em><em>optional</em>) – The memory scope. Defaults to “shared.dyn”</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A TVM buffer object allocated in shared memory</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>T.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
 
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.allocate.alloc_var">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.allocate.</span></span><span class="sig-name descname"><span class="pre">alloc_var</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">scope</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">'local.var'</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.allocate.alloc_var" title="Permalink to this definition">#</a></dt>
-<dd></dd></dl>
+<dd><p>Allocate a single-element variable buffer.</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>dtype</strong> (<em>str</em>) – The data type of the buffer (e.g., ‘float32’, ‘int32’)</p></li>
+<li><p><strong>scope</strong> (<em>str</em><em>, </em><em>optional</em>) – The memory scope. Defaults to “local.var”</p></li>
+</ul>
+</dd>
+<dt class="field-even">Returns<span class="colon">:</span></dt>
+<dd class="field-even"><p>A TVM buffer object allocated as a single-element variable</p>
+</dd>
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>T.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
 
 </section>