Update docs

Author: LeiWang1999

api/tilelang.carver.arch.driver.cuda_driver.html

@@ -875,6 +875,11 @@
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.carver.arch.driver.cuda_driver.get_persisting_l2_cache_max_size">
+<span class="sig-prename descclassname"><span class="pre">tilelang.carver.arch.driver.cuda_driver.</span></span><span class="sig-name descname"><span class="pre">get_persisting_l2_cache_max_size</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">device_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</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">0</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">int</span></span></span><a class="headerlink" href="#tilelang.carver.arch.driver.cuda_driver.get_persisting_l2_cache_max_size" title="Permalink to this definition">#</a></dt>
+<dd></dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.carver.arch.driver.cuda_driver.get_shared_memory_per_block">
 <span class="sig-prename descclassname"><span class="pre">tilelang.carver.arch.driver.cuda_driver.</span></span><span class="sig-name descname"><span class="pre">get_shared_memory_per_block</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">device_id</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">int</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">0</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">format</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">'bytes'</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">Optional</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">]</span></span></span></span><a class="headerlink" href="#tilelang.carver.arch.driver.cuda_driver.get_shared_memory_per_block" title="Permalink to this definition">#</a></dt>
@@ -1015,6 +1020,7 @@
 <li><a class="reference internal" href="#tilelang.carver.arch.driver.cuda_driver.get_device_name"><code class="docutils literal notranslate"><span class="pre">get_device_name()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.carver.arch.driver.cuda_driver.get_max_dynamic_shared_size_bytes"><code class="docutils literal notranslate"><span class="pre">get_max_dynamic_shared_size_bytes()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.carver.arch.driver.cuda_driver.get_num_sms"><code class="docutils literal notranslate"><span class="pre">get_num_sms()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.carver.arch.driver.cuda_driver.get_persisting_l2_cache_max_size"><code class="docutils literal notranslate"><span class="pre">get_persisting_l2_cache_max_size()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.carver.arch.driver.cuda_driver.get_shared_memory_per_block"><code class="docutils literal notranslate"><span class="pre">get_shared_memory_per_block()</span></code></a></li>
 </ul>
 </li>

api/tilelang.carver.arch.driver.html

@@ -496,6 +496,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li class="toctree-l2"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_device_name"><code class="docutils literal notranslate"><span class="pre">get_device_name()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_max_dynamic_shared_size_bytes"><code class="docutils literal notranslate"><span class="pre">get_max_dynamic_shared_size_bytes()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_num_sms"><code class="docutils literal notranslate"><span class="pre">get_num_sms()</span></code></a></li>
+<li class="toctree-l2"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_persisting_l2_cache_max_size"><code class="docutils literal notranslate"><span class="pre">get_persisting_l2_cache_max_size()</span></code></a></li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_shared_memory_per_block"><code class="docutils literal notranslate"><span class="pre">get_shared_memory_per_block()</span></code></a></li>
 </ul>
 </li>

api/tilelang.carver.arch.html

@@ -427,6 +427,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 <li class="toctree-l4"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_device_name"><code class="docutils literal notranslate"><span class="pre">get_device_name()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_max_dynamic_shared_size_bytes"><code class="docutils literal notranslate"><span class="pre">get_max_dynamic_shared_size_bytes()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_num_sms"><code class="docutils literal notranslate"><span class="pre">get_num_sms()</span></code></a></li>
+<li class="toctree-l4"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_persisting_l2_cache_max_size"><code class="docutils literal notranslate"><span class="pre">get_persisting_l2_cache_max_size()</span></code></a></li>
 <li class="toctree-l4"><a class="reference internal" href="tilelang.carver.arch.driver.cuda_driver.html#tilelang.carver.arch.driver.cuda_driver.get_shared_memory_per_block"><code class="docutils literal notranslate"><span class="pre">get_shared_memory_per_block()</span></code></a></li>
 </ul>
 </li>

api/tilelang.html

@@ -898,6 +898,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
 </ul>
 </li>
 <li class="toctree-l2"><a class="reference internal" href="tilelang.language.html#module-tilelang.language">Module contents</a><ul>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.annotate_l2_hit_ratio"><code class="docutils literal notranslate"><span class="pre">annotate_l2_hit_ratio()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.annotate_layout"><code class="docutils literal notranslate"><span class="pre">annotate_layout()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.annotate_padding"><code class="docutils literal notranslate"><span class="pre">annotate_padding()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.import_source"><code class="docutils literal notranslate"><span class="pre">import_source()</span></code></a></li>
@@ -1049,6 +1050,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.LoopVectorizeDynamic"><code class="docutils literal notranslate"><span class="pre">LoopVectorizeDynamic()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerDeviceStorageAccessInfo"><code class="docutils literal notranslate"><span class="pre">LowerDeviceStorageAccessInfo()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerHopperIntrin"><code class="docutils literal notranslate"><span class="pre">LowerHopperIntrin()</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerL2Persistent"><code class="docutils literal notranslate"><span class="pre">LowerL2Persistent()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.LowerTileOp"><code class="docutils literal notranslate"><span class="pre">LowerTileOp()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.MakePackedAPI"><code class="docutils literal notranslate"><span class="pre">MakePackedAPI()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.transform.html#tilelang.transform.MergeIfStmt"><code class="docutils literal notranslate"><span class="pre">MergeIfStmt()</span></code></a></li>

api/tilelang.jit.adapter.html

@@ -517,6 +517,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.backend"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.backend</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.create_dispatch_func"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.create_dispatch_func()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.device_mod"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.device_mod</span></code></a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_l2_persistent_map"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.generate_l2_persistent_map()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_tma_descriptor_args"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.generate_tma_descriptor_args()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.get_dynamic_symbolic_set"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.get_dynamic_symbolic_set()</span></code></a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang.jit.adapter.wrapper.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.get_init_func"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.get_init_func()</span></code></a></li>

api/tilelang.jit.adapter.wrapper.html

@@ -510,6 +510,11 @@
 <span class="sig-name descname"><span class="pre">device_mod</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">Optional</span><span class="p"><span class="pre">[</span></span><span class="pre">IRModule</span><span class="p"><span class="pre">]</span></span></em><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">None</span></em><a class="headerlink" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.device_mod" title="Permalink to this definition">#</a></dt>
 <dd></dd></dl>
 
+<dl class="py method">
+<dt class="sig sig-object py" id="tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_l2_persistent_map">
+<span class="sig-name descname"><span class="pre">generate_l2_persistent_map</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">function_name</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</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">str</span></span></span><a class="headerlink" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_l2_persistent_map" title="Permalink to this definition">#</a></dt>
+<dd></dd></dl>
+
 <dl class="py method">
 <dt class="sig sig-object py" id="tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_tma_descriptor_args">
 <span class="sig-name descname"><span class="pre">generate_tma_descriptor_args</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">desc_name_map</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Dict</span><span class="p"><span class="pre">[</span></span><span class="pre">str</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="pre">str</span><span class="p"><span class="pre">]</span></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">str</span></span></span><a class="headerlink" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_tma_descriptor_args" title="Permalink to this definition">#</a></dt>
@@ -717,6 +722,7 @@
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.backend"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.backend</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.create_dispatch_func"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.create_dispatch_func()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.device_mod"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.device_mod</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_l2_persistent_map"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.generate_l2_persistent_map()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.generate_tma_descriptor_args"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.generate_tma_descriptor_args()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.get_dynamic_symbolic_set"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.get_dynamic_symbolic_set()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.get_init_func"><code class="docutils literal notranslate"><span class="pre">TLCUDASourceWrapper.get_init_func()</span></code></a></li>

api/tilelang.language.html

@@ -771,6 +771,21 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <section id="module-tilelang.language">
 <span id="module-contents"></span><h2>Module contents<a class="headerlink" href="#module-tilelang.language" title="Permalink to this heading">#</a></h2>
 <p>The language interface for tl programs.</p>
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.language.annotate_l2_hit_ratio">
+<span class="sig-prename descclassname"><span class="pre">tilelang.language.</span></span><span class="sig-name descname"><span class="pre">annotate_l2_hit_ratio</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">l2_hit_ratio_map</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Dict</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.annotate_l2_hit_ratio" title="Permalink to this definition">#</a></dt>
+<dd><p>Annotate the L2 hit ratio of the buffer, detailed explanation please refer to:
+<a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/#l2-policy-for-persisting-accesses">https://docs.nvidia.com/cuda/cuda-c-programming-guide/#l2-policy-for-persisting-accesses</a></p>
+<dl class="field-list simple">
+<dt class="field-odd">Parameters<span class="colon">:</span></dt>
+<dd class="field-odd"><p><strong>l2_hit_ratio_map</strong> (<em>dict</em>) – a dictionary of buffer to L2 hit ratio value</p>
+</dd>
+</dl>
+<p class="rubric">Example</p>
+<p># 0.5 is the hit ratio
+T.annotate_l2_hit_ratio({A: 0.5})</p>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.annotate_layout">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.</span></span><span class="sig-name descname"><span class="pre">annotate_layout</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">layout_map</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Dict</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.annotate_layout" title="Permalink to this definition">#</a></dt>
@@ -932,6 +947,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
 <li><a class="reference internal" href="#subpackages">Subpackages</a></li>
 <li><a class="reference internal" href="#submodules">Submodules</a></li>
 <li><a class="reference internal" href="#module-tilelang.language">Module contents</a><ul>
+<li><a class="reference internal" href="#tilelang.language.annotate_l2_hit_ratio"><code class="docutils literal notranslate"><span class="pre">annotate_l2_hit_ratio()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.annotate_layout"><code class="docutils literal notranslate"><span class="pre">annotate_layout()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.annotate_padding"><code class="docutils literal notranslate"><span class="pre">annotate_padding()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.import_source"><code class="docutils literal notranslate"><span class="pre">import_source()</span></code></a></li>