p19647583
/
mindspore
forked from mindspore-Ecosystem/mindspore
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

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

This commit is contained in:
looop5 2020-08-03 14:35:44 +08:00
parent 1091cdeedf
commit c2b4bad040
62 changed files with 2 additions and 3205 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
akg

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

7
cmake/external_libs/dlpack.cmake
View File

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

7
cmake/external_libs/dmlc_core.cmake
View File

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

7
cmake/external_libs/rang.cmake
View File

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

15
cmake/external_libs/tvm_gpu.cmake
View File

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

5
cmake/mind_expression.cmake
View File

@ -46,11 +46,6 @@ if (ENABLE_CPU)
endif()
if (ENABLE_GPU)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/dlpack.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/dmlc_core.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/rang.cmake)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/tvm_gpu.cmake)
if (ENABLE_MPI)
include(${CMAKE_SOURCE_DIR}/cmake/external_libs/nccl.cmake)
endif()

23
cmake/package.cmake
View File

@ -230,29 +230,6 @@ install(
COMPONENT mindspore
)
if (ENABLE_GPU)
install(
DIRECTORY ${CMAKE_SOURCE_DIR}/mindspore/_akg
DESTINATION ${INSTALL_PY_DIR}/../
COMPONENT mindspore
)
if (EXISTS ${incubator_tvm_gpu_ROOT})
file(GLOB_RECURSE GLOG_LIB_LIST ${incubator_tvm_gpu_LIBPATH}/lib*)
install(
FILES ${GLOG_LIB_LIST}
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
install(
DIRECTORY
${incubator_tvm_gpu_ROOT}/topi/python/topi
${incubator_tvm_gpu_ROOT}/python/tvm
DESTINATION ${INSTALL_PY_DIR}/../_akg
COMPONENT mindspore
)
endif ()
endif ()
if ((ENABLE_D OR ENABLE_GPU) AND ENABLE_AKG)
set (AKG_PATH ${CMAKE_SOURCE_DIR}/build/mindspore/akg)
install(

18
mindspore/_akg/__init__.py
View File

@ -1,18 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""__init__"""
from . import add_path
from .op_build import op_build
from .message import compilewithjson

62
mindspore/_akg/add_path.py
View File

@ -1,62 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""add tvm path"""
import sys
import os
def AKGAddPath():
"""_akg add path."""
pwd = os.path.dirname(os.path.realpath(__file__))
tvm_path = os.path.realpath(pwd)
if tvm_path not in sys.path:
sys.path.insert(0, tvm_path)
else:
sys.path.remove(tvm_path)
sys.path.insert(0, tvm_path)
class AKGMetaPathFinder:
"""class AKGMetaPath finder."""
def find_module(self, fullname, path=None):
"""method _akg find module."""
_ = path
if fullname.startswith("_akg.tvm"):
rname = fullname[5:]
return AKGMetaPathLoader(rname)
if fullname.startswith("_akg.topi"):
rname = fullname[5:]
return AKGMetaPathLoader(rname)
return None
class AKGMetaPathLoader:
"""class AKGMetaPathLoader loader."""
def __init__(self, rname):
self.__rname = rname
def load_module(self, fullname):
if self.__rname in sys.modules:
sys.modules.pop(self.__rname)
AKGAddPath()
__import__(self.__rname, globals(), locals())
self.__target_module = sys.modules[self.__rname]
sys.modules[fullname] = self.__target_module
return self.__target_module
sys.meta_path.insert(0, AKGMetaPathFinder())

39
mindspore/_akg/gpu/__init__.py
View File

@ -1,39 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""__init__"""
from .equal import Equal
from .equal import gpu_schedule_Equal
from .tile import Tile
from .tile import gpu_schedule_Tile
from .cast import Cast
from .cast import gpu_schedule_Cast
from .relu6 import ReLU6, gpu_schedule_ReLU6
from .relu6_grad import ReLU6Grad, gpu_schedule_ReLU6Grad
from .squeeze import Squeeze, gpu_schedule_Squeeze
from .squeeze_grad import SqueezeGrad, gpu_schedule_SqueezeGrad
from .mean import SimpleMean, gpu_schedule_SimpleMean
from .mean_grad import SimpleMeanGrad, gpu_schedule_SimpleMeanGrad
from .mul import Mul, gpu_schedule_Mul
from .hsigmoid import HSigmoid, gpu_schedule_HSigmoid
from .hsigmoid_grad import HSigmoidGrad, gpu_schedule_HSigmoidGrad
from .hswish import HSwish, gpu_schedule_HSwish
from .hswish_grad import HSwishGrad, gpu_schedule_HSwishGrad
from .logical_or import LogicalOr, gpu_schedule_LogicalOr
from .logical_not import LogicalNot, gpu_schedule_LogicalNot
from .logical_and import LogicalAnd, gpu_schedule_LogicalAnd
from .sub import Sub, gpu_schedule_Sub
from .less_equal import LessEqual, gpu_schedule_LessEqual
from .notequal import NotEqual, gpu_schedule_NotEqual
from .greater_equal import GreaterEqual, gpu_schedule_GreaterEqual

45
mindspore/_akg/gpu/cast.py
View File

@ -1,45 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""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

56
mindspore/_akg/gpu/default_schedule.py
View File

@ -1,56 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""default schedule function for GPU"""
from queue import Queue
import _akg.tvm as tvm
DEFAULT_GPU_THREAD = 1024
def default_schedule(outs):
"""
default schedule function.
Args:
outs (Union[tvm.tensor.Tensor, list[tvm.tensor.Tensor]]): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
if not isinstance(outs, tvm.tensor.Tensor) and not isinstance(outs, list):
raise ValueError("outs should be list of _akg.tvm.tensor.Tensor or _akg.tvm.tensor.Tensor")
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
outs_list = [outs] if isinstance(outs, tvm.tensor.Tensor) else outs
with tvm.target.create(device):
sch = tvm.create_schedule(outs_list[0].op)
outputs_tensor = Queue()
outputs_tensor.put(outs_list[0])
op_list = []
while not outputs_tensor.empty():
out = outputs_tensor.get()
if out.op not in op_list and isinstance(out.op, tvm.tensor.ComputeOp):
op_list.append(out.op)
for input_tensor in out.op.input_tensors:
outputs_tensor.put(input_tensor)
for op in op_list:
stage = sch[op.output(0)]
bx, tx = stage.split(op.axis[0], factor=DEFAULT_GPU_THREAD)
stage.bind(bx, tvm.thread_axis("blockIdx.x"))
stage.bind(tx, tvm.thread_axis("threadIdx.x"))
return sch

40
mindspore/_akg/gpu/equal.py
View File

@ -1,40 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""equal"""
import _akg.tvm
from _akg.ops.math import equal
from _akg.topi.generic import schedule_elemwise
def Equal(x, y):
"""equal."""
return equal.equal(x, y)
def gpu_schedule_Equal(outs):
"""
gpu schedule for Equal.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

41
mindspore/_akg/gpu/greater_equal.py
View File

@ -1,41 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""greater_equal"""
import _akg.tvm
from _akg.ops.math import greater_equal
from _akg.topi.generic import schedule_elemwise
def GreaterEqual(x, y):
"""GreaterEqual."""
return greater_equal.greater_equal(x, y)
def gpu_schedule_GreaterEqual(outs):
"""
GPU schedule for GreaterEqual.
Args:
outs (tvm.tensor.Tensor): Outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

63
mindspore/_akg/gpu/hsigmoid.py
View File

@ -1,63 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""hsigmoid"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag
@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_hsigmoid(x):
"""
topi hsigmoid
Args:
x:
Returns:
"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, 1,
(x(*i) + 3) / 6)))
def HSigmoid(x):
"""
HSigmoid
Args:
x:
Returns:
"""
return topi_nn_hsigmoid(x)
def gpu_schedule_HSigmoid(outs):
"""
gpu schedule HSigmoid
Args:
outs:
Returns:
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

51
mindspore/_akg/gpu/hsigmoid_grad.py
View File

@ -1,51 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""HSigmoid grad"""
import _akg.topi as topi
import _akg.tvm as tvm
def HSigmoidGrad(y_grad, x):
"""
HSigmoidGrad
Args:
y_grad:
x:
Returns:
"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, 0,
y_grad(*i) / 6)))
def gpu_schedule_HSigmoidGrad(outs):
"""
gpu schedule ReLU6Grad
Args:
outs:
Returns:
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

63
mindspore/_akg/gpu/hswish.py
View File

@ -1,63 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""HSwish"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag
@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_HSwish(x):
"""
topi HSwish
Args:
x:
Returns:
"""
return tvm.compute(x.shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0,
tvm.if_then_else(x(*i) >= 3, x(*i),
x(*i) * (x(*i) + 3) / 6)))
def HSwish(x):
"""
HSwish
Args:
x:
Returns:
"""
return topi_nn_HSwish(x)
def gpu_schedule_HSwish(outs):
"""
gpu schedule HSwish
Args:
outs:
Returns:
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

53
mindspore/_akg/gpu/hswish_grad.py
View File

@ -1,53 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""HSwishGrad"""
import _akg.topi as topi
import _akg.tvm as tvm
def HSwishGrad(y_grad, x):
"""
HSwishGrad
Args:
y_grad:
x:
Returns:
"""
shape = x.shape
res0 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) <= -3, 0, y_grad(*i) * (2 * x(*i) + 3) / 6))
res6 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= 3, y_grad(*i), res0(*i)))
return res6
def gpu_schedule_HSwishGrad(outs):
"""
gpu schedule HSwishGrad
Args:
outs:
Returns:
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

40
mindspore/_akg/gpu/less_equal.py
View File

@ -1,40 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""less_equal"""
import _akg.tvm
from _akg.ops.math import less_equal
from _akg.topi.generic import schedule_elemwise
def LessEqual(x, y):
"""LessEqual."""
return less_equal.less_equal(x, y)
def gpu_schedule_LessEqual(outs):
"""
GPU schedule for LessEqual.
Args:
outs (tvm.tensor.Tensor): Outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

40
mindspore/_akg/gpu/logical_and.py
View File

@ -1,40 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""logical_and"""
import _akg.tvm
from _akg.ops.math import logical_and
from _akg.topi.generic import schedule_elemwise
def LogicalAnd(x, y):
"""LogicalAnd."""
return logical_and.logical_and(x, y)
def gpu_schedule_LogicalAnd(outs):
"""
GPU schedule for LogicalAnd.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

40
mindspore/_akg/gpu/logical_not.py
View File

@ -1,40 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""logical_not"""
import _akg.tvm
from _akg.ops.math import logical_not
from _akg.topi.generic import schedule_elemwise
def LogicalNot(x):
"""LogicalNot."""
return logical_not.logical_not(x)
def gpu_schedule_LogicalNot(outs):
"""
GPU schedule for LogicalNot.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

40
mindspore/_akg/gpu/logical_or.py
View File

@ -1,40 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""logical_or"""
import _akg.tvm
from _akg.ops.math import logical_or
from _akg.topi.generic import schedule_elemwise
def LogicalOr(x, y):
"""LogicalOr."""
return logical_or.logical_or(x, y)
def gpu_schedule_LogicalOr(outs):
"""
GPU schedule for LogicalOr.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

80
mindspore/_akg/gpu/mean.py
View File

@ -1,80 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""mean op compute and schedule"""
import _akg.tvm as tvm
from _akg.ops.math.mean import mean
from .default_schedule import DEFAULT_GPU_THREAD
def Mean(x, axis=None, keepdims=True):
"""mean."""
outs = mean(x, axis, keepdims)
# remove useless mean_output
if isinstance(outs, tuple):
outs = outs[0]
if outs.op.name == "mean_output":
outs = outs.op.input_tensors[0]
return outs
def gpu_schedule_Mean(outs):
"""
gpu schedule function for mean.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
out = outs[0] if isinstance(outs, list) else outs
device = "cuda"
with tvm.target.create(device):
sch = tvm.create_schedule(out.op)
if out.op.name == "T_divide":
tensor_c = out
else: # squeeze
tensor_c = out.op.input_tensors[0]
tensor_b = tensor_c.op.input_tensors[0]
if len(tensor_c.op.axis) >= 2:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[1])
else:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[0])
bx, tx = sch[tensor_c].split(tensor_c.op.axis[0], factor=DEFAULT_GPU_THREAD)
sch[tensor_c].bind(bx, tvm.thread_axis("blockIdx.x"))
sch[tensor_c].bind(tx, tvm.thread_axis("threadIdx.x"))
return sch
def SimpleMean(x):
"""
SimpleMean compute the mean of the input 4D Tensor over last two axises and keep reduced dimensions.
Args:
x (tvm.tensor.Tensor): Tensor of type float16, float32.
Returns:
tvm.tensor.Tensor, has the same type as x, output shape will be (a, b, 1, 1) if input Tensor x is (a, b, c, d).
"""
axis = (2, 3)
keepdims = True
return Mean(x, axis, keepdims)
def gpu_schedule_SimpleMean(outs):
"""gpu schedule function for SimpleMean."""
return gpu_schedule_Mean(outs)

90
mindspore/_akg/gpu/mean_grad.py
View File

@ -1,90 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""mean_grad"""
import _akg.tvm as tvm
import _akg
from _akg.ops.math import mean
from .default_schedule import DEFAULT_GPU_THREAD
def mean_ad(head, input_shape, axis, keepdims):
"""mean autodiff."""
tensor_a = tvm.placeholder(input_shape, head.dtype, "A")
tensor_b = mean.mean(tensor_a, axis, keepdims)
# remove useless mean_output
if isinstance(tensor_b, tuple):
tensor_b = tensor_b[0]
if tensor_b.op.name == "mean_output":
tensor_b = tensor_b.op.input_tensors[0]
jacs = list(_akg.differentiate(tensor_b, [tensor_a], head))
return jacs[0]
def MeanGrad(y_grad, input_shape, axis=None, keepdims=True):
"""Mean Grad."""
if axis is None and not keepdims:
raise ValueError("Mean not support (axis=None && keepdims=False) now")
return mean_ad(y_grad, input_shape, axis, keepdims)
def gpu_schedule_MeanGrad(outs):
"""gpu schedule MeanGrad."""
out = outs[0] if isinstance(outs, list) else outs
device = "cuda"
with tvm.target.create(device):
sch = tvm.create_schedule(out.op)
tensor_c = out
tensor_b = tensor_c.op.input_tensors[0]
if len(tensor_c.op.axis) >= 2:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[1])
else:
sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[0])
bx, tx = sch[tensor_c].split(tensor_c.op.axis[0], factor=DEFAULT_GPU_THREAD)
sch[tensor_c].bind(bx, tvm.thread_axis("blockIdx.x"))
sch[tensor_c].bind(tx, tvm.thread_axis("threadIdx.x"))
return sch
def SimpleMeanGrad(HEAD, input_shape):
"""
Compute Simple Mean Grad.
Args:
HEAD (tvm.tensor.Tensor): output gradient, dy, defined in Primitive.
input_shape (Union[list[int], tuple[int]]): shape of mean input, x.shape.
Returns:
tvm.tensor.Tensor, gradient of mean input.
"""
axis = (2, 3)
keepdims = True
return MeanGrad(HEAD, input_shape, axis, keepdims)
def gpu_schedule_SimpleMeanGrad(outs):
"""
gpu schedule SimpleMeanGrad.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
return gpu_schedule_MeanGrad(outs)

41
mindspore/_akg/gpu/mul.py
View File

@ -1,41 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""mul"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.ops.math import mul
def Mul(x, y):
"""mul."""
return mul.mul(x, y)
def gpu_schedule_Mul(outs):
"""
gpu schedule for mul.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_broadcast(outs)
return sch

41
mindspore/_akg/gpu/notequal.py
View File

@ -1,41 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""notequal"""
import _akg.tvm
from _akg.ops.math import notequal
from _akg.topi.generic import schedule_elemwise
def NotEqual(x, y):
"""notequal."""
return notequal.notequal(x, y)
def gpu_schedule_NotEqual(outs):
"""
gpu schedule for NotEqual.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

54
mindspore/_akg/gpu/relu6.py
View File

@ -1,54 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""relu6"""
import _akg.topi as topi
import _akg.tvm as tvm
from _akg.topi import tag
@tvm.tag_scope(tag=tag.ELEMWISE)
def topi_nn_relu6(x):
"""topi nn relu6."""
return tvm.compute(x.shape, lambda *i: tvm.min(tvm.max(x(*i), tvm.const(0, x.dtype)), tvm.const(6, x.dtype)))
def ReLU6(x):
"""
Compute elementwise with function: min(max(x, 0), 6).
Args:
x (tvm.tensor.Tensor): Tensor of type float16, float32.
Returns:
tvm.tensor.Tensor, has same type and shape as input.
"""
return topi_nn_relu6(x)
def gpu_schedule_ReLU6(outs):
"""
gpu schedule ReLU6.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

59
mindspore/_akg/gpu/relu6_grad.py
View File

@ -1,59 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""relu6 grad"""
import _akg.topi as topi
import _akg.tvm as tvm
def ReLU6Grad(y_grad, x):
"""
Computes Gradients of Rectified Linear 6.
Args:
y_grad (tvm.tensor.Tensor): Tensor of type float16, float32, gradients backpropagated to the ReLU6 op.
x (tvm.tensor.Tensor): Tensor of type float16/float32, inputs that where passed to the ReLU6 op, or its outputs.
Returns:
tvm.tensor.Tensor, has same type and shape as x.
"""
shape = x.shape
dtype = x.dtype
zero = tvm.const(0, dtype)
six = tvm.const(6, dtype)
res0 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= zero, x(*i), zero))
res6 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= six, zero, res0(*i)))
res = tvm.compute(shape, lambda *i: tvm.if_then_else(res6(*i) == zero, zero, y_grad(*i)))
return res
def gpu_schedule_ReLU6Grad(outs):
"""
gpu schedule ReLU6Grad.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_elemwise(outs)
return sch

50
mindspore/_akg/gpu/squeeze.py
View File

@ -1,50 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""squeeze"""
import _akg.topi as topi
import _akg.tvm as tvm
def Squeeze(x, axis=None):
"""
Remove the dimensions which have shape size 1.
Args:
x (tvm.tensor.Tensor): Tensor, input whose shape is to be squeeze.
axis (Union[list, tuple, int, None]): specify which size 1 dimension to be removed.
Returns:
tvm.tensor.Tensor, has the same type and element as x, but some size 1 dimensions are removed.
"""
return topi.squeeze(x, axis)
def gpu_schedule_Squeeze(outs):
"""
gpu schedule Squeeze.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with tvm.target.create(device):
sch = topi.cuda.schedule_injective(outs)
return sch

44
mindspore/_akg/gpu/squeeze_grad.py
View File

@ -1,44 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""squeeze grad"""
import _akg.topi as topi
def SqueezeGrad(y_grad, x_shape):
"""
Computes gradients for squeeze op.
Args:
y_grad (tvm.tensor.Tensor): the gradient needed to be propagation.
x_shape (Union[list, tuple]): output Tensor shape.
Returns:
tvm.tensor.Tensor: output gradient.
"""
return topi.reshape(y_grad, x_shape)
def gpu_schedule_SqueezeGrad(outs):
"""
gpu schedule SqueezeGrad.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
from .default_schedule import default_schedule
return default_schedule(outs)

40
mindspore/_akg/gpu/sub.py
View File

@ -1,40 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""sub"""
import _akg.tvm
from _akg.ops.math import sub
from _akg.topi.generic import schedule_elemwise
def Sub(x, y):
"""Sub."""
return sub.sub(x, y)
def gpu_schedule_Sub(outs):
"""
GPU schedule for Sub.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
sch = schedule_elemwise(outs)
return sch

39
mindspore/_akg/gpu/tile.py
View File

@ -1,39 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""tile"""
import _akg.tvm
from _akg.ops.array import tile
from _akg.topi.generic import schedule_elemwise
def Tile(x, multiples):
"""tile."""
return tile.tile(x, multiples)
def gpu_schedule_Tile(outs):
"""
gpu schedule for tile.
Args:
outs (tvm.tensor.Tensor): outputs of compute.
Returns:
sch (schedule.Schedule): The created schedule.
"""
device = 'cuda'
ctx = _akg.tvm.context(device, 0)
if not ctx.exist:
raise SystemError("Skip because %s is not enabled" % device)
with _akg.tvm.target.create(device):
s = schedule_elemwise(outs)
return s

104
mindspore/_akg/message.py
View File

@ -1,104 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""message"""
import importlib.util
import json
import json.decoder as jd
import logging
import traceback
import os.path
from pathlib import Path
import _akg.tvm
from _akg.utils import validation_check as vc_util
from _akg.utils.dsl_create import TensorUtils
from . import gpu
from . import op_build
@vc_util.check_input_type(str)
def compilewithjson(json_str):
"""compile with json."""
try:
kernel_info = json.loads(json_str)
except jd.JSONDecodeError:
logging.error(traceback.format_exc())
return False
op_name = kernel_info['name']
op_func = None
processor = 'aicore'
if 'process' in kernel_info:
processor = kernel_info['process']
# get custom ops implementation first.
if 'impl_path' in kernel_info and kernel_info['impl_path'] is not None:
impl_path = os.path.realpath(kernel_info['impl_path'])
if os.path.isfile(impl_path):
custom_mod_name = Path(impl_path).resolve().stem
mod_spec = importlib.util.spec_from_file_location(
custom_mod_name, impl_path)
custom_mod = importlib.util.module_from_spec(mod_spec)
mod_spec.loader.exec_module(custom_mod)
op_func = getattr(custom_mod, op_name, None)
# get built-in ops.
if op_func is None:
if processor == 'cuda':
op_func = getattr(gpu, op_name, None)
if op_func is None:
logging.error(
"this op not supported, please check op name %s", str(op_name))
return False
args = {}
tsr = []
for input_desc in kernel_info['input_desc']:
if len(input_desc) == 1:
tensor_shape = input_desc[0]['shape']
tensor_shape = (1,) if not tensor_shape else tensor_shape
vc_util.shape_dtype_max_size_check(tensor_shape)
args[input_desc[0]['name']] = _akg.tvm.placeholder(
shape=tensor_shape, name=input_desc[0]['tensor_name'], dtype=input_desc[0]['data_type'])
tsr.append(args[input_desc[0]['name']])
else:
tmp_input = []
for tmp_desc in input_desc:
tensor_shape = tmp_desc['shape']
tensor_shape = (1,) if not tensor_shape else tensor_shape
vc_util.shape_dtype_max_size_check(tensor_shape)
tmp_input.append(_akg.tvm.placeholder(
shape=tensor_shape, name=tmp_desc['tensor_name'], dtype=tmp_desc['data_type']))
args[input_desc[0]['name']] = tmp_input
tsr = tsr + tmp_input
if kernel_info['attr']:
for ext_arg in kernel_info['attr']:
args[ext_arg['name']] = ext_arg['value']
output = op_func(**args)
if isinstance(output, (list, tuple)):
from inspect import isfunction
tmp_outputs = []
for elem in output:
if not isfunction(elem) or isinstance(elem, dict):
tmp_outputs.append(elem)
output = tmp_outputs
else:
output = [output]
tsr = tsr + [i for i in output if TensorUtils.is_output_value(i)]
return op_build([op_name], output, tsr, processor, kernel_info['op'])

69
mindspore/_akg/op_build.py
View File

@ -1,69 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""op_build"""
import os
import fcntl
import types
import typing
import logging
import traceback
import _akg.tvm
import _akg
from _akg import save_gpu_param as gpu_utils
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(list, (list, tuple), (list, tuple), str, str)
def op_build(opnames, computes, args, device, kernel_name):
"""op_build"""
kernel_meta_path = "./cuda_meta_" + str(os.getpid()) + "/"
if device == "cuda":
cuda_path = os.path.realpath(kernel_meta_path)
if not os.path.isdir(cuda_path):
os.makedirs(cuda_path)
if not opnames:
logging.error("no opname given.")
return None
schedule_name = 'gpu_schedule_' + opnames[0]
schedule_func = getattr(_akg.gpu, schedule_name)
if not isinstance(schedule_func, (types.FunctionType, typing.Callable)):
logging.error("no schedule func found %s", str(schedule_name))
return None
ptx_file = os.path.realpath(kernel_meta_path + kernel_name + ".ptx")
if os.path.exists(ptx_file):
os.chmod(ptx_file, 0o600)
try:
with open(ptx_file, 'at') as file:
fcntl.flock(file.fileno(), fcntl.LOCK_EX)
file.seek(0, 2)
if file.tell() == 0:
s = schedule_func(computes)
foo = _akg.tvm.build(s, args, device, name=kernel_name)
ptx_code = foo.imported_modules[0].get_source("ptx")
file.write(ptx_code)
json_file = os.path.realpath(
kernel_meta_path + kernel_name + ".json")
kernel_info = (ptx_code, json_file, kernel_name)
gpu_utils.save_gpu_params(s, args, kernel_info)
os.chmod(ptx_file, 0o400)
except IOError:
logging.error(traceback.format_exc())
return None
return True
logging.error("Not support device %s.", device)
return None

0
mindspore/_akg/ops/__init__.py
View File

0
mindspore/_akg/ops/array/__init__.py
View File

36
mindspore/_akg/ops/array/tile.py
View File

@ -1,36 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: tile"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple))
def tile(data, multiples):
"""
Repeats the data in the specified dimensions according to the multiples.
Args:
data (tvm.tensor.Tensor): Tensor.
multiples (Union[list, tuple]): Elements must be int. The number of repetitions.
Returns:
tvm.tensor.Tensor, has the same dtype as data.
"""
vc_util.check_shape(data.shape)
vc_util.check_int_list(multiples, "multiples")
output = _akg.topi.tile(data, multiples)
return output

0
mindspore/_akg/ops/math/__init__.py
View File

36
mindspore/_akg/ops/math/cast.py
View File

@ -1,36 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: cast"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, str)
def cast(data, dst_type):
"""
cast data to target type.
Args:
data (tvm.tensor.Tensor): Tensor to be casted.
dst_type (str): target cast type.
Returns:
tvm.tensor.Tensor, type is dst_type.
"""
vc_util.check_shape(data.shape)
out = _akg.topi.cast(data, dst_type)
return out

54
mindspore/_akg/ops/math/equal.py
View File

@ -1,54 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: equal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def equal(input1, input2):
"""
check whether input1 equals to input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. If input1 equal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
shape1, shape2, shape = produce_shapes(shape1, shape2)
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype
# get equal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")
input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] == input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
return res

54
mindspore/_akg/ops/math/greater_equal.py
View File

@ -1,54 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: greaterequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def greater_equal(input1, input2):
"""
Check whether input1 greaterquals to input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. If input1 greaterquals to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
shape1, shape2, shape = produce_shapes(shape1, shape2)
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype
# get greaterquals compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")
input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] >= input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
return res

54
mindspore/_akg/ops/math/less_equal.py
View File

@ -1,54 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: lessequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def less_equal(input1, input2):
"""
Check whether input1 lessequals to input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. If input1 lessequal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
shape1, shape2, shape = produce_shapes(shape1, shape2)
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype
# get lessequal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")
input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] <= input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
return res

41
mindspore/_akg/ops/math/logical_and.py
View File

@ -1,41 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: logical_and"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def logical_and(input1, input2):
"""
Compute logical_and of input1 and input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. LogicalAnd of input1 and input2.
"""
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
res = _akg.topi.logical_and(input1, input2)
return res

32
mindspore/_akg/ops/math/logical_not.py
View File

@ -1,32 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: logical_not"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor)
def logical_not(input1):
"""
Compute logical_not of input1.
Args:
input1 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor.
"""
res = _akg.topi.logical_not(input1)
return res

41
mindspore/_akg/ops/math/logical_or.py
View File

@ -1,41 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: logical_or"""
import _akg.tvm
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def logical_or(input1, input2):
"""
Compute logical_or of input1 and input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. LogicalOr of input1 and input2.
"""
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
res = _akg.topi.logical_or(input1, input2)
return res

47
mindspore/_akg/ops/math/mean.py
View File

@ -1,47 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: mean"""
import _akg.topi
import _akg.tvm
from _akg.utils import format_transform as ft_util
from _akg.utils import validation_check as vc_util
from _akg.ops.math import sum_value
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
def mean(data, axis=None, keepdims=False):
"""
Computes the mean of the values of a Tensor over the whole dataset.
Args:
data (tvm.tensor.Tensor): Tensor.
axis (Union[list, tuple, int, None]): If the tuple is empty, the axis equal to None.
keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.
Returns:
tvm.tensor.Tensor, has the same type as data. If keepdims equal to True, all reduced dimensions are
retained with length 1. else these reduced axis will be eliminate.
"""
shape = [x.value for x in data.shape]
vc_util.reduce_axis_check(shape, axis)
axis = ft_util.refine_reduce_axis(data, axis)
count = 1
for i in axis:
count *= shape[i]
output, _ = sum_value.sum_value(data, axis, keepdims)
res = _akg.topi.divide(output, count)
return res

43
mindspore/_akg/ops/math/mul.py
View File

@ -1,43 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: mul"""
import _akg.topi
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def mul(l_input, r_input):
"""
Calculate x * y element-wise.
Note:
mul supports broadcasting.
Args:
l_input (tvm.tensor.Tensor): Tensor.
r_input (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor, has the same type as l_input and r_input.
"""
shape1 = [x.value for x in l_input.shape]
shape2 = [x.value for x in r_input.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
vc_util.auto_broadcast_check(shape1, shape2)
vc_util.elemwise_dtype_check(l_input.dtype, r_input.dtype)
output = _akg.topi.multiply(l_input, r_input)
return output

54
mindspore/_akg/ops/math/notequal.py
View File

@ -1,54 +0,0 @@
# Copyright 2020 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: notequal"""
import _akg.tvm
import _akg.topi
from _akg.utils.dsl_create import produce_shapes
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def notequal(input1, input2):
"""
check whether input1 notequals to input2.
Args:
input1 (tvm.tensor.Tensor): Tensor.
input2 (tvm.tensor.Tensor): Tensor.
Returns:
tvm.tensor.Tensor. If input1 notequal to input2 return True, else return False.
"""
shape1 = [x.value for x in input1.shape]
shape2 = [x.value for x in input2.shape]
vc_util.check_shape(shape1)
vc_util.check_shape(shape2)
shape1, shape2, shape = produce_shapes(shape1, shape2)
vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
dtype = input1.dtype
# get notequal compute
t_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(1, dtype), "T")
f_value = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.const(0, dtype), "F")
input1_bro = _akg.topi.broadcast_to(input1, shape)
input2_bro = _akg.topi.broadcast_to(input2, shape)
c_out = _akg.tvm.compute(shape, lambda *indice: _akg.tvm.expr.Select(input1_bro[indice] != input2_bro[indice],
t_value[indice], f_value[indice]), name="C")
res = _akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
return res

40
mindspore/_akg/ops/math/sub.py
View File

@ -1,40 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: sub"""
import _akg.topi
import _akg.tvm
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, _akg.tvm.tensor.Tensor)
def sub(data1, data2):
"""
Computes data1 - data2 elementwise, broadcast is supported.
Args:
data1 (tvm.tensor.Tensor): Tensor.
data2 (tvm.tensor.Tensor): Tensor of same type as data1, if shape(data2) != shape(data1), broadcast will happen.
Returns:
tvm.tensor.Tensor, subtracted result, with same type as input tensors and broadcasted shape of data1 and data2.
"""
vc_util.elemwise_dtype_check(data1.dtype, data2.dtype)
vc_util.check_shape(data1.shape)
vc_util.check_shape(data2.shape)
vc_util.auto_broadcast_check(data1.shape, data2.shape)
res = _akg.topi.subtract(data1, data2)
return res

45
mindspore/_akg/ops/math/sum_value.py
View File

@ -1,45 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""operator dsl function: sum"""
import _akg.topi
import _akg.tvm
from _akg.utils import format_transform as ft_util
from _akg.utils import validation_check as vc_util
@vc_util.check_input_type(_akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
def sum_value(inputs, axis=None, keepdims=False):
"""
Compute the sum of elements across dimensions of a tensor.
Args:
inputs (tvm.tensor.Tensor): Tensor.
axis (Union[list, tuple, int, None]): If the list or tuple is empty, the axis equal to None.
keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.
Returns:
tvm.tensor.Tensor, has same type as input. If keepdims is True, all reduced dimensions are retained
with length 1, else these reduced axis will be eliminate.
"""
axis = ft_util.refine_reduce_axis(inputs, axis)
vc_util.check_shape(inputs.shape)
if not axis:
output = _akg.topi.identity(inputs)
else:
output = _akg.topi.sum(inputs, axis=axis, keepdims=keepdims)
return output

87
mindspore/_akg/save_gpu_param.py
View File

@ -1,87 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""save gpu param"""
import os
import hashlib
import _akg.tvm
from _akg.tvm import schedule
from _akg.utils import validation_check as vc_util
def get_dim(dim, axis=True):
"""get dim info"""
dims_str = {
"grid_dim0": "// attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = ",
"grid_dim1": "// attr [iter_var(blockIdx.y, , blockIdx.y)] thread_extent = ",
"grid_dim2": "// attr [iter_var(blockIdx.z, , blockIdx.z)] thread_extent = ",
"block_dim0": "// attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = ",
"block_dim1": "// attr [iter_var(threadIdx.y, , threadIdx.y)] thread_extent = ",
"block_dim2": "// attr [iter_var(threadIdx.z, , threadIdx.z)] thread_extent = "
}
dim_to_axis = {
"grid_dim0": '"blockIdx.x" : ',
"grid_dim1": '"blockIdx.y" : ',
"grid_dim2": '"blockIdx.z" : ',
"block_dim0": '"threadIdx.x" : ',
"block_dim1": '"threadIdx.y" : ',
"block_dim2": '"threadIdx.z" : '
}
if axis:
return dim_to_axis.get(dim)
return dims_str.get(dim)
def parse_params(file, dim, ir):
"""parse parameters"""
dim_str = get_dim(dim, axis=False)
pos = ir.find(dim_str)
if pos != -1:
index = pos + len(dim_str)
param_temp = get_dim(dim)
while ir[index].isdigit():
param_temp += ir[index]
index += 1
file.write(param_temp + ",\n")
else:
param_temp = get_dim(dim) + '1'
file.write(param_temp + ",\n")
@vc_util.check_input_type(schedule.Schedule, (list, tuple), tuple)
def save_gpu_params(s, args, kernel_info):
"""save gpu parameters"""
ptx_code = kernel_info[0]
file_name = kernel_info[1]
kernel_name = kernel_info[2]
ir = str(_akg.tvm.lower(s, args, simple_mode=True))
file_path = os.path.realpath(file_name)
if os.path.exists(file_path):
os.remove(file_path)
sha256 = hashlib.sha256()
sha256.update(ptx_code.encode("utf-8"))
hash_str = sha256.hexdigest()
with os.fdopen(os.open(file_path, os.O_WRONLY | os.O_CREAT, 0o400), 'w') as fo:
fo.write("{\n")
fo.write('"kernelName" : ' + '"' + kernel_name + "_kernel0" + '",\n')
parse_params(fo, "grid_dim0", ir)
parse_params(fo, "grid_dim1", ir)
parse_params(fo, "grid_dim2", ir)
parse_params(fo, "block_dim0", ir)
parse_params(fo, "block_dim1", ir)
parse_params(fo, "block_dim2", ir)
fo.write('"sha256" : ' + '"' + hash_str + '"\n')
fo.write("}\n")

0
mindspore/_akg/utils/__init__.py
View File

122
mindspore/_akg/utils/dsl_create.py
View File

@ -1,122 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""dsl create helping function"""
import _akg
from _akg.utils import format_transform as ft_util
class TensorUtils:
"""Class for creating tensor."""
CREATE_SCH_ONLY = 'create_sch_only'
@classmethod
def get_tensor_attrs(cls, tensor):
"""get tensor attrs."""
tensor_attrs = dict()
if "attrs" in dir(tensor.op):
tensor_attrs = dict(tensor.op.attrs.items())
return tensor_attrs
@classmethod
def update_tensor_attrs(cls, tensor, attrs):
"""update tensor attrs."""
tensor_attrs = cls.get_tensor_attrs(tensor)
tensor_attrs.update(attrs)
tensor = _akg.tvm.compute(tensor.shape,
lambda *indice: tensor[indice],
name=tensor.op.name,
tag=tensor.op.tag,
attrs=tensor_attrs)
return tensor
@classmethod
def is_create_sch_only(cls, tensor):
tensor_attrs = cls.get_tensor_attrs(tensor)
if cls.CREATE_SCH_ONLY in tensor_attrs.keys():
return True
return False
@classmethod
def is_output_value(cls, tensor):
"""check output value."""
return not cls.is_create_sch_only(tensor)
@classmethod
def inplace_set(cls, input_tensor, output_tensor, buffer_name="data_buf"):
"""inplace set."""
input_tensor_shape = ft_util.get_shape(input_tensor)
output_tensor_shape = ft_util.get_shape(output_tensor)
if not input_tensor_shape == output_tensor_shape:
raise RuntimeError("Shape of the input_tensor and the output_tensor should be equal, "
"but got %s and %s"%(input_tensor_shape, output_tensor_shape))
output_tensor = cls.update_tensor_attrs(output_tensor, {cls.CREATE_SCH_ONLY: 1})
data_buf = _akg.tvm.decl_buffer(input_tensor.shape, input_tensor.dtype, name=buffer_name)
binds_info = {input_tensor: data_buf, output_tensor: data_buf}
return output_tensor, binds_info
@classmethod
def inplace_set_tensors(cls, input_tensors, output_tensors, buffer_names=None):
"""
inplace set for tensors
Args:
in_tensors (Union[list, tuple]): Origin input tensors.
out_tensors (Union[list, tuple]): Origin output tensors.
buffer_names (Union[list, tuple] or None): Buffer names used to bind.
Return:
inplace_tensors (list): Output tensors with the inplace info.
binds_infos (dict): Dictionary that maps the input tensor and the output
tensor to buffer.
"""
if not buffer_names:
buffer_names = ["data_buf_%s" % i for i in range(len(input_tensors))]
for arg in (input_tensors, output_tensors, buffer_names):
if not isinstance(arg, (tuple, list)):
raise RuntimeError("arg must be tuple or list!")
if len(input_tensors) != len(output_tensors) or len(input_tensors) != len(buffer_names):
raise RuntimeError("length of the input_tensors, output_tensors and buffer_names must be equal!")
inplace_tensors = []
binds_infos = dict()
for input_tensor, output_tensor, buffer_name in zip(input_tensors, output_tensors, buffer_names):
inplace_tensor, binds_info = cls.inplace_set(input_tensor, output_tensor, buffer_name)
inplace_tensors.append(inplace_tensor)
binds_infos.update(binds_info)
return inplace_tensors, binds_infos
def produce_shapes(shape1, shape2):
"""two input shapes produce three output shape."""
shape1 = list(shape1)
shape2 = list(shape2)
flag = 0
if len(shape1) < len(shape2):
shape1, shape2 = shape2, shape1
flag = 1
output_shape_len = len(shape1)
dec = output_shape_len - len(shape2)
for i in range(dec):
shape2 = [1] + shape2
out_shape = []
for i in range(output_shape_len):
if (shape1[i] != shape2[i]) and (shape1[i] != 1) and (shape2[i] != 1):
raise RuntimeError("input shapes not match!")
out_shape.append(shape1[i] if shape1[i] > shape2[i] else shape2[i])
if flag == 1:
shape1, shape2 = shape2, shape1
return shape1, shape2, out_shape

80
mindspore/_akg/utils/format_transform.py
View File

@ -1,80 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""format transform function"""
import _akg
def refine_reduce_axis(input_content, axis):
"""make reduce axis legal."""
shape = get_shape(input_content)
if axis is None:
axis = [i for i in range(len(shape))]
elif isinstance(axis, int):
axis = [axis]
elif not isinstance(axis, (tuple, list)):
raise TypeError("axis must be one of the type int,tuple,list or None")
if len(axis) > len(shape):
raise ValueError("axis size must not larger than shape size")
axis = list(axis)
for i, _ in enumerate(axis):
if axis[i] < 0:
axis[i] += len(shape)
if axis[i] >= len(shape):
raise ValueError(("axis value-{} exceeds len(axis) which is invalid".format(axis[i])))
axis.sort(reverse=True)
return axis
def get_shape_from_tensor(data):
"""translate _akg.tvm.shape to list type in python."""
tvm_shape = data.shape
py_shape = []
for i in tvm_shape:
if isinstance(i, _akg.tvm.expr.Var):
py_shape.append(i)
else:
py_shape.append(i.value)
return py_shape
def tvm_shape_to_list(tvm_shape):
"""translate _akg.tvm.shape to list type in python."""
py_shape = []
for i in tvm_shape:
if isinstance(i, _akg.tvm.expr.Var):
py_shape.append(i)
else:
py_shape.append(i.value)
return py_shape
def get_shape(data):
"""get shape and save it as list."""
if isinstance(data, _akg.tvm.tensor.Tensor):
shape = get_shape_from_tensor(data)
elif isinstance(data, _akg.tvm.container.Array):
shape = tvm_shape_to_list(data)
elif isinstance(data, int):
shape = [data]
elif isinstance(data, (tuple, list)):
shape = list(data)
else:
raise TypeError("Refine axis does not support type {} for now.".format(type(data)))
return shape

233
mindspore/_akg/utils/validation_check.py
View File

@ -1,233 +0,0 @@
# Copyright 2019 Huawei Technologies Co., Ltd
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""validation check functions"""
from functools import wraps, reduce
from _akg.utils.format_transform import get_shape
MAX_DATA_SIZE = 2 ** 31
def check_input_type_dict(input_dict, input_key, input_name):
"""
check input parameter type for new type: dict.
Note:
rule1: key of input_dict should be in the input_key
rule2: type of input_dict[shape] should be in (list, tuple), if have shape
rule3: type of input_dict[dtype] should be in (str), if have dtype
Args:
input_dict (dict): input_dict
input_key (list or tuple): all input key list, the key of input must in input_key
input_name (str): input param name, only used for error print
Returns:
None
"""
def _check_input_type(input_key, input_type):
if not isinstance(input_dict[input_key], input_type):
raise RuntimeError(
"the input parameter %s[%s] must be %s, while type of input is %s" %
(input_name, input_key, input_type, type(input_dict[input_key])))
for key in input_dict.keys():
if key not in input_key:
raise RuntimeError(
"the input parameter %s must have arrt <%s>" %
(input_name, key))
# check shape's type of input_dict, if have shape
if key == "shape":
_check_input_type(key, (list, tuple))
# check dtype's type of input_dict, if have dtype
if key == "dtype":
_check_input_type(key, (str,))
def check_input_type_list_tuple(inputs, expect):
"""check inputs by a list or tuple of expected types."""
if not isinstance(inputs, expect[1][0]):
raise RuntimeError("the input parameter %s must be (list, tuple), while"
" type of input is %s" % (expect[0], type(inputs)))
for inp in inputs:
if not isinstance(inp, expect[1][1]):
raise RuntimeError("The element in parameter %s must be %s, while "
"type of input is %s" % (
expect[0], expect[1][1], type(inp)))
def check_input_type(*type_args, **_type_kwargs):
"""check input parameter type."""
def out_wrapper(func):
"""outer wrapper function."""
formal_parameter = func.__code__.co_varnames
formal_parameter_list = list(zip(formal_parameter, type_args))
@wraps(func)
def in_wrapper(*args, **kwargs):
"""inner wrapper function."""
for i, arg_v in enumerate(args):
# add for new input dict, if dict, will check shape and dtype
if isinstance(arg_v, dict):
check_input_type_dict(arg_v, arg_v.keys(),
formal_parameter_list[i][0])
if isinstance(formal_parameter_list[i][1], tuple):
if isinstance(formal_parameter_list[i][1][0], tuple) \
and len(formal_parameter_list[i][1]) == 2:
check_input_type_list_tuple(arg_v, formal_parameter_list[i])
continue
if not isinstance(arg_v, formal_parameter_list[i][1]):
raise RuntimeError("the %sth input parameter %s must be %s, "
"while type of input is %s" % (str(i), formal_parameter_list[i][0],
formal_parameter_list[i][1],
type(arg_v)))
for i in kwargs:
for j in formal_parameter_list:
if i in j:
if not isinstance(kwargs[i], j[1]):
raise RuntimeError("the input parameter %s must be "
"%s, while type of input is %s"
"" % (i, j[1], type(kwargs[i])))
break
return func(*args, **kwargs)
return in_wrapper
return out_wrapper
def shape_dtype_max_size_check(shape):
"""check validation of tensor's shape."""
if shape:
mul = int(reduce(lambda x, y: int(x) * int(y), shape))
if mul > MAX_DATA_SIZE:
error_msg = "*".join([str(sh) for sh in shape])
raise RuntimeError("Invalid shape, data is {} bytes ({}), which "
"exceed max data size {} bytes"
.format(mul, error_msg, MAX_DATA_SIZE))
def check_shape(tensor, length=None, tensor_name=""):
"""The common check rule for placeholder data."""
shape = get_shape(tensor)
if not shape:
raise RuntimeError("The ndim of input tensor {} must more than 0, "
"actual input is {}".format(tensor_name, len(shape)))
for shape_v in shape:
if not isinstance(shape_v, int) or shape_v <= 0:
raise RuntimeError("The type of tensor {} axis value must be "
"positive int and value more than 0,"
"actual input is ({}) {}".
format(tensor_name, type(shape_v), shape_v))
if length and len(shape) != length:
raise ValueError('The length of {} should be {}, while actual length is {}'.
format(tensor_name, length, len(shape)))
def ops_dtype_check(dtype, args):
"""check validation of op's dtype."""
expected_dtype = list()
def _get_expect_dtype(expected_dtype, arg):
if isinstance(arg, str):
expected_dtype.append(arg)
elif isinstance(arg, (list, tuple)):
for t in arg:
_get_expect_dtype(expected_dtype, t)
else:
raise TypeError("arg should be either a string, "
"or a list/tuple of string, "
"while current is {}".format(type(arg)))
_get_expect_dtype(expected_dtype, args)
if isinstance(dtype, (list, tuple)):
checking_dtype = [d.lower() for d in dtype]
elif isinstance(dtype, str):
checking_dtype = [dtype.lower()]
else:
raise TypeError("dtype should be either a string or a tuple/list of string")
error_msg = "Supported dtype: {}, while received dtype: {}"
if not set(checking_dtype).issubset(set(expected_dtype)):
raise RuntimeError(error_msg.format(expected_dtype, checking_dtype))
def reduce_axis_check(reduce_shape, reduce_axis):
"""check validation of reduce axis for certain reduce shape."""
dim = len(reduce_shape)
if dim == 1 and int(reduce_shape[0]) == 1:
raise RuntimeError("Error, reduce shape is 1. Scalar is not supported "
"for reduction, please input a vector.")
if isinstance(reduce_axis, int):
if reduce_axis not in range(-dim, dim):
raise RuntimeError("Reduce axis should be in range [%d. %d)"
"" % (-dim, dim))
elif isinstance(reduce_axis, (tuple, list)):
if len(reduce_axis) > len(reduce_shape):
raise RuntimeError("Reduce axis list exceed reduce shape length: "
"%d vs %d, error" % (len(reduce_axis), len(reduce_shape)))
processed_axis = []
for axis in reduce_axis:
processed_axis.append(int(axis + dim) if axis < 0 else int(axis))
if len(set(processed_axis)) < len(processed_axis):
raise RuntimeError("Reduce axis list contains %d duplicated element, please check"
% (len(processed_axis) - len(set(processed_axis))))
for axis in processed_axis:
if axis >= dim:
raise RuntimeError("Invalid reduce axis, axis should less than %d" % dim)
elif reduce_axis is not None:
raise RuntimeError("axis should be a list, tuple or int.")
def elemwise_dtype_check(dtype_a, dtype_b, supported_type=None):
"""check validation of tensor's dtype for element-wise op."""
if supported_type:
ops_dtype_check(dtype_a, supported_type)
ops_dtype_check(dtype_b, supported_type)
if dtype_a.lower() != dtype_b.lower():
raise RuntimeError("Element-wise operation needs same data type, while "
"current is %s vs %s" % (dtype_a.lower(), dtype_b.lower()))
def auto_broadcast_check(shape_a, shape_b):
"""automatic broadcast check."""
shape_l = get_shape(shape_a)
shape_r = get_shape(shape_b)
if len(shape_l) <= len(shape_r):
shape_short = shape_l
shape_long = shape_r
else:
shape_short = shape_r
shape_long = shape_l
dim_diff = len(shape_long) - len(shape_short)
for i in range(dim_diff):
shape_short.insert(0, 1)
for i, shp in enumerate(shape_short):
if int(shp) != int(shape_long[i]) and 1 not in [int(shp), int(shape_long[i])]:
raise RuntimeError("Invalid auto broadcast, dim %d should be 1 or equal, "
"while now is %d vs %d" % (i, shp, shape_long[i]))
def check_int_list(array, array_name):
"""check whether all the elements are integers."""
for num in array:
if not isinstance(num, int):
raise RuntimeError("Type of value in %s should be int, but got type %s" % (array_name, type(num)))

2
setup.py
View File

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

91
third_party/apply_patches.sh vendored
View File

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

100
third_party/patch/incubator-tvm/CMakeLists.txt vendored
View File

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

201
third_party/patch/incubator-tvm/cmake.patch vendored
View File

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

71
third_party/patch/incubator-tvm/find_library.patch vendored
View File

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

55
third_party/patch/incubator-tvm/include.patch vendored
View File

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

120
third_party/patch/incubator-tvm/src_pass.patch vendored
View File

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