Update docs

Author: LeiWang1999

api/tilelang.html

@@ -795,8 +795,13 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 </ul>
 </li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.jit.html#module-tilelang.jit">Module contents</a><ul>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.Kernel"><code class="docutils literal notranslate"><span class="pre">Kernel</span></code></a><ul>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.Kernel.get_kernel_source"><code class="docutils literal notranslate"><span class="pre">Kernel.get_kernel_source()</span></code></a></li>
+</ul>
+</li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.Program"><code class="docutils literal notranslate"><span class="pre">Program</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.compile"><code class="docutils literal notranslate"><span class="pre">compile()</span></code></a></li>
-<li class="toctree-l3"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.jit"><code class="docutils literal notranslate"><span class="pre">jit()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.jit.html#tilelang.jit.jit"><code class="docutils literal notranslate"><span class="pre">jit</span></code></a></li>
 </ul>
 </li>
 </ul>

api/tilelang.jit.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">0x7ff9308ac310&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">0x7fa1465c7220&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">0x7ff9308ac310&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">0x7fa1465c7220&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.logical.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">0x7ff9308aeb00&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">0x7fa1465c5540&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.language.proxy.html

@@ -1345,6 +1345,8 @@ <h2>G</h2>
           <li><a href="api/tilelang.jit.adapter.ctypes.adapter.html#tilelang.jit.adapter.ctypes.adapter.CtypesKernelAdapter.get_kernel_source">(tilelang.jit.adapter.ctypes.adapter.CtypesKernelAdapter method)</a>
 </li>
           <li><a href="api/tilelang.jit.adapter.cython.adapter.html#tilelang.jit.adapter.cython.adapter.CythonKernelAdapter.get_kernel_source">(tilelang.jit.adapter.cython.adapter.CythonKernelAdapter method)</a>
+</li>
+          <li><a href="api/tilelang.jit.html#tilelang.jit.Kernel.get_kernel_source">(tilelang.jit.Kernel method)</a>
 </li>
           <li><a href="api/tilelang.jit.kernel.html#tilelang.jit.kernel.JITKernel.get_kernel_source">(tilelang.jit.kernel.JITKernel method)</a>
 </li>
@@ -1813,12 +1815,10 @@ <h2>I</h2>
   <h2>J</h2>
   <table style="width: 100%" class="indextable genindextable"><tr>
     <td style="width: 33%; vertical-align: top;"><ul>
+        <li><a href="api/tilelang.jit.html#tilelang.jit.jit">jit (class in tilelang.jit)</a>
+</li>
         <li><a href="api/tilelang.autotuner.html#tilelang.autotuner.jit">jit() (in module tilelang.autotuner)</a>
-
-        <ul>
-          <li><a href="api/tilelang.jit.html#tilelang.jit.jit">(in module tilelang.jit)</a>
 </li>
-        </ul></li>
     </ul></td>
     <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="api/tilelang.autotuner.html#tilelang.autotuner.JITContext">JITContext (class in tilelang.autotuner)</a>
@@ -1847,6 +1847,8 @@ <h2>K</h2>
           <li><a href="api/tilelang.primitives.gemm.base.html#tilelang.primitives.gemm.base.GemmBaseParams.k_pack">(tilelang.primitives.gemm.base.GemmBaseParams attribute)</a>
 </li>
         </ul></li>
+        <li><a href="api/tilelang.jit.html#tilelang.jit.Kernel">Kernel (class in tilelang.jit)</a>
+</li>
         <li><a href="api/tilelang.autotuner.html#id2">kernel (tilelang.autotuner.AutotuneResult attribute)</a>, <a href="api/tilelang.autotuner.html#tilelang.autotuner.AutotuneResult.kernel">[1]</a>
 
         <ul>
@@ -2665,6 +2667,8 @@ <h2>P</h2>
         <li><a href="api/tilelang.carver.roller.shape_inference.tir.html#tilelang.carver.roller.shape_inference.tir.DependencyAnalysis.print_dependencies">print_dependencies() (tilelang.carver.roller.shape_inference.tir.DependencyAnalysis method)</a>
 </li>
         <li><a href="api/tilelang.profiler.html#tilelang.profiler.Profiler">Profiler (class in tilelang.profiler)</a>
+</li>
+        <li><a href="api/tilelang.jit.html#tilelang.jit.Program">Program (class in tilelang.jit)</a>
 </li>
         <li><a href="api/tilelang.carver.roller.node.html#tilelang.carver.roller.node.PrimFuncNode.propagate">propagate() (tilelang.carver.roller.node.PrimFuncNode method)</a>
 </li>