!3869 remove redundant compilation of libtvm.so and mindspore/_akg directory

Merge pull request !3869 from looop5/remove_redundant_code_in_gpu

akg

@@ -1 +1 @@
-Subproject commit 949a45538ccb7ae94ad73386b5e3e77005112eea
+Subproject commit 5fe7e5c8377dccfd35c9f661e10ed3dc136208c5

cmake/external_libs/dlpack.cmake

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

cmake/external_libs/dmlc_core.cmake

@@ -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)
-
-

cmake/external_libs/rang.cmake

@@ -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)
-
-

cmake/external_libs/tvm_gpu.cmake

@@ -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)

cmake/mind_expression.cmake

@@ -47,11 +47,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()

cmake/package.cmake

@@ -245,29 +245,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(

mindspore/_akg/__init__.py

@@ -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

mindspore/_akg/add_path.py

@@ -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())

mindspore/_akg/gpu/__init__.py

@@ -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

mindspore/_akg/gpu/cast.py

@@ -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

mindspore/_akg/gpu/default_schedule.py

@@ -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

mindspore/_akg/gpu/equal.py

@@ -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

mindspore/_akg/gpu/greater_equal.py

@@ -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

mindspore/_akg/gpu/hsigmoid.py

@@ -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

mindspore/_akg/gpu/hsigmoid_grad.py

@@ -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

mindspore/_akg/gpu/hswish.py

@@ -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

mindspore/_akg/gpu/hswish_grad.py

@@ -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

mindspore/_akg/gpu/less_equal.py

@@ -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

mindspore/_akg/gpu/logical_and.py

@@ -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

mindspore/_akg/gpu/logical_not.py

@@ -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

mindspore/_akg/gpu/logical_or.py

@@ -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

mindspore/_akg/gpu/mean.py

@@ -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)

mindspore/_akg/gpu/mean_grad.py

@@ -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)

mindspore/_akg/gpu/mul.py

@@ -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

mindspore/_akg/gpu/notequal.py

@@ -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

mindspore/_akg/gpu/relu6.py

@@ -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

mindspore/_akg/gpu/relu6_grad.py

@@ -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

mindspore/_akg/gpu/squeeze.py

@@ -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

mindspore/_akg/gpu/squeeze_grad.py

@@ -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)

mindspore/_akg/gpu/sub.py

@@ -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

mindspore/_akg/gpu/tile.py

@@ -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

mindspore/_akg/message.py

@@ -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'])

mindspore/_akg/op_build.py

@@ -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

mindspore/_akg/ops/array/tile.py

@@ -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

mindspore/_akg/ops/math/cast.py

@@ -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

mindspore/_akg/ops/math/equal.py

@@ -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

mindspore/_akg/ops/math/greater_equal.py

@@ -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

mindspore/_akg/ops/math/less_equal.py

@@ -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

mindspore/_akg/ops/math/logical_and.py

@@ -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

mindspore/_akg/ops/math/logical_not.py

@@ -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

mindspore/_akg/ops/math/logical_or.py

@@ -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

mindspore/_akg/ops/math/mean.py

@@ -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

mindspore/_akg/ops/math/mul.py

@@ -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

mindspore/_akg/ops/math/notequal.py

@@ -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

mindspore/_akg/ops/math/sub.py

@@ -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

mindspore/_akg/ops/math/sum_value.py

@@ -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

mindspore/_akg/save_gpu_param.py

@@ -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")

mindspore/_akg/utils/dsl_create.py

@@ -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

mindspore/_akg/utils/format_transform.py

@@ -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

mindspore/_akg/utils/validation_check.py

@@ -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)))

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)
 
 

third_party/apply_patches.sh

@@ -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}

third_party/patch/incubator-tvm/CMakeLists.txt

@@ -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
-        $<$<COMPILE_LANGUAGE:C>:${CMAKE_C_FLAGS_AKG}>
-        $<$<COMPILE_LANGUAGE:CXX>:${CMAKE_CXX_FLAGS_AKG}>)
-target_include_directories(tvm PRIVATE "${TVM_DIR}/topi/include")
-install(TARGETS tvm)

third_party/patch/incubator-tvm/cmake.patch

@@ -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})

third_party/patch/incubator-tvm/find_library.patch

@@ -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])
- 

third_party/patch/incubator-tvm/include.patch

@@ -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<Var> args;
-+  Array<Var> args_real;
-   /*!
-    * \brief The IterVar axis of threads
-    *  Each axis need host function to specify a size.

third_party/patch/incubator-tvm/src_pass.patch

@@ -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 <tvm/ir_pass.h>
- #include <tvm/ir.h>
- #include <tvm/ir_visitor.h>
-@@ -40,6 +45,17 @@ inline Stmt MakeAssertEQ(Expr lhs, Expr
-   return AssertStmt::make(lhs == rhs, msg, Evaluate::make(0));
- }
- 
-+Array<Var> Param ( Array<NodeRef> api_args,Array<Var> args_real) {
-+    int num_args = static_cast<int>(api_args.size());
-+    for (int i = 0; i < num_args; i++) {
-+        const BufferNode *v = api_args[i].as<BufferNode>();
-+        if(v) {
-+            args_real.push_back(v->data);
-+        }
-+    }
-+    return args_real;
-+}
-+
- LoweredFunc MakeAPI(Stmt body,
-                     std::string name,
-                     Array<NodeRef> api_args,
-@@ -47,6 +63,8 @@ LoweredFunc MakeAPI(Stmt body,
-                     bool is_restricted) {
-   const Stmt nop = Evaluate::make(0);
-   int num_args = static_cast<int>(api_args.size());
-+  Array<Var> 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<LoweredFuncNode> n = make_node<LoweredFuncNode>();
-   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 <tvm/ir.h>
- #include <tvm/lowered_func.h>
- #include <tvm/channel.h>
-@@ -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<IterVar>(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<LoweredFuncNode> n =
-         make_node<LoweredFuncNode>(*f.operator->());
-+    args_real = n->args_real;
-     n->body = this->Mutate(f->body);
-     n->func_type = kHostFunc;
-     Array<LoweredFunc> ret{LoweredFunc(n)};
-@@ -196,6 +203,7 @@ class HostDeviceSplitter : public IRMuta
-   }
- 
-  private:
-+  Array<Var> 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<LoweredFuncNode> na = std::make_shared<LoweredFuncNode>();
-+    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<Expr> call_args;
-     call_args.push_back(StringImm::make(f_device->name));