已合并
docs: inplace_add_atomic README 修正引用块加粗渲染并新增数据流转 SVG 配图 #6223
qixingkai创建于 17 天前
docs: inplace_add_atomic README 修正引用块加粗渲染并新增数据流转 SVG 配图 #6223
已合并
qixingkai创建于 17 天前
共 14 个文件变更+811-10
Rexamples/03_simt_api/03_best_practices/03_instruction_optimizations/inplace_add_atomic/figures/inplace_add_atomic_dataflow.png→docs/zh/guide/figures/inplace_add_atomic_dataflow_atomic.png+2-2
@@ -1,3 +1,3 @@
1version https://git-lfs.github.com/spec/v11version https://git-lfs.github.com/spec/v1
2-oid sha256:768ba4632756a8436f2db49e134285d82505afb4a3141ef110eb9b8c96bd6f0b2+oid sha256:a81bfb0d193409c827ac4d2efc78394aa2dc14c51aa7dd1d3e1ca3a7f20cba0e
3-size 1172903+size 37693
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:8a0102684dbdfc6e65cad83e014daadc1fcd4278b22410151aa415d424a15fe1
3+size 29334
@@ -0,0 +1,4 @@
1+<?xml version="1.0" encoding="UTF-8"?>
2+<!-- Do not edit this file with editors other than draw.io -->
3+<!DOCTYPE svg PUBLIC "-//W3C//DTD SVG 1.1//EN" "http://www.w3.org/Graphics/SVG/1.1/DTD/svg11.dtd">
4+<svg xmlns="http://www.w3.org/2000/svg" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="951px" height="626px" viewBox="-0.5 -0.5 951 626" content="&lt;mxfile host=&quot;Electron&quot; agent=&quot;Mozilla/5.0 (X11; Linux x86_64) AppleWebKit/537.36 (KHTML, like Gecko) draw.io/24.7.17 Chrome/128.0.6613.36 Electron/32.0.1 Safari/537.36&quot; version=&quot;24.7.17&quot;&gt;&#10; &lt;diagram id=&quot;griddim-config-small-shape&quot; name=&quot;Page-1&quot;&gt;&#10; &lt;mxGraphModel dx=&quot;-549&quot; dy=&quot;-198&quot; grid=&quot;0&quot; gridSize=&quot;10&quot; guides=&quot;1&quot; tooltips=&quot;1&quot; connect=&quot;1&quot; arrows=&quot;1&quot; fold=&quot;0&quot; page=&quot;0&quot; pageScale=&quot;1&quot; pageWidth=&quot;1169&quot; pageHeight=&quot;827&quot; math=&quot;0&quot; shadow=&quot;0&quot;&gt;&#10; &lt;root&gt;&#10; &lt;mxCell id=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;1&quot; parent=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;title&quot; value=&quot;Case 2 小数据量下(16384) Task Duration 随线程块数量的变化&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=20;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;100&quot; y=&quot;30&quot; width=&quot;900&quot; height=&quot;35&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;plot&quot; value=&quot;&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=#FFFFFF;strokeColor=#000000;strokeWidth=2;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;100&quot; y=&quot;90&quot; width=&quot;820&quot; height=&quot;520&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;xaxis&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#000000;strokeWidth=2;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;610&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;920&quot; y=&quot;610&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;yaxis&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#000000;strokeWidth=2;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;610&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;90&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;grid6&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;dashed=1;dashPattern=4 4;strokeColor=#D9D9D9;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;520&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;920&quot; y=&quot;520&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;grid8&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;dashed=1;dashPattern=4 4;strokeColor=#D9D9D9;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;415&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;920&quot; y=&quot;415&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;grid10&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;dashed=1;dashPattern=4 4;strokeColor=#D9D9D9;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;310&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;920&quot; y=&quot;310&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;grid12&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;dashed=1;dashPattern=4 4;strokeColor=#D9D9D9;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;100&quot; y=&quot;205&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;920&quot; y=&quot;205&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;curve-4-8&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#2878B5;strokeWidth=4;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;135&quot; y=&quot;155&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;325&quot; y=&quot;405&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;curve-8-16&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#2878B5;strokeWidth=4;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;325&quot; y=&quot;405&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;515&quot; y=&quot;494&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;curve-16-32&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#2878B5;strokeWidth=4;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;515&quot; y=&quot;494&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;705&quot; y=&quot;536&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;curve-32-64&quot; value=&quot;&quot; style=&quot;endArrow=none;html=1;rounded=0;strokeColor=#2878B5;strokeWidth=4;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;705&quot; y=&quot;536&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;895&quot; y=&quot;513&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;point4&quot; value=&quot;&quot; style=&quot;ellipse;whiteSpace=wrap;html=1;aspect=fixed;fillColor=#2878B5;strokeColor=#2878B5;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;127&quot; y=&quot;147&quot; width=&quot;16&quot; height=&quot;16&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;point8&quot; value=&quot;&quot; style=&quot;ellipse;whiteSpace=wrap;html=1;aspect=fixed;fillColor=#2878B5;strokeColor=#2878B5;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;317&quot; y=&quot;397&quot; width=&quot;16&quot; height=&quot;16&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;point16&quot; value=&quot;&quot; style=&quot;ellipse;whiteSpace=wrap;html=1;aspect=fixed;fillColor=#2878B5;strokeColor=#2878B5;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;507&quot; y=&quot;486&quot; width=&quot;16&quot; height=&quot;16&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;point32&quot; value=&quot;&quot; style=&quot;ellipse;whiteSpace=wrap;html=1;aspect=fixed;fillColor=#28AE67;strokeColor=#28AE67;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;694&quot; y=&quot;525&quot; width=&quot;22&quot; height=&quot;22&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;point64&quot; value=&quot;&quot; style=&quot;ellipse;whiteSpace=wrap;html=1;aspect=fixed;fillColor=#2878B5;strokeColor=#2878B5;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;887&quot; y=&quot;505&quot; width=&quot;16&quot; height=&quot;16&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;label4&quot; value=&quot;10.831&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;95&quot; y=&quot;120&quot; width=&quot;80&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;label8&quot; value=&quot;6.569&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;285&quot; y=&quot;370&quot; width=&quot;80&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;label16&quot; value=&quot;4.500&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;475&quot; y=&quot;460&quot; width=&quot;80&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;label32&quot; value=&quot;3.696&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;665&quot; y=&quot;490&quot; width=&quot;80&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;label64&quot; value=&quot;4.133&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;855&quot; y=&quot;480&quot; width=&quot;80&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;ylabel&quot; value=&quot;Task Duration (us)&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=18;rotation=270;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;35&quot; y=&quot;280&quot; width=&quot;140&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;xlabel&quot; value=&quot;gridDim&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=18;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;470&quot; y=&quot;630&quot; width=&quot;100&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-x4&quot; value=&quot;4&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;115&quot; y=&quot;610&quot; width=&quot;40&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-x8&quot; value=&quot;8&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;305&quot; y=&quot;610&quot; width=&quot;40&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-x16&quot; value=&quot;16&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;495&quot; y=&quot;610&quot; width=&quot;40&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-x32&quot; value=&quot;32&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;685&quot; y=&quot;610&quot; width=&quot;40&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-x64&quot; value=&quot;64&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;875&quot; y=&quot;610&quot; width=&quot;40&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-y12&quot; value=&quot;12&quot; style=&quot;text;html=1;align=right;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;50&quot; y=&quot;80&quot; width=&quot;35&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-y10&quot; value=&quot;10&quot; style=&quot;text;html=1;align=right;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;50&quot; y=&quot;190&quot; width=&quot;35&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-y8&quot; value=&quot;8&quot; style=&quot;text;html=1;align=right;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;50&quot; y=&quot;295&quot; width=&quot;35&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-y6&quot; value=&quot;6&quot; style=&quot;text;html=1;align=right;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;50&quot; y=&quot;400&quot; width=&quot;35&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;tick-y4&quot; value=&quot;4&quot; style=&quot;text;html=1;align=right;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=16;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;50&quot; y=&quot;505&quot; width=&quot;35&quot; height=&quot;25&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;legend&quot; value=&quot;● 性能最优点&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#FFFFFF;strokeColor=#CCCCCC;fontSize=16;fontColor=#222222;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;735&quot; y=&quot;105&quot; width=&quot;165&quot; height=&quot;40&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;legend-dot&quot; value=&quot;●&quot; style=&quot;text;html=1;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=26;fontColor=#28AE67;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;744&quot; y=&quot;111&quot; width=&quot;28&quot; height=&quot;26&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;/root&gt;&#10; &lt;/mxGraphModel&gt;&#10; &lt;/diagram&gt;&#10;&lt;/mxfile&gt;&#10;"><defs/><g><g data-cell-id="0"><g data-cell-id="1"><g data-cell-id="title"><g><rect x="50" y="0" width="900" height="35" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 898px; height: 1px; padding-top: 18px; margin-left: 51px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 20px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">Case 2 小数据量下(16384) Task Duration 随线程块数量的变化</div></div></div></foreignObject><text x="500" y="24" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="20px" text-anchor="middle">Case 2 小数据量下(16384) Task Duration 随线程块数量的变化</text></switch></g></g></g><g data-cell-id="plot"><g><rect x="50" y="60" width="820" height="520" fill="#ffffff" stroke="#000000" stroke-width="2" pointer-events="all"/></g></g><g data-cell-id="xaxis"><g><path d="M 50 580 L 870 580" fill="none" stroke="#000000" stroke-width="2" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="yaxis"><g><path d="M 50 580 L 50 60" fill="none" stroke="#000000" stroke-width="2" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="grid6"><g><path d="M 50 490 L 870 490" fill="none" stroke="#d9d9d9" stroke-miterlimit="10" stroke-dasharray="4 4" pointer-events="stroke"/></g></g><g data-cell-id="grid8"><g><path d="M 50 385 L 870 385" fill="none" stroke="#d9d9d9" stroke-miterlimit="10" stroke-dasharray="4 4" pointer-events="stroke"/></g></g><g data-cell-id="grid10"><g><path d="M 50 280 L 870 280" fill="none" stroke="#d9d9d9" stroke-miterlimit="10" stroke-dasharray="4 4" pointer-events="stroke"/></g></g><g data-cell-id="grid12"><g><path d="M 50 175 L 870 175" fill="none" stroke="#d9d9d9" stroke-miterlimit="10" stroke-dasharray="4 4" pointer-events="stroke"/></g></g><g data-cell-id="curve-4-8"><g><path d="M 85 125 L 275 375" fill="none" stroke="#2878b5" stroke-width="4" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="curve-8-16"><g><path d="M 275 375 L 465 464" fill="none" stroke="#2878b5" stroke-width="4" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="curve-16-32"><g><path d="M 465 464 L 655 506" fill="none" stroke="#2878b5" stroke-width="4" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="curve-32-64"><g><path d="M 655 506 L 845 483" fill="none" stroke="#2878b5" stroke-width="4" stroke-miterlimit="10" pointer-events="stroke"/></g></g><g data-cell-id="point4"><g><ellipse cx="85" cy="125" rx="8" ry="8" fill="#2878b5" stroke="#2878b5" pointer-events="all"/></g></g><g data-cell-id="point8"><g><ellipse cx="275" cy="375" rx="8" ry="8" fill="#2878b5" stroke="#2878b5" pointer-events="all"/></g></g><g data-cell-id="point16"><g><ellipse cx="465" cy="464" rx="8" ry="8" fill="#2878b5" stroke="#2878b5" pointer-events="all"/></g></g><g data-cell-id="point32"><g><ellipse cx="655" cy="506" rx="11" ry="11" fill="#28ae67" stroke="#28ae67" pointer-events="all"/></g></g><g data-cell-id="point64"><g><ellipse cx="845" cy="483" rx="8" ry="8" fill="#2878b5" stroke="#2878b5" pointer-events="all"/></g></g><g data-cell-id="label4"><g><rect x="45" y="90" width="80" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 78px; height: 1px; padding-top: 103px; margin-left: 46px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">10.831</div></div></div></foreignObject><text x="85" y="107" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">10.831</text></switch></g></g></g><g data-cell-id="label8"><g><rect x="235" y="340" width="80" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 78px; height: 1px; padding-top: 353px; margin-left: 236px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">6.569</div></div></div></foreignObject><text x="275" y="357" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">6.569</text></switch></g></g></g><g data-cell-id="label16"><g><rect x="425" y="430" width="80" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 78px; height: 1px; padding-top: 443px; margin-left: 426px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">4.500</div></div></div></foreignObject><text x="465" y="447" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">4.500</text></switch></g></g></g><g data-cell-id="label32"><g><rect x="615" y="460" width="80" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 78px; height: 1px; padding-top: 473px; margin-left: 616px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">3.696</div></div></div></foreignObject><text x="655" y="477" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">3.696</text></switch></g></g></g><g data-cell-id="label64"><g><rect x="805" y="450" width="80" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 78px; height: 1px; padding-top: 463px; margin-left: 806px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">4.133</div></div></div></foreignObject><text x="845" y="467" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">4.133</text></switch></g></g></g><g data-cell-id="ylabel"><g><rect x="-15" y="250" width="140" height="25" fill="none" stroke="none" transform="rotate(270,55,262.5)" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)rotate(270 55 262.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 138px; height: 1px; padding-top: 263px; margin-left: -14px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 18px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">Task Duration (us)</div></div></div></foreignObject><text x="55" y="268" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="18px" text-anchor="middle">Task Duration (u...</text></switch></g></g></g><g data-cell-id="xlabel"><g><rect x="420" y="600" width="100" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 98px; height: 1px; padding-top: 613px; margin-left: 421px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 18px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">gridDim</div></div></div></foreignObject><text x="470" y="618" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="18px" text-anchor="middle">gridDim</text></switch></g></g></g><g data-cell-id="tick-x4"><g><rect x="65" y="580" width="40" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 38px; height: 1px; padding-top: 593px; margin-left: 66px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">4</div></div></div></foreignObject><text x="85" y="597" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">4</text></switch></g></g></g><g data-cell-id="tick-x8"><g><rect x="255" y="580" width="40" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 38px; height: 1px; padding-top: 593px; margin-left: 256px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">8</div></div></div></foreignObject><text x="275" y="597" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">8</text></switch></g></g></g><g data-cell-id="tick-x16"><g><rect x="445" y="580" width="40" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 38px; height: 1px; padding-top: 593px; margin-left: 446px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">16</div></div></div></foreignObject><text x="465" y="597" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">16</text></switch></g></g></g><g data-cell-id="tick-x32"><g><rect x="635" y="580" width="40" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 38px; height: 1px; padding-top: 593px; margin-left: 636px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">32</div></div></div></foreignObject><text x="655" y="597" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">32</text></switch></g></g></g><g data-cell-id="tick-x64"><g><rect x="825" y="580" width="40" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 38px; height: 1px; padding-top: 593px; margin-left: 826px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">64</div></div></div></foreignObject><text x="845" y="597" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">64</text></switch></g></g></g><g data-cell-id="tick-y12"><g><rect x="0" y="50" width="35" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-end; width: 33px; height: 1px; padding-top: 63px; margin-left: 0px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: right;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">12</div></div></div></foreignObject><text x="33" y="67" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="end">12</text></switch></g></g></g><g data-cell-id="tick-y10"><g><rect x="0" y="160" width="35" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-end; width: 33px; height: 1px; padding-top: 173px; margin-left: 0px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: right;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">10</div></div></div></foreignObject><text x="33" y="177" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="end">10</text></switch></g></g></g><g data-cell-id="tick-y8"><g><rect x="0" y="265" width="35" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-end; width: 33px; height: 1px; padding-top: 278px; margin-left: 0px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: right;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">8</div></div></div></foreignObject><text x="33" y="282" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="end">8</text></switch></g></g></g><g data-cell-id="tick-y6"><g><rect x="0" y="370" width="35" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-end; width: 33px; height: 1px; padding-top: 383px; margin-left: 0px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: right;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">6</div></div></div></foreignObject><text x="33" y="387" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="end">6</text></switch></g></g></g><g data-cell-id="tick-y4"><g><rect x="0" y="475" width="35" height="25" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-end; width: 33px; height: 1px; padding-top: 488px; margin-left: 0px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: right;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">4</div></div></div></foreignObject><text x="33" y="492" fill="rgb(0, 0, 0)" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="end">4</text></switch></g></g></g><g data-cell-id="legend"><g><rect x="685" y="75" width="165" height="40" rx="6" ry="6" fill="#ffffff" stroke="#cccccc" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 163px; height: 1px; padding-top: 95px; margin-left: 686px;"><div data-drawio-colors="color: #222222; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 16px; font-family: Helvetica; color: rgb(34, 34, 34); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">● 性能最优点</div></div></div></foreignObject><text x="768" y="100" fill="#222222" font-family="&quot;Helvetica&quot;" font-size="16px" text-anchor="middle">● 性能最优点</text></switch></g></g></g><g data-cell-id="legend-dot"><g><rect x="694" y="81" width="28" height="26" fill="none" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 26px; height: 1px; padding-top: 94px; margin-left: 695px;"><div data-drawio-colors="color: #28AE67; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 26px; font-family: Helvetica; color: rgb(40, 174, 103); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">●</div></div></div></foreignObject><text x="708" y="102" fill="#28AE67" font-family="&quot;Helvetica&quot;" font-size="26px" text-anchor="middle">●</text></switch></g></g></g></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>
@@ -0,0 +1,4 @@
1+<?xml version="1.0" encoding="UTF-8"?>
2+<!-- Do not edit this file with editors other than draw.io -->
3+<!DOCTYPE svg PUBLIC "-//W3C//DTD SVG 1.1//EN" "http://www.w3.org/Graphics/SVG/1.1/DTD/svg11.dtd">
4+<svg xmlns="http://www.w3.org/2000/svg" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="662px" height="302px" viewBox="-0.5 -0.5 662 302" content="&lt;mxfile host=&quot;Electron&quot; agent=&quot;Mozilla/5.0 (X11; Linux x86_64) AppleWebKit/537.36 (KHTML, like Gecko) draw.io/24.7.17 Chrome/128.0.6613.36 Electron/32.0.1 Safari/537.36&quot; version=&quot;24.7.17&quot;&gt;&#10; &lt;diagram id=&quot;inplace-add-atomic&quot; name=&quot;原子路径&quot;&gt;&#10; &lt;mxGraphModel dx=&quot;-424&quot; dy=&quot;-242&quot; grid=&quot;0&quot; gridSize=&quot;10&quot; guides=&quot;1&quot; tooltips=&quot;1&quot; connect=&quot;1&quot; arrows=&quot;1&quot; fold=&quot;0&quot; page=&quot;0&quot; pageScale=&quot;1&quot; pageWidth=&quot;800&quot; pageHeight=&quot;520&quot; math=&quot;0&quot; shadow=&quot;0&quot;&gt;&#10; &lt;root&gt;&#10; &lt;mxCell id=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;1&quot; parent=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;gm_at&quot; value=&quot;GM&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;70&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_at_a&quot; value=&quot;a[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;40&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_at_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_at_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_at&quot; value=&quot;L2 Cache&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;315&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_at_a&quot; value=&quot;a[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;40&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_at_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_at_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_at&quot; value=&quot;寄存器&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;560&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_at_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#E0E2E6,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;reg_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_at_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#E0E2E6,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;reg_at&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;atomic_unit&quot; value=&quot;L2 侧完成原子加&amp;lt;br&amp;gt;读旧值 + 计算加法 + 写回&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#FEE2BA;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=12;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;311&quot; y=&quot;155&quot; width=&quot;180&quot; height=&quot;48&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_b_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_at_b&quot; target=&quot;l2_at_b&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_c_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_at_c&quot; target=&quot;l2_at_c&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_b_l2_reg&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_at_b&quot; target=&quot;reg_at_b&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_c_l2_reg&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_at_c&quot; target=&quot;reg_at_c&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_a_read_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.35;exitDx=0;exitDy=0;entryX=0;entryY=0.35;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_at_a&quot; target=&quot;l2_at_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_a_wb_l2_gm&quot; style=&quot;html=1;rounded=0;strokeColor=#686D75;strokeWidth=1;endArrow=blockThin;endFill=1;exitX=0;exitY=0.75;exitDx=0;exitDy=0;entryX=1;entryY=0.75;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_at_a&quot; target=&quot;gm_at_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_at_result_reg_unit&quot; value=&quot;计算结果&quot; style=&quot;edgeStyle=orthogonalEdgeStyle;html=1;rounded=0;strokeColor=#5377FB;dashed=1;dashPattern=3 3;strokeWidth=1;endArrow=blockThin;endFill=1;fontSize=9;fontFamily=FZLanTingHei-R-GBK;fontColor=#191919;labelBackgroundColor=#FFFFFF;exitX=0;exitY=0.2;exitDx=0;exitDy=0;entryX=0.999;entryY=0.203;entryDx=0;entryDy=0;entryPerimeter=0;&quot; parent=&quot;1&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot;&gt;&#10; &lt;mxPoint x=&quot;560&quot; y=&quot;344.9999999999999&quot; as=&quot;sourcePoint&quot; /&gt;&#10; &lt;mxPoint x=&quot;484.8299999999999&quot; y=&quot;345.72&quot; as=&quot;targetPoint&quot; /&gt;&#10; &lt;Array as=&quot;points&quot;&gt;&#10; &lt;mxPoint x=&quot;522&quot; y=&quot;345&quot; /&gt;&#10; &lt;/Array&gt;&#10; &lt;/mxGeometry&gt;&#10; &lt;/mxCell&gt;&#10; &lt;/root&gt;&#10; &lt;/mxGraphModel&gt;&#10; &lt;/diagram&gt;&#10;&lt;/mxfile&gt;&#10;"><defs/><g><g data-cell-id="0"><g data-cell-id="1"><g data-cell-id="gm_at"><g><rect x="0" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 1px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">GM</div></div></div></foreignObject><text x="85" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">GM</text></switch></g></g><g data-cell-id="gm_at_a"><g><rect x="20" y="100" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 118px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">a[index]</div></div></div></foreignObject><text x="85" y="121" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">a[index]</text></switch></g></g></g><g data-cell-id="gm_at_b"><g><rect x="20" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="85" y="169" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="gm_at_c"><g><rect x="20" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="85" y="217" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="l2_at"><g><rect x="245" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 246px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">L2 Cache</div></div></div></foreignObject><text x="330" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">L2 Cache</text></switch></g></g><g data-cell-id="l2_at_a"><g><rect x="265" y="100" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 118px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">a[index]</div></div></div></foreignObject><text x="330" y="121" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">a[index]</text></switch></g></g></g><g data-cell-id="l2_at_b"><g><rect x="265" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="330" y="169" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="l2_at_c"><g><rect x="265" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="330" y="217" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="reg_at"><g><rect x="490" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 491px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">寄存器</div></div></div></foreignObject><text x="575" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">寄存器</text></switch></g></g><g data-cell-id="reg_at_b"><g><rect x="510" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#e0e2e6,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 511px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="575" y="169" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="reg_at_c"><g><rect x="510" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#e0e2e6,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 511px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="575" y="217" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="atomic_unit"><g><rect x="241" y="0" width="180" height="48" rx="7.2" ry="7.2" fill="#fee2ba" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 178px; height: 1px; padding-top: 24px; margin-left: 242px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 12px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">L2 侧完成原子加<br />读旧值 + 计算加法 + 写回</div></div></div></foreignObject><text x="331" y="28" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="12px" text-anchor="middle">L2 侧完成原子加...</text></switch></g></g></g><g data-cell-id="e_at_b_gm_l2"><g><path d="M 150.5 172 L 150.5 160 L 248.5 160 L 248.5 155 L 264.5 166 L 248.5 177 L 248.5 172 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_at_c_gm_l2"><g><path d="M 150.5 220 L 150.5 208 L 248.5 208 L 248.5 203 L 264.5 214 L 248.5 225 L 248.5 220 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_at_b_l2_reg"><g><path d="M 395.5 172 L 395.5 160 L 493.5 160 L 493.5 155 L 509.5 166 L 493.5 177 L 493.5 172 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_at_c_l2_reg"><g><path d="M 395.5 220 L 395.5 208 L 493.5 208 L 493.5 203 L 509.5 214 L 493.5 225 L 493.5 220 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_at_a_read_gm_l2"><g><path d="M 150.5 118.6 L 150.5 106.6 L 248.5 106.6 L 248.5 101.6 L 264.5 112.6 L 248.5 123.6 L 248.5 118.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_at_a_wb_l2_gm"><g><path d="M 265 127 L 158.12 127" fill="none" stroke="#686d75" stroke-miterlimit="10" pointer-events="stroke"/><path d="M 151.12 127 L 158.12 124.67 L 158.12 129.33 Z" fill="#686d75" stroke="#686d75" stroke-miterlimit="10" pointer-events="all"/></g></g><g data-cell-id="e_at_result_reg_unit"><g><path d="M 490 190 L 452 190 L 422.95 190.56" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke"/><path d="M 415.95 190.7 L 422.9 188.23 L 422.99 192.9 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><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: 190px; margin-left: 452px;"><div data-drawio-colors="color: #191919; background-color: #FFFFFF; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 9px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; background-color: rgb(255, 255, 255); white-space: nowrap;">计算结果</div></div></div></foreignObject><text x="452" y="193" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">计算结果</text></switch></g></g></g></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>
@@ -0,0 +1,4 @@
1+<?xml version="1.0" encoding="UTF-8"?>
2+<!-- Do not edit this file with editors other than draw.io -->
3+<!DOCTYPE svg PUBLIC "-//W3C//DTD SVG 1.1//EN" "http://www.w3.org/Graphics/SVG/1.1/DTD/svg11.dtd">
4+<svg xmlns="http://www.w3.org/2000/svg" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="662px" height="242px" viewBox="-0.5 -0.5 662 242" content="&lt;mxfile host=&quot;Electron&quot; agent=&quot;Mozilla/5.0 (X11; Linux x86_64) AppleWebKit/537.36 (KHTML, like Gecko) draw.io/24.7.17 Chrome/128.0.6613.36 Electron/32.0.1 Safari/537.36&quot; version=&quot;24.7.17&quot;&gt;&#10; &lt;diagram id=&quot;inplace-add-plain&quot; name=&quot;普通路径&quot;&gt;&#10; &lt;mxGraphModel dx=&quot;-424&quot; dy=&quot;-287&quot; grid=&quot;0&quot; gridSize=&quot;10&quot; guides=&quot;1&quot; tooltips=&quot;1&quot; connect=&quot;1&quot; arrows=&quot;1&quot; fold=&quot;0&quot; page=&quot;0&quot; pageScale=&quot;1&quot; pageWidth=&quot;800&quot; pageHeight=&quot;520&quot; math=&quot;0&quot; shadow=&quot;0&quot;&gt;&#10; &lt;root&gt;&#10; &lt;mxCell id=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;1&quot; parent=&quot;0&quot; /&gt;&#10; &lt;mxCell id=&quot;gm_pl&quot; value=&quot;GM&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;70&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_pl_a&quot; value=&quot;a[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;40&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_pl_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;gm_pl_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=#C2D5FE;strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;gm_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_pl&quot; value=&quot;L2 Cache&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;315&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_pl_a&quot; value=&quot;a[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;40&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_pl_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;l2_pl_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#B9E9FC,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;l2_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_pl&quot; value=&quot;寄存器&quot; style=&quot;rounded=0;whiteSpace=wrap;html=1;fillColor=none;strokeColor=#686D75;strokeWidth=0.75;container=1;collapsible=0;verticalAlign=top;align=center;spacingTop=6;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;fontStyle=1;fontColor=#191919;&quot; parent=&quot;1&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;560&quot; y=&quot;215&quot; width=&quot;170&quot; height=&quot;240&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_pl_a&quot; value=&quot;a[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#E0E2E6,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;reg_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;40&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_pl_b&quot; value=&quot;b[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#E0E2E6,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;reg_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;88&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;reg_pl_c&quot; value=&quot;c[index]&quot; style=&quot;rounded=1;whiteSpace=wrap;html=1;fillColor=light-dark(#E0E2E6,#121212);strokeColor=none;fontFamily=FZLanTingHei-R-GBK;fontSize=10.5;&quot; parent=&quot;reg_pl&quot; vertex=&quot;1&quot;&gt;&#10; &lt;mxGeometry x=&quot;20&quot; y=&quot;136&quot; width=&quot;130&quot; height=&quot;36&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_b_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_pl_b&quot; target=&quot;l2_pl_b&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_c_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_pl_c&quot; target=&quot;l2_pl_c&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_b_l2_reg&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_pl_b&quot; target=&quot;reg_pl_b&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_c_l2_reg&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.5;exitDx=0;exitDy=0;entryX=0;entryY=0.5;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_pl_c&quot; target=&quot;reg_pl_c&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_a_read_gm_l2&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.35;exitDx=0;exitDy=0;entryX=0;entryY=0.35;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;gm_pl_a&quot; target=&quot;l2_pl_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_a_wb_l2_gm&quot; value=&quot;计算结果&quot; style=&quot;html=1;rounded=0;strokeColor=#5377FB;dashed=1;dashPattern=3 3;strokeWidth=1;endArrow=blockThin;endFill=1;fontSize=9;fontFamily=FZLanTingHei-R-GBK;fontColor=#191919;labelBackgroundColor=#FFFFFF;exitX=0;exitY=0.75;exitDx=0;exitDy=0;entryX=1;entryY=0.75;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_pl_a&quot; target=&quot;gm_pl_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_a_read_l2_reg&quot; style=&quot;shape=flexArrow;endArrow=classic;html=1;rounded=0;fillColor=#E0E2E6;strokeColor=none;width=12;endSize=5;endWidth=9;exitX=1;exitY=0.35;exitDx=0;exitDy=0;entryX=0;entryY=0.35;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;l2_pl_a&quot; target=&quot;reg_pl_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;mxCell id=&quot;e_pl_a_wb_reg_l2&quot; value=&quot;计算结果&quot; style=&quot;html=1;rounded=0;strokeColor=#5377FB;dashed=1;dashPattern=3 3;strokeWidth=1;endArrow=blockThin;endFill=1;fontSize=9;fontFamily=FZLanTingHei-R-GBK;fontColor=#191919;labelBackgroundColor=#FFFFFF;exitX=0;exitY=0.75;exitDx=0;exitDy=0;entryX=1;entryY=0.75;entryDx=0;entryDy=0;&quot; parent=&quot;1&quot; source=&quot;reg_pl_a&quot; target=&quot;l2_pl_a&quot; edge=&quot;1&quot;&gt;&#10; &lt;mxGeometry relative=&quot;1&quot; as=&quot;geometry&quot; /&gt;&#10; &lt;/mxCell&gt;&#10; &lt;/root&gt;&#10; &lt;/mxGraphModel&gt;&#10; &lt;/diagram&gt;&#10;&lt;/mxfile&gt;&#10;"><defs/><g><g data-cell-id="0"><g data-cell-id="1"><g data-cell-id="gm_pl"><g><rect x="0" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 1px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">GM</div></div></div></foreignObject><text x="85" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">GM</text></switch></g></g><g data-cell-id="gm_pl_a"><g><rect x="20" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">a[index]</div></div></div></foreignObject><text x="85" y="61" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">a[index]</text></switch></g></g></g><g data-cell-id="gm_pl_b"><g><rect x="20" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="85" y="109" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="gm_pl_c"><g><rect x="20" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 21px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="85" y="157" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="l2_pl"><g><rect x="245" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 246px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">L2 Cache</div></div></div></foreignObject><text x="330" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">L2 Cache</text></switch></g></g><g data-cell-id="l2_pl_a"><g><rect x="265" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">a[index]</div></div></div></foreignObject><text x="330" y="61" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">a[index]</text></switch></g></g></g><g data-cell-id="l2_pl_b"><g><rect x="265" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="330" y="109" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="l2_pl_c"><g><rect x="265" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#b9e9fc,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 266px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="330" y="157" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="reg_pl"><g><rect x="490" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 491px;"><div data-drawio-colors="color: #191919; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; overflow-wrap: normal;">寄存器</div></div></div></foreignObject><text x="575" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">寄存器</text></switch></g></g><g data-cell-id="reg_pl_a"><g><rect x="510" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#e0e2e6,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 511px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">a[index]</div></div></div></foreignObject><text x="575" y="61" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">a[index]</text></switch></g></g></g><g data-cell-id="reg_pl_b"><g><rect x="510" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#e0e2e6,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 511px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">b[index]</div></div></div></foreignObject><text x="575" y="109" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">b[index]</text></switch></g></g></g><g data-cell-id="reg_pl_c"><g><rect x="510" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="light-dark(#e0e2e6,#121212)" stroke="none" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 511px;"><div data-drawio-colors="color: rgb(0, 0, 0); " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 10.5px; font-family: FZLanTingHei-R-GBK; color: rgb(0, 0, 0); line-height: 1.2; pointer-events: all; white-space: normal; overflow-wrap: normal;">c[index]</div></div></div></foreignObject><text x="575" y="157" fill="rgb(0, 0, 0)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">c[index]</text></switch></g></g></g></g><g data-cell-id="e_pl_b_gm_l2"><g><path d="M 150.5 112 L 150.5 100 L 248.5 100 L 248.5 95 L 264.5 106 L 248.5 117 L 248.5 112 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_c_gm_l2"><g><path d="M 150.5 160 L 150.5 148 L 248.5 148 L 248.5 143 L 264.5 154 L 248.5 165 L 248.5 160 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_b_l2_reg"><g><path d="M 395.5 112 L 395.5 100 L 493.5 100 L 493.5 95 L 509.5 106 L 493.5 117 L 493.5 112 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_c_l2_reg"><g><path d="M 395.5 160 L 395.5 148 L 493.5 148 L 493.5 143 L 509.5 154 L 493.5 165 L 493.5 160 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_a_read_gm_l2"><g><path d="M 150.5 58.6 L 150.5 46.6 L 248.5 46.6 L 248.5 41.6 L 264.5 52.6 L 248.5 63.6 L 248.5 58.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_a_wb_l2_gm"><g><path d="M 265 67 L 158.12 67" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke"/><path d="M 151.12 67 L 158.12 64.67 L 158.12 69.33 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><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: 67px; margin-left: 207px;"><div data-drawio-colors="color: #191919; background-color: #FFFFFF; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 9px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; background-color: rgb(255, 255, 255); white-space: nowrap;">计算结果</div></div></div></foreignObject><text x="207" y="70" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">计算结果</text></switch></g></g></g><g data-cell-id="e_pl_a_read_l2_reg"><g><path d="M 395.5 58.6 L 395.5 46.6 L 493.5 46.6 L 493.5 41.6 L 509.5 52.6 L 493.5 63.6 L 493.5 58.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all"/></g></g><g data-cell-id="e_pl_a_wb_reg_l2"><g><path d="M 510 67 L 403.12 67" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke"/><path d="M 396.12 67 L 403.12 64.67 L 403.12 69.33 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all"/></g><g><g transform="translate(-0.5 -0.5)"><switch><foreignObject pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility" style="overflow: visible; text-align: left;"><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: 67px; margin-left: 452px;"><div data-drawio-colors="color: #191919; background-color: #FFFFFF; " style="box-sizing: border-box; font-size: 0px; text-align: center;"><div style="display: inline-block; font-size: 9px; font-family: FZLanTingHei-R-GBK; color: rgb(25, 25, 25); line-height: 1.2; pointer-events: all; background-color: rgb(255, 255, 255); white-space: nowrap;">计算结果</div></div></div></foreignObject><text x="452" y="70" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">计算结果</text></switch></g></g></g></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>
@@ -332,6 +332,7 @@
332 - [同一Warp内避免分支跳转](operator_practice/simt_operator_optimization/control_flow/avoid_warp_branch.md)332 - [同一Warp内避免分支跳转](operator_practice/simt_operator_optimization/control_flow/avoid_warp_branch.md)
333 - [指令优化](operator_practice/simt_operator_optimization/instruction_optimization/atomic_instruction_optimization.md)333 - [指令优化](operator_practice/simt_operator_optimization/instruction_optimization/atomic_instruction_optimization.md)
334 - [原子操作指令优化](operator_practice/simt_operator_optimization/instruction_optimization/atomic_instruction_optimization.md)334 - [原子操作指令优化](operator_practice/simt_operator_optimization/instruction_optimization/atomic_instruction_optimization.md)
335+ - [原地加法指令优化](operator_practice/simt_operator_optimization/instruction_optimization/inplace_add_instruction_optimization.md)
335 336 
336 - [SIMD与SIMT混合算子性能优化](operator_practice/simd_simt_hybrid_optimization/simd_simt_hybrid_optimization.md)337 - [SIMD与SIMT混合算子性能优化](operator_practice/simd_simt_hybrid_optimization/simd_simt_hybrid_optimization.md)
337 - [概述](operator_practice/simd_simt_hybrid_optimization/overview.md)338 - [概述](operator_practice/simd_simt_hybrid_optimization/overview.md)
@@ -0,0 +1,133 @@
1+# 原地加法指令优化
2+ 
3+【优先级】中
4+ 
5+【描述】如`a[index] += b[index]`、`a[index] += b[index] * c[index]`,原地加法是算子中常见的数据更新形式。在上述使用加法运算符的程序中,累加目标`a[index]`的旧值必须先从L2 Cache读入寄存器才能参与计算,完成加法运算后,结果值从寄存器写回L2 Cache,每个元素都经历一次旧值读入、新值写出的寄存器往返。改用`asc_atomic_add(&a[index], b[index])`这类[原子加接口](../../../../api/SIMT-API/atomic_operations/atomic_operations_intro.md)后,“读出旧值—计算新值—写回新值”三步作为一个不可分割的整体在L2 Cache侧完成,旧值不再进入寄存器,省去其进出寄存器的数据往返。
6+ 
7+【样例介绍】本节以[inplace_add_atomic样例](../../../../../../examples/03_simt_api/03_best_practices/03_instruction_optimizations/inplace_add_atomic/README.md)为载体说明上述优化方向,样例对三个形状均为[4194304]的int32_t数组执行逐元素累加,包含两组对照,一组为三操作数乘加`a[index] += b[index] * c[index]`,另一组为两操作数加法`a[index] += b[index]`。样例启动64个线程块、每个线程块2048个线程,线程块数量为运行时查询到的AIV核数,各组对照仅改变累加语句的代码,即普通原地加或`asc_atomic_add`,其余保持一致。
8+ 
9+<!-- npu="950" id1 -->
10+下文性能数据均使用msOpProf工具在Ascend 950PR系列产品上采集,指标为Task Duration,即Task整体耗时、算子端到端耗时。
11+<!-- end id1 -->
12+ 
13+**优化原理**
14+ 
15+两组对照的原理相同:普通原地加中,累加目标的旧值需经L2 Cache读入寄存器,加法完成后新值再从寄存器写回L2 Cache;改用原子加后,旧值的读出、加法与写回均在L2 Cache侧作为一个不可分割的整体完成,旧值不再经“L2 Cache→寄存器→L2 Cache”的往返。两种累加方式的数据流转分别如下图所示。
16+ 
17+**图 1** 普通原地加的数据流转
18+ 
19+![](../../../figures/inplace_add_atomic_dataflow_plain.png)
20+ 
21+**图 2** 原子加的数据流转
22+ 
23+![](../../../figures/inplace_add_atomic_dataflow_atomic.png)
24+ 
25+以三操作数乘加为例:普通原地加中`a[index]`、`b[index]`、`c[index]`三者都经L2 Cache读入寄存器,乘加完成后`a[index]`的新值再从寄存器写回L2 Cache,每个元素产生3次读、1次写;原子加中只有`b[index]`与`c[index]`读入寄存器并算出乘积,乘积作为原子加的操作数下发,`a[index]`不进入寄存器,每个元素的读从3次降为2次。`b`、`c`的读取与乘法计算完全相同,性能差异仅来自省去的`a[index]`旧值数据往返。
26+ 
27+**三操作数乘加**
28+ 
29+【反例】累加目标的旧值读入寄存器,在核内完成读改写。
30+ 
31+```cpp
32+__global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_muladd_plain(
33+ int32_t* a, const int32_t* b, const int32_t* c, uint32_t element_count)
34+{
35+ uint32_t index = blockIdx.x * blockDim.x + threadIdx.x;
36+ const uint32_t stride = gridDim.x * blockDim.x;
37+ for (; index < element_count; index += stride) {
38+ a[index] += b[index] * c[index];
39+ }
40+}
41+```
42+ 
43+上述实现中,各线程按grid-stride方式遍历数组,即每个线程以总线程数为固定步长,分次处理数组中不同位置的元素。每个元素的乘加都需将`a[index]`旧值读入寄存器、加法完成后写回。该实现的性能数据如下:
44+ 
45+| 累加方式 | Task Duration(μs) | 耗时相对基线 |
46+| :---: | :---: | :---: |
47+| a[index] += b[index] * c[index] | 51.62 | 1× |
48+ 
49+【正例】改用原子加,将累加下发到L2 Cache侧完成。
50+ 
51+```cpp
52+__global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_muladd_atomic(
53+ int32_t* a, const int32_t* b, const int32_t* c, uint32_t element_count)
54+{
55+ uint32_t index = blockIdx.x * blockDim.x + threadIdx.x;
56+ const uint32_t stride = gridDim.x * blockDim.x;
57+ for (; index < element_count; index += stride) {
58+ asc_atomic_add(&a[index], b[index] * c[index]);
59+ }
60+}
61+```
62+ 
63+上述实现的功能与反例一致,线程块数量、线程数、输入读取方式均相同,唯一变化是累加语句的代码。本样例中各线程更新的目标地址互不重叠,不存在同地址竞争,算子结果的正确性并不依赖原子性;此处使用原子加是为了省去旧值进出寄存器的数据往返,获取性能收益。该实现的性能数据如下:
64+ 
65+| 累加方式 | Task Duration(μs) | 耗时相对基线 |
66+| :---: | :---: | :---: |
67+| asc_atomic_add | 47.51 | **0.92×** |
68+ 
69+根据Task Duration数据,改用原子加后执行耗时从51.62μs下降至47.51μs,降幅约8.0%。
70+ 
71+**两操作数加法**
72+ 
73+【反例】累加目标的旧值读入寄存器,在核内完成读改写。
74+ 
75+```cpp
76+__global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_add_plain(
77+ int32_t* a, const int32_t* b, uint32_t element_count)
78+{
79+ uint32_t index = blockIdx.x * blockDim.x + threadIdx.x;
80+ const uint32_t stride = gridDim.x * blockDim.x;
81+ for (; index < element_count; index += stride) {
82+ a[index] += b[index];
83+ }
84+}
85+```
86+ 
87+该实现的性能数据如下:
88+ 
89+| 累加方式 | Task Duration(μs) | 耗时相对基线 |
90+| :---: | :---: | :---: |
91+| a[index] += b[index] | 39.18 | 1× |
92+ 
93+【正例】改用原子加,将累加下发到L2 Cache侧完成。
94+ 
95+```cpp
96+__global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_add_atomic(
97+ int32_t* a, const int32_t* b, uint32_t element_count)
98+{
99+ uint32_t index = blockIdx.x * blockDim.x + threadIdx.x;
100+ const uint32_t stride = gridDim.x * blockDim.x;
101+ for (; index < element_count; index += stride) {
102+ asc_atomic_add(&a[index], b[index]);
103+ }
104+}
105+```
106+ 
107+该实现的核函数遍历方式、线程规模与反例完全一致,仅将累加语句改写为原子加。两操作数加法中每个元素的读从2次降为1次,数据往返的节省比例与三操作数乘加相当。该实现的性能数据如下:
108+ 
109+| 累加方式 | Task Duration(μs) | 耗时相对基线 |
110+| :---: | :---: | :---: |
111+| asc_atomic_add | 36.52 | **0.93×** |
112+ 
113+根据Task Duration数据,改用原子加后执行耗时从39.18μs下降至36.52μs,降幅约6.8%。
114+ 
115+两组对照的性能数据汇总如下:
116+ 
117+| 计算形式 | 累加方式 | Task Duration(μs) | 耗时相对基线 |
118+| :---: | :---: | :---: | :---: |
119+| **三操作数乘加** | **asc_atomic_add** | **47.51** | **0.92×** |
120+| 三操作数乘加 | a[index] += b[index] * c[index] | 51.62 | 1× |
121+| **两操作数加法** | **asc_atomic_add** | **36.52** | **0.93×** |
122+| 两操作数加法 | a[index] += b[index] | 39.18 | 1× |
123+ 
124+**适用边界**
125+ 
126+该优化仅适用于逐元素原地更新,使用前需确认各线程原子操作的目标地址互不重叠:
127+ 
128+- **适用的逐元素原地更新**:每个线程只更新自己负责的元素,目标地址在线程间天然不重叠,如本样例的grid-stride遍历。
129+- **不适用的归约累加**:多个线程向同一地址累加,如直方图统计,此时同一地址上的原子操作只能排队串行执行,直接使用原子加会因地址竞争显著劣化性能,应采用[分层归约](./atomic_instruction_optimization.md)——先在各线程块的UB中局部累加,再向GM汇总。
130+ 
131+此外,原子加支持的数据类型必须是支持指令优化的数据类型;对int64_t等类型,该执行路径的收益需实测确认,不使用返回值时能否生成更优原子指令与数据类型相关,参见[asc_atomic_add约束说明](../../../../api/SIMT-API/atomic_operations/asc_atomic_add.md#约束说明)。
132+ 
133+【总结】原地累加场景中,若累加目标的旧值仅用于累加、不参与其他计算,且各线程目标地址互不重叠,可优先考虑用原子加完成累加:旧值的读出、加法与写回均在L2 Cache侧作为一个不可分割的整体完成,省去旧值进出寄存器的数据往返,实测收益约7%~8%。
@@ -97,3 +97,4 @@ SIMT编程模式以线程(Thread)为基本编程单位,由开发者编写
97| 优化建议 | 优先级 | 样例 |97| 优化建议 | 优先级 | 样例 |
98| --- | :---: | --- |98| --- | :---: | --- |
99| [原子操作指令优化](./instruction_optimization/atomic_instruction_optimization.md) | 高 | [histogram](../../../../../examples/03_simt_api/03_best_practices/03_instruction_optimizations/atomic_histogram)<br>[atomic_add_perf](../../../../../examples/03_simt_api/02_features/01_api_features/02_atomic_operation/atomic_add_perf) |99| [原子操作指令优化](./instruction_optimization/atomic_instruction_optimization.md) | 高 | [histogram](../../../../../examples/03_simt_api/03_best_practices/03_instruction_optimizations/atomic_histogram)<br>[atomic_add_perf](../../../../../examples/03_simt_api/02_features/01_api_features/02_atomic_operation/atomic_add_perf) |
100+| [原地加法指令优化](./instruction_optimization/inplace_add_instruction_optimization.md) | 中 | [inplace_add_atomic](../../../../../examples/03_simt_api/03_best_practices/03_instruction_optimizations/inplace_add_atomic) |
@@ -106,9 +106,15 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_muladd_atomic(
106 106 
107性能数据呈现出一个现象:改用原子加后,Task Duration下降8.0%。差异只来自累加目标`a[index]`旧值的数据路径。107性能数据呈现出一个现象:改用原子加后,Task Duration下降8.0%。差异只来自累加目标`a[index]`旧值的数据路径。
108 108 
109-本样例两种累加方式的数据流转对比如下图所示,左侧为场景0(普通原地加),右侧为场景1(原子加),两侧的GM均存放`a`、`b`、`c`三个数组。109+本样例两种累加方式的数据流转分别如下图所示,两图的GM均存放`a`、`b`、`c`三个数组。
110 110 
111-![两种累加方式的数据流转对比](figures/inplace_add_atomic_dataflow.png)111+普通原地加的数据流转如下图:
112+ 
113+![普通原地加(场景0)的数据流转](figures/inplace_add_atomic_dataflow_plain.png)
114+ 
115+原子加的数据流转如下图:
116+ 
117+![原子加(场景1)的数据流转](figures/inplace_add_atomic_dataflow_atomic.png)
112 118 
113- **原地加法(场景0)**:`a[index]`、`b[index]`、`c[index]`三者都经L2 Cache读入寄存器,乘加完成后,`a[index]`的新值再从寄存器写回L2 Cache。119- **原地加法(场景0)**:`a[index]`、`b[index]`、`c[index]`三者都经L2 Cache读入寄存器,乘加完成后,`a[index]`的新值再从寄存器写回L2 Cache。
114- **原子加法(场景1)**:只有`b[index]`与`c[index]`读入寄存器并算出乘积,乘积作为原子加的操作数下发;读旧值、加法、写回三步作为一个不可分割的整体在L2 Cache侧完成,`a[index]`不进入寄存器,上述往返不再发生。120- **原子加法(场景1)**:只有`b[index]`与`c[index]`读入寄存器并算出乘积,乘积作为原子加的操作数下发;读旧值、加法、写回三步作为一个不可分割的整体在L2 Cache侧完成,`a[index]`不进入寄存器,上述往返不再发生。
@@ -186,6 +192,7 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_add_atomic(
186 192 
1871. **逐元素原地累加优先考虑用原子加完成**:当目标地址的旧值仅用于累加、不参与其他计算时,改用原子加可使旧值的读出、加法与写回均在L2完成,省去旧值进出寄存器的数据往返。1931. **逐元素原地累加优先考虑用原子加完成**:当目标地址的旧值仅用于累加、不参与其他计算时,改用原子加可使旧值的读出、加法与写回均在L2完成,省去旧值进出寄存器的数据往返。
1882. **注意适用边界**:本优化仅适用于逐元素原地更新,使用前需确认各线程原子操作的目标地址互不重叠。其与归约累加(多个线程向同一地址累加,如[atomic_histogram](../atomic_histogram/README.md)样例中的直方图统计)的适用边界不同:归约累加中同一地址上的原子操作只能串行执行,存在地址竞争时直接使用原子加反而会劣化性能,应采用该样例中分块累加再合并的策略。1942. **注意适用边界**:本优化仅适用于逐元素原地更新,使用前需确认各线程原子操作的目标地址互不重叠。其与归约累加(多个线程向同一地址累加,如[atomic_histogram](../atomic_histogram/README.md)样例中的直方图统计)的适用边界不同:归约累加中同一地址上的原子操作只能串行执行,存在地址竞争时直接使用原子加反而会劣化性能,应采用该样例中分块累加再合并的策略。
195+3. **数据类型需支持指令优化**:原子加支持的数据类型必须是支持指令优化的数据类型;对int64_t等类型,该执行路径的收益需实测确认,不使用返回值时能否生成更优原子指令与数据类型相关,参见[asc_atomic_add约束说明](../../../../../docs/zh/api/SIMT-API/atomic_operations/asc_atomic_add.md#约束说明)。
189 196 
190## 编译运行197## 编译运行
191 198 
@@ -227,7 +234,7 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_add_atomic(
227 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j; # NPU仿真模式234 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j; # NPU仿真模式
228 ```235 ```
229 236 
230- > **注意:**切换编译模式前需清理cmake缓存,可在build目录下执行`rm CMakeCache.txt`后重新cmake。237+ > 注意:切换编译模式前需清理cmake缓存,可在build目录下执行`rm CMakeCache.txt`后重新cmake。
231 238 
232 编译选项说明:239 编译选项说明:
233 240 
@@ -259,7 +266,7 @@ make -j
259msopprof ./inplace_add_atomic266msopprof ./inplace_add_atomic
260```267```
261 268 
262-> **关于性能采集时的Validation failed:**本样例的累加目标`a`在host侧分配时写入初值,kernel仅在其上累加。`msopprof`的warmup+replay会在同一块GM内存上重复执行kernel,`a`被累加多次,因此严格校验模式下会报`Validation failed`。该现象是replay机制与校验逻辑的固有冲突。采集性能时建议先以`-DSKIP_VALIDATION=ON`重新编译再执行`msopprof`,跳过校验。269+> 关于性能采集时的Validation failed:本样例的累加目标`a`在host侧分配时写入初值,kernel仅在其上累加。`msopprof`的warmup+replay会在同一块GM内存上重复执行kernel,`a`被累加多次,因此严格校验模式下会报`Validation failed`。该现象是replay机制与校验逻辑的固有冲突。采集性能时建议先以`-DSKIP_VALIDATION=ON`重新编译再执行`msopprof`,跳过校验。
263 270 
264命令完成后,会在默认目录下生成以"OPPROF_{timestamp}_XXX"命名的文件夹,性能数据文件夹结构示例如下:271命令完成后,会在默认目录下生成以"OPPROF_{timestamp}_XXX"命名的文件夹,性能数据文件夹结构示例如下:
265 272 
@@ -107,9 +107,11 @@ __global__ __launch_bounds__(THREADS_PER_BLOCK) void vector_muladd_atomic(
107 107 
108The performance data shows one phenomenon: switching to atomic add reduces the Task Duration by 8.0%. The difference comes solely from the data path of the old value of the accumulation target `a[index]`.108The performance data shows one phenomenon: switching to atomic add reduces the Task Duration by 8.0%. The difference comes solely from the data path of the old value of the accumulation target `a[index]`.
109 109 
110-The following figure compares the data flow of the two accumulation methods in this sample: the left side is Scenario 0 (plain in-place add), the right side is Scenario 1 (atomic add), and GM on both sides holds the three arrays `a`, `b`, and `c`.110+The following figures show the data flow of the two accumulation methods in this sample respectively; GM in both figures holds the three arrays `a`, `b`, and `c`.
111 111 
112-![Data flow comparison of the two accumulation methods](figures/inplace_add_atomic_dataflow.png)112+![Data flow of plain in-place add (Scenario 0)](figures/inplace_add_atomic_dataflow_plain.png)
113+ 
114+![Data flow of atomic add (Scenario 1)](figures/inplace_add_atomic_dataflow_atomic.png)
113 115 
114- **In-place add (Scenario 0)**: `a[index]`, `b[index]`, and `c[index]` are all read into registers through L2 Cache, and after the multiply-add completes, the new value of `a[index]` is written back from the registers to L2 Cache.116- **In-place add (Scenario 0)**: `a[index]`, `b[index]`, and `c[index]` are all read into registers through L2 Cache, and after the multiply-add completes, the new value of `a[index]` is written back from the registers to L2 Cache.
115- **Atomic add (Scenario 1)**: only `b[index]` and `c[index]` are read into registers to compute the product, which is issued as the operand of the atomic add; reading the old value, the addition, and the write-back are completed as one indivisible unit at the L2 Cache side, so `a[index]` never enters the registers and the round trip above does not happen.117- **Atomic add (Scenario 1)**: only `b[index]` and `c[index]` are read into registers to compute the product, which is issued as the operand of the atomic add; reading the old value, the addition, and the write-back are completed as one indivisible unit at the L2 Cache side, so `a[index]` never enters the registers and the round trip above does not happen.
@@ -187,6 +189,7 @@ The conclusion of both controlled groups is consistent: the scenarios using atom
187 189 
1881. **Prefer atomic add for element-wise in-place accumulation**: When the old value of the target address is used only for accumulation and does not participate in other computations, switching to atomic add completes the read, addition, and write-back of the old value in L2, eliminating the data round trip of the old value in and out of the registers.1901. **Prefer atomic add for element-wise in-place accumulation**: When the old value of the target address is used only for accumulation and does not participate in other computations, switching to atomic add completes the read, addition, and write-back of the old value in L2, eliminating the data round trip of the old value in and out of the registers.
1892. **Note the applicability boundary**: This optimization applies only to element-wise in-place updates; before use, confirm that the target addresses of each thread's atomic operations do not overlap. Its boundary differs from reduction-style accumulation (multiple threads accumulating into the same address, e.g. the histogram counting in the [atomic_histogram](../atomic_histogram/README_en.md) sample): in reduction accumulation, atomic operations on the same address can only execute serially, so under address contention, using atomic add directly degrades performance instead; use the block-local accumulation and merge strategy shown in that sample.1912. **Note the applicability boundary**: This optimization applies only to element-wise in-place updates; before use, confirm that the target addresses of each thread's atomic operations do not overlap. Its boundary differs from reduction-style accumulation (multiple threads accumulating into the same address, e.g. the histogram counting in the [atomic_histogram](../atomic_histogram/README_en.md) sample): in reduction accumulation, atomic operations on the same address can only execute serially, so under address contention, using atomic add directly degrades performance instead; use the block-local accumulation and merge strategy shown in that sample.
192+3. **The data type must support the instruction optimization**: The data type of the atomic add target must be one that supports the instruction optimization; for types such as int64_t, the benefit of this execution path needs to be verified through measurement; whether a better atomic instruction is generated when the return value is unused depends on the data type; see the [constraints of asc_atomic_add](../../../../../docs/zh/api/SIMT-API/atomic_operations/asc_atomic_add.md#约束说明).
190 193 
191## Build and Run194## Build and Run
192 195 
@@ -228,7 +231,7 @@ Run the following steps in the root directory of this sample to build and execut
228 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j; # NPU simulation mode231 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j; # NPU simulation mode
229 ```232 ```
230 233 
231- > **Note:** Clear the cmake cache before switching build modes. Execute `rm CMakeCache.txt` in the build directory, then run cmake again.234+ > Note: Clear the cmake cache before switching build modes. Execute `rm CMakeCache.txt` in the build directory, then run cmake again.
232 235 
233 Build Options Description:236 Build Options Description:
234 237 
@@ -260,7 +263,7 @@ make -j
260msopprof ./inplace_add_atomic263msopprof ./inplace_add_atomic
261```264```
262 265 
263-> **Regarding Validation failed during performance collection:** In this sample, the accumulation target `a` is initialized on the host side at allocation, and the kernel only accumulates on it. The `msopprof` warmup+replay re-executes the kernel on the same GM memory, causing `a` to be accumulated multiple times; therefore, strict validation mode reports `Validation failed`. This is an inherent conflict between the replay mechanism and the validation logic. When collecting performance, it is recommended to rebuild with `-DSKIP_VALIDATION=ON` before running `msopprof` to skip validation.266+> Regarding Validation failed during performance collection: In this sample, the accumulation target `a` is initialized on the host side at allocation, and the kernel only accumulates on it. The `msopprof` warmup+replay re-executes the kernel on the same GM memory, causing `a` to be accumulated multiple times; therefore, strict validation mode reports `Validation failed`. This is an inherent conflict between the replay mechanism and the validation logic. When collecting performance, it is recommended to rebuild with `-DSKIP_VALIDATION=ON` before running `msopprof` to skip validation.
264 267 
265After the command completes, a folder named "OPPROF_{timestamp}_XXX" is generated in the default directory. The performance data folder structure is as follows:268After the command completes, a folder named "OPPROF_{timestamp}_XXX" is generated in the default directory. The performance data folder structure is as follows:
266 269 
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:a81bfb0d193409c827ac4d2efc78394aa2dc14c51aa7dd1d3e1ca3a7f20cba0e
3+size 37693
@@ -0,0 +1,308 @@
1+<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="661px" height="301px" viewBox="-0.5 -0.5 661 301" content="&lt;mxfile&gt;&lt;diagram id=&quot;inplace-add-atomic&quot; name=&quot;原子路径&quot;&gt;7Vphb7M2EP41lrYPrQKEAB9DQjpp7ytVa6V1+xI54IJVB0fGadL9+p3BBFxIlfXNgtq3jdT6zvbZvue5u9gqcmbr/Y3Am+w7TwhD9ijZI2eObDtwXPitFC+Vwrf9SpEKmlSqUaO4o/+QSmnV2i1NSKF1lUpyziTdmMqY5zmJpaHDQvCdOeyRM3PVDU5JR3EXY0aMeUr7J01kpk8xao3+jdA00yu7tZk1rsdqRZHhhO9aKidCzkxwLqvWej8jTLnOdMviSO9hX4Lk8pQJ6XqJZTXpGbOtPt7Nd707+VKfWPBtnhA1a4SccJdRSe42OFa9O4AYdJlcM5AsaD5SxmaccQFyznMYFBZS8CdSK5HtTPzJ3HMPPbUbR9elEoCTmOZEaIsxZwxvCroq96O28EyEpIDIlNE0B53kahNYSzGcX00OC9gkzdN71TufqL2B5QVeU6aot/j7G87voR/Quvrj6ib8XY/QlLNG126t0b6wtNw6ihWoD+i73teAqM2SfUul0bghfE2keIEhutfTxNCBYVs6UnYNzax6TNaimD3WSqypnR5MN/BDQzPgDTYscYcPGLkhBfj3yJ2/wQzrvzEDPDez5+4i6tBDc+Y9UHUhaFP8ZBhsE4bauW0UnB4UnMm5QFh1QFj97CD4/qVBiDsgxD87CFbt3UugwOy+6vRNyTMcZ+SrSg1UpZy6LA1SpkpaXKxMMbX7qwSLp1/Ai2EQBQvY4Ux51FafX//foG3HwEcoXxU4lypfnwCcS5a1CpxLlbVPAM5Fy50gaV+9Q5GLpgvkj8vGHAW+agQBmvpfBXCgAuhOhrynVUQZKMlGo8iOJheMYyMsPkKW1fAMk2Y/JjwXzbNY8jWNl9ucHrlcoGiMwghNvTLjRsifIXApUCgYKY0foWBRJ+NKM0VT2OyEwS7DlYBWqloo8lXmDkM1PfC0QX+EwplaxQ7LtXy1xNRCkacagdcYhFmhgwK3NRi6Jir3qwoQoiA6H42AL4sossPpeehin+umYZlEcXsSrd9DlLH/40QhZZZdwr2Z2R1HFxneqOYjI/upelOGA5M80c15zHBR0Nj0t1meX3lfR+4R79fnLR0L62g/u5VU12VV3sieyge9oGr/VdZrV0vzvV69FF5qIQfXPLSF1iwlNtNKqZ53FOGCb0WsPWW8aEksUtL+hlZ3kKR+cz/CCEEYlvTZfIL/YXjjL3jPAm98DN54UHhXgO0SStQXvu/G14jSNr7m19DB4vcL4LMA3BPA5hfZgQDGgC5OPkeWdt6HsnPGPI2P5Wk8MMq7MlWn6w7GR6E79bHCMkmxYjx+us9oXqkXlLEWUA8tYCrveyeCZpmgeWeMzR7QDDQHAk2QYsvUn7T/MtV3t/HUdShwyhsRaGYduNVZ6lccLmTGU55jFjXa47HcIYTreN4ihJ4EF9nhUqSEWywlEeoNyoHlnfdzpnXnCU6+JPW9SjG8IizE8VNaHqp9LSt/jvLTPjGnXAdBYFK0OvhbHC2lWyIoEKR85XuTuOejoZ56yykscfTVzRmP1aGaH9NgFUfaRkPtjtmxP7727Y6ZwyLutWeblqtA7FgGtuCX1rCNGlC8cR7bfr3UqxisLDYRefBlX5CC2PznUDW8+e8rJ/oX&lt;/diagram&gt;&lt;/mxfile&gt;">
2+ <defs/>
3+ <g>
4+ <g>
5+ <rect x="0" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
6+ </g>
7+ <g>
8+ <g transform="translate(-0.5 -0.5)">
9+ <switch>
10+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
11+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 1px;">
12+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
13+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
14+ GM
15+ </div>
16+ </div>
17+ </div>
18+ </foreignObject>
19+ <text x="85" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
20+ GM
21+ </text>
22+ </switch>
23+ </g>
24+ </g>
25+ <g>
26+ <rect x="20" y="100" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
27+ </g>
28+ <g>
29+ <g transform="translate(-0.5 -0.5)">
30+ <switch>
31+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
32+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 118px; margin-left: 21px;">
33+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
34+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
35+ a[index]
36+ </div>
37+ </div>
38+ </div>
39+ </foreignObject>
40+ <text x="85" y="121" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
41+ a[index]
42+ </text>
43+ </switch>
44+ </g>
45+ </g>
46+ <g>
47+ <rect x="20" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
48+ </g>
49+ <g>
50+ <g transform="translate(-0.5 -0.5)">
51+ <switch>
52+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
53+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 21px;">
54+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
55+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
56+ b[index]
57+ </div>
58+ </div>
59+ </div>
60+ </foreignObject>
61+ <text x="85" y="169" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
62+ b[index]
63+ </text>
64+ </switch>
65+ </g>
66+ </g>
67+ <g>
68+ <rect x="20" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
69+ </g>
70+ <g>
71+ <g transform="translate(-0.5 -0.5)">
72+ <switch>
73+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
74+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 21px;">
75+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
76+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
77+ c[index]
78+ </div>
79+ </div>
80+ </div>
81+ </foreignObject>
82+ <text x="85" y="217" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
83+ c[index]
84+ </text>
85+ </switch>
86+ </g>
87+ </g>
88+ <g>
89+ <rect x="245" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
90+ </g>
91+ <g>
92+ <g transform="translate(-0.5 -0.5)">
93+ <switch>
94+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
95+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 246px;">
96+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
97+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
98+ L2 Cache
99+ </div>
100+ </div>
101+ </div>
102+ </foreignObject>
103+ <text x="330" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
104+ L2 Cache
105+ </text>
106+ </switch>
107+ </g>
108+ </g>
109+ <g>
110+ <rect x="265" y="100" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
111+ </g>
112+ <g>
113+ <g transform="translate(-0.5 -0.5)">
114+ <switch>
115+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
116+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 118px; margin-left: 266px;">
117+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
118+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
119+ a[index]
120+ </div>
121+ </div>
122+ </div>
123+ </foreignObject>
124+ <text x="330" y="121" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
125+ a[index]
126+ </text>
127+ </switch>
128+ </g>
129+ </g>
130+ <g>
131+ <rect x="265" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
132+ </g>
133+ <g>
134+ <g transform="translate(-0.5 -0.5)">
135+ <switch>
136+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
137+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 266px;">
138+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
139+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
140+ b[index]
141+ </div>
142+ </div>
143+ </div>
144+ </foreignObject>
145+ <text x="330" y="169" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
146+ b[index]
147+ </text>
148+ </switch>
149+ </g>
150+ </g>
151+ <g>
152+ <rect x="265" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
153+ </g>
154+ <g>
155+ <g transform="translate(-0.5 -0.5)">
156+ <switch>
157+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
158+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 266px;">
159+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
160+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
161+ c[index]
162+ </div>
163+ </div>
164+ </div>
165+ </foreignObject>
166+ <text x="330" y="217" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
167+ c[index]
168+ </text>
169+ </switch>
170+ </g>
171+ </g>
172+ <g>
173+ <rect x="490" y="60" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
174+ </g>
175+ <g>
176+ <g transform="translate(-0.5 -0.5)">
177+ <switch>
178+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
179+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 73px; margin-left: 491px;">
180+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
181+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
182+ 寄存器
183+ </div>
184+ </div>
185+ </div>
186+ </foreignObject>
187+ <text x="575" y="84" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
188+ 寄存器
189+ </text>
190+ </switch>
191+ </g>
192+ </g>
193+ <g>
194+ <rect x="510" y="148" width="130" height="36" rx="5.4" ry="5.4" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(18, 18, 18));"/>
195+ </g>
196+ <g>
197+ <g transform="translate(-0.5 -0.5)">
198+ <switch>
199+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
200+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 166px; margin-left: 511px;">
201+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
202+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
203+ b[index]
204+ </div>
205+ </div>
206+ </div>
207+ </foreignObject>
208+ <text x="575" y="169" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
209+ b[index]
210+ </text>
211+ </switch>
212+ </g>
213+ </g>
214+ <g>
215+ <rect x="510" y="196" width="130" height="36" rx="5.4" ry="5.4" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(18, 18, 18));"/>
216+ </g>
217+ <g>
218+ <g transform="translate(-0.5 -0.5)">
219+ <switch>
220+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
221+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 214px; margin-left: 511px;">
222+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
223+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
224+ c[index]
225+ </div>
226+ </div>
227+ </div>
228+ </foreignObject>
229+ <text x="575" y="217" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
230+ c[index]
231+ </text>
232+ </switch>
233+ </g>
234+ </g>
235+ <g>
236+ <rect x="241" y="0" width="180" height="48" rx="7.2" ry="7.2" fill="#fee2ba" stroke="none" pointer-events="all" style="fill: light-dark(rgb(254, 226, 186), rgb(62, 38, 4));"/>
237+ </g>
238+ <g>
239+ <g transform="translate(-0.5 -0.5)">
240+ <switch>
241+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
242+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 178px; height: 1px; padding-top: 24px; margin-left: 242px;">
243+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
244+ <div style="display: inline-block; font-size: 12px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
245+ L2 侧完成原子加
246+ <br/>
247+ 读旧值 + 计算加法 + 写回
248+ </div>
249+ </div>
250+ </div>
251+ </foreignObject>
252+ <text x="331" y="28" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="12px" text-anchor="middle">
253+ L2 侧完成原子加...
254+ </text>
255+ </switch>
256+ </g>
257+ </g>
258+ <g>
259+ <path d="M 150.5 172 L 150.5 160 L 248.5 160 L 248.5 155 L 264.5 166 L 248.5 177 L 248.5 172 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
260+ </g>
261+ <g>
262+ <path d="M 150.5 220 L 150.5 208 L 248.5 208 L 248.5 203 L 264.5 214 L 248.5 225 L 248.5 220 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
263+ </g>
264+ <g>
265+ <path d="M 395.5 172 L 395.5 160 L 493.5 160 L 493.5 155 L 509.5 166 L 493.5 177 L 493.5 172 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
266+ </g>
267+ <g>
268+ <path d="M 395.5 220 L 395.5 208 L 493.5 208 L 493.5 203 L 509.5 214 L 493.5 225 L 493.5 220 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
269+ </g>
270+ <g>
271+ <path d="M 150.5 118.6 L 150.5 106.6 L 248.5 106.6 L 248.5 101.6 L 264.5 112.6 L 248.5 123.6 L 248.5 118.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
272+ </g>
273+ <g>
274+ <path d="M 265 127 L 158.12 127" fill="none" stroke="#686d75" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
275+ <path d="M 151.12 127 L 158.12 124.67 L 158.12 129.33 Z" fill="#686d75" stroke="#686d75" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(104, 109, 117), rgb(139, 144, 150)); stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
276+ </g>
277+ <g>
278+ <path d="M 490 190 L 452 190 L 422.95 190.56" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
279+ <path d="M 415.95 190.7 L 422.9 188.23 L 422.99 192.9 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(83, 119, 251), rgb(101, 132, 246)); stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
280+ </g>
281+ <g>
282+ <g transform="translate(-0.5 -0.5)">
283+ <switch>
284+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
285+ <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: 190px; margin-left: 452px;">
286+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; background-color: #FFFFFF; ">
287+ <div style="display: inline-block; font-size: 9px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; background-color: light-dark(#FFFFFF, #121212); white-space: nowrap; ">
288+ 计算结果
289+ </div>
290+ </div>
291+ </div>
292+ </foreignObject>
293+ <text x="452" y="193" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">
294+ 计算结果
295+ </text>
296+ </switch>
297+ </g>
298+ </g>
299+ </g>
300+ <switch>
301+ <g requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility"/>
302+ <a transform="translate(0,-5)" xlink:href="https://www.drawio.com/doc/faq/svg-export-text-problems" target="_blank">
303+ <text text-anchor="middle" font-size="10px" x="50%" y="100%">
304+ Text is not SVG - cannot display
305+ </text>
306+ </a>
307+ </switch>
308+</svg>
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:8a0102684dbdfc6e65cad83e014daadc1fcd4278b22410151aa415d424a15fe1
3+size 29334
@@ -0,0 +1,327 @@
1+<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="661px" height="241px" viewBox="-0.5 -0.5 661 241" content="&lt;mxfile&gt;&lt;diagram id=&quot;inplace-add-plain&quot; name=&quot;普通路径&quot;&gt;7VpRb9sqGP01SNtDp9jUcfwYJ3avtE2a1krr9lIRTG1UYiJMm3S/fmBDYuZk6u3NtdUszUP5DmDDdw4cggLgbLm5EmhVfOYZYcAfZRsA58D3PTiO1D+NPDfIZRQ2QC5o1kCjHXBNfxLT06KPNCOVwRpIcs4kXbkg5mVJsHQwJARfu83uOXPfukI56QDXGDHi9NPoN5rJokEno1brfwjNC/PmwD5miWxbA1QFyvi6BcEEwJngXDal5WZGmM6dm5b0QO12XIKU8iUd8uXdijWdnhB7NNO7+mxGJ5/tjAV/LDOie40AjNcFleR6hbCuXSuOFVbIJVORp4r3lLEZZ1youOSlahRXUvAHYkHgw/FkPA+DbY1N4+hDDSriJKIlEeaJmDOGVhVd1OPRQ3giQlLFyJTRvFSY5HoQyERYzV93jis1SFrmN7p2PtZjU09O0ZIyLb30xydU3qh6xdbF14ur+KNpYSTnjT4EFjG58EzcmooX6Y/Cu9k3hOjBkk0LMmxcEb4kUjyrJqY2NMIwC8P3giZe72Tm2TZFS2L+pQGRkXa+ffSOflUwCviDGu5QRw8IBDFV9G9AMP+DMrx/pwyVuZk/D9KkIw+jmddQ1aWgLfEX0+C7NNjktlmAe1iA42ORsOiQsPjbSZhM+iYBd0jAfzsJns1uHywwf587fdLxDOGCnF1qIJeC1pYGsalaFr3ZFNOjv8iQeHinshhHSZSqEc50Rn39ef//Ltr2GngL9tWQ05d9nQA5fdpaQ05ftnYC5PRqd4Lk+/wOJAGYpmByWRfmIJroQhSB6eRsgAMZYDAe8ntaI5SBHDAZJX4y7nEdO8viLVigoWcYD3yb9PRpgoaeYVzwbdLTqw2SevHcqW+rzO/QUBVopYv3jGym+iZXTZqUmSnOMUNVRbHLhmuKv33xNoQc4MDO1/Ob95hsB01k3VCbCtlQeWteqMvfa5cMTDTfmLfXwbMNSpWa23bQ6qXDXbc6sv0OGljFHwU2mXLukSQSOWmfi2wFyexN9wFBCMKQpE/uxfd/phef6T0KvfgQvXhQeheK2zu185z5fTW/zipt8+ueLgZbv2eCj0LwngXsnk8GIhgpdlF2Grs0fB3L8Ij7NDq0T6OBWV7XW3W+7ByGQTIB0wRMPZCEuhCFuhDHIIJAURUpZNYRxkG+O/cKAQzDNFY1GaqK7bFaB1+QlETo2wGoHg47Vw+eK7YF4/jhpqBlA6eUbQ/hu9Nx9OLj9L77AoYWhMUIP+T1pFpN0vqvJbyRI7zwhcLzXOGFR9xf9gjPUeSw28tpmMig+8tBmt1LoiE3GD0S6yLnHeZ0dhhXYcN5mwp3v5iq61o/O4PJLw==&lt;/diagram&gt;&lt;/mxfile&gt;">
2+ <defs/>
3+ <g>
4+ <g>
5+ <rect x="0" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
6+ </g>
7+ <g>
8+ <g transform="translate(-0.5 -0.5)">
9+ <switch>
10+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
11+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 1px;">
12+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
13+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
14+ GM
15+ </div>
16+ </div>
17+ </div>
18+ </foreignObject>
19+ <text x="85" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
20+ GM
21+ </text>
22+ </switch>
23+ </g>
24+ </g>
25+ <g>
26+ <rect x="20" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
27+ </g>
28+ <g>
29+ <g transform="translate(-0.5 -0.5)">
30+ <switch>
31+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
32+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 21px;">
33+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
34+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
35+ a[index]
36+ </div>
37+ </div>
38+ </div>
39+ </foreignObject>
40+ <text x="85" y="61" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
41+ a[index]
42+ </text>
43+ </switch>
44+ </g>
45+ </g>
46+ <g>
47+ <rect x="20" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
48+ </g>
49+ <g>
50+ <g transform="translate(-0.5 -0.5)">
51+ <switch>
52+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
53+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 21px;">
54+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
55+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
56+ b[index]
57+ </div>
58+ </div>
59+ </div>
60+ </foreignObject>
61+ <text x="85" y="109" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
62+ b[index]
63+ </text>
64+ </switch>
65+ </g>
66+ </g>
67+ <g>
68+ <rect x="20" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="#c2d5fe" stroke="none" pointer-events="all" style="fill: light-dark(rgb(194, 213, 254), rgb(40, 56, 91));"/>
69+ </g>
70+ <g>
71+ <g transform="translate(-0.5 -0.5)">
72+ <switch>
73+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
74+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 21px;">
75+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
76+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
77+ c[index]
78+ </div>
79+ </div>
80+ </div>
81+ </foreignObject>
82+ <text x="85" y="157" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
83+ c[index]
84+ </text>
85+ </switch>
86+ </g>
87+ </g>
88+ <g>
89+ <rect x="245" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
90+ </g>
91+ <g>
92+ <g transform="translate(-0.5 -0.5)">
93+ <switch>
94+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
95+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 246px;">
96+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
97+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
98+ L2 Cache
99+ </div>
100+ </div>
101+ </div>
102+ </foreignObject>
103+ <text x="330" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
104+ L2 Cache
105+ </text>
106+ </switch>
107+ </g>
108+ </g>
109+ <g>
110+ <rect x="265" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
111+ </g>
112+ <g>
113+ <g transform="translate(-0.5 -0.5)">
114+ <switch>
115+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
116+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 266px;">
117+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
118+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
119+ a[index]
120+ </div>
121+ </div>
122+ </div>
123+ </foreignObject>
124+ <text x="330" y="61" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
125+ a[index]
126+ </text>
127+ </switch>
128+ </g>
129+ </g>
130+ <g>
131+ <rect x="265" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
132+ </g>
133+ <g>
134+ <g transform="translate(-0.5 -0.5)">
135+ <switch>
136+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
137+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 266px;">
138+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
139+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
140+ b[index]
141+ </div>
142+ </div>
143+ </div>
144+ </foreignObject>
145+ <text x="330" y="109" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
146+ b[index]
147+ </text>
148+ </switch>
149+ </g>
150+ </g>
151+ <g>
152+ <rect x="265" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="#b9e9fc" stroke="none" pointer-events="all" style="fill: light-dark(rgb(185, 233, 252), rgb(18, 18, 18));"/>
153+ </g>
154+ <g>
155+ <g transform="translate(-0.5 -0.5)">
156+ <switch>
157+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
158+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 266px;">
159+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
160+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
161+ c[index]
162+ </div>
163+ </div>
164+ </div>
165+ </foreignObject>
166+ <text x="330" y="157" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
167+ c[index]
168+ </text>
169+ </switch>
170+ </g>
171+ </g>
172+ <g>
173+ <rect x="490" y="0" width="170" height="240" fill="none" stroke="#686d75" stroke-width="0.75" pointer-events="all" style="stroke: light-dark(rgb(104, 109, 117), rgb(139, 144, 150));"/>
174+ </g>
175+ <g>
176+ <g transform="translate(-0.5 -0.5)">
177+ <switch>
178+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
179+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe flex-start; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 13px; margin-left: 491px;">
180+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; ">
181+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
182+ 寄存器
183+ </div>
184+ </div>
185+ </div>
186+ </foreignObject>
187+ <text x="575" y="24" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle" font-weight="bold">
188+ 寄存器
189+ </text>
190+ </switch>
191+ </g>
192+ </g>
193+ <g>
194+ <rect x="510" y="40" width="130" height="36" rx="5.4" ry="5.4" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(18, 18, 18));"/>
195+ </g>
196+ <g>
197+ <g transform="translate(-0.5 -0.5)">
198+ <switch>
199+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
200+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 58px; margin-left: 511px;">
201+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
202+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
203+ a[index]
204+ </div>
205+ </div>
206+ </div>
207+ </foreignObject>
208+ <text x="575" y="61" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
209+ a[index]
210+ </text>
211+ </switch>
212+ </g>
213+ </g>
214+ <g>
215+ <rect x="510" y="88" width="130" height="36" rx="5.4" ry="5.4" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(18, 18, 18));"/>
216+ </g>
217+ <g>
218+ <g transform="translate(-0.5 -0.5)">
219+ <switch>
220+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
221+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 106px; margin-left: 511px;">
222+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
223+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
224+ b[index]
225+ </div>
226+ </div>
227+ </div>
228+ </foreignObject>
229+ <text x="575" y="109" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
230+ b[index]
231+ </text>
232+ </switch>
233+ </g>
234+ </g>
235+ <g>
236+ <rect x="510" y="136" width="130" height="36" rx="5.4" ry="5.4" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(18, 18, 18));"/>
237+ </g>
238+ <g>
239+ <g transform="translate(-0.5 -0.5)">
240+ <switch>
241+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
242+ <div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 128px; height: 1px; padding-top: 154px; margin-left: 511px;">
243+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
244+ <div style="display: inline-block; font-size: 10.5px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: normal; word-wrap: normal; ">
245+ c[index]
246+ </div>
247+ </div>
248+ </div>
249+ </foreignObject>
250+ <text x="575" y="157" fill="light-dark(#000000, #ffffff)" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="11px" text-anchor="middle">
251+ c[index]
252+ </text>
253+ </switch>
254+ </g>
255+ </g>
256+ <g>
257+ <path d="M 150.5 112 L 150.5 100 L 248.5 100 L 248.5 95 L 264.5 106 L 248.5 117 L 248.5 112 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
258+ </g>
259+ <g>
260+ <path d="M 150.5 160 L 150.5 148 L 248.5 148 L 248.5 143 L 264.5 154 L 248.5 165 L 248.5 160 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
261+ </g>
262+ <g>
263+ <path d="M 395.5 112 L 395.5 100 L 493.5 100 L 493.5 95 L 509.5 106 L 493.5 117 L 493.5 112 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
264+ </g>
265+ <g>
266+ <path d="M 395.5 160 L 395.5 148 L 493.5 148 L 493.5 143 L 509.5 154 L 493.5 165 L 493.5 160 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
267+ </g>
268+ <g>
269+ <path d="M 150.5 58.6 L 150.5 46.6 L 248.5 46.6 L 248.5 41.6 L 264.5 52.6 L 248.5 63.6 L 248.5 58.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
270+ </g>
271+ <g>
272+ <path d="M 265 67 L 158.12 67" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
273+ <path d="M 151.12 67 L 158.12 64.67 L 158.12 69.33 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(83, 119, 251), rgb(101, 132, 246)); stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
274+ </g>
275+ <g>
276+ <g transform="translate(-0.5 -0.5)">
277+ <switch>
278+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
279+ <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: 67px; margin-left: 207px;">
280+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; background-color: #FFFFFF; ">
281+ <div style="display: inline-block; font-size: 9px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; background-color: light-dark(#FFFFFF, #121212); white-space: nowrap; ">
282+ 计算结果
283+ </div>
284+ </div>
285+ </div>
286+ </foreignObject>
287+ <text x="207" y="70" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">
288+ 计算结果
289+ </text>
290+ </switch>
291+ </g>
292+ </g>
293+ <g>
294+ <path d="M 395.5 58.6 L 395.5 46.6 L 493.5 46.6 L 493.5 41.6 L 509.5 52.6 L 493.5 63.6 L 493.5 58.6 Z" fill="#e0e2e6" stroke="none" pointer-events="all" style="fill: light-dark(rgb(224, 226, 230), rgb(41, 43, 47));"/>
295+ </g>
296+ <g>
297+ <path d="M 510 67 L 403.12 67" fill="none" stroke="#5377fb" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
298+ <path d="M 396.12 67 L 403.12 64.67 L 403.12 69.33 Z" fill="#5377fb" stroke="#5377fb" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(83, 119, 251), rgb(101, 132, 246)); stroke: light-dark(rgb(83, 119, 251), rgb(101, 132, 246));"/>
299+ </g>
300+ <g>
301+ <g transform="translate(-0.5 -0.5)">
302+ <switch>
303+ <foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
304+ <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: 67px; margin-left: 452px;">
305+ <div style="box-sizing: border-box; font-size: 0; text-align: center; color: #191919; background-color: #FFFFFF; ">
306+ <div style="display: inline-block; font-size: 9px; font-family: &quot;FZLanTingHei-R-GBK&quot;; color: light-dark(#191919, #d8d8d8); line-height: 1.2; pointer-events: all; background-color: light-dark(#FFFFFF, #121212); white-space: nowrap; ">
307+ 计算结果
308+ </div>
309+ </div>
310+ </div>
311+ </foreignObject>
312+ <text x="452" y="70" fill="#191919" font-family="&quot;FZLanTingHei-R-GBK&quot;" font-size="9px" text-anchor="middle">
313+ 计算结果
314+ </text>
315+ </switch>
316+ </g>
317+ </g>
318+ </g>
319+ <switch>
320+ <g requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility"/>
321+ <a transform="translate(0,-5)" xlink:href="https://www.drawio.com/doc/faq/svg-export-text-problems" target="_blank">
322+ <text text-anchor="middle" font-size="10px" x="50%" y="100%">
323+ Text is not SVG - cannot display
324+ </text>
325+ </a>
326+ </switch>
327+</svg>