Skip to content

Commit 5dfe71b

Browse files
github-actions[bot]github-actions[bot]
committed
Update docs
1 parent 4a7d058 commit 5dfe71b

File tree

9 files changed

+90
-37
lines changed

9 files changed

+90
-37
lines changed

‎_sources/autoapi/tilelang/autotuner/param/index.rst.txt‎

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -91,15 +91,9 @@ Module Contents
9191

9292
Additional keyword arguments to pass to the Compiler PassContext.
9393

94-
.. attribute:: Available options
95-
96-
"tir.disable_vectorize": bool, default: False
97-
"tl.disable_tma_lower": bool, default: False
98-
"tl.disable_warp_specialized": bool, default: False
99-
"tl.config_index_bitwidth": int, default: None
100-
"tl.disable_dynamic_tail_split": bool, default: False
101-
"tl.dynamic_vectorize_size_bits": int, default: 128
102-
"tl.disable_safe_memory_legalize": bool, default: False
94+
.. attribute:: Refer to `tilelang.PassConfigKey` for supported options.
95+
96+
10397

10498

10599
.. py:attribute:: out_idx

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

Lines changed: 1 addition & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -61,14 +61,7 @@ Package Contents
6161
:param verbose: Whether to enable verbose output (default: False).
6262
:type verbose: bool, optional
6363
:param pass_configs: Additional keyword arguments to pass to the Compiler PassContext.
64-
Available options:
65-
"tir.disable_vectorize": bool, default: False
66-
"tl.disable_tma_lower": bool, default: False
67-
"tl.disable_warp_specialized": bool, default: False
68-
"tl.config_index_bitwidth": int, default: None
69-
"tl.disable_dynamic_tail_split": bool, default: False
70-
"tl.dynamic_vectorize_size_bits": int, default: 128
71-
"tl.disable_safe_memory_legalize": bool, default: False
64+
Refer to `tilelang.transform.PassConfigKey` for supported options.
7265
:type pass_configs: dict, optional
7366

7467

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

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -188,6 +188,49 @@ Module Contents
188188
:type: Force TileLang to inline let bindings during simplification. Default
189189

190190

191+
.. py:attribute:: TL_STORAGE_REWRITE_DETECT_INPLACE
192+
:value: 'tl.storage_rewrite_detect_inplace'
193+
194+
195+
Control StorageRewrite inplace detection.
196+
197+
When False (default) StorageRewrite keeps distinct temporaries for patterns
198+
such as `dst[i] = f(src[i])`, avoiding implicit aliasing:
199+
200+
```
201+
read = T.allocate([1], "int32", "local.var")
202+
write = T.allocate([1], "int32", "local.var")
203+
read_buf = T.Buffer((1,), "int32", data=read, scope="local.var")
204+
write_buf = T.Buffer((1,), "int32", data=write, scope="local.var")
205+
write_buf[0] = read_buf[0] * 2
206+
f(write_buf[0])
207+
```
208+
209+
Setting the flag to True allows StorageRewrite to reuse the `read` buffer
210+
for the write when it can prove the update is safely inplace, producing IR
211+
like:
212+
213+
```
214+
read = T.allocate([1], "int32", "local.var")
215+
read_buf = T.Buffer((1,), "int32", data=read, scope="local.var")
216+
read_buf[0] = read_buf[0] * 2
217+
f(read_buf[0])
218+
```
219+
220+
This reduces local memory usage but introduces aliasing between the buffers.
221+
222+
Usage:
223+
224+
```python
225+
from tilelang.transform import PassContext, PassConfigKey
226+
227+
with PassContext(
228+
config={PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE.value: True}
229+
):
230+
mod = tilelang.transform.StorageRewrite()(mod)
231+
```
232+
233+
191234
.. py:attribute:: TIR_ENABLE_EQUIV_TERMS_IN_CSE
192235
:value: 'tir.enable_equiv_terms_in_cse_tir'
193236

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

Lines changed: 2 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -607,15 +607,8 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
607607

608608
<dl class="py attribute">
609609
<dt class="sig sig-object py">
610-
<span class="sig-name descname"><span class="pre">Available</span> <span class="pre">options</span></span></dt>
611-
<dd><p>“tir.disable_vectorize”: bool, default: False
612-
“tl.disable_tma_lower”: bool, default: False
613-
“tl.disable_warp_specialized”: bool, default: False
614-
“tl.config_index_bitwidth”: int, default: None
615-
“tl.disable_dynamic_tail_split”: bool, default: False
616-
“tl.dynamic_vectorize_size_bits”: int, default: 128
617-
“tl.disable_safe_memory_legalize”: bool, default: False</p>
618-
</dd></dl>
610+
<span class="sig-name descname"><span class="pre">Refer</span> <span class="pre">to</span> <span class="pre">`tilelang.PassConfigKey`</span> <span class="pre">for</span> <span class="pre">supported</span> <span class="pre">options.</span></span></dt>
611+
<dd></dd></dl>
619612

620613
<dl class="py attribute">
621614
<dt class="sig sig-object py" id="tilelang.autotuner.param.CompileArgs.out_idx">

‎autoapi/tilelang/jit/index.html‎

Lines changed: 1 addition & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -547,16 +547,7 @@ <h2>Package Contents<a class="headerlink" href="#package-contents" title="Link t
547547
:type verbose: bool, optional
548548
:param pass_configs: Additional keyword arguments to pass to the Compiler PassContext.</p>
549549
<blockquote>
550-
<div><dl class="simple">
551-
<dt>Available options:</dt><dd><p>“tir.disable_vectorize”: bool, default: False
552-
“tl.disable_tma_lower”: bool, default: False
553-
“tl.disable_warp_specialized”: bool, default: False
554-
“tl.config_index_bitwidth”: int, default: None
555-
“tl.disable_dynamic_tail_split”: bool, default: False
556-
“tl.dynamic_vectorize_size_bits”: int, default: 128
557-
“tl.disable_safe_memory_legalize”: bool, default: False</p>
558-
</dd>
559-
</dl>
550+
<div><p>Refer to <cite>tilelang.transform.PassConfigKey</cite> for supported options.</p>
560551
</div></blockquote>
561552
<dl class="field-list simple">
562553
<dt class="field-odd">Parameters<span class="colon">:</span></dt>

‎autoapi/tilelang/transform/pass_config/index.html‎

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -695,6 +695,42 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
695695
</dl>
696696
</dd></dl>
697697

698+
<dl class="py attribute">
699+
<dt class="sig sig-object py" id="tilelang.transform.pass_config.PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE">
700+
<span class="sig-name descname"><span class="pre">TL_STORAGE_REWRITE_DETECT_INPLACE</span></span><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">'tl.storage_rewrite_detect_inplace'</span></em><a class="headerlink" href="#tilelang.transform.pass_config.PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE" title="Link to this definition"></a></dt>
701+
<dd><p>Control StorageRewrite inplace detection.</p>
702+
<p>When False (default) StorageRewrite keeps distinct temporaries for patterns
703+
such as <cite>dst[i] = f(src[i])</cite>, avoiding implicit aliasing:</p>
704+
<p><code class="docutils literal notranslate"><span class="pre">`</span>
705+
<span class="pre">read</span> <span class="pre">=</span> <span class="pre">T.allocate([1],</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">&quot;local.var&quot;)</span>
706+
<span class="pre">write</span> <span class="pre">=</span> <span class="pre">T.allocate([1],</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">&quot;local.var&quot;)</span>
707+
<span class="pre">read_buf</span> <span class="pre">=</span> <span class="pre">T.Buffer((1,),</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">data=read,</span> <span class="pre">scope=&quot;local.var&quot;)</span>
708+
<span class="pre">write_buf</span> <span class="pre">=</span> <span class="pre">T.Buffer((1,),</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">data=write,</span> <span class="pre">scope=&quot;local.var&quot;)</span>
709+
<span class="pre">write_buf[0]</span> <span class="pre">=</span> <span class="pre">read_buf[0]</span> <span class="pre">*</span> <span class="pre">2</span>
710+
<span class="pre">f(write_buf[0])</span>
711+
<span class="pre">`</span></code></p>
712+
<p>Setting the flag to True allows StorageRewrite to reuse the <cite>read</cite> buffer
713+
for the write when it can prove the update is safely inplace, producing IR
714+
like:</p>
715+
<p><code class="docutils literal notranslate"><span class="pre">`</span>
716+
<span class="pre">read</span> <span class="pre">=</span> <span class="pre">T.allocate([1],</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">&quot;local.var&quot;)</span>
717+
<span class="pre">read_buf</span> <span class="pre">=</span> <span class="pre">T.Buffer((1,),</span> <span class="pre">&quot;int32&quot;,</span> <span class="pre">data=read,</span> <span class="pre">scope=&quot;local.var&quot;)</span>
718+
<span class="pre">read_buf[0]</span> <span class="pre">=</span> <span class="pre">read_buf[0]</span> <span class="pre">*</span> <span class="pre">2</span>
719+
<span class="pre">f(read_buf[0])</span>
720+
<span class="pre">`</span></code></p>
721+
<p>This reduces local memory usage but introduces aliasing between the buffers.</p>
722+
<p>Usage:</p>
723+
<p><a href="#id1"><span class="problematic" id="id2">``</span></a><a href="#id3"><span class="problematic" id="id4">`</span></a>python
724+
from tilelang.transform import PassContext, PassConfigKey</p>
725+
<dl class="simple">
726+
<dt>with PassContext(</dt><dd><p>config={PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE.value: True}</p>
727+
</dd>
728+
<dt>):</dt><dd><p>mod = tilelang.transform.StorageRewrite()(mod)</p>
729+
</dd>
730+
</dl>
731+
<p><a href="#id5"><span class="problematic" id="id6">``</span></a><a href="#id7"><span class="problematic" id="id8">`</span></a></p>
732+
</dd></dl>
733+
698734
<dl class="py attribute">
699735
<dt class="sig sig-object py" id="tilelang.transform.pass_config.PassConfigKey.TIR_ENABLE_EQUIV_TERMS_IN_CSE">
700736
<span class="sig-name descname"><span class="pre">TIR_ENABLE_EQUIV_TERMS_IN_CSE</span></span><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">'tir.enable_equiv_terms_in_cse_tir'</span></em><a class="headerlink" href="#tilelang.transform.pass_config.PassConfigKey.TIR_ENABLE_EQUIV_TERMS_IN_CSE" title="Link to this definition"></a></dt>
@@ -898,6 +934,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
898934
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TL_DISABLE_SHUFFLE_ELECT"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TL_DISABLE_SHUFFLE_ELECT</span></code></a></li>
899935
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TL_DISABLE_THREAD_STORAGE_SYNC"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TL_DISABLE_THREAD_STORAGE_SYNC</span></code></a></li>
900936
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TL_FORCE_LET_INLINE"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TL_FORCE_LET_INLINE</span></code></a></li>
937+
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE</span></code></a></li>
901938
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TIR_ENABLE_EQUIV_TERMS_IN_CSE"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TIR_ENABLE_EQUIV_TERMS_IN_CSE</span></code></a></li>
902939
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TIR_DISABLE_CSE"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TIR_DISABLE_CSE</span></code></a></li>
903940
<li><a class="reference internal" href="#tilelang.transform.pass_config.PassConfigKey.TIR_SIMPLIFY"><code class="docutils literal notranslate"><span class="pre">PassConfigKey.TIR_SIMPLIFY</span></code></a></li>

‎genindex.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6024,6 +6024,8 @@ <h2>T</h2>
60246024
<li><a href="autoapi/tilelang/env/index.html#tilelang.env.TL_ROOT">TL_ROOT (in module tilelang.env)</a>
60256025
</li>
60266026
<li><a href="autoapi/tilelang/transform/pass_config/index.html#tilelang.transform.pass_config.PassConfigKey.TL_SIMPLIFY">TL_SIMPLIFY (tilelang.transform.pass_config.PassConfigKey attribute)</a>
6027+
</li>
6028+
<li><a href="autoapi/tilelang/transform/pass_config/index.html#tilelang.transform.pass_config.PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE">TL_STORAGE_REWRITE_DETECT_INPLACE (tilelang.transform.pass_config.PassConfigKey attribute)</a>
60276029
</li>
60286030
<li><a href="autoapi/tilelang/env/index.html#tilelang.env.TL_TEMPLATE_NOT_FOUND_MESSAGE">TL_TEMPLATE_NOT_FOUND_MESSAGE (in module tilelang.env)</a>
60296031
</li>

‎objects.inv‎

19 Bytes
Binary file not shown.

‎searchindex.js‎

Lines changed: 1 addition & 1 deletion
Some generated files are not rendered by default. Learn more about customizing how changed files appear on GitHub.

0 commit comments

Comments
 (0)