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>