janetsc commented on PR #13002:
URL: https://github.com/apache/tvm/pull/13002#issuecomment-1270916029
This does pass on the simulator, but I am unfortunately still seeing a
mismatch on hardware.
cc: @supersat, @nverke
enabled targets: llvm; hexagon
pytest marker:
============================= test session starts
==============================
platform linux -- Python 3.8.10, pytest-7.1.3, pluggy-1.0.0
rootdir: /home/janetsc/src/tvm
plugins: lazy-fixture-0.6.3, profiling-1.7.0, forked-1.4.0,
rerunfailures-10.2, xdist-2.5.0
collected 19 items
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py FFFFFF [
31%]
FFFFFFFFFFFFF
[100%]
=================================== FAILURES
===================================
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529b50>
act_shape = (1, 8, 4, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f36a54c8a90>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 8.492e+00, 7.918e+00, 2.390e+03],
[ 9.430e+00, -2.876e+03, 1.755e+04]],
[[ 7.977e+00, ...+01]],
[[-6.245e+02, 6.301e+00, 4.840e+00],
[-4.466e+01, -2.139e+01, 2.222e+04]]]], dtype=float16)
desired = array([[[[8.4 , 7.656, 6.133],
[9.43 , 9.32 , 6.484]],
[[7.594, 8.23 , 5.414],
[8.6 , 7.6...4],
[8.32 , 7.51 , 5.62 ]],
[[6.055, 6.312, 4.844],
[7.402, 7.234, 6.316]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 8.492e+00, 7.918e+00, 2.390e+03],
E [ 9.430e+00, -2.876e+03, 1.755e+04]],
E ...
E y: array([[[[8.4 , 7.656, 6.133],
E [9.43 , 9.32 , 6.484]],
E ...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x10x14x3-hwio3x3x3x3-stride1x1] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529ac0>
act_shape = (1, 10, 14, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f2303b80>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.1992e+00, 8.0391e+00, 1.5148e+01],
[-4.7000e+03, -1.8930e+03, 7.2617e+00],
[ 8.7422e...,
[-8.0240e+03, 1.8953e+01, 5.3000e+03],
[-3.6156e+01, 7.9492e+00, 8.7891e+00]]]], dtype=float16)
desired = array([[[[ 7.203, 7.832, 7.38 ],
[ 7.887, 7.96 , 7.086],
[ 8.734, 7.984, 8.04 ],
[ 7.... [ 9.516, 8.18 , 8.43 ],
[ 9.336, 8.59 , 8.59 ],
[ 8.805, 7.96 , 8.734]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 7.1992e+00, 8.0391e+00, 1.5148e+01],
E [-4.7000e+03, -1.8930e+03, 7.2617e+00],
E [ 8.7422e+00, -2.8781e+01, 8.0391e+00],...
E y: array([[[[ 7.203, 7.832, 7.38 ],
E [ 7.887, 7.96 , 7.086],
E [ 8.734, 7.984, 8.04 ],...
python/tvm/testing/utils.py:120: AssertionError
__
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x3-stride1x1] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529a30>
act_shape = (1, 14, 6, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f2352bb0>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.2070e+00, -1.4296e+04, 2.0723e+00],
[-3.3856e+04, 7.5880e+03, 5.7852e+00],
[-1.8538e...,
[ 6.3086e+00, 7.5273e+00, 1.0431e+02],
[-2.2944e+04, 7.8867e+00, 8.0391e+00]]]], dtype=float16)
desired = array([[[[7.168, 7.867, 8.32 ],
[8.664, 8.14 , 8.53 ],
[7.395, 8.66 , 7.957],
[7.42 , 7.64 ...707],
[6.39 , 7.574, 7.58 ],
[6.79 , 7.527, 7.543],
[6.805, 7.887, 8.03 ]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 7.2070e+00, -1.4296e+04, 2.0723e+00],
E [-3.3856e+04, 7.5880e+03, 5.7852e+00],
E [-1.8538e+02, -2.3162e+02, 6.8242e+00],...
E y: array([[[[7.168, 7.867, 8.32 ],
E [8.664, 8.14 , 8.53 ],
E [7.395, 8.66 , 7.957],...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride1x1] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529970>
act_shape = (1, 14, 6, 3), wgt_shape = (3, 3, 3, 64), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f36a54f2b50>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 3.472e+04, 5.371e+00, 9.117e+00, ..., 5.352e+00,
nan, 3.123e+00],
[ 5.285e+00...1],
[-2.122e+03, nan, -1.761e+03, ..., -2.266e+03,
6.619e+01, 2.133e+04]]]], dtype=float16)
desired = array([[[[4.72 , 5.363, 4.85 , ..., 5.18 , 4.21 , 4.65 ],
[5.273, 5.85 , 5.812, ..., 5.605, 5.723, 4.42 ],
... , 6.69 , 7.176, ..., 7.71 , 6.74 , 6.73 ],
[6.742, 6.984, 6.973, ..., 7.95 , 7.176, 6.07 ]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 3.472e+04, 5.371e+00, 9.117e+00, ..., 5.352e+00,
E nan, 3.123e+00],
E [ 5.285e+00, 5.852e+00, 4.934e+00, ..., 5.598e+00,...
E y: array([[[[4.72 , 5.363, 4.85 , ..., 5.18 , 4.21 , 4.65 ],
E [5.273, 5.85 , 5.812, ..., 5.605, 5.723, 4.42 ],
E [7.03 , 6.375, 7.207, ..., 8.38 , 7.27 , 6.047],...
python/tvm/testing/utils.py:120: AssertionError
__
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride1x1] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529820>
act_shape = (1, 14, 6, 3), wgt_shape = (5, 5, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f2350c10>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[-2.0162e+02, 1.1550e+03, 7.1328e+00],
[ 2.0672e+01, 1.6859e+01, 1.9141e+01]],
[[ 2.629...
[[ 7.0812e+01, 1.9766e+01, 9.6484e+00],
[ 1.0544e+02, 1.9109e+01, 2.1609e+01]]]], dtype=float16)
desired = array([[[[20.16, 17.67, 19.08],
[20.72, 16.86, 19.16]],
[[20.16, 19.31, 20.36],
[19.22, 18....3],
[21.86, 20. , 20.84]],
[[21.25, 19.58, 22.56],
[21.9 , 18.66, 21.61]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[-2.0162e+02, 1.1550e+03, 7.1328e+00],
E [ 2.0672e+01, 1.6859e+01, 1.9141e+01]],
E ...
E y: array([[[[20.16, 17.67, 19.08],
E [20.72, 16.86, 19.16]],
E ...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529790>
act_shape = (1, 8, 8, 3), wgt_shape = (2, 2, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f2335370>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 8.188e+01, 3.518e+00, 2.768e+00],
[-2.091e+04, 3.684e+00, 2.412e+00],
[ 3.502e+00, -1... nan],
[ 2.370e+01, 3.006e+00, 2.691e+00],
[ nan, -2.922e+03, 2.072e+00]]]], dtype=float16)
desired = array([[[[3.055, 3.516, 2.766],
[2.6 , 3.584, 2.535],
[2.943, 3.822, 2.307],
[2.996, 3.861...29 ],
[3.002, 3.396, 2.72 ],
[2.428, 3.006, 2.377],
[3.203, 3.4 , 2.07 ]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 8.188e+01, 3.518e+00, 2.768e+00],
E [-2.091e+04, 3.684e+00, 2.412e+00],
E [ 3.502e+00, -1.340e+04, 2.309e+00],...
E y: array([[[[3.055, 3.516, 2.766],
E [2.6 , 3.584, 2.535],
E [2.943, 3.822, 2.307],...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x64-hwio3x3x64x3-stride1x1] _
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529760>
act_shape = (1, 14, 6, 64), wgt_shape = (3, 3, 64, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f36a5547a30>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 3.9180e+03, -1.2216e+04, 2.9425e+02],
[ 1.5838e+02, 1.5925e+02, 1.4862e+02],
[-3.9776e...,
[ 1.2731e+02, 1.4712e+02, 1.3975e+02],
[ 1.5062e+02, 1.4712e+02, nan]]]], dtype=float16)
desired = array([[[[148.1, 149.8, 148.5],
[158.6, 159.6, 149.2],
[151.2, 151. , 146.6],
[149.9, 148.2...4.5],
[148. , 144.2, 142.5],
[143.1, 147.1, 140.8],
[150.6, 147.4, 147.6]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 3.9180e+03, -1.2216e+04, 2.9425e+02],
E [ 1.5838e+02, 1.5925e+02, 1.4862e+02],
E [-3.9776e+04, 1.5088e+02, 1.4625e+02],...
E y: array([[[[148.1, 149.8, 148.5],
E [158.6, 159.6, 149.2],
E [151.2, 151. , 146.6],...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x4x4x40-hwio3x3x40x3-stride1x1] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529640>
act_shape = (1, 4, 4, 40), wgt_shape = (3, 3, 40, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f223cdc0>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 91.4 , 97.06, -4108. ],
[64160. , 101.94, -2114. ]],
[[ 91.06, 562.5 , 1923. ],
[ 89.56, 7336. , 96.2 ]]]], dtype=float16)
desired = array([[[[ 91. , 97. , 96.44],
[ 90.56, 101.56, 98.9 ]],
[[ 90.9 , 96.6 , 94.7 ],
[ 89.44, 91.75, 93.44]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E Mismatched elements: 6 / 12 (50%)
E Max absolute difference: 64064.
E Max relative difference: 707.5
E x: array([[[[ 91.4 , 97.06, -4108. ],
E [64160. , 101.94, -2114. ]],
E ...
E y: array([[[[ 91. , 97. , 96.44],
E [ 90.56, 101.56, 98.9 ]],
E ...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x4x4x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a55295b0>
act_shape = (1, 4, 4, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f22f8730>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 6.941e+00, 6.070e+00, 1.399e+02],
[ 6.629e+00, -5.740e+03, 7.556e+01]],
[[ 7.383e+00, -1.960e+04, 7.848e+00],
[ 7.543e+00, nan, 8.086e+00]]]], dtype=float16)
desired = array([[[[6.88 , 5.812, 6.98 ],
[6.62 , 6.93 , 7.254]],
[[6.984, 5.895, 7.855],
[7.55 , 5.883, 8.086]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 6.941e+00, 6.070e+00, 1.399e+02],
E [ 6.629e+00, -5.740e+03, 7.556e+01]],
E ...
E y: array([[[[6.88 , 5.812, 6.98 ],
E [6.62 , 6.93 , 7.254]],
E ...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x5x5x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529520>
act_shape = (1, 5, 5, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f219e970>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.945e+00, 8.492e+00, 8.461e+00],
[-2.592e+02, 9.523e+00, -3.584e+01],
[ 6.072e+03, 9...2e+03],
[ 6.387e+00, -1.614e+02, 7.223e+00],
[ 5.980e+00, 6.789e+00, 6.754e+00]]]], dtype=float16)
desired = array([[[[7.863, 8.49 , 8.49 ],
[7.58 , 8.63 , 7.652],
[8.05 , 9.9 , 8.73 ]],
[[7.08 , 8.8... ]],
[[6.65 , 7.83 , 8.305],
[6.4 , 6.375, 7.223],
[5.992, 6.89 , 6.75 ]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 7.945e+00, 8.492e+00, 8.461e+00],
E [-2.592e+02, 9.523e+00, -3.584e+01],
E [ 6.072e+03, 9.016e+00, 9.523e+00]],...
E y: array([[[[7.863, 8.49 , 8.49 ],
E [7.58 , 8.63 , 7.652],
E [8.05 , 9.9 , 8.73 ]],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x6x6x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5529460>
act_shape = (1, 6, 6, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f2358940>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.168e+00, 6.574e+00, 5.840e+00],
[-9.477e+00, 5.971e+04, -3.816e+01],
[ 6.040e+03, 4...4e+04],
[ 3.758e+02, 7.660e+00, -2.836e+01],
[ 1.526e+01, 1.173e+01, 5.914e+00]]]], dtype=float16)
desired = array([[[[6.51 , 6.574, 5.875],
[6.707, 6.727, 5.37 ],
[5.887, 5.01 , 4.33 ],
[5.277, 4.715...52 ],
[6.15 , 6.652, 5.78 ],
[7.715, 7.65 , 6.52 ],
[7.195, 7.45 , 5.906]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 7.168e+00, 6.574e+00, 5.840e+00],
E [-9.477e+00, 5.971e+04, -3.816e+01],
E [ 6.040e+03, 4.137e+00, 5.121e+00],...
E y: array([[[[6.51 , 6.574, 5.875],
E [6.707, 6.727, 5.37 ],
E [5.887, 5.01 , 4.33 ],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x7x7x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a55293d0>
act_shape = (1, 7, 7, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f234edf0>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 8.4141e+00, 6.3398e+00, 6.5391e+00],
[ 2.3422e+01, nan, -3.4812e+01],
[ 6.0400e...,
[ 5.3867e+00, 5.3320e+00, 6.8672e+00],
[-3.1088e+04, 2.0453e+01, 6.9805e+00]]]], dtype=float16)
desired = array([[[[7.11 , 6.34 , 6.58 ],
[7.25 , 6.254, 6.76 ],
[6.066, 5.797, 7.14 ],
[6.93 , 5.883...953],
[7.754, 5.81 , 6.723],
[5.387, 5.324, 6.863],
[7.195, 5.066, 6.992]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 8.4141e+00, 6.3398e+00, 6.5391e+00],
E [ 2.3422e+01, nan, -3.4812e+01],
E [ 6.0400e+03, 4.9141e+00, 5.2256e+04],...
E y: array([[[[7.11 , 6.34 , 6.58 ],
E [7.25 , 6.254, 6.76 ],
E [6.066, 5.797, 7.14 ],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x8x3-hwio3x3x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5547040>
act_shape = (1, 8, 8, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37e964a190>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.9180e+00, 7.2930e+00, 6.7383e+00],
[ 2.3094e+01, 7.1602e+00, -3.3938e+01],
[ 8.0720e...,
[ 6.9805e+00, -1.7094e+01, 1.9053e+00],
[ 6.6641e+00, 1.1898e+01, 1.0305e+01]]]], dtype=float16)
desired = array([[[[6.598, 7.297, 6.78 ],
[6.46 , 6.246, 7.78 ],
[5.73 , 5.46 , 5.758],
[6.926, 6.883...805],
[7.14 , 7.105, 7.36 ],
[6.99 , 7.254, 8.375],
[6.656, 6.66 , 7. ]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 7.9180e+00, 7.2930e+00, 6.7383e+00],
E [ 2.3094e+01, 7.1602e+00, -3.3938e+01],
E [ 8.0720e+03, 4.5703e+00, 6.5508e+00],...
E y: array([[[[6.598, 7.297, 6.78 ],
E [6.46 , 6.246, 7.78 ],
E [5.73 , 5.46 , 5.758],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x8x3-hwio5x5x3x3-stride1x1]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a55470d0>
act_shape = (1, 8, 8, 3), wgt_shape = (5, 5, 3, 3), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f223c2b0>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[-8.5050e+02, 3.2464e+04, 2.3184e+04],
[ 1.7266e+01, 1.8141e+01, 1.9016e+01],
[ 1.6781e...,
[ 2.5449e+00, 1.7031e+01, 1.7156e+01],
[ 1.7953e+01, 1.5945e+01, nan]]]], dtype=float16)
desired = array([[[[18.6 , 18.11, 18.31],
[17.28, 18.16, 17.9 ],
[16.78, 16.66, 17. ],
[18.12, 16.69....45],
[17.03, 17.56, 17.66],
[18. , 17. , 17.55],
[17.97, 15.97, 17.73]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[-8.5050e+02, 3.2464e+04, 2.3184e+04],
E [ 1.7266e+01, 1.8141e+01, 1.9016e+01],
E [ 1.6781e+01, -7.4200e+03, 1.7641e+01],...
E y: array([[[[18.6 , 18.11, 18.31],
E [17.28, 18.16, 17.9 ],
E [16.78, 16.66, 17. ],...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x8x64-hwio2x2x64x64-stride1x1] _
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5547160>
act_shape = (1, 8, 8, 64), wgt_shape = (2, 2, 64, 64), inp_stride = (1, 1)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37e95ca490>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 6.1344e+01, 6.7375e+01, -6.1560e+03, ..., 6.7875e+01,
6.8312e+01, 6.4938e+01],
[ 3.2... [-3.4020e+03, 6.7812e+01, 6.1219e+01, ...,
6.3219e+01,
8.4938e+01, 6.4062e+01]]]], dtype=float16)
desired = array([[[[60.47, 67.25, 62.1 , ..., 66.25, 68.25, 66.5 ],
[61.94, 66.56, 62.78, ..., 66.06, 65.9 , 63.94],
... , 67.6 , 60.12, ..., 65.1 , 66.6 , 63.38],
[62.03, 67.7 , 61.28, ..., 63.2 , 65.06, 64.25]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 6.1344e+01, 6.7375e+01, -6.1560e+03, ...,
6.7875e+01,
E 6.8312e+01, 6.4938e+01],
E [ 3.2096e+04, 6.6688e+01, -2.0500e+03, ..., 6.6188e+01,...
E y: array([[[[60.47, 67.25, 62.1 , ..., 66.25, 68.25, 66.5 ],
E [61.94, 66.56, 62.78, ..., 66.06, 65.9 , 63.94],
E [65.44, 71. , 62.22, ..., 67.2 , 70.56, 66.8 ],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride2x2]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a55471f0>
act_shape = (1, 8, 4, 3), wgt_shape = (3, 3, 3, 3), inp_stride = (2, 2)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37e95c2550>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 7.527e+00, 2.549e+04, 1.552e+02]],
[[ 7.996e+00, -1.960e+04, 6.648e+00]],
[[-2.634e+03, 1.578e+03, -5.585e+02]]]], dtype=float16)
desired = array([[[[7.594, 7.38 , 5.793]],
[[7.617, 7.445, 6.645]],
[[7.63 , 8.266, 6.027]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E Mismatched elements: 6 / 9 (66.7%)
E Max absolute difference: 25488.
E Max relative difference: 3454.
E x: array([[[[ 7.527e+00, 2.549e+04, 1.552e+02]],
E
E [[ 7.996e+00, -1.960e+04, 6.648e+00]],...
E y: array([[[[7.594, 7.38 , 5.793]],
E
E [[7.617, 7.445, 6.645]],...
python/tvm/testing/utils.py:120: AssertionError
_
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride2x2] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5547280>
act_shape = (1, 14, 6, 3), wgt_shape = (3, 3, 3, 64), inp_stride = (2, 2)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f36a54c8970>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 3.4752e+04, 5.7773e+00, 1.2820e+01, 5.6367e+00,
7.7148e+00, 2.5260e+03, 4.3960e+03, 9.1484...0e+03,
7.3789e+00, 7.8477e+00,
1.0350e+03, 6.3008e+00, 7.9648e+00, -4.3242e+00]]]],
dtype=float16)
desired = array([[[[ 9.805, 5.79 , 6.57 , 5.492, 7.74 , 7.57 , 6.43 ,
9.125, 6.582, 7.223, 8.9 , 6.92 , 7... 8.16 , 7.34 ,
7.984, 5.848, 7.383, 7.848, 6.15 , 6.277, 5.125,
7.754]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 3.4752e+04, 5.7773e+00, 1.2820e+01, 5.6367e+00,
E 7.7148e+00, 2.5260e+03, 4.3960e+03, 9.1484e+00,
E 6.5898e+00, 7.2305e+00, 8.6719e+00, 2.2928e+04,...
E y: array([[[[ 9.805, 5.79 , 6.57 , 5.492, 7.74 , 7.57 , 6.43 ,
E 9.125, 6.582, 7.223, 8.9 , 6.92 , 7.695, 8.125,
E 7.047, 7.996, 7.742, 6.64 , 7.438, 5.66 , 8.37 ,...
python/tvm/testing/utils.py:120: AssertionError
__
TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride2x2] __
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a5547310>
act_shape = (1, 14, 6, 3), wgt_shape = (5, 5, 3, 3), inp_stride = (2, 2)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37f23524f0>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[-2.0075e+02, 1.1590e+03, 2.5020e+00]],
[[ 2.5453e+01, nan, 1.9094e+01]],
[[ 1.5... [[ 7.6938e+01, 1.8109e+01, 1.5680e+01]],
[[ 1.7031e+01, 1.9328e+01, 1.8370e+03]]]], dtype=float16)
desired = array([[[[17.23, 19.39, 18.44]],
[[19.38, 20.77, 19.11]],
[[16.98, 18.27, 16.89]],
[[16.48, 18.12, 15.69]],
[[17.02, 19.3 , 18.06]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[-2.0075e+02, 1.1590e+03, 2.5020e+00]],
E
E [[ 2.5453e+01, nan, 1.9094e+01]],...
E y: array([[[[17.23, 19.39, 18.44]],
E
E [[19.38, 20.77, 19.11]],...
python/tvm/testing/utils.py:120: AssertionError
__ TestConv2dIntrin.test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride2x2]
___
self = <test_hexagon.topi.test_conv2d_fp16_intrin.TestConv2dIntrin object at
0x7f36a55473a0>
act_shape = (1, 8, 8, 3), wgt_shape = (2, 2, 3, 3), inp_stride = (2, 2)
inp_offset = (0, 0)
hexagon_session = <tvm.contrib.hexagon.session.Session object at
0x7f37e967b070>
@tvm.testing.requires_hexagon
def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset,
hexagon_session):
"""Test conv2d intrinsic implementation"""
assert act_shape[3] == wgt_shape[2]
# Currently, input offset does not affect the output shape
def get_out_shape(ash, wsh, inp_stride):
assert ash[3] == wsh[2]
osh = (
ash[0],
(ash[1] - wsh[0]) // inp_stride[0] + 1,
(ash[2] - wsh[1]) // inp_stride[1] + 1,
wsh[3],
)
assert tvm.tir.all([x > 0 for x in osh])
return osh
act = np.random.rand(*act_shape).astype("float16")
wgt = np.random.rand(*wgt_shape).astype("float16")
module = build_conv2d(get_hexagon_target("v68"))
mod = hexagon_session.load_module(module)
output = tvm.nd.array(
np.zeros(get_out_shape(act_shape, wgt_shape, inp_stride),
dtype="float16"),
device=hexagon_session.device,
)
mod(
tvm.nd.array(act, device=hexagon_session.device),
tvm.nd.array(wgt, device=hexagon_session.device),
inp_offset[0], # off_t
inp_offset[1], # off_l
inp_stride[0], # stride_height
inp_stride[1], # stride_width
output,
)
out = output.numpy()
# Generate reference output and compare:
ref_out = conv2d_nhwc_python(
act.astype("float32"), wgt.astype("float32"), stride=inp_stride,
padding="VALID"
).astype("float16")
> tvm.testing.assert_allclose(out, ref_out, rtol=5e-2, atol=5e-2)
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:243:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _
actual = array([[[[ 1.2730e+03, 2.9004e+00, 3.0898e+00],
[-8.8375e+01, 2.9941e+00, 2.8320e+00],
[ 3.4473e...,
[ 2.4551e+00, 2.9156e+01, -2.3975e-01],
[-1.5162e+02, 2.5762e+00, 2.9590e+00]]]], dtype=float16)
desired = array([[[[2.393, 2.896, 3.092],
[2.686, 2.898, 2.957],
[2.916, 2.799, 3.514],
[3.285, 3.65 ...086],
[2.457, 2.709, 2.902],
[2.455, 2.64 , 2.27 ],
[2.844, 2.576, 2.965]]]], dtype=float16)
rtol = 0.05, atol = 0.05
def assert_allclose(actual, desired, rtol=1e-7, atol=1e-7):
"""Version of np.testing.assert_allclose with `atol` and `rtol`
fields set
in reasonable defaults.
Arguments `actual` and `desired` are not interchangeable, since the
function
compares the `abs(actual-desired)` with `atol+rtol*abs(desired)`.
Since we
often allow `desired` to be close to zero, we generally want
non-zero `atol`.
"""
actual = np.asanyarray(actual)
desired = np.asanyarray(desired)
np.testing.assert_allclose(actual.shape, desired.shape)
> np.testing.assert_allclose(actual, desired, rtol=rtol, atol=atol,
verbose=True)
E AssertionError:
E Not equal to tolerance rtol=0.05, atol=0.05
E
E x and y nan location mismatch:
E x: array([[[[ 1.2730e+03, 2.9004e+00, 3.0898e+00],
E [-8.8375e+01, 2.9941e+00, 2.8320e+00],
E [ 3.4473e+00, -1.3368e+04, 3.5176e+00],...
E y: array([[[[2.393, 2.896, 3.092],
E [2.686, 2.898, 2.957],
E [2.916, 2.799, 3.514],...
python/tvm/testing/utils.py:120: AssertionError
=========================== short test summary info
============================
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x10x14x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x64-hwio3x3x64x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x4x4x40-hwio3x3x40x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x4x4x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x5x5x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x6x6x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x7x7x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio3x3x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio5x5x3x3-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x64-hwio2x2x64x64-stride1x1]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride2x2]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride2x2]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride2x2]
FAILED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride2x2]
============================= 19 failed in 19.71s
==============================
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
