Update docs

Author: github-actions[bot]

_sources/autoapi/tilelang/engine/phase/index.rst.txt

@@ -11,6 +11,7 @@ Functions
 
    tilelang.engine.phase.allow_warp_specialized
    tilelang.engine.phase.allow_tma_and_warp_specialized
+   tilelang.engine.phase.allow_tma_lower
    tilelang.engine.phase.allow_fence_proxy
    tilelang.engine.phase.allow_vectorize
    tilelang.engine.phase.allow_global_thread_synchronization
@@ -33,6 +34,15 @@ Module Contents
 
 .. py:function:: allow_tma_and_warp_specialized(pass_ctx = None, target = None)
 
+.. py:function:: allow_tma_lower(pass_ctx = None, target = None)
+
+   Return True when TMA lowering is enabled for the given target.
+
+   This is intentionally decoupled from warp specialization so Hopper TMA can
+   be used in a non-warp-specialized pipeline (e.g., no-WS kernels still need
+   mbarrier allocation/init and expect_tx injection).
+
+
 .. py:function:: allow_fence_proxy(target = None)
 
 .. py:function:: allow_vectorize(pass_ctx = None)

autoapi/tilelang/engine/phase/index.html

@@ -568,43 +568,46 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 <tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_tma_and_warp_specialized" title="tilelang.engine.phase.allow_tma_and_warp_specialized"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_tma_and_warp_specialized</span></code></a>([pass_ctx, target])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_fence_proxy" title="tilelang.engine.phase.allow_fence_proxy"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_fence_proxy</span></code></a>([target])</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_tma_lower" title="tilelang.engine.phase.allow_tma_lower"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_tma_lower</span></code></a>([pass_ctx, target])</p></td>
+<td><p>Return True when TMA lowering is enabled for the given target.</p></td>
+</tr>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_fence_proxy" title="tilelang.engine.phase.allow_fence_proxy"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_fence_proxy</span></code></a>([target])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_vectorize" title="tilelang.engine.phase.allow_vectorize"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_vectorize</span></code></a>([pass_ctx])</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_vectorize" title="tilelang.engine.phase.allow_vectorize"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_vectorize</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_global_thread_synchronization" title="tilelang.engine.phase.allow_global_thread_synchronization"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_global_thread_synchronization</span></code></a>([pass_ctx])</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.allow_global_thread_synchronization" title="tilelang.engine.phase.allow_global_thread_synchronization"><code class="xref py py-obj docutils literal notranslate"><span class="pre">allow_global_thread_synchronization</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_aggressive_merge" title="tilelang.engine.phase.should_enable_aggressive_merge"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_aggressive_merge</span></code></a>([pass_ctx, target])</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_aggressive_merge" title="tilelang.engine.phase.should_enable_aggressive_merge"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_aggressive_merge</span></code></a>([pass_ctx, target])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_force_let_inline" title="tilelang.engine.phase.should_force_let_inline"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_force_let_inline</span></code></a>([pass_ctx])</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_force_let_inline" title="tilelang.engine.phase.should_force_let_inline"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_force_let_inline</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_ast_print" title="tilelang.engine.phase.should_enable_ast_print"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_ast_print</span></code></a>([pass_ctx])</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_ast_print" title="tilelang.engine.phase.should_enable_ast_print"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_ast_print</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_layout_visual" title="tilelang.engine.phase.should_enable_layout_visual"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_layout_visual</span></code></a>([pass_ctx])</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_layout_visual" title="tilelang.engine.phase.should_enable_layout_visual"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_layout_visual</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_race_check" title="tilelang.engine.phase.should_enable_race_check"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_race_check</span></code></a>([pass_ctx])</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.should_enable_race_check" title="tilelang.engine.phase.should_enable_race_check"><code class="xref py py-obj docutils literal notranslate"><span class="pre">should_enable_race_check</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.get_layout_visual_formats" title="tilelang.engine.phase.get_layout_visual_formats"><code class="xref py py-obj docutils literal notranslate"><span class="pre">get_layout_visual_formats</span></code></a>([pass_ctx])</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.get_layout_visual_formats" title="tilelang.engine.phase.get_layout_visual_formats"><code class="xref py py-obj docutils literal notranslate"><span class="pre">get_layout_visual_formats</span></code></a>([pass_ctx])</p></td>
 <td><p></p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.LayoutVisual" title="tilelang.engine.phase.LayoutVisual"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LayoutVisual</span></code></a>(mod)</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.LayoutVisual" title="tilelang.engine.phase.LayoutVisual"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LayoutVisual</span></code></a>(mod)</p></td>
 <td><p>Apply layout visualization pass if enabled.</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.PreLowerSemanticCheck" title="tilelang.engine.phase.PreLowerSemanticCheck"><code class="xref py py-obj docutils literal notranslate"><span class="pre">PreLowerSemanticCheck</span></code></a>(mod)</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.PreLowerSemanticCheck" title="tilelang.engine.phase.PreLowerSemanticCheck"><code class="xref py py-obj docutils literal notranslate"><span class="pre">PreLowerSemanticCheck</span></code></a>(mod)</p></td>
 <td><p>Check whether the module is valid before lowering. If not, raise a user-friendly error</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.LowerAndLegalize" title="tilelang.engine.phase.LowerAndLegalize"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LowerAndLegalize</span></code></a>(mod, target)</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.LowerAndLegalize" title="tilelang.engine.phase.LowerAndLegalize"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LowerAndLegalize</span></code></a>(mod, target)</p></td>
 <td><p>Bind target information and progressively legalize and lower frontend Tile IR into a form suitable for downstream optimization and codegen.</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.engine.phase.OptimizeForTarget" title="tilelang.engine.phase.OptimizeForTarget"><code class="xref py py-obj docutils literal notranslate"><span class="pre">OptimizeForTarget</span></code></a>(mod, target)</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.engine.phase.OptimizeForTarget" title="tilelang.engine.phase.OptimizeForTarget"><code class="xref py py-obj docutils literal notranslate"><span class="pre">OptimizeForTarget</span></code></a>(mod, target)</p></td>
 <td><p></p></td>
 </tr>
 </tbody>
@@ -645,6 +648,26 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="tilelang.engine.phase.allow_tma_lower">
+<span class="sig-prename descclassname"><span class="pre">tilelang.engine.phase.</span></span><span class="sig-name descname"><span class="pre">allow_tma_lower</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">pass_ctx</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">target</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.engine.phase.allow_tma_lower" title="Link to this definition">¶</a></dt>
+<dd><p>Return True when TMA lowering is enabled for the given target.</p>
+<p>This is intentionally decoupled from warp specialization so Hopper TMA can
+be used in a non-warp-specialized pipeline (e.g., no-WS kernels still need
+mbarrier allocation/init and expect_tx injection).</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>pass_ctx</strong> (<em>tilelang.transform.PassContext</em><em> | </em><em>None</em>)</p></li>
+<li><p><strong>target</strong> (<em>tvm.target.Target</em><em> | </em><em>None</em>)</p></li>
+</ul>
+</dd>
+<dt class="field-even">Return type<span class="colon">:</span></dt>
+<dd class="field-even"><p><a class="reference internal" href="../../language/dtypes/index.html#tilelang.language.dtypes.bool" title="tilelang.language.dtypes.bool">bool</a></p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.engine.phase.allow_fence_proxy">
 <span class="sig-prename descclassname"><span class="pre">tilelang.engine.phase.</span></span><span class="sig-name descname"><span class="pre">allow_fence_proxy</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">target</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.engine.phase.allow_fence_proxy" title="Link to this definition">¶</a></dt>
@@ -903,6 +926,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <li><a class="reference internal" href="#module-contents">Module Contents</a><ul>
 <li><a class="reference internal" href="#tilelang.engine.phase.allow_warp_specialized"><code class="docutils literal notranslate"><span class="pre">allow_warp_specialized()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.engine.phase.allow_tma_and_warp_specialized"><code class="docutils literal notranslate"><span class="pre">allow_tma_and_warp_specialized()</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.engine.phase.allow_tma_lower"><code class="docutils literal notranslate"><span class="pre">allow_tma_lower()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.engine.phase.allow_fence_proxy"><code class="docutils literal notranslate"><span class="pre">allow_fence_proxy()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.engine.phase.allow_vectorize"><code class="docutils literal notranslate"><span class="pre">allow_vectorize()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.engine.phase.allow_global_thread_synchronization"><code class="docutils literal notranslate"><span class="pre">allow_global_thread_synchronization()</span></code></a></li>

genindex.html

@@ -921,6 +921,8 @@ <h2>A</h2>
         <li><a href="autoapi/tilelang/autodd/index.html#tilelang.autodd.ParTaskManager.allow_larger">allow_larger (tilelang.autodd.ParTaskManager attribute)</a>
 </li>
         <li><a href="autoapi/tilelang/engine/phase/index.html#tilelang.engine.phase.allow_tma_and_warp_specialized">allow_tma_and_warp_specialized() (in module tilelang.engine.phase)</a>
+</li>
+        <li><a href="autoapi/tilelang/engine/phase/index.html#tilelang.engine.phase.allow_tma_lower">allow_tma_lower() (in module tilelang.engine.phase)</a>
 </li>
         <li><a href="autoapi/tilelang/engine/phase/index.html#tilelang.engine.phase.allow_vectorize">allow_vectorize() (in module tilelang.engine.phase)</a>
 </li>