Update docs

Author: LeiWang1999

_sources/autoapi/tilelang/carver/arch/cdna/index.rst.txt

@@ -30,6 +30,9 @@ Module Contents
    Bases: :py:obj:`tilelang.carver.arch.arch_base.TileDevice`
 
 
+   Represents the architecture of a computing device, capturing various hardware specifications.
+
+
    .. py:attribute:: target
 
 

_sources/autoapi/tilelang/carver/arch/cuda/index.rst.txt

@@ -92,6 +92,9 @@ Module Contents
    Bases: :py:obj:`tilelang.carver.arch.arch_base.TileDevice`
 
 
+   Represents the architecture of a computing device, capturing various hardware specifications.
+
+
    .. py:attribute:: target
 
 

_sources/autoapi/tilelang/carver/roller/hint/index.rst.txt

@@ -307,12 +307,12 @@ Module Contents
 
 
    .. py:property:: raxis_order
-      :type: tilelang.carver.roller.rasterization.List[int]
+      :type: List[int]
 
 
 
    .. py:property:: step
-      :type: tilelang.carver.roller.rasterization.List[int]
+      :type: List[int]
 
 
 

_sources/autoapi/tilelang/carver/roller/policy/tensorcore/index.rst.txt

@@ -30,15 +30,11 @@ Module Contents
 
 .. py:data:: logger
 
-.. py:class:: TensorCorePolicy(arch, tags = None)
+.. py:class:: TensorCorePolicy
 
    Bases: :py:obj:`tilelang.carver.roller.policy.default.DefaultPolicy`
 
 
-   Default Policy for fastdlight, a heuristic plan that tries to
-   minimize memory traffic and maximize parallelism.for BitBLAS Schedule.
-
-
    .. py:attribute:: wmma_k
       :type:  int
       :value: 16
@@ -65,67 +61,16 @@ Module Contents
 
    .. py:method:: infer_node_smem_usage(td, node)
 
-      Infers the shared memory usage of a node given a TileDict configuration.
-
-      :param td: The TileDict object containing the tile configuration.
-      :type td: TileDict
-      :param node: The node for which to infer the shared memory usage.
-      :type node: PrimFuncNode
-
-      :returns: The estimated amount of shared memory used by the node.
-      :rtype: int
-
-
 
    .. py:method:: get_node_reduce_step_candidates(node)
 
-      Calculates reduction step candidates for each reduction axis in a PrimFuncNode. General idea : use factor first, since it does not require extra boundary check. for large prime number, which is rare case, use power of 2.
-
-      :param node: The node for which to calculate reduction step candidates. It contains reduction axes (raxis)
-                   with their domains (dom.extent).
-      :type node: PrimFuncNode
-
-      :returns: A dictionary mapping axis variable names to lists of step candidates. For each axis in the node,
-                this function calculates possible step sizes. For axes with a large prime domain, it uses powers of 2
-                as step candidates; for others, it uses all factors of the domain.
-      :rtype: Dict[str, List[int]]
-
-
 
    .. py:method:: check_tile_shape_isvalid(td)
 
-      Checks if the tile shapes in the TileDict are valid for the nodes in this context.
-
-      Parameters:
-      - td (TileDict): The TileDict object containing tile shapes and other configurations.
-
-      Returns:
-      - bool: True if all tile shapes are valid, False otherwise.
-
-
 
    .. py:method:: compute_node_stride_map(node, td)
 
-      Computes the stride map for a given node based on the TileDict configuration.
-
-      :param node: The node for which to compute the stride map.
-      :type node: PrimFuncNode
-      :param td: The TileDict object containing the tile configuration.
-      :type td: TileDict
-
-      :returns: A tuple of dictionaries containing the output strides and tensor strides.
-      :rtype: Tuple[Dict, Dict]
-
-
 
    .. py:method:: plan_rasterization(td)
 
-      Plans the rasterization for the given TileDict. This function is not implemented yet.
-
-      :param td: The TileDict object to plan rasterization for.
-      :type td: TileDict
-
-      :raises RasterRationPlan: This function is not implemented yet.
-
-
 

_sources/autoapi/tilelang/jit/adapter/dlpack/index.rst.txt

@@ -20,8 +20,12 @@ Classes
 Module Contents
 ---------------
 
-.. py:class:: TorchDLPackKernelAdapter
+.. py:class:: TorchDLPackKernelAdapter(mod, params, result_idx)
 
    Bases: :py:obj:`tilelang.jit.adapter.base.BaseKernelAdapter`
 
 
+   Helper class that provides a standard way to create an ABC using
+   inheritance.
+
+

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

@@ -14,6 +14,10 @@ Functions
 
 .. autoapisummary::
 
+   tilelang.language.customize.region
+   tilelang.language.customize.buffer_to_tile_region
+   tilelang.language.customize.buffer_load_to_tile_region
+   tilelang.language.customize.buffer_region_to_tile_region
    tilelang.language.customize.atomic_add
    tilelang.language.customize.atomic_addx2
    tilelang.language.customize.atomic_addx4
@@ -26,6 +30,62 @@ Functions
 Module Contents
 ---------------
 
+.. py:function:: region(buffer, access_type, *args)
+
+   Create a memory region descriptor for tile operations.
+
+   :param buffer: The buffer to create a region for
+   :type buffer: tir.BufferLoad
+   :param access_type: Type of access - 'r' for read, 'w' for write, 'rw' for read-write
+   :type access_type: str
+   :param \*args: Extent expressions defining the region size
+   :type \*args: tir.PrimExpr
+
+   :returns: A region descriptor for tile operations
+   :rtype: tir.Call
+
+
+.. 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
+
+
+.. 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
+
+
+.. 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
+
+
 .. py:function:: atomic_add(dst, value)
 
    Perform an atomic addition operation.
@@ -54,14 +114,14 @@ Module Contents
 
 .. py:function:: atomic_addx4(dst, value)
 
-   Perform an atomic addition operation with double-width operands.
+   Perform an atomic addition operation with quad-width operands.
 
    :param dst: Destination buffer where the atomic addition will be performed
    :type dst: Buffer
-   :param value: Value to be atomically added (double-width)
+   :param value: Value to be atomically added (quad-width)
    :type value: PrimExpr
 
-   :returns: Handle to the double-width atomic addition operation
+   :returns: Handle to the quad-width atomic addition operation
    :rtype: PrimExpr
 
 

autoapi/tilelang/carver/arch/cdna/index.html

@@ -465,7 +465,7 @@ <h2>Classes<a class="headerlink" href="#classes" 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.carver.arch.cdna.CDNA" title="tilelang.carver.arch.cdna.CDNA"><code class="xref py py-obj docutils literal notranslate"><span class="pre">CDNA</span></code></a></p></td>
-<td><p></p></td>
+<td><p>Represents the architecture of a computing device, capturing various hardware specifications.</p></td>
 </tr>
 </tbody>
 </table>
@@ -502,6 +502,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dt class="sig sig-object py" id="tilelang.carver.arch.cdna.CDNA">
 <em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-prename descclassname"><span class="pre">tilelang.carver.arch.cdna.</span></span><span class="sig-name descname"><span class="pre">CDNA</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">target</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.arch.cdna.CDNA" title="Link to this definition">¶</a></dt>
 <dd><p>Bases: <a class="reference internal" href="../arch_base/index.html#tilelang.carver.arch.arch_base.TileDevice" title="tilelang.carver.arch.arch_base.TileDevice"><code class="xref py py-obj docutils literal notranslate"><span class="pre">tilelang.carver.arch.arch_base.TileDevice</span></code></a></p>
+<p>Represents the architecture of a computing device, capturing various hardware specifications.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><p><strong>target</strong> (<em>Union</em><em>[</em><em>tvm.target.Target</em><em>, </em><em>str</em><em>]</em>)</p>

autoapi/tilelang/carver/arch/cuda/index.html

@@ -489,7 +489,7 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">¶
 <td><p></p></td>
 </tr>
 <tr class="row-even"><td><p><a class="reference internal" href="#tilelang.carver.arch.cuda.CUDA" title="tilelang.carver.arch.cuda.CUDA"><code class="xref py py-obj docutils literal notranslate"><span class="pre">CUDA</span></code></a></p></td>
-<td><p></p></td>
+<td><p>Represents the architecture of a computing device, capturing various hardware specifications.</p></td>
 </tr>
 </tbody>
 </table>
@@ -686,6 +686,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dt class="sig sig-object py" id="tilelang.carver.arch.cuda.CUDA">
 <em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-prename descclassname"><span class="pre">tilelang.carver.arch.cuda.</span></span><span class="sig-name descname"><span class="pre">CUDA</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">target</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.arch.cuda.CUDA" title="Link to this definition">¶</a></dt>
 <dd><p>Bases: <a class="reference internal" href="../arch_base/index.html#tilelang.carver.arch.arch_base.TileDevice" title="tilelang.carver.arch.arch_base.TileDevice"><code class="xref py py-obj docutils literal notranslate"><span class="pre">tilelang.carver.arch.arch_base.TileDevice</span></code></a></p>
+<p>Represents the architecture of a computing device, capturing various hardware specifications.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><p><strong>target</strong> (<em>Union</em><em>[</em><em>tvm.target.Target</em><em>, </em><em>str</em><em>]</em>)</p>

autoapi/tilelang/carver/roller/hint/index.html

@@ -565,10 +565,10 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <span class="sig-name descname"><span class="pre">compute_strides_from_shape</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.hint.Stride.compute_strides_from_shape" title="Link to this definition">¶</a></dt>
 <dd><dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
-<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.List</em><em>[</em><em>int</em><em>]</em>)</p>
+<dd class="field-odd"><p><strong>shape</strong> (<em>List</em><em>[</em><em>int</em><em>]</em>)</p>
 </dd>
 <dt class="field-even">Return type<span class="colon">:</span></dt>
-<dd class="field-even"><p>tilelang.carver.roller.rasterization.List[int]</p>
+<dd class="field-even"><p>List[int]</p>
 </dd>
 </dl>
 </dd></dl>
@@ -578,7 +578,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <span class="sig-name descname"><span class="pre">compute_elements_from_shape</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.hint.Stride.compute_elements_from_shape" title="Link to this definition">¶</a></dt>
 <dd><dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
-<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.List</em><em>[</em><em>int</em><em>]</em>)</p>
+<dd class="field-odd"><p><strong>shape</strong> (<em>List</em><em>[</em><em>int</em><em>]</em>)</p>
 </dd>
 <dt class="field-even">Return type<span class="colon">:</span></dt>
 <dd class="field-even"><p>int</p>
@@ -677,7 +677,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <span class="sig-name descname"><span class="pre">get_tile</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">func</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.hint.TileDict.get_tile" 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>tilelang.carver.roller.rasterization.List[int]</p>
+<dd class="field-odd"><p>List[int]</p>
 </dd>
 </dl>
 </dd></dl>
@@ -946,20 +946,20 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 
 <dl class="py property">
 <dt class="sig sig-object py" id="tilelang.carver.roller.hint.Hint.raxis_order">
-<em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">raxis_order</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">tilelang.carver.roller.rasterization.List</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">]</span></span></em><a class="headerlink" href="#tilelang.carver.roller.hint.Hint.raxis_order" title="Link to this definition">¶</a></dt>
+<em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">raxis_order</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">List</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">]</span></span></em><a class="headerlink" href="#tilelang.carver.roller.hint.Hint.raxis_order" 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>tilelang.carver.roller.rasterization.List[int]</p>
+<dd class="field-odd"><p>List[int]</p>
 </dd>
 </dl>
 </dd></dl>
 
 <dl class="py property">
 <dt class="sig sig-object py" id="tilelang.carver.roller.hint.Hint.step">
-<em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">step</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">tilelang.carver.roller.rasterization.List</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">]</span></span></em><a class="headerlink" href="#tilelang.carver.roller.hint.Hint.step" title="Link to this definition">¶</a></dt>
+<em class="property"><span class="pre">property</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">step</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">List</span><span class="p"><span class="pre">[</span></span><span class="pre">int</span><span class="p"><span class="pre">]</span></span></em><a class="headerlink" href="#tilelang.carver.roller.hint.Hint.step" 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>tilelang.carver.roller.rasterization.List[int]</p>
+<dd class="field-odd"><p>List[int]</p>
 </dd>
 </dl>
 </dd></dl>