diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md"
similarity index 82%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md"
index b1b3aa8e14f27c713023fef5e9ef23d7bee37a58..aa64a03016355cc1099f4563782576dd26ebde36 100644
--- "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227.md"
+++ "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227.md"
@@ -1,18 +1,18 @@
-# PyTorch API开发指南
+# PyTorch 算子开发指南
- [简介](#简介md)
- [算子开发流程](#算子开发流程md)
- [算子开发准备](#算子开发准备md)
- [环境准备](#环境准备md)
- - [API速查](#API速查md)
-- [API适配开发](#API适配开发md)
+ - [算子速查](#算子速查md)
+- [算子适配开发](#算子适配开发md)
- [前提条件](#前提条件md)
- [获取PyTorch源码](#获取PyTorch源码md)
- - [注册API开发](#注册API开发md)
+ - [注册算子开发](#注册算子开发md)
- [概述](#概述md)
- - [PyTorch1.8.1 注册API开发](#PyTorch1-8-1-注册API开发md)
- - [API适配插件开发](#API适配插件开发md)
+ - [PyTorch1.8.1 注册算子开发](#PyTorch1-8-1-注册算子开发md)
+ - [算子适配插件开发](#算子适配插件开发md)
- [编译和安装PyTorch框架](#编译和安装PyTorch框架md)
-- [API功能验证](#API功能验证md)
+- [算子功能验证](#算子功能验证md)
- [概述](#概述-0md)
- [实现过程](#实现过程md)
- [FAQ](#FAQmd)
@@ -26,7 +26,7 @@
- [PyTorch编译失败，提示“error: call of overload ....”](#PyTorch编译失败-提示-error-call-of-overloadmd)
- [附录](#附录md)
- [CMake安装方法](#CMake安装方法md)
- - [自定义API导出方法](#自定义API导出方法md)
+ - [自定义算子导出方法](#自定义算子导出方法md)
简介
### 概述
@@ -35,17 +35,17 @@
算子开发流程
-PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
+PyTorch算子开发包含TBE算子开发和PyTorch框架下的算子适配。
-1. TBE算子开发：昇腾AI软件栈中不包含相应的算子，需要先完成TBE算子的开发，再进行PyTorch框架下的API适配。
+1. TBE算子开发：昇腾AI软件栈中不包含相应的算子，需要先完成TBE算子的开发，再进行PyTorch框架下的算子适配。
TBE算子开发流程及方法请参见《CANN TBE自定义算子开发指南》。
-2. PyTorch框架下的API适配：昇腾AI软件栈中已实现了相应的TBE算子，可直接进行PyTorch框架适配。
+2. PyTorch框架下的算子适配：昇腾AI软件栈中已实现了相应的TBE算子，可直接进行PyTorch框架适配。
- PyTorch框架下的API适配流程如下所示。
+ PyTorch框架下的算子适配流程如下所示。
- **图 1** PyTorch框架下的API适配流程
+ **图 1** PyTorch框架下的算子适配流程
**表 1** 算子开发步骤详解
@@ -68,15 +68,15 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
准备算子开发及运行验证所依赖的开发环境与运行环境。
|
-算子开发准备
+ | 算子开发准备
|
2
|
-API速查
+ | 算子速查
|
-查看TBE算子支持列表和PyTorchAPI适配列表。
-- 当前昇腾AI处理器支持的算子列表及支持的算子的详细规格约束。
- 当前PyTorch适配的API列表。
+ | 查看TBE算子支持列表和PyTorch算子适配列表。
+- 当前昇腾AI处理器支持的算子列表及支持的算子的详细规格约束。
- 当前PyTorch适配的算子列表。
|
3
@@ -85,7 +85,7 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
|
获取昇腾社区PyTorch源码。
|
-API适配开发
+ | 算子适配开发
|
4
@@ -97,9 +97,9 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
|
5
|
-API适配层开发
+ | 算子适配层开发
|
-API适配层开发，将基于第三方框架的算子属性映射成适配昇腾AI处理器的算子属性。
+ | 算子适配层开发，将基于第三方框架的算子属性映射成适配昇腾AI处理器的算子属性。
|
6
@@ -111,11 +111,11 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
|
7
|
-API功能验证
+ | 算子功能验证
|
-在真实的硬件环境中验证API功能。
+ | 在真实的硬件环境中验证算子功能。
|
-API功能验证
+ | 算子功能验证
|
@@ -123,11 +123,12 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
+
算子开发准备
- **[环境准备](#环境准备md)**
-- **[API速查](#API速查md)**
+- **[算子速查](#算子速查md)**
环境准备
@@ -153,30 +154,30 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
-API速查
+算子速查
-进行API开发时，您可以查询当前昇腾AI处理器中支持的算子列表和当前PyTorch适配的API列表。根据查询结果进行算子开发或PyTorchAPI适配。
+进行算子开发时，您可以查询当前昇腾AI处理器中支持的算子列表和当前PyTorch适配的算子列表。根据查询结果进行算子开发或PyTorch算子适配。
-- 若昇腾AI处理器不支持算子，则需要进行TBE算子开发和PyTorch框架API适配。
-- 若昇腾AI处理器支持算子但PyTorch框架没有API适配，仅需要进行PyTorch框架下API适配。
-- 若PyTorch框架已经适配API，直接使用API即可，不需要开发和适配。
+- 若昇腾AI处理器不支持算子，则需要进行TBE算子开发和PyTorch框架算子适配。
+- 若昇腾AI处理器支持算子但PyTorch框架没有算子适配，仅需要进行PyTorch框架下算子适配。
+- 若PyTorch框架已经适配算子，直接使用算子即可，不需要开发和适配。
-昇腾AI处理器和PyTorch适配的API查询方式如下。
+昇腾AI处理器和PyTorch适配的算子查询方式如下。
-- 当前昇腾AI处理器中支持的算子以及对应的API约束可以通过以下两种方式查询。
+- 当前昇腾AI处理器中支持的算子以及对应的算子约束可以通过以下两种方式查询。
- 命令行开发方式下，您可以参见《CANN 算子清单 \(Ascend 910\)》进行离线查询。
- MindStudio开发方式下，您可以通过MindStudio进行在线查询，详细查看方法可参见《MindStudio 用户指南》中的“算子&模型速查”章节。
-- 当前PyTorch适配的API列表可以参见[《PyTorch API支持清单》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%20API%E6%94%AF%E6%8C%81%E6%B8%85%E5%8D%95.md)。
+- 当前PyTorch适配的算子列表可以参见[《PyTorch API支持清单》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%20API%E6%94%AF%E6%8C%81%E6%B8%85%E5%8D%95.md)。
-API适配开发
+算子适配开发
- **[前提条件](#前提条件md)**
- **[获取PyTorch源码](#获取PyTorch源码md)**
-- **[注册API开发](#注册API开发md)**
+- **[注册算子开发](#注册算子开发md)**
-- **[API适配插件开发](#API适配插件开发md)**
+- **[算子适配插件开发](#算子适配插件开发md)**
- **[编译和安装PyTorch框架](#编译和安装PyTorch框架md)**
@@ -188,52 +189,52 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
获取PyTorch源码
-针对pytorch1.8.1版本，PyTorch源码获取请参见[《PyTorch安装指南》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97.md)中“安装PyTorch框架”章节，完成在"pytorch/pytorch_v1.8.1"目录生成适配昇腾AI处理器的全量代码步骤。将在pytorch/pytorch目录中进行PyTorch API适配开发。
+针对pytorch1.8.1版本，PyTorch源码获取请参见[《PyTorch安装指南》](https://gitee.com/ascend/pytorch/blob/master/docs/zh/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97/PyTorch%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97.md)中“安装PyTorch框架”章节，完成在"pytorch/pytorch_v1.8.1"目录生成适配昇腾AI处理器的全量代码步骤。将在pytorch/pytorch目录中进行PyTorch 算子适配开发。
-注册API开发
+注册算子开发
- **[概述](#概述md)**
-- **[PyTorch1.8.1 注册API开发](#PyTorch1-8-1-注册API开发md)**
+- **[PyTorch1.8.1 注册算子开发](#PyTorch1-8-1-注册算子开发md)**
概述
-当前制定的NPU适配派发原则是NPUAPI的派发不经过框架公共函数，直接派发成NPU适配的函数，即API执行调用栈中只包含NPU适配的函数调用，不包含框架公共函数。PyTorch框架在编译时，会根据 native\_functions.yaml 的定义，按框架中定义的类型和设备分发原则，生成相应的新API的中间层的调用说明。对于NPU，会生成在 build/aten/src/ATen/NPUType.cpp。
+当前制定的NPU适配派发原则是NPU算子的派发不经过框架公共函数，直接派发成NPU适配的函数，即算子执行调用栈中只包含NPU适配的函数调用，不包含框架公共函数。PyTorch框架在编译时，会根据 native\_functions.yaml 的定义，按框架中定义的类型和设备分发原则，生成相应的新算子的中间层的调用说明。对于NPU，会生成在 build/aten/src/ATen/NPUType.cpp。
-PyTorch1.8.1 注册API开发
+PyTorch1.8.1 注册算子开发
-##### 注册API开发方法
+##### 注册算子开发方法
1. 打开native\_functions.yaml文件。
- native\_functions.yaml 文件中，定义了所有API函数原型，包括函数名称和参数等信息，每个API函数支持不同硬件平台的派发信息。该文件所在路径为pytorch/aten/src/ATen/native/native\_functions.yaml。
+ native\_functions.yaml 文件中，定义了所有算子函数原型，包括函数名称和参数等信息，每个算子函数支持不同硬件平台的派发信息。该文件所在路径为pytorch/aten/src/ATen/native/native\_functions.yaml。
2. 确定需要派发函数。
- - yaml 中已存在的API
+ - yaml 中已存在的算子
- 将所有与待适配API相关的函数进行派发。
+ 将所有与待适配算子相关的函数进行派发。
- - yaml中未存在的自定义API
+ - yaml中未存在的自定义算子
- 由于yaml中没有相关API的信息，需要手动添加相关函数，包括函数名，参数信息，返回类型信息。添加规则及方法请参见“pytorch/aten/src/ATen/native/README.md“。
+ 由于yaml中没有相关算子的信息，需要手动添加相关函数，包括函数名，参数信息，返回类型信息。添加规则及方法请参见“pytorch/aten/src/ATen/native/README.md“。
```
- - func：适配API名称(输入参数信息) -> 返回类型
+ - func：适配算子名称(输入参数信息) -> 返回类型
```
##### 示例
-以torch.add\(\)API为例介绍注册API开发过程。
+以torch.add\(\)算子为例介绍注册算子开发过程。
1. 打开native\_functions.yaml文件。
2. 搜索相关函数。
- 在yaml中搜索add，找到与addAPI相关的函数描述func。由于add是PyTorch内置API，不需要手动添加func。若是自定义API，需要手动添加func。
+ 在yaml中搜索add，找到与add算子相关的函数描述func。由于add是PyTorch内置算子，不需要手动添加func。若是自定义算子，需要手动添加func。
-3. 确定API相关函数名称及其类型的func描述。
+3. 确定算子相关函数名称及其类型的func描述。
- add.Tensor 的函数分发描述。
```
@@ -289,18 +290,18 @@ PyTorch算子开发包含TBE算子开发和PyTorch框架下的API适配。
-API适配插件开发
+算子适配插件开发
#### 简介
-用户通过开发API适配插件，实现PyTorch原生API的输入参数、输出参数和属性的格式转换，使转换后的格式与TBE算子的输入参数、输出参数和属性的格式相同。适配昇腾AI处理器的PyTorch源代码中提供了适配关联、类型转换和判别、处理动态shape等相关的方法供用户使用。
+用户通过开发算子适配插件，实现PyTorch原生算子的输入参数、输出参数和属性的格式转换，使转换后的格式与TBE算子的输入参数、输出参数和属性的格式相同。适配昇腾AI处理器的PyTorch源代码中提供了适配关联、类型转换和判别、处理动态shape等相关的方法供用户使用。
#### npu_native_functions.yaml文件介绍
```
backend: NPU # Backend类型
-cpp_namespace: at_npu::native # 插件中开发API的命名空间
-supported: # 已支持的和PyTorch Native Functions对齐的API
+cpp_namespace: at_npu::native # 插件中开发算子的命名空间
+supported: # 已支持的和PyTorch Native Functions对齐的算子
- add.Tensor
- add.Scalar
- slow_conv3d.out
@@ -313,10 +314,10 @@ supported: # 已支持的和PyTorch Native Functions对齐的API
- addcdiv_
- addcdiv.out
-autograd: # 已支持的和PyTorch Native Functions对齐的继承自Function的具有前反向操作的API
+autograd: # 已支持的和PyTorch Native Functions对齐的继承自Function的具有前反向操作的算子
- maxpool2d
-custom: # 自定义API，需要提供API格式定义
+custom: # 自定义算子，需要提供算子格式定义
- func: npu_dtype_cast(Tensor self, ScalarType dtype) -> Tensor
variants: function, method
- func: npu_dtype_cast_(Tensor(a!) self, Tensor src) -> Tensor(a!)
@@ -326,21 +327,21 @@ custom: # 自定义API，需要提供API格式定义
- func: npu_get_float_status(Tensor self) -> Tensor
variants: function, method
-custom_autograd: # 自定义继承自Function的自定义API
+custom_autograd: # 自定义继承自Function的自定义算子
- func: npu_convolution(Tensor input, Tensor weight, Tensor? bias, ...) -> Tensor
```
-官方提供的native_functions.yaml文件定义了PyTorch Native Functions的具体API定义和分发细节，在NPU设备上适配官方已定义API，我们不需要重新定义，只需要注册NPU分发即可。由于我们可以根据已支持的API（supported，autograd）对应解析官方yaml文件得到每个函数的具体格式，所以对应的函数声明和注册分发可以自动化完成，API迁移和开发的时候只需要关注对应的实现细节即可。对于自定义API，由于没有具体的API定义，我们需要在npu_native_functions.yaml文件中给出定义，以便对API进行结构化解析从而实现自动化注册和Python接口绑定。
+官方提供的native_functions.yaml文件定义了PyTorch Native Functions的具体算子定义和分发细节，在NPU设备上适配官方已定义算子，我们不需要重新定义，只需要注册NPU分发即可。由于我们可以根据已支持的算子（supported，autograd）对应解析官方yaml文件得到每个函数的具体格式，所以对应的函数声明和注册分发可以自动化完成，算子迁移和开发的时候只需要关注对应的实现细节即可。对于自定义算子，由于没有具体的算子定义，我们需要在npu_native_functions.yaml文件中给出定义，以便对算子进行结构化解析从而实现自动化注册和Python接口绑定。
#### 适配插件实现
-1. 注册API。
+1. 注册算子。
- 根据npu_native_functions.yaml文件介绍，添加相应API信息。
+ 根据npu_native_functions.yaml文件介绍，添加相应算子信息。
2. 创建适配插件文件。
- NPU TBE算子适配文件保存在pytorch/torch_npu/csrc/aten/ops目录下，命名风格采用大驼峰，命名格式： + .cpp，如：AddKernelNpu.cpp。
+ NPU TBE算子适配文件保存在pytorch/torch_npu/csrc/aten/ops目录下，命名风格采用大驼峰，命名格式：<算子名\> + .cpp，如：AddKernelNpu.cpp。
3. 引入依赖头文件。
@@ -349,25 +350,25 @@ custom_autograd: # 自定义继承自Function的自定义API
> **说明：**
>工具的功能和使用方法，可查看头文件和源码获得。
-4. 定义API适配主体函数。
+4. 定义算子适配主体函数。
- 根据注册API开发中的分发函数确定自定义API适配主体函数。
+ 根据注册算子开发中的分发函数确定自定义算子适配主体函数。
5. 分别实现适配主体函数。
- 实现API适配主题函数，根据TBE算子原型构造得到对应的input、output、attr。
+ 实现算子适配主题函数，根据TBE算子原型构造得到对应的input、output、attr。
#### 示例
-以torch.add\(\)API为例介绍API适配开发过程。
+以torch.add\(\)算子为例介绍算子适配开发过程。
-1. 注册API。
+1. 注册算子。
- 将addAPI添加到npu_native_functions.yaml文件的对应位置，用于自动化声明和注册。
+ 将add算子添加到npu_native_functions.yaml文件的对应位置，用于自动化声明和注册。
```
- supported: #已支持的Pytorch Native Functions 对齐的API
+ supported: #已支持的Pytorch Native Functions 对齐的算子
add.Tensor
add_.Tensor
add.out
@@ -397,7 +398,7 @@ custom_autograd: # 自定义继承自Function的自定义API
>" CalcuOpUtil.h "中主要包含类型转换和判别的函数。
>" OpAdapter.h"文件中主要包含适配关联的头文件。
-4. 定义API适配主体函数。
+4. 定义算子适配主体函数。
```
at::Tensor NPUNativeFunctions::add(const at::Tensor &self, const at::Tensor &other, at::Scalar alpha)
@@ -409,7 +410,7 @@ custom_autograd: # 自定义继承自Function的自定义API
>  **说明：**
>
- > NPUNativeFunctions是API定义需要添加的命名空间约束。
+ > NPUNativeFunctions是算子定义需要添加的命名空间约束。
5. 分别实现适配主体函数。
@@ -596,7 +597,7 @@ pip3 install --upgrade torch_npu-1.8.1rc1-cp37-cp37m-linux_{arch}.whl
修改代码之后，需要重新执行“编译”和“安装”PyTorch插件过程。
-API功能验证
+算子功能验证
- **[概述](#概述-0md)**
@@ -607,13 +608,13 @@ pip3 install --upgrade torch_npu-1.8.1rc1-cp37-cp37m-linux_{arch}.whl
#### 简介
-完成API适配开发后，可通过运行适配昇腾处理器的PyTorchAPI，验证API运行结果是否正确。
+完成算子适配开发后，可通过运行适配昇腾处理器的PyTorch算子，验证算子运行结果是否正确。
-API功能验证会覆盖API开发的所有交付件，包含实现文件，API原型定义、API信息库以及API适配插件。本节仅对验证的方法做介绍。
+算子功能验证会覆盖算子开发的所有交付件，包含实现文件，算子原型定义、算子信息库以及算子适配插件。本节仅对验证的方法做介绍。
#### 测试用例及测试工具
-进行自定义API功能验证，通过PyTorch前端构造自定义API的函数并运行验证。
+进行自定义算子功能验证，通过PyTorch前端构造自定义算子的函数并运行验证。
在https://gitee.com/ascend/pytorch 中 "pytorch/test/test_network_ops"目录下提供了测试用例及测试工具，供用户参考。
@@ -621,7 +622,7 @@ API功能验证会覆盖API开发的所有交付件，包含实现文件，API
#### 简介
-本章通过具体例子，完成PyTorchAPI的功能测试步骤。
+本章通过具体例子，完成PyTorch算子的功能测试步骤。
#### 操作步骤
@@ -634,7 +635,7 @@ API功能验证会覆盖API开发的所有交付件，包含实现文件，API
${HOME}/Ascend/ascend-toolkit/set_env.sh
```
-2. 编写测试脚本。以addAPI为例，在“pytorch/test/test\_network\_ops“路径下编写测试脚本文件： test\_add.py。
+2. 编写测试脚本。以add算子为例，在“pytorch/test/test\_network\_ops“路径下编写测试脚本文件： test\_add.py。
```
# 引入依赖库
@@ -795,7 +796,7 @@ pip3.7 install torchvision --no-deps
#### 现象描述
-完成“自定义TBE算子”开发，和“PyTorch”适配开发，但执行测试用例，发现无法调用到新开发的API。
+完成“自定义TBE算子”开发，和“PyTorch”适配开发，但执行测试用例，发现无法调用到新开发的算子。
#### 可能原因
@@ -805,15 +806,15 @@ pip3.7 install torchvision --no-deps
#### 处理方法
-1. 参考“[API功能验证](#API功能验证md)”章节，完成运行环境设置，特别注意：
+1. 参考“[算子功能验证](#算子功能验证md)”章节，完成运行环境设置，特别注意：
```
. /home/HwHiAiUser/Ascend/ascend-toolkit/set_env.sh
```
-2. 检查yaml文件中对应API的分发配置，是否正确的完整分发；
+2. 检查yaml文件中对应算子的分发配置，是否正确的完整分发；
3. 分析排查代码实现，建议手段：
- 1. 修改"pytorch"中的API适配实现，让“test\_add.py”可以调用到“自定义API包”中的TBE算子；
+ 1. 修改"pytorch"中的算子适配实现，让“test\_add.py”可以调用到“自定义算子包”中的TBE算子；
"pytorch/aten/src/ATen/native/npu/AddKernelNpu.cpp"
@@ -829,44 +830,44 @@ pip3.7 install torchvision --no-deps
至此步骤，不应该有错误，屏幕应该输出 "add" 中增加的日志打印。若出错，请完成代码清理排查，保证无新开发的代码影响测试。
- 3. 将新开发的“自定义TBE算子”合并到"cann"中，在对应的在API入口增加日志打印，作为运行标识。
+ 3. 将新开发的“自定义TBE算子”合并到"cann"中，在对应的在算子入口增加日志打印，作为运行标识。
4. 完成上面的"cann"编译、安装，调用“python3.7.5 test\_add.py”进行测试。
> **说明：**
- >根据Ascend的设计逻辑，用户开发安装的“custom”API包优先级高于“built-in”的内置API包，在运行加载时，会优先加载调度“custom”包中的API。过程中，若解析“custom”中的API信息文件失败，则会跳过“custom”API包，不加载调度任何“custom”API包中的任何API。
+ >根据Ascend的设计逻辑，用户开发安装的“custom”算子包优先级高于“built-in”的内置算子包，在运行加载时，会优先加载调度“custom”包中的算子。过程中，若解析“custom”中的算子信息文件失败，则会跳过“custom”算子包，不加载调度任何“custom”算子包中的任何算子。
>
- >- 若此步骤出错，或屏幕未输出 "add" 中增加的日志打印，则说明新开发的“自定义TBE算子”有错误，影响了“自定义API包”的加载，建议**优先排查新开发的“自定义TBE算子”中的“API信息定义”是否正确**。
- >- 若此步骤正确，至少说明 **新开发的“自定义TBE算子”中的“API信息定义”不影响运行**。
+ >- 若此步骤出错，或屏幕未输出 "add" 中增加的日志打印，则说明新开发的“自定义TBE算子”有错误，影响了“自定义算子包”的加载，建议**优先排查新开发的“自定义TBE算子”中的“算子信息定义”是否正确**。
+ >- 若此步骤正确，至少说明 **新开发的“自定义TBE算子”中的“算子信息定义”不影响运行**。
5. 调用“python3.7.5 xxx\_testcase.py”进行测试；
> **说明：**
- >- 若屏幕正常输出新开发的“自定义TBE算子”中增加的日志打印，则至少说明调度到了新开发的API。
+ >- 若屏幕正常输出新开发的“自定义TBE算子”中增加的日志打印，则至少说明调度到了新开发的算子。
>- 若屏幕未输出新开发的“自定义TBE算子”中增加的日志打印，则问题可能出现在“PyTorch适配”中，需要排查这一部分的实现代码，较多的可能会出现在“XxxxKernelNpu.cpp”中的输入、输出未能正确适配。
如何确定“TBE算子”是否被“PyTorch适配”正确调用
-不管是“custom”API，还是“built-in”API，在安装后，都是以\*.py源码的方式存放在安装目录下，所以我们可以通过编辑源码，在API入口增加日志的方式，打印输出入参，确定输入的参数是否正确。
+不管是“custom”算子，还是“built-in”算子，在安装后，都是以\*.py源码的方式存放在安装目录下，所以我们可以通过编辑源码，在API入口增加日志的方式，打印输出入参，确定输入的参数是否正确。
> **注意：**
>该操作存在风险，建议在操作前备份计划修改的文件。若未备份，损坏后无法恢复，请及时联系支持人员。
-下面以"built-in"API中的"zn\_2\_nchw"API为例：
+下面以"built-in"算子中的"zn\_2\_nchw"算子为例：
-1. 打开安装在用户目录下的API包安装目录。
+1. 打开安装在用户目录下的算子包安装目录。
```
cd ~/.local/Ascend/opp/op_impl/built-in/ai_core/tbe/impl
ll
```
- 可以看到相应的API\*.py源码文件均为“只读”，即此时不可编辑。
+ 可以看到相应的算子\*.py源码文件均为“只读”，即此时不可编辑。
-2. 修改API\*.py源码文件属性，增加“可写”权限。
+2. 修改算子\*.py源码文件属性，增加“可写”权限。
```
sudo chmod +w zn_2_nchw.py
@@ -875,7 +876,7 @@ pip3.7 install torchvision --no-deps
-3. 打开API\*.py源码文件，增加日志，保存退出。
+3. 打开算子\*.py源码文件，增加日志，保存退出。
```
vi zn_2_nchw.py
@@ -886,8 +887,8 @@ pip3.7 install torchvision --no-deps
上面例子只加了个标识，实际调测时，可以增加打印输入参数信息。
4. 调用脚本执行测试用例，通过打印日志分析“输入参数信息”。
-5. 完成测试分析后，重新打开API\*.py源码文件，删除增加的日志，保存退出。
-6. 修改API\*.py源码文件属性，删除“可写”权限。
+5. 完成测试分析后，重新打开算子\*.py源码文件，删除增加的日志，保存退出。
+6. 修改算子\*.py源码文件属性，删除“可写”权限。
```
sudo chmod -w zn_2_nchw.py
@@ -906,7 +907,7 @@ PyTorch编译失败，提示“ error: ld returned 1 exit status”。
#### 可能原因
-通过日志分析，大概原因为XxxxKernelNpu.cpp中实现的适配函数，与PyTorch框架API要求的分发实现接口参数不匹配。在上面的例子中，是“binary\_cross\_entropy\_npu”，打开对应的XxxxKernelNpu.cpp文件，找到相应的适配函数。
+通过日志分析，大概原因为XxxxKernelNpu.cpp中实现的适配函数，与PyTorch框架算子要求的分发实现接口参数不匹配。在上面的例子中，是“binary\_cross\_entropy\_npu”，打开对应的XxxxKernelNpu.cpp文件，找到相应的适配函数。
@@ -944,7 +945,7 @@ PyTorch编译失败，提示“error: call of overload ....”。
- **[CMake安装方法](#CMake安装方法md)**
-- **[自定义API导出方法](#自定义API导出方法md)**
+- **[自定义算子导出方法](#自定义算子导出方法md)**
CMake安装方法
@@ -990,7 +991,7 @@ PyTorch编译失败，提示“error: call of overload ....”。
#### 简介
-PyTorch训练模型中包含自定义API，将自定义API导出成onnx单API模型，方便转移到其他AI框架中使用。自定义API导出有三种形式，适配NPU的TBE算子导出、C++API导出、纯PythonAPI导出。
+PyTorch训练模型中包含自定义算子，将自定义算子导出成onnx单算子模型，方便转移到其他AI框架中使用。自定义算子导出有三种形式，适配NPU的TBE算子导出、C++算子导出、纯Python算子导出。
#### 前提条件
@@ -1002,10 +1003,10 @@ TBE算子导出有两种方式。
方式一：
-1. 定义和注册API
+1. 定义和注册算子
```
- # 定义API
+ # 定义算子
@parse_args('v', 'v', 'f', 'i', 'i', 'i', 'i')
def symbolic_npu_roi_align(g, input, rois, spatial_scale, pooled_height, pooled_width, sample_num, roi_end_mode):
args = [input, rois]
@@ -1017,7 +1018,7 @@ TBE算子导出有两种方式。
return g.op('torch::npu_roi_align',*args, **kwargs)
- # 注册API
+ # 注册算子
import torch.onnx.symbolic_registry as sym_registry
def register_onnx_sym_npu_roi_align():
sym_registry.register_op('npu_roi_align', symbolic_npu_roi_align, '', 11)
@@ -1072,7 +1073,7 @@ TBE算子导出有两种方式。
1. 定义方法类
```
- # 实现API方法类及符号导出实现的方法
+ # 实现算子方法类及符号导出实现的方法
class CustomClassOp_Func_npu_roi_align(Function):
@staticmethod
def forward(ctx, input, rois, spatial_scale, pooled_height, pooled_width , sample_num, roi_end_mode):
@@ -1090,10 +1091,10 @@ TBE算子导出有两种方式。
return g.op('torch::npu_roi_align',*args, **kwargs)
```
-2. 自定义API模型
+2. 自定义算子模型
```
- # 实现API模型
+ # 实现算子模型
class NpuOp_npu_roi_align_Module(torch.nn.Module):
def __init__(self):
super(NpuOp_npu_roi_align_Module, self).__init__()
@@ -1151,15 +1152,15 @@ TBE算子导出有两种方式。
> **说明：**
>详细实现代码请参见[test\_custom\_ops\_npu\_demo.py](https://gitee.com/ascend/pytorch/blob/master/test/test_npu/test_onnx/torch.onnx/custom_ops_demo/test_custom_ops_npu_demo.py)，如无权限获取代码，请联系华为技术支持申请加入“Ascend”组织。
-#### C++API导出
+#### C++算子导出
-1. 自定义API
+1. 自定义算子
```
import torch
import torch_npu
import torch.utils.cpp_extension
- # 定义C++实现的API
+ # 定义C++实现的算子
def test_custom_add():
op_source = """
#include
@@ -1180,10 +1181,10 @@ TBE算子导出有两种方式。
test_custom_add()
```
-2. 注册自定义API
+2. 注册自定义算子
```
- # 定义API注册方法并注册API
+ # 定义算子注册方法并注册算子
from torch.onnx import register_custom_op_symbolic
def symbolic_custom_add(g, self, other):
@@ -1195,16 +1196,16 @@ TBE算子导出有两种方式。
3. 建立模型
```
- # 建立API模型
+ # 建立算子模型
class CustomAddModel(torch.nn.Module):
def forward(self, a, b):
return torch.ops.custom_namespace.custom_add(a, b)
```
-4. 导出APIonnx模型
+4. 导出算子onnx模型
```
- # 将API导出onnx模型
+ # 将算子导出onnx模型
def do_export(model, inputs, *args, **kwargs):
out = torch.onnx._export(model, inputs, "custom_demo.onnx", *args, **kwargs)
@@ -1218,9 +1219,9 @@ TBE算子导出有两种方式。
> **说明：**
>详细实现代码请参见[test\_custom\_ops\_demo.py](https://gitee.com/ascend/pytorch/blob/master/test/test_npu/test_onnx/torch.onnx/custom_ops_demo/test_custom_ops_demo.py)，如无权限获取代码，请联系华为技术支持申请加入“Ascend”组织。
-#### 纯PythonAPI导出
+#### 纯Python算子导出
-1. 自定义API
+1. 自定义算子
```
import torch
@@ -1237,7 +1238,7 @@ TBE算子导出有两种方式。
import math
from torch.nn import init
- # 定义API类方法
+ # 定义算子类方法
class CustomClassOp_Add_F(Function):
@staticmethod
def forward(ctx, input1,input2):
@@ -1257,7 +1258,7 @@ TBE算子导出有两种方式。
2. 建立模型
```
- # 注册API并建立模型
+ # 注册算子并建立模型
class CustomClassOp_Add(torch.nn.Module):
def __init__(self):
super(CustomClassOp_Add, self).__init__()
@@ -1283,7 +1284,7 @@ TBE算子导出有两种方式。
return rtn
```
-3. 导出APIonnx模型
+3. 导出算子onnx模型
```
ONNX_FILE_NAME = "./custom_python_module_demo.onnx"
@@ -1298,7 +1299,7 @@ TBE算子导出有两种方式。
output = model(input)
do_export(model, input, opset_version=11)
- # 将API导出到onnx模型
+ # 将算子导出到onnx模型
test_class_export()
```
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/PyTorch\346\241\206\346\236\266\344\270\213\347\232\204\347\256\227\345\255\220\351\200\202\351\205\215\346\265\201\347\250\213.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082048.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082056.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082064.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082072.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144082088.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001144241896.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081791.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190081803.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201935.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201951.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/figures/zh-cn_image_0000001190201973.png"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-caution.gif"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-danger.gif"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-note.gif"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-notice.gif"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-tip.gif"
diff --git "a/docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif" "b/docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif"
similarity index 100%
rename from "docs/zh/PyTorch API\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif"
rename to "docs/zh/PyTorch\347\256\227\345\255\220\345\274\200\345\217\221\346\214\207\345\215\227/public_sys-resources/icon-warning.gif"
diff --git a/scripts/codegen/gen_python_functions.py b/scripts/codegen/gen_python_functions.py
index 64337fd148ffeccc34cff6517f8fee10efd1ce88..40cd3e8e29895962b5173c949abb1140a9e643ee 100644
--- a/scripts/codegen/gen_python_functions.py
+++ b/scripts/codegen/gen_python_functions.py
@@ -218,21 +218,36 @@ def create_python_device_bindings(
) -> None:
"""Generates Python bindings to ATen functions"""
py_device_method_defs: List[str] = []
+ device_methods_def_py_dispatch: List[str] = []
grouped = group_filter_overloads(pairs, pred)
PY_DEVICE_METHOD_DEF = CodeTemplate("""\
- torch.${name} = torch_npu.${name}
+ torch.${name} = _${name}
+""")
+
+ PY_DEVICE_METHOD_DEF_DISPATCH = CodeTemplate("""\
+
+@torch_device_guard
+def _${name}(*args, **kwargs):
+ return torch_npu.${name}(*args, **kwargs)
+
""")
def method_device_def(name):
return PY_DEVICE_METHOD_DEF.substitute(name=name)
+ def method_device_def_dispatch(name):
+ return PY_DEVICE_METHOD_DEF_DISPATCH.substitute(name=name)
+
for name in sorted(grouped.keys(), key=lambda x: str(x)):
py_device_method_defs.append(method_device_def(name))
+ device_methods_def_py_dispatch.append(method_device_def_dispatch(name))
+
fm.write_with_template(filename, filename, lambda: {
'device_methods_def_py': py_device_method_defs,
+ 'device_methods_def_py_dispatch': device_methods_def_py_dispatch
})
def query_methods(filepath):
@@ -249,7 +264,6 @@ def create_python_device_bindings(
raise RuntimeError("In device methods file " +
str(fm.install_dir + filename) + " has multi-definition function.")
-
def load_signatures(
native_functions: List[NativeFunction],
*,
diff --git a/scripts/codegen/templates/torch_funcs.py b/scripts/codegen/templates/torch_funcs.py
index 33c9dfe371d3da48649ac6e3f515496ba2913714..d2bb566feda2af112b33a6a1ff72bdbd38e0c5fd 100644
--- a/scripts/codegen/templates/torch_funcs.py
+++ b/scripts/codegen/templates/torch_funcs.py
@@ -13,23 +13,70 @@
# See the License for the specific language governing permissions and
# limitations under the License.
-
import torch
import torch_npu
+from torch_npu.utils.device_guard import torch_device_guard, device
+
+
+@torch_device_guard
+def _device(*args, **kwargs):
+ new_device = torch_npu.new_device(*args, **kwargs)
+ if new_device.type == torch.npu.native_device:
+ return device(type=torch_npu.npu.npu_device, index=new_device.index)
+ return new_device
+
+
+@torch_device_guard
+def _tensor(*args, **kwargs):
+ return torch_npu.tensor(*args, **kwargs)
+
+
+@torch_device_guard
+def _full(*args, **kwargs):
+ return torch_npu.full(*args, **kwargs)
+
+
+@torch_device_guard
+def _randint(*args, **kwargs):
+ return torch_npu.randint(*args, **kwargs)
+
+
+@torch_device_guard
+def _range(*args, **kwargs):
+ return torch_npu.range(*args, **kwargs)
+
+
+@torch_device_guard
+def _arange(*args, **kwargs):
+ return torch_npu.arange(*args, **kwargs)
+
+
+@torch_device_guard
+def _empty_with_format(*args, **kwargs):
+ return torch_npu.empty_with_format(*args, **kwargs)
+
+
+@torch_device_guard
+def _npu_dropout_gen_mask(*args, **kwargs):
+ return torch_npu.npu_dropout_gen_mask(*args, **kwargs)
def jit_script(obj, optimize=None, _frames_up=0, _rcb=None):
# (Ascend) Disable extension of torch.jit.script
return obj
+
+${device_methods_def_py_dispatch}
+
def add_torch_funcs():
- torch.device = torch_npu.new_device
- torch.tensor = torch_npu.tensor
- torch.full = torch_npu.full
- torch.randint = torch_npu.randint
- torch.range = torch_npu.range
- torch.arange = torch_npu.arange
- torch.empty_with_format = torch_npu.empty_with_format
+ torch.device = _device
+ torch.tensor = _tensor
+ torch.full = _full
+ torch.randint = _randint
+ torch.range = _range
+ torch.arange = _arange
+ torch.empty_with_format = _empty_with_format
+ torch.npu_dropout_gen_mask = _npu_dropout_gen_mask
torch.jit.script = jit_script
${device_methods_def_py}
\ No newline at end of file
diff --git a/test/test_device.py b/test/test_device.py
new file mode 100644
index 0000000000000000000000000000000000000000..2d4c0f5975e109b80f24457a937ae8217c2be2c9
--- /dev/null
+++ b/test/test_device.py
@@ -0,0 +1,72 @@
+# Copyright (c) 2021, Huawei Technologies.All rights reserved.
+#
+# Licensed under the BSD 3-Clause License (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# https://opensource.org/licenses/BSD-3-Clause
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+import torch
+import torch_npu
+
+from torch_npu.testing.testcase import TestCase, run_tests
+
+class TestDevice(TestCase):
+ def device_monitor(func):
+ def wrapper(self, *args, **kwargs):
+ device_id = 0
+ torch.npu.set_device(device_id)
+ npu_device = torch.randn(2).npu(device_id).device
+ device_types = [
+ "npu",
+ "npu:"+str(device_id),
+ torch.device("npu:"+str(device_id)),
+ torch.device("npu:"+str(device_id)).type,
+ npu_device
+ ]
+ for device_type in device_types:
+ kwargs["device"] = device_type
+ npu_tensor = func(self, *args, **kwargs)
+ self.assertEqual(npu_tensor.device.type, "npu")
+ self.assertEqual(npu_tensor.device.index, device_id)
+ return wrapper
+
+
+ @device_monitor
+ def test_torch_tensor_to_device(self, device=None):
+ cpu_tensor = torch.randn(2, 3)
+ return cpu_tensor.to(device, torch.int64)
+
+
+ @device_monitor
+ def test_torch_tensor_new_empty_with_device_input(self, device=None):
+ npu_tensor = torch.ones(2, 3).to(device)
+ return npu_tensor.new_empty((2, 3), dtype=torch.float16, device=device)
+
+
+ @device_monitor
+ def test_torch_func_arange_with_device_input(self, device=None):
+ return torch.arange(5, dtype=torch.float32, device=device)
+
+
+ @device_monitor
+ def test_torch_func_zeros_with_device_input(self, device=None):
+ return torch.zeros((2, 3), dtype=torch.int8, device=device)
+
+
+ @device_monitor
+ def test_tensor_method_npu_with_device_input(self, device=None):
+ if isinstance(device, str):
+ device = torch.device(device)
+ cpu_input = torch.randn(2,3)
+ return cpu_input.npu(device)
+
+
+if __name__ == '__main__':
+ run_tests()
\ No newline at end of file
diff --git a/test/test_distributed/run_test.py b/test/test_distributed/run_test.py
index 8032a89b6ec96907d69fbfaa14b668a12ac78e61..2323f4d1f822badffb5d2b97852bb61cd8e92aa2 100644
--- a/test/test_distributed/run_test.py
+++ b/test/test_distributed/run_test.py
@@ -1,25 +1,25 @@
#!/usr/bin/env python
-from __future__ import print_function
import argparse
from datetime import datetime
-import modulefinder
import os
import shutil
import signal
-import subprocess
import sys
import tempfile
import time
import unittest
import torch
-import torch_npu
import torch._six
-from torch.utils import cpp_extension
-from torch.testing._internal.common_utils import TEST_WITH_ROCM, shell
import torch.distributed as dist
+
+from torch.testing._internal.common_utils import TEST_WITH_ROCM, shell
+
+import torch_npu
+
+
PY2 = sys.version_info <= (3,)
PY33 = sys.version_info >= (3, 3)
PY36 = sys.version_info >= (3, 6)
diff --git a/test/test_network_ops/test_all.py b/test/test_network_ops/test_all.py
index 33bcd3d5744821b2294d7fefabe8562ed0a9bed1..d9de4a6001bcd660b8d67472420e660ddbbdb0bf 100644
--- a/test/test_network_ops/test_all.py
+++ b/test/test_network_ops/test_all.py
@@ -37,7 +37,7 @@ class TestAll(TestCase):
output = output.numpy()
return output
- def test_all_shape_format(self, device="npu"):
+ def test_all_shape_format(self):
shape_list = [[1024], [32, 1024], [32, 8, 1024], [128, 32, 8, 1024], [2, 0, 2]]
for item in shape_list:
cpu_input, npu_input = self.create_bool_tensor(item, 0, 1)
@@ -71,7 +71,7 @@ class TestAll(TestCase):
output1 = output1.to("cpu").numpy()
return output0, output1
- def test_alld_shape_format(self, device="npu"):
+ def test_alld_shape_format(self):
shape_list = [[1024], [32, 1024], [32, 8, 1024], [128, 32, 8, 1024]]
for item in shape_list:
cpu_input, npu_input = self.create_bool_tensor(item, 0, 1)
@@ -82,6 +82,15 @@ class TestAll(TestCase):
self.assertRtolEqual(cpu_output.astype(np.int32), npu_out0.astype(np.int32))
self.assertRtolEqual(cpu_output.astype(np.int32), npu_out1.astype(np.int32))
+ def test_all_tensor_numel_0(self):
+ ca = torch.rand(1, 2, 0, 3, 4).bool()
+ na = ca.npu()
+ cout = ca.all(2)
+ nout = na.all(2)
+ cout = cout.numpy()
+ nout = nout.to("cpu").numpy()
+ self.assertRtolEqual(cout, nout)
+
if __name__ == "__main__":
run_tests()
\ No newline at end of file
diff --git a/test/test_network_ops/test_unique_consecutive.py b/test/test_network_ops/test_unique_consecutive.py
new file mode 100644
index 0000000000000000000000000000000000000000..e1f4e04077926972c5e3a9102d8b844d86d7b3ab
--- /dev/null
+++ b/test/test_network_ops/test_unique_consecutive.py
@@ -0,0 +1,65 @@
+# Copyright (c) 2020 Huawei Technologies Co., Ltd
+# Copyright (c) 2019, Facebook CORPORATION.
+# All rights reserved.
+#
+# Licensed under the BSD 3-Clause License (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# https://opensource.org/licenses/BSD-3-Clause
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+import torch
+import numpy as np
+import torch_npu
+
+from torch_npu.testing.testcase import TestCase, run_tests
+
+class TestUniqueConsecutive(TestCase):
+ def test_unique_consecutive(self, device="npu"):
+ shape_format = [
+ [[torch.int32, (2, 3)], 0],
+ [[torch.long, (2, 3)], 1],
+ [[torch.float32, (2, 3)], 0],
+ [[torch.float16, (2, 3)], 1],
+ [[torch.int32, (2, 3)], None],
+ [[torch.long, (2, 3)], None],
+ [[torch.float32, (2, 3)], None],
+ [[torch.float16, (2, 3)], None]
+ ]
+
+ for item in shape_format:
+ cpu_input = torch.rand(item[0][1]).random_(0, 3).to(item[0][0])
+ npu_input = cpu_input.npu()
+ if item[0][0] == torch.float16:
+ cpu_input = cpu_input.float()
+
+ cpu_output, cpu_idx, cpu_counts = torch.unique_consecutive(cpu_input, return_inverse=True,
+ return_counts=True, dim=item[1])
+ npu_output, npu_idx, npu_counts = torch.unique_consecutive(npu_input, return_inverse=True,
+ return_counts=True, dim=item[1])
+
+ if item[0][0] == torch.float16:
+ cpu_output = cpu_output.half()
+ self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())
+ self.assertRtolEqual(cpu_idx.numpy(), npu_idx.cpu().numpy())
+ self.assertRtolEqual(cpu_counts.numpy(), npu_counts.cpu().numpy())
+
+ def test_unique_consecutive_case_in_dino(self, device="npu"):
+ input_list = [
+ torch.tensor([224, 224, 96, 96, 96, 96, 96, 96, 96, 96]),
+ torch.tensor([224, 224])
+ ]
+ for i in input_list:
+ cpu_output, cpu_counts = torch.unique_consecutive(i, return_counts=True)
+ npu_output, npu_counts = torch.unique_consecutive(i.npu(), return_counts=True)
+ self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())
+ self.assertRtolEqual(cpu_counts.numpy(), npu_counts.cpu().numpy())
+
+if __name__ == "__main__":
+ run_tests()
\ No newline at end of file
diff --git a/third_party/acl/inc/acl/acl_op_compiler.h b/third_party/acl/inc/acl/acl_op_compiler.h
index 0807efbd8ed85ed77531ee78904e6ba9515f4103..338ec40e4c8bf0c71cb49b018eb1e95bf8e0fc82 100644
--- a/third_party/acl/inc/acl/acl_op_compiler.h
+++ b/third_party/acl/inc/acl/acl_op_compiler.h
@@ -93,6 +93,12 @@ ACL_FUNC_VISIBILITY aclError aclopCompileAndExecute(const char *opType,
const aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag,
const char *opPath, aclrtStream stream);
+ACL_FUNC_VISIBILITY aclError aclopCompileAndExecuteV2(const char *opType,
+ int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[],
+ int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[],
+ aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag,
+ const char *opPath, aclrtStream stream);
+
/**
* @ingroup AscendCL
* @brief set compile option
diff --git a/third_party/acl/libs/acl_op_compiler.cpp b/third_party/acl/libs/acl_op_compiler.cpp
index ea5feee96822a19c6482e48e9a57c5a0d05a2385..aaa4886cd7194625e27fb10f96f694e6e8f24516 100644
--- a/third_party/acl/libs/acl_op_compiler.cpp
+++ b/third_party/acl/libs/acl_op_compiler.cpp
@@ -43,6 +43,22 @@ aclError aclopCompileAndExecute(
return 0;
}
+aclError aclopCompileAndExecuteV2(
+ const char* opType,
+ int numInputs,
+ aclTensorDesc *inputDesc[],
+ aclDataBuffer *inputs[],
+ int numOutputs,
+ aclTensorDesc *outputDesc[],
+ aclDataBuffer *outputs[],
+ aclopAttr* attr,
+ aclopEngineType engineType,
+ aclopCompileType compileFlag,
+ const char* opPath,
+ aclrtStream stream) {
+ return 0;
+}
+
void GeGeneratorFinalize() {
return;
}
diff --git a/torch_npu/csrc/aten/ops/AllKernelNpu.cpp b/torch_npu/csrc/aten/ops/AllKernelNpu.cpp
index d303b5cddbdbca2acf8658e61c5ab6e06e81a873..d11a357364c9862d62f43cd8006c8c5b67a99536 100644
--- a/torch_npu/csrc/aten/ops/AllKernelNpu.cpp
+++ b/torch_npu/csrc/aten/ops/AllKernelNpu.cpp
@@ -62,8 +62,19 @@ at::Tensor& NPUNativeFunctions::all_out(
at::Tensor NPUNativeFunctions::all(const at::Tensor& self, int64_t dim, bool keepdim) {
TORCH_CHECK(self.scalar_type() == at::ScalarType::Bool || self.scalar_type() == at::ScalarType::Byte,
"all only supports torch.uint8 and torch.bool dtypes");
+ TORCH_CHECK(dim >= -(self.dim()) && dim < self.dim(),
+ "The value of dim must be greater than or equal to -self.dim() and less than self.dim()");
if (self.numel() == 0) {
- at::Tensor res = OpPreparation::ApplyTensor({}, self.options().dtype(at::kInt), self).fill_(1).to(at::ScalarType::Bool);
+ c10::SmallVector outputSize;
+ for(int64_t i = 0; i < self.dim(); i++){
+ if(dim != i){
+ outputSize.emplace_back(self.size(i));
+ }
+ }
+ at::Tensor res = OpPreparation::ApplyTensor(
+ outputSize,
+ self.options().dtype(at::kInt),
+ self).fill_(1).to(at::ScalarType::Bool);
return res;
}
diff --git a/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp b/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp
index 380c8d63f9ce5b98e6d7d53c21bd20e2d3e10bc6..8f03431a8c4464fff660bfca146af133fc30472e 100644
--- a/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp
+++ b/torch_npu/csrc/aten/ops/CdistKernelNpu.cpp
@@ -33,7 +33,7 @@ at::Tensor NPUNativeFunctions::_cdist_forward(
// Since double is not supported in NPU, the type of P needs to be converted from double to float.
float p_float;
if (std::isinf(p)) {
- p_float = std::numeric_limits::infinity();
+ p_float = -1;
}
else {
TORCH_CHECK(p <= std::numeric_limits::max(), "npu dose not support float64" );
diff --git a/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp b/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..de1eba533a5d441b983de4f6eb61409db2057c71
--- /dev/null
+++ b/torch_npu/csrc/aten/ops/UniqueConsecutiveKernelNpu.cpp
@@ -0,0 +1,74 @@
+// Copyright (c) 2020 Huawei Technologies Co., Ltd
+// All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "torch_npu/csrc/framework/utils/OpAdapter.h"
+#include "torch_npu/csrc/aten/NPUNativeFunctions.h"
+
+namespace at_npu {
+namespace native {
+
+std::tuple unique_consecutive_out_npu(
+ at::Tensor& output,
+ at::Tensor& inverse_indices,
+ at::Tensor& counts,
+ const at::Tensor& self,
+ const bool return_inverse,
+ const bool return_counts,
+ c10::optional dim) {
+ at::Tensor selfCopy = self;
+ if (self.scalar_type() == at::ScalarType::Half) {
+ selfCopy = NPUNativeFunctions::npu_dtype_cast(self, at::ScalarType::Float);
+ output = NPUNativeFunctions::npu_dtype_cast(output, at::ScalarType::Float);
+ }
+ c10::SmallVector output_sync_idx = {0, 2};
+ OpCommand cmd;
+ cmd.Sync(output_sync_idx)
+ .Name("UniqueConsecutive")
+ .Input(selfCopy)
+ .Output(output)
+ .Output(inverse_indices)
+ .Output(counts)
+ .Attr("return_idx", return_inverse)
+ .Attr("return_counts", return_counts);
+ if (dim.has_value()) {
+ cmd.Attr("axis", dim.value());
+ }
+ cmd.Run();
+ if (self.scalar_type() == at::ScalarType::Half) {
+ output = NPUNativeFunctions::npu_dtype_cast(output, at::ScalarType::Half);
+ }
+ return std::tie(output, inverse_indices, counts);
+}
+
+std::tuple NPUNativeFunctions::unique_consecutive(
+ const at::Tensor& self,
+ const bool return_inverse,
+ const bool return_counts,
+ c10::optional dim) {
+ at::Tensor output = (dim.has_value()) ?
+ OpPreparation::ApplyTensor(self) : OpPreparation::ApplyTensor(self, {self.numel()});
+ at::Tensor inverse_indices = (dim.has_value()) ?
+ OpPreparation::ApplyTensorWithFormat(self.size(dim.value()), self.options().dtype(at::kLong), ACL_FORMAT_ND) :
+ OpPreparation::ApplyTensorWithFormat(self.sizes(), self.options().dtype(at::kLong), ACL_FORMAT_ND);
+ at::Tensor counts = (dim.has_value()) ?
+ OpPreparation::ApplyTensorWithFormat(self.size(dim.value()), self.options().dtype(at::kLong), ACL_FORMAT_ND) :
+ OpPreparation::ApplyTensorWithFormat({self.numel()}, self.options().dtype(at::kLong), ACL_FORMAT_ND);
+ unique_consecutive_out_npu(output, inverse_indices, counts, self, return_inverse, return_counts, dim);
+ return std::tie(output, inverse_indices, counts);
+}
+
+
+} // namespace native
+} // namespace at
\ No newline at end of file
diff --git a/torch_npu/csrc/core/npu/NPUGuard.h b/torch_npu/csrc/core/npu/NPUGuard.h
index 84e81482bfc2b4263121b10e67af5bb437bbd141..cbbb774972b261603eb535dcfcd54d04e661f7de 100644
--- a/torch_npu/csrc/core/npu/NPUGuard.h
+++ b/torch_npu/csrc/core/npu/NPUGuard.h
@@ -197,7 +197,7 @@ struct NPUStreamGuard {
return guard_.original_device();
}
- private:
+private:
c10::impl::InlineStreamGuard guard_;
};
diff --git a/torch_npu/csrc/framework/OpCommand.cpp b/torch_npu/csrc/framework/OpCommand.cpp
index 9077b7465ed70ac258219c74523efe3dcab54fea..f818a73c6be69c42754fab7b13453eb172f97498 100644
--- a/torch_npu/csrc/framework/OpCommand.cpp
+++ b/torch_npu/csrc/framework/OpCommand.cpp
@@ -12,6 +12,7 @@
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
+
#include
#include "torch_npu/csrc/framework/OpCommand.h"
#include "torch_npu/csrc/core/npu/register/OptionsManager.h"
@@ -127,26 +128,35 @@ OpCommand& OpCommand::Output(
if (!resultTypeDefined && commonType.has_value() &&
output.scalar_type() != commonType.value()) {
output = NPUNativeFunctions::npu_dtype_cast(output, commonType.value());
- }
+ }
)
+ outputTensor.emplace_back(output);
return AddOutput(output, realType);
}
void OpCommand::Run() {
IF_GRAPH_MODE_THEN_RUN(return;)
- if (c10_npu::option::OptionsManager::CheckQueueEnable()) {
+ if (c10_npu::option::OptionsManager::CheckQueueEnable() && !sync) {
ExecuteParas execParams;
aclCmd->ExportParams(execParams);
c10_npu::queue::QueueParas params(c10_npu::queue::COMPILE_AND_EXECUTE, sizeof(ExecuteParas), &execParams);
c10_npu::enCurrentNPUStream(¶ms);
aclCmd->releaseSource(false);
} else {
- aclCmd->Run();
+ aclCmd->Run(sync, sync_index, outputTensor);
aclCmd->releaseSource();
- }
+ }
aclCmds->Pop();
}
+OpCommand& OpCommand::Sync(c10::SmallVector &index) {
+ sync_index = index;
+ if (!index.empty()) {
+ sync = true;
+ }
+ return *this;
+}
+
OpCommand& OpCommand::AddTensorInput(at::Tensor &tensor,
at::ScalarType forceScaleType,
const string &descName,
diff --git a/torch_npu/csrc/framework/OpCommand.h b/torch_npu/csrc/framework/OpCommand.h
index 2ec73f62cb6fc7fa8d8b3ee30b4a5c05b4cf55d2..a7e6eaa76bf4444ec6671dafb1cc38033fabab9b 100644
--- a/torch_npu/csrc/framework/OpCommand.h
+++ b/torch_npu/csrc/framework/OpCommand.h
@@ -117,6 +117,8 @@ public:
// Run a single op
void Run();
+ OpCommand& Sync(c10::SmallVector &sync_index);
+
private:
OpCommand& AddTensorInput(at::Tensor &tensor,
at::ScalarType forceScaleType = at::ScalarType::Undefined,
@@ -150,6 +152,10 @@ private:
c10::optional commonType = c10::nullopt;
c10::optional commonShape = c10::nullopt;
bool resultTypeDefined = false;
+ bool sync = false;
+ c10::SmallVector sync_index;
+ c10::SmallVector outputTensor;
+
}; // class OpCommand
} // namespace native
} // namespace at_npu
diff --git a/torch_npu/csrc/framework/OpParamMaker.cpp b/torch_npu/csrc/framework/OpParamMaker.cpp
index 5d20035d48d4d26e45261211ec6cbc72e46e2939..a84e33201fef2a6753fc4f4350eeebf081ae0bad 100644
--- a/torch_npu/csrc/framework/OpParamMaker.cpp
+++ b/torch_npu/csrc/framework/OpParamMaker.cpp
@@ -107,22 +107,28 @@ namespace at_npu
attrValue.data());
}
- void OpCommandImpl::Run()
- {
+ void OpCommandImpl::Run(
+ bool sync,
+ c10::SmallVector &sync_index,
+ c10::SmallVector &outputTensor) {
NPU_LOGD("Op %s Run.", opName.c_str());
RECORD_FUNCTION(opName, std::vector({}));
if (PyGILState_Check()) {
// we need to release GIL for NPU to compile op.
Py_BEGIN_ALLOW_THREADS
- ACL_REQUIRE_OK_OP(InnerRun(opName, execParam), opName.c_str());
+ ACL_REQUIRE_OK_OP(InnerRun(opName, execParam, sync, sync_index, outputTensor), opName.c_str());
Py_END_ALLOW_THREADS
} else {
- ACL_REQUIRE_OK_OP(InnerRun(opName, execParam), opName.c_str());
+ ACL_REQUIRE_OK_OP(InnerRun(opName, execParam, sync, sync_index, outputTensor), opName.c_str());
}
}
- aclError OpCommandImpl::InnerRun(string name, AclExecParam ¶ms)
- {
+ aclError OpCommandImpl::InnerRun(
+ string name,
+ AclExecParam ¶ms,
+ bool sync,
+ c10::SmallVector &sync_index,
+ c10::SmallVector &outputTensor) {
auto stream = c10_npu::getCurrentNPUStream();
auto inputSize = params.inBuffer.size();
auto outputSize = params.outBuffer.size();
@@ -155,19 +161,45 @@ namespace at_npu
TORCH_CHECK(false, "In aoe mode, AclGenGraphAndDumpForOp failed!");
}
}
- ret = aclopCompileAndExecute(
- name.c_str(),
- inputSize,
- params.inDesc.data(),
- params.inBuffer.data(),
- outputSize,
- params.outDesc.data(),
- params.outBuffer.data(),
- params.attr,
- ACL_ENGINE_SYS,
- ACL_COMPILE_SYS,
- NULL,
- stream);
+ if (!sync) {
+ ret = aclopCompileAndExecute(
+ name.c_str(),
+ inputSize,
+ params.inDesc.data(),
+ params.inBuffer.data(),
+ outputSize,
+ params.outDesc.data(),
+ params.outBuffer.data(),
+ params.attr,
+ ACL_ENGINE_SYS,
+ ACL_COMPILE_SYS,
+ NULL,
+ stream);
+ } else {
+ int64_t dimSize;
+ ret = AclopCompileAndExecuteV2(
+ name.c_str(),
+ inputSize,
+ const_cast(params.inDesc.data()),
+ const_cast(params.inBuffer.data()),
+ outputSize,
+ const_cast(params.outDesc.data()),
+ params.outBuffer.data(),
+ params.attr,
+ ACL_ENGINE_SYS,
+ ACL_COMPILE_SYS,
+ NULL,
+ stream);
+
+ for (size_t i = 0; i < sync_index.size(); i++) {
+ c10::SmallVector real_shape;
+ for (int64_t j = 0; j < outputTensor[sync_index[i]].dim(); j++) {
+ C10_NPU_CHECK(aclGetTensorDescDimV2(params.outDesc[sync_index[i]], j, &dimSize));
+ real_shape.emplace_back(dimSize);
+ }
+ outputTensor[sync_index[i]].resize_(real_shape);
+ }
+ }
++index;
} while (NpuUtils::IsOomError(ret, index) && (index < NPU_MAX_OP_EXEC_TRY_NUM));
if (reset_flag)
diff --git a/torch_npu/csrc/framework/OpParamMaker.h b/torch_npu/csrc/framework/OpParamMaker.h
index d1b02a0dea19c20a27abbea53681926494553ace..58b12eba0a9eda15c99136281830b8614aef1920 100644
--- a/torch_npu/csrc/framework/OpParamMaker.h
+++ b/torch_npu/csrc/framework/OpParamMaker.h
@@ -303,7 +303,7 @@ namespace at_npu
}
}
- void Run();
+ void Run(bool sync, c10::SmallVector &sync_index, c10::SmallVector &outputTensor);
void releaseSource(bool no_blocking = true)
{
@@ -364,7 +364,13 @@ namespace at_npu
}
}
- aclError InnerRun(string name, AclExecParam ¶ms);
+ aclError InnerRun(
+ string name,
+ AclExecParam ¶ms,
+ bool sync,
+ c10::SmallVector &sync_index,
+ c10::SmallVector &outputTensor
+ );
private:
string opName;
diff --git a/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp b/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp
index dd464b3a68c873e4a0d7b97a0aaa2b9e16aec45b..c8c67af34a1c96feaa00c9c354035f057537bf46 100644
--- a/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp
+++ b/torch_npu/csrc/framework/interface/AclOpCompileInterface.cpp
@@ -35,6 +35,7 @@ namespace at_npu
LOAD_FUNCTION(aclGenGraphAndDumpForOp)
LOAD_FUNCTION(aclCreateGraphDumpOpt)
LOAD_FUNCTION(aclDestroyGraphDumpOpt)
+ LOAD_FUNCTION(aclopCompileAndExecuteV2)
aclError AclopSetCompileFlag(aclOpCompileFlag flag) {
typedef aclError (*aclopSetCompileFlagFunc)(aclOpCompileFlag);
@@ -87,5 +88,25 @@ aclError AclDestroyGraphDumpOpt(aclGraphDumpOption* aclGraphDumpOpt) {
return func(aclGraphDumpOpt);
}
+aclError AclopCompileAndExecuteV2(const char *opType,
+ int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[],
+ int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[],
+ aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag,
+ const char *opPath, aclrtStream stream) {
+ typedef aclError(*AclopCompileAndExecuteV2Func)(const char *,
+ int, aclTensorDesc * [], aclDataBuffer * [],
+ int, aclTensorDesc * [], aclDataBuffer * [],
+ aclopAttr *, aclopEngineType, aclopCompileType,
+ const char *, aclrtStream);
+ static AclopCompileAndExecuteV2Func func = nullptr;
+ if (func == nullptr) {
+ func = (AclopCompileAndExecuteV2Func)GET_FUNC(aclopCompileAndExecuteV2);
+ }
+ TORCH_CHECK(func, "Failed to find function ", "aclopCompileAndExecuteV2");
+ auto ret = func(opType, numInputs, inputDesc, inputs, numOutputs,
+ outputDesc, outputs, attr, engineType, compileFlag, opPath, stream);
+ return ret;
+}
+
} // namespace native
} // namespace at_npu
\ No newline at end of file
diff --git a/torch_npu/csrc/framework/interface/AclOpCompileInterface.h b/torch_npu/csrc/framework/interface/AclOpCompileInterface.h
index f59f7fe841cf652a82b50755a34378e2b7c730e4..85f783dd3225a2a6d388f7e57288138a02bc043f 100644
--- a/torch_npu/csrc/framework/interface/AclOpCompileInterface.h
+++ b/torch_npu/csrc/framework/interface/AclOpCompileInterface.h
@@ -72,6 +72,33 @@ aclGraphDumpOption* AclCreateGraphDumpOpt();
*/
aclError AclDestroyGraphDumpOpt(aclGraphDumpOption* aclGraphDumpOpt);
+/**
+ * @ingroup AscendCL
+ * @brief compile and execute op
+ *
+ * @param opType [IN] op type
+ * @param numInputs [IN] number of inputs
+ * @param inputDesc [IN] pointer to array of input tensor descriptions
+ * @param inputs [IN] pointer to array of input buffers
+ * @param numOutputs [IN] number of outputs
+ * @param outputDesc [IN] pointer to array of output tensor descriptions
+ * @param outputs [IN] pointer to array of outputs buffers
+ * @param attr [IN] pointer to instance of aclopAttr.
+ * may pass nullptr if the op has no attribute
+ * @param engineType [IN] engine type
+ * @param compileFlag [IN] compile flag
+ * @param opPath [IN] path of op
+ * @param stream [IN] stream handle
+ * @retval ACL_ERROR_NONE The function is successfully executed.
+ * @retval OtherValues Failure
+ */
+aclError AclopCompileAndExecuteV2(const char *opType,
+ int numInputs, aclTensorDesc *inputDesc[], aclDataBuffer *inputs[],
+ int numOutputs, aclTensorDesc *outputDesc[], aclDataBuffer *outputs[],
+ aclopAttr *attr, aclopEngineType engineType, aclopCompileType compileFlag,
+ const char *opPath, aclrtStream stream);
+
+
} // namespace native
} // namespace at_npu
diff --git a/torch_npu/npu/__init__.py b/torch_npu/npu/__init__.py
index 957a8d77799ae7af3e0e0c1929e2f84be42d2b3f..3b6a8dd890a59cdc35835794c2de5bc3c47684be 100644
--- a/torch_npu/npu/__init__.py
+++ b/torch_npu/npu/__init__.py
@@ -15,7 +15,7 @@
# limitations under the License.
__all__ = [
- "native_device", "is_initialized", "_lazy_call", "_lazy_init", "init", "set_dump",
+ "native_device", "npu_device", "is_initialized", "_lazy_call", "_lazy_init", "init", "set_dump",
"synchronize", "device_count", "set_device", "current_device",
"_get_device_index", "is_available", "device", "device_of",
"stream", "current_stream", "default_stream", "init_dump",
@@ -37,6 +37,7 @@ __all__ = [
import torch
from .device import __device__ as native_device
+from .device import __npu_device__ as npu_device
from .utils import (is_initialized, _lazy_call, _lazy_init, init, set_dump,
synchronize, device_count, set_device, current_device,
_get_device_index, is_available, device, device_of,
diff --git a/torch_npu/npu/device.py b/torch_npu/npu/device.py
index a9607e5f30f6b918271b002ddbdcfbb6c022f367..fbfa35d8228f92b5d7ef0154d09f4315463e3f4a 100644
--- a/torch_npu/npu/device.py
+++ b/torch_npu/npu/device.py
@@ -13,4 +13,5 @@
# See the License for the specific language governing permissions and
# limitations under the License.
-__device__ = 'xla'
\ No newline at end of file
+__device__ = 'xla'
+__npu_device__ = 'npu'
\ No newline at end of file
diff --git a/torch_npu/utils/_tensor_str.py b/torch_npu/utils/_tensor_str.py
index f79ee013833009ca80ee178fb475c4f33d6d8a90..584111668108ee65dcc751df74b15aa05e6a9976 100644
--- a/torch_npu/utils/_tensor_str.py
+++ b/torch_npu/utils/_tensor_str.py
@@ -14,10 +14,14 @@
# limitations under the License.
import torch
+
from torch._tensor_str import _Formatter as SrcFormatter
from torch._tensor_str import PRINT_OPTS, _tensor_str_with_formatter, _add_suffixes, get_summarized_data
from torch.overrides import has_torch_function_unary, handle_torch_function
+import torch_npu
+
+
class _Formatter(SrcFormatter):
def __init__(self, tensor):
self.floating_dtype = tensor.dtype.is_floating_point
@@ -60,8 +64,8 @@ def _str_intern(inp):
self, tangent = torch.autograd.forward_ad.unpack_dual(inp)
if self.device.type != torch._C._get_default_device()\
- or (self.device.type == 'npu' and torch.npu.current_device() != self.device.index):
- suffixes.append('device=\'' + str(self.device) + '\'')
+ or (self.device.type == torch_npu.npu.native_device and torch.npu.current_device() != self.device.index):
+ suffixes.append('device=\'' + str(torch_npu.npu.npu_device) + ':'+ str(torch.npu.current_device())+'\'')
_default_complex_dtype = torch.cdouble if torch.get_default_dtype() == torch.double else torch.cfloat
has_default_dtype = self.dtype in (torch.get_default_dtype(), _default_complex_dtype, torch.int64, torch.bool)
diff --git a/torch_npu/utils/device_guard.py b/torch_npu/utils/device_guard.py
new file mode 100644
index 0000000000000000000000000000000000000000..997239c9c6b64933a563a624420a47d1991baa31
--- /dev/null
+++ b/torch_npu/utils/device_guard.py
@@ -0,0 +1,36 @@
+# Copyright (c) 2020 Huawei Technologies Co., Ltd
+# All rights reserved.
+#
+# Licensed under the BSD 3-Clause License (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# https://opensource.org/licenses/BSD-3-Clause
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+import collections
+import torch_npu
+
+device = collections.namedtuple('device', ['type', 'index'])
+
+def torch_device_guard(func):
+ # Parse args/kwargs from namedtuple(torch.device) to matched torch.device objects
+ def wrapper(*args, **kwargs):
+ if args:
+ args_list = list(args)
+ for index, arg in enumerate(args_list):
+ if isinstance(arg, tuple) and "type='npu'" in str(arg):
+ args_list[index] = torch_npu.new_device(type=torch_npu.npu.native_device, index=arg.index)
+ break
+ args = tuple(args_list)
+ if kwargs and isinstance(kwargs.get("device"), tuple):
+ namedtuple_device = kwargs.get("device")
+ if "type='npu'" in str(namedtuple_device):
+ kwargs['device'] = torch_npu.new_device(type=torch_npu.npu.native_device, index=namedtuple_device.index)
+ return func(*args, **kwargs)
+ return wrapper
\ No newline at end of file
diff --git a/torch_npu/utils/module.py b/torch_npu/utils/module.py
index 5c2f210d3649be6082e491585ec53682dc742094..41495dad803568ad58b0d8d6cf92451ae605a42b 100644
--- a/torch_npu/utils/module.py
+++ b/torch_npu/utils/module.py
@@ -19,6 +19,7 @@ import logging
import torch
from torch.nn.modules._functions import SyncBatchNorm as sync_batch_norm
+from torch_npu.utils.tensor_methods import torch_device_guard
import torch_npu
@@ -50,7 +51,7 @@ def npu(self, device=None):
torch_npu.npu.enable_graph_mode()
return self._apply(lambda t: t.npu(device))
-
+@torch_device_guard
def to(self, *args, **kwargs):
if args and isinstance(args[0], str) and 'npu' in args[0]:
args = tuple([list(args)[0].replace('npu', torch_npu.npu.native_device)])
@@ -173,7 +174,8 @@ def ddp_forward(self, *inputs, **kwargs):
if self.ddp_uneven_inputs_config.ddp_join_enabled:
# Notify joined ranks whether they should sync in backwards pass or not.
self._check_global_requires_backward_grad_sync(is_joined_rank=False)
- if self.device_ids and self.device_type != torch_npu.npu.native_device:
+ # Note: module.device_type was builded from device.type("npu") inside Class Module
+ if self.device_ids and self.device_type != torch_npu.npu.npu_device:
if len(self.device_ids) == 1:
inputs, kwargs = self.to_kwargs(inputs, kwargs, self.device_ids[0])
output = self.module(*inputs[0], **kwargs[0])
diff --git a/torch_npu/utils/tensor_methods.py b/torch_npu/utils/tensor_methods.py
index b427068b9757851982108f51a6cc0953869fd7c3..916483b47726d6f5d650a7739858fdead1e7c6e0 100644
--- a/torch_npu/utils/tensor_methods.py
+++ b/torch_npu/utils/tensor_methods.py
@@ -17,6 +17,8 @@
import warnings
import torch
import torch_npu
+from torch_npu.utils.device_guard import torch_device_guard, device
+
warnings.filterwarnings(action="once")
warning_str = "The tensor methods of custom operators would cause performance drop." + \
@@ -58,13 +60,12 @@ def npu_confusion_transpose(self, perm, shape, transpose_first):
return torch_npu.npu_confusion_transpose(self, perm, shape, transpose_first)
+@torch_device_guard
def _npu(self, *args, **kwargs):
- warnings.warn(warning_str.format("npu"))
return torch_npu._C.npu(self, *args, **kwargs)
@property
def _is_npu(self):
- warnings.warn(warning_str.format("is_npu"))
return torch_npu._C.is_npu(self)
@@ -72,6 +73,7 @@ def _type(self, *args, **kwargs):
return torch_npu._C.type(self, *args, **kwargs)
+@torch_device_guard
def _to(self, *args, **kwargs):
return torch_npu._C.to(self, *args, **kwargs)
@@ -98,6 +100,7 @@ def _storage(self):
return storage_impl(self)
+@torch_device_guard
def _new_empty(self, *args, **kwargs):
if isinstance(args[0], int):
list_args = list(args)
@@ -106,10 +109,18 @@ def _new_empty(self, *args, **kwargs):
return torch_npu._C.new_empty(self, *args, **kwargs)
+@torch_device_guard
def _new_empty_strided(self, *args, **kwargs):
return torch_npu._C.new_empty_strided(self, *args, **kwargs)
+@property
+def _device(self):
+ if torch_npu._C.is_npu(self):
+ return device(type='npu', index=self.get_device())
+ return torch.device("cpu")
+
+
def add_tensor_methods():
torch.Tensor.npu_format_cast_ = npu_format_cast_
torch.Tensor.npu_format_cast = npu_format_cast
@@ -126,3 +137,4 @@ def add_tensor_methods():
torch.Tensor.storage = _storage
torch.Tensor.new_empty = _new_empty
torch.Tensor.new_empty_strided = _new_empty_strided
+ torch.Tensor.device = _device
\ No newline at end of file