!11524 modify model_zoo frcnn's import and numpy version rely

From: @changzherui
Reviewed-by: @kingxian,@zh_qh
Signed-off-by: @kingxian

README.md

@@ -69,6 +69,8 @@ MindSpore offers build options across multiple backends:
 |  | Ubuntu-aarch64 | ✔️ |
 |  | EulerOS-x86 | ✔️ |
 |  | EulerOS-aarch64 | ✔️ |
+|  | CentOS-x86 | ✔️ |
+|  | CentOS-aarch64 | ✔️ |
 | GPU CUDA 10.1 | Ubuntu-x86 | ✔️ |
 | CPU | Ubuntu-x86 | ✔️ |
 |  | Ubuntu-aarch64 | ✔️ |
@@ -79,7 +81,7 @@ For installation using `pip`, take `CPU` and `Ubuntu-x86` build version as an ex
 1. Download whl from [MindSpore download page](https://www.mindspore.cn/versions/en), and install the package.
 
     ```
-    pip install https://ms-release.obs.cn-north-4.myhuaweicloud.com/0.7.0-beta/MindSpore/cpu/ubuntu_x86/mindspore-0.7.0-cp37-cp37m-linux_x86_64.whl
+    pip install https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.1/MindSpore/cpu/ubuntu_x86/mindspore-1.0.1-cp37-cp37m-linux_x86_64.whl
     ```
 
 2. Run the following command to verify the install.
@@ -136,8 +138,8 @@ currently the containerized build options are supported as follows:
 
     For `CPU` backend, you can directly pull and run the latest stable image using the below command:
     ```
-    docker pull mindspore/mindspore-cpu:0.7.0-beta
-    docker run -it mindspore/mindspore-cpu:0.7.0-beta /bin/bash
+    docker pull mindspore/mindspore-cpu:1.0.1
+    docker run -it mindspore/mindspore-cpu:1.0.1 /bin/bash
     ```
 
 * GPU
@@ -170,8 +172,8 @@ currently the containerized build options are supported as follows:
     ```
     Then you can pull and run the latest stable image using the below command:
     ```
-    docker pull mindspore/mindspore-gpu:0.7.0-beta
-    docker run -it --runtime=nvidia --privileged=true mindspore/mindspore-gpu:0.7.0-beta /bin/bash
+    docker pull mindspore/mindspore-gpu:1.0.0
+    docker run -it -v /dev/shm:/dev/shm --runtime=nvidia --privileged=true mindspore/mindspore-gpu:1.0.0 /bin/bash
     ```
 
     To test if the docker image works, please execute the python code below and check the output:
@@ -206,7 +208,7 @@ please check out [docker](docker/README.md) repo for the details.
 
 ## Quickstart
 
-See the [Quick Start](https://www.mindspore.cn/tutorial/training/en/master/quick_start/quick_start.html)
+See the [Quick Start](https://www.mindspore.cn/tutorial/training/en/master/quick_start/quick_start.html) 
 to implement the image classification.
 
 ## Docs
@@ -232,6 +234,34 @@ Check out how MindSpore Open Governance [works](https://gitee.com/mindspore/comm
 Welcome contributions. See our [Contributor Wiki](CONTRIBUTING.md) for
 more details.
 
+## Maintenance phases
+Project stable branches will be in one of the following states:
+| **State**       | **Time frame**    | **Summary**                                          |
+|-------------|---------------|--------------------------------------------------|
+| Planning    | 1 - 3 months  | Features are under planning.                     |
+| Development | 3 months      | Features are under development.                  |
+| Maintained  | 6 - 12 months | All bugfixes are appropriate. Releases produced. |
+| Unmaintained| 0 - 3 months  | All bugfixes are appropriate. No Maintainers and No Releases produced.                                                 |
+| End Of Life (EOL) |  N/A |  Branch no longer accepting changes.    |
+
+## Maintenance status
+| **Branch** | **Status**       | **Initial Release Date** | **Next Phase**                        | **EOL Date**   |
+|--------|--------------|----------------------|-----------------------------------|------------|
+| **r1.1**   | Development  | 2020-12-31 estimated | Maintained <br> 2020-12-31 estimated   |            |
+| **r1.0**   | Maintained   | 2020-09-24           | Unmaintained <br> 2021-03-30 estimated |            |
+| **r0.7**   | Maintained   | 2020-08-31           | Unmaintained <br> 2020-11-30 estimated |            |
+| **r0.6**   | Unmaintained | 2020-07-31           | End Of Life <br> 2020-12-30 estimated  |            |
+| **r0.5**   | Maintained   | 2020-06-30           | Unmaintained <br> 2021-06-30 estimated |            |
+| **r0.3**   | End Of Life  | 2020-05-31           |                                   | 2020-09-30 |
+| **r0.2**   | End Of Life  | 2020-04-30           |                                   | 2020-08-31 |
+| **r0.1**   | End Of Life  | 2020-03-28            |                                   | 2020-06-30 |
+
+
+
+
+
+
+
 ## Release Notes
 
 The release notes, see our [RELEASE](RELEASE.md).

README_CN.md

@@ -66,6 +66,8 @@ MindSpore提供跨多个后端的构建选项：
 |               | Ubuntu-aarch64  | ✔️   |
 |               | EulerOS-x86     | ✔️   |
 |               | EulerOS-aarch64 | ✔️   |
+|               | CentOS-x86      | ✔️   |
+|               | CentOS-aarch64  | ✔️   |
 | GPU CUDA 10.1 | Ubuntu-x86      | ✔️   |
 | CPU           | Ubuntu-x86      | ✔️   |
 |               | Ubuntu-aarch64  | ✔️   |
@@ -76,7 +78,7 @@ MindSpore提供跨多个后端的构建选项：
 1. 请从[MindSpore下载页面](https://www.mindspore.cn/versions)下载并安装whl包。
 
     ```
-    pip install https://ms-release.obs.cn-north-4.myhuaweicloud.com/0.7.0-beta/MindSpore/cpu/ubuntu_x86/mindspore-0.7.0-cp37-cp37m-linux_x86_64.whl
+    pip install https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.1/MindSpore/cpu/ubuntu_x86/mindspore-1.0.1-cp37-cp37m-linux_x86_64.whl
     ```
 
 2. 执行以下命令，验证安装结果。
@@ -132,8 +134,8 @@ MindSpore的Docker镜像托管在[Docker Hub](https://hub.docker.com/r/mindspore
 
     对于`CPU`后端，可以直接使用以下命令获取并运行最新的稳定镜像：
     ```
-    docker pull mindspore/mindspore-cpu:0.7.0-beta
-    docker run -it mindspore/mindspore-cpu:0.7.0-beta /bin/bash
+    docker pull mindspore/mindspore-cpu:1.0.1
+    docker run -it mindspore/mindspore-cpu:1.0.1 /bin/bash
     ```
 
 * GPU
@@ -166,8 +168,8 @@ MindSpore的Docker镜像托管在[Docker Hub](https://hub.docker.com/r/mindspore
     ```
     使用以下命令获取并运行最新的稳定镜像：
     ```
-    docker pull mindspore/mindspore-gpu:0.7.0-beta
-    docker run -it --runtime=nvidia --privileged=true mindspore/mindspore-gpu:0.7.0-beta /bin/bash
+    docker pull mindspore/mindspore-gpu:1.0.0
+    docker run -it -v /dev/shm:/dev/shm --runtime=nvidia --privileged=true mindspore/mindspore-gpu:1.0.0 /bin/bash
     ```
 
     要测试Docker是否正常工作，请运行下面的Python代码并检查输出：
@@ -225,6 +227,34 @@ MindSpore的Docker镜像托管在[Docker Hub](https://hub.docker.com/r/mindspore
 
 欢迎参与贡献。更多详情，请参阅我们的[贡献者Wiki](CONTRIBUTING.md)。
 
+## 分支维护策略
+MindSpore的版本分支有以下几种维护阶段：
+| **状态**       | **持续时间**    | **说明**                                          |
+|-------------|---------------|--------------------------------------------------|
+| Planning    | 1 - 3 months  | 特性规划。                     |
+| Development | 3 months      | 特性开发。                  |
+| Maintained  | 6 - 12 months | 允许所有问题修复的合入，并发布版本。 |
+| Unmaintained| 0 - 3 months  | 允许所有问题修复的合入，无专人维护，不再发布版本。                                                 |
+| End Of Life (EOL) |  N/A |  不再接受修改合入该分支。    |
+
+## 现有分支维护状态
+| **分支名** | **当前状态**       | **上线时间** | **后续状态**                        | **EOL 日期**   |
+|--------|--------------|----------------------|-----------------------------------|------------|
+| **r1.1**   | Development  | 2020-12-31 estimated | Maintained <br> 2020-12-31 estimated   |            |
+| **r1.0**   | Maintained   | 2020-09-24           | Unmaintained <br> 2021-03-30 estimated |            |
+| **r0.7**   | Maintained   | 2020-08-31           | Unmaintained <br> 2020-11-30 estimated |            |
+| **r0.6**   | Unmaintained | 2020-07-31           | End Of Life <br> 2020-12-30 estimated  |            |
+| **r0.5**   | Maintained   | 2020-06-30           | Unmaintained <br> 2021-06-30 estimated |            |
+| **r0.3**   | End Of Life  | 2020-05-31           |                                   | 2020-09-30 |
+| **r0.2**   | End Of Life  | 2020-04-30           |                                   | 2020-08-31 |
+| **r0.1**   | End Of Life  | 2020-03-28            |                                   | 2020-06-30 |
+
+
+
+
+
+
+
 ## 版本说明
 
 版本说明请参阅[RELEASE](RELEASE.md)。

RELEASE.md

@@ -1,3 +1,130 @@
+# Release 1.0.1
+
+## Major Features and Improvements
+## Bugfixes
+* Models
+    * fix some ssd gpu loss is inf.([!7448](https://gitee.com/mindspore/mindspore/pulls/7448))
+    * add Softplus, SoftplusGrad for old backend.([!7480](https://gitee.com/mindspore/mindspore/pulls/7480))
+    * fix unsupported op print and caffe multiple inputs.([!7344](https://gitee.com/mindspore/mindspore/pulls/7344))
+
+## Contributors
+Thanks goes to these wonderful people:
+
+Adel, AGroupofProbiotocs, anthonyaje, anzhengqi, askmiao, baihuawei, baiyangfan, bai-yangfan, bingyaweng, BowenK, buxue, caifubi, CaoJian, caojian05, caozhou, Cathy, changzherui, chenfei, chengxianbin, chenhaozhe, chenjianping, chenzomi, chenzupeng, chujinjin, cj, cjh9368, Corleone, danish, Danish, dayschan, eric, Eric, fary86, fuzhiye, Gaoxiong, gengdongjie, gongdaguo, gukecai, guoqi, gzhcv, hangq, hanhuifeng2020, Harshvardhan, He, heleiwang, hexia, Hoai, HuangBingjian, huangdongrun, huanghui, huangxinjing, huzhifeng, hwjiaorui, Jesse, jianghui58, jiangzhiwen, Jiaqi, jin-xiulang, jinyaohui, jjfeing, John, Jonathan, jonyguo, jzg, kai00, kingfo, kingxian, kpy, kswang, laiyongqiang, leonwanghui, Li, liangchenghui, liangzelang, lichen_101010, lichenever, lihongkang, lilei, limingqi107, ling, linqingke, liubuyu, liuwenhao4, liuxiao78, liuxiao93, liuyang_655, liuzhongkai, Lixia, lixian, liyanliu, liyong, lizhenyu, luoyang, lvchangquan, lvliang, lz, mahdi, Mahdi, maning202007, Margaret_wangrui, mayang, mengyuanli, nhussain, ougongchang, panfengfeng, panyifeng, Payne, Peilin, peixu_ren, Pengyongrong, qianlong, r1chardf1d0, riemann_penn, root, Sheng, shenwei41, simson, Simson, Su, sunsuodong, tao_yunhao, tinazhang, VectorSL, , Wan, wandongdong, wangdongxu, wangmin, wangnan39@huawei.com, wangyue01, wangzhe, wanyiming, Wei, wenchunjiang, wilfChen, WilliamLian, wsc, wukesong, wuweikang, wuxuejian, Xiaoda, xiefangqi, xuanyue, xulei2020, Xun, xuyongfei, yanghaitao, yanghaitao1, yanghaoran, YangLuo, yangruoqi713, yankai, yanzhenxiang2020, yao_yf, yepei6, yeyunpeng, Yi, yoni, yoonlee666, yuchaojie, yujianfeng, yuximiao, zengzitao, Zhang, zhanghaibo5@huawei.com, zhanghuiyao, zhangyihui, zhangz0911gm, zhanke, zhanyuan, zhaodezan, zhaojichen, zhaoting, zhaozhenlong, zhengjun10, zhoufeng, zhousiyi, zhouyaqiang, Zichun, Zirui, Ziyan, zjun, ZPaC
+
+Contributions of any kind are welcome!
+# Release 1.0.0
+
+## Major Features and Improvements
+### MindSpore Training and Inference Framework
+#### Ascend 910
+* New models
+    * DenseNet121: a dense convolutional neural network, which connects each layer to every other layer in a feed-forward fashion for object recognition on ImageNet dataset.
+    * UNet2D-Medical: Unet Medical model for 2D image segmentation, Convolutional Networks for Biomedical Image Segmentation on ISBI Challenge database.
+* Frontend and user interface
+    * Second-Order Optimization
+      * Enable second-order optimization for Bert on Ascend 910, which can achieve a masked lm accuracy of 71.3% in 800 seconds using 8 Ascend 910 (Bert-Large @MLPerf v0.7 dataset).
+    * New GNN model BGCF
+      * Bayesian Graph Convolutional Filtering network which naturally incorporate the uncertainty in the user-item interaction graph shows excellent recommendation performance on Amazon-Beauty dataset.
+    * Add append interface for SequentialCell.
+    * Add a level `auto` for AMP.
+* Executor and performance optimization
+    * Support quantitative network (Resnet50 & YoloV3 & MobileNetV2).
+    * Project ease of use optimization: project compilation time optimization, CMakelist regularization, cudnn, cuda independent compilation and installation independent.
+* Data processing, augmentation, and save format
+    * Support GeneratorDataset return string type
+
+#### Other Hardware Support
+* GPU platform
+    * Enable second-order optimization for resnet50 on GPU, which achieve 30% improvement on training time compared to SGD with Momentum (Resnet50 @ImageNet).
+* CPU platform
+    * ...
+#### User interfaces change log
+* Remove global object GradOperation in Autodiff([!5011](https://gitee.com/mindspore/mindspore/pulls/5011))
+* Remove useless attribute 'name' in Autodiff([!5172](https://gitee.com/mindspore/mindspore/pulls/5172))
+* Rectification distributed init([!5350](https://gitee.com/mindspore/mindspore/pulls/5350))
+* Move the setting of ParalleMode from train.parallel_utils to context([!5351](https://gitee.com/mindspore/mindspore/pulls/5351))
+* Modification of save_checkpoint([!5482](https://gitee.com/mindspore/mindspore/pulls/5482))
+* Wrap numpy random seed into an api([!5634](https://gitee.com/mindspore/mindspore/pulls/5634))
+* Delete enable_fused_layernorm in some modelzoo scripts([!5665](https://gitee.com/mindspore/mindspore/pulls/5665))
+* Move 'multi-subgraphs' interface to internal([!5696](https://gitee.com/mindspore/mindspore/pulls/5696))
+* Rename mirror_mean to gradient_mean([!5700](https://gitee.com/mindspore/mindspore/pulls/5700))
+* Remove default value of 'group' of DepthWiseConv2d([!5865](https://gitee.com/mindspore/mindspore/pulls/5865))
+* Modify interface for function and remove duplicated def([!5958](https://gitee.com/mindspore/mindspore/pulls/5958))
+* Unify Conv2d and DepthwiseConv2d([!5916](https://gitee.com/mindspore/mindspore/pulls/5916))
+* Modification of SoftmaxCrossEntropyWithLogits([!5502](https://gitee.com/mindspore/mindspore/pulls/5502))
+* Change API set_strategy() to shard()([!5991](https://gitee.com/mindspore/mindspore/pulls/5991))
+* Move batch_size from bert_cfg_cfg to cfg([!6233](https://gitee.com/mindspore/mindspore/pulls/6233))
+* Remove unused parameters from SummaryRecord __init__([!5548](https://gitee.com/mindspore/mindspore/pulls/5548))
+* remove sens parameter of TrainOneStepWithLossScaleCell([!5753](https://gitee.com/mindspore/mindspore/pulls/5753))
+* optimize the TrainOneStepCell for user's define([!6159](https://gitee.com/mindspore/mindspore/pulls/6159))
+* delete seed0 and seed1 of nn.Dropout([!5735](https://gitee.com/mindspore/mindspore/pulls/5735))
+* delete DataWrapper([!6101](https://gitee.com/mindspore/mindspore/pulls/6101))
+* LSTM API optimization([!6374](https://gitee.com/mindspore/mindspore/pulls/6374))
+* Merge P\C\F of ops([!5645](https://gitee.com/mindspore/mindspore/pulls/5645))
+* delete SoftmaxCrossEntropyExpand interface([!6607](https://gitee.com/mindspore/mindspore/pulls/6607))
+* Adjust GroupNorm interface([!6329](https://gitee.com/mindspore/mindspore/pulls/6329))
+* Modify init interface to internal interface([!6651](https://gitee.com/mindspore/mindspore/pulls/6651))
+* Log optimization([!5842](https://gitee.com/mindspore/mindspore/pulls/5842))
+* Remove useless API dataset.set_dataset_size（[!5806](https://gitee.com/mindspore/mindspore/pulls/5806))
+* Some of Dataset API add usage parameter（[!5605](https://gitee.com/mindspore/mindspore/pulls/5605))
+* Change the import path, such as from mindspore.dataset.transforms.vision to mindspore.dataset.vision.transforms（[!5384](https://gitee.com/mindspore/mindspore/pulls/5384))
+* Rename ImageFolderDatasetV2 to ImageFolderDataset（[!5384](https://gitee.com/mindspore/mindspore/pulls/5384))
+* Dataset.map parameter optimization（[!5384](https://gitee.com/mindspore/mindspore/pulls/5384))
+* Add new api dataset.get_col_names（[!5384](https://gitee.com/mindspore/mindspore/pulls/5384))
+* Add new api dataset.get_col_names（[!5384](https://gitee.com/mindspore/mindspore/pulls/5384))
+* Remove useless API MindRecord finish（[!5580](https://gitee.com/mindspore/mindspore/pulls/5580))
+
+### MindSpore Lite
+* Converter
+    * Add 6 TFLite op, 7 Caffe op, 1 ONNX op.
+    * Add support for Windows.
+    * Support parallel inference of multiple sessions to adapt to more scenarios
+    * Support 8bits only weight-quantization, most main-stream models has small accuracy loss (less than 0.5%) when compared to non-qunantized fp32 model.
+
+* CPU & GPU
+    * Add 20 CPU ops，include FP32, int8/uint8, FP16 and int32 ops.
+    * Add supporting FP16 for GPU, add 14 GPU ops include FP32/FP16.
+    * Add Buffer/Image2D transform op for GPU
+    * Performance optimization for CPU ops focus on ARM32.
+    * Performance optimization for GPU Convolution using winograd.
+
+* Tool & example
+    * Add object detection Android Demo.
+
+## Bugfixes
+* Models
+    * fix the constant folding problem in multiply.([!6092](https://gitee.com/mindspore/mindspore/pulls/6092))
+    * move batch_size from bert_net_cfg to cfg in bert scripts.([!6233](https://gitee.com/mindspore/mindspore/pulls/6233))
+    * modify the checkpoint file path.([!6137](https://gitee.com/mindspore/mindspore/pulls/6137))
+* Python API
+    * fix semi auto parallel parameter of reshape has another user([!5722](https://gitee.com/mindspore/mindspore/pulls/5722))
+    * raise ValueError when call hook function in graph mode([!5831](https://gitee.com/mindspore/mindspore/pulls/5831))
+* Executor
+    * fix pynative mode to build temporary nn objects.（[!6189](https://gitee.com/mindspore/mindspore/pulls/6189))
+    * fix the accuracy problem of multiple inputs of multi-card communication operator broadcast.([!6522](https://gitee.com/mindspore/mindspore/pulls/5622))
+    * fix the problem that the sample distribution interface categorical does not support graph mode.([!5772](https://gitee.com/mindspore/mindspore/pulls/5772))
+    * fix the random seed failure problem of the polynomial downsampling distribution operator.([!5948](https://gitee.com/mindspore/mindspore/pulls/5948))
+    * fix unnecessary address binding issues in GPU heterogeneous scenarios.([!6232](https://gitee.com/mindspore/mindspore/pulls/6232))
+* GPU platform
+    * fix for kernel resource leak([!5315](https://gitee.com/mindspore/mindspore/pulls/5315))
+    * fix for insufficient memory for continuous unit test running([!5617](https://gitee.com/mindspore/mindspore/pulls/5617))
+    * fix for the memory leak in the sparse slicer([!5578](https://gitee.com/mindspore/mindspore/pulls/5578))
+* Data processing
+    * fix hang when use pyfunc([!6346](https://gitee.com/mindspore/mindspore/pulls/6346))
+    * fix GPU device queue does not release GIL during resource clean up([!5964](https://gitee.com/mindspore/mindspore/pulls/5964))
+    * fix hang if scripte exit unnormally([!6441](https://gitee.com/mindspore/mindspore/pulls/6441))
+* Third party
+    * Sqlite : Update sqlite to 3.32.2 to handle [CVE-2020-11656](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11656), [CVE-2020-13871](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13871), [CVE-2020-11655](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655), [CVE-2020-9327](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-9327), [CVE-2020-13630](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13630), [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-15358), [CVE-2020-13631](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13631), [CVE-2020-13632](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13632), [CVE-2020-13434](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13434), [CVE-2020-13435](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13435), and [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655).
+    * Libjpeg-turbo : Update libjpeg-turbo to 2.0.4 to handle [CVE-2020-13790](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13790).
+
+## Contributors
+Thanks goes to these wonderful people:
+
+Adel, AGroupofProbiotocs, anthonyaje, anzhengqi, askmiao, baihuawei, baiyangfan, bai-yangfan, bingyaweng, BowenK, buxue, caifubi, CaoJian, caojian05, caozhou, Cathy, changzherui, chenfei, chengxianbin, chenhaozhe, chenjianping, chenzomi, chenzupeng, chujinjin, cj, cjh9368, Corleone, danish, Danish, dayschan, eric, Eric, fary86, fuzhiye, Gaoxiong, gengdongjie, gongdaguo, gukecai, guoqi, gzhcv, hangq, hanhuifeng2020, Harshvardhan, He, heleiwang, hexia, Hoai, HuangBingjian, huangdongrun, huanghui, huangxinjing, huzhifeng, hwjiaorui, Jesse, jianghui58, jiangzhiwen, Jiaqi, jin-xiulang, jinyaohui, jjfeing, John, Jonathan, jonyguo, jzg, kai00, kingfo, kingxian, kpy, kswang, laiyongqiang, leonwanghui, Li, liangchenghui, liangzelang, lichen_101010, lichenever, lihongkang, lilei, limingqi107, ling, linqingke, liubuyu, liuwenhao4, liuxiao78, liuxiao93, liuyang_655, liuzhongkai, Lixia, lixian, liyanliu, liyong, lizhenyu, luoyang, lvchangquan, lvliang, lz, mahdi, Mahdi, maning202007, Margaret_wangrui, mayang, mengyuanli, nhussain, ougongchang, panfengfeng, panyifeng, Payne, Peilin, peixu_ren, Pengyongrong, qianlong, r1chardf1d0, riemann_penn, root, Sheng, shenwei41, simson, Simson, Su, sunsuodong, tao_yunhao, tinazhang, VectorSL, , Wan, wandongdong, wangdongxu, wangmin, wangnan39@huawei.com, wangyue01, wangzhe, wanyiming, Wei, wenchunjiang, wilfChen, WilliamLian, wsc, wukesong, wuweikang, wuxuejian, Xiaoda, xiefangqi, xuanyue, xulei2020, Xun, xuyongfei, yanghaitao, yanghaitao1, yanghaoran, YangLuo, yangruoqi713, yankai, yanzhenxiang2020, yao_yf, yepei6, yeyunpeng, Yi, yoni, yoonlee666, yuchaojie, yujianfeng, yuximiao, zengzitao, Zhang, zhanghaibo5@huawei.com, zhanghuiyao, zhangyihui, zhangz0911gm, zhanke, zhanyuan, zhaodezan, zhaojichen, zhaoting, zhaozhenlong, zhengjun10, zhoufeng, zhousiyi, zhouyaqiang, Zichun, Zirui, Ziyan, zjun, ZPaC
+
+Contributions of any kind are welcome!
+
 # Release 0.7.0-beta
 
 ## Major Features and Improvements
@@ -97,6 +224,9 @@
   * fix GeneratorDataset time out([!3624](https://gitee.com/mindspore/mindspore/pulls/3624))
   * fix concat operator get_dataset_size error([!4701](https://gitee.com/mindspore/mindspore/pulls/4701))
   * fixing python validator for Repeat Op([!4366](https://gitee.com/mindspore/mindspore/pulls/4366))
+* Third party
+    * Sqlite : Update sqlite to 3.32.2 to handle [CVE-2020-11656](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11656), [CVE-2020-13871](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13871), [CVE-2020-11655](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655), [CVE-2020-9327](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-9327), [CVE-2020-13630](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13630), [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-15358), [CVE-2020-13631](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13631), [CVE-2020-13632](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13632), [CVE-2020-13434](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13434), [CVE-2020-13435](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13435), and [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655).
+    * Libjpeg-turbo : Update libjpeg-turbo to 2.0.4 to handle [CVE-2020-13790](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13790).
 
 ## Contributors
 Thanks goes to these wonderful people:
@@ -168,6 +298,9 @@ Contributions of any kind are welcome!
 * Data processing
     * Fix bug of RandomColor and RandomSharpness default parameter checking  ([!2833](https://gitee.com/mindspore/mindspore/pulls/2833))
     * Fix process hung when training and eval  ([!3469](https://gitee.com/mindspore/mindspore/pulls/3469))
+* Third party
+    * Sqlite : Update sqlite to 3.32.2 to handle [CVE-2020-11656](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11656), [CVE-2020-13871](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13871), [CVE-2020-11655](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655), [CVE-2020-9327](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-9327), [CVE-2020-13630](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13630), [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-15358), [CVE-2020-13631](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13631), [CVE-2020-13632](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13632), [CVE-2020-13434](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13434), [CVE-2020-13435](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13435), and [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655).
+    * Libjpeg-turbo : Update libjpeg-turbo to 2.0.4 to handle [CVE-2020-13790](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13790).
 
 ## Contributors
 Thanks goes to these wonderful people:
@@ -261,6 +394,8 @@ Contributions of any kind are welcome!
     * Fix bug of Cifar dataset reading([!2096](https://gitee.com/mindspore/mindspore/pulls/2096))
     * Fix bug of C++ behavior in RandomCropAndResize([!2026](https://gitee.com/mindspore/mindspore/pulls/2026))
     * Fix the bug of mindrecord shuffle([!2420](https://gitee.com/mindspore/mindspore/pulls/2420))
+* Third party
+    * Sqlite : Update sqlite to 3.32.2 to handle [CVE-2020-11656](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11656), [CVE-2020-13871](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13871), [CVE-2020-11655](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655), [CVE-2020-9327](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-9327), [CVE-2020-13630](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13630), [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-15358), [CVE-2020-13631](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13631), [CVE-2020-13632](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13632), [CVE-2020-13434](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13434), [CVE-2020-13435](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-13435), and [CVE-2020-15358](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2020-11655).
 
 ## Contributors
 Thanks goes to these wonderful people:

akg

@@ -1 +1 @@
-Subproject commit 4d897c23fc41c5f7013efd0c517796233671518a
+Subproject commit 6c492a8c9d9730ad11ffc5481cc532ae500b0da5

build.bat

@@ -16,12 +16,25 @@
 @title mindspore_build
 
 SET BASEPATH=%CD%
-IF NOT EXIST "%BASEPATH%/build" (
+SET BUILD_PATH=%BASEPATH%/build
+
+find "const int ms_version_major =" mindspore\lite\include\version.h > version.txt
+for /f "delims=\= tokens=2" %%a in ('findstr "const int ms_version_major = " version.txt') do (set x=%%a)
+set VERSION_MAJOR=%x:~1,1%
+find "const int ms_version_minor =" mindspore\lite\include\version.h > version.txt
+for /f "delims=\= tokens=2" %%b in ('findstr "const int ms_version_minor = " version.txt') do (set y=%%b)
+set VERSION_MINOR=%y:~1,1%
+find "const int ms_version_revision =" mindspore\lite\include\version.h > version.txt
+for /f "delims=\= tokens=2" %%c in ('findstr "const int ms_version_revision = " version.txt') do (set z=%%c)
+set VERSION_REVISION=%z:~1,1%
+del version.txt
+echo "======Start building MindSpore Lite %VERSION_MAJOR%.%VERSION_MINOR%.%VERSION_REVISION%======"
+
+IF NOT EXIST "%BUILD_PATH%" (
     md "build"
 )
 
-cd %BASEPATH%/build
-set BUILD_PATH=%CD%
+cd %BUILD_PATH%
 
 IF NOT EXIST "%BUILD_PATH%/mindspore" (
     md "mindspore"
@@ -85,7 +98,7 @@ goto run_eof
     cd %BUILD_PATH%/mindspore
     cmake -DBUILD_DEVICE=on -DBUILD_CONVERTER=on -DPLATFORM_ARM64=off -DSUPPORT_TRAIN=off ^
     -DCMAKE_BUILD_TYPE=Release -DSUPPORT_GPU=off -DBUILD_MINDDATA=off -DOFFLINE_COMPILE=off ^
-    -DMS_VERSION_MAJOR=0 -DMS_VERSION_MINOR=7 -DMS_VERSION_REVISION=0 ^
+    -DMS_VERSION_MAJOR=%VERSION_MAJOR% -DMS_VERSION_MINOR=%VERSION_MINOR% -DMS_VERSION_REVISION=%VERSION_REVISION% ^
     -G "CodeBlocks - MinGW Makefiles" "%BASEPATH%/mindspore/lite"
 GOTO:EOF
 
@@ -123,6 +136,6 @@ GOTO:EOF
 :run_fail
     cd %BASEPATH%
     set errorlevel=1
-    EXIT
+    exit /b %errorlevel%
 
 :run_eof

build.sh

@@ -25,7 +25,7 @@ usage()
   echo "bash build.sh [-d] [-r] [-v] [-c on|off] [-t on|off] [-g on|off] [-h] [-b ge] [-m infer|train] \\"
   echo "              [-a on|off] [-p on|off] [-i] [-L] [-R] [-D on|off] [-j[n]] [-e gpu|d|cpu] \\"
   echo "              [-P on|off] [-z [on|off]] [-M on|off] [-V 9.2|10.1] [-I arm64|arm32|x86_64] [-K] \\"
-  echo "              [-B on|off] [-w on|off] [-E] [-l on|off] [-n full|lite|off] [-T on|off]"
+  echo "              [-B on|off] [-w on|off] [-E] [-l on|off] [-n full|lite|off]"
   echo ""
   echo "Options:"
   echo "    -d Debug mode"
@@ -58,7 +58,6 @@ usage()
   echo "    -B Enable debugger, default on"
   echo "    -E Enable IBVERBS for parameter server, default off"
   echo "    -l Compile with python dependency, default on"
-  echo "    -T Enable on-device training, default off"
 }
 
 # check value of input is 'on' or 'off'
@@ -237,7 +236,7 @@ checkopts()
         ;;
       z)
         eval ARG=\$\{$OPTIND\}
-        if [[ -n $ARG && $ARG != -* ]]; then
+        if [[ -n "$ARG" && "$ARG" != -* ]]; then
           OPTARG="$ARG"
           check_on_off $OPTARG z
           OPTIND=$((OPTIND + 1))
@@ -652,7 +651,7 @@ build_lite()
         checkndk
         cmake -DCMAKE_TOOLCHAIN_FILE="${ANDROID_NDK}/build/cmake/android.toolchain.cmake" -DANDROID_NATIVE_API_LEVEL="19"      \
               -DANDROID_NDK="${ANDROID_NDK}" -DANDROID_ABI="arm64-v8a" -DANDROID_TOOLCHAIN_NAME="aarch64-linux-android-clang"  \
-              -DANDROID_STL="c++_shared" -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSUPPORT_TRAIN=${SUPPORT_TRAIN}                     \
+              -DANDROID_STL="c++_static" -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSUPPORT_TRAIN=${SUPPORT_TRAIN}                     \
               -DBUILD_DEVICE=on -DPLATFORM_ARM64=on -DBUILD_CONVERTER=off -DENABLE_NEON=on -DENABLE_FP16="off"      \
               -DSUPPORT_GPU=${ENABLE_GPU} -DOFFLINE_COMPILE=${OPENCL_OFFLINE_COMPILE} -DBUILD_MINDDATA=${COMPILE_MINDDATA_LITE} \
               -DCMAKE_INSTALL_PREFIX=${BASEPATH}/output/tmp -DMS_VERSION_MAJOR=${VERSION_MAJOR}                           \
@@ -662,7 +661,7 @@ build_lite()
         checkndk
         cmake -DCMAKE_TOOLCHAIN_FILE="${ANDROID_NDK}/build/cmake/android.toolchain.cmake" -DANDROID_NATIVE_API_LEVEL="19"      \
               -DANDROID_NDK="${ANDROID_NDK}" -DANDROID_ABI="armeabi-v7a" -DANDROID_TOOLCHAIN_NAME="clang"                      \
-              -DANDROID_STL="c++_shared" -DCMAKE_BUILD_TYPE=${BUILD_TYPE}                                                      \
+              -DANDROID_STL="c++_static" -DCMAKE_BUILD_TYPE=${BUILD_TYPE}                                                      \
               -DBUILD_DEVICE=on -DPLATFORM_ARM32=on -DENABLE_NEON=on -DSUPPORT_TRAIN=${SUPPORT_TRAIN} -DBUILD_CONVERTER=off    \
               -DSUPPORT_GPU=${ENABLE_GPU} -DOFFLINE_COMPILE=${OPENCL_OFFLINE_COMPILE} -DBUILD_MINDDATA=${COMPILE_MINDDATA_LITE} \
               -DCMAKE_INSTALL_PREFIX=${BASEPATH}/output/tmp -DMS_VERSION_MAJOR=${VERSION_MAJOR}                           \

cmake/check_requirements.cmake

@@ -19,9 +19,14 @@ set(Python3_USE_STATIC_LIBS FALSE)
 find_package(Python3 COMPONENTS Interpreter Development)
 if (Python3_FOUND)
     message("Python3 found, version: ${Python3_VERSION}")
-    message("Python3 library path: ${Python3_LIBRARY_DIRS}")
+    message("Python3 library path: ${Python3_LIBRARY}")
     message("Python3 interpreter: ${Python3_EXECUTABLE}")
-    else()
+elseif (Python3_LIBRARY AND Python3_EXECUTABLE AND
+        ${Python3_VERSION} VERSION_GREATER_EQUAL "3.7.0" AND ${Python3_VERSION} VERSION_LESS "3.8.0")
+    message(WARNING "Maybe python3 environment is broken.")
+    message("Python3 library path: ${Python3_LIBRARY}")
+    message("Python3 interpreter: ${Python3_EXECUTABLE}")
+else ()
     message(FATAL_ERROR "Python3 not found, please install Python>=3.7.5, and set --enable-shared "
             "if you are building Python locally")
 endif ()

cmake/dependency_graphengine.cmake

@@ -56,6 +56,15 @@ else()
     find_library(register libregister.so ${ASCEND_RUNTIME_PATH})
     find_library(resource libresource.so ${ASCEND_RUNTIME_PATH})
     find_library(error_manager liberror_manager.so ${ASCEND_RUNTIME_PATH})
+    # for Atlas env
+    set(ASCEND_TOOLKIT_RUNTIME_PATH ${ASCEND_PATH}/ascend-toolkit/latest/fwkacllib/lib64)
+    find_library(cce libcce.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(hccl libhccl.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(runtime libruntime.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(msprof libmsprof.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(register libregister.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(resource libresource.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(error_manager liberror_manager.so ${ASCEND_TOOLKIT_RUNTIME_PATH})
 endif()
 
 # compile libraries from following directories

cmake/external_libs/nccl.cmake

@@ -1,10 +1,10 @@
 
 set(nccl_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
 mindspore_add_pkg(nccl
-        VER 2.4.8-1
+        VER 2.7.6-1
         LIBS nccl
-        URL https://github.com/NVIDIA/nccl/archive/v2.4.8-1.tar.gz
-        MD5 f14b37d6af1c79db5f57cb029a753727
+        URL https://github.com/NVIDIA/nccl/archive/v2.7.6-1.tar.gz
+        MD5 073b19899f374c5ba07d2db02dc38f9f
         BUILD_OPTION src.build NVCC_GENCODE="-gencode=arch=compute_70,code=sm_70"
         INSTALL_INCS build/include/*
         INSTALL_LIBS build/lib/*)

cmake/package.cmake

@@ -193,6 +193,11 @@ if (NOT ENABLE_GE)
             DESTINATION ${INSTALL_LIB_DIR}
             COMPONENT mindspore
         )
+        install(
+            TARGETS ms_profile
+            DESTINATION ${INSTALL_LIB_DIR}
+            COMPONENT mindspore
+        )
     elseif (ENABLE_TESTCASES)
         install(
             FILES

cmake/package_lite.cmake

@@ -56,7 +56,7 @@ if (PLATFORM_ARM64)
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite.so DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite.a DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/mindspore/core/ir/dtype/type_id.h DESTINATION ${INC_DIR}/ir/dtype COMPONENT ${COMPONENT_NAME})
-    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR} COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h")
+    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR} COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "train_model.h" EXCLUDE PATTERN "train_session.h" EXCLUDE)
     install(DIRECTORY ${TOP_DIR}/mindspore/lite/schema/ DESTINATION ${INC_DIR}/schema COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "inner" EXCLUDE)
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite-optimize.so DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite-fp16.so DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
@@ -65,7 +65,7 @@ elseif (PLATFORM_ARM32)
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite.so DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/mindspore/lite/build/src/libmindspore-lite.a DESTINATION ${LIB_DIR} COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/mindspore/core/ir/dtype/type_id.h DESTINATION ${INC_DIR}/ir/dtype COMPONENT ${COMPONENT_NAME})
-    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR} COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h")
+    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR} COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "train_model.h" EXCLUDE PATTERN "train_session.h" EXCLUDE)
     install(DIRECTORY ${TOP_DIR}/mindspore/lite/schema/ DESTINATION ${INC_DIR}/schema COMPONENT ${COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "inner" EXCLUDE)
     install(DIRECTORY ${TOP_DIR}/third_party/flatbuffers/include DESTINATION ${FLATBF_DIR} COMPONENT ${COMPONENT_NAME})
 elseif (CMAKE_SYSTEM_NAME MATCHES "Windows")
@@ -75,7 +75,7 @@ elseif (CMAKE_SYSTEM_NAME MATCHES "Windows")
     install(FILES ${LIB_LIST} DESTINATION ${TOP_DIR}/build/mindspore/package COMPONENT ${COMPONENT_NAME})
     install(FILES ${TOP_DIR}/build/mindspore/tools/converter/libconverter_parser.a DESTINATION ${TOP_DIR}/build/mindspore/package COMPONENT ${PARSER_NAME})
 else ()
-    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR_RUN_X86} COMPONENT ${RUN_X86_COMPONENT_NAME} FILES_MATCHING PATTERN "*.h")
+    install(DIRECTORY ${TOP_DIR}/mindspore/lite/include/ DESTINATION ${INC_DIR_RUN_X86} COMPONENT ${RUN_X86_COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "train_model.h" EXCLUDE PATTERN "train_session.h" EXCLUDE)
     install(DIRECTORY ${TOP_DIR}/mindspore/lite/schema/ DESTINATION ${INC_DIR_RUN_X86}/schema COMPONENT ${RUN_X86_COMPONENT_NAME} FILES_MATCHING PATTERN "*.h" PATTERN "inner" EXCLUDE)
     install(FILES ${TOP_DIR}/mindspore/core/ir/dtype/type_id.h DESTINATION ${INC_DIR_RUN_X86}/ir/dtype COMPONENT ${RUN_X86_COMPONENT_NAME})
     install(DIRECTORY ${TOP_DIR}/third_party/flatbuffers/include DESTINATION ${FLATBF_DIR_RUN_X86} COMPONENT ${RUN_X86_COMPONENT_NAME})

docker/mindspore-cpu/1.0.0/Dockerfile

@@ -0,0 +1,67 @@
+FROM ubuntu:18.04
+
+MAINTAINER leonwanghui <leon.wanghui@huawei.com>
+
+# Set env
+ENV PYTHON_ROOT_PATH /usr/local/python-3.7.5
+ENV PATH /usr/local/bin:$PATH
+
+# Install base tools
+RUN apt update \
+    && DEBIAN_FRONTEND=noninteractive apt install -y \
+    vim \
+    wget \
+    curl \
+    xz-utils \
+    net-tools \
+    openssh-client \
+    git \
+    ntpdate \
+    tzdata \
+    tcl \
+    sudo \
+    bash-completion
+
+# Install compile tools
+RUN DEBIAN_FRONTEND=noninteractive apt install -y \
+    gcc \
+    g++ \
+    zlibc \
+    make \
+    libgmp-dev \
+    patch \
+    autoconf \
+    libtool \
+    automake \
+    flex
+
+# Set bash
+RUN echo "dash dash/sh boolean false" | debconf-set-selections
+RUN DEBIAN_FRONTEND=noninteractive dpkg-reconfigure dash
+
+# Install python (v3.7.5)
+RUN apt install -y libffi-dev libssl-dev zlib1g-dev libbz2-dev libncurses5-dev \
+    libgdbm-dev libgdbm-compat-dev liblzma-dev libreadline-dev libsqlite3-dev \
+    && cd /tmp \
+    && wget https://github.com/python/cpython/archive/v3.7.5.tar.gz \
+    && tar -xvf v3.7.5.tar.gz \
+    && cd /tmp/cpython-3.7.5 \
+    && mkdir -p ${PYTHON_ROOT_PATH} \
+    && ./configure --prefix=${PYTHON_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -f /usr/local/bin/python \
+    && rm -f /usr/local/bin/pip \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/python3.7 /usr/local/bin/python \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/pip3.7 /usr/local/bin/pip \
+    && rm -rf /tmp/cpython-3.7.5 \
+    && rm -f /tmp/v3.7.5.tar.gz
+
+# Set pip source
+RUN mkdir -pv /root/.pip \
+    && echo "[global]" > /root/.pip/pip.conf \
+    && echo "trusted-host=mirrors.aliyun.com" >> /root/.pip/pip.conf \
+    && echo "index-url=http://mirrors.aliyun.com/pypi/simple/" >> /root/.pip/pip.conf
+
+# Install MindSpore cpu whl package
+RUN pip install --no-cache-dir https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.0/MindSpore/cpu/ubuntu_x86/mindspore-1.0.0-cp37-cp37m-linux_x86_64.whl

docker/mindspore-cpu/1.0.1/Dockerfile

@@ -0,0 +1,67 @@
+FROM ubuntu:18.04
+
+MAINTAINER leonwanghui <leon.wanghui@huawei.com>
+
+# Set env
+ENV PYTHON_ROOT_PATH /usr/local/python-3.7.5
+ENV PATH /usr/local/bin:$PATH
+
+# Install base tools
+RUN apt update \
+    && DEBIAN_FRONTEND=noninteractive apt install -y \
+    vim \
+    wget \
+    curl \
+    xz-utils \
+    net-tools \
+    openssh-client \
+    git \
+    ntpdate \
+    tzdata \
+    tcl \
+    sudo \
+    bash-completion
+
+# Install compile tools
+RUN DEBIAN_FRONTEND=noninteractive apt install -y \
+    gcc \
+    g++ \
+    zlibc \
+    make \
+    libgmp-dev \
+    patch \
+    autoconf \
+    libtool \
+    automake \
+    flex
+
+# Set bash
+RUN echo "dash dash/sh boolean false" | debconf-set-selections
+RUN DEBIAN_FRONTEND=noninteractive dpkg-reconfigure dash
+
+# Install python (v3.7.5)
+RUN apt install -y libffi-dev libssl-dev zlib1g-dev libbz2-dev libncurses5-dev \
+    libgdbm-dev libgdbm-compat-dev liblzma-dev libreadline-dev libsqlite3-dev \
+    && cd /tmp \
+    && wget https://github.com/python/cpython/archive/v3.7.5.tar.gz \
+    && tar -xvf v3.7.5.tar.gz \
+    && cd /tmp/cpython-3.7.5 \
+    && mkdir -p ${PYTHON_ROOT_PATH} \
+    && ./configure --prefix=${PYTHON_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -f /usr/local/bin/python \
+    && rm -f /usr/local/bin/pip \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/python3.7 /usr/local/bin/python \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/pip3.7 /usr/local/bin/pip \
+    && rm -rf /tmp/cpython-3.7.5 \
+    && rm -f /tmp/v3.7.5.tar.gz
+
+# Set pip source
+RUN mkdir -pv /root/.pip \
+    && echo "[global]" > /root/.pip/pip.conf \
+    && echo "trusted-host=mirrors.aliyun.com" >> /root/.pip/pip.conf \
+    && echo "index-url=http://mirrors.aliyun.com/pypi/simple/" >> /root/.pip/pip.conf
+
+# Install MindSpore cpu whl package
+RUN pip install --no-cache-dir https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.1/MindSpore/cpu/ubuntu_x86/mindspore-1.0.1-cp37-cp37m-linux_x86_64.whl

docker/mindspore-gpu/1.0.0/Dockerfile

@@ -0,0 +1,81 @@
+FROM nvidia/cuda:10.1-cudnn7-devel-ubuntu18.04
+
+MAINTAINER leonwanghui <leon.wanghui@huawei.com>
+
+# Set env
+ENV PYTHON_ROOT_PATH /usr/local/python-3.7.5
+ENV OMPI_ROOT_PATH /usr/local/openmpi-3.1.5
+ENV PATH ${OMPI_ROOT_PATH}/bin:/usr/local/bin:$PATH
+ENV LD_LIBRARY_PATH ${OMPI_ROOT_PATH}/lib:$LD_LIBRARY_PATH
+
+# Install base tools
+RUN apt update \
+    && DEBIAN_FRONTEND=noninteractive apt install -y \
+    vim \
+    wget \
+    curl \
+    xz-utils \
+    net-tools \
+    openssh-client \
+    git \
+    ntpdate \
+    tzdata \
+    tcl \
+    sudo \
+    bash-completion
+
+# Install compile tools
+RUN DEBIAN_FRONTEND=noninteractive apt install -y \
+    gcc \
+    g++ \
+    zlibc \
+    make \
+    libgmp-dev \
+    patch \
+    autoconf \
+    libtool \
+    automake \
+    flex
+
+# Set bash
+RUN echo "dash dash/sh boolean false" | debconf-set-selections
+RUN DEBIAN_FRONTEND=noninteractive dpkg-reconfigure dash
+
+# Install python (v3.7.5)
+RUN apt install -y libffi-dev libssl-dev zlib1g-dev libbz2-dev libncurses5-dev \
+    libgdbm-dev libgdbm-compat-dev liblzma-dev libreadline-dev libsqlite3-dev \
+    && cd /tmp \
+    && wget https://github.com/python/cpython/archive/v3.7.5.tar.gz \
+    && tar -xvf v3.7.5.tar.gz \
+    && cd /tmp/cpython-3.7.5 \
+    && mkdir -p ${PYTHON_ROOT_PATH} \
+    && ./configure --prefix=${PYTHON_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -f /usr/local/bin/python \
+    && rm -f /usr/local/bin/pip \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/python3.7 /usr/local/bin/python \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/pip3.7 /usr/local/bin/pip \
+    && rm -rf /tmp/cpython-3.7.5 \
+    && rm -f /tmp/v3.7.5.tar.gz
+
+# Set pip source
+RUN mkdir -pv /root/.pip \
+    && echo "[global]" > /root/.pip/pip.conf \
+    && echo "trusted-host=mirrors.aliyun.com" >> /root/.pip/pip.conf \
+    && echo "index-url=http://mirrors.aliyun.com/pypi/simple/" >> /root/.pip/pip.conf
+
+# Install openmpi (v3.1.5)
+RUN cd /tmp \
+    && wget https://download.open-mpi.org/release/open-mpi/v3.1/openmpi-3.1.5.tar.gz \
+    && tar -xvf openmpi-3.1.5.tar.gz \
+    && cd /tmp/openmpi-3.1.5 \
+    && mkdir -p ${OMPI_ROOT_PATH} \
+    && ./configure --prefix=${OMPI_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -rf /tmp/openmpi-3.1.5 \
+    && rm -f /tmp/openmpi-3.1.5.tar.gz
+
+# Install MindSpore cuda-10.1 whl package
+RUN pip install --no-cache-dir https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.0/MindSpore/gpu/ubuntu_x86/cuda-10.1/mindspore_gpu-1.0.0-cp37-cp37m-linux_x86_64.whl

docker/mindspore-gpu/1.0.1/Dockerfile

@@ -0,0 +1,81 @@
+FROM nvidia/cuda:10.1-cudnn7-devel-ubuntu18.04
+
+MAINTAINER leonwanghui <leon.wanghui@huawei.com>
+
+# Set env
+ENV PYTHON_ROOT_PATH /usr/local/python-3.7.5
+ENV OMPI_ROOT_PATH /usr/local/openmpi-3.1.5
+ENV PATH ${OMPI_ROOT_PATH}/bin:/usr/local/bin:$PATH
+ENV LD_LIBRARY_PATH ${OMPI_ROOT_PATH}/lib:$LD_LIBRARY_PATH
+
+# Install base tools
+RUN apt update \
+    && DEBIAN_FRONTEND=noninteractive apt install -y \
+    vim \
+    wget \
+    curl \
+    xz-utils \
+    net-tools \
+    openssh-client \
+    git \
+    ntpdate \
+    tzdata \
+    tcl \
+    sudo \
+    bash-completion
+
+# Install compile tools
+RUN DEBIAN_FRONTEND=noninteractive apt install -y \
+    gcc \
+    g++ \
+    zlibc \
+    make \
+    libgmp-dev \
+    patch \
+    autoconf \
+    libtool \
+    automake \
+    flex
+
+# Set bash
+RUN echo "dash dash/sh boolean false" | debconf-set-selections
+RUN DEBIAN_FRONTEND=noninteractive dpkg-reconfigure dash
+
+# Install python (v3.7.5)
+RUN apt install -y libffi-dev libssl-dev zlib1g-dev libbz2-dev libncurses5-dev \
+    libgdbm-dev libgdbm-compat-dev liblzma-dev libreadline-dev libsqlite3-dev \
+    && cd /tmp \
+    && wget https://github.com/python/cpython/archive/v3.7.5.tar.gz \
+    && tar -xvf v3.7.5.tar.gz \
+    && cd /tmp/cpython-3.7.5 \
+    && mkdir -p ${PYTHON_ROOT_PATH} \
+    && ./configure --prefix=${PYTHON_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -f /usr/local/bin/python \
+    && rm -f /usr/local/bin/pip \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/python3.7 /usr/local/bin/python \
+    && ln -s ${PYTHON_ROOT_PATH}/bin/pip3.7 /usr/local/bin/pip \
+    && rm -rf /tmp/cpython-3.7.5 \
+    && rm -f /tmp/v3.7.5.tar.gz
+
+# Set pip source
+RUN mkdir -pv /root/.pip \
+    && echo "[global]" > /root/.pip/pip.conf \
+    && echo "trusted-host=mirrors.aliyun.com" >> /root/.pip/pip.conf \
+    && echo "index-url=http://mirrors.aliyun.com/pypi/simple/" >> /root/.pip/pip.conf
+
+# Install openmpi (v3.1.5)
+RUN cd /tmp \
+    && wget https://download.open-mpi.org/release/open-mpi/v3.1/openmpi-3.1.5.tar.gz \
+    && tar -xvf openmpi-3.1.5.tar.gz \
+    && cd /tmp/openmpi-3.1.5 \
+    && mkdir -p ${OMPI_ROOT_PATH} \
+    && ./configure --prefix=${OMPI_ROOT_PATH} \
+    && make -j4 \
+    && make install -j4 \
+    && rm -rf /tmp/openmpi-3.1.5 \
+    && rm -f /tmp/openmpi-3.1.5.tar.gz
+
+# Install MindSpore cuda-10.1 whl package
+RUN pip install --no-cache-dir https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.0.1/MindSpore/gpu/ubuntu_x86/cuda-10.1/mindspore_gpu-1.0.1-cp37-cp37m-linux_x86_64.whl

graphengine

@@ -1 +1 @@
-Subproject commit 4f6ebe0f92da8c758b1b71a2bc7ff33f9344e18a
+Subproject commit d2ab1c2819c90e61d98b5f88699bcd6b6de4423c

include/infer_tensor.h

@@ -74,7 +74,11 @@ class InferTensorBase {
     if (data_len == 0) {
       return true;
     }
-    memcpy_s(mutable_data(), data_size(), data, data_len);
+    auto ret = memcpy_s(mutable_data(), data_size(), data, data_len);
+    if (ret != 0) {
+      MSI_LOG_ERROR << "Set data memcpy_s failed";
+      return false;
+    }
     return true;
   }
 

include/inference.h

@@ -81,7 +81,6 @@ class MS_API InferSession {
   }
   static std::shared_ptr<InferSession> CreateSession(const std::string &device, uint32_t device_id);
 };
-
 }  // namespace inference
 }  // namespace mindspore
 #endif  // MINDSPORE_INCLUDE_MS_SESSION_H

mindspore/_extends/parallel_compile/tbe_compiler/common.py

@@ -13,7 +13,6 @@
 # limitations under the License.
 # ============================================================================
 """tbe common"""
-import json
 import os
 
 class TBEException(Exception):
@@ -27,23 +26,6 @@ class TBEException(Exception):
         return self.__error_msg
 
 
-def get_ddk_version():
-    """get ddk version"""
-    ddk_version = os.environ.get("DDK_VERSION")
-    if ddk_version is None:
-        default_ddk_info_file = '/usr/local/HiAI/runtime/ddk_info'
-        backup_ddk_info_file = '/usr/local/Ascend/fwkacllib/ddk_info'
-        if os.path.exists(default_ddk_info_file):
-            with open(default_ddk_info_file, "r") as fp:
-                ddk_version = json.load(fp)["VERSION"]
-        elif os.path.exists(backup_ddk_info_file):
-            with open(backup_ddk_info_file, "r") as fp:
-                ddk_version = json.load(fp)["VERSION"]
-        else:
-            ddk_version = "Ascend910"
-    return ddk_version
-
-
 def get_build_in_impl_path():
     """get build-in tbe implement path"""
     tbe_impl_path = os.environ.get("TBE_IMPL_PATH")

mindspore/_extends/parallel_compile/tbe_compiler/compiler.py

@@ -18,9 +18,8 @@ import os
 import sys
 from te.platform.cce_conf import te_set_version
 from te.platform.fusion_util import fusion_op
-from common import check_kernel_info, get_args, get_build_in_impl_path, get_ddk_version
+from common import check_kernel_info, get_args, get_build_in_impl_path
 
-ddk_version = get_ddk_version()
 build_in_impl_path = get_build_in_impl_path()
 
 # op function list
@@ -30,7 +29,6 @@ fusion_pattern_end_flag = "fusion_pattern_end"
 
 def _initialize(impl_path):
     """Initialize"""
-    te_set_version(ddk_version)
     if impl_path == "":
         op_module_name = build_in_impl_path
     else:
@@ -53,7 +51,7 @@ def build_op(build_type, json_str):
     """
     kernel_info = json.loads(json_str)
     check_kernel_info(kernel_info)
-
+    te_set_version(kernel_info["op_info"]["socVersion"])
     op_name = kernel_info['op_info']['name']
 
     try:
@@ -111,7 +109,7 @@ def compile_fusion_op(json_str):
         Exception: If specific keyword is not found.
     """
     args = json.loads(json_str)
-    te_set_version(ddk_version)
+    te_set_version(args['fusion_op']["socVersion"])
     if 'fusion_op' not in args or not args['fusion_op']:
         raise ValueError("Json string Errors, key:fusion_op not found.")
     fusion_op_arg = args['fusion_op']

mindspore/_extends/parallel_compile/tbe_compiler/tbe_process.py

@@ -104,7 +104,7 @@ class TbeProcess:
     def __init__(self):
         self.__processe_num = multiprocessing.cpu_count()
         # max_processes_num: Set the maximum number of concurrent processes for compiler
-        max_processes_num = 16
+        max_processes_num = 24
         if self.__processe_num > max_processes_num:
             self.__processe_num = max_processes_num
         self.__pool = None

mindspore/_extends/parse/standard_method.py

@@ -30,7 +30,6 @@ trans = P.Transpose()
 shape_ = P.Shape()
 dtype_ = P.DType()
 
-
 def all_(x, axis=(), keep_dims=False):
     """
     Check all array elements along a given axis evaluate to True.
@@ -144,12 +143,16 @@ def bool_(x):
 
 
 def enumerate_(x, start=0):
-    """Enumerate list or tuple."""
+    """Enumerate list or tuple or tensor."""
     x_type = F.typeof(x)
     ret = ()
     op_name = "enumerate"
-    if check_is_tuple_or_list(x_type, op_name, "first input") and check_is_const_int(start, op_name, "start"):
-        ret = zip(range(start, start + len(x)), x)
+    if check_is_tuple_or_list_or_tensor(x_type, op_name, "first input") and check_is_const_int(start, op_name, "start"):
+        if check_is_tensor(x_type):
+            for i in range(x.shape[0]):
+                ret += ((start + i, x[i]),)
+        else:
+            ret = zip(range(start, start + len(x)), x)
     return ret
 
 
@@ -177,11 +180,19 @@ def check_type_same(x_type, base_type):
 
 
 @constexpr
-def check_is_tuple_or_list(x, op_name, arg_name):
+def check_is_tensor(x):
     """check whether x is list or tuple."""
-    if isinstance(x, (mstype.list_type, mstype.tuple_type)):
+    if isinstance(x, mstype.tensor_type):
         return True
-    raise TypeError(f"For '{op_name}', the '{arg_name}' should be tuple or list, but got {x}.")
+    return False
+
+
+@constexpr
+def check_is_tuple_or_list_or_tensor(x, op_name, arg_name):
+    """check whether x is list or tuple or tensor."""
+    if isinstance(x, (mstype.list_type, mstype.tuple_type, mstype.tensor_type)):
+        return True
+    raise TypeError(f"For '{op_name}', the '{arg_name}' should be tuple or list or tensor, but got {x}.")
 
 
 @constexpr

mindspore/_version_check.py

@@ -78,8 +78,7 @@ class GPUEnvChecker(EnvChecker):
         if Path(self.cuda_bin).is_dir():
             os.environ['PATH'] = self.cuda_bin + ":" + os.environ['PATH']
         else:
-            raise EnvironmentError(
-                f"No such directory: {self.cuda_bin}, please check if cuda is installed correctly.")
+            logger.warning(f"No such directory: {self.cuda_bin}, please check if cuda is installed correctly.")
 
     def check_version(self):
         if not Path(self.cuda_version).is_file():
@@ -92,8 +91,8 @@ class GPUEnvChecker(EnvChecker):
         v = version.parse(v)
         v_str = str(v.major) + "." + str(v.minor)
         if v_str not in self.version:
-            raise EnvironmentError(f"MindSpore version {__version__} and cuda version {v_str} does not match, "
-                                   "reference to the match info on: https://www.mindspore.cn/install")
+            logger.warning(f"MindSpore version {__version__} and cuda version {v_str} does not match, "
+                           "reference to the match info on: https://www.mindspore.cn/install")
 
     def _check_env(self):
         """gpu cuda path check"""
@@ -122,15 +121,25 @@ class AscendEnvChecker(EnvChecker):
 
     def __init__(self):
         self.version = ["1.75.22.0.220"]
-        atlas_fwk_version = "/usr/local/Ascend/nnae/latest/fwkacllib/version.info"
+        atlas_nnae_version = "/usr/local/Ascend/nnae/latest/fwkacllib/version.info"
+        atlas_toolkit_version = "/usr/local/Ascend/ascend-toolkit/latest/fwkacllib/version.info"
         hisi_fwk_version = "/usr/local/Ascend/fwkacllib/version.info"
-        if os.path.exists(atlas_fwk_version):
+        if os.path.exists(atlas_nnae_version):
             # atlas default path
             self.fwk_path = "/usr/local/Ascend/nnae/latest/fwkacllib"
             self.op_impl_path = "/usr/local/Ascend/nnae/latest/opp/op_impl/built-in/ai_core/tbe"
             self.tbe_path = self.fwk_path + "/lib64"
             self.cce_path = self.fwk_path + "/ccec_compiler/bin"
-            self.fwk_version = atlas_fwk_version
+            self.fwk_version = atlas_nnae_version
+            self.op_path = "/usr/local/Ascend/nnae/latest/opp"
+        elif os.path.exists(atlas_toolkit_version):
+            # atlas default path
+            self.fwk_path = "/usr/local/Ascend/ascend-toolkit/latest/fwkacllib"
+            self.op_impl_path = "/usr/local/Ascend/ascend-toolkit/latest/opp/op_impl/built-in/ai_core/tbe"
+            self.tbe_path = self.fwk_path + "/lib64"
+            self.cce_path = self.fwk_path + "/ccec_compiler/bin"
+            self.fwk_version = atlas_toolkit_version
+            self.op_path = "/usr/local/Ascend/ascend-toolkit/latest/opp"
         elif os.path.exists(hisi_fwk_version):
             # hisi default path
             self.fwk_path = "/usr/local/Ascend/fwkacllib"
@@ -138,6 +147,7 @@ class AscendEnvChecker(EnvChecker):
             self.tbe_path = self.fwk_path + "/lib64"
             self.cce_path = self.fwk_path + "/ccec_compiler/bin"
             self.fwk_version = hisi_fwk_version
+            self.op_path = ""
         else:
             # custom or unknown environment
             self.fwk_path = ""
@@ -145,17 +155,20 @@ class AscendEnvChecker(EnvChecker):
             self.tbe_path = ""
             self.cce_path = ""
             self.fwk_version = ""
+            self.op_path = ""
 
         # env
         self.path = os.getenv("PATH")
         self.python_path = os.getenv("PYTHONPATH")
         self.ld_lib_path = os.getenv("LD_LIBRARY_PATH")
+        self.ascend_opp_path = os.getenv("ASCEND_OPP_PATH")
 
         # check content
         self.path_check = "/fwkacllib/ccec_compiler/bin/"
-        self.python_path_check = "opp/op_impl/built_in/ai_core/tbe/"
+        self.python_path_check = "opp/op_impl/built-in/ai_core/tbe/"
         self.ld_lib_path_check_fwk = "/fwkacllib/lib64/"
         self.ld_lib_path_check_addons = "/add-ons/"
+        self.ascend_opp_path_check = "/op"
         self.v = ""
 
     def check_env(self, e):
@@ -171,7 +184,7 @@ class AscendEnvChecker(EnvChecker):
 
         v = self._read_version(self.fwk_version)
         if v not in self.version:
-            raise EnvironmentError(
+            logger.warning(
                 f"MindSpore version {__version__} and Ascend 910 AI software package version {v} does not "
                 "match, reference to the match info on: https://www.mindspore.cn/install")
 
@@ -187,24 +200,33 @@ class AscendEnvChecker(EnvChecker):
             if Path(self.tbe_path).is_dir():
                 os.environ['LD_LIBRARY_PATH'] = self.tbe_path
             else:
-                raise EnvironmentError(
+                logger.warning(
                     f"No such directory: {self.tbe_path}, Please check if Ascend 910 AI software package is "
                     "installed correctly.")
 
         if Path(self.op_impl_path).is_dir():
             sys.path.append(self.op_impl_path)
         else:
-            raise EnvironmentError(
+            logger.warning(
                 f"No such directory: {self.op_impl_path}, Please check if Ascend 910 AI software package is "
                 "installed correctly.")
 
         if Path(self.cce_path).is_dir():
             os.environ['PATH'] = self.cce_path + ":" + os.environ['PATH']
         else:
-            raise EnvironmentError(
+            logger.warning(
                 f"No such directory: {self.cce_path}, Please check if Ascend 910 AI software package is "
                 "installed correctly.")
 
+        if self.op_path is None:
+            pass
+        elif Path(self.op_path).is_dir():
+            os.environ['ASCEND_OPP_PATH'] = self.op_path
+        else:
+            logger.warning(
+                f"No such directory: {self.op_path}, Please check if Ascend 910 AI software package is "
+                "installed correctly.")
+
     def _check_env(self):
         """ascend dependence path check"""
         if self.path is None or self.path_check not in self.path:
@@ -223,6 +245,11 @@ class AscendEnvChecker(EnvChecker):
                            "LD_LIBRARY_PATH, you can reference to the installation guidelines "
                            "https://www.mindspore.cn/install")
 
+        if self.ascend_opp_path is None or self.ascend_opp_path_check not in self.ascend_opp_path:
+            logger.warning(
+                "Can not find opp path (need by mindspore-ascend), please check if you have set env ASCEND_OPP_PATH, "
+                "you can reference to the installation guidelines https://www.mindspore.cn/install")
+
     def _read_version(self, file_path):
         """get ascend version info"""
         with open(file_path, 'r') as f:
@@ -233,6 +260,7 @@ class AscendEnvChecker(EnvChecker):
                     return self.v
         return self.v
 
+
 def check_version_and_env_config():
     """check version and env config"""
     if __package_name__.lower() == "mindspore-ascend":
@@ -253,4 +281,5 @@ def check_version_and_env_config():
     except ImportError as e:
         env_checker.check_env(e)
 
+
 check_version_and_env_config()

mindspore/ccsrc/CMakeLists.txt

@@ -132,7 +132,7 @@ if (ENABLE_D)
     include_directories("${CMAKE_BINARY_DIR}/backend/kernel_compiler/aicpu")
     file(GLOB_RECURSE PROTO_IN RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "backend/kernel_compiler/aicpu/proto/*.proto")
     ms_protobuf_generate(PROTOSRCS PROTOHDRS ${PROTO_IN})
-    
+
     file(GLOB_RECURSE PROTO_DUMP RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "runtime/device/ascend/dump/proto/*.proto")
     ms_protobuf_generate(DUMP_PROTOSRCS PROTOHDRS ${PROTO_DUMP})
 
@@ -176,8 +176,6 @@ endforeach ()
 set_property(SOURCE ${SUB_OBJECTS_SRC} PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_ME)
 add_library(mindspore STATIC ${SUB_OBJECTS_SRC})
 
-target_link_libraries(proto_input mindspore::protobuf)
-
 target_link_libraries(mindspore mindspore_core)
 
 if (ENABLE_DEBUGGER)
@@ -185,7 +183,6 @@ if (ENABLE_DEBUGGER)
     target_link_libraries(proto_input mindspore::grpc++)
 endif()
 
-target_link_libraries(mindspore proto_input)
 target_link_libraries(mindspore securec mindspore::flatbuffers)
 
 if (NOT WIN32)
@@ -222,6 +219,7 @@ if (ENABLE_D)
         set(ASCEND_DRIVER_PATH ${ASCEND_PATH}/driver/lib64/common)
         set(ASCEND_DRIVER_BACK_PATH ${ASCEND_PATH}/driver/lib64/driver)
         set(ASCEND_RUNTIME_PATH ${ASCEND_PATH}/fwkacllib/lib64)
+        set(ASCEND_TOOLKIT_RUNTIME_PATH ${ASCEND_PATH}/ascend-toolkit/latest/fwkacllib/lib64)
     endif()
 
     MESSAGE("USE DAV LIB PATH: ${ASCEND_PATH}")
@@ -230,13 +228,25 @@ if (ENABLE_D)
     find_library(RUNTIME_LIB runtime ${ASCEND_RUNTIME_PATH})
     find_library(TSDCLIENT tsdclient HINTS ${ASCEND_RUNTIME_PATH} ${ASCEND_DRIVER_BACK_PATH})
     find_library(DATATRANSFER datatransfer HINTS ${ASCEND_RUNTIME_PATH} ${ASCEND_DRIVER_BACK_PATH})
-    find_library(PROFILING msprof ${ASCEND_DRIVER_PATH})
-    target_link_libraries(mindspore ge_runtime ${CCE_LIB} ${RUNTIME_LIB} ${TSDCLIENT} ${PROFILING} ${HCCL} ${DATATRANSFER})
-endif()
+    find_library(PROFILING msprofiler ${ASCEND_RUNTIME_PATH})
+    # for atlas env
+    find_library(HCCL hccl ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(CCE_LIB cce ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(RUNTIME_LIB runtime ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    find_library(TSDCLIENT tsdclient HINTS ${ASCEND_TOOLKIT_RUNTIME_PATH} ${ASCEND_DRIVER_BACK_PATH})
+    find_library(DATATRANSFER datatransfer HINTS ${ASCEND_TOOLKIT_RUNTIME_PATH} ${ASCEND_DRIVER_BACK_PATH})
+    find_library(PROFILING msprofiler ${ASCEND_TOOLKIT_RUNTIME_PATH})
 
-# link protobuf
-if (ENABLE_D)
-    target_link_libraries(mindspore mindspore::protobuf)
+    find_library(PROFILING_SHARED msprof ${ASCEND_DRIVER_PATH})
+    add_library(ms_profile SHARED ${PROFILING})
+    set_target_properties(ms_profile PROPERTIES LINKER_LANGUAGE CXX)
+    target_link_libraries(ms_profile -Wl,--start-group ${PROFILING_SHARED} ${PROFILING} mindspore::protobuf -Wl,--end-group)
+    target_link_libraries(mindspore ms_profile ge_runtime ${CCE_LIB} ${RUNTIME_LIB} ${TSDCLIENT} ${HCCL} ${DATATRANSFER} ${OPTILING})
+    target_link_libraries(mindspore -Wl,--start-group proto_input mindspore::protobuf -Wl,--end-group)
+elseif (CMAKE_SYSTEM_NAME MATCHES "Windows")
+    target_link_libraries(mindspore -Wl,--start-group proto_input mindspore::protobuf mindspore::sentencepiece -Wl,--end-group)
+else ()
+    target_link_libraries(mindspore -Wl,--start-group proto_input mindspore::protobuf -Wl,--end-group)
 endif()
 
 # set c_expression building
@@ -259,6 +269,7 @@ endif ()
 set(MINDSPORE_RPATH ${ORIGIN_PATH}/lib)
 if (ENABLE_D)
     set(MINDSPORE_RPATH ${MINDSPORE_RPATH}:/usr/local/Ascend/nnae/latest/fwkacllib/lib64)
+    set(MINDSPORE_RPATH ${MINDSPORE_RPATH}:/usr/local/Ascend/ascend-toolkit/latest/fwkacllib/lib64)
     set(MINDSPORE_RPATH ${MINDSPORE_RPATH}:/usr/local/Ascend/fwkacllib/lib64)
     set(MINDSPORE_RPATH ${MINDSPORE_RPATH}:/usr/local/Ascend/add-ons)
 elseif (ENABLE_GPU)
@@ -278,7 +289,7 @@ else ()
             target_link_libraries(mindspore ibverbs rdmacm)
         endif()
     endif()
-    target_link_libraries(_c_expression PRIVATE -Wl,--whole-archive mindspore -Wl,--no-whole-archive)
+    target_link_libraries(_c_expression PRIVATE -Wl,--whole-archive mindspore proto_input -Wl,--no-whole-archive)
     target_link_libraries(_c_expression PRIVATE mindspore::pybind11_module)
     target_link_libraries(_c_expression PRIVATE mindspore_gvar)
 endif ()
@@ -287,11 +298,6 @@ if (USE_GLOG)
     target_link_libraries(_c_expression PRIVATE mindspore::glog)
 endif ()
 
-if (ENABLE_DUMP_PROTO)
-    message("add protobuf lib to c_expression")
-    target_link_libraries(_c_expression PRIVATE mindspore::protobuf)
-endif ()
-
 if (ENABLE_GPU)
     message("add gpu lib to c_expression")
     target_link_libraries(_c_expression PRIVATE gpu_cuda_lib gpu_queue cublas
@@ -324,7 +330,13 @@ add_library(inference SHARED
         ${LOAD_ONNX_SRC}
         )
 target_link_libraries(inference PRIVATE ${PYTHON_LIBRARIES} ${SECUREC_LIBRARY}
-        -Wl,--whole-archive mindspore -Wl,--no-whole-archive mindspore_gvar mindspore::protobuf)
+        -Wl,--whole-archive mindspore proto_input -Wl,--no-whole-archive mindspore_gvar)
+
+if (ENABLE_D)
+    find_library(adump_server libadump_server.a ${ASCEND_RUNTIME_PATH} ${ASCEND_TOOLKIT_RUNTIME_PATH})
+    target_link_libraries(_c_expression PRIVATE ${adump_server})
+    target_link_libraries(inference PRIVATE ${adump_server})
+endif()
 
 if (ENABLE_CPU)
     target_link_libraries(inference PRIVATE mindspore::dnnl mindspore::mkldnn)
@@ -339,4 +351,3 @@ if (CMAKE_SYSTEM_NAME MATCHES "Linux")
 elseif (CMAKE_SYSTEM_NAME MATCHES "Darwin")
     set_target_properties(inference PROPERTIES MACOSX_RPATH ON)
 endif ()
-

mindspore/ccsrc/backend/kernel_compiler/cpu/adam_cpu_kernel.cc

@@ -0,0 +1,108 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/kernel_compiler/cpu/adam_cpu_kernel.h"
+
+#include <cmath>
+#include <thread>
+#include "backend/kernel_compiler/cpu/mkldnn/mkl_kernel_engine.h"
+#include "runtime/device/cpu/cpu_device_address.h"
+#include "utils/ms_utils.h"
+
+namespace mindspore {
+namespace kernel {
+
+template <typename T>
+void AdamCPUKernel::LaunchAdam(T *var, T *m, T *v, float lr, float beta1, float beta2, float epsilon, const T *gradient,
+                               size_t start, size_t end) {
+  for (size_t i = start; i < end; i++) {
+    m[i] += (gradient[i] - m[i]) * (1 - beta1);
+    v[i] += (gradient[i] * gradient[i] - v[i]) * (1 - beta2);
+    if (use_nesterov) {
+      var[i] -= lr * (m[i] * beta1 + (1 - beta1) * gradient[i]) / (std::sqrt(v[i]) + epsilon);
+    } else {
+      var[i] -= lr * m[i] / (std::sqrt(v[i]) + epsilon);
+    }
+  }
+}
+
+void AdamCPUKernel::InitKernel(const CNodePtr &kernel_node) {
+  MS_EXCEPTION_IF_NULL(kernel_node);
+  size_t input_num = AnfAlgo::GetInputTensorNum(kernel_node);
+  if (input_num != 10) {
+    MS_LOG(EXCEPTION) << "Input number is " << input_num << ", but Adam needs 10 inputs.";
+  }
+  size_t output_num = AnfAlgo::GetOutputTensorNum(kernel_node);
+  if (output_num != 3) {
+    MS_LOG(EXCEPTION) << "Output number is " << output_num << ", but Adam needs 3 outputs.";
+  }
+  use_nesterov = AnfAlgo::GetNodeAttr<bool>(kernel_node, "use_nesterov");
+}
+
+bool AdamCPUKernel::Launch(const std::vector<kernel::AddressPtr> &inputs,
+                           const std::vector<kernel::AddressPtr> & /*workspace*/,
+                           const std::vector<kernel::AddressPtr> &outputs) {
+  if (inputs.size() != 10) {
+    MS_LOG(EXCEPTION) << "Input number is " << inputs.size() << ", but Adam needs 10 inputs.";
+  }
+  if (outputs.size() != 3) {
+    MS_LOG(EXCEPTION) << "Output number is " << outputs.size() << ", but Adam needs 3 outputs.";
+  }
+  if (inputs[0]->size != inputs[1]->size || inputs[0]->size != inputs[2]->size || inputs[0]->size != inputs[9]->size) {
+    MS_LOG(EXCEPTION) << "Error input data size!";
+  }
+  size_t f_size = sizeof(float);
+  if (inputs[3]->size != f_size || inputs[4]->size != f_size || inputs[5]->size != f_size ||
+      inputs[6]->size != f_size || inputs[7]->size != f_size || inputs[8]->size != f_size) {
+    MS_LOG(EXCEPTION) << "The attribute beta_power, beta, lr and epsilon must be float!";
+  }
+  auto var = reinterpret_cast<float *>(inputs[0]->addr);
+  auto m = reinterpret_cast<float *>(inputs[1]->addr);
+  auto v = reinterpret_cast<float *>(inputs[2]->addr);
+  float beta1_power = reinterpret_cast<float *>(inputs[3]->addr)[0];
+  float beta2_power = reinterpret_cast<float *>(inputs[4]->addr)[0];
+  float lr = reinterpret_cast<float *>(inputs[5]->addr)[0];
+  float beta1 = reinterpret_cast<float *>(inputs[6]->addr)[0];
+  float beta2 = reinterpret_cast<float *>(inputs[7]->addr)[0];
+  float epsilon = reinterpret_cast<float *>(inputs[8]->addr)[0];
+  auto gradient = reinterpret_cast<float *>(inputs[9]->addr);
+  if (beta1_power == 1) {
+    MS_LOG(EXCEPTION) << "The beta1_power can't be set 1.";
+  }
+  float new_lr = lr * std::sqrt(1.0 - beta2_power) / (1 - beta1_power);
+
+  // multithreading
+  size_t lens = inputs[0]->size > 0 ? static_cast<size_t>(inputs[0]->size / sizeof(float)) : 1;
+  auto max_thread_num = std::thread::hardware_concurrency();
+  size_t thread_num = lens < 128 * max_thread_num ? std::ceil(lens / 128.0) : max_thread_num;
+  MS_LOG(INFO) << "lens=" << lens << "; use thread_num=" << thread_num << "; max_thread_num: " << max_thread_num;
+  std::vector<std::thread> threads;
+  threads.reserve(thread_num);
+  size_t start = 0;
+  size_t once_compute_size = (lens + thread_num - 1) / thread_num;
+  while (start < lens) {
+    size_t end = (start + once_compute_size) > lens ? lens : (start + once_compute_size);
+    threads.emplace_back(std::thread(&AdamCPUKernel::LaunchAdam<float>, this, var, m, v, new_lr, beta1, beta2, epsilon,
+                                     gradient, start, end));
+    start += once_compute_size;
+  }
+  for (size_t i = 0; i < threads.size(); ++i) {
+    threads[i].join();
+  }
+
+  return true;
+}
+}  // namespace kernel
+}  // namespace mindspore

mindspore/ccsrc/backend/kernel_compiler/cpu/adam_cpu_kernel.h

@@ -0,0 +1,61 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_ADAM_CPU_KERNEL_H_
+#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_ADAM_CPU_KERNEL_H_
+
+#include <vector>
+#include <memory>
+#include "backend/kernel_compiler/cpu/cpu_kernel.h"
+#include "backend/kernel_compiler/cpu/cpu_kernel_factory.h"
+
+namespace mindspore {
+namespace kernel {
+class AdamCPUKernel : public CPUKernel {
+ public:
+  AdamCPUKernel() = default;
+  ~AdamCPUKernel() override = default;
+  template <typename T>
+  void LaunchAdam(T *var, T *m, T *v, float lr, float beta1, float beta2, float epsilon, const T *gradient,
+                  size_t start, size_t end);
+  void InitKernel(const CNodePtr &kernel_node) override;
+
+  bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
+              const std::vector<AddressPtr> &outputs) override;
+
+ private:
+  bool use_nesterov{false};
+};
+
+MS_REG_CPU_KERNEL(Adam,
+                  KernelAttr()
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddInputAttr(kNumberTypeFloat32)
+                    .AddOutputAttr(kNumberTypeFloat32)
+                    .AddOutputAttr(kNumberTypeFloat32)
+                    .AddOutputAttr(kNumberTypeFloat32),
+                  AdamCPUKernel)
+}  // namespace kernel
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_ADAM_CPU_KERNEL_H_

mindspore/ccsrc/backend/kernel_compiler/cpu/assign_cpu_kernel.cc

@@ -0,0 +1,80 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/kernel_compiler/cpu/assign_cpu_kernel.h"
+#include <string>
+#include "runtime/device/cpu/cpu_device_address.h"
+
+namespace mindspore {
+namespace kernel {
+void AssignCPUKernel::InitKernel(const CNodePtr &kernel_node) {
+  auto input_x_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
+  auto input_y_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
+
+  for (size_t i = 0; i < input_x_shape.size(); ++i) {
+    batch_size_ *= input_x_shape[i];
+  }
+
+  if (input_x_shape.size() != input_y_shape.size()) MS_LOG(EXCEPTION) << "x y must be same shape";
+  for (size_t i = 0; i < input_x_shape.size(); ++i) {
+    if (input_x_shape[i] != input_y_shape[i]) {
+      MS_LOG(EXCEPTION) << "x y must be same shape";
+    }
+  }
+  input_x_dtype_ = AnfAlgo::GetPrevNodeOutputInferDataType(kernel_node, 0);
+
+  if (input_x_dtype_ == kNumberTypeFloat32 || input_x_dtype_ == kNumberTypeInt32) {
+    input_x_dtype_size_ = 4;
+  } else if (input_x_dtype_ == kNumberTypeFloat64 || input_x_dtype_ == kNumberTypeInt64) {
+    input_x_dtype_size_ = 8;
+  } else {
+    MS_LOG(EXCEPTION) << "input_x dtype only support float32, float64, int32, int64";
+  }
+}
+
+bool AssignCPUKernel::Launch(const std::vector<kernel::AddressPtr> &inputs,
+                             const std::vector<kernel::AddressPtr> & /*workspace*/,
+                             const std::vector<kernel::AddressPtr> &outputs) {
+  if (input_x_dtype_ == kNumberTypeInt32) {
+    LaunchKernel<int>(inputs, outputs);
+  } else if (input_x_dtype_ == kNumberTypeInt64) {
+    LaunchKernel<int64_t>(inputs, outputs);
+  } else if (input_x_dtype_ == kNumberTypeFloat32) {
+    LaunchKernel<float>(inputs, outputs);
+  } else if (input_x_dtype_ == kNumberTypeFloat64) {
+    LaunchKernel<double>(inputs, outputs);
+  } else {
+    MS_LOG(ERROR) << "indices dtype only support float32, float64, int32, int64";
+    return false;
+  }
+  return true;
+}
+
+template <typename T>
+void AssignCPUKernel::LaunchKernel(const std::vector<AddressPtr> &inputs,
+                                   const std::vector<kernel::AddressPtr> &outputs) {
+  T *input_x = reinterpret_cast<T *>(inputs[0]->addr);
+  T *input_y = reinterpret_cast<T *>(inputs[1]->addr);
+  size_t total_size = input_x_dtype_size_ * batch_size_;
+
+  int ret = memcpy_s(input_x, total_size, input_y, total_size);
+  if (ret != 0) {
+    MS_LOG(EXCEPTION) << "memcpy_s error, errorno" << ret;
+  }
+}
+
+}  // namespace kernel
+}  // namespace mindspore

mindspore/ccsrc/backend/kernel_compiler/cpu/assign_cpu_kernel.h

@@ -0,0 +1,67 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_ASSIGN_CPU_KERNEL_H_
+#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_ASSIGN_CPU_KERNEL_H_
+
+#include <vector>
+#include <memory>
+#include <unordered_map>
+#include "backend/kernel_compiler/cpu/cpu_kernel.h"
+#include "backend/kernel_compiler/cpu/cpu_kernel_factory.h"
+
+namespace mindspore {
+namespace kernel {
+class AssignCPUKernel : public CPUKernel {
+ public:
+  AssignCPUKernel() = default;
+  ~AssignCPUKernel() override = default;
+
+  void InitKernel(const CNodePtr &kernel_node) override;
+
+  bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
+              const std::vector<AddressPtr> &outputs) override;
+
+  template <typename T>
+  void LaunchKernel(const std::vector<AddressPtr> &inputs, const std::vector<kernel::AddressPtr> &outputs);
+
+ private:
+  size_t batch_size_{1};
+  TypeId input_x_dtype_{kTypeUnknown};
+  size_t input_x_dtype_size_ = 4;
+};
+
+MS_REG_CPU_KERNEL(
+  Assign, KernelAttr().AddInputAttr(kNumberTypeInt32).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeInt32),
+  AssignCPUKernel);
+
+MS_REG_CPU_KERNEL(
+  Assign, KernelAttr().AddInputAttr(kNumberTypeInt64).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeInt64),
+  AssignCPUKernel);
+
+MS_REG_CPU_KERNEL(
+  Assign,
+  KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32),
+  AssignCPUKernel);
+
+MS_REG_CPU_KERNEL(
+  Assign,
+  KernelAttr().AddInputAttr(kNumberTypeFloat64).AddInputAttr(kNumberTypeFloat64).AddOutputAttr(kNumberTypeFloat64),
+  AssignCPUKernel);
+
+}  // namespace kernel
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_CPU_UPDATE_CACHE_CPU_KERNEL_H_

mindspore/ccsrc/backend/kernel_compiler/cpu/reduce_cpu_kernel.cc

@@ -39,6 +39,7 @@ void ReduceCPUKernel::InitKernel(const CNodePtr &kernel_node) {
   }
   shape_ = AnfAlgo::GetInputDeviceShape(kernel_node, 0);
   auto axis_addr = AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr(AXIS);
+
   if (axis_addr->isa<ValueTuple>()) {
     auto attr_axis = AnfAlgo::GetNodeAttr<std::vector<int>>(kernel_node, AXIS);
     if (attr_axis.size() > shape_.size()) {
@@ -47,18 +48,24 @@ void ReduceCPUKernel::InitKernel(const CNodePtr &kernel_node) {
       axis_.push_back(shape_.size() - 1);
     } else {
       for (auto axis : attr_axis) {
+        while (axis < 0) {
+          axis += SizeToInt(shape_.size());
+        }
         if (IntToSize(axis) >= (shape_.size())) {
           MS_LOG(EXCEPTION) << "axis value is oversize.";
         }
-        axis < 0 ? axis_.push_back(axis + shape_.size()) : axis_.push_back(axis);
+        axis_.push_back(IntToSize(axis));
       }
     }
   } else if (axis_addr->isa<Int32Imm>()) {
     int axis = AnfAlgo::GetNodeAttr<int>(kernel_node, AXIS);
-    if (axis >= 0 && IntToSize(axis) >= shape_.size()) {
+    while (axis < 0) {
+      axis += SizeToInt(shape_.size());
+    }
+    if (IntToSize(axis) >= shape_.size()) {
       MS_LOG(EXCEPTION) << "axis value is oversize.";
     }
-    axis < 0 ? axis_.push_back(axis + shape_.size()) : axis_.push_back(axis);
+    axis_.push_back(IntToSize(axis));
   } else {
     MS_LOG(EXCEPTION) << "Attribute axis type is invalid.";
   }

mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/ctcloss_impl.cu

@@ -0,0 +1,446 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 <limits>
+#include "ctcloss_impl.cuh"
+#include "runtime/device/gpu/cuda_common.h"
+template <typename T>
+__device__ T LogSumExp(const T logprob1, const T logprob2) {
+  if (logprob1 == logprob2 && logprob1 == -std::numeric_limits<T>::infinity()) {
+    return logprob1;
+  } else {
+    return (logprob1 > logprob2) ? logprob1 + log1pf(expf(logprob2 - logprob1))
+                                 : logprob2 + log1pf(expf(logprob1 - logprob2));
+  }
+}
+
+template <typename T>
+__global__ void CalculateFwdVarKernel(T *log_alpha_b, int *label_value_with_blank, T *softmax_probs,
+                                      const int *sequence_length, bool ctc_merge_repeated, int batch, int SOffSet,
+                                      int maxtime, int blank, int *label_squence_length, int *cum_labels_length,
+                                      bool ignore_longer_outputs_than_inputs) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    if (sequence_length[i] == 0 ||
+        (ignore_longer_outputs_than_inputs && label_squence_length[i] > sequence_length[i])) {
+    } else {
+      T *log_alpha_b_cur = &log_alpha_b[i * SOffSet * maxtime];
+      int *label_value_with_blank_cur = &label_value_with_blank[0];
+      if (i > 0) {
+        label_value_with_blank_cur = &label_value_with_blank[2 * cum_labels_length[i - 1] + i];
+      }
+      int numclass = blank + 1;
+      int U = 2 * label_squence_length[i] + 1;
+      int Ti = sequence_length[i];
+      int low = 0;
+      int high = 0;
+      log_alpha_b_cur[0] = log(softmax_probs[i * numclass + blank]);
+      int label0 = blank;
+      if (U > 1) {
+        label0 = label_value_with_blank_cur[1];
+        log_alpha_b_cur[maxtime] = log(softmax_probs[i * numclass + label0]);
+      }
+      for (int t = 1; t < Ti; ++t) {
+        low = 0;
+        high = U;
+        int low_limit = U - (2 * (Ti - t));
+        int high_limit = 2 * (t + 1);
+        if (low_limit > low) {
+          low = low_limit;
+        }
+        if (high_limit < U) {
+          high = high_limit;
+        }
+        for (int u = low; u < high; ++u) {
+          T sum_log_alpha = -std::numeric_limits<T>::infinity();
+          if (ctc_merge_repeated || label_value_with_blank_cur[u] == blank) {
+            sum_log_alpha = log_alpha_b_cur[u * maxtime + t - 1];
+          }
+          if (u > 0) {
+            sum_log_alpha = LogSumExp(sum_log_alpha, log_alpha_b_cur[(u - 1) * maxtime + t - 1]);
+          }
+          if (u > 1) {
+            const bool matching_labels_merge =
+              ctc_merge_repeated && (label_value_with_blank_cur[u] == label_value_with_blank_cur[u - 2]);
+            if (label_value_with_blank_cur[u] != blank && !matching_labels_merge) {
+              sum_log_alpha = LogSumExp(sum_log_alpha, log_alpha_b_cur[(u - 2) * maxtime + t - 1]);
+            }
+          }
+          log_alpha_b_cur[u * maxtime + t] =
+            log(softmax_probs[i * numclass + label_value_with_blank_cur[u] + t * numclass * batch]) + sum_log_alpha;
+        }
+      }
+    }
+  }
+}
+
+template <typename T>
+__global__ void CalculateBwdVarKernel(T *log_beta_b, int *label_value_with_blank, T *softmax_probs,
+                                      const int *sequence_length, bool ctc_merge_repeated, int batch, int SOffSet,
+                                      int maxtime, int blank, int *label_squence_length, int *cum_labels_length,
+                                      bool ignore_longer_outputs_than_inputs) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    if (sequence_length[i] == 0 ||
+        (ignore_longer_outputs_than_inputs && label_squence_length[i] > sequence_length[i])) {
+    } else {
+      T *log_beta_b_cur = &log_beta_b[i * SOffSet * maxtime];
+      int *label_value_with_blank_cur = &label_value_with_blank[0];
+      if (i > 0) {
+        label_value_with_blank_cur = &label_value_with_blank[2 * cum_labels_length[i - 1] + i];
+      }
+      int numclass = blank + 1;
+      int U = 2 * label_squence_length[i] + 1;
+      int Ti = sequence_length[i];
+      int low = 0;
+      int high = 0;
+      if (U > 1) {
+        for (int u = U - 2; u < U; ++u) {
+          log_beta_b_cur[u * maxtime + Ti - 1] = 0;
+        }
+      } else {
+        log_beta_b_cur[Ti - 1] = 0;
+        log_beta_b_cur[Ti - 2] = 0;
+      }
+      for (int t = Ti - 2; t >= 0; --t) {
+        low = 0;
+        high = U;
+        int low_limit = U - (2 * (Ti - t));
+        int high_limit = 2 * (t + 1);
+        if (low_limit > low) {
+          low = low_limit;
+        }
+        if (high_limit < U) {
+          high = high_limit;
+        }
+        for (int u = low; u < high; ++u) {
+          if (ctc_merge_repeated || label_value_with_blank_cur[u] == blank) {
+            log_beta_b_cur[u * maxtime + t] = LogSumExp(
+              log_beta_b_cur[u * maxtime + t],
+              log_beta_b_cur[u * maxtime + t + 1] +
+                log(softmax_probs[i * numclass + label_value_with_blank_cur[u] + (t + 1) * numclass * batch]));
+          }
+          if (u + 1 < U) {
+            log_beta_b_cur[u * maxtime + t] = LogSumExp(
+              log_beta_b_cur[u * maxtime + t],
+              log_beta_b_cur[(u + 1) * maxtime + t + 1] +
+                log(softmax_probs[i * numclass + label_value_with_blank_cur[u + 1] + (t + 1) * numclass * batch]));
+          }
+          if (u + 2 < U) {
+            const bool matching_labels_merge =
+              ctc_merge_repeated && (label_value_with_blank_cur[u] == label_value_with_blank_cur[u + 2]);
+            if (label_value_with_blank_cur[u] != blank && !matching_labels_merge) {
+              log_beta_b_cur[u * maxtime + t] = LogSumExp(
+                log_beta_b_cur[u * maxtime + t],
+                log_beta_b_cur[(u + 2) * maxtime + t + 1] +
+                  log(softmax_probs[i * numclass + label_value_with_blank_cur[u + 2] + (t + 1) * numclass * batch]));
+            }
+          }
+        }
+      }
+    }
+  }
+}
+
+template <typename T>
+__global__ void ProbInitKernel(T *prob_num, int size) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += blockDim.x * gridDim.x) {
+    prob_num[i] = -std::numeric_limits<T>::infinity();
+  }
+}
+template <typename T>
+__global__ void LogBInitKernel(T *log_b, int log_prob_size) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < log_prob_size; i += blockDim.x * gridDim.x) {
+    log_b[i] = -std::numeric_limits<T>::infinity();
+  }
+}
+
+template <typename T>
+__global__ void CTCLossKernel(T *log_alpha_b, T *log_beta_b, T *softmax_probs, int *label_value_with_blank, int batch,
+                              int SOffSet, int maxtime, int numclass, const int *sequence_length,
+                              int *label_squence_length, int *cum_labels_length, T *cost, T *grads, T *prob_num,
+                              bool ignore_longer_outputs_than_inputs) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    if (sequence_length[i] == 0 ||
+        (ignore_longer_outputs_than_inputs && label_squence_length[i] > sequence_length[i])) {
+    } else {
+      T *grad_cur = &grads[i * numclass];
+      const T *softmax_probs_cur = &softmax_probs[i * numclass];
+      T *prob_num_cur = &prob_num[i * numclass];
+      int U = 2 * label_squence_length[i] + 1;
+      T log_pzx = -std::numeric_limits<T>::infinity();
+      const T *log_alpha_b_cur = &log_alpha_b[i * SOffSet * maxtime];
+      const T *log_beta_b_cur = &log_beta_b[i * SOffSet * maxtime];
+      int *label_value_with_blank_cur = &label_value_with_blank[0];
+      if (i > 0) {
+        label_value_with_blank_cur = &label_value_with_blank[2 * cum_labels_length[i - 1] + i];
+      }
+      for (int u = 0; u < U; ++u) {
+        log_pzx = LogSumExp(log_pzx, log_alpha_b_cur[u * maxtime] + log_beta_b_cur[u * maxtime]);
+      }
+      cost[i] = -log_pzx;
+      // grad
+      int L = numclass;
+      int Ti = sequence_length[i];
+      if (log_pzx == -std::numeric_limits<T>::infinity()) {
+        for (int t = 0; t < Ti; ++t) {
+          for (int l = 0; l < L; ++l) {
+            grad_cur[t * numclass * batch + l] = softmax_probs_cur[t * numclass * batch + l];
+          }
+        }
+      } else {
+        for (int t = 0; t < Ti; ++t) {
+          for (int u = 0; u < U; ++u) {
+            int l = label_value_with_blank_cur[u];
+            prob_num_cur[t * batch * numclass + l] =
+              LogSumExp(prob_num_cur[t * batch * numclass + l],
+                        log_alpha_b_cur[u * maxtime + t] + log_beta_b_cur[u * maxtime + t]);
+          }
+          for (int l = 0; l < L; ++l) {
+            grad_cur[t * numclass * batch + l] =
+              softmax_probs_cur[t * numclass * batch + l] - expf(prob_num_cur[t * batch * numclass + l] - log_pzx);
+          }
+        }
+      }
+    }
+  }
+}
+
+template <typename T>
+__global__ void InnerSoftMaxKernel(const T *probs, T *softmax_probs, const int *sequence_length, int max_time,
+                                   int batch, int numclass) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch * max_time; i += blockDim.x * gridDim.x) {
+    int k = i / batch;
+    int m = i % batch;
+    if (k < sequence_length[m]) {
+      T maxCoeff = 0.;
+      T sumCoeff = 0.;
+      for (int j = i * numclass; j < (i + 1) * numclass; ++j) {
+        if (probs[j] > maxCoeff) {
+          maxCoeff = probs[j];
+        }
+      }
+      for (int j = i * numclass; j < (i + 1) * numclass; ++j) {
+        sumCoeff += exp(probs[j] - maxCoeff);
+        softmax_probs[j] = exp(probs[j] - maxCoeff);
+      }
+      for (int j = i * numclass; j < (i + 1) * numclass; ++j) {
+        softmax_probs[j] /= sumCoeff;
+      }
+    }
+  }
+}
+
+__global__ void GenLabelValuePCRKernel(int *label_value_sp, int *label_value_pcr, int *label_squence_length,
+                                       int *cum_labels_length, int batch) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    int L = label_squence_length[i];
+    label_squence_length[i] = 0;
+    int offset = 0;
+    if (i > 0) {
+      offset = cum_labels_length[i - 1];
+    }
+    for (int l = offset; l < L; ++l) {
+      if (l == offset || label_value_sp[l] != label_value_sp[l - 1]) {
+        label_value_pcr[offset + label_squence_length[i]++] = label_value_sp[l];
+      }
+    }
+  }
+}
+
+__global__ void UpdateLengthKernel(int *label_squence_length, int *cum_labels_length, int *max_labels_length,
+                                   int batch) {
+  max_labels_length[0] = 0;
+  for (int i = 0; i < batch; ++i) {
+    if (label_squence_length[i] > max_labels_length[0]) {
+      max_labels_length[0] = label_squence_length[i];
+    }
+    if (i == 0) {
+      cum_labels_length[i] = label_squence_length[i];
+    } else {
+      cum_labels_length[i] = label_squence_length[i] + cum_labels_length[i - 1];
+    }
+  }
+}
+
+template <typename T>
+void CalculateBwdVar(T *log_beta_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
+                     bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
+                     int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream) {
+  int log_prob_size = SOffSet * batch * maxtime;
+  LogBInitKernel<<<GET_BLOCKS(log_prob_size), GET_THREADS, 0, stream>>>(log_beta_b, log_prob_size);
+  CalculateBwdVarKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(
+    log_beta_b, label_value_with_blank, softmax_probs, sequence_length, ctc_merge_repeated, batch, SOffSet, maxtime,
+    blank, label_squence_length, cum_labels_length, ignore_longer_outputs_than_inputs);
+}
+
+template <typename T>
+void CalculateFwdVar(T *log_alpha_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
+                     bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
+                     int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream) {
+  int log_prob_size = SOffSet * batch * maxtime;
+  LogBInitKernel<<<GET_BLOCKS(log_prob_size), GET_THREADS, 0, stream>>>(log_alpha_b, log_prob_size);
+  CalculateFwdVarKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(
+    log_alpha_b, label_value_with_blank, softmax_probs, sequence_length, ctc_merge_repeated, batch, SOffSet, maxtime,
+    blank, label_squence_length, cum_labels_length, ignore_longer_outputs_than_inputs);
+}
+
+template <typename T>
+void InnerSoftMax(const T *probs, T *softmax_probs, const int *sequence_length, int max_time, int batch, int numclass,
+                  cudaStream_t stream) {
+  InnerSoftMaxKernel<<<GET_BLOCKS(batch * max_time), GET_THREADS, 0, stream>>>(probs, softmax_probs, sequence_length,
+                                                                               max_time, batch, numclass);
+}
+
+__global__ void GenLabelWithBlankKernel(int *label_value, int *label_value_with_blank, int *label_squence_length,
+                                        int *precum_labels_length, int *cum_labels_length, int batch, int blank) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    int offset = 0;
+    int offset1 = 0;
+    if (i > 0) {
+      offset = 2 * cum_labels_length[i - 1] + i;
+      offset1 = precum_labels_length[i - 1];
+    }
+    for (int j = 0; j < label_squence_length[i]; ++j) {
+      label_value_with_blank[offset + 2 * j] = blank;
+      label_value_with_blank[offset + 2 * j + 1] = label_value[offset1 + j];
+    }
+    label_value_with_blank[offset + 2 * label_squence_length[i]] = blank;
+  }
+}
+
+void GenLabelWithBlank(int *label_value, int *label_value_with_blank, int *label_squence_length,
+                       int *precum_labels_length, int *cum_labels_length, int batch, int blank, cudaStream_t stream) {
+  GenLabelWithBlankKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(
+    label_value, label_value_with_blank, label_squence_length, precum_labels_length, cum_labels_length, batch, blank);
+}
+
+void GenLabelValuePCR(int *label_value_sp, int *label_value_pcr, int *label_squence_length, int *cum_labels_length,
+                      int *max_labels_length, int batch, cudaStream_t stream) {
+  GenLabelValuePCRKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(label_value_sp, label_value_pcr,
+                                                                        label_squence_length, cum_labels_length, batch);
+  UpdateLengthKernel<<<1, 1, 0, stream>>>(label_squence_length, cum_labels_length, max_labels_length, batch);
+}
+
+__global__ void GenLabelValueKernel(int *label_value_sp, const int64_t *label_indices, const int *label_values,
+                                    int *label_squence_length, int *cum_labels_length, int size) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += blockDim.x * gridDim.x) {
+    int64_t b = label_indices[i * 2];
+    int offset = 0;
+    if (b > 0) {
+      offset = cum_labels_length[b - 1];
+    }
+    int64_t index = offset + label_indices[i * 2 + 1];
+    label_value_sp[index] = label_values[i];
+  }
+}
+__global__ void LabelValueInitKernel(int *label_value_sp, int size, int blank) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += blockDim.x * gridDim.x) {
+    label_value_sp[i] = blank;
+  }
+}
+__global__ void RecalculateLengthKernel(int *label_value_sp, int *label_squence_length, int *cum_labels_length,
+                                        int batch, int blank) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch; i += blockDim.x * gridDim.x) {
+    int offset = 0;
+    if (i > 0) {
+      offset = cum_labels_length[i - 1];
+    }
+    int L = label_squence_length[i];
+    label_squence_length[i] = 0;
+    for (int j = offset; j < offset + L; ++j) {
+      if (label_value_sp[j] >= blank) {
+        break;
+      } else {
+        label_squence_length[i]++;
+      }
+    }
+  }
+}
+void GenLabelValue(int *label_value_sp, const int64_t *label_indices, const int *label_values,
+                   int *label_squence_length, int *cum_labels_length, int *max_labels_length, int size, int blank,
+                   int batch, cudaStream_t stream) {
+  LabelValueInitKernel<<<GET_BLOCKS(size), GET_THREADS, 0, stream>>>(label_value_sp, size, blank);
+  GenLabelValueKernel<<<GET_BLOCKS(size), GET_THREADS, 0, stream>>>(label_value_sp, label_indices, label_values,
+                                                                    label_squence_length, cum_labels_length, size);
+  RecalculateLengthKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(label_value_sp, label_squence_length,
+                                                                         cum_labels_length, batch, blank);
+  UpdateLengthKernel<<<1, 1, 0, stream>>>(label_squence_length, cum_labels_length, max_labels_length, batch);
+}
+
+__global__ void CalculatePreLengthKernel(int *label_squence_length, int *precum_labels_length, int *cum_labels_length,
+                                         int *max_labels_length, const int64_t *label_indices, int batch, int size) {
+  max_labels_length[0] = 0;
+  for (int i = 0; i < size; ++i) {
+    label_squence_length[label_indices[i * 2]]++;
+    if (max_labels_length[0] < label_indices[i * 2]) {
+      max_labels_length[0] = label_indices[i * 2];
+    }
+  }
+  precum_labels_length[0] = label_squence_length[0];
+  cum_labels_length[0] = label_squence_length[0];
+  for (int i = 1; i < batch; ++i) {
+    cum_labels_length[i] = cum_labels_length[i - 1] + label_squence_length[i];
+    precum_labels_length[i] = precum_labels_length[i - 1] + label_squence_length[i];
+  }
+}
+
+__global__ void CalculateMaxSequenceKernel(const int *sequence_length, int *max_labels_length, int batch) {
+  max_labels_length[0] = 0;
+  for (int i = 0; i < batch; ++i) {
+    if (sequence_length[i] > max_labels_length[0]) {
+      max_labels_length[0] = sequence_length[i];
+    }
+  }
+}
+
+void CalculateMaxSequence(const int *sequence_length, int *max_labels_length, int batch, cudaStream_t stream) {
+  CalculateMaxSequenceKernel<<<1, 1, 0, stream>>>(sequence_length, max_labels_length, batch);
+}
+
+void CalculatePreLength(int *label_squence_length, int *precum_labels_length, int *cum_labels_length,
+                        int *max_labels_length, const int64_t *label_indices, int batch, int size,
+                        cudaStream_t stream) {
+  CalculatePreLengthKernel<<<1, 1, 0, stream>>>(label_squence_length, precum_labels_length, cum_labels_length,
+                                                max_labels_length, label_indices, batch, size);
+}
+
+template <typename T>
+void CTCLoss(T *log_alpha_b, T *log_beta_b, T *softmax_probs, int *label_value_with_blank, int batch, int SOffSet,
+             int maxtime, int numclass, const int *sequence_length, int *label_squence_length, int *cum_labels_length,
+             T *cost, T *grads, T *prob_num, bool ignore_longer_outputs_than_inputs, cudaStream_t stream) {
+  ProbInitKernel<<<GET_BLOCKS(maxtime * batch * numclass), GET_THREADS, 0, stream>>>(prob_num,
+                                                                                     maxtime * batch * numclass);
+  CTCLossKernel<<<GET_BLOCKS(batch), GET_THREADS, 0, stream>>>(
+    log_alpha_b, log_beta_b, softmax_probs, label_value_with_blank, batch, SOffSet, maxtime, numclass, sequence_length,
+    label_squence_length, cum_labels_length, cost, grads, prob_num, ignore_longer_outputs_than_inputs);
+}
+
+template void CalculateFwdVar<float>(float *log_alpha_b, int *label_value_with_blank, float *softmax_probs,
+                                     const int *sequence_length, bool ctc_merge_repeated, int batch, int SOffSet,
+                                     int maxtime, int blank, int *label_squence_length, int *cum_labels_length,
+                                     bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
+
+template void CalculateBwdVar<float>(float *log_beta_b, int *label_value_with_blank, float *softmax_probs,
+                                     const int *sequence_length, bool ctc_merge_repeated, int batch, int SOffSet,
+                                     int maxtime, int blank, int *label_squence_length, int *cum_labels_length,
+                                     bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
+
+template void InnerSoftMax<float>(const float *probs, float *softmax_probs, const int *sequence_length, int max_time,
+                                  int batch, int numclass, cudaStream_t stream);
+
+template void CTCLoss<float>(float *log_alpha_b, float *log_beta_b, float *softmax_probs, int *label_value_with_blank,
+                             int batch, int SOffSet, int maxtime, int numclass, const int *sequence_length,
+                             int *label_squence_length, int *cum_labels_length, float *cost, float *grads,
+                             float *prob_num, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);

mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/ctcloss_impl.cuh

@@ -0,0 +1,51 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+
+#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH
+#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH
+
+template <typename T>
+void CalculateFwdVar(T *log_alpha_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
+                     bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
+                     int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
+
+template <typename T>
+void CalculateBwdVar(T *log_beta_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
+                     bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
+                     int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
+
+template <typename T>
+void InnerSoftMax(const T *probs, T *softmax_cost, const int *sequence_length, int max_time, int batch, int numclass,
+                  cudaStream_t stream);
+
+void GenLabelValuePCR(int *label_value_sp, int *label_value_pcr, int *label_squence_length, int *cum_labels_length,
+                      int *max_labels_length, int batch, cudaStream_t stream);
+
+void GenLabelWithBlank(int *label_value, int *label_value_with_blank, int *label_squence_length,
+                       int *precum_labels_length, int *cum_labels_length, int batch, int blank, cudaStream_t stream);
+
+void GenLabelValue(int *label_value_sp, const int64_t *label_indices, const int *label_values,
+                   int *label_squence_length, int *cum_labels_length, int *max_labels_length, int size, int blank,
+                   int batch, cudaStream_t stream);
+
+void CalculatePreLength(int *label_squence_length, int *precum_labels_length, int *cum_labels_length,
+                        int *max_labels_length, const int64_t *label_indices, int batch, int size, cudaStream_t stream);
+void CalculateMaxSequence(const int *sequence_length, int *max_labels_length, int batch, cudaStream_t stream);
+template <typename T>
+void CTCLoss(T *log_alpha_b, T *log_beta_b, T *softmax_probs, int *label_value_with_blank, int batch, int SOffSet,
+             int maxtime, int numclass, const int *sequence_length, int *label_squence_length, int *cum_labels_length,
+             T *cost, T *grads, T *prob_num, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
+#endif  // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH

mindspore/ccsrc/backend/kernel_compiler/gpu/math/unary_op_gpu_kernel.cc

@@ -44,8 +44,12 @@ MS_REG_GPU_KERNEL_ONE(Square, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddO
                       UnaryOpGpuKernel, half)
 MS_REG_GPU_KERNEL_ONE(Sqrt, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32),
                       UnaryOpGpuKernel, float)
+MS_REG_GPU_KERNEL_ONE(Sqrt, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),
+                      UnaryOpGpuKernel, half)
 MS_REG_GPU_KERNEL_ONE(Rsqrt, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32),
                       UnaryOpGpuKernel, float)
+MS_REG_GPU_KERNEL_ONE(Rsqrt, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),
+                      UnaryOpGpuKernel, half)
 MS_REG_GPU_KERNEL_ONE(Sin, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32),
                       UnaryOpGpuKernel, float)
 MS_REG_GPU_KERNEL_ONE(Sin, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/ctcloss_gpu_kernel.cc

@@ -1,31 +1,31 @@
-/**
- * Copyright 2020 Huawei Technologies Co., Ltd
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * 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 "backend/kernel_compiler/gpu/nn/ctcloss_gpu_kernel.h"
-
-namespace mindspore {
-namespace kernel {
-MS_REG_GPU_KERNEL_ONE(CTCLossV2,
-                      KernelAttr()
-                        .AddInputAttr(kNumberTypeFloat32)
-                        .AddInputAttr(kNumberTypeInt32)
-                        .AddInputAttr(kNumberTypeInt32)
-                        .AddInputAttr(kNumberTypeInt32)
-                        .AddOutputAttr(kNumberTypeFloat32)
-                        .AddOutputAttr(kNumberTypeFloat32),
-                      CtcLossGpuKernel, float)
-}  // namespace kernel
-}  // namespace mindspore
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/kernel_compiler/gpu/nn/ctcloss_gpu_kernel.h"
+
+namespace mindspore {
+namespace kernel {
+MS_REG_GPU_KERNEL_ONE(CTCLoss,
+                      KernelAttr()
+                        .AddInputAttr(kNumberTypeFloat32)
+                        .AddInputAttr(kNumberTypeInt64)
+                        .AddInputAttr(kNumberTypeInt32)
+                        .AddInputAttr(kNumberTypeInt32)
+                        .AddOutputAttr(kNumberTypeFloat32)
+                        .AddOutputAttr(kNumberTypeFloat32),
+                      CtcLossGpuKernel, float)
+}  // namespace kernel
+}  // namespace mindspore

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/ctcloss_gpu_kernel.h

@@ -1,192 +1,233 @@
-/**
- * Copyright 2020 Huawei Technologies Co., Ltd
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * 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.
- */
-
-#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_
-#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_
-
-#include <cuda_runtime_api.h>
-#include <vector>
-#include "backend/kernel_compiler/gpu/gpu_kernel.h"
-#include "backend/kernel_compiler/gpu/gpu_kernel_factory.h"
-#include "runtime/device/gpu/gpu_memory_allocator.h"
-
-namespace mindspore {
-namespace kernel {
-template <typename T>
-class CtcLossGpuKernel : public GpuKernel {
- public:
-  CtcLossGpuKernel()
-      : cudnn_handle_(nullptr),
-        probs_desc_(nullptr),
-        ctcloss_desc_(nullptr),
-        label_size_(0),
-        input_lengths_size_(0),
-        label_lengths_size_(0) {}
-  ~CtcLossGpuKernel() override { DestroyResource(); }
-
-  const std::vector<size_t> &GetInputSizeList() const override { return input_size_list_; }
-  const std::vector<size_t> &GetOutputSizeList() const override { return output_size_list_; }
-  const std::vector<size_t> &GetWorkspaceSizeList() const override { return workspace_size_list_; }
-
-  bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,
-              const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
-    float *probs = GetDeviceAddress<float>(inputs, 0);
-    float *costs = GetDeviceAddress<float>(outputs, 0);
-    float *grads = GetDeviceAddress<float>(outputs, 1);
-
-    // Copy labels/input_lengths/label_length to host as cudnn7.x.x requires
-    int *labels_host = nullptr;
-    int *no_blank_labels_host = nullptr;
-    void *input_lengths_host = nullptr;
-    void *label_lengths_host = nullptr;
-    cudaStream_t stream = reinterpret_cast<cudaStream_t>(stream_ptr);
-    AllocHostMem(&labels_host, &no_blank_labels_host, &input_lengths_host, &label_lengths_host, inputs);
-    CopyToHostSync(labels_host, no_blank_labels_host, input_lengths_host, label_lengths_host, inputs, stream);
-
-    size_t workspace_size = 0;
-    CHECK_CUDNN_RET_WITH_EXCEPT(
-      cudnnGetCTCLossWorkspaceSize(
-        cudnn_handle_, probs_desc_, probs_desc_, reinterpret_cast<int *>(no_blank_labels_host),
-        reinterpret_cast<int *>(label_lengths_host), reinterpret_cast<int *>(input_lengths_host),
-        CUDNN_CTC_LOSS_ALGO_DETERMINISTIC, ctcloss_desc_, &workspace_size),
-      "cudnnGetCTCLossWorkspaceSize failed.");
-    void *workspace = device::gpu::GPUMemoryAllocator::GetInstance().AllocTensorMem(workspace_size);
-    if (workspace == nullptr) {
-      MS_LOG(EXCEPTION) << "Failed to alloc workspace, size: " << workspace_size;
-    }
-
-    CHECK_CUDNN_RET_WITH_EXCEPT(
-      cudnnCTCLoss(cudnn_handle_, probs_desc_, probs, reinterpret_cast<int *>(no_blank_labels_host),
-                   reinterpret_cast<int *>(label_lengths_host), reinterpret_cast<int *>(input_lengths_host), costs,
-                   probs_desc_, grads, CUDNN_CTC_LOSS_ALGO_DETERMINISTIC, ctcloss_desc_, workspace, workspace_size),
-      "cudnnCtcLoss failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
-
-    device::gpu::GPUMemoryAllocator::GetInstance().FreeTensorMem(workspace);
-    FreeHostMem(labels_host, no_blank_labels_host, input_lengths_host, label_lengths_host);
-    return true;
-  }
-  bool Init(const CNodePtr &kernel_node) override {
-    InitResource();
-    auto probs_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
-    if (probs_shape.size() != 3) {
-      MS_LOG(EXCEPTION) << "probs dims: " << probs_shape.size() << " not support.";
-    }
-    probs_dims_[0] = probs_shape[0];
-    probs_dims_[1] = probs_shape[1];
-    probs_dims_[2] = probs_shape[2];
-
-    auto labels_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
-    if (labels_dims.size() != 1 && labels_dims.size() != 2) {
-      MS_LOG(EXCEPTION) << "labels dims: " << labels_dims.size() << " not support.";
-    }
-    label_size_ = sizeof(int);
-    for (auto i : labels_dims) {
-      label_size_ *= i;
-    }
-
-    auto input_length_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 2);
-    input_lengths_size_ = input_length_dims[0] * sizeof(int);
-    auto label_length_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 3);
-    label_lengths_size_ = label_length_dims[0] * sizeof(int);
-    CHECK_CUDNN_RET_WITH_EXCEPT(
-      cudnnSetTensorNdDescriptorEx(probs_desc_, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, 3, probs_dims_),
-      "cudnnSetTensorNdDescriptorEx failed.");
-    CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetCTCLossDescriptorEx(ctcloss_desc_, CUDNN_DATA_FLOAT,
-                                                            CUDNN_LOSS_NORMALIZATION_SOFTMAX, CUDNN_PROPAGATE_NAN),
-                                "cudnnSetCTCLossDescriptorEx failed.");
-    InitSizeLists();
-    return true;
-  }
-
- protected:
-  void InitResource() override {
-    cudnn_handle_ = device::gpu::GPUDeviceManager::GetInstance().GetCudnnHandle();
-    CHECK_CUDNN_RET_WITH_EXCEPT(cudnnCreateTensorDescriptor(&probs_desc_), "cudnnCreateTensorDescriptor failed.");
-    CHECK_CUDNN_RET_WITH_EXCEPT(cudnnCreateCTCLossDescriptor(&ctcloss_desc_), "cudnnCreateCTCLossDescriptor failed.");
-  }
-
-  void InitSizeLists() override {
-    input_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(float));
-    input_size_list_.push_back(label_size_);
-    input_size_list_.push_back(input_lengths_size_);
-    input_size_list_.push_back(label_lengths_size_);
-
-    output_size_list_.push_back(probs_dims_[1] * sizeof(float));
-    output_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(float));
-  }
-
- private:
-  void DestroyResource() noexcept {
-    CHECK_CUDNN_RET_WITH_ERROR(cudnnDestroyCTCLossDescriptor(ctcloss_desc_), "cudnnDestroyCTCLossDescriptor failed.");
-    CHECK_CUDNN_RET_WITH_ERROR(cudnnDestroyTensorDescriptor(probs_desc_), "cudnnDestroyTensorDescriptor failed.");
-  }
-
-  void AllocHostMem(int **labels_host, int **no_blank_labels_host, void **input_lengths_host, void **label_lengths_host,
-                    const std::vector<AddressPtr> &inputs) {
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaMallocHost(labels_host, inputs[1]->size), "cudaMallocHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaMallocHost(no_blank_labels_host, inputs[1]->size), "cudaMallocHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaMallocHost(input_lengths_host, inputs[2]->size), "cudaMallocHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaMallocHost(label_lengths_host, inputs[3]->size), "cudaMallocHost failed.");
-  }
-
-  void FreeHostMem(int *labels_host, int *no_blank_labels_host, void *input_lengths_host, void *label_lengths_host) {
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaFreeHost(label_lengths_host), "cudaFreeHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaFreeHost(input_lengths_host), "cudaFreeHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaFreeHost(labels_host), "cudaFreeHost failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaFreeHost(no_blank_labels_host), "cudaFreeHost failed.");
-  }
-
-  void CopyToHostSync(int *labels_host, int *no_blank_labels_host, void *input_lengths_host, void *label_lengths_host,
-                      const std::vector<AddressPtr> &inputs, cudaStream_t stream) {
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(
-      cudaMemcpyAsync(labels_host, inputs[1]->addr, inputs[1]->size, cudaMemcpyDeviceToHost, stream),
-      "cudaMemcpyAsync failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(
-      cudaMemcpyAsync(input_lengths_host, inputs[2]->addr, inputs[2]->size, cudaMemcpyDeviceToHost, stream),
-      "cudaMemcpyAsync failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(
-      cudaMemcpyAsync(label_lengths_host, inputs[3]->addr, inputs[3]->size, cudaMemcpyDeviceToHost, stream),
-      "cudaMemcpyAsync failed.");
-    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
-
-    // remove blank element
-    size_t j = 0;
-    for (size_t i = 0; i < inputs[1]->size / sizeof(int); i++) {
-      if (labels_host[i] != 0) {
-        no_blank_labels_host[j] = labels_host[i];
-        j++;
-      }
-    }
-  }
-
-  std::vector<size_t> input_size_list_;
-  std::vector<size_t> output_size_list_;
-  std::vector<size_t> workspace_size_list_;
-
-  cudnnHandle_t cudnn_handle_;
-  cudnnTensorDescriptor_t probs_desc_;
-  cudnnCTCLossDescriptor_t ctcloss_desc_;
-  int probs_dims_[3] = {0};
-  int label_size_;
-  int input_lengths_size_;
-  int label_lengths_size_;
-};
-}  // namespace kernel
-}  // namespace mindspore
-
-#endif  // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+
+#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_
+#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_
+
+#include <cuda_runtime_api.h>
+#include <vector>
+#include <limits>
+#include "backend/kernel_compiler/gpu/gpu_kernel.h"
+#include "backend/kernel_compiler/gpu/gpu_kernel_factory.h"
+#include "runtime/device/gpu/gpu_memory_allocator.h"
+#include "backend/kernel_compiler/gpu/cuda_impl/ctcloss_impl.cuh"
+namespace mindspore {
+namespace kernel {
+template <typename T>
+class CtcLossGpuKernel : public GpuKernel {
+ public:
+  CtcLossGpuKernel()
+      : label_indice_size_(0),
+        label_size_(0),
+        squence_lengths_size_(0),
+        preprocess_collapse_repeated_(false),
+        ctc_merge_repeated_(true),
+        ignore_longer_outputs_than_inputs_(false) {}
+  ~CtcLossGpuKernel() override = default;
+
+  const std::vector<size_t> &GetInputSizeList() const override { return input_size_list_; }
+  const std::vector<size_t> &GetOutputSizeList() const override { return output_size_list_; }
+  const std::vector<size_t> &GetWorkspaceSizeList() const override { return workspace_size_list_; }
+
+  bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
+              const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
+    cudaStream_t stream = reinterpret_cast<cudaStream_t>(stream_ptr);
+    const T *probs = GetDeviceAddress<T>(inputs, 0);
+    const int64_t *label_indices = GetDeviceAddress<int64_t>(inputs, 1);
+    const int *label_values = GetDeviceAddress<int>(inputs, 2);
+    const int *sequence_length = GetDeviceAddress<int>(inputs, 3);
+    T *costs = GetDeviceAddress<T>(outputs, 0);
+    T *grads = GetDeviceAddress<T>(outputs, 1);
+    T *softmax_probs = GetDeviceAddress<T>(workspace, 0);
+    int *cum_labels_length = GetDeviceAddress<int>(workspace, 1);
+    int *label_squence_length = GetDeviceAddress<int>(workspace, 2);
+    int *label_value_sp = GetDeviceAddress<int>(workspace, 3);
+    int *label_value_pcr = GetDeviceAddress<int>(workspace, 4);
+    T *prob_num = GetDeviceAddress<T>(workspace, 5);
+    int *precum_labels_length = GetDeviceAddress<int>(workspace, 6);
+    int *max_labels_length = GetDeviceAddress<int>(workspace, 7);
+    int numclass = SizeToInt(probs_dims_[2]);
+    int batch = SizeToInt(probs_dims_[1]);
+    int max_time = SizeToInt(probs_dims_[0]);
+    int max_sequence = 0;
+    CalculateMaxSequence(sequence_length, max_labels_length, batch, stream);
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMemcpyAsync(&max_sequence, max_labels_length, sizeof(int), cudaMemcpyDeviceToHost, stream),
+      "cudaMemcpyAsync failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
+    if (max_time < max_sequence) {
+      MS_LOG(EXCEPTION) << "max_time should be greater than sequence length.";
+    }
+    InnerSoftMax(probs, softmax_probs, sequence_length, max_time, batch, numclass, stream);
+    MemsetForWS(label_value_pcr, cum_labels_length, label_squence_length, costs, grads, stream);
+    int max_labels_length_host = 0;
+    int batch_label = 0;
+    int *label_value_with_blank = nullptr;
+    T *log_alpha_b = nullptr;
+    T *log_beta_b = nullptr;
+    CalculatePreLength(label_squence_length, precum_labels_length, cum_labels_length, max_labels_length, label_indices,
+                       batch, label_size_ / sizeof(int), stream);
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMemcpyAsync(&batch_label, max_labels_length, sizeof(int), cudaMemcpyDeviceToHost, stream),
+      "cudaMemcpyAsync failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
+    if (batch != batch_label + 1) {
+      MS_LOG(EXCEPTION) << "label batch should be equal to input batch.";
+    }
+    GenLabelValue(label_value_sp, label_indices, label_values, label_squence_length, cum_labels_length,
+                  max_labels_length, label_size_ / sizeof(int), numclass - 1, batch, stream);
+    if (preprocess_collapse_repeated_) {
+      GenLabelValuePCR(label_value_sp, label_value_pcr, label_squence_length, cum_labels_length, max_labels_length,
+                       batch, stream);
+    }
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMemcpyAsync(&max_labels_length_host, max_labels_length, sizeof(int), cudaMemcpyDeviceToHost, stream),
+      "cudaMemcpyAsync failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
+    int SOffSet = 2 * max_labels_length_host + 1;
+    int log_prob_size = batch * SOffSet * max_time;
+    if (!ignore_longer_outputs_than_inputs_ && max_labels_length_host > max_time) {
+      MS_LOG(EXCEPTION) << "output size is greater than input size.";
+    }
+    MemManageForCus(&log_alpha_b, &log_beta_b, &label_value_with_blank, cum_labels_length, log_prob_size, batch,
+                    stream);
+
+    if (preprocess_collapse_repeated_) {
+      GenLabelWithBlank(label_value_pcr, label_value_with_blank, label_squence_length, precum_labels_length,
+                        cum_labels_length, batch, numclass - 1, stream);
+    } else {
+      GenLabelWithBlank(label_value_sp, label_value_with_blank, label_squence_length, precum_labels_length,
+                        cum_labels_length, batch, numclass - 1, stream);
+    }
+
+    CalculateFwdVar(log_alpha_b, label_value_with_blank, softmax_probs, sequence_length, ctc_merge_repeated_, batch,
+                    SOffSet, max_time, numclass - 1, label_squence_length, cum_labels_length,
+                    ignore_longer_outputs_than_inputs_, stream);
+    CalculateBwdVar(log_beta_b, label_value_with_blank, softmax_probs, sequence_length, ctc_merge_repeated_, batch,
+                    SOffSet, max_time, numclass - 1, label_squence_length, cum_labels_length,
+                    ignore_longer_outputs_than_inputs_, stream);
+    CTCLoss(log_alpha_b, log_beta_b, softmax_probs, label_value_with_blank, batch, SOffSet, max_time, numclass,
+            sequence_length, label_squence_length, cum_labels_length, costs, grads, prob_num,
+            ignore_longer_outputs_than_inputs_, stream);
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
+    FreeMem(label_value_with_blank, log_alpha_b, log_beta_b);
+    return true;
+  }
+  bool Init(const CNodePtr &kernel_node) override {
+    InitResource();
+    auto probs_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
+    if (probs_shape.size() != 3) {
+      MS_LOG(EXCEPTION) << "probs dims: " << probs_shape.size() << " not support.";
+    }
+    probs_dims_[0] = probs_shape[0];
+    probs_dims_[1] = probs_shape[1];
+    probs_dims_[2] = probs_shape[2];
+    auto indice_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
+    auto labels_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 2);
+    if (labels_dims.size() != 1) {
+      MS_LOG(EXCEPTION) << "labels dims: " << labels_dims.size() << " not support.";
+    }
+    if (indice_dims.size() != 2) {
+      MS_LOG(EXCEPTION) << "labels indice dims: " << indice_dims.size() << " not support.";
+    }
+    label_size_ = sizeof(int);
+    for (auto i : labels_dims) {
+      label_size_ *= i;
+    }
+    label_indice_size_ = sizeof(int64_t);
+    for (auto i : indice_dims) {
+      label_indice_size_ *= i;
+    }
+    auto squence_length_dims = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 3);
+    squence_lengths_size_ = squence_length_dims[0] * sizeof(int);
+    preprocess_collapse_repeated_ = GetAttr<bool>(kernel_node, "preprocess_collapse_repeated");
+    ctc_merge_repeated_ = GetAttr<bool>(kernel_node, "ctc_merge_repeated");
+    ignore_longer_outputs_than_inputs_ = GetAttr<bool>(kernel_node, "ignore_longer_outputs_than_inputs");
+    InitSizeLists();
+    return true;
+  }
+
+ protected:
+  void InitSizeLists() override {
+    input_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(T));
+    input_size_list_.push_back(label_indice_size_);
+    input_size_list_.push_back(label_size_);
+    input_size_list_.push_back(squence_lengths_size_);
+    workspace_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(T));
+    workspace_size_list_.push_back(squence_lengths_size_);
+    workspace_size_list_.push_back(squence_lengths_size_);
+    workspace_size_list_.push_back(label_size_);
+    workspace_size_list_.push_back(label_size_);
+    workspace_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(T));
+    workspace_size_list_.push_back(squence_lengths_size_);
+    workspace_size_list_.push_back(sizeof(int));
+    output_size_list_.push_back(probs_dims_[1] * sizeof(T));
+    output_size_list_.push_back(probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(T));
+  }
+  void MemsetForWS(int *label_value_pcr, int *cum_labels_length, int *label_squence_length, T *costs, T *grads,
+                   cudaStream_t stream) {
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMemsetAsync(label_value_pcr, static_cast<int>(0), label_size_, stream),
+                               "cudaMemSet failed in CtcLossGpuKernel::Launch.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMemsetAsync(cum_labels_length, static_cast<int>(0), squence_lengths_size_, stream),
+                               "cudaMemSet failed in CtcLossGpuKernel::Launch.");
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMemsetAsync(label_squence_length, static_cast<int>(0), squence_lengths_size_, stream),
+      "cudaMemSet failed in CtcLossGpuKernel::Launch.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMemsetAsync(costs, static_cast<T>(0), probs_dims_[1] * sizeof(T), stream),
+                               "cudaMemSet failed in CtcLossGpuKernel::Launch.");
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMemsetAsync(grads, static_cast<T>(0), probs_dims_[0] * probs_dims_[1] * probs_dims_[2] * sizeof(T), stream),
+      "cudaMemSet failed in CtcLossGpuKernel::Launch.");
+  }
+  void MemManageForCus(T **log_alpha_b, T **log_beta_b, int **label_value_with_blank, int *cum_labels_length,
+                       int log_prob_size, int batch, cudaStream_t stream) {
+    int total_labels_size_host = 0;
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMalloc(reinterpret_cast<void **>(log_alpha_b), sizeof(T) * log_prob_size),
+                               "cudaMalloc failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMalloc(reinterpret_cast<void **>(log_beta_b), sizeof(T) * log_prob_size),
+                               "cudaMalloc failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaMemcpyAsync(&total_labels_size_host, cum_labels_length + batch - 1, sizeof(int),
+                                               cudaMemcpyDeviceToHost, stream),
+                               "cudaMemcpyAsync failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamSynchronize(stream), "cudaStreamSynchronize failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(
+      cudaMalloc(reinterpret_cast<void **>(label_value_with_blank), sizeof(int) * (2 * total_labels_size_host + batch)),
+      "cudaMalloc failed.");
+  }
+
+  void FreeMem(int *label_value_with_blank, T *log_alpha_b, T *log_beta_b) {
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaFree(label_value_with_blank), "cudaFree failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaFree(log_alpha_b), "cudaFree failed.");
+    CHECK_CUDA_RET_WITH_EXCEPT(cudaFree(log_beta_b), "cudaFree failed.");
+  }
+
+  std::vector<size_t> input_size_list_;
+  std::vector<size_t> output_size_list_;
+  std::vector<size_t> workspace_size_list_;
+
+  size_t probs_dims_[3] = {0};
+  int label_indice_size_;
+  int label_size_;
+  int squence_lengths_size_;
+  bool preprocess_collapse_repeated_;
+  bool ctc_merge_repeated_;
+  bool ignore_longer_outputs_than_inputs_;
+  T kLogZero_ = -std::numeric_limits<T>::infinity();
+};  // namespace kernel
+}  // namespace kernel
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_CTCLOSS_GPU_KERNEL_H_

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/mirror_pad_gpu_kernel.h

@@ -66,7 +66,6 @@ class MirrorPadGpuFwdKernel : public GpuKernel {
     }
 
     string mode = GetValue<string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("mode"));
-
     if (mode == "REFLECT") {
       mode_ = 0;  // reflected mirroring
     } else {

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/mirror_pad_grad_gpu_kernel.h

@@ -66,7 +66,6 @@ class MirrorPadGpuBackKernel : public GpuKernel {
     }
 
     string mode = GetValue<string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("mode"));
-
     if (mode == "REFLECT") {
       mode_ = 0;  // reflected mirroring
     } else {

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/roi_align_gpu_kernel.cc

@@ -27,6 +27,5 @@ MS_REG_GPU_KERNEL_ONE(
   ROIAlign,
   KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),
   ROIAlignGpuFwdKernel, half)
-
 }  // namespace kernel
 }  // namespace mindspore

mindspore/ccsrc/backend/kernel_compiler/gpu/nn/roi_align_grad_gpu_kernel.cc

@@ -27,6 +27,5 @@ MS_REG_GPU_KERNEL_ONE(
   ROIAlignGrad,
   KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),
   ROIAlignGradGpuFwdKernel, half)
-
 }  // namespace kernel
 }  // namespace mindspore

mindspore/ccsrc/backend/kernel_compiler/kernel_fusion.cc

@@ -25,9 +25,29 @@
 #include "backend/kernel_compiler/tbe/tbe_convert_utils.h"
 #include "utils/ms_context.h"
 
-namespace mindspore {
-namespace kernel {
+namespace mindspore::kernel {
 using mindspore::kernel::tbe::TbeUtils;
+static size_t GenFusionJsonHash(const nlohmann::json &fusion_json) {
+  // get an copy
+  nlohmann::json fusion_json_copy = fusion_json;
+  auto &op_lists = fusion_json_copy["op_list"];
+  for (auto &op : op_lists) {
+    op.erase("name");
+    for (auto &output_desc : op["output_desc"]) {
+      output_desc.erase("name");
+    }
+    if (op["type"] != "Data") {
+      for (auto &input_desc : op["input_desc"]) {
+        input_desc.erase("name");
+      }
+      for (auto &list_arg : op["prebuild_output_attrs"]["list_args"]) {
+        list_arg.erase("name");
+      }
+    }
+  }
+  return std::hash<std::string>()(fusion_json_copy.dump());
+}
+
 std::map<int32_t, KernelModPtr> KernelFusion(const std::vector<FusionScopeInfo> &fusion_scopes) {
   MS_LOG(INFO) << "kernel fusion build start, scope size:" << fusion_scopes.size();
   std::map<int32_t, KernelModPtr> kernel_mod_ret;
@@ -41,8 +61,8 @@ std::map<int32_t, KernelModPtr> KernelFusion(const std::vector<FusionScopeInfo>
       continue;
     }
     // gen kernel_name & check cache
-    std::string json_str = fusion_op.dump();
-    size_t hash_id = std::hash<std::string>()(json_str);
+    size_t hash_id = GenFusionJsonHash(fusion_op);
+    MS_LOG(INFO) << "Fusion op hash id: " << hash_id;
     auto context_ptr = MsContext::GetInstance();
     MS_EXCEPTION_IF_NULL(context_ptr);
     auto device_id = context_ptr->get_param<uint32_t>(MS_CTX_DEVICE_ID);
@@ -102,5 +122,4 @@ std::map<int32_t, KernelModPtr> KernelFusion(const std::vector<FusionScopeInfo>
   MS_LOG(INFO) << "Build Fusion Kernel Failed Num: " << build_failed_num;
   return kernel_mod_ret;
 }
-}  // namespace kernel
-}  // namespace mindspore
+}  // namespace mindspore::kernel

mindspore/ccsrc/backend/kernel_compiler/tbe/tbe_adapter.cc

@@ -87,6 +87,7 @@ static std::map<string, string> tbe_func_adapter_map = {
   {"apply_centered_rms_prop", "apply_centered_rms_prop_d"},
   {"transpose", "transpose_d"},
   {"fill", "fill_d"},
+  {"i_fmr", "ifmr"},
   {"unsorted_segment_sum", "unsorted_segment_sum_d"},
   {"unsorted_segment_prod", "unsorted_segment_prod_d"},
   {"concat", "concat_d"},
@@ -139,7 +140,8 @@ static std::map<string, string> tbe_func_adapter_map = {
   {"inplace_update", "inplace_update_d"},
   {"matrix_diag", "matrix_diag_d"},
   {"matrix_diag_part", "matrix_diag_part_d"},
-  {"matrix_set_diag", "matrix_set_diag_d"}};
+  {"matrix_set_diag", "matrix_set_diag_d"},
+  {"l_stm_input_grad", "lstm_input_grad"}};
 
 void TbeAdapter::NormalizeFuncName(std::string *func_name) {
   if (func_name == nullptr) {

mindspore/ccsrc/backend/kernel_compiler/tbe/tbe_kernel_build.cc

@@ -25,6 +25,7 @@
 #include "backend/kernel_compiler/tbe/tbe_convert_utils.h"
 #include "backend/kernel_compiler/tbe/tbe_utils.h"
 #include "utils/ms_context.h"
+#include "runtime/dev.h"
 
 namespace mindspore {
 namespace kernel {
@@ -86,6 +87,8 @@ constexpr auto kJPyModulePath = "py_module_path";
 constexpr auto kJPreBuildOutsAttrs = "prebuild_outs_attrs";
 constexpr auto kJKwdArgs = "kwds_args";
 constexpr auto kJListArgs = "list_args";
+constexpr auto kJSocVersion = "socVersion";
+constexpr auto kSOC_VERSION = "SOC_VERSION";
 
 bool TbeKernelJsonCreator::GenTbeSingleKernelJson(const std::shared_ptr<mindspore::AnfNode> &anf_node,
                                                   nlohmann::json *kernel_json) {
@@ -122,6 +125,8 @@ bool TbeKernelJsonCreator::GenTbeSingleKernelJson(const std::shared_ptr<mindspor
   nlohmann::json attrs_json;
   (void)GenTbeAttrJson(anf_node, op_info_ptr, &attrs_json);
   op_info_json[kJAttrs] = attrs_json;
+  auto soc_version = TbeKernelJsonCreator::GetSocVersion();
+  op_info_json[kJSocVersion] = soc_version;
   std::string json_str = op_info_json.dump();
   size_t hash_id = std::hash<std::string>()(json_str);
   auto context_ptr = MsContext::GetInstance();
@@ -150,7 +155,13 @@ bool TbeKernelJsonCreator::GenInputDescJson(const std::shared_ptr<AnfNode> &anf_
   MS_EXCEPTION_IF_NULL(input_ptr);
   MS_EXCEPTION_IF_NULL(input_list);
   std::string op_name = AnfAlgo::GetCNodeName(anf_node);
-  if (input_ptr->name() == "input_indices" && op_name == kTopKOpName) {
+  if (op_name == kDynamicRNNOpName && input_ptr->name() == "seq_length") {
+    nlohmann::json input_desc_json;
+    auto in_name = input_ptr->name();
+    input_desc_json[kJName] = in_name + std::to_string(input_i);
+    input_desc_json[kJValid] = false;
+    input_list->emplace_back(input_desc_json);
+  } else if (input_ptr->name() == "input_indices" && op_name == kTopKOpName) {
     TbeAdapter::GenTopKV2IndicesTensorInfo(anf_node, real_input_index, input_list, creater_type_);
   } else {
     auto dtype = GetDeviceInputType(anf_node, real_input_index);
@@ -408,6 +419,30 @@ bool TbeKernelJsonCreator::GenTbeAttrJson(const std::shared_ptr<AnfNode> &anf_no
   return true;
 }
 
+string TbeKernelJsonCreator::GetSocVersion() {
+  // Get default soc version.
+  const int kSocVersionLen = 50;
+  char soc_version[kSocVersionLen] = {0};
+  auto ret = rtGetSocVersion(soc_version, kSocVersionLen);
+  if (ret != RT_ERROR_NONE) {
+    MS_LOG(EXCEPTION) << "GetSocVersion failed.";
+  }
+  MS_LOG(INFO) << "Default SocVersion is " << soc_version;
+  // Get soc version from env value.
+  const char *soc_version_env = getenv(kSOC_VERSION);
+  if (soc_version_env != nullptr) {
+    if (std::strcmp(soc_version, soc_version_env) != 0) {
+      MS_LOG(WARNING) << "SocVerison change to " << soc_version_env;
+      ret = rtSetSocVersion(soc_version_env);
+      if (ret != RT_ERROR_NONE) {
+        MS_LOG(EXCEPTION) << "SetSocVersion to " << soc_version_env << " failed, errorno: " << ret;
+      }
+      return soc_version_env;
+    }
+  }
+  return soc_version;
+}
+
 void TbeKernelJsonCreator::ParseAttrValue(const std::string &type, const mindspore::ValuePtr &value,
                                           nlohmann::json *attr_obj) {
   MS_EXCEPTION_IF_NULL(value);
@@ -624,6 +659,8 @@ bool TbeKernelBuild::GenFusionScopeJson(const std::vector<mindspore::AnfNodePtr>
   index = 0;
   data_list.insert(data_list.end(), compute_list.begin(), compute_list.end());
   (*fusion_json)[kFusionOpList] = data_list;
+  auto soc_version = TbeKernelJsonCreator::GetSocVersion();
+  (*fusion_json)[kJSocVersion] = soc_version;
   return true;
 }
 
@@ -853,6 +890,7 @@ bool TbeKernelBuild::GenFusionDataInputJson(const std::shared_ptr<mindspore::Anf
     (*data_str)[kJName] = name;
     nlohmann::json output_desc;
     output_desc[kJName] = name;
+    output_desc[kJDataType] = 0;
     output_desc[kJShape] = "NULL";
     output_desc_list.push_back(output_desc);
     (*index)++;
@@ -985,6 +1023,7 @@ bool TbeKernelBuild::GenFusionComputeInputJson(const mindspore::CNodePtr &cnode,
     for (size_t i = 0; i < optional_num; ++i) {
       nlohmann::json optional_input_desc;
       optional_input_desc[kJName] = std::string(kOptional) + std::to_string(*index);
+      optional_input_desc[kJShape] = "NULL";
       (*index)++;
       (*layer_iter)->emplace_back(nullptr);
       input_desc_list_tmp.emplace_back(optional_input_desc);

mindspore/ccsrc/backend/kernel_compiler/tbe/tbe_kernel_build.h

@@ -92,6 +92,7 @@ class TbeKernelJsonCreator {
   std::string json_name() { return json_name_; }
   bool GenTbeAttrJson(const std::shared_ptr<AnfNode> &anf_node, const std::shared_ptr<OpInfo> &op_info,
                       nlohmann::json *attrs_json);
+  static string GetSocVersion();
 
  private:
   bool GenTbeInputsJson(const std::shared_ptr<AnfNode> &anf_node, const std::shared_ptr<OpInfo> &op_info,

mindspore/ccsrc/backend/optimizer/ascend/ascend_backend_optimization.cc

@@ -19,6 +19,7 @@
 #include <memory>
 #include <string>
 #include "backend/optimizer/common/optimizer.h"
+#include "backend/optimizer/ascend/ir_fission/dynamic_rnn_grad_fission.h"
 #include "backend/optimizer/ascend/ir_fission/bn_split.h"
 #include "backend/optimizer/ascend/ir_fission/bn_grad_split.h"
 #include "backend/optimizer/ascend/ir_fission/batch_norm_grad_split.h"
@@ -64,6 +65,7 @@
 #include "backend/optimizer/ascend/format_type/rectify_do_mask_kernel_info.h"
 #include "backend/optimizer/ascend/format_type/chang_axis_of_reduce_kernel.h"
 #include "backend/optimizer/ascend/format_type/split_unsupported_transdata.h"
+#include "backend/optimizer/ascend/format_type/insert_reshape_for_extract_image_patches_op.h"
 #include "backend/optimizer/pass/getitem_tuple.h"
 #include "backend/optimizer/pass/optimize_dependence.h"
 #include "backend/optimizer/pass/erase_visit_attr.h"
@@ -106,6 +108,7 @@
 #include "backend/optimizer/ascend/ir_fission/concat_fission.h"
 #include "backend/optimizer/ascend/ir_fission/pack_fission.h"
 #include "backend/optimizer/ascend/enhancer/concat_outputs_for_all_gather.h"
+#include "backend/optimizer/ascend/enhancer/add_placeholder_for_dynamic_rnn.h"
 #include "utils/ms_context.h"
 #include "backend/optimizer/graph_kernel/composite_ops_fusion.h"
 #include "backend/optimizer/graph_kernel/basic_ops_fusion.h"
@@ -231,6 +234,7 @@ void AscendMixPrecision(const std::shared_ptr<session::KernelGraph> &kernel_grap
   auto optimizer = std::make_shared<GraphOptimizer>();
   auto mixed_precision_pm = std::make_shared<PassManager>("cast_pm");
   mixed_precision_pm->AddPass(std::make_shared<InsertCast>());
+  mixed_precision_pm->AddPass(std::make_shared<InsertReshapeForExtractImagePatchesOp>());
   mixed_precision_pm->AddPass(std::make_shared<GetitemTuple>());
   mixed_precision_pm->AddPass(std::make_shared<CommonSubexpressionElimination>());
   mixed_precision_pm->AddPass(std::make_shared<EliminateRedundantOp>());
@@ -276,6 +280,8 @@ void AscendBackendIRFusionOptimization(const std::shared_ptr<session::KernelGrap
   }
   ir_fusion_pm->AddPass(std::make_shared<LayerNormGradSplit>());
   ir_fusion_pm->AddPass(std::make_shared<InsertPadForNMSWithMask>());
+  ir_fusion_pm->AddPass(std::make_shared<InsertPlaceholderForDynamicRNN>());
+  ir_fusion_pm->AddPass(std::make_shared<DynamicRNNGradFission>());
   AddAscendIRFusionRulesPass(ir_fusion_pm.get());
   AddAscendIRFusionPass(ir_fusion_pm.get());
 

mindspore/ccsrc/backend/optimizer/ascend/enhancer/add_placeholder_for_dynamic_rnn.cc

@@ -0,0 +1,77 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/optimizer/ascend/enhancer/add_placeholder_for_dynamic_rnn.h"
+#include <vector>
+#include <memory>
+#include "backend/optimizer/common/helper.h"
+#include "backend/session/anf_runtime_algorithm.h"
+#include "utils/utils.h"
+#include "abstract/abstract_value.h"
+#include "base/core_ops.h"
+
+namespace mindspore {
+namespace opt {
+const BaseRef InsertPlaceholderForDynamicRNN::DefinePattern() const {
+  std::shared_ptr<Var> V = std::make_shared<CondVar>(UnVisited);
+  std::shared_ptr<Var> Xs = std::make_shared<SeqVar>();
+  return VectorRef({V, Xs});
+}
+
+const AnfNodePtr InsertPlaceholderForDynamicRNN::Process(const FuncGraphPtr &func_graph, const AnfNodePtr &node,
+                                                         const EquivPtr &) const {
+  MS_EXCEPTION_IF_NULL(func_graph);
+  MS_EXCEPTION_IF_NULL(node);
+  auto cnode = node->cast<CNodePtr>();
+  MS_EXCEPTION_IF_NULL(cnode);
+  auto op_name = AnfAlgo::GetCNodeName(cnode);
+  if (op_name != kDynamicRNNOpName) {
+    return nullptr;
+  }
+  AnfAlgo::SetNodeAttr(kAttrVisited, MakeValue(true), node);
+  auto kernel_graph = func_graph->cast<std::shared_ptr<session::KernelGraph>>();
+  MS_EXCEPTION_IF_NULL(kernel_graph);
+  size_t input_num = AnfAlgo::GetInputTensorNum(node);
+  if (input_num == 0) {
+    return nullptr;
+  }
+
+  std::vector<AnfNodePtr> new_inputs = {AnfAlgo::GetCNodePrimitiveNode(cnode)};
+  for (size_t in_idx = 0; in_idx < input_num; in_idx++) {
+    auto input_node = AnfAlgo::GetInputNode(cnode, in_idx);
+    if (in_idx == 3) {
+      auto value = std::make_shared<None>();
+      auto value_node = NewValueNode(value);
+      value_node->set_abstract(std::make_shared<abstract::AbstractNone>());
+      auto new_node = kernel_graph->NewValueNode(value_node);
+      kernel_graph->AddValueNodeToGraph(new_node);
+      new_inputs.push_back(new_node);
+    }
+    new_inputs.push_back(input_node);
+  }
+
+  CNodePtr new_node = nullptr;
+  if (kernel_graph == nullptr) {
+    new_node = std::make_shared<CNode>(*cnode);
+  } else {
+    new_node = kernel_graph->NewCNode(cnode);
+  }
+  MS_EXCEPTION_IF_NULL(new_node);
+  new_node->set_inputs(new_inputs);
+  return new_node;
+}
+}  // namespace opt
+}  // namespace mindspore

mindspore/ccsrc/backend/optimizer/ascend/enhancer/add_placeholder_for_dynamic_rnn.h

@@ -0,0 +1,37 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+#ifndef MINDSPORE_CCSRC_PRE_ACTIVATE_ASCEND_ENHANCER_ADD_PLACEHOLDER_FOR_DYNAMIC_RNN_H
+#define MINDSPORE_CCSRC_PRE_ACTIVATE_ASCEND_ENHANCER_ADD_PLACEHOLDER_FOR_DYNAMIC_RNN_H
+
+#include <memory>
+#include <vector>
+#include "backend/optimizer/common/optimizer.h"
+#include "backend/optimizer/ascend/ascend_helper.h"
+
+namespace mindspore {
+namespace opt {
+class InsertPlaceholderForDynamicRNN : public PatternProcessPass {
+ public:
+  explicit InsertPlaceholderForDynamicRNN(bool multigraph = true)
+      : PatternProcessPass("add_placeholder_for_dynamic_rnn", multigraph) {}
+  ~InsertPlaceholderForDynamicRNN() override = default;
+  const BaseRef DefinePattern() const override;
+  const AnfNodePtr Process(const FuncGraphPtr &, const AnfNodePtr &, const EquivPtr &) const override;
+};
+}  // namespace opt
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_PRE_ACTIVATE_ASCEND_ENHANCER_ADD_PLACEHOLDER_FOR_DYNAMIC_RNN_H

mindspore/ccsrc/backend/optimizer/ascend/format_type/insert_reshape_for_extract_image_patches_op.cc

@@ -0,0 +1,65 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/optimizer/ascend/format_type/insert_reshape_for_extract_image_patches_op.h"
+#include <memory>
+#include "backend/optimizer/ascend/ascend_helper.h"
+#include "backend/session/anf_runtime_algorithm.h"
+#include "utils/utils.h"
+#include "base/core_ops.h"
+
+namespace mindspore {
+namespace opt {
+const BaseRef InsertReshapeForExtractImagePatchesOp::DefinePattern() const {
+  VarPtr Xs = std::make_shared<SeqVar>();
+  return VectorRef({prim::kPrimExtractImagePatches, Xs});
+}
+
+const AnfNodePtr InsertReshapeForExtractImagePatchesOp::Process(const FuncGraphPtr &func_graph, const AnfNodePtr &node,
+                                                                const EquivPtr &equiv) const {
+  MS_EXCEPTION_IF_NULL(func_graph);
+  MS_EXCEPTION_IF_NULL(equiv);
+  auto extract = CheckAnfNodeIfCNodeAndInputSize(node, 2);
+  MS_EXCEPTION_IF_NULL(extract);
+  auto in_node = extract->input(1);
+  MS_EXCEPTION_IF_NULL(in_node);
+  auto extract_kernel_build_info = AnfAlgo::GetSelectKernelBuildInfo(extract);
+  auto in_node_kernel_build_info = AnfAlgo::GetSelectKernelBuildInfo(in_node);
+  MS_EXCEPTION_IF_NULL(extract_kernel_build_info);
+  MS_EXCEPTION_IF_NULL(in_node_kernel_build_info);
+  std::vector<AnfNodePtr> reshape_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimReshape->name())),
+                                            in_node};
+  auto reshape_builder = std::make_shared<kernel::KernelBuildInfo::KernelBuildInfoBuilder>();
+  reshape_builder->SetInputsFormat({kOpFormat_NC1HWC0});
+  reshape_builder->SetOutputsFormat({kOpFormat_NC1HWC0});
+  reshape_builder->SetInputsDeviceType({AnfAlgo::GetOutputDeviceDataType(in_node, 0)});
+  reshape_builder->SetOutputsDeviceType({AnfAlgo::GetOutputDeviceDataType(in_node, 0)});
+  reshape_builder->SetKernelType(in_node_kernel_build_info->kernel_type());
+  reshape_builder->SetFusionType(in_node_kernel_build_info->fusion_type());
+  reshape_builder->SetProcessor(in_node_kernel_build_info->processor());
+
+  auto reshape = func_graph->NewCNode(reshape_inputs);
+  reshape->set_scope(in_node->scope());
+  auto shape_tmp = AnfAlgo::GetOutputInferShape(in_node, 0);
+  AnfAlgo::SetOutputInferTypeAndShape({AnfAlgo::GetOutputDeviceDataType(in_node, 0)},
+                                      {{shape_tmp[0], shape_tmp[2], shape_tmp[3], shape_tmp[1]}}, reshape.get());
+  AnfAlgo::SetSelectKernelBuildInfo(reshape_builder->Build(), reshape.get());
+  AnfAlgo::SetNodeAttr("nop_op", MakeValue(true), reshape);
+  AnfAlgo::SetNodeInput(extract, reshape, 0);
+  return extract;
+}
+}  // namespace opt
+}  // namespace mindspore

mindspore/ccsrc/backend/optimizer/ascend/format_type/insert_reshape_for_extract_image_patches_op.h

@@ -0,0 +1,41 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+
+#ifndef MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_FORMAT_TYPE_INSERT_RESHAPE_FOR_EXTRACT_IMAGE_PATCHES_OP_H
+#define MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_FORMAT_TYPE_INSERT_RESHAPE_FOR_EXTRACT_IMAGE_PATCHES_OP_H
+#include <vector>
+#include <string>
+#include <utility>
+#include <memory>
+#include "ir/anf.h"
+#include "backend/optimizer/common/pattern_engine.h"
+#include "backend/optimizer/common/helper.h"
+#include "backend/optimizer/common/optimizer.h"
+
+namespace mindspore {
+namespace opt {
+class InsertReshapeForExtractImagePatchesOp : public PatternProcessPass {
+ public:
+  explicit InsertReshapeForExtractImagePatchesOp(bool multigraph = true)
+      : PatternProcessPass("insert_reshape_for_extract_image_patches_op", multigraph) {}
+  ~InsertReshapeForExtractImagePatchesOp() override = default;
+  const BaseRef DefinePattern() const override;
+  const AnfNodePtr Process(const FuncGraphPtr &, const AnfNodePtr &, const EquivPtr &) const override;
+};
+}  // namespace opt
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_FORMAT_TYPE_INSERT_RESHAPE_FOR_EXTRACT_IMAGE_PATCHES_OP_H

mindspore/ccsrc/backend/optimizer/ascend/ir_fission/dynamic_rnn_grad_fission.cc

@@ -0,0 +1,250 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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 "backend/optimizer/ascend/ir_fission/dynamic_rnn_grad_fission.h"
+#include <vector>
+#include <memory>
+#include <algorithm>
+#include "backend/session/anf_runtime_algorithm.h"
+#include "backend/optimizer/common/helper.h"
+
+namespace mindspore {
+namespace opt {
+constexpr size_t kDynamicRNNGradInputNum = 16;
+constexpr size_t kLSTMInputGradOutputNum = 4;
+const BaseRef DynamicRNNGradFission::DefinePattern() const {
+  VarPtr Xs = std::make_shared<SeqVar>();
+  return VectorRef({prim::kPrimDynamicRNNGrad, Xs});
+}
+
+AnfNodePtr CreateSplitVD(const FuncGraphPtr &graph, const AnfNodePtr &node) {
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node);
+  // SplitV
+  std::vector<AnfNodePtr> splitvd_input = {NewValueNode(std::make_shared<Primitive>(prim::kPrimSplitV->name())), node};
+  auto split_vd = graph->NewCNode(splitvd_input);
+  MS_EXCEPTION_IF_NULL(split_vd);
+  auto dtypes = {AnfAlgo::GetOutputInferDataType(node, 0), AnfAlgo::GetOutputInferDataType(node, 0)};
+  std::vector<size_t> shape = {AnfAlgo::GetOutputInferShape(node, 0)[0] - 1, AnfAlgo::GetOutputInferShape(node, 0)[1],
+                               AnfAlgo::GetOutputInferShape(node, 0)[2]};
+  auto shape2 = {IntToSize(1), AnfAlgo::GetOutputInferShape(node, 0)[1], AnfAlgo::GetOutputInferShape(node, 0)[2]};
+  std::vector<std::vector<size_t>> shapes = {shape, shape2};
+  AnfAlgo::SetOutputInferTypeAndShape(dtypes, shapes, split_vd.get());
+  AnfAlgo::SetNodeAttr("split_dim", MakeValue(0), split_vd);
+  AnfAlgo::SetNodeAttr("num_split", MakeValue(2), split_vd);
+  int tmp = SizeToInt(AnfAlgo::GetOutputInferShape(node, 0)[0]) - 1;
+  AnfAlgo::SetNodeAttr("size_splits", MakeValue(std::vector<int>{tmp, 1}), split_vd);
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), split_vd);
+  return split_vd;
+}
+
+AnfNodePtr CreateLSTMInputGrad(const FuncGraphPtr &graph, const AnfNodePtr &node) {
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node);
+  auto cnode = node->cast<CNodePtr>();
+  MS_EXCEPTION_IF_NULL(cnode);
+  const auto &dynamic_rnn_grad_inputs = cnode->inputs();
+  std::vector<AnfNodePtr> lstm_input_grad_inputs = {NewValueNode(std::make_shared<Primitive>(kLSTMInputGradOpName)),
+                                                    dynamic_rnn_grad_inputs[2],
+                                                    dynamic_rnn_grad_inputs[6],
+                                                    dynamic_rnn_grad_inputs[8],
+                                                    dynamic_rnn_grad_inputs[9],
+                                                    dynamic_rnn_grad_inputs[10],
+                                                    dynamic_rnn_grad_inputs[11],
+                                                    dynamic_rnn_grad_inputs[12],
+                                                    dynamic_rnn_grad_inputs[13],
+                                                    dynamic_rnn_grad_inputs[14],
+                                                    dynamic_rnn_grad_inputs[15],
+                                                    dynamic_rnn_grad_inputs[16]};
+  std::vector<AnfNodePtr> ori_outputs;
+  CreateMultipleOutputsOfAnfNode(graph, node, 5, &ori_outputs);
+  auto lstm_op = graph->NewCNode(lstm_input_grad_inputs);
+  MS_EXCEPTION_IF_NULL(lstm_op);
+  auto ori_type = AnfAlgo::GetOutputInferDataType(dynamic_rnn_grad_inputs[8], 0);
+  auto types = {AnfAlgo::GetOutputInferDataType(ori_outputs[2], 0), AnfAlgo::GetOutputInferDataType(ori_outputs[3], 0),
+                AnfAlgo::GetOutputInferDataType(ori_outputs[4], 0), ori_type};
+  std::vector<size_t> ori_shape = {AnfAlgo::GetOutputInferShape(dynamic_rnn_grad_inputs[8], 0)[0],
+                                   AnfAlgo::GetOutputInferShape(dynamic_rnn_grad_inputs[8], 0)[1],
+                                   4 * AnfAlgo::GetOutputInferShape(dynamic_rnn_grad_inputs[8], 0)[2]};
+  auto shapes = {AnfAlgo::GetOutputInferShape(ori_outputs[2], 0), AnfAlgo::GetOutputInferShape(ori_outputs[3], 0),
+                 AnfAlgo::GetOutputInferShape(ori_outputs[4], 0), ori_shape};
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, lstm_op.get());
+  return lstm_op;
+}
+
+AnfNodePtr CreateBatchMatMul(const FuncGraphPtr &graph, const AnfNodePtr &node1, const AnfNodePtr &node2) {
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node1);
+  MS_EXCEPTION_IF_NULL(node2);
+  // BatchMatMul
+  std::vector<AnfNodePtr> matmul_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimBatchMatMul->name())),
+                                           node2, node1};
+  auto batch_matmul = graph->NewCNode(matmul_inputs);
+  MS_EXCEPTION_IF_NULL(batch_matmul);
+  auto types = {AnfAlgo::GetOutputInferDataType(node1, 0)};
+  std::vector<size_t> shape = {AnfAlgo::GetOutputInferShape(node2, 0)[0], AnfAlgo::GetOutputInferShape(node2, 0)[2],
+                               AnfAlgo::GetOutputInferShape(node1, 0)[2]};
+  auto shapes = {shape};
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), batch_matmul);
+  AnfAlgo::SetNodeAttr("transpose_x1", MakeValue(true), batch_matmul);
+  AnfAlgo::SetNodeAttr("transpose_x2", MakeValue(false), batch_matmul);
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, batch_matmul.get());
+  return batch_matmul;
+}
+
+AnfNodePtr AddHConcatD(const FuncGraphPtr &graph, const AnfNodePtr &node1, const AnfNodePtr &node2) {
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node1);
+  MS_EXCEPTION_IF_NULL(node2);
+  std::vector<AnfNodePtr> ori_outputs;
+  CreateMultipleOutputsOfAnfNode(graph, node2, 2, &ori_outputs);
+  auto ori_shape = AnfAlgo::GetOutputInferShape(node1, 0);
+  std::vector<std::vector<size_t>> shape_tmp;
+  if (ori_shape.size() == 3) {
+    shape_tmp = {ori_shape};
+  } else {
+    shape_tmp = {{IntToSize(1), ori_shape[0], ori_shape[1]}};
+  }
+  auto ori_dtype = {AnfAlgo::GetOutputInferDataType(node1, 0)};
+  // reshape
+  std::vector<AnfNodePtr> reshape_input = {NewValueNode(std::make_shared<Primitive>(prim::kPrimReshape->name())),
+                                           node1};
+  auto reshape = graph->NewCNode(reshape_input);
+  AnfAlgo::SetNodeAttr(kAttrVisited, MakeValue(true), reshape);
+  AnfAlgo::SetOutputInferTypeAndShape(ori_dtype, shape_tmp, reshape.get());
+
+  // concatd --> concat
+  std::vector<AnfNodePtr> concat_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimConcat->name())),
+                                           reshape, ori_outputs[0]};
+  auto concat_op = graph->NewCNode(concat_inputs);
+  MS_EXCEPTION_IF_NULL(concat_op);
+  std::vector<size_t> input = {AnfAlgo::GetOutputInferShape(node2, 0)[0] + 1, AnfAlgo::GetOutputInferShape(node2, 0)[1],
+                               AnfAlgo::GetOutputInferShape(node2, 0)[2]};
+  auto types = {AnfAlgo::GetOutputInferDataType(node1, 0)};
+  auto shapes = {input};
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, concat_op.get());
+  AnfAlgo::SetNodeAttr(kAttrN, MakeValue(2), concat_op);
+  AnfAlgo::SetNodeAttr(kAttrDynInputSizes, MakeValue(std::vector<int>{2}), concat_op);
+  AnfAlgo::SetNodeAttr("axis", MakeValue(0), concat_op);
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), concat_op);
+  return concat_op;
+}
+
+AnfNodePtr AddConcatD(const FuncGraphPtr &graph, const AnfNodePtr &node1, const AnfNodePtr &node2) {
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node1);
+  MS_EXCEPTION_IF_NULL(node2);
+  // concatd --> concat
+  std::vector<AnfNodePtr> concat_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimConcat->name())), node1,
+                                           node2};
+  auto concat_op = graph->NewCNode(concat_inputs);
+  MS_EXCEPTION_IF_NULL(concat_op);
+  std::vector<size_t> input = {AnfAlgo::GetOutputInferShape(node1, 0)[0], AnfAlgo::GetOutputInferShape(node1, 0)[1],
+                               AnfAlgo::GetOutputInferShape(node1, 0)[2] + AnfAlgo::GetOutputInferShape(node2, 0)[2]};
+  auto types = {AnfAlgo::GetOutputInferDataType(node1, 0)};
+  auto shapes = {input};
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, concat_op.get());
+  AnfAlgo::SetNodeAttr(kAttrN, MakeValue(2), concat_op);
+  AnfAlgo::SetNodeAttr(kAttrDynInputSizes, MakeValue(std::vector<int>{2}), concat_op);
+  AnfAlgo::SetNodeAttr("axis", MakeValue(2), concat_op);
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), concat_op);
+  return concat_op;
+}
+
+AnfNodePtr AddDwReduceSum(const FuncGraphPtr &graph, const AnfNodePtr &node1, const AnfNodePtr &node2) {
+  // node1 : dynamic output
+  // node2 : matmul
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node1);
+  MS_EXCEPTION_IF_NULL(node2);
+  std::vector<AnfNodePtr> ori_outputs;
+  CreateMultipleOutputsOfAnfNode(graph, node1, 5, &ori_outputs);
+  // ReduceSumd
+  std::vector<AnfNodePtr> reducesum_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimReduceSum->name())),
+                                              node2};
+  auto reduce_sumd = graph->NewCNode(reducesum_inputs);
+  MS_EXCEPTION_IF_NULL(reduce_sumd);
+  auto types = {AnfAlgo::GetOutputInferDataType(ori_outputs[0], 0)};
+  auto shapes = {AnfAlgo::GetOutputInferShape(ori_outputs[0], 0)};
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, reduce_sumd.get());
+  AnfAlgo::SetNodeAttr(kAttrAxis, MakeValue(std::vector<int>{0}), reduce_sumd);
+  AnfAlgo::SetNodeAttr("keep_dims", MakeValue(false), reduce_sumd);
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), reduce_sumd);
+  return reduce_sumd;
+}
+
+AnfNodePtr AddDbReduceSum(const FuncGraphPtr &graph, const AnfNodePtr &node1, const AnfNodePtr &node2) {
+  // node1 lstm output
+  // node2 // dynamic output
+  MS_EXCEPTION_IF_NULL(graph);
+  MS_EXCEPTION_IF_NULL(node1);
+  MS_EXCEPTION_IF_NULL(node2);
+  std::vector<AnfNodePtr> ori_outputs;
+  CreateMultipleOutputsOfAnfNode(graph, node2, 5, &ori_outputs);
+  // ReduceSumd --> ReduceSum
+  std::vector<AnfNodePtr> reducerum_inputs = {NewValueNode(std::make_shared<Primitive>(prim::kPrimReduceSum->name())),
+                                              node1};
+  auto reduce_sumd = graph->NewCNode(reducerum_inputs);
+  MS_EXCEPTION_IF_NULL(reduce_sumd);
+  auto types = {AnfAlgo::GetOutputInferDataType(ori_outputs[1], 0)};
+  auto shapes = {AnfAlgo::GetOutputInferShape(ori_outputs[1], 0)};
+  AnfAlgo::SetOutputInferTypeAndShape(types, shapes, reduce_sumd.get());
+  AnfAlgo::SetNodeAttr(kAttrAxis, MakeValue(std::vector<int>{0, 1}), reduce_sumd);
+  AnfAlgo::SetNodeAttr("keep_dims", MakeValue(false), reduce_sumd);
+  AnfAlgo::SetNodeAttr("is_backend_insert", MakeValue(true), reduce_sumd);
+  return reduce_sumd;
+}
+
+const AnfNodePtr DynamicRNNGradFission::Process(const FuncGraphPtr &func_graph, const AnfNodePtr &node,
+                                                const EquivPtr &) const {
+  MS_EXCEPTION_IF_NULL(func_graph);
+  MS_EXCEPTION_IF_NULL(node);
+  auto cnode = node->cast<CNodePtr>();
+  MS_EXCEPTION_IF_NULL(cnode);
+  if (cnode->size() < kDynamicRNNGradInputNum + 1) {
+    MS_LOG(INFO) << "The input num of DynamicRNNGrad less than" << kDynamicRNNGradInputNum
+                 << ". The node should not be changed";
+    return nullptr;
+  }
+  // input_list of dynamic_rnn_grad
+  const auto &ori_inputs = cnode->inputs();
+  // create split_vd
+  auto split_vd = CreateSplitVD(func_graph, ori_inputs[7]);
+  // create concat_1
+  auto h_concat = AddHConcatD(func_graph, ori_inputs[5], split_vd);
+  // create concat_2
+  auto concat = AddConcatD(func_graph, ori_inputs[1], h_concat);
+  // create lsym_input_grad
+  auto lstm_input_grad = CreateLSTMInputGrad(func_graph, cnode);
+  std::vector<AnfNodePtr> lstm_outputs;
+  CreateMultipleOutputsOfAnfNode(func_graph, lstm_input_grad, kLSTMInputGradOutputNum, &lstm_outputs);
+  // create matmul
+  auto batch_matmul = CreateBatchMatMul(func_graph, lstm_outputs[3], concat);
+  // create reduce_sum_1
+  auto dw_reduce_sum = AddDwReduceSum(func_graph, node, batch_matmul);
+  // create reduce_sum_2
+  auto db_reduce_sum = AddDbReduceSum(func_graph, lstm_outputs[3], node);
+  std::vector<AnfNodePtr> make_tuple_inputs = {NewValueNode(prim::kPrimMakeTuple),
+                                               dw_reduce_sum,
+                                               db_reduce_sum,
+                                               lstm_outputs[0],
+                                               lstm_outputs[1],
+                                               lstm_outputs[2]};
+  auto make_tuple = func_graph->NewCNode(make_tuple_inputs);
+  MS_EXCEPTION_IF_NULL(make_tuple);
+  return make_tuple;
+}
+}  // namespace opt
+}  // namespace mindspore

mindspore/ccsrc/backend/optimizer/ascend/ir_fission/dynamic_rnn_grad_fission.h

@@ -0,0 +1,33 @@
+/**
+ * Copyright 2020 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * 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.
+ */
+#ifndef MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_IR_FISSION_DYNAMIC_RNN_GRAD_FISSION_H_
+#define MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_IR_FISSION_DYNAMIC_RNN_GRAD_FISSION_H_
+
+#include "backend/optimizer/common/optimizer.h"
+
+namespace mindspore {
+namespace opt {
+class DynamicRNNGradFission : public PatternProcessPass {
+ public:
+  explicit DynamicRNNGradFission(bool multigraph = true) : PatternProcessPass("dynamic_rnn_grad_fission", multigraph) {}
+  ~DynamicRNNGradFission() override = default;
+  const BaseRef DefinePattern() const override;
+  const AnfNodePtr Process(const FuncGraphPtr &, const AnfNodePtr &, const EquivPtr &) const override;
+};
+}  // namespace opt
+}  // namespace mindspore
+
+#endif  // MINDSPORE_CCSRC_BACKEND_OPTIMIZER_ASCEND_IR_FISSION_DYNAMIC_RNN_GRAD_FISSION_H_

mindspore/ccsrc/backend/optimizer/ascend/ir_fission/transdata_split.cc

@@ -21,11 +21,10 @@
 
 namespace mindspore {
 namespace opt {
-const std::set<std::pair<string, string>> invalid_formats_pair = {{kOpFormat_C1HWNCoC0, kOpFormat_NCHW},
-                                                                  {kOpFormat_NCHW, kOpFormat_C1HWNCoC0},
-                                                                  {kOpFormat_C1HWNCoC0, kOpFormat_DEFAULT},
-                                                                  {kOpFormat_DEFAULT, kOpFormat_FRACTAL_ZN_LSTM},
-                                                                  {kOpFormat_DEFAULT, kOpFormat_C1HWNCoC0}};
+const std::set<std::pair<string, string>> invalid_formats_pair = {
+  {kOpFormat_C1HWNCoC0, kOpFormat_NCHW},          {kOpFormat_NCHW, kOpFormat_C1HWNCoC0},
+  {kOpFormat_C1HWNCoC0, kOpFormat_DEFAULT},       {kOpFormat_DEFAULT, kOpFormat_FRACTAL_ZN_LSTM},
+  {kOpFormat_FRACTAL_ZN_LSTM, kOpFormat_DEFAULT}, {kOpFormat_DEFAULT, kOpFormat_C1HWNCoC0}};
 
 bool TransDataSplit::Run(const FuncGraphPtr &func_graph) {
   MS_EXCEPTION_IF_NULL(func_graph);
@@ -83,6 +82,9 @@ bool TransDataSplit::DoSplit(const FuncGraphPtr &func_graph, const AnfNodePtr &n
     new_transpose_node = NewTransOpNode(func_graph, AnfAlgo::GetInputNode(node->cast<CNodePtr>(), 0), kernel_select_,
                                         false, prim::kPrimTranspose->name());
     AnfAlgo::SetNodeAttr(kAttrPerm, MakeValue(std::vector<int>{2, 3, 1, 0}), new_transpose_node);
+    if (output_format == kOpFormat_FRACTAL_ZN_LSTM) {
+      AnfAlgo::SetNodeAttr("nop_op", MakeValue(true), new_transpose_node);
+    }
     RefreshKernelBuildInfo(input_format, kOpFormat_HWCN, new_transpose_node);
 
     // trans hwcn to output_format

mindspore/ccsrc/backend/optimizer/common/helper.cc

@@ -404,7 +404,11 @@ bool IsNopNode(const AnfNodePtr &node) {
   }
   CNodePtr cnode = node->cast<CNodePtr>();
   MS_EXCEPTION_IF_NULL(cnode);
-  if (nop_nodes.find(AnfAlgo::GetCNodeName(cnode)) == nop_nodes.end()) {
+  bool is_nop_node = false;
+  if (AnfAlgo::HasNodeAttr("nop_op", cnode)) {
+    is_nop_node = AnfAlgo::GetNodeAttr<bool>(cnode, "nop_op");
+  }
+  if (nop_nodes.find(AnfAlgo::GetCNodeName(cnode)) == nop_nodes.end() && !is_nop_node) {
     return false;
   }
   return true;

mindspore/ccsrc/backend/session/anf_runtime_algorithm.cc

@@ -516,6 +516,10 @@ std::vector<size_t> AnfRuntimeAlgorithm::GetInputDeviceShape(const AnfNodePtr &n
   if (trans::IsNeedPadding(format, infer_shape.size())) {
     infer_shape = trans::PaddingShapeTo4d(infer_shape, GetInputReshapeType(node, input_idx));
   }
+  if (node->isa<CNode>() && GetCNodeName(node) == kExtractImagePatchesOpName) {
+    auto shape_tmp = {infer_shape[0], infer_shape[3], infer_shape[1], infer_shape[2]};
+    return trans::TransShapeToDevice(shape_tmp, format);
+  }
   return trans::TransShapeToDevice(infer_shape, format);
 }
 

mindspore/ccsrc/backend/session/ascend_session.cc

@@ -44,6 +44,7 @@
 #include "debug/data_dump/e2e_dump_util.h"
 #include "debug/anf_ir_dump.h"
 #include "debug/dump_proto.h"
+#include "toolchain/adx_datadump_server.h"
 
 namespace mindspore {
 namespace session {
@@ -179,6 +180,7 @@ GraphId AscendSession::CompileGraph(NotNull<FuncGraphPtr> func_graph) {
     debugger_->PreExecute(root_graph);
   }
 #endif
+  SetSummaryNodes(root_graph.get());
   // alloc mem
   MemoryAlloc(root_graph.get());
   // generate and load task into device
@@ -424,6 +426,16 @@ void AscendSession::SelectKernel(const KernelGraph &kernel_graph) const {
   MS_LOG(INFO) << "Finish!";
 }
 
+void DumpInit() {
+  auto &json_parser = DumpJsonParser::GetInstance();
+  json_parser.Parse();
+  if (json_parser.async_dump_enabled()) {
+    if (AdxDataDumpServerInit() != 0) {
+      MS_LOG(EXCEPTION) << "Adx data dump server init failed";
+    }
+  }
+}
+
 void AscendSession::InitRuntimeResource() {
   MS_LOG(INFO) << "Start!";
   auto runtime_instance = device::KernelRuntimeManager::Instance().GetKernelRuntime(kAscendDevice, device_id_);
@@ -431,7 +443,7 @@ void AscendSession::InitRuntimeResource() {
   if (!runtime_instance->Init()) {
     MS_LOG(EXCEPTION) << "Kernel runtime init error.";
   }
-  DumpJsonParser::GetInstance().Parse();
+  DumpInit();
   MS_LOG(INFO) << "Finish!";
 }
 

mindspore/ccsrc/backend/session/executor.cc

@@ -102,7 +102,7 @@ void Executor::CheckException() {
 }
 
 void Executor::WorkerJoin() {
-  if (worker_->joinable()) {
+  if (worker_->joinable() && worker_->get_id() != std::this_thread::get_id()) {
     {
       std::unique_lock<std::mutex> lock(task_mutex_);
       auto task = std::make_shared<ExitTask>();

mindspore/ccsrc/backend/session/executor.h

@@ -125,7 +125,7 @@ class CreateCommGroupTask : public Task {
   void Run() override;
   std::string group_name_;
   std::vector<uint32_t> ranks_;
-  bool result_;
+  bool result_{false};
 };
 
 class DestroyCommGroupTask : public Task {
@@ -134,7 +134,7 @@ class DestroyCommGroupTask : public Task {
   ~DestroyCommGroupTask() override = default;
   void Run() override;
   std::string group_name_;
-  bool result_;
+  bool result_{false};
 };
 
 class ExitTask : public Task {

mindspore/ccsrc/debug/data_dump/dump_json_parser.cc

@@ -169,7 +169,7 @@ void DumpJsonParser::ParseAsyncDumpSetting(const nlohmann::json &content) {
 }
 
 void DumpJsonParser::ParseE2eDumpSetting(const nlohmann::json &content) {
-  auto e2e_dump_setting = CheckJsonKeyExist(content, kE2eDumpSettings);
+  auto e2e_dump_setting = content.find(kE2eDumpSettings);
   if (e2e_dump_setting == content.end()) {
     MS_LOG(INFO) << "No e2e_dump_settings";
     return;

mindspore/ccsrc/debug/data_dump/e2e_dump_util.cc

@@ -64,8 +64,11 @@ void E2eDumpUtil::DumpGPUMemToFile(const std::string &file_path, const std::stri
                                    const ShapeVector &int_shapes, const TypeId &type, size_t slot, Debugger *debugger) {
 #ifdef ENABLE_DEBUGGER
   auto format = kOpFormat_DEFAULT;
+  MS_EXCEPTION_IF_NULL(debugger);
   DebugServices *debug_services = debugger->debug_services();
+  MS_EXCEPTION_IF_NULL(debug_services);
   TensorLoader *tensor_loader = debug_services->tensor_loader();
+  MS_EXCEPTION_IF_NULL(tensor_loader);
   auto ret = tensor_loader->DumpTensorToFile(original_kernel_name, trans_flag, file_path, format, int_shapes, type,
                                              addr->type_id(), addr->format(), slot);
   if (!ret) {

mindspore/ccsrc/debug/debugger/debugger.cc

@@ -309,7 +309,7 @@ void Debugger::PostDebugOp() {
   }
 }
 
-std::map<std::pair<uint32_t, uint32_t>, std::string> &Debugger::GetStreamTaskToOpnameMap() {
+std::map<std::pair<uint32_t, uint32_t>, std::string> Debugger::GetStreamTaskToOpnameMap() const {
   return stream_task_to_opname_;
 }
 
@@ -412,6 +412,7 @@ void Debugger::CommandLoop() {
         MS_LOG(ERROR) << "Failed to connect to MindInsight debugger server. Please check the config "
                          "of debugger host and port.";
         Exit();
+        run = true;
       }
       MS_LOG(ERROR) << "Number of consecutive WaitForCommand fail:" << num_wait_fail << "; Retry after "
                     << num_wait_fail << "s";
@@ -434,6 +435,8 @@ void Debugger::CommandLoop() {
       case DebuggerCommand::kExitCMD:
         MS_LOG(INFO) << "ExitCMD";
         Exit();
+        // Used for debugger termination
+        run = true;
         break;
       case DebuggerCommand::kRunCMD:
         MS_LOG(INFO) << "RunCMD";
@@ -583,9 +586,9 @@ std::list<TensorProto> Debugger::LoadTensors(const ProtoVector<TensorProto> &ten
 }
 
 void Debugger::Exit() {
-  // clear resource before exit
-  pipeline::ClearResAtexit();
-  std::exit(EXIT_FAILURE);
+  // Notify main thread to terminate after clearing resources
+  pipeline::ExecutorPy::DebugTerminate(true);
+  debugger_enabled_ = false;
 }
 
 std::list<WatchpointHit> Debugger::CheckWatchpoints(const std::string &watchnode) {
@@ -754,7 +757,8 @@ uint64_t BytestoInt64(const std::vector<char> &buffer) {
   uint64_t ret;
 
   ret = ((uint64_t)buffer[7] << 56) | ((uint64_t)buffer[6] << 48) | ((uint64_t)buffer[5] << 40) |
-        ((uint64_t)buffer[4] << 32) | (buffer[3] << 24) | (buffer[2] << 16) | (buffer[1] << 8) | buffer[0];
+        ((uint64_t)buffer[4] << 32) | ((uint64_t)buffer[3] << 24) | ((uint64_t)buffer[2] << 16) |
+        ((uint64_t)buffer[1] << 8) | ((uint64_t)buffer[0]);
 
   return ret;
 }

mindspore/ccsrc/debug/debugger/debugger.h

@@ -94,7 +94,7 @@ class Debugger : public std::enable_shared_from_this<Debugger> {
 
   int32_t step_num() const;
 
-  std::map<std::pair<uint32_t, uint32_t>, std::string> &GetStreamTaskToOpnameMap();
+  std::map<std::pair<uint32_t, uint32_t>, std::string> GetStreamTaskToOpnameMap() const;
 
   // check if any feature that uses the debugger backend is enabled
   bool DebuggerBackendEnabled();

mindspore/ccsrc/debug/debugger/grpc_client.cc

@@ -79,9 +79,9 @@ std::vector<std::string> ChunkString(std::string str, int graph_size) {
     }
     std::string buffer;
     buffer.resize(chunk_size);
-    errno_t err = memcpy_s(reinterpret_cast<char *>(buffer.data()), chunk_size, str.data() + size_iter, chunk_size);
+    auto err = memcpy_s(reinterpret_cast<char *>(buffer.data()), chunk_size, str.data() + size_iter, chunk_size);
     if (err != 0) {
-      MS_LOG(ERROR) << "memcpy_s failed. err code is: " << err;
+      MS_LOG(EXCEPTION) << "memcpy_s failed. errorno is: " << err;
     }
     buf.push_back(buffer);
     size_iter += CHUNK_SIZE;

mindspore/ccsrc/frontend/optimizer/irpass.cc

@@ -25,7 +25,7 @@
 #include "frontend/optimizer/irpass/inline.h"
 #include "frontend/optimizer/irpass/incorporate_call.h"
 #include "frontend/optimizer/irpass/incorporate_getitem.h"
-#include "frontend/optimizer/irpass/item_tuple_eliminate.h"
+#include "frontend/optimizer/irpass/item_tuple_or_list_eliminate.h"
 #include "frontend/optimizer/irpass/mark_interface_fusion.h"
 #include "frontend/optimizer/irpass/merge_addn.h"
 #include "frontend/optimizer/irpass/minmax_grad.h"
@@ -66,8 +66,9 @@ OptimizeIRPassLib::OptimizeIRPassLib() {
     MakeSubstitution(std::make_shared<AdjustAllReduceMulAdd>(), "adjust_all_reduce_mul_add", prim::kPrimAddN);
 
   // ops eliminate
-  item_tuple_eliminate_ = MakeSubstitution(std::make_shared<ItemTupleEliminater>(), "item_tuple_eliminate",
-                                           {prim::kPrimTupleGetItem, prim::kPrimTupleSetItem, prim::kPrimListGetItem});
+  item_tuple_or_list_eliminate_ = MakeSubstitution(
+    std::make_shared<ItemTupleOrListEliminater>(), "item_tuple_or_list_eliminate",
+    {prim::kPrimTupleGetItem, prim::kPrimTupleSetItem, prim::kPrimListGetItem, prim::kPrimListSetItem});
   tile_eliminate_ = MakeSubstitution(std::make_shared<TileMultiplyByOne>(), "tile_eliminate", prim::kPrimTile);
   cast_eliminate_ = MakeSubstitution(std::make_shared<CastEliminater>(), "cast_eliminate", prim::kPrimCast);
   reshape_eliminate_ = MakeSubstitution(std::make_shared<ReshapeEliminater>(), "reshape_eliminate", prim::kPrimReshape);

mindspore/ccsrc/frontend/optimizer/irpass.h

@@ -39,7 +39,7 @@ class OptimizeIRPassLib {
   SubstitutionPtr adjust_all_reduce_mul_add_;
 
   //  ops eliminate
-  SubstitutionPtr item_tuple_eliminate_;
+  SubstitutionPtr item_tuple_or_list_eliminate_;
   SubstitutionPtr tile_eliminate_;
   SubstitutionPtr cast_eliminate_;
   SubstitutionPtr reshape_eliminate_;

mindspore/ccsrc/frontend/optimizer/irpass/item_tuple_or_list_eliminate.h

@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#ifndef MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_ELIMINATE_H_
-#define MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_ELIMINATE_H_
+#ifndef MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_OR_LIST_ELIMINATE_H_
+#define MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_OR_LIST_ELIMINATE_H_
 
 #include <algorithm>
 #include <memory>
@@ -33,6 +33,7 @@ namespace irpass {
 // (a, b, c, ...)[0] => a
 // (a, b, c, ...)[1] => b
 // {prim::kPrimTupleGetItem, {prim::kPrimMakeTuple, Xs}, C}
+// {prim::kPrimListGetItem, {prim::kPrimMakeList, Xs}, C}
 class GetitemEliminater : public AnfVisitor {
  public:
   AnfNodePtr operator()(const OptimizerPtr &, const AnfNodePtr &node) override {
@@ -80,6 +81,7 @@ class GetitemEliminater : public AnfVisitor {
 // (a, b, c, ...)[0] => a
 // (a, b, c, ...)[1] => b
 // {prim::kPrimTupleGetItem, C1, C}
+// {prim::kPrimListGetItem, C1, C}
 class GetitemConstEliminater : public AnfVisitor {
  public:
   AnfNodePtr operator()(const OptimizerPtr &, const AnfNodePtr &node) override {
@@ -124,11 +126,13 @@ class GetitemConstEliminater : public AnfVisitor {
 // setitem((a, b, c, ...), 0, z) => (z, b, c, ...)
 // setitem((a, b, c, ...), 1, z) => (a, z, c, ...)
 // {prim::kPrimTupleSetItem, {prim::kPrimMakeTuple, Xs}, C, Z}
+// {prim::kPrimListSetItem, {prim::kPrimMakeList, Xs}, C, Z}
 class SetitemEliminater : public AnfVisitor {
  public:
   AnfNodePtr operator()(const OptimizerPtr &, const AnfNodePtr &node) override {
     Reset();
     AnfVisitor::Match(prim::kPrimTupleSetItem, {IsCNode, IsVNode, IsNode})(node);
+    AnfVisitor::Match(prim::kPrimListSetItem, {IsCNode, IsVNode, IsNode})(node);
 
     auto fg = node->func_graph();
     if (fg != nullptr && z_ != nullptr) {
@@ -178,11 +182,13 @@ class SetitemEliminater : public AnfVisitor {
 };
 
 // {prim::kPrimTupleGetItem, {prim::kPrimTupleSetItem, Y, C1, X}, C2}
+// {prim::kPrimListGetItem, {prim::kPrimListSetItem, Y, C1, X}, C2}
 class GetSetitemEliminater : public AnfVisitor {
  public:
   AnfNodePtr operator()(const OptimizerPtr &, const AnfNodePtr &node) override {
     Reset();
     AnfVisitor::Match(prim::kPrimTupleGetItem, {IsCNode, IsVNode})(node);
+    AnfVisitor::Match(prim::kPrimListGetItem, {IsCNode, IsVNode})(node);
 
     auto fg = node->func_graph();
     if (fg != nullptr && key1_ >= 0 && key2_ >= 0) {
@@ -195,7 +201,7 @@ class GetSetitemEliminater : public AnfVisitor {
   }
 
   void Visit(const CNodePtr &cnode) override {
-    if (IsPrimitiveCNode(cnode, prim::kPrimTupleSetItem)) {
+    if (IsPrimitiveCNode(cnode, prim::kPrimTupleSetItem) || IsPrimitiveCNode(cnode, prim::kPrimListSetItem)) {
       if (cnode->size() < 4) {
         return;
       }
@@ -239,6 +245,8 @@ class GetSetitemEliminater : public AnfVisitor {
 
 // {prim::kPrimTupleGetItem, {prim::kPrimDepend, X, Y}, C} ->
 // {prim::kPrimDepend, {prim::kPrimTupleGetItem, X, C}, Y}
+// {prim::kPrimListGetItem, {prim::kPrimDepend, X, Y}, C} ->
+// {prim::kPrimDepend, {prim::kPrimListGetItem, X, C}, Y}
 class GetitemDependReorder : public AnfVisitor {
  public:
   AnfNodePtr operator()(const OptimizerPtr &, const AnfNodePtr &node) override {
@@ -274,9 +282,9 @@ class GetitemDependReorder : public AnfVisitor {
   AnfNodePtr x_{nullptr}, y_{nullptr}, c_{nullptr};
 };
 
-class ItemTupleEliminater : public OptimizerCaller {
+class ItemTupleOrListEliminater : public OptimizerCaller {
  public:
-  ItemTupleEliminater()
+  ItemTupleOrListEliminater()
       : get_item_eliminater_(std::make_shared<GetitemEliminater>()),
         get_item_const_eliminater_(std::make_shared<GetitemConstEliminater>()),
         set_item_eliminater_(std::make_shared<SetitemEliminater>()),
@@ -288,7 +296,7 @@ class ItemTupleEliminater : public OptimizerCaller {
     eliminaters_.emplace_back(get_set_item_eliminater_);
     eliminaters_.emplace_back(get_item_depend_reorder_);
   }
-  ~ItemTupleEliminater() = default;
+  ~ItemTupleOrListEliminater() = default;
 
   AnfNodePtr operator()(const OptimizerPtr &optimizer, const AnfNodePtr &node) override {
     AnfNodePtr new_node;
@@ -309,4 +317,4 @@ class ItemTupleEliminater : public OptimizerCaller {
 }  // namespace irpass
 }  // namespace opt
 }  // namespace mindspore
-#endif  // MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_ELIMINATE_H_
+#endif  // MINDSPORE_CCSRC_FRONTEND_OPTIMIZER_IRPASS_ITEM_TUPLE_OR_LIST_ELIMINATE_H_

mindspore/ccsrc/frontend/optimizer/irpass/special_op_eliminate.h

@@ -255,7 +255,7 @@ class PynativeEliminater : public OptimizerCaller {
     MS_LOG(DEBUG) << "Start FillZero";
     ValuePtr out = nullptr;
     if (value->isa<Int32Imm>()) {
-      return value;
+      return MakeValue(value->cast<Int32ImmPtr>()->value());
     }
 
     if (value->isa<tensor::Tensor>()) {
@@ -298,9 +298,10 @@ class PynativeEliminater : public OptimizerCaller {
       if (rep != nullptr) {
         if (rep->isa<ValueNode>()) {
           auto value_node = rep->cast<ValueNodePtr>();
-          value_node->set_value(FillZero(value_node->value()));
+          auto new_value_node = NewValueNode(FillZero(value_node->value()));
+          new_value_node->set_has_new_value(value_node->has_new_value());
           MS_LOG(DEBUG) << "Zeros_like replace ok " << rep->DebugString(4);
-          return rep;
+          return new_value_node;
         }
       }
     }
@@ -315,9 +316,10 @@ class PynativeEliminater : public OptimizerCaller {
       if (rep != nullptr) {
         if (rep->isa<ValueNode>()) {
           auto value_node = rep->cast<ValueNodePtr>();
-          value_node->set_value(FillZero(value_node->value()));
+          auto new_value_node = NewValueNode(FillZero(value_node->value()));
+          new_value_node->set_has_new_value(value_node->has_new_value());
           MS_LOG(DEBUG) << "Zeros_like replace ok 2 " << rep->DebugString(4);
-          return rep;
+          return new_value_node;
         }
       }
     }

mindspore/ccsrc/minddata/dataset/CMakeLists.txt

@@ -139,20 +139,21 @@ endif()
 target_link_libraries(_c_dataengine PRIVATE mindspore mindspore_gvar)
 if (${CMAKE_SYSTEM_NAME} MATCHES "Windows")
     if (ENABLE_PYTHON)
-        target_link_libraries(_c_dataengine PRIVATE mindspore::pybind11_module ${PYTHON_LIBRARIES} mindspore::protobuf ${SECUREC_LIBRARY})
+        target_link_libraries(_c_dataengine PRIVATE mindspore::pybind11_module ${PYTHON_LIBRARIES} ${SECUREC_LIBRARY})
     else()
-        target_link_libraries(_c_dataengine PRIVATE mindspore::protobuf ${SECUREC_LIBRARY})
+        target_link_libraries(_c_dataengine PRIVATE ${SECUREC_LIBRARY})
     endif()
 else()
     set(ICU_LIB mindspore::icuuc mindspore::icudata mindspore::icui18n)
     if (ENABLE_PYTHON)
-        target_link_libraries(_c_dataengine PRIVATE mindspore::pybind11_module -ldl mindspore::protobuf ${SECUREC_LIBRARY})
+        target_link_libraries(_c_dataengine PRIVATE mindspore::pybind11_module -ldl ${SECUREC_LIBRARY})
     else()
-        target_link_libraries(_c_dataengine PRIVATE -ldl mindspore::protobuf ${SECUREC_LIBRARY})
+        target_link_libraries(_c_dataengine PRIVATE -ldl ${SECUREC_LIBRARY})
     endif()
+    target_link_libraries(_c_dataengine PUBLIC mindspore::sentencepiece)
 endif()
 target_link_libraries(_c_dataengine PUBLIC mindspore::jpeg_turbo mindspore::turbojpeg mindspore::opencv_core mindspore::opencv_imgcodecs
-        mindspore::opencv_imgproc mindspore::tinyxml2 mindspore::sentencepiece mindspore::sentencepiece_train ${ICU_LIB})
+                       mindspore::opencv_imgproc mindspore::tinyxml2 mindspore::sentencepiece_train ${ICU_LIB})
 if (ENABLE_GPUQUE)
     target_link_libraries(_c_dataengine PRIVATE gpu_queue
                                      ${CUDNN_LIBRARY_PATH}
@@ -171,7 +172,7 @@ if (${CMAKE_SYSTEM_NAME} MATCHES "Windows")
 else()
     target_link_libraries(_c_dataengine PRIVATE _c_mindrecord)
     if (ENABLE_CPU AND (ENABLE_D OR ENABLE_GPU))
-        target_link_libraries(_c_dataengine PRIVATE mindspore::pslite mindspore::protobuf ${zeromq_DIRPATH}/zmq_install/lib/libzmq.a)
+        target_link_libraries(_c_dataengine PRIVATE mindspore::pslite ${zeromq_DIRPATH}/zmq_install/lib/libzmq.a)
         if (${ENABLE_IBVERBS} STREQUAL "ON")
             target_link_libraries(_c_dataengine PRIVATE ibverbs rdmacm)
         endif()

mindspore/ccsrc/minddata/dataset/api/python/bindings.cc

@@ -76,6 +76,12 @@ PYBIND_REGISTER(
              THROW_IF_ERROR(de.GetOutputTypes(&out));
              return out;
            })
+      .def("GetDataInfo",
+           [](DEPipeline &de) {
+             py::list types, shapes;
+             THROW_IF_ERROR(de.GetDataInfo(&types, &shapes));
+             return py::make_tuple(types, shapes);
+           })
       .def("GetDatasetSize", &DEPipeline::GetDatasetSize)
       .def("GetBatchSize", &DEPipeline::GetBatchSize)
       .def("GetNumClasses", &DEPipeline::GetNumClasses)

mindspore/ccsrc/minddata/dataset/api/python/bindings/dataset/core/bindings.cc

@@ -69,7 +69,7 @@ PYBIND_REGISTER(Tensor, 0, ([](const py::module *m) {
                       auto &tensor = py::cast<Tensor &>(t);
                       if (tensor.type() == DataType::DE_STRING) {
                         py::array res;
-                        tensor.GetDataAsNumpyStrings(&res);
+                        THROW_IF_ERROR(tensor.GetDataAsNumpyStrings(&res));
                         return res;
                       }
                       py::buffer_info info;

mindspore/ccsrc/minddata/dataset/api/python/de_pipeline.cc

@@ -241,6 +241,30 @@ Status DEPipeline::GetNextAsList(py::list *output) {
   return Status::OK();
 }
 
+Status DEPipeline::GetDataInfo(py::list *types, py::list *shapes) {
+  Status s;
+  DATA_INFO data_info;
+  // tree_.root() must be DeviceQueueOp
+  DeviceQueueOp *op = dynamic_cast<DeviceQueueOp *>(tree_->root().get());
+  if (op == nullptr) {
+    return Status(StatusCode::kUnexpectedError, __LINE__, __FILE__, "GetDataInfo only supported by DeviceQueueOp");
+  }
+  {
+    py::gil_scoped_release gil_release;
+    s = op->GetDataInfo(&data_info);
+  }
+  RETURN_IF_NOT_OK(s);
+  for (auto el : data_info) {
+    types->append(el.first.AsNumpyType());
+    py::list shape;
+    for (auto dim : el.second.AsVector()) {
+      shape.append(dim);
+    }
+    shapes->append(shape);
+  }
+  return Status::OK();
+}
+
 Status DEPipeline::GetOutputShapes(py::list *output) {
   std::vector<TensorShape> shapes;
   Status s;
@@ -1052,6 +1076,8 @@ Status DEPipeline::ParseDeviceQueueOp(const py::dict &args, std::shared_ptr<Data
         (void)builder->SetDeviceId(ToInt(value));
       } else if (key == "send_epoch_end") {
         (void)builder->SetSendEpochEnd(ToBool(value));
+      } else if (key == "create_data_info_queue") {
+        (void)builder->SetCreateDataInfoQueue(ToBool(value));
       }
     }
   }

mindspore/ccsrc/minddata/dataset/api/python/de_pipeline.h

@@ -111,6 +111,8 @@ class DEPipeline {
 
   Status GetOutputTypes(py::list *output);
 
+  Status GetDataInfo(py::list *types, py::list *shapes);
+
   Status SaveDataset(const std::vector<std::string> &file_names, const std::string &file_type);
 
   int GetDatasetSize() const;

mindspore/ccsrc/minddata/dataset/engine/datasetops/batch_op.cc

@@ -26,6 +26,7 @@
 #include "minddata/dataset/engine/db_connector.h"
 #include "minddata/dataset/engine/opt/pass.h"
 #include "minddata/dataset/kernels/data/data_utils.h"
+#include "minddata/dataset/util/status.h"
 
 namespace mindspore {
 namespace dataset {
@@ -129,6 +130,15 @@ Status BatchOp::operator()() {
       worker_queues_[cnt++ % num_workers_]->EmplaceBack(std::make_pair(nullptr, CBatchInfo(batchCtrl::kEOE))));
     RETURN_IF_NOT_OK(GetBatchSize(&cur_batch_size, CBatchInfo(epoch_num, batch_num, cnt - epoch_num)));
     RETURN_IF_NOT_OK(child_iterator_->FetchNextTensorRow(&new_row));
+
+#if !defined(_WIN32) && !defined(_WIN64)
+    if (batch_num % 10 == 0 && (num_workers_ > 1 || batch_map_func_) && GetMemoryUsage() > MAX_MEMORY_USAGE_THRESHOLD) {
+      MS_LOG(WARNING) << "Memory consumption is more than " << MAX_MEMORY_USAGE_THRESHOLD * 100 << "%, "
+                      << "which may cause oom error. Please reduce num_parallel_workers size / "
+                      << "optimize per_batch_map function / other python data preprocess function to "
+                      << "reduce memory usage.";
+    }
+#endif
   }  // end of eof_handled() == false
   RETURN_IF_NOT_OK(
     worker_queues_[cnt++ % num_workers_]->EmplaceBack(std::make_pair(nullptr, CBatchInfo(batchCtrl::kEOF))));

mindspore/ccsrc/minddata/dataset/engine/datasetops/device_queue_op.cc

@@ -32,14 +32,15 @@
 namespace mindspore {
 namespace dataset {
 DeviceQueueOp::DeviceQueueOp(std::string channel_name, DeviceType device_type, int32_t device_id, int32_t prefetch_size,
-                             bool send_epoch_end)
+                             bool send_epoch_end, bool create_data_info_queue)
     : PipelineOp(1),
       channel_name_(channel_name),
       device_type_(device_type),
       device_id_(device_id),
       prefetch_size_(prefetch_size),
       send_epoch_end_(send_epoch_end),
-      stop_send_(false) {
+      stop_send_(false),
+      create_data_info_queue_(create_data_info_queue) {
 #ifdef ENABLE_TDTQUE
   ascend_keep_waiting_ = true;
 #endif
@@ -84,6 +85,10 @@ Status DeviceQueueOp::operator()() {
 
   if (device_type_ == DeviceType::Ascend) {
 #ifdef ENABLE_TDTQUE
+    if (create_data_info_queue_) {
+      data_info_queue_ptr_ = std::make_unique<DATA_INFO_QUEUE>(kDataInfoQueueCapacity);
+      RETURN_IF_NOT_OK(data_info_queue_ptr_->Register(tree_->AllTasks()));
+    }
     RETURN_IF_NOT_OK(SendDataToAscend());
 #endif
   } else if (device_type_ == DeviceType::GPU) {
@@ -137,6 +142,13 @@ Status DeviceQueueOp::SendDataToAscend() {
             return Status(StatusCode::kTDTPushFailure, "TDT Push Failed");
           }
         }
+        if (create_data_info_queue_) {
+          DATA_INFO data_info;
+          (void)std::transform(
+            currRow.begin(), currRow.end(), std::back_inserter(data_info),
+            [](const std::shared_ptr<Tensor> &ts) { return std::make_pair(ts->type(), ts->shape()); });
+          RETURN_IF_NOT_OK(data_info_queue_ptr_->Add(data_info));
+        }
 
         if (isProfilingEnable) {
           end_time = ProfilingTime::GetCurMilliSecond();
@@ -187,6 +199,21 @@ Status DeviceQueueOp::SendDataToAscend() {
 
   return Status::OK();
 }
+
+#endif
+
+#ifdef ENABLE_TDTQUE
+Status DeviceQueueOp::GetDataInfo(DATA_INFO *data_info) {
+  if (!create_data_info_queue_) {
+    return Status(StatusCode::kUnexpectedError, __LINE__, __FILE__, "DataInfo queue is not created.");
+  }
+  RETURN_IF_NOT_OK(data_info_queue_ptr_->PopFront(data_info));
+  return Status::OK();
+}
+#else
+Status DeviceQueueOp::GetDataInfo(DATA_INFO *data_info) {
+  return Status(StatusCode::kUnexpectedError, __LINE__, __FILE__, "GetDataInfo is not supported yet.");
+}
 #endif
 
 #ifdef ENABLE_GPUQUE

mindspore/ccsrc/minddata/dataset/engine/datasetops/device_queue_op.h

@@ -18,6 +18,7 @@
 
 #include <memory>
 #include <string>
+#include <utility>
 #include <vector>
 
 #include "minddata/dataset/engine/datasetops/pipeline_op.h"
@@ -25,6 +26,7 @@
 #include "minddata/dataset/util/status.h"
 
 #ifdef ENABLE_TDTQUE
+#include "minddata/dataset/util/queue.h"
 #include "minddata/dataset/engine/tdt/tdt_plugin.h"
 #endif
 
@@ -37,6 +39,10 @@ using mindspore::device::GpuBufferMgr;
 
 namespace mindspore {
 namespace dataset {
+
+using DATA_INFO = std::vector<std::pair<DataType, TensorShape>>;
+using DATA_INFO_QUEUE = Queue<DATA_INFO>;
+const int kDataInfoQueueCapacity = 128;
 class DeviceQueueOp : public PipelineOp {
  public:
   static const uint32_t INVALID_HANDLE = 0xffffffffUL;
@@ -86,13 +92,18 @@ class DeviceQueueOp : public PipelineOp {
       return *this;
     }
 
+    Builder &SetCreateDataInfoQueue(bool create_data_info_queue) {
+      builder_create_data_info_queue_ = create_data_info_queue;
+      return *this;
+    }
     //  Name: Build()
     //  Description: The final step for building a DeviceQueueOp via the Builder is
     //              to call this Build() method.  It will instantiate the DeviceQueueOp
     //              and return it to caller as a shared pointer.
     Status Build(std::shared_ptr<DeviceQueueOp> *ptr) {
       *ptr = std::make_shared<DeviceQueueOp>(builder_channel_name_, builder_device_type_, builder_device_id_,
-                                             builder_prefetch_size_, builder_send_epoch_end_);
+                                             builder_prefetch_size_, builder_send_epoch_end_,
+                                             builder_create_data_info_queue_);
       return Status::OK();
     }
 
@@ -102,12 +113,13 @@ class DeviceQueueOp : public PipelineOp {
     DeviceType builder_device_type_;
     std::string builder_channel_name_;
     bool builder_send_epoch_end_;
+    bool builder_create_data_info_queue_;
   };
 
   //  Name: constructor
   //  Description
   DeviceQueueOp(std::string channel_name, DeviceType device_type, int32_t device_id, int32_t prefetch_size,
-                bool send_epoch_end);
+                bool send_epoch_end, bool create_data_info_queue);
 
   //  Name: destructor
   //  Description
@@ -132,6 +144,8 @@ class DeviceQueueOp : public PipelineOp {
   void StopWaiting() { ascend_keep_waiting_ = false; }
 #endif
 
+  Status GetDataInfo(DATA_INFO *data_info);
+
   // Name: Print()
   // Description: A function that prints info about the node
   void Print(std::ostream &out,              // In: The output stream to print to
@@ -164,6 +178,7 @@ class DeviceQueueOp : public PipelineOp {
 #ifdef ENABLE_TDTQUE
   Status SendDataToAscend();
   bool ascend_keep_waiting_;
+
 #endif
 
 #ifdef ENABLE_GPUQUE
@@ -182,6 +197,8 @@ class DeviceQueueOp : public PipelineOp {
   const int32_t prefetch_size_;
   const bool send_epoch_end_;
   bool stop_send_;
+  bool create_data_info_queue_;
+  std::unique_ptr<DATA_INFO_QUEUE> data_info_queue_ptr_;
 
 #ifdef ENABLE_TDTQUE
   std::shared_ptr<TdtPlugin> tdtInstancePtr;

mindspore/ccsrc/minddata/dataset/kernels/image/image_utils.cc

@@ -97,6 +97,8 @@ Status Resize(const std::shared_ptr<Tensor> &input, std::shared_ptr<Tensor> *out
   if (input_cv->Rank() != 3 && input_cv->Rank() != 2) {
     RETURN_STATUS_UNEXPECTED("Input Tensor is not in shape of <H,W,C> or <H,W>");
   }
+  // OpenCv lanuch too many threads default
+  cv::setNumThreads(0);
   cv::Mat in_image = input_cv->mat();
   // resize image too large or too small
   if (output_height == 0 || output_height > in_image.rows * 1000 || output_width == 0 ||

mindspore/ccsrc/minddata/dataset/kernels/image/lite_cv/image_process.cc

@@ -294,6 +294,9 @@ static bool ConvertRGBAToGRAY(const unsigned char *data, LDataType data_type, in
 }
 
 bool InitFromPixel(const unsigned char *data, LPixelType pixel_type, LDataType data_type, int w, int h, LiteMat &m) {
+  if (data == nullptr) {
+    return false;
+  }
   if (w <= 0 || h <= 0) {
     return false;
   }
@@ -517,6 +520,9 @@ bool Pad(const LiteMat &src, LiteMat &dst, int top, int bottom, int left, int ri
   if (top <= 0 || bottom <= 0 || left <= 0 || right <= 0) {
     return false;
   }
+  if (src.IsEmpty()) {
+    return false;
+  }
   if (pad_type == PADD_BORDER_CONSTANT && src.data_type_ == LDataType::FLOAT32) {
     PadWithConstant<float>(src, dst, top, bottom, left, right, pad_type, fill_b_or_gray, fill_g, fill_r);
   } else if (pad_type == PADD_BORDER_CONSTANT && src.data_type_ == LDataType::UINT8) {

mindspore/ccsrc/minddata/dataset/util/arena.cc

@@ -37,7 +37,8 @@ struct MemHdr {
 ArenaImpl::ArenaImpl(void *ptr, size_t sz) : size_in_bytes_(sz), ptr_(ptr) {
   // Divide the memory into blocks. Ignore the last partial block.
   uint64_t num_blks = size_in_bytes_ / ARENA_BLK_SZ;
-  MS_LOG(DEBUG) << "Size of memory pool is " << num_blks << ", number of blocks of size is " << ARENA_BLK_SZ << ".";
+  MS_LOG(DEBUG) << "Arena memory pool is created. Number of blocks : " << num_blks << ". Block size : " << ARENA_BLK_SZ
+                << ".";
   tr_.Insert(0, num_blks);
 }
 
@@ -233,16 +234,16 @@ std::ostream &operator<<(std::ostream &os, const ArenaImpl &s) {
 
 Status Arena::Init() {
   try {
-    auto sz = size_in_MB_ * 1048576L;
-    mem_ = std::make_unique<uint8_t[]>(sz);
-    impl_ = std::make_unique<ArenaImpl>(mem_.get(), sz);
+    int64_t sz = size_in_MB_ * 1048576L;
+    RETURN_IF_NOT_OK(DeMalloc(sz, &ptr_, false));
+    impl_ = std::make_unique<ArenaImpl>(ptr_, sz);
   } catch (std::bad_alloc &e) {
     return Status(StatusCode::kOutOfMemory);
   }
   return Status::OK();
 }
 
-Arena::Arena(size_t val_in_MB) : size_in_MB_(val_in_MB) {}
+Arena::Arena(size_t val_in_MB) : ptr_(nullptr), size_in_MB_(val_in_MB) {}
 
 Status Arena::CreateArena(std::shared_ptr<Arena> *p_ba, size_t val_in_MB) {
   RETURN_UNEXPECTED_IF_NULL(p_ba);

mindspore/ccsrc/minddata/dataset/util/arena.h

@@ -19,6 +19,7 @@
 #include <memory>
 #include <mutex>
 #include <utility>
+#include "minddata/dataset/util/allocator.h"
 #include "minddata/dataset/util/memory_pool.h"
 #include "minddata/dataset/util/treap.h"
 
@@ -103,7 +104,12 @@ class Arena : public MemoryPool {
   // Disable copy and assignment constructor
   Arena(const Arena &) = delete;
   Arena &operator=(const Arena &) = delete;
-  ~Arena() override = default;
+  ~Arena() override {
+    if (ptr_ != nullptr) {
+      free(ptr_);
+    }
+    ptr_ = nullptr;
+  }
 
   /// As a derived class of MemoryPool, we have to implement the following.
   /// But we simply transfer the call to the implementation class
@@ -140,7 +146,7 @@ class Arena : public MemoryPool {
  protected:
   mutable std::mutex mux_;
   std::unique_ptr<ArenaImpl> impl_;
-  std::unique_ptr<uint8_t[]> mem_;
+  void *ptr_;
   size_t size_in_MB_;
 
   explicit Arena(size_t val_in_MB = 4096);

mindspore/ccsrc/minddata/dataset/util/services.cc

@@ -30,6 +30,8 @@ namespace mindspore {
 namespace dataset {
 std::unique_ptr<Services> Services::instance_ = nullptr;
 std::once_flag Services::init_instance_flag_;
+std::set<std::string> Services::unique_id_list_ = {};
+std::mutex Services::unique_id_mutex_;
 
 #if !defined(_WIN32) && !defined(_WIN64) && !defined(__ANDROID__) && !defined(ANDROID)
 std::string Services::GetUserName() {
@@ -52,8 +54,23 @@ std::string Services::GetUniqueID() {
   std::mt19937 gen = GetRandomDevice();
   std::uniform_int_distribution<uint32_t> dist(0, kStr.size() - 1);
   char buffer[UNIQUEID_LEN];
-  for (int i = 0; i < UNIQUEID_LEN; i++) {
-    buffer[i] = kStr[dist(gen)];
+  {
+    std::unique_lock<std::mutex> lock(unique_id_mutex_);
+    while (true) {
+      auto ret = memset_s(buffer, UNIQUEID_LEN, 0, UNIQUEID_LEN);
+      if (ret != 0) {
+        MS_LOG(ERROR) << "memset_s error, errorno(" << ret << ")";
+        return std::string("");
+      }
+      for (int i = 0; i < UNIQUEID_LEN; i++) {
+        buffer[i] = kStr[dist(gen)];
+      }
+      if (unique_id_list_.find(std::string(buffer, UNIQUEID_LEN)) != unique_id_list_.end()) {
+        continue;
+      }
+      unique_id_list_.insert(std::string(buffer, UNIQUEID_LEN));
+      break;
+    }
   }
   return std::string(buffer, UNIQUEID_LEN);
 }

mindspore/ccsrc/minddata/dataset/util/services.h

@@ -19,6 +19,7 @@
 #include <algorithm>
 #include <memory>
 #include <mutex>
+#include <set>
 #include <string>
 #include <vector>
 #include "minddata/dataset/util/memory_pool.h"
@@ -97,6 +98,8 @@ class Services {
  private:
   static std::once_flag init_instance_flag_;
   static std::unique_ptr<Services> instance_;
+  static std::set<std::string> unique_id_list_;
+  static std::mutex unique_id_mutex_;
   // A small pool used for small objects that last until the
   // Services Manager shuts down. Used by all sub-services.
   std::shared_ptr<MemoryPool> pool_;

mindspore/ccsrc/minddata/dataset/util/sig_handler.cc

@@ -35,6 +35,7 @@ void RegisterHandlers() {
   (void)sigemptyset(&new_int_action.sa_mask);
   new_int_action.sa_flags = SA_RESTART | SA_SIGINFO;
   (void)sigaction(SIGINT, &new_int_action, nullptr);
+  (void)sigaction(SIGTERM, &new_int_action, nullptr);
 }
 
 extern void IntHandler(int sig_num,          // The signal that was raised

mindspore/ccsrc/minddata/dataset/util/status.cc

@@ -15,7 +15,13 @@
  */
 #include "minddata/dataset/util/status.h"
 #include <sstream>
+#include <string>
+#include <memory.h>
+#include <stdio.h>
+#include <stdlib.h>
 #include "utils/ms_utils.h"
+#include "./securec.h"
+
 #include "minddata/dataset/util/task_manager.h"
 
 namespace mindspore {
@@ -119,5 +125,55 @@ std::ostream &operator<<(std::ostream &os, const Status &s) {
 std::string Status::ToString() const { return err_msg_; }
 
 StatusCode Status::get_code() const { return code_; }
+
+#if !defined(_WIN32) && !defined(_WIN64)
+float GetMemoryUsage() {
+  char buf[128] = {0};
+
+  FILE *fd;
+  fd = fopen("/proc/meminfo", "r");
+  if (fd == nullptr) {
+    MS_LOG(WARNING) << "The meminfo file: /proc/meminfo is opened failed.";
+    return 0.0;
+  }
+
+  uint32_t status_count = 0;
+  uint64_t mem_total = 0L;
+  uint64_t mem_available = 0L;
+  while (fgets(buf, sizeof(buf), fd)) {
+    if (status_count == 2) {  // get MemTotal and MemAvailable yet
+      break;
+    }
+
+    // get title
+    std::string line(buf);
+    std::string::size_type position = line.find(":");
+    std::string title = line.substr(0, position);
+
+    // get the value when MemTotal or MemAvailable
+    if (title == "MemTotal") {
+      std::string::size_type pos1 = line.find_last_of(" ");
+      std::string::size_type pos2 = line.find_last_of(" ", pos1 - 1);
+      mem_total = atol(line.substr(pos2, pos1 - pos2).c_str());
+      status_count++;
+    } else if (title == "MemAvailable") {
+      std::string::size_type pos1 = line.find_last_of(" ");
+      std::string::size_type pos2 = line.find_last_of(" ", pos1 - 1);
+      mem_available = atol(line.substr(pos2, pos1 - pos2).c_str());
+      status_count++;
+    }
+
+    (void)memset_s(buf, sizeof(buf), 0, sizeof(buf));
+  }
+  fclose(fd);
+
+  if (status_count != 2 || mem_total == 0 || mem_available > mem_total) {
+    MS_LOG(WARNING) << "Get memory usage failed.";
+    return 0.0;
+  }
+
+  return (1.0 - static_cast<float>(static_cast<double>(mem_available) / static_cast<double>(mem_total)));
+}
+#endif
 }  // namespace dataset
 }  // namespace mindspore

mindspore/ccsrc/minddata/dataset/util/status.h

@@ -141,6 +141,12 @@ class Status {
   StatusCode code_;
   std::string err_msg_;
 };
+
+#if !defined(_WIN32) && !defined(_WIN64)
+const float MAX_MEMORY_USAGE_THRESHOLD = 0.95;
+
+float GetMemoryUsage();
+#endif
 }  // namespace dataset
 }  // namespace mindspore
 #endif  // MINDSPORE_CCSRC_MINDDATA_DATASET_UTIL_STATUS_H_

mindspore/ccsrc/pipeline/jit/parse/parse.cc

@@ -1501,7 +1501,7 @@ void Parser::HandleAssignSubscript(const FunctionBlockPtr &block, const py::obje
   AnfNodePtr slice_node = ParseExprNode(block, slice_obj);
   CNodePtr setitem_app = block->func_graph()->NewCNode({op_setitem, value_node, slice_node, assigned_node});
   // getitem apply should return the sequence data structure itself
-  std::string var_name = "";
+  std::string var_name;
   if (ast_->IsClassMember(value_obj)) {
     std::string attr_name = value_obj.attr("attr").cast<std::string>();
     var_name = "self." + attr_name;
@@ -1515,9 +1515,18 @@ void Parser::HandleAssignSubscript(const FunctionBlockPtr &block, const py::obje
                               << py::str(obj).cast<std::string>() << "' with type '"
                               << py::str(obj_type).cast<std::string>() << "'.";
     }
-  } else {
-    var_name = value_obj.attr("id").cast<std::string>();
+    block->WriteVariable(var_name, setitem_app);
+    return;
   }
+  if (AstSubType(py::cast<int32_t>(ast_->CallParserObjMethod(PYTHON_PARSE_GET_AST_TYPE, value_obj))) ==
+      AST_SUB_TYPE_SUBSCRIPT) {
+    HandleAssignSubscript(block, value_obj, setitem_app);
+    return;
+  }
+  if (!py::hasattr(value_obj, "id")) {
+    MS_EXCEPTION(TypeError) << "Attribute id not found in " << py::str(value_obj).cast<std::string>();
+  }
+  var_name = value_obj.attr("id").cast<std::string>();
   block->WriteVariable(var_name, setitem_app);
 }
 
@@ -1790,38 +1799,92 @@ bool UpdateFuncGraphFlags(py::object obj, const FuncGraphPtr &func_graph) {
   return true;
 }
 
+// Generate and copy a ValueNode, or a CNode with its child nodes
+static AnfNodePtr CopyNodesFromParamDefaultValue(const FuncGraphPtr func_graph, const AnfNodePtr &param_node) {
+  MS_EXCEPTION_IF_NULL(param_node);
+  if (param_node->isa<ValueNode>()) {
+    return std::make_shared<ValueNode>(param_node->cast<ValueNodePtr>()->value());
+  }
+
+  // Parameter default value is CNode.
+  std::size_t index = 0;
+  std::vector<AnfNodePtr> old_cnodes;
+  old_cnodes.emplace_back(param_node);
+  auto res = func_graph->NewCNode({});
+  std::vector<CNodePtr> new_cnodes;
+  new_cnodes.emplace_back(res);
+  while (index < old_cnodes.size()) {
+    auto current = old_cnodes[index];
+    auto current_new_cnode = new_cnodes[index];
+    index++;
+    MS_EXCEPTION_IF_NULL(current);
+    if (current->isa<CNode>()) {
+      auto &inputs = current->cast<CNodePtr>()->inputs();
+      for (auto it = inputs.begin(); it != inputs.end(); it++) {
+        AnfNodePtr input = *it;
+        if (input != nullptr && input->isa<CNode>()) {
+          old_cnodes.emplace_back(input);
+          auto new_cnode = func_graph->NewCNode({});
+          new_cnodes.emplace_back(new_cnode);
+          current_new_cnode->add_input(new_cnode);
+        } else if (input->isa<ValueNode>()) {
+          current_new_cnode->add_input(std::make_shared<ValueNode>(input->cast<ValueNodePtr>()->value()));
+        } else {
+          MS_LOG(EXCEPTION) << "Wrong type item in default parameters: " << input->ToString();
+        }
+      }
+    }
+  }
+  return res;
+}
+
 FuncGraphPtr MakeTopGraph(const py::object &cell, const ValuePtr &cell_ptr) {
+  auto current_graph = dyn_cast<FuncGraph>(cell_ptr);
+  if (current_graph == nullptr) {
+    MS_LOG(EXCEPTION) << "Current graph cast failed from " << cell_ptr->ToString();
+  }
+
   auto func_graph = std::make_shared<FuncGraph>();
-  func_graph->debug_info()->set_name("top");
+  func_graph->debug_info()->set_name(current_graph->debug_info()->name() + "_wrapper");
 
-  // def top(*arg, *kwargs):
-  auto param_vargs = func_graph->add_parameter();
-  auto args_name = "args";
-  param_vargs->set_name(args_name);
-  param_vargs->debug_info()->set_name(args_name);
-
-  auto param_vkwargs = func_graph->add_parameter();
-  args_name = "kwargs";
-  param_vkwargs->set_name(args_name);
-  param_vkwargs->debug_info()->set_name(args_name);
-
-  func_graph->set_has_vararg(true);
-  func_graph->set_has_kwarg(true);
-  func_graph->set_kwonlyargs_count(0);
+  // Copy all parameters information
+  for (auto &para : current_graph->parameters()) {
+    auto param = func_graph->add_parameter();
+    auto orig_param = para->cast<ParameterPtr>();
+    auto name = orig_param->name();
+    param->set_name(name);
+    param->debug_info()->set_name(name);
+  }
+  func_graph->set_has_vararg(current_graph->has_vararg());
+  func_graph->set_has_kwarg(current_graph->has_kwarg());
+  func_graph->set_kwonlyargs_count(current_graph->kwonlyargs_count());
+  // Copy all default values
+  for (auto &d : current_graph->parameter_default_value()) {
+    func_graph->set_param_default_value(d.first, CopyNodesFromParamDefaultValue(func_graph, d.second));
+  }
 
   // cell_obj
+  MS_LOG(DEBUG) << "add Flag for " << std::string(py::str(cell));
   parse::UpdateFuncGraphFlags(cell, func_graph);
   // top graph's construct flag
   if (py::hasattr(cell, "construct")) {
     parse::UpdateFuncGraphFlags(cell.attr("construct"), func_graph);
   }
 
-  // ret = cell_obj(*arg, *kwargs)
-  auto call_fn = MakeUnpackCall(func_graph, NewValueNode(cell_ptr), {param_vargs, param_vkwargs});
-
-  // return ret
-  func_graph->set_output(call_fn);
-  MS_LOG(DEBUG) << "add Flag for " << std::string(py::str(cell));
+  auto unpacking = func_graph->has_vararg() || func_graph->has_kwarg();
+  if (!unpacking) {
+    std::vector<AnfNodePtr> inputs;
+    inputs.emplace_back(NewValueNode(cell_ptr));
+    auto &params = func_graph->parameters();
+    (void)std::transform(params.begin(), params.end(), std::back_inserter(inputs),
+                         [](AnfNodePtr node) -> AnfNodePtr { return node; });
+    func_graph->set_output(func_graph->NewCNode(inputs));
+  } else {
+    // ret = cell_obj(*arg, *kwargs)
+    auto call_fn = MakeUnpackCall(func_graph, NewValueNode(cell_ptr), func_graph->parameters());
+    // return ret
+    func_graph->set_output(call_fn);
+  }
   return func_graph;
 }
 }  // namespace parse

mindspore/ccsrc/pipeline/jit/pass.cc

@@ -99,7 +99,7 @@ OptPassGroupMap GetOptPassesA(const opt::irpass::OptimizeIRPassLib &irpass) {
     irpass.specialize_transform_,
 
     // Miscellaneous
-    irpass.item_tuple_eliminate_,
+    irpass.item_tuple_or_list_eliminate_,
     irpass.env_get_item_eliminate_,
     irpass.cast_eliminate_,
     irpass.reshape_eliminate_,
@@ -186,8 +186,9 @@ OptPassGroupMap GetOptPassesAfterCconv(const opt::irpass::OptimizeIRPassLib &irp
 }
 
 OptPassGroupMap GetOptPassesTransformGraph(const opt::irpass::OptimizeIRPassLib &irpass) {
-  opt::OptPassConfig d_1 = opt::OptPassConfig({// Safe inlining
-                                               irpass.call_graph_tuple_transform_, irpass.item_tuple_eliminate_});
+  opt::OptPassConfig d_1 =
+    opt::OptPassConfig({// Safe inlining
+                        irpass.call_graph_tuple_transform_, irpass.item_tuple_or_list_eliminate_});
 
   OptPassGroupMap map_a({{"d_1", d_1}, {"renormalize", opt::OptPassConfig::Renormalize()}});
 
@@ -196,7 +197,7 @@ OptPassGroupMap GetOptPassesTransformGraph(const opt::irpass::OptimizeIRPassLib
 
 OptPassGroupMap GetOptPassesB(const opt::irpass::OptimizeIRPassLib &irpass) {
   opt::OptPassConfig b_1 = opt::OptPassConfig(
-    {irpass.zero_like_fill_zero_, irpass.item_tuple_eliminate_, irpass.float_tuple_getitem_switch_,
+    {irpass.zero_like_fill_zero_, irpass.item_tuple_or_list_eliminate_, irpass.float_tuple_getitem_switch_,
      irpass.reset_defer_inline_, irpass.inline_, irpass.special_op_eliminate_, irpass.get_make_ref_eliminate_,
      irpass.incorporate_env_getitem_, irpass.incorporate_env_getitem_switch_, irpass.env_get_item_eliminate_,
      irpass.incorporate_env_getitem_switch_layer_, irpass.value_based_eliminate_});

mindspore/ccsrc/pipeline/jit/pipeline.cc

@@ -76,6 +76,7 @@ const char IR_TYPE_MINDIR[] = "mind_ir";
 
 ExecutorPyPtr ExecutorPy::executor_ = nullptr;
 std::mutex ExecutorPy::instance_lock_;
+bool ExecutorPy::debugger_terminate_ = false;
 
 std::unordered_map<abstract::AbstractBasePtrList, int, abstract::AbstractBasePtrListHasher,
                    abstract::AbstractBasePtrListEqual>
@@ -748,7 +749,17 @@ void ExecutorPy::ProcessVmArg(const py::tuple &args, const std::string &phase, V
   ProcessVmArgInner(args, GetResource(phase), arg_list);
 }
 
+void ExecutorPy::TerminateDebugger() {
+  if (debugger_terminate_) {
+    MS_LOG(INFO) << "Terminate debugger and clear resources!";
+    ClearResAtexit();
+    exit(1);
+  }
+}
+
 py::object ExecutorPy::Run(const py::tuple &args, const py::object &phase) {
+  // Mindspore debugger notify main thread to exit after one step, and will not run next step
+  TerminateDebugger();
   std::size_t size = args.size();
   if (!py::isinstance<py::str>(phase)) {
     MS_LOG(EXCEPTION) << "Run failed, phase input is not a str";

mindspore/ccsrc/pipeline/jit/pipeline.h

@@ -97,6 +97,9 @@ class ExecutorPy : public std::enable_shared_from_this<ExecutorPy> {
   void DelNetRes(const std::string &id);
   void ReleaseResource(const py::object &phase);
   static void ClearRes();
+  static bool GetDebugTerminate() { return debugger_terminate_; }
+  static void DebugTerminate(bool val) { debugger_terminate_ = val; }
+  void TerminateDebugger();
 
   std::map<std::string, std::pair<PrimitivePyPtr, std::string>> FetchInfoForQuantExport(const std::string &phase_s);
 
@@ -111,6 +114,7 @@ class ExecutorPy : public std::enable_shared_from_this<ExecutorPy> {
   std::map<std::string, ExecutorInfoPtr> info_;
   static std::shared_ptr<ExecutorPy> executor_;
   static std::mutex instance_lock_;
+  static bool debugger_terminate_;
 };
 using ExecutorPyPtr = std::shared_ptr<ExecutorPy>;
 

mindspore/ccsrc/pipeline/pynative/pynative_execute.cc

@@ -260,7 +260,7 @@ py::object DoAutoCast(const py::object &arg, const TypeId &type_id) {
 py::object DoParamMixPrecisionCast(bool *is_cast, const py::object obj) {
   auto tensor = py::cast<tensor::TensorPtr>(obj);
   auto cast_type = tensor->cast_dtype();
-  py::object cast_output;
+  py::object cast_output = obj;
   if (cast_type != nullptr) {
     auto source_element = tensor->Dtype();
     if (source_element != nullptr && IsSubType(source_element, kFloat) && *source_element != *cast_type) {
@@ -282,12 +282,15 @@ py::object DoParamMixPrecisionCastTuple(bool *is_cast, const py::tuple tuple) {
       result[i] = DoParamMixPrecisionCast(is_cast, tuple[i]);
     } else if (py::isinstance<py::tuple>(tuple[i])) {
       result[i] = DoParamMixPrecisionCastTuple(is_cast, tuple[i]);
+    } else {
+      result[i] = tuple[i];
     }
   }
   return result;
 }
 
 bool GetSignatureType(const PrimitivePyPtr &prim, std::vector<SignatureEnumDType> *dtypes) {
+  MS_EXCEPTION_IF_NULL(dtypes);
   auto signature = prim->signatures();
   bool has_sig_dtype = false;
   (void)std::transform(signature.begin(), signature.end(), std::back_inserter(*dtypes),
@@ -733,20 +736,29 @@ ValuePtr PynativeExecutor::GetForwardValue(const OpExecInfoPtr &op_exec_info) {
 
 AnfNodePtr PynativeExecutor::MakeCNode(const OpExecInfoPtr &op_exec_info, std::vector<bool> *op_masks,
                                        abstract::AbstractBasePtrList *args_spec_list) {
+  MS_EXCEPTION_IF_NULL(op_masks);
+  MS_EXCEPTION_IF_NULL(args_spec_list);
   CNodePtr cnode = nullptr;
   std::vector<AnfNodePtr> inputs;
+
   auto prim = op_exec_info->py_primitive;
+  const auto &signature = prim->signatures();
+
   inputs.push_back(NewValueNode(prim));
 
   size_t size = op_exec_info->op_inputs.size();
+  auto sig_size = signature.size();
   // ignore signature for cast op
+  if (sig_size > 0 && sig_size != size) {
+    MS_EXCEPTION(ValueError) << op_exec_info->op_name << " inputs size " << size << " does not match the requires "
+                             << "inputs size " << sig_size;
+  }
   bool is_cast_op = (op_exec_info->op_name == "Cast");
   if (!is_cast_op) {
-    const auto &signature = prim->signatures();
     for (size_t i = 0; i < size; i++) {
       auto obj = op_exec_info->op_inputs[i];
       auto sig = SignatureEnumRW::kRWDefault;
-      if (signature.size() > 0) {
+      if (sig_size > 0) {
         sig = signature[i].rw;
       }
       MS_LOG(DEBUG) << "check mix precision " << op_exec_info->op_name << " input " << i << " "
@@ -1173,10 +1185,6 @@ py::tuple RunOp(const py::args &args) {
     std::ostringstream oss;
     trace::TraceGraphEval();
     trace::GetEvalStackInfo(oss);
-    // call py::print to output function call stack to STDOUT, in case of output the log to file, the user can see
-    // these info from screen, no need to open log file to find these info
-    py::print(oss.str());
-    MS_LOG(ERROR) << oss.str();
     PynativeExecutor::GetInstance()->Clean();
     // re-throw this exception to Python interpreter to handle it
     throw(py::error_already_set(ex));

mindspore/ccsrc/profiler/device/gpu/gpu_profiling.cc

@@ -406,8 +406,10 @@ void GPUProfiler::OpDataProducerBegin(const std::string op_name, void *stream) {
     CHECK_CUDA_RET_WITH_ERROR(cudaEventRecord(op_event_start_, (CUstream)stream_),
                               "cudaEventRecord op event start failed");
     op_host_time_start_ = GetHostTimeStamp();
+    op_cupti_time_start_ = GetCUPTITimeStamp();
   } else {
     op_host_time_start_ = GetHostTimeStamp();
+    op_cupti_time_start_ = GetCUPTITimeStamp();
   }
   SetRunTimeData(op_name, stream);
 }
@@ -431,7 +433,7 @@ void GPUProfiler::OpDataProducerEnd() {
   }
   MS_LOG(DEBUG) << "Host Time Elapsed(us)," << op_name_ << "," << op_time_elapsed;
   SetRunTimeData(op_name_, op_time_elapsed);
-  SetRunTimeData(op_name_, op_host_time_start_, op_time_elapsed);
+  SetRunTimeData(op_name_, op_cupti_time_start_, op_time_elapsed);
 }
 
 void GPUProfiler::StopCUPTI() {

mindspore/ccsrc/profiler/device/gpu/gpu_profiling.h

@@ -170,6 +170,7 @@ class GPUProfiler {
   cudaEvent_t op_event_stop_;
   uint64_t op_host_time_start_;
   uint64_t op_host_time_stop_;
+  uint64_t op_cupti_time_start_;
   std::string profile_data_path_;
 };
 }  // namespace gpu

mindspore/ccsrc/pybind_api/ir/tensor_py.cc

@@ -455,7 +455,7 @@ REGISTER_PYBIND_DEFINE(Tensor, ([](const py::module *m) {
                                   >>> data.set_dtype(mindspore.int32)
                                   mindspore.int32
                               )mydelimiter")
-                           .def("set_cast_dtype", &Tensor::set_cast_dtype)
+                           .def("set_cast_dtype", &Tensor::set_cast_dtype, py::arg("dtype") = nullptr)
                            .def("__str__", &Tensor::ToString)
                            .def("__repr__", &Tensor::ToStringRepr)
                            .def(py::pickle(

mindspore/ccsrc/runtime/device/CMakeLists.txt

@@ -14,12 +14,12 @@ endif ()
 
 if (ENABLE_CPU)
     file(GLOB_RECURSE CPU_SRC_LIST RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "cpu/*.cc")
-    list(REMOVE_ITEM CPU_SRC_LIST "cpu/mpi/mpi_adapter.cc", "cpu/mpi/mpi_export.cc")
+    list(REMOVE_ITEM CPU_SRC_LIST "cpu/mpi/mpi_adapter.cc" "cpu/mpi/mpi_export.cc")
 endif ()
 
 if (ENABLE_MPI)
     if (ENABLE_CPU)
-        file(GLOB_RECURSE MPI_SRC_LIST "cpu/mpi/mpi_adapter.cc", "cpu/mpi/mpi_export.cc")
+        file(GLOB_RECURSE MPI_SRC_LIST "cpu/mpi/mpi_adapter.cc" "cpu/mpi/mpi_export.cc")
         set_property(SOURCE ${MPI_SRC_LIST}
             PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_DEVICE)
         add_library(mpi_adapter SHARED ${MPI_SRC_LIST})

mindspore/ccsrc/runtime/device/ascend/ascend_kernel_runtime.cc

@@ -40,6 +40,7 @@
 #include "runtime/device/ascend/ascend_memory_manager.h"
 #include "debug/tensor_load.h"
 #include "debug/data_dump/dump_json_parser.h"
+#include "toolchain/adx_datadump_server.h"
 #include "utils/shape_utils.h"
 #ifdef MEM_REUSE_DEBUG
 #include "backend/optimizer/mem_reuse/mem_reuse_checker.h"
@@ -169,6 +170,14 @@ bool AscendKernelRuntime::NeedDestroyHccl() {
   return true;
 }
 
+void AsyncDataDumpUninit() {
+  if (DumpJsonParser::GetInstance().async_dump_enabled()) {
+    if (AdxDataDumpServerUnInit() != 0) {
+      MS_LOG(ERROR) << "Adx data dump server uninit failed";
+    }
+  }
+}
+
 void AscendKernelRuntime::ReleaseDeviceRes() {
   MS_LOG(INFO) << "Ascend finalize start";
 #ifdef ENABLE_DEBUGGER
@@ -184,6 +193,8 @@ void AscendKernelRuntime::ReleaseDeviceRes() {
   // release ge runtime
   ClearGraphModelMap();
 
+  AsyncDataDumpUninit();
+
   auto context_ptr = MsContext::GetInstance();
   MS_EXCEPTION_IF_NULL(context_ptr);
   auto ret = rtSetDevice(context_ptr->get_param<uint32_t>(MS_CTX_DEVICE_ID));

mindspore/ccsrc/runtime/device/ascend/dump/data_dumper.cc

@@ -57,7 +57,6 @@ constexpr const char *kOpTypeOpDebug = "Opdebug";
 namespace mindspore {
 namespace device {
 namespace ascend {
-
 DataDumper::~DataDumper() {
   ReleaseDevMem(&dev_load_mem_);
   ReleaseDevMem(&dev_unload_mem_);
@@ -116,7 +115,7 @@ void DataDumper::LoadDumpInfo() {
   auto debugger = mindspore::Debugger::GetInstance();
   MS_EXCEPTION_IF_NULL(debugger);
   if (debugger->DebuggerBackendEnabled()) {
-    std::map<std::pair<uint32_t, uint32_t>, std::string> &stream_task_to_opname = debugger->GetStreamTaskToOpnameMap();
+    std::map<std::pair<uint32_t, uint32_t>, std::string> stream_task_to_opname = debugger->GetStreamTaskToOpnameMap();
     // extract stream id, task id and opname from runtime_info_map for overflow detection
     std::transform(runtime_info_map_.begin(), runtime_info_map_.end(),
                    std::inserter(stream_task_to_opname, stream_task_to_opname.end()),

mindspore/ccsrc/runtime/device/ascend/kernel_select_ascend.cc

@@ -52,8 +52,12 @@ const int kUnSupportMixedDataTypeIndex = -1;
 bool MatchInferOutputDataType(const CNodePtr &cnode, const kernel::KernelBuildInfo &kernel_build_info) {
   MS_EXCEPTION_IF_NULL(cnode);
   // Check input data type
+  auto name = AnfAlgo::GetCNodeName(cnode);
   for (size_t input_index = 0; input_index < kernel_build_info.GetInputNum(); ++input_index) {
     TypeId input_origin_type = AnfAlgo::GetPrevNodeOutputInferDataType(cnode, input_index);
+    if (name == kDynamicRNNOpName && input_origin_type == kMetaTypeNone) {
+      continue;
+    }
     if (kernel_build_info.GetInputDeviceType(input_index) != input_origin_type) {
       return false;
     }
@@ -478,6 +482,9 @@ void SetTensorDeviceInfo(const kernel::KernelBuildInfo &selected_kernel_info, co
       AnfAlgo::SetSelectKernelBuildInfo(builder->Build(), input_kernel_node.get());
       continue;
     }
+    if (selected_kernel_info.GetInputFormat(input_index) == kOpFormat_FRACTAL_ZN_LSTM) {
+      continue;
+    }
     // we set special device info of a input tensor.
     bool is_ref = false;
     auto op_info = kernel::OpLib::FindOp(AnfAlgo::GetCNodeName(kernel_node), kernel::kTBE);

mindspore/ccsrc/runtime/device/ascend/tasksink/task_generator.cc

@@ -127,8 +127,12 @@ bool TaskGenerator::LaunchKernel(const CNodePtr &anf_node_ptr, uint32_t stream_i
   auto kernel_mod = AnfAlgo::GetKernelMod(anf_node_ptr);
   MS_EXCEPTION_IF_NULL(kernel_mod);
   kernel_mod->set_kernel_name(anf_node_ptr->fullname_with_scope());
+  auto op_name = AnfAlgo::GetCNodeName(anf_node_ptr);
   if (AnfAlgo::GetCNodeName(anf_node_ptr) != kAtomicAddrCleanOpName) {
     for (size_t i = 0; i < AnfAlgo::GetInputTensorNum(anf_node_ptr); ++i) {
+      if (op_name == kDynamicRNNOpName && i == 3) {
+        continue;
+      }
       auto real_input_index = AnfAlgo::GetRealInputIndex(anf_node_ptr, i);
       auto device_address = AnfAlgo::GetPrevNodeOutputAddr(anf_node_ptr, real_input_index);
       AddressPtr input = std::make_shared<Address>();

mindspore/ccsrc/runtime/device/gpu/gpu_common.h

@@ -141,7 +141,6 @@ inline bool CheckNullInput(std::vector<size_t> input_shape) {
       MS_LOG(EXCEPTION) << "CUAD curand Error: " << message << " | curandStatus: " << status; \
     }                                                                                         \
   }
-
 }  // namespace gpu
 }  // namespace device
 }  // namespace mindspore

mindspore/ccsrc/runtime/device/gpu/mpi/mpi_initializer.cc

@@ -25,7 +25,6 @@
 namespace mindspore {
 namespace device {
 namespace gpu {
-
 MPIInitializer &MPIInitializer::GetInstance() {
   static MPIInitializer instance;
   return instance;

mindspore/ccsrc/transform/graph_ir/op_adapter_map.h

@@ -78,6 +78,8 @@ constexpr const char kNameApplyAdam[] = "Adam";
 constexpr const char kNameExtractImagePatches[] = "ExtractImagePatches";
 constexpr const char kNameReLU6[] = "ReLU6";
 constexpr const char kNameReLU6Grad[] = "ReLU6Grad";
+constexpr const char kNameSoftplus[] = "Softplus";
+constexpr const char kNameSoftplusGrad[] = "SoftplusGrad";
 constexpr const char kNameElu[] = "Elu";
 constexpr const char kNameEluGrad[] = "EluGrad";
 constexpr const char kNameTensorScatterUpdate[] = "TensorScatterUpdate";

mindspore/ccsrc/transform/graph_ir/op_declare/nonlinear_fuc_ops_declare.cc

@@ -71,6 +71,18 @@ ATTR_MAP(Relu6Grad) = EMPTY_ATTR_MAP;
 OUTPUT_MAP(Relu6Grad) = {{0, OUTPUT_DESC(backprops)}};
 REG_ADPT_DESC(Relu6Grad, kNameReLU6Grad, ADPT_DESC(Relu6Grad))
 
+// Softplus
+INPUT_MAP(Softplus) = {{1, INPUT_DESC(x)}};
+ATTR_MAP(Softplus) = EMPTY_ATTR_MAP;
+OUTPUT_MAP(Softplus) = {{0, OUTPUT_DESC(y)}};
+REG_ADPT_DESC(Softplus, kNameSoftplus, ADPT_DESC(Softplus))
+
+// SoftplusGrad
+INPUT_MAP(SoftplusGrad) = {{1, INPUT_DESC(gradients)}, {2, INPUT_DESC(features)}};
+ATTR_MAP(SoftplusGrad) = EMPTY_ATTR_MAP;
+OUTPUT_MAP(SoftplusGrad) = {{0, OUTPUT_DESC(backprops)}};
+REG_ADPT_DESC(SoftplusGrad, kNameSoftplusGrad, ADPT_DESC(SoftplusGrad))
+
 // ReluGrad
 INPUT_MAP(ReluGrad) = {{1, INPUT_DESC(gradients)}, {2, INPUT_DESC(features)}};
 ATTR_MAP(ReluGrad) = EMPTY_ATTR_MAP;

mindspore/ccsrc/transform/graph_ir/op_declare/nonlinear_fuc_ops_declare.h

@@ -19,8 +19,8 @@
 
 #include <string>
 #include <unordered_map>
-#include "transform/graph_ir/op_declare/op_declare_macro.h"
 #include "ops/nonlinear_fuc_ops.h"
+#include "transform/graph_ir/op_declare/op_declare_macro.h"
 
 namespace mindspore::transform {
 DECLARE_OP_ADAPTER(ReluGrad)
@@ -32,6 +32,12 @@ DECLARE_OP_USE_OUTPUT(Relu6)
 DECLARE_OP_ADAPTER(Relu6Grad)
 DECLARE_OP_USE_OUTPUT(Relu6Grad)
 
+DECLARE_OP_ADAPTER(Softplus)
+DECLARE_OP_USE_OUTPUT(Softplus)
+
+DECLARE_OP_ADAPTER(SoftplusGrad)
+DECLARE_OP_USE_OUTPUT(SoftplusGrad)
+
 DECLARE_OP_ADAPTER(Tanh)
 DECLARE_OP_USE_OUTPUT(Tanh)
 

mindspore/ccsrc/utils/context/context_extends.cc

@@ -76,7 +76,7 @@ bool OpenTsd(const std::shared_ptr<MsContext> &ms_context_ptr) {
   MS_LOG(INFO) << "Device id = " << device_id << ", rank size = " << rank_size << ".";
   TDT_StatusT status = TsdOpen(device_id, rank_size);
   if (status != TDT_OK) {
-    MS_LOG(EXCEPTION) << "Device " << device_id << " is occupied, open tsd failed, status = " << status << ".";
+    MS_LOG(EXCEPTION) << "Device " << device_id << " open tsd failed, status = " << status << ".";
     return false;
   }
   ms_context_ptr->increase_param<uint32_t>(MS_CTX_TSD_REF);