mirror of
https://github.com/deepseek-ai/FlashMLA
synced 2025-06-26 18:15:54 +00:00
Add the deep-dive blog
This commit is contained in:
Shengyu Liu
parent
c7996e951d
commit
984059a7cf
77
docs/20250422-new-kernel-deep-dive.md
Normal file
77
docs/20250422-new-kernel-deep-dive.md
Normal file
@ -0,0 +1,77 @@
|
||||
# A Deep-Dive Into the New Flash MLA Kernel
|
||||
|
||||
In the [previous version](https://github.com/deepseek-ai/FlashMLA/tree/b31bfe72a83ea205467b3271a5845440a03ed7cb) of the Flash MLA kernel, we have achieved impressive performance: 3000 GB/s in memory-intensive settings and 580 TFlops in compute-bound settings. Now, we're pushing these numbers even further, reaching up to 660 TFlops.
|
||||
|
||||
In this blog, we present a deep dive into the new kernel, explaining the optimizations and techniques behind this performance boost. We'll first explain why the MLA kernel is compute-bound despite being a decoding-stage attention kernel, then discuss our high-level kernel schedule design, and finally cover the technical details of the new kernel.
|
||||
|
||||
## A Theoretical Analysis of the MLA Algorithm
|
||||
|
||||
GPU kernels can be classified as either compute-bound (limited by floating-point operations per second, FLOPs) or memory-bound (limited by memory bandwidth). To identify the kernel's bottleneck, we calculate the ratio of FLOPs to memory bandwidth (FLOPs/byte) and compare it with the GPU's capacity.
|
||||
|
||||
Assume the number of q heads is $h_q$, the number of q tokens per request is $s_q$ (should be 1 if MTP / speculative decoding is disabled), the number of kv tokens per request is $s_k\ (s_k \gg h_q s_q)$, and the head dimensions of K and V are $d_k$ and $d_v$ respectively. The number of FLOPs is roughly $2 (h_q s_q \cdot d_k \cdot s_k + h_q s_q \cdot s_k \cdot d_v) = 2 h_q s_q s_k (d_k+d_v)$, and the memory access volume (in bytes) is $\mathop{\text{sizeof}}(\text{bfloat16}) \times (h_q s_q d_k + s_k d_k + h_q s_q d_v) \approx 2s_k d_k$. Thus, the compute-memory ratio is $h_q s_q \cdot \frac{d_k+d_v}{d_k} \approx 2 h_q s_q$.
|
||||
|
||||
An NVIDIA H800 SXM5 GPU has a peak memory bandwidth of 3.35 TB/s and peak FLOPs of 990 TFlops. However, due to throttling (reducing to ~1600 MHz in our case), the practical peak FLOPs drops to ~865 TFlops. Therefore, when $h_qs_q \ge \frac{1}{2} \cdot \frac{865}{3.35} = 128$, the kernel is compute-bound; otherwise, it's memory-bound.
|
||||
|
||||
According to [the overview of DeepSeek's Online Inference System](https://github.com/deepseek-ai/open-infra-index/blob/main/202502OpenSourceWeek/day_6_one_more_thing_deepseekV3R1_inference_system_overview.md), we don't use Tensor Parallel for decoding instances, meaning $h_q$ is 128 and the kernel is compute-bound. Thus, we need to optimize the kernel for compute-bound settings.
|
||||
|
||||
## High-Level Design of the New Kernel
|
||||
|
||||
To fully utilize GPU compute resources, we need to overlap CUDA Core operations with Tensor Core operations and memory access with computation, keeping the Tensor Core constantly busy. This requires redesigning the kernel's "schedule."
|
||||
|
||||
[FlashAttention-3's paper](https://arxiv.org/abs/2205.14135) introduces ping-pong scheduling and intra-warpgroup GEMM-softmax pipelining to overlap block-wise matmul and CUDA Core operations. However, these techniques can't be directly applied here due to resource constraints. The output matrix (scaled and accumulated during each mainloop round, similar to [FlashAttention's algorithm](https://arxiv.org/abs/2205.14135)) must be stored in registers due to [WGMMA instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions) requirements. Each $64 \times 512$ output matrix occupies 32,768 32-bit registers. With only 65,536 32-bit registers per SM, we can store only one output matrix per SM. This eliminates the possiblility of having two output matrices and letting them use CUDA Core and Tensor Core in a interleaved manner. We need to find another clever way to overlap CUDA Core and Tensor Core computation.
|
||||
|
||||
(You might pause here to ponder - perhaps you can find a better solution than ours!)
|
||||
|
||||
Our solution involves an additional mathematical transformation beyond FlashAttention's online softmax and accumulation approach. In each step, we take two KV blocks (called $K_0$, $K_1$, $V_0$, and $V_1$). Since the output matrix occupies 32,768 registers (too many for one warpgroup), we split it vertically into $O_L$ and $O_R$ (each $64 \times 256$). We similarly split $V_0$ and $V_1$ into $V_{0L}$, $V_{0R}$, $V_{1L}$, and $V_{1R}$ (each $64 \times 256$). The output matrix is then computed as follows:
|
||||
|
||||
0. Maintain a running max $m$ (initialized to $-\infty$, shared between the two warpgroups) and output matrices $\vec o_L, \vec o_R$ (initialized to 0).
|
||||
1. [0] Compute $\vec p_0 = \vec q K_0^\intercal / qk\_scale$.
|
||||
2. [1] Compute $\vec p_1 = \vec q K_1^\intercal / qk\_scale$.
|
||||
3. [0] Compute $mp_0 = \max(\vec p_0)$, $m\_new_0 = \max(m, mp_0)$, and $scale_0 = \exp(m\_new_0 - m)$. Update $m \gets m\_new_0$.
|
||||
4. [0] Perform softmax on $\vec p_0$: $\vec p_0 \gets \exp(\vec p_0 - m\_new_0)$.
|
||||
5. [0] Update $\vec o_L \gets \vec o_L \cdot scale_0 + \vec p_0 V_{0L}$.
|
||||
6. [1] Compute $mp_1 = \max(\vec p_1)$, $m\_new_1 = \max(m, mp_1)$, and $scale_1 = \exp(m\_new_1 - m)$. Update $m \gets m\_new_1$.
|
||||
7. [1] Perform softmax on $\vec p_1$: $\vec p_1 \gets \exp(\vec p_1 - m\_new_1)$.
|
||||
8. [1] Update $\vec o_R \gets \vec o_R \cdot (scale_0 \cdot scale_1) + \vec p_1 V_{1R}$.
|
||||
9. [0] Update $\vec p_0 \gets \vec p_0 \cdot scale_1$.
|
||||
10. [1] Update $\vec o_R \gets \vec o_R + \vec p_0 V_{0R}$.
|
||||
11. [0] Update $\vec o_L \gets \vec o_L \cdot scale_1 + \vec p_1 V_{1L}$.
|
||||
|
||||
Note: We assume one q head for simplicity, so $\vec q$ and $\vec o$ are vectors. Bracketed numbers indicate the warpgroup performing the operation. Assume $\vec o_L$ resides in warpgroup 0's register and $\vec o_R$ resides in warpgroup 1's register.
|
||||
|
||||
This schedule can be viewed as a "ping-pong" variant using one output matrix—we call it "seesaw" scheduling. It's mathematically equivalent to FlashAttention's online softmax algorithm. This schedule allows us to overlap CUDA Core and Tensor Core operations by interleaving the two warpgroups, and also allows us to overlap memory access with computation since we can launch the corresponding Tensor Memory Accelerator (TMA) instructions right after data is no longer needed.
|
||||
|
||||
The complete schedule is shown below (remember that in MLA, $K$ and $V$ are the same with different names):
|
||||
|
||||
|
||||
|
||||
## Discussion of Technical Details
|
||||
|
||||
This section covers technical details of the new kernel.
|
||||
|
||||
First, although the kernel targets compute-bound scenarios (where memory bandwidth isn't the bottleneck), we can't ignore memory latency. If the data is not ready when we want to use it, we have to wait. To solve this problem, we employ the following techniques:
|
||||
|
||||
- **Fine-grained TMA copy - GEMM pipelining:** For a $64 \times 576$ K block, we launch 9 TMA copies (each moving a $64 \times 64$ block). GEMM operations begin as soon as each TMA copy completes (When the first TMA copy is done, we can start the first GEMM operation, and so on), improving memory latency tolerance.
|
||||
- **Cache hints:** Using `cute::TMA::CacheHintSm90::EVICT_FIRST` for TMA copies improves L2 cache hit rates, as shown by experiments.
|
||||
|
||||
These optimizations achieve up to 80% Tensor Core utilization (of the throttled theoretical peak) and 3 TB/s memory bandwidth on an H800 SXM5 GPU. While slightly slower (~2%) than the old ping-pong buffer version in memory-bound settings, this is acceptable.
|
||||
|
||||
Other performance improvements include:
|
||||
- **Programmatic Dependent Launch.** We use [programmatic dependent launch](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programmatic-dependent-launch-and-synchronization) to overlap `splitkv_mla` and `combine` kernels.
|
||||
- **Tile Scheduler.** We implement a tile scheduler to allocate jobs (requests and blocks) to SMs. This ensures a balanced load across SMs.
|
||||
|
||||
## Acknowledgements
|
||||
|
||||
FlashMLA's algorithm and scheduling is inspired by [FlashAttention](https://github.com/dao-AILab/flash-attention/), [Flash-Decoding](https://crfm.stanford.edu/2023/10/12/flashdecoding.html), and [CUTLASS](https://github.com/nvidia/cutlass), as well as many projects behind them. We thank the authors for their great work.
|
||||
|
||||
## Citation
|
||||
|
||||
```bibtex
|
||||
@misc{flashmla2025,
|
||||
title={FlashMLA: Efficient MLA decoding kernels},
|
||||
author={Jiashi Li, Shengyu Liu},
|
||||
year={2025},
|
||||
publisher = {GitHub},
|
||||
howpublished = {\url{https://github.com/deepseek-ai/FlashMLA}},
|
||||
}
|
||||
```
|
||||
852
docs/assets/MLA Kernel Sched.drawio.svg
Normal file
852
docs/assets/MLA Kernel Sched.drawio.svg
Normal file
@ -0,0 +1,852 @@
|
||||
<svg host="65bd71144e" xmlns="http://www.w3.org/2000/svg" style="background: transparent; background-color: transparent;" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="511px" height="808px" viewBox="-0.5 -0.5 511 808" content="<mxfile scale="1" border="0"><diagram name="Page-1" id="t0OnRo4-AS1naIkeEJg6">7V1bk6K6Fv41VnU/2EW4qY99OTNz6vTU7rN7395OoUTlDAI7YNszv34nkCgkKIKJQ2umq0ZIIISs9a1bVsLAely9f0Zesvwa+zAcmIb/PrCeBqYJRsDBP6TkOy1xHbMoWaDAp2W7gtfgB6SFBi1dBz5MKxdmcRxmQVItnMVRBGdZpcxDKN5UL5vHYfWpibegTzR2Ba8zL4TCZX8GfrYsSsdO6eovMFgs2ZOBQWtWHruYFqRLz483pSLrXwPrEcVxVhyt3h9hSEaPjUtx36c9tduOIRhlx9xAx/3NC9f03Wi/su/sZWHk35Mxw2dTL4VPXopf4WGZrUJcAvAhiteRD0mDBj77/3qVvNK7ozjCPw/zMN7cRwF++SCO6GXbF89P4jWawReIghXMIHpNvFkQLVhl5qGM8oDNzmt6lJd/CsKQ3of7XT0rtSEOEx056C/KNKaD9hnGuGPoO75gs6M3I/eyRGpWhmCI3/atyi8eZbvFtrntE17iAPfENChGhq5NG6IQGZquUW2EjllxX5nCjU3ZNtcUHrgFzISm8EHp1XdFOQvVs5MlsBN6wQ8z8mMj/S85JE830v8YIqPhsX/2plhYVNjLC4MF4ZoZphVEuOANoizAYLynFavA90kbDwimwQ9vmrdHaJ6QN8rf0XkYOE+1dDcrkoLevMMneRZ8P8gRtNa4A7Y9OY367JJ4Pk/hqbSwpUObR2IVU1W8HYvS3iPRNow7d2Ls/jlVNI274rKpYQHx8mDq1MAU1MMU9AOm9sXC1JUP04pqzVD8DT7GYYxp8uTDubcOs1bkmsdR9slbBSFhzC8wfIPkOlpB8Q8APRcfFBJWefBm3xZ5L8ULagyGw4Km1pxotCBaWwy9l0yOyyl20FkWCU0pFD6jGoZ3c07wg7cK47t/r+OchYIIDtkI3uNLJrhzTk4idgk+WtDfXVPVIsKetc2TimGa81feePKe32W0eOxnmDeduwgGeyAeiOKZ1X7g4prenbnDvye+l0HS+tdO3VXZN/Qy3eqisvn4Gs+zlfd+g8tuGzp97h7/Uuol4YEbXNLUx37xAXru1F2+9Py2gmr1JKoBV6YtYox6ZIuM1bgMrVXvYf3/ARVzZ+e9WcfLU8wTgfr/TlN82AsHYCwLdEPjzrB7hDkW09QhuJ6h2OXDZN0jbuowC4DAPdQYMR+YBTXdufN/GM/9gDPj+gt06IH8mLr26K/Do7fHvBveObjIt6RQBIlB/5/s0B9o3tjjMx1y6IFyR05Ch5U59BL6hl7AzqPfefHbcU0TL6p99nQrZIazQsrc5+4t7svQ99C3m4FpzfN/A/MxZ0F0MxwuYF5Jb8kr8HXAJH+3t0XP65tDi+mNUdzBfm6LX1JjOk5xUj6gDfKvXHb+i9cTiHBq3EIGYX7ZRf1Z3ALcXjBZulHlaIBf8IidRXYdE4TaQzcdm6oYnxcYnALtJrRnoZemwayjWY1xNnFHludemWE9x9ZyaQwg8B04wuWzNXrbjpuPje38BHwEo9pxOFO4+zSZPTmfVa1gXvgaY7G205nYPN/wzphEYotzoj0KvQJpEx69i73Kn/DQsVcpqDWq2DslyUoZaMX5EupP7UKvoBx7Bb/2BM/S5lJ6F3vdvoKOvSozES809srPtE66ihy+oZEyCWTunfypRHSI638Nvi8Dvwy5ZkwsUKGiKYXLzOoNcmReu/km7Rdrv7iQU0aDf3O8xON9LIkiTv4qBm2hy+AePsfBlJYtISRLSeSmmoUP1WyJtFgHwSz2nmRLXPDyBwXrHzTAJbngd3ZpdVLVHOqe0djk2ktEuxhV4xz0tMguZ7lRPfHP5SYY9wvt7QJuJ9uq07FDwrhXbavOxzM4mx20VcnJi5fhlqK8xDQ+hgXLJ0fb3RdJW3xTljq5JAYONQg0CLrq6RHHuRborJoB7xGqc+wYvDQINAgkaAIeBICPunbXBHwgWCIGxPitxoDGQFdFwK86OV0PKGD5dsFqPUGnJ+hK1v4dmOz+jSvc3jn8eLhZdTk/lrhsIi0vdi8m716uZfKOCYYrm7yzxPmN32CUknRvjur47TN1BPfWWZxSugpilAVRS5qUFZ0/wbktOwxdPkrgMJ4qCS63RnDxEYAyh1RIfoi+4ozD4+9P95q6sqjrjJxm6tapJSnUFScvNHVlUldIuD4rdcX5Ay2bJdPX/JmyWZyM+NNDCXnfZGCyaSpNZgkqmI+pOLZI5okqMovh9jKZxQ0ONZm7oplfdH5OMrNHHRlZoMOqowrHRhX2Mkhf4gOCnW93TuG1GpuSFwewzxwE1umbF5K+yS1P6xwA421sdbuD23U7hZwY/L3KNY2cnu2eizU+G+nFOFePVjgyxry8FY62GH9Sql30FONFTzFy0+zd11ZzmSbqkkDtdvnF2rzS5lWtVdR9B1fOTlO3HNA+c3Kt705d58pZfT6fmw2yvves7vLhl+N5WyEzt0uS1b7CkTnSoHNWUNMCaYmfpxFDeX1yFihnXp6z4LQLRWnMHanyT9iECDd1xzY+27amzlpm32EsccBLkECyx5+Pi3/7Srbx23hBznoRKQr6A0zGvhcIzHaBM+3Fay/+kHjivHjA75HUOVde4dJtR/5GAFepm4Q5Wqsz9cWm1K2Zc8Q45gdSTNKW7/dNMbmipa5ReQoqTwehuox1V3QP0tJnQ4vd7vqAN8aVvcAba4n3u+XgT7TXBRJgUy8hh9iOCb8/IGxjQWJQbZZBBokFROo2yEv2gRIw8+wlTgNqTIVwnpWo9sxVb6nH6IwKyOwn816Etsky45MJt/AsoRgcQPEp6UeuaKJvFsZwuo5my+EHSSWUQYSRsKvUGXPA3COMRBlwmIdB8mUPNHhWb8QGhyUl0LD5BMzxGZEhmm6bBbg+ZAg0GJ0RGAqi6TpB8lqWXVbYtvsugQo/Kb13woB6hjfRO+nKH8bzlSyndKVNUnyo5ZQjBVEiLeeuU85J+2avwrTKkWhb1Qs+cC2CbyQt1vaxBN+Z09x07s+F5f6cvI0vP6ukUuqJC3M1s2tmb46yN2QdHx90Hx1uSCKry/9C0XVPsJw85823pJD2YtBmj3H3608x7mqML2lfFurbROdYjDBoHMrAobSJT4W5J2Nx4nNPeKkvQGRjc4FANFsBUdt+2vartf1O+IABnz6nUPCcefWmZvYWzM5xdg3398arH54QzOR3VhaaksjuOoil2V2CbBc+y91Ztiv8wPdYDGKlf/01sIg9+ZrFKM9rxSTF3L/0irMVXMVkINycVfzgDR8uyCGqvxHBRZAShjYNTHXIbsQ9K90rIGx/AkZTpsqJ6Rch4Yyh76FvN5gxZwb5G5h4IPH/1mhO/m45BqYQOZCVjz1SY7x9TcHuPYqNt9sC8bsRbr83Uk6jqcvh4GOhXXI4xu1CQeIGVxzyfQeOfbtOLo7NqeW6HwH5YrYZSxRvr+gam5KI/XbpOFdBSmH8R10paTa1JI+QDOuakG2A1B2TCilZt/a10JfElKuQ1P17HbOKYaHOiPKdJO85fVg107DPcZzg6imx+LycEgHRzngQ8giNX1bMxbOYZhYUvawu3SThj5zU89w6yOK8J6sVJPqa9m/gGpsF+F+6nmJbNQsw57pkM/qDXe1sWvzknTWZ5b7LVJFsKWyX3ZdAuJVTJ1oK+BTFhLw71kdesvyKOYtc8Q8=</diagram></mxfile>">
|
||||
<defs/>
|
||||
<g>
|
||||
<g>
|
||||
<path d="M 60 280 L 60 100" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 57 280 L 63 280" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 63 100 L 57 100" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 178px; margin-left: 60px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rP0 = sQ @ sK0
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="60" y="181" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rP0 = sQ @ sK0
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 299.31 460 L 299.31 280" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 296.31 460 L 302.31 460" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 302.31 280 L 296.31 280" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 358px; margin-left: 300px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rP1 = sQ @ sK1
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="300" y="361" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rP1 = sQ @ sK1
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 360 L 140 280" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 360 L 143 360" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 280 L 137 280" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 315px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
<div style="line-height: 90%;">
|
||||
<div>
|
||||
<font style="font-size: 9px; line-height: 90%;">
|
||||
Get sScale0
|
||||
</font>
|
||||
</div>
|
||||
<div>
|
||||
<font style="font-size: 9px; line-height: 90%;">
|
||||
Update sM
|
||||
</font>
|
||||
</div>
|
||||
<font style="font-size: 9px; line-height: 90%;">
|
||||
rPb = rP0 = Softmax(rP0)
|
||||
</font>
|
||||
<div>
|
||||
<font style="font-size: 9px; line-height: 90%;">
|
||||
rO0 = Scale(rO0)
|
||||
</font>
|
||||
</div>
|
||||
<div>
|
||||
<font style="font-size: 9px; line-height: 90%;">
|
||||
Update rL
|
||||
</font>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="318" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Get sScale0...
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 380 L 140 360" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 380 L 143 380" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 360 L 137 360" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 371px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="374" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 60 540 L 60 460" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 57 540 L 63 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 63 460 L 57 460" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 495px; margin-left: 60px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rO0 += rPb @ sV0L
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="60" y="498" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rO0 += rPb @ sV0L
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 520 L 220 460" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 520 L 223 520" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 460 L 217 460" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 486px; margin-left: 220px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
<div style="line-height: 90%;">
|
||||
<div>
|
||||
<font style="line-height: 90%; font-size: 9px;">
|
||||
Get sScale1
|
||||
</font>
|
||||
</div>
|
||||
<div>
|
||||
<font style="line-height: 90%; font-size: 9px;">
|
||||
Update sM
|
||||
</font>
|
||||
</div>
|
||||
<font style="line-height: 90%; font-size: 9px;">
|
||||
rP1b = Softmax(rP1
|
||||
<span style="background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); color: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));">
|
||||
)
|
||||
</span>
|
||||
</font>
|
||||
<div>
|
||||
<font style="line-height: 90%; font-size: 9px;">
|
||||
rO1 = Scale(rO1)
|
||||
<span style="background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); color: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"></span>
|
||||
</font>
|
||||
</div>
|
||||
<div>
|
||||
<span style="background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); color: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));">
|
||||
<font style="line-height: 90%; font-size: 9px;">
|
||||
Update rL
|
||||
</font>
|
||||
</span>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="220" y="489" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Get sScale1...
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 150 360 L 207.5 455.83" fill="none" stroke="#9673a6" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
<path d="M 209.42 459.04 L 204.71 456.04 L 207.5 455.83 L 209 453.47 Z" fill="#9673a6" stroke="#9673a6" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(150, 115, 166), rgb(149, 119, 163)); stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 250 540 L 250 520" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 247 540 L 253 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 253 520 L 247 520" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 531px; margin-left: 250px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="250" y="534" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 300 620 L 300 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 297 620 L 303 620" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 303 540 L 297 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 575px; margin-left: 300px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rO1 += rP1b @ sV1R
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="300" y="578" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rO1 += rP1b @ sV1R
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 630 L 140 610" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 630 L 143 630" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 610 L 137 610" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 619px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rO0 = Scale(rO0)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="623" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rO0 = Scale(rO0)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 200 520 L 154.52 538.19" fill="none" stroke="#9673a6" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
<path d="M 151.04 539.58 L 154.75 535.41 L 154.52 538.19 L 156.61 540.05 Z" fill="#9673a6" stroke="#9673a6" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(150, 115, 166), rgb(149, 119, 163)); stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 60 780 L 60 700" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 57 780 L 63 780" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 63 700 L 57 700" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 735px; margin-left: 60px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rO0 += sP1 @ sV1L
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="60" y="738" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rO0 += sP1 @ sV1L
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 299.6 700 L 300 620" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 296.6 699.99 L 302.6 700.01" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 303 620.01 L 297 619.99" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 655px; margin-left: 300px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
rO1 += sP0 @ sV0R
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="300" y="658" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
rO1 += sP0 @ sV0R
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 80 L 74.47 108.08" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 71.03 109.56 L 74.64 105.29 L 74.47 108.08 L 76.61 109.89 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 230 230 L 286.56 286.56" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 289.21 289.21 L 283.91 287.44 L 286.56 286.56 L 287.44 283.91 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 130 370 L 72.92 446.11" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 70.67 449.11 L 71.67 443.61 L 72.92 446.11 L 75.67 446.61 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 260 530 L 285.38 538.46" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 288.94 539.65 L 283.41 540.44 L 285.38 538.46 L 284.99 535.69 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 139.8 580 L 139.8 560" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 136.8 580 L 142.8 580" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 142.8 560 L 136.8 560" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 569px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
sP0 = Scale(rP0)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="573" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
sP0 = Scale(rP0)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="40" y="20" width="60" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 35px; margin-left: 70px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Tensor
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="70" y="38" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Tensor
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="125" y="20" width="50" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 35px; margin-left: 150px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
CUDA
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="150" y="38" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
CUDA
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="210" y="20" width="50" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 35px; margin-left: 235px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
CUDA
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="235" y="38" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
CUDA
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="280" y="20" width="60" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 35px; margin-left: 310px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Tensor
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="310" y="38" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Tensor
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="70" y="0" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 15px; margin-left: 115px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Warpgroup 0
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="115" y="18" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Warpgroup 0
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="220" y="0" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 15px; margin-left: 265px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Warpgroup 1
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="265" y="18" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Warpgroup 1
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 40 50 L 340 50" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 150 580 L 205.13 580" fill="none" stroke="#9673a6" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
<path d="M 208.88 580 L 203.88 582.5 L 205.13 580 L 203.88 577.5 Z" fill="#9673a6" stroke="#9673a6" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(150, 115, 166), rgb(149, 119, 163)); stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 600 L 220 580" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 600 L 223 600" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 580 L 217 580" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 591px; margin-left: 220px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="220" y="594" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 230 590 L 285.65 617.82" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 289 619.5 L 283.41 619.5 L 285.65 617.82 L 285.65 615.03 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 210 600 L 154.8 609.2" fill="none" stroke="#9673a6" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
<path d="M 151.1 609.82 L 155.62 606.53 L 154.8 609.2 L 156.45 611.46 Z" fill="#9673a6" stroke="#9673a6" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(150, 115, 166), rgb(149, 119, 163)); stroke: light-dark(rgb(150, 115, 166), rgb(149, 119, 163));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 70 540 L 125.13 540" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 128.88 540 L 123.88 542.5 L 125.13 540 L 123.88 537.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 650 L 140 630" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 650 L 143 650" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 630 L 137 630" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 641px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="644" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 150 90 L 149.86 60" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 147 90.01 L 153 89.99" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 152.86 59.99 L 146.86 60.01" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 76px; margin-left: 150px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Pipelined TMA wait and issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="150" y="79" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Pipelined TMA wait and issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 130 640 L 73.44 696.56" fill="none" stroke="#b85450" stroke-miterlimit="10" stroke-dasharray="1 1" pointer-events="stroke" style="stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
<path d="M 70.79 699.21 L 72.56 693.91 L 73.44 696.56 L 76.09 697.44 Z" fill="#b85450" stroke="#b85450" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(184, 84, 80), rgb(215, 129, 126)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 240 L 220 210" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 240 L 223 240" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 210 L 217 210" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 226px; margin-left: 220px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Pipelined TMA wait and issue
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="220" y="229" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Pipelined TMA wait and issue
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 540 L 220 520" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 540 L 223 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 520 L 217 520" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 531px; margin-left: 210px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
sP1 = rP1b
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="210" y="534" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
sP1 = rP1b
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 90 290 L 87.5 290 Q 85 290 85 300 L 85 307.5 Q 85 315 82.5 315 L 81.25 315 Q 80 315 82.5 315 L 83.75 315 Q 85 315 85 325 L 85 332.5 Q 85 340 87.5 340 L 90 340" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="0" y="300" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 315px; margin-left: 45px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
|
||||
<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
|
||||
wg0-bunch-0
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="45" y="319" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">
|
||||
wg0-bunch-0
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 280 460 L 277.5 460 Q 275 460 275 470 L 275 477.5 Q 275 485 272.5 485 L 271.25 485 Q 270 485 272.5 485 L 273.75 485 Q 275 485 275 495 L 275 502.5 Q 275 510 277.5 510 L 280 510" fill="none" stroke="#000000" stroke-miterlimit="10" transform="translate(275,0)scale(-1,1)translate(-275,0)" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="270" y="470" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 485px; margin-left: 315px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
|
||||
<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
|
||||
wg1-bunch-0
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="315" y="489" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">
|
||||
wg1-bunch-0
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 560 L 140 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 560 L 143 560" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 540 L 137 540" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 549px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue TMA (nxt V0L)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="553" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue TMA (nxt V0L)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 140 800 L 140 780" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 137 800 L 143 800" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 143 780 L 137 780" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 789px; margin-left: 140px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue TMA (nxt V1L)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="140" y="793" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue TMA (nxt V1L)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 70 780 L 125.13 780" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 128.88 780 L 123.88 782.5 L 125.13 780 L 123.88 777.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 290 620 L 234.87 620" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 231.12 620 L 236.12 617.5 L 234.87 620 L 236.12 622.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 640 L 220 620" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 640 L 223 640" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 620 L 217 620" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 631px; margin-left: 220px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue TMA (nxt V1R)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="220" y="634" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue TMA (nxt V1R)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 220 720 L 220 700" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 217 720 L 223 720" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
<path d="M 223 700 L 217 700" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 711px; margin-left: 220px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; background-color: #ffffff; ">
|
||||
<div style="display: inline-block; font-size: 11px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; background-color: light-dark(#ffffff, var(--ge-dark-color, #121212)); white-space: nowrap; ">
|
||||
Issue TMA (nxt V0R)
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="220" y="714" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="11px" text-anchor="middle">
|
||||
Issue TMA (nxt V0R)
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 290 700 L 234.87 700" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 231.12 700 L 236.12 697.5 L 234.87 700 L 236.12 702.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 70 280 L 125.13 280" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 128.88 280 L 123.88 282.5 L 125.13 280 L 123.88 277.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 290 460 L 234.87 460" fill="none" stroke="#d6b656" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
<path d="M 231.12 460 L 236.12 457.5 L 234.87 460 L 236.12 462.5 Z" fill="#d6b656" stroke="#d6b656" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(214, 182, 86), rgb(109, 81, 0)); stroke: light-dark(rgb(214, 182, 86), rgb(109, 81, 0));"/>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="310" y="80" width="190" height="40" fill="none" stroke="#c0c0c0" stroke-dasharray="8 8" pointer-events="all" style="stroke: light-dark(rgb(192, 192, 192), rgb(127, 127, 127));"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 100px; margin-left: 312px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">
|
||||
<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
|
||||
sXX: Stored on shared memory
|
||||
<div>
|
||||
rXX: Stored on register file
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="312" y="104" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px">
|
||||
sXX: Stored on shared memory...
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 0 285 L 200 285" fill="none" stroke="#82b366" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(130, 179, 102), rgb(68, 110, 44));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 200 465 L 500 465" fill="none" stroke="#82b366" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(130, 179, 102), rgb(68, 110, 44));"/>
|
||||
</g>
|
||||
<g>
|
||||
<path d="M 200 285 L 200 465" fill="none" stroke="#82b366" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(130, 179, 102), rgb(68, 110, 44));"/>
|
||||
</g>
|
||||
<g>
|
||||
<rect x="310" y="430" width="200" height="40" fill="none" stroke="none" pointer-events="all"/>
|
||||
</g>
|
||||
<g>
|
||||
<g transform="translate(-0.5 -0.5)">
|
||||
<switch>
|
||||
<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
|
||||
<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 450px; margin-left: 410px;">
|
||||
<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
|
||||
<div style="display: inline-block; font-size: 10px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
|
||||
<font style="font-size: 9px;">
|
||||
Loop boundary in our code
|
||||
</font>
|
||||
<div>
|
||||
<font style="font-size: 9px;">
|
||||
(plz refer to comments in `wg1_subroutine`)
|
||||
</font>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</div>
|
||||
</foreignObject>
|
||||
<text x="410" y="453" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="10px" text-anchor="middle">
|
||||
Loop boundary in our code...
|
||||
</text>
|
||||
</switch>
|
||||
</g>
|
||||
</g>
|
||||
</g>
|
||||
<switch>
|
||||
<g requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility"/>
|
||||
<a transform="translate(0,-5)" xlink:href="https://www.drawio.com/doc/faq/svg-export-text-problems" target="_blank">
|
||||
<text text-anchor="middle" font-size="10px" x="50%" y="100%">
|
||||
Text is not SVG - cannot display
|
||||
</text>
|
||||
</a>
|
||||
</switch>
|
||||
</svg>
|
||||
|
After Width: | Height: | Size: 73 KiB |
Loading…
Reference in New Issue
Block a user