Skip to content

Commit 50bbed2

Browse files
committed
Update docs
1 parent 54fb845 commit 50bbed2

File tree

17 files changed

+167
-44
lines changed

17 files changed

+167
-44
lines changed

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

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

3232

33+
Represents the architecture of a computing device, capturing various hardware specifications.
34+
35+
3336
.. py:attribute:: target
3437
3538

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,9 @@ 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+
9598
.. py:attribute:: target
9699
97100

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -307,12 +307,12 @@ Module Contents
307307
308308
309309
.. py:property:: raxis_order
310-
:type: tilelang.carver.roller.rasterization.List[int]
310+
:type: List[int]
311311

312312

313313

314314
.. py:property:: step
315-
:type: tilelang.carver.roller.rasterization.List[int]
315+
:type: List[int]
316316

317317

318318

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

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

3131
.. py:data:: logger
3232
33-
.. py:class:: TensorCorePolicy
33+
.. py:class:: TensorCorePolicy(arch, tags = None)
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+
3842
.. py:attribute:: wmma_k
3943
:type: int
4044
:value: 16
@@ -61,16 +65,67 @@ Module Contents
6165

6266
.. py:method:: infer_node_smem_usage(td, node)
6367
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+
6479

6580
.. py:method:: get_node_reduce_step_candidates(node)
6681
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+
6794

6895
.. py:method:: check_tile_shape_isvalid(td)
6996
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+
70106

71107
.. py:method:: compute_node_stride_map(node, td)
72108
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+
73120

74121
.. py:method:: plan_rasterization(td)
75122
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+
76131

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,11 @@ 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+
2328
.. py:attribute:: batch_size
2429
:type: int
2530
:value: 1

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

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -20,12 +20,8 @@ Classes
2020
Module Contents
2121
---------------
2222

23-
.. py:class:: TorchDLPackKernelAdapter(mod, params, result_idx)
23+
.. py:class:: TorchDLPackKernelAdapter
2424
2525
Bases: :py:obj:`tilelang.jit.adapter.base.BaseKernelAdapter`
2626

2727

28-
Helper class that provides a standard way to create an ABC using
29-
inheritance.
30-
31-

_sources/autoapi/tilelang/transform/pass_config/index.rst.txt

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -68,6 +68,14 @@ Module Contents
6868
:type: Disable fast math optimization. Default
6969

7070

71+
.. py:attribute:: TL_PTXAS_REGISTER_USAGE_LEVEL
72+
:value: 'tl.ptxas_register_usage_level'
73+
74+
75+
The PTXAS register usage level in [0, 10], which controls the
76+
aggressiveness of optimizations that affect register usage. Default: None
77+
78+
7179
.. py:attribute:: TL_ENABLE_PTXAS_VERBOSE_OUTPUT
7280
:value: 'tl.enable_ptxas_verbose_output'
7381

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -465,7 +465,7 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">¶
465465
<table class="autosummary longtable docutils align-default">
466466
<tbody>
467467
<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>
468-
<td><p></p></td>
468+
<td><p>Represents the architecture of a computing device, capturing various hardware specifications.</p></td>
469469
</tr>
470470
</tbody>
471471
</table>
@@ -502,6 +502,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
502502
<dt class="sig sig-object py" id="tilelang.carver.arch.cdna.CDNA">
503503
<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>
504504
<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>
505+
<p>Represents the architecture of a computing device, capturing various hardware specifications.</p>
505506
<dl class="field-list simple">
506507
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
507508
<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

Lines changed: 2 additions & 1 deletion
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></p></td>
492+
<td><p>Represents the architecture of a computing device, capturing various hardware specifications.</p></td>
493493
</tr>
494494
</tbody>
495495
</table>
@@ -686,6 +686,7 @@ <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>
689690
<dl class="field-list simple">
690691
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
691692
<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

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -565,10 +565,10 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
565565
<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>
566566
<dd><dl class="field-list simple">
567567
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
568-
<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.List</em><em>[</em><em>int</em><em>]</em>)</p>
568+
<dd class="field-odd"><p><strong>shape</strong> (<em>List</em><em>[</em><em>int</em><em>]</em>)</p>
569569
</dd>
570570
<dt class="field-even">Return type<span class="colon">:</span></dt>
571-
<dd class="field-even"><p>tilelang.carver.roller.rasterization.List[int]</p>
571+
<dd class="field-even"><p>List[int]</p>
572572
</dd>
573573
</dl>
574574
</dd></dl>
@@ -578,7 +578,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
578578
<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>
579579
<dd><dl class="field-list simple">
580580
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
581-
<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.List</em><em>[</em><em>int</em><em>]</em>)</p>
581+
<dd class="field-odd"><p><strong>shape</strong> (<em>List</em><em>[</em><em>int</em><em>]</em>)</p>
582582
</dd>
583583
<dt class="field-even">Return type<span class="colon">:</span></dt>
584584
<dd class="field-even"><p>int</p>
@@ -677,7 +677,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
677677
<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>
678678
<dd><dl class="field-list simple">
679679
<dt class="field-odd">Return type<span class="colon">:</span></dt>
680-
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.List[int]</p>
680+
<dd class="field-odd"><p>List[int]</p>
681681
</dd>
682682
</dl>
683683
</dd></dl>
@@ -946,20 +946,20 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
946946

947947
<dl class="py property">
948948
<dt class="sig sig-object py" id="tilelang.carver.roller.hint.Hint.raxis_order">
949-
<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>
949+
<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>
950950
<dd><dl class="field-list simple">
951951
<dt class="field-odd">Return type<span class="colon">:</span></dt>
952-
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.List[int]</p>
952+
<dd class="field-odd"><p>List[int]</p>
953953
</dd>
954954
</dl>
955955
</dd></dl>
956956

957957
<dl class="py property">
958958
<dt class="sig sig-object py" id="tilelang.carver.roller.hint.Hint.step">
959-
<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>
959+
<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>
960960
<dd><dl class="field-list simple">
961961
<dt class="field-odd">Return type<span class="colon">:</span></dt>
962-
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.List[int]</p>
962+
<dd class="field-odd"><p>List[int]</p>
963963
</dd>
964964
</dl>
965965
</dd></dl>

0 commit comments

Comments
 (0)