Skip to content

Conversation

@csullivan
Copy link
Contributor

@csullivan csullivan commented Nov 10, 2022

This PR shows how to utilize the sch.rolling_buffer schedule primitive to avoid copying repeated rows when loading to fast memory. The example below is like conv2d with filter=[2, 1] and unit weights.

Before applying a rolling buffer to the height axis:

# from tvm.script import tir as T
@tvm.script.ir_module
class Module:
    @T.prim_func
    def main(A: T.Buffer[(6, 32), "float16"], B: T.Buffer[(5, 32), "float16"]):
        # function attr dict
        T.func_attr({"global_symbol": "main"})
        # body
        # with T.block("root")
        A_global_vtcm = T.alloc_buffer([6, 32], dtype="float16", scope="global.vtcm")
        for output_row in T.serial(5):
            for ax0, ax1 in T.grid(2, 32):
                with T.block("A_global.vtcm"):
                    v0 = T.axis.spatial(6, output_row + ax0)
                    v1 = T.axis.spatial(32, ax1)
                    T.reads(A[v0, v1])
                    T.writes(A_global_vtcm[v0, v1])
                    A_global_vtcm[v0, v1] = A[v0, v1]
            for i, rh in T.grid(32, 2):
                with T.block("compute"):
                    vrow, vi, vrh = T.axis.remap("SSR", [output_row, i, rh])
                    T.reads(A_global_vtcm[vrow + vrh, vi])
                    T.writes(B[vrow, vi])
                    with T.init():
                        B[vrow, vi] = T.cast(0, "float16")
                    B[vrow, vi] = B[vrow, vi] + A_global_vtcm[vrow + vrh, vi]

After applying sch.rolling_buffer to the cache_read block along the height axis:

# from tvm.script import tir as T
@tvm.script.ir_module
class Module:
    @T.prim_func
    def main(A: T.Buffer[(6, 32), "float16"], B: T.Buffer[(5, 32), "float16"]):
        # function attr dict
        T.func_attr({"global_symbol": "main"})
        # body
        # with T.block("root")
        A_global_vtcm = T.alloc_buffer([2, 32], dtype="float16", scope="global.vtcm")
        for output_row in T.serial(5):
            for ax0, ax1 in T.grid(2, 32):
                with T.block("A_global.vtcm"):
                    T.where(output_row < 1 or 1 <= ax0)
                    v0 = T.axis.opaque(6, output_row + ax0)
                    v1 = T.axis.spatial(32, ax1)
                    T.reads(A[v0, v1])
                    T.writes(A_global_vtcm[v0 % 2, v1])
                    A_global_vtcm[v0 % 2, v1] = A[v0, v1]
            for i, rh in T.grid(32, 2):
                with T.block("compute"):
                    vrow = T.axis.opaque(5, output_row)
                    vi, vrh = T.axis.remap("SR", [i, rh])
                    T.reads(A_global_vtcm[(vrow + vrh) % 2, vi])
                    T.writes(B[vrow, vi])
                    with T.init():
                        B[vrow, vi] = T.cast(0, "float16")
                    B[vrow, vi] = B[vrow, vi] + A_global_vtcm[(vrow + vrh) % 2, vi]

@tvm-bot
Copy link
Collaborator

tvm-bot commented Nov 10, 2022

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

Generated by tvm-bot

to avoid repeated loads on input rows needed by vertically
output rows.
@csullivan csullivan marked this pull request as ready for review November 10, 2022 04:55
@csullivan
Copy link
Contributor Author

cc @kparzysz-quic

Many thanks (!!) to @mbaret @NicolaLancellotti for the initial contribution of rolling buffer [1, 2], and to @liangW-intellif and @wrongtest-intellif for their recent translation of rolling buffer to a TIR schedule primitive [3], as well as all others involved in this work. It is a fantastic example of the benefit of a strong open source community!

Copy link
Member

@mehrdadh mehrdadh left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM! just small changes

return conv2d_2x1_filter


def test_rolling_buffer_conv2d_2x1(hexagon_session):
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add @tvm.testing.requires_hexagon

a = tvm.nd.array(np.ones((6, 32), dtype=dtype), device=hexagon_session.device)
b = tvm.nd.array(np.zeros((5, 32), dtype=dtype), device=hexagon_session.device)
mod(a, b)
tvm.testing.assert_allclose(b.numpy(), np.full((5, 32), 2, dtype=dtype), atol=1e-3, rtol=1e-3)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add this to the end:

if __name__ == "__main__":
    tvm.testing.main()

@csullivan
Copy link
Contributor Author

@tvm-bot rerun

2 similar comments
@masahi
Copy link
Member

masahi commented Nov 28, 2022

@tvm-bot rerun

@kparzysz-quic
Copy link
Contributor

@tvm-bot rerun

@github-actions
Copy link
Contributor

github-actions bot commented Jan 9, 2023

Failed to re-run CI in https://github.com/apache/tvm/actions/runs/3876284649

Details
Traceback (most recent call last):
  File "ci/scripts/github/github_tvmbot.py", line 593, in comment_failure
    raise item
  File "ci/scripts/github/github_tvmbot.py", line 699, in run
    pr.rerun_jenkins_ci()
  File "ci/scripts/github/github_tvmbot.py", line 552, in rerun_jenkins_ci
    post(url, auth=("tvm-bot", TVM_BOT_JENKINS_TOKEN))
  File "/home/runner/work/tvm/tvm/ci/scripts/jenkins/git_utils.py", line 53, in post
    with request.urlopen(req, data) as response:
  File "/usr/lib/python3.8/urllib/request.py", line 222, in urlopen
    return opener.open(url, data, timeout)
  File "/usr/lib/python3.8/urllib/request.py", line 531, in open
    response = meth(req, response)
  File "/usr/lib/python3.8/urllib/request.py", line 640, in http_response
    response = self.parent.error(
  File "/usr/lib/python3.8/urllib/request.py", line 569, in error
    return self._call_chain(*args)
  File "/usr/lib/python3.8/urllib/request.py", line 502, in _call_chain
    result = func(*args)
  File "/usr/lib/python3.8/urllib/request.py", line 649, in http_error_default
    raise HTTPError(req.full_url, code, msg, hdrs, fp)
urllib.error.HTTPError: HTTP Error 500: Server Error

with response


  
  <!DOCTYPE html><html><head resURL="/static/e3b9d568" data-rooturl="" data-resurl="/static/e3b9d568" data-extensions-available="true" data-unit-test="false" data-imagesurl="/static/e3b9d568/images" data-crumb-header="Jenkins-Crumb" data-crumb-value="feab3f95df923bbb057f5a8022c738110a9136ff18248f08e656785c09abf5dea00d43488d85fd2e6238225d794f8499d78b0df72517f4d6a7ebf5087f1ae73f">
    
    

    <title>Jenkins [Jenkins]</title><link rel="stylesheet" href="/static/e3b9d568/jsbundles/styles.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/css/responsive-grid.css" type="text/css"><link rel="shortcut icon" href="/static/e3b9d568/favicon.ico" type="image/vnd.microsoft.icon"><script src="/static/e3b9d568/scripts/prototype.js" type="text/javascript"></script><script src="/static/e3b9d568/scripts/behavior.js" type="text/javascript"></script><script src='/adjuncts/e3b9d568/org/kohsuke/stapler/bind.js' type='text/javascript'></script><script src="/static/e3b9d568/scripts/yui/yahoo/yahoo-min.js"></script><script src="/static/e3b9d568/scripts/yui/dom/dom-min.js"></script><script src="/static/e3b9d568/scripts/yui/event/event-min.js"></script><script src="/static/e3b9d568/scripts/yui/animation/animation-min.js"></script><script src="/static/e3b9d568/scripts/yui/dragdrop/dragdrop-min.js"></script><script src="/static/e3b9d568/scripts/yui/container/container-min.js"></script><script src="/static/e3b9d568/scripts/yui/connection/connection-min.js"></script><script src="/static/e3b9d568/scripts/yui/datasource/datasource-min.js"></script><script src="/static/e3b9d568/scripts/yui/autocomplete/autocomplete-min.js"></script><script src="/static/e3b9d568/scripts/yui/menu/menu-min.js"></script><script src="/static/e3b9d568/scripts/yui/element/element-min.js"></script><script src="/static/e3b9d568/scripts/yui/button/button-min.js"></script><script src="/static/e3b9d568/scripts/yui/storage/storage-min.js"></script><script src="/static/e3b9d568/scripts/hudson-behavior.js" type="text/javascript"></script><script src="/static/e3b9d568/scripts/sortable.js" type="text/javascript"></script><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/container/assets/container.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/container/assets/skins/sam/container.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/menu/assets/skins/sam/menu.css" type="text/css"><link rel="search" href="/opensearch.xml" type="application/opensearchdescription+xml" title="Jenkins"><meta name="ROBOTS" content="INDEX,NOFOLLOW"><meta name="viewport" content="width=device-width, initial-scale=1"><script src="/static/e3b9d568/jsbundles/vendors.js" type="text/javascript"></script><script src="/static/e3b9d568/jsbundles/page-init.js" type="text/javascript"></script><script src="/static/e3b9d568/jsbundles/sortable-drag-drop.js" type="text/javascript"></script></head><body data-model-type="hudson.model.Hudson" id="jenkins" class="yui-skin-sam one-column jenkins-2.361.2" data-version="2.361.2"><a href="#skip2content" class="skiplink">Skip to content</a><header id="page-header" class="page-header"><div class="page-header__brand"><div class="logo"><a id="jenkins-home-link" href="/"><img src="/static/e3b9d568/images/svgs/logo.svg" alt="[Jenkins]" id="jenkins-head-icon"><img src="/static/e3b9d568/images/title.svg" alt="Jenkins" width="139" id="jenkins-name-icon" height="34"></a></div><a href="/" class="page-header__brand-link"><img src="/static/e3b9d568/images/svgs/logo.svg" alt="[Jenkins]" class="page-header__brand-image"><span class="page-header__brand-name">Jenkins</span></a></div><div class="searchbox hidden-xs"><form role="search" method="get" name="search" action="/search/" style="position:relative;" class="no-json"><div id="search-box-sizer"></div><div id="searchform"><input role="searchbox" name="q" placeholder="Search" id="search-box" class="main-search__input"><span class="main-search__icon-leading"><svg class="" class="" aria-hidden="true" xmlns="http://www.w3.org/2000/svg" class="" viewBox="0 0 512 512"><title></title><path d="M221.09 64a157.09 157.09 0 10157.09 157.09A157.1 157.1 0 00221.09 64z" fill="none" stroke="currentColor" stroke-miterlimit="10" stroke-width="32"/><path fill="none" stroke="currentColor" stroke-linecap="round" stroke-miterlimit="10" stroke-width="32" d="M338.29 338.29L448 448"/></svg></span><a href="https://www.jenkins.io/redirect/search-box" class="main-search__icon-trailing"><svg class="" class="" aria-hidden="true" xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><path d="M256 40a216 216 0 10216 216A216 216 0 00256 40z" fill="none" stroke="currentColor" stroke-miterlimit="10" stroke-width="38"/><path d="M200 202.29s.84-17.5 19.57-32.57C230.68 160.77 244 158.18 256 158c10.93-.14 20.69 1.67 26.53 4.45 10 4.76 29.47 16.38 29.47 41.09 0 26-17 37.81-36.37 50.8S251 281.43 251 296" fill="none" stroke="currentColor" stroke-linecap="round" stroke-miterlimit="10" stroke-width="38"/><circle cx="250" cy="360" r="25" fill="currentColor"/></svg></a><div id="search-box-completion" data-search-url="/search/"></div><script src='/adjuncts/e3b9d568/jenkins/views/JenkinsHeader/search-box.js' type='text/javascript'></script></div></form></div><div class="login page-header__hyperlinks"><div id="visible-am-insertion" class="page-header__am-wrapper"></div><div id="visible-sec-am-insertion" class="page-header__am-wrapper"></div><a href="/securityRealm/commenceLogin?from=%2Fjob%2Ftvm-arm%2Fjob%2FPR-13340%2FbuildWithParameters"><b>log in</b></a></div></header><script src="/static/e3b9d568/jsbundles/keyboard-shortcuts.js" type="text/javascript"></script><div id="breadcrumbBar"><script src='/adjuncts/e3b9d568/lib/layout/breadcrumbs.js' type='text/javascript'></script><div class="top-sticker noedge"><div class="top-sticker-inner"><div class="jenkins-breadcrumbs"><ul id="breadcrumbs"><li class="item"><a href="/" class="model-link">Dashboard</a></li><li href="/" class="children"></li></ul><div id="breadcrumb-menu-target"></div></div></div></div></div><div id="page-body" class="clear"><div id="main-panel"><a name="skip2content"></a><h1 style="text-align: center"><img src="/static/e3b9d568/images/rage.svg" width="154" height="179"><span style="font-size:50px"> Oops!</span></h1><div id="error-description"><h2 style="text-align: center">A problem occurred while processing the request.</h2><p style="text-align: center">Logging ID=8759cb03-29c2-401f-af83-67c328129544</div></div></div><footer class="page-footer"><div class="container-fluid"><div class="page-footer__flex-row"><div class="page-footer__footer-id-placeholder" id="footer"></div><div class="page-footer__links rest_api hidden-xs"><a href="api/">REST API</a></div><div class="page-footer__links page-footer__links--white jenkins_ver"><a rel="noopener noreferrer" href="https://www.jenkins.io/" target="_blank">Jenkins 2.361.2</a></div></div></div></footer></body></html>

@csullivan
Copy link
Contributor Author

@tvm-bot rerun

@github-actions
Copy link
Contributor

Failed to re-run CI in https://github.com/apache/tvm/actions/runs/4008831418

Details
Traceback (most recent call last):
  File "ci/scripts/github/github_tvmbot.py", line 594, in comment_failure
    raise item
  File "ci/scripts/github/github_tvmbot.py", line 700, in run
    pr.rerun_jenkins_ci()
  File "ci/scripts/github/github_tvmbot.py", line 553, in rerun_jenkins_ci
    post(url, auth=("tvm-bot", TVM_BOT_JENKINS_TOKEN))
  File "/home/runner/work/tvm/tvm/ci/scripts/jenkins/git_utils.py", line 53, in post
    with request.urlopen(req, data) as response:
  File "/usr/lib/python3.8/urllib/request.py", line 222, in urlopen
    return opener.open(url, data, timeout)
  File "/usr/lib/python3.8/urllib/request.py", line 531, in open
    response = meth(req, response)
  File "/usr/lib/python3.8/urllib/request.py", line 640, in http_response
    response = self.parent.error(
  File "/usr/lib/python3.8/urllib/request.py", line 569, in error
    return self._call_chain(*args)
  File "/usr/lib/python3.8/urllib/request.py", line 502, in _call_chain
    result = func(*args)
  File "/usr/lib/python3.8/urllib/request.py", line 649, in http_error_default
    raise HTTPError(req.full_url, code, msg, hdrs, fp)
urllib.error.HTTPError: HTTP Error 500: Server Error

with response


  
  <!DOCTYPE html><html><head resURL="/static/e3b9d568" data-rooturl="" data-resurl="/static/e3b9d568" data-extensions-available="true" data-unit-test="false" data-imagesurl="/static/e3b9d568/images" data-crumb-header="Jenkins-Crumb" data-crumb-value="c78f18bdf9977cf79191d32d8f2e5149434a1dd1b6b42998f3edf086ecb2a21312e7f72a3e6e13e6357cedbfd05aa98de50cdc83eca9747cd931af89f7ff2997">
    
    

    <title>Jenkins [Jenkins]</title><link rel="stylesheet" href="/static/e3b9d568/jsbundles/styles.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/css/responsive-grid.css" type="text/css"><link rel="shortcut icon" href="/static/e3b9d568/favicon.ico" type="image/vnd.microsoft.icon"><script src="/static/e3b9d568/scripts/prototype.js" type="text/javascript"></script><script src="/static/e3b9d568/scripts/behavior.js" type="text/javascript"></script><script src='/adjuncts/e3b9d568/org/kohsuke/stapler/bind.js' type='text/javascript'></script><script src="/static/e3b9d568/scripts/yui/yahoo/yahoo-min.js"></script><script src="/static/e3b9d568/scripts/yui/dom/dom-min.js"></script><script src="/static/e3b9d568/scripts/yui/event/event-min.js"></script><script src="/static/e3b9d568/scripts/yui/animation/animation-min.js"></script><script src="/static/e3b9d568/scripts/yui/dragdrop/dragdrop-min.js"></script><script src="/static/e3b9d568/scripts/yui/container/container-min.js"></script><script src="/static/e3b9d568/scripts/yui/connection/connection-min.js"></script><script src="/static/e3b9d568/scripts/yui/datasource/datasource-min.js"></script><script src="/static/e3b9d568/scripts/yui/autocomplete/autocomplete-min.js"></script><script src="/static/e3b9d568/scripts/yui/menu/menu-min.js"></script><script src="/static/e3b9d568/scripts/yui/element/element-min.js"></script><script src="/static/e3b9d568/scripts/yui/button/button-min.js"></script><script src="/static/e3b9d568/scripts/yui/storage/storage-min.js"></script><script src="/static/e3b9d568/scripts/hudson-behavior.js" type="text/javascript"></script><script src="/static/e3b9d568/scripts/sortable.js" type="text/javascript"></script><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/container/assets/container.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/container/assets/skins/sam/container.css" type="text/css"><link rel="stylesheet" href="/static/e3b9d568/scripts/yui/menu/assets/skins/sam/menu.css" type="text/css"><link rel="search" href="/opensearch.xml" type="application/opensearchdescription+xml" title="Jenkins"><meta name="ROBOTS" content="INDEX,NOFOLLOW"><meta name="viewport" content="width=device-width, initial-scale=1"><script src="/static/e3b9d568/jsbundles/vendors.js" type="text/javascript"></script><script src="/static/e3b9d568/jsbundles/page-init.js" type="text/javascript"></script><script src="/static/e3b9d568/jsbundles/sortable-drag-drop.js" type="text/javascript"></script></head><body data-model-type="hudson.model.Hudson" id="jenkins" class="yui-skin-sam one-column jenkins-2.361.2" data-version="2.361.2"><a href="#skip2content" class="skiplink">Skip to content</a><header id="page-header" class="page-header"><div class="page-header__brand"><div class="logo"><a id="jenkins-home-link" href="/"><img src="/static/e3b9d568/images/svgs/logo.svg" alt="[Jenkins]" id="jenkins-head-icon"><img src="/static/e3b9d568/images/title.svg" alt="Jenkins" width="139" id="jenkins-name-icon" height="34"></a></div><a href="/" class="page-header__brand-link"><img src="/static/e3b9d568/images/svgs/logo.svg" alt="[Jenkins]" class="page-header__brand-image"><span class="page-header__brand-name">Jenkins</span></a></div><div class="searchbox hidden-xs"><form role="search" method="get" name="search" action="/search/" style="position:relative;" class="no-json"><div id="search-box-sizer"></div><div id="searchform"><input role="searchbox" name="q" placeholder="Search" id="search-box" class="main-search__input"><span class="main-search__icon-leading"><svg class="" class="" aria-hidden="true" xmlns="http://www.w3.org/2000/svg" class="" viewBox="0 0 512 512"><title></title><path d="M221.09 64a157.09 157.09 0 10157.09 157.09A157.1 157.1 0 00221.09 64z" fill="none" stroke="currentColor" stroke-miterlimit="10" stroke-width="32"/><path fill="none" stroke="currentColor" stroke-linecap="round" stroke-miterlimit="10" stroke-width="32" d="M338.29 338.29L448 448"/></svg></span><a href="https://www.jenkins.io/redirect/search-box" class="main-search__icon-trailing"><svg class="" class="" aria-hidden="true" xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><path d="M256 40a216 216 0 10216 216A216 216 0 00256 40z" fill="none" stroke="currentColor" stroke-miterlimit="10" stroke-width="38"/><path d="M200 202.29s.84-17.5 19.57-32.57C230.68 160.77 244 158.18 256 158c10.93-.14 20.69 1.67 26.53 4.45 10 4.76 29.47 16.38 29.47 41.09 0 26-17 37.81-36.37 50.8S251 281.43 251 296" fill="none" stroke="currentColor" stroke-linecap="round" stroke-miterlimit="10" stroke-width="38"/><circle cx="250" cy="360" r="25" fill="currentColor"/></svg></a><div id="search-box-completion" data-search-url="/search/"></div><script src='/adjuncts/e3b9d568/jenkins/views/JenkinsHeader/search-box.js' type='text/javascript'></script></div></form></div><div class="login page-header__hyperlinks"><div id="visible-am-insertion" class="page-header__am-wrapper"></div><div id="visible-sec-am-insertion" class="page-header__am-wrapper"></div><a href="/securityRealm/commenceLogin?from=%2Fjob%2Ftvm-arm%2Fjob%2FPR-13340%2FbuildWithParameters"><b>log in</b></a></div></header><script src="/static/e3b9d568/jsbundles/keyboard-shortcuts.js" type="text/javascript"></script><div id="breadcrumbBar"><script src='/adjuncts/e3b9d568/lib/layout/breadcrumbs.js' type='text/javascript'></script><div class="top-sticker noedge"><div class="top-sticker-inner"><div class="jenkins-breadcrumbs"><ul id="breadcrumbs"><li class="item"><a href="/" class="model-link">Dashboard</a></li><li href="/" class="children"></li></ul><div id="breadcrumb-menu-target"></div></div></div></div></div><div id="page-body" class="clear"><div id="main-panel"><a name="skip2content"></a><h1 style="text-align: center"><img src="/static/e3b9d568/images/rage.svg" width="154" height="179"><span style="font-size:50px"> Oops!</span></h1><div id="error-description"><h2 style="text-align: center">A problem occurred while processing the request.</h2><p style="text-align: center">Logging ID=dc135f49-40b0-4791-abe1-5e5a8cf06b51</div></div></div><footer class="page-footer"><div class="container-fluid"><div class="page-footer__flex-row"><div class="page-footer__footer-id-placeholder" id="footer"></div><div class="page-footer__links rest_api hidden-xs"><a href="api/">REST API</a></div><div class="page-footer__links page-footer__links--white jenkins_ver"><a rel="noopener noreferrer" href="https://www.jenkins.io/" target="_blank">Jenkins 2.361.2</a></div></div></div></footer></body></html>

@csullivan csullivan closed this by deleting the head repository Feb 1, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants