{ "type": "excalidraw", "version": 2, "source": "generated", "elements": [ { "id": "el0001", "type": "text", "x": 20, "y": 10, "width": 822.8000000000001, "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": "K4.2: both matmuls on the tensor cores, P never leaves the registers", "fontSize": 22, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "K4.2: both matmuls on the tensor cores, P never leaves the registers", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0002", "type": "rectangle", "x": 20, "y": 60, "width": 120, "height": 70, "angle": 0, "strokeColor": "#f08c00", "backgroundColor": "#ffec99", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1695753999, "version": 1, "versionNonce": 207388625, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0003", "type": "text", "x": 33.525, "y": 78.75, "width": 92.95, "height": 32.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": 311111476, "version": 1, "versionNonce": 404285458, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "sQ (64, d)\nswizzled SMEM", "fontSize": 13, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "sQ (64, d)\nswizzled SMEM", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0004", "type": "rectangle", "x": 20, "y": 150, "width": 120, "height": 70, "angle": 0, "strokeColor": "#f08c00", "backgroundColor": "#ffec99", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1570621945, "version": 1, "versionNonce": 249103478, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0005", "type": "text", "x": 33.525, "y": 168.75, "width": 92.95, "height": 32.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": 922121677, "version": 1, "versionNonce": 161042649, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "sK_j (64, d)\nswizzled SMEM", "fontSize": 13, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "sK_j (64, d)\nswizzled SMEM", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0006", "type": "arrow", "x": 145, "y": 95, "width": 70, "height": 35, "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": 369140571, "version": 1, "versionNonce": 1862494043, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 70, 35 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0007", "type": "arrow", "x": 145, "y": 185, "width": 70, "height": -35, "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": 1796035740, "version": 1, "versionNonce": 300026768, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 70, -35 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0008", "type": "rectangle", "x": 220, "y": 105, "width": 150, "height": 60, "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": 1033639717, "version": 1, "versionNonce": 389609434, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0009", "type": "text", "x": 260.35, "y": 117.5, "width": 69.30000000000001, "height": 35.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": 1823296039, "version": 1, "versionNonce": 253877687, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "WGMMA\nS = Q K^T", "fontSize": 14, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "WGMMA\nS = Q K^T", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0010", "type": "text", "x": 234.5, "y": 168, "width": 121.00000000000001, "height": 27.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": 531725348, "version": 1, "versionNonce": 958804058, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "d / 16 instructions,\n64 x 64 x 16 each", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "d / 16 instructions,\n64 x 64 x 16 each", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0011", "type": "arrow", "x": 375, "y": 135, "width": 60, "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": 265695474, "version": 1, "versionNonce": 1703729685, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 60, 0 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0012", "type": "rectangle", "x": 440, "y": 90, "width": 170, "height": 90, "angle": 0, "strokeColor": "#e03131", "backgroundColor": "#ffc9c9", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 212984477, "version": 1, "versionNonce": 949539217, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0013", "type": "text", "x": 460.65, "y": 110.625, "width": 128.70000000000002, "height": 48.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": 200071089, "version": 1, "versionNonce": 571981486, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "acc_qk\n32 FP32 per thread\n(C fragment)", "fontSize": 13, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "acc_qk\n32 FP32 per thread\n(C fragment)", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0014", "type": "arrow", "x": 525, "y": 185, "width": 0, "height": 40, "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": 1243862423, "version": 1, "versionNonce": 1800188483, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, 40 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0015", "type": "rectangle", "x": 440, "y": 230, "width": 170, "height": 70, "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": 619570853, "version": 1, "versionNonce": 505913793, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0016", "type": "text", "x": 455.7, "y": 242.5, "width": 138.60000000000002, "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": 1324919353, "version": 1, "versionNonce": 776213900, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "softmax in place:\nscale, mask, row max,\nexp, row sum", "fontSize": 12, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "softmax in place:\nscale, mask, row max,\nexp, row sum", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0017", "type": "arrow", "x": 525, "y": 305, "width": 0, "height": 40, "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": 442620899, "version": 1, "versionNonce": 806899910, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, 40 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0018", "type": "rectangle", "x": 440, "y": 350, "width": 170, "height": 60, "angle": 0, "strokeColor": "#6741d9", "backgroundColor": "#d0bfff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "index": null, "roundness": { "type": 3 }, "seed": 1599435268, "version": 1, "versionNonce": 418461139, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0019", "type": "text", "x": 462.3, "y": 365.0, "width": 125.4, "height": 30.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": 269676600, "version": 1, "versionNonce": 255985077, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "P as BF16, same\nregisters, A layout", "fontSize": 12, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "P as BF16, same\nregisters, A layout", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0020", "type": "arrow", "x": 615, "y": 380, "width": 60, "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": 884585952, "version": 1, "versionNonce": 2132084005, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 60, 0 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0021", "type": "rectangle", "x": 680, "y": 340, "width": 160, "height": 80, "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": 1836494975, "version": 1, "versionNonce": 1349251824, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0022", "type": "text", "x": 721.5, "y": 362.5, "width": 77.0, "height": 35.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": 1999744785, "version": 1, "versionNonce": 1946412081, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "WGMMA\nO += P V_j", "fontSize": 14, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "WGMMA\nO += P V_j", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0023", "type": "rectangle", "x": 680, "y": 230, "width": 160, "height": 70, "angle": 0, "strokeColor": "#f08c00", "backgroundColor": "#ffec99", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "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": 717.1, "y": 250.0, "width": 85.80000000000001, "height": 30.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": 1066984056, "version": 1, "versionNonce": 772092315, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "sV_j (d, 64)\nMN-major SMEM", "fontSize": 12, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "sV_j (d, 64)\nMN-major SMEM", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0025", "type": "arrow", "x": 760, "y": 305, "width": 0, "height": 30, "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": 1048386556, "version": 1, "versionNonce": 351564608, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 0, 30 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "el0026", "type": "text", "x": 635.975, "y": 425, "width": 248.05, "height": 41.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": 1289560150, "version": 1, "versionNonce": 2126508551, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "acc_pv: 64 x d / 128 FP32 per thread,\nrescaled by exp(m_old - m_new) each tile,\ndivided by l once at the end", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "acc_pv: 64 x d / 128 FP32 per thread,\nrescaled by exp(m_old - m_new) each tile,\ndivided by l once at the end", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0027", "type": "rectangle", "x": 20, "y": 475, "width": 820, "height": 125, "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": 1475216846, "version": 1, "versionNonce": 1927728187, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0028", "type": "text", "x": 202.29999999999998, "y": 481, "width": 455.40000000000003, "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": 1236683273, "version": 1, "versionNonce": 314395343, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "who holds what in a 64 x 64 accumulator (PTX wgmma D fragment layout)", "fontSize": 12, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "who holds what in a 64 x 64 accumulator (PTX wgmma D fragment layout)", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0029", "type": "rectangle", "x": 40, "y": 504, "width": 190, "height": 28, "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": 507088657, "version": 1, "versionNonce": 1795823849, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0030", "type": "text", "x": 80.55, "y": 511.125, "width": 108.9, "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": "warp 0: rows 0..15", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "warp 0: rows 0..15", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0031", "type": "rectangle", "x": 238, "y": 504, "width": 190, "height": 28, "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": 652768598, "version": 1, "versionNonce": 2100080515, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0032", "type": "text", "x": 275.525, "y": 511.125, "width": 114.95, "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": "warp 1: rows 16..31", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "warp 1: rows 16..31", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0033", "type": "rectangle", "x": 436, "y": 504, "width": 190, "height": 28, "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": 333377415, "version": 1, "versionNonce": 1347535309, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0034", "type": "text", "x": 473.525, "y": 511.125, "width": 114.95, "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": "warp 2: rows 32..47", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "warp 2: rows 32..47", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0035", "type": "rectangle", "x": 634, "y": 504, "width": 190, "height": 28, "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": 2133201996, "version": 1, "versionNonce": 1959386987, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "el0036", "type": "text", "x": 671.525, "y": 511.125, "width": 114.95, "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": 295334610, "version": 1, "versionNonce": 401991736, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "warp 3: rows 48..63", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "warp 3: rows 48..63", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0037", "type": "text", "x": 57.924999999999955, "y": 540, "width": 744.1500000000001, "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": 1159380354, "version": 1, "versionNonce": 2036236842, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "thread t of a warp holds rows 16w + t/4 and 16w + t/4 + 8, columns 8i + 2(t%4) and 8i + 2(t%4) + 1 for i = 0..7: 32 values.", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "thread t of a warp holds rows 16w + t/4 and 16w + t/4 + 8, columns 8i + 2(t%4) and 8i + 2(t%4) + 1 for i = 0..7: 32 values.", "autoResize": true, "lineHeight": 1.25 }, { "id": "el0038", "type": "text", "x": 91.19999999999999, "y": 560, "width": 677.6, "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": 279172787, "version": 1, "versionNonce": 260573196, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "the four lanes with equal t/4 share a row, so the row max and row sum are reduced across them with two shuffles.", "fontSize": 11, "fontFamily": 1, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "the four lanes with equal t/4 share a row, so the row max and row sum are reduced across them with two shuffles.", "autoResize": true, "lineHeight": 1.25 } ], "appState": { "gridSize": 20, "gridStep": 5, "gridModeEnabled": false, "viewBackgroundColor": "#fdf8f6" }, "files": {} }