{ "type": "excalidraw", "version": 2, "source": "gen_five_levels.py", "elements": [ { "id": "e0001", "type": "text", "x": 20, "y": 14, "width": 660.0, "height": 25.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 7919, "version": 1, "versionNonce": 104729, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "The five levels of CUTLASS 3.x, nested as the templates nest", "fontSize": 20, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "The five levels of CUTLASS 3.x, nested as the templates nest", "lineHeight": 1.25, "autoResize": true, "baseline": 20 }, { "id": "e0002", "type": "rectangle", "x": 20, "y": 60, "width": 600, "height": 470, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#f1f3f5", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 15838, "version": 1, "versionNonce": 209458, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0003", "type": "text", "x": 34, "y": 70, "width": 361.90000000000003, "height": 17.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 23757, "version": 1, "versionNonce": 314187, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Device (host): GemmUniversalAdapter", "fontSize": 14, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Device (host): GemmUniversalAdapter", "lineHeight": 1.25, "autoResize": true, "baseline": 14 }, { "id": "e0004", "type": "text", "x": 34, "y": 90, "width": 453.75000000000006, "height": 13.75, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 31676, "version": 1, "versionNonce": 418916, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "gemm/device/gemm_universal_adapter.h. Arguments to Params, then the launch.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "gemm/device/gemm_universal_adapter.h. Arguments to Params, then the launch.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0005", "type": "rectangle", "x": 40, "y": 116, "width": 560, "height": 398, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 39595, "version": 1, "versionNonce": 523645, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0006", "type": "text", "x": 54, "y": 126, "width": 469.70000000000005, "height": 17.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 47514, "version": 1, "versionNonce": 628374, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Kernel (GPU): GemmUniversal", "fontSize": 14, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Kernel (GPU): GemmUniversal", "lineHeight": 1.25, "autoResize": true, "baseline": 14 }, { "id": "e0007", "type": "text", "x": 54, "y": 146, "width": 344.85, "height": 13.75, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 55433, "version": 1, "versionNonce": 733103, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0008", "type": "text", "x": 54, "y": 162, "width": 326.70000000000005, "height": 13.75, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 63352, "version": 1, "versionNonce": 837832, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Warp roles, and the persistent loop over output tiles.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Warp roles, and the persistent loop over output tiles.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0009", "type": "rectangle", "x": 56, "y": 190, "width": 272, "height": 308, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#d0bfff", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 71271, "version": 1, "versionNonce": 942561, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0010", "type": "text", "x": 68, "y": 200, "width": 154.0, "height": 17.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 79190, "version": 1, "versionNonce": 1047290, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Collective: mainloop", "fontSize": 14, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Collective: mainloop", "lineHeight": 1.25, "autoResize": true, "baseline": 14 }, { "id": "e0011", "type": "text", "x": 68, "y": 220, "width": 205.70000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 87109, "version": 1, "versionNonce": 1152019, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "CollectiveMma", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "CollectiveMma", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0012", "type": "text", "x": 68, "y": 236, "width": 209.00000000000003, "height": 11.875, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 95028, "version": 1, "versionNonce": 1256748, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "sm90_mma_tma_gmma_ss_warpspecialized.hpp", "fontSize": 9.5, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "sm90_mma_tma_gmma_ss_warpspecialized.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 9.5 }, { "id": "e0013", "type": "text", "x": 68, "y": 252, "width": 223.85000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 102947, "version": 1, "versionNonce": 1361477, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "The loop over K tiles (not unrolled).", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "The loop over K tiles (not unrolled).", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0014", "type": "rectangle", "x": 70, "y": 282, "width": 244, "height": 202, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#ffec99", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 110866, "version": 1, "versionNonce": 1466206, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0015", "type": "text", "x": 82, "y": 292, "width": 128.70000000000002, "height": 16.25, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 118785, "version": 1, "versionNonce": 1570935, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Tiled MMA and copy", "fontSize": 13, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Tiled MMA and copy", "lineHeight": 1.25, "autoResize": true, "baseline": 13 }, { "id": "e0016", "type": "text", "x": 82, "y": 311, "width": 181.50000000000003, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 126704, "version": 1, "versionNonce": 1675664, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "TiledMMA, TiledCopy (TMA load)", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "TiledMMA, TiledCopy (TMA load)", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0017", "type": "text", "x": 82, "y": 327, "width": 133.10000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 134623, "version": 1, "versionNonce": 1780393, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "cute::gemm, cute::copy", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "cute::gemm, cute::copy", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0018", "type": "text", "x": 82, "y": 343, "width": 242.00000000000003, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 142542, "version": 1, "versionNonce": 1885122, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Loops inside one call: static, unrolled.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Loops inside one call: static, unrolled.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0019", "type": "rectangle", "x": 84, "y": 374, "width": 216, "height": 96, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 150461, "version": 1, "versionNonce": 1989851, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0020", "type": "text", "x": 96, "y": 384, "width": 28.6, "height": 16.25, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 158380, "version": 1, "versionNonce": 2094580, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Atom", "fontSize": 13, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Atom", "lineHeight": 1.25, "autoResize": true, "baseline": 13 }, { "id": "e0021", "type": "text", "x": 96, "y": 404, "width": 114.95, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 166299, "version": 1, "versionNonce": 2199309, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "MMA_Atom: one wgmma", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "MMA_Atom: one wgmma", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0022", "type": "text", "x": 96, "y": 420, "width": 139.15, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 174218, "version": 1, "versionNonce": 2304038, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Copy_Atom: one TMA load", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Copy_Atom: one TMA load", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0023", "type": "text", "x": 96, "y": 438, "width": 121.00000000000001, "height": 12.5, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 182137, "version": 1, "versionNonce": 2408767, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "cute/atom/mma_atom.hpp", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "cute/atom/mma_atom.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0024", "type": "text", "x": 96, "y": 452, "width": 126.50000000000001, "height": 12.5, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 190056, "version": 1, "versionNonce": 2513496, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "cute/atom/copy_atom.hpp", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "cute/atom/copy_atom.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0025", "type": "rectangle", "x": 340, "y": 190, "width": 246, "height": 308, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#d0bfff", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 197975, "version": 1, "versionNonce": 2618225, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0026", "type": "text", "x": 352, "y": 200, "width": 154.0, "height": 17.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 205894, "version": 1, "versionNonce": 2722954, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Collective: epilogue", "fontSize": 14, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Collective: epilogue", "lineHeight": 1.25, "autoResize": true, "baseline": 14 }, { "id": "e0027", "type": "text", "x": 352, "y": 220, "width": 114.95, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 213813, "version": 1, "versionNonce": 2827683, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "CollectiveEpilogue<", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "CollectiveEpilogue<", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0028", "type": "text", "x": 352, "y": 236, "width": 211.75000000000003, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 221732, "version": 1, "versionNonce": 2932412, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": " Sm90TmaWarpSpecialized<...>, ...>", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": " Sm90TmaWarpSpecialized<...>, ...>", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0029", "type": "text", "x": 352, "y": 252, "width": 193.32500000000002, "height": 11.875, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 229651, "version": 1, "versionNonce": 3037141, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "sm90_epilogue_tma_warpspecialized.hpp", "fontSize": 9.5, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "sm90_epilogue_tma_warpspecialized.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 9.5 }, { "id": "e0030", "type": "text", "x": 352, "y": 268, "width": 139.15, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 237570, "version": 1, "versionNonce": 3141870, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Writes one output tile.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Writes one output tile.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0031", "type": "rectangle", "x": 354, "y": 298, "width": 218, "height": 186, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#ffec99", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 245489, "version": 1, "versionNonce": 3246599, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0032", "type": "text", "x": 366, "y": 308, "width": 71.5, "height": 16.25, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 253408, "version": 1, "versionNonce": 3351328, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Tiled copy", "fontSize": 13, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Tiled copy", "lineHeight": 1.25, "autoResize": true, "baseline": 13 }, { "id": "e0033", "type": "text", "x": 366, "y": 327, "width": 127.05000000000001, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 261327, "version": 1, "versionNonce": 3456057, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "TiledCopy (TMA store)", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "TiledCopy (TMA store)", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0034", "type": "rectangle", "x": 368, "y": 374, "width": 190, "height": 96, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#b2f2bb", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 269246, "version": 1, "versionNonce": 3560786, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0035", "type": "text", "x": 380, "y": 384, "width": 28.6, "height": 16.25, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 277165, "version": 1, "versionNonce": 3665515, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Atom", "fontSize": 13, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Atom", "lineHeight": 1.25, "autoResize": true, "baseline": 13 }, { "id": "e0036", "type": "text", "x": 380, "y": 404, "width": 145.20000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 285084, "version": 1, "versionNonce": 3770244, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Copy_Atom: one TMA store", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Copy_Atom: one TMA store", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0037", "type": "text", "x": 380, "y": 438, "width": 126.50000000000001, "height": 12.5, "angle": 0, "strokeColor": "#868e96", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 293003, "version": 1, "versionNonce": 3874973, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "cute/atom/copy_atom.hpp", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "cute/atom/copy_atom.hpp", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0038", "type": "text", "x": 660, "y": 66, "width": 184.8, "height": 17.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 300922, "version": 1, "versionNonce": 3979702, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "How the kernel is picked", "fontSize": 14, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "How the kernel is picked", "lineHeight": 1.25, "autoResize": true, "baseline": 14 }, { "id": "e0039", "type": "rectangle", "x": 660, "y": 96, "width": 290, "height": 70, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#f1f3f5", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 308841, "version": 1, "versionNonce": 4084431, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0040", "type": "text", "x": 672, "y": 104, "width": 138.60000000000002, "height": 15.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 316760, "version": 1, "versionNonce": 4189160, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Dense mainloop policy", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Dense mainloop policy", "lineHeight": 1.25, "autoResize": true, "baseline": 12 }, { "id": "e0041", "type": "text", "x": 672, "y": 122, "width": 114.95, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 324679, "version": 1, "versionNonce": 4293889, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "MainloopSm90TmaGmma", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "MainloopSm90TmaGmma", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0042", "type": "text", "x": 672, "y": 138, "width": 242.00000000000003, "height": 12.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 332598, "version": 1, "versionNonce": 4398618, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": " WarpSpecialized", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": " WarpSpecialized", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0043", "type": "rectangle", "x": 660, "y": 190, "width": 290, "height": 70, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#ffc9c9", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 340517, "version": 1, "versionNonce": 4503347, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0044", "type": "text", "x": 672, "y": 198, "width": 132.0, "height": 15.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 348436, "version": 1, "versionNonce": 4608076, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "W4A8 mainloop policy", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "W4A8 mainloop policy", "lineHeight": 1.25, "autoResize": true, "baseline": 12 }, { "id": "e0045", "type": "text", "x": 672, "y": 216, "width": 145.20000000000002, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 356355, "version": 1, "versionNonce": 4712805, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "MainloopSm90TmaGmmaRmemA", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "MainloopSm90TmaGmmaRmemA", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0046", "type": "text", "x": 672, "y": 232, "width": 176.0, "height": 12.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 364274, "version": 1, "versionNonce": 4817534, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": " WarpSpecializedMixedInput<...>", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": " WarpSpecializedMixedInput<...>", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0047", "type": "arrow", "x": 950, "y": 131, "width": 18, "height": 191, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 372193, "version": 1, "versionNonce": 4922263, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 18, 0 ], [ 18, 191 ], [ 2, 191 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "e0048", "type": "arrow", "x": 805, "y": 260, "width": 0, "height": 40, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 380112, "version": 1, "versionNonce": 5026992, "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": "e0049", "type": "rectangle", "x": 660, "y": 300, "width": 290, "height": 58, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "#a5d8ff", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 388031, "version": 1, "versionNonce": 5131721, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "e0050", "type": "text", "x": 672, "y": 308, "width": 145.20000000000002, "height": 15.0, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 395950, "version": 1, "versionNonce": 5236450, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Both have ::Schedule =", "fontSize": 12, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "Both have ::Schedule =", "lineHeight": 1.25, "autoResize": true, "baseline": 12 }, { "id": "e0051", "type": "text", "x": 672, "y": 326, "width": 211.75000000000003, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 403869, "version": 1, "versionNonce": 5341179, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "KernelTmaWarpSpecializedCooperative", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "KernelTmaWarpSpecializedCooperative", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0052", "type": "arrow", "x": 660, "y": 329, "width": 58, "height": 195, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 411788, "version": 1, "versionNonce": 5445908, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ -30, 0 ], [ -30, -195 ], [ -58, -195 ] ], "lastCommittedPoint": null, "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow", "elbowed": false }, { "id": "e0053", "type": "text", "x": 660, "y": 372, "width": 217.8, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 419707, "version": 1, "versionNonce": 5550637, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "GemmUniversal's specialization tests", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "GemmUniversal's specialization tests", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0054", "type": "text", "x": 660, "y": 388, "width": 258.5, "height": 12.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 427626, "version": 1, "versionNonce": 5655366, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "is_base_of", "fontSize": 10, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": " Mainloop::DispatchPolicy::Schedule>", "lineHeight": 1.25, "autoResize": true, "baseline": 10 }, { "id": "e0056", "type": "text", "x": 660, "y": 424, "width": 242.00000000000003, "height": 13.75, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 443464, "version": 1, "versionNonce": 5864824, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "So both get the same cooperative kernel.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "So both get the same cooperative kernel.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 }, { "id": "e0057", "type": "text", "x": 660, "y": 440, "width": 199.65, "height": 27.5, "angle": 0, "strokeColor": "#1e1e1e", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1.5, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "frameId": null, "roundness": null, "seed": 451383, "version": 1, "versionNonce": 5969553, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "The epilogue class is shared too.\nThe mainloop collective differs.", "fontSize": 11, "fontFamily": 1, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "The epilogue class is shared too.\nThe mainloop collective differs.", "lineHeight": 1.25, "autoResize": true, "baseline": 11 } ], "appState": { "gridSize": null, "viewBackgroundColor": "#fdf8f6" }, "files": {} }