| |
| |
| |
| |
| |
| |
| <!DOCTYPE html> |
| <html class="writer-html5" lang="en" > |
| <head> |
| <meta charset="utf-8"> |
| |
| <meta name="viewport" content="width=device-width, initial-scale=1.0"> |
| |
| <title>Cross Compilation and RPC — tvm 0.17.dev0 documentation</title> |
| |
| |
| |
| <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/css/bootstrap.min.css" integrity="sha384-Gn5384xqQ1aoWXA+058RXPxPg6fy4IWvTNh0E263XmFcJlSAwiGgFAW/dAiS6JXm" crossorigin="anonymous"> |
| <link rel="stylesheet" href="../_static/css/theme.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/pygments.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/css/theme.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/sg_gallery.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/sg_gallery-binder.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/sg_gallery-dataframe.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/sg_gallery-rendered-html.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/pygments.css" type="text/css" /> |
| <link rel="stylesheet" href="../_static/css/tlcpack_theme.css" type="text/css" /> |
| |
| |
| |
| <link rel="shortcut icon" href="../_static/tvm-logo-square.png"/> |
| |
| |
| |
| |
| |
| |
| |
| <script type="text/javascript" id="documentation_options" data-url_root="../" src="../_static/documentation_options.js"></script> |
| <script data-url_root="../" id="documentation_options" src="../_static/documentation_options.js"></script> |
| <script src="../_static/jquery.js"></script> |
| <script src="../_static/underscore.js"></script> |
| <script src="../_static/doctools.js"></script> |
| |
| <script type="text/javascript" src="../_static/js/theme.js"></script> |
| |
| |
| <script type="text/javascript" src="../_static/js/tlcpack_theme.js"></script> |
| <link rel="index" title="Index" href="../genindex.html" /> |
| <link rel="search" title="Search" href="../search.html" /> |
| <link rel="next" title="Quick Start Tutorial for Compiling Deep Learning Models" href="relay_quick_start.html" /> |
| <link rel="prev" title="Blitz Course to TensorIR" href="tensor_ir_blitz_course.html" /> |
| </head> |
| |
| <body class="wy-body-for-nav"> |
| |
| |
| <div class="wy-grid-for-nav"> |
| |
| |
| <header class="header"> |
| <div class="innercontainer"> |
| <div class="headerInner d-flex justify-content-between align-items-center"> |
| <div class="headerLogo"> |
| <a href="https://tvm.apache.org/"><img src=https://tvm.apache.org/assets/images/logo.svg alt="logo"></a> |
| </div> |
| |
| <div id="headMenu" class="headerNav"> |
| <button type="button" id="closeHeadMenu" class="navCloseBtn"><img src="../_static/img/close-icon.svg" alt="Close"></button> |
| <ul class="nav"> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvm.apache.org/community>Community</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvm.apache.org/download>Download</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvm.apache.org/vta>VTA</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvm.apache.org/blog>Blog</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvm.apache.org/docs>Docs</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://tvmconf.org>Conference</a> |
| </li> |
| <li class="nav-item"> |
| <a class="nav-link" href=https://github.com/apache/tvm/>Github</a> |
| </li> |
| </ul> |
| <div class="responsivetlcdropdown"> |
| <button type="button" class="btn-link"> |
| ASF |
| </button> |
| <ul> |
| <li> |
| <a href=https://apache.org/>Apache Homepage</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/licenses/>License</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/security/>Security</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/foundation/thanks.html>Thanks</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/events/current-event>Events</a> |
| </li> |
| </ul> |
| </div> |
| </div> |
| <div class="responsiveMenuIcon"> |
| <button type="button" id="menuBtn" class="btn-menu"><img src="../_static/img/menu-icon.svg" alt="Menu Icon"></button> |
| </div> |
| |
| <div class="tlcDropdown"> |
| <div class="dropdown"> |
| <button type="button" class="btn-link dropdown-toggle" data-toggle="dropdown" aria-haspopup="true" aria-expanded="false"> |
| ASF |
| </button> |
| <div class="dropdown-menu dropdown-menu-right"> |
| <ul> |
| <li> |
| <a href=https://apache.org/>Apache Homepage</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/licenses/>License</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/security/>Security</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/foundation/thanks.html>Thanks</a> |
| </li> |
| <li> |
| <a href=https://www.apache.org/events/current-event>Events</a> |
| </li> |
| </ul> |
| </div> |
| </div> |
| </div> |
| </div> |
| </div> |
| </header> |
| |
| <nav data-toggle="wy-nav-shift" class="wy-nav-side fixed"> |
| <div class="wy-side-scroll"> |
| <div class="wy-side-nav-search" > |
| |
| |
| |
| <a href="../index.html"> |
| |
| |
| |
| |
| <img src="../_static/tvm-logo-small.png" class="logo" alt="Logo"/> |
| |
| </a> |
| |
| |
| |
| |
| <input type="checkbox" class="version-toggle-box" hidden id="version-toggle"> |
| <label for="version-toggle" class="version-toggle-label"> |
| <div tabindex="0" class="version version-selector version-selector-show"> |
| 0.17.dev0 <span class="chevron versions-hidden"><svg fill="none" height="24" viewBox="0 0 24 24" width="24" xmlns="http://www.w3.org/2000/svg"><path d="m8 4 8 8-8 8" stroke="#000" stroke-linecap="round" stroke-linejoin="round" stroke-width="2"/></svg></span><span class="chevron versions-shown"><svg fill="none" height="24" viewBox="0 0 24 24" width="24" xmlns="http://www.w3.org/2000/svg"><path d="m4 8 8 8 8-8" stroke="#000" stroke-linecap="round" stroke-linejoin="round" stroke-width="2"/></svg></span> |
| </div> |
| </label> |
| <div class="version-details wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation"> |
| <p class="caption" role="heading"><span class="caption-text">Versions</span></p> |
| <ol style="text-align: left"> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="/">0.17.dev0 (main)</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.8.0/">v0.8.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.9.0/">v0.9.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.10.0/">v0.10.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.11.0/">v0.11.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.12.0/">v0.12.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.13.0/">v0.13.0</a></div></li> |
| |
| |
| |
| |
| <li><div class="version"><a style="font-size: 0.8em; padding: 4px" href="v0.14.0/">v0.14.0</a></div></li> |
| |
| </ol> |
| </div> |
| |
| |
| |
| |
| <div role="search"> |
| <form id="rtd-search-form" class="wy-form" action="../search.html" method="get"> |
| <input type="text" name="q" placeholder="Search docs" aria-label="Search docs" /> |
| <input type="hidden" name="check_keywords" value="yes" /> |
| <input type="hidden" name="area" value="default" /> |
| </form> |
| </div> |
| |
| |
| </div> |
| |
| |
| <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation"> |
| |
| |
| |
| |
| |
| |
| <p class="caption" role="heading"><span class="caption-text">Getting Started</span></p> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../install/index.html">Installing TVM</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../contribute/index.html">Contributor Guide</a></li> |
| </ul> |
| <p class="caption" role="heading"><span class="caption-text">User Guide</span></p> |
| <ul class="current"> |
| <li class="toctree-l1 current"><a class="reference internal" href="index.html">User Tutorial</a><ul class="current"> |
| <li class="toctree-l2"><a class="reference internal" href="introduction.html">Introduction</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="introduction.html#an-overview-of-tvm-and-model-optimization">An Overview of TVM and Model Optimization</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="install.html">Installing TVM</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="tvmc_command_line_driver.html">Compiling and Optimizing a Model with TVMC</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="tvmc_python.html">Getting Starting using TVMC Python: a high-level API for TVM</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="autotvm_relay_x86.html">Compiling and Optimizing a Model with the Python Interface (AutoTVM)</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="tensor_expr_get_started.html">Working with Operators Using Tensor Expression</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="autotvm_matmul_x86.html">Optimizing Operators with Schedule Templates and AutoTVM</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="auto_scheduler_matmul_x86.html">Optimizing Operators with Auto-scheduling</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="tensor_ir_blitz_course.html">Blitz Course to TensorIR</a></li> |
| <li class="toctree-l2 current"><a class="current reference internal" href="#">Cross Compilation and RPC</a><ul> |
| <li class="toctree-l3"><a class="reference internal" href="#build-tvm-runtime-on-device">Build TVM Runtime on Device</a></li> |
| <li class="toctree-l3"><a class="reference internal" href="#set-up-rpc-server-on-device">Set Up RPC Server on Device</a></li> |
| <li class="toctree-l3"><a class="reference internal" href="#declare-and-cross-compile-kernel-on-local-machine">Declare and Cross Compile Kernel on Local Machine</a></li> |
| <li class="toctree-l3"><a class="reference internal" href="#run-cpu-kernel-remotely-by-rpc">Run CPU Kernel Remotely by RPC</a></li> |
| <li class="toctree-l3"><a class="reference internal" href="#run-opencl-kernel-remotely-by-rpc">Run OpenCL Kernel Remotely by RPC</a></li> |
| <li class="toctree-l3"><a class="reference internal" href="#summary">Summary</a></li> |
| </ul> |
| </li> |
| <li class="toctree-l2"><a class="reference internal" href="relay_quick_start.html">Quick Start Tutorial for Compiling Deep Learning Models</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="uma.html">Making your Hardware Accelerator TVM-ready with UMA</a></li> |
| <li class="toctree-l2"><a class="reference internal" href="intro_topi.html">Introduction to TOPI</a></li> |
| </ul> |
| </li> |
| <li class="toctree-l1"><a class="reference internal" href="../how_to/index.html">How To Guides</a></li> |
| </ul> |
| <p class="caption" role="heading"><span class="caption-text">Developer Guide</span></p> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../dev/tutorial/index.html">Developer Tutorial</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../dev/how_to/how_to.html">Developer How-To Guide</a></li> |
| </ul> |
| <p class="caption" role="heading"><span class="caption-text">Architecture Guide</span></p> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../arch/index.html">Design and Architecture</a></li> |
| </ul> |
| <p class="caption" role="heading"><span class="caption-text">Topic Guides</span></p> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../topic/microtvm/index.html">microTVM: TVM on bare-metal</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../topic/vta/index.html">VTA: Versatile Tensor Accelerator</a></li> |
| </ul> |
| <p class="caption" role="heading"><span class="caption-text">Reference Guide</span></p> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../reference/langref/index.html">Language Reference</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../reference/api/python/index.html">Python API</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../reference/api/links.html">Other APIs</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../reference/publications.html">Publications</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../genindex.html">Index</a></li> |
| </ul> |
| |
| |
| |
| </div> |
| |
| </div> |
| </nav> |
| |
| <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap"> |
| |
| <nav class="wy-nav-top" aria-label="top navigation" data-toggle="wy-nav-top"> |
| |
| <div class="togglemenu"> |
| |
| </div> |
| <div class="nav-content"> |
| <!-- tvm --> |
| Table of Contents |
| </div> |
| |
| </nav> |
| |
| |
| <div class="wy-nav-content"> |
| |
| <div class="rst-content"> |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| <div role="navigation" aria-label="breadcrumbs navigation"> |
| |
| <ul class="wy-breadcrumbs"> |
| |
| <li><a href="../index.html">Docs</a> <span class="br-arrow">></span></li> |
| |
| <li><a href="index.html">User Tutorial</a> <span class="br-arrow">></span></li> |
| |
| <li>Cross Compilation and RPC</li> |
| |
| |
| |
| |
| |
| |
| |
| |
| <li class="wy-breadcrumbs-aside"> |
| |
| |
| |
| <a href="https://github.com/apache/tvm/edit/main/gallery/tutorial/cross_compilation_and_rpc.py" class="fa fa-github"> Edit on GitHub</a> |
| |
| |
| |
| </li> |
| |
| </ul> |
| |
| |
| <hr/> |
| </div> |
| <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article"> |
| <div itemprop="articleBody"> |
| |
| <div class="sphx-glr-download-link-note admonition note"> |
| <p class="admonition-title">Note</p> |
| <p>This tutorial can be used interactively with Google Colab! You can also click |
| <a class="reference internal" href="#sphx-glr-download-tutorial-cross-compilation-and-rpc-py"><span class="std std-ref">here</span></a> to run the Jupyter notebook locally.</p> |
| <a class="reference external image-reference" href="https://colab.research.google.com/github/apache/tvm-site/blob/asf-site/docs/_downloads/f289ca2466fcf79c024068c1f8642bd0/cross_compilation_and_rpc.ipynb"><img alt="https://raw.githubusercontent.com/tlc-pack/web-data/main/images/utilities/colab_button.svg" class="align-center" src="https://raw.githubusercontent.com/tlc-pack/web-data/main/images/utilities/colab_button.svg" width="300px" /></a> |
| </div> |
| <div class="sphx-glr-example-title section" id="cross-compilation-and-rpc"> |
| <span id="tutorial-cross-compilation-and-rpc"></span><span id="sphx-glr-tutorial-cross-compilation-and-rpc-py"></span><h1>Cross Compilation and RPC<a class="headerlink" href="#cross-compilation-and-rpc" title="Permalink to this headline">¶</a></h1> |
| <p><strong>Author</strong>: <a class="reference external" href="https://github.com/ZihengJiang/">Ziheng Jiang</a>, <a class="reference external" href="https://github.com/merrymercy/">Lianmin Zheng</a></p> |
| <p>This tutorial introduces cross compilation and remote device |
| execution with RPC in TVM.</p> |
| <p>With cross compilation and RPC, you can <strong>compile a program on your |
| local machine then run it on the remote device</strong>. It is useful when |
| the remote device resource are limited, like Raspberry Pi and mobile |
| platforms. In this tutorial, we will use the Raspberry Pi for a CPU example |
| and the Firefly-RK3399 for an OpenCL example.</p> |
| <div class="section" id="build-tvm-runtime-on-device"> |
| <h2>Build TVM Runtime on Device<a class="headerlink" href="#build-tvm-runtime-on-device" title="Permalink to this headline">¶</a></h2> |
| <p>The first step is to build the TVM runtime on the remote device.</p> |
| <div class="admonition note"> |
| <p class="admonition-title">Note</p> |
| <p>All instructions in both this section and the next section should be |
| executed on the target device, e.g. Raspberry Pi. We assume the target |
| is running Linux.</p> |
| </div> |
| <p>Since we do compilation on the local machine, the remote device is only used |
| for running the generated code. We only need to build the TVM runtime on |
| the remote device.</p> |
| <div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>git<span class="w"> </span>clone<span class="w"> </span>--recursive<span class="w"> </span>https://github.com/apache/tvm<span class="w"> </span>tvm |
| <span class="nb">cd</span><span class="w"> </span>tvm |
| make<span class="w"> </span>runtime<span class="w"> </span>-j2 |
| </pre></div> |
| </div> |
| <p>After building the runtime successfully, we need to set environment variables |
| in <code class="code docutils literal notranslate"><span class="pre">~/.bashrc</span></code> file. We can edit <code class="code docutils literal notranslate"><span class="pre">~/.bashrc</span></code> |
| using <code class="code docutils literal notranslate"><span class="pre">vi</span> <span class="pre">~/.bashrc</span></code> and add the line below (Assuming your TVM |
| directory is in <code class="code docutils literal notranslate"><span class="pre">~/tvm</span></code>):</p> |
| <div class="highlight-bash notranslate"><div class="highlight"><pre><span></span><span class="nb">export</span><span class="w"> </span><span class="nv">PYTHONPATH</span><span class="o">=</span><span class="nv">$PYTHONPATH</span>:~/tvm/python |
| </pre></div> |
| </div> |
| <p>To update the environment variables, execute <code class="code docutils literal notranslate"><span class="pre">source</span> <span class="pre">~/.bashrc</span></code>.</p> |
| </div> |
| <div class="section" id="set-up-rpc-server-on-device"> |
| <h2>Set Up RPC Server on Device<a class="headerlink" href="#set-up-rpc-server-on-device" title="Permalink to this headline">¶</a></h2> |
| <p>To start an RPC server, run the following command on your remote device |
| (Which is Raspberry Pi in this example).</p> |
| <blockquote> |
| <div><div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>python<span class="w"> </span>-m<span class="w"> </span>tvm.exec.rpc_server<span class="w"> </span>--host<span class="w"> </span><span class="m">0</span>.0.0.0<span class="w"> </span>--port<span class="o">=</span><span class="m">9090</span> |
| </pre></div> |
| </div> |
| </div></blockquote> |
| <p>If you see the line below, it means the RPC server started |
| successfully on your device.</p> |
| <blockquote> |
| <div><div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>INFO:root:RPCServer:<span class="w"> </span><span class="nb">bind</span><span class="w"> </span>to<span class="w"> </span><span class="m">0</span>.0.0.0:9090 |
| </pre></div> |
| </div> |
| </div></blockquote> |
| </div> |
| <div class="section" id="declare-and-cross-compile-kernel-on-local-machine"> |
| <h2>Declare and Cross Compile Kernel on Local Machine<a class="headerlink" href="#declare-and-cross-compile-kernel-on-local-machine" title="Permalink to this headline">¶</a></h2> |
| <div class="admonition note"> |
| <p class="admonition-title">Note</p> |
| <p>Now we go back to the local machine, which has a full TVM installed |
| (with LLVM).</p> |
| </div> |
| <p>Here we will declare a simple kernel on the local machine:</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span> |
| |
| <span class="kn">import</span> <span class="nn">tvm</span> |
| <span class="kn">from</span> <span class="nn">tvm</span> <span class="kn">import</span> <span class="n">te</span> |
| <span class="kn">from</span> <span class="nn">tvm</span> <span class="kn">import</span> <span class="n">rpc</span> |
| <span class="kn">from</span> <span class="nn">tvm.contrib</span> <span class="kn">import</span> <span class="n">utils</span> |
| |
| <a href="../reference/api/python/tir.html#tvm.tir.IntImm" title="tvm.tir.IntImm" class="sphx-glr-backref-module-tvm-tir sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">n</span></a> <span class="o">=</span> <a href="../reference/api/python/runtime.html#tvm.runtime.convert" title="tvm.runtime.convert" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">runtime</span><span class="o">.</span><span class="n">convert</span></a><span class="p">(</span><span class="mi">1024</span><span class="p">)</span> |
| <a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span></a> <span class="o">=</span> <a href="../reference/api/python/te.html#tvm.te.placeholder" title="tvm.te.placeholder" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">placeholder</span></a><span class="p">((</span><a href="../reference/api/python/tir.html#tvm.tir.IntImm" title="tvm.tir.IntImm" class="sphx-glr-backref-module-tvm-tir sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">n</span></a><span class="p">,),</span> <span class="n">name</span><span class="o">=</span><span class="s2">"A"</span><span class="p">)</span> |
| <a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a> <span class="o">=</span> <a href="../reference/api/python/te.html#tvm.te.compute" title="tvm.te.compute" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">compute</span></a><span class="p">((</span><a href="../reference/api/python/tir.html#tvm.tir.IntImm" title="tvm.tir.IntImm" class="sphx-glr-backref-module-tvm-tir sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">n</span></a><span class="p">,),</span> <span class="k">lambda</span> <span class="n">i</span><span class="p">:</span> <a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span></a><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="o">+</span> <span class="mf">1.0</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">"B"</span><span class="p">)</span> |
| <a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a> <span class="o">=</span> <a href="../reference/api/python/te.html#tvm.te.create_schedule" title="tvm.te.create_schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span></a><span class="p">(</span><a href="../reference/api/python/te.html#tvm.te.Tensor.op" title="tvm.te.Tensor.op" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-property"><span class="n">B</span><span class="o">.</span><span class="n">op</span></a><span class="p">)</span> |
| </pre></div> |
| </div> |
| <p>Then we cross compile the kernel. |
| The target should be ‘llvm -mtriple=armv7l-linux-gnueabihf’ for |
| Raspberry Pi 3B, but we use ‘llvm’ here to make this tutorial runnable |
| on our webpage building server. See the detailed note in the following block.</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><a href="https://docs.python.org/3/library/functions.html#bool" title="builtins.bool" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">local_demo</span></a> <span class="o">=</span> <span class="kc">True</span> |
| |
| <span class="k">if</span> <a href="https://docs.python.org/3/library/functions.html#bool" title="builtins.bool" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">local_demo</span></a><span class="p">:</span> |
| <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a> <span class="o">=</span> <span class="s2">"llvm"</span> |
| <span class="k">else</span><span class="p">:</span> |
| <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a> <span class="o">=</span> <span class="s2">"llvm -mtriple=armv7l-linux-gnueabihf"</span> |
| |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a> <span class="o">=</span> <a href="../reference/api/python/driver.html#tvm.build" title="tvm.build" class="sphx-glr-backref-module-tvm sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">build</span></a><span class="p">(</span><a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a><span class="p">,</span> <span class="p">[</span><a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span></a><span class="p">,</span> <a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a><span class="p">],</span> <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a><span class="o">=</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">"add_one"</span><span class="p">)</span> |
| <span class="c1"># save the lib at a local temp folder</span> |
| <a href="../reference/api/python/contrib.html#tvm.contrib.utils.TempDirectory" title="tvm.contrib.utils.TempDirectory" class="sphx-glr-backref-module-tvm-contrib-utils sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">temp</span></a> <span class="o">=</span> <a href="../reference/api/python/contrib.html#tvm.contrib.utils.tempdir" title="tvm.contrib.utils.tempdir" class="sphx-glr-backref-module-tvm-contrib-utils sphx-glr-backref-type-py-function"><span class="n">utils</span><span class="o">.</span><span class="n">tempdir</span></a><span class="p">()</span> |
| <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a> <span class="o">=</span> <a href="../reference/api/python/contrib.html#tvm.contrib.utils.TempDirectory.relpath" title="tvm.contrib.utils.TempDirectory.relpath" class="sphx-glr-backref-module-tvm-contrib-utils sphx-glr-backref-type-py-method"><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span></a><span class="p">(</span><span class="s2">"lib.tar"</span><span class="p">)</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module.export_library" title="tvm.runtime.Module.export_library" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-method"><span class="n">func</span><span class="o">.</span><span class="n">export_library</span></a><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a><span class="p">)</span> |
| </pre></div> |
| </div> |
| <div class="admonition note"> |
| <p class="admonition-title">Note</p> |
| <p>To run this tutorial with a real remote device, change <code class="code docutils literal notranslate"><span class="pre">local_demo</span></code> |
| to False and replace <code class="code docutils literal notranslate"><span class="pre">target</span></code> in <code class="code docutils literal notranslate"><span class="pre">build</span></code> with the appropriate |
| target triple for your device. The target triple which might be |
| different for different devices. For example, it is |
| <code class="code docutils literal notranslate"><span class="pre">'llvm</span> <span class="pre">-mtriple=armv7l-linux-gnueabihf'</span></code> for Raspberry Pi 3B and |
| <code class="code docutils literal notranslate"><span class="pre">'llvm</span> <span class="pre">-mtriple=aarch64-linux-gnu'</span></code> for RK3399.</p> |
| <p>Usually, you can query the target by running <code class="code docutils literal notranslate"><span class="pre">gcc</span> <span class="pre">-v</span></code> on your |
| device, and looking for the line starting with <code class="code docutils literal notranslate"><span class="pre">Target:</span></code> |
| (Though it may still be a loose configuration.)</p> |
| <p>Besides <code class="code docutils literal notranslate"><span class="pre">-mtriple</span></code>, you can also set other compilation options |
| like:</p> |
| <ul> |
| <li><dl class="simple"> |
| <dt>-mcpu=<cpuname></dt><dd><p>Specify a specific chip in the current architecture to generate code for. By default this is inferred from the target triple and autodetected to the current architecture.</p> |
| </dd> |
| </dl> |
| </li> |
| <li><dl> |
| <dt>-mattr=a1,+a2,-a3,…</dt><dd><p>Override or control specific attributes of the target, such as whether SIMD operations are enabled or not. The default set of attributes is set by the current CPU. |
| To get the list of available attributes, you can do:</p> |
| <div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>llc<span class="w"> </span>-mtriple<span class="o">=</span><your<span class="w"> </span>device<span class="w"> </span>target<span class="w"> </span>triple><span class="w"> </span>-mattr<span class="o">=</span><span class="nb">help</span> |
| </pre></div> |
| </div> |
| </dd> |
| </dl> |
| </li> |
| </ul> |
| <p>These options are consistent with <a class="reference external" href="http://llvm.org/docs/CommandGuide/llc.html">llc</a>. |
| It is recommended to set target triple and feature set to contain specific |
| feature available, so we can take full advantage of the features of the |
| board. |
| You can find more details about cross compilation attributes from |
| <a class="reference external" href="https://clang.llvm.org/docs/CrossCompilation.html">LLVM guide of cross compilation</a>.</p> |
| </div> |
| </div> |
| <div class="section" id="run-cpu-kernel-remotely-by-rpc"> |
| <h2>Run CPU Kernel Remotely by RPC<a class="headerlink" href="#run-cpu-kernel-remotely-by-rpc" title="Permalink to this headline">¶</a></h2> |
| <p>We show how to run the generated CPU kernel on the remote device. |
| First we obtain an RPC session from remote device.</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="k">if</span> <a href="https://docs.python.org/3/library/functions.html#bool" title="builtins.bool" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">local_demo</span></a><span class="p">:</span> |
| <a href="../reference/api/python/rpc.html#tvm.rpc.LocalSession" title="tvm.rpc.LocalSession" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">remote</span></a> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.LocalSession" title="tvm.rpc.LocalSession" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-class"><span class="n">rpc</span><span class="o">.</span><span class="n">LocalSession</span></a><span class="p">()</span> |
| <span class="k">else</span><span class="p">:</span> |
| <span class="c1"># The following is my environment, change this to the IP address of your target device</span> |
| <span class="n">host</span> <span class="o">=</span> <span class="s2">"10.77.1.162"</span> |
| <span class="n">port</span> <span class="o">=</span> <span class="mi">9090</span> |
| <a href="../reference/api/python/rpc.html#tvm.rpc.LocalSession" title="tvm.rpc.LocalSession" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">remote</span></a> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.connect" title="tvm.rpc.connect" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-function"><span class="n">rpc</span><span class="o">.</span><span class="n">connect</span></a><span class="p">(</span><span class="n">host</span><span class="p">,</span> <span class="n">port</span><span class="p">)</span> |
| </pre></div> |
| </div> |
| <p>Upload the lib to the remote device, then invoke a device local |
| compiler to relink them. Now <cite>func</cite> is a remote module object.</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.upload" title="tvm.rpc.RPCSession.upload" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">upload</span></a><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a><span class="p">)</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.load_module" title="tvm.rpc.RPCSession.load_module" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">load_module</span></a><span class="p">(</span><span class="s2">"lib.tar"</span><span class="p">)</span> |
| |
| <span class="c1"># create arrays on the remote device</span> |
| <span class="n">dev</span> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.cpu" title="tvm.rpc.RPCSession.cpu" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">cpu</span></a><span class="p">()</span> |
| <span class="n">a</span> <span class="o">=</span> <a href="../reference/api/python/ndarray.html#tvm.nd.array" title="tvm.nd.array" class="sphx-glr-backref-module-tvm-nd sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span></a><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="mi">1024</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span><span class="o">.</span><span class="n">dtype</span></a><span class="p">),</span> <span class="n">dev</span><span class="p">)</span> |
| <span class="n">b</span> <span class="o">=</span> <a href="../reference/api/python/ndarray.html#tvm.nd.array" title="tvm.nd.array" class="sphx-glr-backref-module-tvm-nd sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span></a><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="mi">1024</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span><span class="o">.</span><span class="n">dtype</span></a><span class="p">),</span> <span class="n">dev</span><span class="p">)</span> |
| <span class="c1"># the function will run on the remote device</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">)</span> |
| <span class="n">np</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_equal</span><span class="p">(</span><span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="mi">1</span><span class="p">)</span> |
| </pre></div> |
| </div> |
| <p>When you want to evaluate the performance of the kernel on the remote |
| device, it is important to avoid the overhead of network. |
| <code class="code docutils literal notranslate"><span class="pre">time_evaluator</span></code> will returns a remote function that runs the |
| function over number times, measures the cost per run on the remote |
| device and returns the measured cost. Network overhead is excluded.</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">time_f</span> <span class="o">=</span> <a href="../reference/api/python/runtime.html#tvm.runtime.Module.time_evaluator" title="tvm.runtime.Module.time_evaluator" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-method"><span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span></a><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span></a><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">number</span><span class="o">=</span><span class="mi">10</span><span class="p">)</span> |
| <span class="n">cost</span> <span class="o">=</span> <span class="n">time_f</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">)</span><span class="o">.</span><span class="n">mean</span> |
| <span class="nb">print</span><span class="p">(</span><span class="s2">"</span><span class="si">%g</span><span class="s2"> secs/op"</span> <span class="o">%</span> <span class="n">cost</span><span class="p">)</span> |
| </pre></div> |
| </div> |
| <div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>1.2e-07 secs/op |
| </pre></div> |
| </div> |
| </div> |
| <div class="section" id="run-opencl-kernel-remotely-by-rpc"> |
| <h2>Run OpenCL Kernel Remotely by RPC<a class="headerlink" href="#run-opencl-kernel-remotely-by-rpc" title="Permalink to this headline">¶</a></h2> |
| <p>For remote OpenCL devices, the workflow is almost the same as above. |
| You can define the kernel, upload files, and run via RPC.</p> |
| <div class="admonition note"> |
| <p class="admonition-title">Note</p> |
| <p>Raspberry Pi does not support OpenCL, the following code is tested on |
| Firefly-RK3399. You may follow this <a class="reference external" href="https://gist.github.com/mli/585aed2cec0b5178b1a510f9f236afa2">tutorial</a> |
| to setup the OS and OpenCL driver for RK3399.</p> |
| <p>Also we need to build the runtime with OpenCL enabled on rk3399 board. In the TVM |
| root directory, execute</p> |
| </div> |
| <div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>cp<span class="w"> </span>cmake/config.cmake<span class="w"> </span>. |
| sed<span class="w"> </span>-i<span class="w"> </span><span class="s2">"s/USE_OPENCL OFF/USE_OPENCL ON/"</span><span class="w"> </span>config.cmake |
| make<span class="w"> </span>runtime<span class="w"> </span>-j4 |
| </pre></div> |
| </div> |
| <p>The following function shows how we run an OpenCL kernel remotely</p> |
| <div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="k">def</span> <span class="nf">run_opencl</span><span class="p">():</span> |
| <span class="c1"># NOTE: This is the setting for my rk3399 board. You need to modify</span> |
| <span class="c1"># them according to your environment.</span> |
| <span class="n">opencl_device_host</span> <span class="o">=</span> <span class="s2">"10.77.1.145"</span> |
| <span class="n">opencl_device_port</span> <span class="o">=</span> <span class="mi">9090</span> |
| <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a> <span class="o">=</span> <a href="../reference/api/python/target.html#tvm.target.Target" title="tvm.target.Target" class="sphx-glr-backref-module-tvm-target sphx-glr-backref-type-py-class"><span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span></a><span class="p">(</span><span class="s2">"opencl"</span><span class="p">,</span> <span class="n">host</span><span class="o">=</span><span class="s2">"llvm -mtriple=aarch64-linux-gnu"</span><span class="p">)</span> |
| |
| <span class="c1"># create schedule for the above "add one" compute declaration</span> |
| <a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a> <span class="o">=</span> <a href="../reference/api/python/te.html#tvm.te.create_schedule" title="tvm.te.create_schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span></a><span class="p">(</span><a href="../reference/api/python/te.html#tvm.te.Tensor.op" title="tvm.te.Tensor.op" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-property"><span class="n">B</span><span class="o">.</span><span class="n">op</span></a><span class="p">)</span> |
| <span class="n">xo</span><span class="p">,</span> <span class="n">xi</span> <span class="o">=</span> <a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a><span class="p">[</span><a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><a href="../reference/api/python/te.html#tvm.te.Tensor.op" title="tvm.te.Tensor.op" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-property"><span class="n">B</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span></a><span class="p">[</span><span class="mi">0</span><span class="p">],</span> <span class="n">factor</span><span class="o">=</span><span class="mi">32</span><span class="p">)</span> |
| <a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a><span class="p">[</span><a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">xo</span><span class="p">,</span> <a href="../reference/api/python/te.html#tvm.te.thread_axis" title="tvm.te.thread_axis" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span></a><span class="p">(</span><span class="s2">"blockIdx.x"</span><span class="p">))</span> |
| <a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a><span class="p">[</span><a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">xi</span><span class="p">,</span> <a href="../reference/api/python/te.html#tvm.te.thread_axis" title="tvm.te.thread_axis" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-function"><span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span></a><span class="p">(</span><span class="s2">"threadIdx.x"</span><span class="p">))</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a> <span class="o">=</span> <a href="../reference/api/python/driver.html#tvm.build" title="tvm.build" class="sphx-glr-backref-module-tvm sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">build</span></a><span class="p">(</span><a href="../reference/api/python/te.html#tvm.te.Schedule" title="tvm.te.Schedule" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">s</span></a><span class="p">,</span> <span class="p">[</span><a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span></a><span class="p">,</span> <a href="../reference/api/python/te.html#tvm.te.Tensor" title="tvm.te.Tensor" class="sphx-glr-backref-module-tvm-te sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">B</span></a><span class="p">],</span> <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a><span class="o">=</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">target</span></a><span class="p">)</span> |
| |
| <a href="../reference/api/python/rpc.html#tvm.rpc.LocalSession" title="tvm.rpc.LocalSession" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">remote</span></a> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.connect" title="tvm.rpc.connect" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-function"><span class="n">rpc</span><span class="o">.</span><span class="n">connect</span></a><span class="p">(</span><span class="n">opencl_device_host</span><span class="p">,</span> <span class="n">opencl_device_port</span><span class="p">)</span> |
| |
| <span class="c1"># export and upload</span> |
| <a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a> <span class="o">=</span> <a href="../reference/api/python/contrib.html#tvm.contrib.utils.TempDirectory.relpath" title="tvm.contrib.utils.TempDirectory.relpath" class="sphx-glr-backref-module-tvm-contrib-utils sphx-glr-backref-type-py-method"><span class="n">temp</span><span class="o">.</span><span class="n">relpath</span></a><span class="p">(</span><span class="s2">"lib_cl.tar"</span><span class="p">)</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module.export_library" title="tvm.runtime.Module.export_library" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-method"><span class="n">func</span><span class="o">.</span><span class="n">export_library</span></a><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a><span class="p">)</span> |
| <a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.upload" title="tvm.rpc.RPCSession.upload" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">upload</span></a><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">path</span></a><span class="p">)</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.load_module" title="tvm.rpc.RPCSession.load_module" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">load_module</span></a><span class="p">(</span><span class="s2">"lib_cl.tar"</span><span class="p">)</span> |
| |
| <span class="c1"># run</span> |
| <span class="n">dev</span> <span class="o">=</span> <a href="../reference/api/python/rpc.html#tvm.rpc.RPCSession.cl" title="tvm.rpc.RPCSession.cl" class="sphx-glr-backref-module-tvm-rpc sphx-glr-backref-type-py-method"><span class="n">remote</span><span class="o">.</span><span class="n">cl</span></a><span class="p">()</span> |
| <span class="n">a</span> <span class="o">=</span> <a href="../reference/api/python/ndarray.html#tvm.nd.array" title="tvm.nd.array" class="sphx-glr-backref-module-tvm-nd sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span></a><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="mi">1024</span><span class="p">)</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span><span class="o">.</span><span class="n">dtype</span></a><span class="p">),</span> <span class="n">dev</span><span class="p">)</span> |
| <span class="n">b</span> <span class="o">=</span> <a href="../reference/api/python/ndarray.html#tvm.nd.array" title="tvm.nd.array" class="sphx-glr-backref-module-tvm-nd sphx-glr-backref-type-py-function"><span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span></a><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">(</span><span class="mi">1024</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><a href="https://docs.python.org/3/library/stdtypes.html#str" title="builtins.str" class="sphx-glr-backref-module-builtins sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">A</span><span class="o">.</span><span class="n">dtype</span></a><span class="p">),</span> <span class="n">dev</span><span class="p">)</span> |
| <a href="../reference/api/python/runtime.html#tvm.runtime.Module" title="tvm.runtime.Module" class="sphx-glr-backref-module-tvm-runtime sphx-glr-backref-type-py-class sphx-glr-backref-instance"><span class="n">func</span></a><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">)</span> |
| <span class="n">np</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_equal</span><span class="p">(</span><span class="n">b</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">a</span><span class="o">.</span><span class="n">numpy</span><span class="p">()</span> <span class="o">+</span> <span class="mi">1</span><span class="p">)</span> |
| <span class="nb">print</span><span class="p">(</span><span class="s2">"OpenCL test passed!"</span><span class="p">)</span> |
| </pre></div> |
| </div> |
| </div> |
| <div class="section" id="summary"> |
| <h2>Summary<a class="headerlink" href="#summary" title="Permalink to this headline">¶</a></h2> |
| <p>This tutorial provides a walk through of cross compilation and RPC |
| features in TVM.</p> |
| <ul class="simple"> |
| <li><p>Set up an RPC server on the remote device.</p></li> |
| <li><p>Set up the target device configuration to cross compile the kernels on the |
| local machine.</p></li> |
| <li><p>Upload and run the kernels remotely via the RPC API.</p></li> |
| </ul> |
| <div class="sphx-glr-footer sphx-glr-footer-example docutils container" id="sphx-glr-download-tutorial-cross-compilation-and-rpc-py"> |
| <div class="sphx-glr-download sphx-glr-download-python docutils container"> |
| <p><a class="reference download internal" download="" href="../_downloads/766206ab8f1fd80ac34d9816cb991a0d/cross_compilation_and_rpc.py"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Python</span> <span class="pre">source</span> <span class="pre">code:</span> <span class="pre">cross_compilation_and_rpc.py</span></code></a></p> |
| </div> |
| <div class="sphx-glr-download sphx-glr-download-jupyter docutils container"> |
| <p><a class="reference download internal" download="" href="../_downloads/f289ca2466fcf79c024068c1f8642bd0/cross_compilation_and_rpc.ipynb"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Jupyter</span> <span class="pre">notebook:</span> <span class="pre">cross_compilation_and_rpc.ipynb</span></code></a></p> |
| </div> |
| </div> |
| <p class="sphx-glr-signature"><a class="reference external" href="https://sphinx-gallery.github.io">Gallery generated by Sphinx-Gallery</a></p> |
| </div> |
| </div> |
| |
| |
| </div> |
| |
| </div> |
| |
| |
| <footer> |
| |
| <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation"> |
| |
| <a href="relay_quick_start.html" class="btn btn-neutral float-right" title="Quick Start Tutorial for Compiling Deep Learning Models" accesskey="n" rel="next">Next <span class="fa fa-arrow-circle-right"></span></a> |
| |
| |
| <a href="tensor_ir_blitz_course.html" class="btn btn-neutral float-left" title="Blitz Course to TensorIR" accesskey="p" rel="prev"><span class="fa fa-arrow-circle-left"></span> Previous</a> |
| |
| </div> |
| |
| <div id="button" class="backtop"><img src="../_static/img/right.svg" alt="backtop"/> </div> |
| <section class="footerSec"> |
| <div class="footerHeader"> |
| <div class="d-flex align-md-items-center justify-content-between flex-column flex-md-row"> |
| <div class="copywrite d-flex align-items-center"> |
| <h5 id="copy-right-info">© 2023 Apache Software Foundation | All rights reserved</h5> |
| </div> |
| </div> |
| |
| </div> |
| |
| <div> |
| <div class="footernote">Copyright © 2023 The Apache Software Foundation. Apache TVM, Apache, the Apache feather, and the Apache TVM project logo are either trademarks or registered trademarks of the Apache Software Foundation.</div> |
| </div> |
| |
| </section> |
| </footer> |
| </div> |
| </div> |
| |
| </section> |
| |
| </div> |
| |
| |
| <script src="https://cdnjs.cloudflare.com/ajax/libs/popper.js/1.12.9/umd/popper.min.js" integrity="sha384-ApNbgh9B+Y1QKtv3Rn7W3mgPxhU9K/ScQsAP7hUibX39j7fakFPskvXusvfa0b4Q" crossorigin="anonymous"></script> |
| <script src="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/js/bootstrap.min.js" integrity="sha384-JZR6Spejh4U02d8jOt6vLEHfe/JQGiRRSQQxSfFWpi1MquVdAyjUar5+76PVCmYl" crossorigin="anonymous"></script> |
| |
| </body> |
| <script type="text/javascript"> |
| jQuery(function () { |
| SphinxRtdTheme.Navigation.enable(true); |
| }); |
| </script> |
| |
| |
| |
| |
| <!-- Theme Analytics --> |
| <script> |
| (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){ |
| (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o), |
| m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m) |
| })(window,document,'script','https://www.google-analytics.com/analytics.js','ga'); |
| |
| ga('create', 'UA-75982049-2', 'auto'); |
| ga('send', 'pageview'); |
| </script> |
| |
| |
| |
| |
| </body> |
| </html> |
