Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[Hexagon] Add trivial conv2d operator to Hexagon relay strategy #8915

Merged
merged 1 commit into from
Sep 3, 2021
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
[Hexagon] Add trivial conv2d operator to Hexagon relay strategy
This is just to enable relay codegen for conv2d on Hexagon, not for
performance.
Krzysztof Parzyszek committed Sep 2, 2021
commit c464e06de99e2ec9de13357ebcc7bab849f12628
1 change: 1 addition & 0 deletions python/tvm/relay/op/strategy/__init__.py
Original file line number Diff line number Diff line change
@@ -28,3 +28,4 @@
from . import bifrost
from . import rocm
from . import intel_graphics
from . import hexagon
44 changes: 44 additions & 0 deletions python/tvm/relay/op/strategy/hexagon.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,44 @@
# 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.

"""Definition of Hexagon operator strategy."""

# pylint: disable=invalid-name,unused-argument,wildcard-import,unused-wildcard-import

from tvm import topi
from .generic import *
from .. import op as _op


@conv2d_strategy.register("hexagon")
def conv2d_strategy_hexagon(attrs, inputs, out_type, target):
"""Conv2d strategy for Hexagon"""
strategy = _op.OpStrategy()
data_layout = attrs.data_layout
kernel_layout = attrs.kernel_layout

if data_layout == "NHWC" and kernel_layout == "HWIO":
strategy.add_implementation(
wrap_compute_conv2d(topi.nn.conv2d_nhwc),
wrap_topi_schedule(topi.hexagon.schedule_conv2d_nhwc),
name="conv2d.hexagon",
)
return strategy

raise RuntimeError(
"Unsupported layouts: data_layout:{}, kernel_layout:{}".format(data_layout, kernel_layout)
)
1 change: 1 addition & 0 deletions python/tvm/topi/__init__.py
Original file line number Diff line number Diff line change
@@ -61,6 +61,7 @@
from . import sparse
from . import hls
from . import random
from . import hexagon

# error reporting
from .utils import InvalidShapeError
22 changes: 22 additions & 0 deletions python/tvm/topi/hexagon/__init__.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
# 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.

""" Schedules for Hexagon. """

# pylint: disable=wildcard-import

from .conv2d import *
26 changes: 26 additions & 0 deletions python/tvm/topi/hexagon/conv2d.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,26 @@
# 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.

""" Schedules for conv2d. """

import tvm


def schedule_conv2d_nhwc(outs):
"""Schedule for Conv2d NHWC operator."""
s = tvm.te.create_schedule([x.op for x in outs])
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For Hexagon, don't we need mark pipeline attribute no longer anymore?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We're moving towards running entire graphs on Hexagon, so no offloading is necessary. You can still write an op that will use pipeline, but for relay-based compilation we will use ops that run entirely on Hexagon.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we add a possibility to appoint a layer to be offloaded to CPU? I have compared TFLite case and SNPE one for object detection models and it happened that it's better to execute Detection Output on CPU instead of DSP. It might be an exception but such use case is quite widely used

Copy link
Contributor Author

@kparzysz-quic kparzysz-quic Sep 3, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Qualcomm ships products that do this sort of a thing, but I'm not familiar with how the separation between CPU and DSP happens there. In these products, the communication between CPU and DSP is handled by the runtime that is a part of these products, and the offloading does not occur within an individual op. That means that even there there is no need to offload anything explicitly in TVM, because the entire op will run either on CPU or on DSP.

Edit: I think I misunderstood your question---yes SNPE does allow that, but we are not planning to implement this in TVM in the whole-graph compilation mode, at least not in the near term.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I might have misunderstood your message We're moving towards running entire graphs on Hexagon, so no offloading is necessary. My point that sometimes it might be more efficient to offload some parts and give control to user and I got impression that we are going to cancel this ability.

Copy link
Contributor Author

@kparzysz-quic kparzysz-quic Sep 3, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, but we're not planning to support this in the near term.

Edit: we don't have that ability in TVM (for Hexagon) right now, so we're not canceling anything.

return s