Star1085
540
代码
代码
Issues243
Pull Requests185
Actions
项目讨论
Wiki
分析
Star1085
540
已合并
issue问题修复 #685
jiangxiuhan创建于 8月13日
issue问题修复 #685
已合并
jiangxiuhan创建于 8月13日
从jiangxiuhan1/cann-learning-hub_6676test合入到cann/cann-learning-hubtest
共 23 个文件变更+57-40
Mtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.01_chapter_intro.ipynb+2-2
@@ -95,7 +95,7 @@
95 " </tr>\n",95 " </tr>\n",
96 " <tr>\n",96 " <tr>\n",
97 " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\">02.10</td>\n",97 " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\">02.10</td>\n",
98- " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\"><a href=\"./02.10_quant_dispatch_practice.ipynb\">量化 Dispatch 实践</a></td>\n",98+ " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\"><a href=\"./02.10_chapter_test.ipynb\">章节测试</a></td>\n",
99 " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\">通过完整实战项目，动手实现支持 MXFP8 动态量化的高性能 Dispatch 算子</td>\n",99 " <td style=\"border:1px solid #ccc; padding:8px; text-align:left;\">通过完整实战项目，动手实现支持 MXFP8 动态量化的高性能 Dispatch 算子</td>\n",
100 " </tr>\n",100 " </tr>\n",
101 "</table>\n",101 "</table>\n",
@@ -107,7 +107,7 @@
107 "\n",107 "\n",
108 "## 实践项目\n",108 "## 实践项目\n",
109 "\n",109 "\n",
110- "本教程提供 **量化 Dispatch 算子实践项目**，位于 [`scaffold_02_10/`](./scaffold_02_10/) 目录：\n",110+ "本教程提供 **量化 Dispatch 算子实践项目**，位于 [`src_02_10/`](./src_02_10/) 目录：\n",
111 "\n",111 "\n",
112 "- **项目目标**：实现支持 MXFP8 动态量化的 Dispatch 算子，完成 Token 分发、跨卡通信与数据整理\n",112 "- **项目目标**：实现支持 MXFP8 动态量化的 Dispatch 算子，完成 Token 分发、跨卡通信与数据整理\n",
113 "- **技术要点**：量化计算、共享内存通信、流水并行、多核同步\n",113 "- **技术要点**：量化计算、共享内存通信、流水并行、多核同步\n",
Mtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.07_win_memory_layout.ipynb+4-4
@@ -11,7 +11,7 @@
11 "\n",11 "\n",
12 "**学习目标**：掌握 Win 区中「Win数据区 / Win状态区」「ping-pong 翻转」「Dispatch / Combine 子区」三层划分；理解 token 在 Win 上为何切成 480 + 32 字节的小块。\n",12 "**学习目标**：掌握 Win 区中「Win数据区 / Win状态区」「ping-pong 翻转」「Dispatch / Combine 子区」三层划分；理解 token 在 Win 上为何切成 480 + 32 字节的小块。\n",
13 "\n",13 "\n",
14- "**关键常量速查表**（来自实践代码 [moe_distribute_comm.h](./answer/scaffold_02_10/include/moe_distribute_comm.h) / [moe_distribute_dispatch_non_quant.h](./answer/scaffold_02_10/include/moe_distribute_dispatch_non_quant.h) / [moe_distribute_combine.h](./answer/scaffold_02_10/include/moe_distribute_combine.h)）：\n",14+ "**关键常量速查表**（来自实践代码 [moe_distribute_comm.h](./src_02_10/include/moe_distribute_comm.h) / [moe_distribute_dispatch_non_quant.h](./src_02_10/include/moe_distribute_dispatch_non_quant.h) / [moe_distribute_combine.h](./src_02_10/include/moe_distribute_combine.h)）：\n",
15 "\n",15 "\n",
16 "<table style=\"margin: 0; margin-right: auto; border-collapse: collapse;\">\n",16 "<table style=\"margin: 0; margin-right: auto; border-collapse: collapse;\">\n",
17 " <tr>\n",17 " <tr>\n",
@@ -144,8 +144,8 @@
144 "\n",144 "\n",
145 "Win数据区的每一个 ping-pong half 都被 Dispatch 与 Combine 共用，二者切法如下（图 G）：\n",145 "Win数据区的每一个 ping-pong half 都被 Dispatch 与 Combine 共用，二者切法如下（图 G）：\n",
146 "\n",146 "\n",
147- "- **Combine 预留段**：位于 half 起点 `0 ~ axisBS × axisK × hSizeAlignCombine` 字节，其中 `hSizeAlignCombine = ceil(h * sizeof(X), 512) × 512`，即按 512 B 对齐每条 token 的 K 个 topk 槽。[Combine 2.1 阶段回送](./01.06_dispatch_combine_core_flow.ipynb)的 token 就写到这一段：第 i 条本地 token 的第 k 个 topk 槽位于 `(i × K + k) × hSizeAlignCombine`。\n",147+ "- **Combine 预留段**：位于 half 起点 `0 ~ axisBS × axisK × hSizeAlignCombine` 字节，其中 `hSizeAlignCombine = ceil(h * sizeof(X), 512) × 512`，即按 512 B 对齐每条 token 的 K 个 topk 槽。[Combine 2.1 阶段回送](./02.06_dispatch_combine_core_flow.ipynb)的 token 就写到这一段：第 i 条本地 token 的第 k 个 topk 槽位于 `(i × K + k) × hSizeAlignCombine`。\n",
148- "- **Dispatch 数据段**：紧接在 Combine 预留段之后，从 `axisBS × axisK × hSizeAlignCombine` 一直占到 half 末尾。[Dispatch 1.1阶段写到对端](./01.06_dispatch_combine_core_flow.ipynb)阶段写到对端的 token（按 (本地专家, 源卡, token 序号) 排列）落在这里。\n",148+ "- **Dispatch 数据段**：紧接在 Combine 预留段之后，从 `axisBS × axisK × hSizeAlignCombine` 一直占到 half 末尾。[Dispatch 1.1阶段写到对端](./02.06_dispatch_combine_core_flow.ipynb)阶段写到对端的 token（按 (本地专家, 源卡, token 序号) 排列）落在这里。\n",
149 "\n",149 "\n",
150 "代码对应（`SetDataStatus`）：\n",150 "代码对应（`SetDataStatus`）：\n",
151 "\n",151 "\n",
@@ -180,7 +180,7 @@
180 "- `blockCntPerToken_ = ceil(hOutSizeAlign_, 480)`（一条 token 切成多少个 480 B 子块）；\n",180 "- `blockCntPerToken_ = ceil(hOutSizeAlign_, 480)`（一条 token 切成多少个 480 B 子块）；\n",
181 "- `hCommuSize_ = blockCntPerToken_ × 512 B`（一条 token 在 Win 上的完整占位 = 子块数 × 512）。\n",181 "- `hCommuSize_ = blockCntPerToken_ × 512 B`（一条 token 在 Win 上的完整占位 = 子块数 × 512）。\n",
182 "\n",182 "\n",
183- "**每个 32 B token-flag 槽位的语义**：当 Dispatch 把一个数据子块成功推送到对端 Win 后，会紧跟着把一个 UB 对齐块（首位 `1.0`，其余 0）写到该子块尾部的 32 B 区域。[Dispatch 1.4 阶段](./01.06_dispatch_combine_core_flow.ipynb)在搬运对端 token 之前，先 `Sum` 这些 token-flag——求和等于「期望的子块数」即代表该 token 全部到齐。这就是 1.4 中提到的「token 粒度 flag」的物理形式。\n",183+ "**每个 32 B token-flag 槽位的语义**：当 Dispatch 把一个数据子块成功推送到对端 Win 后，会紧跟着把一个 UB 对齐块（首位 `1.0`，其余 0）写到该子块尾部的 32 B 区域。[Dispatch 1.4 阶段](./02.06_dispatch_combine_core_flow.ipynb)在搬运对端 token 之前，先 `Sum` 这些 token-flag——求和等于「期望的子块数」即代表该 token 全部到齐。这就是 1.4 中提到的「token 粒度 flag」的物理形式。\n",
184 "\n",184 "\n",
185 "**图 H · 一条 token 在 Win数据区上的物理排布**：\n",185 "**图 H · 一条 token 在 Win数据区上的物理排布**：\n",
186 "\n",186 "\n",
Mtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.08_tiling_guide.ipynb+5-5
@@ -16,7 +16,7 @@
16 "本实践示例代码核心目录如下：\n",16 "本实践示例代码核心目录如下：\n",
17 "\n",17 "\n",
18 "```\n",18 "```\n",
19- "answer/scaffold_01_10/ # 完整脚手架（含量化版与非量化版）\n",19+ "src_02_10/ # 完整脚手架（含量化版与非量化版）\n",
20 "├─ include/ # 头文件目录\n",20 "├─ include/ # 头文件目录\n",
21 "│ ├─ moe_distribute_dispatch.h # ★ Dispatch 算子核心头文件\n",21 "│ ├─ moe_distribute_dispatch.h # ★ Dispatch 算子核心头文件\n",
22 "│ │ # - DispatchTilingData 结构体定义\n",22 "│ │ # - DispatchTilingData 结构体定义\n",
@@ -92,7 +92,7 @@
92 "outputs": [],92 "outputs": [],
93 "source": [93 "source": [
94 "# 查看 Host 主程序中的 Tiling 函数定义\n",94 "# 查看 Host 主程序中的 Tiling 函数定义\n",
95- "!cat scaffold_02_10/src/dispatch_and_combine_final.asc | head -100"95+ "!cat src_02_10/src/dispatch_and_combine_final.asc | head -100"
96 ]96 ]
97 },97 },
98 {98 {
@@ -107,7 +107,7 @@
107 "outputs": [],107 "outputs": [],
108 "source": [108 "source": [
109 "# 查看 Dispatch TilingData 结构体定义（在头文件中搜索 DispatchTilingData）\n",109 "# 查看 Dispatch TilingData 结构体定义（在头文件中搜索 DispatchTilingData）\n",
110- "!grep -A 15 'struct.*DispatchTilingData' scaffold_02_10/include/moe_distribute_dispatch.h"110+ "!grep -A 15 'struct.*DispatchTilingData' src_02_10/include/moe_distribute_dispatch.h"
111 ]111 ]
112 },112 },
113 {113 {
@@ -122,7 +122,7 @@
122 "outputs": [],122 "outputs": [],
123 "source": [123 "source": [
124 "# 查看 Combine TilingData 结构体定义\n",124 "# 查看 Combine TilingData 结构体定义\n",
125- "!grep -A 15 'struct.*MoeDistributeCombineShmemTilingData' scaffold_02_10/include/moe_distribute_combine.h"125+ "!grep -A 15 'struct.*MoeDistributeCombineShmemTilingData' src_02_10/include/moe_distribute_combine.h"
126 ]126 ]
127 },127 },
128 {128 {
@@ -232,7 +232,7 @@
232 "\n",232 "\n",
233 "Host 端 Tiling 函数的核心动作就是 **按值填充结构体**——sample 中并没有复杂派生计算，所有派生量都留给 kernel 端算（见下一节）。\n",233 "Host 端 Tiling 函数的核心动作就是 **按值填充结构体**——sample 中并没有复杂派生计算，所有派生量都留给 kernel 端算（见下一节）。\n",
234 "\n",234 "\n",
235- "**最小示例（节选自 [src/0_non_quant_naive.asc](./answer/scaffold_01_10/src/0_non_quant_naive.asc)）**：\n",235+ "**最小示例（节选自 [src/0_non_quant_naive.asc](./src_02_10/src/0_non_quant_naive.asc)）**：\n",
236 "\n",236 "\n",
237 "```cpp\n",237 "```cpp\n",
238 "void SetDispatchTilingData(DispatchTilingData& dispatchTilingData,\n",238 "void SetDispatchTilingData(DispatchTilingData& dispatchTilingData,\n",
Mtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.09_kernel_stage_guide.ipynb+8-8
@@ -21,7 +21,7 @@
21 "本节涉及的代码文件与上一节（Tiling）相同，但视角从 Host 侧转向 Kernel 侧：\n",21 "本节涉及的代码文件与上一节（Tiling）相同，但视角从 Host 侧转向 Kernel 侧：\n",
22 "\n",22 "\n",
23 "```\n",23 "```\n",
24- "answer/scaffold_02_10/\n",24+ "src_02_10/\n",
25 "├─ include/ # ★ 本节重点：Kernel 算子实现\n",25 "├─ include/ # ★ 本节重点：Kernel 算子实现\n",
26 "│ ├─ moe_distribute_dispatch.h # Dispatch 算子类\n",26 "│ ├─ moe_distribute_dispatch.h # Dispatch 算子类\n",
27 "│ │ # - Init()：绑定 GM、计算派生量、分配 UB\n",27 "│ │ # - Init()：绑定 GM、计算派生量、分配 UB\n",
@@ -93,7 +93,7 @@
93 "outputs": [],93 "outputs": [],
94 "source": [94 "source": [
95 "# 查看 Kernel 入口函数（DispatchKernel / CombineKernel）\n",95 "# 查看 Kernel 入口函数（DispatchKernel / CombineKernel）\n",
96- "!grep -A 20 '__global__.*DispatchKernel' scaffold_02_10/src/dispatch_and_combine_final.asc"96+ "!grep -A 20 '__global__.*DispatchKernel' src_02_10/src/dispatch_and_combine_final.asc"
97 ]97 ]
98 },98 },
99 {99 {
@@ -108,7 +108,7 @@
108 "outputs": [],108 "outputs": [],
109 "source": [109 "source": [
110 "# 查看 Dispatch Process 阶段拆解\n",110 "# 查看 Dispatch Process 阶段拆解\n",
111- "!grep -A 15 'void.*Process()' scaffold_02_10/include/moe_distribute_dispatch.h | head -20"111+ "!grep -A 15 'void.*Process()' src_02_10/include/moe_distribute_dispatch.h | head -20"
112 ]112 ]
113 },113 },
114 {114 {
@@ -123,7 +123,7 @@
123 "outputs": [],123 "outputs": [],
124 "source": [124 "source": [
125 "# 查看 Combine Process 阶段拆解\n",125 "# 查看 Combine Process 阶段拆解\n",
126- "!grep -A 15 'void.*Process()' scaffold_02_10/include/moe_distribute_combine.h | head -20"126+ "!grep -A 15 'void.*Process()' src_02_10/include/moe_distribute_combine.h | head -20"
127 ]127 ]
128 },128 },
129 {129 {
@@ -138,7 +138,7 @@
138 "outputs": [],138 "outputs": [],
139 "source": [139 "source": [
140 "# 查看通信公共头（Win区偏移常量）\n",140 "# 查看通信公共头（Win区偏移常量）\n",
141- "!cat scaffold_02_10/include/moe_distribute_comm.h | head -60"141+ "!cat src_02_10/include/moe_distribute_comm.h | head -60"
142 ]142 ]
143 },143 },
144 {144 {
@@ -201,7 +201,7 @@
201 "source": [201 "source": [
202 "## **1. Kernel 入口模板**\n",202 "## **1. Kernel 入口模板**\n",
203 "\n",203 "\n",
204- "两个 kernel 入口遵循统一实现结构（节选自 [src/0_non_quant_naive.asc](./scaffold_01_10/src/dispatch_and_combine_final.asc)）：\n",204+ "两个 kernel 入口遵循统一实现结构（节选自 [src/0_non_quant_naive.asc](./src_02_10/src/dispatch_and_combine_final.asc)）：\n",
205 "\n",205 "\n",
206 "```cpp\n",206 "```cpp\n",
207 "__global__ __aicore__ __vector__ void DispatchKernel(\n",207 "__global__ __aicore__ __vector__ void DispatchKernel(\n",
@@ -463,9 +463,9 @@
463 "\n",463 "\n",
464 "---\n",464 "---\n",
465 "\n",465 "\n",
466- "**导航**：[← 上一章：Tiling开发指引](02.08_tiling_guide.ipynb) | [返回课程主页](02.01_chapter_intro.ipynb) | [下一章：量化实践 →](02.10_quant_dispatch_practice.ipynb)\n",466+ "**导航**：[← 上一章：Tiling开发指引](02.08_tiling_guide.ipynb) | [返回课程主页](02.01_chapter_intro.ipynb) | [下一章：章节测试 →](02.10_chapter_test.ipynb)\n",
467 "\n",467 "\n",
468- "> **下一步**：完成理论学习后，进入 [02.10 量化实践](02.10_quant_dispatch_practice.ipynb) 进行实际代码开发\n",468+ "> **下一步**：完成理论学习后，进入 [02.10 章节测试](02.10_chapter_test.ipynb) 进行实际代码开发\n",
469 "---"469 "---"
470 ]470 ]
471 }471 }
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.10_quant_dispatch_practice.ipynb→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/02.10_chapter_test.ipynb+10-11
@@ -9,7 +9,7 @@
9 "\n",9 "\n",
10 "## 背景\n",10 "## 背景\n",
11 "\n",11 "\n",
12- "在 [02.05 算子语义概览](02.05_operator_logic_overview.ipynb)、[02.06 核心流程拆解](02.06_dispatch_combine_core_flow.ipynb)、[02.07 Win 区布局](01.07_win_memory_layout.ipynb)、[02.08 Tiling 开发指引](02.08_tiling_guide.ipynb) 与 [02.09 Kernel Stage 开发指引](02.09_kernel_stage_guide.ipynb) 中，已经完整呈现了非量化版本 MoE Dispatch 算子的语义、Win 区布局、Tiling 字段与 Kernel 阶段划分。\n",12+ "在 [02.05 算子语义概览](02.05_operator_logic_overview.ipynb)、[02.06 核心流程拆解](02.06_dispatch_combine_core_flow.ipynb)、[02.07 Win 区布局](02.07_win_memory_layout.ipynb)、[02.08 Tiling 开发指引](02.08_tiling_guide.ipynb) 与 [02.09 Kernel Stage 开发指引](02.09_kernel_stage_guide.ipynb) 中，已经完整呈现了非量化版本 MoE Dispatch 算子的语义、Win 区布局、Tiling 字段与 Kernel 阶段划分。\n",
13 "\n",13 "\n",
14 "在大模型推理场景下，token 数据从 BF16/FP16 量化到 FP8 可显著降低跨卡通信量。本章实践要求在已有非量化版本理解的基础上，**实现 MX 量化版本 Dispatch 算子的关键路径**。MX 量化（Modular Exponent，每 32 个值共用一个 8-bit 指数 scale，输出 FP8_E5M2）的算子核函数已在实践代码中提供完整 `Quant::ComputeMaxExp` / `ComputeScale` / `ComputeData` 实现，**不需要重复实现量化算法本身**，本章实践聚焦的是把量化嵌入到 Dispatch 的数据流中。\n",14 "在大模型推理场景下，token 数据从 BF16/FP16 量化到 FP8 可显著降低跨卡通信量。本章实践要求在已有非量化版本理解的基础上，**实现 MX 量化版本 Dispatch 算子的关键路径**。MX 量化（Modular Exponent，每 32 个值共用一个 8-bit 指数 scale，输出 FP8_E5M2）的算子核函数已在实践代码中提供完整 `Quant::ComputeMaxExp` / `ComputeScale` / `ComputeData` 实现，**不需要重复实现量化算法本身**，本章实践聚焦的是把量化嵌入到 Dispatch 的数据流中。\n",
15 "\n",15 "\n",
@@ -142,7 +142,7 @@
142 "- `output/`：算子输出 bin（按 `chip_{rankId}` 分目录）\n",142 "- `output/`：算子输出 bin（按 `chip_{rankId}` 分目录）\n",
143 "\n",143 "\n",
144 "\n",144 "\n",
145- "本练习提供完整自包含脚手架 `scaffold_02_10/`（含 `CMakeLists.txt`、`include/`、`src/`、`cmake/`、`third_party/shmem/` 占位等）。运行下方第一个代码 cell 即把脚手架复制到 `practice_workspace/`，开发者只需编辑 `include/moe_distribute_dispatch.h` 中标有 `===== TODO N =====` 的 3 处函数体。"145+ "本练习提供完整自包含脚手架 `src_02_10/`（含 `CMakeLists.txt`、`include/`、`src/`、`cmake/`、`third_party/shmem/` 占位等）。运行下方第一个代码 cell 即把脚手架复制到 `practice_workspace/`，开发者只需编辑 `include/moe_distribute_dispatch.h` 中标有 `===== TODO N =====` 的 3 处函数体。"
146 ]146 ]
147 },147 },
148 {148 {
@@ -152,11 +152,11 @@
152 "metadata": {},152 "metadata": {},
153 "outputs": [],153 "outputs": [],
154 "source": [154 "source": [
155- "# 工程初始化：把自包含脚手架 scaffold_01_10/ 复制到 practice_workspace/\n",155+ "# 工程初始化：把自包含脚手架 src_02_10/ 复制到 practice_workspace/\n",
156 "import shutil\n",156 "import shutil\n",
157 "from pathlib import Path\n",157 "from pathlib import Path\n",
158 "\n",158 "\n",
159- "SCAFFOLD = Path('./scaffold_02_10')\n",159+ "SCAFFOLD = Path('./src_02_10')\n",
160 "WORKSPACE = Path('practice_workspace')\n",160 "WORKSPACE = Path('practice_workspace')\n",
161 "\n",161 "\n",
162 "assert SCAFFOLD.is_dir(), f'未找到脚手架目录：{SCAFFOLD.resolve()}'\n",162 "assert SCAFFOLD.is_dir(), f'未找到脚手架目录：{SCAFFOLD.resolve()}'\n",
@@ -215,7 +215,7 @@
215 "source": [215 "source": [
216 "## TODO 1：Buffer 大小计算\n",216 "## TODO 1：Buffer 大小计算\n",
217 "\n",217 "\n",
218- "> 🔗 [跳转到代码 TODO 1](./scaffold_02_10/include/moe_distribute_dispatch.h#L261)\n",218+ "> 🔗 [跳转到代码 TODO 1](./src_02_10/include/moe_distribute_dispatch.h#L261)\n",
219 "\n",219 "\n",
220 "**目标函数：** `MoeDistributeDispatch::SetTilingDataAndCal`（位于 `practice_workspace/include/moe_distribute_dispatch.h`）\n",220 "**目标函数：** `MoeDistributeDispatch::SetTilingDataAndCal`（位于 `practice_workspace/include/moe_distribute_dispatch.h`）\n",
221 "\n",221 "\n",
@@ -333,7 +333,7 @@
333 "source": [333 "source": [
334 "## TODO 2：发送流程\n",334 "## TODO 2：发送流程\n",
335 "\n",335 "\n",
336- "> 🔗 [跳转到代码 TODO 2](./scaffold_02_10/include/moe_distribute_dispatch.h#L357)\n",336+ "> 🔗 [跳转到代码 TODO 2](./src_02_10/include/moe_distribute_dispatch.h#L357)\n",
337 "\n",337 "\n",
338 "**目标函数：** `MoeDistributeDispatch::TokenToExpertInQuant`\n",338 "**目标函数：** `MoeDistributeDispatch::TokenToExpertInQuant`\n",
339 "\n",339 "\n",
@@ -389,7 +389,7 @@
389 "source": [389 "source": [
390 "## TODO 3：接收流程\n",390 "## TODO 3：接收流程\n",
391 "\n",391 "\n",
392- "> 🔗 [跳转到代码 TODO 3](./scaffold_02_10/include/moe_distribute_dispatch.h#L934)\n",392+ "> 🔗 [跳转到代码 TODO 3](./src_02_10/include/moe_distribute_dispatch.h#L934)\n",
393 "\n",393 "\n",
394 "**目标函数：** `MoeDistributeDispatch::CopyInAndOut`\n",394 "**目标函数：** `MoeDistributeDispatch::CopyInAndOut`\n",
395 "\n",395 "\n",
@@ -531,10 +531,9 @@
531 "outputs": [],531 "outputs": [],
532 "source": [532 "source": [
533 "import os\n",533 "import os\n",
534- "CANN_SAMPLES_ROOT = os.path.abspath(os.path.join('..', '..', '..', '..', 'cann-samples'))\n",534+ "PRACTICE_DIR = os.path.abspath('practice_workspace')\n",
535- "BUILD_DIR = os.path.join(CANN_SAMPLES_ROOT, 'build_practice')\n",535+ "BUILD_DIR = os.path.join(PRACTICE_DIR, 'build')\n",
536- "EXE = os.path.join(BUILD_DIR, 'Samples', '2_Performance',\n",536+ "EXE = os.path.join(BUILD_DIR, 'moe_dispatch_and_combine_story')\n",
537- " 'moe_dispatch_and_combine_story', 'moe_dispatch_and_combine_story')\n",
538 "!source /usr/local/Ascend/ascend-toolkit/set_env.sh && cd practice_workspace && {EXE} 2 8"537 "!source /usr/local/Ascend/ascend-toolkit/set_env.sh && cd practice_workspace && {EXE} 2 8"
539 ]538 ]
540 },539 },
Mtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/README.md+28-10
@@ -7,16 +7,34 @@
7 7 
8| Notebook | Link | 状态 |8| Notebook | Link | 状态 |
9|--|--|--|9|--|--|--|
10-| 2.1 章节介绍 | [在线体验](./02.01_chapter_intro.ipynb) | ✅ 已发布 |10+| 2.1 章节介绍 | - | ✅ 已发布 |
11-| 2.2 MoE 架构概述 | [在线体验](./02.02_moe_architexture_overview.ipynb) | ✅ 已发布 |11+| 2.2 MoE 架构概述 | - | ✅ 已发布 |
12-| 2.3 并行策略 | [在线体验](./02.03_parallel_strategy.ipynb) | ✅ 已发布 |12+| 2.3 并行策略 | - | ✅ 已发布 |
13-| 2.4 Dispatch/Combine 算子 | [在线体验](./02.04_dispatch_combine_operator.ipynb) | ✅ 已发布 |13+| 2.4 Dispatch/Combine 算子 | - | ✅ 已发布 |
14-| 2.5 算子逻辑概述 | [在线体验](./02.05_operator_logic_overview.ipynb) | ✅ 已发布 |14+| 2.5 算子逻辑概述 | - | ✅ 已发布 |
15-| 2.6 核心流程拆解 | [在线体验](./02.06_dispatch_combine_core_flow.ipynb) | ✅ 已发布 |15+| 2.6 核心流程拆解 | - | ✅ 已发布 |
16-| 2.7 Win 区内存布局 | [在线体验](./02.07_win_memory_layout.ipynb) | ✅ 已发布 |16+| 2.7 Win 区内存布局 | - | ✅ 已发布 |
17-| 2.8 Tiling 指南 | [在线体验](./02.08_tiling_guide.ipynb) | ✅ 已发布 |17+| 2.8 Tiling 指南 | - | ✅ 已发布 |
18-| 2.9 Kernel 阶段指南 | [在线体验](./02.09_kernel_stage_guide.ipynb) | ✅ 已发布 |18+| 2.9 Kernel 阶段指南 | - | ✅ 已发布 |
19-| 2.10 量化 Dispatch 实践 | [在线体验](./02.10_quant_dispatch_practice.ipynb) | ✅ 已发布 |19+| 2.10 章节测试 | - | ✅ 已发布 |
20+ 
21+## 软硬件配套说明
22+ 
23+| 项目 | 要求 |
24+| --- | --- |
25+| 支持硬件 | Atlas A5 训练/推理系列产品|
26+| CANN 版本 | 9.0.0 及以上 |
27+| Python | 3.11 |
28+ 
29+## 在线体验环境
30+ 
31+本教程支持以下在线体验环境：
32+ 
33+| 体验环境 | 镜像模板 / 版本 | Python 内核 | 说明 |
34+| --- | --- | --- | --- |
35+| CANNLab 950尝鲜体验 | cann_9.0.0_py3.11-A5-arm | Python 3.11.4 |参考 [CANNLab 环境体验指南](https://gitcode.com/cann/cann-learning-hub/blob/master/docs/CANNLab_env_experience_guide.md)创建CANNLab环境运行notebook |
36+ 
37+> **注意：** 如在本地环境离线体验，需自行安装配套的 CANN 软件，具体请参考 [CANN 安装指南](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/600alpha003/softwareinstall/instg/atlasdeploy_03_0001.html)，并选择对应CANN版本的文档。
20 38 
21 39 
22## 运行环境与约束40## 运行环境与约束
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/anwser_02_10/01_buffer_size_calc.txt→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/answer_02_10/01_buffer_size_calc.txt+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/anwser_02_10/02_send_flow.txt→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/answer_02_10/02_send_flow.txt+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/anwser_02_10/03_receive_flow.txt→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/answer/answer_02_10/03_receive_flow.txt+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/CMakeLists.txt→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/CMakeLists.txt+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/cmake/ascend.cmake→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/cmake/ascend.cmake+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/cmake/shmem.cmake→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/cmake/shmem.cmake+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/include/moe_distribute_combine.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/include/moe_distribute_combine.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/include/moe_distribute_comm.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/include/moe_distribute_comm.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/include/moe_distribute_dispatch.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/include/moe_distribute_dispatch.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/include/moe_distribute_dispatch_non_quant.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/include/moe_distribute_dispatch_non_quant.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/include/moe_distribute_dispatch_quant.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/include/moe_distribute_dispatch_quant.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/scripts/gen_data.py→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/scripts/gen_data.py+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/scripts/verify_result.py→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/scripts/verify_result.py+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/src/0_non_quant_naive.asc→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/src/0_non_quant_naive.asc+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/src/dispatch_and_combine_final.asc→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/src/dispatch_and_combine_final.asc+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/src/utils.h→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/src/utils.h+0-0
文件重命名但无更改。
Rtutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/scaffold_02_10/third_party/shmem/.gitkeep→tutorials/MC2_fused_operator_development/02_MoE_basic_operator_tutorial/src_02_10/third_party/shmem/.gitkeep+0-0
文件重命名但无更改。