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

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


The following commit(s) were added to refs/heads/main by this push:
     new 8a4a446368c [DOCS] Add self-contained HTML diagrams for TIRx matrix 
guides (#79)
8a4a446368c is described below

commit 8a4a446368ca483de17667142fc9e3dad0455ae8
Author: Shushi Hong <[email protected]>
AuthorDate: Thu Nov 5 22:38:53 2026 -0500

    [DOCS] Add self-contained HTML diagrams for TIRx matrix guides (#79)
    
    * [DOCS] Add HTML diagrams for TIRx matrix instruction guides
    
    * [DOCS] Make matrix diagrams self-contained HTML
---
 images/docs/tirx/blackwell_tensor_cores.html | 370 +++++++++++++++++++++++++++
 images/docs/tirx/warp_matrix.html            | 350 +++++++++++++++++++++++++
 2 files changed, 720 insertions(+)

diff --git a/images/docs/tirx/blackwell_tensor_cores.html 
b/images/docs/tirx/blackwell_tensor_cores.html
new file mode 100644
index 00000000000..d50b2582691
--- /dev/null
+++ b/images/docs/tirx/blackwell_tensor_cores.html
@@ -0,0 +1,370 @@
+<!DOCTYPE html>
+<!--
+  Licensed to the Apache Software Foundation (ASF) under one
+  or more contributor license agreements. See the NOTICE file
+  distributed with this work for additional information
+  regarding copyright ownership. The ASF licenses this file
+  to you under the Apache License, Version 2.0 (the
+  "License"); you may not use this file except in compliance
+  with the License. You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+  Unless required by applicable law or agreed to in writing,
+  software distributed under the License is distributed on an
+  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+  KIND, either express or implied. See the License for the
+  specific language governing permissions and limitations
+  under the License.
+-->
+<html lang="en">
+<head>
+  <meta charset="utf-8">
+  <meta name="viewport" content="width=device-width, initial-scale=1">
+  <title>Blackwell matrix multiplication and TMEM lifetime</title>
+  <style>
+    /* Match the typography, blue cards, and neutral frame in 
authoring_layers.html. */
+    :root {
+      color-scheme: light;
+      --surface: #fff;
+      --text: #222;
+      --muted: #536276;
+      --frame: #b6bdc6;
+      --frame-bg: #f8fafc;
+      --tag: #475569;
+      --line: #64748b;
+      --layer-border: #5b9bd5;
+      --layer-bg: #eef5fc;
+      --a-bg: #e2effc;
+      --a-border: #7a9fc7;
+      --b-bg: #eee7f8;
+      --b-border: #a38bc2;
+    }
+    :root[data-theme="dark"] {
+      color-scheme: dark;
+      --surface: #14181e;
+      --text: #e5e9ef;
+      --muted: #b1bdcd;
+      --frame: #627084;
+      --frame-bg: #1c232c;
+      --line: #a1afc1;
+      --layer-bg: #1d3043;
+      --a-bg: #243c54;
+      --a-border: #7cadd7;
+      --b-bg: #3c304e;
+      --b-border: #b09bd0;
+    }
+    * { box-sizing: border-box; }
+    html, body { margin: 0; }
+    body {
+      padding: 14px 4px 4px;
+      background: var(--surface);
+      color: var(--text);
+      font: 13px/1.5 'Inter', 'SF Pro', 'Segoe UI', system-ui, sans-serif;
+    }
+    h1, h2, p { margin: 0; }
+    h2 { font-size: 14px; font-weight: 700; }
+    code { font: 12px/1.5 'SF Mono', 'Fira Code', Consolas, monospace; }
+    .diagram {
+      position: relative;
+      max-width: 920px;
+      margin: 0 auto;
+      padding: 26px 18px 16px;
+      border: 2px solid var(--frame);
+      border-radius: 12px;
+      background: var(--frame-bg);
+    }
+    .compact-diagram { max-width: 680px; }
+    .scope {
+      position: absolute;
+      top: -11px;
+      left: 16px;
+      padding: 1px 10px;
+      border-radius: 4px;
+      background: var(--tag);
+      color: #fff;
+      font-size: 11px;
+      font-weight: 700;
+      letter-spacing: .02em;
+    }
+    .card {
+      padding: 10px 14px;
+      border: 2px solid var(--layer-border);
+      border-radius: 8px;
+      background: var(--layer-bg);
+      text-align: center;
+    }
+    .card h2 + p, .card h2 + .pair { margin-top: 4px; }
+    .endpoint { border-color: var(--frame); background: var(--surface); }
+    .pair { display: grid; grid-template-columns: 1fr 1fr; gap: 0; }
+    .pair > :first-child { padding-right: 12px; }
+    .pair > :last-child { padding-left: 12px; border-left: 1px solid 
var(--frame); }
+    .card:not(.endpoint) .pair:not(.operands) > :last-child { border-color: 
var(--layer-border); }
+    .card h2 + .pair.operands { margin-top: 8px; }
+    .pair.operands > p {
+      margin: 0 6px;
+      padding: 6px 8px;
+      border: 1px solid var(--operand-border);
+      border-radius: 6px;
+      background: var(--operand-bg);
+    }
+    .operand-a { --operand-bg: var(--a-bg); --operand-border: var(--a-border); 
}
+    .operand-b { --operand-bg: var(--b-bg); --operand-border: var(--b-border); 
}
+    .operand-detail { display: block; }
+    .note { color: var(--muted); font-size: 11px; }
+    .note code { font-size: inherit; }
+    .packing, .coverage { margin-top: 6px; }
+    .equation { margin-bottom: 14px; text-align: center; }
+    .equation code { font-size: 15px; font-weight: 600; }
+    .parallel {
+      display: grid;
+      grid-template-columns: minmax(0, 1fr) minmax(0, 1fr);
+      padding: 0 16px;
+    }
+    .sync-note {
+      grid-column: 1 / -1;
+      position: relative;
+      padding: 10px 0 0;
+      text-align: center;
+    }
+    .sync-note::before, .sync-note::after {
+      content: '';
+      position: absolute;
+      top: 0;
+      bottom: 0;
+      border-left: 1.5px solid var(--line);
+    }
+    .sync-note::before { left: 25%; }
+    .sync-note::after { left: 75%; }
+    .sync-note .operation { display: inline-block; z-index: 1; }
+    .merge {
+      height: 20px;
+      margin: 0 calc(25% + 8px);
+      border: solid var(--line);
+      border-width: 0 1.5px 1.5px;
+    }
+    .connector {
+      position: relative;
+      display: flex;
+      flex-direction: column;
+      align-items: center;
+      justify-content: center;
+      min-height: 50px;
+      padding: 8px 0 13px;
+      text-align: center;
+    }
+    .connector::before {
+      content: '';
+      position: absolute;
+      left: 50%;
+      top: 0;
+      bottom: 3px;
+      border-left: 1.5px solid var(--line);
+    }
+    .connector::after {
+      content: '';
+      position: absolute;
+      left: calc(50% - 3.25px);
+      bottom: 0;
+      border-left: 4px solid transparent;
+      border-right: 4px solid transparent;
+      border-top: 6px solid var(--line);
+    }
+    .operation {
+      position: relative;
+      max-width: 100%;
+      padding: 1px 10px;
+      background: var(--frame-bg);
+    }
+    .operation > code, .operation > strong { font-weight: 600; }
+    .handoff { gap: 12px; }
+    .lifetime {
+      display: grid;
+      grid-template-columns: 1fr 1fr;
+      gap: 8px 12px;
+      margin-bottom: 8px;
+      text-align: center;
+    }
+    .lifetime > section {
+      padding: 5px 8px;
+      border: 1px solid var(--frame);
+      border-radius: 6px;
+      background: var(--surface);
+    }
+    .lifetime h2 { font-size: 12px; }
+    .blackwell-diagram .handoff { gap: 10px; }
+    .release {
+      margin-top: 14px;
+      padding-top: 10px;
+      border-top: 1px solid var(--frame);
+      text-align: center;
+    }
+    .sr-only {
+      position: absolute;
+      width: 1px;
+      height: 1px;
+      padding: 0;
+      overflow: hidden;
+      clip-path: inset(50%);
+      white-space: nowrap;
+    }
+    /* Use the same compact vertical layout for both matrix examples. */
+    @media (min-width: 640px) {
+      .compact-diagram { padding: 20px 10px 12px; }
+      .compact-diagram .equation { margin-bottom: 8px; }
+      .compact-diagram .card { padding-block: 6px; }
+      .compact-diagram .card h2 + .pair.operands { margin-top: 4px; }
+      .compact-diagram .pair.operands > p { padding-block: 4px; }
+      .compact-diagram .operand-detail { display: inline-block; }
+      .compact-diagram .operand-detail::before { content: ' · '; }
+      .compact-diagram .connector { min-height: 34px; padding: 3px 0 10px; }
+      .compact-diagram .sync-note { padding-top: 6px; }
+      .compact-diagram .compact-operation {
+        display: inline-flex;
+        flex-wrap: wrap;
+        align-items: baseline;
+        justify-content: center;
+        gap: 0 8px;
+      }
+      .compact-diagram .merge { height: 12px; }
+      .compact-diagram .packing, .compact-diagram .coverage { margin-top: 4px; 
}
+      .compact-diagram .single-output h2,
+      .compact-diagram .single-output > p:not(.note) { display: inline; }
+      .compact-diagram .single-output > p:not(.note) { margin-left: 8px; }
+      .blackwell-diagram .release { margin-top: 8px; padding-top: 6px; }
+      .blackwell-diagram .release h2 { font-size: 12px; }
+    }
+    @media (max-width: 480px) {
+      .diagram { padding: 24px 10px 12px; }
+      .card { padding: 9px 8px; }
+      .pair > :first-child { padding-right: 8px; }
+      .pair > :last-child { padding-left: 8px; }
+      .operation { padding: 2px 6px; }
+      .lifetime { gap: 8px 12px; }
+      .parallel { padding: 0 10px; }
+      .merge { margin-inline: calc(25% + 5px); }
+    }
+  </style>
+</head>
+<body>
+  <main class="diagram compact-diagram blackwell-diagram" 
aria-labelledby="scope" aria-describedby="description">
+    <h1 class="scope" id="scope">One CTA · 128 threads</h1>
+    <p class="sr-only" id="description">
+      The CTA computes Out = A @ B.T. Warp 0 allocates TMEM and thread 0
+      initializes the completion barrier mbar. All 128 threads copy A to shared
+      memory a; threads 0 through 63 copy B to shared memory b. All threads
+      execute a proxy fence and synchronize. Thread 0 issues four MMA
+      instructions that read a and b directly and accumulate into TMEM d,
+      then commits them to mbar. All threads wait and execute the thread-sync
+      fence. All four warps load TMEM into register fragments frag and wait for
+      the loads. Each thread stores one row of Out. A CTA barrier precedes
+      TMEM release by warp 0.
+    </p>
+    <p class="equation"><code>Out = A @ B.T</code></p>
+    <div class="lifetime">
+      <section>
+        <h2>Allocate TMEM</h2>
+        <p class="note">Warp 0 · 64 columns</p>
+      </section>
+      <section>
+        <h2>Initialize barrier</h2>
+        <p class="note">Thread 0 · <code>mbar</code></p>
+      </section>
+    </div>
+    <section class="card endpoint">
+      <h2>Inputs · global memory</h2>
+      <div class="pair operands">
+        <p class="operand-a"><code>A</code> · 128 × 64 · float16</p>
+        <p class="operand-b"><code>B</code> · 64 × 64 · float16</p>
+      </div>
+    </section>
+    <div class="parallel" aria-label="Cooperative copy of A and B to shared 
memory">
+      <div class="connector">
+        <div class="operation compact-operation">Copy A <span class="note">All 
128 threads</span></div>
+      </div>
+      <div class="connector">
+        <div class="operation compact-operation">Copy B <span 
class="note">Threads 0–63</span></div>
+      </div>
+    </div>
+    <section class="card">
+      <h2>Shared memory</h2>
+      <div class="pair operands">
+        <p class="operand-a"><strong>Shared copy of A 
(<code>a</code>)</strong><span class="operand-detail">128 × 64 · 
float16</span></p>
+        <p class="operand-b"><strong>Shared copy of B 
(<code>b</code>)</strong><span class="operand-detail">64 × 64 · 
float16</span></p>
+      </div>
+    </section>
+    <div class="parallel">
+      <div class="sync-note">
+        <div class="operation compact-operation">
+          <strong>Proxy fence + CTA barrier</strong>
+          <p class="note">All 128 threads · shared inputs ready for MMA</p>
+        </div>
+      </div>
+    </div>
+    <div class="merge" aria-hidden="true"></div>
+    <div class="connector">
+      <div class="operation compact-operation">
+        <code>tcgen05.mma</code>
+        <p class="note">Thread 0 · four K = 16 slices</p>
+      </div>
+    </div>
+    <section class="card single-output">
+      <h2>Accumulator · TMEM</h2>
+      <p><code>d</code> · 128 × 64 · float32</p>
+      <p class="note coverage">The first slice starts a fresh result; the next 
three accumulate.</p>
+    </section>
+    <div class="connector handoff">
+      <div class="operation">
+        <strong>Wait for MMA completion</strong>
+        <p class="note">Thread 0 commits to <code>mbar</code>; all 128 threads 
wait, then fence.</p>
+        <p class="note"><code>tcgen05.fence::after_thread_sync</code></p>
+      </div>
+      <div class="operation">
+        <p><code>tcgen05.ld</code> + <code>tcgen05.wait::ld</code></p>
+        <p class="note">All four warps · load TMEM into registers, then wait 
for the loads.</p>
+      </div>
+    </div>
+    <section class="card single-output">
+      <h2>Output register fragments</h2>
+      <p><code>frag</code> · 64 float32 values per thread</p>
+      <p class="note coverage">128 threads × 64 values = one 128 × 64 output 
tile</p>
+    </section>
+    <div class="connector">
+      <div class="operation">Each thread stores one output row</div>
+    </div>
+    <section class="card endpoint single-output">
+      <h2>Out · global memory</h2>
+      <p>128 × 64 · float32</p>
+    </section>
+    <section class="release">
+      <h2>Release TMEM</h2>
+      <p class="note">After a CTA barrier, warp 0 relinquishes the permit and 
deallocates 64 columns.</p>
+    </section>
+  </main>
+  <script>
+    // Report the diagram's natural height and receive the documentation theme.
+    (() => {
+      const root = document.documentElement;
+      const dark = matchMedia('(prefers-color-scheme: dark)');
+      root.dataset.theme = dark.matches ? 'dark' : 'light';
+      if (parent === window) {
+        dark.addEventListener('change', () => { root.dataset.theme = 
dark.matches ? 'dark' : 'light'; });
+        return;
+      }
+      const parentOrigin = document.referrer ? new 
URL(document.referrer).origin : '*';
+      const report = () => parent.postMessage({
+        type: 'tirx-matrix-dataflow:height', height: 
Math.ceil(document.body.getBoundingClientRect().height)
+      }, parentOrigin);
+      new ResizeObserver(report).observe(document.body);
+      addEventListener('message', (event) => {
+        if (event.source !== parent || (parentOrigin !== '*' && event.origin 
!== parentOrigin)) return;
+        if (event.data?.type !== 'tirx-matrix-dataflow:theme') return;
+        if (['light', 'dark'].includes(event.data.theme)) root.dataset.theme = 
event.data.theme;
+        report();
+      });
+      addEventListener('load', report);
+      report();
+    })();
+  </script>
+</body>
+</html>
diff --git a/images/docs/tirx/warp_matrix.html 
b/images/docs/tirx/warp_matrix.html
new file mode 100644
index 00000000000..7cb84465278
--- /dev/null
+++ b/images/docs/tirx/warp_matrix.html
@@ -0,0 +1,350 @@
+<!DOCTYPE html>
+<!--
+  Licensed to the Apache Software Foundation (ASF) under one
+  or more contributor license agreements. See the NOTICE file
+  distributed with this work for additional information
+  regarding copyright ownership. The ASF licenses this file
+  to you under the Apache License, Version 2.0 (the
+  "License"); you may not use this file except in compliance
+  with the License. You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+  Unless required by applicable law or agreed to in writing,
+  software distributed under the License is distributed on an
+  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+  KIND, either express or implied. See the License for the
+  specific language governing permissions and limitations
+  under the License.
+-->
+<html lang="en">
+<head>
+  <meta charset="utf-8">
+  <meta name="viewport" content="width=device-width, initial-scale=1">
+  <title>Warp matrix multiplication data flow</title>
+  <style>
+    /* Match the typography, blue cards, and neutral frame in 
authoring_layers.html. */
+    :root {
+      color-scheme: light;
+      --surface: #fff;
+      --text: #222;
+      --muted: #536276;
+      --frame: #b6bdc6;
+      --frame-bg: #f8fafc;
+      --tag: #475569;
+      --line: #64748b;
+      --layer-border: #5b9bd5;
+      --layer-bg: #eef5fc;
+      --a-bg: #e2effc;
+      --a-border: #7a9fc7;
+      --b-bg: #eee7f8;
+      --b-border: #a38bc2;
+    }
+    :root[data-theme="dark"] {
+      color-scheme: dark;
+      --surface: #14181e;
+      --text: #e5e9ef;
+      --muted: #b1bdcd;
+      --frame: #627084;
+      --frame-bg: #1c232c;
+      --line: #a1afc1;
+      --layer-bg: #1d3043;
+      --a-bg: #243c54;
+      --a-border: #7cadd7;
+      --b-bg: #3c304e;
+      --b-border: #b09bd0;
+    }
+    * { box-sizing: border-box; }
+    html, body { margin: 0; }
+    body {
+      padding: 14px 4px 4px;
+      background: var(--surface);
+      color: var(--text);
+      font: 13px/1.5 'Inter', 'SF Pro', 'Segoe UI', system-ui, sans-serif;
+    }
+    h1, h2, p { margin: 0; }
+    h2 { font-size: 14px; font-weight: 700; }
+    code { font: 12px/1.5 'SF Mono', 'Fira Code', Consolas, monospace; }
+    .diagram {
+      position: relative;
+      max-width: 920px;
+      margin: 0 auto;
+      padding: 26px 18px 16px;
+      border: 2px solid var(--frame);
+      border-radius: 12px;
+      background: var(--frame-bg);
+    }
+    .compact-diagram { max-width: 680px; }
+    .scope {
+      position: absolute;
+      top: -11px;
+      left: 16px;
+      padding: 1px 10px;
+      border-radius: 4px;
+      background: var(--tag);
+      color: #fff;
+      font-size: 11px;
+      font-weight: 700;
+      letter-spacing: .02em;
+    }
+    .card {
+      padding: 10px 14px;
+      border: 2px solid var(--layer-border);
+      border-radius: 8px;
+      background: var(--layer-bg);
+      text-align: center;
+    }
+    .card h2 + p, .card h2 + .pair { margin-top: 4px; }
+    .endpoint { border-color: var(--frame); background: var(--surface); }
+    .pair { display: grid; grid-template-columns: 1fr 1fr; gap: 0; }
+    .pair > :first-child { padding-right: 12px; }
+    .pair > :last-child { padding-left: 12px; border-left: 1px solid 
var(--frame); }
+    .card:not(.endpoint) .pair:not(.operands) > :last-child { border-color: 
var(--layer-border); }
+    .card h2 + .pair.operands { margin-top: 8px; }
+    .pair.operands > p {
+      margin: 0 6px;
+      padding: 6px 8px;
+      border: 1px solid var(--operand-border);
+      border-radius: 6px;
+      background: var(--operand-bg);
+    }
+    .operand-a { --operand-bg: var(--a-bg); --operand-border: var(--a-border); 
}
+    .operand-b { --operand-bg: var(--b-bg); --operand-border: var(--b-border); 
}
+    .operand-detail { display: block; }
+    .note { color: var(--muted); font-size: 11px; }
+    .note code { font-size: inherit; }
+    .packing, .coverage { margin-top: 6px; }
+    .equation { margin-bottom: 14px; text-align: center; }
+    .equation code { font-size: 15px; font-weight: 600; }
+    .parallel {
+      display: grid;
+      grid-template-columns: minmax(0, 1fr) minmax(0, 1fr);
+      padding: 0 16px;
+    }
+    .sync-note {
+      grid-column: 1 / -1;
+      position: relative;
+      padding: 10px 0 0;
+      text-align: center;
+    }
+    .sync-note::before, .sync-note::after {
+      content: '';
+      position: absolute;
+      top: 0;
+      bottom: 0;
+      border-left: 1.5px solid var(--line);
+    }
+    .sync-note::before { left: 25%; }
+    .sync-note::after { left: 75%; }
+    .sync-note .operation { display: inline-block; z-index: 1; }
+    .merge {
+      height: 20px;
+      margin: 0 calc(25% + 8px);
+      border: solid var(--line);
+      border-width: 0 1.5px 1.5px;
+    }
+    .connector {
+      position: relative;
+      display: flex;
+      flex-direction: column;
+      align-items: center;
+      justify-content: center;
+      min-height: 50px;
+      padding: 8px 0 13px;
+      text-align: center;
+    }
+    .connector::before {
+      content: '';
+      position: absolute;
+      left: 50%;
+      top: 0;
+      bottom: 3px;
+      border-left: 1.5px solid var(--line);
+    }
+    .connector::after {
+      content: '';
+      position: absolute;
+      left: calc(50% - 3.25px);
+      bottom: 0;
+      border-left: 4px solid transparent;
+      border-right: 4px solid transparent;
+      border-top: 6px solid var(--line);
+    }
+    .operation {
+      position: relative;
+      max-width: 100%;
+      padding: 1px 10px;
+      background: var(--frame-bg);
+    }
+    .operation > code, .operation > strong { font-weight: 600; }
+    .handoff { gap: 12px; }
+    .lifetime {
+      display: grid;
+      grid-template-columns: 1fr 1fr;
+      gap: 8px 12px;
+      margin-bottom: 8px;
+      text-align: center;
+    }
+    .lifetime > section {
+      padding: 5px 8px;
+      border: 1px solid var(--frame);
+      border-radius: 6px;
+      background: var(--surface);
+    }
+    .lifetime h2 { font-size: 12px; }
+    .blackwell-diagram .handoff { gap: 10px; }
+    .release {
+      margin-top: 14px;
+      padding-top: 10px;
+      border-top: 1px solid var(--frame);
+      text-align: center;
+    }
+    .sr-only {
+      position: absolute;
+      width: 1px;
+      height: 1px;
+      padding: 0;
+      overflow: hidden;
+      clip-path: inset(50%);
+      white-space: nowrap;
+    }
+    /* Use the same compact vertical layout for both matrix examples. */
+    @media (min-width: 640px) {
+      .compact-diagram { padding: 20px 10px 12px; }
+      .compact-diagram .equation { margin-bottom: 8px; }
+      .compact-diagram .card { padding-block: 6px; }
+      .compact-diagram .card h2 + .pair.operands { margin-top: 4px; }
+      .compact-diagram .pair.operands > p { padding-block: 4px; }
+      .compact-diagram .operand-detail { display: inline-block; }
+      .compact-diagram .operand-detail::before { content: ' · '; }
+      .compact-diagram .connector { min-height: 34px; padding: 3px 0 10px; }
+      .compact-diagram .sync-note { padding-top: 6px; }
+      .compact-diagram .compact-operation {
+        display: inline-flex;
+        flex-wrap: wrap;
+        align-items: baseline;
+        justify-content: center;
+        gap: 0 8px;
+      }
+      .compact-diagram .merge { height: 12px; }
+      .compact-diagram .packing, .compact-diagram .coverage { margin-top: 4px; 
}
+      .compact-diagram .single-output h2,
+      .compact-diagram .single-output > p:not(.note) { display: inline; }
+      .compact-diagram .single-output > p:not(.note) { margin-left: 8px; }
+      .blackwell-diagram .release { margin-top: 8px; padding-top: 6px; }
+      .blackwell-diagram .release h2 { font-size: 12px; }
+    }
+    @media (max-width: 480px) {
+      .diagram { padding: 24px 10px 12px; }
+      .card { padding: 9px 8px; }
+      .pair > :first-child { padding-right: 8px; }
+      .pair > :last-child { padding-left: 8px; }
+      .operation { padding: 2px 6px; }
+      .lifetime { gap: 8px 12px; }
+      .parallel { padding: 0 10px; }
+      .merge { margin-inline: calc(25% + 5px); }
+    }
+  </style>
+</head>
+<body>
+  <main class="diagram compact-diagram warp-diagram" aria-labelledby="scope" 
aria-describedby="description">
+    <h1 class="scope" id="scope">One warp · 32 lanes</h1>
+    <p class="sr-only" id="description">
+      The warp computes Out = A @ B.T. All 32 lanes copy float16 inputs A and B
+      from global to shared memory, then synchronize at a CTA barrier. Separate
+      ldmatrix instructions load A into ar and B into br. These register 
fragments
+      feed one warp-wide mma.sync operation. Each lane stores four float32 
results,
+      together forming the complete 16 by 8 output tile in global memory.
+    </p>
+    <p class="equation"><code>Out = A @ B.T</code></p>
+    <section class="card endpoint">
+      <h2>Inputs · global memory</h2>
+      <div class="pair operands">
+        <p class="operand-a"><code>A</code> · 16 × 16 · float16</p>
+        <p class="operand-b"><code>B</code> · 8 × 16 · float16</p>
+      </div>
+    </section>
+    <div class="parallel" aria-label="Cooperative copy of A and B to shared 
memory">
+      <div class="connector">
+        <div class="operation">Copy A</div>
+      </div>
+      <div class="connector">
+        <div class="operation">Copy B</div>
+      </div>
+    </div>
+    <section class="card">
+      <h2>Shared memory</h2>
+      <div class="pair operands">
+        <p class="operand-a"><strong>Shared copy of A 
(<code>a</code>)</strong><span class="operand-detail">16 × 16 · 
float16</span></p>
+        <p class="operand-b"><strong>Shared copy of B 
(<code>b</code>)</strong><span class="operand-detail">8 × 16 · 
float16</span></p>
+      </div>
+    </section>
+    <div class="parallel">
+      <div class="sync-note">
+        <div class="operation compact-operation">
+          <strong>CTA barrier</strong>
+          <p class="note">All 32 lanes synchronize before loading</p>
+        </div>
+      </div>
+      <div class="connector">
+        <div class="operation"><code>ldmatrix.x4</code></div>
+      </div>
+      <div class="connector">
+        <div class="operation"><code>ldmatrix.x2</code></div>
+      </div>
+    </div>
+    <section class="card">
+      <h2>Input register fragments</h2>
+      <div class="pair operands">
+        <p class="operand-a"><strong>A fragment 
(<code>ar</code>)</strong><span class="operand-detail">4 uint32 words per 
lane</span></p>
+        <p class="operand-b"><strong>B fragment 
(<code>br</code>)</strong><span class="operand-detail">2 uint32 words per 
lane</span></p>
+      </div>
+      <p class="note packing">Each uint32 word packs two float16 elements.</p>
+    </section>
+    <div class="merge" aria-hidden="true"></div>
+    <div class="connector">
+      <div class="operation compact-operation">
+        <code>mma.sync · m16n8k16</code>
+        <p class="note">All 32 lanes participate · initial accumulator = 0</p>
+      </div>
+    </div>
+    <section class="card single-output">
+      <h2>Output register fragments</h2>
+      <p><code>dr</code> · 4 float32 values per lane</p>
+      <p class="note coverage">32 lanes × 4 values = one 16 × 8 output tile</p>
+    </section>
+    <div class="connector">
+      <div class="operation">Each lane stores its four results</div>
+    </div>
+    <section class="card endpoint single-output">
+      <h2>Out · global memory</h2>
+      <p>16 × 8 · float32</p>
+    </section>
+  </main>
+  <script>
+    // Report the diagram's natural height and receive the documentation theme.
+    (() => {
+      const root = document.documentElement;
+      const dark = matchMedia('(prefers-color-scheme: dark)');
+      root.dataset.theme = dark.matches ? 'dark' : 'light';
+      if (parent === window) {
+        dark.addEventListener('change', () => { root.dataset.theme = 
dark.matches ? 'dark' : 'light'; });
+        return;
+      }
+      const parentOrigin = document.referrer ? new 
URL(document.referrer).origin : '*';
+      const report = () => parent.postMessage({
+        type: 'tirx-matrix-dataflow:height', height: 
Math.ceil(document.body.getBoundingClientRect().height)
+      }, parentOrigin);
+      new ResizeObserver(report).observe(document.body);
+      addEventListener('message', (event) => {
+        if (event.source !== parent || (parentOrigin !== '*' && event.origin 
!== parentOrigin)) return;
+        if (event.data?.type !== 'tirx-matrix-dataflow:theme') return;
+        if (['light', 'dark'].includes(event.data.theme)) root.dataset.theme = 
event.data.theme;
+        report();
+      });
+      addEventListener('load', report);
+      report();
+    })();
+  </script>
+</body>
+</html>

Reply via email to