Skip to content

Commit 686d7a6

Browse files
Update docs
1 parent 21ab6bd commit 686d7a6

File tree

3 files changed

+21
-17
lines changed

3 files changed

+21
-17
lines changed

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

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -53,14 +53,16 @@ Module Contents
5353
- Accepts `Buffer`/`BufferRegion`/`BufferLoad` on either side. Extents are
5454
derived as follows: `Buffer -> shape`, `BufferRegion -> [r.extent]`,
5555
`BufferLoad -> extents from its inferred/encoded region`.
56-
- If both `src` and `dst` are scalar `BufferLoad` without region extents,
57-
lowers to a direct store: `dst[...] = src`.
58-
- If one side is missing extents, it is treated as all-ones with the other
59-
side's rank to enable broadcasting.
60-
- Extents are right-aligned and legalized via `legalize_pairwise_extents`:
61-
per tail-dimension, equal keeps as-is, a `1` broadcasts to the other,
62-
otherwise a conservative `tir.max` is used to remain safe for dynamic
63-
shapes.
56+
- Normally, we require the extents of both sides to be the same. If they
57+
differ, the copy instruction follows an internal rule to select one side
58+
as the base range and create iteration space. This may generate unexpected
59+
code. And if some dimensions are 1, unexpected errors may happen.
60+
- Small Optimization: If both `src` and `dst` are scalar `BufferLoad` without
61+
region extents, lowers to a direct store: `dst[...] = src[...]`.
62+
- Syntactic Sugar: TileLang supports passing the head address of a buffer to represent
63+
the whole buffer if there are no ambiguity. For example, T.copy(A, A_shared[i, j]).
64+
To support this, we need some special shape checking. But remember currently we don't
65+
support something like "broadcast".
6466
- The finalized extents are encoded with `tl.region` via `to_buffer_region`
6567
and passed through to the backend; low-level loop construction and any
6668
scope-specific decisions happen during lowering.

autoapi/tilelang/language/copy_op/index.html

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -612,14 +612,16 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
612612
<cite>BufferLoad -&gt; extents from its inferred/encoded region</cite>.</p>
613613
</div></blockquote>
614614
<ul class="simple">
615-
<li><p>If both <cite>src</cite> and <cite>dst</cite> are scalar <cite>BufferLoad</cite> without region extents,
616-
lowers to a direct store: <cite>dst[…] = src</cite>.</p></li>
617-
<li><p>If one side is missing extents, it is treated as all-ones with the other
618-
side’s rank to enable broadcasting.</p></li>
619-
<li><p>Extents are right-aligned and legalized via <cite>legalize_pairwise_extents</cite>:
620-
per tail-dimension, equal keeps as-is, a <cite>1</cite> broadcasts to the other,
621-
otherwise a conservative <cite>tir.max</cite> is used to remain safe for dynamic
622-
shapes.</p></li>
615+
<li><p>Normally, we require the extents of both sides to be the same. If they
616+
differ, the copy instruction follows an internal rule to select one side
617+
as the base range and create iteration space. This may generate unexpected
618+
code. And if some dimensions are 1, unexpected errors may happen.</p></li>
619+
<li><p>Small Optimization: If both <cite>src</cite> and <cite>dst</cite> are scalar <cite>BufferLoad</cite> without
620+
region extents, lowers to a direct store: <cite>dst[…] = src[…]</cite>.</p></li>
621+
<li><p>Syntactic Sugar: TileLang supports passing the head address of a buffer to represent
622+
the whole buffer if there are no ambiguity. For example, T.copy(A, A_shared[i, j]).
623+
To support this, we need some special shape checking. But remember currently we don’t
624+
support something like “broadcast”.</p></li>
623625
<li><p>The finalized extents are encoded with <cite>tl.region</cite> via <cite>to_buffer_region</cite>
624626
and passed through to the backend; low-level loop construction and any
625627
scope-specific decisions happen during lowering.</p></li>

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)