{ "type": "excalidraw", "version": 2, "source": "https://excalidraw.com", "elements": [ { "id": "title", "type": "text", "x": 300, "y": 20, "width": 900, "height": 40, "angle": 0, "strokeColor": "#1e40af", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 1000, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "ldmatrix 指令与 CuTe TV Layout 完整映射图", "fontSize": 32, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "ldmatrix 指令与 CuTe TV Layout 完整映射图", "autoResize": false, "lineHeight": 1.25 }, { "id": "subtitle", "type": "text", "x": 200, "y": 68, "width": 1100, "height": 24, "angle": 0, "strokeColor": "#6b7280", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 1001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "核心：理解 (Thread, Value) → (Row, Col) 的映射，以及 ldmatrix 如何将 SMEM 数据直接放入 MMA 期望的寄存器格式", "fontSize": 16, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "核心：理解 (Thread, Value) → (Row, Col) 的映射，以及 ldmatrix 如何将 SMEM 数据直接放入 MMA 期望的寄存器格式", "autoResize": false, "lineHeight": 1.25 }, { "id": "section1_bg", "type": "rectangle", "x": 30, "y": 110, "width": 700, "height": 530, "angle": 0, "strokeColor": "#3b82f6", "backgroundColor": "#dbeafe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 30, "groupIds": [], "roundness": { "type": 3 }, "seed": 1100, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "section1_title", "type": "text", "x": 50, "y": 120, "width": 660, "height": 28, "angle": 0, "strokeColor": "#1e40af", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 1101, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Part 1: ldmatrix.m8n8.x1 线程→矩阵元素 映射", "fontSize": 22, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Part 1: ldmatrix.m8n8.x1 线程→矩阵元素 映射", "autoResize": false, "lineHeight": 1.25 }, { "id": "matrix_8x8_bg", "type": "rectangle", "x": 50, "y": 160, "width": 480, "height": 460, "angle": 0, "strokeColor": "#334155", "backgroundColor": "#f8fafc", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 1102, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "matrix_label", "type": "text", "x": 60, "y": 168, "width": 460, "height": 20, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 1103, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "8×8 矩阵（每格=1个fp16）— 颜色=线程组（4线程/行）", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "8×8 矩阵（每格=1个fp16）— 颜色=线程组（4线程/行）", "autoResize": false, "lineHeight": 1.25 }, { "id": "row0_c01", "type": "rectangle", "x": 100, "y": 200, "width": 100, "height": 44, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "#fecaca", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": { "type": 3 }, "seed": 2000, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row0_c01_text", "type": "text", "x": 108, "y": 204, "width": 90, "height": 36, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": null, "seed": 2001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T0\n(0,0)(0,1)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T0\n(0,0)(0,1)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row0_c23", "type": "rectangle", "x": 200, "y": 200, "width": 100, "height": 44, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "#fde8e8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": { "type": 3 }, "seed": 2002, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row0_c23_text", "type": "text", "x": 208, "y": 204, "width": 90, "height": 36, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": null, "seed": 2003, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T1\n(0,2)(0,3)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T1\n(0,2)(0,3)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row0_c45", "type": "rectangle", "x": 300, "y": 200, "width": 100, "height": 44, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "#fecaca", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": { "type": 3 }, "seed": 2004, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row0_c45_text", "type": "text", "x": 308, "y": 204, "width": 90, "height": 36, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": null, "seed": 2005, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T2\n(0,4)(0,5)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T2\n(0,4)(0,5)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row0_c67", "type": "rectangle", "x": 400, "y": 200, "width": 100, "height": 44, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "#fde8e8", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": { "type": 3 }, "seed": 2006, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row0_c67_text", "type": "text", "x": 408, "y": 204, "width": 90, "height": 36, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row0" ], "roundness": null, "seed": 2007, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T3\n(0,6)(0,7)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T3\n(0,6)(0,7)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row1_c01", "type": "rectangle", "x": 100, "y": 248, "width": 100, "height": 44, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "#fed7aa", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": { "type": 3 }, "seed": 2100, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row1_c01_text", "type": "text", "x": 108, "y": 252, "width": 90, "height": 36, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": null, "seed": 2101, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T4\n(1,0)(1,1)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T4\n(1,0)(1,1)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row1_c23", "type": "rectangle", "x": 200, "y": 248, "width": 100, "height": 44, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "#ffedd5", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": { "type": 3 }, "seed": 2102, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row1_c23_text", "type": "text", "x": 208, "y": 252, "width": 90, "height": 36, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": null, "seed": 2103, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T5\n(1,2)(1,3)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T5\n(1,2)(1,3)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row1_c45", "type": "rectangle", "x": 300, "y": 248, "width": 100, "height": 44, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "#fed7aa", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": { "type": 3 }, "seed": 2104, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row1_c45_text", "type": "text", "x": 308, "y": 252, "width": 90, "height": 36, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": null, "seed": 2105, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T6\n(1,4)(1,5)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T6\n(1,4)(1,5)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row1_c67", "type": "rectangle", "x": 400, "y": 248, "width": 100, "height": 44, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "#ffedd5", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": { "type": 3 }, "seed": 2106, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row1_c67_text", "type": "text", "x": 408, "y": 252, "width": 90, "height": 36, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row1" ], "roundness": null, "seed": 2107, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T7\n(1,6)(1,7)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T7\n(1,6)(1,7)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row2_c01", "type": "rectangle", "x": 100, "y": 296, "width": 100, "height": 44, "angle": 0, "strokeColor": "#ca8a04", "backgroundColor": "#fef9c3", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row2" ], "roundness": { "type": 3 }, "seed": 2200, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row2_text", "type": "text", "x": 108, "y": 300, "width": 90, "height": 36, "angle": 0, "strokeColor": "#ca8a04", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row2" ], "roundness": null, "seed": 2201, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T8\n(2,0)(2,1)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T8\n(2,0)(2,1)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row2_dots", "type": "text", "x": 230, "y": 306, "width": 270, "height": 28, "angle": 0, "strokeColor": "#6b7280", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 2202, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T9(2,2)(2,3) T10(2,4)(2,5) T11(2,6)(2,7)", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "T9(2,2)(2,3) T10(2,4)(2,5) T11(2,6)(2,7)", "autoResize": false, "lineHeight": 1.25 }, { "id": "dots_middle", "type": "text", "x": 230, "y": 348, "width": 200, "height": 20, "angle": 0, "strokeColor": "#9ca3af", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 2300, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "... 同理 Row3~Row6 ...", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "... 同理 Row3~Row6 ...", "autoResize": false, "lineHeight": 1.25 }, { "id": "row7_c01", "type": "rectangle", "x": 100, "y": 378, "width": 100, "height": 44, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#ede9fe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": { "type": 3 }, "seed": 2400, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row7_c01_text", "type": "text", "x": 105, "y": 382, "width": 90, "height": 36, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": null, "seed": 2401, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T28\n(7,0)(7,1)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T28\n(7,0)(7,1)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row7_c23", "type": "rectangle", "x": 200, "y": 378, "width": 100, "height": 44, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#f5f3ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": { "type": 3 }, "seed": 2402, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row7_c23_text", "type": "text", "x": 205, "y": 382, "width": 90, "height": 36, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": null, "seed": 2403, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T29\n(7,2)(7,3)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T29\n(7,2)(7,3)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row7_c45", "type": "rectangle", "x": 300, "y": 378, "width": 100, "height": 44, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#ede9fe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": { "type": 3 }, "seed": 2404, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row7_c45_text", "type": "text", "x": 305, "y": 382, "width": 90, "height": 36, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": null, "seed": 2405, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T30\n(7,4)(7,5)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T30\n(7,4)(7,5)", "autoResize": false, "lineHeight": 1.25 }, { "id": "row7_c67", "type": "rectangle", "x": 400, "y": 378, "width": 100, "height": 44, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#f5f3ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": { "type": 3 }, "seed": 2406, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "row7_c67_text", "type": "text", "x": 405, "y": 382, "width": 90, "height": 36, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "row7" ], "roundness": null, "seed": 2407, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "T31\n(7,6)(7,7)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "T31\n(7,6)(7,7)", "autoResize": false, "lineHeight": 1.25 }, { "id": "formula_box", "type": "rectangle", "x": 60, "y": 440, "width": 450, "height": 100, "angle": 0, "strokeColor": "#059669", "backgroundColor": "#d1fae5", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 2500, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "formula_text", "type": "text", "x": 75, "y": 448, "width": 430, "height": 84, "angle": 0, "strokeColor": "#065f46", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 2501, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "映射公式 (ldmatrix.m8n8.x1 非转置):\n row = thread_id / 4\n col_start = (thread_id % 4) * 2\n reg[u32] = {mat[row][col_start], mat[row][col_start+1]}", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "映射公式 (ldmatrix.m8n8.x1 非转置):\n row = thread_id / 4\n col_start = (thread_id % 4) * 2\n reg[u32] = {mat[row][col_start], mat[row][col_start+1]}", "autoResize": false, "lineHeight": 1.25 }, { "id": "addr_note", "type": "text", "x": 535, "y": 165, "width": 185, "height": 270, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 2600, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "SMEM 地址提供:\n\nT0 → Row0 addr\n ↳ 16B (8×fp16)\n\nT1 → Row1 addr\n ↳ 16B\n\n...\n\nT7 → Row7 addr\n ↳ 16B\n\nT8~T31 地址忽略\n（但必须参与 sync）\n\n硬件自动广播到\n4线程/行", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "SMEM 地址提供:\n\nT0 → Row0 addr\n ↳ 16B (8×fp16)\n\nT1 → Row1 addr\n ↳ 16B\n\n...\n\nT7 → Row7 addr\n ↳ 16B\n\nT8~T31 地址忽略\n（但必须参与 sync）\n\n硬件自动广播到\n4线程/行", "autoResize": false, "lineHeight": 1.25 }, { "id": "section2_bg", "type": "rectangle", "x": 770, "y": 110, "width": 780, "height": 530, "angle": 0, "strokeColor": "#9333ea", "backgroundColor": "#f3e8ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 25, "groupIds": [], "roundness": { "type": 3 }, "seed": 3000, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "section2_title", "type": "text", "x": 790, "y": 120, "width": 740, "height": 28, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 3001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Part 2: ldmatrix.x4 — 加载 4 个 8×8 块（覆盖 MMA 16×16 A 操作数）", "fontSize": 20, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Part 2: ldmatrix.x4 — 加载 4 个 8×8 块（覆盖 MMA 16×16 A 操作数）", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat0_bg", "type": "rectangle", "x": 800, "y": 166, "width": 160, "height": 160, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "#fecaca", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 60, "groupIds": [ "x4grid" ], "roundness": { "type": 3 }, "seed": 3100, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "mat0_label", "type": "text", "x": 830, "y": 175, "width": 100, "height": 48, "angle": 0, "strokeColor": "#991b1b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3101, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Matrix 0\n8×8 (M0-7,K0-7)\n→ reg d0", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Matrix 0\n8×8 (M0-7,K0-7)\n→ reg d0", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat0_addr", "type": "text", "x": 808, "y": 228, "width": 145, "height": 90, "angle": 0, "strokeColor": "#991b1b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3102, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "地址: T0~T7\nT0→Row0, T1→Row1\n...\nT7→Row7", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "地址: T0~T7\nT0→Row0, T1→Row1\n...\nT7→Row7", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat1_bg", "type": "rectangle", "x": 970, "y": 166, "width": 160, "height": 160, "angle": 0, "strokeColor": "#2563eb", "backgroundColor": "#bfdbfe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 60, "groupIds": [ "x4grid" ], "roundness": { "type": 3 }, "seed": 3200, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "mat1_label", "type": "text", "x": 995, "y": 175, "width": 110, "height": 48, "angle": 0, "strokeColor": "#1e3a5f", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3201, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Matrix 1\n8×8 (M0-7,K8-15)\n→ reg d1", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Matrix 1\n8×8 (M0-7,K8-15)\n→ reg d1", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat1_addr", "type": "text", "x": 978, "y": 228, "width": 145, "height": 90, "angle": 0, "strokeColor": "#1e3a5f", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3202, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "地址: T8~T15\nT8→Row0, T9→Row1\n...\nT15→Row7", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "地址: T8~T15\nT8→Row0, T9→Row1\n...\nT15→Row7", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat2_bg", "type": "rectangle", "x": 800, "y": 336, "width": 160, "height": 160, "angle": 0, "strokeColor": "#16a34a", "backgroundColor": "#bbf7d0", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 60, "groupIds": [ "x4grid" ], "roundness": { "type": 3 }, "seed": 3300, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "mat2_label", "type": "text", "x": 823, "y": 345, "width": 115, "height": 48, "angle": 0, "strokeColor": "#14532d", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3301, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Matrix 2\n8×8 (M8-15,K0-7)\n→ reg d2", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Matrix 2\n8×8 (M8-15,K0-7)\n→ reg d2", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat2_addr", "type": "text", "x": 808, "y": 400, "width": 145, "height": 90, "angle": 0, "strokeColor": "#14532d", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3302, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "地址: T16~T23\nT16→Row0,T17→Row1\n...\nT23→Row7", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "地址: T16~T23\nT16→Row0,T17→Row1\n...\nT23→Row7", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat3_bg", "type": "rectangle", "x": 970, "y": 336, "width": 160, "height": 160, "angle": 0, "strokeColor": "#d97706", "backgroundColor": "#fde68a", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 60, "groupIds": [ "x4grid" ], "roundness": { "type": 3 }, "seed": 3400, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "mat3_label", "type": "text", "x": 993, "y": 345, "width": 115, "height": 48, "angle": 0, "strokeColor": "#78350f", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3401, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Matrix 3\n8×8(M8-15,K8-15)\n→ reg d3", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Matrix 3\n8×8(M8-15,K8-15)\n→ reg d3", "autoResize": false, "lineHeight": 1.25 }, { "id": "mat3_addr", "type": "text", "x": 978, "y": 400, "width": 145, "height": 90, "angle": 0, "strokeColor": "#78350f", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "x4grid" ], "roundness": null, "seed": 3402, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "地址: T24~T31\nT24→Row0,T25→Row1\n...\nT31→Row7", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "地址: T24~T31\nT24→Row0,T25→Row1\n...\nT31→Row7", "autoResize": false, "lineHeight": 1.25 }, { "id": "x4_axis_m", "type": "text", "x": 777, "y": 240, "width": 20, "height": 120, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 3500, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "M\n‖\n1\n6\n行", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "M\n‖\n1\n6\n行", "autoResize": false, "lineHeight": 1.25 }, { "id": "x4_axis_k", "type": "text", "x": 880, "y": 148, "width": 160, "height": 18, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 3501, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "K = 16 列 →", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "K = 16 列 →", "autoResize": false, "lineHeight": 1.25 }, { "id": "x4_thread_example", "type": "rectangle", "x": 1160, "y": 166, "width": 370, "height": 340, "angle": 0, "strokeColor": "#1e40af", "backgroundColor": "#eff6ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 3600, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "x4_example_title", "type": "text", "x": 1175, "y": 174, "width": 340, "height": 18, "angle": 0, "strokeColor": "#1e40af", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 3601, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Thread 0 的 4 个寄存器 (ldmatrix.x4):", "fontSize": 15, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Thread 0 的 4 个寄存器 (ldmatrix.x4):", "autoResize": false, "lineHeight": 1.25 }, { "id": "x4_example_regs", "type": "text", "x": 1175, "y": 200, "width": 340, "height": 290, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 3602, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "d0 = {A[0][0], A[0][1]} ← Mat0\n Row 0, Col 0-1\n\nd1 = {A[0][8], A[0][9]} ← Mat1\n Row 0, Col 8-9\n\nd2 = {A[8][0], A[8][1]} ← Mat2\n Row 8, Col 0-1\n\nd3 = {A[8][8], A[8][9]} ← Mat3\n Row 8, Col 8-9\n\n─────────────────────────\nThread 5 (t=5):\n row=5/4=1, col_base=(5%4)*2=2\n d0={A[1][2],A[1][3]}\n d1={A[1][10],A[1][11]}\n d2={A[9][2],A[9][3]}\n d3={A[9][10],A[9][11]}", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "d0 = {A[0][0], A[0][1]} ← Mat0\n Row 0, Col 0-1\n\nd1 = {A[0][8], A[0][9]} ← Mat1\n Row 0, Col 8-9\n\nd2 = {A[8][0], A[8][1]} ← Mat2\n Row 8, Col 0-1\n\nd3 = {A[8][8], A[8][9]} ← Mat3\n Row 8, Col 8-9\n\n─────────────────────────\nThread 5 (t=5):\n row=5/4=1, col_base=(5%4)*2=2\n d0={A[1][2],A[1][3]}\n d1={A[1][10],A[1][11]}\n d2={A[9][2],A[9][3]}\n d3={A[9][10],A[9][11]}", "autoResize": false, "lineHeight": 1.25 }, { "id": "section3_bg", "type": "rectangle", "x": 30, "y": 680, "width": 1520, "height": 470, "angle": 0, "strokeColor": "#0891b2", "backgroundColor": "#cffafe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 25, "groupIds": [], "roundness": { "type": 3 }, "seed": 4000, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "section3_title", "type": "text", "x": 50, "y": 690, "width": 800, "height": 28, "angle": 0, "strokeColor": "#0e7490", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Part 3: CuTe TV Layout — (Thread, Value) → (Row, Col) 映射原理", "fontSize": 22, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Part 3: CuTe TV Layout — (Thread, Value) → (Row, Col) 映射原理", "autoResize": false, "lineHeight": 1.25 }, { "id": "tv_concept_box", "type": "rectangle", "x": 50, "y": 730, "width": 460, "height": 400, "angle": 0, "strokeColor": "#0e7490", "backgroundColor": "#ecfeff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 4100, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "tv_concept_title", "type": "text", "x": 65, "y": 738, "width": 420, "height": 20, "angle": 0, "strokeColor": "#0e7490", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4101, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "SM80_16x8x16 A操作数 TV Layout", "fontSize": 16, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "SM80_16x8x16 A操作数 TV Layout", "autoResize": false, "lineHeight": 1.25 }, { "id": "tv_code", "type": "text", "x": 60, "y": 766, "width": 440, "height": 355, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4102, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "ALayout: (T32, V8) → (M16, K16)\n\nLayout<\n Shape , Shape<_2,_2,_2>>,\n Stride,Stride<_16,_8,_128>>\n>\n\n解读:\n Thread = t_outer(4) * 8 + t_inner(8)\n t_inner stride=1 → K方向连续\n t_outer stride=32 → M方向跳2行\n\n Value = v2(2)*4 + v1(2)*2 + v0(2)\n v0 stride=16 → M方向+1行\n v1 stride=8 → K方向+8列(后半段)\n v2 stride=128 → M方向+8行(下半段)\n\nThread 0 持有 (16×16矩阵中):\n v0:(0,0) v1:(1,0) → d0寄存器\n v2:(0,8) v3:(1,8) → d1寄存器\n v4:(8,0) v5:(9,0) → d2寄存器\n v6:(8,8) v7:(9,8) → d3寄存器", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "ALayout: (T32, V8) → (M16, K16)\n\nLayout<\n Shape , Shape<_2,_2,_2>>,\n Stride,Stride<_16,_8,_128>>\n>\n\n解读:\n Thread = t_outer(4) * 8 + t_inner(8)\n t_inner stride=1 → K方向连续\n t_outer stride=32 → M方向跳2行\n\n Value = v2(2)*4 + v1(2)*2 + v0(2)\n v0 stride=16 → M方向+1行\n v1 stride=8 → K方向+8列(后半段)\n v2 stride=128 → M方向+8行(下半段)\n\nThread 0 持有 (16×16矩阵中):\n v0:(0,0) v1:(1,0) → d0寄存器\n v2:(0,8) v3:(1,8) → d1寄存器\n v4:(8,0) v5:(9,0) → d2寄存器\n v6:(8,8) v7:(9,8) → d3寄存器", "autoResize": false, "lineHeight": 1.25 }, { "id": "tv_visual_box", "type": "rectangle", "x": 530, "y": 730, "width": 500, "height": 400, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#faf5ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 4200, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "tv_visual_title", "type": "text", "x": 545, "y": 738, "width": 470, "height": 20, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4201, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Thread 0 在 A[16×16] 中的元素位置（用●标记）", "fontSize": 15, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Thread 0 在 A[16×16] 中的元素位置（用●标记）", "autoResize": false, "lineHeight": 1.25 }, { "id": "tv_visual_grid", "type": "text", "x": 545, "y": 766, "width": 470, "height": 350, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4202, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": " K→ 0 1 2 3 4 5 6 7 │ 8 9 10 11 12 13 14 15\n M ─────────────────────────────┼──────────────────────\n 0 │ ●v0 . . . . . . . │ ●v2 . . . . . . .\n 1 │ ●v1 . . . . . . . │ ●v3 . . . . . . .\n 2 │ . . . . . . . . │ . . . . . . . .\n 3 │ . . . . . . . . │ . . . . . . . .\n 4 │ . . . . . . . . │ . . . . . . . .\n 5 │ . . . . . . . . │ . . . . . . . .\n 6 │ . . . . . . . . │ . . . . . . . .\n 7 │ . . . . . . . . │ . . . . . . . .\n ──┼─────────────────────────────┼──────────────────────\n 8 │ ●v4 . . . . . . . │ ●v6 . . . . . . .\n 9 │ ●v5 . . . . . . . │ ●v7 . . . . . . .\n 10 │ . . . . . . . . │ . . . . . . . .\n 11 │ . . . . . . . . │ . . . . . . . .\n ...\n\n ● = Thread 0 持有的元素\n v0,v1 打包在 d0 寄存器 (u32)\n v2,v3 打包在 d1 寄存器 (u32)\n v4,v5 打包在 d2 寄存器 (u32)\n v6,v7 打包在 d3 寄存器 (u32)", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": " K→ 0 1 2 3 4 5 6 7 │ 8 9 10 11 12 13 14 15\n M ─────────────────────────────┼──────────────────────\n 0 │ ●v0 . . . . . . . │ ●v2 . . . . . . .\n 1 │ ●v1 . . . . . . . │ ●v3 . . . . . . .\n 2 │ . . . . . . . . │ . . . . . . . .\n 3 │ . . . . . . . . │ . . . . . . . .\n 4 │ . . . . . . . . │ . . . . . . . .\n 5 │ . . . . . . . . │ . . . . . . . .\n 6 │ . . . . . . . . │ . . . . . . . .\n 7 │ . . . . . . . . │ . . . . . . . .\n ──┼─────────────────────────────┼──────────────────────\n 8 │ ●v4 . . . . . . . │ ●v6 . . . . . . .\n 9 │ ●v5 . . . . . . . │ ●v7 . . . . . . .\n 10 │ . . . . . . . . │ . . . . . . . .\n 11 │ . . . . . . . . │ . . . . . . . .\n ...\n\n ● = Thread 0 持有的元素\n v0,v1 打包在 d0 寄存器 (u32)\n v2,v3 打包在 d1 寄存器 (u32)\n v4,v5 打包在 d2 寄存器 (u32)\n v6,v7 打包在 d3 寄存器 (u32)", "autoResize": false, "lineHeight": 1.25 }, { "id": "pipeline_box", "type": "rectangle", "x": 1050, "y": 730, "width": 480, "height": 400, "angle": 0, "strokeColor": "#059669", "backgroundColor": "#ecfdf5", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 4300, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "pipeline_title", "type": "text", "x": 1065, "y": 738, "width": 450, "height": 20, "angle": 0, "strokeColor": "#059669", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4301, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "关键洞察：ldmatrix DstLayout ≈ MMA ALayout", "fontSize": 16, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "关键洞察：ldmatrix DstLayout ≈ MMA ALayout", "autoResize": false, "lineHeight": 1.25 }, { "id": "pipeline_content", "type": "text", "x": 1065, "y": 768, "width": 450, "height": 350, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 4302, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "ldmatrix.trans.x4 的 DstLayout:\n Shape , Shape<_16,_2,_2>>\n Stride,Stride<_1,_128,_1024>>\n\nMMA SM80 的 ALayout:\n Shape , Shape<_2,_2,_2>>\n Stride,Stride<_16,_8,_128>>\n\n二者的 Thread 维度分组方式完全一致:\n ► 4组 × 8线程 = 32线程\n ► 内层8线程对应 K 方向\n ► 外层4组对应 M 方向\n\n这就是为什么 ldmatrix 的输出可以\n直接作为 mma.sync 的输入！\n\n─────────────────────────────\nCuTe 代码:\n make_tiled_copy_A(\n Copy_Atom{},\n tiled_mma\n );\n→ 自动组合 ldmatrix DstLayout\n 与 MMA ALayout，保证兼容性", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "ldmatrix.trans.x4 的 DstLayout:\n Shape , Shape<_16,_2,_2>>\n Stride,Stride<_1,_128,_1024>>\n\nMMA SM80 的 ALayout:\n Shape , Shape<_2,_2,_2>>\n Stride,Stride<_16,_8,_128>>\n\n二者的 Thread 维度分组方式完全一致:\n ► 4组 × 8线程 = 32线程\n ► 内层8线程对应 K 方向\n ► 外层4组对应 M 方向\n\n这就是为什么 ldmatrix 的输出可以\n直接作为 mma.sync 的输入！\n\n─────────────────────────────\nCuTe 代码:\n make_tiled_copy_A(\n Copy_Atom{},\n tiled_mma\n );\n→ 自动组合 ldmatrix DstLayout\n 与 MMA ALayout，保证兼容性", "autoResize": false, "lineHeight": 1.25 }, { "id": "section4_bg", "type": "rectangle", "x": 30, "y": 1190, "width": 1520, "height": 360, "angle": 0, "strokeColor": "#be185d", "backgroundColor": "#fce7f3", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 25, "groupIds": [], "roundness": { "type": 3 }, "seed": 5000, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "section4_title", "type": "text", "x": 50, "y": 1200, "width": 700, "height": 28, "angle": 0, "strokeColor": "#be185d", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 5001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Part 4: 全流程 — GMEM → SMEM → Register → MMA", "fontSize": 22, "fontFamily": 5, "textAlign": "left", "verticalAlign": "middle", "containerId": null, "originalText": "Part 4: 全流程 — GMEM → SMEM → Register → MMA", "autoResize": false, "lineHeight": 1.25 }, { "id": "stage1_box", "type": "rectangle", "x": 60, "y": 1250, "width": 280, "height": 130, "angle": 0, "strokeColor": "#4f46e5", "backgroundColor": "#eef2ff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": { "type": 3 }, "seed": 5100, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "stage1_text", "type": "text", "x": 75, "y": 1258, "width": 250, "height": 115, "angle": 0, "strokeColor": "#3730a3", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5101, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Global Memory\nA[M×K], B[N×K]\n\nCopy Atom:\n SM80_CP_ASYNC_CACHEGLOBAL\n ", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Global Memory\nA[M×K], B[N×K]\n\nCopy Atom:\n SM80_CP_ASYNC_CACHEGLOBAL\n ", "autoResize": false, "lineHeight": 1.25 }, { "id": "arrow1", "type": "arrow", "x": 345, "y": 1315, "width": 80, "height": 0, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 3, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5200, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 80, 0 ] ], "lastCommittedPoint": [ 80, 0 ], "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow" }, { "id": "arrow1_label", "type": "text", "x": 348, "y": 1290, "width": 70, "height": 18, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5201, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "cp.async", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "cp.async", "autoResize": false, "lineHeight": 1.25 }, { "id": "stage2_box", "type": "rectangle", "x": 430, "y": 1250, "width": 340, "height": 130, "angle": 0, "strokeColor": "#0891b2", "backgroundColor": "#ecfeff", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": { "type": 3 }, "seed": 5300, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "stage2_text", "type": "text", "x": 445, "y": 1258, "width": 310, "height": 115, "angle": 0, "strokeColor": "#155e75", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5301, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Shared Memory (+ Swizzle)\nSmemLayoutAtom + Swizzle<3,3,3>\n\n消除 Bank Conflict:\nXOR 高位→低位，8行映射到8个bank\n无冲突 = 最大带宽", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Shared Memory (+ Swizzle)\nSmemLayoutAtom + Swizzle<3,3,3>\n\n消除 Bank Conflict:\nXOR 高位→低位，8行映射到8个bank\n无冲突 = 最大带宽", "autoResize": false, "lineHeight": 1.25 }, { "id": "arrow2", "type": "arrow", "x": 775, "y": 1315, "width": 80, "height": 0, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 3, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5400, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 80, 0 ] ], "lastCommittedPoint": [ 80, 0 ], "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow" }, { "id": "arrow2_label", "type": "text", "x": 770, "y": 1282, "width": 90, "height": 28, "angle": 0, "strokeColor": "#dc2626", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5401, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "ldmatrix\n(本图核心)", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "ldmatrix\n(本图核心)", "autoResize": false, "lineHeight": 1.25 }, { "id": "stage3_box", "type": "rectangle", "x": 860, "y": 1250, "width": 300, "height": 130, "angle": 0, "strokeColor": "#ea580c", "backgroundColor": "#fff7ed", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": { "type": 3 }, "seed": 5500, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "stage3_text", "type": "text", "x": 875, "y": 1258, "width": 270, "height": 115, "angle": 0, "strokeColor": "#9a3412", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5501, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Register File\n(MMA TV Layout 格式)\n\nA: 4 u32 regs/thread (8 fp16)\nB: 2 u32 regs/thread (4 fp16)\n已经是 mma.sync 期望的格式!", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Register File\n(MMA TV Layout 格式)\n\nA: 4 u32 regs/thread (8 fp16)\nB: 2 u32 regs/thread (4 fp16)\n已经是 mma.sync 期望的格式!", "autoResize": false, "lineHeight": 1.25 }, { "id": "arrow3", "type": "arrow", "x": 1165, "y": 1315, "width": 80, "height": 0, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 3, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5600, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "points": [ [ 0, 0 ], [ 80, 0 ] ], "lastCommittedPoint": [ 80, 0 ], "startBinding": null, "endBinding": null, "startArrowhead": null, "endArrowhead": "arrow" }, { "id": "arrow3_label", "type": "text", "x": 1168, "y": 1290, "width": 75, "height": 18, "angle": 0, "strokeColor": "#374151", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5601, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "mma.sync", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "middle", "containerId": null, "originalText": "mma.sync", "autoResize": false, "lineHeight": 1.25 }, { "id": "stage4_box", "type": "rectangle", "x": 1250, "y": 1250, "width": 280, "height": 130, "angle": 0, "strokeColor": "#059669", "backgroundColor": "#d1fae5", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": { "type": 3 }, "seed": 5700, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "stage4_text", "type": "text", "x": 1265, "y": 1258, "width": 250, "height": 115, "angle": 0, "strokeColor": "#065f46", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [ "pipeline" ], "roundness": null, "seed": 5701, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "Accumulator\nC/D: 4 f32 regs/thread\n\nCLayout: SM80_16x8_Row\n(T32,V4) → (M16,N8)\n\nmma.sync.m16n8k16", "fontSize": 14, "fontFamily": 5, "textAlign": "center", "verticalAlign": "top", "containerId": null, "originalText": "Accumulator\nC/D: 4 f32 regs/thread\n\nCLayout: SM80_16x8_Row\n(T32,V4) → (M16,N8)\n\nmma.sync.m16n8k16", "autoResize": false, "lineHeight": 1.25 }, { "id": "cute_code_box", "type": "rectangle", "x": 60, "y": 1400, "width": 1470, "height": 130, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "#f1f5f9", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": { "type": 3 }, "seed": 5800, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false }, { "id": "cute_code_text", "type": "text", "x": 75, "y": 1408, "width": 1440, "height": 115, "angle": 0, "strokeColor": "#1e293b", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 0, "opacity": 100, "groupIds": [], "roundness": null, "seed": 5801, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "对应 CuTe 代码:\n using TiledMMA = TiledMMA>;\n auto copy_A = make_tiled_copy_A(Copy_Atom{}, tiled_mma); // A: ldmatrix.trans (行主序SMEM → 列消费)\n auto copy_B = make_tiled_copy_B(Copy_Atom{}, tiled_mma); // B: ldmatrix.normal (行主序SMEM → 行消费)\n cute::copy(copy_A, thr_copy.partition_S(sA), thr_copy.partition_D(rA)); // 底层发射 ldmatrix PTX\n cute::gemm(tiled_mma, rA, rB, rC); // 底层发射 mma.sync PTX", "fontSize": 14, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "对应 CuTe 代码:\n using TiledMMA = TiledMMA>;\n auto copy_A = make_tiled_copy_A(Copy_Atom{}, tiled_mma); // A: ldmatrix.trans (行主序SMEM → 列消费)\n auto copy_B = make_tiled_copy_B(Copy_Atom{}, tiled_mma); // B: ldmatrix.normal (行主序SMEM → 行消费)\n cute::copy(copy_A, thr_copy.partition_S(sA), thr_copy.partition_D(rA)); // 底层发射 ldmatrix PTX\n cute::gemm(tiled_mma, rA, rB, rC); // 底层发射 mma.sync PTX", "autoResize": false, "lineHeight": 1.25 }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-bg", "type": "rectangle", "x": 80, "y": 1630.0, "width": 1470.0, "height": 300, "angle": 0, "strokeColor": "#64748b", "backgroundColor": "#e2e8f0", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 70, "groupIds": [], "seed": 93001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "roundness": { "type": 3 } }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-title", "type": "text", "x": 108, "y": 1652.0, "width": 1414.0, "height": 38, "angle": 0, "strokeColor": "#1e40af", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "seed": 91001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "练习题与面试题补充：ldmatrix / tensor core layout", "fontSize": 24, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "练习题与面试题补充：ldmatrix / tensor core layout", "autoResize": false, "lineHeight": 1.25 }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-practice-box", "type": "rectangle", "x": 108, "y": 1708.0, "width": 693.0, "height": 190, "angle": 0, "strokeColor": "#2563eb", "backgroundColor": "#bfdbfe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "seed": 93002, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "roundness": { "type": 3 } }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-practice-text", "type": "text", "x": 130, "y": 1728.0, "width": 670.0, "height": 150, "angle": 0, "strokeColor": "#243042", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "seed": 91001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "练习题\n1. 画出 8x8 tile 从 shared memory 到 register 的 lane 分布\n2. 解释 transpose variant 对矩阵布局的影响\n3. 对照 mma 指令需要的 fragment layout 做一次手算", "fontSize": 17, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "练习题\n1. 画出 8x8 tile 从 shared memory 到 register 的 lane 分布\n2. 解释 transpose variant 对矩阵布局的影响\n3. 对照 mma 指令需要的 fragment layout 做一次手算", "autoResize": false, "lineHeight": 1.25 }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-interview-box", "type": "rectangle", "x": 829.0, "y": 1708.0, "width": 693.0, "height": 190, "angle": 0, "strokeColor": "#7c3aed", "backgroundColor": "#ddd6fe", "fillStyle": "solid", "strokeWidth": 2, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "seed": 93003, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "roundness": { "type": 3 } }, { "id": "practice-interview-panel-ldmatrix-tv-layout-详解-interview-text", "type": "text", "x": 851.0, "y": 1728.0, "width": 670.0, "height": 150, "angle": 0, "strokeColor": "#243042", "backgroundColor": "transparent", "fillStyle": "solid", "strokeWidth": 1, "strokeStyle": "solid", "roughness": 1, "opacity": 100, "groupIds": [], "seed": 91001, "version": 1, "isDeleted": false, "boundElements": null, "updated": 1, "link": null, "locked": false, "text": "面试题\n1. ldmatrix 解决了 tensor core 前的数据装载什么问题\n2. shared memory bank conflict 如何影响 ldmatrix\n3. TV layout 和 MMA fragment layout 如何对应", "fontSize": 17, "fontFamily": 5, "textAlign": "left", "verticalAlign": "top", "containerId": null, "originalText": "面试题\n1. ldmatrix 解决了 tensor core 前的数据装载什么问题\n2. shared memory bank conflict 如何影响 ldmatrix\n3. TV layout 和 MMA fragment layout 如何对应", "autoResize": false, "lineHeight": 1.25 } ], "appState": { "gridSize": null, "viewBackgroundColor": "#ffffff" }, "files": {} }