Update docs

Author: github-actions[bot]

_sources/autoapi/tilelang/language/utils/index.rst.txt

@@ -10,7 +10,6 @@ Functions
 .. autoapisummary::
 
    tilelang.language.utils.region
-   tilelang.language.utils.buffer_to_tile_region
    tilelang.language.utils.buffer_load_to_tile_region
    tilelang.language.utils.buffer_region_to_tile_region
    tilelang.language.utils.index_to_coordinates
@@ -22,63 +21,17 @@ Module Contents
 
 .. py:function:: region(buffer, access_type, *args)
 
-   Create a tile memory-region descriptor for a BufferLoad.
-
-   Maps access_type ('r', 'w', 'rw') to the numeric codes expected by the `tl.region` intrinsic
-   (1, 2, 3 respectively) and returns a tir.Call representing the region with the provided extents.
-
-   :param buffer: The BufferLoad that identifies the underlying buffer and indices.
-   :type buffer: tir.BufferLoad
-   :param access_type: One of 'r', 'w', or 'rw' indicating read, write, or read-write access.
-   :type access_type: str
-   :param \*args: Extent expressions for each region dimension.
-   :type \*args: tir.PrimExpr
-
-   :returns: A call to the `tl.region` intrinsic describing the memory region.
-   :rtype: tir.Call
-
-   :raises KeyError: If access_type is not one of 'r', 'w', or 'rw'.
-
-
-.. py:function:: buffer_to_tile_region(buffer, access_type)
-
-   Convert a TVM buffer to a tile region descriptor.
-
-   :param buffer: The buffer to convert
-   :type buffer: tir.Buffer
-   :param access_type: Type of access - 'r' for read, 'w' for write, 'rw' for read-write
-   :type access_type: str
-
-   :returns: A region descriptor covering the entire buffer
-   :rtype: tir.Call
+   Create a tl.region call for a BufferLoad and extents.
 
 
 .. py:function:: buffer_load_to_tile_region(load, access_type, extents)
 
-   Convert a buffer load operation to a tile region descriptor.
-
-   :param load: The buffer load operation
-   :type load: tir.BufferLoad
-   :param access_type: Type of access - 'r' for read, 'w' for write, 'rw' for read-write
-   :type access_type: str
-   :param extents: List of expressions defining the region size
-   :type extents: List[tir.PrimExpr]
-
-   :returns: A region descriptor for the loaded area
-   :rtype: tir.Call
+   Convert a BufferLoad to a tl.region call with explicit extents.
 
 
 .. py:function:: buffer_region_to_tile_region(buffer_region, access_type, extents)
 
-   Convert a buffer region to a tile region descriptor.
-
-   :param buffer_region: The buffer region to convert
-   :type buffer_region: tir.BufferRegion
-   :param access_type: Type of access - 'r' for read, 'w' for write, 'rw' for read-write
-   :type access_type: str
-
-   :returns: A region descriptor for the specified buffer region
-   :rtype: tir.Call
+   Clamp extents and return a tl.region call.
 
 
 .. py:function:: index_to_coordinates(index, shape)

_sources/autoapi/tilelang/tileop/gemm/gemm_base/index.rst.txt

@@ -152,6 +152,11 @@ Module Contents
 
 
 
+   .. py:property:: mbar
+      :type: tvm.tir.Buffer
+
+
+
    .. py:property:: C_coords
 
 

_sources/autoapi/tilelang/utils/language/index.rst.txt

@@ -125,7 +125,7 @@ Module Contents
    :raises AssertionError: If the module contains more than one global function.
 
 
-.. py:function:: get_buffer_region_from_load(buffer_load)
+.. py:function:: get_buffer_region_from_load(buffer_load, extents = None)
 
    Get the buffer region from a buffer load.
 
@@ -134,14 +134,12 @@ Module Contents
    convert load to region
 
 
-.. py:function:: to_buffer_region(obj)
+.. py:function:: to_buffer_region(obj, access_type = 'rw', extents = None)
 
-   Convert Buffer/BufferRegion/BufferLoad to a BufferRegion.
+   Convert to/from the tl.region representation.
 
-   - Buffer -> full-region BufferRegion covering entire shape
-   - BufferRegion -> returned as-is
-   - BufferLoad -> best-effort convert via get_buffer_region_from_load;
-     if scalar, fall back to 1-sized ranges at given indices
+   - Buffer/BufferLoad/BufferRegion -> returns a tl.region call (PrimExpr)
+   - tl.region Call -> returns the decoded BufferRegion for analysis
 
 
 .. py:function:: retrieve_shape(obj)

autoapi/tilelang/language/utils/index.html

@@ -507,21 +507,18 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 <table class="autosummary longtable docutils align-default">
 <tbody>
 <tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.utils.region" title="tilelang.language.utils.region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">region</span></code></a>(buffer, access_type, *args)</p></td>
-<td><p>Create a tile memory-region descriptor for a BufferLoad.</p></td>
+<td><p>Create a tl.region call for a BufferLoad and extents.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.utils.buffer_to_tile_region" title="tilelang.language.utils.buffer_to_tile_region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">buffer_to_tile_region</span></code></a>(buffer, access_type)</p></td>
-<td><p>Convert a TVM buffer to a tile region descriptor.</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.utils.buffer_load_to_tile_region" title="tilelang.language.utils.buffer_load_to_tile_region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">buffer_load_to_tile_region</span></code></a>(load, access_type, extents)</p></td>
+<td><p>Convert a BufferLoad to a tl.region call with explicit extents.</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.utils.buffer_load_to_tile_region" title="tilelang.language.utils.buffer_load_to_tile_region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">buffer_load_to_tile_region</span></code></a>(load, access_type, extents)</p></td>
-<td><p>Convert a buffer load operation to a tile region descriptor.</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.utils.buffer_region_to_tile_region" title="tilelang.language.utils.buffer_region_to_tile_region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">buffer_region_to_tile_region</span></code></a>(buffer_region, ...)</p></td>
+<td><p>Clamp extents and return a tl.region call.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.utils.buffer_region_to_tile_region" title="tilelang.language.utils.buffer_region_to_tile_region"><code class="xref py py-obj docutils literal notranslate"><span class="pre">buffer_region_to_tile_region</span></code></a>(buffer_region, ...)</p></td>
-<td><p>Convert a buffer region to a tile region descriptor.</p></td>
-</tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.utils.index_to_coordinates" title="tilelang.language.utils.index_to_coordinates"><code class="xref py py-obj docutils literal notranslate"><span class="pre">index_to_coordinates</span></code></a>(index, shape)</p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.utils.index_to_coordinates" title="tilelang.language.utils.index_to_coordinates"><code class="xref py py-obj docutils literal notranslate"><span class="pre">index_to_coordinates</span></code></a>(index, shape)</p></td>
 <td><p>Convert a flat (linear) index into multi-dimensional coordinates for a given shape.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.language.utils.linear_index" title="tilelang.language.utils.linear_index"><code class="xref py py-obj docutils literal notranslate"><span class="pre">linear_index</span></code></a>(*args)</p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.language.utils.linear_index" title="tilelang.language.utils.linear_index"><code class="xref py py-obj docutils literal notranslate"><span class="pre">linear_index</span></code></a>(*args)</p></td>
 <td><p>Compute a flat (linear) index from multi-dimensional coordinates and strides.</p></td>
 </tr>
 </tbody>
@@ -533,88 +530,45 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dl class="py function">
 <dt class="sig sig-object py" id="tilelang.language.utils.region">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.utils.</span></span><span class="sig-name descname"><span class="pre">region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">buffer</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span></em>, <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.utils.region" title="Link to this definition">¶</a></dt>
-<dd><p>Create a tile memory-region descriptor for a BufferLoad.</p>
-<p>Maps access_type (‘r’, ‘w’, ‘rw’) to the numeric codes expected by the <cite>tl.region</cite> intrinsic
-(1, 2, 3 respectively) and returns a tir.Call representing the region with the provided extents.</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>buffer</strong> (<em>tir.BufferLoad</em>) – The BufferLoad that identifies the underlying buffer and indices.</p></li>
-<li><p><strong>access_type</strong> (<em>str</em>) – One of ‘r’, ‘w’, or ‘rw’ indicating read, write, or read-write access.</p></li>
-<li><p><strong>*args</strong> (<em>tir.PrimExpr</em>) – Extent expressions for each region dimension.</p></li>
-</ul>
-</dd>
-<dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A call to the <cite>tl.region</cite> intrinsic describing the memory region.</p>
-</dd>
-<dt class="field-odd">Return type<span class="colon">:</span></dt>
-<dd class="field-odd"><p>tir.Call</p>
-</dd>
-<dt class="field-even">Raises<span class="colon">:</span></dt>
-<dd class="field-even"><p><strong>KeyError</strong> – If access_type is not one of ‘r’, ‘w’, or ‘rw’.</p>
-</dd>
-</dl>
-</dd></dl>
-
-<dl class="py function">
-<dt class="sig sig-object py" id="tilelang.language.utils.buffer_to_tile_region">
-<span class="sig-prename descclassname"><span class="pre">tilelang.language.utils.</span></span><span class="sig-name descname"><span class="pre">buffer_to_tile_region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">buffer</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.utils.buffer_to_tile_region" title="Link to this definition">¶</a></dt>
-<dd><p>Convert a TVM buffer to a tile region descriptor.</p>
+<dd><p>Create a tl.region call for a BufferLoad and extents.</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>buffer</strong> (<em>tir.Buffer</em>) – The buffer to convert</p></li>
-<li><p><strong>access_type</strong> (<em>str</em>) – Type of access - ‘r’ for read, ‘w’ for write, ‘rw’ for read-write</p></li>
+<li><p><strong>buffer</strong> (<em>tvm.tir.BufferLoad</em>)</p></li>
+<li><p><strong>access_type</strong> (<em>str</em>)</p></li>
+<li><p><strong>args</strong> (<em>tvm.tir.PrimExpr</em>)</p></li>
 </ul>
 </dd>
-<dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A region descriptor covering the entire buffer</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.utils.buffer_load_to_tile_region">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.utils.</span></span><span class="sig-name descname"><span class="pre">buffer_load_to_tile_region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">load</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">extents</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.utils.buffer_load_to_tile_region" title="Link to this definition">¶</a></dt>
-<dd><p>Convert a buffer load operation to a tile region descriptor.</p>
+<dd><p>Convert a BufferLoad to a tl.region call with explicit extents.</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>load</strong> (<em>tir.BufferLoad</em>) – The buffer load operation</p></li>
-<li><p><strong>access_type</strong> (<em>str</em>) – Type of access - ‘r’ for read, ‘w’ for write, ‘rw’ for read-write</p></li>
-<li><p><strong>extents</strong> (<em>List</em><em>[</em><em>tir.PrimExpr</em><em>]</em>) – List of expressions defining the region size</p></li>
+<li><p><strong>load</strong> (<em>tvm.tir.BufferLoad</em>)</p></li>
+<li><p><strong>access_type</strong> (<em>str</em>)</p></li>
+<li><p><strong>extents</strong> (<em>list</em><em>[</em><em>tvm.tir.PrimExpr</em><em>]</em>)</p></li>
 </ul>
 </dd>
-<dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A region descriptor for the loaded area</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.utils.buffer_region_to_tile_region">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.utils.</span></span><span class="sig-name descname"><span class="pre">buffer_region_to_tile_region</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">buffer_region</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">access_type</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">extents</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.utils.buffer_region_to_tile_region" title="Link to this definition">¶</a></dt>
-<dd><p>Convert a buffer region to a tile region descriptor.</p>
+<dd><p>Clamp extents and return a tl.region call.</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>buffer_region</strong> (<em>tir.BufferRegion</em>) – The buffer region to convert</p></li>
-<li><p><strong>access_type</strong> (<em>str</em>) – Type of access - ‘r’ for read, ‘w’ for write, ‘rw’ for read-write</p></li>
+<li><p><strong>buffer_region</strong> (<em>tvm.tir.BufferRegion</em>)</p></li>
+<li><p><strong>access_type</strong> (<em>str</em>)</p></li>
 <li><p><strong>extents</strong> (<em>list</em><em>[</em><em>tvm.tir.PrimExpr</em><em>]</em>)</p></li>
 </ul>
 </dd>
-<dt class="field-even">Returns<span class="colon">:</span></dt>
-<dd class="field-even"><p>A region descriptor for the specified buffer region</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>
 
@@ -740,7 +694,6 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <li><a class="reference internal" href="#functions">Functions</a></li>
 <li><a class="reference internal" href="#module-contents">Module Contents</a><ul>
 <li><a class="reference internal" href="#tilelang.language.utils.region"><code class="docutils literal notranslate"><span class="pre">region()</span></code></a></li>
-<li><a class="reference internal" href="#tilelang.language.utils.buffer_to_tile_region"><code class="docutils literal notranslate"><span class="pre">buffer_to_tile_region()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.utils.buffer_load_to_tile_region"><code class="docutils literal notranslate"><span class="pre">buffer_load_to_tile_region()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.utils.buffer_region_to_tile_region"><code class="docutils literal notranslate"><span class="pre">buffer_region_to_tile_region()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.language.utils.index_to_coordinates"><code class="docutils literal notranslate"><span class="pre">index_to_coordinates()</span></code></a></li>

autoapi/tilelang/tileop/gemm/gemm_base/index.html

@@ -805,6 +805,16 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 </dl>
 </dd></dl>
 
+<dl class="py property">
+<dt class="sig sig-object py" id="tilelang.tileop.gemm.gemm_base.GemmBase.mbar">
+<em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">mbar</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">tvm.tir.Buffer</span></em><a class="headerlink" href="#tilelang.tileop.gemm.gemm_base.GemmBase.mbar" title="Link to this definition">¶</a></dt>
+<dd><dl class="field-list simple">
+<dt class="field-odd">Return type<span class="colon">:</span></dt>
+<dd class="field-odd"><p>tvm.tir.Buffer</p>
+</dd>
+</dl>
+</dd></dl>
+
 <dl class="py property">
 <dt class="sig sig-object py" id="tilelang.tileop.gemm.gemm_base.GemmBase.C_coords">
 <em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">C_coords</span></span><a class="headerlink" href="#tilelang.tileop.gemm.gemm_base.GemmBase.C_coords" title="Link to this definition">¶</a></dt>
@@ -938,6 +948,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.wg_wait"><code class="docutils literal notranslate"><span class="pre">GemmBase.wg_wait</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.policy"><code class="docutils literal notranslate"><span class="pre">GemmBase.policy</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.mbarptr"><code class="docutils literal notranslate"><span class="pre">GemmBase.mbarptr</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.mbar"><code class="docutils literal notranslate"><span class="pre">GemmBase.mbar</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.C_coords"><code class="docutils literal notranslate"><span class="pre">GemmBase.C_coords</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.get_region_base_offsets"><code class="docutils literal notranslate"><span class="pre">GemmBase.get_region_base_offsets()</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.tileop.gemm.gemm_base.GemmBase.A_base_offsets"><code class="docutils literal notranslate"><span class="pre">GemmBase.A_base_offsets</span></code></a></li>