(tvm-site) branch asf-site updated: deploying tirx docs (40d297af8c9b13a5bc890e0e5622a578c21469b2)

[tqchen.apache.org via commits]

This is an automated email from the ASF dual-hosted git repository.

tqchen pushed a commit to branch asf-site
in repository https://gitbox.apache.org/repos/asf/tvm-site.git


The following commit(s) were added to refs/heads/asf-site by this push:
     new b564f736cdd deploying tirx docs 
(40d297af8c9b13a5bc890e0e5622a578c21469b2)
b564f736cdd is described below

commit b564f736cdd34273144459f43c611f7b8b29141f
Author: tvm-bot <95660001+tvm-...@users.noreply.github.com>
AuthorDate: Fri Nov 6 05:15:18 2026 +0000

    deploying tirx docs (40d297af8c9b13a5bc890e0e5622a578c21469b2)
---
 .../tile/blackwell_tensor_cores.rst                | 141 ++++++++++++---------
 .../programming_guides/tile/warp_matrix.rst        | 114 +++++++++--------
 .../cpp/generated/file_tvm_ir_attrs.h.rst          |   4 +-
 .../cpp/generated/file_tvm_ir_prim_op.h.rst        |  52 ++++----
 .../generated/file_tvm_ir_prim_pattern_match.h.rst |  12 +-
 .../cpp/generated/file_tvm_tirx_op_math.h.rst      |   4 +-
 .../file_tvm_tirx_script_ir_builder_ir.h.rst       |   4 +-
 .../reference/cpp/generated/namespace_tvm.rst      |  24 ++--
 .../namespace_tvm__script__ir_builder__tirx.rst    |   4 +-
 tirx/commit_hash                                   |   2 +-
 .../tile/blackwell_tensor_cores.html               | 129 ++++++++++---------
 tirx/programming_guides/tile/warp_matrix.html      | 108 ++++++++--------
 .../cpp/generated/file_tvm_ir_attrs.h.html         |   2 +-
 .../cpp/generated/file_tvm_ir_prim_op.h.html       |  34 ++---
 .../file_tvm_ir_prim_pattern_match.h.html          |   6 +-
 .../cpp/generated/file_tvm_tirx_op_math.h.html     |   2 +-
 .../file_tvm_tirx_script_ir_builder_ir.h.html      |   2 +-
 tirx/reference/cpp/generated/namespace_tvm.html    |  16 +--
 .../namespace_tvm__script__ir_builder__tirx.html   |   2 +-
 .../generated/tvm_tirx.tirx.tile.Operand.html      |   2 +-
 tirx/searchindex.js                                |   2 +-
 21 files changed, 343 insertions(+), 323 deletions(-)

diff --git a/tirx/_sources/programming_guides/tile/blackwell_tensor_cores.rst 
b/tirx/_sources/programming_guides/tile/blackwell_tensor_cores.rst
index 6ed98cbcbad..fd5d899a67c 100644
--- a/tirx/_sources/programming_guides/tile/blackwell_tensor_cores.rst
+++ b/tirx/_sources/programming_guides/tile/blackwell_tensor_cores.rst
@@ -27,38 +27,49 @@ shared memory.
 The Accumulator's Lifetime
 --------------------------
 
-The example computes ``Out = A @ B.T`` with A of shape ``(128, 64)`` and
-B of shape ``(64, 64)``. One CTA of 128 threads performs these stages:
-
-.. code-block:: text
-
-    allocate TMEM -> stage A/B in shared memory -> issue MMA
-                                                    |
-                                    commit + wait for completion
-                                                    |
-    release TMEM <- write Out <- load TMEM into registers + wait
-
-Different stages require different participants. The conditions in the
-kernel select them explicitly:
-
-.. list-table::
-   :header-rows: 1
-
-   * - Participants
-     - Operations
-   * - ``warp == 0``: all 32 threads of warp 0
-     - Allocate TMEM, relinquish the allocation permit, and deallocate TMEM.
-   * - ``tid == 0``: thread 0
-     - Initialize the barrier, issue the MMA instructions, and commit them.
-   * - All four warps
-     - Stage inputs, synchronize, load the completed result from TMEM, and 
write Out.
-
-``mma_shared_layout`` constructs the swizzled shared layouts for A and B.
-Mode 3 selects a 128-byte swizzle; its 1024-byte period determines the
-shared allocation alignment used below.
-``tmem_mma_operand_layout`` maps the float32 accumulator to TMEM lanes
-and columns for the chosen instruction. The ``128 x 64`` accumulator
-uses 64 columns, so the allocation and deallocation both specify 64.
+The example computes ``Out = A @ B.T`` with A of shape ``(128, 64)``,
+B of shape ``(64, 64)``, and Out of shape ``(128, 64)``. One CTA of 128
+threads performs these stages:
+
+.. raw:: html
+
+   <iframe data-tirx-matrix-dataflow
+           
src="https://tvm.apache.org/images/docs/tirx/blackwell_tensor_cores.html";
+           title="Blackwell matrix multiplication: TMEM allocation, MMA 
completion, and output"
+           style="display: block; width: 100%; height: 740px; border: 0;"
+           loading="lazy"></iframe>
+   <script>
+     // Keep both matrix diagrams sized and themed with their documentation 
pages.
+     (() => {
+       const dark = matchMedia('(prefers-color-scheme: dark)');
+       
document.querySelectorAll('iframe[data-tirx-matrix-dataflow]').forEach((frame) 
=> {
+         if (frame.dataset.ready) return;
+         frame.dataset.ready = 'true';
+         const origin = new URL(frame.src, document.baseURI).origin;
+         const syncTheme = () => {
+           const selected = document.documentElement.dataset.theme;
+           const theme = ['light', 'dark'].includes(selected) ? selected : 
(dark.matches ? 'dark' : 'light');
+           frame.contentWindow.postMessage({ type: 
'tirx-matrix-dataflow:theme', theme }, origin);
+         };
+         addEventListener('message', (event) => {
+           if (event.source !== frame.contentWindow || event.origin !== 
origin) return;
+           const data = event.data;
+           if (data?.type !== 'tirx-matrix-dataflow:height') return;
+           if (!Number.isFinite(data.height) || data.height < 100 || 
data.height > 3000) return;
+           frame.style.height = Math.ceil(data.height) + 'px';
+         });
+         new MutationObserver(syncTheme).observe(document.documentElement, {
+           attributes: true, attributeFilter: ['data-theme']
+         });
+         dark.addEventListener('change', syncTheme);
+         frame.addEventListener('load', syncTheme);
+         syncTheme();
+       });
+     })();
+   </script>
+
+The kernel uses ``d`` for the TMEM accumulator and ``frag`` for the output
+values held in each thread's registers:
 
 .. code-block:: python
 
@@ -133,35 +144,38 @@ uses 64 columns, so the allocation and deallocation both 
specify 64.
             T.ptx.tcgen05.relinquish_alloc_permit.cta_group__1.sync.aligned()
             T.ptx.tcgen05.dealloc.cta_group__1.sync.aligned.b32(address[0], 
T.uint32(64))
 
-The allocation writes a TMEM base address to the shared ``address`` slot.
-After the CTA barrier publishes the address, ``decl_tmem`` gives the allocated
-storage a tensor shape and layout. It does not perform another allocation.
-
-Each thread stages input values in shared memory and executes a proxy
-fence. The following CTA barrier lets the issuing thread proceed only
-after all writers have finished this handoff to the asynchronous reader.
-The A/B views split K=64 into four K=16 slices. For example,
-``av`` has shape ``(1, 4, 128, 16)``: one M tile, four K repeats, and a
-``128 x 16`` A slice per issue. B likewise has one N tile and four K
-repeats. Consequently, the tile call issues four MMA instructions.
-``accum=False`` starts a fresh result on the first slice; subsequent
-slices accumulate into D.
-
-The elected thread commits these operations to ``mbar``, initialized with
-one expected arrival. All threads wait for phase 0 and execute
-``tcgen05.fence::after_thread_sync`` before reading TMEM. This completion
-protocol uses a barrier arrival, rather than the expected-byte count
-used by :doc:`TMA copies <tma>`.
-
-The register fragment assigns one output row to each thread. Its
-``(4, 32, 64)`` view lets each of the four warps load its 32 TMEM lanes.
-``tcgen05.wait::ld`` then makes the loaded registers ready for the output
-stores. It waits for these loads; the preceding barrier and fence are
-what establish that MMA has finished. The final CTA barrier keeps TMEM
-alive until every warp has finished using it, after which warp 0 gives
-up further allocations and releases the columns.
-
-On a supported GPU, we can check the product as follows:
+All 32 threads of warp 0 collectively allocate 64 TMEM columns for the
+``128 x 64`` float32 accumulator. The allocation writes its base address
+to ``address[0]``.
+
+The inputs are staged in shared-memory buffers ``a`` and ``b``.
+``mma_shared_layout`` mode ``3`` selects a 128-byte swizzle with
+1024-byte base alignment. Each thread executes ``fence.proxy.async``
+after its copies; the CTA barrier then ensures all inputs are ready
+for MMA and the TMEM address is visible to all threads.
+``T.cuda.decl_tmem`` binds that allocation to tensor ``d``.
+
+The views ``av`` and ``bv`` split K=64 into four K=16 slices, so thread 0's
+single ``T.cuda.tile.apply`` call expands into four MMA instructions.
+With ``accum=False``, the first instruction writes its product into ``d``;
+the next three add their products to it.
+
+Thread 0 calls ``tcgen05.commit`` to signal ``mbar`` once all four MMA
+instructions complete, matching the barrier's arrival count of 1.
+All threads wait for phase 0, then execute
+``tcgen05.fence::after_thread_sync`` before loading the result from TMEM.
+
+``register_layout`` assigns one output row of 64 float32 values to each
+thread. All four warps use ``tcgen05.ld`` to load their rows into ``frag``.
+``tcgen05.wait::ld`` waits for these loads to complete before each thread
+writes its row to ``Out``.
+
+The final CTA barrier ensures every warp has finished using TMEM.
+Warp 0 then collectively relinquishes the CTA's allocation permit and
+deallocates the 64 columns.
+
+On a supported GPU, we can compare the result with a PyTorch float32
+matrix product:
 
 .. code-block:: python
 
@@ -202,6 +216,11 @@ have different participants and completion rules:
      - One elected thread issues the copy and commits it to an
        ``mbarrier``; consumers wait and use the thread-sync fence.
 
+For ``tcgen05.mma`` and ``tcgen05.cp``, a commit signals completion through
+a barrier arrival. :doc:`TMA loads <tma>` instead track completion using
+an expected-byte count. The barrier setup must match the operation being
+tracked.
+
 ``tcgen05.cp`` is useful for operands that must reside in TMEM, such as
 the scale factors of a block-scaled MMA. Its shared source must be ready
 for asynchronous access, just as the MMA inputs above must be. The
diff --git a/tirx/_sources/programming_guides/tile/warp_matrix.rst 
b/tirx/_sources/programming_guides/tile/warp_matrix.rst
index bb5225099d7..ca035f7d574 100644
--- a/tirx/_sources/programming_guides/tile/warp_matrix.rst
+++ b/tirx/_sources/programming_guides/tile/warp_matrix.rst
@@ -30,46 +30,46 @@ From Shared Memory to a Matrix Product
 All 32 lanes cooperate on one ``16 x 8`` output tile. The data follows
 this path:
 
-.. code-block:: text
-
-    float16 inputs -> shared memory -> ldmatrix -> register fragments
-                                                       |
-                                                   mma.sync
-                                                       |
-                                      float32 registers -> Out
-
-Each lane holds part of a matrix, called its register fragment. In this
-example, the fragments have the following sizes:
-
-.. list-table::
-   :header-rows: 1
-
-   * - Registers in each lane
-     - Contents
-     - Complete matrix across the warp
-   * - ``ar``: four uint32 words
-     - Eight float16 elements, packed in pairs
-     - A: ``(16, 16)``
-   * - ``br``: two uint32 words
-     - Four float16 elements, packed in pairs
-     - B: ``(8, 16)``
-   * - ``dr``: four float32 values
-     - Four output elements
-     - Out: ``(16, 8)``
-
-The instruction fixes which elements belong to each lane. The
-``ldmatrix`` views use coordinates ``(matrix, row, packed_word)``: each
-``8 x 8`` half matrix has eight rows of four words. A contains four such
-matrices and B contains two, giving view shapes ``(4, 8, 4)`` and
-``(2, 8, 4)``. The load layouts place their words in ``ar`` and ``br``.
-
-The ``*_mma_layout`` views describe those same registers with MMA operand
-coordinates. For example, A's ``(1, 1, 16, 8)`` view has one M repeat,
-one K repeat, 16 rows, and eight packed words per row. These views let
-both instructions use the same fragments without moving values between
-lanes. The
-:ref:`warp MMA reference <cuda-tile-contract-mma-sync>` gives the fixed
-lane mappings represented by these layouts.
+.. raw:: html
+
+   <iframe data-tirx-matrix-dataflow
+           src="https://tvm.apache.org/images/docs/tirx/warp_matrix.html";
+           title="Warp matrix multiplication: shared memory, register 
fragments, and output"
+           style="display: block; width: 100%; height: 600px; border: 0;"
+           loading="lazy"></iframe>
+   <script>
+     // Keep both matrix diagrams sized and themed with their documentation 
pages.
+     (() => {
+       const dark = matchMedia('(prefers-color-scheme: dark)');
+       
document.querySelectorAll('iframe[data-tirx-matrix-dataflow]').forEach((frame) 
=> {
+         if (frame.dataset.ready) return;
+         frame.dataset.ready = 'true';
+         const origin = new URL(frame.src, document.baseURI).origin;
+         const syncTheme = () => {
+           const selected = document.documentElement.dataset.theme;
+           const theme = ['light', 'dark'].includes(selected) ? selected : 
(dark.matches ? 'dark' : 'light');
+           frame.contentWindow.postMessage({ type: 
'tirx-matrix-dataflow:theme', theme }, origin);
+         };
+         addEventListener('message', (event) => {
+           if (event.source !== frame.contentWindow || event.origin !== 
origin) return;
+           const data = event.data;
+           if (data?.type !== 'tirx-matrix-dataflow:height') return;
+           if (!Number.isFinite(data.height) || data.height < 100 || 
data.height > 3000) return;
+           frame.style.height = Math.ceil(data.height) + 'px';
+         });
+         new MutationObserver(syncTheme).observe(document.documentElement, {
+           attributes: true, attributeFilter: ['data-theme']
+         });
+         dark.addEventListener('change', syncTheme);
+         frame.addEventListener('load', syncTheme);
+         syncTheme();
+       });
+     })();
+   </script>
+
+A register fragment is the part of a matrix held in one lane's registers.
+The kernel below uses ``ar`` and ``br`` for the input fragments and ``dr``
+for the output fragment:
 
 .. code-block:: python
 
@@ -126,22 +126,24 @@ lane mappings represented by these layouts.
         for r in range(4):
             Out[lane // 4 + 8 * (r // 2), 2 * (lane % 4) + r % 2] = dr[r]
 
-Each thread first copies part of A and B to shared memory. The CTA barrier
-ensures those stores are visible before ``ldmatrix`` reads them. The
-``uint32`` views pack two float16 elements into each instruction operand;
-they preserve the bits rather than convert the values to integers.
-
-``ldmatrix.x4`` loads the four A matrices, and ``ldmatrix.x2`` loads the
-two B matrices. Both instructions require all 32 lanes to participate.
-
-The MMA call takes ``D, A, B, C`` in that order. Its views include repeat
-axes of length one before the matrix dimensions. The selected
-``m16n8k16`` instruction computes the whole product in one issue, with
-``T.float32(0)`` as the initial C value. Each lane receives four output
-elements in ``dr``; the final loop writes them to their matrix positions.
-For lane 0, ``dr[0]`` through ``dr[3]`` belong to ``Out[0, 0]``,
-``Out[0, 1]``, ``Out[8, 0]``, and ``Out[8, 1]``. The other lanes use the
-same coordinate formula to cover the remaining output elements.
+The threads copy ``A`` and ``B`` into shared-memory buffers ``a`` and ``b``,
+then synchronize before ``ldmatrix`` reads them.
+
+The ``uint32`` views reinterpret pairs of float16 values as 32-bit words.
+``rearrange`` expresses the shared-memory views in
+``(matrix, row, packed_word)`` coordinates for ``ldmatrix``.
+All 32 lanes participate in ``ldmatrix.x4`` to load A into ``ar`` and
+``ldmatrix.x2`` to load B into ``br``.
+
+The load and MMA layouts describe the same register fragments in each
+instruction's coordinates. This lets ``mma.sync`` reuse the loaded values
+without moving them between lanes. The :ref:`warp MMA reference
+<cuda-tile-contract-mma-sync>` gives the fixed lane mappings.
+
+The MMA call takes ``D, A, B, C``, with ``T.float32(0)`` as the initial C
+value. The ``m16n8k16`` instruction computes the product in one issue,
+leaving four float32 results per lane in ``dr``. The final loop stores
+them at their corresponding positions in ``Out``.
 
 We can compile the kernel and compare its result with a float32 matrix
 product:
diff --git a/tirx/_sources/reference/cpp/generated/file_tvm_ir_attrs.h.rst 
b/tirx/_sources/reference/cpp/generated/file_tvm_ir_attrs.h.rst
index 9d9b525598e..8ea33e9396b 100644
--- a/tirx/_sources/reference/cpp/generated/file_tvm_ir_attrs.h.rst
+++ b/tirx/_sources/reference/cpp/generated/file_tvm_ir_attrs.h.rst
@@ -116,7 +116,7 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a31d64ff4720a3ed13273776208ab1391`
 
-- :ref:`exhale_function_namespacetvm_1a7e2bc626db8be997b1562c79df3d9e11`
-
 - :ref:`exhale_function_namespacetvm_1a011c96fb07f51688c56daf97a96c0302`
 
+- :ref:`exhale_function_namespacetvm_1a7e2bc626db8be997b1562c79df3d9e11`
+
diff --git a/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_op.h.rst 
b/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_op.h.rst
index d92d12d7668..90ead4bdf06 100644
--- a/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_op.h.rst
+++ b/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_op.h.rst
@@ -143,6 +143,8 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a6a699844e347e400050c1656008de099`
 
+- :ref:`exhale_function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7`
+
 - :ref:`exhale_function_namespacetvm_1af977c9b717c435e3b04fd61333204084`
 
 - :ref:`exhale_function_namespacetvm_1a8c90cd58c6eb3d6b54b92d067ff3ce3a`
@@ -151,19 +153,17 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1aef784ffcb982740d996747683b6e6bad`
 
-- :ref:`exhale_function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7`
+- :ref:`exhale_function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa`
 
-- :ref:`exhale_function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c`
+- :ref:`exhale_function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b`
 
 - :ref:`exhale_function_namespacetvm_1a1871d5104013b80f203905b787ef5508`
 
-- :ref:`exhale_function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b`
-
-- :ref:`exhale_function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef`
+- :ref:`exhale_function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb`
 
-- :ref:`exhale_function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa`
+- :ref:`exhale_function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c`
 
-- :ref:`exhale_function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb`
+- :ref:`exhale_function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef`
 
 - :ref:`exhale_function_namespacetvm_1adca2cf79ae266f512eac9686d152093b`
 
@@ -185,6 +185,10 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a691f5d8db08495baefbe9d4b09a9da46`
 
+- :ref:`exhale_function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120`
+
+- :ref:`exhale_function_namespacetvm_1a534a6e697773ff5607038c8925146d25`
+
 - :ref:`exhale_function_namespacetvm_1aa24e148562e149f345023f01a94fa9e2`
 
 - :ref:`exhale_function_namespacetvm_1acdfbd53bb9384d52e6af0bbb440b9f11`
@@ -193,14 +197,10 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1ae359d9b51f7f8869e5dbaf61c8661cbb`
 
-- :ref:`exhale_function_namespacetvm_1a534a6e697773ff5607038c8925146d25`
-
-- :ref:`exhale_function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120`
+- :ref:`exhale_function_namespacetvm_1a7743df0828a14fc0db059a00df04db37`
 
 - :ref:`exhale_function_namespacetvm_1a88ccf8598f756f100c10500969c21719`
 
-- :ref:`exhale_function_namespacetvm_1a7743df0828a14fc0db059a00df04db37`
-
 - :ref:`exhale_function_namespacetvm_1a5e369b48cc28774b18e3fda6172de78c`
 
 - :ref:`exhale_function_namespacetvm_1aac95b43cceaf788735406b50c9b09f48`
@@ -283,6 +283,8 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a3fc66f0c6cd2135365cd24917d931abe`
 
+- :ref:`exhale_function_namespacetvm_1a5c5034de2993b9130b7bd9d593a11bb5`
+
 - :ref:`exhale_function_namespacetvm_1af70bb4a982810d795dbd17ce73c6b124`
 
 - :ref:`exhale_function_namespacetvm_1ace5dbde3bde1ba48d14a3f9064a45aee`
@@ -293,8 +295,6 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a1815d8b152819885a5733554f374a9ca`
 
-- :ref:`exhale_function_namespacetvm_1a5c5034de2993b9130b7bd9d593a11bb5`
-
 - :ref:`exhale_function_namespacetvm_1a9233d7b436ae52ac070442ba19caed18`
 
 - :ref:`exhale_function_namespacetvm_1ac86d7f61b9f0899a08685145f6bb5051`
@@ -311,6 +311,8 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a96061840d12f84eeecc8fae11e245242`
 
+- :ref:`exhale_function_namespacetvm_1abc417454badf61b154d6a8d87cd8f171`
+
 - :ref:`exhale_function_namespacetvm_1a679ff94dec26779d8769231abb229647`
 
 - :ref:`exhale_function_namespacetvm_1aef861fe5325bc0b415a905a24c42f10a`
@@ -323,16 +325,16 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1abde487c0197942c4ebb1b47277b89dac`
 
-- :ref:`exhale_function_namespacetvm_1abc417454badf61b154d6a8d87cd8f171`
-
 - :ref:`exhale_function_namespacetvm_1a164d454c519fecaa5611c86b65469acc`
 
-- :ref:`exhale_function_namespacetvm_1a136427374941fbf8e50f53b1cab39e38`
-
 - :ref:`exhale_function_namespacetvm_1a18256ba1213ce5ff3cf8037a314354b7`
 
+- :ref:`exhale_function_namespacetvm_1a136427374941fbf8e50f53b1cab39e38`
+
 - :ref:`exhale_function_namespacetvm_1a51dc569142bf8ce8ea55f73029d3807d`
 
+- :ref:`exhale_function_namespacetvm_1a1f98476c3a413f6cdfc7b7e490f3221b`
+
 - :ref:`exhale_function_namespacetvm_1a4c5092e248ab7daa5de5c22717670d8e`
 
 - :ref:`exhale_function_namespacetvm_1abc5d3aba4f3f15098d5ac2fb0c3dfd39`
@@ -343,13 +345,13 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a0854363590c38f5479b1da5e70c4f002`
 
-- :ref:`exhale_function_namespacetvm_1a1f98476c3a413f6cdfc7b7e490f3221b`
+- :ref:`exhale_function_namespacetvm_1af682776c3609284f1bc3ea436e21a67a`
 
 - :ref:`exhale_function_namespacetvm_1a9d8412e5f401f59f5ca85ed556d70810`
 
 - :ref:`exhale_function_namespacetvm_1ad0449d28f23318cc5163159a58c80ba3`
 
-- :ref:`exhale_function_namespacetvm_1af682776c3609284f1bc3ea436e21a67a`
+- :ref:`exhale_function_namespacetvm_1a598f8139c469abc4066dbdd0a0a0845d`
 
 - :ref:`exhale_function_namespacetvm_1a06d97bd5ee2c12e8547be0cc42f6b300`
 
@@ -361,10 +363,10 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a6eea8276bcc178425bc14f3d878970ff`
 
-- :ref:`exhale_function_namespacetvm_1a598f8139c469abc4066dbdd0a0a0845d`
-
 - :ref:`exhale_function_namespacetvm_1a2ea3b45c96d3980227e418f7158ce5c3`
 
+- :ref:`exhale_function_namespacetvm_1ad93d00f7b080dc3f905f5c34c170a041`
+
 - :ref:`exhale_function_namespacetvm_1a9cea8f3789d8f3dc78acae43e9a6aad6`
 
 - :ref:`exhale_function_namespacetvm_1a6d0ad14c882c11311836138a2c164cf3`
@@ -375,7 +377,7 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a7e2181bca182f90533ec35537714d09d`
 
-- :ref:`exhale_function_namespacetvm_1ad93d00f7b080dc3f905f5c34c170a041`
+- :ref:`exhale_function_namespacetvm_1a5530417da455bd46f5dc55f27d69bcdf`
 
 - :ref:`exhale_function_namespacetvm_1a35961a6074b72fae0dfc48ee395e0673`
 
@@ -387,14 +389,12 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1aae1dcfef78728c5490d3c107b4abac5a`
 
-- :ref:`exhale_function_namespacetvm_1a5530417da455bd46f5dc55f27d69bcdf`
+- :ref:`exhale_function_namespacetvm_1a1ce1eb32fc9d76ebe5a6b8d185024d41`
 
 - :ref:`exhale_function_namespacetvm_1abccc3d3e96ee608022b148ffa05034eb`
 
 - :ref:`exhale_function_namespacetvm_1a550c2a251b2a6fd2a72172fe3db75d40`
 
-- :ref:`exhale_function_namespacetvm_1a1ce1eb32fc9d76ebe5a6b8d185024d41`
-
 - :ref:`exhale_function_namespacetvm_1abd7d1b3232218b25e2e0cf6ef699a65f`
 
 - :ref:`exhale_function_namespacetvm_1a82dc2fe21e7a64be5a1b11c2a8775d31`
diff --git 
a/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.rst 
b/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.rst
index d578db62aa6..816794f43e6 100644
--- a/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.rst
+++ b/tirx/_sources/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.rst
@@ -172,12 +172,12 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a5eabe48483942dd9209e0f1970d92966`
 
+- :ref:`exhale_function_namespacetvm_1a90215dca2b70c507d5e32ed71704b54c`
+
 - :ref:`exhale_function_namespacetvm_1ac157e28ac096e2178d5af4f5ac124090`
 
 - :ref:`exhale_function_namespacetvm_1a3beb6ede9dc1c8d2ac91c0ffed0b32f3`
 
-- :ref:`exhale_function_namespacetvm_1a90215dca2b70c507d5e32ed71704b54c`
-
 - :ref:`exhale_function_namespacetvm_1a460572123ac43346571969c7c474d9b3`
 
 - :ref:`exhale_function_namespacetvm_1a7d182ece0a5fe24bf0479f38063f0829`
@@ -238,12 +238,12 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1af6d0ad1d909999b19bf03a951f6c0083`
 
+- :ref:`exhale_function_namespacetvm_1aabc6cf9e961e38f1c2b335bd36277487`
+
 - :ref:`exhale_function_namespacetvm_1a5bea550dd56905267e2e2b0d72a34c62`
 
 - :ref:`exhale_function_namespacetvm_1a81851f002c71fea8695a2e2bd29eefc7`
 
-- :ref:`exhale_function_namespacetvm_1aabc6cf9e961e38f1c2b335bd36277487`
-
 - :ref:`exhale_function_namespacetvm_1a93f1c7a184054d6af9aab47a636e912f`
 
 - :ref:`exhale_function_namespacetvm_1ab480e963072a0d618a1f10d597167545`
@@ -262,12 +262,12 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1adfe0027298faaa905130ff80023b4804`
 
-- :ref:`exhale_function_namespacetvm_1aa99dbf24a4d6e8c09587296bda16ff0d`
-
 - :ref:`exhale_function_namespacetvm_1ac58e284059c962f5a96e46d98fe24e8d`
 
 - :ref:`exhale_function_namespacetvm_1ab042ef46f754a15ae6c332babcd67907`
 
+- :ref:`exhale_function_namespacetvm_1aa99dbf24a4d6e8c09587296bda16ff0d`
+
 - :ref:`exhale_function_namespacetvm_1a34d3774087ee1b5cddd400921762235f`
 
 - :ref:`exhale_function_namespacetvm_1aafc5e5fa4393becb77c4b6d7da3ace2d`
diff --git a/tirx/_sources/reference/cpp/generated/file_tvm_tirx_op_math.h.rst 
b/tirx/_sources/reference/cpp/generated/file_tvm_tirx_op_math.h.rst
index c3d3a5d8832..70041408a6c 100644
--- a/tirx/_sources/reference/cpp/generated/file_tvm_tirx_op_math.h.rst
+++ b/tirx/_sources/reference/cpp/generated/file_tvm_tirx_op_math.h.rst
@@ -57,9 +57,9 @@ Functions
 
 - 
:ref:`exhale_function_namespacetvm_1_1tirx_1aed615c151b18908d314e2f4fd62ae9f7`
 
+- 
:ref:`exhale_function_namespacetvm_1_1tirx_1a4f5fb501032540c31b8249d4b63160c2`
+
 - 
:ref:`exhale_function_namespacetvm_1_1tirx_1ac66dc986673e79d2d9f6bbc4a7dc2d6f`
 
 - 
:ref:`exhale_function_namespacetvm_1_1tirx_1a232ca4bffb36db0d70b967ce516639cd`
 
-- 
:ref:`exhale_function_namespacetvm_1_1tirx_1a4f5fb501032540c31b8249d4b63160c2`
-
diff --git 
a/tirx/_sources/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.rst
 
b/tirx/_sources/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.rst
index 83340d6a3c2..23478a29b42 100644
--- 
a/tirx/_sources/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.rst
+++ 
b/tirx/_sources/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.rst
@@ -86,10 +86,10 @@ Functions
 ---------
 
 
-- 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31`
-
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a74518f928669160b1414e0828eaafbb6`
 
+- 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31`
+
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a91682000d93e5bd2f20173780c529c14`
 
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a774eac1d7ce7a993b13f5e3947a8c8a4`
diff --git a/tirx/_sources/reference/cpp/generated/namespace_tvm.rst 
b/tirx/_sources/reference/cpp/generated/namespace_tvm.rst
index 984529cc6c5..1bf3271a0a5 100644
--- a/tirx/_sources/reference/cpp/generated/namespace_tvm.rst
+++ b/tirx/_sources/reference/cpp/generated/namespace_tvm.rst
@@ -494,6 +494,8 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a6a699844e347e400050c1656008de099`
 
+- :ref:`exhale_function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7`
+
 - :ref:`exhale_function_namespacetvm_1af977c9b717c435e3b04fd61333204084`
 
 - :ref:`exhale_function_namespacetvm_1a8c90cd58c6eb3d6b54b92d067ff3ce3a`
@@ -502,19 +504,17 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1aef784ffcb982740d996747683b6e6bad`
 
-- :ref:`exhale_function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7`
+- :ref:`exhale_function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa`
 
-- :ref:`exhale_function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c`
+- :ref:`exhale_function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b`
 
 - :ref:`exhale_function_namespacetvm_1a1871d5104013b80f203905b787ef5508`
 
-- :ref:`exhale_function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b`
-
-- :ref:`exhale_function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef`
+- :ref:`exhale_function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb`
 
-- :ref:`exhale_function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa`
+- :ref:`exhale_function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c`
 
-- :ref:`exhale_function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb`
+- :ref:`exhale_function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef`
 
 - :ref:`exhale_function_namespacetvm_1a2c34094a0d90d2a554336564b1fba3a2`
 
@@ -542,6 +542,10 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1a43cd8da29b0be251d8d9fc4e93a5db3b`
 
+- :ref:`exhale_function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120`
+
+- :ref:`exhale_function_namespacetvm_1a534a6e697773ff5607038c8925146d25`
+
 - :ref:`exhale_function_namespacetvm_1aa24e148562e149f345023f01a94fa9e2`
 
 - :ref:`exhale_function_namespacetvm_1acdfbd53bb9384d52e6af0bbb440b9f11`
@@ -550,14 +554,10 @@ Functions
 
 - :ref:`exhale_function_namespacetvm_1ae359d9b51f7f8869e5dbaf61c8661cbb`
 
-- :ref:`exhale_function_namespacetvm_1a534a6e697773ff5607038c8925146d25`
-
-- :ref:`exhale_function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120`
+- :ref:`exhale_function_namespacetvm_1a7743df0828a14fc0db059a00df04db37`
 
 - :ref:`exhale_function_namespacetvm_1a88ccf8598f756f100c10500969c21719`
 
-- :ref:`exhale_function_namespacetvm_1a7743df0828a14fc0db059a00df04db37`
-
 - :ref:`exhale_function_namespacetvm_1a5e369b48cc28774b18e3fda6172de78c`
 
 - :ref:`exhale_function_namespacetvm_1aac95b43cceaf788735406b50c9b09f48`
diff --git 
a/tirx/_sources/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.rst
 
b/tirx/_sources/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.rst
index d3aa0624792..ec57e5350b7 100644
--- 
a/tirx/_sources/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.rst
+++ 
b/tirx/_sources/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.rst
@@ -32,10 +32,10 @@ Functions
 ---------
 
 
-- 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31`
-
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a74518f928669160b1414e0828eaafbb6`
 
+- 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31`
+
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a91682000d93e5bd2f20173780c529c14`
 
 - 
:ref:`exhale_function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a774eac1d7ce7a993b13f5e3947a8c8a4`
diff --git a/tirx/commit_hash b/tirx/commit_hash
index 0272750636c..82346eb76b5 100644
--- a/tirx/commit_hash
+++ b/tirx/commit_hash
@@ -1 +1 @@
-a6e9f0ced84f9da1896afb3151926160a7617647
+40d297af8c9b13a5bc890e0e5622a578c21469b2
diff --git a/tirx/programming_guides/tile/blackwell_tensor_cores.html 
b/tirx/programming_guides/tile/blackwell_tensor_cores.html
index c656a52b9cc..f0970614f21 100644
--- a/tirx/programming_guides/tile/blackwell_tensor_cores.html
+++ b/tirx/programming_guides/tile/blackwell_tensor_cores.html
@@ -403,42 +403,44 @@ TMEM holds the accumulator while the Tensor Core reads 
the input tiles from
 shared memory.</p>
 <section id="the-accumulator-s-lifetime">
 <h2>The Accumulator’s Lifetime<a class="headerlink" 
href="#the-accumulator-s-lifetime" title="Link to this heading">#</a></h2>
-<p>The example computes <code class="docutils literal notranslate"><span 
class="pre">Out</span> <span class="pre">=</span> <span class="pre">A</span> 
<span class="pre">&#64;</span> <span class="pre">B.T</span></code> with A of 
shape <code class="docutils literal notranslate"><span class="pre">(128,</span> 
<span class="pre">64)</span></code> and
-B of shape <code class="docutils literal notranslate"><span 
class="pre">(64,</span> <span class="pre">64)</span></code>. One CTA of 128 
threads performs these stages:</p>
-<div class="highlight-text notranslate"><div 
class="highlight"><pre><span></span>allocate TMEM -&gt; stage A/B in shared 
memory -&gt; issue MMA
-                                                |
-                                commit + wait for completion
-                                                |
-release TMEM &lt;- write Out &lt;- load TMEM into registers + wait
-</pre></div>
-</div>
-<p>Different stages require different participants. The conditions in the
-kernel select them explicitly:</p>
-<div class="pst-scrollable-table-container"><table class="table">
-<thead>
-<tr class="row-odd"><th class="head"><p>Participants</p></th>
-<th class="head"><p>Operations</p></th>
-</tr>
-</thead>
-<tbody>
-<tr class="row-even"><td><p><code class="docutils literal notranslate"><span 
class="pre">warp</span> <span class="pre">==</span> <span 
class="pre">0</span></code>: all 32 threads of warp 0</p></td>
-<td><p>Allocate TMEM, relinquish the allocation permit, and deallocate 
TMEM.</p></td>
-</tr>
-<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span 
class="pre">tid</span> <span class="pre">==</span> <span 
class="pre">0</span></code>: thread 0</p></td>
-<td><p>Initialize the barrier, issue the MMA instructions, and commit 
them.</p></td>
-</tr>
-<tr class="row-even"><td><p>All four warps</p></td>
-<td><p>Stage inputs, synchronize, load the completed result from TMEM, and 
write Out.</p></td>
-</tr>
-</tbody>
-</table>
-</div>
-<p><code class="docutils literal notranslate"><span 
class="pre">mma_shared_layout</span></code> constructs the swizzled shared 
layouts for A and B.
-Mode 3 selects a 128-byte swizzle; its 1024-byte period determines the
-shared allocation alignment used below.
-<code class="docutils literal notranslate"><span 
class="pre">tmem_mma_operand_layout</span></code> maps the float32 accumulator 
to TMEM lanes
-and columns for the chosen instruction. The <code class="docutils literal 
notranslate"><span class="pre">128</span> <span class="pre">x</span> <span 
class="pre">64</span></code> accumulator
-uses 64 columns, so the allocation and deallocation both specify 64.</p>
+<p>The example computes <code class="docutils literal notranslate"><span 
class="pre">Out</span> <span class="pre">=</span> <span class="pre">A</span> 
<span class="pre">&#64;</span> <span class="pre">B.T</span></code> with A of 
shape <code class="docutils literal notranslate"><span class="pre">(128,</span> 
<span class="pre">64)</span></code>,
+B of shape <code class="docutils literal notranslate"><span 
class="pre">(64,</span> <span class="pre">64)</span></code>, and Out of shape 
<code class="docutils literal notranslate"><span class="pre">(128,</span> <span 
class="pre">64)</span></code>. One CTA of 128
+threads performs these stages:</p>
+<iframe data-tirx-matrix-dataflow
+        
src="https://tvm.apache.org/images/docs/tirx/blackwell_tensor_cores.html";
+        title="Blackwell matrix multiplication: TMEM allocation, MMA 
completion, and output"
+        style="display: block; width: 100%; height: 740px; border: 0;"
+        loading="lazy"></iframe>
+<script>
+  // Keep both matrix diagrams sized and themed with their documentation pages.
+  (() => {
+    const dark = matchMedia('(prefers-color-scheme: dark)');
+    
document.querySelectorAll('iframe[data-tirx-matrix-dataflow]').forEach((frame) 
=> {
+      if (frame.dataset.ready) return;
+      frame.dataset.ready = 'true';
+      const origin = new URL(frame.src, document.baseURI).origin;
+      const syncTheme = () => {
+        const selected = document.documentElement.dataset.theme;
+        const theme = ['light', 'dark'].includes(selected) ? selected : 
(dark.matches ? 'dark' : 'light');
+        frame.contentWindow.postMessage({ type: 'tirx-matrix-dataflow:theme', 
theme }, origin);
+      };
+      addEventListener('message', (event) => {
+        if (event.source !== frame.contentWindow || event.origin !== origin) 
return;
+        const data = event.data;
+        if (data?.type !== 'tirx-matrix-dataflow:height') return;
+        if (!Number.isFinite(data.height) || data.height < 100 || data.height 
> 3000) return;
+        frame.style.height = Math.ceil(data.height) + 'px';
+      });
+      new MutationObserver(syncTheme).observe(document.documentElement, {
+        attributes: true, attributeFilter: ['data-theme']
+      });
+      dark.addEventListener('change', syncTheme);
+      frame.addEventListener('load', syncTheme);
+      syncTheme();
+    });
+  })();
+</script><p>The kernel uses <code class="docutils literal notranslate"><span 
class="pre">d</span></code> for the TMEM accumulator and <code class="docutils 
literal notranslate"><span class="pre">frag</span></code> for the output
+values held in each thread’s registers:</p>
 <div class="highlight-python notranslate"><div 
class="highlight"><pre><span></span><span class="kn">from</span><span 
class="w"> </span><span class="nn">tvm_tirx.script</span><span class="w"> 
</span><span class="kn">import</span> <span class="n">tirx</span> <span 
class="k">as</span> <span class="n">T</span>
 <span class="kn">from</span><span class="w"> </span><span 
class="nn">tvm_tirx.backend.cuda.axis</span><span class="w"> </span><span 
class="kn">import</span> <span class="n">tid_in_wg</span>
 <span class="kn">from</span><span class="w"> </span><span 
class="nn">tvm_tirx.backend.cuda.layout</span><span class="w"> </span><span 
class="kn">import</span> <span class="n">mma_shared_layout</span><span 
class="p">,</span> <span class="n">tmem_mma_operand_layout</span>
@@ -511,31 +513,32 @@ uses 64 columns, so the allocation and deallocation both 
specify 64.</p>
         <span class="n">T</span><span class="o">.</span><span 
class="n">ptx</span><span class="o">.</span><span class="n">tcgen05</span><span 
class="o">.</span><span class="n">dealloc</span><span class="o">.</span><span 
class="n">cta_group__1</span><span class="o">.</span><span 
class="n">sync</span><span class="o">.</span><span 
class="n">aligned</span><span class="o">.</span><span class="n">b32</span><span 
class="p">(</span><span class="n">address</span><span class="p">[</span><span 
clas [...]
 </pre></div>
 </div>
-<p>The allocation writes a TMEM base address to the shared <code 
class="docutils literal notranslate"><span class="pre">address</span></code> 
slot.
-After the CTA barrier publishes the address, <code class="docutils literal 
notranslate"><span class="pre">decl_tmem</span></code> gives the allocated
-storage a tensor shape and layout. It does not perform another allocation.</p>
-<p>Each thread stages input values in shared memory and executes a proxy
-fence. The following CTA barrier lets the issuing thread proceed only
-after all writers have finished this handoff to the asynchronous reader.
-The A/B views split K=64 into four K=16 slices. For example,
-<code class="docutils literal notranslate"><span class="pre">av</span></code> 
has shape <code class="docutils literal notranslate"><span 
class="pre">(1,</span> <span class="pre">4,</span> <span 
class="pre">128,</span> <span class="pre">16)</span></code>: one M tile, four K 
repeats, and a
-<code class="docutils literal notranslate"><span class="pre">128</span> <span 
class="pre">x</span> <span class="pre">16</span></code> A slice per issue. B 
likewise has one N tile and four K
-repeats. Consequently, the tile call issues four MMA instructions.
-<code class="docutils literal notranslate"><span 
class="pre">accum=False</span></code> starts a fresh result on the first slice; 
subsequent
-slices accumulate into D.</p>
-<p>The elected thread commits these operations to <code class="docutils 
literal notranslate"><span class="pre">mbar</span></code>, initialized with
-one expected arrival. All threads wait for phase 0 and execute
-<code class="docutils literal notranslate"><span 
class="pre">tcgen05.fence::after_thread_sync</span></code> before reading TMEM. 
This completion
-protocol uses a barrier arrival, rather than the expected-byte count
-used by <a class="reference internal" href="tma.html"><span class="doc">TMA 
copies</span></a>.</p>
-<p>The register fragment assigns one output row to each thread. Its
-<code class="docutils literal notranslate"><span class="pre">(4,</span> <span 
class="pre">32,</span> <span class="pre">64)</span></code> view lets each of 
the four warps load its 32 TMEM lanes.
-<code class="docutils literal notranslate"><span 
class="pre">tcgen05.wait::ld</span></code> then makes the loaded registers 
ready for the output
-stores. It waits for these loads; the preceding barrier and fence are
-what establish that MMA has finished. The final CTA barrier keeps TMEM
-alive until every warp has finished using it, after which warp 0 gives
-up further allocations and releases the columns.</p>
-<p>On a supported GPU, we can check the product as follows:</p>
+<p>All 32 threads of warp 0 collectively allocate 64 TMEM columns for the
+<code class="docutils literal notranslate"><span class="pre">128</span> <span 
class="pre">x</span> <span class="pre">64</span></code> float32 accumulator. 
The allocation writes its base address
+to <code class="docutils literal notranslate"><span 
class="pre">address[0]</span></code>.</p>
+<p>The inputs are staged in shared-memory buffers <code class="docutils 
literal notranslate"><span class="pre">a</span></code> and <code 
class="docutils literal notranslate"><span class="pre">b</span></code>.
+<code class="docutils literal notranslate"><span 
class="pre">mma_shared_layout</span></code> mode <code class="docutils literal 
notranslate"><span class="pre">3</span></code> selects a 128-byte swizzle with
+1024-byte base alignment. Each thread executes <code class="docutils literal 
notranslate"><span class="pre">fence.proxy.async</span></code>
+after its copies; the CTA barrier then ensures all inputs are ready
+for MMA and the TMEM address is visible to all threads.
+<code class="docutils literal notranslate"><span 
class="pre">T.cuda.decl_tmem</span></code> binds that allocation to tensor 
<code class="docutils literal notranslate"><span 
class="pre">d</span></code>.</p>
+<p>The views <code class="docutils literal notranslate"><span 
class="pre">av</span></code> and <code class="docutils literal 
notranslate"><span class="pre">bv</span></code> split K=64 into four K=16 
slices, so thread 0’s
+single <code class="docutils literal notranslate"><span 
class="pre">T.cuda.tile.apply</span></code> call expands into four MMA 
instructions.
+With <code class="docutils literal notranslate"><span 
class="pre">accum=False</span></code>, the first instruction writes its product 
into <code class="docutils literal notranslate"><span 
class="pre">d</span></code>;
+the next three add their products to it.</p>
+<p>Thread 0 calls <code class="docutils literal notranslate"><span 
class="pre">tcgen05.commit</span></code> to signal <code class="docutils 
literal notranslate"><span class="pre">mbar</span></code> once all four MMA
+instructions complete, matching the barrier’s arrival count of 1.
+All threads wait for phase 0, then execute
+<code class="docutils literal notranslate"><span 
class="pre">tcgen05.fence::after_thread_sync</span></code> before loading the 
result from TMEM.</p>
+<p><code class="docutils literal notranslate"><span 
class="pre">register_layout</span></code> assigns one output row of 64 float32 
values to each
+thread. All four warps use <code class="docutils literal notranslate"><span 
class="pre">tcgen05.ld</span></code> to load their rows into <code 
class="docutils literal notranslate"><span class="pre">frag</span></code>.
+<code class="docutils literal notranslate"><span 
class="pre">tcgen05.wait::ld</span></code> waits for these loads to complete 
before each thread
+writes its row to <code class="docutils literal notranslate"><span 
class="pre">Out</span></code>.</p>
+<p>The final CTA barrier ensures every warp has finished using TMEM.
+Warp 0 then collectively relinquishes the CTA’s allocation permit and
+deallocates the 64 columns.</p>
+<p>On a supported GPU, we can compare the result with a PyTorch float32
+matrix product:</p>
 <div class="highlight-python notranslate"><div 
class="highlight"><pre><span></span><span class="kn">import</span><span 
class="w"> </span><span class="nn">torch</span>
 <span class="kn">import</span><span class="w"> </span><span 
class="nn">tvm_tirx</span>
 
@@ -580,6 +583,10 @@ before a dependent access or reuse of the affected 
TMEM.</p></td>
 </tbody>
 </table>
 </div>
+<p>For <code class="docutils literal notranslate"><span 
class="pre">tcgen05.mma</span></code> and <code class="docutils literal 
notranslate"><span class="pre">tcgen05.cp</span></code>, a commit signals 
completion through
+a barrier arrival. <a class="reference internal" href="tma.html"><span 
class="doc">TMA loads</span></a> instead track completion using
+an expected-byte count. The barrier setup must match the operation being
+tracked.</p>
 <p><code class="docutils literal notranslate"><span 
class="pre">tcgen05.cp</span></code> is useful for operands that must reside in 
TMEM, such as
 the scale factors of a block-scaled MMA. Its shared source must be ready
 for asynchronous access, just as the MMA inputs above must be. The
diff --git a/tirx/programming_guides/tile/warp_matrix.html 
b/tirx/programming_guides/tile/warp_matrix.html
index b48ba3be2e2..84c446e7c22 100644
--- a/tirx/programming_guides/tile/warp_matrix.html
+++ b/tirx/programming_guides/tile/warp_matrix.html
@@ -405,50 +405,42 @@ example computes <code class="docutils literal 
notranslate"><span class="pre">Ou
 <h2>From Shared Memory to a Matrix Product<a class="headerlink" 
href="#from-shared-memory-to-a-matrix-product" title="Link to this 
heading">#</a></h2>
 <p>All 32 lanes cooperate on one <code class="docutils literal 
notranslate"><span class="pre">16</span> <span class="pre">x</span> <span 
class="pre">8</span></code> output tile. The data follows
 this path:</p>
-<div class="highlight-text notranslate"><div 
class="highlight"><pre><span></span>float16 inputs -&gt; shared memory -&gt; 
ldmatrix -&gt; register fragments
-                                                   |
-                                               mma.sync
-                                                   |
-                                  float32 registers -&gt; Out
-</pre></div>
-</div>
-<p>Each lane holds part of a matrix, called its register fragment. In this
-example, the fragments have the following sizes:</p>
-<div class="pst-scrollable-table-container"><table class="table">
-<thead>
-<tr class="row-odd"><th class="head"><p>Registers in each lane</p></th>
-<th class="head"><p>Contents</p></th>
-<th class="head"><p>Complete matrix across the warp</p></th>
-</tr>
-</thead>
-<tbody>
-<tr class="row-even"><td><p><code class="docutils literal notranslate"><span 
class="pre">ar</span></code>: four uint32 words</p></td>
-<td><p>Eight float16 elements, packed in pairs</p></td>
-<td><p>A: <code class="docutils literal notranslate"><span 
class="pre">(16,</span> <span class="pre">16)</span></code></p></td>
-</tr>
-<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span 
class="pre">br</span></code>: two uint32 words</p></td>
-<td><p>Four float16 elements, packed in pairs</p></td>
-<td><p>B: <code class="docutils literal notranslate"><span 
class="pre">(8,</span> <span class="pre">16)</span></code></p></td>
-</tr>
-<tr class="row-even"><td><p><code class="docutils literal notranslate"><span 
class="pre">dr</span></code>: four float32 values</p></td>
-<td><p>Four output elements</p></td>
-<td><p>Out: <code class="docutils literal notranslate"><span 
class="pre">(16,</span> <span class="pre">8)</span></code></p></td>
-</tr>
-</tbody>
-</table>
-</div>
-<p>The instruction fixes which elements belong to each lane. The
-<code class="docutils literal notranslate"><span 
class="pre">ldmatrix</span></code> views use coordinates <code class="docutils 
literal notranslate"><span class="pre">(matrix,</span> <span 
class="pre">row,</span> <span class="pre">packed_word)</span></code>: each
-<code class="docutils literal notranslate"><span class="pre">8</span> <span 
class="pre">x</span> <span class="pre">8</span></code> half matrix has eight 
rows of four words. A contains four such
-matrices and B contains two, giving view shapes <code class="docutils literal 
notranslate"><span class="pre">(4,</span> <span class="pre">8,</span> <span 
class="pre">4)</span></code> and
-<code class="docutils literal notranslate"><span class="pre">(2,</span> <span 
class="pre">8,</span> <span class="pre">4)</span></code>. The load layouts 
place their words in <code class="docutils literal notranslate"><span 
class="pre">ar</span></code> and <code class="docutils literal 
notranslate"><span class="pre">br</span></code>.</p>
-<p>The <code class="docutils literal notranslate"><span 
class="pre">*_mma_layout</span></code> views describe those same registers with 
MMA operand
-coordinates. For example, A’s <code class="docutils literal notranslate"><span 
class="pre">(1,</span> <span class="pre">1,</span> <span class="pre">16,</span> 
<span class="pre">8)</span></code> view has one M repeat,
-one K repeat, 16 rows, and eight packed words per row. These views let
-both instructions use the same fragments without moving values between
-lanes. The
-<a class="reference internal" 
href="../../reference/python/backend.cuda.tile.instructions.html#cuda-tile-contract-mma-sync"><span
 class="std std-ref">warp MMA reference</span></a> gives the fixed
-lane mappings represented by these layouts.</p>
+<iframe data-tirx-matrix-dataflow
+        src="https://tvm.apache.org/images/docs/tirx/warp_matrix.html";
+        title="Warp matrix multiplication: shared memory, register fragments, 
and output"
+        style="display: block; width: 100%; height: 600px; border: 0;"
+        loading="lazy"></iframe>
+<script>
+  // Keep both matrix diagrams sized and themed with their documentation pages.
+  (() => {
+    const dark = matchMedia('(prefers-color-scheme: dark)');
+    
document.querySelectorAll('iframe[data-tirx-matrix-dataflow]').forEach((frame) 
=> {
+      if (frame.dataset.ready) return;
+      frame.dataset.ready = 'true';
+      const origin = new URL(frame.src, document.baseURI).origin;
+      const syncTheme = () => {
+        const selected = document.documentElement.dataset.theme;
+        const theme = ['light', 'dark'].includes(selected) ? selected : 
(dark.matches ? 'dark' : 'light');
+        frame.contentWindow.postMessage({ type: 'tirx-matrix-dataflow:theme', 
theme }, origin);
+      };
+      addEventListener('message', (event) => {
+        if (event.source !== frame.contentWindow || event.origin !== origin) 
return;
+        const data = event.data;
+        if (data?.type !== 'tirx-matrix-dataflow:height') return;
+        if (!Number.isFinite(data.height) || data.height < 100 || data.height 
> 3000) return;
+        frame.style.height = Math.ceil(data.height) + 'px';
+      });
+      new MutationObserver(syncTheme).observe(document.documentElement, {
+        attributes: true, attributeFilter: ['data-theme']
+      });
+      dark.addEventListener('change', syncTheme);
+      frame.addEventListener('load', syncTheme);
+      syncTheme();
+    });
+  })();
+</script><p>A register fragment is the part of a matrix held in one lane’s 
registers.
+The kernel below uses <code class="docutils literal notranslate"><span 
class="pre">ar</span></code> and <code class="docutils literal 
notranslate"><span class="pre">br</span></code> for the input fragments and 
<code class="docutils literal notranslate"><span class="pre">dr</span></code>
+for the output fragment:</p>
 <div class="highlight-python notranslate"><div 
class="highlight"><pre><span></span><span class="kn">from</span><span 
class="w"> </span><span class="nn">tvm_tirx.script</span><span class="w"> 
</span><span class="kn">import</span> <span class="n">tirx</span> <span 
class="k">as</span> <span class="n">T</span>
 <span class="kn">from</span><span class="w"> </span><span 
class="nn">tvm_tirx.backend.cuda.axis</span><span class="w"> </span><span 
class="kn">import</span> <span class="n">laneid</span>
 <span class="kn">from</span><span class="w"> </span><span 
class="nn">tvm_tirx.tirx.layout</span><span class="w"> </span><span 
class="kn">import</span> <span class="n">S</span><span class="p">,</span> <span 
class="n">TileLayout</span>
@@ -503,20 +495,20 @@ lane mappings represented by these layouts.</p>
         <span class="n">Out</span><span class="p">[</span><span 
class="n">lane</span> <span class="o">//</span> <span class="mi">4</span> <span 
class="o">+</span> <span class="mi">8</span> <span class="o">*</span> <span 
class="p">(</span><span class="n">r</span> <span class="o">//</span> <span 
class="mi">2</span><span class="p">),</span> <span class="mi">2</span> <span 
class="o">*</span> <span class="p">(</span><span class="n">lane</span> <span 
class="o">%</span> <span class="mi">4</span [...]
 </pre></div>
 </div>
-<p>Each thread first copies part of A and B to shared memory. The CTA barrier
-ensures those stores are visible before <code class="docutils literal 
notranslate"><span class="pre">ldmatrix</span></code> reads them. The
-<code class="docutils literal notranslate"><span 
class="pre">uint32</span></code> views pack two float16 elements into each 
instruction operand;
-they preserve the bits rather than convert the values to integers.</p>
-<p><code class="docutils literal notranslate"><span 
class="pre">ldmatrix.x4</span></code> loads the four A matrices, and <code 
class="docutils literal notranslate"><span 
class="pre">ldmatrix.x2</span></code> loads the
-two B matrices. Both instructions require all 32 lanes to participate.</p>
-<p>The MMA call takes <code class="docutils literal notranslate"><span 
class="pre">D,</span> <span class="pre">A,</span> <span class="pre">B,</span> 
<span class="pre">C</span></code> in that order. Its views include repeat
-axes of length one before the matrix dimensions. The selected
-<code class="docutils literal notranslate"><span 
class="pre">m16n8k16</span></code> instruction computes the whole product in 
one issue, with
-<code class="docutils literal notranslate"><span 
class="pre">T.float32(0)</span></code> as the initial C value. Each lane 
receives four output
-elements in <code class="docutils literal notranslate"><span 
class="pre">dr</span></code>; the final loop writes them to their matrix 
positions.
-For lane 0, <code class="docutils literal notranslate"><span 
class="pre">dr[0]</span></code> through <code class="docutils literal 
notranslate"><span class="pre">dr[3]</span></code> belong to <code 
class="docutils literal notranslate"><span class="pre">Out[0,</span> <span 
class="pre">0]</span></code>,
-<code class="docutils literal notranslate"><span class="pre">Out[0,</span> 
<span class="pre">1]</span></code>, <code class="docutils literal 
notranslate"><span class="pre">Out[8,</span> <span 
class="pre">0]</span></code>, and <code class="docutils literal 
notranslate"><span class="pre">Out[8,</span> <span 
class="pre">1]</span></code>. The other lanes use the
-same coordinate formula to cover the remaining output elements.</p>
+<p>The threads copy <code class="docutils literal notranslate"><span 
class="pre">A</span></code> and <code class="docutils literal 
notranslate"><span class="pre">B</span></code> into shared-memory buffers <code 
class="docutils literal notranslate"><span class="pre">a</span></code> and 
<code class="docutils literal notranslate"><span class="pre">b</span></code>,
+then synchronize before <code class="docutils literal notranslate"><span 
class="pre">ldmatrix</span></code> reads them.</p>
+<p>The <code class="docutils literal notranslate"><span 
class="pre">uint32</span></code> views reinterpret pairs of float16 values as 
32-bit words.
+<code class="docutils literal notranslate"><span 
class="pre">rearrange</span></code> expresses the shared-memory views in
+<code class="docutils literal notranslate"><span class="pre">(matrix,</span> 
<span class="pre">row,</span> <span class="pre">packed_word)</span></code> 
coordinates for <code class="docutils literal notranslate"><span 
class="pre">ldmatrix</span></code>.
+All 32 lanes participate in <code class="docutils literal notranslate"><span 
class="pre">ldmatrix.x4</span></code> to load A into <code class="docutils 
literal notranslate"><span class="pre">ar</span></code> and
+<code class="docutils literal notranslate"><span 
class="pre">ldmatrix.x2</span></code> to load B into <code class="docutils 
literal notranslate"><span class="pre">br</span></code>.</p>
+<p>The load and MMA layouts describe the same register fragments in each
+instruction’s coordinates. This lets <code class="docutils literal 
notranslate"><span class="pre">mma.sync</span></code> reuse the loaded values
+without moving them between lanes. The <a class="reference internal" 
href="../../reference/python/backend.cuda.tile.instructions.html#cuda-tile-contract-mma-sync"><span
 class="std std-ref">warp MMA reference</span></a> gives the fixed lane 
mappings.</p>
+<p>The MMA call takes <code class="docutils literal notranslate"><span 
class="pre">D,</span> <span class="pre">A,</span> <span class="pre">B,</span> 
<span class="pre">C</span></code>, with <code class="docutils literal 
notranslate"><span class="pre">T.float32(0)</span></code> as the initial C
+value. The <code class="docutils literal notranslate"><span 
class="pre">m16n8k16</span></code> instruction computes the product in one 
issue,
+leaving four float32 results per lane in <code class="docutils literal 
notranslate"><span class="pre">dr</span></code>. The final loop stores
+them at their corresponding positions in <code class="docutils literal 
notranslate"><span class="pre">Out</span></code>.</p>
 <p>We can compile the kernel and compare its result with a float32 matrix
 product:</p>
 <div class="highlight-python notranslate"><div 
class="highlight"><pre><span></span><span class="kn">import</span><span 
class="w"> </span><span class="nn">torch</span>
diff --git a/tirx/reference/cpp/generated/file_tvm_ir_attrs.h.html 
b/tirx/reference/cpp/generated/file_tvm_ir_attrs.h.html
index eb3cfee7797..4326df14700 100644
--- a/tirx/reference/cpp/generated/file_tvm_ir_attrs.h.html
+++ b/tirx/reference/cpp/generated/file_tvm_ir_attrs.h.html
@@ -449,8 +449,8 @@
 <h2>Functions<a class="headerlink" href="#functions" title="Link to this 
heading">#</a></h2>
 <ul class="simple">
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a31d64ff4720a3ed13273776208ab1391.html#exhale-function-namespacetvm-1a31d64ff4720a3ed13273776208ab1391"><span
 class="std std-ref">Template Function tvm::WithAttrs</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a7e2bc626db8be997b1562c79df3d9e11.html#exhale-function-namespacetvm-1a7e2bc626db8be997b1562c79df3d9e11"><span
 class="std std-ref">Template Function tvm::WithoutAttr(TFunc, const 
std::string&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a011c96fb07f51688c56daf97a96c0302.html#exhale-function-namespacetvm-1a011c96fb07f51688c56daf97a96c0302"><span
 class="std std-ref">Template Function tvm::WithoutAttr(TFunc, 
ffi::reflection::AttrKey&lt;T&gt;)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a7e2bc626db8be997b1562c79df3d9e11.html#exhale-function-namespacetvm-1a7e2bc626db8be997b1562c79df3d9e11"><span
 class="std std-ref">Template Function tvm::WithoutAttr(TFunc, const 
std::string&amp;)</span></a></p></li>
 </ul>
 </section>
 </section>
diff --git a/tirx/reference/cpp/generated/file_tvm_ir_prim_op.h.html 
b/tirx/reference/cpp/generated/file_tvm_ir_prim_op.h.html
index ac864202455..ffaa85f8bce 100644
--- a/tirx/reference/cpp/generated/file_tvm_ir_prim_op.h.html
+++ b/tirx/reference/cpp/generated/file_tvm_ir_prim_op.h.html
@@ -456,17 +456,17 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a5e7d2093417e92997774363760ee3632.html#exhale-function-namespacetvm-1a5e7d2093417e92997774363760ee3632"><span
 class="std std-ref">Function tvm::floormod(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac61810f050d2bf1b8e0f6a9b81ebfa8e.html#exhale-function-namespacetvm-1ac61810f050d2bf1b8e0f6a9b81ebfa8e"><span
 class="std std-ref">Function tvm::floormod(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6a699844e347e400050c1656008de099.html#exhale-function-namespacetvm-1a6a699844e347e400050c1656008de099"><span
 class="std std-ref">Function tvm::greater(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7.html#exhale-function-namespacetvm-1ae24221fc6f06b441fcfbe2ad8e2ecac7"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af977c9b717c435e3b04fd61333204084.html#exhale-function-namespacetvm-1af977c9b717c435e3b04fd61333204084"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a8c90cd58c6eb3d6b54b92d067ff3ce3a.html#exhale-function-namespacetvm-1a8c90cd58c6eb3d6b54b92d067ff3ce3a"><span
 class="std std-ref">Function tvm::greater(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3f318d5322769e405e1ef38a5a70b8fe.html#exhale-function-namespacetvm-1a3f318d5322769e405e1ef38a5a70b8fe"><span
 class="std std-ref">Function tvm::greater(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aef784ffcb982740d996747683b6e6bad.html#exhale-function-namespacetvm-1aef784ffcb982740d996747683b6e6bad"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7.html#exhale-function-namespacetvm-1ae24221fc6f06b441fcfbe2ad8e2ecac7"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c.html#exhale-function-namespacetvm-1afdb60a9a5c7aee496daf0b96e150566c"><span
 class="std std-ref">Function tvm::greater_equal(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a1871d5104013b80f203905b787ef5508.html#exhale-function-namespacetvm-1a1871d5104013b80f203905b787ef5508"><span
 class="std std-ref">Function tvm::greater_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b.html#exhale-function-namespacetvm-1a566a2df8159832da07ec2b0e68459e8b"><span
 class="std std-ref">Function tvm::greater_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef.html#exhale-function-namespacetvm-1a6dfb155a76270e35150565a0d5b48aef"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa.html#exhale-function-namespacetvm-1a8a8f498f0643b9b9651efdff92af9ffa"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b.html#exhale-function-namespacetvm-1a566a2df8159832da07ec2b0e68459e8b"><span
 class="std std-ref">Function tvm::greater_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a1871d5104013b80f203905b787ef5508.html#exhale-function-namespacetvm-1a1871d5104013b80f203905b787ef5508"><span
 class="std std-ref">Function tvm::greater_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb.html#exhale-function-namespacetvm-1a7744c56beeaca3e5ebefd33b47c3d5cb"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c.html#exhale-function-namespacetvm-1afdb60a9a5c7aee496daf0b96e150566c"><span
 class="std std-ref">Function tvm::greater_equal(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef.html#exhale-function-namespacetvm-1a6dfb155a76270e35150565a0d5b48aef"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1adca2cf79ae266f512eac9686d152093b.html#exhale-function-namespacetvm-1adca2cf79ae266f512eac9686d152093b"><span
 class="std std-ref">Function tvm::if_then_else(PrimExpr, PrimExpr, PrimExpr, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3835e89801a9e13a6e7ec0f965c338f8.html#exhale-function-namespacetvm-1a3835e89801a9e13a6e7ec0f965c338f8"><span
 class="std std-ref">Function tvm::indexdiv(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a66b47fb09c9239e1f21b48c3658c7ae3.html#exhale-function-namespacetvm-1a66b47fb09c9239e1f21b48c3658c7ae3"><span
 class="std std-ref">Function tvm::indexdiv(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
@@ -477,14 +477,14 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a43cd8da29b0be251d8d9fc4e93a5db3b.html#exhale-function-namespacetvm-1a43cd8da29b0be251d8d9fc4e93a5db3b"><span
 class="std std-ref">Function tvm::left_shift(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a41ee509ee2a4a8d9ae91969a2af01468.html#exhale-function-namespacetvm-1a41ee509ee2a4a8d9ae91969a2af01468"><span
 class="std std-ref">Function tvm::left_shift(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a691f5d8db08495baefbe9d4b09a9da46.html#exhale-function-namespacetvm-1a691f5d8db08495baefbe9d4b09a9da46"><span
 class="std std-ref">Function tvm::left_shift(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120.html#exhale-function-namespacetvm-1a6f2b313be57b820bf8954df496a2e120"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a534a6e697773ff5607038c8925146d25.html#exhale-function-namespacetvm-1a534a6e697773ff5607038c8925146d25"><span
 class="std std-ref">Function tvm::less(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aa24e148562e149f345023f01a94fa9e2.html#exhale-function-namespacetvm-1aa24e148562e149f345023f01a94fa9e2"><span
 class="std std-ref">Function tvm::less(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1acdfbd53bb9384d52e6af0bbb440b9f11.html#exhale-function-namespacetvm-1acdfbd53bb9384d52e6af0bbb440b9f11"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a73232ab0487607c7e99939cd0120e9ea.html#exhale-function-namespacetvm-1a73232ab0487607c7e99939cd0120e9ea"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ae359d9b51f7f8869e5dbaf61c8661cbb.html#exhale-function-namespacetvm-1ae359d9b51f7f8869e5dbaf61c8661cbb"><span
 class="std std-ref">Function tvm::less(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a534a6e697773ff5607038c8925146d25.html#exhale-function-namespacetvm-1a534a6e697773ff5607038c8925146d25"><span
 class="std std-ref">Function tvm::less(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120.html#exhale-function-namespacetvm-1a6f2b313be57b820bf8954df496a2e120"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a88ccf8598f756f100c10500969c21719.html#exhale-function-namespacetvm-1a88ccf8598f756f100c10500969c21719"><span
 class="std std-ref">Function tvm::less_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7743df0828a14fc0db059a00df04db37.html#exhale-function-namespacetvm-1a7743df0828a14fc0db059a00df04db37"><span
 class="std std-ref">Function tvm::less_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a88ccf8598f756f100c10500969c21719.html#exhale-function-namespacetvm-1a88ccf8598f756f100c10500969c21719"><span
 class="std std-ref">Function tvm::less_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a5e369b48cc28774b18e3fda6172de78c.html#exhale-function-namespacetvm-1a5e369b48cc28774b18e3fda6172de78c"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aac95b43cceaf788735406b50c9b09f48.html#exhale-function-namespacetvm-1aac95b43cceaf788735406b50c9b09f48"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a01ca142e2f7055eaedc956cc814c90ff.html#exhale-function-namespacetvm-1a01ca142e2f7055eaedc956cc814c90ff"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
@@ -526,12 +526,12 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1afc83b50366be0862a6c8f88c9a4e62d5.html#exhale-function-namespacetvm-1afc83b50366be0862a6c8f88c9a4e62d5"><span
 class="std std-ref">Function tvm::operator&amp;&amp;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aca62d5095a23e0e24db1d6611d118eab.html#exhale-function-namespacetvm-1aca62d5095a23e0e24db1d6611d118eab"><span
 class="std std-ref">Function tvm::operator&amp;&amp;(const PrimExpr&amp;, 
bool)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3fc66f0c6cd2135365cd24917d931abe.html#exhale-function-namespacetvm-1a3fc66f0c6cd2135365cd24917d931abe"><span
 class="std std-ref">Function tvm::operator&amp;&amp;(bool, const 
PrimExpr&amp;)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a5c5034de2993b9130b7bd9d593a11bb5.html#exhale-function-namespacetvm-1a5c5034de2993b9130b7bd9d593a11bb5"><span
 class="std std-ref">Function tvm::operator*(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af70bb4a982810d795dbd17ce73c6b124.html#exhale-function-namespacetvm-1af70bb4a982810d795dbd17ce73c6b124"><span
 class="std std-ref">Function tvm::operator*(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ace5dbde3bde1ba48d14a3f9064a45aee.html#exhale-function-namespacetvm-1ace5dbde3bde1ba48d14a3f9064a45aee"><span
 class="std std-ref">Function tvm::operator*(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aca621e1d2df8562819bc021c1410b741.html#exhale-function-namespacetvm-1aca621e1d2df8562819bc021c1410b741"><span
 class="std std-ref">Function tvm::operator*(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6823188ec16be854223bbffe349c975d.html#exhale-function-namespacetvm-1a6823188ec16be854223bbffe349c975d"><span
 class="std std-ref">Function tvm::operator*(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a1815d8b152819885a5733554f374a9ca.html#exhale-function-namespacetvm-1a1815d8b152819885a5733554f374a9ca"><span
 class="std std-ref">Function tvm::operator*(const PrimExpr&amp;, 
double)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a5c5034de2993b9130b7bd9d593a11bb5.html#exhale-function-namespacetvm-1a5c5034de2993b9130b7bd9d593a11bb5"><span
 class="std std-ref">Function tvm::operator*(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a9233d7b436ae52ac070442ba19caed18.html#exhale-function-namespacetvm-1a9233d7b436ae52ac070442ba19caed18"><span
 class="std std-ref">Function tvm::operator*=</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac86d7f61b9f0899a08685145f6bb5051.html#exhale-function-namespacetvm-1ac86d7f61b9f0899a08685145f6bb5051"><span
 class="std std-ref">Function tvm::operator+(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ad728a6c2c3d21242a4df808aadb722eb.html#exhale-function-namespacetvm-1ad728a6c2c3d21242a4df808aadb722eb"><span
 class="std std-ref">Function tvm::operator+(int, const 
PrimExpr&amp;)</span></a></p></li>
@@ -540,48 +540,48 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac225d9ad74343d0e97a5e943e7f6a7d1.html#exhale-function-namespacetvm-1ac225d9ad74343d0e97a5e943e7f6a7d1"><span
 class="std std-ref">Function tvm::operator+(const PrimExpr&amp;, 
double)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af246f441d4ac21b110185b77240b2dcc.html#exhale-function-namespacetvm-1af246f441d4ac21b110185b77240b2dcc"><span
 class="std std-ref">Function tvm::operator+(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a96061840d12f84eeecc8fae11e245242.html#exhale-function-namespacetvm-1a96061840d12f84eeecc8fae11e245242"><span
 class="std std-ref">Function tvm::operator+=</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1abc417454badf61b154d6a8d87cd8f171.html#exhale-function-namespacetvm-1abc417454badf61b154d6a8d87cd8f171"><span
 class="std std-ref">Function tvm::operator-(PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a679ff94dec26779d8769231abb229647.html#exhale-function-namespacetvm-1a679ff94dec26779d8769231abb229647"><span
 class="std std-ref">Function tvm::operator-(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aef861fe5325bc0b415a905a24c42f10a.html#exhale-function-namespacetvm-1aef861fe5325bc0b415a905a24c42f10a"><span
 class="std std-ref">Function tvm::operator-(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a4f40ad3340a853d58664bc864dc10d47.html#exhale-function-namespacetvm-1a4f40ad3340a853d58664bc864dc10d47"><span
 class="std std-ref">Function tvm::operator-(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af7c46ff33a2727f48b10d7d563f4a746.html#exhale-function-namespacetvm-1af7c46ff33a2727f48b10d7d563f4a746"><span
 class="std std-ref">Function tvm::operator-(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a028ba217f99b6cb1592a6a56b2bc9ee5.html#exhale-function-namespacetvm-1a028ba217f99b6cb1592a6a56b2bc9ee5"><span
 class="std std-ref">Function tvm::operator-(const PrimExpr&amp;, 
double)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1abde487c0197942c4ebb1b47277b89dac.html#exhale-function-namespacetvm-1abde487c0197942c4ebb1b47277b89dac"><span
 class="std std-ref">Function tvm::operator-(PrimExpr, 
PrimExpr)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1abc417454badf61b154d6a8d87cd8f171.html#exhale-function-namespacetvm-1abc417454badf61b154d6a8d87cd8f171"><span
 class="std std-ref">Function tvm::operator-(PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a164d454c519fecaa5611c86b65469acc.html#exhale-function-namespacetvm-1a164d454c519fecaa5611c86b65469acc"><span
 class="std std-ref">Function tvm::operator-=</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a136427374941fbf8e50f53b1cab39e38.html#exhale-function-namespacetvm-1a136427374941fbf8e50f53b1cab39e38"><span
 class="std std-ref">Template Function tvm::operator/(const PrimExpr&amp;, 
const TB&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a18256ba1213ce5ff3cf8037a314354b7.html#exhale-function-namespacetvm-1a18256ba1213ce5ff3cf8037a314354b7"><span
 class="std std-ref">Function tvm::operator/(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a136427374941fbf8e50f53b1cab39e38.html#exhale-function-namespacetvm-1a136427374941fbf8e50f53b1cab39e38"><span
 class="std std-ref">Template Function tvm::operator/(const PrimExpr&amp;, 
const TB&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a51dc569142bf8ce8ea55f73029d3807d.html#exhale-function-namespacetvm-1a51dc569142bf8ce8ea55f73029d3807d"><span
 class="std std-ref">Template Function tvm::operator/=</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a1f98476c3a413f6cdfc7b7e490f3221b.html#exhale-function-namespacetvm-1a1f98476c3a413f6cdfc7b7e490f3221b"><span
 class="std std-ref">Function tvm::operator&lt;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a4c5092e248ab7daa5de5c22717670d8e.html#exhale-function-namespacetvm-1a4c5092e248ab7daa5de5c22717670d8e"><span
 class="std std-ref">Function tvm::operator&lt;(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1abc5d3aba4f3f15098d5ac2fb0c3dfd39.html#exhale-function-namespacetvm-1abc5d3aba4f3f15098d5ac2fb0c3dfd39"><span
 class="std std-ref">Function tvm::operator&lt;(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a46877235265ab97544ec2e561f521b0f.html#exhale-function-namespacetvm-1a46877235265ab97544ec2e561f521b0f"><span
 class="std std-ref">Function tvm::operator&lt;(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aa672271dbd566a0e7b9e4c87664bccb4.html#exhale-function-namespacetvm-1aa672271dbd566a0e7b9e4c87664bccb4"><span
 class="std std-ref">Function tvm::operator&lt;(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a0854363590c38f5479b1da5e70c4f002.html#exhale-function-namespacetvm-1a0854363590c38f5479b1da5e70c4f002"><span
 class="std std-ref">Function tvm::operator&lt;(const PrimExpr&amp;, 
double)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a1f98476c3a413f6cdfc7b7e490f3221b.html#exhale-function-namespacetvm-1a1f98476c3a413f6cdfc7b7e490f3221b"><span
 class="std std-ref">Function tvm::operator&lt;(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1af682776c3609284f1bc3ea436e21a67a.html#exhale-function-namespacetvm-1af682776c3609284f1bc3ea436e21a67a"><span
 class="std std-ref">Function tvm::operator&lt;&lt;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a9d8412e5f401f59f5ca85ed556d70810.html#exhale-function-namespacetvm-1a9d8412e5f401f59f5ca85ed556d70810"><span
 class="std std-ref">Function tvm::operator&lt;&lt;(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ad0449d28f23318cc5163159a58c80ba3.html#exhale-function-namespacetvm-1ad0449d28f23318cc5163159a58c80ba3"><span
 class="std std-ref">Function tvm::operator&lt;&lt;(int, const 
PrimExpr&amp;)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1af682776c3609284f1bc3ea436e21a67a.html#exhale-function-namespacetvm-1af682776c3609284f1bc3ea436e21a67a"><span
 class="std std-ref">Function tvm::operator&lt;&lt;(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a598f8139c469abc4066dbdd0a0a0845d.html#exhale-function-namespacetvm-1a598f8139c469abc4066dbdd0a0a0845d"><span
 class="std std-ref">Function tvm::operator&lt;=(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a06d97bd5ee2c12e8547be0cc42f6b300.html#exhale-function-namespacetvm-1a06d97bd5ee2c12e8547be0cc42f6b300"><span
 class="std std-ref">Function tvm::operator&lt;=(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a872f50bd7175eccf440865311aa75232.html#exhale-function-namespacetvm-1a872f50bd7175eccf440865311aa75232"><span
 class="std std-ref">Function tvm::operator&lt;=(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ad5dbec0c48b8644c5c6e9d773ddc106b.html#exhale-function-namespacetvm-1ad5dbec0c48b8644c5c6e9d773ddc106b"><span
 class="std std-ref">Function tvm::operator&lt;=(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af94a56db543e741a23bbf2f51c49091a.html#exhale-function-namespacetvm-1af94a56db543e741a23bbf2f51c49091a"><span
 class="std std-ref">Function tvm::operator&lt;=(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6eea8276bcc178425bc14f3d878970ff.html#exhale-function-namespacetvm-1a6eea8276bcc178425bc14f3d878970ff"><span
 class="std std-ref">Function tvm::operator&lt;=(const PrimExpr&amp;, 
double)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a598f8139c469abc4066dbdd0a0a0845d.html#exhale-function-namespacetvm-1a598f8139c469abc4066dbdd0a0a0845d"><span
 class="std std-ref">Function tvm::operator&lt;=(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a2ea3b45c96d3980227e418f7158ce5c3.html#exhale-function-namespacetvm-1a2ea3b45c96d3980227e418f7158ce5c3"><span
 class="std std-ref">Function tvm::operator==(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1ad93d00f7b080dc3f905f5c34c170a041.html#exhale-function-namespacetvm-1ad93d00f7b080dc3f905f5c34c170a041"><span
 class="std std-ref">Function tvm::operator&gt;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a9cea8f3789d8f3dc78acae43e9a6aad6.html#exhale-function-namespacetvm-1a9cea8f3789d8f3dc78acae43e9a6aad6"><span
 class="std std-ref">Function tvm::operator&gt;(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6d0ad14c882c11311836138a2c164cf3.html#exhale-function-namespacetvm-1a6d0ad14c882c11311836138a2c164cf3"><span
 class="std std-ref">Function tvm::operator&gt;(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6aeb6ed068c5de8ab908ff234337aeeb.html#exhale-function-namespacetvm-1a6aeb6ed068c5de8ab908ff234337aeeb"><span
 class="std std-ref">Function tvm::operator&gt;(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1acc92dcd3d81981e983ddf05347bc9371.html#exhale-function-namespacetvm-1acc92dcd3d81981e983ddf05347bc9371"><span
 class="std std-ref">Function tvm::operator&gt;(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7e2181bca182f90533ec35537714d09d.html#exhale-function-namespacetvm-1a7e2181bca182f90533ec35537714d09d"><span
 class="std std-ref">Function tvm::operator&gt;(const PrimExpr&amp;, 
double)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1ad93d00f7b080dc3f905f5c34c170a041.html#exhale-function-namespacetvm-1ad93d00f7b080dc3f905f5c34c170a041"><span
 class="std std-ref">Function tvm::operator&gt;(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a5530417da455bd46f5dc55f27d69bcdf.html#exhale-function-namespacetvm-1a5530417da455bd46f5dc55f27d69bcdf"><span
 class="std std-ref">Function tvm::operator&gt;=(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a35961a6074b72fae0dfc48ee395e0673.html#exhale-function-namespacetvm-1a35961a6074b72fae0dfc48ee395e0673"><span
 class="std std-ref">Function tvm::operator&gt;=(const PrimExpr&amp;, 
float)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7a94a354cd62137652e09fa887a96100.html#exhale-function-namespacetvm-1a7a94a354cd62137652e09fa887a96100"><span
 class="std std-ref">Function tvm::operator&gt;=(float, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af7dee311b945dfc5a821a119c1db9ad1.html#exhale-function-namespacetvm-1af7dee311b945dfc5a821a119c1db9ad1"><span
 class="std std-ref">Function tvm::operator&gt;=(int, const 
PrimExpr&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac194836fc11a8ba34e44738da17fd116.html#exhale-function-namespacetvm-1ac194836fc11a8ba34e44738da17fd116"><span
 class="std std-ref">Function tvm::operator&gt;=(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aae1dcfef78728c5490d3c107b4abac5a.html#exhale-function-namespacetvm-1aae1dcfef78728c5490d3c107b4abac5a"><span
 class="std std-ref">Function tvm::operator&gt;=(const PrimExpr&amp;, 
double)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a5530417da455bd46f5dc55f27d69bcdf.html#exhale-function-namespacetvm-1a5530417da455bd46f5dc55f27d69bcdf"><span
 class="std std-ref">Function tvm::operator&gt;=(PrimExpr, 
PrimExpr)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a1ce1eb32fc9d76ebe5a6b8d185024d41.html#exhale-function-namespacetvm-1a1ce1eb32fc9d76ebe5a6b8d185024d41"><span
 class="std std-ref">Function tvm::operator&gt;&gt;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1abccc3d3e96ee608022b148ffa05034eb.html#exhale-function-namespacetvm-1abccc3d3e96ee608022b148ffa05034eb"><span
 class="std std-ref">Function tvm::operator&gt;&gt;(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a550c2a251b2a6fd2a72172fe3db75d40.html#exhale-function-namespacetvm-1a550c2a251b2a6fd2a72172fe3db75d40"><span
 class="std std-ref">Function tvm::operator&gt;&gt;(int, const 
PrimExpr&amp;)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a1ce1eb32fc9d76ebe5a6b8d185024d41.html#exhale-function-namespacetvm-1a1ce1eb32fc9d76ebe5a6b8d185024d41"><span
 class="std std-ref">Function tvm::operator&gt;&gt;(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1abd7d1b3232218b25e2e0cf6ef699a65f.html#exhale-function-namespacetvm-1abd7d1b3232218b25e2e0cf6ef699a65f"><span
 class="std std-ref">Function tvm::operator^(PrimExpr, 
PrimExpr)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a82dc2fe21e7a64be5a1b11c2a8775d31.html#exhale-function-namespacetvm-1a82dc2fe21e7a64be5a1b11c2a8775d31"><span
 class="std std-ref">Function tvm::operator^(const PrimExpr&amp;, 
int)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6f638564e5e4d1023096523800f2579e.html#exhale-function-namespacetvm-1a6f638564e5e4d1023096523800f2579e"><span
 class="std std-ref">Function tvm::operator^(int, const 
PrimExpr&amp;)</span></a></p></li>
diff --git a/tirx/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.html 
b/tirx/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.html
index fac324f26a7..3239522e2f1 100644
--- a/tirx/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.html
+++ b/tirx/reference/cpp/generated/file_tvm_ir_prim_pattern_match.h.html
@@ -487,9 +487,9 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac9d50c060c54980dff8f0f62736f7e1f.html#exhale-function-namespacetvm-1ac9d50c060c54980dff8f0f62736f7e1f"><span
 class="std std-ref">Template Function tvm::max(const Pattern&lt;TA&gt;&amp;, 
const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a752f6f3c02b9ae41ce24db2d1a769105.html#exhale-function-namespacetvm-1a752f6f3c02b9ae41ce24db2d1a769105"><span
 class="std std-ref">Template Function tvm::max(const Pattern&lt;TA&gt;&amp;, 
int64_t)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a5eabe48483942dd9209e0f1970d92966.html#exhale-function-namespacetvm-1a5eabe48483942dd9209e0f1970d92966"><span
 class="std std-ref">Template Function tvm::max(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a90215dca2b70c507d5e32ed71704b54c.html#exhale-function-namespacetvm-1a90215dca2b70c507d5e32ed71704b54c"><span
 class="std std-ref">Template Function tvm::min(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac157e28ac096e2178d5af4f5ac124090.html#exhale-function-namespacetvm-1ac157e28ac096e2178d5af4f5ac124090"><span
 class="std std-ref">Template Function tvm::min(const Pattern&lt;TA&gt;&amp;, 
const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3beb6ede9dc1c8d2ac91c0ffed0b32f3.html#exhale-function-namespacetvm-1a3beb6ede9dc1c8d2ac91c0ffed0b32f3"><span
 class="std std-ref">Template Function tvm::min(const Pattern&lt;TA&gt;&amp;, 
int64_t)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a90215dca2b70c507d5e32ed71704b54c.html#exhale-function-namespacetvm-1a90215dca2b70c507d5e32ed71704b54c"><span
 class="std std-ref">Template Function tvm::min(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a460572123ac43346571969c7c474d9b3.html#exhale-function-namespacetvm-1a460572123ac43346571969c7c474d9b3"><span
 class="std std-ref">Template Function tvm::operator!(const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7d182ece0a5fe24bf0479f38063f0829.html#exhale-function-namespacetvm-1a7d182ece0a5fe24bf0479f38063f0829"><span
 class="std std-ref">Template Function tvm::operator!=(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ab5df457f526fb099d0d8de6a51639307.html#exhale-function-namespacetvm-1ab5df457f526fb099d0d8de6a51639307"><span
 class="std std-ref">Template Function tvm::operator!=(const 
Pattern&lt;TA&gt;&amp;, int64_t)</span></a></p></li>
@@ -520,9 +520,9 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1afd9566424968e15c462fb9bc2326c568.html#exhale-function-namespacetvm-1afd9566424968e15c462fb9bc2326c568"><span
 class="std std-ref">Template Function tvm::operator&lt;=(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a509224ee70e43a87c9a9f90050130067.html#exhale-function-namespacetvm-1a509224ee70e43a87c9a9f90050130067"><span
 class="std std-ref">Template Function tvm::operator&lt;=(const 
Pattern&lt;TA&gt;&amp;, int64_t)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af6d0ad1d909999b19bf03a951f6c0083.html#exhale-function-namespacetvm-1af6d0ad1d909999b19bf03a951f6c0083"><span
 class="std std-ref">Template Function tvm::operator&lt;=(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1aabc6cf9e961e38f1c2b335bd36277487.html#exhale-function-namespacetvm-1aabc6cf9e961e38f1c2b335bd36277487"><span
 class="std std-ref">Template Function tvm::operator==(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a5bea550dd56905267e2e2b0d72a34c62.html#exhale-function-namespacetvm-1a5bea550dd56905267e2e2b0d72a34c62"><span
 class="std std-ref">Template Function tvm::operator==(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a81851f002c71fea8695a2e2bd29eefc7.html#exhale-function-namespacetvm-1a81851f002c71fea8695a2e2bd29eefc7"><span
 class="std std-ref">Template Function tvm::operator==(const 
Pattern&lt;TA&gt;&amp;, int64_t)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1aabc6cf9e961e38f1c2b335bd36277487.html#exhale-function-namespacetvm-1aabc6cf9e961e38f1c2b335bd36277487"><span
 class="std std-ref">Template Function tvm::operator==(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a93f1c7a184054d6af9aab47a636e912f.html#exhale-function-namespacetvm-1a93f1c7a184054d6af9aab47a636e912f"><span
 class="std std-ref">Template Function tvm::operator&gt;(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ab480e963072a0d618a1f10d597167545.html#exhale-function-namespacetvm-1ab480e963072a0d618a1f10d597167545"><span
 class="std std-ref">Template Function tvm::operator&gt;(const 
Pattern&lt;TA&gt;&amp;, int64_t)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a1bb911ef90e2f328ba2ae63387b9e936.html#exhale-function-namespacetvm-1a1bb911ef90e2f328ba2ae63387b9e936"><span
 class="std std-ref">Template Function tvm::operator&gt;(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
@@ -532,9 +532,9 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a8c7a3555362a34c57e7eda751b699d6a.html#exhale-function-namespacetvm-1a8c7a3555362a34c57e7eda751b699d6a"><span
 class="std std-ref">Template Function tvm::operator&gt;&gt;(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7fdc1742f400de487cb807737b16aacb.html#exhale-function-namespacetvm-1a7fdc1742f400de487cb807737b16aacb"><span
 class="std std-ref">Template Function tvm::operator^(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1adfe0027298faaa905130ff80023b4804.html#exhale-function-namespacetvm-1adfe0027298faaa905130ff80023b4804"><span
 class="std std-ref">Template Function tvm::operator|(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1aa99dbf24a4d6e8c09587296bda16ff0d.html#exhale-function-namespacetvm-1aa99dbf24a4d6e8c09587296bda16ff0d"><span
 class="std std-ref">Template Function tvm::operator||(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ac58e284059c962f5a96e46d98fe24e8d.html#exhale-function-namespacetvm-1ac58e284059c962f5a96e46d98fe24e8d"><span
 class="std std-ref">Template Function tvm::operator||(const 
Pattern&lt;TA&gt;&amp;, const Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ab042ef46f754a15ae6c332babcd67907.html#exhale-function-namespacetvm-1ab042ef46f754a15ae6c332babcd67907"><span
 class="std std-ref">Template Function tvm::operator||(const 
Pattern&lt;TA&gt;&amp;, int64_t)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1aa99dbf24a4d6e8c09587296bda16ff0d.html#exhale-function-namespacetvm-1aa99dbf24a4d6e8c09587296bda16ff0d"><span
 class="std std-ref">Template Function tvm::operator||(int64_t, const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a34d3774087ee1b5cddd400921762235f.html#exhale-function-namespacetvm-1a34d3774087ee1b5cddd400921762235f"><span
 class="std std-ref">Template Function tvm::operator~(const 
Pattern&lt;TA&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aafc5e5fa4393becb77c4b6d7da3ace2d.html#exhale-function-namespacetvm-1aafc5e5fa4393becb77c4b6d7da3ace2d"><span
 class="std std-ref">Template Function tvm::ramp(const 
Pattern&lt;TBase&gt;&amp;, const Pattern&lt;TStride&gt;&amp;, const 
Pattern&lt;TLanes&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a98593b4c0d2622859925df160ca32f3d.html#exhale-function-namespacetvm-1a98593b4c0d2622859925df160ca32f3d"><span
 class="std std-ref">Template Function tvm::ramp(const 
Pattern&lt;TBase&gt;&amp;, int, int)</span></a></p></li>
diff --git a/tirx/reference/cpp/generated/file_tvm_tirx_op_math.h.html 
b/tirx/reference/cpp/generated/file_tvm_tirx_op_math.h.html
index 54e6935745b..c5f3f7f45c4 100644
--- a/tirx/reference/cpp/generated/file_tvm_tirx_op_math.h.html
+++ b/tirx/reference/cpp/generated/file_tvm_tirx_op_math.h.html
@@ -408,9 +408,9 @@
 <h2>Functions<a class="headerlink" href="#functions" title="Link to this 
heading">#</a></h2>
 <ul class="simple">
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1tirx_1aed615c151b18908d314e2f4fd62ae9f7.html#exhale-function-namespacetvm-1-1tirx-1aed615c151b18908d314e2f4fd62ae9f7"><span
 class="std std-ref">Function tvm::tirx::fast_erf_float_expr</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1_1tirx_1a4f5fb501032540c31b8249d4b63160c2.html#exhale-function-namespacetvm-1-1tirx-1a4f5fb501032540c31b8249d4b63160c2"><span
 class="std std-ref">Function tvm::tirx::logaddexp(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1tirx_1ac66dc986673e79d2d9f6bbc4a7dc2d6f.html#exhale-function-namespacetvm-1-1tirx-1ac66dc986673e79d2d9f6bbc4a7dc2d6f"><span
 class="std std-ref">Function tvm::tirx::logaddexp(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1tirx_1a232ca4bffb36db0d70b967ce516639cd.html#exhale-function-namespacetvm-1-1tirx-1a232ca4bffb36db0d70b967ce516639cd"><span
 class="std std-ref">Function tvm::tirx::logaddexp(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1_1tirx_1a4f5fb501032540c31b8249d4b63160c2.html#exhale-function-namespacetvm-1-1tirx-1a4f5fb501032540c31b8249d4b63160c2"><span
 class="std std-ref">Function tvm::tirx::logaddexp(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 </ul>
 </section>
 </section>
diff --git 
a/tirx/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.html 
b/tirx/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.html
index e2f1c3bf13e..464f4315251 100644
--- a/tirx/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.html
+++ b/tirx/reference/cpp/generated/file_tvm_tirx_script_ir_builder_ir.h.html
@@ -427,8 +427,8 @@
 <section id="functions">
 <h2>Functions<a class="headerlink" href="#functions" title="Link to this 
heading">#</a></h2>
 <ul class="simple">
-<li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1ac307c2206e6fcacb36fe08289e5b5a31"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
Var)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a74518f928669160b1414e0828eaafbb6.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a74518f928669160b1414e0828eaafbb6"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
TensorVar)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1ac307c2206e6fcacb36fe08289e5b5a31"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
Var)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a91682000d93e5bd2f20173780c529c14.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a91682000d93e5bd2f20173780c529c14"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncAttrs</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a774eac1d7ce7a993b13f5e3947a8c8a4.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a774eac1d7ce7a993b13f5e3947a8c8a4"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncName</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1afd5dc6b469056bd40917e93b6def788c.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1afd5dc6b469056bd40917e93b6def788c"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncRet</span></a></p></li>
diff --git a/tirx/reference/cpp/generated/namespace_tvm.html 
b/tirx/reference/cpp/generated/namespace_tvm.html
index 0cdfdc57270..a3bc1c9dc29 100644
--- a/tirx/reference/cpp/generated/namespace_tvm.html
+++ b/tirx/reference/cpp/generated/namespace_tvm.html
@@ -632,17 +632,17 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a9a25946825952db35e05fecf60b06be9.html#exhale-function-namespacetvm-1a9a25946825952db35e05fecf60b06be9"><span
 class="std std-ref">Function tvm::ForKind2String</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a0ebca0080421e4b6d76d0259da74fa52.html#exhale-function-namespacetvm-1a0ebca0080421e4b6d76d0259da74fa52"><span
 class="std std-ref">Function tvm::GetRegionBodyParams</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a6a699844e347e400050c1656008de099.html#exhale-function-namespacetvm-1a6a699844e347e400050c1656008de099"><span
 class="std std-ref">Function tvm::greater(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7.html#exhale-function-namespacetvm-1ae24221fc6f06b441fcfbe2ad8e2ecac7"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1af977c9b717c435e3b04fd61333204084.html#exhale-function-namespacetvm-1af977c9b717c435e3b04fd61333204084"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a8c90cd58c6eb3d6b54b92d067ff3ce3a.html#exhale-function-namespacetvm-1a8c90cd58c6eb3d6b54b92d067ff3ce3a"><span
 class="std std-ref">Function tvm::greater(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3f318d5322769e405e1ef38a5a70b8fe.html#exhale-function-namespacetvm-1a3f318d5322769e405e1ef38a5a70b8fe"><span
 class="std std-ref">Function tvm::greater(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aef784ffcb982740d996747683b6e6bad.html#exhale-function-namespacetvm-1aef784ffcb982740d996747683b6e6bad"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1ae24221fc6f06b441fcfbe2ad8e2ecac7.html#exhale-function-namespacetvm-1ae24221fc6f06b441fcfbe2ad8e2ecac7"><span
 class="std std-ref">Function tvm::greater(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c.html#exhale-function-namespacetvm-1afdb60a9a5c7aee496daf0b96e150566c"><span
 class="std std-ref">Function tvm::greater_equal(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a1871d5104013b80f203905b787ef5508.html#exhale-function-namespacetvm-1a1871d5104013b80f203905b787ef5508"><span
 class="std std-ref">Function tvm::greater_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b.html#exhale-function-namespacetvm-1a566a2df8159832da07ec2b0e68459e8b"><span
 class="std std-ref">Function tvm::greater_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef.html#exhale-function-namespacetvm-1a6dfb155a76270e35150565a0d5b48aef"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a8a8f498f0643b9b9651efdff92af9ffa.html#exhale-function-namespacetvm-1a8a8f498f0643b9b9651efdff92af9ffa"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a566a2df8159832da07ec2b0e68459e8b.html#exhale-function-namespacetvm-1a566a2df8159832da07ec2b0e68459e8b"><span
 class="std std-ref">Function tvm::greater_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a1871d5104013b80f203905b787ef5508.html#exhale-function-namespacetvm-1a1871d5104013b80f203905b787ef5508"><span
 class="std std-ref">Function tvm::greater_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7744c56beeaca3e5ebefd33b47c3d5cb.html#exhale-function-namespacetvm-1a7744c56beeaca3e5ebefd33b47c3d5cb"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1afdb60a9a5c7aee496daf0b96e150566c.html#exhale-function-namespacetvm-1afdb60a9a5c7aee496daf0b96e150566c"><span
 class="std std-ref">Function tvm::greater_equal(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a6dfb155a76270e35150565a0d5b48aef.html#exhale-function-namespacetvm-1a6dfb155a76270e35150565a0d5b48aef"><span
 class="std std-ref">Function tvm::greater_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a2c34094a0d90d2a554336564b1fba3a2.html#exhale-function-namespacetvm-1a2c34094a0d90d2a554336564b1fba3a2"><span
 class="std std-ref">Template Function tvm::if_then_else(const 
Pattern&lt;TCond&gt;&amp;, const Pattern&lt;TA&gt;&amp;, const 
Pattern&lt;TB&gt;&amp;)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1adca2cf79ae266f512eac9686d152093b.html#exhale-function-namespacetvm-1adca2cf79ae266f512eac9686d152093b"><span
 class="std std-ref">Function tvm::if_then_else(PrimExpr, PrimExpr, PrimExpr, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a3835e89801a9e13a6e7ec0f965c338f8.html#exhale-function-namespacetvm-1a3835e89801a9e13a6e7ec0f965c338f8"><span
 class="std std-ref">Function tvm::indexdiv(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
@@ -656,14 +656,14 @@
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a41ee509ee2a4a8d9ae91969a2af01468.html#exhale-function-namespacetvm-1a41ee509ee2a4a8d9ae91969a2af01468"><span
 class="std std-ref">Function tvm::left_shift(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a691f5d8db08495baefbe9d4b09a9da46.html#exhale-function-namespacetvm-1a691f5d8db08495baefbe9d4b09a9da46"><span
 class="std std-ref">Function tvm::left_shift(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a43cd8da29b0be251d8d9fc4e93a5db3b.html#exhale-function-namespacetvm-1a43cd8da29b0be251d8d9fc4e93a5db3b"><span
 class="std std-ref">Function tvm::left_shift(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120.html#exhale-function-namespacetvm-1a6f2b313be57b820bf8954df496a2e120"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a534a6e697773ff5607038c8925146d25.html#exhale-function-namespacetvm-1a534a6e697773ff5607038c8925146d25"><span
 class="std std-ref">Function tvm::less(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aa24e148562e149f345023f01a94fa9e2.html#exhale-function-namespacetvm-1aa24e148562e149f345023f01a94fa9e2"><span
 class="std std-ref">Function tvm::less(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1acdfbd53bb9384d52e6af0bbb440b9f11.html#exhale-function-namespacetvm-1acdfbd53bb9384d52e6af0bbb440b9f11"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a73232ab0487607c7e99939cd0120e9ea.html#exhale-function-namespacetvm-1a73232ab0487607c7e99939cd0120e9ea"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1ae359d9b51f7f8869e5dbaf61c8661cbb.html#exhale-function-namespacetvm-1ae359d9b51f7f8869e5dbaf61c8661cbb"><span
 class="std std-ref">Function tvm::less(PrimExpr, PrimExpr, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a534a6e697773ff5607038c8925146d25.html#exhale-function-namespacetvm-1a534a6e697773ff5607038c8925146d25"><span
 class="std std-ref">Function tvm::less(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a6f2b313be57b820bf8954df496a2e120.html#exhale-function-namespacetvm-1a6f2b313be57b820bf8954df496a2e120"><span
 class="std std-ref">Function tvm::less(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
-<li><p><a class="reference internal" 
href="function_namespacetvm_1a88ccf8598f756f100c10500969c21719.html#exhale-function-namespacetvm-1a88ccf8598f756f100c10500969c21719"><span
 class="std std-ref">Function tvm::less_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a7743df0828a14fc0db059a00df04db37.html#exhale-function-namespacetvm-1a7743df0828a14fc0db059a00df04db37"><span
 class="std std-ref">Function tvm::less_equal(float, const PrimExpr&amp;, 
Location)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1a88ccf8598f756f100c10500969c21719.html#exhale-function-namespacetvm-1a88ccf8598f756f100c10500969c21719"><span
 class="std std-ref">Function tvm::less_equal(int, const PrimExpr&amp;, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a5e369b48cc28774b18e3fda6172de78c.html#exhale-function-namespacetvm-1a5e369b48cc28774b18e3fda6172de78c"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, float, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1aac95b43cceaf788735406b50c9b09f48.html#exhale-function-namespacetvm-1aac95b43cceaf788735406b50c9b09f48"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, int, 
Location)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1a01ca142e2f7055eaedc956cc814c90ff.html#exhale-function-namespacetvm-1a01ca142e2f7055eaedc956cc814c90ff"><span
 class="std std-ref">Function tvm::less_equal(const PrimExpr&amp;, double, 
Location)</span></a></p></li>
diff --git 
a/tirx/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.html 
b/tirx/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.html
index 10d7254d393..d063a5b7394 100644
--- a/tirx/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.html
+++ b/tirx/reference/cpp/generated/namespace_tvm__script__ir_builder__tirx.html
@@ -390,8 +390,8 @@
 <section id="functions">
 <h2>Functions<a class="headerlink" href="#functions" title="Link to this 
heading">#</a></h2>
 <ul class="simple">
-<li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1ac307c2206e6fcacb36fe08289e5b5a31"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
Var)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a74518f928669160b1414e0828eaafbb6.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a74518f928669160b1414e0828eaafbb6"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
TensorVar)</span></a></p></li>
+<li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1ac307c2206e6fcacb36fe08289e5b5a31.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1ac307c2206e6fcacb36fe08289e5b5a31"><span
 class="std std-ref">Function tvm::script::ir_builder::tirx::Arg(ffi::String, 
Var)</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a91682000d93e5bd2f20173780c529c14.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a91682000d93e5bd2f20173780c529c14"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncAttrs</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1a774eac1d7ce7a993b13f5e3947a8c8a4.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1a774eac1d7ce7a993b13f5e3947a8c8a4"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncName</span></a></p></li>
 <li><p><a class="reference internal" 
href="function_namespacetvm_1_1script_1_1ir__builder_1_1tirx_1afd5dc6b469056bd40917e93b6def788c.html#exhale-function-namespacetvm-1-1script-1-1ir-builder-1-1tirx-1afd5dc6b469056bd40917e93b6def788c"><span
 class="std std-ref">Function 
tvm::script::ir_builder::tirx::FuncRet</span></a></p></li>
diff --git a/tirx/reference/python/generated/tvm_tirx.tirx.tile.Operand.html 
b/tirx/reference/python/generated/tvm_tirx.tirx.tile.Operand.html
index 87edbc24c3d..28a872dcdf0 100644
--- a/tirx/reference/python/generated/tvm_tirx.tirx.tile.Operand.html
+++ b/tirx/reference/python/generated/tvm_tirx.tirx.tile.Operand.html
@@ -430,7 +430,7 @@
 <h1>tvm_tirx.tirx.tile.Operand<a class="headerlink" 
href="#tvm-tirx-tirx-tile-operand" title="Link to this heading">#</a></h1>
 <dl class="py class">
 <dt class="sig sig-object py" id="tvm_tirx.tirx.tile.Operand">
-<span class="property"><span class="k"><span 
class="pre">class</span></span><span class="w"> </span></span><span 
class="sig-prename descclassname"><span 
class="pre">tvm_tirx.tirx.tile.</span></span><span class="sig-name 
descname"><span class="pre">Operand</span></span><span 
class="sig-paren">(</span><em class="sig-param"><span class="n"><span 
class="pre">name</span></span><span class="p"><span 
class="pre">:</span></span><span class="w"> </span><span class="n"><span 
class="pre">str</span> [...]
+<span class="property"><span class="k"><span 
class="pre">class</span></span><span class="w"> </span></span><span 
class="sig-prename descclassname"><span 
class="pre">tvm_tirx.tirx.tile.</span></span><span class="sig-name 
descname"><span class="pre">Operand</span></span><span 
class="sig-paren">(</span><em class="sig-param"><span class="n"><span 
class="pre">name</span></span><span class="p"><span 
class="pre">:</span></span><span class="w"> </span><span class="n"><span 
class="pre">str</span> [...]
 <dd><p>Bases: <code class="xref py py-class docutils literal 
notranslate"><span class="pre">object</span></code></p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
diff --git a/tirx/searchindex.js b/tirx/searchindex.js
index e6d9bf71200..cc333c71d94 100644
--- a/tirx/searchindex.js
+++ b/tirx/searchindex.js
@@ -1 +1 @@
-Search.setIndex({"alltitles":{"A CUDA Function":[[23,"a-cuda-function"]],"A 
Global-to-Shared Copy":[[30,"a-global-to-shared-copy"]],"A Load, Compute, and 
Store Example":[[34,"a-load-compute-and-store-example"]],"A Shared-Memory 
Exchange":[[32,"a-shared-memory-exchange"]],"A Tile Load, Computation, and 
Store":[[35,"a-tile-load-computation-and-store"]],"API 
Reference":[[10,null]],"Absolute value and 
negation":[[1537,"absolute-value-and-negation"]],"Accumulation, transpose and 
alpha/beta":[ [...]
\ No newline at end of file
+Search.setIndex({"alltitles":{"A CUDA Function":[[23,"a-cuda-function"]],"A 
Global-to-Shared Copy":[[30,"a-global-to-shared-copy"]],"A Load, Compute, and 
Store Example":[[34,"a-load-compute-and-store-example"]],"A Shared-Memory 
Exchange":[[32,"a-shared-memory-exchange"]],"A Tile Load, Computation, and 
Store":[[35,"a-tile-load-computation-and-store"]],"API 
Reference":[[10,null]],"Absolute value and 
negation":[[1537,"absolute-value-and-negation"]],"Accumulation, transpose and 
alpha/beta":[ [...]
\ No newline at end of file