{ "type": "excalidraw", "version": 2, "source": "generated", "elements": [ { "id": "el0001", "type": "text", "x": 20, "y": 10, "width": 895.4000000000001, "height": 27.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1390851129, "version": 1, "versionNonce": 647892280, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "K5: warp specialization, one producer warpgroup and one consumer warpgroup", "fontSize": 22, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "K5: warp specialization, one producer warpgroup and one consumer warpgroup", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0002", "type": "text", "x": 700, "y": 50, "width": 28.6, "height": 16.25, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1695753999, "version": 1, "versionNonce": 207388625, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "time", "fontSize": 13, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "time", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0003", "type": "arrow", "x": 600, "y": 58, "width": 240, "height": 0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 311111476, "version": 1, "versionNonce": 404285458, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 240, 0 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0004", "type": "text", "x": 20, "y": 80, "width": 165.0, "height": 45.0, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1570621945, "version": 1, "versionNonce": 249103478, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "WG 0, producer\n128 threads, 24 regs each\nwarp 0 issues TMA", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "WG 0, producer\n128 threads, 24 regs each\nwarp 0 issues TMA", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0005", "type": "rectangle", "x": 200, "y": 78, "width": 50, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 922121677, "version": 1, "versionNonce": 161042649, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0006", "type": "text", "x": 221.975, "y": 86.125, "width": 6.050000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 369140571, "version": 1, "versionNonce": 1862494043, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Q", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Q", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0007", "type": "rectangle", "x": 255, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1796035740, "version": 1, "versionNonce": 300026768, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0008", "type": "text", "x": 267.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1033639717, "version": 1, "versionNonce": 389609434, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV0", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV0", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0009", "type": "rectangle", "x": 303, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1823296039, "version": 1, "versionNonce": 253877687, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0010", "type": "text", "x": 315.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 531725348, "version": 1, "versionNonce": 958804058, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV1", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV1", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0011", "type": "rectangle", "x": 351, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 265695474, "version": 1, "versionNonce": 1703729685, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0012", "type": "text", "x": 363.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 212984477, "version": 1, "versionNonce": 949539217, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV2", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV2", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0013", "type": "rectangle", "x": 399, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 200071089, "version": 1, "versionNonce": 571981486, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0014", "type": "text", "x": 411.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1243862423, "version": 1, "versionNonce": 1800188483, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV3", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV3", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0015", "type": "rectangle", "x": 447, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 619570853, "version": 1, "versionNonce": 505913793, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0016", "type": "text", "x": 459.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1324919353, "version": 1, "versionNonce": 776213900, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV4", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV4", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0017", "type": "rectangle", "x": 505, "y": 78, "width": 120, "height": 30, "angle": 0, "strokeColor": "#495057", "backgroundColor": "#e9ecef", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 442620899, "version": 1, "versionNonce": 806899910, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0018", "type": "text", "x": 512.75, "y": 80.5, "width": 104.50000000000001, "height": 25.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1599435268, "version": 1, "versionNonce": 418461139, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "acquire: wait for a\nfree slot", "fontSize": 10, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "acquire: wait for a\nfree slot", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0019", "type": "rectangle", "x": 635, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "dashed", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 269676600, "version": 1, "versionNonce": 255985077, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0020", "type": "text", "x": 647.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 884585952, "version": 1, "versionNonce": 2132084005, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV5", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV5", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0021", "type": "rectangle", "x": 683, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "dashed", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1836494975, "version": 1, "versionNonce": 1349251824, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0022", "type": "text", "x": 695.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1999744785, "version": 1, "versionNonce": 1946412081, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV6", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV6", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0023", "type": "rectangle", "x": 731, "y": 78, "width": 44, "height": 30, "angle": 0, "strokeColor": "#1971c2", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "dashed", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1552984409, "version": 1, "versionNonce": 1287489454, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0024", "type": "text", "x": 743.925, "y": 86.125, "width": 18.150000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1066984056, "version": 1, "versionNonce": 772092315, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KV7", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "KV7", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0025", "type": "text", "x": 20, "y": 125, "width": 738.1, "height": 13.75, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1048386556, "version": 1, "versionNonce": 351564608, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "the producer runs ahead until all 5 slots are in flight, then blocks in producer_acquire until a release (green) frees one", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "the producer runs ahead until all 5 slots are in flight, then blocks in producer_acquire until a release (green) frees one", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0026", "type": "text", "x": 20, "y": 160, "width": 171.60000000000002, "height": 45.0, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1289560150, "version": 1, "versionNonce": 2126508551, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "WG 1, consumer\n128 threads, 240 regs each\nWGMMA + softmax + epilogue", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "WG 1, consumer\n128 threads, 240 regs each\nWGMMA + softmax + epilogue", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0027", "type": "rectangle", "x": 200, "y": 158, "width": 60, "height": 30, "angle": 0, "strokeColor": "#495057", "backgroundColor": "#e9ecef", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1475216846, "version": 1, "versionNonce": 1927728187, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0028", "type": "text", "x": 213.5, "y": 166.75, "width": 33.0, "height": 12.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1236683273, "version": 1, "versionNonce": 314395343, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "wait Q", "fontSize": 10, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "wait Q", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0029", "type": "rectangle", "x": 265, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 507088657, "version": 1, "versionNonce": 1795823849, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0030", "type": "text", "x": 278.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 708506837, "version": 1, "versionNonce": 1469118511, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "QK", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "QK", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0031", "type": "rectangle", "x": 307, "y": 158, "width": 50, "height": 30, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 652768598, "version": 1, "versionNonce": 2100080515, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0032", "type": "text", "x": 310.825, "y": 166.125, "width": 42.35, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1811180650, "version": 1, "versionNonce": 168393880, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "softmax", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "softmax", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0033", "type": "rectangle", "x": 359, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 333377415, "version": 1, "versionNonce": 1347535309, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0034", "type": "text", "x": 372.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1460814403, "version": 1, "versionNonce": 1504004732, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "PV", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "PV", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0035", "type": "arrow", "x": 401, "y": 188, "width": 0, "height": -53, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 2133201996, "version": 1, "versionNonce": 1959386987, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, -53 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0036", "type": "rectangle", "x": 409, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 295334610, "version": 1, "versionNonce": 401991736, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0037", "type": "text", "x": 422.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1159380354, "version": 1, "versionNonce": 2036236842, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "QK", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "QK", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0038", "type": "rectangle", "x": 451, "y": 158, "width": 50, "height": 30, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 279172787, "version": 1, "versionNonce": 260573196, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0039", "type": "text", "x": 454.825, "y": 166.125, "width": 42.35, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1329753548, "version": 1, "versionNonce": 1914012529, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "softmax", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "softmax", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0040", "type": "rectangle", "x": 503, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1222328496, "version": 1, "versionNonce": 1656961616, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0041", "type": "text", "x": 516.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1490376254, "version": 1, "versionNonce": 96907016, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "PV", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "PV", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0042", "type": "arrow", "x": 545, "y": 188, "width": 0, "height": -53, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 1982966163, "version": 1, "versionNonce": 1526706730, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, -53 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0043", "type": "rectangle", "x": 553, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 721762279, "version": 1, "versionNonce": 502922617, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0044", "type": "text", "x": 566.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 2120395275, "version": 1, "versionNonce": 253207297, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "QK", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "QK", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0045", "type": "rectangle", "x": 595, "y": 158, "width": 50, "height": 30, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 937195260, "version": 1, "versionNonce": 1234510746, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0046", "type": "text", "x": 598.825, "y": 166.125, "width": 42.35, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 555512016, "version": 1, "versionNonce": 1063497604, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "softmax", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "softmax", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0047", "type": "rectangle", "x": 647, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1708957521, "version": 1, "versionNonce": 1679116189, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0048", "type": "text", "x": 660.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 2132480061, "version": 1, "versionNonce": 346094056, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "PV", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "PV", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0049", "type": "arrow", "x": 689, "y": 188, "width": 0, "height": -53, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 714537755, "version": 1, "versionNonce": 1929245187, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, -53 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0050", "type": "rectangle", "x": 697, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1725048951, "version": 1, "versionNonce": 1193309984, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0051", "type": "text", "x": 710.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 588093311, "version": 1, "versionNonce": 1849076401, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "QK", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "QK", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0052", "type": "rectangle", "x": 739, "y": 158, "width": 50, "height": 30, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1195809358, "version": 1, "versionNonce": 1783684942, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0053", "type": "text", "x": 742.825, "y": 166.125, "width": 42.35, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1540910401, "version": 1, "versionNonce": 1633982922, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "softmax", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "softmax", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0054", "type": "rectangle", "x": 791, "y": 158, "width": 40, "height": 30, "angle": 0, "strokeColor": "#e8590c", "backgroundColor": "#ffd8a8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 991070208, "version": 1, "versionNonce": 648200382, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0055", "type": "text", "x": 804.95, "y": 166.125, "width": 12.100000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 356416555, "version": 1, "versionNonce": 756849393, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "PV", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "PV", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0056", "type": "arrow", "x": 833, "y": 188, "width": 0, "height": -53, "angle": 0, "strokeColor": "#2f9e44", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 649821630, "version": 1, "versionNonce": 996247159, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, -53 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0057", "type": "text", "x": 200, "y": 195, "width": 732.0500000000001, "height": 13.75, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1002170859, "version": 1, "versionNonce": 51810463, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "each tile: consumer_wait on the full barrier, compute, consumer_release (green) which lets the producer refill that slot.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "each tile: consumer_wait on the full barrier, compute, consumer_release (green) which lets the producer refill that slot.", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0058", "type": "rectangle", "x": 20, "y": 240, "width": 820, "height": 120, "angle": 0, "strokeColor": "#495057", "backgroundColor": "#f8f9fa", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 2082899072, "version": 1, "versionNonce": 783156688, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0059", "type": "text", "x": 341.45, "y": 246, "width": 177.10000000000002, "height": 17.5, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1128488134, "version": 1, "versionNonce": 1210883261, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "registers per SM: 65536", "fontSize": 14, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "registers per SM: 65536", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0060", "type": "rectangle", "x": 40, "y": 272, "width": 180, "height": 70, "angle": 0, "strokeColor": "#495057", "backgroundColor": "#e9ecef", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 17581914, "version": 1, "versionNonce": 625675343, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0061", "type": "text", "x": 54.099999999999994, "y": 284.5, "width": 151.8, "height": 45.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1799361520, "version": 1, "versionNonce": 1585932014, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "K4.2: 128 threads,\nlaunch-time allocation,\nspills at d = 128", "fontSize": 12, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "K4.2: 128 threads,\nlaunch-time allocation,\nspills at d = 128", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0062", "type": "arrow", "x": 225, "y": 307, "width": 40, "height": 0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 2 }, "seed": 1368426742, "version": 1, "versionNonce": 538981927, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 40, 0 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0063", "type": "rectangle", "x": 270, "y": 272, "width": 550, "height": 70, "angle": 0, "strokeColor": "#495057", "backgroundColor": "#e9ecef", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 231897702, "version": 1, "versionNonce": 1961269854, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0064", "type": "text", "x": 287.875, "y": 293.25, "width": 514.25, "height": 27.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1685254564, "version": 1, "versionNonce": 1709696036, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "K5: setmaxnreg.dec to 24 in the producer, .inc to 240 in the consumer\n128 x 24 + 128 x 240 = 33792 per CTA; acc_pv (64 at d = 128), acc_qk (32), P (16) fit", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "K5: setmaxnreg.dec to 24 in the producer, .inc to 240 in the consumer\n128 x 24 + 128 x 240 = 33792 per CTA; acc_pv (64 at d = 128), acc_qk (32), P (16) fit", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0065", "type": "text", "x": 20, "y": 375, "width": 1128.6000000000001, "height": 15.0, "angle": 0, "strokeColor": "#495057", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": null, "seed": 1713601029, "version": 1, "versionNonce": 1692732590, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "What changed vs K4.2: the same consumer body, but the TMA issue loop moved to its own warpgroup, so loads are no longer interleaved with compute in one instruction stream.", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "What changed vs K4.2: the same consumer body, but the TMA issue loop moved to its own warpgroup, so loads are no longer interleaved with compute in one instruction stream.", "autoResize": true, "lineHeight": 1.25 } ], "appState": { "gridSize": 20, "gridStep": 5, "gridModeEnabled": false, "viewBackgroundColor": "#fdf8f6" }, "files": {} }