From c2b4bad0403d8c0c778933490aafdcbb97166147 Mon Sep 17 00:00:00 2001 From: looop5 Date: Mon, 3 Aug 2020 14:35:44 +0800 Subject: [PATCH] remove redundant compilation of libtvm.so and mindspore/_akg directory --- akg | 2 +- cmake/external_libs/dlpack.cmake | 7 - cmake/external_libs/dmlc_core.cmake | 7 - cmake/external_libs/rang.cmake | 7 - cmake/external_libs/tvm_gpu.cmake | 15 -- cmake/mind_expression.cmake | 5 - cmake/package.cmake | 23 -- mindspore/_akg/__init__.py | 18 -- mindspore/_akg/add_path.py | 62 ----- mindspore/_akg/gpu/__init__.py | 39 --- mindspore/_akg/gpu/cast.py | 45 ---- mindspore/_akg/gpu/default_schedule.py | 56 ----- mindspore/_akg/gpu/equal.py | 40 --- mindspore/_akg/gpu/greater_equal.py | 41 --- mindspore/_akg/gpu/hsigmoid.py | 63 ----- mindspore/_akg/gpu/hsigmoid_grad.py | 51 ---- mindspore/_akg/gpu/hswish.py | 63 ----- mindspore/_akg/gpu/hswish_grad.py | 53 ---- mindspore/_akg/gpu/less_equal.py | 40 --- mindspore/_akg/gpu/logical_and.py | 40 --- mindspore/_akg/gpu/logical_not.py | 40 --- mindspore/_akg/gpu/logical_or.py | 40 --- mindspore/_akg/gpu/mean.py | 80 ------ mindspore/_akg/gpu/mean_grad.py | 90 ------- mindspore/_akg/gpu/mul.py | 41 --- mindspore/_akg/gpu/notequal.py | 41 --- mindspore/_akg/gpu/relu6.py | 54 ---- mindspore/_akg/gpu/relu6_grad.py | 59 ----- mindspore/_akg/gpu/squeeze.py | 50 ---- mindspore/_akg/gpu/squeeze_grad.py | 44 ---- mindspore/_akg/gpu/sub.py | 40 --- mindspore/_akg/gpu/tile.py | 39 --- mindspore/_akg/message.py | 104 -------- mindspore/_akg/op_build.py | 69 ------ mindspore/_akg/ops/__init__.py | 0 mindspore/_akg/ops/array/__init__.py | 0 mindspore/_akg/ops/array/tile.py | 36 --- mindspore/_akg/ops/math/__init__.py | 0 mindspore/_akg/ops/math/cast.py | 36 --- mindspore/_akg/ops/math/equal.py | 54 ---- mindspore/_akg/ops/math/greater_equal.py | 54 ---- mindspore/_akg/ops/math/less_equal.py | 54 ---- mindspore/_akg/ops/math/logical_and.py | 41 --- mindspore/_akg/ops/math/logical_not.py | 32 --- mindspore/_akg/ops/math/logical_or.py | 41 --- mindspore/_akg/ops/math/mean.py | 47 ---- mindspore/_akg/ops/math/mul.py | 43 ---- mindspore/_akg/ops/math/notequal.py | 54 ---- mindspore/_akg/ops/math/sub.py | 40 --- mindspore/_akg/ops/math/sum_value.py | 45 ---- mindspore/_akg/save_gpu_param.py | 87 ------- mindspore/_akg/utils/__init__.py | 0 mindspore/_akg/utils/dsl_create.py | 122 --------- mindspore/_akg/utils/format_transform.py | 80 ------ mindspore/_akg/utils/validation_check.py | 233 ------------------ setup.py | 2 +- third_party/apply_patches.sh | 91 ------- .../patch/incubator-tvm/CMakeLists.txt | 100 -------- third_party/patch/incubator-tvm/cmake.patch | 201 --------------- .../patch/incubator-tvm/find_library.patch | 71 ------ third_party/patch/incubator-tvm/include.patch | 55 ----- .../patch/incubator-tvm/src_pass.patch | 120 --------- 62 files changed, 2 insertions(+), 3205 deletions(-) delete mode 100644 cmake/external_libs/dlpack.cmake delete mode 100644 cmake/external_libs/dmlc_core.cmake delete mode 100644 cmake/external_libs/rang.cmake delete mode 100644 cmake/external_libs/tvm_gpu.cmake delete mode 100644 mindspore/_akg/__init__.py delete mode 100644 mindspore/_akg/add_path.py delete mode 100644 mindspore/_akg/gpu/__init__.py delete mode 100644 mindspore/_akg/gpu/cast.py delete mode 100644 mindspore/_akg/gpu/default_schedule.py delete mode 100644 mindspore/_akg/gpu/equal.py delete mode 100644 mindspore/_akg/gpu/greater_equal.py delete mode 100644 mindspore/_akg/gpu/hsigmoid.py delete mode 100644 mindspore/_akg/gpu/hsigmoid_grad.py delete mode 100644 mindspore/_akg/gpu/hswish.py delete mode 100644 mindspore/_akg/gpu/hswish_grad.py delete mode 100644 mindspore/_akg/gpu/less_equal.py delete mode 100644 mindspore/_akg/gpu/logical_and.py delete mode 100644 mindspore/_akg/gpu/logical_not.py delete mode 100644 mindspore/_akg/gpu/logical_or.py delete mode 100644 mindspore/_akg/gpu/mean.py delete mode 100644 mindspore/_akg/gpu/mean_grad.py delete mode 100644 mindspore/_akg/gpu/mul.py delete mode 100644 mindspore/_akg/gpu/notequal.py delete mode 100644 mindspore/_akg/gpu/relu6.py delete mode 100644 mindspore/_akg/gpu/relu6_grad.py delete mode 100644 mindspore/_akg/gpu/squeeze.py delete mode 100644 mindspore/_akg/gpu/squeeze_grad.py delete mode 100644 mindspore/_akg/gpu/sub.py delete mode 100644 mindspore/_akg/gpu/tile.py delete mode 100644 mindspore/_akg/message.py delete mode 100644 mindspore/_akg/op_build.py delete mode 100644 mindspore/_akg/ops/__init__.py delete mode 100644 mindspore/_akg/ops/array/__init__.py delete mode 100644 mindspore/_akg/ops/array/tile.py delete mode 100644 mindspore/_akg/ops/math/__init__.py delete mode 100644 mindspore/_akg/ops/math/cast.py delete mode 100644 mindspore/_akg/ops/math/equal.py delete mode 100644 mindspore/_akg/ops/math/greater_equal.py delete mode 100644 mindspore/_akg/ops/math/less_equal.py delete mode 100644 mindspore/_akg/ops/math/logical_and.py delete mode 100644 mindspore/_akg/ops/math/logical_not.py delete mode 100644 mindspore/_akg/ops/math/logical_or.py delete mode 100644 mindspore/_akg/ops/math/mean.py delete mode 100644 mindspore/_akg/ops/math/mul.py delete mode 100644 mindspore/_akg/ops/math/notequal.py delete mode 100644 mindspore/_akg/ops/math/sub.py delete mode 100644 mindspore/_akg/ops/math/sum_value.py delete mode 100644 mindspore/_akg/save_gpu_param.py delete mode 100644 mindspore/_akg/utils/__init__.py delete mode 100644 mindspore/_akg/utils/dsl_create.py delete mode 100644 mindspore/_akg/utils/format_transform.py delete mode 100644 mindspore/_akg/utils/validation_check.py delete mode 100755 third_party/apply_patches.sh delete mode 100644 third_party/patch/incubator-tvm/CMakeLists.txt delete mode 100644 third_party/patch/incubator-tvm/cmake.patch delete mode 100644 third_party/patch/incubator-tvm/find_library.patch delete mode 100644 third_party/patch/incubator-tvm/include.patch delete mode 100644 third_party/patch/incubator-tvm/src_pass.patch diff --git a/akg b/akg index 949a45538cc..5fe7e5c8377 160000 --- a/akg +++ b/akg @@ -1 +1 @@ -Subproject commit 949a45538ccb7ae94ad73386b5e3e77005112eea +Subproject commit 5fe7e5c8377dccfd35c9f661e10ed3dc136208c5 diff --git a/cmake/external_libs/dlpack.cmake b/cmake/external_libs/dlpack.cmake deleted file mode 100644 index a2375c7d350..00000000000 --- a/cmake/external_libs/dlpack.cmake +++ /dev/null @@ -1,7 +0,0 @@ -mindspore_add_pkg(dlpack - VER 0.2 - HEAD_ONLY ./ - URL https://github.com/dmlc/dlpack/archive/0acb731e0e43d15deee27b66f10e4c5b4e667913.zip - MD5 6b8093f17ad4e830d3c63eb3171c4b45) - - diff --git a/cmake/external_libs/dmlc_core.cmake b/cmake/external_libs/dmlc_core.cmake deleted file mode 100644 index e07df83fd6e..00000000000 --- a/cmake/external_libs/dmlc_core.cmake +++ /dev/null @@ -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) - - diff --git a/cmake/external_libs/rang.cmake b/cmake/external_libs/rang.cmake deleted file mode 100644 index 45ea375cb56..00000000000 --- a/cmake/external_libs/rang.cmake +++ /dev/null @@ -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) - - diff --git a/cmake/external_libs/tvm_gpu.cmake b/cmake/external_libs/tvm_gpu.cmake deleted file mode 100644 index 834e2d159df..00000000000 --- a/cmake/external_libs/tvm_gpu.cmake +++ /dev/null @@ -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) \ No newline at end of file diff --git a/cmake/mind_expression.cmake b/cmake/mind_expression.cmake index e7aa6ac3f4a..4d88bea4ac4 100644 --- a/cmake/mind_expression.cmake +++ b/cmake/mind_expression.cmake @@ -46,11 +46,6 @@ if (ENABLE_CPU) 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() diff --git a/cmake/package.cmake b/cmake/package.cmake index 1b0f349f159..3dbe3e21456 100644 --- a/cmake/package.cmake +++ b/cmake/package.cmake @@ -230,29 +230,6 @@ install( 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 OR ENABLE_GPU) AND ENABLE_AKG) set (AKG_PATH ${CMAKE_SOURCE_DIR}/build/mindspore/akg) install( diff --git a/mindspore/_akg/__init__.py b/mindspore/_akg/__init__.py deleted file mode 100644 index d0c1f0ffe46..00000000000 --- a/mindspore/_akg/__init__.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/add_path.py b/mindspore/_akg/add_path.py deleted file mode 100644 index d1e50f81771..00000000000 --- a/mindspore/_akg/add_path.py +++ /dev/null @@ -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()) diff --git a/mindspore/_akg/gpu/__init__.py b/mindspore/_akg/gpu/__init__.py deleted file mode 100644 index 4c114995947..00000000000 --- a/mindspore/_akg/gpu/__init__.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/cast.py b/mindspore/_akg/gpu/cast.py deleted file mode 100644 index dd7517f0a5c..00000000000 --- a/mindspore/_akg/gpu/cast.py +++ /dev/null @@ -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. - -"""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.""" - if x.dtype == "int64" and dst_type == "float16": - x = cast.cast(x, "float32") - 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 diff --git a/mindspore/_akg/gpu/default_schedule.py b/mindspore/_akg/gpu/default_schedule.py deleted file mode 100644 index 811cc2d7108..00000000000 --- a/mindspore/_akg/gpu/default_schedule.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/equal.py b/mindspore/_akg/gpu/equal.py deleted file mode 100644 index 3321c10b2c7..00000000000 --- a/mindspore/_akg/gpu/equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/greater_equal.py b/mindspore/_akg/gpu/greater_equal.py deleted file mode 100644 index 0212cac03c2..00000000000 --- a/mindspore/_akg/gpu/greater_equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/hsigmoid.py b/mindspore/_akg/gpu/hsigmoid.py deleted file mode 100644 index b313c2fd5a2..00000000000 --- a/mindspore/_akg/gpu/hsigmoid.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/hsigmoid_grad.py b/mindspore/_akg/gpu/hsigmoid_grad.py deleted file mode 100644 index bdde4ed3caf..00000000000 --- a/mindspore/_akg/gpu/hsigmoid_grad.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/hswish.py b/mindspore/_akg/gpu/hswish.py deleted file mode 100644 index 3def3c4b353..00000000000 --- a/mindspore/_akg/gpu/hswish.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/hswish_grad.py b/mindspore/_akg/gpu/hswish_grad.py deleted file mode 100644 index cadbf0f663d..00000000000 --- a/mindspore/_akg/gpu/hswish_grad.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/less_equal.py b/mindspore/_akg/gpu/less_equal.py deleted file mode 100644 index c58346e929c..00000000000 --- a/mindspore/_akg/gpu/less_equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/logical_and.py b/mindspore/_akg/gpu/logical_and.py deleted file mode 100644 index 6453901458d..00000000000 --- a/mindspore/_akg/gpu/logical_and.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/logical_not.py b/mindspore/_akg/gpu/logical_not.py deleted file mode 100644 index 0a381071872..00000000000 --- a/mindspore/_akg/gpu/logical_not.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/logical_or.py b/mindspore/_akg/gpu/logical_or.py deleted file mode 100644 index 1bd49bedbca..00000000000 --- a/mindspore/_akg/gpu/logical_or.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/mean.py b/mindspore/_akg/gpu/mean.py deleted file mode 100644 index e9cdb6d5517..00000000000 --- a/mindspore/_akg/gpu/mean.py +++ /dev/null @@ -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) diff --git a/mindspore/_akg/gpu/mean_grad.py b/mindspore/_akg/gpu/mean_grad.py deleted file mode 100644 index 9d91ee3f403..00000000000 --- a/mindspore/_akg/gpu/mean_grad.py +++ /dev/null @@ -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) diff --git a/mindspore/_akg/gpu/mul.py b/mindspore/_akg/gpu/mul.py deleted file mode 100644 index 5c289a62a6f..00000000000 --- a/mindspore/_akg/gpu/mul.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/notequal.py b/mindspore/_akg/gpu/notequal.py deleted file mode 100644 index 3e3a6561a1b..00000000000 --- a/mindspore/_akg/gpu/notequal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/relu6.py b/mindspore/_akg/gpu/relu6.py deleted file mode 100644 index 9a0a3d7a451..00000000000 --- a/mindspore/_akg/gpu/relu6.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/relu6_grad.py b/mindspore/_akg/gpu/relu6_grad.py deleted file mode 100644 index 62aeabb4c07..00000000000 --- a/mindspore/_akg/gpu/relu6_grad.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/squeeze.py b/mindspore/_akg/gpu/squeeze.py deleted file mode 100644 index b5f55facaac..00000000000 --- a/mindspore/_akg/gpu/squeeze.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/squeeze_grad.py b/mindspore/_akg/gpu/squeeze_grad.py deleted file mode 100644 index ae31de8e84d..00000000000 --- a/mindspore/_akg/gpu/squeeze_grad.py +++ /dev/null @@ -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) diff --git a/mindspore/_akg/gpu/sub.py b/mindspore/_akg/gpu/sub.py deleted file mode 100644 index 611e4228fd3..00000000000 --- a/mindspore/_akg/gpu/sub.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/gpu/tile.py b/mindspore/_akg/gpu/tile.py deleted file mode 100644 index 1eb6979b099..00000000000 --- a/mindspore/_akg/gpu/tile.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/message.py b/mindspore/_akg/message.py deleted file mode 100644 index 3d1f81f914b..00000000000 --- a/mindspore/_akg/message.py +++ /dev/null @@ -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']) diff --git a/mindspore/_akg/op_build.py b/mindspore/_akg/op_build.py deleted file mode 100644 index 92101f657ec..00000000000 --- a/mindspore/_akg/op_build.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/__init__.py b/mindspore/_akg/ops/__init__.py deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/mindspore/_akg/ops/array/__init__.py b/mindspore/_akg/ops/array/__init__.py deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/mindspore/_akg/ops/array/tile.py b/mindspore/_akg/ops/array/tile.py deleted file mode 100644 index 2fa485ea36c..00000000000 --- a/mindspore/_akg/ops/array/tile.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/__init__.py b/mindspore/_akg/ops/math/__init__.py deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/mindspore/_akg/ops/math/cast.py b/mindspore/_akg/ops/math/cast.py deleted file mode 100644 index 78140bfe271..00000000000 --- a/mindspore/_akg/ops/math/cast.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/equal.py b/mindspore/_akg/ops/math/equal.py deleted file mode 100644 index 2dbb1ba7333..00000000000 --- a/mindspore/_akg/ops/math/equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/greater_equal.py b/mindspore/_akg/ops/math/greater_equal.py deleted file mode 100644 index 00ad0166435..00000000000 --- a/mindspore/_akg/ops/math/greater_equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/less_equal.py b/mindspore/_akg/ops/math/less_equal.py deleted file mode 100644 index 5a566fbbcab..00000000000 --- a/mindspore/_akg/ops/math/less_equal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/logical_and.py b/mindspore/_akg/ops/math/logical_and.py deleted file mode 100644 index 480d4e17416..00000000000 --- a/mindspore/_akg/ops/math/logical_and.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/logical_not.py b/mindspore/_akg/ops/math/logical_not.py deleted file mode 100644 index 9befe7e8163..00000000000 --- a/mindspore/_akg/ops/math/logical_not.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/logical_or.py b/mindspore/_akg/ops/math/logical_or.py deleted file mode 100644 index 8fb0b80567c..00000000000 --- a/mindspore/_akg/ops/math/logical_or.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/mean.py b/mindspore/_akg/ops/math/mean.py deleted file mode 100644 index e8300f22fcb..00000000000 --- a/mindspore/_akg/ops/math/mean.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/mul.py b/mindspore/_akg/ops/math/mul.py deleted file mode 100644 index a690089da2d..00000000000 --- a/mindspore/_akg/ops/math/mul.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/notequal.py b/mindspore/_akg/ops/math/notequal.py deleted file mode 100644 index 16d5e4a0f43..00000000000 --- a/mindspore/_akg/ops/math/notequal.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/sub.py b/mindspore/_akg/ops/math/sub.py deleted file mode 100644 index 6ae2ee51efc..00000000000 --- a/mindspore/_akg/ops/math/sub.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/ops/math/sum_value.py b/mindspore/_akg/ops/math/sum_value.py deleted file mode 100644 index b9720469a60..00000000000 --- a/mindspore/_akg/ops/math/sum_value.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/save_gpu_param.py b/mindspore/_akg/save_gpu_param.py deleted file mode 100644 index ed2c9fe23ab..00000000000 --- a/mindspore/_akg/save_gpu_param.py +++ /dev/null @@ -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") diff --git a/mindspore/_akg/utils/__init__.py b/mindspore/_akg/utils/__init__.py deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/mindspore/_akg/utils/dsl_create.py b/mindspore/_akg/utils/dsl_create.py deleted file mode 100644 index 9d27039b28a..00000000000 --- a/mindspore/_akg/utils/dsl_create.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/utils/format_transform.py b/mindspore/_akg/utils/format_transform.py deleted file mode 100644 index c7a69b26cc1..00000000000 --- a/mindspore/_akg/utils/format_transform.py +++ /dev/null @@ -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 diff --git a/mindspore/_akg/utils/validation_check.py b/mindspore/_akg/utils/validation_check.py deleted file mode 100644 index 1231b3110ed..00000000000 --- a/mindspore/_akg/utils/validation_check.py +++ /dev/null @@ -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))) diff --git a/setup.py b/setup.py index 46a6e04a4e7..2836a24c31c 100644 --- a/setup.py +++ b/setup.py @@ -146,7 +146,7 @@ class BuildPy(build_py): super().run() mindspore_dir = os.path.join(pkg_dir, 'build', 'lib', 'mindspore') update_permissions(mindspore_dir) - mindspore_dir = os.path.join(pkg_dir, 'build', 'lib', '_akg') + mindspore_dir = os.path.join(pkg_dir, 'build', 'lib', 'akg') update_permissions(mindspore_dir) diff --git a/third_party/apply_patches.sh b/third_party/apply_patches.sh deleted file mode 100755 index fbd06b68b67..00000000000 --- a/third_party/apply_patches.sh +++ /dev/null @@ -1,91 +0,0 @@ -#!/bin/bash -# 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. -# ============================================================================ - -PWD_PATH=`pwd` -THIRD_PARTY_PATH=$(cd "$(dirname $0)"; pwd) -if [ $# -lt 1 ]; then - echo "Usage: sh apply_patches.sh [build_dir]" - echo " build_dir is the directory where you type \"cmake\"" - echo " Open source software incubator-tvm will be copied to build_dir" - echo " where patches will be applied on." - exit 1 -fi -BUILD_PATH=$1 - -if [ -d ${BUILD_PATH}/incubator-tvm ]; then - rm -rf ${BUILD_PATH}/incubator-tvm -fi -DLPACK_PATH=$2 -DMLC_PATH=$3 -RANG_PATH=$4 -TVM_PATH=$5 -mkdir ${BUILD_PATH}/incubator-tvm -cp -rf ${TVM_PATH}/* ${BUILD_PATH}/incubator-tvm/ -cp -rf ${DLPACK_PATH}/* ${BUILD_PATH}/incubator-tvm/3rdparty/dlpack/ -cp -rf ${DMLC_PATH}/* ${BUILD_PATH}/incubator-tvm/3rdparty/dmlc-core/ -cp -rf ${RANG_PATH}/* ${BUILD_PATH}/incubator-tvm/3rdparty/rang/ - -check_dir_not_empty() -{ - if [ ! $# -eq 1 ]; then - echo "Usage: check_dir_not_empty dir_path" - exit 1 - fi - - if [ ! -d $1 ]; then - echo "Directory $1 does not exist." - exit 1 - fi - - fileCounts=`ls $1 | wc -l` - if [ ${fileCounts} -eq 0 ]; then - echo "Directory $1 is empty." - exit 1 - fi -} - -apply_patch() -{ - if [ ! $# -eq 1 ]; then - echo "Usage: apply_patch patch_name" - exit 1 - fi - - if [ ! -f $1 ]; then - echo "Patch $1 does not exist." - exit 1 - fi - - patch -p1 < $1 - if [ $? -eq 0 ]; then - echo "Patch $1 applied successfully." - else - echo "Patch $1 not applied." - fi -} - -# apply patches on tvm -TVM_PATH=${BUILD_PATH}/incubator-tvm -TVM_PATCH_PATH=${THIRD_PARTY_PATH}/patch/incubator-tvm -check_dir_not_empty "${TVM_PATH}" -check_dir_not_empty "${TVM_PATCH_PATH}" -cd ${TVM_PATH} -apply_patch "${TVM_PATCH_PATH}/cmake.patch" -apply_patch "${TVM_PATCH_PATH}/find_library.patch" -apply_patch "${TVM_PATCH_PATH}/include.patch" -apply_patch "${TVM_PATCH_PATH}/src_pass.patch" - -cd ${PWD_PATH} diff --git a/third_party/patch/incubator-tvm/CMakeLists.txt b/third_party/patch/incubator-tvm/CMakeLists.txt deleted file mode 100644 index d8964579cdc..00000000000 --- a/third_party/patch/incubator-tvm/CMakeLists.txt +++ /dev/null @@ -1,100 +0,0 @@ -cmake_minimum_required(VERSION 3.2) -project(tvm C CXX) -set(TVM_DIR ${CMAKE_CURRENT_SOURCE_DIR}) -# Utility functions -include(${TVM_DIR}/cmake/util/Util.cmake) -include(${TVM_DIR}/cmake/util/FindCUDA.cmake) - -# include directories -include_directories(AFTER "${TVM_DIR}/include") -include_directories(AFTER "${TVM_DIR}/src") -include_directories(AFTER "${TVM_DIR}") -include_directories(AFTER "${TVM_DIR}/src/schedule") - -include_directories(AFTER "${TVM_DIR}/3rdparty/dmlc-core/include") -include_directories(AFTER "${TVM_DIR}/3rdparty/dlpack/include") -include_directories(AFTER "${TVM_DIR}/3rdparty/compiler-rt") -include_directories(AFTER "${TVM_DIR}/3rdparty/rang/include") - -# lib contain dlopen and dlclose -set(TVM_RUNTIME_LINKER_LIBS ${CMAKE_DL_LIBS}) - -# add source group -file(GLOB_RECURSE GROUP_SOURCE "${TVM_DIR}/src/*.cc" "src/*.cc") -file(GLOB_RECURSE GROUP_INCLUDE "${TVM_DIR}/src/*.h" - "${TVM_DIR}/include/*.h" "src/*.h" "include/*.h") -assign_source_group("Source" ${GROUP_SOURCE}) -assign_source_group("Include" ${GROUP_INCLUDE}) - -file(GLOB COMPILER_SRCS - "pre_activate/gpu/*.cc" - ${TVM_DIR}/src/api/*.cc - ${TVM_DIR}/src/arithmetic/*.cc - ${TVM_DIR}/src/autotvm/*.cc - ${TVM_DIR}/src/codegen/*.cc - ${TVM_DIR}/src/lang/*.cc - ${TVM_DIR}/src/pass/*.cc - ${TVM_DIR}/src/op/*.cc - ${TVM_DIR}/src/node/*.cc - ${TVM_DIR}/src/schedule/*.cc - ${TVM_DIR}/src/runtime/*.cc - ${TVM_DIR}/src/runtime/vm/*.cc - ${TVM_DIR}/src/runtime/vm/profiler/*.cc - ${TVM_DIR}/src/codegen/stackvm/*.cc) - -file(GLOB_RECURSE RELAY_SRCS ${TVM_DIR}/src/relay/*.cc) -list(APPEND COMPILER_SRCS ${RELAY_SRCS}) - -file(GLOB DATATYPE_SRCS ${TVM_DIR}/src/codegen/datatype/*.cc) -list(APPEND COMPILER_SRCS ${DATATYPE_SRCS}) - -file(GLOB COMPILER_VERILOG_SRCS ${TVM_DIR}/src/codegen/verilog/*.cc) -list(APPEND COMPILER_SRCS ${COMPILER_VERILOG_SRCS}) - -file(GLOB TOPI_SRCS ${TVM_DIR}/topi/src/*.cc) - -file(GLOB RUNTIME_SRCS - ${TVM_DIR}/src/runtime/*.cc - ${TVM_DIR}/src/runtime/vm/*.cc - ${TVM_DIR}/src/runtime/stub/*.cc - ${TVM_DIR}/src/runtime/stackvm/*.cc) - - -file(GLOB COMPILER_OFF_SRCS - ${TVM_DIR}/src/codegen/opt/build_*_off.cc) - -list(REMOVE_ITEM COMPILER_OFF_SRCS - ${TVM_DIR}/src/codegen/opt/build_cuda_off.cc) -set(USE_CUDA "ON") -list(APPEND COMPILER_SRCS ${COMPILER_OFF_SRCS}) -# Module rules -include(${TVM_DIR}/cmake/modules/CUDA.cmake) - -set(CMAKE_C_FLAGS_AKG -pipe -Wall -fPIC -fstack-protector-all) -set(CMAKE_C_FLAGS_AKG ${CMAKE_C_FLAGS_AKG} -Wl,-z,relro,-z,now,-z,noexecstack) - -set(CMAKE_CXX_FLAGS_AKG -std=c++11 -pipe -Wall -fPIC -fstack-protector-all) -set(CMAKE_CXX_FLAGS_AKG ${CMAKE_CXX_FLAGS_AKG} -Wl,-z,relro,-z,now,-z,noexecstack) - -if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug") - message("-- Build in Debug mode") - set(CMAKE_C_FLAGS_AKG ${CMAKE_C_FLAGS_AKG} -O0 -g -rdynamic) - set(CMAKE_CXX_FLAGS_AKG ${CMAKE_CXX_FLAGS_AKG} -O0 -g -rdynamic) -else() - message("-- Build in Release mode") - set(CMAKE_C_FLAGS_AKG ${CMAKE_C_FLAGS_AKG} -O2 -Werror) - set(CMAKE_CXX_FLAGS_AKG ${CMAKE_CXX_FLAGS_AKG} -O2 -Werror) -endif() -if(CMAKE_CXX_COMPILER_ID MATCHES "GNU" AND CMAKE_CXX_COMPILER_VERSION - VERSION_GREATER 7.0) - set(CMAKE_CXX_FLAGS_AKG ${CMAKE_CXX_FLAGS_AKG} -faligned-new) -endif() - -add_library(tvm SHARED ${COMPILER_SRCS} ${RUNTIME_SRCS} ${TOPI_SRCS}) - -target_link_libraries(tvm ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) -target_compile_options(tvm PRIVATE - $<$:${CMAKE_C_FLAGS_AKG}> - $<$:${CMAKE_CXX_FLAGS_AKG}>) -target_include_directories(tvm PRIVATE "${TVM_DIR}/topi/include") -install(TARGETS tvm) \ No newline at end of file diff --git a/third_party/patch/incubator-tvm/cmake.patch b/third_party/patch/incubator-tvm/cmake.patch deleted file mode 100644 index 820c7e24fd8..00000000000 --- a/third_party/patch/incubator-tvm/cmake.patch +++ /dev/null @@ -1,201 +0,0 @@ -diff -Npur tvm/cmake/modules/ANTLR.cmake tvm_new/cmake/modules/ANTLR.cmake ---- tvm/cmake/modules/ANTLR.cmake 2019-12-14 15:11:37.562418441 +0800 -+++ tvm_new/cmake/modules/ANTLR.cmake 2019-12-14 11:28:49.161977599 +0800 -@@ -14,12 +14,15 @@ - # KIND, either express or implied. See the License for the - # specific language governing permissions and limitations - # under the License. -+ -+# 2019.12.30 - Modify current directory of tvm. -+ - if(USE_ANTLR) - find_antlr(${USE_ANTLR}) - if(ANTLR4) - - set(RELAY_PARSER_DIR -- ${CMAKE_CURRENT_SOURCE_DIR}/python/tvm/relay/grammar) -+ ${TVM_DIR}/python/tvm/relay/grammar) - - set(RELAY_PARSER - ${RELAY_PARSER_DIR}/py3/RelayVisitor.py -diff -Npur tvm/cmake/modules/CUDA.cmake tvm_new/cmake/modules/CUDA.cmake ---- tvm/cmake/modules/CUDA.cmake 2019-12-14 15:11:37.562418441 +0800 -+++ tvm_new/cmake/modules/CUDA.cmake 2019-12-14 11:28:49.161977599 +0800 -@@ -15,6 +15,8 @@ - # specific language governing permissions and limitations - # under the License. - -+# 2019.12.30 - Modify current directory of tvm. -+ - # CUDA Module - find_cuda(${USE_CUDA}) - -@@ -29,9 +31,9 @@ if(USE_CUDA) - message(FATAL_ERROR "Cannot find CUDA, USE_CUDA=" ${USE_CUDA}) - endif() - message(STATUS "Build with CUDA support") -- file(GLOB RUNTIME_CUDA_SRCS src/runtime/cuda/*.cc) -+ file(GLOB RUNTIME_CUDA_SRCS ${TVM_DIR}/src/runtime/cuda/*.cc) - list(APPEND RUNTIME_SRCS ${RUNTIME_CUDA_SRCS}) -- list(APPEND COMPILER_SRCS src/codegen/opt/build_cuda_on.cc) -+ list(APPEND COMPILER_SRCS ${TVM_DIR}/src/codegen/opt/build_cuda_on.cc) - - list(APPEND TVM_LINKER_LIBS ${CUDA_NVRTC_LIBRARY}) - list(APPEND TVM_RUNTIME_LINKER_LIBS ${CUDA_CUDART_LIBRARY}) -@@ -40,18 +42,18 @@ if(USE_CUDA) - - if(USE_CUDNN) - message(STATUS "Build with cuDNN support") -- file(GLOB CONTRIB_CUDNN_SRCS src/runtime/contrib/cudnn/*.cc) -+ file(GLOB CONTRIB_CUDNN_SRCS ${TVM_DIR}/src/runtime/contrib/cudnn/*.cc) - list(APPEND RUNTIME_SRCS ${CONTRIB_CUDNN_SRCS}) - list(APPEND TVM_RUNTIME_LINKER_LIBS ${CUDA_CUDNN_LIBRARY}) - endif(USE_CUDNN) - - if(USE_CUBLAS) - message(STATUS "Build with cuBLAS support") -- file(GLOB CONTRIB_CUBLAS_SRCS src/runtime/contrib/cublas/*.cc) -+ file(GLOB CONTRIB_CUBLAS_SRCS ${TVM_DIR}/src/runtime/contrib/cublas/*.cc) - list(APPEND RUNTIME_SRCS ${CONTRIB_CUBLAS_SRCS}) - list(APPEND TVM_RUNTIME_LINKER_LIBS ${CUDA_CUBLAS_LIBRARY}) - endif(USE_CUBLAS) - - else(USE_CUDA) -- list(APPEND COMPILER_SRCS src/codegen/opt/build_cuda_off.cc) -+ list(APPEND COMPILER_SRCS ${TVM_DIR}/src/codegen/opt/build_cuda_off.cc) - endif(USE_CUDA) -diff -Npur tvm/cmake/modules/LLVM.cmake tvm_new/cmake/modules/LLVM.cmake ---- tvm/cmake/modules/LLVM.cmake 2019-12-14 15:11:37.562418441 +0800 -+++ tvm_new/cmake/modules/LLVM.cmake 2019-12-14 11:28:49.161977599 +0800 -@@ -15,6 +15,8 @@ - # specific language governing permissions and limitations - # under the License. - -+# 2019.12.30 - Modify current directory of tvm. -+ - # LLVM rules - add_definitions(-DDMLC_USE_FOPEN64=0) - -@@ -26,7 +28,7 @@ if(NOT USE_LLVM STREQUAL "OFF") - message(STATUS "Set TVM_LLVM_VERSION=" ${TVM_LLVM_VERSION}) - # Set flags that are only needed for LLVM target - add_definitions(-DTVM_LLVM_VERSION=${TVM_LLVM_VERSION}) -- file(GLOB COMPILER_LLVM_SRCS src/codegen/llvm/*.cc) -+ file(GLOB COMPILER_LLVM_SRCS ${TVM_DIR}/src/codegen/llvm/*.cc) - list(APPEND TVM_LINKER_LIBS ${LLVM_LIBS}) - list(APPEND COMPILER_SRCS ${COMPILER_LLVM_SRCS}) - if(NOT MSVC) -diff -Npur tvm/cmake/modules/Micro.cmake tvm_new/cmake/modules/Micro.cmake ---- tvm/cmake/modules/Micro.cmake 2019-12-14 15:11:37.562418441 +0800 -+++ tvm_new/cmake/modules/Micro.cmake 2019-12-14 11:28:49.161977599 +0800 -@@ -15,8 +15,10 @@ - # specific language governing permissions and limitations - # under the License. - -+# 2019.12.30 - Modify current directory of tvm. -+ - if(USE_MICRO) - message(STATUS "Build with Micro support") -- file(GLOB RUNTIME_MICRO_SRCS src/runtime/micro/*.cc) -+ file(GLOB RUNTIME_MICRO_SRCS ${TVM_DIR}/src/runtime/micro/*.cc) - list(APPEND RUNTIME_SRCS ${RUNTIME_MICRO_SRCS}) - endif(USE_MICRO) -diff -Npur tvm/cmake/modules/VTA.cmake tvm_new/cmake/modules/VTA.cmake ---- tvm/cmake/modules/VTA.cmake 2019-12-14 15:11:37.562418441 +0800 -+++ tvm_new/cmake/modules/VTA.cmake 2019-12-14 14:42:32.358381133 +0800 -@@ -15,17 +15,19 @@ - # specific language governing permissions and limitations - # under the License. - -+# 2019.12.30 - Modify current directory of tvm. -+ - # CMake Build rules for VTA - find_program(PYTHON NAMES python python3 python3.6) - - if(MSVC) - message(STATUS "VTA build is skipped in Windows..") - elseif(PYTHON) -- set(VTA_CONFIG ${PYTHON} ${CMAKE_CURRENT_SOURCE_DIR}/vta/config/vta_config.py) -+ set(VTA_CONFIG ${PYTHON} ${TVM_DIR}/vta/config/vta_config.py) - - if(EXISTS ${CMAKE_CURRENT_BINARY_DIR}/vta_config.json) - message(STATUS "Use VTA config " ${CMAKE_CURRENT_BINARY_DIR}/vta_config.json) -- set(VTA_CONFIG ${PYTHON} ${CMAKE_CURRENT_SOURCE_DIR}/vta/config/vta_config.py -+ set(VTA_CONFIG ${PYTHON} ${TVM_DIR}/vta/config/vta_config.py - --use-cfg=${CMAKE_CURRENT_BINARY_DIR}/vta_config.json) - endif() - -@@ -40,18 +42,18 @@ elseif(PYTHON) - # Fast simulator driver build - if(USE_VTA_FSIM) - # Add fsim driver sources -- file(GLOB FSIM_RUNTIME_SRCS vta/src/*.cc) -- list(APPEND FSIM_RUNTIME_SRCS vta/src/sim/sim_driver.cc) -- list(APPEND FSIM_RUNTIME_SRCS vta/src/vmem/virtual_memory.cc vta/src/vmem/virtual_memory.h) -- list(APPEND FSIM_RUNTIME_SRCS vta/src/sim/sim_tlpp.cc) -+ file(GLOB FSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/*.cc) -+ list(APPEND FSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/sim/sim_driver.cc) -+ list(APPEND FSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/vmem/virtual_memory.cc ${TVM_DIR}/vta/src/vmem/virtual_memory.h) -+ list(APPEND FSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/sim/sim_tlpp.cc) - # Target lib: vta_fsim - add_library(vta_fsim SHARED ${FSIM_RUNTIME_SRCS}) -- target_include_directories(vta_fsim PUBLIC vta/include) -+ target_include_directories(vta_fsim PUBLIC ${TVM_DIR}/vta/include) - foreach(__def ${VTA_DEFINITIONS}) - string(SUBSTRING ${__def} 3 -1 __strip_def) - target_compile_definitions(vta_fsim PUBLIC ${__strip_def}) - endforeach() -- include_directories("vta/include") -+ include_directories("${TVM_DIR}/vta/include") - if(APPLE) - set_target_properties(vta_fsim PROPERTIES LINK_FLAGS "-undefined dynamic_lookup") - endif(APPLE) -@@ -61,18 +63,18 @@ elseif(PYTHON) - # Cycle accurate simulator driver build - if(USE_VTA_TSIM) - # Add tsim driver sources -- file(GLOB TSIM_RUNTIME_SRCS vta/src/*.cc) -- list(APPEND TSIM_RUNTIME_SRCS vta/src/tsim/tsim_driver.cc) -- list(APPEND TSIM_RUNTIME_SRCS vta/src/dpi/module.cc) -- list(APPEND TSIM_RUNTIME_SRCS vta/src/vmem/virtual_memory.cc vta/src/vmem/virtual_memory.h) -+ file(GLOB TSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/*.cc) -+ list(APPEND TSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/tsim/tsim_driver.cc) -+ list(APPEND TSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/dpi/module.cc) -+ list(APPEND TSIM_RUNTIME_SRCS ${TVM_DIR}/vta/src/vmem/virtual_memory.cc ${TVM_DIR}/vta/src/vmem/virtual_memory.h) - # Target lib: vta_tsim - add_library(vta_tsim SHARED ${TSIM_RUNTIME_SRCS}) -- target_include_directories(vta_tsim PUBLIC vta/include) -+ target_include_directories(vta_tsim PUBLIC ${TVM_DIR}/vta/include) - foreach(__def ${VTA_DEFINITIONS}) - string(SUBSTRING ${__def} 3 -1 __strip_def) - target_compile_definitions(vta_tsim PUBLIC ${__strip_def}) - endforeach() -- include_directories("vta/include") -+ include_directories("${TVM_DIR}/vta/include") - if(APPLE) - set_target_properties(vta_tsim PROPERTIES LINK_FLAGS "-undefined dynamic_lookup") - endif(APPLE) -@@ -80,19 +82,19 @@ elseif(PYTHON) - - # VTA FPGA driver sources - if(USE_VTA_FPGA) -- file(GLOB FPGA_RUNTIME_SRCS vta/src/*.cc) -+ file(GLOB FPGA_RUNTIME_SRCS ${TVM_DIR}/vta/src/*.cc) - # Rules for Zynq-class FPGAs with pynq OS support (see pynq.io) - if(${VTA_TARGET} STREQUAL "pynq" OR - ${VTA_TARGET} STREQUAL "ultra96") -- list(APPEND FPGA_RUNTIME_SRCS vta/src/pynq/pynq_driver.cc) -+ list(APPEND FPGA_RUNTIME_SRCS ${TVM_DIR}/vta/src/pynq/pynq_driver.cc) - # Rules for Pynq v2.4 - find_library(__cma_lib NAMES cma PATH /usr/lib) - elseif(${VTA_TARGET} STREQUAL "de10nano") # DE10-Nano rules -- file(GLOB FPGA_RUNTIME_SRCS vta/src/de10nano/*.cc vta/src/*.cc) -+ file(GLOB FPGA_RUNTIME_SRCS ${TVM_DIR}/vta/src/de10nano/*.cc ${TVM_DIR}/vta/src/*.cc) - endif() - # Target lib: vta - add_library(vta SHARED ${FPGA_RUNTIME_SRCS}) -- target_include_directories(vta PUBLIC vta/include) -+ target_include_directories(vta PUBLIC ${TVM_DIR}/vta/include) - foreach(__def ${VTA_DEFINITIONS}) - string(SUBSTRING ${__def} 3 -1 __strip_def) - target_compile_definitions(vta PUBLIC ${__strip_def}) diff --git a/third_party/patch/incubator-tvm/find_library.patch b/third_party/patch/incubator-tvm/find_library.patch deleted file mode 100644 index f7b2f9af0a9..00000000000 --- a/third_party/patch/incubator-tvm/find_library.patch +++ /dev/null @@ -1,71 +0,0 @@ ---- tvm/python/tvm/_ffi/base.py 2020-03-12 16:17:39.089828527 +0800 -+++ tvm_new/python/tvm/_ffi/base.py 2020-03-12 16:17:16.829829558 +0800 -@@ -16,6 +16,9 @@ - # under the License. - # coding: utf-8 - # pylint: disable=invalid-name -+ -+# 2019.12.30 - Modify _load_lib function. -+ - """Base library for TVM FFI.""" - from __future__ import absolute_import - -@@ -47,8 +50,18 @@ else: - - - def _load_lib(): -- """Load libary by searching possible path.""" -- lib_path = libinfo.find_lib_path() -+ """Load library by searching possible path.""" -+ pwd = os.path.dirname(os.path.realpath(__file__)) -+ path = os.path.realpath(pwd+"/../../../mindspore/lib") -+ lib_path = [] -+ files = os.listdir(path) -+ for f in files: -+ if f.startswith("libtvm.") and f.endswith(".so"): -+ lib_path.append(path+"/"+f) -+ break -+ if not lib_path: -+ raise RuntimeError("mindspore library cannot find.") -+ - lib = ctypes.CDLL(lib_path[0], ctypes.RTLD_GLOBAL) - # DMatrix functions - lib.TVMGetLastError.restype = ctypes.c_char_p -diff -Npur tvm/topi/python/topi/cpp/impl.py tvm_new/topi/python/topi/cpp/impl.py ---- tvm/topi/python/topi/cpp/impl.py 2020-03-12 16:17:39.129828525 +0800 -+++ tvm_new/topi/python/topi/cpp/impl.py 2020-03-12 16:17:16.873829556 +0800 -@@ -14,6 +14,9 @@ - # KIND, either express or implied. See the License for the - # specific language governing permissions and limitations - # under the License. -+ -+# 2019.12.30 - Modify _load_lib function. -+ - """Load Lib for C++ TOPI ops and schedules""" - import sys - import os -@@ -30,12 +33,18 @@ def _get_lib_names(): - return ['libtvm_topi.so', 'tvm_topi.so'] - - def _load_lib(): -- """Load libary by searching possible path.""" -- curr_path = os.path.dirname(os.path.realpath(os.path.expanduser(__file__))) -- lib_search = curr_path -- lib_path = libinfo.find_lib_path(_get_lib_names(), lib_search, optional=True) -- if lib_path is None: -- return None, None -+ """Load library by searching possible path.""" -+ pwd = os.path.dirname(os.path.realpath(__file__)) -+ path = os.path.realpath(pwd+"/../../../mindspore/lib") -+ lib_path = [] -+ files = os.listdir(path) -+ for f in files: -+ if f.startswith("libtvm.") and f.endswith(".so"): -+ lib_path.append(path+"/"+f) -+ break -+ if not lib_path: -+ raise RuntimeError("mindspore library cannot find.") -+ - lib = ctypes.CDLL(lib_path[0], ctypes.RTLD_GLOBAL) - return lib, os.path.basename(lib_path[0]) - diff --git a/third_party/patch/incubator-tvm/include.patch b/third_party/patch/incubator-tvm/include.patch deleted file mode 100644 index 270c7a0f399..00000000000 --- a/third_party/patch/incubator-tvm/include.patch +++ /dev/null @@ -1,55 +0,0 @@ -diff -Npur tvm/include/tvm/expr_operator.h tvm_new/include/tvm/expr_operator.h ---- tvm/include/tvm/expr_operator.h 2019-12-28 10:11:27.369814744 +0800 -+++ tvm_new/include/tvm/expr_operator.h 2019-12-28 10:11:27.209812391 +0800 -@@ -25,6 +25,11 @@ - * when the type is int32 or int64 for simplifying the index expressions. - */ - // Acknowledgement: Most operator APIs originate from Halide. -+ -+/* -+ * 2019.12.30 - Add new operator for expr. -+ */ -+ - #ifndef TVM_EXPR_OPERATOR_H_ - #define TVM_EXPR_OPERATOR_H_ - -@@ -217,6 +222,16 @@ TVM_DLL Expr operator*(Expr a, Expr b); - */ - TVM_DLL Expr operator/(Expr a, Expr b); - /*! -+ * \brief mod operator -+ * -+ * \param a left operand -+ * \param b right operand -+ * \return The result expression. -+ * \note this function does eager constant folding for -+ * index types(int32, int64) when possible. -+ */ -+TVM_DLL Expr operator%(Expr a, Expr b); -+/*! - * \brief left shift operator - * - * \param a left operand -diff -Npur tvm/include/tvm/lowered_func.h tvm_new/include/tvm/lowered_func.h ---- tvm/include/tvm/lowered_func.h 2019-12-28 10:11:27.369814744 +0800 -+++ tvm_new/include/tvm/lowered_func.h 2019-12-28 10:11:27.209812391 +0800 -@@ -22,6 +22,11 @@ - * \brief Information about a lowered TVM function. - * This data structure is final step toward codegen. - */ -+ -+/* -+ * 2019.12.30 - Add new var array for args_real. -+ */ -+ - #ifndef TVM_LOWERED_FUNC_H_ - #define TVM_LOWERED_FUNC_H_ - -@@ -74,6 +79,7 @@ class LoweredFuncNode : public ir::Funct - * This function can only take pod type(int, float) and void* as arguments. - */ - Array args; -+ Array args_real; - /*! - * \brief The IterVar axis of threads - * Each axis need host function to specify a size. diff --git a/third_party/patch/incubator-tvm/src_pass.patch b/third_party/patch/incubator-tvm/src_pass.patch deleted file mode 100644 index 5450ca142ec..00000000000 --- a/third_party/patch/incubator-tvm/src_pass.patch +++ /dev/null @@ -1,120 +0,0 @@ -diff -Npur tvm/src/pass/make_api.cc tvm_new/src/pass/make_api.cc ---- tvm/src/pass/make_api.cc 2019-12-14 15:11:37.626419432 +0800 -+++ tvm_new/src/pass/make_api.cc 2019-12-14 14:58:46.562493287 +0800 -@@ -20,6 +20,11 @@ - /*! - * \file make_api.cc Build API function. - */ -+ -+/* -+ * 2019.12.30 - Define new function to push buffer node from api_args to args_real. -+ */ -+ - #include - #include - #include -@@ -40,6 +45,17 @@ inline Stmt MakeAssertEQ(Expr lhs, Expr - return AssertStmt::make(lhs == rhs, msg, Evaluate::make(0)); - } - -+Array Param ( Array api_args,Array args_real) { -+ int num_args = static_cast(api_args.size()); -+ for (int i = 0; i < num_args; i++) { -+ const BufferNode *v = api_args[i].as(); -+ if(v) { -+ args_real.push_back(v->data); -+ } -+ } -+ return args_real; -+} -+ - LoweredFunc MakeAPI(Stmt body, - std::string name, - Array api_args, -@@ -47,6 +63,8 @@ LoweredFunc MakeAPI(Stmt body, - bool is_restricted) { - const Stmt nop = Evaluate::make(0); - int num_args = static_cast(api_args.size()); -+ Array args_real; -+ args_real = Param (api_args, args_real); - CHECK_LE(num_unpacked_args, num_args); - int num_packed_args = num_args - num_unpacked_args; - // Data field definitions -@@ -170,6 +188,7 @@ LoweredFunc MakeAPI(Stmt body, - NodePtr n = make_node(); - n->name = name; - n->args = args; -+ n->args_real = args_real; - n->handle_data_type = binder.def_handle_dtype(); - n->is_packed_func = num_unpacked_args == 0; - n->is_restricted = is_restricted; -diff -Npur tvm/src/pass/split_host_device.cc tvm_new/src/pass/split_host_device.cc ---- tvm/src/pass/split_host_device.cc 2019-12-14 15:11:37.626419432 +0800 -+++ tvm_new/src/pass/split_host_device.cc 2019-12-14 11:28:49.293979656 +0800 -@@ -21,6 +21,11 @@ - * \file split_host_device.cc - * \brief Split device function from host. - */ -+ -+/* -+ * 2019.12.30 - Add new implements for host device splitter. -+ */ -+ - #include - #include - #include -@@ -38,6 +43,7 @@ class IRUseDefAnalysis : public IRMutato - Stmt Mutate_(const AttrStmt *op, const Stmt& s) final { - if (op->attr_key == attr::thread_extent) { - IterVar iv = Downcast(op->node); -+ iv = IterVarNode::make(Range(0, op->value), iv->var, iv->iter_type, iv->thread_tag); - CHECK_NE(iv->thread_tag.length(), 0U); - // thread_extent can appear multiple times - // use the first appearance as def. -@@ -186,6 +192,7 @@ class HostDeviceSplitter : public IRMuta - name_ = f->name; - NodePtr n = - make_node(*f.operator->()); -+ args_real = n->args_real; - n->body = this->Mutate(f->body); - n->func_type = kHostFunc; - Array ret{LoweredFunc(n)}; -@@ -196,6 +203,7 @@ class HostDeviceSplitter : public IRMuta - } - - private: -+ Array args_real; - Stmt SplitDeviceFunc(Stmt body) { - std::ostringstream os; - os << name_ << "_kernel" << device_funcs_.size(); -@@ -223,6 +231,30 @@ class HostDeviceSplitter : public IRMuta - n->args.push_back(v); - } - } -+std::shared_ptr na = std::make_shared(); -+ for (unsigned i = 0; i < (unsigned)args_real.size(); i++) { -+ bool match = false; -+ for (unsigned j = 0; j < (unsigned)n->args.size(); j++) { -+ if (strcmp(args_real[i].get()->name_hint.c_str(), n->args[j].get()->name_hint.c_str()) == 0) { -+ na->args.push_back(n->args[j]); -+ match = true; -+ break; -+ } else { -+ continue; -+ } -+ } -+ -+ if (!match) { -+ na->args.push_back(args_real[i]); -+ // mark handle data type. -+ for (auto kv : handle_data_type_) { -+ if (strcmp(args_real[i].get()->name_hint.c_str(), kv.first->name_hint.c_str()) == 0) { -+ n->handle_data_type.Set(args_real[i], kv.second); -+ } -+ } -+ } -+ } -+ n->args = na->args; - LoweredFunc f_device(n); - Array call_args; - call_args.push_back(StringImm::make(f_device->name));