Skip to content

Commit 76d6668

Browse files
github-actions[bot]github-actions[bot]
committed
Update docs
1 parent d010279 commit 76d6668

File tree

170 files changed

+1725
-277
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

170 files changed

+1725
-277
lines changed

‎_sources/autoapi/tilelang/language/annotations/index.rst.txt‎

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
tilelang.language.annotations
2+
=============================
3+
4+
.. py:module:: tilelang.language.annotations
5+
6+
.. autoapi-nested-parse::
7+
8+
Annotation helpers exposed on the TileLang language surface.
9+
10+
11+
12+
Functions
13+
---------
14+
15+
.. autoapisummary::
16+
17+
tilelang.language.annotations.use_swizzle
18+
tilelang.language.annotations.annotate_layout
19+
tilelang.language.annotations.annotate_safe_value
20+
tilelang.language.annotations.annotate_l2_hit_ratio
21+
22+
23+
Module Contents
24+
---------------
25+
26+
.. py:function:: use_swizzle(panel_size, order = 'row', enable = True)
27+
28+
Annotate a kernel to use a specific threadblock swizzle pattern.
29+
30+
31+
.. py:function:: annotate_layout(layout_map)
32+
33+
Annotate the layout of the buffer.
34+
35+
36+
.. py:function:: annotate_safe_value(safe_value_map)
37+
38+
Annotate the safe value of the buffer.
39+
40+
41+
.. py:function:: annotate_l2_hit_ratio(l2_hit_ratio_map)
42+
43+
Annotate the L2 hit ratio of the buffer.
44+
45+

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

Lines changed: 2 additions & 99 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@ Submodules
1616
:maxdepth: 1
1717

1818
/autoapi/tilelang/language/allocate/index
19+
/autoapi/tilelang/language/annotations/index
1920
/autoapi/tilelang/language/atomic/index
2021
/autoapi/tilelang/language/builtin/index
2122
/autoapi/tilelang/language/copy/index
@@ -35,6 +36,7 @@ Submodules
3536
/autoapi/tilelang/language/print/index
3637
/autoapi/tilelang/language/proxy/index
3738
/autoapi/tilelang/language/reduce/index
39+
/autoapi/tilelang/language/symbolics/index
3840
/autoapi/tilelang/language/tir/index
3941
/autoapi/tilelang/language/utils/index
4042
/autoapi/tilelang/language/warpgroup/index
@@ -45,110 +47,11 @@ Functions
4547

4648
.. autoapisummary::
4749

48-
tilelang.language.symbolic
49-
tilelang.language.use_swizzle
50-
tilelang.language.annotate_layout
51-
tilelang.language.annotate_safe_value
52-
tilelang.language.annotate_l2_hit_ratio
5350
tilelang.language.import_source
5451

5552

5653
Package Contents
5754
----------------
5855

59-
.. py:function:: symbolic(name, dtype = 'int32')
60-
61-
Create a TIR symbolic variable.
62-
63-
:param name: Identifier for the variable in generated TIR.
64-
:type name: str
65-
:param dtype: Data type string for the variable (e.g., "int32"). Defaults to "int32".
66-
:type dtype: str
67-
68-
:returns: A TIR variable with the given name and dtype for use in TIR/TensorIR kernels.
69-
:rtype: tir.Var
70-
71-
72-
.. py:function:: use_swizzle(panel_size, order = 'row', enable = True)
73-
74-
.. py:function:: annotate_layout(layout_map)
75-
76-
Annotate the layout of the buffer
77-
78-
:param layout_map: a dictionary of buffer to layout
79-
:type layout_map: Dict
80-
81-
:returns: a block attribute
82-
:rtype: block_attr
83-
84-
.. rubric:: Example
85-
86-
@T.prim_func
87-
def main(
88-
A: T.Tensor((M, N), dtype),
89-
B: T.Tensor((M, N), dtype),
90-
):
91-
# Initialize Kernel Context
92-
with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
93-
A_shared = T.alloc_shared((block_M, block_N), dtype)
94-
95-
T.annotate_layout({A_shared: layout})
96-
for i, j in T.Parallel(block_M, block_N):
97-
A_shared[i, j] = A[by * block_M + i, bx * block_N + j]
98-
99-
for i, j in T.Parallel(block_M, block_N):
100-
B[by * block_M + i, bx * block_N + j] = A_shared[i, j]
101-
102-
return main
103-
104-
105-
.. py:function:: annotate_safe_value(safe_value_map)
106-
107-
Annotate the safe value of the buffer.
108-
109-
A safe value of a buffer is the value that will be used when the
110-
buffer is accessed out of bounds.
111-
112-
:param safe_value_map: a dictionary of buffer to safe value
113-
:type safe_value_map: dict
114-
115-
:returns: a block attribute
116-
:rtype: block_attr
117-
118-
.. rubric:: Example
119-
120-
@T.prim_func
121-
def main(
122-
A: T.Tensor((M, N), dtype),
123-
B: T.Tensor((M, N), dtype),
124-
):
125-
# Initialize Kernel Context
126-
with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
127-
A_shared = T.alloc_shared((block_M, block_N), dtype)
128-
129-
T.annotate_safe_value({A: safe_value})
130-
for i, j in T.Parallel(block_M, block_N):
131-
A_shared[i, j] = A[by * block_M + i - 10, bx * block_N + j]
132-
133-
for i, j in T.Parallel(block_M, block_N):
134-
B[by * block_M + i, bx * block_N + j] = A_shared[i, j]
135-
136-
return main
137-
138-
139-
.. py:function:: annotate_l2_hit_ratio(l2_hit_ratio_map)
140-
141-
Annotate the L2 hit ratio of the buffer, detailed explanation please refer to:
142-
https://docs.nvidia.com/cuda/cuda-c-programming-guide/#l2-policy-for-persisting-accesses
143-
144-
:param l2_hit_ratio_map: a dictionary of buffer to L2 hit ratio value
145-
:type l2_hit_ratio_map: dict
146-
147-
.. rubric:: Example
148-
149-
# 0.5 is the hit ratio
150-
T.annotate_l2_hit_ratio({A: 0.5})
151-
152-
15356
.. py:function:: import_source(source = None)
15457

‎_sources/autoapi/tilelang/language/symbolics/index.rst.txt‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
tilelang.language.symbolics
2+
===========================
3+
4+
.. py:module:: tilelang.language.symbolics
5+
6+
.. autoapi-nested-parse::
7+
8+
Symbolic variable helpers exposed on the TileLang language surface.
9+
10+
11+
12+
Functions
13+
---------
14+
15+
.. autoapisummary::
16+
17+
tilelang.language.symbolics.dynamic
18+
tilelang.language.symbolics.symbolic
19+
20+
21+
Module Contents
22+
---------------
23+
24+
.. py:function:: dynamic(name, dtype = 'int32')
25+
26+
Create a TIR dynamic symbolic variable.
27+
28+
:param name: Identifier for the variable in generated TIR.
29+
:type name: str
30+
:param dtype: Data type string for the variable (e.g., "int32"). Defaults to "int32".
31+
:type dtype: str
32+
33+
:returns: A TIR variable with the given name and dtype for use in TIR/TensorIR kernels.
34+
:rtype: tir.Var
35+
36+
37+
.. py:function:: symbolic(name, dtype = 'int32')
38+
39+
Deprecated alias for `T.dynamic`.
40+
41+

‎_sources/deeplearning_operators/elementwise.md.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -89,7 +89,7 @@ def elementwise_add(
8989
In the compilation process above, a fixed shape was used. However, in practical usage, we often want the kernel to support dynamic shapes. So, how can we compile a kernel in TileLang to handle dynamic shapes? In TileLang, we can replace the target size with a dynamic symbolic value, making the dimension dynamic. The following example illustrates this:
9090

9191
```python
92-
program = elementwise_add(T.symbolic("N"), threads=256, dtype="bfloat16")
92+
program = elementwise_add(T.dynamic("N"), threads=256, dtype="bfloat16")
9393
kernel = tilelang.compile(program, out_idx=-1, target="cuda", execution_backend="cython")
9494
```
9595

‎autoapi/index.html‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -343,6 +343,7 @@
343343
</li>
344344
<li class="toctree-l2 has-children"><a class="reference internal" href="tilelang/language/index.html">tilelang.language</a><input aria-label="Toggle navigation of tilelang.language" class="toctree-checkbox" id="toctree-checkbox-21" name="toctree-checkbox-21" role="switch" type="checkbox"/><label for="toctree-checkbox-21"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
345345
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/allocate/index.html">tilelang.language.allocate</a></li>
346+
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/annotations/index.html">tilelang.language.annotations</a></li>
346347
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/atomic/index.html">tilelang.language.atomic</a></li>
347348
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/builtin/index.html">tilelang.language.builtin</a></li>
348349
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/copy/index.html">tilelang.language.copy</a></li>
@@ -365,6 +366,7 @@
365366
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/print/index.html">tilelang.language.print</a></li>
366367
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/proxy/index.html">tilelang.language.proxy</a></li>
367368
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/reduce/index.html">tilelang.language.reduce</a></li>
369+
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/symbolics/index.html">tilelang.language.symbolics</a></li>
368370
<li class="toctree-l3 has-children"><a class="reference internal" href="tilelang/language/tir/index.html">tilelang.language.tir</a><input aria-label="Toggle navigation of tilelang.language.tir" class="toctree-checkbox" id="toctree-checkbox-23" name="toctree-checkbox-23" role="switch" type="checkbox"/><label for="toctree-checkbox-23"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
369371
<li class="toctree-l4"><a class="reference internal" href="tilelang/language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
370372
<li class="toctree-l4"><a class="reference internal" href="tilelang/language/tir/ir/index.html">tilelang.language.tir.ir</a></li>
@@ -602,6 +604,7 @@ <h1>API Reference<a class="headerlink" href="#api-reference" title="Link to this
602604
</li>
603605
<li class="toctree-l2"><a class="reference internal" href="tilelang/language/index.html">tilelang.language</a><ul>
604606
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/allocate/index.html">tilelang.language.allocate</a></li>
607+
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/annotations/index.html">tilelang.language.annotations</a></li>
605608
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/atomic/index.html">tilelang.language.atomic</a></li>
606609
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/builtin/index.html">tilelang.language.builtin</a></li>
607610
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/copy/index.html">tilelang.language.copy</a></li>
@@ -624,6 +627,7 @@ <h1>API Reference<a class="headerlink" href="#api-reference" title="Link to this
624627
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/print/index.html">tilelang.language.print</a></li>
625628
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/proxy/index.html">tilelang.language.proxy</a></li>
626629
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/reduce/index.html">tilelang.language.reduce</a></li>
630+
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/symbolics/index.html">tilelang.language.symbolics</a></li>
627631
<li class="toctree-l3"><a class="reference internal" href="tilelang/language/tir/index.html">tilelang.language.tir</a><ul>
628632
<li class="toctree-l4"><a class="reference internal" href="tilelang/language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
629633
<li class="toctree-l4"><a class="reference internal" href="tilelang/language/tir/ir/index.html">tilelang.language.tir.ir</a></li>

‎autoapi/tilelang/autotuner/capture/index.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -343,6 +343,7 @@
343343
</li>
344344
<li class="toctree-l2 has-children"><a class="reference internal" href="../../language/index.html">tilelang.language</a><input aria-label="Toggle navigation of tilelang.language" class="toctree-checkbox" id="toctree-checkbox-21" name="toctree-checkbox-21" role="switch" type="checkbox"/><label for="toctree-checkbox-21"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
345345
<li class="toctree-l3"><a class="reference internal" href="../../language/allocate/index.html">tilelang.language.allocate</a></li>
346+
<li class="toctree-l3"><a class="reference internal" href="../../language/annotations/index.html">tilelang.language.annotations</a></li>
346347
<li class="toctree-l3"><a class="reference internal" href="../../language/atomic/index.html">tilelang.language.atomic</a></li>
347348
<li class="toctree-l3"><a class="reference internal" href="../../language/builtin/index.html">tilelang.language.builtin</a></li>
348349
<li class="toctree-l3"><a class="reference internal" href="../../language/copy/index.html">tilelang.language.copy</a></li>
@@ -365,6 +366,7 @@
365366
<li class="toctree-l3"><a class="reference internal" href="../../language/print/index.html">tilelang.language.print</a></li>
366367
<li class="toctree-l3"><a class="reference internal" href="../../language/proxy/index.html">tilelang.language.proxy</a></li>
367368
<li class="toctree-l3"><a class="reference internal" href="../../language/reduce/index.html">tilelang.language.reduce</a></li>
369+
<li class="toctree-l3"><a class="reference internal" href="../../language/symbolics/index.html">tilelang.language.symbolics</a></li>
368370
<li class="toctree-l3 has-children"><a class="reference internal" href="../../language/tir/index.html">tilelang.language.tir</a><input aria-label="Toggle navigation of tilelang.language.tir" class="toctree-checkbox" id="toctree-checkbox-23" name="toctree-checkbox-23" role="switch" type="checkbox"/><label for="toctree-checkbox-23"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
369371
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
370372
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/ir/index.html">tilelang.language.tir.ir</a></li>

‎autoapi/tilelang/autotuner/index.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -343,6 +343,7 @@
343343
</li>
344344
<li class="toctree-l2 has-children"><a class="reference internal" href="../language/index.html">tilelang.language</a><input aria-label="Toggle navigation of tilelang.language" class="toctree-checkbox" id="toctree-checkbox-21" name="toctree-checkbox-21" role="switch" type="checkbox"/><label for="toctree-checkbox-21"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
345345
<li class="toctree-l3"><a class="reference internal" href="../language/allocate/index.html">tilelang.language.allocate</a></li>
346+
<li class="toctree-l3"><a class="reference internal" href="../language/annotations/index.html">tilelang.language.annotations</a></li>
346347
<li class="toctree-l3"><a class="reference internal" href="../language/atomic/index.html">tilelang.language.atomic</a></li>
347348
<li class="toctree-l3"><a class="reference internal" href="../language/builtin/index.html">tilelang.language.builtin</a></li>
348349
<li class="toctree-l3"><a class="reference internal" href="../language/copy/index.html">tilelang.language.copy</a></li>
@@ -365,6 +366,7 @@
365366
<li class="toctree-l3"><a class="reference internal" href="../language/print/index.html">tilelang.language.print</a></li>
366367
<li class="toctree-l3"><a class="reference internal" href="../language/proxy/index.html">tilelang.language.proxy</a></li>
367368
<li class="toctree-l3"><a class="reference internal" href="../language/reduce/index.html">tilelang.language.reduce</a></li>
369+
<li class="toctree-l3"><a class="reference internal" href="../language/symbolics/index.html">tilelang.language.symbolics</a></li>
368370
<li class="toctree-l3 has-children"><a class="reference internal" href="../language/tir/index.html">tilelang.language.tir</a><input aria-label="Toggle navigation of tilelang.language.tir" class="toctree-checkbox" id="toctree-checkbox-23" name="toctree-checkbox-23" role="switch" type="checkbox"/><label for="toctree-checkbox-23"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
369371
<li class="toctree-l4"><a class="reference internal" href="../language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
370372
<li class="toctree-l4"><a class="reference internal" href="../language/tir/ir/index.html">tilelang.language.tir.ir</a></li>

‎autoapi/tilelang/autotuner/param/index.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -343,6 +343,7 @@
343343
</li>
344344
<li class="toctree-l2 has-children"><a class="reference internal" href="../../language/index.html">tilelang.language</a><input aria-label="Toggle navigation of tilelang.language" class="toctree-checkbox" id="toctree-checkbox-21" name="toctree-checkbox-21" role="switch" type="checkbox"/><label for="toctree-checkbox-21"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
345345
<li class="toctree-l3"><a class="reference internal" href="../../language/allocate/index.html">tilelang.language.allocate</a></li>
346+
<li class="toctree-l3"><a class="reference internal" href="../../language/annotations/index.html">tilelang.language.annotations</a></li>
346347
<li class="toctree-l3"><a class="reference internal" href="../../language/atomic/index.html">tilelang.language.atomic</a></li>
347348
<li class="toctree-l3"><a class="reference internal" href="../../language/builtin/index.html">tilelang.language.builtin</a></li>
348349
<li class="toctree-l3"><a class="reference internal" href="../../language/copy/index.html">tilelang.language.copy</a></li>
@@ -365,6 +366,7 @@
365366
<li class="toctree-l3"><a class="reference internal" href="../../language/print/index.html">tilelang.language.print</a></li>
366367
<li class="toctree-l3"><a class="reference internal" href="../../language/proxy/index.html">tilelang.language.proxy</a></li>
367368
<li class="toctree-l3"><a class="reference internal" href="../../language/reduce/index.html">tilelang.language.reduce</a></li>
369+
<li class="toctree-l3"><a class="reference internal" href="../../language/symbolics/index.html">tilelang.language.symbolics</a></li>
368370
<li class="toctree-l3 has-children"><a class="reference internal" href="../../language/tir/index.html">tilelang.language.tir</a><input aria-label="Toggle navigation of tilelang.language.tir" class="toctree-checkbox" id="toctree-checkbox-23" name="toctree-checkbox-23" role="switch" type="checkbox"/><label for="toctree-checkbox-23"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
369371
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
370372
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/ir/index.html">tilelang.language.tir.ir</a></li>

‎autoapi/tilelang/autotuner/tuner/index.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -343,6 +343,7 @@
343343
</li>
344344
<li class="toctree-l2 has-children"><a class="reference internal" href="../../language/index.html">tilelang.language</a><input aria-label="Toggle navigation of tilelang.language" class="toctree-checkbox" id="toctree-checkbox-21" name="toctree-checkbox-21" role="switch" type="checkbox"/><label for="toctree-checkbox-21"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
345345
<li class="toctree-l3"><a class="reference internal" href="../../language/allocate/index.html">tilelang.language.allocate</a></li>
346+
<li class="toctree-l3"><a class="reference internal" href="../../language/annotations/index.html">tilelang.language.annotations</a></li>
346347
<li class="toctree-l3"><a class="reference internal" href="../../language/atomic/index.html">tilelang.language.atomic</a></li>
347348
<li class="toctree-l3"><a class="reference internal" href="../../language/builtin/index.html">tilelang.language.builtin</a></li>
348349
<li class="toctree-l3"><a class="reference internal" href="../../language/copy/index.html">tilelang.language.copy</a></li>
@@ -365,6 +366,7 @@
365366
<li class="toctree-l3"><a class="reference internal" href="../../language/print/index.html">tilelang.language.print</a></li>
366367
<li class="toctree-l3"><a class="reference internal" href="../../language/proxy/index.html">tilelang.language.proxy</a></li>
367368
<li class="toctree-l3"><a class="reference internal" href="../../language/reduce/index.html">tilelang.language.reduce</a></li>
369+
<li class="toctree-l3"><a class="reference internal" href="../../language/symbolics/index.html">tilelang.language.symbolics</a></li>
368370
<li class="toctree-l3 has-children"><a class="reference internal" href="../../language/tir/index.html">tilelang.language.tir</a><input aria-label="Toggle navigation of tilelang.language.tir" class="toctree-checkbox" id="toctree-checkbox-23" name="toctree-checkbox-23" role="switch" type="checkbox"/><label for="toctree-checkbox-23"><span class="icon"><svg><use href="#svg-arrow-right"></use></svg></span></label><ul>
369371
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/entry/index.html">tilelang.language.tir.entry</a></li>
370372
<li class="toctree-l4"><a class="reference internal" href="../../language/tir/ir/index.html">tilelang.language.tir.ir</a></li>

0 commit comments

Comments
 (0)