From 441d485f63487fb2c71ae9d8d5a72b94664cafc3 Mon Sep 17 00:00:00 2001 From: farshidsp Date: Tue, 26 Apr 2022 14:35:54 -0700 Subject: [PATCH 1/7] Add test for registered scheduales - depthwise_conv2d --- .../topi/test_depthwise_conv2d.py | 310 ++++++++++++++++++ 1 file changed, 310 insertions(+) create mode 100644 tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py new file mode 100644 index 000000000000..202b4c46971e --- /dev/null +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -0,0 +1,310 @@ +# 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. + +import sys + +import numpy as np +import pytest + +import tvm +import tvm.testing +import tvm.topi.testing + +from tvm import te, topi +from tvm.topi.utils import get_const_tuple +from tvm.topi.nn.utils import get_pad_tuple +from ..conftest import requires_hexagon_toolchain + + +random_seed = tvm.testing.parameter(0) + +in_dtype, out_dtype = tvm.testing.parameters( + ("float32", "float32"), +) + +@tvm.testing.fixture +def input_shape(layout, batch, in_channel, in_size, filter_shape): + if layout == "NCHW": + return (batch, in_channel, in_size, in_size) + elif layout == "NHWC": + return (batch, in_size, in_size, in_channel) + elif layout == "NCHWc": + oc_block = filter_shape[-1] + ic_block = next(bn for bn in range(oc_block, 0, -1) if in_channel % bn == 0) + return (batch, in_channel // ic_block, in_size, in_size, ic_block) + + +@tvm.testing.fixture +def filter_shape(layout, in_channel, channel_multiplier, kernel): + filter_channel = in_channel + if layout == "NCHW": + return (filter_channel, channel_multiplier, kernel, kernel) + elif layout == "NHWC": + return (kernel, kernel, filter_channel, channel_multiplier) + elif layout == "NCHWc": + out_channel = in_channel * channel_multiplier + # For testing the functionality, we choose an arbitrary block + # size that can divide out_channel, regardless of the + # performance. + oc_block = next(bn for bn in range(16, 0, -1) if out_channel % bn == 0) + return (out_channel // oc_block, 1, kernel, kernel, 1, oc_block) + + +@tvm.testing.fixture +def scale_shape(layout, in_channel, channel_multiplier, filter_shape): + out_channel = in_channel * channel_multiplier + + if layout in ("NCHW", "NHWC"): + return (out_channel,) + + if layout == "NCHWc": + oc_block = filter_shape[-1] + return (out_channel // oc_block, oc_block) + + raise ValueError("Unknown layout {}".format(layout)) + + +@tvm.testing.fixture +def shift_shape(scale_shape): + return scale_shape + + +@tvm.testing.fixture(cache_return_value=True) +def ref_data( + random_seed, + in_dtype, + out_dtype, + layout, + input_shape, + filter_shape, + dilation, + stride, + padding, + scale_shape, + shift_shape, + use_scale_shift, + apply_relu, +): + np.random.seed(random_seed) + + print(input_shape) + + # scipy.signal.convolve2d does not support float16 data types, and + # the python fallback is too slow for general use. Computing + # ref_data in float32 will have fewer rounding errors than the TVM + # float16 compute, but those vary based on schedule anyways. + conv_dtype = "float32" if in_dtype == "float16" else in_dtype + + input_np = np.random.uniform(size=input_shape).astype(in_dtype) + filter_np = np.random.uniform(size=filter_shape).astype(in_dtype) + scale_np = np.random.uniform(size=scale_shape).astype(out_dtype) + shift_np = np.random.uniform(size=shift_shape).astype(out_dtype) + if layout == "NCHW": + np_depthwise_conv2d = tvm.topi.testing.depthwise_conv2d_python_nchw + dilation = (1, 1, dilation, dilation) + reshape = (1, -1, 1, 1) + elif layout == "NHWC": + np_depthwise_conv2d = tvm.topi.testing.depthwise_conv2d_python_nhwc + dilation = (dilation, dilation, 1, 1) + reshape = (1, 1, 1, -1) + elif layout == "NCHWc": + np_depthwise_conv2d = tvm.topi.testing.depthwise_conv2d_python_nchwc + dilation = (1, 1, dilation, dilation, 1, 1) + reshape = (1, scale_shape[0], 1, 1, scale_shape[1]) + + dilated_filter_np = tvm.topi.testing.dilate_python(filter_np, dilation) + output_np = np_depthwise_conv2d( + input_np.astype(conv_dtype), dilated_filter_np.astype(conv_dtype), stride, padding + ).astype(out_dtype) + + if use_scale_shift: + output_np = output_np * scale_np.reshape(reshape) + shift_np.reshape(reshape) + if apply_relu: + output_np = np.maximum(output_np, 0) + + return ( + input_np, + filter_np, + scale_np, + shift_np, + output_np, + ) + + +class BaseDepthwiseConv2D: + """Provides the test_conv2d test function, to be used by other test classes. + + Test parameter sets are split out into different classes for + readability (e.g. used for mobilenet), and for restrictions + (e.g. implemented only for llvm). + """ + + # layout = tvm.testing.parameter("NCHW", "NHWC") + + # (batch, in_channel, in_size, channel_multiplier, kernel, stride) = tvm.testing.parameters( + # (1, 728, 32, 1, 3, 1), + # (4, 256, 64, 2, 5, 2), + # ) + # padding = tvm.testing.parameter("SAME", "VALID") + # dilation = tvm.testing.parameter(1, 2) + + use_scale_shift = tvm.testing.parameter(False, ids=["no_scale_shift"]) + apply_relu = tvm.testing.parameter(False, ids=["no_relu"]) + + run_after_compile = True + + @requires_hexagon_toolchain + def test_conv2d( + self, + hexagon_session, + target, + dev, + in_dtype, + out_dtype, + layout, + input_shape, + filter_shape, + scale_shape, + shift_shape, + use_scale_shift, + apply_relu, + batch, + in_channel, + channel_multiplier, + kernel, + stride, + padding, + dilation, + ref_data, + ): + target_hexagon = tvm.target.hexagon("v68") + + + # Transform the padding argument from 'str' to 'tuple' to + # match the "workload" tuple in TopHub. Which padding_args to + # use for each layout chosen to reproduce previous behavior. + if dilation == 1: + padding_args = get_pad_tuple(padding, (kernel, kernel)) + padding_args_i = [0, 1, 2, 3] if layout == "NCHW" else [0, 1] + padding_args = [padding_args[i] for i in padding_args_i] + else: + padding_args = padding + + # placeholder + Input = te.placeholder(input_shape, name="Input", dtype=in_dtype) + Filter = te.placeholder(filter_shape, name="Filter", dtype=in_dtype) + Scale = te.placeholder(scale_shape, name="Scale", dtype=out_dtype) + Shift = te.placeholder(shift_shape, name="Shift", dtype=out_dtype) + + if layout == "NCHW": + topi_scale_shift = topi.nn.scale_shift_nchw + fcompute_args = (Input, Filter, stride, padding_args, dilation, out_dtype) + + elif layout == "NHWC": + topi_scale_shift = topi.nn.scale_shift_nhwc + fcompute_args = (Input, Filter, stride, padding_args, dilation, out_dtype) + + elif layout == "NCHWc": + topi_scale_shift = topi.nn.scale_shift_nchwc + in_layout = "NCHW{}c".format(input_shape[-1]) + out_layout = "NCHW{}c".format(filter_shape[-1]) + fcompute_args = ( + Input, + Filter, + stride, + padding, + dilation, + in_layout, + out_layout, + out_dtype, + ) + + + with tvm.target.Target(target_hexagon): + # Declare, build schedule + if layout == "NCHW": + fcompute = topi.nn.depthwise_conv2d_nchw + fschedule = topi.hexagon.schedule_depthwise_conv2d_nchw + elif layout == "NHWC": + fcompute = topi.nn.depthwise_conv2d_nhwc + fschedule = topi.hexagon.schedule_depthwise_conv2d_nhwc + C = fcompute(*fcompute_args) + if use_scale_shift: + C = topi_scale_shift(C, Scale, Shift) + if apply_relu: + C = topi.nn.relu(C) + + s = fschedule([C]) + + # Build and run + f = tvm.build(s, [Input, Filter, Scale, Shift, C], tvm.target.Target(target_hexagon, host=target_hexagon)) + mod = hexagon_session.load_module(f) + + input_np, filter_np, scale_np, shift_np, output_np = ref_data + + dev = hexagon_session.device + input_tvm = tvm.nd.array(input_np, dev) + filter_tvm = tvm.nd.array(filter_np, dev) + scale_tvm = tvm.nd.array(scale_np, dev) + shift_tvm = tvm.nd.array(shift_np, dev) + output_tvm = tvm.nd.array( + np.zeros(shape=get_const_tuple(C.shape), dtype=C.dtype), + dev, + ) + + mod(input_tvm, filter_tvm, scale_tvm, shift_tvm, output_tvm) + + tol = {"rtol": 1e-4, "atol": 1e-5} + tvm.testing.assert_allclose(output_np, output_tvm.numpy(), **tol) + + +class TestDepthwiseConv2D_MobilenetWorkloads_small(BaseDepthwiseConv2D): + """Extra tests to verify functionality for workloads used by mobilenet.""" + + layout = tvm.testing.parameter("NCHW", "NHWC") + + batch = tvm.testing.parameter(1) + channel_multiplier = tvm.testing.parameter(1) + kernel = tvm.testing.parameter(3) + padding = tvm.testing.parameter("SAME") + dilation = tvm.testing.parameter(1) + + in_channel, in_size, stride = tvm.testing.parameters( + (32, 112, 1), + (64, 112, 2), + (128, 56, 1), + (128, 56, 2), + ) + +# class TestDepthwiseConv2D_MobilenetWorkloads_large(BaseDepthwiseConv2D): +# """Extra tests to verify functionality for workloads used by mobilenet.""" + +# layout = tvm.testing.parameter("NCHW", "NHWC") + +# batch = tvm.testing.parameter(1) +# channel_multiplier = tvm.testing.parameter(1) +# kernel = tvm.testing.parameter(3) +# padding = tvm.testing.parameter("SAME") +# dilation = tvm.testing.parameter(1) + +# in_channel, in_size, stride = tvm.testing.parameters( +# (256, 28, 1), +# (256, 28, 2), +# (512, 14, 1), +# (512, 14, 2), +# (1024, 7, 1), +# ) \ No newline at end of file From 59dd7e5d8208722de2d8b05dd13992d8db432a2a Mon Sep 17 00:00:00 2001 From: farshidsp Date: Tue, 26 Apr 2022 16:59:20 -0700 Subject: [PATCH 2/7] added more test to depthwise_conv2 --- .../topi/test_depthwise_conv2d.py | 42 ++++++------------- 1 file changed, 12 insertions(+), 30 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index 202b4c46971e..ebd03b9984c0 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -153,26 +153,13 @@ class BaseDepthwiseConv2D: (e.g. implemented only for llvm). """ - # layout = tvm.testing.parameter("NCHW", "NHWC") - - # (batch, in_channel, in_size, channel_multiplier, kernel, stride) = tvm.testing.parameters( - # (1, 728, 32, 1, 3, 1), - # (4, 256, 64, 2, 5, 2), - # ) - # padding = tvm.testing.parameter("SAME", "VALID") - # dilation = tvm.testing.parameter(1, 2) - use_scale_shift = tvm.testing.parameter(False, ids=["no_scale_shift"]) apply_relu = tvm.testing.parameter(False, ids=["no_relu"]) - run_after_compile = True - @requires_hexagon_toolchain def test_conv2d( self, hexagon_session, - target, - dev, in_dtype, out_dtype, layout, @@ -272,7 +259,7 @@ def test_conv2d( tvm.testing.assert_allclose(output_np, output_tvm.numpy(), **tol) -class TestDepthwiseConv2D_MobilenetWorkloads_small(BaseDepthwiseConv2D): +class TestDepthwiseConv2D_MobilenetWorkloads(BaseDepthwiseConv2D): """Extra tests to verify functionality for workloads used by mobilenet.""" layout = tvm.testing.parameter("NCHW", "NHWC") @@ -288,23 +275,18 @@ class TestDepthwiseConv2D_MobilenetWorkloads_small(BaseDepthwiseConv2D): (64, 112, 2), (128, 56, 1), (128, 56, 2), + (256, 28, 1), ) -# class TestDepthwiseConv2D_MobilenetWorkloads_large(BaseDepthwiseConv2D): -# """Extra tests to verify functionality for workloads used by mobilenet.""" +class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): -# layout = tvm.testing.parameter("NCHW", "NHWC") - -# batch = tvm.testing.parameter(1) -# channel_multiplier = tvm.testing.parameter(1) -# kernel = tvm.testing.parameter(3) -# padding = tvm.testing.parameter("SAME") -# dilation = tvm.testing.parameter(1) + layout = tvm.testing.parameter("NCHW", "NHWC") + use_scale_shift = tvm.testing.parameter(True, False, ids=["with_scale_shift", "no_scale_shift"]) + apply_relu = tvm.testing.parameter(True, False, ids=["with_relu", "no_relu"]) -# in_channel, in_size, stride = tvm.testing.parameters( -# (256, 28, 1), -# (256, 28, 2), -# (512, 14, 1), -# (512, 14, 2), -# (1024, 7, 1), -# ) \ No newline at end of file + (batch, in_channel, in_size, channel_multiplier, kernel, stride) = tvm.testing.parameters( + (1, 64, 32, 1, 3, 1), + (1, 128, 64, 2, 5, 2), + ) + padding = tvm.testing.parameter("VALID") + dilation = tvm.testing.parameter(1) \ No newline at end of file From ad15633d5d0fb2d0e55f5e0aeb1213addf401610 Mon Sep 17 00:00:00 2001 From: farshidsp Date: Tue, 26 Apr 2022 18:05:23 -0700 Subject: [PATCH 3/7] adding new line at the end of the file --- .../python/contrib/test_hexagon/topi/test_depthwise_conv2d.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index ebd03b9984c0..c8462cbdddcd 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -289,4 +289,5 @@ class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): (1, 128, 64, 2, 5, 2), ) padding = tvm.testing.parameter("VALID") - dilation = tvm.testing.parameter(1) \ No newline at end of file + dilation = tvm.testing.parameter(1) + \ No newline at end of file From ab8b529c4ed6c6e4ecddbaacc6373757373eb27d Mon Sep 17 00:00:00 2001 From: farshidsp Date: Wed, 27 Apr 2022 10:16:04 -0700 Subject: [PATCH 4/7] reformatted the file --- .../test_hexagon/topi/test_depthwise_conv2d.py | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index c8462cbdddcd..786f50c536e1 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -36,6 +36,7 @@ ("float32", "float32"), ) + @tvm.testing.fixture def input_shape(layout, batch, in_channel, in_size, filter_shape): if layout == "NCHW": @@ -180,7 +181,6 @@ def test_conv2d( ): target_hexagon = tvm.target.hexagon("v68") - # Transform the padding argument from 'str' to 'tuple' to # match the "workload" tuple in TopHub. Which padding_args to # use for each layout chosen to reproduce previous behavior. @@ -220,7 +220,6 @@ def test_conv2d( out_dtype, ) - with tvm.target.Target(target_hexagon): # Declare, build schedule if layout == "NCHW": @@ -238,7 +237,11 @@ def test_conv2d( s = fschedule([C]) # Build and run - f = tvm.build(s, [Input, Filter, Scale, Shift, C], tvm.target.Target(target_hexagon, host=target_hexagon)) + f = tvm.build( + s, + [Input, Filter, Scale, Shift, C], + tvm.target.Target(target_hexagon, host=target_hexagon), + ) mod = hexagon_session.load_module(f) input_np, filter_np, scale_np, shift_np, output_np = ref_data @@ -278,6 +281,7 @@ class TestDepthwiseConv2D_MobilenetWorkloads(BaseDepthwiseConv2D): (256, 28, 1), ) + class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): layout = tvm.testing.parameter("NCHW", "NHWC") @@ -290,4 +294,3 @@ class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): ) padding = tvm.testing.parameter("VALID") dilation = tvm.testing.parameter(1) - \ No newline at end of file From e5b567095e6b88640ffc345cfc79f742f31f6424 Mon Sep 17 00:00:00 2001 From: farshidsp Date: Wed, 27 Apr 2022 11:51:31 -0700 Subject: [PATCH 5/7] resolve comments --- .../contrib/test_hexagon/topi/test_depthwise_conv2d.py | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index 786f50c536e1..6343a10f1f77 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -154,9 +154,6 @@ class BaseDepthwiseConv2D: (e.g. implemented only for llvm). """ - use_scale_shift = tvm.testing.parameter(False, ids=["no_scale_shift"]) - apply_relu = tvm.testing.parameter(False, ids=["no_relu"]) - @requires_hexagon_toolchain def test_conv2d( self, @@ -266,6 +263,8 @@ class TestDepthwiseConv2D_MobilenetWorkloads(BaseDepthwiseConv2D): """Extra tests to verify functionality for workloads used by mobilenet.""" layout = tvm.testing.parameter("NCHW", "NHWC") + use_scale_shift = tvm.testing.parameter(False, ids=["no_scale_shift"]) + apply_relu = tvm.testing.parameter(False, ids=["no_relu"]) batch = tvm.testing.parameter(1) channel_multiplier = tvm.testing.parameter(1) @@ -282,7 +281,7 @@ class TestDepthwiseConv2D_MobilenetWorkloads(BaseDepthwiseConv2D): ) -class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): +class TestDepthwiseConv2D(BaseDepthwiseConv2D): layout = tvm.testing.parameter("NCHW", "NHWC") use_scale_shift = tvm.testing.parameter(True, False, ids=["with_scale_shift", "no_scale_shift"]) @@ -294,3 +293,6 @@ class TestDepthwiseConv2D_More(BaseDepthwiseConv2D): ) padding = tvm.testing.parameter("VALID") dilation = tvm.testing.parameter(1) + + +# TODO(hexagon-team): add TestDepthwiseConv2D_NCHWc test. From 2cde8dbcee49c330376fd46b0beacfcaf16c42af Mon Sep 17 00:00:00 2001 From: farshidsp Date: Thu, 28 Apr 2022 18:11:27 -0700 Subject: [PATCH 6/7] add schedule and tests for conv2d_transpose_nchw --- python/tvm/topi/hexagon/conv2d.py | 26 +++ .../topi/test_conv2d_transpose.py | 177 ++++++++++++++++++ 2 files changed, 203 insertions(+) create mode 100644 tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py diff --git a/python/tvm/topi/hexagon/conv2d.py b/python/tvm/topi/hexagon/conv2d.py index 4f564faa0ab4..d8f44d663843 100644 --- a/python/tvm/topi/hexagon/conv2d.py +++ b/python/tvm/topi/hexagon/conv2d.py @@ -18,6 +18,7 @@ """Schedule for conv2d""" import tvm +from ..utils import traverse_inline def schedule_conv2d_nhwc(outs): @@ -60,3 +61,28 @@ def schedule_depthwise_conv2d_nchw(outs): def schedule_depthwise_conv2d_nhwc(out): return schedule_conv2d_nhwc(out) + + +def schedule_conv2d_transpose_nchw(outs): + """Create schedule for tensors""" + outs = [outs] if isinstance(outs, tvm.te.tensor.Tensor) else outs + s = schedule_conv2d_nchw(outs) + + def _callback(op): + if "unpack_nchwc" in op.tag: + conv_out = op.input_tensors[0] + # retrieve data + data_vec = conv_out.op.input_tensors[0] + if isinstance(data_vec, tvm.te.ComputeOp): + data_pad = data_vec.op.input_tensors[0] + data_dilate = data_pad.op.input_tensors[0] + s[data_dilate].compute_inline() + s[data_pad].compute_inline() + # retrieve kernel + kernel_vec = conv_out.op.input_tensors[1] + if isinstance(kernel_vec, tvm.te.ComputeOp): + kernel_transform = kernel_vec.op.input_tensors[0] + s[kernel_transform].compute_inline() + + traverse_inline(s, outs[0].op, _callback) + return s diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py new file mode 100644 index 000000000000..0afa407ee03e --- /dev/null +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py @@ -0,0 +1,177 @@ +# 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. +"""Test code for transposed convolution.""" +import numpy as np +import tvm +import tvm.testing +from tvm import te +from tvm import topi +import tvm.topi.testing +from tvm.contrib.pickle_memoize import memoize +from tvm.topi.utils import get_const_tuple +from ..conftest import requires_hexagon_toolchain + + +# TODO Should add kernal to tvm.testing.fixture + +random_seed = tvm.testing.parameter(0) + + +@tvm.testing.fixture +def shift_shape(batch): + return batch + + +@tvm.testing.fixture +def shift_shape(in_channel): + return in_channel + + +@tvm.testing.fixture +def shift_shape(in_size): + return in_size + + +@tvm.testing.fixture +def shift_shape(num_filter): + return num_filter + + +@tvm.testing.fixture +def shift_shape(stride): + return stride + + +@tvm.testing.fixture +def shift_shape(padding): + return padding + + +@tvm.testing.fixture +def shift_shape(output_padding): + return output_padding + + +class BaseConv2DTransposeTests: + @requires_hexagon_toolchain + def test_conv2d( + self, + hexagon_session, + batch, + in_channel, + in_size, + num_filter, + stride, + padding, + output_padding, + random_seed, + ): + + target_hexagon = tvm.target.hexagon("v68") + + in_height, in_width = in_size + kernel_height, kernel_width = (1, 1) + stride_height, stride_width = stride + pad_top, pad_left, pad_bottom, pad_right = padding + + A = te.placeholder((batch, in_channel, in_height, in_width), name="A") + W = te.placeholder((in_channel, num_filter, kernel_height, kernel_width), name="W") + + a_shape = get_const_tuple(A.shape) + w_shape = get_const_tuple(W.shape) + dtype = A.dtype + + def get_ref_data(): + + np.random.seed(random_seed) + a_np = np.random.uniform(size=a_shape).astype(dtype) + w_np = np.random.uniform(size=w_shape).astype(dtype) + b_np = tvm.topi.testing.conv2d_transpose_nchw_python( + a_np, w_np, stride, padding, output_padding + ) + c_np = np.maximum(b_np, 0) + return a_np, w_np, b_np, c_np + + a_np, w_np, b_np, c_np = get_ref_data() + + fcompute_args = ( + A, + W, + [stride_height, stride_width], + [pad_top, pad_left, pad_bottom, pad_right], + A.dtype, + output_padding, + ) + + with tvm.target.Target(target_hexagon): + fcompute = topi.nn.conv2d_transpose_nchw + fschedule = topi.hexagon.schedule_conv2d_transpose_nchw + B = fcompute(*fcompute_args) + C = topi.nn.relu(B) + s1 = fschedule([B]) + s2 = fschedule([C]) + + dev = hexagon_session.device + + a = tvm.nd.array(a_np, dev) + w = tvm.nd.array(w_np, dev) + b = tvm.nd.array(np.zeros(get_const_tuple(B.shape), dtype=B.dtype), dev) + c = tvm.nd.array(np.zeros(get_const_tuple(C.shape), dtype=C.dtype), dev) + + func1 = tvm.build(s1, [A, W, B], tvm.target.Target(target_hexagon, host=target_hexagon)) + func2 = tvm.build(s2, [A, W, C], tvm.target.Target(target_hexagon, host=target_hexagon)) + + mod1 = hexagon_session.load_module(func1) + mod2 = hexagon_session.load_module(func2) + + mod1(a, w, b) + mod2(a, w, c) + tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-5) + tvm.testing.assert_allclose(c.numpy(), c_np, rtol=1e-5) + + +class TestConv2DTranspose_1(BaseConv2DTransposeTests): + + batch = tvm.testing.parameter(1) + in_channel = tvm.testing.parameter(3, 8) + in_size = tvm.testing.parameter((224, 224)) + num_filter = tvm.testing.parameter(1, 8) + stride = tvm.testing.parameter((1, 1)) + padding = tvm.testing.parameter((0, 0, 0, 0)) + output_padding = tvm.testing.parameter((0, 0)) + + +class TestConv2DTranspose_2(BaseConv2DTransposeTests): + + batch = tvm.testing.parameter(1) + in_channel = tvm.testing.parameter(512) + in_size = tvm.testing.parameter((8, 1)) + num_filter = tvm.testing.parameter(128) + stride = tvm.testing.parameter((31, 1)) + padding = tvm.testing.parameter((0, 0, 0, 0)) + output_padding = tvm.testing.parameter((0, 0)) + + +class TestConv2DTranspose_3(BaseConv2DTransposeTests): + + batch = tvm.testing.parameter(1) + in_channel = tvm.testing.parameter(32) + in_size = tvm.testing.parameter((8192, 1)) + num_filter = tvm.testing.parameter(1) + stride = tvm.testing.parameter((1, 1)) + padding = tvm.testing.parameter((0, 0, 0, 0)) + output_padding = tvm.testing.parameter((0, 0)) From 1f9f97efc4e004b95b2088096c381fcd78b40a76 Mon Sep 17 00:00:00 2001 From: farshidsp Date: Tue, 3 May 2022 10:07:22 -0700 Subject: [PATCH 7/7] registering conv2d_transpose strategy and clean up test --- python/tvm/relay/op/strategy/hexagon.py | 20 +++++++++++ .../topi/test_conv2d_transpose.py | 34 ++++--------------- 2 files changed, 27 insertions(+), 27 deletions(-) diff --git a/python/tvm/relay/op/strategy/hexagon.py b/python/tvm/relay/op/strategy/hexagon.py index cfd9a8b5ddc2..da15a5412517 100644 --- a/python/tvm/relay/op/strategy/hexagon.py +++ b/python/tvm/relay/op/strategy/hexagon.py @@ -112,6 +112,26 @@ def softmax_strategy_hexagon(attrs, inputs, out_type, target): return strategy +@conv2d_transpose_strategy.register("hexagon") +def conv2d_transpose_strategy_hexagon(attrs, inputs, out_type, target): + """conv2d_transpose hexagon strategy""" + layout = attrs.data_layout + dilation = get_const_tuple(attrs.dilation) + groups = attrs.groups + assert layout == "NCHW", "only support nchw for now" + assert dilation == (1, 1), "not support dilate now" + strategy = _op.OpStrategy() + if groups == 1: + strategy.add_implementation( + wrap_compute_conv2d_transpose(topi.nn.conv2d_transpose_nchw), + wrap_topi_schedule(topi.hexagon.schedule_conv2d_transpose_nchw), + name="conv2d_transpose_nchw.generic", + ) + else: + raise RuntimeError("Unsupported conv2d_transpose layout {}".format(layout)) + return strategy + + # --- Op schedule registration diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py index 0afa407ee03e..1dbac67aeb76 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py @@ -144,34 +144,14 @@ def get_ref_data(): tvm.testing.assert_allclose(c.numpy(), c_np, rtol=1e-5) -class TestConv2DTranspose_1(BaseConv2DTransposeTests): +class TestConv2DTranspose(BaseConv2DTransposeTests): - batch = tvm.testing.parameter(1) - in_channel = tvm.testing.parameter(3, 8) - in_size = tvm.testing.parameter((224, 224)) - num_filter = tvm.testing.parameter(1, 8) - stride = tvm.testing.parameter((1, 1)) - padding = tvm.testing.parameter((0, 0, 0, 0)) - output_padding = tvm.testing.parameter((0, 0)) - - -class TestConv2DTranspose_2(BaseConv2DTransposeTests): - - batch = tvm.testing.parameter(1) - in_channel = tvm.testing.parameter(512) - in_size = tvm.testing.parameter((8, 1)) - num_filter = tvm.testing.parameter(128) - stride = tvm.testing.parameter((31, 1)) - padding = tvm.testing.parameter((0, 0, 0, 0)) - output_padding = tvm.testing.parameter((0, 0)) - - -class TestConv2DTranspose_3(BaseConv2DTransposeTests): + (batch, in_channel, in_size, num_filter, stride) = tvm.testing.parameters( + (1, 3, (224, 224), 1, (1, 1)), + (1, 8, (224, 224), 1, (1, 1)), + (1, 512, (8, 1), 128, (31, 1)), + (1, 32, (8192, 1), 1, (1, 1)), + ) - batch = tvm.testing.parameter(1) - in_channel = tvm.testing.parameter(32) - in_size = tvm.testing.parameter((8192, 1)) - num_filter = tvm.testing.parameter(1) - stride = tvm.testing.parameter((1, 1)) padding = tvm.testing.parameter((0, 0, 0, 0)) output_padding = tvm.testing.parameter((0, 0))