def test_alloc_vtcm():
    if not check_prereq_and_setup():
        return
    target = tvm.target.hexagon("v66")

    buf_len = 2048
    A = tvm.te.placeholder((buf_len, ), name="A", dtype="int8")
    B = tvm.te.placeholder((buf_len, ), name="B", dtype="int8")

    A_buf = tvm.te.compute((buf_len, ), lambda *i: A(*i), "A_buf")
    B_buf = tvm.te.compute((buf_len, ), lambda *i: B(*i), "B_buf")
    C = tvm.te.compute((buf_len, ), lambda *i: A_buf(*i) + B_buf(*i), name="C")
    s = tvm.te.create_schedule(C.op)

    # Use VTCM for each buffer.
    s[A_buf].set_scope("local.vtcm")
    s[B_buf].set_scope("local.vtcm")

    config = {"tir.add_lower_pass": hexagon.ir_lower_vtcm_pass()}
    with tvm.transform.PassContext(config=config):
        irmod = tvm.lower(s, [A, B, C], name="alloc_vtcm")

    calls = re.findall("HexagonBackend[A-Za-z]*VTCM", str(irmod["alloc_vtcm"]))
    assert "HexagonBackendAllocateVTCM" in calls
    assert "HexagonBackendFreeVTCM" in calls
Exemplo n.º 2
0
def test_alloc_vtcm():
    if not check_prereq_and_setup():
        return
    target = tvm.target.hexagon('v66')

    buf_len = 2048
    A = tvm.te.placeholder((buf_len, ), name='A', dtype='int8')
    B = tvm.te.placeholder((buf_len, ), name='B', dtype='int8')

    A_buf = tvm.te.compute((buf_len, ), lambda *i: A(*i), 'A_buf')
    B_buf = tvm.te.compute((buf_len, ), lambda *i: B(*i), 'B_buf')
    C = tvm.te.compute((buf_len, ), lambda *i: A_buf(*i) + B_buf(*i), name='C')
    s = tvm.te.create_schedule(C.op)

    # Use VTCM for each buffer.
    s[A_buf].set_scope("local.vtcm")
    s[B_buf].set_scope("local.vtcm")

    config = {'tir.add_lower_pass': hexagon.ir_lower_vtcm_pass()}
    with tvm.transform.PassContext(config=config):
        irmod = tvm.lower(s, [A, B, C], name='alloc_vtcm')

    calls = re.findall('HexagonBackend[A-Za-z]*VTCM', str(irmod['alloc_vtcm']))
    assert 'HexagonBackendAllocateVTCM' in calls
    assert 'HexagonBackendFreeVTCM' in calls