Skip to content

Commit d69a650

Browse files
committed
Update docs
1 parent 6ed8cb4 commit d69a650

File tree

7 files changed

+21
-155
lines changed

7 files changed

+21
-155
lines changed

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

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -92,9 +92,6 @@ Module Contents
9292
Bases: :py:obj:`tilelang.carver.arch.arch_base.TileDevice`
9393

9494

95-
Represents the architecture of a computing device, capturing various hardware specifications.
96-
97-
9895
.. py:attribute:: target
9996
10097

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

Lines changed: 1 addition & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -30,15 +30,11 @@ Module Contents
3030

3131
.. py:data:: logger
3232
33-
.. py:class:: TensorCorePolicy(arch, tags = None)
33+
.. py:class:: TensorCorePolicy
3434
3535
Bases: :py:obj:`tilelang.carver.roller.policy.default.DefaultPolicy`
3636

3737

38-
Default Policy for fastdlight, a heuristic plan that tries to
39-
minimize memory traffic and maximize parallelism.for BitBLAS Schedule.
40-
41-
4238
.. py:attribute:: wmma_k
4339
:type: int
4440
:value: 16
@@ -65,67 +61,16 @@ Module Contents
6561

6662
.. py:method:: infer_node_smem_usage(td, node)
6763
68-
Infers the shared memory usage of a node given a TileDict configuration.
69-
70-
:param td: The TileDict object containing the tile configuration.
71-
:type td: TileDict
72-
:param node: The node for which to infer the shared memory usage.
73-
:type node: PrimFuncNode
74-
75-
:returns: The estimated amount of shared memory used by the node.
76-
:rtype: int
77-
78-
7964
8065
.. py:method:: get_node_reduce_step_candidates(node)
8166
82-
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.
83-
84-
:param node: The node for which to calculate reduction step candidates. It contains reduction axes (raxis)
85-
with their domains (dom.extent).
86-
:type node: PrimFuncNode
87-
88-
:returns: A dictionary mapping axis variable names to lists of step candidates. For each axis in the node,
89-
this function calculates possible step sizes. For axes with a large prime domain, it uses powers of 2
90-
as step candidates; for others, it uses all factors of the domain.
91-
:rtype: Dict[str, List[int]]
92-
93-
9467
9568
.. py:method:: check_tile_shape_isvalid(td)
9669
97-
Checks if the tile shapes in the TileDict are valid for the nodes in this context.
98-
99-
Parameters:
100-
- td (TileDict): The TileDict object containing tile shapes and other configurations.
101-
102-
Returns:
103-
- bool: True if all tile shapes are valid, False otherwise.
104-
105-
10670
10771
.. py:method:: compute_node_stride_map(node, td)
10872
109-
Computes the stride map for a given node based on the TileDict configuration.
110-
111-
:param node: The node for which to compute the stride map.
112-
:type node: PrimFuncNode
113-
:param td: The TileDict object containing the tile configuration.
114-
:type td: TileDict
115-
116-
:returns: A tuple of dictionaries containing the output strides and tensor strides.
117-
:rtype: Tuple[Dict, Dict]
118-
119-
12073
12174
.. py:method:: plan_rasterization(td)
12275
123-
Plans the rasterization for the given TileDict. This function is not implemented yet.
124-
125-
:param td: The TileDict object to plan rasterization for.
126-
:type td: TileDict
127-
128-
:raises RasterRationPlan: This function is not implemented yet.
129-
130-
13176

_sources/autoapi/tilelang/carver/template/general_reduce/index.rst.txt

Lines changed: 0 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -20,11 +20,6 @@ Module Contents
2020
Bases: :py:obj:`tilelang.carver.template.base.BaseTemplate`
2121

2222

23-
Base class template for hardware-aware configurations.
24-
This serves as an abstract base class (ABC) that defines the structure
25-
for subclasses implementing hardware-specific optimizations.
26-
27-
2823
.. py:attribute:: structure
2924
:type: Union[str, List[str]]
3025
:value: None
@@ -45,19 +40,6 @@ Module Contents
4540

4641
.. py:method:: get_hardware_aware_configs(arch = None, topk = 10)
4742
48-
Abstract method that must be implemented by subclasses.
49-
It should return a list of hardware-aware configurations (hints)
50-
based on the specified architecture.
51-
52-
:param arch: The target architecture. Defaults to None.
53-
:type arch: TileDevice, optional
54-
:param topk: Number of top configurations to return. Defaults to 10.
55-
:type topk: int, optional
56-
57-
:returns: A list of recommended hardware-aware configurations.
58-
:rtype: List[Hint]
59-
60-
6143
6244
.. py:method:: initialize_function()
6345

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -489,7 +489,7 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">¶
489489
<td><p></p></td>
490490
</tr>
491491
<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>
492-
<td><p>Represents the architecture of a computing device, capturing various hardware specifications.</p></td>
492+
<td><p></p></td>
493493
</tr>
494494
</tbody>
495495
</table>
@@ -686,7 +686,6 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
686686
<dt class="sig sig-object py" id="tilelang.carver.arch.cuda.CUDA">
687687
<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>
688688
<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>
689-
<p>Represents the architecture of a computing device, capturing various hardware specifications.</p>
690689
<dl class="field-list simple">
691690
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
692691
<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/policy/tensorcore/index.html

Lines changed: 12 additions & 60 deletions
Original file line numberDiff line numberDiff line change
@@ -478,7 +478,7 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">¶
478478
<table class="autosummary longtable docutils align-default">
479479
<tbody>
480480
<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy" title="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy"><code class="xref py py-obj docutils literal notranslate"><span class="pre">TensorCorePolicy</span></code></a></p></td>
481-
<td><p>Default Policy for fastdlight, a heuristic plan that tries to</p></td>
481+
<td><p></p></td>
482482
</tr>
483483
</tbody>
484484
</table>
@@ -493,18 +493,8 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
493493

494494
<dl class="py class">
495495
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy">
496-
<em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-prename descclassname"><span class="pre">tilelang.carver.roller.policy.tensorcore.</span></span><span class="sig-name descname"><span class="pre">TensorCorePolicy</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">arch</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">tags</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.carver.roller.policy.tensorcore.TensorCorePolicy" title="Link to this definition"></a></dt>
496+
<em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-prename descclassname"><span class="pre">tilelang.carver.roller.policy.tensorcore.</span></span><span class="sig-name descname"><span class="pre">TensorCorePolicy</span></span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy" title="Link to this definition"></a></dt>
497497
<dd><p>Bases: <a class="reference internal" href="../default/index.html#tilelang.carver.roller.policy.default.DefaultPolicy" title="tilelang.carver.roller.policy.default.DefaultPolicy"><code class="xref py py-obj docutils literal notranslate"><span class="pre">tilelang.carver.roller.policy.default.DefaultPolicy</span></code></a></p>
498-
<p>Default Policy for fastdlight, a heuristic plan that tries to
499-
minimize memory traffic and maximize parallelism.for BitBLAS Schedule.</p>
500-
<dl class="field-list simple">
501-
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
502-
<dd class="field-odd"><ul class="simple">
503-
<li><p><strong>arch</strong> (<em>tilelang.carver.arch.TileDevice</em>)</p></li>
504-
<li><p><strong>tags</strong> (<em>Optional</em><em>[</em><em>Dict</em><em>]</em>)</p></li>
505-
</ul>
506-
</dd>
507-
</dl>
508498
<dl class="py attribute">
509499
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.wmma_k">
510500
<span class="sig-name descname"><span class="pre">wmma_k</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">int</span></em><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">16</span></em><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.wmma_k" title="Link to this definition"></a></dt>
@@ -528,52 +518,25 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
528518
<dl class="py method">
529519
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.infer_node_smem_usage">
530520
<span class="sig-name descname"><span class="pre">infer_node_smem_usage</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">td</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">node</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.infer_node_smem_usage" title="Link to this definition"></a></dt>
531-
<dd><p>Infers the shared memory usage of a node given a TileDict configuration.</p>
532-
<dl class="field-list simple">
521+
<dd><dl class="field-list simple">
533522
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
534523
<dd class="field-odd"><ul class="simple">
535-
<li><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>TileDict</em></a>) – The TileDict object containing the tile configuration.</p></li>
536-
<li><p><strong>node</strong> (<a class="reference internal" href="../../node/index.html#tilelang.carver.roller.node.PrimFuncNode" title="tilelang.carver.roller.node.PrimFuncNode"><em>PrimFuncNode</em></a>) – The node for which to infer the shared memory usage.</p></li>
524+
<li><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>tilelang.carver.roller.hint.TileDict</em></a>)</p></li>
525+
<li><p><strong>node</strong> (<a class="reference internal" href="../../node/index.html#tilelang.carver.roller.node.PrimFuncNode" title="tilelang.carver.roller.node.PrimFuncNode"><em>tilelang.carver.roller.node.PrimFuncNode</em></a>)</p></li>
537526
</ul>
538527
</dd>
539-
<dt class="field-even">Returns<span class="colon">:</span></dt>
540-
<dd class="field-even"><p>The estimated amount of shared memory used by the node.</p>
541-
</dd>
542-
<dt class="field-odd">Return type<span class="colon">:</span></dt>
543-
<dd class="field-odd"><p>int</p>
544-
</dd>
545528
</dl>
546529
</dd></dl>
547530

548531
<dl class="py method">
549532
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.get_node_reduce_step_candidates">
550533
<span class="sig-name descname"><span class="pre">get_node_reduce_step_candidates</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">node</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.get_node_reduce_step_candidates" title="Link to this definition"></a></dt>
551-
<dd><p>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.</p>
552-
<dl class="field-list simple">
553-
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
554-
<dd class="field-odd"><p><strong>node</strong> (<a class="reference internal" href="../../node/index.html#tilelang.carver.roller.node.PrimFuncNode" title="tilelang.carver.roller.node.PrimFuncNode"><em>PrimFuncNode</em></a>) – The node for which to calculate reduction step candidates. It contains reduction axes (raxis)
555-
with their domains (dom.extent).</p>
556-
</dd>
557-
<dt class="field-even">Returns<span class="colon">:</span></dt>
558-
<dd class="field-even"><p>A dictionary mapping axis variable names to lists of step candidates. For each axis in the node,
559-
this function calculates possible step sizes. For axes with a large prime domain, it uses powers of 2
560-
as step candidates; for others, it uses all factors of the domain.</p>
561-
</dd>
562-
<dt class="field-odd">Return type<span class="colon">:</span></dt>
563-
<dd class="field-odd"><p>Dict[str, List[int]]</p>
564-
</dd>
565-
</dl>
566-
</dd></dl>
534+
<dd></dd></dl>
567535

568536
<dl class="py method">
569537
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.check_tile_shape_isvalid">
570538
<span class="sig-name descname"><span class="pre">check_tile_shape_isvalid</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">td</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.check_tile_shape_isvalid" title="Link to this definition"></a></dt>
571-
<dd><p>Checks if the tile shapes in the TileDict are valid for the nodes in this context.</p>
572-
<p>Parameters:
573-
- td (TileDict): The TileDict object containing tile shapes and other configurations.</p>
574-
<p>Returns:
575-
- bool: True if all tile shapes are valid, False otherwise.</p>
576-
<dl class="field-list simple">
539+
<dd><dl class="field-list simple">
577540
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
578541
<dd class="field-odd"><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>tilelang.carver.roller.hint.TileDict</em></a>)</p>
579542
</dd>
@@ -583,33 +546,22 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
583546
<dl class="py method">
584547
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.compute_node_stride_map">
585548
<span class="sig-name descname"><span class="pre">compute_node_stride_map</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">node</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">td</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.compute_node_stride_map" title="Link to this definition"></a></dt>
586-
<dd><p>Computes the stride map for a given node based on the TileDict configuration.</p>
587-
<dl class="field-list simple">
549+
<dd><dl class="field-list simple">
588550
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
589551
<dd class="field-odd"><ul class="simple">
590-
<li><p><strong>node</strong> (<a class="reference internal" href="../../node/index.html#tilelang.carver.roller.node.PrimFuncNode" title="tilelang.carver.roller.node.PrimFuncNode"><em>PrimFuncNode</em></a>) – The node for which to compute the stride map.</p></li>
591-
<li><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>TileDict</em></a>) – The TileDict object containing the tile configuration.</p></li>
552+
<li><p><strong>node</strong> (<a class="reference internal" href="../../node/index.html#tilelang.carver.roller.node.PrimFuncNode" title="tilelang.carver.roller.node.PrimFuncNode"><em>tilelang.carver.roller.node.PrimFuncNode</em></a>)</p></li>
553+
<li><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>tilelang.carver.roller.hint.TileDict</em></a>)</p></li>
592554
</ul>
593555
</dd>
594-
<dt class="field-even">Returns<span class="colon">:</span></dt>
595-
<dd class="field-even"><p>A tuple of dictionaries containing the output strides and tensor strides.</p>
596-
</dd>
597-
<dt class="field-odd">Return type<span class="colon">:</span></dt>
598-
<dd class="field-odd"><p>Tuple[Dict, Dict]</p>
599-
</dd>
600556
</dl>
601557
</dd></dl>
602558

603559
<dl class="py method">
604560
<dt class="sig sig-object py" id="tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.plan_rasterization">
605561
<span class="sig-name descname"><span class="pre">plan_rasterization</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">td</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.carver.roller.policy.tensorcore.TensorCorePolicy.plan_rasterization" title="Link to this definition"></a></dt>
606-
<dd><p>Plans the rasterization for the given TileDict. This function is not implemented yet.</p>
607-
<dl class="field-list simple">
562+
<dd><dl class="field-list simple">
608563
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
609-
<dd class="field-odd"><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>TileDict</em></a>) – The TileDict object to plan rasterization for.</p>
610-
</dd>
611-
<dt class="field-even">Raises<span class="colon">:</span></dt>
612-
<dd class="field-even"><p><strong>RasterRationPlan</strong> – This function is not implemented yet.</p>
564+
<dd class="field-odd"><p><strong>td</strong> (<a class="reference internal" href="../../hint/index.html#tilelang.carver.roller.hint.TileDict" title="tilelang.carver.roller.hint.TileDict"><em>tilelang.carver.roller.hint.TileDict</em></a>)</p>
613565
</dd>
614566
</dl>
615567
</dd></dl>

0 commit comments

Comments
 (0)