Skip to content

Commit 5e11d0d

Browse files
committed
Update docs
1 parent d69a650 commit 5e11d0d

File tree

11 files changed

+178
-42
lines changed

11 files changed

+178
-42
lines changed

_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: List[int]
310+
:type: tilelang.carver.roller.rasterization.List[int]
311311

312312

313313

314314
.. py:property:: step
315-
:type: List[int]
315+
:type: tilelang.carver.roller.rasterization.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: 0 additions & 5 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:: batch_size
2924
:type: int
3025
:value: 1

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

Lines changed: 18 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:: structure
2429
:type: Union[str, List[str]]
2530
:value: None
@@ -40,6 +45,19 @@ Module Contents
4045

4146
.. py:method:: get_hardware_aware_configs(arch = None, topk = 10)
4247
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+
4361

4462
.. py:method:: initialize_function()
4563

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

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

23-
.. py:class:: TorchDLPackKernelAdapter
23+
.. py:class:: TorchDLPackKernelAdapter(mod, params, result_idx)
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+

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>List</em><em>[</em><em>int</em><em>]</em>)</p>
568+
<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.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>List[int]</p>
571+
<dd class="field-even"><p>tilelang.carver.roller.rasterization.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>List</em><em>[</em><em>int</em><em>]</em>)</p>
581+
<dd class="field-odd"><p><strong>shape</strong> (<em>tilelang.carver.roller.rasterization.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>List[int]</p>
680+
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.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">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">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>
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>List[int]</p>
952+
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.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">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">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>
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>List[int]</p>
962+
<dd class="field-odd"><p>tilelang.carver.roller.rasterization.List[int]</p>
963963
</dd>
964964
</dl>
965965
</dd></dl>

0 commit comments

Comments
 (0)