diff --git a/CMakeLists.txt b/CMakeLists.txt
index 099e13c67c2bb8987cf8b30a97c704e2d92084a0..0f4736d7e9b25bfd00e1f07a522b3a1feaa2ed7a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -99,8 +99,9 @@ add_definitions(-DDMLC_LOG_CUSTOMIZE=1)
 if(USE_AKG_LOG)
   add_definitions(-DUSE_AKG_LOG=1)
 endif()
-if(NOT USE_CUDA)
-  add_definitions("-DBACKEND_D")
+if(NOT USE_CUDA
+   OR ENABLE_AKG)
+  add_definitions("-DFIX_INPUT_ORDER_TVM")
 endif()
 
 # Generic compilation options
diff --git a/python/akg/ms/gpu/__init__.py b/python/akg/ms/gpu/__init__.py
index 0cf7937dcc47befed4f8649adab93bbc0bc3791a..76f2e39aa4b29442ef33262f2f388a623b83715b 100644
--- a/python/akg/ms/gpu/__init__.py
+++ b/python/akg/ms/gpu/__init__.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -15,16 +13,23 @@
 # limitations under the License.
 
 """__init__"""
+from .notequal import NotEqual
 from .equal import Equal
-from .equal import gpu_schedule_Equal
+from .greater_equal import GreaterEqual
+from .less_equal import LessEqual
 from .tile import Tile
-from .tile import gpu_schedule_Tile
 from .cast import Cast
-from .relu6 import ReLU6, gpu_schedule_ReLU6
-from .relu6_grad import ReLU6Grad, gpu_schedule_ReLU6Grad
-from .squeeze import Squeeze, gpu_schedule_Squeeze
+from .relu6 import ReLU6
+from .logical_and import LogicalAnd
+from .logical_not import LogicalNot
+from .logical_or import LogicalOr
+from .relu6_grad import ReLU6Grad
+from .squeeze import 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 .mean import SimpleMean
+from .sub import Sub
 from .mul import Mul
+from .hsigmoid import HSigmoid
+from .hsigmoid_grad import HSigmoidGrad
+from .hswish import HSwish
+from .hswish_grad import HSwishGrad
diff --git a/python/akg/ms/gpu/cast.py b/python/akg/ms/gpu/cast.py
index 4ef744d3a31bd03a2ac5072c77fd42f30d9d6c73..0478a20f94a1d16709617f70aea313cab0daf735 100644
--- a/python/akg/ms/gpu/cast.py
+++ b/python/akg/ms/gpu/cast.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -15,13 +13,13 @@
 # limitations under the License.
 
 """cast"""
-import logging
-import akg.tvm
-from akg.ops.math import cast
-from akg.topi.generic import schedule_elemwise
+import akg
+from akg.ops.math_gpu import cast
 import akg.topi as topi
 
 @akg.schedule(topi.cuda.schedule_injective)
 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)
diff --git a/python/akg/ms/gpu/default_schedule.py b/python/akg/ms/gpu/default_schedule.py
index 0af76cf313e76ec135fa952b887b3a03293bb8b2..a2e5d15afa7bfc698a2d3f7b1cc487cce5e8a08b 100644
--- a/python/akg/ms/gpu/default_schedule.py
+++ b/python/akg/ms/gpu/default_schedule.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
diff --git a/python/akg/ms/gpu/equal.py b/python/akg/ms/gpu/equal.py
index cd44f01b3dd1341f3a2be4c6f1eafb10ada99c2a..5ad7984303a500d24d7083e6838ab898520210e6 100644
--- a/python/akg/ms/gpu/equal.py
+++ b/python/akg/ms/gpu/equal.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -13,30 +13,13 @@
 # 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
+import akg
+import akg.topi as topi
+from akg.ops.math_gpu import equal
 
+@akg.schedule(topi.cuda.schedule_injective)
 def Equal(x, y):
-    """equal."""
+    """Equal"""
     return equal.equal(x, y)
-
-
-def gpu_schedule_Equal(outs):
-    """
-    gpu schedule for Equal.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    device = 'cuda'
-    ctx = akg.tvm.context(device, 0)
-    if not ctx.exist:
-        raise SystemError("Skip because %s is not enabled" % device)
-    with akg.tvm.target.create(device):
-        sch = schedule_elemwise(outs)
-    return sch
diff --git a/python/akg/ms/gpu/greater_equal.py b/python/akg/ms/gpu/greater_equal.py
new file mode 100644
index 0000000000000000000000000000000000000000..12da66cdf7cfa84d6338f78e5e2f191dc19ed8b3
--- /dev/null
+++ b/python/akg/ms/gpu/greater_equal.py
@@ -0,0 +1,23 @@
+# 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
+import akg.topi as topi
+from akg.ops.math_gpu import greater_equal
+
+@akg.schedule(topi.cuda.schedule_injective)
+def GreaterEqual(x, y):
+    """GreaterEqual"""
+    return greater_equal.greater_equal(x, y)
diff --git a/python/akg/ms/gpu/hsigmoid.py b/python/akg/ms/gpu/hsigmoid.py
new file mode 100644
index 0000000000000000000000000000000000000000..4ac3bdb1a5c8abf5d5527edbb3baeafc832c6a38
--- /dev/null
+++ b/python/akg/ms/gpu/hsigmoid.py
@@ -0,0 +1,45 @@
+# 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.
+
+"""hsigmoid"""
+import akg.topi as topi
+import akg.tvm as tvm
+from akg.topi import tag
+import akg
+
+@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)))
+
+@akg.schedule(topi.cuda.schedule_injective)
+def HSigmoid(x):
+    """
+    HSigmoid
+    Args:
+        x:
+
+    Returns:
+
+    """
+    return topi_nn_hsigmoid(x)
diff --git a/python/akg/ms/gpu/hsigmoid_grad.py b/python/akg/ms/gpu/hsigmoid_grad.py
new file mode 100644
index 0000000000000000000000000000000000000000..4d6d021d5ea9f2eb62d5d8eda038554d0011c309
--- /dev/null
+++ b/python/akg/ms/gpu/hsigmoid_grad.py
@@ -0,0 +1,34 @@
+# 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.
+
+"""HSigmoid grad"""
+import akg.topi as topi
+import akg.tvm as tvm
+import akg
+
+@akg.schedule(topi.cuda.schedule_injective)
+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)))
+
diff --git a/python/akg/ms/gpu/hswish.py b/python/akg/ms/gpu/hswish.py
new file mode 100644
index 0000000000000000000000000000000000000000..956ab5e8e408caacac535c880ecd397c317ac5dd
--- /dev/null
+++ b/python/akg/ms/gpu/hswish.py
@@ -0,0 +1,46 @@
+# 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.
+
+"""HSwish"""
+import akg.topi as topi
+import akg.tvm as tvm
+from akg.topi import tag
+import akg
+
+
+@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)))
+
+@akg.schedule(topi.cuda.schedule_injective)
+def HSwish(x):
+    """
+    HSwish
+    Args:
+        x:
+
+    Returns:
+
+    """
+    return topi_nn_HSwish(x)
diff --git a/python/akg/ms/gpu/hswish_grad.py b/python/akg/ms/gpu/hswish_grad.py
new file mode 100644
index 0000000000000000000000000000000000000000..cdf6d0a81f450f5083beeb08fdaa2ec0f48bf5f7
--- /dev/null
+++ b/python/akg/ms/gpu/hswish_grad.py
@@ -0,0 +1,35 @@
+# 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.
+
+"""HSwishGrad"""
+import akg.topi as topi
+import akg.tvm as tvm
+import akg
+
+@akg.schedule(topi.cuda.schedule_injective)
+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
diff --git a/python/akg/ms/gpu/less_equal.py b/python/akg/ms/gpu/less_equal.py
new file mode 100644
index 0000000000000000000000000000000000000000..73e6b225fecef76961ebc4a978e52e95d9a0ac0c
--- /dev/null
+++ b/python/akg/ms/gpu/less_equal.py
@@ -0,0 +1,22 @@
+# 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
+import akg.topi
+from akg.ops.math_gpu import less_equal
+
+@akg.schedule(akg.topi.cuda.schedule_injective)
+def LessEqual(x, y):
+    return less_equal.less_equal(x, y)
diff --git a/python/akg/ms/gpu/logical_and.py b/python/akg/ms/gpu/logical_and.py
new file mode 100644
index 0000000000000000000000000000000000000000..713b19fa505fb0926e5d08e2bc8f266c2d6b1fdd
--- /dev/null
+++ b/python/akg/ms/gpu/logical_and.py
@@ -0,0 +1,24 @@
+# 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
+import akg.topi as topi
+from akg.ops.math_gpu import logical_and
+
+@akg.schedule(topi.cuda.schedule_injective)
+def LogicalAnd(x, y):
+    """LogicalAnd."""
+    return logical_and.logical_and(x, y)
+
diff --git a/python/akg/ms/gpu/logical_not.py b/python/akg/ms/gpu/logical_not.py
new file mode 100644
index 0000000000000000000000000000000000000000..05adc6ed39a667a146acb8f8a51a5c9dc38f95bb
--- /dev/null
+++ b/python/akg/ms/gpu/logical_not.py
@@ -0,0 +1,24 @@
+# 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
+import akg.topi as topi
+from akg.ops.math_gpu import logical_not
+
+@akg.schedule(topi.cuda.schedule_injective)
+def LogicalNot(x):
+    """LogicalNot."""
+    return logical_not.logical_not(x)
+
diff --git a/python/akg/ms/gpu/logical_or.py b/python/akg/ms/gpu/logical_or.py
new file mode 100644
index 0000000000000000000000000000000000000000..deb6e7fa39c3e360d9a89a002d17c1bcbd99be3f
--- /dev/null
+++ b/python/akg/ms/gpu/logical_or.py
@@ -0,0 +1,23 @@
+# 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
+import akg.topi as topi
+from akg.ops.math_gpu import logical_or
+
+@akg.schedule(topi.cuda.schedule_injective)
+def LogicalOr(x, y):
+    """LogicalOr."""
+    return logical_or.logical_or(x, y)
diff --git a/python/akg/ms/gpu/mean.py b/python/akg/ms/gpu/mean.py
index 257f4105e66ac5f702a018c746b3d58f399e337c..cb0ae9a20ea92c689cc4893c877ebad3a663fe14 100644
--- a/python/akg/ms/gpu/mean.py
+++ b/python/akg/ms/gpu/mean.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -15,21 +15,10 @@
 # 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
-
+from akg.ops.math_gpu.sum_value import sum_value
+import akg
+from akg.ops.math_gpu.mean import mean
 
 def gpu_schedule_Mean(outs):
     """
@@ -43,25 +32,28 @@ def gpu_schedule_Mean(outs):
     """
     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]
+    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])
+    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"))
+    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
 
+@akg.schedule(gpu_schedule_Mean)
+def Mean(data, axis=None, keepdims=False):
+    return mean(data, axis, keepdims)
+
+@akg.schedule(gpu_schedule_Mean)
 def SimpleMean(x):
     """
     SimpleMean compute the mean of the input 4D Tensor over last two axises and keep reduced dimensions.
@@ -74,9 +66,4 @@ def SimpleMean(x):
     """
     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)
+    return mean(x, axis, keepdims)
diff --git a/python/akg/ms/gpu/mean_grad.py b/python/akg/ms/gpu/mean_grad.py
deleted file mode 100644
index 701d2654c7926c0f63a76c1442bdb6a7851aa66b..0000000000000000000000000000000000000000
--- a/python/akg/ms/gpu/mean_grad.py
+++ /dev/null
@@ -1,92 +0,0 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-"""mean_grad"""
-import akg.tvm as tvm
-import akg
-from akg.ops.math import mean
-from .default_schedule import DEFAULT_GPU_THREAD
-
-
-def mean_ad(head, input_shape, axis, keepdims):
-    """mean autodiff."""
-    tensor_a = tvm.placeholder(input_shape, head.dtype, "A")
-    tensor_b = mean.mean(tensor_a, axis, keepdims)
-
-    # remove useless mean_output
-    if isinstance(tensor_b, tuple):
-        tensor_b = tensor_b[0]
-    if tensor_b.op.name == "mean_output":
-        tensor_b = tensor_b.op.input_tensors[0]
-
-    jacs = list(akg.differentiate(tensor_b, [tensor_a], head))
-    return jacs[0]
-
-
-def MeanGrad(y_grad, input_shape, axis=None, keepdims=True):
-    """Mean Grad."""
-    if axis is None and not keepdims:
-        raise ValueError("Mean not support (axis=None && keepdims=False)  now")
-    return mean_ad(y_grad, input_shape, axis, keepdims)
-
-
-def gpu_schedule_MeanGrad(outs):
-    """gpu schedule MeanGrad."""
-    out = outs[0] if isinstance(outs, list) else outs
-
-    device = "cuda"
-    with tvm.target.create(device):
-        sch = tvm.create_schedule(out.op)
-        tensor_c = out
-        tensor_b = tensor_c.op.input_tensors[0]
-        if len(tensor_c.op.axis) >= 2:
-            sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[1])
-        else:
-            sch[tensor_b].compute_at(sch[tensor_c], tensor_c.op.axis[0])
-
-        bx, tx = sch[tensor_c].split(tensor_c.op.axis[0], factor=DEFAULT_GPU_THREAD)
-        sch[tensor_c].bind(bx, tvm.thread_axis("blockIdx.x"))
-        sch[tensor_c].bind(tx, tvm.thread_axis("threadIdx.x"))
-
-    return sch
-
-def SimpleMeanGrad(HEAD, input_shape):
-    """
-    Compute Simple Mean Grad.
-
-    Args:
-        HEAD (tvm.tensor.Tensor): output gradient, dy, defined in Primitive.
-        input_shape (Union[list[int], tuple[int]]): shape of mean input, x.shape.
-
-    Returns:
-        tvm.tensor.Tensor, gradient of mean input.
-    """
-    axis = (2, 3)
-    keepdims = True
-    return MeanGrad(HEAD, input_shape, axis, keepdims)
-
-
-def gpu_schedule_SimpleMeanGrad(outs):
-    """
-    gpu schedule SimpleMeanGrad.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    return gpu_schedule_MeanGrad(outs)
diff --git a/python/akg/ms/gpu/mul.py b/python/akg/ms/gpu/mul.py
index 4dda614e766a660aed565f793c77774b2a989a8e..334674fa0129638442117657bd5a47c1117c5156 100644
--- a/python/akg/ms/gpu/mul.py
+++ b/python/akg/ms/gpu/mul.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -17,8 +17,7 @@
 """mul"""
 import akg
 import akg.topi as topi
-import akg.tvm as tvm
-from akg.ops.math import mul
+from akg.ops.math_gpu import mul
 
 @akg.schedule(topi.cuda.schedule_injective)
 def Mul(x, y):
diff --git a/python/akg/ms/gpu/notequal.py b/python/akg/ms/gpu/notequal.py
new file mode 100644
index 0000000000000000000000000000000000000000..1f2147ee23d3a85c649096ac0f1e57a93f912377
--- /dev/null
+++ b/python/akg/ms/gpu/notequal.py
@@ -0,0 +1,23 @@
+# 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
+import akg.topi
+from akg.ops.math_gpu import notequal
+
+@akg.schedule(akg.topi.cuda.schedule_injective)
+def NotEqual(x, y):
+    """notequal."""
+    return notequal.notequal(x, y)
diff --git a/python/akg/ms/gpu/relu6.py b/python/akg/ms/gpu/relu6.py
index efa45805f5ee121cd80d0240b0f401e38533245e..a65677a42ac2daeead2bde036b49186952111649 100644
--- a/python/akg/ms/gpu/relu6.py
+++ b/python/akg/ms/gpu/relu6.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -16,14 +14,16 @@
 
 """relu6"""
 import akg.topi as topi
-import akg.tvm as tvm
 from akg.topi import tag
+import akg
+import akg.tvm as tvm
 
 @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)))
 
+@akg.schedule(topi.cuda.schedule_injective)
 def ReLU6(x):
     """
     Compute elementwise with function: min(max(x, 0), 6).
@@ -35,22 +35,3 @@ def ReLU6(x):
         tvm.tensor.Tensor, has same type and shape as input.
     """
     return topi_nn_relu6(x)
-
-
-def gpu_schedule_ReLU6(outs):
-    """
-    gpu schedule ReLU6.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    device = 'cuda'
-    ctx = tvm.context(device, 0)
-    if not ctx.exist:
-        raise SystemError("Skip because %s is not enabled" % device)
-    with tvm.target.create(device):
-        sch = topi.cuda.schedule_elemwise(outs)
-    return sch
diff --git a/python/akg/ms/gpu/relu6_grad.py b/python/akg/ms/gpu/relu6_grad.py
index 45fb000bebca764502a2c57f0fae3c1f496ec69b..f321b1e22df7a0126653fe9b1b9b9efe074e3375 100644
--- a/python/akg/ms/gpu/relu6_grad.py
+++ b/python/akg/ms/gpu/relu6_grad.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -17,7 +15,9 @@
 """relu6 grad"""
 import akg.topi as topi
 import akg.tvm as tvm
+import akg
 
+@akg.schedule(topi.cuda.schedule_injective)
 def ReLU6Grad(y_grad, x):
     """
     Computes Gradients of Rectified Linear 6.
@@ -39,23 +39,3 @@ def ReLU6Grad(y_grad, x):
     res6 = tvm.compute(shape, lambda *i: tvm.if_then_else(x(*i) >= six, zero, res0(*i)))
     res = tvm.compute(shape, lambda *i: tvm.if_then_else(res6(*i) == zero, zero, y_grad(*i)))
     return res
-
-
-def gpu_schedule_ReLU6Grad(outs):
-    """
-    gpu schedule ReLU6Grad.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    device = 'cuda'
-    ctx = tvm.context(device, 0)
-    if not ctx.exist:
-        raise SystemError("Skip because %s is not enabled" % device)
-
-    with tvm.target.create(device):
-        sch = topi.cuda.schedule_elemwise(outs)
-    return sch
diff --git a/python/akg/ms/gpu/squeeze.py b/python/akg/ms/gpu/squeeze.py
index 678e30170e1453040bb30d7cf0e78ec50e63049b..1c0fa7f58e57bf621e86cfc89e956c1ebc0e9e49 100644
--- a/python/akg/ms/gpu/squeeze.py
+++ b/python/akg/ms/gpu/squeeze.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -16,8 +14,9 @@
 
 """squeeze"""
 import akg.topi as topi
-import akg.tvm as tvm
+import akg
 
+@akg.schedule(topi.cuda.schedule_injective)
 def Squeeze(x, axis=None):
     """
     Remove the dimensions which have shape size 1.
@@ -30,23 +29,3 @@ def Squeeze(x, axis=None):
         tvm.tensor.Tensor, has the same type and element as x, but some size 1 dimensions are removed.
     """
     return topi.squeeze(x, axis)
-
-
-def gpu_schedule_Squeeze(outs):
-    """
-    gpu schedule Squeeze.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    device = 'cuda'
-    ctx = tvm.context(device, 0)
-    if not ctx.exist:
-        raise SystemError("Skip because %s is not enabled" % device)
-
-    with tvm.target.create(device):
-        sch = topi.cuda.schedule_injective(outs)
-    return sch
diff --git a/python/akg/ms/gpu/squeeze_grad.py b/python/akg/ms/gpu/squeeze_grad.py
index 20626ebfff41a43014f20f9adb5ac399dea2ecf4..86bbb0042373cc7a63d7112179768f104a93d938 100644
--- a/python/akg/ms/gpu/squeeze_grad.py
+++ b/python/akg/ms/gpu/squeeze_grad.py
@@ -1,6 +1,4 @@
-#!/usr/bin/env python3
-# coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -16,31 +14,31 @@
 
 """squeeze grad"""
 import akg.topi as topi
+import akg
 
-def SqueezeGrad(y_grad, x_shape, axis=None):
+def gpu_schedule_SqueezeGrad(outs):
     """
-    Computes gradients for squeeze op.
+    gpu schedule SqueezeGrad.
 
     Args:
-        y_grad (tvm.tensor.Tensor): the gradient needed to be propagation.
-        x_shape (Union[list, tuple]): output Tensor shape.
-        axis (Union[list, tuple, int, None], optional): eliminated axis by squeeze.
+        outs (tvm.tensor.Tensor): outputs of compute.
 
     Returns:
-        tvm.tensor.Tensor: output gradient.
+        sch (schedule.Schedule): The created schedule.
     """
-    return topi.reshape(y_grad, x_shape)
-
+    from .default_schedule import default_schedule
+    return default_schedule(outs)
 
-def gpu_schedule_SqueezeGrad(outs):
+@akg.schedule(gpu_schedule_SqueezeGrad)
+def SqueezeGrad(y_grad, x_shape):
     """
-    gpu schedule SqueezeGrad.
+    Computes gradients for squeeze op.
 
     Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
+        y_grad (tvm.tensor.Tensor): the gradient needed to be propagation.
+        x_shape (Union[list, tuple]): output Tensor shape.
 
     Returns:
-        sch (schedule.Schedule): The created schedule.
+        tvm.tensor.Tensor: output gradient.
     """
-    from .default_schedule import default_schedule
-    return default_schedule(outs)
+    return topi.reshape(y_grad, x_shape)
diff --git a/python/akg/ms/gpu/sub.py b/python/akg/ms/gpu/sub.py
new file mode 100644
index 0000000000000000000000000000000000000000..47c92dceb55ce3f1a3aaaee0a318588debbc133b
--- /dev/null
+++ b/python/akg/ms/gpu/sub.py
@@ -0,0 +1,23 @@
+# 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
+import akg.topi as topi
+from akg.ops.math_gpu import sub
+
+@akg.schedule(topi.cuda.schedule_injective)
+def Sub(x, y):
+    """Sub."""
+    return sub.sub(x, y)
diff --git a/python/akg/ms/gpu/tile.py b/python/akg/ms/gpu/tile.py
index 2ac548dc51ad4e3cd035b8aef1a64e06ca788c63..d98707cdbc6ac090c57fb9fe56ef44ef8e2603ed 100644
--- a/python/akg/ms/gpu/tile.py
+++ b/python/akg/ms/gpu/tile.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# 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.
@@ -13,29 +13,13 @@
 # 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
+from akg.ops.array_gpu import tile
+import akg.topi as topi
+import akg
 
+@akg.schedule(topi.cuda.schedule_injective)
 def Tile(x, multiples):
     """tile."""
     return tile.tile(x, multiples)
-
-def gpu_schedule_Tile(outs):
-    """
-    gpu schedule for tile.
-
-    Args:
-        outs (tvm.tensor.Tensor): outputs of compute.
-
-    Returns:
-        sch (schedule.Schedule): The created schedule.
-    """
-    device = 'cuda'
-    ctx = akg.tvm.context(device, 0)
-    if not ctx.exist:
-        raise SystemError("Skip because %s is not enabled" % device)
-    with akg.tvm.target.create(device):
-        s = schedule_elemwise(outs)
-    return s
diff --git a/python/akg/ops/array_gpu/__init__.py b/python/akg/ops/array_gpu/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/python/akg/ops/array_gpu/tile.py b/python/akg/ops/array_gpu/tile.py
new file mode 100644
index 0000000000000000000000000000000000000000..e60fcc4ffb094595584b42253a3ca32f699fb0a5
--- /dev/null
+++ b/python/akg/ops/array_gpu/tile.py
@@ -0,0 +1,36 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: tile"""
+import akg.tvm
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, (list, tuple))
+def tile(data, multiples):
+    """
+    Repeats the data in the specified dimensions according to the multiples.
+
+    Args:
+        data (tvm.tensor.Tensor): Tensor.
+        multiples (Union[list, tuple]): Elements must be int. The number of repetitions.
+
+    Returns:
+        tvm.tensor.Tensor, has the same dtype as data.
+    """
+    vc_util.check_shape(data.shape)
+    vc_util.check_int_list(multiples, "multiples")
+    output = akg.topi.tile(data, multiples)
+    return output
diff --git a/python/akg/ops/math_gpu/__init__.py b/python/akg/ops/math_gpu/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/python/akg/ops/math_gpu/cast.py b/python/akg/ops/math_gpu/cast.py
new file mode 100644
index 0000000000000000000000000000000000000000..7266fd60c1daad5009a7b9a57d7d0f883af29c26
--- /dev/null
+++ b/python/akg/ops/math_gpu/cast.py
@@ -0,0 +1,36 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: cast"""
+import akg.tvm
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, str)
+def cast(data, dst_type):
+    """
+    cast data to target type.
+
+    Args:
+        data (tvm.tensor.Tensor): Tensor to be casted.
+        dst_type (str): target cast type.
+
+    Returns:
+        tvm.tensor.Tensor, type is dst_type.
+    """
+    vc_util.check_shape(data.shape)
+    out = akg.topi.cast(data, dst_type)
+
+    return out
diff --git a/python/akg/ops/math_gpu/equal.py b/python/akg/ops/math_gpu/equal.py
new file mode 100644
index 0000000000000000000000000000000000000000..ba06b05524d4d4267a5cc5e74aa3a4b1fb518e5e
--- /dev/null
+++ b/python/akg/ops/math_gpu/equal.py
@@ -0,0 +1,54 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: equal"""
+import akg.tvm
+import akg.topi
+from akg.utils.dsl_create import produce_shapes
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def equal(input1, input2):
+    """
+    check whether input1 equals to input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. If input1 equal to input2 return True, else return False.
+    """
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    shape1, shape2, shape = produce_shapes(shape1, shape2)
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+    dtype = input1.dtype
+
+    # get equal compute
+    t_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(1, dtype), "T")
+    f_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(0, dtype), "F")
+
+    input1_bro = akg.topi.broadcast_to(input1, shape)
+    input2_bro = akg.topi.broadcast_to(input2, shape)
+    c_out = akg.tvm.compute(shape, lambda *indice: akg.tvm.expr.Select(input1_bro[indice] == input2_bro[indice],
+                                                                         t_value[indice], f_value[indice]), name="C")
+    res = akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
+
+    return res
diff --git a/python/akg/ops/math_gpu/greater_equal.py b/python/akg/ops/math_gpu/greater_equal.py
new file mode 100644
index 0000000000000000000000000000000000000000..f149e2849b878d395d22f361a30caf0a74985f2f
--- /dev/null
+++ b/python/akg/ops/math_gpu/greater_equal.py
@@ -0,0 +1,54 @@
+# Copyright 2020 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: greaterequal"""
+import akg.tvm
+import akg.topi
+from akg.utils.dsl_create import produce_shapes
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def greater_equal(input1, input2):
+    """
+    Check whether input1 greaterquals to input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. If input1 greaterquals to input2 return True, else return False.
+    """
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    shape1, shape2, shape = produce_shapes(shape1, shape2)
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+    dtype = input1.dtype
+
+    # get greaterquals compute
+    t_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(1, dtype), "T")
+    f_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(0, dtype), "F")
+
+    input1_bro = akg.topi.broadcast_to(input1, shape)
+    input2_bro = akg.topi.broadcast_to(input2, shape)
+    c_out = akg.tvm.compute(shape, lambda *indice: akg.tvm.expr.Select(input1_bro[indice] >= input2_bro[indice],
+                                                                         t_value[indice], f_value[indice]), name="C")
+    res = akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
+
+    return res
diff --git a/python/akg/ops/math_gpu/less_equal.py b/python/akg/ops/math_gpu/less_equal.py
new file mode 100644
index 0000000000000000000000000000000000000000..521df7db608c42ab33abfae3e508f5abd40eb6a8
--- /dev/null
+++ b/python/akg/ops/math_gpu/less_equal.py
@@ -0,0 +1,54 @@
+# Copyright 2020 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: lessequal"""
+import akg.tvm
+import akg.topi
+from akg.utils.dsl_create import produce_shapes
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def less_equal(input1, input2):
+    """
+    Check whether input1 lessequals to input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. If input1 lessequal to input2 return True, else return False.
+    """
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    shape1, shape2, shape = produce_shapes(shape1, shape2)
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+    dtype = input1.dtype
+
+    # get lessequal compute
+    t_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(1, dtype), "T")
+    f_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(0, dtype), "F")
+
+    input1_bro = akg.topi.broadcast_to(input1, shape)
+    input2_bro = akg.topi.broadcast_to(input2, shape)
+    c_out = akg.tvm.compute(shape, lambda *indice: akg.tvm.expr.Select(input1_bro[indice] <= input2_bro[indice],
+                                                                         t_value[indice], f_value[indice]), name="C")
+    res = akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
+
+    return res
diff --git a/python/akg/ops/math_gpu/logical_and.py b/python/akg/ops/math_gpu/logical_and.py
new file mode 100644
index 0000000000000000000000000000000000000000..8472370c2c2ccd566ed163c4a66a689eaf4b439c
--- /dev/null
+++ b/python/akg/ops/math_gpu/logical_and.py
@@ -0,0 +1,41 @@
+# Copyright 2020 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: logical_and"""
+import akg.tvm
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def logical_and(input1, input2):
+    """
+    Compute logical_and of input1 and input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. LogicalAnd of input1 and input2.
+    """
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    res = akg.topi.logical_and(input1, input2)
+    return res
diff --git a/python/akg/ops/math_gpu/logical_not.py b/python/akg/ops/math_gpu/logical_not.py
new file mode 100644
index 0000000000000000000000000000000000000000..e9401398384af0c6eb617a8e2a98f590913bcad7
--- /dev/null
+++ b/python/akg/ops/math_gpu/logical_not.py
@@ -0,0 +1,32 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: logical_not"""
+import akg.tvm
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor)
+def logical_not(input1):
+    """
+    Compute logical_not of input1.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor.
+    """
+    res = akg.topi.logical_not(input1)
+    return res
diff --git a/python/akg/ops/math_gpu/logical_or.py b/python/akg/ops/math_gpu/logical_or.py
new file mode 100644
index 0000000000000000000000000000000000000000..b1487e6bcd6405cd16675003306af3c85596f21f
--- /dev/null
+++ b/python/akg/ops/math_gpu/logical_or.py
@@ -0,0 +1,41 @@
+# Copyright 2020 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: logical_or"""
+import akg.tvm
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def logical_or(input1, input2):
+    """
+    Compute logical_or of input1 and input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. LogicalOr of input1 and input2.
+    """
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    res = akg.topi.logical_or(input1, input2)
+    return res
diff --git a/python/akg/ops/math_gpu/mean.py b/python/akg/ops/math_gpu/mean.py
new file mode 100644
index 0000000000000000000000000000000000000000..5c738df40faa3b83d52e5cb28dbf729ca0aa4a7e
--- /dev/null
+++ b/python/akg/ops/math_gpu/mean.py
@@ -0,0 +1,47 @@
+# 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_gpu import sum_value
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
+def mean(data, axis=None, keepdims=False):
+    """
+    Computes the mean of the values of a Tensor over the whole dataset.
+
+    Args:
+        data (tvm.tensor.Tensor): Tensor.
+        axis (Union[list, tuple, int, None]): If the tuple is empty, the axis equal to None.
+        keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.
+
+    Returns:
+            tvm.tensor.Tensor, has the same type as data. If keepdims equal to True, all reduced dimensions are
+            retained with length 1. else these reduced axis will be eliminate.
+    """
+    shape = [x.value for x in data.shape]
+    vc_util.reduce_axis_check(shape, axis)
+    axis = ft_util.refine_reduce_axis(data, axis)
+
+    count = 1
+    for i in axis:
+        count *= shape[i]
+    output, _ = sum_value.sum_value(data, axis, keepdims)
+    res = akg.topi.divide(output, count)
+
+    return res
diff --git a/python/akg/ops/math_gpu/mul.py b/python/akg/ops/math_gpu/mul.py
new file mode 100644
index 0000000000000000000000000000000000000000..8377a63d6979a0bbaf56c74531d4f26853d4a838
--- /dev/null
+++ b/python/akg/ops/math_gpu/mul.py
@@ -0,0 +1,43 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: mul"""
+import akg.topi
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def mul(l_input, r_input):
+    """
+    Calculate x * y element-wise.
+
+    Note:
+        mul supports broadcasting.
+
+    Args:
+        l_input (tvm.tensor.Tensor): Tensor.
+        r_input (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor, has the same type as l_input and r_input.
+    """
+    shape1 = [x.value for x in l_input.shape]
+    shape2 = [x.value for x in r_input.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+    vc_util.auto_broadcast_check(shape1, shape2)
+    vc_util.elemwise_dtype_check(l_input.dtype, r_input.dtype)
+    output = akg.topi.multiply(l_input, r_input)
+
+    return output
diff --git a/python/akg/ops/math_gpu/notequal.py b/python/akg/ops/math_gpu/notequal.py
new file mode 100644
index 0000000000000000000000000000000000000000..2ec7ab3432b7e3a809b2f523f8f1f9e7a35b523e
--- /dev/null
+++ b/python/akg/ops/math_gpu/notequal.py
@@ -0,0 +1,54 @@
+# Copyright 2020 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: notequal"""
+import akg.tvm
+import akg.topi
+from akg.utils.dsl_create import produce_shapes
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def notequal(input1, input2):
+    """
+    check whether input1 notequals to input2.
+
+    Args:
+        input1 (tvm.tensor.Tensor): Tensor.
+        input2 (tvm.tensor.Tensor): Tensor.
+
+    Returns:
+        tvm.tensor.Tensor. If input1 notequal to input2 return True, else return False.
+    """
+    shape1 = [x.value for x in input1.shape]
+    shape2 = [x.value for x in input2.shape]
+    vc_util.check_shape(shape1)
+    vc_util.check_shape(shape2)
+
+    shape1, shape2, shape = produce_shapes(shape1, shape2)
+
+    vc_util.elemwise_dtype_check(input1.dtype, input2.dtype)
+    dtype = input1.dtype
+
+    # get notequal compute
+    t_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(1, dtype), "T")
+    f_value = akg.tvm.compute(shape, lambda *indice: akg.tvm.const(0, dtype), "F")
+
+    input1_bro = akg.topi.broadcast_to(input1, shape)
+    input2_bro = akg.topi.broadcast_to(input2, shape)
+    c_out = akg.tvm.compute(shape, lambda *indice: akg.tvm.expr.Select(input1_bro[indice] != input2_bro[indice],
+                                                                         t_value[indice], f_value[indice]), name="C")
+    res = akg.tvm.compute(shape, lambda *indice: c_out(*indice).astype("bool"), name="res")
+
+    return res
diff --git a/python/akg/ops/math_gpu/sub.py b/python/akg/ops/math_gpu/sub.py
new file mode 100644
index 0000000000000000000000000000000000000000..a4a85b0a09feefa37565d90321fe31ddba48626f
--- /dev/null
+++ b/python/akg/ops/math_gpu/sub.py
@@ -0,0 +1,40 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: sub"""
+import akg.topi
+import akg.tvm
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, akg.tvm.tensor.Tensor)
+def sub(data1, data2):
+    """
+    Computes data1 - data2 elementwise, broadcast is supported.
+
+    Args:
+        data1 (tvm.tensor.Tensor): Tensor.
+        data2 (tvm.tensor.Tensor): Tensor of same type as data1, if shape(data2) != shape(data1), broadcast will happen.
+
+    Returns:
+        tvm.tensor.Tensor, subtracted result, with same type as input tensors and broadcasted shape of data1 and data2.
+    """
+    vc_util.elemwise_dtype_check(data1.dtype, data2.dtype)
+    vc_util.check_shape(data1.shape)
+    vc_util.check_shape(data2.shape)
+    vc_util.auto_broadcast_check(data1.shape, data2.shape)
+
+    res = akg.topi.subtract(data1, data2)
+
+    return res
diff --git a/python/akg/ops/math_gpu/sum_value.py b/python/akg/ops/math_gpu/sum_value.py
new file mode 100644
index 0000000000000000000000000000000000000000..ea71bab9c48b78dcc9bed47b42cc5898c36d91d7
--- /dev/null
+++ b/python/akg/ops/math_gpu/sum_value.py
@@ -0,0 +1,45 @@
+# Copyright 2019 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""operator dsl function: sum"""
+
+import akg.topi
+import akg.tvm
+from akg.utils import format_transform as ft_util
+from akg.utils import validation_check as vc_util
+
+
+@vc_util.check_input_type(akg.tvm.tensor.Tensor, (list, tuple, int, type(None)), (bool, type(None)))
+def sum_value(inputs, axis=None, keepdims=False):
+    """
+    Compute the sum of elements across dimensions of a tensor.
+
+    Args:
+        inputs (tvm.tensor.Tensor): Tensor.
+        axis (Union[list, tuple, int, None]): If the list or tuple is empty, the axis equal to None.
+        keepdims (bool): If keepdims equal to True, the result shape length is same to input shape length.
+
+    Returns:
+        tvm.tensor.Tensor, has same type as input. If keepdims is True, all reduced dimensions are retained
+        with length 1, else these reduced axis will be eliminate.
+    """
+    axis = ft_util.refine_reduce_axis(inputs, axis)
+    vc_util.check_shape(inputs.shape)
+
+    if not axis:
+        output = akg.topi.identity(inputs)
+    else:
+        output = akg.topi.sum(inputs, axis=axis, keepdims=keepdims)
+
+    return output
diff --git a/third_party/incubator-tvm/src/pass/split_host_device.cc b/third_party/incubator-tvm/src/pass/split_host_device.cc
index 4f62c7c56b38c8251ac0498a40ccf7407f672e3b..37e9cb4e2838630ee5a87fc7b7b2078f4e475f8c 100644
--- a/third_party/incubator-tvm/src/pass/split_host_device.cc
+++ b/third_party/incubator-tvm/src/pass/split_host_device.cc
@@ -234,7 +234,7 @@ class HostDeviceSplitter : public IRMutator {
       }
     }
 
-#ifdef BACKEND_D
+#ifdef FIX_INPUT_ORDER_TVM
     std::shared_ptr<LoweredFuncNode> na = std::make_shared<LoweredFuncNode>();
     for (unsigned i = 0; i < (unsigned)args_real.size(); i++) {
       bool match = false;