forked from metax-maca/op_optimization
修改fused_moe教程
This commit is contained in:
Xinyi Wu (i26343) - Application Ecology
parent
af2909cf63
commit
911d79c1ac
|
|
@ -232,13 +232,13 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
|
||||
**常见问题：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `Python.h: No such file or directory` | Python 头文件路径未找到 | 确认 `PYTHON_BIN` 路径正确，脚本自动探测 `sysconfig.get_path('include')` |
|
||||
| `libpython3.x.so: cannot find` | 链接时找不到 Python 库 | 1、执行 `find $CONDA_PREFIX -name "libpython3*.so*"`查找绝对路径<br>2、将该路径赋值给 `LIBPYTHON_PATH` |
|
||||
| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
|
||||
| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
|
||||
| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `Python.h: No such file or directory` | Python 头文件路径未找到 | 确认 `PYTHON_BIN` 路径正确，脚本自动探测 `sysconfig.get_path('include')` |
|
||||
| `libpython3.x.so: cannot find` | 链接时找不到 Python 库 | 1、执行 `find $CONDA_PREFIX -name "libpython3*.so*"`查找绝对路径<br>2、将该路径赋值给 `LIBPYTHON_PATH` |
|
||||
| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
|
||||
| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
|
||||
| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
|
||||
|
||||
#### Step 4：正确性验证
|
||||
|
||||
|
|
@ -248,62 +248,62 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all # 运行全部计算方式
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all # 运行全部计算方式
|
||||
|
||||
# --backend：选择计算方式
|
||||
# 只测 pybind：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
|
||||
# 只测 triton：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
|
||||
# 只测 reference:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend reference
|
||||
```
|
||||
# --backend：选择计算方式
|
||||
# 只测 pybind：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
|
||||
# 只测 triton：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
|
||||
# 只测 reference:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend reference
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
编译成功无报错，输出示例如下：
|
||||
编译成功无报错，输出示例如下：
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69531, C\[last\]=-0.44531
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69531, C\[last\]=-0.44531
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57813, C\[last\]=-0.49805
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57813, C\[last\]=-0.49805
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08594, C\[last\]=-0.33594
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08594, C\[last\]=-0.33594
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.6934, C\[last\]=-0.4451
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.6934, C\[last\]=-0.4451
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.5768, C\[last\]=-0.4975
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.5768, C\[last\]=-0.4975
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.0875, C\[last\]=-0.3362
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.0875, C\[last\]=-0.3362
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69337, C\[last\]=-0.44513
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69337, C\[last\]=-0.44513
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57678, C\[last\]=-0.49749
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57678, C\[last\]=-0.49749
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08748, C\[last\]=-0.33618
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08748, C\[last\]=-0.33618
|
||||
|
||||
|
||||
**结果解释：**
|
||||
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
|
||||
|
||||
* ”rows=... , cols=...“：输出 Tensor 的形状；
|
||||
* ”rows=... , cols=...“：输出 Tensor 的形状；
|
||||
|
||||
* ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
|
||||
* ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
|
||||
|
||||
|
||||
|
||||
**常见问题：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| `FAILED: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误；确认 TopK 索引与权重是否一致 |
|
||||
| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
|
||||
| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
|
||||
| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| `FAILED: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误；确认 TopK 索引与权重是否一致 |
|
||||
| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
|
||||
| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
|
||||
| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
|
||||
|
||||
#### Step 5：性能测试
|
||||
|
||||
|
|
@ -313,59 +313,60 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
|
||||
# --backend：选择计算方式
|
||||
# --warmup：设置预热次数
|
||||
# --iters：设置迭代次数
|
||||
```
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
|
||||
# --backend：选择计算方式
|
||||
# --warmup：设置预热次数
|
||||
# --iters：设置迭代次数
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
编译成功无报错，输出示例如下：
|
||||
编译成功无报错，输出示例如下：
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=0.308978, TOPS=0.027149, warmup=5, iters=20
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=0.308978, TOPS=0.027149, warmup=5, iters=20
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=0.304500, TOPS=0.055098, warmup=5, iters=20
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=0.304500, TOPS=0.055098, warmup=5, iters=20
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=0.297775, TOPS=0.042256, warmup=5, iters=20
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=0.297775, TOPS=0.042256, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=1685.43, TOPS=0.000005, warmup=5, iters=20
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=1685.43, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=3384.52, TOPS=0.000005, warmup=5, iters=20
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=3384.52, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=2532.14, TOPS=0.000005, warmup=5, iters=20
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=2532.14, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=19.013421, TOPS=0.000441, warmup=5, iters=20
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=19.013421, TOPS=0.000441, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=16.745914, TOPS=0.001002, warmup=5, iters=20
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=16.745914, TOPS=0.001002, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=19.630328, TOPS=0.000641, warmup=5, iters=20
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=19.630328, TOPS=0.000641, warmup=5, iters=20
|
||||
|
||||
|
||||
**结果解释：**
|
||||
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3”：分别对应选择前 1 / 2 / 3 个专家场景下的 MoE 算子；
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3”：分别对应选择前 1 / 2 / 3 个专家场景下的 MoE 算子；
|
||||
|
||||
* “avg\_ms”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的时间；
|
||||
* “avg\_ms”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的时间；
|
||||
|
||||
* “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗时；
|
||||
* “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗时；
|
||||
|
||||
* “warmup=5, iters=20”：预热轮数和正式迭代数。
|
||||
* “warmup=5, iters=20”：预热轮数和正式迭代数。
|
||||
|
||||
|
||||
**常见错误：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape；重启容器或设备 |
|
||||
| avg\_ms 异常抖动（±50%） | 其他进程占用 GPU | 关闭其他占用显存的进程，单机单任务运行 |
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape；重启容器或设备 |
|
||||
| avg\_ms 异常抖动（±50%） | 其他进程占用 GPU | 关闭其他占用显存的进程，单机单任务运行 |
|
||||
|
||||
### 6.2 在 XPU-OJ 平台进行提交
|
||||
|
||||
平台链接：[https://xpuoj.com/](https://xpuoj.com/)
|
||||
平台链接：[https://xpuoj.com/](https://xpuoj.com/)
|
||||
|
||||
#### Step 6：从 Benchmark 到 XPU-OJ 提交
|
||||
|
||||
|
|
@ -375,16 +376,16 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
|
||||
1、厘清 Benchmark 和 XPU-OJ 的区别：
|
||||
|
||||
赛事镜像中的 Benchmark 脚本用于理解目标算子的调用方式、输入输出 shape 和性能基线；XPU-OJ 题包用于定义最终评测接口、数据范围、参考输出和精度要求。
|
||||
赛事镜像中的 Benchmark 脚本用于理解目标算子的调用方式、输入输出 shape 和性能基线；XPU-OJ 题包用于定义最终评测接口、数据范围、参考输出和精度要求。
|
||||
|
||||
| **维度** | **Benchmark 脚本** | **XPU-OJ 提交** |
|
||||
| --- | --- | --- |
|
||||
| **目的** | 理解算子接口、建立性能基线 | 统一环境下的正确性+性能评测 |
|
||||
| **接口形式** | Python API | 三种接口供选择：<br>* CUDA C（`extern "C"`）<br> <br>* TileLang（Python `@jit`）<br> <br>* Triton（Python `@triton.jit`） |
|
||||
| **函数签名** | `backend_fn(...)` | `run_kernel(...)` |
|
||||
| **数据范围** | 多种 head\_dim / batch\_size / seq\_len 组合 | 固定参数范围（以题包为准） |
|
||||
| **验证** | 无自动正确性校验，人工对比输出数值 | 强制通过 `torch.allclose(rtol=2e-2, atol=5e-3)` |
|
||||
| **输出** | 终端直接输出 | 排行榜得分 |
|
||||
| **维度** | **Benchmark 脚本** | **XPU-OJ 提交** |
|
||||
| --- | --- | --- |
|
||||
| **目的** | 理解算子接口、建立性能基线 | 统一环境下的正确性+性能评测 |
|
||||
| **接口形式** | Python API | 三种接口供选择：<br>* CUDA C（`extern "C"`）<br> <br>* TileLang（Python `@jit`）<br> <br>* Triton（Python `@triton.jit`） |
|
||||
| **函数签名** | `backend_fn(...)` | `run_kernel(...)` |
|
||||
| **数据范围** | 多种 head\_dim / batch\_size / seq\_len 组合 | 固定参数范围（以题包为准） |
|
||||
| **验证** | 无自动正确性校验，人工对比输出数值 | 强制通过 `torch.allclose(rtol=2e-2, atol=5e-3)` |
|
||||
| **输出** | 终端直接输出 | 排行榜得分 |
|
||||
|
||||
2、理解完成 benchmark 验证并成功建立性能基线后，需要完成以下转换：
|
||||
|
||||
|
|
@ -526,25 +527,28 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
* OJ 平台对单测试点的评分下公式：
|
||||
|
||||
|
||||
$S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_k - T\_h}{T\_b - T\_h}}$
|
||||
$$
|
||||
S(T_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T_k - T_h}{T_b - T_h}}
|
||||
$$
|
||||
|
||||
其中，$T\_k$代表你提交的 kernel 平均执行时间；$T\_b$代表基准算子的的平均执行时间，对应 50 分；$T\_h$代表硬件理论下限耗时，对应100分 。
|
||||
其中，$T_k$代表你提交的 kernel 平均执行时间；$T_b$代表基准算子的的平均执行时间，对应 50 分；$T_h$代表硬件理论下限耗时，对应100分 。
|
||||
|
||||
当单测试点得分超过 150 分时，平台会按对数压缩规则显示：
|
||||
当单测试点得分超过 150 分时，平台会按对数压缩规则显示：
|
||||
|
||||
$$
|
||||
$S\_{display}=150+10\*log\_{10}(S/150)$
|
||||
$$
|
||||
|
||||
总得分为各测试点得分的算术平均，总耗时为各测试点$T\_k$的求和。
|
||||
总得分为各测试点得分的算术平均，总耗时为各测试点$T_k$的求和。
|
||||
|
||||
* 关键分数节点：
|
||||
|
||||
|
||||
| **性能** | **得分** | **含义** |
|
||||
| --- | --- | --- |
|
||||
| $T\_k=T\_b$ | 50 分 | 与 Baseline 等速 |
|
||||
| $T\_k=T\_h$ | 100 分 | 达到硬件理论上限 |
|
||||
| $T\_k<T\_h$ | 大于 100 分 | 超越理论估算，可能因估算偏保守 |
|
||||
| $T\_k≫T\_b$ | 接近 0 分 | 远慢于 Baseline |
|
||||
| **性能** | **得分** | **含义** |
|
||||
| --- | --- | --- |
|
||||
| $T_k=T_b$ | 50 分 | 与 Baseline 等速 |
|
||||
| $T_k=T_h$ | 100 分 | 达到硬件理论上限 |
|
||||
| $T_k<T_h$ | 大于 100 分 | 超越理论估算，可能因估算偏保守 |
|
||||
| $T_k≫T_b$ | 接近 0 分 | 远慢于 Baseline |
|
||||
|
||||
#### Step 13：榜单查看与优化方向
|
||||
|
||||
|
|
@ -567,18 +571,17 @@ PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
|||
|
||||
[](https://www.picgo.net/image/image9.4t3S4J)
|
||||
|
||||
2. 制定优化方向
|
||||
|
||||
2. 制定优化方向
|
||||
|
||||
| **优化方向** | **具体说明** |
|
||||
| --- | --- |
|
||||
| 算子融合 | 将矩阵乘、scale、softmax 等步骤合并为单个内核，减少 HBM 往返。 |
|
||||
| 并行策略调整 | 在 Decode 阶段采用分块或 Split-K 思路，提升长 KV 序列并行度。 |
|
||||
| 在线 Softmax | 引入局部最大值与局部求和动态缩放，保证数值稳定并减少访存。 |
|
||||
| 显存访问合并 | 保证相邻线程访问相邻地址，按 HeadDim 等连续维度向量化加载。 |
|
||||
| 利用内存层级 | 将频繁更新的标量放入寄存器，块内复用数据放入共享内存。 |
|
||||
| 软件流水线 | 在当前分块计算时预取下一分块数据，隐藏内存加载延迟。 |
|
||||
| 减少冗余计算 | 提取循环不变量，处理变长序列时减少复杂分支。 |
|
||||
| **优化方向** | **具体说明** |
|
||||
| --- | --- |
|
||||
| 算子融合 | 将矩阵乘、scale、softmax 等步骤合并为单个内核，减少 HBM 往返。 |
|
||||
| 并行策略调整 | 在 Decode 阶段采用分块或 Split-K 思路，提升长 KV 序列并行度。 |
|
||||
| 在线 Softmax | 引入局部最大值与局部求和动态缩放，保证数值稳定并减少访存。 |
|
||||
| 显存访问合并 | 保证相邻线程访问相邻地址，按 HeadDim 等连续维度向量化加载。 |
|
||||
| 利用内存层级 | 将频繁更新的标量放入寄存器，块内复用数据放入共享内存。 |
|
||||
| 软件流水线 | 在当前分块计算时预取下一分块数据，隐藏内存加载延迟。 |
|
||||
| 减少冗余计算 | 提取循环不变量，处理变长序列时减少复杂分支。 |
|
||||
|
||||
## 七、项目实践2-Kernel Swift 智能算子迁移系统自动调优
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue