Skip to content

Commit 16f37f6

Browse files
committed
Update docs
1 parent 78c882a commit 16f37f6

File tree

7 files changed

+89
-6
lines changed

7 files changed

+89
-6
lines changed

api/tilelang.html

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -945,6 +945,7 @@ <h2>Subpackages<a class="headerlink" href="#subpackages" title="Permalink to thi
945945
</li>
946946
<li class="toctree-l2"><a class="reference internal" href="tilelang.language.html#module-tilelang.language">Module contents</a><ul>
947947
<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.annotate_layout"><code class="docutils literal notranslate"><span class="pre">annotate_layout()</span></code></a></li>
948+
<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.annotate_padding"><code class="docutils literal notranslate"><span class="pre">annotate_padding()</span></code></a></li>
948949
<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.import_source"><code class="docutils literal notranslate"><span class="pre">import_source()</span></code></a></li>
949950
<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.symbolic"><code class="docutils literal notranslate"><span class="pre">symbolic()</span></code></a></li>
950951
<li class="toctree-l3"><a class="reference internal" href="tilelang.language.html#tilelang.language.use_swizzle"><code class="docutils literal notranslate"><span class="pre">use_swizzle()</span></code></a></li>

api/tilelang.language.html

Lines changed: 82 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -766,8 +766,87 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
766766
<p>The language interface for tl programs.</p>
767767
<dl class="py function">
768768
<dt class="sig sig-object py" id="tilelang.language.annotate_layout">
769-
<span class="sig-prename descclassname"><span class="pre">tilelang.language.</span></span><span class="sig-name descname"><span class="pre">annotate_layout</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">layout_map</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.annotate_layout" title="Permalink to this definition">#</a></dt>
770-
<dd></dd></dl>
769+
<span class="sig-prename descclassname"><span class="pre">tilelang.language.</span></span><span class="sig-name descname"><span class="pre">annotate_layout</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">layout_map</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Dict</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.annotate_layout" title="Permalink to this definition">#</a></dt>
770+
<dd><p>Annotate the layout of the buffer</p>
771+
<dl class="field-list simple">
772+
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
773+
<dd class="field-odd"><p><strong>layout_map</strong> (<em>Dict</em>) – a dictionary of buffer to layout</p>
774+
</dd>
775+
<dt class="field-even">Returns<span class="colon">:</span></dt>
776+
<dd class="field-even"><p>a block attribute</p>
777+
</dd>
778+
<dt class="field-odd">Return type<span class="colon">:</span></dt>
779+
<dd class="field-odd"><p>block_attr</p>
780+
</dd>
781+
</dl>
782+
<p class="rubric">Example</p>
783+
<p>&#64;T.prim_func
784+
def main(</p>
785+
<blockquote>
786+
<div><p>A: T.Tensor((M, N), dtype),
787+
B: T.Tensor((M, N), dtype),</p>
788+
</div></blockquote>
789+
<dl>
790+
<dt>):</dt><dd><p># Initialize Kernel Context
791+
with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):</p>
792+
<blockquote>
793+
<div><p>A_shared = T.alloc_shared((block_M, block_N), dtype)</p>
794+
<p>T.annotate_layout({A_shared: layout})
795+
for i, j in T.Parallel(block_M, block_N):</p>
796+
<blockquote>
797+
<div><p>A_shared[i, j] = A[by * block_M + i, bx * block_N + j]</p>
798+
</div></blockquote>
799+
<dl class="simple">
800+
<dt>for i, j in T.Parallel(block_M, block_N):</dt><dd><p>B[by * block_M + i, bx * block_N + j] = A_shared[i, j]</p>
801+
</dd>
802+
</dl>
803+
</div></blockquote>
804+
</dd>
805+
</dl>
806+
<p>return main</p>
807+
</dd></dl>
808+
809+
<dl class="py function">
810+
<dt class="sig sig-object py" id="tilelang.language.annotate_padding">
811+
<span class="sig-prename descclassname"><span class="pre">tilelang.language.</span></span><span class="sig-name descname"><span class="pre">annotate_padding</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">padding_map</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Dict</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.annotate_padding" title="Permalink to this definition">#</a></dt>
812+
<dd><p>Annotate the padding of the buffer</p>
813+
<dl class="field-list simple">
814+
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
815+
<dd class="field-odd"><p><strong>padding_map</strong> (<em>dict</em>) – a dictionary of buffer to padding value</p>
816+
</dd>
817+
<dt class="field-even">Returns<span class="colon">:</span></dt>
818+
<dd class="field-even"><p>a block attribute</p>
819+
</dd>
820+
<dt class="field-odd">Return type<span class="colon">:</span></dt>
821+
<dd class="field-odd"><p>block_attr</p>
822+
</dd>
823+
</dl>
824+
<p class="rubric">Example</p>
825+
<p>&#64;T.prim_func
826+
def main(</p>
827+
<blockquote>
828+
<div><p>A: T.Tensor((M, N), dtype),
829+
B: T.Tensor((M, N), dtype),</p>
830+
</div></blockquote>
831+
<dl>
832+
<dt>):</dt><dd><p># Initialize Kernel Context
833+
with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):</p>
834+
<blockquote>
835+
<div><p>A_shared = T.alloc_shared((block_M, block_N), dtype)</p>
836+
<p>T.annotate_padding({A_shared: pad_value})
837+
for i, j in T.Parallel(block_M, block_N):</p>
838+
<blockquote>
839+
<div><p>A_shared[i, j] = A[by * block_M + i - 10, bx * block_N + j]</p>
840+
</div></blockquote>
841+
<dl class="simple">
842+
<dt>for i, j in T.Parallel(block_M, block_N):</dt><dd><p>B[by * block_M + i, bx * block_N + j] = A_shared[i, j]</p>
843+
</dd>
844+
</dl>
845+
</div></blockquote>
846+
</dd>
847+
</dl>
848+
<p>return main</p>
849+
</dd></dl>
771850

772851
<dl class="py function">
773852
<dt class="sig sig-object py" id="tilelang.language.import_source">
@@ -847,6 +926,7 @@ <h2>Submodules<a class="headerlink" href="#submodules" title="Permalink to this
847926
<li><a class="reference internal" href="#submodules">Submodules</a></li>
848927
<li><a class="reference internal" href="#module-tilelang.language">Module contents</a><ul>
849928
<li><a class="reference internal" href="#tilelang.language.annotate_layout"><code class="docutils literal notranslate"><span class="pre">annotate_layout()</span></code></a></li>
929+
<li><a class="reference internal" href="#tilelang.language.annotate_padding"><code class="docutils literal notranslate"><span class="pre">annotate_padding()</span></code></a></li>
850930
<li><a class="reference internal" href="#tilelang.language.import_source"><code class="docutils literal notranslate"><span class="pre">import_source()</span></code></a></li>
851931
<li><a class="reference internal" href="#tilelang.language.symbolic"><code class="docutils literal notranslate"><span class="pre">symbolic()</span></code></a></li>
852932
<li><a class="reference internal" href="#tilelang.language.use_swizzle"><code class="docutils literal notranslate"><span class="pre">use_swizzle()</span></code></a></li>

api/tilelang.language.logical.html

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -412,7 +412,7 @@
412412
<p>The language interface for tl programs.</p>
413413
<dl class="py function">
414414
<dt class="sig sig-object py" id="tilelang.language.logical.all_of">
415-
<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">all_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7fceea39b6d0&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.all_of" title="Permalink to this definition">#</a></dt>
415+
<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">all_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f5873842bf0&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.all_of" title="Permalink to this definition">#</a></dt>
416416
<dd><p>Check if all elements in the buffer are true.</p>
417417
<dl class="field-list simple">
418418
<dt class="field-odd">Parameters<span class="colon">:</span></dt>
@@ -426,7 +426,7 @@
426426

427427
<dl class="py function">
428428
<dt class="sig sig-object py" id="tilelang.language.logical.any_of">
429-
<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">any_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7fceea39b6d0&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.any_of" title="Permalink to this definition">#</a></dt>
429+
<span class="sig-prename descclassname"><span class="pre">tilelang.language.logical.</span></span><span class="sig-name descname"><span class="pre">any_of</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="pre">buffer:</span> <span class="pre">~typing.Union[&lt;tilelang.language.proxy.TensorProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f5873842bf0&gt;,</span> <span class="pre">~tvm.tir.stmt.BufferRegion]</span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.logical.any_of" title="Permalink to this definition">#</a></dt>
430430
<dd><p>Check if any element in the buffer is true.</p>
431431
<dl class="field-list simple">
432432
<dt class="field-odd">Parameters<span class="colon">:</span></dt>

api/tilelang.language.proxy.html

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -460,7 +460,7 @@
460460
<p>Buffer proxy class for constructing tir buffer.</p>
461461
<dl class="py method">
462462
<dt class="sig sig-object py" id="tilelang.language.proxy.BufferProxy.from_ptr">
463-
<span class="sig-name descname"><span class="pre">from_ptr</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">pointer_var</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Var</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">shape</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">tuple</span><span class="p"><span class="pre">[</span></span><span class="pre">tvm.ir.expr.PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="p"><span class="pre">...</span></span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</span></span><span class="w"> </span><span class="o"><span class="pre">=</span></span><span class="w"> </span><span class="default_value"><span class="pre">'float32'</span></span></em><span class="sig-paren">)</span> <span class="sig-return"><span class="sig-return-icon">&#x2192;</span> <span class="sig-return-typehint"><span class="pre">&lt;tilelang.language.proxy.BufferProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7fceea39b700&gt;</span></span></span><a class="headerlink" href="#tilelang.language.proxy.BufferProxy.from_ptr" title="Permalink to this definition">#</a></dt>
463+
<span class="sig-name descname"><span class="pre">from_ptr</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">pointer_var</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">Var</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">shape</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">tuple</span><span class="p"><span class="pre">[</span></span><span class="pre">tvm.ir.expr.PrimExpr</span><span class="p"><span class="pre">,</span></span><span class="w"> </span><span class="p"><span class="pre">...</span></span><span class="p"><span class="pre">]</span></span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><span class="pre">str</span></span><span class="w"> </span><span class="o"><span class="pre">=</span></span><span class="w"> </span><span class="default_value"><span class="pre">'float32'</span></span></em><span class="sig-paren">)</span> <span class="sig-return"><span class="sig-return-icon">&#x2192;</span> <span class="sig-return-typehint"><span class="pre">&lt;tilelang.language.proxy.BufferProxy</span> <span class="pre">object</span> <span class="pre">at</span> <span class="pre">0x7f5873842f20&gt;</span></span></span><a class="headerlink" href="#tilelang.language.proxy.BufferProxy.from_ptr" title="Permalink to this definition">#</a></dt>
464464
<dd><p>Create a buffer from a pointer, shape, and data type.</p>
465465
<dl class="field-list simple">
466466
<dt class="field-odd">Parameters<span class="colon">:</span></dt>

genindex.html

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -476,6 +476,8 @@ <h2>A</h2>
476476
<li><a href="api/tilelang.tools.Analyzer.html#tilelang.tools.Analyzer.Analyzer">Analyzer (class in tilelang.tools.Analyzer)</a>
477477
</li>
478478
<li><a href="api/tilelang.language.html#tilelang.language.annotate_layout">annotate_layout() (in module tilelang.language)</a>
479+
</li>
480+
<li><a href="api/tilelang.language.html#tilelang.language.annotate_padding">annotate_padding() (in module tilelang.language)</a>
479481
</li>
480482
<li><a href="api/tilelang.transform.html#tilelang.transform.AnnotateDeviceRegions">AnnotateDeviceRegions() (in module tilelang.transform)</a>
481483
</li>

objects.inv

8 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)