[Hexagon] Fix chunk address table generation in chunkify_hwio_16b by kparzysz-quic · Pull Request #13002 · apache/tvm

kparzysz-quic · 2022-10-06T22:01:12Z

The filter data is contiguous, the iteration is over 16-bit elements, not pointers.

kparzysz-quic · 2022-10-06T22:02:05Z

This should fix the failure on hardware in the conv2d fp16 test.
cc: @janetsc

Edit: apparently it still fails on hardware, so the failing test was left disabled.

janetsc · 2022-10-07T00:42:34Z

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 ==============================

janetsc · 2022-10-07T11:34:08Z

@kparzysz-quic - I’m out today, but I think @supersat did a quick experiment of zeroing memory, which produced closer answers on hardware. Maybe the combination of this fix, and zeroing the buffers will do it.

janetsc

Looks good on the simulator.

supersat · 2022-10-19T19:54:03Z

All tests pass on hardware after zeroing out memory before returning from VTCMAllocation::VTCMAllocation():

platform linux -- Python 3.8.10, pytest-6.2.5, py-1.11.0, pluggy-1.0.0 -- /usr/bin/python3
cachedir: .pytest_cache
rootdir: /home/kkoscher/Projects/tvm
plugins: xdist-2.5.0, forked-1.4.0
collected 19 items                                                                                                                                                                 

tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x10x14x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x64-hwio3x3x64x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x4x4x40-hwio3x3x40x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x4x4x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x5x5x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x6x6x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x7x7x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio3x3x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio5x5x3x3-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x64-hwio2x2x64x64-stride1x1] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x4x3-hwio3x3x3x3-stride2x2] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio3x3x3x64-stride2x2] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x14x6x3-hwio5x5x3x3-stride2x2] PASSED
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py::TestConv2dIntrin::test_conv2d[offset0x0-nhwc1x8x8x3-hwio2x2x3x3-stride2x2] PASSED

=============================================================================== 19 passed in 40.70s ================================================================================```

Properly zeroing memory should likely be done in the conv2d kernel rather than relying on VTCMAllocator to do that every time, but that can be addressed in another PR, so this one LGTM.

…ache#13002) * [Hexagon] Fix chunk address table generation in chunkify_hwio_16b The filter data is contiguous, the iteration is over 16-bit elements, not pointers. * Re-disable the conv2d_fp16 test, since it still fails on hardware

[Hexagon] Fix chunk address table generation in chunkify_hwio_16b

3d5e1dd

The filter data is contiguous, the iteration is over 16-bit elements, not pointers.

Re-disable the conv2d_fp16 test, since it still fails on hardware

3d581f6

janetsc approved these changes Oct 10, 2022

View reviewed changes

areusch added needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it and removed needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it labels Oct 19, 2022

supersat mentioned this pull request Oct 19, 2022

[Bug] [Hexagon] HVX conv2d kernel incorrectly relies on VTCM being zeroed #13144

Closed

masahi approved these changes Oct 21, 2022

View reviewed changes

masahi merged commit 0e21840 into apache:main Oct 21, 2022

kparzysz-quic deleted the fix-chunking branch January 10, 2023 16:30

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Hexagon] Fix chunk address table generation in chunkify_hwio_16b#13002

[Hexagon] Fix chunk address table generation in chunkify_hwio_16b#13002
masahi merged 2 commits intoapache:mainfrom
kparzysz-quic:fix-chunking

kparzysz-quic commented Oct 6, 2022

Uh oh!

kparzysz-quic commented Oct 6, 2022 •

edited

Loading

Uh oh!

janetsc commented Oct 7, 2022

Uh oh!

janetsc commented Oct 7, 2022

Uh oh!

janetsc left a comment

Uh oh!

supersat commented Oct 19, 2022

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

kparzysz-quic commented Oct 6, 2022

Uh oh!

kparzysz-quic commented Oct 6, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

janetsc commented Oct 7, 2022

Uh oh!

janetsc commented Oct 7, 2022

Uh oh!

janetsc left a comment

Choose a reason for hiding this comment

Uh oh!

supersat commented Oct 19, 2022

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

kparzysz-quic commented Oct 6, 2022 •

edited

Loading