Initial for lwc

cmake/config.cmake

@@ -133,7 +133,7 @@ set(USE_MICRO_STANDALONE_RUNTIME OFF)
 # - OFF: disable llvm, note this will disable CPU codegen
 #        which is needed for most cases
 # - /path/to/llvm-config: enable specific LLVM when multiple llvm-dev is available.
-set(USE_LLVM OFF)
+set(USE_LLVM ON)
 
 #---------------------------------------------
 # Contrib libraries

python/tvm/relay/op/strategy/lwc.py

@@ -0,0 +1,18 @@
+import logging
+import re
+
+from tvm import tir, topi
+from tvm.auto_scheduler import is_auto_scheduler_enabled
+from tvm.meta_schedule import is_meta_schedule_enabled
+from tvm.relay.ty import is_dynamic
+from tvm.target import Target
+from tvm.te import SpecializedCondition
+
+from .. import op as _op
+from .generic import *
+
+logger = logging.getLogger("strategy")
+@schedule_pool.register("x330")
+def schedule_pool_lwc(attrs, outs, target):
+    with target:
+        return topi.lwc.schedule_pool(outs, attrs.layout)

python/tvm/topi/__init__.py

@@ -49,7 +49,7 @@
 from .stft import *
 from . import generic
 from . import nn
-from . import x86
+#from . import x86
 from . import cuda
 from . import gpu
 from . import arm_cpu
@@ -65,7 +65,7 @@
 from . import random
 from . import hexagon
 from . import adreno
-
+from . import lwc
 # error reporting
 from .utils import InvalidShapeError
 

python/tvm/topi/generic/default.py

@@ -24,7 +24,7 @@ def default_schedule(outs, auto_inline):
     """Default schedule for llvm."""
     target = tvm.target.Target.current(allow_none=False)
     outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
-    if target.kind.name not in ("llvm", "c"):
+    if target.kind.name not in ("llvm", "c", "lwc"):
         raise RuntimeError("schedule not registered for '%s'" % target)
     s = te.create_schedule([x.op for x in outs])
     if auto_inline:

python/tvm/topi/generic/injective.py

@@ -56,7 +56,7 @@ def schedule_injective(outs):
         The computation schedule for the op.
     """
     target = tvm.target.Target.current(allow_none=False)
-    if target.kind.name != "llvm":
+    if target.kind.name not in ("llvm", "lwc"):
         raise RuntimeError("schedule_injective not registered for '%s'" % target)
     outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
     x = outs[0]

python/tvm/topi/lwc/__init__.py

@@ -0,0 +1,35 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you 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.
+
+# pylint: disable=wildcard-import
+"""LWC declaration and schedules.
+
+This is a recommended way of using TOPI API.
+To use the generic schedule function, user must set
+the current target scope using with block. See also :any:`tvm.target`
+
+Example
+-------
+.. code-block:: python
+
+  # create schedule that dispatches to topi.cuda.schedule_injective
+  with tvm.target.Target("cuda"):
+    s = tvm.tir.generic.schedule_injective(outs)
+"""
+from __future__ import absolute_import as _abs
+
+from .nn import *

python/tvm/topi/lwc/default.py

@@ -0,0 +1,34 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you 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.
+# pylint: disable=invalid-name,unused-argument
+"""The default schedule used by various operators"""
+import tvm
+from tvm import te
+
+
+def default_schedule(outs, auto_inline):
+    """Default schedule for lwc."""
+    target = tvm.target.Target.current(allow_none=False)
+    outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
+    if target.kind.name not in ("lwc"):
+        raise RuntimeError("schedule not registered for '%s'" % target)
+    s = te.create_schedule([x.op for x in outs])
+    if auto_inline:
+        x = outs[0]
+        te.schedule.AutoInlineInjective(s)
+        s[x].fuse(s[x].op.axis)
+    return s