{
  "type": "excalidraw",
  "version": 2,
  "source": "https://excalidraw.com",
  "elements": [
    {
      "id": "bg-top",
      "type": "rectangle",
      "x": 30,
      "y": 20,
      "width": 1740,
      "height": 86,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f8fafc",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8081,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "title",
      "type": "text",
      "x": 60,
      "y": 39,
      "width": 1006.5000000000006,
      "height": 37.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8082,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "FlashAttention v3  Hopper asynchrony: TMA + WGMMA + ping-pong",
      "fontSize": 30,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "FlashAttention v3  Hopper asynchrony: TMA + WGMMA + ping-pong",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "subtitle",
      "type": "text",
      "x": 60,
      "y": 75,
      "width": 671.5,
      "height": 21.25,
      "angle": 0,
      "strokeColor": "#374151",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8083,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "目标：面向 H100，把搬运、QK、softmax、PV 交错起来，让 Tensor Core 尽量不停。",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "目标：面向 H100，把搬运、QK、softmax、PV 交错起来，让 Tensor Core 尽量不停。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "core-panel",
      "type": "rectangle",
      "x": 60,
      "y": 135,
      "width": 510,
      "height": 350,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8084,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "core-title",
      "type": "text",
      "x": 85,
      "y": 160,
      "width": 194.7,
      "height": 27.5,
      "angle": 0,
      "strokeColor": "#6d28d9",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8085,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Hopper 上的新瓶颈",
      "fontSize": 22,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Hopper 上的新瓶颈",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "core-body",
      "type": "text",
      "x": 85,
      "y": 205,
      "width": 508.3000000000003,
      "height": 255,
      "angle": 0,
      "strokeColor": "#4c1d95",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8086,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "v2 在 H100 上仍未吃满：\nTensor Core 更快后，数据搬运和 softmax 更容易挡路。\n\nHopper 给了三个关键能力：\nTMA：异步搬 tile 到 SMEM\nWGMMA：warpgroup 级 Tensor Core MMA\nFP8：更高吞吐，但要控制误差\n\nv3 的核心不是少算，而是异步流水：\nproducer 负责预取下一块 K/V\nconsumer 负责当前块 QK/softmax/PV\n两个 consumer warpgroup ping-pong，互相遮住 softmax 延迟",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "v2 在 H100 上仍未吃满：\nTensor Core 更快后，数据搬运和 softmax 更容易挡路。\n\nHopper 给了三个关键能力：\nTMA：异步搬 tile 到 SMEM\nWGMMA：warpgroup 级 Tensor Core MMA\nFP8：更高吞吐，但要控制误差\n\nv3 的核心不是少算，而是异步流水：\nproducer 负责预取下一块 K/V\nconsumer 负责当前块 QK/softmax/PV\n两个 consumer warpgroup ping-pong，互相遮住 softmax 延迟",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pseudo-panel",
      "type": "rectangle",
      "x": 620,
      "y": 135,
      "width": 560,
      "height": 430,
      "angle": 0,
      "strokeColor": "#2563eb",
      "backgroundColor": "#dbeafe",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8087,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pseudo-title",
      "type": "text",
      "x": 645,
      "y": 160,
      "width": 337.70000000000016,
      "height": 27.5,
      "angle": 0,
      "strokeColor": "#1e40af",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8088,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "概念伪代码：producer-consumer",
      "fontSize": 22,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "概念伪代码：producer-consumer",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pseudo-code",
      "type": "text",
      "x": 645,
      "y": 205,
      "width": 325.6000000000002,
      "height": 400,
      "angle": 0,
      "strokeColor": "#1e3a8a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8089,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "producer warpgroup:\n  for k_block j:\n    TMA async load K[j],V[j]\n    write into circular SMEM stage\n    signal stage ready\n\nconsumer warpgroup A/B:\n  wait stage j ready\n  S_j = WGMMA(Q, K_j.T)\n  update rowmax and online softmax\n  O_partial = WGMMA(P_j, V_j)\n  ping-pong with other consumer\n\ninterleaving idea:\n  while A does softmax(j)\n  B can launch WGMMA(j+1)\n\nFP8 path:\n  block scale Q/K/V\n  incoherent processing reduces error",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "producer warpgroup:\n  for k_block j:\n    TMA async load K[j],V[j]\n    write into circular SMEM stage\n    signal stage ready\n\nconsumer warpgroup A/B:\n  wait stage j ready\n  S_j = WGMMA(Q, K_j.T)\n  update rowmax and online softmax\n  O_partial = WGMMA(P_j, V_j)\n  ping-pong with other consumer\n\ninterleaving idea:\n  while A does softmax(j)\n  B can launch WGMMA(j+1)\n\nFP8 path:\n  block scale Q/K/V\n  incoherent processing reduces error",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipeline-panel",
      "type": "rectangle",
      "x": 60,
      "y": 525,
      "width": 1120,
      "height": 340,
      "angle": 0,
      "strokeColor": "#64748b",
      "backgroundColor": "#f8fafc",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8090,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipeline-title",
      "type": "text",
      "x": 85,
      "y": 550,
      "width": 308,
      "height": 27.5,
      "angle": 0,
      "strokeColor": "#334155",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8091,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "流水线时间轴：搬运和计算重叠",
      "fontSize": 22,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "流水线时间轴：搬运和计算重叠",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "time-0",
      "type": "text",
      "x": 120,
      "y": 590,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8092,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "t0",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "t0",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "time-1",
      "type": "text",
      "x": 330,
      "y": 590,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8093,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "t1",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "t1",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "time-2",
      "type": "text",
      "x": 540,
      "y": 590,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8094,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "t2",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "t2",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "time-3",
      "type": "text",
      "x": 750,
      "y": 590,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8095,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "t3",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "t3",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "time-4",
      "type": "text",
      "x": 960,
      "y": 590,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8096,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "t4",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "t4",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "row-l-0",
      "type": "text",
      "x": 85,
      "y": 638,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8097,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Producer",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Producer",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-0-1",
      "type": "rectangle",
      "x": 120,
      "y": 620,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8098,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-0-1",
      "type": "text",
      "x": 132,
      "y": 635,
      "width": 100.09999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8099,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA load K/V0",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA load K/V0",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-0-2",
      "type": "rectangle",
      "x": 330,
      "y": 620,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8100,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-0-2",
      "type": "text",
      "x": 342,
      "y": 635,
      "width": 100.09999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8101,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA load K/V1",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA load K/V1",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-0-3",
      "type": "rectangle",
      "x": 540,
      "y": 620,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8102,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-0-3",
      "type": "text",
      "x": 552,
      "y": 635,
      "width": 100.09999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8103,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA load K/V2",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA load K/V2",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-0-4",
      "type": "rectangle",
      "x": 750,
      "y": 620,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8104,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-0-4",
      "type": "text",
      "x": 762,
      "y": 635,
      "width": 100.09999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8105,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA load K/V3",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA load K/V3",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-0-5",
      "type": "rectangle",
      "x": 960,
      "y": 620,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8106,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-0-5",
      "type": "text",
      "x": 972,
      "y": 635,
      "width": 100.09999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8107,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA load K/V4",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA load K/V4",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "row-l-1",
      "type": "text",
      "x": 85,
      "y": 708,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8108,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Consumer A",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Consumer A",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-1-1",
      "type": "rectangle",
      "x": 120,
      "y": 690,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8109,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-1-1",
      "type": "text",
      "x": 132,
      "y": 705,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8110,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "WGMMA QK0",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "WGMMA QK0",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-1-2",
      "type": "rectangle",
      "x": 330,
      "y": 690,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8111,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-1-2",
      "type": "text",
      "x": 342,
      "y": 705,
      "width": 107.79999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8112,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "softmax0 + PV0",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "softmax0 + PV0",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-1-3",
      "type": "rectangle",
      "x": 540,
      "y": 690,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8113,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-1-3",
      "type": "text",
      "x": 552,
      "y": 705,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8114,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "WGMMA QK2",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "WGMMA QK2",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-1-4",
      "type": "rectangle",
      "x": 750,
      "y": 690,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8115,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-1-4",
      "type": "text",
      "x": 762,
      "y": 705,
      "width": 107.79999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8116,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "softmax2 + PV2",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "softmax2 + PV2",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-1-5",
      "type": "rectangle",
      "x": 960,
      "y": 690,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8117,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-1-5",
      "type": "text",
      "x": 972,
      "y": 705,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8118,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "write O",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "write O",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "row-l-2",
      "type": "text",
      "x": 85,
      "y": 778,
      "width": 90,
      "height": 20,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8119,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Consumer B",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Consumer B",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-2-1",
      "type": "rectangle",
      "x": 120,
      "y": 760,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8120,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-2-1",
      "type": "text",
      "x": 132,
      "y": 775,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8121,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "idle",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "idle",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-2-2",
      "type": "rectangle",
      "x": 330,
      "y": 760,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8122,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-2-2",
      "type": "text",
      "x": 342,
      "y": 775,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8123,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "WGMMA QK1",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "WGMMA QK1",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-2-3",
      "type": "rectangle",
      "x": 540,
      "y": 760,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8124,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-2-3",
      "type": "text",
      "x": 552,
      "y": 775,
      "width": 107.79999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8125,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "softmax1 + PV1",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "softmax1 + PV1",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-2-4",
      "type": "rectangle",
      "x": 750,
      "y": 760,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8126,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-2-4",
      "type": "text",
      "x": 762,
      "y": 775,
      "width": 90,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8127,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "WGMMA QK3",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "WGMMA QK3",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipe-2-5",
      "type": "rectangle",
      "x": 960,
      "y": 760,
      "width": 170,
      "height": 50,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8128,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "pipe-t-2-5",
      "type": "text",
      "x": 972,
      "y": 775,
      "width": 107.79999999999998,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8129,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "softmax3 + PV3",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "softmax3 + PV3",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "pipeline-note",
      "type": "text",
      "x": 85,
      "y": 830,
      "width": 1183.1999999999987,
      "height": 20,
      "angle": 0,
      "strokeColor": "#334155",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8130,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "关键讲法：v3 把原先串行的 load -> QK -> softmax -> PV，拆成异步阶段。softmax 不是 Tensor Core 任务，所以用另一个 warpgroup 的 GEMM 来遮住它。",
      "fontSize": 16,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "关键讲法：v3 把原先串行的 load -> QK -> softmax -> PV，拆成异步阶段。softmax 不是 Tensor Core 任务，所以用另一个 warpgroup 的 GEMM 来遮住它。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "move-panel",
      "type": "rectangle",
      "x": 1220,
      "y": 135,
      "width": 490,
      "height": 730,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8131,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "move-title",
      "type": "text",
      "x": 1245,
      "y": 160,
      "width": 154,
      "height": 27.5,
      "angle": 0,
      "strokeColor": "#15803d",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8132,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "线程与数据搬运",
      "fontSize": 22,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "线程与数据搬运",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "v3-lane-lane-0",
      "type": "rectangle",
      "x": 1250,
      "y": 215,
      "width": 245,
      "height": 72,
      "angle": 0,
      "strokeColor": "#2563eb",
      "backgroundColor": "#dbeafe",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8133,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "v3-lane-lane-text-0",
      "type": "text",
      "x": 1268,
      "y": 232,
      "width": 90,
      "height": 42.5,
      "angle": 0,
      "strokeColor": "#2563eb",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8134,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "HBM\nK/V tiles",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "HBM\nK/V tiles",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "v3-lane-lane-1",
      "type": "rectangle",
      "x": 1527,
      "y": 215,
      "width": 245,
      "height": 72,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "#dcfce7",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8135,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "v3-lane-lane-text-1",
      "type": "text",
      "x": 1545,
      "y": 232,
      "width": 93.49999999999999,
      "height": 42.5,
      "angle": 0,
      "strokeColor": "#16a34a",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8136,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "TMA\nasync copy",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "TMA\nasync copy",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "v3-lane-lane-2",
      "type": "rectangle",
      "x": 1804,
      "y": 215,
      "width": 245,
      "height": 72,
      "angle": 0,
      "strokeColor": "#059669",
      "backgroundColor": "#d1fae5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8137,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "v3-lane-lane-text-2",
      "type": "text",
      "x": 1822,
      "y": 232,
      "width": 102.84999999999998,
      "height": 42.5,
      "angle": 0,
      "strokeColor": "#059669",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8138,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "SMEM ring\nstage 0/1/2",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "SMEM ring\nstage 0/1/2",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "v3-lane-arrow-0",
      "type": "arrow",
      "x": 1495,
      "y": 251,
      "width": 32,
      "height": 0,
      "angle": 0,
      "strokeColor": "#3b82f6",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8139,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "points": [
        [
          0,
          0
        ],
        [
          32,
          0
        ]
      ],
      "startBinding": null,
      "endBinding": null,
      "startArrowhead": null,
      "endArrowhead": "arrow"
    },
    {
      "id": "v3-lane-arrow-1",
      "type": "arrow",
      "x": 1772,
      "y": 251,
      "width": 32,
      "height": 0,
      "angle": 0,
      "strokeColor": "#3b82f6",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8140,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "points": [
        [
          0,
          0
        ],
        [
          32,
          0
        ]
      ],
      "startBinding": null,
      "endBinding": null,
      "startArrowhead": null,
      "endArrowhead": "arrow"
    },
    {
      "id": "v3-smem-to-wgmma",
      "type": "arrow",
      "x": 1490,
      "y": 330,
      "width": 0,
      "height": 75,
      "angle": 0,
      "strokeColor": "#3b82f6",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8141,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "points": [
        [
          0,
          0
        ],
        [
          0,
          75
        ]
      ],
      "startBinding": null,
      "endBinding": null,
      "startArrowhead": null,
      "endArrowhead": "arrow"
    },
    {
      "id": "wgmma-box",
      "type": "rectangle",
      "x": 1310,
      "y": 420,
      "width": 340,
      "height": 85,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#f3e8ff",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8142,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "wgmma-text",
      "type": "text",
      "x": 1330,
      "y": 445,
      "width": 168.3,
      "height": 42.5,
      "angle": 0,
      "strokeColor": "#6d28d9",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8143,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Consumer warpgroup\nWGMMA: QK 和 PV",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Consumer warpgroup\nWGMMA: QK 和 PV",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "v3-wgmma-soft",
      "type": "arrow",
      "x": 1480,
      "y": 505,
      "width": 0,
      "height": 65,
      "angle": 0,
      "strokeColor": "#3b82f6",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8144,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "points": [
        [
          0,
          0
        ],
        [
          0,
          65
        ]
      ],
      "startBinding": null,
      "endBinding": null,
      "startArrowhead": null,
      "endArrowhead": "arrow"
    },
    {
      "id": "soft-box",
      "type": "rectangle",
      "x": 1310,
      "y": 585,
      "width": 340,
      "height": 85,
      "angle": 0,
      "strokeColor": "#ea580c",
      "backgroundColor": "#ffedd5",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8145,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "soft-text",
      "type": "text",
      "x": 1330,
      "y": 610,
      "width": 233.75000000000009,
      "height": 42.5,
      "angle": 0,
      "strokeColor": "#9a3412",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8146,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "CUDA cores\nmax / exp / sum / rescale",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "CUDA cores\nmax / exp / sum / rescale",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "move-note",
      "type": "text",
      "x": 1245,
      "y": 710,
      "width": 772.6500000000001,
      "height": 21.25,
      "angle": 0,
      "strokeColor": "#14532d",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8147,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "Producer 和 consumer 分工后，线程不是一起等内存，而是有人持续搬下一块，有人持续算当前块。",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "Producer 和 consumer 分工后，线程不是一起等内存，而是有人持续搬下一块，有人持续算当前块。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "script-panel",
      "type": "rectangle",
      "x": 60,
      "y": 910,
      "width": 1650,
      "height": 150,
      "angle": 0,
      "strokeColor": "#f59e0b",
      "backgroundColor": "#fffbeb",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8148,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "script-title",
      "type": "text",
      "x": 85,
      "y": 935,
      "width": 132,
      "height": 27.5,
      "angle": 0,
      "strokeColor": "#b45309",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8149,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "视频讲解顺序",
      "fontSize": 22,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "视频讲解顺序",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "script-body",
      "type": "text",
      "x": 85,
      "y": 980,
      "width": 1753.1999999999969,
      "height": 22.5,
      "angle": 0,
      "strokeColor": "#78350f",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8150,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "先回顾 v2 仍是同步主循环；再画出 producer/consumer；然后用时间轴解释 TMA 预取、WGMMA 计算、softmax 交错；最后补 FP8：吞吐更高，但要用 block quantization 和 incoherent processing 保精度。",
      "fontSize": 18,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "先回顾 v2 仍是同步主循环；再画出 producer/consumer；然后用时间轴解释 TMA 预取、WGMMA 计算、softmax 交错；最后补 FP8：吞吐更高，但要用 block quantization 和 incoherent processing 保精度。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "legend",
      "type": "rectangle",
      "x": 60,
      "y": 1110,
      "width": 1030,
      "height": 70,
      "angle": 0,
      "strokeColor": "#64748b",
      "backgroundColor": "#f8fafc",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "dashed",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8151,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "legend-text",
      "type": "text",
      "x": 80,
      "y": 1128,
      "width": 1109.999999999999,
      "height": 18.75,
      "angle": 0,
      "strokeColor": "#334155",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8152,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "读图方法：蓝色表示 HBM/全局内存，绿色表示 SMEM/SRAM，紫色表示寄存器/累加器，橙色表示 Tensor Core 或 CUDA Core 计算。箭头表示数据搬运或流水依赖。",
      "fontSize": 15,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "读图方法：蓝色表示 HBM/全局内存，绿色表示 SMEM/SRAM，紫色表示寄存器/累加器，橙色表示 Tensor Core 或 CUDA Core 计算。箭头表示数据搬运或流水依赖。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "source",
      "type": "rectangle",
      "x": 1220,
      "y": 1110,
      "width": 520,
      "height": 70,
      "angle": 0,
      "strokeColor": "#94a3b8",
      "backgroundColor": "#f8fafc",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "dashed",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "roundness": {
        "type": 3
      },
      "seed": 8153,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false
    },
    {
      "id": "source-text",
      "type": "text",
      "x": 1240,
      "y": 1128,
      "width": 632.0999999999995,
      "height": 17.5,
      "angle": 0,
      "strokeColor": "#475569",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 8154,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "来源：FlashAttention-3 arXiv:2407.08608；NVIDIA Next Generation of FlashAttention。",
      "fontSize": 14,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "来源：FlashAttention-3 arXiv:2407.08608；NVIDIA Next Generation of FlashAttention。",
      "autoResize": true,
      "lineHeight": 1.25
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-bg",
      "type": "rectangle",
      "x": 80,
      "y": 1260.0,
      "width": 1580,
      "height": 300,
      "angle": 0,
      "strokeColor": "#64748b",
      "backgroundColor": "#e2e8f0",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 70,
      "groupIds": [],
      "seed": 93001,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "roundness": {
        "type": 3
      }
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-title",
      "type": "text",
      "x": 108,
      "y": 1282.0,
      "width": 1524,
      "height": 38,
      "angle": 0,
      "strokeColor": "#1e40af",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 1,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 91001,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "练习题与面试题补充：FlashAttention IO-aware attention",
      "fontSize": 24,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "练习题与面试题补充：FlashAttention IO-aware attention",
      "autoResize": false,
      "lineHeight": 1.25
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-practice-box",
      "type": "rectangle",
      "x": 108,
      "y": 1338.0,
      "width": 748.0,
      "height": 190,
      "angle": 0,
      "strokeColor": "#2563eb",
      "backgroundColor": "#bfdbfe",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 93002,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "roundness": {
        "type": 3
      }
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-practice-text",
      "type": "text",
      "x": 130,
      "y": 1358.0,
      "width": 725.0,
      "height": 150,
      "angle": 0,
      "strokeColor": "#243042",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 1,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 91001,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "练习题\n1. 手推一个 Q block 扫两个 K/V tile 的 online softmax 更新\n2. 画出 S tile、P tile、O tile 的生命周期\n3. 比较 v1/v2/v3/v4 在 work partition 和 pipeline 上的变化",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "练习题\n1. 手推一个 Q block 扫两个 K/V tile 的 online softmax 更新\n2. 画出 S tile、P tile、O tile 的生命周期\n3. 比较 v1/v2/v3/v4 在 work partition 和 pipeline 上的变化",
      "autoResize": false,
      "lineHeight": 1.25
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-interview-box",
      "type": "rectangle",
      "x": 884.0,
      "y": 1338.0,
      "width": 748.0,
      "height": 190,
      "angle": 0,
      "strokeColor": "#7c3aed",
      "backgroundColor": "#ddd6fe",
      "fillStyle": "solid",
      "strokeWidth": 2,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 93003,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "roundness": {
        "type": 3
      }
    },
    {
      "id": "practice-interview-panel-flashattention-v3-hopper-async-pipeline-原理详解-interview-text",
      "type": "text",
      "x": 906.0,
      "y": 1358.0,
      "width": 725.0,
      "height": 150,
      "angle": 0,
      "strokeColor": "#243042",
      "backgroundColor": "transparent",
      "fillStyle": "solid",
      "strokeWidth": 1,
      "strokeStyle": "solid",
      "roughness": 1,
      "opacity": 100,
      "groupIds": [],
      "seed": 91001,
      "version": 1,
      "isDeleted": false,
      "boundElements": null,
      "updated": 1,
      "link": null,
      "locked": false,
      "text": "面试题\n1. FlashAttention 为什么降低 HBM traffic\n2. online softmax 的 m/l/O 三个状态如何保证数值正确\n3. v2 的 work partitioning 相比 v1 改善了什么",
      "fontSize": 17,
      "fontFamily": 5,
      "textAlign": "left",
      "verticalAlign": "top",
      "containerId": null,
      "originalText": "面试题\n1. FlashAttention 为什么降低 HBM traffic\n2. online softmax 的 m/l/O 三个状态如何保证数值正确\n3. v2 的 work partitioning 相比 v1 改善了什么",
      "autoResize": false,
      "lineHeight": 1.25
    }
  ],
  "appState": {
    "gridSize": null,
    "viewBackgroundColor": "#ffffff"
  },
  "files": {}
}