Translate TMA chapter into Chinese

[tlopex]

• Translate the TMA chapter into Chinese and enable it in the Chinese documentation build.
• Explain TMA tile copies, tensor-map descriptors, shared-memory swizzling, 3D TMA tiling, load/store completion, and pipelining in a tutorial-oriented flow.
• Localize and refine the three interactive TMA demos, including explicit 16-byte sector units and swizzle-atom boundaries.
• Add a Chinese diagram illustrating TMA load completion through mbarrier transaction-byte tracking.

[gemini-code-assist]

Code Review

This pull request translates and expands the documentation for the Tensor Memory Accelerator (TMA) chapter, including updating interactive HTML demos and adding a synchronization flow diagram. The reviewer provided critical feedback regarding the asynchronous load synchronization flow: in both the SVG diagram and the markdown text, registering the expected transaction bytes (expect_tx) must be performed before initiating the asynchronous copy (copy_async) to prevent potential race conditions where the hardware transfer completes before the bytes are registered.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

[gemini-code-assist]

在 TMA 异步 Load 的同步流程中，登记 expected transaction bytes（即 expect_tx）必须在发起异步拷贝（即 copy_async）之前完成。如果在发起 copy_async 之后才执行 expect_tx，可能会因为硬件传输极快而导致 complete-tx 信号先于 expect_tx 到达，从而引发 barrier 状态错误或未定义行为。

建议将步骤 1（发起 TMA load）与步骤 2（登记 byte count）的顺序进行调换，使流程在学术和实践上都保持严谨。

  <g>
    <circle cx="64" cy="245" r="24" fill="#4a90d9" stroke="#111827" stroke-width="2"/>
    <text x="64" y="252" text-anchor="middle" fill="#ffffff" class="step">1</text>
    <text x="102" y="251" class="step">登记 byte count</text>
    <line x1="265" y1="245" x2="965" y2="245" stroke="#d97706" stroke-width="4" marker-end="url(#arrow-orange)"/>
    <text x="615" y="230" text-anchor="middle" class="label mono" fill="#d97706">arrive.expect_tx(4096)</text>
    <text x="965" y="278" text-anchor="middle" class="detail">arrival = 0, pending bytes = 4096</text>
  </g>

  <g>
    <circle cx="64" cy="365" r="24" fill="#4a90d9" stroke="#111827" stroke-width="2"/>
    <text x="64" y="372" text-anchor="middle" fill="#ffffff" class="step">2</text>
    <text x="102" y="371" class="step">发起 TMA load</text>
    <line x1="265" y1="345" x2="615" y2="345" stroke="#2563eb" stroke-width="4" marker-end="url(#arrow-blue)"/>
    <text x="440" y="333" text-anchor="middle" class="label mono" fill="#2563eb">copy_async(A)</text>
    <line x1="265" y1="395" x2="615" y2="395" stroke="#2563eb" stroke-width="4" marker-end="url(#arrow-blue)"/>
    <text x="440" y="383" text-anchor="middle" class="label mono" fill="#2563eb">copy_async(B)</text>
  </g>

[gemini-code-assist]

为了与修正后的同步流程图保持一致，并避免读者在编写代码时引入潜在的竞争冒险（Race Condition），建议在文本中明确指出应当先执行 expect_tx 登记字节数，再发起 TMA load 拷贝操作。

建议修改为：

假设初始化 mbarrier 时将 expected arrival count 设为 1。发起 copy 的 thread 先执行 mbarrier.arrive.expect_tx(4096) 登记字节数，然后再发起 A、B 两次 TMA load 并将它们关联 to 这个 barrier。这一步既完成该 thread 的一次 arrival，也把 pending bytes 设为 4096：