From 6dd1566002bb089e6c84edd7cb5c542a912d32d6 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:11:32 -0700 Subject: [PATCH 01/35] fix a lot of initial tests --- python/tvm/relay/op/nn/nn.py | 2 +- python/tvm/relay/qnn/op/layout_conversions.py | 2 +- python/tvm/relay/qnn/op/qnn.py | 6 +- python/tvm/relay/testing/dcgan.py | 3 +- .../transform/fake_quantization_to_integer.py | 18 +++++ python/tvm/topi/nn/conv2d_transpose.py | 76 ++++++++++++------- python/tvm/topi/nn/utils.py | 2 + src/relay/op/nn/convolution.h | 6 +- .../relay/test_autotvm_task_extraction.py | 3 +- tests/python/relay/test_op_level2.py | 60 +++++++++------ .../relay/test_pass_convert_op_layout.py | 16 ++-- .../test_pass_fake_quantization_to_integer.py | 24 +++++- 12 files changed, 147 insertions(+), 71 deletions(-) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 5a17db745b3e..8baa96e9bab0 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -522,7 +522,7 @@ def conv2d_transpose( channels=None, kernel_size=None, data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_layout="", output_padding=(0, 0), out_dtype="", diff --git a/python/tvm/relay/qnn/op/layout_conversions.py b/python/tvm/relay/qnn/op/layout_conversions.py index 1a3b1771d6ce..24c787e0a00a 100644 --- a/python/tvm/relay/qnn/op/layout_conversions.py +++ b/python/tvm/relay/qnn/op/layout_conversions.py @@ -119,7 +119,7 @@ def convert_qnn_conv2d_transpose(attrs, inputs, tinfos, desired_layouts): # Handle default kernel layouts if desired_data_layout == "NCHW": - new_attrs["kernel_layout"] = "OIHW" + new_attrs["kernel_layout"] = "IOHW" return relay.qnn.op.conv2d_transpose(*inputs, **new_attrs) if desired_data_layout == "NHWC": new_attrs["kernel_layout"] = "HWIO" diff --git a/python/tvm/relay/qnn/op/qnn.py b/python/tvm/relay/qnn/op/qnn.py index 83b5cf0a831c..d1121f7b90fc 100644 --- a/python/tvm/relay/qnn/op/qnn.py +++ b/python/tvm/relay/qnn/op/qnn.py @@ -18,13 +18,15 @@ """QNN dialect operators.""" from __future__ import absolute_import as _abs + from tvm import relay from tvm.relay.expr import Tuple, TupleWrapper from tvm.relay.op.nn.utils import get_pad_tuple2d from tvm.topi.nn.qnn import SQNN_DTYPE_TO_CODE -from . import _make + from ... import op as reg from ...op import OpPattern +from . import _make def requantize( @@ -382,7 +384,7 @@ def conv2d_transpose( channels=None, kernel_size=None, data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_layout="", output_padding=(0, 0), out_dtype="", diff --git a/python/tvm/relay/testing/dcgan.py b/python/tvm/relay/testing/dcgan.py index fc531b765a88..acc478330dc9 100644 --- a/python/tvm/relay/testing/dcgan.py +++ b/python/tvm/relay/testing/dcgan.py @@ -27,6 +27,7 @@ arXiv preprint arXiv:1511.06434 (2015). """ from tvm import relay + from . import layers from .init import create_workload @@ -41,7 +42,7 @@ def deconv2d(data, ishape, oshape, kshape, layout, name, stride=(2, 2)): adj_x = (target_shape[1] + 2 * pad_x - kshape[1]) % stride[1] if layout == "NCHW": - kernel_layout = "OIHW" + kernel_layout = "IOHW" elif layout == "NHWC": kernel_layout = "HWOI" else: diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index 1adde9a4a430..ecf5d9581bd8 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -19,6 +19,7 @@ from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType from tvm.tir import bijective_layout + from ..op import register_fake_quantization_to_integer @@ -156,6 +157,23 @@ def conv2d(expr, type_map): return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] +@register_fake_quantization_to_integer("nn.conv2d_transpose") +def conv2d_transpose(expr, type_map): + attrs = {**expr.attrs} + attrs.pop("out_dtype") + x, weight = expr.args + x_t = type_map[x] + w_t = type_map[weight] + conv_scale = fold_constant(x_t.scale * w_t.scale) + conv_zp = get_zeros(conv_scale) + out = relay.qnn.op.conv2d_transpose( + x, weight, x_t.zero_point, w_t.zero_point, x_t.zero_point, w_t.zero_point, **attrs + ) + out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] + out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] + return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] + + @register_fake_quantization_to_integer("nn.dense") def dense(expr, type_map): """Rewrite a dense op""" diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index 22188bcd45a4..8291bfed4ea6 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -17,12 +17,12 @@ # pylint: disable=invalid-name, unused-variable, unused-argument """Transposed 2D convolution operators (sometimes called Deconvolution).""" import tvm -from tvm import te -from tvm import relay +from tvm import relay, te + +from ..utils import simplify from .dilate import dilate from .pad import pad from .utils import get_pad_tuple -from ..utils import simplify def conv2d_transpose_nchw(Input, Filter, strides, padding, out_dtype, output_padding): @@ -116,6 +116,41 @@ def declaration_conv2d_transpose_impl(data, kernel, strides, padding, out_dtype, return Output +def layout_transform(tensor: "relay.Expr", current_layout: str, desired_layout: str): + """Transform a tensor with the current layout to the desired layout. + + E.g. layout_transform(t, "NCHW", "CNHW") --> relay.transpose(t, [1, 0, 2, 3]) + + Parameters + ---------- + tensor: relay.Expr + The Tensor to transpose + + current_layout: str + The current layout e.g. NCHW or OIHW + + desired_layout: str + The desired layout, must be compatible with current_layout + + Returns + ------- + The layout_transformed tensor. + """ + if sorted(current_layout) != sorted(desired_layout): + raise ValueError(f"Incompatible layouts: {current_layout} vs {desired_layout}") + + if current_layout == desired_layout: + return tensor + + current_layout_map = {c: i for i, c in enumerate(current_layout)} + desired_layout_map = {c: i for i, c in enumerate(desired_layout)} + + axes = [None] * len(current_layout) + for c, i in current_layout_map.items(): + axes[i] = desired_layout_map[c] + return relay.transpose(tensor, axes=axes) + + @tvm.target.generic_func def conv2d_transpose_legalize(attrs, inputs, types): """Legalizes Transposed 2D convolution op. @@ -134,36 +169,16 @@ def conv2d_transpose_legalize(attrs, inputs, types): result : tvm.relay.Expr The legalized expr """ + data, kernel = inputs + kernel_layout = attrs["kernel_layout"] if attrs["data_layout"] == "NHWC": - data, kernel = inputs - kernel_layout = attrs["kernel_layout"] - # Convert Kernel layout to IOHW - # kernel_layout is different from input kernel layout - IO is swapped - if kernel_layout == "HWIO": - # input kernel layout is swapped to HWOI - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(3, 2, 0, 1)) - elif kernel_layout == "HWOI": - # input kernel layout is swapped to HWIO - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(2, 3, 0, 1)) - elif kernel_layout == "IOHW": - # input kernel layout is swapped to OIHW - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(1, 0, 2, 3)) - elif kernel_layout == "OIHW": - # input kernel layout is swapped to IOHW - # output kernel layout will be IOHW - pass - else: - # Skip legalize. Let relay.nn.conv2d_transpose to handle the case - return None + kernel = layout_transform(kernel, kernel_layout, "IOHW") # Set new attrs for conv2d_transpose. new_attrs = {k: attrs[k] for k in attrs.keys()} new_attrs["data_layout"] = "NCHW" # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW - new_attrs["kernel_layout"] = "OIHW" + new_attrs["kernel_layout"] = "IOHW" # Convert data to NCHW. data = relay.transpose(data, axes=(0, 3, 1, 2)) @@ -171,5 +186,12 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Convert back to original NHWC layout. out = relay.transpose(deconv, axes=(0, 2, 3, 1)) return out + elif attrs["data_layout"] == "NCHW": + kernel = layout_transform(kernel, kernel_layout, "IOHW") + new_attrs = {k: attrs[k] for k in attrs.keys()} + + # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + new_attrs["kernel_layout"] = "IOHW" + return relay.nn.conv2d_transpose(data, kernel, **new_attrs) return None diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index ff00441e9850..eea35c814c11 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,6 +19,8 @@ from __future__ import absolute_import import tvm +from tvm import relay + from ..utils import get_const_int diff --git a/src/relay/op/nn/convolution.h b/src/relay/op/nn/convolution.h index c27227b2eb73..d9958076adc1 100644 --- a/src/relay/op/nn/convolution.h +++ b/src/relay/op/nn/convolution.h @@ -1044,7 +1044,7 @@ bool Conv2DTransposeRel(const Array& types, int num_inputs, const Attrs& a if (data == nullptr) return false; static const Layout kNCHW("NCHW"); - static const Layout kOIHW("OIHW"); + static const Layout kIOHW("IOHW"); const Conv2DTransposeAttrs* param = attrs.as(); ICHECK(param != nullptr); @@ -1056,9 +1056,9 @@ bool Conv2DTransposeRel(const Array& types, int num_inputs, const Attrs& a << "Conv only support input layouts that are convertible from NCHW." << " But got " << in_layout; - const auto trans_kernel_layout = tir::BijectiveLayout(kernel_layout, kOIHW); + const auto trans_kernel_layout = tir::BijectiveLayout(kernel_layout, kIOHW); ICHECK(trans_kernel_layout.defined()) - << "Conv only support kernel layouts that are convertible from OIHW." + << "Conv only support kernel layouts that are convertible from IOHW." << " But got " << kernel_layout; Layout out_layout(param->out_layout == "" ? param->data_layout : param->out_layout); diff --git a/tests/python/relay/test_autotvm_task_extraction.py b/tests/python/relay/test_autotvm_task_extraction.py index 83480a044f45..f83a23d38adb 100644 --- a/tests/python/relay/test_autotvm_task_extraction.py +++ b/tests/python/relay/test_autotvm_task_extraction.py @@ -16,8 +16,7 @@ # under the License. """Test task extraction for autotvm""" import tvm.relay.testing -from tvm import relay -from tvm import autotvm +from tvm import autotvm, relay def get_network(name, batch_size): diff --git a/tests/python/relay/test_op_level2.py b/tests/python/relay/test_op_level2.py index da2877063c45..db712be4262e 100644 --- a/tests/python/relay/test_op_level2.py +++ b/tests/python/relay/test_op_level2.py @@ -20,13 +20,12 @@ import numpy as np import pytest - import tvm import tvm.testing import tvm.topi.testing - from tvm import autotvm, relay, te from tvm.contrib import utils +from tvm.ir.module import IRModule from tvm.relay import transform from tvm.relay.testing import run_infer_type from tvm.topi.cuda.conv3d_winograd import _infer_tile_size @@ -838,25 +837,42 @@ def test_conv2d_transpose_infer_type(): @tvm.testing.uses_gpu def test_conv2d_transpose_nchw_run(): - dshape = (1, 3, 18, 18) - kshape = (3, 10, 3, 3) - oshape = (1, 10, 36, 36) - x = relay.var("x", shape=dshape) - w = relay.var("w") - y = relay.nn.conv2d_transpose( - x, w, channels=10, kernel_size=(3, 3), strides=(2, 2), padding=(1, 1), output_padding=(1, 1) - ) - func = relay.Function([x, w], y) - dtype = "float32" - data = np.random.uniform(size=dshape).astype(dtype) - kernel = np.random.uniform(size=kshape).astype(dtype) - ref_res = tvm.topi.testing.conv2d_transpose_nchw_python(data, kernel, 2, 1, (1, 1)) + k_layouts = {"OIHW": (10, 3, 3, 3), "IOHW": (3, 10, 3, 3)} - for target, dev in tvm.testing.enabled_targets(): - op_res1 = relay.create_executor("graph", device=dev, target=target).evaluate(func)( - data, kernel + for k_layout, kshape in k_layouts.items(): + dshape = (1, 3, 18, 18) + oshape = (1, 10, 36, 36) + x = relay.var("x", shape=dshape) + w = relay.var("w") + y = relay.nn.conv2d_transpose( + x, + w, + channels=10, + kernel_size=(3, 3), + strides=(2, 2), + padding=(1, 1), + output_padding=(1, 1), + kernel_layout=k_layout, + data_layout="NCHW", ) - tvm.testing.assert_allclose(op_res1.numpy(), ref_res, rtol=1e-5, atol=1e-5) + func = relay.Function([x, w], y) + dtype = "float32" + data = np.random.uniform(size=dshape).astype(dtype) + kernel = np.random.uniform(size=kshape).astype(dtype) + + if k_layout != "IOHW": + # Must be OIHW so switch + kernel_iohw = np.transpose(kernel, [1, 0, 2, 3]) + else: + kernel_iohw = kernel + + ref_res = tvm.topi.testing.conv2d_transpose_nchw_python(data, kernel_iohw, 2, 1, (1, 1)) + + for target, dev in tvm.testing.enabled_targets(): + op_res1 = relay.create_executor("graph", device=dev, target=target).evaluate(func)( + data, kernel + ) + tvm.testing.assert_allclose(op_res1.numpy(), ref_res, rtol=1e-5, atol=1e-5) @tvm.testing.uses_gpu @@ -866,8 +882,7 @@ def test_conv2d_transpose_nhwc_run(): oshape_nhwc = (1, 36, 36, 10) x = relay.var("x", shape=dshape_nhwc) w = relay.var("w") - # kshape and kernel_layout should have swapped IO. - # kshape is HWOI and kernel_layout is HWIO + y = relay.nn.conv2d_transpose( x, w, @@ -877,13 +892,12 @@ def test_conv2d_transpose_nhwc_run(): padding=(1, 1), output_padding=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", ) func = relay.Function([x, w], y) dtype = "float32" data = np.random.uniform(size=dshape_nhwc).astype(dtype) kernel = np.random.uniform(size=kshape_hwoi).astype(dtype) - # use true kshape layout here - HWOI ref_res = tvm.topi.testing.conv2d_transpose_nhwc_python( data, kernel, "HWOI", 2, 1, output_padding=(1, 1) diff --git a/tests/python/relay/test_pass_convert_op_layout.py b/tests/python/relay/test_pass_convert_op_layout.py index 9b4d154360b2..c6a5c53a1db4 100644 --- a/tests/python/relay/test_pass_convert_op_layout.py +++ b/tests/python/relay/test_pass_convert_op_layout.py @@ -16,15 +16,12 @@ # under the License. """Test alter op layout pass""" import pytest - import tvm -from tvm import te - -from tvm import relay +from tvm import relay, te +from tvm.relay import analysis, transform +from tvm.relay.op import op as reg from tvm.relay.op import register_alter_op_layout -from tvm.relay import transform, analysis from tvm.relay.transform.infer_layout_utils import InferCorrectLayoutOutput -from tvm.relay.op import op as reg def run_opt_pass(expr, passes): @@ -182,7 +179,7 @@ def expected(): x = relay.var("x", shape=(1, 56, 56, 64)) weight = relay.var("weight", shape=(3, 3, 64, 64)) x = relay.layout_transform(x, "NHWC", "NCHW") - weight = relay.layout_transform(weight, "HWIO", "OIHW") + weight = relay.layout_transform(weight, "HWIO", "IOHW") y = relay.nn.conv2d_transpose(x, weight, channels=64, kernel_size=(3, 3), padding=(1, 1)) y = relay.nn.relu(y) y = relay.layout_transform(y, "NCHW", "NHWC") @@ -190,7 +187,7 @@ def expected(): return y a = before() - a = run_opt_pass(a, transform.ConvertLayout({"nn.conv2d_transpose": ["NCHW", "OIHW"]})) + a = run_opt_pass(a, transform.ConvertLayout({"nn.conv2d_transpose": ["NCHW", "IOHW"]})) b = run_opt_pass(expected(), transform.InferType()) assert tvm.ir.structural_equal(a, b), "Actual = \n" + str(a) @@ -1134,7 +1131,7 @@ def expected(): x = relay.var("x", shape=(1, 56, 56, 64), dtype="int8") weight = relay.var("weight", shape=(3, 3, 64, 64), dtype="int8") x = relay.layout_transform(x, "NHWC", "NCHW") - weight = relay.layout_transform(weight, "HWIO", "OIHW") + weight = relay.layout_transform(weight, "HWIO", "IOHW") y = relay.qnn.op.conv2d_transpose( x, weight, @@ -1164,7 +1161,6 @@ def expected(): a = before() a = run_opt_pass(a, transform.ConvertLayout({"qnn.conv2d_transpose": ["NCHW", "default"]})) b = run_opt_pass(expected(), transform.InferType()) - assert tvm.ir.structural_equal(a, b), "Actual = \n" + str(a) diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index c49d837ed920..5f01189b7e33 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -17,13 +17,13 @@ # pylint: disable=unused-wildcard-import import numpy as np import pytest - import tvm from tvm import relay def compare_fq_to_int(expr, args, allow_rounding_error=False): mod = tvm.IRModule.from_expr(expr) + breakpoint() mod = tvm.relay.transform.InferType()(mod) mod_int = tvm.relay.transform.FakeQuantizationToInteger()(mod) @@ -89,6 +89,28 @@ def test_fake_quantize_conv_per_channel(): compare_fq_to_int(op, [x_np, w_np], allow_rounding_error=True) +def test_fake_quantize_transposeconv(): + for out_dtype in ["int8", "uint8"]: + x = relay.var("x", shape=[1, 3, 224, 224], dtype="int8") + w = relay.var("w", shape=[3, 16, 5, 5], dtype="int8") + one = relay.const(1.0) + zero = relay.const(0) + + op = relay.op.nn.conv2d_transpose( + relay.qnn.op.dequantize(x, relay.const(2.0), zero), + relay.qnn.op.dequantize(w, relay.const(0.5), zero), + kernel_size=[5, 5], + data_layout="NCHW", + kernel_layout="IOHW", + ) + op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) + + x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") + w_np = np.random.randint(-128, 127, size=[16, 3, 5, 5], dtype="int8") + + compare_fq_to_int(op, [x_np, w_np]) + + def test_fake_quantize_dense(): for out_dtype in ["int8", "uint8"]: x = relay.var("x", shape=[128, 64], dtype="int8") From 6b00883c380fc049de67bff35d1e0dcbcb11f5f4 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:18:24 -0700 Subject: [PATCH 02/35] make pytorch tests pass --- python/tvm/relay/frontend/pytorch.py | 11 ++++--- tests/python/frontend/pytorch/test_forward.py | 30 +------------------ 2 files changed, 8 insertions(+), 33 deletions(-) diff --git a/python/tvm/relay/frontend/pytorch.py b/python/tvm/relay/frontend/pytorch.py index 3fc202a7cc91..b3aecb589352 100644 --- a/python/tvm/relay/frontend/pytorch.py +++ b/python/tvm/relay/frontend/pytorch.py @@ -19,8 +19,8 @@ # pylint: disable=import-outside-toplevel, simplifiable-if-expression, cell-var-from-loop, unnecessary-lambda # pylint: disable=missing-function-docstring """PT: PyTorch frontend.""" -import itertools import functools +import itertools import logging import math import sys @@ -40,11 +40,11 @@ from ..prelude import Prelude, StaticTensorArrayOps from ..ty import Any, TensorType, TupleType from . import qnn_torch -from .common import AttrCvt, get_relay_op, unbind, lstm_cell, gru_cell -from .common import infer_value as _infer_value +from .common import AttrCvt, get_relay_op, gru_cell from .common import infer_shape as _infer_shape +from .common import infer_value as _infer_value from .common import infer_value_simulated as _infer_value_simulated -from .common import try_infer_value +from .common import lstm_cell, try_infer_value, unbind from .pytorch_utils import is_version_greater_than __all__ = ["from_pytorch"] @@ -1022,6 +1022,9 @@ def convolution(self, inputs, input_types): elif len(kernel_size) == 2: data_layout = "NCHW" kernel_layout = "OIHW" + if use_transpose: + # Transposed convolutions have IOHW layout. + kernel_layout = "IOHW" else: data_layout = "NCW" kernel_layout = "OIW" diff --git a/tests/python/frontend/pytorch/test_forward.py b/tests/python/frontend/pytorch/test_forward.py index 0031f4143fab..a05c66d2c974 100644 --- a/tests/python/frontend/pytorch/test_forward.py +++ b/tests/python/frontend/pytorch/test_forward.py @@ -21,6 +21,7 @@ from time import time import numpy as np +import pytest import torch import torchvision import tvm @@ -32,7 +33,6 @@ from tvm import relay from tvm.contrib import graph_executor from tvm.contrib.nvcc import have_fp16 -import pytest sys.setrecursionlimit(10000) @@ -994,20 +994,6 @@ def test_forward_conv_transpose( # opt to make the stride 1 + output padding stride = output_padding + 1 - # Conv 3D Transpose Tests - conv3d_input_shape = [1, in_channels, 16, 16, 16] - conv3d_input_data = torch.rand(conv3d_input_shape).float() - conv3d_transpose = torch.nn.ConvTranspose3d( - in_channels=in_channels, - out_channels=out_channels, - kernel_size=kernel_size, - stride=stride, - output_padding=output_padding, - groups=groups, - bias=bias, - ).eval() - verify_model(conv3d_transpose, conv3d_input_data) - # Conv 2D Transpose Tests conv2d_input_shape = [1, in_channels, 128, 256] conv2d_input_data = torch.rand(conv2d_input_shape).float() @@ -1022,20 +1008,6 @@ def test_forward_conv_transpose( ).eval() verify_model(conv2d_transpose, conv2d_input_data) - # # Conv 1D Transpose Tests - conv1d_input_shape = [1, in_channels, 10] - conv1d_input_data = torch.rand(conv1d_input_shape).float() - conv1d_transpose = torch.nn.ConvTranspose1d( - in_channels=in_channels, - out_channels=out_channels, - kernel_size=kernel_size, - stride=stride, - output_padding=output_padding, - groups=groups, - bias=bias, - ).eval() - verify_model(conv1d_transpose, conv1d_input_data) - def test_forward_deform_conv(): torch.set_grad_enabled(False) From 0447f108f1cd5791d4f765aafa4dc945df9252b1 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:42:48 -0700 Subject: [PATCH 03/35] lint --- python/tvm/relay/transform/fake_quantization_to_integer.py | 1 + python/tvm/topi/nn/conv2d_transpose.py | 3 ++- python/tvm/topi/nn/utils.py | 1 - 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index ecf5d9581bd8..febccde83dcf 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -159,6 +159,7 @@ def conv2d(expr, type_map): @register_fake_quantization_to_integer("nn.conv2d_transpose") def conv2d_transpose(expr, type_map): + """Rewrite a conv2d_transpose op""" attrs = {**expr.attrs} attrs.pop("out_dtype") x, weight = expr.args diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index 8291bfed4ea6..e668ed0b20d4 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -186,7 +186,8 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Convert back to original NHWC layout. out = relay.transpose(deconv, axes=(0, 2, 3, 1)) return out - elif attrs["data_layout"] == "NCHW": + + if attrs["data_layout"] == "NCHW": kernel = layout_transform(kernel, kernel_layout, "IOHW") new_attrs = {k: attrs[k] for k in attrs.keys()} diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index eea35c814c11..d1c219d8eca0 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,7 +19,6 @@ from __future__ import absolute_import import tvm -from tvm import relay from ..utils import get_const_int From 2425b41174c6825d07a45360f8219066ca189bd9 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 21:57:15 -0700 Subject: [PATCH 04/35] add test --- tests/python/relay/test_pass_fake_quantization_to_integer.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index 5f01189b7e33..8ac0f91bb066 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -23,7 +23,6 @@ def compare_fq_to_int(expr, args, allow_rounding_error=False): mod = tvm.IRModule.from_expr(expr) - breakpoint() mod = tvm.relay.transform.InferType()(mod) mod_int = tvm.relay.transform.FakeQuantizationToInteger()(mod) From a1c8f6b5cf80fe8c9e473ed1dce4a39ea54de08c Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 22:23:03 -0700 Subject: [PATCH 05/35] fix bug with layout transform --- python/tvm/topi/nn/conv2d_transpose.py | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index e668ed0b20d4..99c7442240c7 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -146,8 +146,8 @@ def layout_transform(tensor: "relay.Expr", current_layout: str, desired_layout: desired_layout_map = {c: i for i, c in enumerate(desired_layout)} axes = [None] * len(current_layout) - for c, i in current_layout_map.items(): - axes[i] = desired_layout_map[c] + for c, i in desired_layout_map.items(): + axes[i] = current_layout_map[c] return relay.transpose(tensor, axes=axes) @@ -169,6 +169,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): result : tvm.relay.Expr The legalized expr """ + data, kernel = inputs kernel_layout = attrs["kernel_layout"] if attrs["data_layout"] == "NHWC": @@ -177,7 +178,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Set new attrs for conv2d_transpose. new_attrs = {k: attrs[k] for k in attrs.keys()} new_attrs["data_layout"] = "NCHW" - # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + # layout of kernel should be IOHW, but kernel_layout will be swapped - OIHW new_attrs["kernel_layout"] = "IOHW" # Convert data to NCHW. @@ -191,7 +192,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): kernel = layout_transform(kernel, kernel_layout, "IOHW") new_attrs = {k: attrs[k] for k in attrs.keys()} - # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + # layout of kernel should be IOHW, but kernel_layout will be swapped - OIHW new_attrs["kernel_layout"] = "IOHW" return relay.nn.conv2d_transpose(data, kernel, **new_attrs) From 26e1462ee29cdf00e26d95c2ead45b6f928ba9b1 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 14:26:41 -0700 Subject: [PATCH 06/35] change layouts for conv2d_transpose too --- python/tvm/relay/op/contrib/vitis_ai.py | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index 7b9324cdc671..31a7bb555367 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -18,13 +18,13 @@ """Vitis-AI codegen annotation of supported operators""" import warnings -import numpy as np -from tvm import relay +import numpy as np import tvm._ffi +from tvm import relay from tvm.relay import transform -from tvm.relay.expr import Tuple, TupleGetItem from tvm.relay.build_module import bind_params_by_name +from tvm.relay.expr import Tuple, TupleGetItem from tvm.relay.op.annotation import compiler_begin, compiler_end # Placeholder for PyXIR module @@ -164,11 +164,13 @@ def partition_for_vitis_ai(mod, params=None, dpu=None, **opts): desired_layouts_in_partition = { "nn.conv2d": ["NHWC", "default"], + "nn.conv2d_transpose": ["NHWC", "default"], "nn.upsampling": ["NHWC"], "image.resize2d": ["NHWC"], } desired_layouts_in_main = { "nn.conv2d": ["NCHW", "default"], + "nn.conv2d_transpose": ["NCHW", "default"], "nn.upsampling": ["NCHW"], "image.resize2d": ["NCHW"], } From 8b93a0e46a4132759783373a129f8c3b1d985bfb Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Thu, 21 Oct 2021 15:03:50 -0700 Subject: [PATCH 07/35] fix vitis tests --- python/tvm/relay/op/contrib/vitis_ai.py | 2 -- .../contrib/test_vitis_ai/test_vitis_ai_codegen.py | 13 ++++++++++++- 2 files changed, 12 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index 31a7bb555367..fe6ee817bc36 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -164,13 +164,11 @@ def partition_for_vitis_ai(mod, params=None, dpu=None, **opts): desired_layouts_in_partition = { "nn.conv2d": ["NHWC", "default"], - "nn.conv2d_transpose": ["NHWC", "default"], "nn.upsampling": ["NHWC"], "image.resize2d": ["NHWC"], } desired_layouts_in_main = { "nn.conv2d": ["NCHW", "default"], - "nn.conv2d_transpose": ["NCHW", "default"], "nn.upsampling": ["NCHW"], "image.resize2d": ["NCHW"], } diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index c5c9cc7f818c..c50bed737d6f 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -241,6 +241,10 @@ def test_upsampling(dpu_target): verify_codegen(mod, dpu_target=dpu_target) +@pytest.mark.skip( + reason="I and O used to be mixed up in kernel layouts in TVM." + "This is fixed, but vitis needs to adopt the new convention" +) @pytest.mark.parametrize( "dpu_target", ["DPUCADF8H", "DPUCAHX8H-u50", "DPUCAHX8L", "DPUCVDX8H", "DPUCVDX8G", "DPUCZDX8G-zcu104"], @@ -253,7 +257,14 @@ def test_conv2d_transpose(dpu_target): x = relay.var("x", shape=dshape) w = relay.const(np.zeros(kshape, dtype="float32")) y = relay.nn.conv2d_transpose( - x, w, channels=10, kernel_size=(3, 3), strides=(1, 1), padding=(1, 1) + x, + w, + channels=10, + kernel_size=(3, 3), + strides=(1, 1), + padding=(1, 1), + data_layout="NCHW", + kernel_layout="IOHW", ) func = relay.Function([x], y) params = {} From aa3daf986e8db96f36bfd51e8c07e89ce6823943 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 15:29:22 -0700 Subject: [PATCH 08/35] fix qnn conv2d transpose tests --- .../relay/test_op_qnn_conv2_transpose.py | 53 +++++++++---------- 1 file changed, 26 insertions(+), 27 deletions(-) diff --git a/tests/python/relay/test_op_qnn_conv2_transpose.py b/tests/python/relay/test_op_qnn_conv2_transpose.py index 9fd3d1b84537..a7c3c9110f6d 100644 --- a/tests/python/relay/test_op_qnn_conv2_transpose.py +++ b/tests/python/relay/test_op_qnn_conv2_transpose.py @@ -15,13 +15,12 @@ # specific language governing permissions and limitations # under the License. -import tvm -from tvm import te import numpy as np -from tvm import relay +import tvm +from tvm import relay, te +from tvm.contrib import graph_executor from tvm.relay import transform from tvm.relay.testing import run_infer_type -from tvm.contrib import graph_executor from tvm.relay.testing.temp_op_attr import TempOpAttr @@ -224,7 +223,7 @@ def test_no_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -248,7 +247,7 @@ def test_no_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -274,7 +273,7 @@ def test_kernel_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -298,7 +297,7 @@ def test_kernel_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -324,7 +323,7 @@ def test_input_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -348,7 +347,7 @@ def test_input_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -374,7 +373,7 @@ def test_both_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -398,7 +397,7 @@ def test_both_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -424,7 +423,7 @@ def test_different_dtype(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], ) @@ -449,7 +448,7 @@ def test_different_dtype(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], ) @@ -460,7 +459,7 @@ def test_layout(): # uint8 input data_shape = (2, 2, 4, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -476,14 +475,14 @@ def test_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) data_shape = (2, 2, 4, 3) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 1, 3) # HWIO + kernel_shape = (2, 2, 1, 3) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -499,7 +498,7 @@ def test_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -525,7 +524,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -533,7 +532,7 @@ def test_padding(): # Try different layout data_shape = (2, 2, 4, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -549,7 +548,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -557,7 +556,7 @@ def test_padding(): # Try asymmetric padding data_shape = (2, 8, 6, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -573,7 +572,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -600,7 +599,7 @@ def test_const_folding(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], groups=1, @@ -614,7 +613,7 @@ def test_broadcast_layout(): # Test broadcast support for NHWC layout. data_shape = (1, 229, 229, 3) # NHWC data_dtype = "uint8" - kernel_shape = (7, 7, 64, 3) # HWIO + kernel_shape = (7, 7, 64, 3) # HWOI kernel_dtype = "int8" _, qnn_func = get_funcs( data_shape=data_shape, @@ -630,7 +629,7 @@ def test_broadcast_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) func = qnn_func["main"].body @@ -670,7 +669,7 @@ def test_per_channel_kernel_scale(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) From 5bb9e808e5b1756d4d6eb8207d31de89c02d64c1 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:04:10 -0700 Subject: [PATCH 09/35] fix fake quantization pass --- python/tvm/relay/qnn/op/qnn.py | 2 +- python/tvm/relay/transform/fake_quantization_to_integer.py | 4 +++- tests/python/relay/test_pass_fake_quantization_to_integer.py | 2 +- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/qnn/op/qnn.py b/python/tvm/relay/qnn/op/qnn.py index d1121f7b90fc..4749b63b858f 100644 --- a/python/tvm/relay/qnn/op/qnn.py +++ b/python/tvm/relay/qnn/op/qnn.py @@ -387,7 +387,7 @@ def conv2d_transpose( kernel_layout="IOHW", out_layout="", output_padding=(0, 0), - out_dtype="", + out_dtype="int32", ): """This operator deconvolves quantized data with quantized kernel. The scale of the output quantized tensor is the product of the kernel_scale and diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index febccde83dcf..b8c63f4aa5ef 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -18,6 +18,7 @@ import tvm from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType +from tvm.ir.module import IRModule from tvm.tir import bijective_layout from ..op import register_fake_quantization_to_integer @@ -167,8 +168,9 @@ def conv2d_transpose(expr, type_map): w_t = type_map[weight] conv_scale = fold_constant(x_t.scale * w_t.scale) conv_zp = get_zeros(conv_scale) + out = relay.qnn.op.conv2d_transpose( - x, weight, x_t.zero_point, w_t.zero_point, x_t.zero_point, w_t.zero_point, **attrs + x, weight, x_t.zero_point, w_t.zero_point, x_t.scale, w_t.scale, **attrs ) out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index 8ac0f91bb066..f99c15639dd5 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -105,7 +105,7 @@ def test_fake_quantize_transposeconv(): op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") - w_np = np.random.randint(-128, 127, size=[16, 3, 5, 5], dtype="int8") + w_np = np.random.randint(-128, 127, size=[3, 16, 5, 5], dtype="int8") compare_fq_to_int(op, [x_np, w_np]) From 8c526e28553519f20dd61c00e65f8496ebcf7114 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:08:06 -0700 Subject: [PATCH 10/35] add todo --- .../contrib/test_vitis_ai/test_vitis_ai_codegen.py | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index c50bed737d6f..b89cc37d29b8 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -19,24 +19,23 @@ """Vitis-AI codegen tests""" import sys -import numpy as np +import numpy as np import pytest pytest.importorskip("pyxir") import pyxir.contrib.target.DPUCADF8H import pyxir.contrib.target.DPUCAHX8H import pyxir.contrib.target.DPUCAHX8L -import pyxir.contrib.target.DPUCVDX8H import pyxir.contrib.target.DPUCVDX8G +import pyxir.contrib.target.DPUCVDX8H import pyxir.contrib.target.DPUCZDX8G - import tvm from tvm import relay +from tvm.contrib.target import vitis_ai from tvm.relay import transform -from tvm.relay.op.contrib.vitis_ai import annotation from tvm.relay.build_module import bind_params_by_name -from tvm.contrib.target import vitis_ai +from tvm.relay.op.contrib.vitis_ai import annotation from .infrastructure import skip_test, verify_codegen @@ -243,7 +242,10 @@ def test_upsampling(dpu_target): @pytest.mark.skip( reason="I and O used to be mixed up in kernel layouts in TVM." - "This is fixed, but vitis needs to adopt the new convention" + "This is fixed, but vitis needs to adopt the new convention." + "To change, simply remove this line:" + "https://github.com/Xilinx/pyxir/blob/bef661d6d77adcdbd2cf4163f2cf3a1d31d40406/" + "python/pyxir/frontend/tvm/relay_tools/relay_l2_convolution.py#L380" ) @pytest.mark.parametrize( "dpu_target", From 03bad0848546a353de5352c639880e5f1741cb3d Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:21:32 -0700 Subject: [PATCH 11/35] lint --- python/tvm/relay/transform/fake_quantization_to_integer.py | 1 - 1 file changed, 1 deletion(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index b8c63f4aa5ef..2bb9cf463a64 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -18,7 +18,6 @@ import tvm from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType -from tvm.ir.module import IRModule from tvm.tir import bijective_layout from ..op import register_fake_quantization_to_integer From 962766c21c6a33629d98f279b45827eb850dd451 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:36:40 -0700 Subject: [PATCH 12/35] undo just formatting changes --- python/tvm/relay/op/contrib/vitis_ai.py | 6 ++-- python/tvm/topi/nn/utils.py | 1 - tests/python/frontend/pytorch/test_forward.py | 30 ++++++++++++++++++- 3 files changed, 32 insertions(+), 5 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index fe6ee817bc36..7b9324cdc671 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -18,13 +18,13 @@ """Vitis-AI codegen annotation of supported operators""" import warnings - import numpy as np -import tvm._ffi + from tvm import relay +import tvm._ffi from tvm.relay import transform -from tvm.relay.build_module import bind_params_by_name from tvm.relay.expr import Tuple, TupleGetItem +from tvm.relay.build_module import bind_params_by_name from tvm.relay.op.annotation import compiler_begin, compiler_end # Placeholder for PyXIR module diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index d1c219d8eca0..ff00441e9850 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,7 +19,6 @@ from __future__ import absolute_import import tvm - from ..utils import get_const_int diff --git a/tests/python/frontend/pytorch/test_forward.py b/tests/python/frontend/pytorch/test_forward.py index a05c66d2c974..0031f4143fab 100644 --- a/tests/python/frontend/pytorch/test_forward.py +++ b/tests/python/frontend/pytorch/test_forward.py @@ -21,7 +21,6 @@ from time import time import numpy as np -import pytest import torch import torchvision import tvm @@ -33,6 +32,7 @@ from tvm import relay from tvm.contrib import graph_executor from tvm.contrib.nvcc import have_fp16 +import pytest sys.setrecursionlimit(10000) @@ -994,6 +994,20 @@ def test_forward_conv_transpose( # opt to make the stride 1 + output padding stride = output_padding + 1 + # Conv 3D Transpose Tests + conv3d_input_shape = [1, in_channels, 16, 16, 16] + conv3d_input_data = torch.rand(conv3d_input_shape).float() + conv3d_transpose = torch.nn.ConvTranspose3d( + in_channels=in_channels, + out_channels=out_channels, + kernel_size=kernel_size, + stride=stride, + output_padding=output_padding, + groups=groups, + bias=bias, + ).eval() + verify_model(conv3d_transpose, conv3d_input_data) + # Conv 2D Transpose Tests conv2d_input_shape = [1, in_channels, 128, 256] conv2d_input_data = torch.rand(conv2d_input_shape).float() @@ -1008,6 +1022,20 @@ def test_forward_conv_transpose( ).eval() verify_model(conv2d_transpose, conv2d_input_data) + # # Conv 1D Transpose Tests + conv1d_input_shape = [1, in_channels, 10] + conv1d_input_data = torch.rand(conv1d_input_shape).float() + conv1d_transpose = torch.nn.ConvTranspose1d( + in_channels=in_channels, + out_channels=out_channels, + kernel_size=kernel_size, + stride=stride, + output_padding=output_padding, + groups=groups, + bias=bias, + ).eval() + verify_model(conv1d_transpose, conv1d_input_data) + def test_forward_deform_conv(): torch.set_grad_enabled(False) From 83d4fa6f9d0ba777606fc90beb455b5b213419e9 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:41:30 -0700 Subject: [PATCH 13/35] remove formatting only change --- tests/python/relay/test_autotvm_task_extraction.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/python/relay/test_autotvm_task_extraction.py b/tests/python/relay/test_autotvm_task_extraction.py index f83a23d38adb..83480a044f45 100644 --- a/tests/python/relay/test_autotvm_task_extraction.py +++ b/tests/python/relay/test_autotvm_task_extraction.py @@ -16,7 +16,8 @@ # under the License. """Test task extraction for autotvm""" import tvm.relay.testing -from tvm import autotvm, relay +from tvm import relay +from tvm import autotvm def get_network(name, batch_size): From b2f049d956b90e6f73f8366b283a4de45b992bc7 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:43:38 -0700 Subject: [PATCH 14/35] remove f2qi for later pr --- .../transform/fake_quantization_to_integer.py | 20 ---------------- .../test_pass_fake_quantization_to_integer.py | 23 +------------------ 2 files changed, 1 insertion(+), 42 deletions(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index 2bb9cf463a64..1adde9a4a430 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -19,7 +19,6 @@ from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType from tvm.tir import bijective_layout - from ..op import register_fake_quantization_to_integer @@ -157,25 +156,6 @@ def conv2d(expr, type_map): return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] -@register_fake_quantization_to_integer("nn.conv2d_transpose") -def conv2d_transpose(expr, type_map): - """Rewrite a conv2d_transpose op""" - attrs = {**expr.attrs} - attrs.pop("out_dtype") - x, weight = expr.args - x_t = type_map[x] - w_t = type_map[weight] - conv_scale = fold_constant(x_t.scale * w_t.scale) - conv_zp = get_zeros(conv_scale) - - out = relay.qnn.op.conv2d_transpose( - x, weight, x_t.zero_point, w_t.zero_point, x_t.scale, w_t.scale, **attrs - ) - out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] - out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] - return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] - - @register_fake_quantization_to_integer("nn.dense") def dense(expr, type_map): """Rewrite a dense op""" diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index f99c15639dd5..c49d837ed920 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -17,6 +17,7 @@ # pylint: disable=unused-wildcard-import import numpy as np import pytest + import tvm from tvm import relay @@ -88,28 +89,6 @@ def test_fake_quantize_conv_per_channel(): compare_fq_to_int(op, [x_np, w_np], allow_rounding_error=True) -def test_fake_quantize_transposeconv(): - for out_dtype in ["int8", "uint8"]: - x = relay.var("x", shape=[1, 3, 224, 224], dtype="int8") - w = relay.var("w", shape=[3, 16, 5, 5], dtype="int8") - one = relay.const(1.0) - zero = relay.const(0) - - op = relay.op.nn.conv2d_transpose( - relay.qnn.op.dequantize(x, relay.const(2.0), zero), - relay.qnn.op.dequantize(w, relay.const(0.5), zero), - kernel_size=[5, 5], - data_layout="NCHW", - kernel_layout="IOHW", - ) - op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) - - x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") - w_np = np.random.randint(-128, 127, size=[3, 16, 5, 5], dtype="int8") - - compare_fq_to_int(op, [x_np, w_np]) - - def test_fake_quantize_dense(): for out_dtype in ["int8", "uint8"]: x = relay.var("x", shape=[128, 64], dtype="int8") From 525ab244408786033e0def5b55d20823d66f4e0a Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 21:32:08 -0700 Subject: [PATCH 15/35] more frontend tests fixes --- python/tvm/relay/frontend/mxnet.py | 46 ++++++++++++++++---------- python/tvm/relay/frontend/qnn_torch.py | 8 ++--- python/tvm/relay/frontend/tflite.py | 39 +++++++++++----------- 3 files changed, 50 insertions(+), 43 deletions(-) diff --git a/python/tvm/relay/frontend/mxnet.py b/python/tvm/relay/frontend/mxnet.py index 59b4e99de999..1b1d60119967 100644 --- a/python/tvm/relay/frontend/mxnet.py +++ b/python/tvm/relay/frontend/mxnet.py @@ -18,40 +18,50 @@ """MXNet symbol frontend.""" import json import math + import numpy as np import tvm -from tvm.ir import IRModule - from tvm import relay +from tvm.ir import IRModule from tvm.topi.utils import get_const_tuple + +from ... import nd as _nd from .. import analysis from .. import expr as _expr from .. import function as _function from .. import op as _op from .. import scope_builder as _scope_builder -from ... import nd as _nd - from .common import StrAttrsDict -from .common import infer_type as _infer_type +from .common import get_name as _get_name from .common import infer_shape as _infer_shape +from .common import infer_type as _infer_type from .common import infer_value as _infer_value -from .common import get_name as _get_name -from .nnvm_common import _rename, _binop_scalar, _rbinop_scalar, _reduce -from .nnvm_common import _arg_reduce, _init_op, _softmax_op, _cast -from .nnvm_common import _clip, _transpose, _upsampling -from .nnvm_common import _elemwise_sum, _reshape -from .nnvm_common import _warn_not_used from .mxnet_qnn_op_utils import ( - quantize_mxnet_min_max, - quantize_conv_weights_bias_channel_mkldnn_from_var, - quantize_conv_bias_mkldnn_from_var, - get_conv_mkldnn_requantized_scale_outDtype, dequantize_mxnet_min_max, + get_conv_mkldnn_requantized_scale_outDtype, get_mkldnn_int8_scale, - get_mkldnn_uint8_scale, get_mkldnn_requantize_scale_outDtype, + get_mkldnn_uint8_scale, + quantize_conv_bias_mkldnn_from_var, + quantize_conv_weights_bias_channel_mkldnn_from_var, + quantize_mxnet_min_max, +) +from .nnvm_common import ( + _arg_reduce, + _binop_scalar, + _cast, + _clip, + _elemwise_sum, + _init_op, + _rbinop_scalar, + _reduce, + _rename, + _reshape, + _softmax_op, + _transpose, + _upsampling, + _warn_not_used, ) - __all__ = ["from_mxnet"] @@ -329,7 +339,7 @@ def _mx_conv2d_transpose(inputs, attrs): if "kernel_layout" in attrs.attrs: kernel_layout = attrs.get_str("kernel_layout") else: - kernel_layout = "HWIO" if data_layout == "NHWC" else "OIHW" + kernel_layout = "HWIO" if data_layout == "NHWC" else "IOHW" new_attrs = {} new_attrs["channels"] = attrs.get_int("num_filter") diff --git a/python/tvm/relay/frontend/qnn_torch.py b/python/tvm/relay/frontend/qnn_torch.py index 172ab1e41268..5772313fb97f 100644 --- a/python/tvm/relay/frontend/qnn_torch.py +++ b/python/tvm/relay/frontend/qnn_torch.py @@ -19,7 +19,6 @@ import logging import numpy as np - import tvm from tvm import relay from tvm.relay import expr as _expr @@ -1043,11 +1042,8 @@ def _impl(inputs, _): weight_shape = list(infer_shape(weight)) - # Swap I and O dims to match shape relay expects for OIHW - weight_shape[0], weight_shape[1] = weight_shape[1], weight_shape[0] - kernel_size = (weight_shape[2], weight_shape[3]) - out_channels = weight_shape[0] + out_channels = weight_shape[1] conv_out = relay.qnn.op.conv2d_transpose( inputs[0], @@ -1064,7 +1060,7 @@ def _impl(inputs, _): channels=out_channels, output_padding=output_padding, out_dtype="int32", - kernel_layout="OIHW", + kernel_layout="IOHW", ) return _do_bias_and_requantize( diff --git a/python/tvm/relay/frontend/tflite.py b/python/tvm/relay/frontend/tflite.py index 3688ff5ff4e5..5d3681ebe132 100644 --- a/python/tvm/relay/frontend/tflite.py +++ b/python/tvm/relay/frontend/tflite.py @@ -16,24 +16,25 @@ # under the License. # pylint: disable=invalid-name, unused-argument, too-many-lines, import-outside-toplevel """Tensorflow lite frontend.""" -import math import itertools +import math + import numpy as np import tvm +from tvm import relay from tvm.ir import IRModule -from tvm import relay +from ... import nd as _nd from .. import analysis from .. import expr as _expr from .. import function as _function from .. import op as _op from .. import qnn as _qnn -from ... import nd as _nd from .common import ExprTable -from .common import infer_shape as _infer_shape, to_int_list +from .common import infer_shape as _infer_shape +from .common import to_int_list from .tflite_flexbuffer import FlexBufferDecoder - __all__ = ["from_tflite"] @@ -53,9 +54,9 @@ class OperatorConverter(object): def __init__(self, model, subgraph, exp_tab): try: + from tflite.ActivationFunctionType import ActivationFunctionType from tflite.BuiltinOperator import BuiltinOperator from tflite.BuiltinOptions import BuiltinOptions - from tflite.ActivationFunctionType import ActivationFunctionType except ImportError: raise ImportError("The tflite package must be installed") @@ -1061,8 +1062,8 @@ def convert_log_softmax(self, op): def convert_concatenation(self, op): """Convert TFLite concatenation""" try: - from tflite.ConcatenationOptions import ConcatenationOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.ConcatenationOptions import ConcatenationOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1248,10 +1249,10 @@ def _convert_elemwise(self, relay_op, op, ignore_qnn_params=False): """Generic method to Convert TFLite elemwise""" try: from tflite.AddOptions import AddOptions - from tflite.SubOptions import SubOptions - from tflite.MulOptions import MulOptions - from tflite.DivOptions import DivOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.DivOptions import DivOptions + from tflite.MulOptions import MulOptions + from tflite.SubOptions import SubOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1810,9 +1811,9 @@ def convert_reduce_any(self, op): def _convert_arg_min_max(self, relay_op, op): """Generic method converting TFLite arg_min_max""" try: - from tflite.BuiltinOptions import BuiltinOptions - from tflite.ArgMinOptions import ArgMinOptions from tflite.ArgMaxOptions import ArgMaxOptions + from tflite.ArgMinOptions import ArgMinOptions + from tflite.BuiltinOptions import BuiltinOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1859,8 +1860,8 @@ def convert_arg_max(self, op): def convert_fully_connected(self, op): """Convert TFLite fully connected""" try: - from tflite.FullyConnectedOptions import FullyConnectedOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.FullyConnectedOptions import FullyConnectedOptions from tflite.TensorType import TensorType except ImportError: raise ImportError("The tflite package must be installed") @@ -2030,10 +2031,10 @@ def convert_conv(self, op, conv_type): """convolution implementation.""" try: from tflite.BuiltinOptions import BuiltinOptions - from tflite.TensorType import TensorType from tflite.Conv2DOptions import Conv2DOptions from tflite.DepthwiseConv2DOptions import DepthwiseConv2DOptions from tflite.Padding import Padding + from tflite.TensorType import TensorType except ImportError: raise ImportError("The tflite package must be installed") @@ -2440,8 +2441,8 @@ def convert_pool2d(self, op, pool_type): """pool2d implementation.""" try: from tflite.BuiltinOptions import BuiltinOptions - from tflite.Pool2DOptions import Pool2DOptions from tflite.Padding import Padding + from tflite.Pool2DOptions import Pool2DOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -2856,9 +2857,9 @@ def convert_transpose_conv(self, op): """Convert TFLite TRANSPOSE_CONV""" try: from tflite.BuiltinOptions import BuiltinOptions + from tflite.Padding import Padding from tflite.TensorType import TensorType from tflite.TransposeConvOptions import TransposeConvOptions - from tflite.Padding import Padding except ImportError: raise ImportError("The tflite package must be installed") @@ -2952,7 +2953,7 @@ def convert_transpose_conv(self, op): channels=int(out_channels), kernel_size=(int(kernel_h), int(kernel_w)), data_layout="NHWC", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) else: @@ -2964,7 +2965,7 @@ def convert_transpose_conv(self, op): channels=int(out_channels), kernel_size=(int(kernel_h), int(kernel_w)), data_layout="NHWC", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype=output_tensor_type_str, ) @@ -3723,8 +3724,8 @@ def from_tflite(model, shape_dict=None, dtype_dict=None, op_converter=OperatorCo The parameter dict to be used by relay """ try: - import tflite.SubGraph import tflite.BuiltinOperator + import tflite.SubGraph except ImportError: raise ImportError("The tflite package must be installed") From 306ac98cb37b2087c692f97055a3fd640d8f84dd Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:11:32 -0700 Subject: [PATCH 16/35] fix a lot of initial tests --- python/tvm/relay/op/nn/nn.py | 2 +- python/tvm/relay/qnn/op/layout_conversions.py | 2 +- python/tvm/relay/qnn/op/qnn.py | 6 +- python/tvm/relay/testing/dcgan.py | 3 +- .../transform/fake_quantization_to_integer.py | 18 +++++ python/tvm/topi/nn/conv2d_transpose.py | 76 ++++++++++++------- python/tvm/topi/nn/utils.py | 2 + src/relay/op/nn/convolution.h | 6 +- .../relay/test_autotvm_task_extraction.py | 3 +- tests/python/relay/test_op_level2.py | 60 +++++++++------ .../relay/test_pass_convert_op_layout.py | 16 ++-- .../test_pass_fake_quantization_to_integer.py | 24 +++++- 12 files changed, 147 insertions(+), 71 deletions(-) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 1821ff17258a..c7b376ec3d64 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -522,7 +522,7 @@ def conv2d_transpose( channels=None, kernel_size=None, data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_layout="", output_padding=(0, 0), out_dtype="", diff --git a/python/tvm/relay/qnn/op/layout_conversions.py b/python/tvm/relay/qnn/op/layout_conversions.py index 1a3b1771d6ce..24c787e0a00a 100644 --- a/python/tvm/relay/qnn/op/layout_conversions.py +++ b/python/tvm/relay/qnn/op/layout_conversions.py @@ -119,7 +119,7 @@ def convert_qnn_conv2d_transpose(attrs, inputs, tinfos, desired_layouts): # Handle default kernel layouts if desired_data_layout == "NCHW": - new_attrs["kernel_layout"] = "OIHW" + new_attrs["kernel_layout"] = "IOHW" return relay.qnn.op.conv2d_transpose(*inputs, **new_attrs) if desired_data_layout == "NHWC": new_attrs["kernel_layout"] = "HWIO" diff --git a/python/tvm/relay/qnn/op/qnn.py b/python/tvm/relay/qnn/op/qnn.py index 83b5cf0a831c..d1121f7b90fc 100644 --- a/python/tvm/relay/qnn/op/qnn.py +++ b/python/tvm/relay/qnn/op/qnn.py @@ -18,13 +18,15 @@ """QNN dialect operators.""" from __future__ import absolute_import as _abs + from tvm import relay from tvm.relay.expr import Tuple, TupleWrapper from tvm.relay.op.nn.utils import get_pad_tuple2d from tvm.topi.nn.qnn import SQNN_DTYPE_TO_CODE -from . import _make + from ... import op as reg from ...op import OpPattern +from . import _make def requantize( @@ -382,7 +384,7 @@ def conv2d_transpose( channels=None, kernel_size=None, data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_layout="", output_padding=(0, 0), out_dtype="", diff --git a/python/tvm/relay/testing/dcgan.py b/python/tvm/relay/testing/dcgan.py index fc531b765a88..acc478330dc9 100644 --- a/python/tvm/relay/testing/dcgan.py +++ b/python/tvm/relay/testing/dcgan.py @@ -27,6 +27,7 @@ arXiv preprint arXiv:1511.06434 (2015). """ from tvm import relay + from . import layers from .init import create_workload @@ -41,7 +42,7 @@ def deconv2d(data, ishape, oshape, kshape, layout, name, stride=(2, 2)): adj_x = (target_shape[1] + 2 * pad_x - kshape[1]) % stride[1] if layout == "NCHW": - kernel_layout = "OIHW" + kernel_layout = "IOHW" elif layout == "NHWC": kernel_layout = "HWOI" else: diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index 1adde9a4a430..ecf5d9581bd8 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -19,6 +19,7 @@ from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType from tvm.tir import bijective_layout + from ..op import register_fake_quantization_to_integer @@ -156,6 +157,23 @@ def conv2d(expr, type_map): return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] +@register_fake_quantization_to_integer("nn.conv2d_transpose") +def conv2d_transpose(expr, type_map): + attrs = {**expr.attrs} + attrs.pop("out_dtype") + x, weight = expr.args + x_t = type_map[x] + w_t = type_map[weight] + conv_scale = fold_constant(x_t.scale * w_t.scale) + conv_zp = get_zeros(conv_scale) + out = relay.qnn.op.conv2d_transpose( + x, weight, x_t.zero_point, w_t.zero_point, x_t.zero_point, w_t.zero_point, **attrs + ) + out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] + out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] + return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] + + @register_fake_quantization_to_integer("nn.dense") def dense(expr, type_map): """Rewrite a dense op""" diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index 22188bcd45a4..8291bfed4ea6 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -17,12 +17,12 @@ # pylint: disable=invalid-name, unused-variable, unused-argument """Transposed 2D convolution operators (sometimes called Deconvolution).""" import tvm -from tvm import te -from tvm import relay +from tvm import relay, te + +from ..utils import simplify from .dilate import dilate from .pad import pad from .utils import get_pad_tuple -from ..utils import simplify def conv2d_transpose_nchw(Input, Filter, strides, padding, out_dtype, output_padding): @@ -116,6 +116,41 @@ def declaration_conv2d_transpose_impl(data, kernel, strides, padding, out_dtype, return Output +def layout_transform(tensor: "relay.Expr", current_layout: str, desired_layout: str): + """Transform a tensor with the current layout to the desired layout. + + E.g. layout_transform(t, "NCHW", "CNHW") --> relay.transpose(t, [1, 0, 2, 3]) + + Parameters + ---------- + tensor: relay.Expr + The Tensor to transpose + + current_layout: str + The current layout e.g. NCHW or OIHW + + desired_layout: str + The desired layout, must be compatible with current_layout + + Returns + ------- + The layout_transformed tensor. + """ + if sorted(current_layout) != sorted(desired_layout): + raise ValueError(f"Incompatible layouts: {current_layout} vs {desired_layout}") + + if current_layout == desired_layout: + return tensor + + current_layout_map = {c: i for i, c in enumerate(current_layout)} + desired_layout_map = {c: i for i, c in enumerate(desired_layout)} + + axes = [None] * len(current_layout) + for c, i in current_layout_map.items(): + axes[i] = desired_layout_map[c] + return relay.transpose(tensor, axes=axes) + + @tvm.target.generic_func def conv2d_transpose_legalize(attrs, inputs, types): """Legalizes Transposed 2D convolution op. @@ -134,36 +169,16 @@ def conv2d_transpose_legalize(attrs, inputs, types): result : tvm.relay.Expr The legalized expr """ + data, kernel = inputs + kernel_layout = attrs["kernel_layout"] if attrs["data_layout"] == "NHWC": - data, kernel = inputs - kernel_layout = attrs["kernel_layout"] - # Convert Kernel layout to IOHW - # kernel_layout is different from input kernel layout - IO is swapped - if kernel_layout == "HWIO": - # input kernel layout is swapped to HWOI - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(3, 2, 0, 1)) - elif kernel_layout == "HWOI": - # input kernel layout is swapped to HWIO - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(2, 3, 0, 1)) - elif kernel_layout == "IOHW": - # input kernel layout is swapped to OIHW - # output kernel layout will be IOHW - kernel = relay.transpose(kernel, axes=(1, 0, 2, 3)) - elif kernel_layout == "OIHW": - # input kernel layout is swapped to IOHW - # output kernel layout will be IOHW - pass - else: - # Skip legalize. Let relay.nn.conv2d_transpose to handle the case - return None + kernel = layout_transform(kernel, kernel_layout, "IOHW") # Set new attrs for conv2d_transpose. new_attrs = {k: attrs[k] for k in attrs.keys()} new_attrs["data_layout"] = "NCHW" # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW - new_attrs["kernel_layout"] = "OIHW" + new_attrs["kernel_layout"] = "IOHW" # Convert data to NCHW. data = relay.transpose(data, axes=(0, 3, 1, 2)) @@ -171,5 +186,12 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Convert back to original NHWC layout. out = relay.transpose(deconv, axes=(0, 2, 3, 1)) return out + elif attrs["data_layout"] == "NCHW": + kernel = layout_transform(kernel, kernel_layout, "IOHW") + new_attrs = {k: attrs[k] for k in attrs.keys()} + + # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + new_attrs["kernel_layout"] = "IOHW" + return relay.nn.conv2d_transpose(data, kernel, **new_attrs) return None diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index ff00441e9850..eea35c814c11 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,6 +19,8 @@ from __future__ import absolute_import import tvm +from tvm import relay + from ..utils import get_const_int diff --git a/src/relay/op/nn/convolution.h b/src/relay/op/nn/convolution.h index c27227b2eb73..d9958076adc1 100644 --- a/src/relay/op/nn/convolution.h +++ b/src/relay/op/nn/convolution.h @@ -1044,7 +1044,7 @@ bool Conv2DTransposeRel(const Array& types, int num_inputs, const Attrs& a if (data == nullptr) return false; static const Layout kNCHW("NCHW"); - static const Layout kOIHW("OIHW"); + static const Layout kIOHW("IOHW"); const Conv2DTransposeAttrs* param = attrs.as(); ICHECK(param != nullptr); @@ -1056,9 +1056,9 @@ bool Conv2DTransposeRel(const Array& types, int num_inputs, const Attrs& a << "Conv only support input layouts that are convertible from NCHW." << " But got " << in_layout; - const auto trans_kernel_layout = tir::BijectiveLayout(kernel_layout, kOIHW); + const auto trans_kernel_layout = tir::BijectiveLayout(kernel_layout, kIOHW); ICHECK(trans_kernel_layout.defined()) - << "Conv only support kernel layouts that are convertible from OIHW." + << "Conv only support kernel layouts that are convertible from IOHW." << " But got " << kernel_layout; Layout out_layout(param->out_layout == "" ? param->data_layout : param->out_layout); diff --git a/tests/python/relay/test_autotvm_task_extraction.py b/tests/python/relay/test_autotvm_task_extraction.py index 83480a044f45..f83a23d38adb 100644 --- a/tests/python/relay/test_autotvm_task_extraction.py +++ b/tests/python/relay/test_autotvm_task_extraction.py @@ -16,8 +16,7 @@ # under the License. """Test task extraction for autotvm""" import tvm.relay.testing -from tvm import relay -from tvm import autotvm +from tvm import autotvm, relay def get_network(name, batch_size): diff --git a/tests/python/relay/test_op_level2.py b/tests/python/relay/test_op_level2.py index da2877063c45..db712be4262e 100644 --- a/tests/python/relay/test_op_level2.py +++ b/tests/python/relay/test_op_level2.py @@ -20,13 +20,12 @@ import numpy as np import pytest - import tvm import tvm.testing import tvm.topi.testing - from tvm import autotvm, relay, te from tvm.contrib import utils +from tvm.ir.module import IRModule from tvm.relay import transform from tvm.relay.testing import run_infer_type from tvm.topi.cuda.conv3d_winograd import _infer_tile_size @@ -838,25 +837,42 @@ def test_conv2d_transpose_infer_type(): @tvm.testing.uses_gpu def test_conv2d_transpose_nchw_run(): - dshape = (1, 3, 18, 18) - kshape = (3, 10, 3, 3) - oshape = (1, 10, 36, 36) - x = relay.var("x", shape=dshape) - w = relay.var("w") - y = relay.nn.conv2d_transpose( - x, w, channels=10, kernel_size=(3, 3), strides=(2, 2), padding=(1, 1), output_padding=(1, 1) - ) - func = relay.Function([x, w], y) - dtype = "float32" - data = np.random.uniform(size=dshape).astype(dtype) - kernel = np.random.uniform(size=kshape).astype(dtype) - ref_res = tvm.topi.testing.conv2d_transpose_nchw_python(data, kernel, 2, 1, (1, 1)) + k_layouts = {"OIHW": (10, 3, 3, 3), "IOHW": (3, 10, 3, 3)} - for target, dev in tvm.testing.enabled_targets(): - op_res1 = relay.create_executor("graph", device=dev, target=target).evaluate(func)( - data, kernel + for k_layout, kshape in k_layouts.items(): + dshape = (1, 3, 18, 18) + oshape = (1, 10, 36, 36) + x = relay.var("x", shape=dshape) + w = relay.var("w") + y = relay.nn.conv2d_transpose( + x, + w, + channels=10, + kernel_size=(3, 3), + strides=(2, 2), + padding=(1, 1), + output_padding=(1, 1), + kernel_layout=k_layout, + data_layout="NCHW", ) - tvm.testing.assert_allclose(op_res1.numpy(), ref_res, rtol=1e-5, atol=1e-5) + func = relay.Function([x, w], y) + dtype = "float32" + data = np.random.uniform(size=dshape).astype(dtype) + kernel = np.random.uniform(size=kshape).astype(dtype) + + if k_layout != "IOHW": + # Must be OIHW so switch + kernel_iohw = np.transpose(kernel, [1, 0, 2, 3]) + else: + kernel_iohw = kernel + + ref_res = tvm.topi.testing.conv2d_transpose_nchw_python(data, kernel_iohw, 2, 1, (1, 1)) + + for target, dev in tvm.testing.enabled_targets(): + op_res1 = relay.create_executor("graph", device=dev, target=target).evaluate(func)( + data, kernel + ) + tvm.testing.assert_allclose(op_res1.numpy(), ref_res, rtol=1e-5, atol=1e-5) @tvm.testing.uses_gpu @@ -866,8 +882,7 @@ def test_conv2d_transpose_nhwc_run(): oshape_nhwc = (1, 36, 36, 10) x = relay.var("x", shape=dshape_nhwc) w = relay.var("w") - # kshape and kernel_layout should have swapped IO. - # kshape is HWOI and kernel_layout is HWIO + y = relay.nn.conv2d_transpose( x, w, @@ -877,13 +892,12 @@ def test_conv2d_transpose_nhwc_run(): padding=(1, 1), output_padding=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", ) func = relay.Function([x, w], y) dtype = "float32" data = np.random.uniform(size=dshape_nhwc).astype(dtype) kernel = np.random.uniform(size=kshape_hwoi).astype(dtype) - # use true kshape layout here - HWOI ref_res = tvm.topi.testing.conv2d_transpose_nhwc_python( data, kernel, "HWOI", 2, 1, output_padding=(1, 1) diff --git a/tests/python/relay/test_pass_convert_op_layout.py b/tests/python/relay/test_pass_convert_op_layout.py index 2359dcdf93d9..5360018dcd67 100644 --- a/tests/python/relay/test_pass_convert_op_layout.py +++ b/tests/python/relay/test_pass_convert_op_layout.py @@ -16,15 +16,12 @@ # under the License. """Test alter op layout pass""" import pytest - import tvm -from tvm import te - -from tvm import relay +from tvm import relay, te +from tvm.relay import analysis, transform +from tvm.relay.op import op as reg from tvm.relay.op import register_alter_op_layout -from tvm.relay import transform, analysis from tvm.relay.transform.infer_layout_utils import InferCorrectLayoutOutput -from tvm.relay.op import op as reg def run_opt_pass(expr, passes): @@ -182,7 +179,7 @@ def expected(): x = relay.var("x", shape=(1, 56, 56, 64)) weight = relay.var("weight", shape=(3, 3, 64, 64)) x = relay.layout_transform(x, "NHWC", "NCHW") - weight = relay.layout_transform(weight, "HWIO", "OIHW") + weight = relay.layout_transform(weight, "HWIO", "IOHW") y = relay.nn.conv2d_transpose(x, weight, channels=64, kernel_size=(3, 3), padding=(1, 1)) y = relay.nn.relu(y) y = relay.layout_transform(y, "NCHW", "NHWC") @@ -190,7 +187,7 @@ def expected(): return y a = before() - a = run_opt_pass(a, transform.ConvertLayout({"nn.conv2d_transpose": ["NCHW", "OIHW"]})) + a = run_opt_pass(a, transform.ConvertLayout({"nn.conv2d_transpose": ["NCHW", "IOHW"]})) b = run_opt_pass(expected(), transform.InferType()) assert tvm.ir.structural_equal(a, b), "Actual = \n" + str(a) @@ -1373,7 +1370,7 @@ def expected(): x = relay.var("x", shape=(1, 56, 56, 64), dtype="int8") weight = relay.var("weight", shape=(3, 3, 64, 64), dtype="int8") x = relay.layout_transform(x, "NHWC", "NCHW") - weight = relay.layout_transform(weight, "HWIO", "OIHW") + weight = relay.layout_transform(weight, "HWIO", "IOHW") y = relay.qnn.op.conv2d_transpose( x, weight, @@ -1403,7 +1400,6 @@ def expected(): a = before() a = run_opt_pass(a, transform.ConvertLayout({"qnn.conv2d_transpose": ["NCHW", "default"]})) b = run_opt_pass(expected(), transform.InferType()) - assert tvm.ir.structural_equal(a, b), "Actual = \n" + str(a) diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index c49d837ed920..5f01189b7e33 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -17,13 +17,13 @@ # pylint: disable=unused-wildcard-import import numpy as np import pytest - import tvm from tvm import relay def compare_fq_to_int(expr, args, allow_rounding_error=False): mod = tvm.IRModule.from_expr(expr) + breakpoint() mod = tvm.relay.transform.InferType()(mod) mod_int = tvm.relay.transform.FakeQuantizationToInteger()(mod) @@ -89,6 +89,28 @@ def test_fake_quantize_conv_per_channel(): compare_fq_to_int(op, [x_np, w_np], allow_rounding_error=True) +def test_fake_quantize_transposeconv(): + for out_dtype in ["int8", "uint8"]: + x = relay.var("x", shape=[1, 3, 224, 224], dtype="int8") + w = relay.var("w", shape=[3, 16, 5, 5], dtype="int8") + one = relay.const(1.0) + zero = relay.const(0) + + op = relay.op.nn.conv2d_transpose( + relay.qnn.op.dequantize(x, relay.const(2.0), zero), + relay.qnn.op.dequantize(w, relay.const(0.5), zero), + kernel_size=[5, 5], + data_layout="NCHW", + kernel_layout="IOHW", + ) + op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) + + x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") + w_np = np.random.randint(-128, 127, size=[16, 3, 5, 5], dtype="int8") + + compare_fq_to_int(op, [x_np, w_np]) + + def test_fake_quantize_dense(): for out_dtype in ["int8", "uint8"]: x = relay.var("x", shape=[128, 64], dtype="int8") From b200eeb81a1705317ff7b075ad326884c04b7691 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:18:24 -0700 Subject: [PATCH 17/35] make pytorch tests pass --- python/tvm/relay/frontend/pytorch.py | 11 ++++--- tests/python/frontend/pytorch/test_forward.py | 30 +------------------ 2 files changed, 8 insertions(+), 33 deletions(-) diff --git a/python/tvm/relay/frontend/pytorch.py b/python/tvm/relay/frontend/pytorch.py index 3fc202a7cc91..b3aecb589352 100644 --- a/python/tvm/relay/frontend/pytorch.py +++ b/python/tvm/relay/frontend/pytorch.py @@ -19,8 +19,8 @@ # pylint: disable=import-outside-toplevel, simplifiable-if-expression, cell-var-from-loop, unnecessary-lambda # pylint: disable=missing-function-docstring """PT: PyTorch frontend.""" -import itertools import functools +import itertools import logging import math import sys @@ -40,11 +40,11 @@ from ..prelude import Prelude, StaticTensorArrayOps from ..ty import Any, TensorType, TupleType from . import qnn_torch -from .common import AttrCvt, get_relay_op, unbind, lstm_cell, gru_cell -from .common import infer_value as _infer_value +from .common import AttrCvt, get_relay_op, gru_cell from .common import infer_shape as _infer_shape +from .common import infer_value as _infer_value from .common import infer_value_simulated as _infer_value_simulated -from .common import try_infer_value +from .common import lstm_cell, try_infer_value, unbind from .pytorch_utils import is_version_greater_than __all__ = ["from_pytorch"] @@ -1022,6 +1022,9 @@ def convolution(self, inputs, input_types): elif len(kernel_size) == 2: data_layout = "NCHW" kernel_layout = "OIHW" + if use_transpose: + # Transposed convolutions have IOHW layout. + kernel_layout = "IOHW" else: data_layout = "NCW" kernel_layout = "OIW" diff --git a/tests/python/frontend/pytorch/test_forward.py b/tests/python/frontend/pytorch/test_forward.py index 0031f4143fab..a05c66d2c974 100644 --- a/tests/python/frontend/pytorch/test_forward.py +++ b/tests/python/frontend/pytorch/test_forward.py @@ -21,6 +21,7 @@ from time import time import numpy as np +import pytest import torch import torchvision import tvm @@ -32,7 +33,6 @@ from tvm import relay from tvm.contrib import graph_executor from tvm.contrib.nvcc import have_fp16 -import pytest sys.setrecursionlimit(10000) @@ -994,20 +994,6 @@ def test_forward_conv_transpose( # opt to make the stride 1 + output padding stride = output_padding + 1 - # Conv 3D Transpose Tests - conv3d_input_shape = [1, in_channels, 16, 16, 16] - conv3d_input_data = torch.rand(conv3d_input_shape).float() - conv3d_transpose = torch.nn.ConvTranspose3d( - in_channels=in_channels, - out_channels=out_channels, - kernel_size=kernel_size, - stride=stride, - output_padding=output_padding, - groups=groups, - bias=bias, - ).eval() - verify_model(conv3d_transpose, conv3d_input_data) - # Conv 2D Transpose Tests conv2d_input_shape = [1, in_channels, 128, 256] conv2d_input_data = torch.rand(conv2d_input_shape).float() @@ -1022,20 +1008,6 @@ def test_forward_conv_transpose( ).eval() verify_model(conv2d_transpose, conv2d_input_data) - # # Conv 1D Transpose Tests - conv1d_input_shape = [1, in_channels, 10] - conv1d_input_data = torch.rand(conv1d_input_shape).float() - conv1d_transpose = torch.nn.ConvTranspose1d( - in_channels=in_channels, - out_channels=out_channels, - kernel_size=kernel_size, - stride=stride, - output_padding=output_padding, - groups=groups, - bias=bias, - ).eval() - verify_model(conv1d_transpose, conv1d_input_data) - def test_forward_deform_conv(): torch.set_grad_enabled(False) From 3223e8814ebafe7246c3af4ec698d2f6e9e61b44 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 18:42:48 -0700 Subject: [PATCH 18/35] lint --- python/tvm/relay/transform/fake_quantization_to_integer.py | 1 + python/tvm/topi/nn/conv2d_transpose.py | 3 ++- python/tvm/topi/nn/utils.py | 1 - 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index ecf5d9581bd8..febccde83dcf 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -159,6 +159,7 @@ def conv2d(expr, type_map): @register_fake_quantization_to_integer("nn.conv2d_transpose") def conv2d_transpose(expr, type_map): + """Rewrite a conv2d_transpose op""" attrs = {**expr.attrs} attrs.pop("out_dtype") x, weight = expr.args diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index 8291bfed4ea6..e668ed0b20d4 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -186,7 +186,8 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Convert back to original NHWC layout. out = relay.transpose(deconv, axes=(0, 2, 3, 1)) return out - elif attrs["data_layout"] == "NCHW": + + if attrs["data_layout"] == "NCHW": kernel = layout_transform(kernel, kernel_layout, "IOHW") new_attrs = {k: attrs[k] for k in attrs.keys()} diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index eea35c814c11..d1c219d8eca0 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,7 +19,6 @@ from __future__ import absolute_import import tvm -from tvm import relay from ..utils import get_const_int From 3c453406ba615cce85fbd0113931af1b4d978dbd Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 21:57:15 -0700 Subject: [PATCH 19/35] add test --- tests/python/relay/test_pass_fake_quantization_to_integer.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index 5f01189b7e33..8ac0f91bb066 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -23,7 +23,6 @@ def compare_fq_to_int(expr, args, allow_rounding_error=False): mod = tvm.IRModule.from_expr(expr) - breakpoint() mod = tvm.relay.transform.InferType()(mod) mod_int = tvm.relay.transform.FakeQuantizationToInteger()(mod) From 2a95c3f8e535af4942be83402ff933e837b79db0 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Wed, 20 Oct 2021 22:23:03 -0700 Subject: [PATCH 20/35] fix bug with layout transform --- python/tvm/topi/nn/conv2d_transpose.py | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/python/tvm/topi/nn/conv2d_transpose.py b/python/tvm/topi/nn/conv2d_transpose.py index e668ed0b20d4..99c7442240c7 100644 --- a/python/tvm/topi/nn/conv2d_transpose.py +++ b/python/tvm/topi/nn/conv2d_transpose.py @@ -146,8 +146,8 @@ def layout_transform(tensor: "relay.Expr", current_layout: str, desired_layout: desired_layout_map = {c: i for i, c in enumerate(desired_layout)} axes = [None] * len(current_layout) - for c, i in current_layout_map.items(): - axes[i] = desired_layout_map[c] + for c, i in desired_layout_map.items(): + axes[i] = current_layout_map[c] return relay.transpose(tensor, axes=axes) @@ -169,6 +169,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): result : tvm.relay.Expr The legalized expr """ + data, kernel = inputs kernel_layout = attrs["kernel_layout"] if attrs["data_layout"] == "NHWC": @@ -177,7 +178,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): # Set new attrs for conv2d_transpose. new_attrs = {k: attrs[k] for k in attrs.keys()} new_attrs["data_layout"] = "NCHW" - # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + # layout of kernel should be IOHW, but kernel_layout will be swapped - OIHW new_attrs["kernel_layout"] = "IOHW" # Convert data to NCHW. @@ -191,7 +192,7 @@ def conv2d_transpose_legalize(attrs, inputs, types): kernel = layout_transform(kernel, kernel_layout, "IOHW") new_attrs = {k: attrs[k] for k in attrs.keys()} - # layout of kernel should be IOHW, but kernel_layout should be swapped - OIHW + # layout of kernel should be IOHW, but kernel_layout will be swapped - OIHW new_attrs["kernel_layout"] = "IOHW" return relay.nn.conv2d_transpose(data, kernel, **new_attrs) From 737238e7ea0fd2beedde98f1fcb9f18d6b995fab Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 14:26:41 -0700 Subject: [PATCH 21/35] change layouts for conv2d_transpose too --- python/tvm/relay/op/contrib/vitis_ai.py | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index 7b9324cdc671..31a7bb555367 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -18,13 +18,13 @@ """Vitis-AI codegen annotation of supported operators""" import warnings -import numpy as np -from tvm import relay +import numpy as np import tvm._ffi +from tvm import relay from tvm.relay import transform -from tvm.relay.expr import Tuple, TupleGetItem from tvm.relay.build_module import bind_params_by_name +from tvm.relay.expr import Tuple, TupleGetItem from tvm.relay.op.annotation import compiler_begin, compiler_end # Placeholder for PyXIR module @@ -164,11 +164,13 @@ def partition_for_vitis_ai(mod, params=None, dpu=None, **opts): desired_layouts_in_partition = { "nn.conv2d": ["NHWC", "default"], + "nn.conv2d_transpose": ["NHWC", "default"], "nn.upsampling": ["NHWC"], "image.resize2d": ["NHWC"], } desired_layouts_in_main = { "nn.conv2d": ["NCHW", "default"], + "nn.conv2d_transpose": ["NCHW", "default"], "nn.upsampling": ["NCHW"], "image.resize2d": ["NCHW"], } From 08107d9dc33fc1e972e8b6d8650a6aee994a8211 Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Thu, 21 Oct 2021 15:03:50 -0700 Subject: [PATCH 22/35] fix vitis tests --- python/tvm/relay/op/contrib/vitis_ai.py | 2 -- .../contrib/test_vitis_ai/test_vitis_ai_codegen.py | 13 ++++++++++++- 2 files changed, 12 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index 31a7bb555367..fe6ee817bc36 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -164,13 +164,11 @@ def partition_for_vitis_ai(mod, params=None, dpu=None, **opts): desired_layouts_in_partition = { "nn.conv2d": ["NHWC", "default"], - "nn.conv2d_transpose": ["NHWC", "default"], "nn.upsampling": ["NHWC"], "image.resize2d": ["NHWC"], } desired_layouts_in_main = { "nn.conv2d": ["NCHW", "default"], - "nn.conv2d_transpose": ["NCHW", "default"], "nn.upsampling": ["NCHW"], "image.resize2d": ["NCHW"], } diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index c5c9cc7f818c..c50bed737d6f 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -241,6 +241,10 @@ def test_upsampling(dpu_target): verify_codegen(mod, dpu_target=dpu_target) +@pytest.mark.skip( + reason="I and O used to be mixed up in kernel layouts in TVM." + "This is fixed, but vitis needs to adopt the new convention" +) @pytest.mark.parametrize( "dpu_target", ["DPUCADF8H", "DPUCAHX8H-u50", "DPUCAHX8L", "DPUCVDX8H", "DPUCVDX8G", "DPUCZDX8G-zcu104"], @@ -253,7 +257,14 @@ def test_conv2d_transpose(dpu_target): x = relay.var("x", shape=dshape) w = relay.const(np.zeros(kshape, dtype="float32")) y = relay.nn.conv2d_transpose( - x, w, channels=10, kernel_size=(3, 3), strides=(1, 1), padding=(1, 1) + x, + w, + channels=10, + kernel_size=(3, 3), + strides=(1, 1), + padding=(1, 1), + data_layout="NCHW", + kernel_layout="IOHW", ) func = relay.Function([x], y) params = {} From d9de5d0c10b7f162696ec75d2f7d663b203fb873 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 15:29:22 -0700 Subject: [PATCH 23/35] fix qnn conv2d transpose tests --- .../relay/test_op_qnn_conv2_transpose.py | 53 +++++++++---------- 1 file changed, 26 insertions(+), 27 deletions(-) diff --git a/tests/python/relay/test_op_qnn_conv2_transpose.py b/tests/python/relay/test_op_qnn_conv2_transpose.py index 9fd3d1b84537..a7c3c9110f6d 100644 --- a/tests/python/relay/test_op_qnn_conv2_transpose.py +++ b/tests/python/relay/test_op_qnn_conv2_transpose.py @@ -15,13 +15,12 @@ # specific language governing permissions and limitations # under the License. -import tvm -from tvm import te import numpy as np -from tvm import relay +import tvm +from tvm import relay, te +from tvm.contrib import graph_executor from tvm.relay import transform from tvm.relay.testing import run_infer_type -from tvm.contrib import graph_executor from tvm.relay.testing.temp_op_attr import TempOpAttr @@ -224,7 +223,7 @@ def test_no_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -248,7 +247,7 @@ def test_no_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -274,7 +273,7 @@ def test_kernel_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -298,7 +297,7 @@ def test_kernel_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -324,7 +323,7 @@ def test_input_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -348,7 +347,7 @@ def test_input_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -374,7 +373,7 @@ def test_both_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -398,7 +397,7 @@ def test_both_zero_point(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -424,7 +423,7 @@ def test_different_dtype(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], ) @@ -449,7 +448,7 @@ def test_different_dtype(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], ) @@ -460,7 +459,7 @@ def test_layout(): # uint8 input data_shape = (2, 2, 4, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -476,14 +475,14 @@ def test_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) data_shape = (2, 2, 4, 3) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 1, 3) # HWIO + kernel_shape = (2, 2, 1, 3) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -499,7 +498,7 @@ def test_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -525,7 +524,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -533,7 +532,7 @@ def test_padding(): # Try different layout data_shape = (2, 2, 4, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -549,7 +548,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -557,7 +556,7 @@ def test_padding(): # Try asymmetric padding data_shape = (2, 8, 6, 4) # NHWC data_dtype = "uint8" - kernel_shape = (2, 2, 3, 4) # HWIO + kernel_shape = (2, 2, 3, 4) # HWOI kernel_dtype = "uint8" ref_func, qnn_func = get_funcs( data_shape=data_shape, @@ -573,7 +572,7 @@ def test_padding(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) verify(ref_func, qnn_func, data_shape, data_dtype, kernel_shape, kernel_dtype) @@ -600,7 +599,7 @@ def test_const_folding(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", channels=kernel_shape[1], groups=1, @@ -614,7 +613,7 @@ def test_broadcast_layout(): # Test broadcast support for NHWC layout. data_shape = (1, 229, 229, 3) # NHWC data_dtype = "uint8" - kernel_shape = (7, 7, 64, 3) # HWIO + kernel_shape = (7, 7, 64, 3) # HWOI kernel_dtype = "int8" _, qnn_func = get_funcs( data_shape=data_shape, @@ -630,7 +629,7 @@ def test_broadcast_layout(): strides=(1, 1), dilation=(1, 1), data_layout="NHWC", - kernel_layout="HWIO", + kernel_layout="HWOI", out_dtype="int32", ) func = qnn_func["main"].body @@ -670,7 +669,7 @@ def test_per_channel_kernel_scale(): strides=(1, 1), dilation=(1, 1), data_layout="NCHW", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) From 815cc4bf8c1cd704190303109f741100be90d22b Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:04:10 -0700 Subject: [PATCH 24/35] fix fake quantization pass --- python/tvm/relay/qnn/op/qnn.py | 2 +- python/tvm/relay/transform/fake_quantization_to_integer.py | 4 +++- tests/python/relay/test_pass_fake_quantization_to_integer.py | 2 +- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/python/tvm/relay/qnn/op/qnn.py b/python/tvm/relay/qnn/op/qnn.py index d1121f7b90fc..4749b63b858f 100644 --- a/python/tvm/relay/qnn/op/qnn.py +++ b/python/tvm/relay/qnn/op/qnn.py @@ -387,7 +387,7 @@ def conv2d_transpose( kernel_layout="IOHW", out_layout="", output_padding=(0, 0), - out_dtype="", + out_dtype="int32", ): """This operator deconvolves quantized data with quantized kernel. The scale of the output quantized tensor is the product of the kernel_scale and diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index febccde83dcf..b8c63f4aa5ef 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -18,6 +18,7 @@ import tvm from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType +from tvm.ir.module import IRModule from tvm.tir import bijective_layout from ..op import register_fake_quantization_to_integer @@ -167,8 +168,9 @@ def conv2d_transpose(expr, type_map): w_t = type_map[weight] conv_scale = fold_constant(x_t.scale * w_t.scale) conv_zp = get_zeros(conv_scale) + out = relay.qnn.op.conv2d_transpose( - x, weight, x_t.zero_point, w_t.zero_point, x_t.zero_point, w_t.zero_point, **attrs + x, weight, x_t.zero_point, w_t.zero_point, x_t.scale, w_t.scale, **attrs ) out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index 8ac0f91bb066..f99c15639dd5 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -105,7 +105,7 @@ def test_fake_quantize_transposeconv(): op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") - w_np = np.random.randint(-128, 127, size=[16, 3, 5, 5], dtype="int8") + w_np = np.random.randint(-128, 127, size=[3, 16, 5, 5], dtype="int8") compare_fq_to_int(op, [x_np, w_np]) From 9baac547047dc854b8bea1e2c9a4a67e799508b1 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:08:06 -0700 Subject: [PATCH 25/35] add todo --- .../contrib/test_vitis_ai/test_vitis_ai_codegen.py | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index c50bed737d6f..b89cc37d29b8 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -19,24 +19,23 @@ """Vitis-AI codegen tests""" import sys -import numpy as np +import numpy as np import pytest pytest.importorskip("pyxir") import pyxir.contrib.target.DPUCADF8H import pyxir.contrib.target.DPUCAHX8H import pyxir.contrib.target.DPUCAHX8L -import pyxir.contrib.target.DPUCVDX8H import pyxir.contrib.target.DPUCVDX8G +import pyxir.contrib.target.DPUCVDX8H import pyxir.contrib.target.DPUCZDX8G - import tvm from tvm import relay +from tvm.contrib.target import vitis_ai from tvm.relay import transform -from tvm.relay.op.contrib.vitis_ai import annotation from tvm.relay.build_module import bind_params_by_name -from tvm.contrib.target import vitis_ai +from tvm.relay.op.contrib.vitis_ai import annotation from .infrastructure import skip_test, verify_codegen @@ -243,7 +242,10 @@ def test_upsampling(dpu_target): @pytest.mark.skip( reason="I and O used to be mixed up in kernel layouts in TVM." - "This is fixed, but vitis needs to adopt the new convention" + "This is fixed, but vitis needs to adopt the new convention." + "To change, simply remove this line:" + "https://github.com/Xilinx/pyxir/blob/bef661d6d77adcdbd2cf4163f2cf3a1d31d40406/" + "python/pyxir/frontend/tvm/relay_tools/relay_l2_convolution.py#L380" ) @pytest.mark.parametrize( "dpu_target", From 435e3b287d99d0a5b0e74665132760979f1a17c1 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:21:32 -0700 Subject: [PATCH 26/35] lint --- python/tvm/relay/transform/fake_quantization_to_integer.py | 1 - 1 file changed, 1 deletion(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index b8c63f4aa5ef..2bb9cf463a64 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -18,7 +18,6 @@ import tvm from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType -from tvm.ir.module import IRModule from tvm.tir import bijective_layout from ..op import register_fake_quantization_to_integer From 1228fc5f453c08a8f95e66db4bf589b1bc024960 Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:36:40 -0700 Subject: [PATCH 27/35] undo just formatting changes --- python/tvm/relay/op/contrib/vitis_ai.py | 6 ++-- python/tvm/topi/nn/utils.py | 1 - tests/python/frontend/pytorch/test_forward.py | 30 ++++++++++++++++++- 3 files changed, 32 insertions(+), 5 deletions(-) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index fe6ee817bc36..7b9324cdc671 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -18,13 +18,13 @@ """Vitis-AI codegen annotation of supported operators""" import warnings - import numpy as np -import tvm._ffi + from tvm import relay +import tvm._ffi from tvm.relay import transform -from tvm.relay.build_module import bind_params_by_name from tvm.relay.expr import Tuple, TupleGetItem +from tvm.relay.build_module import bind_params_by_name from tvm.relay.op.annotation import compiler_begin, compiler_end # Placeholder for PyXIR module diff --git a/python/tvm/topi/nn/utils.py b/python/tvm/topi/nn/utils.py index d1c219d8eca0..ff00441e9850 100644 --- a/python/tvm/topi/nn/utils.py +++ b/python/tvm/topi/nn/utils.py @@ -19,7 +19,6 @@ from __future__ import absolute_import import tvm - from ..utils import get_const_int diff --git a/tests/python/frontend/pytorch/test_forward.py b/tests/python/frontend/pytorch/test_forward.py index a05c66d2c974..0031f4143fab 100644 --- a/tests/python/frontend/pytorch/test_forward.py +++ b/tests/python/frontend/pytorch/test_forward.py @@ -21,7 +21,6 @@ from time import time import numpy as np -import pytest import torch import torchvision import tvm @@ -33,6 +32,7 @@ from tvm import relay from tvm.contrib import graph_executor from tvm.contrib.nvcc import have_fp16 +import pytest sys.setrecursionlimit(10000) @@ -994,6 +994,20 @@ def test_forward_conv_transpose( # opt to make the stride 1 + output padding stride = output_padding + 1 + # Conv 3D Transpose Tests + conv3d_input_shape = [1, in_channels, 16, 16, 16] + conv3d_input_data = torch.rand(conv3d_input_shape).float() + conv3d_transpose = torch.nn.ConvTranspose3d( + in_channels=in_channels, + out_channels=out_channels, + kernel_size=kernel_size, + stride=stride, + output_padding=output_padding, + groups=groups, + bias=bias, + ).eval() + verify_model(conv3d_transpose, conv3d_input_data) + # Conv 2D Transpose Tests conv2d_input_shape = [1, in_channels, 128, 256] conv2d_input_data = torch.rand(conv2d_input_shape).float() @@ -1008,6 +1022,20 @@ def test_forward_conv_transpose( ).eval() verify_model(conv2d_transpose, conv2d_input_data) + # # Conv 1D Transpose Tests + conv1d_input_shape = [1, in_channels, 10] + conv1d_input_data = torch.rand(conv1d_input_shape).float() + conv1d_transpose = torch.nn.ConvTranspose1d( + in_channels=in_channels, + out_channels=out_channels, + kernel_size=kernel_size, + stride=stride, + output_padding=output_padding, + groups=groups, + bias=bias, + ).eval() + verify_model(conv1d_transpose, conv1d_input_data) + def test_forward_deform_conv(): torch.set_grad_enabled(False) From 35e56177254459faea896d323fc84893ef02548e Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:41:30 -0700 Subject: [PATCH 28/35] remove formatting only change --- tests/python/relay/test_autotvm_task_extraction.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/python/relay/test_autotvm_task_extraction.py b/tests/python/relay/test_autotvm_task_extraction.py index f83a23d38adb..83480a044f45 100644 --- a/tests/python/relay/test_autotvm_task_extraction.py +++ b/tests/python/relay/test_autotvm_task_extraction.py @@ -16,7 +16,8 @@ # under the License. """Test task extraction for autotvm""" import tvm.relay.testing -from tvm import autotvm, relay +from tvm import relay +from tvm import autotvm def get_network(name, batch_size): From 30101fa56be9b1e84dfb65ae6679073a9ab7687e Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 16:43:38 -0700 Subject: [PATCH 29/35] remove f2qi for later pr --- .../transform/fake_quantization_to_integer.py | 20 ---------------- .../test_pass_fake_quantization_to_integer.py | 23 +------------------ 2 files changed, 1 insertion(+), 42 deletions(-) diff --git a/python/tvm/relay/transform/fake_quantization_to_integer.py b/python/tvm/relay/transform/fake_quantization_to_integer.py index 2bb9cf463a64..1adde9a4a430 100644 --- a/python/tvm/relay/transform/fake_quantization_to_integer.py +++ b/python/tvm/relay/transform/fake_quantization_to_integer.py @@ -19,7 +19,6 @@ from tvm import relay from tvm.ir import TensorAffineType, TupleAffineType from tvm.tir import bijective_layout - from ..op import register_fake_quantization_to_integer @@ -157,25 +156,6 @@ def conv2d(expr, type_map): return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] -@register_fake_quantization_to_integer("nn.conv2d_transpose") -def conv2d_transpose(expr, type_map): - """Rewrite a conv2d_transpose op""" - attrs = {**expr.attrs} - attrs.pop("out_dtype") - x, weight = expr.args - x_t = type_map[x] - w_t = type_map[weight] - conv_scale = fold_constant(x_t.scale * w_t.scale) - conv_zp = get_zeros(conv_scale) - - out = relay.qnn.op.conv2d_transpose( - x, weight, x_t.zero_point, w_t.zero_point, x_t.scale, w_t.scale, **attrs - ) - out_layout = attrs["out_layout"] if attrs["out_layout"] != "" else attrs["data_layout"] - out_axis = bijective_layout(out_layout, "NCHW").backward_index(list(range(4)))[1] - return [out, TensorAffineType(conv_scale, conv_zp, out.attrs.out_dtype, out_axis.value)] - - @register_fake_quantization_to_integer("nn.dense") def dense(expr, type_map): """Rewrite a dense op""" diff --git a/tests/python/relay/test_pass_fake_quantization_to_integer.py b/tests/python/relay/test_pass_fake_quantization_to_integer.py index f99c15639dd5..c49d837ed920 100644 --- a/tests/python/relay/test_pass_fake_quantization_to_integer.py +++ b/tests/python/relay/test_pass_fake_quantization_to_integer.py @@ -17,6 +17,7 @@ # pylint: disable=unused-wildcard-import import numpy as np import pytest + import tvm from tvm import relay @@ -88,28 +89,6 @@ def test_fake_quantize_conv_per_channel(): compare_fq_to_int(op, [x_np, w_np], allow_rounding_error=True) -def test_fake_quantize_transposeconv(): - for out_dtype in ["int8", "uint8"]: - x = relay.var("x", shape=[1, 3, 224, 224], dtype="int8") - w = relay.var("w", shape=[3, 16, 5, 5], dtype="int8") - one = relay.const(1.0) - zero = relay.const(0) - - op = relay.op.nn.conv2d_transpose( - relay.qnn.op.dequantize(x, relay.const(2.0), zero), - relay.qnn.op.dequantize(w, relay.const(0.5), zero), - kernel_size=[5, 5], - data_layout="NCHW", - kernel_layout="IOHW", - ) - op = relay.qnn.op.quantize(op, one, zero, out_dtype=out_dtype) - - x_np = np.random.randint(-128, 127, size=[1, 3, 224, 224], dtype="int8") - w_np = np.random.randint(-128, 127, size=[3, 16, 5, 5], dtype="int8") - - compare_fq_to_int(op, [x_np, w_np]) - - def test_fake_quantize_dense(): for out_dtype in ["int8", "uint8"]: x = relay.var("x", shape=[128, 64], dtype="int8") From 0fa4acb53cb1b4f06de3d459b9e0839bdffb549b Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Thu, 21 Oct 2021 21:32:08 -0700 Subject: [PATCH 30/35] more frontend tests fixes --- python/tvm/relay/frontend/mxnet.py | 46 ++++++++++++++++---------- python/tvm/relay/frontend/qnn_torch.py | 8 ++--- python/tvm/relay/frontend/tflite.py | 39 +++++++++++----------- 3 files changed, 50 insertions(+), 43 deletions(-) diff --git a/python/tvm/relay/frontend/mxnet.py b/python/tvm/relay/frontend/mxnet.py index 59b4e99de999..1b1d60119967 100644 --- a/python/tvm/relay/frontend/mxnet.py +++ b/python/tvm/relay/frontend/mxnet.py @@ -18,40 +18,50 @@ """MXNet symbol frontend.""" import json import math + import numpy as np import tvm -from tvm.ir import IRModule - from tvm import relay +from tvm.ir import IRModule from tvm.topi.utils import get_const_tuple + +from ... import nd as _nd from .. import analysis from .. import expr as _expr from .. import function as _function from .. import op as _op from .. import scope_builder as _scope_builder -from ... import nd as _nd - from .common import StrAttrsDict -from .common import infer_type as _infer_type +from .common import get_name as _get_name from .common import infer_shape as _infer_shape +from .common import infer_type as _infer_type from .common import infer_value as _infer_value -from .common import get_name as _get_name -from .nnvm_common import _rename, _binop_scalar, _rbinop_scalar, _reduce -from .nnvm_common import _arg_reduce, _init_op, _softmax_op, _cast -from .nnvm_common import _clip, _transpose, _upsampling -from .nnvm_common import _elemwise_sum, _reshape -from .nnvm_common import _warn_not_used from .mxnet_qnn_op_utils import ( - quantize_mxnet_min_max, - quantize_conv_weights_bias_channel_mkldnn_from_var, - quantize_conv_bias_mkldnn_from_var, - get_conv_mkldnn_requantized_scale_outDtype, dequantize_mxnet_min_max, + get_conv_mkldnn_requantized_scale_outDtype, get_mkldnn_int8_scale, - get_mkldnn_uint8_scale, get_mkldnn_requantize_scale_outDtype, + get_mkldnn_uint8_scale, + quantize_conv_bias_mkldnn_from_var, + quantize_conv_weights_bias_channel_mkldnn_from_var, + quantize_mxnet_min_max, +) +from .nnvm_common import ( + _arg_reduce, + _binop_scalar, + _cast, + _clip, + _elemwise_sum, + _init_op, + _rbinop_scalar, + _reduce, + _rename, + _reshape, + _softmax_op, + _transpose, + _upsampling, + _warn_not_used, ) - __all__ = ["from_mxnet"] @@ -329,7 +339,7 @@ def _mx_conv2d_transpose(inputs, attrs): if "kernel_layout" in attrs.attrs: kernel_layout = attrs.get_str("kernel_layout") else: - kernel_layout = "HWIO" if data_layout == "NHWC" else "OIHW" + kernel_layout = "HWIO" if data_layout == "NHWC" else "IOHW" new_attrs = {} new_attrs["channels"] = attrs.get_int("num_filter") diff --git a/python/tvm/relay/frontend/qnn_torch.py b/python/tvm/relay/frontend/qnn_torch.py index 172ab1e41268..5772313fb97f 100644 --- a/python/tvm/relay/frontend/qnn_torch.py +++ b/python/tvm/relay/frontend/qnn_torch.py @@ -19,7 +19,6 @@ import logging import numpy as np - import tvm from tvm import relay from tvm.relay import expr as _expr @@ -1043,11 +1042,8 @@ def _impl(inputs, _): weight_shape = list(infer_shape(weight)) - # Swap I and O dims to match shape relay expects for OIHW - weight_shape[0], weight_shape[1] = weight_shape[1], weight_shape[0] - kernel_size = (weight_shape[2], weight_shape[3]) - out_channels = weight_shape[0] + out_channels = weight_shape[1] conv_out = relay.qnn.op.conv2d_transpose( inputs[0], @@ -1064,7 +1060,7 @@ def _impl(inputs, _): channels=out_channels, output_padding=output_padding, out_dtype="int32", - kernel_layout="OIHW", + kernel_layout="IOHW", ) return _do_bias_and_requantize( diff --git a/python/tvm/relay/frontend/tflite.py b/python/tvm/relay/frontend/tflite.py index 3688ff5ff4e5..5d3681ebe132 100644 --- a/python/tvm/relay/frontend/tflite.py +++ b/python/tvm/relay/frontend/tflite.py @@ -16,24 +16,25 @@ # under the License. # pylint: disable=invalid-name, unused-argument, too-many-lines, import-outside-toplevel """Tensorflow lite frontend.""" -import math import itertools +import math + import numpy as np import tvm +from tvm import relay from tvm.ir import IRModule -from tvm import relay +from ... import nd as _nd from .. import analysis from .. import expr as _expr from .. import function as _function from .. import op as _op from .. import qnn as _qnn -from ... import nd as _nd from .common import ExprTable -from .common import infer_shape as _infer_shape, to_int_list +from .common import infer_shape as _infer_shape +from .common import to_int_list from .tflite_flexbuffer import FlexBufferDecoder - __all__ = ["from_tflite"] @@ -53,9 +54,9 @@ class OperatorConverter(object): def __init__(self, model, subgraph, exp_tab): try: + from tflite.ActivationFunctionType import ActivationFunctionType from tflite.BuiltinOperator import BuiltinOperator from tflite.BuiltinOptions import BuiltinOptions - from tflite.ActivationFunctionType import ActivationFunctionType except ImportError: raise ImportError("The tflite package must be installed") @@ -1061,8 +1062,8 @@ def convert_log_softmax(self, op): def convert_concatenation(self, op): """Convert TFLite concatenation""" try: - from tflite.ConcatenationOptions import ConcatenationOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.ConcatenationOptions import ConcatenationOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1248,10 +1249,10 @@ def _convert_elemwise(self, relay_op, op, ignore_qnn_params=False): """Generic method to Convert TFLite elemwise""" try: from tflite.AddOptions import AddOptions - from tflite.SubOptions import SubOptions - from tflite.MulOptions import MulOptions - from tflite.DivOptions import DivOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.DivOptions import DivOptions + from tflite.MulOptions import MulOptions + from tflite.SubOptions import SubOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1810,9 +1811,9 @@ def convert_reduce_any(self, op): def _convert_arg_min_max(self, relay_op, op): """Generic method converting TFLite arg_min_max""" try: - from tflite.BuiltinOptions import BuiltinOptions - from tflite.ArgMinOptions import ArgMinOptions from tflite.ArgMaxOptions import ArgMaxOptions + from tflite.ArgMinOptions import ArgMinOptions + from tflite.BuiltinOptions import BuiltinOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -1859,8 +1860,8 @@ def convert_arg_max(self, op): def convert_fully_connected(self, op): """Convert TFLite fully connected""" try: - from tflite.FullyConnectedOptions import FullyConnectedOptions from tflite.BuiltinOptions import BuiltinOptions + from tflite.FullyConnectedOptions import FullyConnectedOptions from tflite.TensorType import TensorType except ImportError: raise ImportError("The tflite package must be installed") @@ -2030,10 +2031,10 @@ def convert_conv(self, op, conv_type): """convolution implementation.""" try: from tflite.BuiltinOptions import BuiltinOptions - from tflite.TensorType import TensorType from tflite.Conv2DOptions import Conv2DOptions from tflite.DepthwiseConv2DOptions import DepthwiseConv2DOptions from tflite.Padding import Padding + from tflite.TensorType import TensorType except ImportError: raise ImportError("The tflite package must be installed") @@ -2440,8 +2441,8 @@ def convert_pool2d(self, op, pool_type): """pool2d implementation.""" try: from tflite.BuiltinOptions import BuiltinOptions - from tflite.Pool2DOptions import Pool2DOptions from tflite.Padding import Padding + from tflite.Pool2DOptions import Pool2DOptions except ImportError: raise ImportError("The tflite package must be installed") @@ -2856,9 +2857,9 @@ def convert_transpose_conv(self, op): """Convert TFLite TRANSPOSE_CONV""" try: from tflite.BuiltinOptions import BuiltinOptions + from tflite.Padding import Padding from tflite.TensorType import TensorType from tflite.TransposeConvOptions import TransposeConvOptions - from tflite.Padding import Padding except ImportError: raise ImportError("The tflite package must be installed") @@ -2952,7 +2953,7 @@ def convert_transpose_conv(self, op): channels=int(out_channels), kernel_size=(int(kernel_h), int(kernel_w)), data_layout="NHWC", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype="int32", ) else: @@ -2964,7 +2965,7 @@ def convert_transpose_conv(self, op): channels=int(out_channels), kernel_size=(int(kernel_h), int(kernel_w)), data_layout="NHWC", - kernel_layout="OIHW", + kernel_layout="IOHW", out_dtype=output_tensor_type_str, ) @@ -3723,8 +3724,8 @@ def from_tflite(model, shape_dict=None, dtype_dict=None, op_converter=OperatorCo The parameter dict to be used by relay """ try: - import tflite.SubGraph import tflite.BuiltinOperator + import tflite.SubGraph except ImportError: raise ImportError("The tflite package must be installed") From 5424dcfee289d7b530adcbf3bc6dcb4b69d17e55 Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Fri, 22 Oct 2021 10:50:21 -0700 Subject: [PATCH 31/35] jostle From 54bd9201754ffdfa6186808b233d942f875c676e Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Fri, 22 Oct 2021 15:11:34 -0700 Subject: [PATCH 32/35] fix keras --- python/tvm/relay/frontend/keras.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/python/tvm/relay/frontend/keras.py b/python/tvm/relay/frontend/keras.py index bf6293a2a90c..4a0336a056b3 100644 --- a/python/tvm/relay/frontend/keras.py +++ b/python/tvm/relay/frontend/keras.py @@ -354,11 +354,14 @@ def _convert_convolution(inexpr, keras_layer, etab): else: kernel_layout = "HWIO" else: - kernel_layout = "OIHW" + if is_deconv: + kernel_layout = "IOHW" + else: + kernel_layout = "OIHW" if is_deconv: kernel_h, kernel_w, n_filters, in_channels = weight.shape - if kernel_layout == "OIHW": + if kernel_layout == "IOHW": weight = weight.transpose([3, 2, 0, 1]) elif is_depthconv: kernel_h, kernel_w, in_channels, depth_mult = weight.shape From dbfa74b501e88f7620f2ffc9a05cdea98e18ff4b Mon Sep 17 00:00:00 2001 From: Andrew Luo Date: Sun, 24 Oct 2021 18:06:17 -0700 Subject: [PATCH 33/35] fix another frontend test --- python/tvm/relay/frontend/caffe.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/python/tvm/relay/frontend/caffe.py b/python/tvm/relay/frontend/caffe.py index b8273b0324c0..f0938def86e2 100644 --- a/python/tvm/relay/frontend/caffe.py +++ b/python/tvm/relay/frontend/caffe.py @@ -21,11 +21,12 @@ import numpy as np import tvm from tvm.ir import IRModule + +from ... import nd as _nd from .. import analysis from .. import expr as _expr from .. import function as _function from .. import op as _op -from ... import nd as _nd from .common import ExprTable from .common import infer_shape as _infer_shape @@ -513,6 +514,9 @@ def convert_deconv(self, op): weight_shape = [-1, conv_params.num_output, kh, kw] weight_value = np.asarray(weight.data, np.float32) weight_value = np.reshape(weight_value, weight_shape) + + # weight shape is in relay's IOHW format rn, we need it to be OIHW + weight_value = np.transpose(weight_value, [1, 0, 2, 3]) else: raise Exception("No weight value of layer {} in caffemodel".format(op.name)) @@ -520,7 +524,6 @@ def convert_deconv(self, op): in_expr = self.exp_tab.get_expr(inputs[0]) out = _op.nn.conv2d_transpose(data=in_expr, weight=weight_expr, **params) if bias: - bias_value = np.asarray(bias.data, np.float32) bias_expr = self.exp_tab.new_const(bias_value, dtype="float32") out = _op.nn.bias_add(out, bias_expr) From 7a2838d9ce4e603b4e0b164d0dfb4826f662f0f3 Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Mon, 25 Oct 2021 10:51:37 -0700 Subject: [PATCH 34/35] fix things --- python/tvm/relay/frontend/tensorflow_ops.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/python/tvm/relay/frontend/tensorflow_ops.py b/python/tvm/relay/frontend/tensorflow_ops.py index a8213d4b1c49..26ea4f4dbc2a 100644 --- a/python/tvm/relay/frontend/tensorflow_ops.py +++ b/python/tvm/relay/frontend/tensorflow_ops.py @@ -461,8 +461,11 @@ def _impl(inputs, attr, params, mod): raise tvm.error.OpAttributeInvalid(msg.format(attr["padding"])) if "kernel_layout" not in attr: - if opname in ["conv", "conv_transpose"]: + if opname == "conv": attr["kernel_layout"] = "HWIO" if attr["data_format"] == "NHWC" else "OIHW" + elif opname == "conv_transpose": + # conv_transpose in TVM has weights be IOHW for NCHW + attr["kernel_layout"] = "HWIO" if attr["data_format"] == "NHWC" else "IOHW" else: attr["kernel_layout"] = "HWOI" if attr["data_format"] == "NHWC" else "OIHW" From 2c2136e8cb9e18dfbb16b8bbfb6cda25028c99e9 Mon Sep 17 00:00:00 2001 From: Andrew Zhao Luo Date: Mon, 25 Oct 2021 14:28:50 -0700 Subject: [PATCH 35/35] jostle ci