Browse Source

!334 sync code 0807

Merge pull request !334 from changzherui/code_sync_0807
tags/v0.7.0-beta
mindspore-ci-bot Gitee 5 years ago
parent
commit
1546d803eb
100 changed files with 1732 additions and 2968 deletions
  1. +18
    -2
      .gitignore
  2. +18
    -0
      .gitmodules
  3. +9
    -2
      CMakeLists.txt
  4. +8
    -8
      README.md
  5. +71
    -0
      RELEASE.md
  6. +1
    -1
      akg
  7. +329
    -128
      build.sh
  8. +0
    -7
      cmake/external_libs/dlpack.cmake
  9. +0
    -7
      cmake/external_libs/dmlc_core.cmake
  10. +1
    -1
      cmake/external_libs/glog.cmake
  11. +1
    -0
      cmake/external_libs/jpeg_turbo.cmake
  12. +6
    -6
      cmake/external_libs/mkl_dnn.cmake
  13. +0
    -7
      cmake/external_libs/rang.cmake
  14. +25
    -0
      cmake/external_libs/sentencepiece.cmake
  15. +0
    -15
      cmake/external_libs/tvm_gpu.cmake
  16. +11
    -11
      cmake/mind_expression.cmake
  17. +4
    -4
      cmake/options.cmake
  18. +38
    -29
      cmake/package.cmake
  19. +5
    -2
      cmake/package_script.cmake
  20. +4
    -2
      config/data_dump.json
  21. +67
    -0
      docker/mindspore-cpu/0.6.0-beta/Dockerfile
  22. +83
    -0
      docker/mindspore-gpu/0.6.0-beta/Dockerfile
  23. +1
    -1
      graphengine
  24. +0
    -0
      hub/docs/.gitkeep
  25. +0
    -0
      hub/images/.gitkeep
  26. +0
    -0
      hub/scripts/.gitkeep
  27. +108
    -0
      include/infer_log.h
  28. +204
    -0
      include/infer_tensor.h
  29. +48
    -10
      include/inference.h
  30. +0
    -69
      include/ms_tensor.h
  31. +0
    -18
      mindspore/_akg/__init__.py
  32. +0
    -62
      mindspore/_akg/add_path.py
  33. +0
    -39
      mindspore/_akg/gpu/__init__.py
  34. +0
    -43
      mindspore/_akg/gpu/cast.py
  35. +0
    -56
      mindspore/_akg/gpu/default_schedule.py
  36. +0
    -40
      mindspore/_akg/gpu/equal.py
  37. +0
    -41
      mindspore/_akg/gpu/greater_equal.py
  38. +0
    -63
      mindspore/_akg/gpu/hsigmoid.py
  39. +0
    -51
      mindspore/_akg/gpu/hsigmoid_grad.py
  40. +0
    -63
      mindspore/_akg/gpu/hswish.py
  41. +0
    -53
      mindspore/_akg/gpu/hswish_grad.py
  42. +0
    -40
      mindspore/_akg/gpu/less_equal.py
  43. +0
    -40
      mindspore/_akg/gpu/logical_and.py
  44. +0
    -40
      mindspore/_akg/gpu/logical_not.py
  45. +0
    -40
      mindspore/_akg/gpu/logical_or.py
  46. +0
    -80
      mindspore/_akg/gpu/mean.py
  47. +0
    -90
      mindspore/_akg/gpu/mean_grad.py
  48. +0
    -41
      mindspore/_akg/gpu/mul.py
  49. +0
    -41
      mindspore/_akg/gpu/notequal.py
  50. +0
    -54
      mindspore/_akg/gpu/relu6.py
  51. +0
    -59
      mindspore/_akg/gpu/relu6_grad.py
  52. +0
    -50
      mindspore/_akg/gpu/squeeze.py
  53. +0
    -44
      mindspore/_akg/gpu/squeeze_grad.py
  54. +0
    -40
      mindspore/_akg/gpu/sub.py
  55. +0
    -39
      mindspore/_akg/gpu/tile.py
  56. +0
    -104
      mindspore/_akg/message.py
  57. +0
    -69
      mindspore/_akg/op_build.py
  58. +0
    -36
      mindspore/_akg/ops/array/tile.py
  59. +0
    -36
      mindspore/_akg/ops/math/cast.py
  60. +0
    -54
      mindspore/_akg/ops/math/equal.py
  61. +0
    -54
      mindspore/_akg/ops/math/greater_equal.py
  62. +0
    -54
      mindspore/_akg/ops/math/less_equal.py
  63. +0
    -41
      mindspore/_akg/ops/math/logical_and.py
  64. +0
    -32
      mindspore/_akg/ops/math/logical_not.py
  65. +0
    -41
      mindspore/_akg/ops/math/logical_or.py
  66. +0
    -47
      mindspore/_akg/ops/math/mean.py
  67. +0
    -43
      mindspore/_akg/ops/math/mul.py
  68. +0
    -54
      mindspore/_akg/ops/math/notequal.py
  69. +0
    -40
      mindspore/_akg/ops/math/sub.py
  70. +0
    -45
      mindspore/_akg/ops/math/sum_value.py
  71. +0
    -87
      mindspore/_akg/save_gpu_param.py
  72. +0
    -122
      mindspore/_akg/utils/dsl_create.py
  73. +0
    -80
      mindspore/_akg/utils/format_transform.py
  74. +0
    -233
      mindspore/_akg/utils/validation_check.py
  75. +1
    -1
      mindspore/_checkparam.py
  76. +19
    -0
      mindspore/_extends/builtin_operations.py
  77. +5
    -0
      mindspore/_extends/parallel_compile/akg_compiler/__init__.py
  78. +88
    -0
      mindspore/_extends/parallel_compile/akg_compiler/akg_process.py
  79. +0
    -71
      mindspore/_extends/parallel_compile/akg_compiler/multi_process_compiler.py
  80. +5
    -6
      mindspore/_extends/parallel_compile/tbe_compiler/tbe_process.py
  81. +2
    -2
      mindspore/_extends/parse/__init__.py
  82. +8
    -1
      mindspore/_extends/parse/namespace.py
  83. +27
    -9
      mindspore/_extends/parse/parser.py
  84. +4
    -2
      mindspore/_extends/parse/resources.py
  85. +57
    -7
      mindspore/_extends/parse/standard_method.py
  86. +2
    -2
      mindspore/_extends/parse/trope.py
  87. +19
    -0
      mindspore/_extends/remote/__init__.py
  88. +174
    -0
      mindspore/_extends/remote/kernel_build_server.py
  89. +148
    -0
      mindspore/_extends/remote/kernel_build_server_ascend.py
  90. +63
    -0
      mindspore/_extends/remote/kernel_build_server_gpu.py
  91. +19
    -22
      mindspore/ccsrc/CMakeLists.txt
  92. +12
    -8
      mindspore/ccsrc/backend/kernel_compiler/CMakeLists.txt
  93. +1
    -1
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_build.cc
  94. +3
    -3
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_build.h
  95. +3
    -3
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_metadata.h
  96. +1
    -1
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_mod.cc
  97. +3
    -3
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_mod.h
  98. +3
    -3
      mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_util.h
  99. +2
    -9
      mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_attrs_process.cc
  100. +3
    -3
      mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_attrs_process.h

+ 18
- 2
.gitignore View File

@@ -4,6 +4,20 @@ mindspore/lib
output
*.ir

# flatbuffer
mindspore/lite/tools/converter/parser/tflite/schema_generated.h
mindspore/lite/tools/converter/parser/caffe/caffe.pb.cc
mindspore/lite/tools/converter/parser/caffe/caffe.pb.h
mindspore/lite/tools/converter/parser/onnx/onnx.pb.h
mindspore/lite/tools/converter/parser/onnx/onnx.pb.h
mindspore/lite/tools/converter/schema/*.h
mindspore/lite/tools/converter/schema/inner
mindspore/lite/schema/*.h
mindspore/lite/schema/inner

mindspore/lite/src/runtime/kernel/opencl/cl/fp16/*.inc
mindspore/lite/src/runtime/kernel/opencl/cl/fp32/*.inc

# Cmake files
CMakeFiles/
cmake_install.cmake
@@ -27,6 +41,7 @@ cmake-build-debug
*.pb.h
*.pb.cc
*.pb
*_grpc.py

# Object files
*.o
@@ -71,5 +86,6 @@ test_temp_summary_event_file/
mindspore/version.py
mindspore/default_config.py
mindspore/.commit_id
onnx.proto
mindspore/ccsrc/onnx.proto

# lite test file
mindspore/lite/test/do_test/

+ 18
- 0
.gitmodules View File

@@ -1,6 +1,7 @@
[submodule "third_party/flatbuffers"]
path = third_party/flatbuffers
url = https://github.com/google/flatbuffers.git
ignore = all
[submodule "third_party/googletest"]
path = third_party/googletest
url = https://github.com/google/googletest.git
@@ -10,9 +11,26 @@
[submodule "third_party/protobuf"]
path = third_party/protobuf
url = https://github.com/protocolbuffers/protobuf.git
ignore = all
[submodule "akg"]
path = akg
url = https://gitee.com/mindspore/akg.git
[submodule "graphengine"]
path = graphengine
url = https://gitee.com/ms-incubator/graphengine.git
[submodule "third_party/OpenCL-CLHPP"]
path = third_party/OpenCL-CLHPP
url = https://github.com/KhronosGroup/OpenCL-CLHPP.git
[submodule "third_party/OpenCL-Headers"]
path = third_party/OpenCL-Headers
url = https://github.com/KhronosGroup/OpenCL-Headers.git
[submodule "third_party/opencv"]
path = third_party/opencv
url = https://github.com/opencv/opencv.git
[submodule "third_party/eigen"]
path = third_party/eigen
url = https://gitlab.com/libeigen/eigen.git
[submodule "third_party/libjpeg-turbo"]
path = third_party/libjpeg-turbo
url = https://github.com/libjpeg-turbo/libjpeg-turbo.git
ignore = dirty

+ 9
- 2
CMakeLists.txt View File

@@ -21,7 +21,7 @@ if (ENABLE_PYTHON)
add_compile_definitions(ENABLE_PYTHON)
endif()

set(CMAKE_CXX_FLAGS_DEBUG "$ENV{CXXFLAGS} -O0 -g2 -ggdb -fno-inline-functions -fno-omit-frame-pointer -Wl,--allow-shlib-undefined -D_LIBCPP_INLINE_VISIBILITY='' -D'_LIBCPP_EXTERN_TEMPLATE(...)=' -DHALF_ENABLE_CPP11_USER_LITERALS=0 -D_FORTIFY_SOURCE=2 -Wno-cpp")
set(CMAKE_CXX_FLAGS_DEBUG "$ENV{CXXFLAGS} -O0 -g2 -ggdb -fno-inline-functions -fno-omit-frame-pointer -Wl,--allow-shlib-undefined -D_LIBCPP_INLINE_VISIBILITY='' -D_LIBCPP_DISABLE_EXTERN_TEMPLATE=1 -DHALF_ENABLE_CPP11_USER_LITERALS=0 -D_FORTIFY_SOURCE=2 -Wno-cpp")

set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -I/usr/local/include -std=c++17 -Werror -Wall -Wno-deprecated-declarations -fPIC")
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
@@ -42,7 +42,7 @@ if (NOT Patch_FOUND)
endif ()
message(PATCH_EXECUTABLE = ${Patch_EXECUTABLE})

if (ENABLE_AKG AND ENABLE_D)
if (ENABLE_AKG AND (ENABLE_D OR ENABLE_GPU))
add_subdirectory("${CMAKE_SOURCE_DIR}/akg")
endif()

@@ -51,6 +51,8 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR})
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/flatbuffers/include)
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/flatbuffers/include/flatbuffers)

if (NOT ENABLE_ACL)

include(${CMAKE_SOURCE_DIR}/cmake/dependency_utils.cmake)
find_package(Python3 3.7 COMPONENTS Interpreter Development)
if(Python3_FOUND)
@@ -100,8 +102,13 @@ if (ENABLE_TESTCASES)
add_subdirectory(tests)
endif()

endif() # NOT ENABLE_ACL

if (ENABLE_SERVING)
add_subdirectory(serving)
add_subdirectory(serving/example/cpp_client)
endif()

if (NOT ENABLE_ACL)
include(cmake/package.cmake)
endif() # NOT ENABLE_ACL

+ 8
- 8
README.md View File

@@ -75,7 +75,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.5.0-beta/MindSpore/cpu/ubuntu_x86/mindspore-0.5.0-cp37-cp37m-linux_x86_64.whl
pip install https://ms-release.obs.cn-north-4.myhuaweicloud.com/0.6.0-beta/MindSpore/cpu/ubuntu_x86/mindspore-0.6.0-cp37-cp37m-linux_x86_64.whl
```

2. Run the following command to verify the install.
@@ -132,8 +132,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.5.0-beta
docker run -it mindspore/mindspore-cpu:0.5.0-beta /bin/bash
docker pull mindspore/mindspore-cpu:0.6.0-beta
docker run -it mindspore/mindspore-cpu:0.6.0-beta /bin/bash
```

* GPU
@@ -150,8 +150,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.5.0-beta
docker run -it --runtime=nvidia --privileged=true mindspore/mindspore-gpu:0.5.0-beta /bin/bash
docker pull mindspore/mindspore-gpu:0.6.0-beta
docker run -it --runtime=nvidia --privileged=true mindspore/mindspore-gpu:0.6.0-beta /bin/bash
```

To test if the docker image works, please execute the python code below and check the output:
@@ -202,10 +202,10 @@ Check out how MindSpore Open Governance [works](https://gitee.com/mindspore/comm

### Communication

- [MindSpore Slack](https://join.slack.com/t/mindspore/shared_invite/enQtOTcwMTIxMDI3NjM0LTNkMWM2MzI5NjIyZWU5ZWQ5M2EwMTQ5MWNiYzMxOGM4OWFhZjI4M2E5OGI2YTg3ODU1ODE2Njg1MThiNWI3YmQ) - Communication platform for developers.
- [MindSpore Slack](https://join.slack.com/t/mindspore/shared_invite/zt-dgk65rli-3ex4xvS4wHX7UDmsQmfu8w) - Communication platform for developers.
- IRC channel at `#mindspore` (only for meeting minutes logging purpose)
- Video Conferencing: https://meet.jit.si
- Mailing-list: https://mailweb.mindspore.cn/postorius/lists
- Video Conferencing: TBD
- Mailing-list: <https://mailweb.mindspore.cn/postorius/lists>

## Contributing



+ 71
- 0
RELEASE.md View File

@@ -1,3 +1,74 @@
# Release 0.6.0-beta

## Major Features and Improvements
### Ascend 910 Training and Inference Framework
* New models
* There are official, research and community under modelzoo.
* Official is maintained with the newest APIs by MindSpore team, MaskRCNN are added.
* Research is uploaded by researchers for official review, and APIs may not be updated in time.
* Community reprints the relevant links of partner research results.
* Hub added on the same level as modelzoo, synchronous storage of materials needed for official hub web pages which will be launched soon.
* Support pre-trained models, few lines of code can be used to download and load pre-trained models, supporting inference or transfer learning.
* Frontend and user interface
* Supports user side operator compilation and graph execution error rendering.
* Uniform definition dynamic learning rate behavior in optimizers.
* Support IndexSlice in sparse expression.
* Support use parent construct method during construct.
* Support asynchronous execution save checkpoint file.
* Support implicit type conversion in pynative mode.
* User interfaces change log
* unform learning rate behavior in optimizers([!2755](https://gitee.com/mindspore/mindspore/pulls/2755))
* rename operator of sparse optimizer([!3217](https://gitee.com/mindspore/mindspore/pulls/3217))
* move profiler module from mindinsight to mindspore([!3075](https://gitee.com/mindspore/mindspore/pulls/3075))
* VOCDataset output change to multi-columns([!3093](https://gitee.com/mindspore/mindspore/pulls/3093))
* GetDatasize feature([!3212](https://gitee.com/mindspore/mindspore/pulls/3212))
* dataset: modify config api([!2936](https://gitee.com/mindspore/mindspore/pulls/2936))
* Executor and performance optimization
* Decouple C++ and python, so make the architecture more extensible.
* Parameter Server for distributed deep learning supported.
* Serving:a flexible service deployment framework for deep learning models.
* Memory reuse is enhanced, and the batch size of Bert large model is increased from 96 to 160 on a single server.
* Data processing, augmentation, and save format
* Support MindRecord save operator after date processing
* Support automatic fusion operator, such as decode/resize/crop
* Support CSV dataset loading
### Other Hardware Support
* GPU platform
* New model supported: ResNext50, WarpCTC and GoogLeNet.
* Support hyperparametric search and data enhanced automl on GPU.
* Support Resnet50 automatic parallel in GPU backend.

## Bugfixes
* Models
* Improved the performance and accuracy on ResNet50([!3456](https://gitee.com/mindspore/mindspore/pulls/3456))
* Fixed the performance test case of bert([!3486](https://gitee.com/mindspore/mindspore/pulls/3486))
* Python API
* Fix assign used in while loop([!2720](https://gitee.com/mindspore/mindspore/pulls/2720))
* Revert optimize the graph output of all nop node.([!2857](https://gitee.com/mindspore/mindspore/pulls/2857))
* Print tensor as numpy.([!2859](https://gitee.com/mindspore/mindspore/pulls/2859))
* Support weight decay for sparse optimizer([!2668](https://gitee.com/mindspore/mindspore/pulls/2668))
* Fix BatchToSpaceND([!2741](https://gitee.com/mindspore/mindspore/pulls/2741))
* Fixing type check mistakes of InplaceAdd and Inplace Sub ops([!2744](https://gitee.com/mindspore/mindspore/pulls/2744]))
* Change order param only equal to group param([!2748](https://gitee.com/mindspore/mindspore/pulls/2748))
* Executor
* The performance of graph whith control flow is optimized([!2931](https://gitee.com/mindspore/mindspore/pulls/2931))
* Fix bug of wrong number of tuple layers([!3390](https://gitee.com/mindspore/mindspore/pulls/3390))
* Fix cpu multi graph memory exception([!3631](https://gitee.com/mindspore/mindspore/pulls/3631))
* Enable data sync when calling operator without defining a cell([!3081](https://gitee.com/mindspore/mindspore/pulls/3081))
* Fix argmaxwith value error in pynative mode on GPU([!3082](https://gitee.com/mindspore/mindspore/pulls/3082))
* Fix precision error with fp16 input on pynative mode([!3196](https://gitee.com/mindspore/mindspore/pulls/3196))
* 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))

## Contributors
Thanks goes to these wonderful people:

Alexey Shevlyakov, avakh, baihuawei, BowenK, buxue, caifubi, caojian05, Cathy Wong, changzherui, chenfei, chengxianbin, chenhaozhe, chenjianping, chentingting, chenzomi, chujinjin, Danish Farid, dayschan, dengwentao, dinghao, etone-chan, fangzehua, fary86, geekun, Giancarlo Colmenares, gong chen, gukecai, guohongzilong, hangangqiang, heleiwang, hesham, He Wei, hexia, hongxing, huangdongrun, huanghui, islam_amin, Jamie Nisbet, Jesse Lee, jiangjinsheng, jiangzhiwen, jinyaohui, jjfeing, jojobugfree, Jonathan Yan, jonyguo, Junhan Hu, Kang, kingfo, kouzhenzhong, kpy, kswang, laiyongqiang, leopz, liangzelang, lichenever, lihongkang, Li Hongzhang, lilei, limingqi107, lirongzhen1, liubuyu, liuchongming74, liuwenhao4, liuxiao, Lixia Chen, liyanliu, liyong, lizhenyu, lvliang, Mahdi, Margaret_wangrui, meixiaowei, ms_yan, nhussain, ougongchang, panfengfeng, panyifeng, peilinwang, Peilin Wang, pkuliuliu, qianlong, rick_sanchez, shibeiji, Shida He, shijianning, simson, sunsuodong, suteng, Tinazhang, Tron Zhang, unknown, VectorSL, wandongdong, wangcong, wangdongxu, wangdongxu6, wanghua, wangnan39, Wei Luning, wenchunjiang, wenkai, wilfChen, WilliamLian, wukesong, Xian Weizhao, Xiaoda Zhang, xiefangqi, xulei2020, xunxue, xutianchun, Yang, yanghaitao, yanghaitao1, yanghaoran, yangjie, yangjie159, YangLuo, Yanjun Peng, yankai, yanzhenxiang2020, yao_yf, Yi Huaijie, yoonlee666, yuchaojie, yujianfeng, zhangzhongpeng, zhangdengcheng, Zhang Qinghua, zhangyinxia, zhangz0911gm, zhaojichen, zhaoting, zhaozhenlong, zhoufeng, zhouneng, zhousiyi, Zirui Wu, Ziyan, zjun, ZPaC, lihongzhang, wangdongxu

Contributions of any kind are welcome!


# Release 0.5.0-beta

## Major Features and Improvements


+ 1
- 1
akg

@@ -1 +1 @@
Subproject commit df57a6cf9450e347d1854687d1fe66a420ee3b35
Subproject commit 5fe7e5c8377dccfd35c9f661e10ed3dc136208c5

+ 329
- 128
build.sh View File

@@ -24,8 +24,9 @@ usage()
{
echo "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] [-Q on|off] [-S 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] [-K] [-B on|off] [-E] [-l on|off]"
echo " [-a on|off] [-Q 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]"
echo ""
echo "Options:"
echo " -d Debug mode"
@@ -48,13 +49,14 @@ usage()
echo " -P Enable dump anf graph to file in ProtoBuffer format, default on"
echo " -Q Enable dump memory, default off"
echo " -D Enable dumping of function graph ir, default on"
echo " -S Enable async data dump, default off"
echo " -z Compile dataset & mindrecord, default on"
echo " -n Compile minddata lite"
echo " -M Enable MPI and NCCL for GPU training, gpu default on"
echo " -V Specify the minimum required cuda version, default CUDA 10.1"
echo " -I Compile predict, default off"
echo " -I Compile lite"
echo " -K Compile with AKG, default on"
echo " -s Enable serving module, default off"
echo " -w Enable acl module, default off"
echo " -B Enable debugger, default off"
echo " -E Enable IBVERBS for parameter server, default off"
echo " -l Compile with python dependency, default on"
@@ -89,28 +91,34 @@ checkopts()
ENABLE_TIMELINE="off"
ENABLE_DUMP2PROTO="on"
ENABLE_DUMPE2E="off"
ENABLE_DATA_DUMP="off"
ENABLE_DUMP_IR="on"
COMPILE_MINDDATA="on"
COMPILE_MINDDATA_LITE="off"
ENABLE_MPI="off"
CUDA_VERSION="10.1"
COMPILE_PREDICT="off"
COMPILE_LITE="off"
LITE_PLATFORM=""
SUPPORT_TRAIN="off"
USE_GLOG="on"
PREDICT_PLATFORM=""
ENABLE_AKG="on"
ENABLE_SERVING="off"
ENABLE_ACL="off"
ENABLE_DEBUGGER="off"
ENABLE_IBVERBS="off"
ENABLE_PYTHON="on"
ENABLE_GPU="off"

# Process the options
while getopts 'drvj:c:t:hsb:a:g:p:ie:m:l:I:LRP:Q:S:D:zM:V:K:sB:E' opt
while getopts 'drvj:c:t:hsb:a:g:p:ie:m:l:I:LRP:Q:D:zM:V:K:swB:En' opt
do
OPTARG=$(echo ${OPTARG} | tr '[A-Z]' '[a-z]')
case "${opt}" in
d)
DEBUG_MODE="on"
;;
n)
COMPILE_MINDDATA_LITE="on"
;;
r)
DEBUG_MODE="off"
;;
@@ -186,6 +194,7 @@ checkopts()
elif [[ "X$OPTARG" == "Xd" || "X$OPTARG" == "Xascend" ]]; then
ENABLE_D="on"
ENABLE_CPU="on"
ENABLE_SERVING="on"
elif [[ "X$OPTARG" == "Xcpu" ]]; then
ENABLE_CPU="on"
else
@@ -220,11 +229,6 @@ checkopts()
ENABLE_DUMPE2E="$OPTARG"
echo "enable dump end to end"
;;
S)
check_on_off $OPTARG S
ENABLE_DATA_DUMP="$OPTARG"
echo "enable data dump"
;;
D)
check_on_off $OPTARG D
ENABLE_DUMP_IR="$OPTARG"
@@ -244,13 +248,16 @@ checkopts()
fi
;;
I)
COMPILE_PREDICT="on"
COMPILE_LITE="on"
if [[ "$OPTARG" == "arm64" ]]; then
PREDICT_PLATFORM="arm64"
LITE_PLATFORM="arm64"
elif [[ "$OPTARG" == "arm32" ]]; then
LITE_PLATFORM="arm32"
elif [[ "$OPTARG" == "x86_64" ]]; then
PREDICT_PLATFORM="x86_64"
ENABLE_CONVERTER="on"
LITE_PLATFORM="x86_64"
else
echo "-I parameter must be arm64 or x86_64"
echo "-I parameter must be arm64、arm32 or x86_64"
exit 1
fi
;;
@@ -262,6 +269,10 @@ checkopts()
ENABLE_SERVING="on"
echo "enable serving"
;;
w)
ENABLE_ACL="on"
echo "enable acl"
;;
B)
check_on_off $OPTARG B
ENABLE_DEBUGGER="on"
@@ -279,10 +290,13 @@ checkopts()
done
}
checkopts "$@"
if [[ "X$ENABLE_GPU" = "Xon" ]] && [[ "X$ENABLE_DUMPE2E" = "Xon" ]]; then
ENABLE_DEBUGGER="on"
fi
echo "---------------- MindSpore: build start ----------------"
mkdir -pv "${BUILD_PATH}/package/mindspore/lib"
git submodule update --init graphengine
if [[ "X$ENABLE_AKG" = "Xon" ]] && [[ "X$ENABLE_D" = "Xon" ]]; then
if [[ "X$ENABLE_AKG" = "Xon" ]] && [[ "X$ENABLE_D" = "Xon" || "X$ENABLE_GPU" = "Xon" ]]; then
git submodule update --init --recursive akg
fi

@@ -328,9 +342,6 @@ build_mindspore()
if [[ "X$ENABLE_DUMPE2E" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_DUMP_E2E=ON"
fi
if [[ "X$ENABLE_DATA_DUMP" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_DATA_DUMP=ON"
fi
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_DUMP_IR=${ENABLE_DUMP_IR}"
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_PYTHON=${ENABLE_PYTHON}"
if [[ "X$ENABLE_MPI" = "Xon" ]]; then
@@ -340,7 +351,7 @@ build_mindspore()
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_D=ON"
fi
if [[ "X$ENABLE_GPU" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_GPU=ON -DCUDA_PATH=$CUDA_PATH -DCUDNN_PATH=$CUDNN_PATH -DMS_REQUIRE_CUDA_VERSION=${CUDA_VERSION}"
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_GPU=ON -DUSE_CUDA=ON -DCUDA_PATH=$CUDA_PATH -DCUDNN_PATH=$CUDNN_PATH -DMS_REQUIRE_CUDA_VERSION=${CUDA_VERSION}"
fi
if [[ "X$ENABLE_CPU" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_CPU=ON"
@@ -351,12 +362,15 @@ build_mindspore()
if [[ "X$USE_GLOG" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DUSE_GLOG=ON"
fi
if [[ "X$ENABLE_AKG" = "Xon" ]] && [[ "X$ENABLE_D" = "Xon" ]]; then
if [[ "X$ENABLE_AKG" = "Xon" ]] && [[ "X$ENABLE_D" = "Xon" || "X$ENABLE_GPU" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_AKG=ON"
fi
if [[ "X$ENABLE_SERVING" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_SERVING=ON"
fi
if [[ "X$ENABLE_ACL" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_ACL=ON"
fi
if [[ "X$ENABLE_DEBUGGER" = "Xon" ]]; then
CMAKE_ARGS="${CMAKE_ARGS} -DENABLE_DEBUGGER=ON"
fi
@@ -371,132 +385,319 @@ build_mindspore()
if [[ -n "$VERBOSE" ]]; then
CMAKE_VERBOSE="--verbose"
fi
if [[ "X$ENABLE_ACL" = "Xon" ]]; then
cmake --build . ${CMAKE_VERBOSE} -j$THREAD_NUM
else
cmake --build . --target package ${CMAKE_VERBOSE} -j$THREAD_NUM
fi
echo "success to build mindspore project!"
}

build_predict()
{
git submodule update --init --recursive third_party/incubator-tvm
echo "start build predict project"

git submodule update --init --recursive third_party/flatbuffers
git submodule update --init --recursive third_party/googletest
git submodule update --init --recursive third_party/protobuf

rm -rf "${BASEPATH}/predict/build"
mkdir -pv "${BASEPATH}/predict/build"
rm -rf "${BASEPATH}/predict/output"
mkdir -pv "${BASEPATH}/predict/output"

if [[ "$PREDICT_PLATFORM" == "arm64" ]]; then
if [ "${ANDROID_NDK}" ]; then
echo -e "\e[31mANDROID_NDK_PATH=$ANDROID_NDK \e[0m"
else
echo -e "\e[31mplease set ANDROID_NDK_PATH in environment variable for example: export ANDROID_NDK=/root/usr/android-ndk-r16b/ \e[0m"
exit 1
fi
checkndk() {
if [ "${ANDROID_NDK}" ]; then
echo -e "\e[31mANDROID_NDK_PATH=$ANDROID_NDK \e[0m"
else
echo -e "\e[31mplease set ANDROID_NDK_PATH in environment variable for example: export ANDROID_NDK=/root/usr/android-ndk-r20b/ \e[0m"
exit 1
fi
}

#build flatbuf
cd "${BASEPATH}/third_party/flatbuffers"
rm -rf build && mkdir -p build && cd build && cmake .. && make -j$THREAD_NUM
FLATC="${BASEPATH}"/third_party/flatbuffers/build/flatc
cd "${BASEPATH}"/predict/schema && mkdir -p "${BASEPATH}"/predict/schema/inner
gene_flatbuffer() {
FLAT_DIR="${BASEPATH}/mindspore/lite/schema"
cd ${FLAT_DIR} && rm -rf "${FLAT_DIR}/inner" && mkdir -p "${FLAT_DIR}/inner"
find . -name "*.fbs" -print0 | xargs -0 "${FLATC}" -c -b
find . -name "*.fbs" -print0 | xargs -0 "${FLATC}" -c -b --reflect-types --gen-mutable --reflect-names --gen-object-api -o ${BASEPATH}/predict/schema/inner

# check LLVM_PATH
if [ "${LLVM_PATH}" == "" ]; then
echo "Please set LLVM_PATH in env for example export LLVM_PATH=/xxxx/bin/llvm-config"
exit
fi

#build tvm
tvm_open_source="${BASEPATH}/third_party/incubator-tvm"
tvm_kernel_build="${BASEPATH}/predict/module/tvm_kernel"
if [ ! -f "${tvm_kernel_build}"/incubator-tvm/build/libtvm.so ]; then
rm -fr "${tvm_kernel_build}"/incubator-tvm
cp -fr "${tvm_open_source}" "${tvm_kernel_build}"
mkdir -p "${tvm_kernel_build}"/incubator-tvm/build
patch -d "${tvm_kernel_build}"/incubator-tvm -p1 < "${BASEPATH}"/third_party/patch/predict/0001-RetBugFix-CustomRuntime_v06.patch
cp "${tvm_kernel_build}"/lite/src/codegen/llvm/lite_rtfunc_reset.cc "${tvm_kernel_build}"/incubator-tvm/src/codegen/llvm/
cp "${tvm_open_source}"/cmake/config.cmake "${tvm_kernel_build}"/incubator-tvm
if [ "${LLVM_PATH}" ]; then
sed -i "s#set(USE_LLVM .*)#set(USE_LLVM \"${LLVM_PATH}\")#g" "${tvm_kernel_build}"/incubator-tvm/config.cmake
else
echo "need set LLVM_PATH in env for example export LLVM_PATH=/xxxx/bin/llvm-config"
find . -name "*.fbs" -print0 | xargs -0 "${FLATC}" -c -b --reflect-types --gen-mutable --reflect-names --gen-object-api -o "${FLAT_DIR}/inner"

FLAT_DIR="${BASEPATH}/mindspore/lite/tools/converter/parser/tflite"
cd ${FLAT_DIR}
find . -name "*.fbs" -print0 | xargs -0 "${FLATC}" -c -b --reflect-types --gen-mutable --reflect-names --gen-object-api -o "${FLAT_DIR}/"
}

build_flatbuffer() {
cd ${BASEPATH}
FLATC="${BASEPATH}"/third_party/flatbuffers/build/flatc
if [[ ! -f "${FLATC}" ]]; then
git submodule update --init --recursive third_party/flatbuffers
cd ${BASEPATH}/third_party/flatbuffers
rm -rf build && mkdir -pv build && cd build && cmake .. && make -j$THREAD_NUM
gene_flatbuffer
fi
if [[ "${INC_BUILD}" == "off" ]]; then
gene_flatbuffer
fi
}

gene_protobuf() {
PROTO_SRC_DIR="${BASEPATH}/mindspore/lite/tools/converter/parser/caffe"
find ${PROTO_SRC_DIR} -name "*.proto" -print0 | xargs -0 "${PROTOC}" -I"${PROTO_SRC_DIR}" --cpp_out="${PROTO_SRC_DIR}"
PROTO_SRC_DIR="${BASEPATH}/mindspore/lite/tools/converter/parser/onnx"
find ${PROTO_SRC_DIR} -name "*.proto" -print0 | xargs -0 "${PROTOC}" -I"${PROTO_SRC_DIR}" --cpp_out="${PROTO_SRC_DIR}"
}

build_protobuf() {
cd ${BASEPATH}
PROTOC="${BASEPATH}"/third_party/protobuf/build/bin/protoc
if [[ ! -f "${PROTOC}" ]]; then
git submodule update --init --recursive third_party/protobuf
cd ${BASEPATH}/third_party/protobuf
rm -rf build && mkdir -pv build && ./autogen.sh
./configure --prefix=${BASEPATH}/third_party/protobuf/build
make clean && make -j$THREAD_NUM && make install
gene_protobuf
fi
if [[ "${INC_BUILD}" == "off" ]]; then
gene_protobuf
fi
}

build_gtest() {
cd ${BASEPATH}
git submodule update --init --recursive third_party/googletest
}

gene_clhpp() {
CL_SRC_DIR="${BASEPATH}/mindspore/lite/src/runtime/kernel/opencl/cl"
for sub_dir in "${CL_SRC_DIR}"/*
do
data_type="$(basename ${sub_dir})"
if [ ! -d ${CL_SRC_DIR}/${data_type} ]; then
continue
fi
cd "${tvm_kernel_build}"/incubator-tvm/build
cmake ..
make -j$THREAD_NUM
else
cd "${tvm_kernel_build}"/incubator-tvm/build
make -j$THREAD_NUM
fi

#gen op
predict_tvm_op_lib_path="${BASEPATH}/predict/module/tvm_kernel/build/lib_x86"
predict_platform="x86"
if [[ "$PREDICT_PLATFORM" == "arm64" ]]; then
predict_tvm_op_lib_path="${BASEPATH}/predict/module/tvm_kernel/build/lib_arm64"
predict_platform="arm64"
fi

need_get_libs=true
if [ -d "${predict_tvm_op_lib_path}" ]; then
file_list=$(ls "${predict_tvm_op_lib_path}")
if [ -n "${file_list}" ]; then
libstime=$(stat -c %Y "${predict_tvm_op_lib_path}"/* | sort -u | tail -n1)
pythontime=$(find "${BASEPATH}"/predict/module/tvm_kernel/lite/python/ -name "*.py" -exec stat -c %Y {} \; |
sort -u | tail -n1)
if [ "${libstime}" -ge "${pythontime}" ]; then
need_get_libs=false
else
rm -fr "${predict_tvm_op_lib_path}"
cd ${CL_SRC_DIR}/${data_type}
rm -rf *.inc
echo "$(cd "$(dirname $0)"; pwd)"
for file_path in "${CL_SRC_DIR}/${data_type}"/*
do
file="$(basename ${file_path})"
inc_file=`echo ${CL_SRC_DIR}/${data_type}/${file} | sed 's/$/.inc/'`
sed 's/^/\"/;s/$/ \\n\" \\/' ${CL_SRC_DIR}/${data_type}/${file} > ${inc_file}
kernel_name=`echo ${file} | sed s'/.\{3\}$//'`
sed -i "1i\static const char *${kernel_name}_source_${data_type} =\"\\n\" \\" ${inc_file}
sed -i '$a\;' ${inc_file}
done
done
}

gene_ocl_program() {
CL_SRC_DIR="${BASEPATH}/mindspore/lite/src/runtime/kernel/opencl/cl"
SPIRV_DIR=build/spirv
rm -rf ${SPIRV_DIR}
mkdir -pv ${SPIRV_DIR}
for sub_dir in "${CL_SRC_DIR}"/*
do
data_type="$(basename ${sub_dir})"
if [ ! -d ${CL_SRC_DIR}/${data_type} ]; then
continue
fi
fi
#echo $(cd "$(dirname $0)"; pwd)
for file_path in "${CL_SRC_DIR}/${data_type}"/*
do
file="$(basename ${file_path})"
if [ "${file##*.}" != "cl" ]; then
continue
fi
clang -Xclang -finclude-default-header -cl-std=CL2.0 --target=spir64-unknown-unknown -emit-llvm \
-c -O0 -o ${SPIRV_DIR}/${file%.*}.bc ${CL_SRC_DIR}/${data_type}/${file}
done
done

bcs=`ls ${SPIRV_DIR}/*.bc`
llvm-link ${bcs} -o ${SPIRV_DIR}/program.bc
llvm-spirv -o ${SPIRV_DIR}/program.spv ${SPIRV_DIR}/program.bc

CL_PROGRAM_PATH="${BASEPATH}/mindspore/lite/src/runtime/kernel/opencl/cl/program.inc"
echo "#include <vector>" > ${CL_PROGRAM_PATH}
echo "std::vector<unsigned char> g_program_binary = {" >> ${CL_PROGRAM_PATH}
#hexdump -v -e '16/1 "0x%02x, " "\n"' ${SPIRV_DIR}/program.spv >> ${CL_PROGRAM_PATH}
hexdump -v -e '1/1 "0x%02x, "' ${SPIRV_DIR}/program.spv >> ${CL_PROGRAM_PATH}
echo "};" >> ${CL_PROGRAM_PATH}
echo "Compile SPIRV done"
}

build_opencl() {
cd ${BASEPATH}
git submodule update --init third_party/OpenCL-Headers
git submodule update --init third_party/OpenCL-CLHPP
if [[ "${OPENCL_OFFLINE_COMPILE}" == "on" ]]; then
gene_ocl_program
else
gene_clhpp
fi
}

build_opencv() {
cd ${BASEPATH}
if [[ "${INC_BUILD}" == "off" ]]; then
git submodule update --init --recursive third_party/opencv
cd ${BASEPATH}/third_party/opencv
rm -rf build && mkdir -p build && cd build && cmake ${CMAKE_MINDDATA_ARGS} -DBUILD_SHARED_LIBS=ON -DBUILD_ANDROID_PROJECTS=OFF \
-DBUILD_LIST=core,imgcodecs,imgproc -DBUILD_ZLIB=ON .. && make -j$THREAD_NUM
fi
}

build_jpeg_turbo() {
cd ${BASEPATH}
if [[ "${INC_BUILD}" == "off" ]]; then
git submodule update --init --recursive third_party/libjpeg-turbo
cd ${BASEPATH}/third_party/libjpeg-turbo
rm -rf build && mkdir -p build && cd build && cmake ${CMAKE_MINDDATA_ARGS} -DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX="${BASEPATH}/third_party/libjpeg-turbo" .. && make -j$THREAD_NUM && make install
fi
}

build_eigen() {
cd ${BASEPATH}
git submodule update --init --recursive third_party/eigen
}

build_minddata_lite_deps()
{
echo "start build minddata lite project"
if [[ "${LITE_PLATFORM}" == "arm64" ]]; then
CMAKE_MINDDATA_ARGS="-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}"
elif [[ "${LITE_PLATFORM}" == "arm32" ]]; then
CMAKE_MINDDATA_ARGS="-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}"
else
CMAKE_MINDDATA_ARGS="-DCMAKE_BUILD_TYPE=${BUILD_TYPE}"
fi
build_opencv
build_eigen
build_jpeg_turbo
}

if $need_get_libs; then
PYTHONPATH_OLD=${PYTHONPATH}
export PYTHONPATH="${tvm_kernel_build}/incubator-tvm/python:${tvm_kernel_build}/incubator-tvm/topi/python:${tvm_kernel_build}/incubator-tvm/nnvm/python:${tvm_kernel_build}/lite/python:"
cd "${BASEPATH}"/predict/module/tvm_kernel/lite/python/at_ops
python3 at_gen_strip.py ${predict_platform}
export PYTHONPATH=${PYTHONPATH_OLD}
build_lite()
{
echo "start build mindspore lite project"

if [ "${ENABLE_GPU}" == "on" ] && [ "${LITE_PLATFORM}" == "arm64" ]; then
echo "start build opencl"
build_opencl
fi
if [[ "${LITE_PLATFORM}" == "x86_64" ]]; then
build_protobuf
fi
build_flatbuffer
build_gtest

cd "${BASEPATH}/predict/build"
if [[ "$PREDICT_PLATFORM" == "arm64" ]]; then
cmake -DCMAKE_TOOLCHAIN_FILE="${ANDROID_NDK}/build/cmake/android.toolchain.cmake" \
-DANDROID_NATIVE_API_LEVEL=android-19 -DANDROID_NDK="${ANDROID_NDK}" \
-DANDROID_TOOLCHAIN_NAME="aarch64-linux-android-clang" -DANDROID_STL="c++_shared" \
-DANDROID_ABI="arm64-v8a" -DENABLE_PREDICT_ARM64=ON -DANDROID_ALLOW_UNDEFINED_SYMBOLS=TRUE ..
elif [[ "$PREDICT_PLATFORM" == "x86_64" ]]; then
cmake ..
if [ "${COMPILE_MINDDATA_LITE}" == "on" ]; then
build_minddata_lite_deps
fi

make ${VERBOSE} -j$THREAD_NUM
if [[ "$PREDICT_PLATFORM" == "x86_64" ]]; then
cd "${BASEPATH}/predict/build/test" && ./run_tests.sh
cd "${BASEPATH}/mindspore/lite"
if [[ "${INC_BUILD}" == "off" ]]; then
rm -rf build
fi
mkdir -pv build
cd build
BUILD_TYPE="Release"
if [[ "${DEBUG_MODE}" == "on" ]]; then
BUILD_TYPE="Debug"
fi

# copy securec include files
mkdir -p "${BASEPATH}/predict/output/include/securec/include"
cp "${BASEPATH}"/third_party/securec/include/* "${BASEPATH}"/predict/output/include/securec/include
if [[ "${LITE_PLATFORM}" == "arm64" ]]; then
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} \
-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} \
"${BASEPATH}/mindspore/lite"
elif [[ "${LITE_PLATFORM}" == "arm32" ]]; then
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} \
-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} \
"${BASEPATH}/mindspore/lite"
else
cmake -DBUILD_DEVICE=on -DPLATFORM_ARM64=off -DBUILD_CONVERTER=${ENABLE_CONVERTER} -DSUPPORT_TRAIN=${SUPPORT_TRAIN} \
-DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSUPPORT_GPU=${ENABLE_GPU} -DBUILD_MINDDATA=${COMPILE_MINDDATA_LITE} \
-DOFFLINE_COMPILE=${OPENCL_OFFLINE_COMPILE} "${BASEPATH}/mindspore/lite"
fi
VERBOSE=2 make -j$THREAD_NUM
COMPILE_RET=$?

cd "${BASEPATH}/predict/output/"
if [[ "$PREDICT_PLATFORM" == "x86_64" ]]; then
tar -cf MSPredict-0.5.0-linux_x86_64.tar.gz include/ lib/ --warning=no-file-changed
elif [[ "$PREDICT_PLATFORM" == "arm64" ]]; then
tar -cf MSPredict-0.5.0-linux_aarch64.tar.gz include/ lib/ --warning=no-file-changed
if [[ "${COMPILE_RET}" -ne 0 ]]; then
echo "---------------- mindspore lite: build failed ----------------"
else
mkdir -pv ${BASEPATH}/mindspore/lite/output/
if [[ "$LITE_PLATFORM" == "x86_64" ]]; then
OUTPUT_DIR=${BASEPATH}/output/MSLite-0.6.0-linux_x86_64
rm -rf ${OUTPUT_DIR} && mkdir -p ${OUTPUT_DIR} && cd ${OUTPUT_DIR}
mkdir -p ${OUTPUT_DIR}/converter && mkdir -p ${OUTPUT_DIR}/time_profile
mkdir -p ${OUTPUT_DIR}/benchmark && mkdir -p ${OUTPUT_DIR}/include && mkdir -p ${OUTPUT_DIR}/lib
mkdir -p ${OUTPUT_DIR}/third_party
cp ${BASEPATH}/mindspore/lite/build/tools/converter/converter_lite ${OUTPUT_DIR}/converter/
cp ${BASEPATH}/mindspore/lite/build/tools/benchmark/benchmark ${OUTPUT_DIR}/benchmark/
cp ${BASEPATH}/mindspore/lite/build/tools/time_profile/timeprofile ${OUTPUT_DIR}/time_profile/
cp ${BASEPATH}/mindspore/lite/include/*.h ${OUTPUT_DIR}/include/
mkdir -p ${OUTPUT_DIR}/include/ir/dtype/
cp ${BASEPATH}/mindspore/core/ir/dtype/type_id.h ${OUTPUT_DIR}/include/ir/dtype/
mkdir -p ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/schema/*.h ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/build/src/libmindspore-lite.so ${OUTPUT_DIR}/lib/
mkdir -p ${OUTPUT_DIR}/third_party/protobuf/lib
cp -r ${BASEPATH}/third_party/protobuf/build/include/ ${OUTPUT_DIR}/third_party/protobuf/
cp -r ${BASEPATH}/third_party/protobuf/build/lib/libprotobuf.so.19 ${OUTPUT_DIR}/third_party/protobuf/lib/
cp -r ${BASEPATH}/third_party/protobuf/build/lib/libprotobuf.so.19.0.0 ${OUTPUT_DIR}/third_party/protobuf/lib/
mkdir -p ${OUTPUT_DIR}/third_party/flatbuffers
cp -r ${BASEPATH}/third_party/flatbuffers/include/ ${OUTPUT_DIR}/third_party/flatbuffers/
cd ..
tar -czf MSLite-0.6.0-linux_x86_64.tar.gz MSLite-0.6.0-linux_x86_64/ --warning=no-file-changed
sha256sum MSLite-0.6.0-linux_x86_64.tar.gz > MSLite-0.6.0-linux_x86_64.tar.gz.sha256
rm -rf MSLite-0.6.0-linux_x86_64/
elif [[ "$LITE_PLATFORM" == "arm64" ]]; then
OUTPUT_DIR=${BASEPATH}/output/MSLite-0.6.0-linux_arm64
rm -rf ${OUTPUT_DIR} && mkdir -p ${OUTPUT_DIR} && cd ${OUTPUT_DIR}
mkdir -p ${OUTPUT_DIR}/time_profile && mkdir -p ${OUTPUT_DIR}/benchmark
mkdir -p ${OUTPUT_DIR}/include && mkdir -p ${OUTPUT_DIR}/lib
mkdir -p ${OUTPUT_DIR}/third_party
cp ${BASEPATH}/mindspore/lite/build/tools/benchmark/benchmark ${OUTPUT_DIR}/benchmark/
cp ${BASEPATH}/mindspore/lite/build/tools/time_profile/timeprofile ${OUTPUT_DIR}/time_profile/
cp ${BASEPATH}/mindspore/lite/include/*.h ${OUTPUT_DIR}/include/
mkdir -p ${OUTPUT_DIR}/include/ir/dtype/
cp ${BASEPATH}/mindspore/core/ir/dtype/type_id.h ${OUTPUT_DIR}/include/ir/dtype/
mkdir -p ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/schema/*.h ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/build/src/libmindspore-lite.so ${OUTPUT_DIR}/lib/
mkdir -p ${OUTPUT_DIR}/third_party/flatbuffers
cp -r ${BASEPATH}/third_party/flatbuffers/include/ ${OUTPUT_DIR}/third_party/flatbuffers/
cd ..
tar -czf MSLite-0.6.0-linux_arm64.tar.gz MSLite-0.6.0-linux_arm64/ --warning=no-file-changed
sha256sum MSLite-0.6.0-linux_arm64.tar.gz > MSLite-0.6.0-linux_arm64.tar.gz.sha256
rm -rf MSLite-0.6.0-linux_arm64/
elif [[ "$LITE_PLATFORM" == "arm32" ]]; then
OUTPUT_DIR=${BASEPATH}/output/MSLite-0.6.0-linux_arm32
rm -rf ${OUTPUT_DIR} && mkdir -p ${OUTPUT_DIR} && cd ${OUTPUT_DIR}
mkdir -p ${OUTPUT_DIR}/time_profile && mkdir -p ${OUTPUT_DIR}/benchmark
mkdir -p ${OUTPUT_DIR}/include && mkdir -p ${OUTPUT_DIR}/lib
mkdir -p ${OUTPUT_DIR}/third_party
cp ${BASEPATH}/mindspore/lite/build/tools/benchmark/benchmark ${OUTPUT_DIR}/benchmark/
cp ${BASEPATH}/mindspore/lite/build/tools/time_profile/timeprofile ${OUTPUT_DIR}/time_profile/
cp ${BASEPATH}/mindspore/lite/include/*.h ${OUTPUT_DIR}/include/
mkdir -p ${OUTPUT_DIR}/include/ir/dtype/
cp ${BASEPATH}/mindspore/core/ir/dtype/type_id.h ${OUTPUT_DIR}/include/ir/dtype/
mkdir -p ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/schema/*.h ${OUTPUT_DIR}/include/schema/
cp ${BASEPATH}/mindspore/lite/build/src/libmindspore-lite.so ${OUTPUT_DIR}/lib/
mkdir -p ${OUTPUT_DIR}/third_party/flatbuffers
cp -r ${BASEPATH}/third_party/flatbuffers/include/ ${OUTPUT_DIR}/third_party/flatbuffers/
cd ..
tar -czf MSLite-0.6.0-linux_arm32.tar.gz MSLite-0.6.0-linux_arm32/ --warning=no-file-changed
sha256sum MSLite-0.6.0-linux_arm32.tar.gz > MSLite-0.6.0-linux_arm32.tar.gz.sha256
rm -rf MSLite-0.6.0-linux_arm32/
fi
echo "---------------- mindspore lite: build success ----------------"
fi
echo "success to build predict project!"
}

if [[ "X$COMPILE_PREDICT" = "Xon" ]]; then
build_predict
echo "---------------- mindspore: build end ----------------"
if [[ "X$COMPILE_LITE" = "Xon" ]]; then
build_lite
exit
else
build_mindspore


+ 0
- 7
cmake/external_libs/dlpack.cmake View File

@@ -1,7 +0,0 @@
mindspore_add_pkg(dlpack
VER 0.2
HEAD_ONLY ./
URL https://github.com/dmlc/dlpack/archive/0acb731e0e43d15deee27b66f10e4c5b4e667913.zip
MD5 6b8093f17ad4e830d3c63eb3171c4b45)



+ 0
- 7
cmake/external_libs/dmlc_core.cmake View File

@@ -1,7 +0,0 @@
mindspore_add_pkg(dmlc-core
VER 0.3
HEAD_ONLY ./
URL https://github.com/dmlc/dmlc-core/archive/808f485387f9a03f78fa9f1159f387d0d91b7a28.zip
MD5 ea36f94c57752bf40fb02dfc362f1ed9)



+ 1
- 1
cmake/external_libs/glog.cmake View File

@@ -1,4 +1,4 @@
set(glog_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2 ${SECURE_CXX_FLAGS}")
set(glog_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2 ${SECURE_CXX_FLAGS} -D_GLIBCXX_USE_CXX11_ABI=0")
set(glog_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
mindspore_add_pkg(glog
VER 0.4.0


+ 1
- 0
cmake/external_libs/jpeg_turbo.cmake View File

@@ -12,6 +12,7 @@ mindspore_add_pkg(jpeg_turbo
URL https://github.com/libjpeg-turbo/libjpeg-turbo/archive/2.0.4.tar.gz
MD5 44c43e4a9fb352f47090804529317c88
CMAKE_OPTION -DCMAKE_BUILD_TYPE=Release -DCMAKE_SKIP_RPATH=TRUE
PATCHES ${CMAKE_SOURCE_DIR}/third_party/patch/jpeg_turbo/jpeg_turbo.patch001
)
include_directories(${jpeg_turbo_INC})
add_library(mindspore::jpeg_turbo ALIAS jpeg_turbo::jpeg)

+ 6
- 6
cmake/external_libs/mkl_dnn.cmake View File

@@ -2,18 +2,18 @@ set(onednn_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2")
set(onednn_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
if (CMAKE_SYSTEM_NAME MATCHES "Windows")
mindspore_add_pkg(onednn
VER 1.1.1
VER 1.5
LIBS dnnl mkldnn
HEAD_ONLY ./include
RELEASE on
URL https://github.com/oneapi-src/oneDNN/releases/download/v1.1.1/dnnl_win_1.1.1_cpu_vcomp.zip
MD5 ecaab9ed549643067699c80e5cea1c23)
URL https://github.com/oneapi-src/oneDNN/releases/download/v1.5/dnnl_win_1.5.0_cpu_vcomp.zip
MD5 17757c84f49edd42d34ae8c9288110a1)
else()
mindspore_add_pkg(onednn
VER 1.1.2
VER 1.5
LIBS dnnl mkldnn
URL https://github.com/oneapi-src/oneDNN/archive/v1.1.2.tar.gz
MD5 ab40d52230f3ad1d7a6f06ce0f6bc17a
URL https://github.com/oneapi-src/oneDNN/archive/v1.5.tar.gz
MD5 5d97e0e8f4c0b37da5f524533b7a644b
CMAKE_OPTION -DDNNL_ARCH_OPT_FLAGS='' -DDNNL_CPU_RUNTIME='SEQ' -DDNNL_BUILD_EXAMPLES=OFF -DDNNL_BUILD_TESTS=OFF)
endif()



+ 0
- 7
cmake/external_libs/rang.cmake View File

@@ -1,7 +0,0 @@
mindspore_add_pkg(rang
VER 3.1.0
HEAD_ONLY ./
URL https://github.com/agauniyal/rang/archive/cabe04d6d6b05356fa8f9741704924788f0dd762.zip
MD5 0c5c9b251fea9ee7ce32f188655be0ea)



+ 25
- 0
cmake/external_libs/sentencepiece.cmake View File

@@ -0,0 +1,25 @@
if (WIN32)
set(sentencepiece_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2 -Wno-unused-result -Wno-stringop-overflow -Wno-format-extra-args -Wno-format")
set(sentencepiece_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
mindspore_add_pkg(sentencepiece
VER 0.1.92
LIBS sentencepiece sentencepiece_train
URL https://github.com/google/sentencepiece/archive/v0.1.92.tar.gz
CMAKE_OPTION -DCMAKE_BUILD_TYPE=Release -DSPM_USE_BUILTIN_PROTOBUF=ON
MD5 5dfd2241914b5598a68b2a8542ed8e91
)
else ()
set(sentencepiece_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2 -Wno-unused-result -Wno-sign-compare")
set(sentencepiece_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
mindspore_add_pkg(sentencepiece
VER 0.1.92
LIBS sentencepiece sentencepiece_train
URL https://github.com/google/sentencepiece/archive/v0.1.92.tar.gz
CMAKE_OPTION -DCMAKE_BUILD_TYPE=Release -DSPM_USE_BUILTIN_PROTOBUF=OFF -DSPM_ENABLE_SHARED=OFF -DPROTOBUF_INC=${protobuf_INC}
MD5 5dfd2241914b5598a68b2a8542ed8e91
PATCHES ${CMAKE_SOURCE_DIR}/third_party/patch/sentencepiece/sentencepiece.patch001
)
endif ()
include_directories(${sentencepiece_INC})
add_library(mindspore::sentencepiece ALIAS sentencepiece::sentencepiece)
add_library(mindspore::sentencepiece_train ALIAS sentencepiece::sentencepiece_train)

+ 0
- 15
cmake/external_libs/tvm_gpu.cmake View File

@@ -1,15 +0,0 @@
set(incubator_tvm_gpu_CXXFLAGS "-D_FORTIFY_SOURCE=2 -O2")
set(incubator_tvm_gpu_CFLAGS "-D_FORTIFY_SOURCE=2 -O2")
mindspore_add_pkg(incubator_tvm_gpu
VER 0.6.0
LIBS tvm
URL https://github.com/apache/incubator-tvm/archive/v0.6.0.tar.gz
MD5 9cbbd32545a776023acabbba270449fe
CUSTOM_CMAKE ${CMAKE_SOURCE_DIR}/third_party/patch/incubator-tvm/
SUBMODULES ${dlpack_DIRPATH} ${dmlc-core_DIRPATH} ${rang_DIRPATH}
SOURCEMODULES topi/python/topi python/tvm
PATCHES ${CMAKE_SOURCE_DIR}/third_party/patch/incubator-tvm/find_library.patch
${CMAKE_SOURCE_DIR}/third_party/patch/incubator-tvm/include.patch
${CMAKE_SOURCE_DIR}/third_party/patch/incubator-tvm/src_pass.patch
CMAKE_OPTION " ")
add_library(mindspore::tvm ALIAS incubator_tvm_gpu::tvm)

+ 11
- 11
cmake/mind_expression.cmake View File

@@ -15,7 +15,7 @@ include(${CMAKE_SOURCE_DIR}/cmake/external_libs/json.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/dependency_securec.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/protobuf.cmake)

if (ENABLE_DEBUGGER OR ENABLE_SERVING)
if (ENABLE_DEBUGGER OR ENABLE_SERVING OR ENABLE_TESTCASES)
# build dependencies of gRPC
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/absl.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/c-ares.cmake)
@@ -30,7 +30,7 @@ include(${CMAKE_SOURCE_DIR}/cmake/external_libs/flatbuffers.cmake)
if(USE_GLOG)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/glog.cmake)
endif()
if (NOT ${CMAKE_SYSTEM_NAME} MATCHES "Windows" AND NOT ENABLE_GE)
if (ENABLE_CPU AND (ENABLE_D OR ENABLE_GPU))
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/zeromq.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/pslite.cmake)
endif()
@@ -38,19 +38,15 @@ endif()
find_package(Python3)
include_directories(${Python3_INCLUDE_DIRS})
include_directories(${CMAKE_SOURCE_DIR}/third_party)
if (ENABLE_MPI)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/ompi.cmake)
endif()

if (ENABLE_CPU)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/mkl_dnn.cmake)
if (ENABLE_MPI)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/ompi.cmake)
endif()
endif()

if (ENABLE_GPU)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/dlpack.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/dmlc_core.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/rang.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/tvm_gpu.cmake)

if (ENABLE_MPI)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/nccl.cmake)
endif()
@@ -69,12 +65,16 @@ endif()

if (ENABLE_MINDDATA)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/icu4c.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/jpeg_turbo.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/libtiff.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/opencv.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/sqlite.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/tinyxml2.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/cppjieba.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/sentencepiece.cmake)
endif()

if (ENABLE_MINDDATA OR ENABLE_SERVING)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/jpeg_turbo.cmake)
endif()

include(${CMAKE_SOURCE_DIR}/cmake/external_libs/gtest.cmake)


+ 4
- 4
cmake/options.cmake View File

@@ -70,6 +70,10 @@ if (ENABLE_GPU)
add_compile_definitions(ENABLE_GPU_COLLECTIVE)
endif()

if (ENABLE_CPU)
add_compile_definitions(ENABLE_CPU)
endif()

if (ENABLE_GE)
add_compile_definitions(ENABLE_GE)
add_compile_definitions(CUSTOM_OP)
@@ -116,10 +120,6 @@ if(ENABLE_DUMP_E2E)
add_compile_definitions(ENABLE_DUMP_E2E)
endif()

if(ENABLE_DATA_DUMP)
add_compile_definitions(ENABLE_DATA_DUMP)
endif()

if(ENABLE_DEBUGGER)
add_compile_definitions(ENABLE_DEBUGGER)
endif()

+ 38
- 29
cmake/package.cmake View File

@@ -40,6 +40,7 @@ if (CMAKE_SYSTEM_NAME MATCHES "Windows")
set(jpeg_turbo_LIBPATH ${jpeg_turbo_LIBPATH}/../bin/)
set(sqlite_LIBPATH ${sqlite_LIBPATH}/../bin/)
set(tinyxml2_LIBPATH ${tinyxml2_LIBPATH}/../bin/)
set(sentencepiece_LIBPATH ${sentencepiece_LIBPATH}/../bin/)
else ()
set(INSTALL_LIB_DIR "lib")
endif ()
@@ -91,6 +92,14 @@ if (ENABLE_MINDDATA)
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
file(GLOB_RECURSE SENTENCEPIECE_LIB_LIST
${sentencepiece_LIBPATH}/libsentencepiece*
)
install(
FILES ${SENTENCEPIECE_LIB_LIST}
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
if (CMAKE_SYSTEM_NAME MATCHES "Windows")
message("icu4c does not support windows system temporarily")
else()
@@ -123,24 +132,39 @@ if (ENABLE_CPU)
endif ()

if (ENABLE_MPI)
install(
TARGETS _ms_mpi
DESTINATION ${INSTALL_BASE_DIR}
COMPONENT mindspore
if (ENABLE_GPU)
install(
TARGETS _ms_mpi
DESTINATION ${INSTALL_BASE_DIR}
COMPONENT mindspore
)
endif ()
if (ENABLE_CPU)
install(
TARGETS mpi_adapter
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
endif ()
file(GLOB_RECURSE MPI_LIB_LIST
${ompi_LIBPATH}/libmpi${CMAKE_SHARED_LIBRARY_SUFFIX}*
${ompi_LIBPATH}/libopen*${CMAKE_SHARED_LIBRARY_SUFFIX}*
)
install(
TARGETS mpi_adapter
FILES ${MPI_LIB_LIST}
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
endif ()

if (ENABLE_GPU)
if (ENABLE_MPI)
install(
TARGETS gpu_collective
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
endif ()
install(
TARGETS gpu_queue
DESTINATION ${INSTALL_LIB_DIR}
@@ -216,34 +240,12 @@ install(
${CMAKE_SOURCE_DIR}/mindspore/common
${CMAKE_SOURCE_DIR}/mindspore/ops
${CMAKE_SOURCE_DIR}/mindspore/communication
${CMAKE_SOURCE_DIR}/mindspore/profiler
DESTINATION ${INSTALL_PY_DIR}
COMPONENT mindspore
)

if (ENABLE_GPU)
install(
DIRECTORY ${CMAKE_SOURCE_DIR}/mindspore/_akg
DESTINATION ${INSTALL_PY_DIR}/../
COMPONENT mindspore
)
if (EXISTS ${incubator_tvm_gpu_ROOT})
file(GLOB_RECURSE GLOG_LIB_LIST ${incubator_tvm_gpu_LIBPATH}/lib*)
install(
FILES ${GLOG_LIB_LIST}
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
install(
DIRECTORY
${incubator_tvm_gpu_ROOT}/topi/python/topi
${incubator_tvm_gpu_ROOT}/python/tvm
DESTINATION ${INSTALL_PY_DIR}/../_akg
COMPONENT mindspore
)
endif ()
endif ()

if (ENABLE_D AND ENABLE_AKG)
if ((ENABLE_D OR ENABLE_GPU) AND ENABLE_AKG)
set (AKG_PATH ${CMAKE_SOURCE_DIR}/build/mindspore/akg)
install(
DIRECTORY
@@ -268,6 +270,13 @@ if (ENABLE_SERVING)
COMPONENT mindspore
)

install(
FILES ${CMAKE_SOURCE_DIR}/build/mindspore/serving/ms_service_pb2.py
${CMAKE_SOURCE_DIR}/build/mindspore/serving/ms_service_pb2_grpc.py
DESTINATION ${INSTALL_PY_DIR}
COMPONENT mindspore
)

install(
TARGETS inference
DESTINATION ${INSTALL_LIB_DIR}


+ 5
- 2
cmake/package_script.cmake View File

@@ -1,13 +1,16 @@
# find exec
find_package(Python3 3.7 COMPONENTS Interpreter Development)
if (NOT Python3_FOUND)
message("No python3 found.")
return ()
message(FATAL_ERROR "No python3 found.")
endif ()

set(PYTHON ${Python3_EXECUTABLE})
set(PYTHON_VERSION ${Python3_VERSION_MAJOR}.${Python3_VERSION_MINOR})

if (NOT PYTHON_VERSION MATCHES "3.7")
message(FATAL_ERROR "FIND PYTHON VERSION ${PYTHON_VERSION} BUT CAN NOT MATCH PYTHON VERSION 3.7")
endif ()

find_package(Git)
if (NOT GIT_FOUND)
message("No git found.")


+ 4
- 2
config/data_dump.json View File

@@ -1,14 +1,16 @@
{
"DumpSettings": {
"net_name": "ResNet50",
"mode": 1,
"dump_mode": 1,
"op_debug_mode": 3,
"iteration": 0,
"kernels": ["Default/Conv2D-op2", "Default/TensorAdd-op10"]
},

"DumpSettingsSpec": {
"net_name": "net name eg:ResNet50",
"mode": "0: dump all kernels, 1: dump kernels in kernels list",
"dump_mode": "0: dump all kernels, 1: dump kernels in kernels list",
"op_debug_mode": "0: close debug, 1: debug ai-core overflow, 2: debug atomic overflow, 3: debug all overflow",
"iteration": "specified iteration ",
"kernels": "op's full scope name which need to be dump"
}

+ 67
- 0
docker/mindspore-cpu/0.6.0-beta/Dockerfile View File

@@ -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/0.6.0-beta/MindSpore/cpu/ubuntu_x86/mindspore-0.6.0-cp37-cp37m-linux_x86_64.whl

+ 83
- 0
docker/mindspore-gpu/0.6.0-beta/Dockerfile View File

@@ -0,0 +1,83 @@
FROM nvidia/cuda:10.1-cudnn7-runtime-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 \
libnccl2=2.4.8-1+cuda10.1 \
libnccl-dev=2.4.8-1+cuda10.1

# 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/0.6.0-beta/MindSpore/gpu/ubuntu_x86/cuda-10.1/mindspore_gpu-0.6.0-cp37-cp37m-linux_x86_64.whl

+ 1
- 1
graphengine

@@ -1 +1 @@
Subproject commit eee707935c066c16e9b9cd207f8125871b6b97cf
Subproject commit e64a1cfc0457c96859bc9be1693443aa14f2e9df

mindspore/_akg/ops/__init__.py → hub/docs/.gitkeep View File


mindspore/_akg/ops/array/__init__.py → hub/images/.gitkeep View File


mindspore/_akg/ops/math/__init__.py → hub/scripts/.gitkeep View File


+ 108
- 0
include/infer_log.h View File

@@ -0,0 +1,108 @@
/**
* Copyright 2019 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_INFERENCE_LOG_H_
#define MINDSPORE_INFERENCE_LOG_H_

#include <stdarg.h>
#include <stdint.h>
#include <string>
#include <sstream>
#include <memory>
#include <iostream>

#ifndef ENABLE_ACL
#include "mindspore/core/utils/log_adapter.h"
#else // ENABLE_ACL
#include "acl/acl.h"
#endif

namespace mindspore::inference {

class LogStream {
public:
LogStream() { sstream_ = std::make_shared<std::stringstream>(); }
~LogStream() = default;

template <typename T>
LogStream &operator<<(const T &val) noexcept {
(*sstream_) << val;
return *this;
}

LogStream &operator<<(std::ostream &func(std::ostream &os)) noexcept {
(*sstream_) << func;
return *this;
}

friend class LogWriter;
friend class Status;

private:
std::shared_ptr<std::stringstream> sstream_;
};

#ifndef ENABLE_ACL
#define MSI_LOG(level) MS_LOG(level)

#define MSI_LOG_DEBUG MSI_LOG(DEBUG)
#define MSI_LOG_INFO MSI_LOG(INFO)
#define MSI_LOG_WARNING MSI_LOG(WARNING)
#define MSI_LOG_ERROR MSI_LOG(ERROR)

#define MSI_ASSERT(item) MS_ASSERT(item)

#else // ENABLE_ACL

class LogWriter {
public:
LogWriter(const char *file, int line, const char *func, aclLogLevel log_level)
: file_(file), line_(line), func_(func), log_level_(log_level) {}
~LogWriter() = default;

void operator<(const LogStream &stream) const noexcept __attribute__((visibility("default"))) {
std::ostringstream msg;
msg << stream.sstream_->rdbuf();
OutputLog(msg);
}

private:
void OutputLog(const std::ostringstream &msg) const { aclAppLog(log_level_, func_, file_, line_, msg.str().c_str()); }

const char *file_;
int line_;
const char *func_;
aclLogLevel log_level_;
};

#define MSILOG_IF(level) inference::LogWriter(__FILE__, __LINE__, __FUNCTION__, ACL_##level) < inference::LogStream()

#define MSI_LOG(level) MSI_LOG_##level

#define MSI_LOG_DEBUG MSILOG_IF(DEBUG)
#define MSI_LOG_INFO MSILOG_IF(INFO)
#define MSI_LOG_WARNING MSILOG_IF(WARNING)
#define MSI_LOG_ERROR MSILOG_IF(ERROR)

#define MSI_ASSERT(item)

#endif // ENABLE_ACL

#define INFER_STATUS(code) inference::Status(code) < inference::LogStream()

} // namespace mindspore::inference

#endif // MINDSPORE_INFERENCE_LOG_H_

+ 204
- 0
include/infer_tensor.h View File

@@ -0,0 +1,204 @@
/**
* 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_INCLUDE_INFER_TENSOR_H_
#define MINDSPORE_INCLUDE_INFER_TENSOR_H_

#include <utility>
#include <vector>
#include <memory>
#include <numeric>
#include <map>
#include <functional>

#include "securec/include/securec.h"
#include "include/infer_log.h"

namespace mindspore {
#define MS_API __attribute__((visibility("default")))
namespace inference {

enum DataType {
kMSI_Unknown = 0,
kMSI_Bool = 1,
kMSI_Int8 = 2,
kMSI_Int16 = 3,
kMSI_Int32 = 4,
kMSI_Int64 = 5,
kMSI_Uint8 = 6,
kMSI_Uint16 = 7,
kMSI_Uint32 = 8,
kMSI_Uint64 = 9,
kMSI_Float16 = 10,
kMSI_Float32 = 11,
kMSI_Float64 = 12,
};

class InferTensorBase {
public:
InferTensorBase() = default;
virtual ~InferTensorBase() = default;

virtual DataType data_type() const = 0;
virtual void set_data_type(DataType type) = 0;
virtual std::vector<int64_t> shape() const = 0;
virtual void set_shape(const std::vector<int64_t> &shape) = 0;
virtual const void *data() const = 0;
virtual size_t data_size() const = 0;
virtual bool resize_data(size_t data_len) = 0;
virtual void *mutable_data() = 0;

bool set_data(const void *data, size_t data_len) {
resize_data(data_len);
if (mutable_data() == nullptr) {
MSI_LOG_ERROR << "set data failed, data len " << data_len;
return false;
}
if (data_size() != data_len) {
MSI_LOG_ERROR << "set data failed, tensor current data size " << data_size() << " not match data len "
<< data_len;
return false;
}
if (data_len == 0) {
return true;
}
memcpy_s(mutable_data(), data_size(), data, data_len);
return true;
}

int64_t ElementNum() const {
std::vector<int64_t> shapex = shape();
return std::accumulate(shapex.begin(), shapex.end(), 1LL, std::multiplies<int64_t>());
}

int GetTypeSize(DataType type) const {
const std::map<DataType, size_t> type_size_map{
{kMSI_Bool, sizeof(bool)}, {kMSI_Float64, sizeof(double)}, {kMSI_Int8, sizeof(int8_t)},
{kMSI_Uint8, sizeof(uint8_t)}, {kMSI_Int16, sizeof(int16_t)}, {kMSI_Uint16, sizeof(uint16_t)},
{kMSI_Int32, sizeof(int32_t)}, {kMSI_Uint32, sizeof(uint32_t)}, {kMSI_Int64, sizeof(int64_t)},
{kMSI_Uint64, sizeof(uint64_t)}, {kMSI_Float16, sizeof(uint16_t)}, {kMSI_Float32, sizeof(float)},
};
auto it = type_size_map.find(type);
if (it != type_size_map.end()) {
return it->second;
}
return 0;
}
};

class InferTensor : public InferTensorBase {
public:
DataType type_;
std::vector<int64_t> shape_;
std::vector<uint8_t> data_;

public:
InferTensor() = default;
InferTensor(DataType type, std::vector<int64_t> shape, const void *data, size_t data_len) {
set_data_type(type);
set_shape(shape);
set_data(data, data_len);
}

void set_data_type(DataType type) override { type_ = type; }
DataType data_type() const override { return type_; }

void set_shape(const std::vector<int64_t> &shape) override { shape_ = shape; }
std::vector<int64_t> shape() const override { return shape_; }

const void *data() const override { return data_.data(); }
size_t data_size() const override { return data_.size(); }

bool resize_data(size_t data_len) override {
data_.resize(data_len);
return true;
}
void *mutable_data() override { return data_.data(); }
};

class InferImagesBase {
public:
virtual size_t batch_size() const = 0;
virtual bool get(size_t index, const void *&pic_buffer, uint32_t &pic_size) const = 0;
virtual size_t input_index() const = 0; // the index of images as input in model
};

class RequestBase {
public:
virtual size_t size() const = 0;
virtual const InferTensorBase *operator[](size_t index) const = 0;
};

class ImagesRequestBase {
public:
virtual size_t size() const = 0;
virtual const InferImagesBase *operator[](size_t index) const = 0;
};

class ReplyBase {
public:
virtual size_t size() const = 0;
virtual InferTensorBase *operator[](size_t index) = 0;
virtual const InferTensorBase *operator[](size_t index) const = 0;
virtual InferTensorBase *add() = 0;
virtual void clear() = 0;
};

class VectorInferTensorWrapReply : public ReplyBase {
public:
explicit VectorInferTensorWrapReply(std::vector<InferTensor> &tensor_list) : tensor_list_(tensor_list) {}

size_t size() const { return tensor_list_.size(); }
InferTensorBase *operator[](size_t index) {
if (index >= tensor_list_.size()) {
MSI_LOG_ERROR << "visit invalid index " << index << " total size " << tensor_list_.size();
return nullptr;
}
return &(tensor_list_[index]);
}
const InferTensorBase *operator[](size_t index) const {
if (index >= tensor_list_.size()) {
MSI_LOG_ERROR << "visit invalid index " << index << " total size " << tensor_list_.size();
return nullptr;
}
return &(tensor_list_[index]);
}
InferTensorBase *add() {
tensor_list_.push_back(InferTensor());
return &(tensor_list_.back());
}
void clear() { tensor_list_.clear(); }
std::vector<InferTensor> &tensor_list_;
};

class VectorInferTensorWrapRequest : public RequestBase {
public:
explicit VectorInferTensorWrapRequest(const std::vector<InferTensor> &tensor_list) : tensor_list_(tensor_list) {}

size_t size() const { return tensor_list_.size(); }
const InferTensorBase *operator[](size_t index) const {
if (index >= tensor_list_.size()) {
MSI_LOG_ERROR << "visit invalid index " << index << " total size " << tensor_list_.size();
return nullptr;
}
return &(tensor_list_[index]);
}
const std::vector<InferTensor> &tensor_list_;
};

} // namespace inference
} // namespace mindspore
#endif // MINDSPORE_INCLUDE_INFER_TENSOR_H_

+ 48
- 10
include/inference.h View File

@@ -20,25 +20,63 @@
#include <memory>
#include <vector>
#include <string>
#include "include/ms_tensor.h"
#include "include/infer_tensor.h"
#include "include/infer_log.h"

namespace mindspore {
class FuncGraph;
namespace inference {
class MS_API MSSession {
public:
MSSession() = default;

static std::shared_ptr<MSSession> CreateSession(const std::string &device, uint32_t device_id);
enum StatusCode { SUCCESS = 0, FAILED, INVALID_INPUTS };

virtual uint32_t CompileGraph(std::shared_ptr<FuncGraph> funcGraphPtr) = 0;
class Status {
public:
Status() : status_code_(FAILED) {}
Status(enum StatusCode status_code, const std::string &status_msg = "")
: status_code_(status_code), status_msg_(status_msg) {}
bool IsSuccess() const { return status_code_ == SUCCESS; }
enum StatusCode StatusCode() const { return status_code_; }
std::string StatusMessage() const { return status_msg_; }
bool operator==(const Status &other) const { return status_code_ == other.status_code_; }
bool operator==(enum StatusCode other_code) const { return status_code_ == other_code; }
bool operator!=(const Status &other) const { return status_code_ != other.status_code_; }
bool operator!=(enum StatusCode other_code) const { return status_code_ != other_code; }
operator bool() const = delete;
Status &operator<(const LogStream &stream) noexcept __attribute__((visibility("default"))) {
status_msg_ = stream.sstream_->str();
return *this;
}

virtual MultiTensor RunGraph(uint32_t graph_id, const std::vector<std::shared_ptr<inference::MSTensor>> &inputs) = 0;
private:
enum StatusCode status_code_;
std::string status_msg_;
};

std::shared_ptr<FuncGraph> MS_API LoadModel(const char *model_buf, size_t size, const std::string &device);
class MS_API InferSession {
public:
InferSession() = default;
virtual ~InferSession() = default;
virtual Status InitEnv(const std::string &device_type, uint32_t device_id) = 0;
virtual Status FinalizeEnv() = 0;
virtual Status LoadModelFromFile(const std::string &file_name, uint32_t &model_id) = 0;
virtual Status UnloadModel(uint32_t model_id) = 0;
// override this method to avoid request/reply data copy
virtual Status ExecuteModel(uint32_t model_id, const RequestBase &request, ReplyBase &reply) = 0;

virtual Status ExecuteModel(uint32_t model_id, const std::vector<InferTensor> &inputs,
std::vector<InferTensor> &outputs) {
VectorInferTensorWrapRequest request(inputs);
VectorInferTensorWrapReply reply(outputs);
return ExecuteModel(model_id, request, reply);
}
// default not support input data preprocess(decode, resize, crop, crop&paste, etc.)
virtual Status ExecuteModel(uint32_t /*model_id*/,
const ImagesRequestBase & /*images_inputs*/, // images for preprocess
const RequestBase & /*request*/, ReplyBase & /*reply*/) {
return FAILED;
}
static std::shared_ptr<InferSession> CreateSession(const std::string &device, uint32_t device_id);
};

void MS_API ExitInference();
} // namespace inference
} // namespace mindspore
#endif // MINDSPORE_INCLUDE_MS_SESSION_H

+ 0
- 69
include/ms_tensor.h View File

@@ -1,69 +0,0 @@
/**
* 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_INCLUDE_MS_TENSOR_H_
#define MINDSPORE_INCLUDE_MS_TENSOR_H_

#include <utility>
#include <vector>
#include <memory>
#include "mindspore/core/ir/dtype/type_id.h"

namespace mindspore {
#define MS_API __attribute__((visibility("default")))
namespace inference {
class MS_API MSTensor {
public:
MSTensor() = default;
// brief Create a MSTensor pointer.
//
// param data_type DataTypeId of tensor to be created.
// param shape Shape of tensor to be created.
// return MSTensor pointer.
static MSTensor *CreateTensor(TypeId data_type, const std::vector<int> &shape);

~MSTensor() = default;

virtual TypeId data_type() const = 0;

virtual TypeId set_data_type(const TypeId data_type) = 0;

virtual std::vector<int> shape() const = 0;

virtual size_t set_shape(const std::vector<int> &shape) = 0;

virtual int DimensionSize(size_t index) const = 0;
// brief Get number of element in MSTensor.
//
// return Number of element in MSTensor.
virtual int ElementsNum() const = 0;

virtual std::size_t hash() const = 0;
// brief Get byte size of data in MSTensor.
//
// return Byte size of data in MSTensor.
virtual size_t Size() const = 0;
// brief Get pointer of data in MSTensor.
//
// The data pointer can be used to both write or read data in MSTensor.
//
// return A pointer points to data in MSTensor.
virtual void *MutableData() const = 0;
};
using MultiTensor = std::vector<std::shared_ptr<inference::MSTensor>>;
} // namespace inference
} // namespace mindspore
#endif // MINDSPORE_INCLUDE_MS_TENSOR_H_

+ 0
- 18
mindspore/_akg/__init__.py View File

@@ -1,18 +0,0 @@
# Copyright 2019 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.

"""__init__"""
from . import add_path
from .op_build import op_build
from .message import compilewithjson

+ 0
- 62
mindspore/_akg/add_path.py View File

@@ -1,62 +0,0 @@
# 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.

"""add tvm path"""
import sys
import os


def AKGAddPath():
"""_akg add path."""
pwd = os.path.dirname(os.path.realpath(__file__))
tvm_path = os.path.realpath(pwd)
if tvm_path not in sys.path:
sys.path.insert(0, tvm_path)
else:
sys.path.remove(tvm_path)
sys.path.insert(0, tvm_path)


class AKGMetaPathFinder:
"""class AKGMetaPath finder."""

def find_module(self, fullname, path=None):
"""method _akg find module."""
_ = path
if fullname.startswith("_akg.tvm"):
rname = fullname[5:]
return AKGMetaPathLoader(rname)
if fullname.startswith("_akg.topi"):
rname = fullname[5:]
return AKGMetaPathLoader(rname)
return None


class AKGMetaPathLoader:
"""class AKGMetaPathLoader loader."""

def __init__(self, rname):
self.__rname = rname

def load_module(self, fullname):
if self.__rname in sys.modules:
sys.modules.pop(self.__rname)
AKGAddPath()
__import__(self.__rname, globals(), locals())
self.__target_module = sys.modules[self.__rname]
sys.modules[fullname] = self.__target_module
return self.__target_module


sys.meta_path.insert(0, AKGMetaPathFinder())

+ 0
- 39
mindspore/_akg/gpu/__init__.py View File

@@ -1,39 +0,0 @@
# Copyright 2019 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.

"""__init__"""
from .equal import Equal
from .equal import gpu_schedule_Equal
from .tile import Tile
from .tile import gpu_schedule_Tile
from .cast import Cast
from .cast import gpu_schedule_Cast
from .relu6 import ReLU6, gpu_schedule_ReLU6
from .relu6_grad import ReLU6Grad, gpu_schedule_ReLU6Grad
from .squeeze import Squeeze, gpu_schedule_Squeeze
from .squeeze_grad import SqueezeGrad, gpu_schedule_SqueezeGrad
from .mean import SimpleMean, gpu_schedule_SimpleMean
from .mean_grad import SimpleMeanGrad, gpu_schedule_SimpleMeanGrad
from .mul import Mul, gpu_schedule_Mul
from .hsigmoid import HSigmoid, gpu_schedule_HSigmoid
from .hsigmoid_grad import HSigmoidGrad, gpu_schedule_HSigmoidGrad
from .hswish import HSwish, gpu_schedule_HSwish
from .hswish_grad import HSwishGrad, gpu_schedule_HSwishGrad
from .logical_or import LogicalOr, gpu_schedule_LogicalOr
from .logical_not import LogicalNot, gpu_schedule_LogicalNot
from .logical_and import LogicalAnd, gpu_schedule_LogicalAnd
from .sub import Sub, gpu_schedule_Sub
from .less_equal import LessEqual, gpu_schedule_LessEqual
from .notequal import NotEqual, gpu_schedule_NotEqual
from .greater_equal import GreaterEqual, gpu_schedule_GreaterEqual

+ 0
- 43
mindspore/_akg/gpu/cast.py View File

@@ -1,43 +0,0 @@
# Copyright 2019 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.

"""cast"""
import logging
import _akg.tvm
from _akg.ops.math import cast
from _akg.topi.generic import schedule_elemwise

def Cast(x, dst_type):
"""cast."""
return cast.cast(x, dst_type)


def gpu_schedule_Cast(outs):
"""
gpu schedule for cast.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
logging.info("Skip because %s is not enabled", device)
return None
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 56
mindspore/_akg/gpu/default_schedule.py View File

@@ -1,56 +0,0 @@
# Copyright 2019 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.

"""default schedule function for GPU"""
from queue import Queue

import _akg.tvm as tvm

DEFAULT_GPU_THREAD = 1024


def default_schedule(outs):
"""
default schedule function.

Args:
outs (Union[tvm.tensor.Tensor, list[tvm.tensor.Tensor]]): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
if not isinstance(outs, tvm.tensor.Tensor) and not isinstance(outs, list):
raise ValueError("outs should be list of _akg.tvm.tensor.Tensor or _akg.tvm.tensor.Tensor")
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
outs_list = [outs] if isinstance(outs, tvm.tensor.Tensor) else outs
with tvm.target.create(device):
sch = tvm.create_schedule(outs_list[0].op)
outputs_tensor = Queue()
outputs_tensor.put(outs_list[0])
op_list = []
while not outputs_tensor.empty():
out = outputs_tensor.get()
if out.op not in op_list and isinstance(out.op, tvm.tensor.ComputeOp):
op_list.append(out.op)
for input_tensor in out.op.input_tensors:
outputs_tensor.put(input_tensor)
for op in op_list:
stage = sch[op.output(0)]
bx, tx = stage.split(op.axis[0], factor=DEFAULT_GPU_THREAD)
stage.bind(bx, tvm.thread_axis("blockIdx.x"))
stage.bind(tx, tvm.thread_axis("threadIdx.x"))
return sch

+ 0
- 40
mindspore/_akg/gpu/equal.py View File

@@ -1,40 +0,0 @@
# Copyright 2019 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.
"""equal"""
import _akg.tvm
from _akg.ops.math import equal
from _akg.topi.generic import schedule_elemwise

def Equal(x, y):
"""equal."""
return equal.equal(x, y)


def gpu_schedule_Equal(outs):
"""
gpu schedule for Equal.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 41
mindspore/_akg/gpu/greater_equal.py View File

@@ -1,41 +0,0 @@
# 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.

"""greater_equal"""
import _akg.tvm
from _akg.ops.math import greater_equal
from _akg.topi.generic import schedule_elemwise

def GreaterEqual(x, y):
"""GreaterEqual."""
return greater_equal.greater_equal(x, y)


def gpu_schedule_GreaterEqual(outs):
"""
GPU schedule for GreaterEqual.

Args:
outs (tvm.tensor.Tensor): Outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 63
mindspore/_akg/gpu/hsigmoid.py View File

@@ -1,63 +0,0 @@
# Copyright 2019 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.

"""hsigmoid"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag


@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_hsigmoid(x):
"""
topi hsigmoid
Args:
x:

Returns:

"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, 1,
(x(*i) + 3) / 6)))


def HSigmoid(x):
"""
HSigmoid
Args:
x:

Returns:

"""
return topi_nn_hsigmoid(x)


def gpu_schedule_HSigmoid(outs):
"""
gpu schedule HSigmoid
Args:
outs:

Returns:

"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 51
mindspore/_akg/gpu/hsigmoid_grad.py View File

@@ -1,51 +0,0 @@
# Copyright 2019 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.

"""HSigmoid grad"""
import _akg.topi as topi
import _akg.tvm as tvm


def HSigmoidGrad(y_grad, x):
"""
HSigmoidGrad
Args:
y_grad:
x:

Returns:

"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, 0,
y_grad(*i) / 6)))


def gpu_schedule_HSigmoidGrad(outs):
"""
gpu schedule ReLU6Grad
Args:
outs:

Returns:

"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)

with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 63
mindspore/_akg/gpu/hswish.py View File

@@ -1,63 +0,0 @@
# Copyright 2019 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.

"""HSwish"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag


@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_HSwish(x):
"""
topi HSwish
Args:
x:

Returns:

"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, x(*i),
x(*i) * (x(*i) + 3) / 6)))


def HSwish(x):
"""
HSwish
Args:
x:

Returns:

"""
return topi_nn_HSwish(x)


def gpu_schedule_HSwish(outs):
"""
gpu schedule HSwish
Args:
outs:

Returns:

"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 53
mindspore/_akg/gpu/hswish_grad.py View File

@@ -1,53 +0,0 @@
# Copyright 2019 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.

"""HSwishGrad"""
import _akg.topi as topi
import _akg.tvm as tvm


def HSwishGrad(y_grad, x):
"""
HSwishGrad
Args:
y_grad:
x:

Returns:

"""
shape = x.shape

res0 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0, y_grad(*i) * (2 * x(*i) + 3) / 6))
res6 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= 3, y_grad(*i), res0(*i)))
return res6


def gpu_schedule_HSwishGrad(outs):
"""
gpu schedule HSwishGrad
Args:
outs:

Returns:

"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)

with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 40
mindspore/_akg/gpu/less_equal.py View File

@@ -1,40 +0,0 @@
# 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.
"""less_equal"""
import _akg.tvm
from _akg.ops.math import less_equal
from _akg.topi.generic import schedule_elemwise

def LessEqual(x, y):
"""LessEqual."""
return less_equal.less_equal(x, y)


def gpu_schedule_LessEqual(outs):
"""
GPU schedule for LessEqual.

Args:
outs (tvm.tensor.Tensor): Outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 40
mindspore/_akg/gpu/logical_and.py View File

@@ -1,40 +0,0 @@
# 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.
"""logical_and"""
import _akg.tvm
from _akg.ops.math import logical_and
from _akg.topi.generic import schedule_elemwise

def LogicalAnd(x, y):
"""LogicalAnd."""
return logical_and.logical_and(x, y)


def gpu_schedule_LogicalAnd(outs):
"""
GPU schedule for LogicalAnd.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 40
mindspore/_akg/gpu/logical_not.py View File

@@ -1,40 +0,0 @@
# 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.
"""logical_not"""
import _akg.tvm
from _akg.ops.math import logical_not
from _akg.topi.generic import schedule_elemwise

def LogicalNot(x):
"""LogicalNot."""
return logical_not.logical_not(x)


def gpu_schedule_LogicalNot(outs):
"""
GPU schedule for LogicalNot.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 40
mindspore/_akg/gpu/logical_or.py View File

@@ -1,40 +0,0 @@
# 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.
"""logical_or"""
import _akg.tvm
from _akg.ops.math import logical_or
from _akg.topi.generic import schedule_elemwise

def LogicalOr(x, y):
"""LogicalOr."""
return logical_or.logical_or(x, y)


def gpu_schedule_LogicalOr(outs):
"""
GPU schedule for LogicalOr.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 80
mindspore/_akg/gpu/mean.py View File

@@ -1,80 +0,0 @@
# Copyright 2019 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.

"""mean op compute and schedule"""
import _akg.tvm as tvm
from _akg.ops.math.mean import mean
from .default_schedule import DEFAULT_GPU_THREAD

def Mean(x, axis=None, keepdims=True):
"""mean."""
outs = mean(x, axis, keepdims)

# remove useless mean_output
if isinstance(outs, tuple):
outs = outs[0]
if outs.op.name == "mean_output":
outs = outs.op.input_tensors[0]
return outs


def gpu_schedule_Mean(outs):
"""
gpu schedule function for mean.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
out = outs[0] if isinstance(outs, list) else outs

device = "cuda"
with tvm.target.create(device):
sch = tvm.create_schedule(out.op)
if out.op.name == "T_divide":
tensor_c = out
else: # squeeze
tensor_c = out.op.input_tensors[0]

tensor_b = tensor_c.op.input_tensors[0]
if len(tensor_c.op.axis) >= 2:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[1])
else:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[0])

bx, tx = sch[tensor_c].split(tensor_c.op.axis[0], factor=DEFAULT_GPU_THREAD)
sch[tensor_c].bind(bx, tvm.thread_axis("blockIdx.x"))
sch[tensor_c].bind(tx, tvm.thread_axis("threadIdx.x"))
return sch

def SimpleMean(x):
"""
SimpleMean compute the mean of the input 4D Tensor over last two axises and keep reduced dimensions.

Args:
x (tvm.tensor.Tensor): Tensor of type float16, float32.

Returns:
tvm.tensor.Tensor, has the same type as x, output shape will be (a, b, 1, 1) if input Tensor x is (a, b, c, d).
"""
axis = (2, 3)
keepdims = True
return Mean(x, axis, keepdims)


def gpu_schedule_SimpleMean(outs):
"""gpu schedule function for SimpleMean."""
return gpu_schedule_Mean(outs)

+ 0
- 90
mindspore/_akg/gpu/mean_grad.py View File

@@ -1,90 +0,0 @@
# Copyright 2019 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.

"""mean_grad"""
import _akg.tvm as tvm
import _akg
from _akg.ops.math import mean
from .default_schedule import DEFAULT_GPU_THREAD


def mean_ad(head, input_shape, axis, keepdims):
"""mean autodiff."""
tensor_a = tvm.placeholder(input_shape, head.dtype, "A")
tensor_b = mean.mean(tensor_a, axis, keepdims)

# remove useless mean_output
if isinstance(tensor_b, tuple):
tensor_b = tensor_b[0]
if tensor_b.op.name == "mean_output":
tensor_b = tensor_b.op.input_tensors[0]

jacs = list(_akg.differentiate(tensor_b, [tensor_a], head))
return jacs[0]


def MeanGrad(y_grad, input_shape, axis=None, keepdims=True):
"""Mean Grad."""
if axis is None and not keepdims:
raise ValueError("Mean not support (axis=None && keepdims=False) now")
return mean_ad(y_grad, input_shape, axis, keepdims)


def gpu_schedule_MeanGrad(outs):
"""gpu schedule MeanGrad."""
out = outs[0] if isinstance(outs, list) else outs

device = "cuda"
with tvm.target.create(device):
sch = tvm.create_schedule(out.op)
tensor_c = out
tensor_b = tensor_c.op.input_tensors[0]
if len(tensor_c.op.axis) >= 2:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[1])
else:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[0])

bx, tx = sch[tensor_c].split(tensor_c.op.axis[0], factor=DEFAULT_GPU_THREAD)
sch[tensor_c].bind(bx, tvm.thread_axis("blockIdx.x"))
sch[tensor_c].bind(tx, tvm.thread_axis("threadIdx.x"))

return sch

def SimpleMeanGrad(HEAD, input_shape):
"""
Compute Simple Mean Grad.

Args:
HEAD (tvm.tensor.Tensor): output gradient, dy, defined in Primitive.
input_shape (Union[list[int], tuple[int]]): shape of mean input, x.shape.

Returns:
tvm.tensor.Tensor, gradient of mean input.
"""
axis = (2, 3)
keepdims = True
return MeanGrad(HEAD, input_shape, axis, keepdims)


def gpu_schedule_SimpleMeanGrad(outs):
"""
gpu schedule SimpleMeanGrad.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
return gpu_schedule_MeanGrad(outs)

+ 0
- 41
mindspore/_akg/gpu/mul.py View File

@@ -1,41 +0,0 @@
# Copyright 2019 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.

"""mul"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.ops.math import mul

def Mul(x, y):
"""mul."""
return mul.mul(x, y)


def gpu_schedule_Mul(outs):
"""
gpu schedule for mul.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_broadcast(outs)
return sch

+ 0
- 41
mindspore/_akg/gpu/notequal.py View File

@@ -1,41 +0,0 @@
# 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.

"""notequal"""
import _akg.tvm
from _akg.ops.math import notequal
from _akg.topi.generic import schedule_elemwise

def NotEqual(x, y):
"""notequal."""
return notequal.notequal(x, y)


def gpu_schedule_NotEqual(outs):
"""
gpu schedule for NotEqual.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 54
mindspore/_akg/gpu/relu6.py View File

@@ -1,54 +0,0 @@
# Copyright 2019 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.

"""relu6"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag

@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_relu6(x):
"""topi nn relu6."""
return tvm.compute(x.shape, lambda *i: tvm.min(tvm.max(x(*i), tvm.const(0, x.dtype)), tvm.const(6, x.dtype)))

def ReLU6(x):
"""
Compute elementwise with function: min(max(x, 0), 6).

Args:
x (tvm.tensor.Tensor): Tensor of type float16, float32.

Returns:
tvm.tensor.Tensor, has same type and shape as input.
"""
return topi_nn_relu6(x)


def gpu_schedule_ReLU6(outs):
"""
gpu schedule ReLU6.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 59
mindspore/_akg/gpu/relu6_grad.py View File

@@ -1,59 +0,0 @@
# Copyright 2019 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.

"""relu6 grad"""
import _akg.topi as topi
import _akg.tvm as tvm

def ReLU6Grad(y_grad, x):
"""
Computes Gradients of Rectified Linear 6.

Args:
y_grad (tvm.tensor.Tensor): Tensor of type float16, float32, gradients backpropagated to the ReLU6 op.
x (tvm.tensor.Tensor): Tensor of type float16/float32, inputs that where passed to the ReLU6 op, or its outputs.

Returns:
tvm.tensor.Tensor, has same type and shape as x.
"""
shape = x.shape
dtype = x.dtype

zero = tvm.const(0, dtype)
six = tvm.const(6, dtype)

res0 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= zero, x(*i), zero))
res6 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= six, zero, res0(*i)))
res = tvm.compute(shape, lambda *i: tvm.if_then_else(res6(*i) == zero, zero, y_grad(*i)))
return res


def gpu_schedule_ReLU6Grad(outs):
"""
gpu schedule ReLU6Grad.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)

with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

+ 0
- 50
mindspore/_akg/gpu/squeeze.py View File

@@ -1,50 +0,0 @@
# Copyright 2019 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.

"""squeeze"""
import _akg.topi as topi
import _akg.tvm as tvm

def Squeeze(x, axis=None):
"""
Remove the dimensions which have shape size 1.

Args:
x (tvm.tensor.Tensor): Tensor, input whose shape is to be squeeze.
axis (Union[list, tuple, int, None]): specify which size 1 dimension to be removed.

Returns:
tvm.tensor.Tensor, has the same type and element as x, but some size 1 dimensions are removed.
"""
return topi.squeeze(x, axis)


def gpu_schedule_Squeeze(outs):
"""
gpu schedule Squeeze.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)

with tvm.target.create(device):
sch = topi.cuda.schedule_injective(outs)
return sch

+ 0
- 44
mindspore/_akg/gpu/squeeze_grad.py View File

@@ -1,44 +0,0 @@
# Copyright 2019 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.

"""squeeze grad"""
import _akg.topi as topi


def SqueezeGrad(y_grad, x_shape):
"""
Computes gradients for squeeze op.

Args:
y_grad (tvm.tensor.Tensor): the gradient needed to be propagation.
x_shape (Union[list, tuple]): output Tensor shape.

Returns:
tvm.tensor.Tensor: output gradient.
"""
return topi.reshape(y_grad, x_shape)


def gpu_schedule_SqueezeGrad(outs):
"""
gpu schedule SqueezeGrad.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
from .default_schedule import default_schedule
return default_schedule(outs)

+ 0
- 40
mindspore/_akg/gpu/sub.py View File

@@ -1,40 +0,0 @@
# 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.
"""sub"""
import _akg.tvm
from _akg.ops.math import sub
from _akg.topi.generic import schedule_elemwise

def Sub(x, y):
"""Sub."""
return sub.sub(x, y)


def gpu_schedule_Sub(outs):
"""
GPU schedule for Sub.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

+ 0
- 39
mindspore/_akg/gpu/tile.py View File

@@ -1,39 +0,0 @@
# Copyright 2019 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.
"""tile"""
import _akg.tvm
from _akg.ops.array import tile
from _akg.topi.generic import schedule_elemwise

def Tile(x, multiples):
"""tile."""
return tile.tile(x, multiples)

def gpu_schedule_Tile(outs):
"""
gpu schedule for tile.

Args:
outs (tvm.tensor.Tensor): outputs of compute.

Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
s = schedule_elemwise(outs)
return s

+ 0
- 104
mindspore/_akg/message.py View File

@@ -1,104 +0,0 @@
# Copyright 2019 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.

"""message"""
import importlib.util
import json
import json.decoder as jd
import logging
import traceback
import os.path
from pathlib import Path
import _akg.tvm
from _akg.utils import validation_check as vc_util
from _akg.utils.dsl_create import TensorUtils
from . import gpu
from . import op_build


@vc_util.check_input_type(str)
def compilewithjson(json_str):
"""compile with json."""
try:
kernel_info = json.loads(json_str)
except jd.JSONDecodeError:
logging.error(traceback.format_exc())
return False

op_name = kernel_info['name']
op_func = None
processor = 'aicore'
if 'process' in kernel_info:
processor = kernel_info['process']
# get custom ops implementation first.
if 'impl_path' in kernel_info and kernel_info['impl_path'] is not None:
impl_path = os.path.realpath(kernel_info['impl_path'])
if os.path.isfile(impl_path):
custom_mod_name = Path(impl_path).resolve().stem
mod_spec = importlib.util.spec_from_file_location(
custom_mod_name, impl_path)
custom_mod = importlib.util.module_from_spec(mod_spec)
mod_spec.loader.exec_module(custom_mod)
op_func = getattr(custom_mod, op_name, None)

# get built-in ops.
if op_func is None:
if processor == 'cuda':
op_func = getattr(gpu, op_name, None)

if op_func is None:
logging.error(
"this op not supported, please check op name %s", str(op_name))
return False

args = {}
tsr = []
for input_desc in kernel_info['input_desc']:
if len(input_desc) == 1:
tensor_shape = input_desc[0]['shape']
tensor_shape = (1,) if not tensor_shape else tensor_shape
vc_util.shape_dtype_max_size_check(tensor_shape)
args[input_desc[0]['name']] = _akg.tvm.placeholder(
shape=tensor_shape, name=input_desc[0]['tensor_name'], dtype=input_desc[0]['data_type'])
tsr.append(args[input_desc[0]['name']])
else:
tmp_input = []
for tmp_desc in input_desc:
tensor_shape = tmp_desc['shape']
tensor_shape = (1,) if not tensor_shape else tensor_shape
vc_util.shape_dtype_max_size_check(tensor_shape)
tmp_input.append(_akg.tvm.placeholder(
shape=tensor_shape, name=tmp_desc['tensor_name'], dtype=tmp_desc['data_type']))
args[input_desc[0]['name']] = tmp_input
tsr = tsr + tmp_input

if kernel_info['attr']:
for ext_arg in kernel_info['attr']:
args[ext_arg['name']] = ext_arg['value']

output = op_func(**args)

if isinstance(output, (list, tuple)):
from inspect import isfunction
tmp_outputs = []
for elem in output:
if not isfunction(elem) or isinstance(elem, dict):
tmp_outputs.append(elem)

output = tmp_outputs
else:
output = [output]

tsr = tsr + [i for i in output if TensorUtils.is_output_value(i)]
return op_build([op_name], output, tsr, processor, kernel_info['op'])

+ 0
- 69
mindspore/_akg/op_build.py View File

@@ -1,69 +0,0 @@
# Copyright 2019 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.

"""op_build"""
import os
import fcntl
import types
import typing
import logging
import traceback
import _akg.tvm
import _akg
from _akg import save_gpu_param as gpu_utils
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(list, (list, tuple), (list, tuple), str, str)
def op_build(opnames, computes, args, device, kernel_name):
"""op_build"""
kernel_meta_path = "./cuda_meta_" + str(os.getpid()) + "/"
if device == "cuda":
cuda_path = os.path.realpath(kernel_meta_path)
if not os.path.isdir(cuda_path):
os.makedirs(cuda_path)
if not opnames:
logging.error("no opname given.")
return None

schedule_name = 'gpu_schedule_' + opnames[0]
schedule_func = getattr(_akg.gpu, schedule_name)
if not isinstance(schedule_func, (types.FunctionType, typing.Callable)):
logging.error("no schedule func found %s", str(schedule_name))
return None

ptx_file = os.path.realpath(kernel_meta_path + kernel_name + ".ptx")
if os.path.exists(ptx_file):
os.chmod(ptx_file, 0o600)
try:
with open(ptx_file, 'at') as file:
fcntl.flock(file.fileno(), fcntl.LOCK_EX)
file.seek(0, 2)
if file.tell() == 0:
s = schedule_func(computes)
foo = _akg.tvm.build(s, args, device, name=kernel_name)
ptx_code = foo.imported_modules[0].get_source("ptx")
file.write(ptx_code)
json_file = os.path.realpath(
kernel_meta_path + kernel_name + ".json")
kernel_info = (ptx_code, json_file, kernel_name)
gpu_utils.save_gpu_params(s, args, kernel_info)
os.chmod(ptx_file, 0o400)
except IOError:
logging.error(traceback.format_exc())
return None
return True

logging.error("Not support device %s.", device)
return None

+ 0
- 36
mindspore/_akg/ops/array/tile.py View File

@@ -1,36 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: tile"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple))
def tile(data, multiples):
"""
Repeats the data in the specified dimensions according to the multiples.

Args:
data (tvm.tensor.Tensor): Tensor.
multiples (Union[list, tuple]): Elements must be int. The number of repetitions.

Returns:
tvm.tensor.Tensor, has the same dtype as data.
"""
vc_util.check_shape(data.shape)
vc_util.check_int_list(multiples, "multiples")
output = _akg.topi.tile(data, multiples)
return output

+ 0
- 36
mindspore/_akg/ops/math/cast.py View File

@@ -1,36 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: cast"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, str)
def cast(data, dst_type):
"""
cast data to target type.

Args:
data (tvm.tensor.Tensor): Tensor to be casted.
dst_type (str): target cast type.

Returns:
tvm.tensor.Tensor, type is dst_type.
"""
vc_util.check_shape(data.shape)
out = _akg.topi.cast(data, dst_type)

return out

+ 0
- 54
mindspore/_akg/ops/math/equal.py View File

@@ -1,54 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: equal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def equal(input1, input2):
"""
check whether input1 equals to input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. If input1 equal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

shape1, shape2, shape = produce_shapes(shape1, shape2)

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype

# get equal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")

input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] == input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")

return res

+ 0
- 54
mindspore/_akg/ops/math/greater_equal.py View File

@@ -1,54 +0,0 @@
# 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.

"""operator dsl function: greaterequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def greater_equal(input1, input2):
"""
Check whether input1 greaterquals to input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. If input1 greaterquals to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

shape1, shape2, shape = produce_shapes(shape1, shape2)

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype

# get greaterquals compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")

input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] >= input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")

return res

+ 0
- 54
mindspore/_akg/ops/math/less_equal.py View File

@@ -1,54 +0,0 @@
# 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.

"""operator dsl function: lessequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def less_equal(input1, input2):
"""
Check whether input1 lessequals to input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. If input1 lessequal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

shape1, shape2, shape = produce_shapes(shape1, shape2)

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype

# get lessequal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")

input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] <= input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")

return res

+ 0
- 41
mindspore/_akg/ops/math/logical_and.py View File

@@ -1,41 +0,0 @@
# 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.

"""operator dsl function: logical_and"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util

@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def logical_and(input1, input2):
"""
Compute logical_and of input1 and input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. LogicalAnd of input1 and input2.
"""

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)

shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

res = _akg.topi.logical_and(input1, input2)
return res

+ 0
- 32
mindspore/_akg/ops/math/logical_not.py View File

@@ -1,32 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: logical_not"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util

@vc_util.check_input_type(_akg.tvm.tensor.Tensor)
def logical_not(input1):
"""
Compute logical_not of input1.

Args:
input1 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor.
"""
res = _akg.topi.logical_not(input1)
return res

+ 0
- 41
mindspore/_akg/ops/math/logical_or.py View File

@@ -1,41 +0,0 @@
# 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.

"""operator dsl function: logical_or"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util

@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def logical_or(input1, input2):
"""
Compute logical_or of input1 and input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. LogicalOr of input1 and input2.
"""

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)

shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

res = _akg.topi.logical_or(input1, input2)
return res

+ 0
- 47
mindspore/_akg/ops/math/mean.py View File

@@ -1,47 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: mean"""
import _akg.topi
import _akg.tvm
from _akg.utils import format_transform as ft_util
from _akg.utils import validation_check as vc_util
from _akg.ops.math import sum_value


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
def mean(data, axis=None, keepdims=False):
"""
Computes the mean of the values of a Tensor over the whole dataset.

Args:
data (tvm.tensor.Tensor): Tensor.
axis (Union[list, tuple, int, None]): If the tuple is empty, the axis equal to None.
keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.

Returns:
tvm.tensor.Tensor, has the same type as data. If keepdims equal to True, all reduced dimensions are
retained with length 1. else these reduced axis will be eliminate.
"""
shape = [x.value for x in data.shape]
vc_util.reduce_axis_check(shape, axis)
axis = ft_util.refine_reduce_axis(data, axis)

count = 1
for i in axis:
count *= shape[i]
output, _ = sum_value.sum_value(data, axis, keepdims)
res = _akg.topi.divide(output, count)

return res

+ 0
- 43
mindspore/_akg/ops/math/mul.py View File

@@ -1,43 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: mul"""
import _akg.topi
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def mul(l_input, r_input):
"""
Calculate x * y element-wise.

Note:
mul supports broadcasting.

Args:
l_input (tvm.tensor.Tensor): Tensor.
r_input (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor, has the same type as l_input and r_input.
"""
shape1 = [x.value for x in l_input.shape]
shape2 = [x.value for x in r_input.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
vc_util.auto_broadcast_check(shape1, shape2)
vc_util.elemwise_dtype_check(l_input.dtype, r_input.dtype)
output = _akg.topi.multiply(l_input, r_input)

return output

+ 0
- 54
mindspore/_akg/ops/math/notequal.py View File

@@ -1,54 +0,0 @@
# 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.

"""operator dsl function: notequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def notequal(input1, input2):
"""
check whether input1 notequals to input2.

Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.

Returns:
tvm.tensor.Tensor. If input1 notequal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)

shape1, shape2, shape = produce_shapes(shape1, shape2)

vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype

# get notequal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")

input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] != input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")

return res

+ 0
- 40
mindspore/_akg/ops/math/sub.py View File

@@ -1,40 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: sub"""
import _akg.topi
import _akg.tvm
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def sub(data1, data2):
"""
Computes data1 - data2 elementwise, broadcast is supported.

Args:
data1 (tvm.tensor.Tensor): Tensor.
data2 (tvm.tensor.Tensor): Tensor of same type as data1, if shape(data2) != shape(data1), broadcast will happen.

Returns:
tvm.tensor.Tensor, subtracted result, with same type as input tensors and broadcasted shape of data1 and data2.
"""
vc_util.elemwise_dtype_check(data1.dtype, data2.dtype)
vc_util.check_shape(data1.shape)
vc_util.check_shape(data2.shape)
vc_util.auto_broadcast_check(data1.shape, data2.shape)

res = _akg.topi.subtract(data1, data2)

return res

+ 0
- 45
mindspore/_akg/ops/math/sum_value.py View File

@@ -1,45 +0,0 @@
# Copyright 2019 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.

"""operator dsl function: sum"""

import _akg.topi
import _akg.tvm
from _akg.utils import format_transform as ft_util
from _akg.utils import validation_check as vc_util


@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
def sum_value(inputs, axis=None, keepdims=False):
"""
Compute the sum of elements across dimensions of a tensor.

Args:
inputs (tvm.tensor.Tensor): Tensor.
axis (Union[list, tuple, int, None]): If the list or tuple is empty, the axis equal to None.
keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.

Returns:
tvm.tensor.Tensor, has same type as input. If keepdims is True, all reduced dimensions are retained
with length 1, else these reduced axis will be eliminate.
"""
axis = ft_util.refine_reduce_axis(inputs, axis)
vc_util.check_shape(inputs.shape)

if not axis:
output = _akg.topi.identity(inputs)
else:
output = _akg.topi.sum(inputs, axis=axis, keepdims=keepdims)

return output

+ 0
- 87
mindspore/_akg/save_gpu_param.py View File

@@ -1,87 +0,0 @@
# Copyright 2019 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.

"""save gpu param"""
import os
import hashlib
import _akg.tvm
from _akg.tvm import schedule
from _akg.utils import validation_check as vc_util


def get_dim(dim, axis=True):
"""get dim info"""
dims_str = {
"grid_dim0": "// attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = ",
"grid_dim1": "// attr [iter_var(blockIdx.y, , blockIdx.y)] thread_extent = ",
"grid_dim2": "// attr [iter_var(blockIdx.z, , blockIdx.z)] thread_extent = ",
"block_dim0": "// attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = ",
"block_dim1": "// attr [iter_var(threadIdx.y, , threadIdx.y)] thread_extent = ",
"block_dim2": "// attr [iter_var(threadIdx.z, , threadIdx.z)] thread_extent = "
}
dim_to_axis = {
"grid_dim0": '"blockIdx.x" : ',
"grid_dim1": '"blockIdx.y" : ',
"grid_dim2": '"blockIdx.z" : ',
"block_dim0": '"threadIdx.x" : ',
"block_dim1": '"threadIdx.y" : ',
"block_dim2": '"threadIdx.z" : '
}
if axis:
return dim_to_axis.get(dim)
return dims_str.get(dim)


def parse_params(file, dim, ir):
"""parse parameters"""
dim_str = get_dim(dim, axis=False)
pos = ir.find(dim_str)
if pos != -1:
index = pos + len(dim_str)
param_temp = get_dim(dim)

while ir[index].isdigit():
param_temp += ir[index]
index += 1
file.write(param_temp + ",\n")
else:
param_temp = get_dim(dim) + '1'
file.write(param_temp + ",\n")


@vc_util.check_input_type(schedule.Schedule, (list, tuple), tuple)
def save_gpu_params(s, args, kernel_info):
"""save gpu parameters"""
ptx_code = kernel_info[0]
file_name = kernel_info[1]
kernel_name = kernel_info[2]
ir = str(_akg.tvm.lower(s, args, simple_mode=True))
file_path = os.path.realpath(file_name)
if os.path.exists(file_path):
os.remove(file_path)

sha256 = hashlib.sha256()
sha256.update(ptx_code.encode("utf-8"))
hash_str = sha256.hexdigest()
with os.fdopen(os.open(file_path, os.O_WRONLY | os.O_CREAT, 0o400), 'w') as fo:
fo.write("{\n")
fo.write('"kernelName" : ' + '"' + kernel_name + "_kernel0" + '",\n')
parse_params(fo, "grid_dim0", ir)
parse_params(fo, "grid_dim1", ir)
parse_params(fo, "grid_dim2", ir)
parse_params(fo, "block_dim0", ir)
parse_params(fo, "block_dim1", ir)
parse_params(fo, "block_dim2", ir)
fo.write('"sha256" : ' + '"' + hash_str + '"\n')
fo.write("}\n")

+ 0
- 122
mindspore/_akg/utils/dsl_create.py View File

@@ -1,122 +0,0 @@
# Copyright 2019 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.

"""dsl create helping function"""
import _akg
from _akg.utils import format_transform as ft_util

class TensorUtils:
"""Class for creating tensor."""
CREATE_SCH_ONLY = 'create_sch_only'

@classmethod
def get_tensor_attrs(cls, tensor):
"""get tensor attrs."""
tensor_attrs = dict()
if "attrs" in dir(tensor.op):
tensor_attrs = dict(tensor.op.attrs.items())
return tensor_attrs

@classmethod
def update_tensor_attrs(cls, tensor, attrs):
"""update tensor attrs."""
tensor_attrs = cls.get_tensor_attrs(tensor)
tensor_attrs.update(attrs)
tensor = _akg.tvm.compute(tensor.shape,
lambda *indice: tensor[indice],
name=tensor.op.name,
tag=tensor.op.tag,
attrs=tensor_attrs)
return tensor

@classmethod
def is_create_sch_only(cls, tensor):
tensor_attrs = cls.get_tensor_attrs(tensor)
if cls.CREATE_SCH_ONLY in tensor_attrs.keys():
return True
return False

@classmethod
def is_output_value(cls, tensor):
"""check output value."""
return not cls.is_create_sch_only(tensor)

@classmethod
def inplace_set(cls, input_tensor, output_tensor, buffer_name="data_buf"):
"""inplace set."""
input_tensor_shape = ft_util.get_shape(input_tensor)
output_tensor_shape = ft_util.get_shape(output_tensor)
if not input_tensor_shape == output_tensor_shape:
raise RuntimeError("Shape of the input_tensor and the output_tensor should be equal, "
"but got %s and %s"%(input_tensor_shape, output_tensor_shape))
output_tensor = cls.update_tensor_attrs(output_tensor, {cls.CREATE_SCH_ONLY: 1})
data_buf = _akg.tvm.decl_buffer(input_tensor.shape, input_tensor.dtype, name=buffer_name)
binds_info = {input_tensor: data_buf, output_tensor: data_buf}
return output_tensor, binds_info

@classmethod
def inplace_set_tensors(cls, input_tensors, output_tensors, buffer_names=None):
"""
inplace set for tensors

Args:
in_tensors (Union[list, tuple]): Origin input tensors.
out_tensors (Union[list, tuple]): Origin output tensors.
buffer_names (Union[list, tuple] or None): Buffer names used to bind.

Return:
inplace_tensors (list): Output tensors with the inplace info.
binds_infos (dict): Dictionary that maps the input tensor and the output
tensor to buffer.
"""
if not buffer_names:
buffer_names = ["data_buf_%s" % i for i in range(len(input_tensors))]
for arg in (input_tensors, output_tensors, buffer_names):
if not isinstance(arg, (tuple, list)):
raise RuntimeError("arg must be tuple or list!")
if len(input_tensors) != len(output_tensors) or len(input_tensors) != len(buffer_names):
raise RuntimeError("length of the input_tensors, output_tensors and buffer_names must be equal!")

inplace_tensors = []
binds_infos = dict()
for input_tensor, output_tensor, buffer_name in zip(input_tensors, output_tensors, buffer_names):
inplace_tensor, binds_info = cls.inplace_set(input_tensor, output_tensor, buffer_name)
inplace_tensors.append(inplace_tensor)
binds_infos.update(binds_info)
return inplace_tensors, binds_infos

def produce_shapes(shape1, shape2):
"""two input shapes produce three output shape."""
shape1 = list(shape1)
shape2 = list(shape2)
flag = 0
if len(shape1) < len(shape2):
shape1, shape2 = shape2, shape1
flag = 1

output_shape_len = len(shape1)
dec = output_shape_len - len(shape2)
for i in range(dec):
shape2 = [1] + shape2

out_shape = []
for i in range(output_shape_len):
if (shape1[i] != shape2[i]) and (shape1[i] != 1) and (shape2[i] != 1):
raise RuntimeError("input shapes not match!")
out_shape.append(shape1[i] if shape1[i] > shape2[i] else shape2[i])

if flag == 1:
shape1, shape2 = shape2, shape1

return shape1, shape2, out_shape

+ 0
- 80
mindspore/_akg/utils/format_transform.py View File

@@ -1,80 +0,0 @@
# Copyright 2019 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.

"""format transform function"""
import _akg

def refine_reduce_axis(input_content, axis):
"""make reduce axis legal."""
shape = get_shape(input_content)
if axis is None:
axis = [i for i in range(len(shape))]
elif isinstance(axis, int):
axis = [axis]
elif not isinstance(axis, (tuple, list)):
raise TypeError("axis must be one of the type int,tuple,list or None")

if len(axis) > len(shape):
raise ValueError("axis size must not larger than shape size")

axis = list(axis)

for i, _ in enumerate(axis):
if axis[i] < 0:
axis[i] += len(shape)

if axis[i] >= len(shape):
raise ValueError(("axis value-{} exceeds len(axis) which is invalid".format(axis[i])))

axis.sort(reverse=True)

return axis


def get_shape_from_tensor(data):
"""translate _akg.tvm.shape to list type in python."""
tvm_shape = data.shape
py_shape = []
for i in tvm_shape:
if isinstance(i, _akg.tvm.expr.Var):
py_shape.append(i)
else:
py_shape.append(i.value)
return py_shape


def tvm_shape_to_list(tvm_shape):
"""translate _akg.tvm.shape to list type in python."""
py_shape = []
for i in tvm_shape:
if isinstance(i, _akg.tvm.expr.Var):
py_shape.append(i)
else:
py_shape.append(i.value)
return py_shape


def get_shape(data):
"""get shape and save it as list."""
if isinstance(data, _akg.tvm.tensor.Tensor):
shape = get_shape_from_tensor(data)
elif isinstance(data, _akg.tvm.container.Array):
shape = tvm_shape_to_list(data)
elif isinstance(data, int):
shape = [data]
elif isinstance(data, (tuple, list)):
shape = list(data)
else:
raise TypeError("Refine axis does not support type {} for now.".format(type(data)))
return shape

+ 0
- 233
mindspore/_akg/utils/validation_check.py View File

@@ -1,233 +0,0 @@
# Copyright 2019 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.

"""validation check functions"""
from functools import wraps, reduce
from _akg.utils.format_transform import get_shape

MAX_DATA_SIZE = 2 ** 31

def check_input_type_dict(input_dict, input_key, input_name):
"""
check input parameter type for new type: dict.

Note:
rule1: key of input_dict should be in the input_key
rule2: type of input_dict[shape] should be in (list, tuple), if have shape
rule3: type of input_dict[dtype] should be in (str), if have dtype

Args:
input_dict (dict): input_dict
input_key (list or tuple): all input key list, the key of input must in input_key
input_name (str): input param name, only used for error print

Returns:
None
"""
def _check_input_type(input_key, input_type):
if not isinstance(input_dict[input_key], input_type):
raise RuntimeError(
"the input parameter %s[%s] must be %s, while type of input is %s" %
(input_name, input_key, input_type, type(input_dict[input_key])))

for key in input_dict.keys():
if key not in input_key:
raise RuntimeError(
"the input parameter %s must have arrt <%s>" %
(input_name, key))

# check shape's type of input_dict, if have shape
if key == "shape":
_check_input_type(key, (list, tuple))

# check dtype's type of input_dict, if have dtype
if key == "dtype":
_check_input_type(key, (str,))


def check_input_type_list_tuple(inputs, expect):
"""check inputs by a list or tuple of expected types."""
if not isinstance(inputs, expect[1][0]):
raise RuntimeError("the input parameter %s must be (list, tuple), while"
" type of input is %s" % (expect[0], type(inputs)))
for inp in inputs:
if not isinstance(inp, expect[1][1]):
raise RuntimeError("The element in parameter %s must be %s, while "
"type of input is %s" % (
expect[0], expect[1][1], type(inp)))


def check_input_type(*type_args, **_type_kwargs):
"""check input parameter type."""
def out_wrapper(func):
"""outer wrapper function."""
formal_parameter = func.__code__.co_varnames
formal_parameter_list = list(zip(formal_parameter, type_args))

@wraps(func)
def in_wrapper(*args, **kwargs):
"""inner wrapper function."""
for i, arg_v in enumerate(args):
# add for new input dict, if dict, will check shape and dtype
if isinstance(arg_v, dict):
check_input_type_dict(arg_v, arg_v.keys(),
formal_parameter_list[i][0])

if isinstance(formal_parameter_list[i][1], tuple):
if isinstance(formal_parameter_list[i][1][0], tuple) \
and len(formal_parameter_list[i][1]) == 2:
check_input_type_list_tuple(arg_v, formal_parameter_list[i])
continue

if not isinstance(arg_v, formal_parameter_list[i][1]):
raise RuntimeError("the %sth input parameter %s must be %s, "
"while type of input is %s" % (str(i), formal_parameter_list[i][0],
formal_parameter_list[i][1],
type(arg_v)))
for i in kwargs:
for j in formal_parameter_list:
if i in j:
if not isinstance(kwargs[i], j[1]):
raise RuntimeError("the input parameter %s must be "
"%s, while type of input is %s"
"" % (i, j[1], type(kwargs[i])))
break
return func(*args, **kwargs)

return in_wrapper

return out_wrapper


def shape_dtype_max_size_check(shape):
"""check validation of tensor's shape."""
if shape:
mul = int(reduce(lambda x, y: int(x) * int(y), shape))
if mul > MAX_DATA_SIZE:
error_msg = "*".join([str(sh) for sh in shape])
raise RuntimeError("Invalid shape, data is {} bytes ({}), which "
"exceed max data size {} bytes"
.format(mul, error_msg, MAX_DATA_SIZE))


def check_shape(tensor, length=None, tensor_name=""):
"""The common check rule for placeholder data."""
shape = get_shape(tensor)
if not shape:
raise RuntimeError("The ndim of input tensor {} must more than 0, "
"actual input is {}".format(tensor_name, len(shape)))

for shape_v in shape:
if not isinstance(shape_v, int) or shape_v <= 0:
raise RuntimeError("The type of tensor {} axis value must be "
"positive int and value more than 0,"
"actual input is ({}) {}".
format(tensor_name, type(shape_v), shape_v))

if length and len(shape) != length:
raise ValueError('The length of {} should be {}, while actual length is {}'.
format(tensor_name, length, len(shape)))


def ops_dtype_check(dtype, args):
"""check validation of op's dtype."""
expected_dtype = list()

def _get_expect_dtype(expected_dtype, arg):
if isinstance(arg, str):
expected_dtype.append(arg)
elif isinstance(arg, (list, tuple)):
for t in arg:
_get_expect_dtype(expected_dtype, t)
else:
raise TypeError("arg should be either a string, "
"or a list/tuple of string, "
"while current is {}".format(type(arg)))

_get_expect_dtype(expected_dtype, args)

if isinstance(dtype, (list, tuple)):
checking_dtype = [d.lower() for d in dtype]
elif isinstance(dtype, str):
checking_dtype = [dtype.lower()]
else:
raise TypeError("dtype should be either a string or a tuple/list of string")
error_msg = "Supported dtype: {}, while received dtype: {}"
if not set(checking_dtype).issubset(set(expected_dtype)):
raise RuntimeError(error_msg.format(expected_dtype, checking_dtype))


def reduce_axis_check(reduce_shape, reduce_axis):
"""check validation of reduce axis for certain reduce shape."""
dim = len(reduce_shape)
if dim == 1 and int(reduce_shape[0]) == 1:
raise RuntimeError("Error, reduce shape is 1. Scalar is not supported "
"for reduction, please input a vector.")
if isinstance(reduce_axis, int):
if reduce_axis not in range(-dim, dim):
raise RuntimeError("Reduce axis should be in range [%d. %d)"
"" % (-dim, dim))
elif isinstance(reduce_axis, (tuple, list)):
if len(reduce_axis) > len(reduce_shape):
raise RuntimeError("Reduce axis list exceed reduce shape length: "
"%d vs %d, error" % (len(reduce_axis), len(reduce_shape)))
processed_axis = []
for axis in reduce_axis:
processed_axis.append(int(axis + dim) if axis < 0 else int(axis))
if len(set(processed_axis)) < len(processed_axis):
raise RuntimeError("Reduce axis list contains %d duplicated element, please check"
% (len(processed_axis) - len(set(processed_axis))))
for axis in processed_axis:
if axis >= dim:
raise RuntimeError("Invalid reduce axis, axis should less than %d" % dim)
elif reduce_axis is not None:
raise RuntimeError("axis should be a list, tuple or int.")


def elemwise_dtype_check(dtype_a, dtype_b, supported_type=None):
"""check validation of tensor's dtype for element-wise op."""
if supported_type:
ops_dtype_check(dtype_a, supported_type)
ops_dtype_check(dtype_b, supported_type)
if dtype_a.lower() != dtype_b.lower():
raise RuntimeError("Element-wise operation needs same data type, while "
"current is %s vs %s" % (dtype_a.lower(), dtype_b.lower()))


def auto_broadcast_check(shape_a, shape_b):
"""automatic broadcast check."""
shape_l = get_shape(shape_a)
shape_r = get_shape(shape_b)

if len(shape_l) <= len(shape_r):
shape_short = shape_l
shape_long = shape_r
else:
shape_short = shape_r
shape_long = shape_l

dim_diff = len(shape_long) - len(shape_short)
for i in range(dim_diff):
shape_short.insert(0, 1)
for i, shp in enumerate(shape_short):
if int(shp) != int(shape_long[i]) and 1 not in [int(shp), int(shape_long[i])]:
raise RuntimeError("Invalid auto broadcast, dim %d should be 1 or equal, "
"while now is %d vs %d" % (i, shp, shape_long[i]))


def check_int_list(array, array_name):
"""check whether all the elements are integers."""
for num in array:
if not isinstance(num, int):
raise RuntimeError("Type of value in %s should be int, but got type %s" % (array_name, type(num)))

+ 1
- 1
mindspore/_checkparam.py View File

@@ -672,7 +672,7 @@ def check_input_data(*data, data_class):

def check_output_data(data):
"""Output data check."""
if not data:
if data is None:
raise RuntimeError('Executor return data ' + str(data) + ', please check your net or input data.')




+ 19
- 0
mindspore/_extends/builtin_operations.py View File

@@ -14,7 +14,10 @@
# ============================================================================
"""builtin_operations"""
import numpy as np
from mindspore.ops import functional as F
from mindspore.ops import composite as C
from mindspore.common.tensor import Tensor
import mindspore.common.dtype as mstype
from mindspore.common.dtype import dtype_to_nptype, get_py_obj_dtype


@@ -113,6 +116,7 @@ def bool_or(x, y):
"""Implement `bool_or`."""
return x or y


def vm_compare(*args):
"""Implement `vm_compare` for tensor."""
obj_str = args[-1]
@@ -141,10 +145,12 @@ def list_len(x):
"""Implement `list_len`."""
return len(x)


def Depend(value, expr):
"""Implement `Depend`."""
return value


# only used in PyNative mode
def make_ref(key, value, ref):
return value
@@ -171,3 +177,16 @@ def tuple_to_array(x):
def stop_gradient(x):
"""Implement `stop_gradient`."""
return x


hyper_map = C.HyperMap()


def mixed_precision_cast(dst_type, x):
"""Implement `mixed_precision_cast`."""
def cast_inner(data):
if isinstance(data, Tensor) and data.dtype in (mstype.float32, mstype.float16):
return F.cast(data, dst_type)
return data

return hyper_map(cast_inner, x)

+ 5
- 0
mindspore/_extends/parallel_compile/akg_compiler/__init__.py View File

@@ -12,3 +12,8 @@
# See the License for the specific language governing permissions and
# limitations under the License.
# ============================================================================
"""
Extension functions.

Python functions that will be called in the c++ parts of MindSpore.
"""

+ 88
- 0
mindspore/_extends/parallel_compile/akg_compiler/akg_process.py View File

@@ -0,0 +1,88 @@
# 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.
# ============================================================================
"""akg process"""
import os
import subprocess
import sys
from multiprocessing import Pool, cpu_count

def _compile_akg_task(*json_strs):
"""
compile func called in single process

Parameters:
json_strs: list. List contains multiple kernel infos, suitable for json compile api.
"""
akg_compiler = os.path.join(os.path.split(
os.path.realpath(__file__))[0], "compiler.py")
for json_str in json_strs:
res = subprocess.run(
[sys.executable, akg_compiler, json_str], text=True)
if res.returncode != 0:
raise ValueError("Failed, args: {}!".format(json_str))

def create_akg_parallel_process(process_num, wait_time):
"""
create AkgParallelCompiler object

Returns:
AkgParallelCompiler
"""
return AkgProcess(process_num, wait_time)

class AkgProcess:
"""akg kernel parallel process"""

def __init__(self, process_num, wait_time):
"""
Args:
process_num: int. processes number
waittime: int. max time the function blocked
"""
if not isinstance(process_num, int):
raise ValueError("process number must be a num")
if not isinstance(wait_time, int):
raise ValueError("wait time must be a num")
if process_num == 0:
process_num = 1
max_proc_num = 16
self.process_num = min([cpu_count(), max_proc_num, process_num])
self.args = [[] for _ in range(self.process_num)]
self.wait_time = wait_time
self.argc = 0

def compile(self):
"""
compile kernel by multi processes
Return:
True for all compile success, False for some failed.
"""
if self.argc == 0:
raise ValueError("json must be not null")
with Pool(processes=self.process_num) as pool:
res = pool.starmap_async(_compile_akg_task, self.args)
res.get(timeout=self.wait_time)
return True

def accept_json(self, json):
"""
accept json data before compile
Args:
json: str. kernel info.
"""
if not isinstance(json, str):
raise ValueError("json must be a str")
self.args[self.argc % self.process_num].append(json)
self.argc += 1

+ 0
- 71
mindspore/_extends/parallel_compile/akg_compiler/multi_process_compiler.py View File

@@ -1,71 +0,0 @@
# 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.
# ============================================================================
"""Providing multi process compile with json"""
import os
import subprocess
import sys
from multiprocessing import Pool, cpu_count


def _compile_akg_task(*json_strs):
"""
compile func called in single process

Parameters:
json_strs: list. List contains multiple kernel infos, suitable for json compile api.
"""
akg_compiler = os.path.join(os.path.split(
os.path.realpath(__file__))[0], "compiler.py")
for json_str in json_strs:
res = subprocess.run(
[sys.executable, akg_compiler, json_str], text=True)
if res.returncode != 0:
raise ValueError("Failed, args: {}!".format(json_str))


def compile_akg_kernel_parallel(json_infos, process, waitime):
"""
compile kernel use multi processes

Parameters:
json_infos: list. list contain kernel info(task id and json str)
process: int. processes num
waittime: int. max time the function blocked

Returns:
True for all compile success, False for some failed.
"""
if not isinstance(json_infos, list):
raise ValueError("json_infos must be a list")
if not isinstance(process, int):
raise ValueError("process must be a num")
if not isinstance(waitime, int):
raise ValueError("waittime must be a num")

if process == 0 and json_infos:
process = 1

cpu_proc_num = cpu_count()
max_proc_num = 16
process = min([cpu_proc_num, max_proc_num, process])

args = [[] for _ in range(process)]
for p, info in enumerate(json_infos):
args[p % process].append(info)

with Pool(processes=process) as pool:
res = pool.starmap_async(_compile_akg_task, args)
res.get(timeout=waitime)
return True

+ 5
- 6
mindspore/_extends/parallel_compile/tbe_compiler/tbe_process.py View File

@@ -22,14 +22,14 @@ import json
from .common import check_kernel_info, TBEException
from .helper import _op_select_format, _check_supported

def create_tbe_parallel_compiler():
def create_tbe_parallel_process():
"""
create TBEParallelCompiler object

Returns:
TBEParallelCompiler
"""
return compile_pool
return tbe_process

def op_select_format(op_json: str):
"""
@@ -98,8 +98,8 @@ def run_compiler(op_json):
except subprocess.CalledProcessError as e:
return "TBEException", "PreCompileProcessFailed:\n" + e.stdout + "\n" + e.stderr + "\ninput_args: " + op_json

class CompilerPool:
"""compiler pool"""
class TbeProcess:
"""tbe process"""

def __init__(self):
self.__processe_num = multiprocessing.cpu_count()
@@ -168,5 +168,4 @@ class CompilerPool:
if self.__running_tasks:
self.__running_tasks.clear()


compile_pool = CompilerPool()
tbe_process = TbeProcess()

+ 2
- 2
mindspore/_extends/parse/__init__.py View File

@@ -21,12 +21,12 @@ from .parser import (Parser, create_obj_instance, generate_scope,
get_class_member_namespace_symbol, create_slice_obj,
get_dataclass_attributes, get_dataclass_methods, get_obj_id,
get_module_namespace, get_obj_type, get_object_key,
get_default_input, get_parse_method_of_class, get_scope_name,
get_parse_method_of_class, get_scope_name,
is_class_member, parse_cb, resolve_symbol)
from .serialize import *

__all__ = ['parse_cb', 'get_parse_method_of_class', 'get_bprop_method_of_class', 'resolve_symbol',
'get_object_key', 'get_default_input', 'get_class_instance_type', 'is_class_member',
'get_object_key', 'get_class_instance_type', 'is_class_member',
'get_obj_type', 'get_obj_id', 'create_obj_instance', 'get_module_namespace',
'get_class_member_namespace_symbol', 'get_obj_id', 'Parser', 'get_dataclass_attributes',
'get_dataclass_methods', 'dump_obj', 'load_obj', 'get_dataclass_methods', 'get_scope_name',


+ 8
- 1
mindspore/_extends/parse/namespace.py View File

@@ -99,12 +99,19 @@ class ClassMemberNamespace(Namespace):
obj (Object): A python class object.
"""
def __init__(self, obj):
self.__class_member_namespace__ = True
label = f'{obj.__module__}..<{obj.__class__.__name__}::{id(obj)}>'
super().__init__(label, obj)

def __getitem__(self, name):
d, = self.dicts
if name == "self":
return d
if name == "namespace":
return self
try:
return getattr(d, name)
if hasattr(d, name):
return getattr(d, name)
return d.__dict__[name]
except ValueError:
raise UnboundLocalError(name)

+ 27
- 9
mindspore/_extends/parse/parser.py View File

@@ -70,6 +70,7 @@ parse_expr_statement_white_list = (
"append",
)


def create_slice_obj(start, end, step):
"""Create slice object"""
return slice(start, end, step)
@@ -201,17 +202,9 @@ def get_object_key(obj):
if isinstance(obj, types.MethodType):
method_instance = obj.__self__
instance_id = "%s_ID%d" % (str(method_instance.__class__.__name__), id(method_instance))
obj_id = instance_id + obj_id
obj_id = instance_id + obj_id + str(obj.__hash__())
return obj_id, obj_key

def get_default_input(obj):
if hasattr(obj, '__parameter__'):
return obj.default_input
if isinstance(obj, tuple):
convert = lambda x: x.default_input if hasattr(x, '__parameter__') else x
args = tuple(convert(x) for x in obj)
return args
return obj

def is_class_member(node):
"""Check the attr is class member variable."""
@@ -224,10 +217,12 @@ def is_class_member(node):
return True
return False


def get_obj_id(obj):
"""Get the obj id."""
return str(id(obj))


def get_obj_type(obj):
"""Get the obj type."""
obj_type = RESOLVE_TYPE_INVALID
@@ -320,6 +315,7 @@ def get_dataclass_methods(cls):
if isinstance(getattr(cls, name), (types.FunctionType,))}
return methods


class Parser:
"""
Parser python code to ast tree.
@@ -453,6 +449,28 @@ class Parser:
logger.debug("ops info = %r", ops_info)
return ops_info

def analyze_super(self, class_type_node, subclass_instance):
"""Analyze super and return a class instance."""
sub_class = type(subclass_instance)
if class_type_node is None:
return super(sub_class, subclass_instance)
if isinstance(class_type_node, ast.Name):
class_name = getattr(class_type_node, 'id')
elif isinstance(class_type_node, ast.Attribute):
class_name = getattr(class_type_node, 'attr')
else:
raise ValueError(f"When call 'super', the first arg should be a class type, "
f"but got {class_type_node.__class__.__name__}.")

target_father_class = None
for class_element in sub_class.mro():
if class_element.__name__ == class_name:
target_father_class = class_element
break
if target_father_class is None:
raise ValueError("When call 'super', the second arg should be an instance of first arg.")
return super(target_father_class, subclass_instance)

def get_location(self, node):
"""
Get location of node start and end line no.


+ 4
- 2
mindspore/_extends/parse/resources.py View File

@@ -17,7 +17,7 @@
"""Resources for ast tree parse."""
import ast
import math
from mindspore import IndexedSlices
from mindspore import RowTensor, SparseTensor
from mindspore.ops.composite import multitype_ops
from mindspore.ops import functional as F, composite as C
from . import standard_method as M
@@ -117,6 +117,7 @@ convert_object_map = {
T.zip: C.zip_operation,
T.print: F.print_,
T.enumerate: M.enumerate_,
T.isinstance: M.isinstance_,

# custom define operation
T.iter: M.ms_iter,
@@ -139,5 +140,6 @@ convert_object_map = {
math.tan: NO_IMPLEMENT,

# user defined
IndexedSlices: F.make_indexed_slices,
RowTensor: F.make_row_tensor,
SparseTensor: F.make_sparse_tensor,
}

+ 57
- 7
mindspore/_extends/parse/standard_method.py View File

@@ -27,6 +27,42 @@ from ...ops.composite.base import _append
__all__ = ['MultitypeFuncGraph', 'env_get', 'hyper_add', 'zeros_like', 'ones_like']

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.

Args:
x (Tensor): A Tensor to be reduced.
axis (Union[None, int, tuple(int)): Dimensions of reduction.
keep_dims (bool): Whether to keep the reduced dimensions.

Returns:
Tensor, has the same data type as x.
"""

reduce_all = P.ReduceAll(keep_dims)
return reduce_all(x, axis)


def any_(x, axis=(), keep_dims=False):
"""
Check any array element along a given axis evaluate to True.

Args:
x (Tensor): A Tensor to be reduced.
axis (Union[None, int, tuple(int)): Dimensions of reduction.
keep_dims (bool): Whether to keep the reduced dimensions.

Returns:
Tensor, has the same data type as x.
"""

reduce_any = P.ReduceAny(keep_dims)
return reduce_any(x, axis)


def transpose(x):
@@ -114,6 +150,12 @@ def enumerate_(x, start=0):
return ret


def isinstance_(x, base_type):
"""Determine whether x is an instance of base_type."""
x_type = F.typeof(x)
return check_type_same(x_type, base_type)


def while_cond(x):
"""For while condtion, if the condition is a tensor, the loop will not be unrolled"""
if F.issubclass_(F.typeof(x), F.typeof(mstype.tensor)):
@@ -123,6 +165,14 @@ def while_cond(x):
return x


@constexpr
def check_type_same(x_type, base_type):
"""Check x_type is same as base_type."""
if mstype.issubclass_(x_type, base_type):
return True
raise TypeError(f"The arg 'x' should be a {base_type}, but got {x_type}.")


@constexpr
def check_is_tuple_or_list(x, op_name, arg_name):
"""check whether x is list or tuple."""
@@ -146,7 +196,8 @@ def check_is_tensor_bool_cond(shp):
"""check if tensor is a bool condition"""
if shp in ((), (1,)):
return True
raise ValueError("tensor as bool condition, its shape should be () or (1,), but got ", shp)
raise ValueError("The truth value of an array with several elements is ambiguous.")


@constexpr
def const_tensor_to_bool(x):
@@ -154,13 +205,12 @@ def const_tensor_to_bool(x):
if x is None:
raise ValueError("Only constant tensor bool can be converted to bool")
x = x.asnumpy()
if x.shape not in ((), (1,)):
raise ValueError("Tensor to bool should input shape () or (1), but got ", x.shape)
if x.shape == ():
value = bool(x)
else:
value = bool(x[0])
return value
return bool(x)
if x.shape == (1,):
return bool(x[0])
raise ValueError("The truth value of an array with several elements is ambiguous.")


def tensor_bool(x):
"""tensor as conditon, if is constant, return immediate bool value"""


+ 2
- 2
mindspore/_extends/parse/trope.py View File

@@ -27,7 +27,7 @@ from operator import ( # noqa

# support system function call
from builtins import ( # noqa
bool, getattr, setattr, len, iter, next, pow, range, map, zip, print, enumerate
bool, getattr, setattr, len, iter, next, pow, range, map, zip, print, enumerate, isinstance
)

# support functools
@@ -44,7 +44,7 @@ __all__ = ['add', 'sub', 'mul', 'truediv', 'floordiv', 'mod', 'eq', 'ne', 'lt',
'not_', 'and_', 'or_', 'xor', 'lshift', 'rshift', 'invert', 'is_', 'is_not', 'contains',
'matmul', 'getitem', 'setitem',
'bool', 'getattr', 'setattr', 'len', 'iter', 'next', 'pow', 'range', 'map', 'zip',
'partial', 'print', 'enumerate',
'partial', 'print', 'enumerate', 'isinstance',
'exp', 'log', 'sin', 'cos', 'tan']




+ 19
- 0
mindspore/_extends/remote/__init__.py View File

@@ -0,0 +1,19 @@
# 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.
# ============================================================================
"""
Server functions.

Python functions that will be called in the c++ client part of MindSpore.
"""

+ 174
- 0
mindspore/_extends/remote/kernel_build_server.py View File

@@ -0,0 +1,174 @@
# 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.
# ============================================================================
"""kernel build server"""
import os
import time

class Messager:
'''Messager'''

def __init__(self, fdin, fdout):
self.fdin = fdin
self.fdout = fdout
self.fin = os.fdopen(fdin, "r")
self.fout = os.fdopen(fdout, "w")
self.message = ''

def __del__(self):
os.close(self.fdin)
os.close(self.fdout)

def get_message(self):
"""
Get message from remote

Returns:
message
"""
try:
# Not read by input() anymore
res = self.fin.readline()
if not res:
logger.debug('[TRACE]', "read nothing...")
self.exit()
if res[len(res) - 1] == '\n':
res = res[0:len(res)-1]
self.message = res
logger.debug('[IN]', self.message)
except EOFError:
self.exit()
finally:
pass
if self.message == '' or self.message == 'FINISH':
self.send_ack()
self.exit()
return self.message

def send_res(self, res, keep_format=True):
"""
Send result to remote

Args:
keep_format: True or False
"""
logger.debug('[OUT]', str(res))
if keep_format:
res_str = str(res).replace('\n', '[LF]').replace('\r', '[CR]').replace(' ', '[SP]')
else:
res_str = str(res).replace('\n', '').replace('\r', '').replace(' ', '')
tag = '[~]' # The same as client kTAG

# Not write by print(tag + res_str, flush=True) any more
try:
self.fout.write(tag + res_str + "\n")
self.fout.flush()
except BrokenPipeError as err:
logger.info('[TRACE]', 'Write, ' + str(err))
self.exit()
finally:
pass

def send_ack(self, success=True):
"""
Send ack to remote

Args:
success: True or False
"""
if success:
self.send_res('ACK')
else:
self.send_res('ERR')

def loop(self):
"""
Messaging loop
"""
while True:
self.handle()

def run(self):
self.loop()

def handle(self):
"""
A interface communicates with remote.

Note:
All subclasses should override this interface.
"""
raise NotImplementedError

def exit(self):
"""
A interface handles the procedure before exit.

Note:
All subclasses should override this interface.
"""
raise NotImplementedError

class Logger:
"""
Replace dummy 'logger' to output log as below:
logger = Logger(0, True, "remote_kernel_build_" + time.strftime("%Y_%m_%d_%H_%M_%S", time.localtime()) + ".log")
"""
def __init__(self, level=1, dumpfile=False, filename='Logger.log'):
"""
Args:
level: 0 for debug and info, 1 for info
dumpfile: if dump log into file
"""
self.level = level
self.dumpfile = dumpfile
if self.dumpfile:
self.log = open(filename, "a")

def write(self, msg):
self.log.write(msg)
self.flush()

def writeline(self, tag, msg):
prefix = tag + ' REMOTE(' + str(os.getpid()) + ',python)'
line = prefix + '\t' + time.strftime("%Y-%m-%d %H:%M:%S", time.localtime()) + ':\t' + msg
print(line, flush=True)
if self.dumpfile:
self.write(line + '\n')

def debug(self, tag, msg):
if self.level == 0:
self.writeline('[DEBUG]' + tag, msg)

def info(self, tag, msg):
self.writeline('[INFO]' + tag, msg)

def flush(self):
self.log.flush()

class DummyLogger:
"""DummyLogger"""
def __init__(self):
pass

def debug(self, tag, msg):
pass

def info(self, tag, msg):
pass

logger = DummyLogger()

def get_logger():
return logger

+ 148
- 0
mindspore/_extends/remote/kernel_build_server_ascend.py View File

@@ -0,0 +1,148 @@
# 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.
# ============================================================================
"""kernel build server for ascend"""
import sys
from mindspore._extends.remote.kernel_build_server import Messager, get_logger
from mindspore._extends.parallel_compile.tbe_compiler.tbe_process import create_tbe_parallel_process, op_select_format, check_supported
from mindspore._extends.parallel_compile.akg_compiler.akg_process import create_akg_parallel_process

class TbeBuilder:
"""Tbe building wrapper"""

def __init__(self):
self.tbe_builder = create_tbe_parallel_process()

def start(self, json):
return self.tbe_builder.start_compile_op(json)

def wait(self):
return self.tbe_builder.wait_one()

def reset(self):
self.tbe_builder.reset_task_info()

def exit(self):
self.tbe_builder.exit()

class AkgBuilder:
"""Akg building wrapper"""

def __init__(self):
pass

def create(self, process_num, waitime):
self.akg_builder = create_akg_parallel_process(process_num, waitime)

def accept_json(self, json):
return self.akg_builder.accept_json(json)

def compile(self):
return self.akg_builder.compile()

class AscendMessager(Messager):
'''
Ascend Messager
It works as a server, communicating with c++ client.
'''

def __init__(self, fdin, fdout):
super().__init__(fdin, fdout)
get_logger().info('[TRACE]', 'Ascend Messager init...')
self.tbe_builder = TbeBuilder()
self.akg_builder = AkgBuilder()

def handle(self):
"""
Communicate with remote client.
Reference protocol between them at PR#3821 and PR#3935
"""
arg = self.get_message()
if arg == 'TBE/START':
self.send_ack()
json = self.get_message()
res = self.tbe_builder.start(json)
self.send_res(res)
elif arg == 'TBE/WAIT':
self.send_ack()
task_id, res, pre = self.tbe_builder.wait()
get_logger().debug('[TRACE]', str(task_id) + '/' + str(res) + '/' + str(pre))
if self.get_message() != 'CONTINUE':
self.send_ack(False)
self.exit()
self.send_res(task_id)
if self.get_message() != 'CONTINUE':
self.send_ack(False)
self.exit()
self.send_res(res)
if self.get_message() != 'CONTINUE':
self.send_ack(False)
self.exit()
self.send_res(pre)
elif arg == 'TBE/RESET':
self.tbe_builder.reset()
self.send_ack()
elif arg == 'AKG/START':
self.send_ack()
process_num_str = self.get_message()
self.send_ack()
wait_time_str = self.get_message()
self.akg_builder.create(int(process_num_str), int(wait_time_str))
self.send_ack()
elif arg == 'AKG/DATA':
self.send_ack()
while True:
req = self.get_message()
if req.startswith('{'):
self.akg_builder.accept_json(req)
self.send_ack()
elif req == 'AKG/WAIT':
res = self.akg_builder.compile()
self.send_res(res)
break
else:
self.send_ack(False)
break
elif arg == 'FORMAT':
self.send_ack()
json = self.get_message()
self.send_res(op_select_format(json))
elif arg == 'SUPPORT':
self.send_ack()
json = self.get_message()
get_logger().debug('[SUPPORT]', json)
try:
res = check_supported(json)
except json.decoder.JSONDecodeError:
self.send_ack(False)
self.exit()
finally:
pass
self.send_res(res)
else:
self.send_ack(False)
self.exit()

def exit(self):
self.tbe_builder.reset()
self.tbe_builder.exit()
get_logger().info('[TRACE]', 'Ascend Messager Exit...')
exit()

if __name__ == '__main__':
if len(sys.argv) != 3:
raise Exception('Incorrect argv: {}'.format(sys.argv))
get_logger().debug('[TRACE]', 'argv: ' + str(sys.argv))
messager = AscendMessager(int(sys.argv[1]), int(sys.argv[2]))
messager.run()

+ 63
- 0
mindspore/_extends/remote/kernel_build_server_gpu.py View File

@@ -0,0 +1,63 @@
# 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.
# ============================================================================
"""kernel build server for gpu"""
import os
import sys
from mindspore._extends.remote.kernel_build_server import Messager, get_logger
from mindspore._extends.parallel_compile.akg_compiler.compiler import run_compiler as akg_compile_single

class GpuMessager(Messager):
'''
GPU Messager
It works as a server, communicating with c++ client.
'''

def __init__(self, fdin, fdout):
super().__init__(fdin, fdout)
get_logger().info('[TRACE]', 'GPU Messager init...')

def handle(self):
"""
Communicate with remote client.
Reference protocol between them at PR#4063
"""
arg = self.get_message()
if arg == 'AKG/PID':
self.send_res(os.getpid())
elif arg == 'AKG/COMPILE':
self.send_ack()
json = self.get_message()
try:
akg_compile_single(json)
except ValueError:
self.send_ack(False)
self.exit()
finally:
pass
self.send_ack()
else:
self.send_ack(False)
self.exit()

def exit(self):
get_logger().info('[TRACE]', 'GPU Messager Exit...')
exit()

if __name__ == '__main__':
if len(sys.argv) != 3:
raise Exception('Incorrect argv: {}'.format(sys.argv))
get_logger().debug('[TRACE]', 'argv: ' + str(sys.argv))
messager = GpuMessager(int(sys.argv[1]), int(sys.argv[2]))
messager.run()

+ 19
- 22
mindspore/ccsrc/CMakeLists.txt View File

@@ -44,7 +44,7 @@ if(ENABLE_GPU)
"backend/kernel_compiler/akg/akg_kernel_attrs_process.cc"
)

list(APPEND CUDA_NVCC_FLAGS -arch=sm_53)
list(APPEND CUDA_NVCC_FLAGS -arch=sm_53 --expt-relaxed-constexpr)
list(REMOVE_ITEM GPU_SRC_LIST "runtime/device/gpu/blocking_queue.cc" "runtime/device/gpu/gpu_buffer_mgr.cc")
list(REMOVE_ITEM GPU_SRC_LIST "runtime/device/gpu/mpi/mpi_initializer.cc"
"runtime/device/gpu/distribution/collective_wrapper.cc"
@@ -60,11 +60,6 @@ if(ENABLE_GPU)
add_compile_definitions(ENABLE_GPU)
endif ()

## make flatuffer files
include_directories("${CMAKE_BINARY_DIR}/predict/schema/inner")
file(GLOB_RECURSE FLATBUFFER_IN RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "predict/schema/*.fbs")
set(FLATBUFFER_OU "${CMAKE_BINARY_DIR}/predict/schema/inner")
ms_build_flatbuffers("${FLATBUFFER_IN}" "${FLATBUFFER_IN}" flat_input "${FLATBUFFER_OU}")

## make protobuf files
file(COPY "${ms_onnx_INC}/onnx/onnx.proto" DESTINATION ${CMAKE_BINARY_DIR}/proto)
@@ -104,13 +99,9 @@ endif ()

if (ENABLE_D)
include_directories("${CMAKE_BINARY_DIR}/backend/kernel_compiler/aicpu")
include_directories("${CMAKE_BINARY_DIR}/predict/generator/ir")
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_INNER RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "predict/proto/*.proto")
ms_protobuf_generate(PREDICT_PROTOSRCS PREDICT_PROTOHDRS ${PROTO_INNER})

file(GLOB_RECURSE PROTO_DUMP RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "runtime/device/ascend/dump/proto/*.proto")
ms_protobuf_generate(DUMP_PROTOSRCS PROTOHDRS ${PROTO_DUMP})

@@ -139,7 +130,7 @@ set(SUB_COMP
frontend/operator
pipeline/jit
pipeline/pynative
common debug gvar predict pybind_api utils vm
common debug gvar pybind_api utils vm
)

foreach (_comp ${SUB_COMP})
@@ -147,16 +138,18 @@ foreach (_comp ${SUB_COMP})
string(REPLACE "/" "_" sub ${_comp})
if (TARGET _mindspore_${sub}_obj)
list(APPEND SUB_OBJECTS_SRC $<TARGET_OBJECTS:_mindspore_${sub}_obj>)
add_dependencies(_mindspore_${sub}_obj proto_input flat_input)
add_dependencies(_mindspore_${sub}_obj proto_input )
endif ()
endforeach ()
add_subdirectory(${CMAKE_SOURCE_DIR}/mindspore/core/base base)
list(APPEND SUB_OBJECTS_SRC $<TARGET_OBJECTS:_mindspore_base_obj>)
add_subdirectory(${CMAKE_SOURCE_DIR}/mindspore/core/abstract abstract)
list(APPEND SUB_OBJECTS_SRC $<TARGET_OBJECTS:_mindspore_abstract_obj>)
add_subdirectory(${CMAKE_SOURCE_DIR}/mindspore/core/utils util)
list(APPEND SUB_OBJECTS_SRC $<TARGET_OBJECTS:_mindspore_core_utils_obj>)
add_subdirectory(${CMAKE_SOURCE_DIR}/mindspore/core/ir ir)
list(APPEND SUB_OBJECTS_SRC $<TARGET_OBJECTS:_mindspore_ir_obj>)
add_dependencies(_mindspore_base_obj _mindspore_ir_obj _mindspore_abstract_obj proto_input flat_input)
add_dependencies(_mindspore_core_utils_obj _mindspore_base_obj _mindspore_ir_obj _mindspore_abstract_obj proto_input )

set_property(SOURCE ${SUB_OBJECTS_SRC} PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_ME)
add_library(mindspore STATIC ${SUB_OBJECTS_SRC})
@@ -169,7 +162,7 @@ if (ENABLE_DEBUGGER)
endif()

target_link_libraries(mindspore proto_input)
if (ENABLE_MPI)
if (ENABLE_MPI AND ENABLE_CPU)
target_link_libraries(mindspore securec mindspore::flatbuffers mpi_adapter)
else ()
target_link_libraries(mindspore securec mindspore::flatbuffers)
@@ -252,15 +245,15 @@ if (CMAKE_SYSTEM_NAME MATCHES "Windows")
target_link_libraries(mindspore mindspore_gvar)
target_link_libraries(_c_expression PRIVATE -Wl,--whole-archive mindspore -Wl,--no-whole-archive)
else ()
target_link_libraries(_c_expression PRIVATE -Wl,--whole-archive mindspore -Wl,--no-whole-archive)
target_link_libraries(_c_expression PRIVATE mindspore::pybind11_module)
target_link_libraries(_c_expression PRIVATE mindspore_gvar)
if (NOT ENABLE_GE)
target_link_libraries(_c_expression PRIVATE mindspore::pslite mindspore::protobuf ${zeromq_DIRPATH}/zmq_install/lib/libzmq.a)
if (ENABLE_CPU AND (ENABLE_D OR ENABLE_GPU))
target_link_libraries(mindspore mindspore::pslite mindspore::protobuf ${zeromq_DIRPATH}/zmq_install/lib/libzmq.a)
if (${ENABLE_IBVERBS} STREQUAL "ON")
target_link_libraries(_c_expression PRIVATE ibverbs rdmacm)
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 mindspore::pybind11_module)
target_link_libraries(_c_expression PRIVATE mindspore_gvar)
endif ()

if (USE_GLOG)
@@ -278,7 +271,11 @@ if (ENABLE_GPU)
${CUDA_PATH}/lib64/libcurand.so
${CUDNN_PATH}/lib64/libcudnn.so
${CUDA_PATH}/lib64/libcudart.so
${CUDA_PATH}/lib64/stubs/libcuda.so)
${CUDA_PATH}/lib64/stubs/libcuda.so
${CUDA_PATH}/lib64/libcusolver.so)
if (ENABLE_MPI)
set_target_properties(_ms_mpi PROPERTIES INSTALL_RPATH ${ORIGIN_PATH})
endif()
endif ()

if (ENABLE_CPU)
@@ -296,7 +293,7 @@ set(LOAD_ONNX_SRC
${CMAKE_CURRENT_SOURCE_DIR}/utils/load_onnx/anf_model_parser.cc
)
add_library(inference SHARED
${CMAKE_CURRENT_SOURCE_DIR}/backend/session/session.cc
${CMAKE_CURRENT_SOURCE_DIR}/backend/session/infer_session.cc
${LOAD_ONNX_SRC}
)
target_link_libraries(inference PRIVATE ${PYTHON_LIBRARIES} ${SECUREC_LIBRARY}


+ 12
- 8
mindspore/ccsrc/backend/kernel_compiler/CMakeLists.txt View File

@@ -26,14 +26,6 @@ if (ENABLE_CPU)
"cpu/*.cc"
)

list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/push_kernel.cc"
"cpu/ps/pull_kernel.cc"
"cpu/ps/embedding_look_up_ps_kernel.cc"
"cpu/ps/embedding_look_up_proxy_kernel.cc"
"cpu/ps/apply_momentum_ps_kernel.cc"
"cpu/ps/sparse_apply_adam_ps_kernel.cc"
"cpu/ps/sparse_apply_ftrl_ps_kernel.cc")

if (NOT ENABLE_MPI)
list(REMOVE_ITEM CPU_SRC_LIST "cpu/allgather_cpu_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/reduce_scatter_cpu_kernel.cc")
@@ -41,6 +33,18 @@ if (ENABLE_CPU)
endif ()
endif ()

if (NOT (ENABLE_CPU AND (ENABLE_D OR ENABLE_GPU)))
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/apply_momentum_ps_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/embedding_look_up_proxy_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/embedding_look_up_ps_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/pserver_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/pull_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/push_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/sparse_apply_adam_ps_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/sparse_apply_ftrl_ps_kernel.cc")
list(REMOVE_ITEM CPU_SRC_LIST "cpu/ps/sparse_apply_lazy_adam_ps_kernel.cc")
endif()

if (ENABLE_GPU)
file(GLOB_RECURSE CUDA_SRC_LIST RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}
"gpu/*.cu"


+ 1
- 1
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_build.cc View File

@@ -30,7 +30,7 @@
#include "proto/attr.pb.h"
#include "proto/node_def.pb.h"
#include "backend/session/anf_runtime_algorithm.h"
#include "common/utils.h"
#include "utils/ms_utils.h"
#include "backend/kernel_compiler/aicpu/aicpu_util.h"
#include "backend/session/kernel_graph.h"
#include "backend/kernel_compiler/common_utils.h"


+ 3
- 3
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_build.h View File

@@ -13,8 +13,8 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_BUILD_H_
#define MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_BUILD_H_
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_BUILD_H_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_BUILD_H_
#include <memory>
#include "backend/kernel_compiler/kernel.h"

@@ -24,4 +24,4 @@ KernelModPtr AicpuOpBuild(const std::shared_ptr<AnfNode> &anf_node);
} // namespace kernel
} // namespace mindspore

#endif // MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_BUILD_H_
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_BUILD_H_

+ 3
- 3
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_metadata.h View File

@@ -14,8 +14,8 @@
* limitations under the License.
*/

#ifndef MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_META_DATA_H_
#define MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_META_DATA_H_
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_META_DATA_H_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_META_DATA_H_

#include <string>
#include <vector>
@@ -27,4 +27,4 @@ namespace kernel {
void AicpuMetadataInfo(const CNodePtr &kernel_node, std::vector<std::shared_ptr<KernelBuildInfo>> *kernel_info_list);
} // namespace kernel
} // namespace mindspore
#endif // MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_META_DATA_H_
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_META_DATA_H_

+ 1
- 1
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_mod.cc View File

@@ -26,7 +26,7 @@
#include "backend/kernel_compiler/aicpu/aicpu_kernel_build.h"
#include "utils/convert_utils.h"
#include "backend/kernel_compiler/aicpu/aicpu_util.h"
#include "utils/context/ms_context.h"
#include "utils/ms_context.h"

using AicpuTaskInfoPtr = std::shared_ptr<ge::model_runner::AicpuTaskInfo>;



+ 3
- 3
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_kernel_mod.h View File

@@ -13,8 +13,8 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_MOD_H_
#define MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_MOD_H_
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_MOD_H_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_MOD_H_
#include <vector>
#include <memory>
#include <string>
@@ -72,4 +72,4 @@ using AicputOpKernelModPtrList = std::vector<AicpuOpKernelModPtr>;
} // namespace kernel
} // namespace mindspore

#endif // MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_KERNEL_MOD_H_
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_KERNEL_MOD_H_

+ 3
- 3
mindspore/ccsrc/backend/kernel_compiler/aicpu/aicpu_util.h View File

@@ -13,8 +13,8 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_UTIL_H_
#define MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_UTIL_H_
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_UTIL_H_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_UTIL_H_

#include <cstdint>
#include <vector>
@@ -61,4 +61,4 @@ class AicpuOpUtil {
} // namespace kernel
} // namespace mindspore

#endif // MINDSPORE_MINDSPORE_CCSRC_KERNEL_AICPU_AICPU_UTIL_H_
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AICPU_AICPU_UTIL_H_

+ 2
- 9
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_attrs_process.cc View File

@@ -18,6 +18,7 @@
#include <algorithm>
#include "backend/session/anf_runtime_algorithm.h"
#include "backend/optimizer/common/helper.h"
#include "backend/kernel_compiler/common_utils.h"

namespace mindspore {
namespace kernel {
@@ -75,15 +76,7 @@ void SetAkgAttrsForCast(const AnfNodePtr &anf_node) {

std::string dst_type;
TypeId output_type = AnfAlgo::GetOutputDeviceDataType(anf_node, 0);
if (output_type == kFloat32->type_id()) {
dst_type = "float32";
} else if (output_type == kFloat16->type_id()) {
dst_type = "float16";
} else if (output_type == kInt32->type_id()) {
dst_type = "int32";
} else {
MS_LOG(WARNING) << "Unknown cast_to type: " << TypeIdToType(output_type)->ToString();
}
dst_type = TypeId2String(output_type);
AnfAlgo::SetNodeAttr("dst_type", MakeValue(dst_type), anf_node);
}



+ 3
- 3
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_attrs_process.h View File

@@ -13,8 +13,8 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_AKG_AKG_KERNEL_ATTRS_PROCESS_H
#define MINDSPORE_CCSRC_KERNEL_AKG_AKG_KERNEL_ATTRS_PROCESS_H
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AKG_AKG_KERNEL_ATTRS_PROCESS_H
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AKG_AKG_KERNEL_ATTRS_PROCESS_H

#include <vector>
#include <memory>
@@ -55,4 +55,4 @@ const std::unordered_map<std::string, std::function<void(const AnfNodePtr &anf_n
};
} // namespace kernel
} // namespace mindspore
#endif // MINDSPORE_CCSRC_KERNEL_AKG_AKG_KERNEL_ATTRS_PROCESS_H
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_AKG_AKG_KERNEL_ATTRS_PROCESS_H

Some files were not shown because too many files changed in this diff

Loading…
Cancel
Save