| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
添加aclgraph+<<<>>> 示例代码 Co-authored-by: wangkai<wangkai579@huawei.com> # message auto-generated for no-merge-commit merge: !4019 merge master into master 添加aclgraph+<<<>>> 示例代码 Created-by: mihudan Commit-by: wangkai Merged-by: ascend-robot Description: <!-- Thanks for sending a pull request! --> **What type of PR is this?** > Uncomment only one /kind <> line, hit enter to put that in a new line, and remove leading whitespaces from that line: > > /kind bug /kind task > /kind feature **What does this PR do / why do we need it**: 添加aclgraph 与<<<>>>联调的demo测试代码，进行验证并供用户参考。 展示了如何使用PyTorch的torch.librar以及pybind两张方式注册自定义算子，通过<<<>>>内核调用符调用核函数，并适配aclgraph使用该自定义算子，以简单的Add算子和三角函数计算的原地算子为例，实现aclgraph下自定义算子的调用。 展示了3种aclgraph的使能方式，通过对比NPU输出与CPU标准加法结果来验证自定义算子的数值正确性。 1. torch.npu.NPUGraph() 2. torch.npu.make_graphed_callables 3. backend="npugraph_ex" ## 算子描述 ### Add算子 - 算子功能： Add算子实现了两个数据相加，返回相加结果的功能。对应的算子原型为： ascendc_add(Tensor x, Tensor y) -> Tensor - 算子规格： <table> <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">add_custom</td></tr> </tr> <tr><td rowspan="3" align="center">算子输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> <tr><td align="center">x</td><td align="center">8 * 2048</td><td align="center">int</td><td align="center">ND</td></tr> <tr><td align="center">y</td><td align="center">8 * 2048</td><td align="center">int</td><td align="center">ND</td></tr> </tr> </tr> <tr><td rowspan="1" align="center">算子输出</td><td align="center">z</td><td align="center">8 * 2048</td><td align="center">int</td><td align="center">ND</td></tr> </tr> </table> ### 原地三角函数算子 - 算子功能： 该算子入参为x, out_sin ,out_cos, 算子调用后，out_sin会被原地修改为sin(x)计算结果，out_cos会被原地修改为cos(x)计算结果，返回值tan(x)计算结果。对应的算子原型为： ascendc_trig(Tensor x, Tensor(a!) out_sin, Tensor(b!) out_cos) -> Tensor - 算子规格： <table> <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">trig_inplace_custom</td></tr> </tr> <tr><td rowspan="4" align="center">算子输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> <tr><td align="center">x</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> <tr><td align="center">out_sin</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> <tr><td align="center">out_cos</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> </tr> </tr> <tr><td rowspan="3" align="center">算子输出</td><td align="center">out_sin</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> <tr><td align="center">out_cos</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> <tr><td align="center">out_tan</td><td align="center">8 * 2048</td><td align="center">float</td><td align="center">ND</td></tr> </tr> </table> ## 代码实现介绍 - 以Add算子为例，样例在*.asc文件中定义了一个名为ascendc_ops的命名空间，并在其中注册了ascendc_add函数。在ascendc_add函数中通过c10_npu::getCurrentNPUStream()函数获取当前NPU上的流，并通过内核调用符<<<>>>调用自定义的Kernel函数add_custom，在NPU上执行算子。 c++ add_custom<<<blockDim, nullptr, aclStream>>>(xGm, yGm, zGm, totalLength); - PyTorch提供TORCH_LIBRARY_FRAGMENT宏作为自定义算子注册的核心接口，用于创建并初始化自定义算子库，注册后在Python侧可以通过torch.ops.namespace.op_name方式进行调用，例如： c++ TORCH_LIBRARY_FRAGMENT(ascendc_ops, m) { m.def(ascendc_add"(Tensor x, Tensor y) -> Tensor"); } - TORCH_LIBRARY_IMPL用于将算子逻辑绑定到特定的DispatchKey（PyTorch设备调度标识）。针对NPU设备，需要将算子实现注册到PrivateUse1这一专属的DispatchKey上，例如： c++ TORCH_LIBRARY_IMPL(ascendc_ops, PrivateUse1, m) { m.impl("ascendc_add", TORCH_FN(ascendc_ops::ascendc_add)); } - 注册Meta函数： 注册Meta函数使faketensor流程正常工作，在使用fx, compile等功能涉及，注册代码如下： c++ TORCH_LIBRARY_IMPL(ascendc_ops, Meta, m) { m.impl("ascendc_add", &add_impl_meta); } **Special notes for your reviewers**: See merge request: Ascend/op-plugin!4019 | 5 个月前 | |
fix <<<>>> sample doc Co-authored-by: wangkai<wangkai579@huawei.com> # message auto-generated for no-merge-commit merge: !4080 merge master into master fix <<<>>> sample doc Created-by: mihudan Commit-by: wangkai Merged-by: ascend-robot Description: <!-- PR描述模板更新日期：20260203 --> # 【合入来源】 > (如有)请关联需求文档/issue链接 - [x] 需求 # 【修改方案】 > 请描述修改内容的具体实现，涉及哪些组件之间进行交互，可以用1、2、3、...进行罗列\ > 如果是需求或者重构类的PR，需要补充详细设计文档（说明上下游组件关系、时序图、类图、DFX能力等内容） <<<>>> aclgraph代码示例 注释描述有误 # 【资料变更】 > 请确认是否涉及资料变更。如涉及，需要在PR中体现，并简要说明修改内容。如不涉及，需填写“不涉及” 不涉及 # 【接口变更】 > 请确认是否涉及跨代码仓或者客户面可见的接口变更。如涉及，需要详细说明接口以及对应的变更内容，同时需要在资料中体现。如不涉及，需填写“不涉及” 不涉及 # 【功能验证】 > 说明测试场景，测试方法。如果本次测试方式与常规单元测试不同，请详细说明您的测试步骤\ > 新增/变更内容是否已新增/适配UT测试用例看护，并补充测试自验证截图 示例demo，需编译安装，无UT # 【CheckList】 > PR提交人对以下CheckList自检项进行全量自检，自检通过或不涉及，均修改 [ ] 为 [x] - [x] 代码注释完备，正确记录错误日志 - [x] 代码实现进行了返回值、空指针等校验 - [x] PR标题正确使用类型标签，如：feat、fix、refactor、docs、test等 - [x] PR持续集成流水线（CI）执行通过，代码检查无异常 See merge request: Ascend/op-plugin!4080 | 4 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 5 个月前 | ||
| 4 个月前 |