From 8903b1a3251370ee1013fc2f9f3ef6004fa0e4b2 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 21 Jul 2020 03:23:18 -0600 Subject: [PATCH 01/22] [BYOC][CONTRIB] VITIS-AI integration --- CMakeLists.txt | 3 +- cmake/config.cmake | 3 + cmake/modules/contrib/VITISAI.cmake | 49 ++ docker/Dockerfile.ci_vai | 55 ++ docker/bash.sh | 23 + docker/install/ubuntu_install_python.sh | 2 +- docker/install/ubuntu_install_vai_core.sh | 37 ++ docs/deploy/vitis_ai.rst | 617 ++++++++++++++++++ python/tvm/contrib/target/vitis_ai.py | 109 ++++ python/tvm/contrib/vitis_ai_runtime.py | 54 ++ python/tvm/relay/op/contrib/vitis_ai.py | 92 +++ .../contrib/vitis_ai/vitis_ai_runtime.cc | 147 +++++ .../contrib/vitis_ai/vitis_ai_runtime.h | 97 +++ tests/python/contrib/test_vitis_ai_codegen.py | 203 ++++++ tests/python/contrib/test_vitis_ai_runtime.py | 104 +++ 15 files changed, 1593 insertions(+), 2 deletions(-) create mode 100644 cmake/modules/contrib/VITISAI.cmake create mode 100644 docker/Dockerfile.ci_vai create mode 100644 docker/install/ubuntu_install_vai_core.sh create mode 100755 docs/deploy/vitis_ai.rst create mode 100644 python/tvm/contrib/target/vitis_ai.py create mode 100644 python/tvm/contrib/vitis_ai_runtime.py create mode 100644 python/tvm/relay/op/contrib/vitis_ai.py create mode 100755 src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc create mode 100755 src/runtime/contrib/vitis_ai/vitis_ai_runtime.h create mode 100644 tests/python/contrib/test_vitis_ai_codegen.py create mode 100644 tests/python/contrib/test_vitis_ai_runtime.py diff --git a/CMakeLists.txt b/CMakeLists.txt index 80581076b925..6090d1d01c51 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -75,7 +75,7 @@ tvm_option(USE_COREML "Build with coreml support" OFF) tvm_option(USE_TARGET_ONNX "Build with ONNX Codegen support" OFF) tvm_option(USE_ARM_COMPUTE_LIB "Build with Arm Compute Library" OFF) tvm_option(USE_ARM_COMPUTE_LIB_GRAPH_RUNTIME "Build with Arm Compute Library graph runtime" OFF) - +tvm_option(USE_VITIS_AI "Build with VITIS-AI Codegen support" OFF) # include directories include_directories(${CMAKE_INCLUDE_PATH}) @@ -323,6 +323,7 @@ include(cmake/modules/contrib/TF_TVMDSOOP.cmake) include(cmake/modules/contrib/CoreML.cmake) include(cmake/modules/contrib/ONNX.cmake) include(cmake/modules/contrib/ArmComputeLib.cmake) +include(cmake/modules/contrib/VITISAI.cmake) include(cmake/modules/Git.cmake) include(cmake/modules/LibInfo.cmake) diff --git a/cmake/config.cmake b/cmake/config.cmake index e7e2a5a8a5c9..9e3094eb011b 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -220,6 +220,9 @@ set(USE_ETHOSN OFF) # otherwise use ETHOSN_HW (OFF) to use the software test infrastructure set(USE_ETHOSN_HW OFF) +# Whether use VITIS-AI codegen +set(USE_VITIS_AI OFF) + # Build ANTLR parser for Relay text format # Possible values: # - ON: enable ANTLR by searching default locations (cmake find_program for antlr4 and /usr/local for jar) diff --git a/cmake/modules/contrib/VITISAI.cmake b/cmake/modules/contrib/VITISAI.cmake new file mode 100644 index 000000000000..6ccdbaa15fdc --- /dev/null +++ b/cmake/modules/contrib/VITISAI.cmake @@ -0,0 +1,49 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +if(USE_VITIS_AI) + set(PYXIR_SHARED_LIB libpyxir.so) + find_package(PythonInterp 3.6 REQUIRED) + if(NOT PYTHON) + find_program(PYTHON NAMES python3 python3.6) + endif() + if(PYTHON) + execute_process(COMMAND "${PYTHON_EXECUTABLE}" "-c" + "import pyxir as px; print(px.get_include_dir()); print(px.get_lib_dir());" + RESULT_VARIABLE __result + OUTPUT_VARIABLE __output + OUTPUT_STRIP_TRAILING_WHITESPACE) + + if(__result MATCHES 0) + string(REGEX REPLACE ";" "\\\\;" __values ${__output}) + string(REGEX REPLACE "\r?\n" ";" __values ${__values}) + list(GET __values 0 PYXIR_INCLUDE_DIR) + list(GET __values 1 PYXIR_LIB_DIR) + endif() + + else() + message(STATUS "To find Pyxir, Python interpreter is required to be found.") + endif() + +message(STATUS "Build with contrib.vitisai") +include_directories(${PYXIR_INCLUDE_DIR}) +file(GLOB VAI_CONTRIB_SRC src/runtime/contrib/vitis_ai/*.cc) +link_directories(${PYXIR_LIB_DIR}) +list(APPEND TVM_RUNTIME_LINKER_LIBS "pyxir") +list(APPEND RUNTIME_SRCS ${VAI_CONTRIB_SRC}) +endif(USE_VITIS_AI) + diff --git a/docker/Dockerfile.ci_vai b/docker/Dockerfile.ci_vai new file mode 100644 index 000000000000..d0a30f20520e --- /dev/null +++ b/docker/Dockerfile.ci_vai @@ -0,0 +1,55 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +# CI docker VAI env +FROM xilinx/vitis-ai:latest + +RUN apt-get update --fix-missing + + +COPY install/ubuntu_install_core.sh /install/ubuntu_install_core.sh +RUN bash /install/ubuntu_install_core.sh + +COPY install/ubuntu_install_python.sh /install/ubuntu_install_python.sh +RUN bash /install/ubuntu_install_python.sh + +COPY install/ubuntu_install_python_package.sh /install/ubuntu_install_python_package.sh +RUN bash /install/ubuntu_install_python_package.sh + +COPY install/ubuntu_install_llvm.sh /install/ubuntu_install_llvm.sh +RUN bash /install/ubuntu_install_llvm.sh + +# NNPACK deps +COPY install/ubuntu_install_nnpack.sh /install/ubuntu_install_nnpack.sh +RUN bash /install/ubuntu_install_nnpack.sh + +ENV PATH $PATH:$CARGO_HOME/bin:/usr/lib/go-1.10/bin + +# ANTLR deps +COPY install/ubuntu_install_java.sh /install/ubuntu_install_java.sh +RUN bash /install/ubuntu_install_java.sh + +# Install Vitis-AI ubuntu dependencies +COPY install/ubuntu_install_vai_core.sh /install/ubuntu_install_vai_core.sh +RUN bash /install/ubuntu_install_vai_core.sh + +# Install dependencies inside vitis-ai-tensorflow conda +RUN . $VAI_ROOT/conda/etc/profile.d/conda.sh && \ + conda activate vitis-ai-tensorflow && \ + pip install --no-cache-dir antlr4-python3-runtime + +ENV USER="root" diff --git a/docker/bash.sh b/docker/bash.sh index 73bfb12268f3..7817adaf8839 100755 --- a/docker/bash.sh +++ b/docker/bash.sh @@ -75,6 +75,27 @@ else CI_PY_ENV="" fi +if [[ "${DOCKER_IMAGE_NAME}" == *"ci_vai"* && -d "/dev/shm" && -d "/opt/xilinx/dsa" && -d "/opt/xilinx/overlaybins" ]]; then + WORKSPACE_VOLUMES="-v /dev/shm:/dev/shm -v /opt/xilinx/dsa:/opt/xilinx/dsa -v /opt/xilinx/overlaybins:/opt/xilinx/overlaybins" + XCLMGMT_DRIVER="$(find /dev -name xclmgmt\*)" + DOCKER_DEVICES="" + for i in ${XCLMGMT_DRIVER} ; + do + DOCKER_DEVICES+="--device=$i " + done + + RENDER_DRIVER="$(find /dev/dri -name renderD\*)" + for i in ${RENDER_DRIVER} ; + do + DOCKER_DEVICES+="--device=$i " + done + +else + DOCKER_DEVICES="" + WORKSPACE_VOLUMES="" +fi + + # Print arguments. echo "WORKSPACE: ${WORKSPACE}" echo "DOCKER CONTAINER NAME: ${DOCKER_IMAGE_NAME}" @@ -95,6 +116,8 @@ fi # and share the PID namespace (--pid=host) so the process inside does not have # pid 1 and SIGKILL is propagated to the process inside (jenkins can kill it). ${DOCKER_BINARY} run --rm --pid=host\ + ${DOCKER_DEVICES}\ + ${WORKSPACE_VOLUMES}\ -v ${WORKSPACE}:/workspace \ -v ${SCRIPT_DIR}:/docker \ "${EXTRA_MOUNTS[@]}" \ diff --git a/docker/install/ubuntu_install_python.sh b/docker/install/ubuntu_install_python.sh index c1f9d5081f57..58d72f327aa6 100755 --- a/docker/install/ubuntu_install_python.sh +++ b/docker/install/ubuntu_install_python.sh @@ -27,7 +27,7 @@ apt-get install -y python-dev # python 3.6 apt-get install -y software-properties-common -add-apt-repository ppa:deadsnakes/ppa +add-apt-repository -y ppa:deadsnakes/ppa apt-get update apt-get install -y python-pip python-dev python3.6 python3.6-dev diff --git a/docker/install/ubuntu_install_vai_core.sh b/docker/install/ubuntu_install_vai_core.sh new file mode 100644 index 000000000000..00189a7c308a --- /dev/null +++ b/docker/install/ubuntu_install_vai_core.sh @@ -0,0 +1,37 @@ +#!/bin/bash +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +set -e +set -u +set -o pipefail + +# install libraries for building Vitis-AI on ubuntu +apt-get update && apt-get install -y --no-install-recommends \ + build-essential\ + ca-certificates\ + cmake\ + sudo\ + wget\ + git\ + vim\ + graphviz\ + python-dev\ + gnupg2 + +apt-get update && apt-get install -y gcc-aarch64-linux-gnu + diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst new file mode 100755 index 000000000000..3e1db774e67c --- /dev/null +++ b/docs/deploy/vitis_ai.rst @@ -0,0 +1,617 @@ +Vitis-AI Integration +==================== + +`Vitis-AI `__ is Xilinx's +development stack for hardware-accelerated AI inference on Xilinx +platforms, including both edge devices and Alveo cards. It consists of +optimized IP, tools, libraries, models, and example designs. It is +designed with high efficiency and ease of use in mind, unleashing the +full potential of AI acceleration on Xilinx FPGA and ACAP. + +The current Vitis-AI Byoc flow inside TVM enables acceleration of Neural +Network model inference on edge and cloud. The identifiers for the +supported edge and cloud Deep Learning Processor Units (DPU's) are +DPUCZDX8G respectively DPUCADX8G. DPUCZDX8G and DPUCADX8G are hardware +accelerators for convolutional neural networks (CNN's) on top of the +Xilinx `Zynq Ultrascale+ +MPSoc `__ +respectively +`Alveo `__ +(U200/U250) platforms. For more information about the DPU identifiers +see the section on `DPU naming information <#dpu-naming-information>`__. + +On this page you will find information on how to +`build <#build-instructions>`__ TVM with Vitis-AI and on how to `get +started <#getting-started>`__ with an example. + +DPU naming information +---------------------- + ++---------------------------------+-----------------+-------------------------------------------------------------------------+------------------------------------------------------------+---------------------------------------------------+--------------------------------------------------------------------------+ +| DPU | Application | HW Platform | Quantization Method | Quantization Bitwidth | Design Target | ++=================================+=================+=========================================================================+============================================================+===================================================+==========================================================================+ +| Deep Learning Processing Unit | C: CNN R: RNN | AD: Alveo DDR AH: Alveo HBM VD: Versal DDR with AIE & PL ZD: Zynq DDR | X: DECENT I: Integer threshold F: Float threshold R: RNN | 4: 4-bit 8: 8-bit 16: 16-bit M: Mixed Precision | G: General purpose H: High throughput L: Low latency C: Cost optimized | ++---------------------------------+-----------------+-------------------------------------------------------------------------+------------------------------------------------------------+---------------------------------------------------+--------------------------------------------------------------------------+ + +Build instructions +------------------ + +This section lists the instructions for building TVM with Vitis-AI for +both `cloud <#cloud-dpucadx8g>`__ and `edge <#edge-dpuczdx8g>`__. + +Cloud (DPUCADX8G) +~~~~~~~~~~~~~~~~~ + +For Vitis-AI acceleration in the cloud TVM has to be built on top of the +Xilinx Alveo platform. + +System requirements +^^^^^^^^^^^^^^^^^^^ + +The following table lists system requirements for running docker +containers as well as Alveo cards. + ++-----------------------------------------------------+----------------------------------------------------------+ +| **Component** | **Requirement** | ++=====================================================+==========================================================+ +| Motherboard | PCI Express 3.0-compliant with one dual-width x16 slot | ++-----------------------------------------------------+----------------------------------------------------------+ +| System Power Supply | 225W | ++-----------------------------------------------------+----------------------------------------------------------+ +| Operating System | Ubuntu 16.04, 18.04 | ++-----------------------------------------------------+----------------------------------------------------------+ +| | CentOS 7.4, 7.5 | ++-----------------------------------------------------+----------------------------------------------------------+ +| | RHEL 7.4, 7.5 | ++-----------------------------------------------------+----------------------------------------------------------+ +| CPU | Intel i3/i5/i7/i9/Xeon 64-bit CPU | ++-----------------------------------------------------+----------------------------------------------------------+ +| GPU (Optional to accelerate quantization) | NVIDIA GPU with a compute capability > 3.0 | ++-----------------------------------------------------+----------------------------------------------------------+ +| CUDA Driver (Optional to accelerate quantization) | nvidia-410 | ++-----------------------------------------------------+----------------------------------------------------------+ +| FPGA | Xilinx Alveo U200 or U250 | ++-----------------------------------------------------+----------------------------------------------------------+ +| Docker Version | 19.03.1 | ++-----------------------------------------------------+----------------------------------------------------------+ + +Hardware setup and docker build +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +1. Clone the Vitis AI repository: + :: + + + git clone --recurse-submodules https://github.com/Xilinx/Vitis-AI + +2. Install Docker, and add the user to the docker group. Link the user + to docker installation instructions from the following docker's + website: + + - https://docs.docker.com/install/linux/docker-ce/ubuntu/ + - https://docs.docker.com/install/linux/docker-ce/centos/ + - https://docs.docker.com/install/linux/linux-postinstall/ + +3. Any GPU instructions will have to be separated from Vitis AI. +4. Set up Vitis AI to target Alveo cards. To target Alveo cards with + Vitis AI for machine learning workloads, you must install the + following software components: + + - Xilinx Runtime (XRT) + - Alveo Deployment Shells (DSAs) + - Xilinx Resource Manager (XRM) (xbutler) + - Xilinx Overlaybins (Accelerators to Dynamically Load - binary + programming files) + + While it is possible to install all of these software components + individually, a script has been provided to automatically install + them at once. To do so: + + - Run the following commands: + :: + + + cd Vitis-AI/alveo/packages + sudo su + ./install.sh + + - Power cycle the system. + +5. Clone tvm repo and pyxir repo + :: + + + git clone --recursive https://github.com/apache/incubator-tvm.git + git clone --recursive https://github.com/Xilinx/pyxir.git + +6. Build and start the tvm runtime Vitis-AI Docker Container. + :: + + + bash incubator-tvm/docker/build.sh ci_vai bash + bash incubator-tvm/docker/bash.sh tvm.ci_vai + + #Setup inside container + source /opt/xilinx/xrt/setup.sh + . $VAI_ROOT/conda/etc/profile.d/conda.sh + conda activate vitis-ai-tensorflow + +7. Install PyXIR + :: + + + + cd pyxir + python3 setup.py install --use_vai_rt_dpucadx8g --user + + +8. Build TVM inside the container with Vitis-AI + :: + + + cd incubator-tvm + mkdir build + cp cmake/config.cmake build + cd build + echo set\(USE_LLVM ON\) >> config.cmake + echo set\(USE_VITIS_AI ON\) >> config.cmake + cmake .. + make -j$(nproc) + +9. Install TVM + :: + cd incubator-tvm/python + pip3 install -e . --user + +Edge (DPUCZDX8G) +^^^^^^^^^^^^^^^^ + + +For edge deployment we make use of two systems referred to as host and +edge. The `host <#host-requirements>`__ system is responsible for +quantization and compilation of the neural network model in a first +offline step. Afterwards, the model will de deployed on the +`edge <#edge-requirements>`__ system. + +Host requirements +^^^^^^^^^^^^^^^^^ + +The following table lists system requirements for running the TVM - +Vitis-AI docker container. + ++-----------------------------------------------------+----------------------------------------------+ +| **Component** | **Requirement** | ++=====================================================+==============================================+ +| Operating System | Ubuntu 16.04, 18.04 | ++-----------------------------------------------------+----------------------------------------------+ +| | CentOS 7.4, 7.5 | ++-----------------------------------------------------+----------------------------------------------+ +| | RHEL 7.4, 7.5 | ++-----------------------------------------------------+----------------------------------------------+ +| CPU | Intel i3/i5/i7/i9/Xeon 64-bit CPU | ++-----------------------------------------------------+----------------------------------------------+ +| GPU (Optional to accelerate quantization) | NVIDIA GPU with a compute capability > 3.0 | ++-----------------------------------------------------+----------------------------------------------+ +| CUDA Driver (Optional to accelerate quantization) | nvidia-410 | ++-----------------------------------------------------+----------------------------------------------+ +| FPGA | Not necessary on host | ++-----------------------------------------------------+----------------------------------------------+ +| Docker Version | 19.03.1 | ++-----------------------------------------------------+----------------------------------------------+ + +Host setup and docker build +^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +1. Clone tvm repo +:: + git clone --recursive https://github.com/apache/incubator-tvm.git +2. Build and start the tvm runtime Vitis-AI Docker Container. +:: + cd incubator-tvm + bash incubator-tvm/docker/build.sh ci_vai bash + bash incubator-tvm/docker/bash.sh tvm.ci_vai + + #Setup inside container + . $VAI_ROOT/conda/etc/profile.d/conda.sh + conda activate vitis-ai-tensorflow + +3. Install PyXIR +:: + + + git clone --recursive https://github.com/Xilinx/pyxir.git + cd pyxir + python3 setup.py install --user + + +4. Build TVM inside the container with Vitis-AI. +:: + cd incubator-tvm + mkdir build + cp cmake/config.cmake build + cd build + echo set\(USE_LLVM ON\) >> config.cmake + echo set\(USE_VITIS_AI ON\) >> config.cmake + cmake .. + make -j$(nproc) + +5. Install TVM +:: + cd incubator-tvm/python + pip3 install -e . --user + +Edge requirements +^^^^^^^^^^^^^^^^^ + +The DPUCZDX8G can be deployed on the `Zynq Ultrascale+ +MPSoc `__ +platform. The following development boards can be used out-of-the-box: + ++--------------------+----------------------+-----------------------------------------------------------------------+ +| **Target board** | **TVM identifier** | **Info** | ++====================+======================+=======================================================================+ +| Ultra96 | DPUCZDX8G-ultra96 | https://www.xilinx.com/products/boards-and-kits/1-vad4rl.html | ++--------------------+----------------------+-----------------------------------------------------------------------+ +| ZCU104 | DPUCZDX8G-zcu104 | https://www.xilinx.com/products/boards-and-kits/zcu104.html | ++--------------------+----------------------+-----------------------------------------------------------------------+ +| ZCU102 | DPUCZDX8G-zcu102 | https://www.xilinx.com/products/boards-and-kits/ek-u1-zcu102-g.html | ++--------------------+----------------------+-----------------------------------------------------------------------+ + +Edge hardware setup +^^^^^^^^^^^^^^^^^^^ + +NOTE: This section provides instructions setting up with the `Pynq `__ platform but Petalinux based flows are also supported. + +1. Download the Pynq v2.5 image for your target (use Z1 or Z2 for + Ultra96 target depending on board version) Link to image: + https://github.com/Xilinx/PYNQ/releases/tag/v2.5 +2. Follow Pynq instructions for setting up the board: `pynq + setup `__ +3. After connecting to the board, make sure to run as root. Execute + ``su`` +4. Set up DPU on Pynq by following the steps here: `DPU Pynq + setup `__ +5. Run the following command to download the DPU bitstream: + + :: + + + python3 -c 'from pynq_dpu import DpuOverlay ; overlay = DpuOverlay("dpu.bit")' + +6. Check whether the DPU kernel is alive: + :: + + + dexplorer -w + +Edge TVM setup +^^^^^^^^^^^^^^ + +NOTE: When working on Petalinux instead of Pynq, the following steps might take more manual work (e.g building hdf5 from source). Also, TVM has a scipy dependency which you then might have to build from source or circumvent. We don't depend on scipy in our flow. + +Building TVM depends on the Xilinx +`PyXIR `__ package. PyXIR acts as an +interface between TVM and Vitis-AI tools. + +1. First install the PyXIR h5py and pydot dependencies: +:: + + + apt-get install libhdf5-dev + pip3 install pydot h5py +2. Install PyXIR +:: + + + git clone --recursive https://github.com/Xilinx/pyxir.git + cd pyxir + sudo python3 setup.py install --use_vai_rt_dpuczdx8g + +3. Build TVM with Vitis-AI +:: + + + git clone --recursive https://github.com/apache/incubator-tvm + cd incubator-tvm + mkdir build + cp cmake/config.cmake build + cd build + echo set\(USE_VITIS_AI ON\) >> config.cmake + cmake .. + make + +4. Install TVM +:: + cd incubator-tvm/python + pip3 install -e . --user + +5. Check whether the setup was successful in the Python shell: +:: + + + python3 -c 'import pyxir; import tvm' + + +Getting started +--------------- + +This section shows how to use TVM with Vitis-AI. For this it's important +to understand that neural network models are quantized for Vitis-AI +execution in fixed point arithmetic. The approach we take here is to +quantize on-the-fly using the first N inputs as explained in the next +section. + +On-the-fly quantization +~~~~~~~~~~~~~~~~~~~~~~~ + +Usually, to be able to accelerate inference of Neural Network models +with Vitis-AI DPU accelerators, those models need to quantized upfront. +In TVM - Vitis-AI flow, we make use of on-the-fly quantization to remove +this additional preprocessing step. In this flow, one doesn't need to +quantize his/her model upfront but can make use of the typical inference +execution calls (module.run) to quantize the model on-the-fly using the +first N inputs that are provided (see more information below). This will +set up and calibrate the Vitis-AI DPU and from that point onwards +inference will be accelerated for all next inputs. Note that the edge +flow deviates slightly from the explained flow in that inference won't +be accelerated after the first N inputs but the model will have been +quantized and compiled and can be moved to the edge device for +deployment. Please check out the `edge <#Edge%20usage>`__ usage +instructions below for more information. + +Config/Settings +~~~~~~~~~~~~~~~ + +A couple of environment variables can be used to customize the Vitis-AI +Byoc flow. + ++----------------------------+----------------------------------------+--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------+ +| **Environment Variable** | **Default if unset** | **Explanation** | ++============================+========================================+============================================================================================================================================================================================================================================================================================================================================+ +| PX\_QUANT\_SIZE | 128 | The number of inputs that will be used for quantization (necessary for Vitis-AI acceleration) | ++----------------------------+----------------------------------------+--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------+ +| PX\_BUILD\_DIR | Use the on-the-fly quantization flow | Loads the quantization and compilation information from the provided build directory and immediately starts Vitis-AI hardware acceleration. This configuration can be used if the model has been executed before using on-the-fly quantization during which the quantization and comilation information was cached in a build directory. | ++----------------------------+----------------------------------------+--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------+ + +Cloud usage +~~~~~~~~~~~ + +This section shows how to accelerate a convolutional neural network +model in TVM with Vitis-AI on the cloud. + +To be able to target the Vitis-AI cloud DPUCADX8G target we first have +to import the target in PyXIR. This PyXIR package is the interface being +used by TVM to integrate with the Vitis-AI stack. Additionaly, import +the typical TVM and Relay modules and the Vitis-AI contrib module inside +TVM. + +:: + + import pyxir + import pyxir.contrib.target.DPUCADX8G + + import tvm + import tvm.relay as relay + from tvm.contrib.target import vitis_ai + from tvm.relay.build_module import bind_params_by_name + from tvm.relay.op.contrib.vitis_ai import annotation + +After importing a convolutional neural network model using the usual +Relay API's, annotate the Relay expression for the given Vitis-AI DPU +target and partition the graph. + +:: + + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, target) + mod = relay.transform.MergeCompilerRegions()(mod) + mod = relay.transform.PartitionGraph()(mod) + +Now, we can build the TVM runtime library for executing the model. The +TVM target is 'llvm' as the operations that can't be handled by the DPU +are executed on the CPU. The Vitis-AI target is DPUCADX8G as we are +targeting the cloud DPU and this target is passed as a config to the TVM +build call. + +:: + + tvm_target = 'llvm' + target='DPUCADX8G' + + with tvm.transform.PassContext(opt_level=3, config= {'target_': target}): + graph, lib, params = relay.build(mod, tvm_target, params=params) + +As one more step before we can accelerate a model with Vitis-AI in TVM +we have to quantize and compile the model for execution on the DPU. We +make use of on-the-fly quantization for this. Using this method one +doesn’t need to quantize their model upfront and can make use of the +typical inference execution calls (module.run) to calibrate the model +on-the-fly using the first N inputs that are provided. After the first N +iterations, computations will be accelerated on the DPU. So now we will +feed N inputs to the TVM runtime module. Note that these first N inputs +will take a substantial amount of time. + +:: + + module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) + module.set_input(**params) + + # First N (default = 128) inputs are used for quantization calibration and will + # be executed on the CPU + # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) + for i in range(128): + module.set_input(input_name, inputs[i]) + module.run() + +Afterwards, inference will be accelerated on the DPU. + +:: + + module.set_input(name, data) + module.run() + +To save and load the built module, one can use the typical TVM API's: + +:: + + # save the graph, lib and params into separate files + from tvm.contrib import util + + temp = util.tempdir() + path_lib = temp.relpath("deploy_lib.so") + lib.export_library(path_lib) + with open(temp.relpath("deploy_graph.json"), "w") as fo: + fo.write(graph) + with open(temp.relpath("deploy_param.params"), "wb") as fo: + fo.write(relay.save_param_dict(params)) + +Load the module from compiled files and run inference + +:: + + # load the module into memory + loaded_json = open(temp.relpath("deploy_graph.json")).read() + loaded_lib = tvm.runtime.load_module(path_lib) + loaded_params = bytearray(open(temp.relpath("deploy_param.params"), "rb").read()) + + module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) + module.load_params(loaded_params) + module.set_input(name, data) + module.run() + +Edge usage +~~~~~~~~~~ + +This section shows how to accelerate a convolutional neural network +model in TVM with Vitis-AI at the edge. The first couple of steps will +have to be run on the host machine and take care of quantization and +compilation for deployment at the edge. + +Host steps +^^^^^^^^^^ + +To be able to target the Vitis-AI cloud DPUCZDX8G target we first have +to import the target in PyXIR. This PyXIR package is the interface being +used by TVM to integrate with the Vitis-AI stack. Additionaly, import +the typical TVM and Relay modules and the Vitis-AI contrib module inside +TVM. + +:: + + import pyxir + import pyxir.contrib.target.DPUCZDX8G + + import tvm + import tvm.relay as relay + from tvm.contrib.target import vitis_ai + from tvm.relay.build_module import bind_params_by_name + from tvm.relay.op.contrib.vitis_ai import annotation + +After importing a convolutional neural network model using the usual +Relay API's, annotate the Relay expression for the given Vitis-AI DPU +target and partition the graph. + +:: + + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, target) + mod = relay.transform.MergeCompilerRegions()(mod) + mod = relay.transform.PartitionGraph()(mod) + +Now, we can build the TVM runtime library for executing the model. The +TVM target is 'llvm' as the operations that can't be handled by the DPU +are executed on the CPU. At this point that means the CPU on the host. +The Vitis-AI target is DPUCZDX8G-zcu104 as we are targeting the edge DPU +on the ZCU104 board and this target is passed as a config to the TVM +build call. Note that different identifiers can be passed for different +targets, see `edge targets info <#edge-requirements>`__. + +:: + + tvm_target = 'llvm' + target='DPUCZDX8G-zcu104' + + with tvm.transform.PassContext(opt_level=3, config= {'target_': target}): + graph, lib, params = relay.build(mod, tvm_target, params=params) + +Additionaly, already build the deployment module for the ARM CPU target +and serialize: + +:: + + # Export lib for aarch64 target + + tvm_target = tvm.target.arm_cpu('ultra96') + lib_kwargs = { + 'fcompile': contrib.cc.create_shared, + 'cc': "/usr/aarch64-linux-gnu/bin/ld" + } + + with tvm.transform.PassContext(opt_level=3, + config={'target_': target, + 'vai_build_dir_': target + '_build'}): + graph_arm, lib_arm, params_arm = relay.build( + mod, tvm_target, params=params) + + lib_dpuv2.export_library('tvm_dpu_arm.so', **lib_kwargs) + with open("tvm_dpu_arm.json","w") as f: + f.write(graph_dpuv2) + with open("tvm_dpu_arm.params", "wb") as f: + f.write(relay.save_param_dict(params_dpuv2)) + +As one more step before we can deploy a model with Vitis-AI in TVM at +the edge we have to quantize and compile the model for execution on the +DPU. We make use of on-the-fly quantization on the host machine for +this. This involves using the TVM inference calls (module.run) to +quantize the model on the host using N inputs. After providing N inputs +we can then move the TVM and Vitis-AI build files to the edge device for +deployment. + +:: + + module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) + module.set_input(**params) + + # First N (default = 128) inputs are used for quantization calibration and will + # be executed on the CPU + # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) + for i in range(128): + module.set_input(input_name, inputs[i]) + module.run() + +Now, move the TVM build files (tvm\_dpu\_arm.json, tvm\_dpu\_arm.so, +tvm\_dpu\_arm.params) and the DPU build directory (e.g. +DPUCZDX8G-zcu104\_build) to the edge device. For information on setting +up the edge device check out the `edge setup <#edge-dpuczdx8g>`__ +section. + +Edge steps +^^^^^^^^^^ + +The following steps will have to be executed on the edge device after +setup and moving the build files from the host. + +Move the target build directory to the same folder where the example +running script is located and explicitly set the path to the build +directory using the PX\_BUILD\_DIR environment variable. + +:: + + export PX_BUILD_DIR={PATH-TO-DPUCZDX8G-BUILD_DIR} + +Then load the TVM runtime module into memory and feed inputs for +inference. + +:: + + # load the module into memory + loaded_json = open(temp.relpath("tvm_dpu_arm.json")).read() + loaded_lib = tvm.runtime.load_module("tvm_dpu_arm.so") + loaded_params = bytearray(open(temp.relpath("tvm_dpu_arm.params"), "rb").read()) + + module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) + module.load_params(loaded_params) + module.set_input(name, data) + module.run() + + diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py new file mode 100644 index 000000000000..932db232f8c8 --- /dev/null +++ b/python/tvm/contrib/target/vitis_ai.py @@ -0,0 +1,109 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=invalid-name, unused-argument, import-outside-toplevel +"""Utility to compile VITISAI models""" + +import os + +from tvm.relay.expr import Tuple, Call +import tvm._ffi + +import pyxir +import pyxir.frontend.tvm + +from .. import vitis_ai_runtime + +class CodegenVitisAI: + """ + Traverse subgraphs and build XGraph + """ + def __init__(self, model_name, function): + + self.model_name = model_name + self.function = function + self.params = {} + + + + def convert_pyxir(self, target): + """ + Convert relay submodule expression to PYXIR(XGRAPH) + """ + xgraph = pyxir.frontend.tvm.from_relay(self.function, + params=self.params, postprocessing=None) + xgraph = pyxir.partition(xgraph, targets=[target]) + return xgraph + + def get_output_names(self): + """ + Get output names from subgraph + """ + func = self.function + output_relay_ids = [] + expr = func.body + if isinstance(expr, Tuple): + for field in expr.fields: + output_relay_ids.append(hash(field)) + elif isinstance(expr, Call): + output_relay_ids.append(hash(expr)) + else: + raise ValueError("does not support {}".format(type(expr))) + return output_relay_ids + +@tvm._ffi.register_func("relay.ext.vai") +def vai_compiler(ref): + """ + Create a VAI runtime from a Relay module. + """ + assert isinstance(ref, tvm.relay.function.Function) + + model_dir = os.getcwd() + out_tensor_names = [] + name = str(ref.attrs.global_symbol) + + pass_context = tvm.get_global_func("transform.GetCurrentPassContext")() + target = str(pass_context.config['target_']) + vai_build_dir = str(pass_context.config['vai_build_dir_']) \ + if 'vai_build_dir_' in pass_context.config else None + if vai_build_dir and not os.path.exists(vai_build_dir): + raise ValueError("Provided Vitis-AI build dir: `{}` could not be found" + .format(vai_build_dir)) + if not vai_build_dir: + builder = CodegenVitisAI(name, ref) + model_dir = target + "_build/" + xgraph = builder.convert_pyxir(target) + output_relay_ids = builder.get_output_names() + layers = xgraph.get_layers() + # get the output tensor names using xgraph and output relay ids + out_tensor_names = [] + for layer in layers: + if not layer.internal: + if layer.attrs['relay_id'][0] in output_relay_ids: + out_tensor_names.append(layer.name) + if len(out_tensor_names) == 0: + raise ValueError("During codegeneration the loading of subexpression \ + failed due to output tensorname mismatch in relay pyxir interface.") + + # Save/serialize XGraph + if not os.path.exists(model_dir): + os.mkdir(model_dir) + xgraph.meta_attrs['tvm_out_tensors'] = out_tensor_names + pyxir.graph.io.xgraph_io.XGraphIO.save(xgraph, model_dir + 'dpu_xgraph') + else: + model_dir = vai_build_dir + + return vitis_ai_runtime.create(name, model_dir, target).module diff --git a/python/tvm/contrib/vitis_ai_runtime.py b/python/tvm/contrib/vitis_ai_runtime.py new file mode 100644 index 000000000000..a03120b301b8 --- /dev/null +++ b/python/tvm/contrib/vitis_ai_runtime.py @@ -0,0 +1,54 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +"""VitisAI runtime that load and run Xgraph.""" +import tvm._ffi + +def create(name, model_dir, target): + """Create a runtime executor module given a xgraph model and context. + Parameters + ---------- + model_dir : str + The directory where the compiled models are located. + target : str + The target for running subgraph. + + Returns + ------- + vai_runtime : VaiModule + Runtime Vai module that can be used to execute xgraph model. + """ + runtime_func = "tvm.vitis_ai_runtime.create" + fcreate = tvm._ffi.get_global_func(runtime_func) + return VitisAIModule(fcreate(name, model_dir, target)) + +class VitisAIModule(object): + """Wrapper runtime module. + + This is a thin wrapper of the underlying TVM module. + you can also directly call set_input, run, and get_output + of underlying module functions + + Parameters + ---------- + module : Module + The internal tvm module that holds the actual vai functions. + + """ + + def __init__(self, module): + self.module = module diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py new file mode 100644 index 000000000000..8734daa39750 --- /dev/null +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -0,0 +1,92 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=invalid-name, unused-argument, no-else-return, E1102 +"""VITISAI codegen supported operators.""" + +import numpy as np + +from tvm import relay +import tvm._ffi +from tvm.relay.expr import Tuple, TupleGetItem +from tvm.relay import transform +from tvm.relay.op.annotation import compiler_begin, compiler_end + +import pyxir +import pyxir.frontend.tvm + + +@transform.function_pass(opt_level=0) +class VitisAIAnnotationPass: + """The explicit pass wrapper around VitisAIAnnotationPass.""" + def __init__(self, compiler, relay_ids): + self.compiler = compiler + self.relay_ids = relay_ids + def transform_function(self, func, mod, ctx): + """Transform func to annotate.""" + annotator = self + class Annotator(tvm.relay.ExprMutator): + """Annotator for VITIS-AI DPU.""" + def visit_tuple(self, tup): + field_list = [] + cond = int(hash(tup)) + for field in tup.fields: + if cond in annotator.relay_ids: + field_list.append(compiler_begin(super().visit(field), annotator.compiler)) + else: + field_list.append(super().visit(field)) + if cond in annotator.relay_ids: + return compiler_end(Tuple(field_list), annotator.compiler) + else: + return Tuple(field_list) + + def visit_tuple_getitem(self, op): + if int(hash(op.tuple_value)) in annotator.relay_ids: + tuple_value = compiler_begin(super().visit(op.tuple_value), + annotator.compiler) + return compiler_end(TupleGetItem(tuple_value, op.index), annotator.compiler) + else: + tuple_value = super().visit(op.tuple_value) + return TupleGetItem(tuple_value, op.index) + def visit_call(self, call): + if int(hash(call)) in annotator.relay_ids: + new_args = [] + for arg in call.args: + ann = compiler_begin(super().visit(arg), + annotator.compiler) + new_args.append(ann) + new_call = relay.Call(call.op, new_args, call.attrs, + call.type_args) + return compiler_end(new_call, annotator.compiler) + + else: + return super().visit_call(call) + return Annotator().visit(func) + + + +def annotation(mod, params, target): + """ + An annotator for VITISAI. + """ + xgraph = pyxir.frontend.tvm.from_relay(mod, params, postprocessing=None) + xgraph = pyxir.partition(xgraph, targets=[target]) + layers = xgraph.get_layers() + relay_ids = [list(np.array(layer.attrs['relay_id']).flatten()) + for layer in layers if layer.target == target] + relay_ids_flatten = [item for sublist in relay_ids for item in sublist] + mod = VitisAIAnnotationPass("vai", relay_ids_flatten)(mod) + return mod diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc new file mode 100755 index 000000000000..a042b260d653 --- /dev/null +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -0,0 +1,147 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +/*! + * \file vitis_ai_runtime.cc + */ +#include +#include + +#include "vitis_ai_runtime.h" + +namespace tvm { +namespace runtime { + +TVM_REGISTER_PASS_CONFIG_OPTION("target_", String); +TVM_REGISTER_PASS_CONFIG_OPTION("vai_build_dir_", String); + +std::shared_ptr load_xgraph_model(const std::string& model_path) { + std::string model_name = model_path + "/" + "dpu_xgraph.json"; + std::string model_weights = model_path + "/" + "dpu_xgraph.h5"; + return pyxir::load(model_name, model_weights); +} + +void VitisAIRuntime::Init(const std::string& model_path, const std::string& target) { + model_path_ = model_path; + target_ = target; + xgraph_ = load_xgraph_model(model_path_); + in_tensor_names_ = xgraph_->get_input_names(); + out_tensor_names_ = xgraph_->get_meta_attr("tvm_out_tensors").get_strings(); + pyxir::partition(xgraph_, std::vector{target}, ""); + pyxir::RunOptionsHolder run_options(new pyxir::runtime::RunOptions()); + run_options->on_the_fly_quantization = true; + rt_mod_ = pyxir::build_rt(xgraph_, target_ , in_tensor_names_, out_tensor_names_, + "vai", run_options); +} + + +Module VitisAIRuntimeCreate(const std::string& name, + const std::string& model_path, + const std::string& target) { + Array const_vars; + auto exec = make_object(name, const_vars); + exec->Init(model_path, target); + return Module(exec); +} + + + +TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.create").set_body([](TVMArgs args, TVMRetValue* rv) { + *rv = VitisAIRuntimeCreate(args[0], args[1], args[2]); +}); + +Module VitisAIRuntimeLoadFromBinary(void* strm ) { + dmlc::Stream* stream = static_cast(strm); + + std::string model_path; + std::string symbol_name; + std::vector const_vars; + std::string target; + stream->Read(&model_path); + stream->Read(&target); + stream->Read(&symbol_name); + stream->Read(&const_vars); + Array const_names; + for (const auto& it : const_vars) { + const_names.push_back(it); + } + auto exec = make_object(symbol_name, const_names); + exec->Init(model_path, target); + return Module(exec); + } + +TVM_REGISTER_GLOBAL("runtime.module.loadbinary_VitisAIRuntime") + .set_body_typed(VitisAIRuntimeLoadFromBinary); + +void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { + stream->Write(this-> model_path_); + stream->Write(this-> target_); + stream->Write(this->symbol_name_); + std::vector consts; + for (const auto& it : const_names_) { + consts.push_back(it); + } + stream->Write(consts); + } + + +PackedFunc VitisAIRuntime::GetFunction(const std::string& name, + const ObjectPtr& sptr_to_self) { + if (name == "get_symbol") { + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->symbol_name_; }); + } else if (name == "get_const_vars") { + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->const_names_; }); + } else if ("__init_" + this->symbol_name_ == name) { + // The function to initialize constant tensors. + return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { + CHECK_EQ(args.size(), 1U); + this->initialized_ = true; + *rv = 0; + }); + } else { + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { + DLTensor* inputs = args[0]; + std::vector in_shape; + for (int i = 0; i < inputs->ndim; ++i) + in_shape.push_back(inputs->shape[i]); + pyxir::XBufferHolder xb_in = std::shared_ptr( + new pyxir::XBuffer(reinterpret_cast(static_cast(inputs->data)), 4, + "f", in_shape.size(), in_shape, false, false)); + std::vector out_tensors; + for (unsigned i = 0; i < out_tensor_names_.size(); ++i) { + DLTensor* output_tensor = args[args.size() - out_tensor_names_.size()+i]; + std::vector out_shape; + for (int i = 0; i < output_tensor->ndim; ++i) + out_shape.push_back(output_tensor->shape[i]); + void* output_data = reinterpret_cast (static_cast(output_tensor->data)); + out_tensors.push_back(std::shared_ptr( + new pyxir::XBuffer(output_data, 4, "f", out_shape.size(), out_shape, + false, false))); + } + std::vector in_tensors{xb_in}; + // Execute the subgraph. + rt_mod_->execute(in_tensors, out_tensors); + }); + } + } +} // namespace runtime +} // namespace tvm diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h new file mode 100755 index 000000000000..7ff69a456612 --- /dev/null +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h @@ -0,0 +1,97 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +/*! + * \brief Vitis-AI runtime that can run model + * containing only tvm PackedFunc. + * \file vitis_ai_runtime.h + */ +#ifndef TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ +#define TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ + + +#include +#include +#include + +#include +#include +#include + +#include +#include + + + +namespace tvm { +namespace runtime { + +/*! + * \brief VAI runtime. + * + * This runtime can be accessed in various language via + * TVM runtime PackedFunc API. + */ +class VitisAIRuntime : public ModuleNode { + public: +VitisAIRuntime(const std::string& symbol_name, const Array const_names): + symbol_name_(symbol_name), const_names_(const_names) {} + /*! + * \brief Get member function to front-end. + * \param name The name of the function. + * \param sptr_to_self The pointer to the module node. + * \return The corresponding member function. + */ + virtual PackedFunc GetFunction(const std::string& name, const ObjectPtr& sptr_to_self); + /*! + * \return The type key of the executor. + */ + const char* type_key() const { return "VitisAIRuntime"; } + + /*! + * \brief Initialize the vai runtime with pyxir. + * \param model_path The compiled model path. + * \param target The name of the target being used + */ +void Init(const std::string& model_path, const std::string& target); + /*! + * \brief Serialize the content of the pyxir directory and save it to + * binary stream. + * \param stream The binary stream to save to. + */ + void SaveToBinary(dmlc::Stream* stream) final; + + private: + /*! \brief The only subgraph name for this module. */ + std::string symbol_name_; + /*! \brief The graph. */ + std::string graph_json_; + /*! \brief The required constant names. */ + Array const_names_; + std::shared_ptr xgraph_; + pyxir::RtModHolder rt_mod_; + std::string model_path_; + std::string target_; + std::vector in_tensor_names_; + std::vector out_tensor_names_; + bool initialized_{false}; +}; +} // namespace runtime +} // namespace tvm +#endif // TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ diff --git a/tests/python/contrib/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai_codegen.py new file mode 100644 index 000000000000..673047f113a2 --- /dev/null +++ b/tests/python/contrib/test_vitis_ai_codegen.py @@ -0,0 +1,203 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611 +"""Vitis-AI codegen tests.""" + +import numpy as np + +import tvm +from tvm import relay +from tvm.relay import transform +from tvm.relay.op.contrib.vitis_ai import annotation +from tvm.contrib.target import vitis_ai + +import pyxir +import pyxir.contrib.target.DPUCADX8G + +def set_func_attr(func, compile_name, symbol_name): + func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) + func = func.with_attr("Inline", tvm.tir.IntImm("int32", 1)) + func = func.with_attr("Compiler", compile_name) + func = func.with_attr("global_symbol", symbol_name) + return func + +def _create_graph(): + shape = (10, 10) + mod = tvm.IRModule() + x = relay.var('x', shape=shape) + y = relay.var('y', shape=shape) + z = x + x + p = y * y + func = relay.Function([x, y], p - z) + mod["main"] = func + params = {} + params["x"] = np.random.rand(10, 10).astype('float32') + params["y"] = np.random.rand(10, 10).astype('float32') + return mod, params + + +def _construct_model(func, params=None): + mod = tvm.IRModule() + mod["main"] = func + if params is None: + params = {} + mod = annotation(mod, params, "DPUCADX8G") + mod = transform.MergeCompilerRegions()(mod) + mod = transform.PartitionGraph()(mod) + fcompile = tvm._ffi.get_global_func("relay.ext.vai") + subgraph_mod = tvm.IRModule() + for _, funcnode in mod.functions.items(): + if funcnode.attrs and 'Compiler' in funcnode.attrs and \ + funcnode.attrs['Compiler'] == 'vai': + subgraph_mod["main"] = funcnode + with tvm.transform.PassContext(opt_level=3, config={'target_':'DPUCADX8G'}): + fcompile(subgraph_mod["main"]) + + +def test_add(): + shape = (10, 10) + x = relay.var('x', shape=shape) + y = x + x + func = relay.Function([x], y) + _construct_model(func) + +def test_relu(): + shape = (10, 10) + x = relay.var('x', shape=shape) + y = relay.nn.relu(x) + func = relay.Function([x], y) + _construct_model(func) + +def test_conv2d(): + x = relay.var('x', shape=(1, 3, 224, 224)) + w = relay.const(np.zeros((16, 3, 3, 3), dtype='float32')) + y = relay.nn.conv2d(x, w, strides=[2, 2], padding=[1, 1, 1, 1], kernel_size=[3, 3]) + func = relay.Function([x], y) + params = {} + params["x"] = np.zeros((16, 3, 3, 3), dtype='float32') + _construct_model(func, params) + + +def test_global_avg_pool2d(): + shape = (10, 10, 10, 10) + x = relay.var('x', shape=shape) + y = relay.nn.global_avg_pool2d(x) + func = relay.Function([x], y) + _construct_model(func) + +def test_annotate(): + """Test annotation with Vitis-AI DP (DPUCADX8G)""" + def partition(): + data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) + weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) + + conv = relay.nn.conv2d( + data=data, + weight=weight, + kernel_size=(3, 3), + channels=16, + padding=(1, 1)) + bn_output = relay.nn.batch_norm(conv, bn_gamma, bn_beta, bn_mmean, + bn_mvar) + + func = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, + bn_mvar], bn_output.astuple()) + mod = tvm.IRModule() + mod["main"] = func + params = {} + params["weight"] = np.random.rand(16, 3, 3, 3).astype('float32') + params["bn_gamma"] = np.random.rand(16).astype('float32') + params["bn_beta"] = np.random.rand(16).astype('float32') + params["bn_mean"] = np.random.rand(16).astype('float32') + params["bn_var"] = np.random.rand(16).astype('float32') + mod = annotation(mod, params, "DPUCADX8G") + + opt_pass = tvm.transform.Sequential([ + transform.InferType(), + transform.PartitionGraph(), + transform.SimplifyInference(), + transform.FoldConstant(), + transform.AlterOpLayout(), + ]) + + with tvm.transform.PassContext(opt_level=3): + mod = opt_pass(mod) + + return mod + + def expected(): + # function for batch_norm + data0 = relay.var("data0", relay.TensorType((1, 16, 224, 224), + "float32")) + mod = tvm.IRModule() + bn_gamma = relay.var("bn_gamma1", relay.TensorType((16, ), "float32")) + bn_beta = relay.var("bn_beta1", relay.TensorType((16, ), "float32")) + bn_mmean = relay.var("bn_mean1", relay.TensorType((16, ), "float32")) + bn_mvar = relay.var("bn_var1", relay.TensorType((16, ), "float32")) + + bn = relay.nn.batch_norm(data0, bn_gamma, bn_beta, bn_mmean, bn_mvar) + func0 = relay.Function([data0, bn_gamma, bn_beta, bn_mmean, bn_mvar], + bn.astuple()) + func0 = set_func_attr(func0, "vai", "vai_2") + gv0 = relay.GlobalVar("vai_2") + mod[gv0] = func0 + + # function for conv2d + data1 = relay.var("data1", relay.TensorType((1, 3, 224, 224), "float32")) + weight1 = relay.var("weight1", relay.TensorType((16, 3, 3, 3), "float32")) + conv = relay.nn.conv2d( + data=data1, + weight=weight1, + kernel_size=(3, 3), + channels=16, + padding=(1, 1)) + func1 = relay.Function([data1, weight1], conv) + func1 = set_func_attr(func1, "vai", "vai_0") + gv1 = relay.GlobalVar("vai_0") + mod[gv1] = func1 + + # main function + data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) + weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) + bn_gamma0 = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) + bn_beta0 = relay.var("bn_beta", relay.TensorType((16, ), "float32")) + bn_mmean0 = relay.var("bn_mean", relay.TensorType((16, ), "float32")) + bn_mvar0 = relay.var("bn_var", relay.TensorType((16, ), "float32")) + + call1 = gv1(data, weight) + call0 = gv0(call1, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0) + mod["main"] = relay.Function([data, weight, bn_gamma0, bn_beta0, bn_mmean0, + bn_mvar0], call0) + mod = transform.InferType()(mod) + return mod + + partitioned = partition() + ref_mod = expected() + + assert tvm.ir.structural_equal(partitioned, ref_mod, map_free_vars=True) + + +if __name__ == "__main__": + test_annotate() + test_add() + test_relu() + test_conv2d() + test_global_avg_pool2d() diff --git a/tests/python/contrib/test_vitis_ai_runtime.py b/tests/python/contrib/test_vitis_ai_runtime.py new file mode 100644 index 000000000000..e3630a58ed94 --- /dev/null +++ b/tests/python/contrib/test_vitis_ai_runtime.py @@ -0,0 +1,104 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611 + +""" Vitis-AI runtime test """ + +import sys +import numpy as np + +import pyxir +import pyxir.contrib.target.DPUCADX8G + +import tvm +import tvm.relay.testing +from tvm import relay +from tvm import runtime +from tvm.relay import transform +from tvm.contrib import util +from tvm.relay.backend import compile_engine +from tvm.relay.build_module import bind_params_by_name +from tvm.relay.op.contrib.vitis_ai import annotation +from tvm.contrib.target import vitis_ai + + + +def check_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", + ctx=tvm.cpu(), params=None): + """Check the result between reference and generated output with vitis-ai byoc flow""" + if sys.platform == "win32": + print("Skip test on Windows for now") + return + + def update_lib(lib): + tmp_path = util.tempdir() + lib_name = 'lib.so' + lib_path = tmp_path.relpath(lib_name) + lib.export_library(lib_path) + lib = runtime.load_module(lib_path) + + return lib + + def check_graph_runtime_result(): + compile_engine.get().clear() + with tvm.transform.PassContext(opt_level=3, config={'target_' : 'DPUCADX8G'}): + json, lib, param = relay.build(mod, target=target, params=params) + lib = update_lib(lib) + rt_mod = tvm.contrib.graph_runtime.create(json, lib, ctx) + + for name, data in map_inputs.items(): + rt_mod.set_input(name, data) + rt_mod.set_input(**param) + rt_mod.run() + + out_shapes = out_shape if isinstance(out_shape, list) else [out_shape] + results = result if isinstance(result, list) else [result] + + for idx, shape in enumerate(out_shapes): + out = tvm.nd.empty(shape, ctx=ctx) + out = rt_mod.get_output(idx, out) + + tvm.testing.assert_allclose(out.asnumpy(), results[idx], rtol=tol, atol=tol) + + check_graph_runtime_result() + + +def test_extern_vai_resnet18(): + """Test resnet18 model using Vitis-AI byoc flow""" + if not tvm.get_global_func("relay.ext.vai", True): + print("skip because VITIS-AI codegen is not available") + return + + dtype = 'float32' + ishape = (1, 3, 224, 224) + + mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, "DPUCADX8G") + mod = transform.MergeCompilerRegions()(mod) + mod = transform.PartitionGraph()(mod) + + ref_mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) + ref_ex = relay.create_executor("graph", mod=ref_mod, ctx=tvm.cpu(0)) + i_data = np.random.uniform(0, 1, ishape).astype(dtype) + + ref_res = ref_ex.evaluate()(i_data, **params) + + check_result(mod, {"data": i_data}, + (1, 1000), ref_res.asnumpy(), tol=1e-5, params=params) +if __name__ == "__main__": + test_extern_vai_resnet18() From 5eebecb22e9a09760838ed01ba8a995ba5348518 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Mon, 31 Aug 2020 23:31:23 -0600 Subject: [PATCH 02/22] Remove environment related files --- docker/Dockerfile.ci_vai | 55 --------- docker/bash.sh | 136 ---------------------- docker/install/ubuntu_install_python.sh | 40 ------- docker/install/ubuntu_install_vai_core.sh | 37 ------ 4 files changed, 268 deletions(-) delete mode 100644 docker/Dockerfile.ci_vai delete mode 100755 docker/bash.sh delete mode 100755 docker/install/ubuntu_install_python.sh delete mode 100644 docker/install/ubuntu_install_vai_core.sh diff --git a/docker/Dockerfile.ci_vai b/docker/Dockerfile.ci_vai deleted file mode 100644 index d0a30f20520e..000000000000 --- a/docker/Dockerfile.ci_vai +++ /dev/null @@ -1,55 +0,0 @@ -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -# CI docker VAI env -FROM xilinx/vitis-ai:latest - -RUN apt-get update --fix-missing - - -COPY install/ubuntu_install_core.sh /install/ubuntu_install_core.sh -RUN bash /install/ubuntu_install_core.sh - -COPY install/ubuntu_install_python.sh /install/ubuntu_install_python.sh -RUN bash /install/ubuntu_install_python.sh - -COPY install/ubuntu_install_python_package.sh /install/ubuntu_install_python_package.sh -RUN bash /install/ubuntu_install_python_package.sh - -COPY install/ubuntu_install_llvm.sh /install/ubuntu_install_llvm.sh -RUN bash /install/ubuntu_install_llvm.sh - -# NNPACK deps -COPY install/ubuntu_install_nnpack.sh /install/ubuntu_install_nnpack.sh -RUN bash /install/ubuntu_install_nnpack.sh - -ENV PATH $PATH:$CARGO_HOME/bin:/usr/lib/go-1.10/bin - -# ANTLR deps -COPY install/ubuntu_install_java.sh /install/ubuntu_install_java.sh -RUN bash /install/ubuntu_install_java.sh - -# Install Vitis-AI ubuntu dependencies -COPY install/ubuntu_install_vai_core.sh /install/ubuntu_install_vai_core.sh -RUN bash /install/ubuntu_install_vai_core.sh - -# Install dependencies inside vitis-ai-tensorflow conda -RUN . $VAI_ROOT/conda/etc/profile.d/conda.sh && \ - conda activate vitis-ai-tensorflow && \ - pip install --no-cache-dir antlr4-python3-runtime - -ENV USER="root" diff --git a/docker/bash.sh b/docker/bash.sh deleted file mode 100755 index 7817adaf8839..000000000000 --- a/docker/bash.sh +++ /dev/null @@ -1,136 +0,0 @@ -#!/usr/bin/env bash - -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -# -# Start a bash, mount /workspace to be current directory. -# -# Usage: docker/bash.sh -# Starts an interactive session -# -# Usage2: docker/bash.sh [COMMAND] -# Execute command in the docker image, non-interactive -# -if [ "$#" -lt 1 ]; then - echo "Usage: docker/bash.sh [COMMAND]" - exit -1 -fi - -DOCKER_IMAGE_NAME=("$1") - -if [ "$#" -eq 1 ]; then - COMMAND="bash" - if [[ $(uname) == "Darwin" ]]; then - # Docker's host networking driver isn't supported on macOS. - # Use default bridge network and expose port for jupyter notebook. - CI_DOCKER_EXTRA_PARAMS=("-it -p 8888:8888") - else - CI_DOCKER_EXTRA_PARAMS=("-it --net=host") - fi -else - shift 1 - COMMAND=("$@") -fi - -SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -WORKSPACE="$(pwd)" - -# Use nvidia-docker if the container is GPU. -if [[ ! -z $CUDA_VISIBLE_DEVICES ]]; then - CUDA_ENV="-e CUDA_VISIBLE_DEVICES=${CUDA_VISIBLE_DEVICES}" -else - CUDA_ENV="" -fi - -if [[ "${DOCKER_IMAGE_NAME}" == *"gpu"* ]]; then - if ! type "nvidia-docker" 1> /dev/null 2> /dev/null - then - DOCKER_BINARY="docker" - CUDA_ENV=" --gpus all "${CUDA_ENV} - else - DOCKER_BINARY="nvidia-docker" - fi -else - DOCKER_BINARY="docker" -fi - -if [[ "${DOCKER_IMAGE_NAME}" == *"ci"* ]]; then - CI_PY_ENV="-e PYTHONPATH=/workspace/python" -else - CI_PY_ENV="" -fi - -if [[ "${DOCKER_IMAGE_NAME}" == *"ci_vai"* && -d "/dev/shm" && -d "/opt/xilinx/dsa" && -d "/opt/xilinx/overlaybins" ]]; then - WORKSPACE_VOLUMES="-v /dev/shm:/dev/shm -v /opt/xilinx/dsa:/opt/xilinx/dsa -v /opt/xilinx/overlaybins:/opt/xilinx/overlaybins" - XCLMGMT_DRIVER="$(find /dev -name xclmgmt\*)" - DOCKER_DEVICES="" - for i in ${XCLMGMT_DRIVER} ; - do - DOCKER_DEVICES+="--device=$i " - done - - RENDER_DRIVER="$(find /dev/dri -name renderD\*)" - for i in ${RENDER_DRIVER} ; - do - DOCKER_DEVICES+="--device=$i " - done - -else - DOCKER_DEVICES="" - WORKSPACE_VOLUMES="" -fi - - -# Print arguments. -echo "WORKSPACE: ${WORKSPACE}" -echo "DOCKER CONTAINER NAME: ${DOCKER_IMAGE_NAME}" -echo "" - -echo "Running '${COMMAND[@]}' inside ${DOCKER_IMAGE_NAME}..." - -# When running from a git worktree, also mount the original git dir. -EXTRA_MOUNTS=( ) -if [ -f "${WORKSPACE}/.git" ]; then - git_dir=$(cd ${WORKSPACE} && git rev-parse --git-common-dir) - if [ "${git_dir}" != "${WORKSPACE}/.git" ]; then - EXTRA_MOUNTS=( "${EXTRA_MOUNTS[@]}" -v "${git_dir}:${git_dir}" ) - fi -fi - -# By default we cleanup - remove the container once it finish running (--rm) -# and share the PID namespace (--pid=host) so the process inside does not have -# pid 1 and SIGKILL is propagated to the process inside (jenkins can kill it). -${DOCKER_BINARY} run --rm --pid=host\ - ${DOCKER_DEVICES}\ - ${WORKSPACE_VOLUMES}\ - -v ${WORKSPACE}:/workspace \ - -v ${SCRIPT_DIR}:/docker \ - "${EXTRA_MOUNTS[@]}" \ - -w /workspace \ - -e "CI_BUILD_HOME=/workspace" \ - -e "CI_BUILD_USER=$(id -u -n)" \ - -e "CI_BUILD_UID=$(id -u)" \ - -e "CI_BUILD_GROUP=$(id -g -n)" \ - -e "CI_BUILD_GID=$(id -g)" \ - -e "CI_PYTEST_ADD_OPTIONS=$CI_PYTEST_ADD_OPTIONS" \ - ${CI_PY_ENV} \ - ${CUDA_ENV} \ - ${CI_DOCKER_EXTRA_PARAMS[@]} \ - ${DOCKER_IMAGE_NAME} \ - bash --login /docker/with_the_same_user \ - ${COMMAND[@]} diff --git a/docker/install/ubuntu_install_python.sh b/docker/install/ubuntu_install_python.sh deleted file mode 100755 index 58d72f327aa6..000000000000 --- a/docker/install/ubuntu_install_python.sh +++ /dev/null @@ -1,40 +0,0 @@ -#!/bin/bash -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -set -e -set -u -set -o pipefail - -# install python and pip, don't modify this, modify install_python_package.sh -apt-get update -apt-get install -y python-dev - -# python 3.6 -apt-get install -y software-properties-common - -add-apt-repository -y ppa:deadsnakes/ppa -apt-get update -apt-get install -y python-pip python-dev python3.6 python3.6-dev - -rm -f /usr/bin/python3 && ln -s /usr/bin/python3.6 /usr/bin/python3 - -# Install pip -cd /tmp && wget -q https://bootstrap.pypa.io/get-pip.py && python2 get-pip.py && python3.6 get-pip.py - -# Pin pip version -pip3 install pip==19.3.1 diff --git a/docker/install/ubuntu_install_vai_core.sh b/docker/install/ubuntu_install_vai_core.sh deleted file mode 100644 index 00189a7c308a..000000000000 --- a/docker/install/ubuntu_install_vai_core.sh +++ /dev/null @@ -1,37 +0,0 @@ -#!/bin/bash -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -set -e -set -u -set -o pipefail - -# install libraries for building Vitis-AI on ubuntu -apt-get update && apt-get install -y --no-install-recommends \ - build-essential\ - ca-certificates\ - cmake\ - sudo\ - wget\ - git\ - vim\ - graphviz\ - python-dev\ - gnupg2 - -apt-get update && apt-get install -y gcc-aarch64-linux-gnu - From 9dce16153188d02d8855c91d87a1906c953d3ed4 Mon Sep 17 00:00:00 2001 From: anilmartha Date: Tue, 1 Sep 2020 15:35:19 +0530 Subject: [PATCH 03/22] Update vitis_ai.rst --- docs/deploy/vitis_ai.rst | 78 +++++++++++++++++++++++++++------------- 1 file changed, 54 insertions(+), 24 deletions(-) diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index 3e1db774e67c..9aa12d8c1d0c 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -1,3 +1,21 @@ +.. Licensed to the Apache Software Foundation (ASF) under one + or more contributor license agreements. See the NOTICE file + distributed with this work for additional information + regarding copyright ownership. The ASF licenses this file + to you under the Apache License, Version 2.0 (the + "License"); you may not use this file except in compliance + with the License. You may obtain a copy of the License at + +.. http://www.apache.org/licenses/LICENSE-2.0 + +.. Unless required by applicable law or agreed to in writing, + software distributed under the License is distributed on an + "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + KIND, either express or implied. See the License for the + specific language governing permissions and limitations + under the License. + + Vitis-AI Integration ==================== @@ -79,14 +97,14 @@ Hardware setup and docker build ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ 1. Clone the Vitis AI repository: - :: - + .. code:: bash git clone --recurse-submodules https://github.com/Xilinx/Vitis-AI 2. Install Docker, and add the user to the docker group. Link the user to docker installation instructions from the following docker's website: + - https://docs.docker.com/install/linux/docker-ce/ubuntu/ - https://docs.docker.com/install/linux/docker-ce/centos/ @@ -108,9 +126,8 @@ Hardware setup and docker build them at once. To do so: - Run the following commands: - :: - - + .. code:: bash + cd Vitis-AI/alveo/packages sudo su ./install.sh @@ -118,15 +135,16 @@ Hardware setup and docker build - Power cycle the system. 5. Clone tvm repo and pyxir repo - :: + + .. code:: bash git clone --recursive https://github.com/apache/incubator-tvm.git git clone --recursive https://github.com/Xilinx/pyxir.git 6. Build and start the tvm runtime Vitis-AI Docker Container. - :: + .. code:: bash bash incubator-tvm/docker/build.sh ci_vai bash bash incubator-tvm/docker/bash.sh tvm.ci_vai @@ -137,16 +155,16 @@ Hardware setup and docker build conda activate vitis-ai-tensorflow 7. Install PyXIR - :: - + .. code:: bash cd pyxir python3 setup.py install --use_vai_rt_dpucadx8g --user 8. Build TVM inside the container with Vitis-AI - :: + + .. code:: bash cd incubator-tvm @@ -159,7 +177,9 @@ Hardware setup and docker build make -j$(nproc) 9. Install TVM - :: + + .. code:: bash + cd incubator-tvm/python pip3 install -e . --user @@ -203,10 +223,12 @@ Host setup and docker build ^^^^^^^^^^^^^^^^^^^^^^^^^^^ 1. Clone tvm repo -:: + +.. code:: bash git clone --recursive https://github.com/apache/incubator-tvm.git 2. Build and start the tvm runtime Vitis-AI Docker Container. -:: + +.. code:: bash cd incubator-tvm bash incubator-tvm/docker/build.sh ci_vai bash bash incubator-tvm/docker/bash.sh tvm.ci_vai @@ -216,7 +238,8 @@ Host setup and docker build conda activate vitis-ai-tensorflow 3. Install PyXIR -:: + +.. code:: bash git clone --recursive https://github.com/Xilinx/pyxir.git @@ -225,7 +248,9 @@ Host setup and docker build 4. Build TVM inside the container with Vitis-AI. -:: + +.. code:: bash + cd incubator-tvm mkdir build cp cmake/config.cmake build @@ -236,7 +261,8 @@ Host setup and docker build make -j$(nproc) 5. Install TVM -:: + +.. code:: bash cd incubator-tvm/python pip3 install -e . --user @@ -273,14 +299,13 @@ NOTE: This section provides instructions setting up with the `Pynq `__ 5. Run the following command to download the DPU bitstream: - :: - + .. code:: bash python3 -c 'from pynq_dpu import DpuOverlay ; overlay = DpuOverlay("dpu.bit")' 6. Check whether the DPU kernel is alive: - :: + .. code:: bash dexplorer -w @@ -294,13 +319,15 @@ Building TVM depends on the Xilinx interface between TVM and Vitis-AI tools. 1. First install the PyXIR h5py and pydot dependencies: -:: + +.. code:: bash: apt-get install libhdf5-dev pip3 install pydot h5py 2. Install PyXIR -:: + +.. code:: bash: git clone --recursive https://github.com/Xilinx/pyxir.git @@ -308,7 +335,8 @@ interface between TVM and Vitis-AI tools. sudo python3 setup.py install --use_vai_rt_dpuczdx8g 3. Build TVM with Vitis-AI -:: + +.. code:: bash: git clone --recursive https://github.com/apache/incubator-tvm @@ -321,13 +349,15 @@ interface between TVM and Vitis-AI tools. make 4. Install TVM -:: + +.. code:: bash: + cd incubator-tvm/python pip3 install -e . --user 5. Check whether the setup was successful in the Python shell: -:: +.. code:: bash: python3 -c 'import pyxir; import tvm' From fb18b1763ec96c50ef1f62fff5a65ed0ca515d91 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Thu, 3 Sep 2020 05:46:50 -0600 Subject: [PATCH 04/22] Add review changes --- cmake/modules/contrib/VITISAI.cmake | 2 +- docs/deploy/vitis_ai.rst | 325 ++++++++++-------- python/tvm/contrib/target/vitis_ai.py | 42 ++- python/tvm/contrib/vitis_ai_runtime.py | 54 --- python/tvm/relay/op/contrib/vitis_ai.py | 40 ++- .../contrib/vitis_ai/vitis_ai_runtime.cc | 4 +- .../test_vitis_ai_codegen.py | 128 ++++--- .../test_vitis_ai_runtime.py | 25 +- 8 files changed, 325 insertions(+), 295 deletions(-) delete mode 100644 python/tvm/contrib/vitis_ai_runtime.py rename tests/python/contrib/{ => test_vitis_ai}/test_vitis_ai_codegen.py (61%) rename tests/python/contrib/{ => test_vitis_ai}/test_vitis_ai_runtime.py (88%) diff --git a/cmake/modules/contrib/VITISAI.cmake b/cmake/modules/contrib/VITISAI.cmake index 6ccdbaa15fdc..13207c27c62e 100644 --- a/cmake/modules/contrib/VITISAI.cmake +++ b/cmake/modules/contrib/VITISAI.cmake @@ -37,7 +37,7 @@ if(USE_VITIS_AI) else() message(STATUS "To find Pyxir, Python interpreter is required to be found.") - endif() +endif() message(STATUS "Build with contrib.vitisai") include_directories(${PYXIR_INCLUDE_DIR}) diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index 9aa12d8c1d0c..adb8d4e7902c 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -97,6 +97,7 @@ Hardware setup and docker build ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ 1. Clone the Vitis AI repository: + .. code:: bash git clone --recurse-submodules https://github.com/Xilinx/Vitis-AI @@ -110,7 +111,23 @@ Hardware setup and docker build - https://docs.docker.com/install/linux/docker-ce/centos/ - https://docs.docker.com/install/linux/linux-postinstall/ -3. Any GPU instructions will have to be separated from Vitis AI. +3. Download the latest Vitis AI Docker with the following command. This container runs on CPU. + + .. code:: bash + + + docker pull xilinx/vitis-ai:latest + + To accelerate the quantization, you can optionally use the Vitis-AI GPU docker image. Use the below commands to build the Vitis-AI GPU docker container: + + + .. code:: bash + + + cd Vitis-AI/docker + ./docker_build_gpu.sh + + 4. Set up Vitis AI to target Alveo cards. To target Alveo cards with Vitis AI for machine learning workloads, you must install the following software components: @@ -146,8 +163,8 @@ Hardware setup and docker build .. code:: bash - bash incubator-tvm/docker/build.sh ci_vai bash - bash incubator-tvm/docker/bash.sh tvm.ci_vai + bash incubator-tvm/docker/build.sh demo_vitis_ai bash + bash incubator-tvm/docker/bash.sh tvm.demo_vitis_ai #Setup inside container source /opt/xilinx/xrt/setup.sh @@ -224,47 +241,50 @@ Host setup and docker build 1. Clone tvm repo -.. code:: bash - git clone --recursive https://github.com/apache/incubator-tvm.git + .. code:: bash + + git clone --recursive https://github.com/apache/incubator-tvm.git 2. Build and start the tvm runtime Vitis-AI Docker Container. -.. code:: bash - cd incubator-tvm - bash incubator-tvm/docker/build.sh ci_vai bash - bash incubator-tvm/docker/bash.sh tvm.ci_vai - - #Setup inside container - . $VAI_ROOT/conda/etc/profile.d/conda.sh - conda activate vitis-ai-tensorflow + .. code:: bash + + cd incubator-tvm + bash incubator-tvm/docker/build.sh demo_vitis_ai bash + bash incubator-tvm/docker/bash.sh tvm.demo_vitis_ai + + #Setup inside container + . $VAI_ROOT/conda/etc/profile.d/conda.sh + conda activate vitis-ai-tensorflow 3. Install PyXIR -.. code:: bash + .. code:: bash - git clone --recursive https://github.com/Xilinx/pyxir.git - cd pyxir - python3 setup.py install --user + git clone --recursive https://github.com/Xilinx/pyxir.git + cd pyxir + python3 setup.py install --user 4. Build TVM inside the container with Vitis-AI. -.. code:: bash + .. code:: bash - cd incubator-tvm - mkdir build - cp cmake/config.cmake build - cd build - echo set\(USE_LLVM ON\) >> config.cmake - echo set\(USE_VITIS_AI ON\) >> config.cmake - cmake .. - make -j$(nproc) + cd incubator-tvm + mkdir build + cp cmake/config.cmake build + cd build + echo set\(USE_LLVM ON\) >> config.cmake + echo set\(USE_VITIS_AI ON\) >> config.cmake + cmake .. + make -j$(nproc) 5. Install TVM -.. code:: bash - cd incubator-tvm/python - pip3 install -e . --user + .. code:: bash + + cd incubator-tvm/python + pip3 install -e . --user Edge requirements ^^^^^^^^^^^^^^^^^ @@ -285,8 +305,10 @@ platform. The following development boards can be used out-of-the-box: Edge hardware setup ^^^^^^^^^^^^^^^^^^^ ++------------------------------------------------------------------------------------------------------------------------------------------------------------+ +| .. note:: This section provides instructions for setting up with the `Pynq `__ platform but Petalinux based flows are also supported. | ++------------------------------------------------------------------------------------------------------------------------------------------------------------+ -NOTE: This section provides instructions setting up with the `Pynq `__ platform but Petalinux based flows are also supported. 1. Download the Pynq v2.5 image for your target (use Z1 or Z2 for Ultra96 target depending on board version) Link to image: @@ -312,7 +334,10 @@ NOTE: This section provides instructions setting up with the `Pynq `__ package. PyXIR acts as an @@ -320,46 +345,46 @@ interface between TVM and Vitis-AI tools. 1. First install the PyXIR h5py and pydot dependencies: -.. code:: bash: + .. code:: bash: - apt-get install libhdf5-dev - pip3 install pydot h5py + apt-get install libhdf5-dev + pip3 install pydot h5py 2. Install PyXIR -.. code:: bash: + .. code:: bash: - git clone --recursive https://github.com/Xilinx/pyxir.git - cd pyxir - sudo python3 setup.py install --use_vai_rt_dpuczdx8g + git clone --recursive https://github.com/Xilinx/pyxir.git + cd pyxir + sudo python3 setup.py install --use_vai_rt_dpuczdx8g 3. Build TVM with Vitis-AI -.. code:: bash: + .. code:: bash: - git clone --recursive https://github.com/apache/incubator-tvm - cd incubator-tvm - mkdir build - cp cmake/config.cmake build - cd build - echo set\(USE_VITIS_AI ON\) >> config.cmake - cmake .. - make + git clone --recursive https://github.com/apache/incubator-tvm + cd incubator-tvm + mkdir build + cp cmake/config.cmake build + cd build + echo set\(USE_VITIS_AI ON\) >> config.cmake + cmake .. + make 4. Install TVM -.. code:: bash: + .. code:: bash: cd incubator-tvm/python pip3 install -e . --user 5. Check whether the setup was successful in the Python shell: -.. code:: bash: + .. code:: bash: - python3 -c 'import pyxir; import tvm' + python3 -c 'import pyxir; import tvm' Getting started @@ -415,27 +440,27 @@ used by TVM to integrate with the Vitis-AI stack. Additionaly, import the typical TVM and Relay modules and the Vitis-AI contrib module inside TVM. -:: +.. code:: bash - import pyxir - import pyxir.contrib.target.DPUCADX8G + import pyxir + import pyxir.contrib.target.DPUCADX8G - import tvm - import tvm.relay as relay - from tvm.contrib.target import vitis_ai - from tvm.relay.build_module import bind_params_by_name - from tvm.relay.op.contrib.vitis_ai import annotation + import tvm + import tvm.relay as relay + from tvm.contrib.target import vitis_ai + from tvm.relay.build_module import bind_params_by_name + from tvm.relay.op.contrib.vitis_ai import annotation After importing a convolutional neural network model using the usual Relay API's, annotate the Relay expression for the given Vitis-AI DPU target and partition the graph. -:: +.. code:: bash - mod["main"] = bind_params_by_name(mod["main"], params) - mod = annotation(mod, params, target) - mod = relay.transform.MergeCompilerRegions()(mod) - mod = relay.transform.PartitionGraph()(mod) + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, target) + mod = relay.transform.MergeCompilerRegions()(mod) + mod = relay.transform.PartitionGraph()(mod) Now, we can build the TVM runtime library for executing the model. The TVM target is 'llvm' as the operations that can't be handled by the DPU @@ -443,13 +468,13 @@ are executed on the CPU. The Vitis-AI target is DPUCADX8G as we are targeting the cloud DPU and this target is passed as a config to the TVM build call. -:: +.. code:: bash - tvm_target = 'llvm' - target='DPUCADX8G' + tvm_target = 'llvm' + target='DPUCADX8G' - with tvm.transform.PassContext(opt_level=3, config= {'target_': target}): - graph, lib, params = relay.build(mod, tvm_target, params=params) + with tvm.transform.PassContext(opt_level=3, config= {'relay.ext.vitis_ai.options.target': target}): + graph, lib, params = relay.build(mod, tvm_target, params=params) As one more step before we can accelerate a model with Vitis-AI in TVM we have to quantize and compile the model for execution on the DPU. We @@ -461,53 +486,53 @@ iterations, computations will be accelerated on the DPU. So now we will feed N inputs to the TVM runtime module. Note that these first N inputs will take a substantial amount of time. -:: +.. code:: bash: - module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) - module.set_input(**params) + module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) + module.set_input(**params) - # First N (default = 128) inputs are used for quantization calibration and will - # be executed on the CPU - # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) - for i in range(128): - module.set_input(input_name, inputs[i]) - module.run() + # First N (default = 128) inputs are used for quantization calibration and will + # be executed on the CPU + # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) + for i in range(128): + module.set_input(input_name, inputs[i]) + module.run() Afterwards, inference will be accelerated on the DPU. -:: +.. code:: bash - module.set_input(name, data) - module.run() + module.set_input(name, data) + module.run() To save and load the built module, one can use the typical TVM API's: -:: +.. code:: bash - # save the graph, lib and params into separate files - from tvm.contrib import util + # save the graph, lib and params into separate files + from tvm.contrib import util - temp = util.tempdir() - path_lib = temp.relpath("deploy_lib.so") - lib.export_library(path_lib) - with open(temp.relpath("deploy_graph.json"), "w") as fo: - fo.write(graph) - with open(temp.relpath("deploy_param.params"), "wb") as fo: - fo.write(relay.save_param_dict(params)) + temp = util.tempdir() + path_lib = temp.relpath("deploy_lib.so") + lib.export_library(path_lib) + with open(temp.relpath("deploy_graph.json"), "w") as fo: + fo.write(graph) + with open(temp.relpath("deploy_param.params"), "wb") as fo: + fo.write(relay.save_param_dict(params)) Load the module from compiled files and run inference -:: +.. code:: bash - # load the module into memory - loaded_json = open(temp.relpath("deploy_graph.json")).read() - loaded_lib = tvm.runtime.load_module(path_lib) - loaded_params = bytearray(open(temp.relpath("deploy_param.params"), "rb").read()) + # load the module into memory + loaded_json = open(temp.relpath("deploy_graph.json")).read() + loaded_lib = tvm.runtime.load_module(path_lib) + loaded_params = bytearray(open(temp.relpath("deploy_param.params"), "rb").read()) - module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) - module.load_params(loaded_params) - module.set_input(name, data) - module.run() + module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) + module.load_params(loaded_params) + module.set_input(name, data) + module.run() Edge usage ~~~~~~~~~~ @@ -526,27 +551,27 @@ used by TVM to integrate with the Vitis-AI stack. Additionaly, import the typical TVM and Relay modules and the Vitis-AI contrib module inside TVM. -:: +.. code:: bash - import pyxir - import pyxir.contrib.target.DPUCZDX8G + import pyxir + import pyxir.contrib.target.DPUCZDX8G - import tvm - import tvm.relay as relay - from tvm.contrib.target import vitis_ai - from tvm.relay.build_module import bind_params_by_name - from tvm.relay.op.contrib.vitis_ai import annotation + import tvm + import tvm.relay as relay + from tvm.contrib.target import vitis_ai + from tvm.relay.build_module import bind_params_by_name + from tvm.relay.op.contrib.vitis_ai import annotation After importing a convolutional neural network model using the usual Relay API's, annotate the Relay expression for the given Vitis-AI DPU target and partition the graph. -:: +.. code:: bash - mod["main"] = bind_params_by_name(mod["main"], params) - mod = annotation(mod, params, target) - mod = relay.transform.MergeCompilerRegions()(mod) - mod = relay.transform.PartitionGraph()(mod) + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, target) + mod = relay.transform.MergeCompilerRegions()(mod) + mod = relay.transform.PartitionGraph()(mod) Now, we can build the TVM runtime library for executing the model. The TVM target is 'llvm' as the operations that can't be handled by the DPU @@ -556,38 +581,38 @@ on the ZCU104 board and this target is passed as a config to the TVM build call. Note that different identifiers can be passed for different targets, see `edge targets info <#edge-requirements>`__. -:: +.. code:: bash - tvm_target = 'llvm' - target='DPUCZDX8G-zcu104' + tvm_target = 'llvm' + target='DPUCZDX8G-zcu104' - with tvm.transform.PassContext(opt_level=3, config= {'target_': target}): - graph, lib, params = relay.build(mod, tvm_target, params=params) + with tvm.transform.PassContext(opt_level=3, config= {'relay.ext.vitis_ai.options.target': target}): + graph, lib, params = relay.build(mod, tvm_target, params=params) Additionaly, already build the deployment module for the ARM CPU target and serialize: -:: +.. code:: bash - # Export lib for aarch64 target + # Export lib for aarch64 target - tvm_target = tvm.target.arm_cpu('ultra96') - lib_kwargs = { + tvm_target = tvm.target.arm_cpu('ultra96') + lib_kwargs = { 'fcompile': contrib.cc.create_shared, 'cc': "/usr/aarch64-linux-gnu/bin/ld" - } + } - with tvm.transform.PassContext(opt_level=3, - config={'target_': target, - 'vai_build_dir_': target + '_build'}): + with tvm.transform.PassContext(opt_level=3, + config={'relay.ext.vitis_ai.options.target': target, + 'relay.ext.vitis_ai.options.build_dir': target + '_build'}): graph_arm, lib_arm, params_arm = relay.build( mod, tvm_target, params=params) - lib_dpuv2.export_library('tvm_dpu_arm.so', **lib_kwargs) - with open("tvm_dpu_arm.json","w") as f: - f.write(graph_dpuv2) - with open("tvm_dpu_arm.params", "wb") as f: - f.write(relay.save_param_dict(params_dpuv2)) + lib_dpuv2.export_library('tvm_dpu_arm.so', **lib_kwargs) + with open("tvm_dpu_arm.json","w") as f: + f.write(graph_dpuv2) + with open("tvm_dpu_arm.params", "wb") as f: + f.write(relay.save_param_dict(params_dpuv2)) As one more step before we can deploy a model with Vitis-AI in TVM at the edge we have to quantize and compile the model for execution on the @@ -597,17 +622,17 @@ quantize the model on the host using N inputs. After providing N inputs we can then move the TVM and Vitis-AI build files to the edge device for deployment. -:: +.. code:: bash - module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) - module.set_input(**params) + module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) + module.set_input(**params) - # First N (default = 128) inputs are used for quantization calibration and will - # be executed on the CPU - # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) - for i in range(128): - module.set_input(input_name, inputs[i]) - module.run() + # First N (default = 128) inputs are used for quantization calibration and will + # be executed on the CPU + # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) + for i in range(128): + module.set_input(input_name, inputs[i]) + module.run() Now, move the TVM build files (tvm\_dpu\_arm.json, tvm\_dpu\_arm.so, tvm\_dpu\_arm.params) and the DPU build directory (e.g. @@ -625,23 +650,21 @@ Move the target build directory to the same folder where the example running script is located and explicitly set the path to the build directory using the PX\_BUILD\_DIR environment variable. -:: +.. code:: bash export PX_BUILD_DIR={PATH-TO-DPUCZDX8G-BUILD_DIR} Then load the TVM runtime module into memory and feed inputs for inference. -:: - - # load the module into memory - loaded_json = open(temp.relpath("tvm_dpu_arm.json")).read() - loaded_lib = tvm.runtime.load_module("tvm_dpu_arm.so") - loaded_params = bytearray(open(temp.relpath("tvm_dpu_arm.params"), "rb").read()) - - module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) - module.load_params(loaded_params) - module.set_input(name, data) - module.run() +.. code:: bash + # load the module into memory + loaded_json = open(temp.relpath("tvm_dpu_arm.json")).read() + loaded_lib = tvm.runtime.load_module("tvm_dpu_arm.so") + loaded_params = bytearray(open(temp.relpath("tvm_dpu_arm.params"), "rb").read()) + module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) + module.load_params(loaded_params) + module.set_input(name, data) + module.run() diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index 932db232f8c8..327887d8e637 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -15,33 +15,29 @@ # specific language governing permissions and limitations # under the License. # pylint: disable=invalid-name, unused-argument, import-outside-toplevel -"""Utility to compile VITISAI models""" +"""Utility to compile Vitis-AI models""" import os -from tvm.relay.expr import Tuple, Call +from tvm.relay.expr import Tuple, Call, TupleGetItem import tvm._ffi import pyxir import pyxir.frontend.tvm -from .. import vitis_ai_runtime class CodegenVitisAI: """ - Traverse subgraphs and build XGraph + Traverse Relay expression and convert into PyXIR XGraph format """ def __init__(self, model_name, function): - self.model_name = model_name self.function = function self.params = {} - - def convert_pyxir(self, target): """ - Convert relay submodule expression to PYXIR(XGRAPH) + Convert Relay expression to PyXIR XGraph """ xgraph = pyxir.frontend.tvm.from_relay(self.function, params=self.params, postprocessing=None) @@ -50,7 +46,7 @@ def convert_pyxir(self, target): def get_output_names(self): """ - Get output names from subgraph + Get output names from Relay expression """ func = self.function output_relay_ids = [] @@ -60,14 +56,16 @@ def get_output_names(self): output_relay_ids.append(hash(field)) elif isinstance(expr, Call): output_relay_ids.append(hash(expr)) + elif isinstance(expr, TupleGetItem): + output_relay_ids.append(hash(expr.tuple_value)) else: - raise ValueError("does not support {}".format(type(expr))) + raise ValueError("Vitis-AI codegen does not support {} as output".format(type(expr))) return output_relay_ids -@tvm._ffi.register_func("relay.ext.vai") -def vai_compiler(ref): +@tvm._ffi.register_func("relay.ext.vitis_ai") +def vitis_ai_compiler(ref): """ - Create a VAI runtime from a Relay module. + Create a Vitis-AI runtime from the provided Relay expression """ assert isinstance(ref, tvm.relay.function.Function) @@ -76,19 +74,24 @@ def vai_compiler(ref): name = str(ref.attrs.global_symbol) pass_context = tvm.get_global_func("transform.GetCurrentPassContext")() - target = str(pass_context.config['target_']) - vai_build_dir = str(pass_context.config['vai_build_dir_']) \ + target = str(pass_context.config['relay.ext.vitis_ai.options.target']) + vai_build_dir = str(pass_context.config['relay.ext.vitis_ai.options.build_dir']) \ if 'vai_build_dir_' in pass_context.config else None if vai_build_dir and not os.path.exists(vai_build_dir): raise ValueError("Provided Vitis-AI build dir: `{}` could not be found" .format(vai_build_dir)) + + # If build directory is not passed as a parameter in transform.PassContext, + # we will build the Vitis-AI PyXIR runtime from scratch if not vai_build_dir: + # Convert Relay expression into XGraph and do partitioning inside PyXIR builder = CodegenVitisAI(name, ref) model_dir = target + "_build/" xgraph = builder.convert_pyxir(target) output_relay_ids = builder.get_output_names() layers = xgraph.get_layers() - # get the output tensor names using xgraph and output relay ids + + # Get the output tensor names using XGraph and output Relay ids out_tensor_names = [] for layer in layers: if not layer.internal: @@ -96,7 +99,7 @@ def vai_compiler(ref): out_tensor_names.append(layer.name) if len(out_tensor_names) == 0: raise ValueError("During codegeneration the loading of subexpression \ - failed due to output tensorname mismatch in relay pyxir interface.") + failed due to output tensor name mismatch in Relay PyXIR interface.") # Save/serialize XGraph if not os.path.exists(model_dir): @@ -106,4 +109,7 @@ def vai_compiler(ref): else: model_dir = vai_build_dir - return vitis_ai_runtime.create(name, model_dir, target).module + # Create Vitis-AI runtime module + runtime_func = "tvm.vitis_ai_runtime.create" + fcreate = tvm._ffi.get_global_func(runtime_func) + return fcreate(name, model_dir, target) diff --git a/python/tvm/contrib/vitis_ai_runtime.py b/python/tvm/contrib/vitis_ai_runtime.py deleted file mode 100644 index a03120b301b8..000000000000 --- a/python/tvm/contrib/vitis_ai_runtime.py +++ /dev/null @@ -1,54 +0,0 @@ -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -"""VitisAI runtime that load and run Xgraph.""" -import tvm._ffi - -def create(name, model_dir, target): - """Create a runtime executor module given a xgraph model and context. - Parameters - ---------- - model_dir : str - The directory where the compiled models are located. - target : str - The target for running subgraph. - - Returns - ------- - vai_runtime : VaiModule - Runtime Vai module that can be used to execute xgraph model. - """ - runtime_func = "tvm.vitis_ai_runtime.create" - fcreate = tvm._ffi.get_global_func(runtime_func) - return VitisAIModule(fcreate(name, model_dir, target)) - -class VitisAIModule(object): - """Wrapper runtime module. - - This is a thin wrapper of the underlying TVM module. - you can also directly call set_input, run, and get_output - of underlying module functions - - Parameters - ---------- - module : Module - The internal tvm module that holds the actual vai functions. - - """ - - def __init__(self, module): - self.module = module diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index 8734daa39750..cca0baa26476 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -15,7 +15,7 @@ # specific language governing permissions and limitations # under the License. # pylint: disable=invalid-name, unused-argument, no-else-return, E1102 -"""VITISAI codegen supported operators.""" +"""Vitis-AI codegen supported operators""" import numpy as np @@ -31,16 +31,33 @@ @transform.function_pass(opt_level=0) class VitisAIAnnotationPass: - """The explicit pass wrapper around VitisAIAnnotationPass.""" + + """ + The VitisAIAnnotationPass is responsible for annotating Relay expressions + in the way that they are supported through Vitis-AI accelerators + """ + def __init__(self, compiler, relay_ids): self.compiler = compiler self.relay_ids = relay_ids + def transform_function(self, func, mod, ctx): - """Transform func to annotate.""" + """ + Transform function for annotating Relay module + """ + annotator = self + class Annotator(tvm.relay.ExprMutator): - """Annotator for VITIS-AI DPU.""" + + """ + Annotator for Vitis-AI DPU accelerators + """ + def visit_tuple(self, tup): + """ + Visit the Tuple expression and add compiler_begin and compiler_end annotations + """ field_list = [] cond = int(hash(tup)) for field in tup.fields: @@ -54,6 +71,10 @@ def visit_tuple(self, tup): return Tuple(field_list) def visit_tuple_getitem(self, op): + """ + Visit the TupleGetItem expression and add compiler_begin and compiler_end + annotations + """ if int(hash(op.tuple_value)) in annotator.relay_ids: tuple_value = compiler_begin(super().visit(op.tuple_value), annotator.compiler) @@ -61,7 +82,12 @@ def visit_tuple_getitem(self, op): else: tuple_value = super().visit(op.tuple_value) return TupleGetItem(tuple_value, op.index) + def visit_call(self, call): + """ + Visit the function Call expression and add compiler_begin and compiler_end + annotations + """ if int(hash(call)) in annotator.relay_ids: new_args = [] for arg in call.args: @@ -80,13 +106,15 @@ def visit_call(self, call): def annotation(mod, params, target): """ - An annotator for VITISAI. + Annotate Relay expression for Vitis-AI DPU accelerators """ xgraph = pyxir.frontend.tvm.from_relay(mod, params, postprocessing=None) xgraph = pyxir.partition(xgraph, targets=[target]) + layers = xgraph.get_layers() relay_ids = [list(np.array(layer.attrs['relay_id']).flatten()) for layer in layers if layer.target == target] relay_ids_flatten = [item for sublist in relay_ids for item in sublist] - mod = VitisAIAnnotationPass("vai", relay_ids_flatten)(mod) + mod = VitisAIAnnotationPass("vitis_ai", relay_ids_flatten)(mod) + return mod diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index a042b260d653..64930f2c5842 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -28,8 +28,8 @@ namespace tvm { namespace runtime { -TVM_REGISTER_PASS_CONFIG_OPTION("target_", String); -TVM_REGISTER_PASS_CONFIG_OPTION("vai_build_dir_", String); +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.target", String); +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.build_dir", String); std::shared_ptr load_xgraph_model(const std::string& model_path) { std::string model_name = model_path + "/" + "dpu_xgraph.json"; diff --git a/tests/python/contrib/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py similarity index 61% rename from tests/python/contrib/test_vitis_ai_codegen.py rename to tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index 673047f113a2..f03dfad7cde4 100644 --- a/tests/python/contrib/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -14,20 +14,22 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611 +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 """Vitis-AI codegen tests.""" - +import sys import numpy as np +import pytest +pytest.importorskip('pyxir') +import pyxir.contrib.target.DPUCADX8G + import tvm from tvm import relay from tvm.relay import transform from tvm.relay.op.contrib.vitis_ai import annotation +from tvm.relay.build_module import bind_params_by_name from tvm.contrib.target import vitis_ai -import pyxir -import pyxir.contrib.target.DPUCADX8G - def set_func_attr(func, compile_name, symbol_name): func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) func = func.with_attr("Inline", tvm.tir.IntImm("int32", 1)) @@ -55,16 +57,19 @@ def _construct_model(func, params=None): mod["main"] = func if params is None: params = {} + mod["main"] = bind_params_by_name(mod["main"], params) mod = annotation(mod, params, "DPUCADX8G") mod = transform.MergeCompilerRegions()(mod) mod = transform.PartitionGraph()(mod) - fcompile = tvm._ffi.get_global_func("relay.ext.vai") + fcompile = tvm._ffi.get_global_func("relay.ext.vitis_ai") subgraph_mod = tvm.IRModule() for _, funcnode in mod.functions.items(): if funcnode.attrs and 'Compiler' in funcnode.attrs and \ - funcnode.attrs['Compiler'] == 'vai': + funcnode.attrs['Compiler'] == 'vitis_ai': subgraph_mod["main"] = funcnode - with tvm.transform.PassContext(opt_level=3, config={'target_':'DPUCADX8G'}): + with tvm.transform.PassContext(opt_level=3, \ + config={'relay.ext.vitis_ai.options.target': + 'DPUCADX8G'}): fcompile(subgraph_mod["main"]) @@ -88,9 +93,27 @@ def test_conv2d(): y = relay.nn.conv2d(x, w, strides=[2, 2], padding=[1, 1, 1, 1], kernel_size=[3, 3]) func = relay.Function([x], y) params = {} - params["x"] = np.zeros((16, 3, 3, 3), dtype='float32') + params["x"] = np.zeros((1, 3, 224, 224), dtype='float32') + params["w"] = np.random.rand(16, 3, 3, 3).astype('float32') _construct_model(func, params) +def test_batchnorm(): + data = relay.var('data', shape=(1, 16, 112, 112)) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) + bn_output = relay.nn.batch_norm(data, bn_gamma, bn_beta, bn_mmean, + bn_mvar) + func = relay.Function([data, bn_gamma, bn_beta, bn_mmean, + bn_mvar], bn_output[0]) + params = {} + params["data"] = np.zeros((1, 16, 112, 112), dtype='float32') + params["bn_gamma"] = np.random.rand(16).astype('float32') + params["bn_beta"] = np.random.rand(16).astype('float32') + params["bn_mean"] = np.random.rand(16).astype('float32') + params["bn_var"] = np.random.rand(16).astype('float32') + _construct_model(func, params) def test_global_avg_pool2d(): shape = (10, 10, 10, 10) @@ -99,6 +122,13 @@ def test_global_avg_pool2d(): func = relay.Function([x], y) _construct_model(func) +def test_avg_pool2d(): + shape = (10, 10, 10, 10) + x = relay.var('x', shape=shape) + y = relay.nn.avg_pool2d(x, pool_size=(3, 3)) + func = relay.Function([x], y) + _construct_model(func) + def test_annotate(): """Test annotation with Vitis-AI DP (DPUCADX8G)""" def partition(): @@ -131,11 +161,8 @@ def partition(): mod = annotation(mod, params, "DPUCADX8G") opt_pass = tvm.transform.Sequential([ - transform.InferType(), + transform.MergeCompilerRegions(), transform.PartitionGraph(), - transform.SimplifyInference(), - transform.FoldConstant(), - transform.AlterOpLayout(), ]) with tvm.transform.PassContext(opt_level=3): @@ -144,49 +171,39 @@ def partition(): return mod def expected(): - # function for batch_norm - data0 = relay.var("data0", relay.TensorType((1, 16, 224, 224), - "float32")) - mod = tvm.IRModule() - bn_gamma = relay.var("bn_gamma1", relay.TensorType((16, ), "float32")) - bn_beta = relay.var("bn_beta1", relay.TensorType((16, ), "float32")) - bn_mmean = relay.var("bn_mean1", relay.TensorType((16, ), "float32")) - bn_mvar = relay.var("bn_var1", relay.TensorType((16, ), "float32")) - - bn = relay.nn.batch_norm(data0, bn_gamma, bn_beta, bn_mmean, bn_mvar) - func0 = relay.Function([data0, bn_gamma, bn_beta, bn_mmean, bn_mvar], - bn.astuple()) - func0 = set_func_attr(func0, "vai", "vai_2") - gv0 = relay.GlobalVar("vai_2") - mod[gv0] = func0 - - # function for conv2d - data1 = relay.var("data1", relay.TensorType((1, 3, 224, 224), "float32")) - weight1 = relay.var("weight1", relay.TensorType((16, 3, 3, 3), "float32")) + # function variables for conv2d + data0 = relay.var("data0", relay.TensorType((1, 3, 224, 224), "float32")) + weight0 = relay.var("weight0", relay.TensorType((16, 3, 3, 3), "float32")) conv = relay.nn.conv2d( - data=data1, - weight=weight1, + data=data0, + weight=weight0, kernel_size=(3, 3), channels=16, padding=(1, 1)) - func1 = relay.Function([data1, weight1], conv) - func1 = set_func_attr(func1, "vai", "vai_0") - gv1 = relay.GlobalVar("vai_0") - mod[gv1] = func1 + + # function variables for batch_norm + bn_gamma0 = relay.var("bn_gamma0", relay.TensorType((16, ), "float32")) + bn_beta0 = relay.var("bn_beta0", relay.TensorType((16, ), "float32")) + bn_mmean0 = relay.var("bn_mean0", relay.TensorType((16, ), "float32")) + bn_mvar0 = relay.var("bn_var0", relay.TensorType((16, ), "float32")) + bn = relay.nn.batch_norm(conv, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0) + func0 = relay.Function([data0, weight0, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0], + bn.astuple()) + func0 = set_func_attr(func0, "vitis_ai", "vitis_ai_0") + gv0 = relay.GlobalVar("vitis_ai_0") + mod = tvm.IRModule() + mod[gv0] = func0 # main function data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) - bn_gamma0 = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) - bn_beta0 = relay.var("bn_beta", relay.TensorType((16, ), "float32")) - bn_mmean0 = relay.var("bn_mean", relay.TensorType((16, ), "float32")) - bn_mvar0 = relay.var("bn_var", relay.TensorType((16, ), "float32")) - - call1 = gv1(data, weight) - call0 = gv0(call1, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0) - mod["main"] = relay.Function([data, weight, bn_gamma0, bn_beta0, bn_mmean0, - bn_mvar0], call0) - mod = transform.InferType()(mod) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) + call0 = gv0(data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar) + mod["main"] = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, + bn_mvar], call0) return mod partitioned = partition() @@ -196,8 +213,13 @@ def expected(): if __name__ == "__main__": - test_annotate() - test_add() - test_relu() - test_conv2d() - test_global_avg_pool2d() + if sys.platform == "win32": + print("Skip test on Windows for now") + else: + test_annotate() + test_add() + test_relu() + test_conv2d() + test_batchnorm() + test_global_avg_pool2d() + test_avg_pool2d() diff --git a/tests/python/contrib/test_vitis_ai_runtime.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py similarity index 88% rename from tests/python/contrib/test_vitis_ai_runtime.py rename to tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py index e3630a58ed94..6b15f2a237a2 100644 --- a/tests/python/contrib/test_vitis_ai_runtime.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py @@ -14,14 +14,15 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611 +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 """ Vitis-AI runtime test """ import sys import numpy as np -import pyxir +import pytest +pytest.importorskip('pyxir') import pyxir.contrib.target.DPUCADX8G import tvm @@ -39,10 +40,7 @@ def check_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", ctx=tvm.cpu(), params=None): - """Check the result between reference and generated output with vitis-ai byoc flow""" - if sys.platform == "win32": - print("Skip test on Windows for now") - return + """ To check the result between reference and byoc vitis-ai flow""" def update_lib(lib): tmp_path = util.tempdir() @@ -50,12 +48,12 @@ def update_lib(lib): lib_path = tmp_path.relpath(lib_name) lib.export_library(lib_path) lib = runtime.load_module(lib_path) - return lib def check_graph_runtime_result(): compile_engine.get().clear() - with tvm.transform.PassContext(opt_level=3, config={'target_' : 'DPUCADX8G'}): + with tvm.transform.PassContext(opt_level=3, + config={'relay.ext.vitis_ai.options.target' : 'DPUCADX8G'}): json, lib, param = relay.build(mod, target=target, params=params) lib = update_lib(lib) rt_mod = tvm.contrib.graph_runtime.create(json, lib, ctx) @@ -79,7 +77,11 @@ def check_graph_runtime_result(): def test_extern_vai_resnet18(): """Test resnet18 model using Vitis-AI byoc flow""" - if not tvm.get_global_func("relay.ext.vai", True): + if sys.platform == "win32": + print("Skip test on Windows for now") + return + + if not tvm.get_global_func("relay.ext.vitis_ai", True): print("skip because VITIS-AI codegen is not available") return @@ -101,4 +103,7 @@ def test_extern_vai_resnet18(): check_result(mod, {"data": i_data}, (1, 1000), ref_res.asnumpy(), tol=1e-5, params=params) if __name__ == "__main__": - test_extern_vai_resnet18() + if sys.platform == "win32": + print("Skip test on Windows for now") + else: + test_extern_vai_resnet18() From 0b572277989032bdabcdd3bd959555451a80fa2a Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Sun, 6 Sep 2020 23:00:14 -0600 Subject: [PATCH 05/22] Remove new lines and note frame in vitis_ai.rst --- docs/deploy/vitis_ai.rst | 15 ++++++++------- python/tvm/contrib/target/vitis_ai.py | 1 + python/tvm/relay/op/contrib/vitis_ai.py | 5 ----- 3 files changed, 9 insertions(+), 12 deletions(-) diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index adb8d4e7902c..9ef1ed2d9ed0 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -305,10 +305,10 @@ platform. The following development boards can be used out-of-the-box: Edge hardware setup ^^^^^^^^^^^^^^^^^^^ -+------------------------------------------------------------------------------------------------------------------------------------------------------------+ -| .. note:: This section provides instructions for setting up with the `Pynq `__ platform but Petalinux based flows are also supported. | -+------------------------------------------------------------------------------------------------------------------------------------------------------------+ +.. note:: + This section provides instructions for setting up with the `Pynq `__ platform but + Petalinux based flows are also supported. 1. Download the Pynq v2.5 image for your target (use Z1 or Z2 for Ultra96 target depending on board version) Link to image: @@ -334,10 +334,11 @@ Edge hardware setup Edge TVM setup ^^^^^^^^^^^^^^ -+----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------+ -| .. note:: When working on Petalinux instead of Pynq, the following steps might take more manual work (e.g building hdf5 from source). Also, TVM has a scipy dependency which you then might have to build from source or circumvent. We don't depend on scipy in our flow. | -+----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------+ - +.. note:: + + When working on Petalinux instead of Pynq, the following steps might take more manual work (e.g building + hdf5 from source). Also, TVM has a scipy dependency which you then might have to build from source or + circumvent. We don't depend on scipy in our flow. Building TVM depends on the Xilinx `PyXIR `__ package. PyXIR acts as an diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index 327887d8e637..06ae827e2697 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -113,3 +113,4 @@ def vitis_ai_compiler(ref): runtime_func = "tvm.vitis_ai_runtime.create" fcreate = tvm._ffi.get_global_func(runtime_func) return fcreate(name, model_dir, target) + diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index cca0baa26476..f2df83692855 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -31,12 +31,10 @@ @transform.function_pass(opt_level=0) class VitisAIAnnotationPass: - """ The VitisAIAnnotationPass is responsible for annotating Relay expressions in the way that they are supported through Vitis-AI accelerators """ - def __init__(self, compiler, relay_ids): self.compiler = compiler self.relay_ids = relay_ids @@ -45,15 +43,12 @@ def transform_function(self, func, mod, ctx): """ Transform function for annotating Relay module """ - annotator = self class Annotator(tvm.relay.ExprMutator): - """ Annotator for Vitis-AI DPU accelerators """ - def visit_tuple(self, tup): """ Visit the Tuple expression and add compiler_begin and compiler_end annotations From 969faf68d0ecc48de3823b980fadf4b33cb8ff47 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Sun, 6 Sep 2020 23:03:05 -0600 Subject: [PATCH 06/22] use sys.exit --- .../test_vitis_ai/test_vitis_ai_codegen.py | 17 +++++++++-------- .../test_vitis_ai/test_vitis_ai_runtime.py | 4 ++-- 2 files changed, 11 insertions(+), 10 deletions(-) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index f03dfad7cde4..9caa8a7ae9a8 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -215,11 +215,12 @@ def expected(): if __name__ == "__main__": if sys.platform == "win32": print("Skip test on Windows for now") - else: - test_annotate() - test_add() - test_relu() - test_conv2d() - test_batchnorm() - test_global_avg_pool2d() - test_avg_pool2d() + sys.exit(0) + + test_annotate() + test_add() + test_relu() + test_conv2d() + test_batchnorm() + test_global_avg_pool2d() + test_avg_pool2d() diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py index 6b15f2a237a2..45ce7aaf34b2 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py @@ -105,5 +105,5 @@ def test_extern_vai_resnet18(): if __name__ == "__main__": if sys.platform == "win32": print("Skip test on Windows for now") - else: - test_extern_vai_resnet18() + sys.exit(0) + test_extern_vai_resnet18() From 680f3732d6a68bbfe9490ac31f15f279448282e2 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 8 Sep 2020 10:36:20 -0600 Subject: [PATCH 07/22] Add condition for vitis_ai runtime exec function --- src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index 64930f2c5842..a033ba0e7e7a 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -116,8 +116,8 @@ PackedFunc VitisAIRuntime::GetFunction(const std::string& name, this->initialized_ = true; *rv = 0; }); - } else { - return PackedFunc( + } else if (this->symbol_name_ == name) { + return PackedFunc( [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { DLTensor* inputs = args[0]; std::vector in_shape; @@ -141,7 +141,9 @@ PackedFunc VitisAIRuntime::GetFunction(const std::string& name, // Execute the subgraph. rt_mod_->execute(in_tensors, out_tensors); }); - } + } else { + return PackedFunc(); + } } } // namespace runtime } // namespace tvm From 20015aef5cc426a8ac94c9b57fdec665fc7f0fdd Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 8 Sep 2020 10:37:11 -0600 Subject: [PATCH 08/22] remove unused graph_json --- src/runtime/contrib/vitis_ai/vitis_ai_runtime.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h index 7ff69a456612..8a1a54e65644 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h @@ -80,8 +80,6 @@ void Init(const std::string& model_path, const std::string& target); private: /*! \brief The only subgraph name for this module. */ std::string symbol_name_; - /*! \brief The graph. */ - std::string graph_json_; /*! \brief The required constant names. */ Array const_names_; std::shared_ptr xgraph_; From fb9d3f25e2fa181b3baf5631dfe151586b8953d1 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 8 Sep 2020 10:38:03 -0600 Subject: [PATCH 09/22] correct indentation --- cmake/modules/contrib/VITISAI.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/modules/contrib/VITISAI.cmake b/cmake/modules/contrib/VITISAI.cmake index 13207c27c62e..36acad5fc3ea 100644 --- a/cmake/modules/contrib/VITISAI.cmake +++ b/cmake/modules/contrib/VITISAI.cmake @@ -36,7 +36,7 @@ if(USE_VITIS_AI) endif() else() - message(STATUS "To find Pyxir, Python interpreter is required to be found.") + message(STATUS "To find Pyxir, Python interpreter is required to be found.") endif() message(STATUS "Build with contrib.vitisai") From 4352c7b2a1c8dd51bc1c6edf3aae46e167cca79a Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 8 Sep 2020 10:43:57 -0600 Subject: [PATCH 10/22] use code python instead of bash --- docs/deploy/vitis_ai.rst | 46 ++++++++++++++++------------------------ 1 file changed, 18 insertions(+), 28 deletions(-) diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index 9ef1ed2d9ed0..1ab0a234bfb4 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -114,19 +114,15 @@ Hardware setup and docker build 3. Download the latest Vitis AI Docker with the following command. This container runs on CPU. .. code:: bash - - + docker pull xilinx/vitis-ai:latest To accelerate the quantization, you can optionally use the Vitis-AI GPU docker image. Use the below commands to build the Vitis-AI GPU docker container: - .. code:: bash - - - cd Vitis-AI/docker - ./docker_build_gpu.sh + cd Vitis-AI/docker + ./docker_build_gpu.sh 4. Set up Vitis AI to target Alveo cards. To target Alveo cards with Vitis AI for machine learning workloads, you must install the @@ -154,8 +150,7 @@ Hardware setup and docker build 5. Clone tvm repo and pyxir repo .. code:: bash - - + git clone --recursive https://github.com/apache/incubator-tvm.git git clone --recursive https://github.com/Xilinx/pyxir.git @@ -183,7 +178,6 @@ Hardware setup and docker build .. code:: bash - cd incubator-tvm mkdir build cp cmake/config.cmake build @@ -260,7 +254,6 @@ Host setup and docker build .. code:: bash - git clone --recursive https://github.com/Xilinx/pyxir.git cd pyxir python3 setup.py install --user @@ -348,14 +341,13 @@ interface between TVM and Vitis-AI tools. .. code:: bash: - apt-get install libhdf5-dev pip3 install pydot h5py + 2. Install PyXIR .. code:: bash: - git clone --recursive https://github.com/Xilinx/pyxir.git cd pyxir sudo python3 setup.py install --use_vai_rt_dpuczdx8g @@ -364,7 +356,6 @@ interface between TVM and Vitis-AI tools. .. code:: bash: - git clone --recursive https://github.com/apache/incubator-tvm cd incubator-tvm mkdir build @@ -441,7 +432,7 @@ used by TVM to integrate with the Vitis-AI stack. Additionaly, import the typical TVM and Relay modules and the Vitis-AI contrib module inside TVM. -.. code:: bash +.. code:: python import pyxir import pyxir.contrib.target.DPUCADX8G @@ -456,7 +447,7 @@ After importing a convolutional neural network model using the usual Relay API's, annotate the Relay expression for the given Vitis-AI DPU target and partition the graph. -.. code:: bash +.. code:: python mod["main"] = bind_params_by_name(mod["main"], params) mod = annotation(mod, params, target) @@ -469,7 +460,7 @@ are executed on the CPU. The Vitis-AI target is DPUCADX8G as we are targeting the cloud DPU and this target is passed as a config to the TVM build call. -.. code:: bash +.. code:: python tvm_target = 'llvm' target='DPUCADX8G' @@ -487,7 +478,7 @@ iterations, computations will be accelerated on the DPU. So now we will feed N inputs to the TVM runtime module. Note that these first N inputs will take a substantial amount of time. -.. code:: bash: +.. code:: python module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) module.set_input(**params) @@ -501,14 +492,14 @@ will take a substantial amount of time. Afterwards, inference will be accelerated on the DPU. -.. code:: bash +.. code:: python module.set_input(name, data) module.run() To save and load the built module, one can use the typical TVM API's: -.. code:: bash +.. code:: python # save the graph, lib and params into separate files from tvm.contrib import util @@ -523,7 +514,7 @@ To save and load the built module, one can use the typical TVM API's: Load the module from compiled files and run inference -.. code:: bash +.. code:: python # load the module into memory loaded_json = open(temp.relpath("deploy_graph.json")).read() @@ -552,7 +543,7 @@ used by TVM to integrate with the Vitis-AI stack. Additionaly, import the typical TVM and Relay modules and the Vitis-AI contrib module inside TVM. -.. code:: bash +.. code:: python import pyxir import pyxir.contrib.target.DPUCZDX8G @@ -567,7 +558,7 @@ After importing a convolutional neural network model using the usual Relay API's, annotate the Relay expression for the given Vitis-AI DPU target and partition the graph. -.. code:: bash +.. code:: python mod["main"] = bind_params_by_name(mod["main"], params) mod = annotation(mod, params, target) @@ -582,7 +573,7 @@ on the ZCU104 board and this target is passed as a config to the TVM build call. Note that different identifiers can be passed for different targets, see `edge targets info <#edge-requirements>`__. -.. code:: bash +.. code:: python tvm_target = 'llvm' target='DPUCZDX8G-zcu104' @@ -593,10 +584,9 @@ targets, see `edge targets info <#edge-requirements>`__. Additionaly, already build the deployment module for the ARM CPU target and serialize: -.. code:: bash +.. code:: python # Export lib for aarch64 target - tvm_target = tvm.target.arm_cpu('ultra96') lib_kwargs = { 'fcompile': contrib.cc.create_shared, @@ -623,7 +613,7 @@ quantize the model on the host using N inputs. After providing N inputs we can then move the TVM and Vitis-AI build files to the edge device for deployment. -.. code:: bash +.. code:: python module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) module.set_input(**params) @@ -658,7 +648,7 @@ directory using the PX\_BUILD\_DIR environment variable. Then load the TVM runtime module into memory and feed inputs for inference. -.. code:: bash +.. code:: python # load the module into memory loaded_json = open(temp.relpath("tvm_dpu_arm.json")).read() From 3c453788caf9def062f20b3fcd796459680e4420 Mon Sep 17 00:00:00 2001 From: "anilm (generated by with_the_same_user script)" Date: Wed, 9 Sep 2020 06:42:19 -0600 Subject: [PATCH 11/22] Rename VITISAI.cmake to VitisAI.cmake --- CMakeLists.txt | 2 +- cmake/modules/contrib/VitisAI.cmake | 44 +++++++++++++++++++++++++++++ 2 files changed, 45 insertions(+), 1 deletion(-) create mode 100644 cmake/modules/contrib/VitisAI.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 6090d1d01c51..fbf421d5191b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -323,7 +323,7 @@ include(cmake/modules/contrib/TF_TVMDSOOP.cmake) include(cmake/modules/contrib/CoreML.cmake) include(cmake/modules/contrib/ONNX.cmake) include(cmake/modules/contrib/ArmComputeLib.cmake) -include(cmake/modules/contrib/VITISAI.cmake) +include(cmake/modules/contrib/VitisAI.cmake) include(cmake/modules/Git.cmake) include(cmake/modules/LibInfo.cmake) diff --git a/cmake/modules/contrib/VitisAI.cmake b/cmake/modules/contrib/VitisAI.cmake new file mode 100644 index 000000000000..896a3c8d1513 --- /dev/null +++ b/cmake/modules/contrib/VitisAI.cmake @@ -0,0 +1,44 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +if(USE_VITIS_AI) + set(PYXIR_SHARED_LIB libpyxir.so) + find_package(PythonInterp 3.6 REQUIRED) + if(NOT PYTHON) + find_program(PYTHON NAMES python3 python3.6) + endif() + execute_process(COMMAND "${PYTHON_EXECUTABLE}" "-c" + "import pyxir as px; print(px.get_include_dir()); print(px.get_lib_dir());" + RESULT_VARIABLE __result + OUTPUT_VARIABLE __output + OUTPUT_STRIP_TRAILING_WHITESPACE) + + if(__result MATCHES 0) + string(REGEX REPLACE ";" "\\\\;" __values ${__output}) + string(REGEX REPLACE "\r?\n" ";" __values ${__values}) + list(GET __values 0 PYXIR_INCLUDE_DIR) + list(GET __values 1 PYXIR_LIB_DIR) + else() + message(FATAL_ERROR "Can't build TVM with Vitis-AI because PyXIR can't be found") + endif() + message(STATUS "Build with contrib.vitisai") + include_directories(${PYXIR_INCLUDE_DIR}) + file(GLOB VAI_CONTRIB_SRC src/runtime/contrib/vitis_ai/*.cc) + link_directories(${PYXIR_LIB_DIR}) + list(APPEND TVM_RUNTIME_LINKER_LIBS "pyxir") + list(APPEND RUNTIME_SRCS ${VAI_CONTRIB_SRC}) +endif(USE_VITIS_AI) From dea81b5a3edbb5b750d96a322e1fec8e1697e54e Mon Sep 17 00:00:00 2001 From: "anilm (generated by with_the_same_user script)" Date: Wed, 9 Sep 2020 06:44:05 -0600 Subject: [PATCH 12/22] use relay.ext.vitis_ai.options.build_dir in comparison --- cmake/modules/contrib/VITISAI.cmake | 49 --------------------------- python/tvm/contrib/target/vitis_ai.py | 2 +- 2 files changed, 1 insertion(+), 50 deletions(-) delete mode 100644 cmake/modules/contrib/VITISAI.cmake diff --git a/cmake/modules/contrib/VITISAI.cmake b/cmake/modules/contrib/VITISAI.cmake deleted file mode 100644 index 36acad5fc3ea..000000000000 --- a/cmake/modules/contrib/VITISAI.cmake +++ /dev/null @@ -1,49 +0,0 @@ -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. - -if(USE_VITIS_AI) - set(PYXIR_SHARED_LIB libpyxir.so) - find_package(PythonInterp 3.6 REQUIRED) - if(NOT PYTHON) - find_program(PYTHON NAMES python3 python3.6) - endif() - if(PYTHON) - execute_process(COMMAND "${PYTHON_EXECUTABLE}" "-c" - "import pyxir as px; print(px.get_include_dir()); print(px.get_lib_dir());" - RESULT_VARIABLE __result - OUTPUT_VARIABLE __output - OUTPUT_STRIP_TRAILING_WHITESPACE) - - if(__result MATCHES 0) - string(REGEX REPLACE ";" "\\\\;" __values ${__output}) - string(REGEX REPLACE "\r?\n" ";" __values ${__values}) - list(GET __values 0 PYXIR_INCLUDE_DIR) - list(GET __values 1 PYXIR_LIB_DIR) - endif() - - else() - message(STATUS "To find Pyxir, Python interpreter is required to be found.") -endif() - -message(STATUS "Build with contrib.vitisai") -include_directories(${PYXIR_INCLUDE_DIR}) -file(GLOB VAI_CONTRIB_SRC src/runtime/contrib/vitis_ai/*.cc) -link_directories(${PYXIR_LIB_DIR}) -list(APPEND TVM_RUNTIME_LINKER_LIBS "pyxir") -list(APPEND RUNTIME_SRCS ${VAI_CONTRIB_SRC}) -endif(USE_VITIS_AI) - diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index 06ae827e2697..f1b65b2b6737 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -76,7 +76,7 @@ def vitis_ai_compiler(ref): pass_context = tvm.get_global_func("transform.GetCurrentPassContext")() target = str(pass_context.config['relay.ext.vitis_ai.options.target']) vai_build_dir = str(pass_context.config['relay.ext.vitis_ai.options.build_dir']) \ - if 'vai_build_dir_' in pass_context.config else None + if 'relay.ext.vitis_ai.options.build_dir' in pass_context.config else None if vai_build_dir and not os.path.exists(vai_build_dir): raise ValueError("Provided Vitis-AI build dir: `{}` could not be found" .format(vai_build_dir)) From 832fae0381fcbae74cb0e61fa94bb80bc94aeac2 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Fri, 11 Sep 2020 03:55:22 -0600 Subject: [PATCH 13/22] Re-add deleted docker related files --- docker/bash.sh | 113 ++++++++++++++++++++++++ docker/install/ubuntu_install_python.sh | 40 +++++++++ 2 files changed, 153 insertions(+) create mode 100644 docker/bash.sh create mode 100644 docker/install/ubuntu_install_python.sh diff --git a/docker/bash.sh b/docker/bash.sh new file mode 100644 index 000000000000..73bfb12268f3 --- /dev/null +++ b/docker/bash.sh @@ -0,0 +1,113 @@ +#!/usr/bin/env bash + +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +# +# Start a bash, mount /workspace to be current directory. +# +# Usage: docker/bash.sh +# Starts an interactive session +# +# Usage2: docker/bash.sh [COMMAND] +# Execute command in the docker image, non-interactive +# +if [ "$#" -lt 1 ]; then + echo "Usage: docker/bash.sh [COMMAND]" + exit -1 +fi + +DOCKER_IMAGE_NAME=("$1") + +if [ "$#" -eq 1 ]; then + COMMAND="bash" + if [[ $(uname) == "Darwin" ]]; then + # Docker's host networking driver isn't supported on macOS. + # Use default bridge network and expose port for jupyter notebook. + CI_DOCKER_EXTRA_PARAMS=("-it -p 8888:8888") + else + CI_DOCKER_EXTRA_PARAMS=("-it --net=host") + fi +else + shift 1 + COMMAND=("$@") +fi + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +WORKSPACE="$(pwd)" + +# Use nvidia-docker if the container is GPU. +if [[ ! -z $CUDA_VISIBLE_DEVICES ]]; then + CUDA_ENV="-e CUDA_VISIBLE_DEVICES=${CUDA_VISIBLE_DEVICES}" +else + CUDA_ENV="" +fi + +if [[ "${DOCKER_IMAGE_NAME}" == *"gpu"* ]]; then + if ! type "nvidia-docker" 1> /dev/null 2> /dev/null + then + DOCKER_BINARY="docker" + CUDA_ENV=" --gpus all "${CUDA_ENV} + else + DOCKER_BINARY="nvidia-docker" + fi +else + DOCKER_BINARY="docker" +fi + +if [[ "${DOCKER_IMAGE_NAME}" == *"ci"* ]]; then + CI_PY_ENV="-e PYTHONPATH=/workspace/python" +else + CI_PY_ENV="" +fi + +# Print arguments. +echo "WORKSPACE: ${WORKSPACE}" +echo "DOCKER CONTAINER NAME: ${DOCKER_IMAGE_NAME}" +echo "" + +echo "Running '${COMMAND[@]}' inside ${DOCKER_IMAGE_NAME}..." + +# When running from a git worktree, also mount the original git dir. +EXTRA_MOUNTS=( ) +if [ -f "${WORKSPACE}/.git" ]; then + git_dir=$(cd ${WORKSPACE} && git rev-parse --git-common-dir) + if [ "${git_dir}" != "${WORKSPACE}/.git" ]; then + EXTRA_MOUNTS=( "${EXTRA_MOUNTS[@]}" -v "${git_dir}:${git_dir}" ) + fi +fi + +# By default we cleanup - remove the container once it finish running (--rm) +# and share the PID namespace (--pid=host) so the process inside does not have +# pid 1 and SIGKILL is propagated to the process inside (jenkins can kill it). +${DOCKER_BINARY} run --rm --pid=host\ + -v ${WORKSPACE}:/workspace \ + -v ${SCRIPT_DIR}:/docker \ + "${EXTRA_MOUNTS[@]}" \ + -w /workspace \ + -e "CI_BUILD_HOME=/workspace" \ + -e "CI_BUILD_USER=$(id -u -n)" \ + -e "CI_BUILD_UID=$(id -u)" \ + -e "CI_BUILD_GROUP=$(id -g -n)" \ + -e "CI_BUILD_GID=$(id -g)" \ + -e "CI_PYTEST_ADD_OPTIONS=$CI_PYTEST_ADD_OPTIONS" \ + ${CI_PY_ENV} \ + ${CUDA_ENV} \ + ${CI_DOCKER_EXTRA_PARAMS[@]} \ + ${DOCKER_IMAGE_NAME} \ + bash --login /docker/with_the_same_user \ + ${COMMAND[@]} diff --git a/docker/install/ubuntu_install_python.sh b/docker/install/ubuntu_install_python.sh new file mode 100644 index 000000000000..c1f9d5081f57 --- /dev/null +++ b/docker/install/ubuntu_install_python.sh @@ -0,0 +1,40 @@ +#!/bin/bash +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +set -e +set -u +set -o pipefail + +# install python and pip, don't modify this, modify install_python_package.sh +apt-get update +apt-get install -y python-dev + +# python 3.6 +apt-get install -y software-properties-common + +add-apt-repository ppa:deadsnakes/ppa +apt-get update +apt-get install -y python-pip python-dev python3.6 python3.6-dev + +rm -f /usr/bin/python3 && ln -s /usr/bin/python3.6 /usr/bin/python3 + +# Install pip +cd /tmp && wget -q https://bootstrap.pypa.io/get-pip.py && python2 get-pip.py && python3.6 get-pip.py + +# Pin pip version +pip3 install pip==19.3.1 From 2a845fda6bb23ef2303811d83a187c531ba91d29 Mon Sep 17 00:00:00 2001 From: Jorn Tuyls Date: Sun, 20 Sep 2020 22:35:20 -0700 Subject: [PATCH 14/22] Make use of PyXIR XGraph and RuntimeModule serialization & refactor flow --- docker/bash.sh | 0 docker/install/ubuntu_install_python.sh | 0 .../ubuntu_install_vitis_ai_packages_ci.sh | 2 +- docs/deploy/vitis_ai.rst | 142 +++++------ python/tvm/contrib/target/vitis_ai.py | 102 +++++--- python/tvm/relay/op/contrib/vitis_ai.py | 43 ++-- .../contrib/vitis_ai/vitis_ai_runtime.cc | 237 ++++++++++++------ .../contrib/vitis_ai/vitis_ai_runtime.h | 69 +++-- .../python/contrib/test_vitis_ai/__init__.py | 18 ++ .../contrib/test_vitis_ai/infrastructure.py | 143 +++++++++++ .../test_vitis_ai/test_vitis_ai_codegen.py | 228 ++++++++++++----- .../test_vitis_ai/test_vitis_ai_runtime.py | 109 -------- .../test_vitis_ai_runtime_cpu_part.py | 74 ++++++ 13 files changed, 754 insertions(+), 413 deletions(-) mode change 100644 => 100755 docker/bash.sh mode change 100644 => 100755 docker/install/ubuntu_install_python.sh create mode 100644 tests/python/contrib/test_vitis_ai/__init__.py create mode 100644 tests/python/contrib/test_vitis_ai/infrastructure.py delete mode 100644 tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py create mode 100644 tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py diff --git a/docker/bash.sh b/docker/bash.sh old mode 100644 new mode 100755 diff --git a/docker/install/ubuntu_install_python.sh b/docker/install/ubuntu_install_python.sh old mode 100644 new mode 100755 diff --git a/docker/install/ubuntu_install_vitis_ai_packages_ci.sh b/docker/install/ubuntu_install_vitis_ai_packages_ci.sh index d4077bc67b44..c34ed3addce2 100644 --- a/docker/install/ubuntu_install_vitis_ai_packages_ci.sh +++ b/docker/install/ubuntu_install_vitis_ai_packages_ci.sh @@ -25,5 +25,5 @@ mkdir "$PYXIR_HOME" pip3 install progressbar -git clone --recursive --branch v0.1.2 https://github.com/Xilinx/pyxir.git "${PYXIR_HOME}" +git clone --recursive --branch v0.1.3 https://github.com/Xilinx/pyxir.git "${PYXIR_HOME}" cd "${PYXIR_HOME}" && python3 setup.py install diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index 1ab0a234bfb4..f5bda10cddea 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -158,8 +158,8 @@ Hardware setup and docker build .. code:: bash - bash incubator-tvm/docker/build.sh demo_vitis_ai bash - bash incubator-tvm/docker/bash.sh tvm.demo_vitis_ai + ./incubator-tvm/docker/build.sh demo_vitis_ai bash + ./incubator-tvm/docker/bash.sh tvm.demo_vitis_ai #Setup inside container source /opt/xilinx/xrt/setup.sh @@ -243,8 +243,8 @@ Host setup and docker build .. code:: bash cd incubator-tvm - bash incubator-tvm/docker/build.sh demo_vitis_ai bash - bash incubator-tvm/docker/bash.sh tvm.demo_vitis_ai + ./incubator-tvm/docker/build.sh demo_vitis_ai bash + ./incubator-tvm/docker/bash.sh tvm.demo_vitis_ai #Setup inside container . $VAI_ROOT/conda/etc/profile.d/conda.sh @@ -440,6 +440,7 @@ TVM. import tvm import tvm.relay as relay from tvm.contrib.target import vitis_ai + from tvm.contrib import util, graph_runtime from tvm.relay.build_module import bind_params_by_name from tvm.relay.op.contrib.vitis_ai import annotation @@ -466,7 +467,7 @@ build call. target='DPUCADX8G' with tvm.transform.PassContext(opt_level=3, config= {'relay.ext.vitis_ai.options.target': target}): - graph, lib, params = relay.build(mod, tvm_target, params=params) + lib = relay.build(mod, tvm_target, params=params) As one more step before we can accelerate a model with Vitis-AI in TVM we have to quantize and compile the model for execution on the DPU. We @@ -480,8 +481,7 @@ will take a substantial amount of time. .. code:: python - module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) - module.set_input(**params) + module = graph_runtime.GraphModule(lib["default"](tvm.cpu())) # First N (default = 128) inputs are used for quantization calibration and will # be executed on the CPU @@ -501,28 +501,17 @@ To save and load the built module, one can use the typical TVM API's: .. code:: python - # save the graph, lib and params into separate files - from tvm.contrib import util - - temp = util.tempdir() - path_lib = temp.relpath("deploy_lib.so") - lib.export_library(path_lib) - with open(temp.relpath("deploy_graph.json"), "w") as fo: - fo.write(graph) - with open(temp.relpath("deploy_param.params"), "wb") as fo: - fo.write(relay.save_param_dict(params)) + lib_path = "deploy_lib.so" + lib.export_library(lib_path) Load the module from compiled files and run inference .. code:: python # load the module into memory - loaded_json = open(temp.relpath("deploy_graph.json")).read() - loaded_lib = tvm.runtime.load_module(path_lib) - loaded_params = bytearray(open(temp.relpath("deploy_param.params"), "rb").read()) + loaded_lib = tvm.runtime.load_module(lib_path) - module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) - module.load_params(loaded_params) + module = graph_runtime.GraphModule(lib["default"](tvm.cpu())) module.set_input(name, data) module.run() @@ -551,6 +540,7 @@ TVM. import tvm import tvm.relay as relay from tvm.contrib.target import vitis_ai + from tvm.contrib import util, graph_runtime from tvm.relay.build_module import bind_params_by_name from tvm.relay.op.contrib.vitis_ai import annotation @@ -567,22 +557,61 @@ target and partition the graph. Now, we can build the TVM runtime library for executing the model. The TVM target is 'llvm' as the operations that can't be handled by the DPU -are executed on the CPU. At this point that means the CPU on the host. +are executed on the CPU. At this point that means the CPU on the host machine. The Vitis-AI target is DPUCZDX8G-zcu104 as we are targeting the edge DPU on the ZCU104 board and this target is passed as a config to the TVM build call. Note that different identifiers can be passed for different -targets, see `edge targets info <#edge-requirements>`__. +targets, see `edge targets info <#edge-requirements>`__. Additionally, we +provide the 'export_runtime_module' config that points to a file to which we +can export the Vitis-AI runtime module. We have to do this because we will +first be compiling and quantizing the model on the host machine before building +the model for edge deployment. As you will see later on, the exported runtime +module will be passed to the edge build so that the Vitis-AI runtime module +can be included. .. code:: python + from tvm.contrib import util + + temp = util.tempdir() + tvm_target = 'llvm' target='DPUCZDX8G-zcu104' + export_rt_mod_file = temp.relpath("vitis_ai.rtmod") + + with tvm.transform.PassContext(opt_level=3, config= {'relay.ext.vitis_ai.options.target': target, + 'relay.ext.vitis_ai.options.export_runtime_module': export_rt_mod_file}): + lib = relay.build(mod, tvm_target, params=params) + +We will quantize and compile the model for execution on the DPU using on-the-fly +quantization on the host machine. This makes use of TVM inference calls +(module.run) to quantize the model on the host with the first N inputs. - with tvm.transform.PassContext(opt_level=3, config= {'relay.ext.vitis_ai.options.target': target}): - graph, lib, params = relay.build(mod, tvm_target, params=params) +.. code:: python -Additionaly, already build the deployment module for the ARM CPU target -and serialize: + module = graph_runtime.GraphModule(lib["default"](tvm.cpu())) + + # First N (default = 128) inputs are used for quantization calibration and will + # be executed on the CPU + # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) + for i in range(128): + module.set_input(input_name, inputs[i]) + module.run() + +Save the TVM lib module so that the Vitis-AI runtime module will also be exported +(to the 'export_runtime_module' path we previously passed as a config). + +.. code:: python + + from tvm.contrib import util + + temp = util.tempdir() + lib.export_library(temp.relpath("tvm_lib.so")) + +After quantizing and compiling the model for Vitis-AI acceleration using the +first N inputs we can build the model for execution on the ARM edge device. +Here we pass the previously exported Vitis-AI runtime module so it can be included +in the TVM build. .. code:: python @@ -594,68 +623,29 @@ and serialize: } with tvm.transform.PassContext(opt_level=3, - config={'relay.ext.vitis_ai.options.target': target, - 'relay.ext.vitis_ai.options.build_dir': target + '_build'}): - graph_arm, lib_arm, params_arm = relay.build( - mod, tvm_target, params=params) + config={'relay.ext.vitis_ai.options.load_runtime_module': export_rt_mod_file}): + lib_arm = relay.build(mod, tvm_target, params=params) lib_dpuv2.export_library('tvm_dpu_arm.so', **lib_kwargs) - with open("tvm_dpu_arm.json","w") as f: - f.write(graph_dpuv2) - with open("tvm_dpu_arm.params", "wb") as f: - f.write(relay.save_param_dict(params_dpuv2)) - -As one more step before we can deploy a model with Vitis-AI in TVM at -the edge we have to quantize and compile the model for execution on the -DPU. We make use of on-the-fly quantization on the host machine for -this. This involves using the TVM inference calls (module.run) to -quantize the model on the host using N inputs. After providing N inputs -we can then move the TVM and Vitis-AI build files to the edge device for -deployment. - -.. code:: python - - module = tvm.contrib.graph_runtime.create(graph, lib, tvm.cpu()) - module.set_input(**params) - - # First N (default = 128) inputs are used for quantization calibration and will - # be executed on the CPU - # This config can be changed by setting the 'PX_QUANT_SIZE' (e.g. export PX_QUANT_SIZE=64) - for i in range(128): - module.set_input(input_name, inputs[i]) - module.run() Now, move the TVM build files (tvm\_dpu\_arm.json, tvm\_dpu\_arm.so, -tvm\_dpu\_arm.params) and the DPU build directory (e.g. -DPUCZDX8G-zcu104\_build) to the edge device. For information on setting +tvm\_dpu\_arm.params) to the edge device. For information on setting up the edge device check out the `edge setup <#edge-dpuczdx8g>`__ section. Edge steps ^^^^^^^^^^ -The following steps will have to be executed on the edge device after -setup and moving the build files from the host. - -Move the target build directory to the same folder where the example -running script is located and explicitly set the path to the build -directory using the PX\_BUILD\_DIR environment variable. - -.. code:: bash - - export PX_BUILD_DIR={PATH-TO-DPUCZDX8G-BUILD_DIR} - -Then load the TVM runtime module into memory and feed inputs for -inference. +After setting up TVM with Vitis-AI on the edge device, you can now load +the TVM runtime module into memory and feed inputs for inference. .. code:: python + ctx = tvm.cpu() + # load the module into memory - loaded_json = open(temp.relpath("tvm_dpu_arm.json")).read() - loaded_lib = tvm.runtime.load_module("tvm_dpu_arm.so") - loaded_params = bytearray(open(temp.relpath("tvm_dpu_arm.params"), "rb").read()) + lib = tvm.runtime.load_module("tvm_dpu_arm.so") - module = tvm.contrib.graph_runtime.create(loaded_json, loaded_lib, ctx) - module.load_params(loaded_params) + module = graph_runtime.GraphModule(lib["default"](tvm.cpu())) module.set_input(name, data) module.run() diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index f1b65b2b6737..4ae35853b2f3 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -15,39 +15,36 @@ # specific language governing permissions and limitations # under the License. # pylint: disable=invalid-name, unused-argument, import-outside-toplevel -"""Utility to compile Vitis-AI models""" -import os +"""Utility to offload (sub-)models to Vitis-AI""" -from tvm.relay.expr import Tuple, Call, TupleGetItem -import tvm._ffi +import warnings import pyxir import pyxir.frontend.tvm +from tvm.relay.expr import Tuple, Call, TupleGetItem +import tvm._ffi + class CodegenVitisAI: - """ - Traverse Relay expression and convert into PyXIR XGraph format - """ + + """Traverse Relay expression and convert into PyXIR XGraph format""" + def __init__(self, model_name, function): self.model_name = model_name self.function = function self.params = {} def convert_pyxir(self, target): - """ - Convert Relay expression to PyXIR XGraph - """ + """Convert Relay expression to PyXIR XGraph""" xgraph = pyxir.frontend.tvm.from_relay(self.function, params=self.params, postprocessing=None) xgraph = pyxir.partition(xgraph, targets=[target]) return xgraph def get_output_names(self): - """ - Get output names from Relay expression - """ + """Get output names from Relay expression""" func = self.function output_relay_ids = [] expr = func.body @@ -62,31 +59,57 @@ def get_output_names(self): raise ValueError("Vitis-AI codegen does not support {} as output".format(type(expr))) return output_relay_ids + @tvm._ffi.register_func("relay.ext.vitis_ai") def vitis_ai_compiler(ref): - """ - Create a Vitis-AI runtime from the provided Relay expression - """ + """Create a Vitis-AI runtime from the provided Relay expression""" assert isinstance(ref, tvm.relay.function.Function) - model_dir = os.getcwd() out_tensor_names = [] name = str(ref.attrs.global_symbol) pass_context = tvm.get_global_func("transform.GetCurrentPassContext")() - target = str(pass_context.config['relay.ext.vitis_ai.options.target']) + + # The target Vitis-AI accelerator device + target = str(pass_context.config['relay.ext.vitis_ai.options.target']) \ + if 'relay.ext.vitis_ai.options.target' in pass_context.config else None + + # (Optional configs) The build and work directories to be used by Vitis-AI vai_build_dir = str(pass_context.config['relay.ext.vitis_ai.options.build_dir']) \ - if 'relay.ext.vitis_ai.options.build_dir' in pass_context.config else None - if vai_build_dir and not os.path.exists(vai_build_dir): - raise ValueError("Provided Vitis-AI build dir: `{}` could not be found" - .format(vai_build_dir)) - - # If build directory is not passed as a parameter in transform.PassContext, - # we will build the Vitis-AI PyXIR runtime from scratch - if not vai_build_dir: + if 'relay.ext.vitis_ai.options.build_dir' in pass_context.config else \ + tvm.contrib.util.tempdir().relpath("") + vai_work_dir = str(pass_context.config['relay.ext.vitis_ai.options.work_dir']) \ + if 'relay.ext.vitis_ai.options.work_dir' in pass_context.config else \ + tvm.contrib.util.tempdir().relpath("") + + # (Optional configs) Export and load PyXIR runtime module to file if provided. This is used to + # compile and quantize a model on the host and deploy it at the edge + export_runtime_module = \ + str(pass_context.config['relay.ext.vitis_ai.options.export_runtime_module']) \ + if 'relay.ext.vitis_ai.options.export_runtime_module' in pass_context.config else "" + load_runtime_module = \ + str(pass_context.config['relay.ext.vitis_ai.options.load_runtime_module']) \ + if 'relay.ext.vitis_ai.options.load_runtime_module' in pass_context.config else "" + + # Config checks + if load_runtime_module != "" and target is not None: + warnings.warn("Both `load_runtime_module` and `target` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal target so the `target` config will be ignored") + if load_runtime_module != "" and 'relay.ext.vitis_ai.options.build_dir' in pass_context.config: + warnings.warn("Both `load_runtime_module` and `build_dir` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal build directory so the `build_dir` config will be ignored") + if load_runtime_module != "" and 'relay.ext.vitis_ai.options.work_dir' in pass_context.config: + warnings.warn("Both `load_runtime_module` and `work_dir` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal work directory so the `work_dir` config will be ignored") + + + # If load_runtime_module is not set, we will build the PyXIR runtime module from scratch + if load_runtime_module == "": # Convert Relay expression into XGraph and do partitioning inside PyXIR builder = CodegenVitisAI(name, ref) - model_dir = target + "_build/" xgraph = builder.convert_pyxir(target) output_relay_ids = builder.get_output_names() layers = xgraph.get_layers() @@ -95,22 +118,21 @@ def vitis_ai_compiler(ref): out_tensor_names = [] for layer in layers: if not layer.internal: - if layer.attrs['relay_id'][0] in output_relay_ids: - out_tensor_names.append(layer.name) + for relay_id in layer.attrs['relay_id']: + if relay_id in output_relay_ids: + out_tensor_names.append(layer.name) + break if len(out_tensor_names) == 0: raise ValueError("During codegeneration the loading of subexpression \ failed due to output tensor name mismatch in Relay PyXIR interface.") - - # Save/serialize XGraph - if not os.path.exists(model_dir): - os.mkdir(model_dir) xgraph.meta_attrs['tvm_out_tensors'] = out_tensor_names - pyxir.graph.io.xgraph_io.XGraphIO.save(xgraph, model_dir + 'dpu_xgraph') - else: - model_dir = vai_build_dir + xgraph_str = pyxir.get_xgraph_str(xgraph) - # Create Vitis-AI runtime module - runtime_func = "tvm.vitis_ai_runtime.create" - fcreate = tvm._ffi.get_global_func(runtime_func) - return fcreate(name, model_dir, target) + runtime_func = "tvm.vitis_ai_runtime.from_xgraph" + fcreate = tvm._ffi.get_global_func(runtime_func) + return fcreate(name, xgraph_str, target, vai_build_dir, vai_work_dir, + export_runtime_module) + runtime_func = "tvm.vitis_ai_runtime.from_rt_mod" + fcreate = tvm._ffi.get_global_func(runtime_func) + return fcreate(name, load_runtime_module, export_runtime_module) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index f2df83692855..f9ba4e1a6916 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -15,44 +15,37 @@ # specific language governing permissions and limitations # under the License. # pylint: disable=invalid-name, unused-argument, no-else-return, E1102 -"""Vitis-AI codegen supported operators""" +"""Vitis-AI codegen annotation of supported operators""" import numpy as np +import pyxir +import pyxir.frontend.tvm + from tvm import relay import tvm._ffi from tvm.relay.expr import Tuple, TupleGetItem from tvm.relay import transform from tvm.relay.op.annotation import compiler_begin, compiler_end -import pyxir -import pyxir.frontend.tvm - @transform.function_pass(opt_level=0) class VitisAIAnnotationPass: - """ - The VitisAIAnnotationPass is responsible for annotating Relay expressions - in the way that they are supported through Vitis-AI accelerators - """ + """Responsible for annotating Relay expressions for Vitis-AI DPU accelerators""" + def __init__(self, compiler, relay_ids): self.compiler = compiler self.relay_ids = relay_ids def transform_function(self, func, mod, ctx): - """ - Transform function for annotating Relay module - """ + """Transform function for annotating Relay module""" annotator = self class Annotator(tvm.relay.ExprMutator): - """ - Annotator for Vitis-AI DPU accelerators - """ + """Annotator for Vitis-AI DPU accelerators""" + def visit_tuple(self, tup): - """ - Visit the Tuple expression and add compiler_begin and compiler_end annotations - """ + """Add compiler_begin and compiler_end annotations to Tuple""" field_list = [] cond = int(hash(tup)) for field in tup.fields: @@ -66,10 +59,7 @@ def visit_tuple(self, tup): return Tuple(field_list) def visit_tuple_getitem(self, op): - """ - Visit the TupleGetItem expression and add compiler_begin and compiler_end - annotations - """ + """Add compiler_begin and compiler_end annotations to TupleGetItem""" if int(hash(op.tuple_value)) in annotator.relay_ids: tuple_value = compiler_begin(super().visit(op.tuple_value), annotator.compiler) @@ -79,10 +69,7 @@ def visit_tuple_getitem(self, op): return TupleGetItem(tuple_value, op.index) def visit_call(self, call): - """ - Visit the function Call expression and add compiler_begin and compiler_end - annotations - """ + """Add compiler_begin and compiler_end annotations to the Call expr""" if int(hash(call)) in annotator.relay_ids: new_args = [] for arg in call.args: @@ -95,14 +82,12 @@ def visit_call(self, call): else: return super().visit_call(call) - return Annotator().visit(func) + return Annotator().visit(func) def annotation(mod, params, target): - """ - Annotate Relay expression for Vitis-AI DPU accelerators - """ + """Annotate Relay expression for Vitis-AI DPU accelerators""" xgraph = pyxir.frontend.tvm.from_relay(mod, params, postprocessing=None) xgraph = pyxir.partition(xgraph, targets=[target]) diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index a033ba0e7e7a..1ef041bc2e7c 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -20,130 +20,207 @@ /*! * \file vitis_ai_runtime.cc */ + +#include +#include + #include #include #include "vitis_ai_runtime.h" +using namespace pyxir::runtime; + namespace tvm { namespace runtime { +/*! \brief The target Vitis-AI accelerator device */ TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.target", String); +/*! \brief (Optional config) The build directory to be used by Vitis-AI */ TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.build_dir", String); - -std::shared_ptr load_xgraph_model(const std::string& model_path) { - std::string model_name = model_path + "/" + "dpu_xgraph.json"; - std::string model_weights = model_path + "/" + "dpu_xgraph.h5"; - return pyxir::load(model_name, model_weights); +/*! \brief (Optional config) The work directory to be used by Vitis-AI */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.work_dir", String); +/*! \brief (Optional config) Export PyXIR runtime module to disk during serialization if provided */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.export_runtime_module", String); +/*! \brief (Optional config) Load PyXIR runtime module from disk */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.load_runtime_module", String); + + +VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, + const Array const_names, + const std::string& serialized_rt_mod, + const std::string& export_rt_mod_path) + : symbol_name_(symbol_name), const_names_(const_names), + export_rt_mod_path_(export_rt_mod_path) { + + std::istringstream sstream(serialized_rt_mod); + rt_mod_.reset(new RuntimeModule()); + rt_mod_->deserialize(sstream); + in_tensor_names_ = rt_mod_->get_in_tensor_names(); + out_tensor_names_ = rt_mod_->get_out_tensor_names(); } -void VitisAIRuntime::Init(const std::string& model_path, const std::string& target) { - model_path_ = model_path; - target_ = target; - xgraph_ = load_xgraph_model(model_path_); - in_tensor_names_ = xgraph_->get_input_names(); - out_tensor_names_ = xgraph_->get_meta_attr("tvm_out_tensors").get_strings(); - pyxir::partition(xgraph_, std::vector{target}, ""); + +VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, + const std::string& xgraph_str, + const Array const_names, + const std::string& target, + const std::string& build_dir, + const std::string& work_dir, + const std::string& export_rt_mod_path) + : symbol_name_(symbol_name), const_names_(const_names), + export_rt_mod_path_(export_rt_mod_path) { + + std::istringstream xgraph_sstream(xgraph_str); + pyxir::XGraphHolder xgraph = std::make_shared(""); + pyxir::read(xgraph, xgraph_sstream); + in_tensor_names_ = xgraph->get_input_names(); + out_tensor_names_ = xgraph->get_meta_attr("tvm_out_tensors").get_strings(); + + pyxir::partition(xgraph, std::vector{target}, ""); + pyxir::RunOptionsHolder run_options(new pyxir::runtime::RunOptions()); run_options->on_the_fly_quantization = true; - rt_mod_ = pyxir::build_rt(xgraph_, target_ , in_tensor_names_, out_tensor_names_, - "vai", run_options); + run_options->build_dir = build_dir; + if (!work_dir.empty()) + run_options->work_dir = work_dir; + rt_mod_ = pyxir::build_rt(xgraph, target, in_tensor_names_, out_tensor_names_, + "vai", run_options); } Module VitisAIRuntimeCreate(const std::string& name, - const std::string& model_path, - const std::string& target) { + const std::string& xgraph_str, + const std::string& target, + const std::string& build_dir, + const std::string& work_dir, + const std::string& export_rt_mod_path) { Array const_vars; - auto exec = make_object(name, const_vars); - exec->Init(model_path, target); + auto exec = make_object(name, xgraph_str, const_vars, target, + build_dir, work_dir, + export_rt_mod_path); return Module(exec); } +TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.from_xgraph").set_body([](TVMArgs args, TVMRetValue* rv) { + *rv = VitisAIRuntimeCreate(args[0], args[1], args[2], + args[3], args[4], args[5]); +}); -TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.create").set_body([](TVMArgs args, TVMRetValue* rv) { - *rv = VitisAIRuntimeCreate(args[0], args[1], args[2]); +Module VitisAIRuntimeCreate(const std::string& name, + const std::string& serialized_rt_mod, + const std::string& export_rt_mod_path) { + Array const_vars; + auto exec = make_object(name, const_vars, serialized_rt_mod, + export_rt_mod_path); + return Module(exec); +} + +TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.from_rt_mod").set_body([](TVMArgs args, TVMRetValue* rv) { + std::string load_rt_mod_path = args[1]; + assert(!load_rt_mod_path.empty()); + std::ifstream in_file(load_rt_mod_path); + std::stringstream buffer; + buffer << in_file.rdbuf(); + std::string serialized_rt_mod = buffer.str(); + in_file.close(); + *rv = VitisAIRuntimeCreate(args[0], serialized_rt_mod, args[2]); }); + Module VitisAIRuntimeLoadFromBinary(void* strm ) { - dmlc::Stream* stream = static_cast(strm); - - std::string model_path; - std::string symbol_name; - std::vector const_vars; - std::string target; - stream->Read(&model_path); - stream->Read(&target); - stream->Read(&symbol_name); - stream->Read(&const_vars); - Array const_names; - for (const auto& it : const_vars) { - const_names.push_back(it); - } - auto exec = make_object(symbol_name, const_names); - exec->Init(model_path, target); - return Module(exec); + dmlc::Stream* stream = static_cast(strm); + std::string symbol_name; + std::vector const_vars; + std::string serialized_rt_mod; + std::string export_rt_mod_path; + stream->Read(&serialized_rt_mod); + stream->Read(&export_rt_mod_path); + stream->Read(&symbol_name); + stream->Read(&const_vars); + Array const_names; + for (const auto& it : const_vars) { + const_names.push_back(it); } + auto exec = make_object(symbol_name, const_names, serialized_rt_mod, + export_rt_mod_path); + return Module(exec); +} TVM_REGISTER_GLOBAL("runtime.module.loadbinary_VitisAIRuntime") .set_body_typed(VitisAIRuntimeLoadFromBinary); void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { - stream->Write(this-> model_path_); - stream->Write(this-> target_); - stream->Write(this->symbol_name_); + std::ostringstream sstream; + rt_mod_->serialize(sstream); + stream->Write(sstream.str()); + stream->Write(export_rt_mod_path_); + stream->Write(symbol_name_); std::vector consts; for (const auto& it : const_names_) { consts.push_back(it); } stream->Write(consts); + + // If export_runtime_module_ member variable is set, we will additionally export the PyXIR + // runtime_module to the specified file + if (!export_rt_mod_path_.empty()) { + std::ofstream out_file(export_rt_mod_path_); + out_file << sstream.str(); + out_file.close(); } +} PackedFunc VitisAIRuntime::GetFunction(const std::string& name, - const ObjectPtr& sptr_to_self) { + const ObjectPtr& sptr_to_self) { if (name == "get_symbol") { - return PackedFunc( - [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->symbol_name_; }); - } else if (name == "get_const_vars") { - return PackedFunc( - [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->const_names_; }); - } else if ("__init_" + this->symbol_name_ == name) { - // The function to initialize constant tensors. - return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { - CHECK_EQ(args.size(), 1U); - this->initialized_ = true; - *rv = 0; + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->symbol_name_; }); + } else if (name == "get_const_vars") { + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { *rv = this->const_names_; }); + } else if ("__init_" + this->symbol_name_ == name) { + // The function to initialize constant tensors. + return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { + CHECK_EQ(args.size(), 1U); + this->initialized_ = true; + *rv = 0; + }); + } else if (this->symbol_name_ == name) { + return PackedFunc( + [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { + // Initialize input tensors + DLTensor* inputs = args[0]; + std::vector in_tensors; + std::vector in_shape; + for (int i = 0; i < inputs->ndim; ++i) + in_shape.push_back(inputs->shape[i]); + in_tensors.push_back( + std::shared_ptr(new pyxir::XBuffer(reinterpret_cast(static_cast(inputs->data)), + 4, "f", in_shape.size(), in_shape, false, false))); + + // Initialize output tensors + std::vector out_tensors; + for (unsigned i = 0; i < out_tensor_names_.size(); ++i) { + DLTensor* output_tensor = args[args.size() - out_tensor_names_.size()+i]; + std::vector out_shape; + for (int i = 0; i < output_tensor->ndim; ++i) + out_shape.push_back(output_tensor->shape[i]); + void* output_data = reinterpret_cast (static_cast(output_tensor->data)); + out_tensors.push_back( + std::shared_ptr(new pyxir::XBuffer(output_data, 4, "f", out_shape.size(), + out_shape, false, false))); + } + + // Execute the subgraph. + rt_mod_->execute(in_tensors, out_tensors); }); - } else if (this->symbol_name_ == name) { - return PackedFunc( - [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { - DLTensor* inputs = args[0]; - std::vector in_shape; - for (int i = 0; i < inputs->ndim; ++i) - in_shape.push_back(inputs->shape[i]); - pyxir::XBufferHolder xb_in = std::shared_ptr( - new pyxir::XBuffer(reinterpret_cast(static_cast(inputs->data)), 4, - "f", in_shape.size(), in_shape, false, false)); - std::vector out_tensors; - for (unsigned i = 0; i < out_tensor_names_.size(); ++i) { - DLTensor* output_tensor = args[args.size() - out_tensor_names_.size()+i]; - std::vector out_shape; - for (int i = 0; i < output_tensor->ndim; ++i) - out_shape.push_back(output_tensor->shape[i]); - void* output_data = reinterpret_cast (static_cast(output_tensor->data)); - out_tensors.push_back(std::shared_ptr( - new pyxir::XBuffer(output_data, 4, "f", out_shape.size(), out_shape, - false, false))); - } - std::vector in_tensors{xb_in}; - // Execute the subgraph. - rt_mod_->execute(in_tensors, out_tensors); - }); - } else { - return PackedFunc(); - } + } else { + return PackedFunc(); } +} + } // namespace runtime } // namespace tvm diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h index 8a1a54e65644..06086083d0e0 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h @@ -25,20 +25,17 @@ #ifndef TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ #define TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ - -#include -#include -#include - #include #include #include -#include +#include +#include +#include +#include #include - namespace tvm { namespace runtime { @@ -50,26 +47,52 @@ namespace runtime { */ class VitisAIRuntime : public ModuleNode { public: -VitisAIRuntime(const std::string& symbol_name, const Array const_names): - symbol_name_(symbol_name), const_names_(const_names) {} + /*! + * \brief Create VitisAI runtime from serialized XGraph + * \param symbol_name The name of the function. + * \param const_names The names of each constant in the sub-graph. + * \param serialized_rt_mod The serialized runtime module. + * \param export_rt_mod_path The path to the file to be used for exporting the + * PyXIR runtime module. + */ + VitisAIRuntime(const std::string& symbol_name, + const Array const_names, + const std::string& serialized_rt_mod, + const std::string& export_rt_mod); + + /*! + * \brief Create VitisAI runtime from serialized XGraph + * \param symbol_name The name of the function. + * \param xgraph_str serialized XGraph representation + * \param const_names The names of each constant in the sub-graph. + * \param target The Vitis-AI device target (e.g. DPUCADX8G, DPUCZDX8G). + * \param build_dir The directory to be used for Vitis-AI build files. + * \param work_dir The directory to be used for Vitis-AI work files. + * \param export_rt_mod_path The path to the file to be used for exporting the + * PyXIR runtime module. + */ + VitisAIRuntime(const std::string& symbol_name, + const std::string& xgraph_str, + const Array const_names, + const std::string& target, + const std::string& build_dir, + const std::string& work_dir, + const std::string& export_runtime_module_path); + /*! * \brief Get member function to front-end. * \param name The name of the function. * \param sptr_to_self The pointer to the module node. * \return The corresponding member function. */ - virtual PackedFunc GetFunction(const std::string& name, const ObjectPtr& sptr_to_self); + virtual PackedFunc GetFunction(const std::string& name, + const ObjectPtr& sptr_to_self); + /*! * \return The type key of the executor. */ const char* type_key() const { return "VitisAIRuntime"; } - /*! - * \brief Initialize the vai runtime with pyxir. - * \param model_path The compiled model path. - * \param target The name of the target being used - */ -void Init(const std::string& model_path, const std::string& target); /*! * \brief Serialize the content of the pyxir directory and save it to * binary stream. @@ -78,18 +101,22 @@ void Init(const std::string& model_path, const std::string& target); void SaveToBinary(dmlc::Stream* stream) final; private: - /*! \brief The only subgraph name for this module. */ + /*! \brief The only subgraph name for this module */ std::string symbol_name_; - /*! \brief The required constant names. */ + /*! \brief The required constant names */ Array const_names_; - std::shared_ptr xgraph_; + /*! \brief The runtime module */ pyxir::RtModHolder rt_mod_; - std::string model_path_; - std::string target_; + /*! \brief The XGraph input tensor names in the order as provided by TVM */ std::vector in_tensor_names_; + /*! \brief The XGraph output tensor names in the order as provided by TVM */ std::vector out_tensor_names_; + /*! \brief The file path for exporting the runtime module if set */ + std::string export_rt_mod_path_; + /*! \brief Whether constant tensors have been initialized */ bool initialized_{false}; }; + } // namespace runtime } // namespace tvm #endif // TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ diff --git a/tests/python/contrib/test_vitis_ai/__init__.py b/tests/python/contrib/test_vitis_ai/__init__.py new file mode 100644 index 000000000000..c5fe1539b059 --- /dev/null +++ b/tests/python/contrib/test_vitis_ai/__init__.py @@ -0,0 +1,18 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +""" Infrastructure and tests for Vitis-AI codegen """ diff --git a/tests/python/contrib/test_vitis_ai/infrastructure.py b/tests/python/contrib/test_vitis_ai/infrastructure.py new file mode 100644 index 000000000000..1799f42cdd59 --- /dev/null +++ b/tests/python/contrib/test_vitis_ai/infrastructure.py @@ -0,0 +1,143 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 + +"""Expose Vitis-AI test functions to the Python frontend""" + +import sys +import numpy as np + +import pytest +pytest.importorskip('pyxir') +import pyxir.contrib.target.DPUCADX8G +import pyxir.contrib.target.DPUCZDX8G + +import tvm +from tvm import relay +from tvm import runtime +from tvm.relay import transform +from tvm.relay.op.contrib.vitis_ai import annotation +from tvm.relay.build_module import bind_params_by_name +from tvm.contrib.target import vitis_ai +from tvm.contrib import graph_runtime +from tvm.contrib import util + + +def get_cpu_op_count(mod): + """Traverse graph counting ops offloaded to TVM.""" + class Counter(tvm.relay.ExprVisitor): + def __init__(self): + super().__init__() + self.count = 0 + + def visit_call(self, call): + if isinstance(call.op, tvm.ir.Op): + self.count += 1 + + super().visit_call(call) + + c = Counter() + c.visit(mod["main"]) + return c.count + +def skip_test(): + """Skip test if it requires the Vitis-AI codegen and it's not present.""" + if not tvm.get_global_func("relay.ext.vitis_ai", True): + print("Skip test because Vitis-AI codegen is not available.") + return True + return False + +def build_module(mod, target, dpu_target='DPUCADX8G', params=None, + enable_vitis_ai=True, tvm_ops=0, + vitis_ai_partitions=1): + """Build module for Vitis-AI codegen.""" + if isinstance(mod, tvm.relay.expr.Call): + mod = tvm.IRModule.from_expr(mod) + if params is None: + params = {} + + with tvm.transform.PassContext(opt_level=3, + config={'relay.ext.vitis_ai.options.target': dpu_target}): + if enable_vitis_ai: + mod["main"] = bind_params_by_name(mod["main"], params) + mod = annotation(mod, params, dpu_target) + mod = transform.MergeCompilerRegions()(mod) + mod = transform.PartitionGraph()(mod) + tvm_op_count = get_cpu_op_count(mod) + assert tvm_op_count == tvm_ops, \ + "Got {} TVM operators, expected {}".format(tvm_op_count, tvm_ops) + partition_count = 0 + for global_var in mod.get_global_vars(): + if "vitis_ai" in global_var.name_hint: + partition_count += 1 + + assert vitis_ai_partitions == partition_count, \ + "Got {} Vitis-AI partitions, expected {}".format( + partition_count, vitis_ai_partitions) + relay.backend.compile_engine.get().clear() + return relay.build(mod, target, params=params) + +def update_lib(lib, cross_compile=None): + tmp_path = util.tempdir() + lib_name = 'lib.so' + lib_path = tmp_path.relpath(lib_name) + if cross_compile: + lib.export_library(lib_path, cc=cross_compile) + else: + lib.export_library(lib_path) + lib = runtime.load_module(lib_path) + return lib + +def extract_vitis_ai_modules(module): + """Get the Vits-AI runtime module from llvm module.""" + return list(filter(lambda mod: mod.type_key == "VitisAIRuntime", + module.get_lib().imported_modules)) + +def verify_codegen(module, num_vitis_ai_modules=1, params=None, + target='llvm', dpu_target='DPUCADX8G'): + """Check Vitis-AI codegen against a known good output.""" + module = build_module(module, target, params=params, dpu_target=dpu_target) + vitis_ai_modules = extract_vitis_ai_modules(module) + + assert len(vitis_ai_modules) == num_vitis_ai_modules, \ + f"The number of Vitis-AI modules produced ({len(vitis_ai_modules)}) does not " \ + f"match the expected value ({num_vitis_ai_modules})." + + +def verify_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", + ctx=tvm.cpu(), params=None, + dpu_target='DPUCADX8G', tvm_ops=0): + """To check the result between reference and byoc vitis-ai flow""" + + lib = build_module(mod, target, params=params, + dpu_target=dpu_target, tvm_ops=tvm_ops) + lib = update_lib(lib) + ctx = tvm.cpu() + rt_mod = graph_runtime.GraphModule(lib["default"](tvm.cpu())) + + for name, data in map_inputs.items(): + rt_mod.set_input(name, data) + rt_mod.set_input(**params) + rt_mod.run() + + out_shapes = out_shape if isinstance(out_shape, list) else [out_shape] + results = result if isinstance(result, list) else [result] + + for idx, shape in enumerate(out_shapes): + out = tvm.nd.empty(shape, ctx=ctx) + out = rt_mod.get_output(idx, out) + tvm.testing.assert_allclose(out.asnumpy(), results[idx], rtol=tol, atol=tol) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index 9caa8a7ae9a8..8528bdd0d7fa 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -15,13 +15,16 @@ # specific language governing permissions and limitations # under the License. # pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 -"""Vitis-AI codegen tests.""" + +"""Vitis-AI codegen tests""" + import sys import numpy as np import pytest pytest.importorskip('pyxir') import pyxir.contrib.target.DPUCADX8G +import pyxir.contrib.target.DPUCZDX8G import tvm from tvm import relay @@ -30,6 +33,8 @@ from tvm.relay.build_module import bind_params_by_name from tvm.contrib.target import vitis_ai +from .infrastructure import skip_test, verify_codegen + def set_func_attr(func, compile_name, symbol_name): func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) func = func.with_attr("Inline", tvm.tir.IntImm("int32", 1)) @@ -37,67 +42,77 @@ def set_func_attr(func, compile_name, symbol_name): func = func.with_attr("global_symbol", symbol_name) return func -def _create_graph(): - shape = (10, 10) +def test_conv2d(): + """Test conv2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + x = relay.var('x', shape=(1, 3, 224, 224)) + w = relay.const(np.zeros((16, 3, 3, 3), dtype='float32')) + y = relay.nn.conv2d(x, w, strides=[2, 2], padding=[1, 1, 1, 1], kernel_size=[3, 3]) + func = relay.Function([x], y) + params = {} + params["x"] = np.zeros((1, 3, 224, 224), dtype='float32') + params["w"] = np.random.rand(16, 3, 3, 3).astype('float32') mod = tvm.IRModule() - x = relay.var('x', shape=shape) - y = relay.var('y', shape=shape) - z = x + x - p = y * y - func = relay.Function([x, y], p - z) mod["main"] = func - params = {} - params["x"] = np.random.rand(10, 10).astype('float32') - params["y"] = np.random.rand(10, 10).astype('float32') - return mod, params + verify_codegen(mod, params=params, dpu_target='DPUCADX8G') + verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') +def test_depthwise_conv(): + """Test depthwise_conv operator for Vitis-AI DPUCZDX8G-zcu104 target""" -def _construct_model(func, params=None): + dtype = 'float32' + ishape = (1, 32, 14, 14) + wshape = (32, 1, 3, 3) + data = relay.var("data", shape=(ishape), dtype=dtype) + weights = relay.var("weights", shape=(wshape), dtype=dtype) + depthwise_conv2d = relay.nn.conv2d(data, + weights, + kernel_size=(3, 3), + padding=(1, 1), + groups=32) + func = relay.Function([data, weights], depthwise_conv2d) + params = {} + params["weights"] = np.random.randn(32, 1, 3, 3).astype(dtype) + params["data"] = np.random.randn(1, 32, 14, 14).astype(dtype) mod = tvm.IRModule() mod["main"] = func - if params is None: - params = {} - mod["main"] = bind_params_by_name(mod["main"], params) - mod = annotation(mod, params, "DPUCADX8G") - mod = transform.MergeCompilerRegions()(mod) - mod = transform.PartitionGraph()(mod) - fcompile = tvm._ffi.get_global_func("relay.ext.vitis_ai") - subgraph_mod = tvm.IRModule() - for _, funcnode in mod.functions.items(): - if funcnode.attrs and 'Compiler' in funcnode.attrs and \ - funcnode.attrs['Compiler'] == 'vitis_ai': - subgraph_mod["main"] = funcnode - with tvm.transform.PassContext(opt_level=3, \ - config={'relay.ext.vitis_ai.options.target': - 'DPUCADX8G'}): - fcompile(subgraph_mod["main"]) + verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') -def test_add(): - shape = (10, 10) - x = relay.var('x', shape=shape) - y = x + x - func = relay.Function([x], y) - _construct_model(func) +def test_bias_add(): + """Test bias_add operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + dtype = 'float32' + ishape = (1, 32, 14, 14) + data = relay.var("data", shape=(ishape), dtype=dtype) + bias = relay.var("bias", relay.TensorType((32, ), dtype)) + out = relay.nn.bias_add(data, bias) + func = relay.Function([data, bias], out) + params = {} + params["bias"] = np.random.randn(32).astype(dtype) + params["data"] = np.random.randn(1, 32, 14, 14).astype(dtype) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, params=params, dpu_target='DPUCADX8G') + verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + def test_relu(): + """Test relu operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + shape = (10, 10) x = relay.var('x', shape=shape) y = relay.nn.relu(x) func = relay.Function([x], y) - _construct_model(func) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') -def test_conv2d(): - x = relay.var('x', shape=(1, 3, 224, 224)) - w = relay.const(np.zeros((16, 3, 3, 3), dtype='float32')) - y = relay.nn.conv2d(x, w, strides=[2, 2], padding=[1, 1, 1, 1], kernel_size=[3, 3]) - func = relay.Function([x], y) - params = {} - params["x"] = np.zeros((1, 3, 224, 224), dtype='float32') - params["w"] = np.random.rand(16, 3, 3, 3).astype('float32') - _construct_model(func, params) def test_batchnorm(): + """Test batchnorm operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + data = relay.var('data', shape=(1, 16, 112, 112)) bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) @@ -113,25 +128,115 @@ def test_batchnorm(): params["bn_beta"] = np.random.rand(16).astype('float32') params["bn_mean"] = np.random.rand(16).astype('float32') params["bn_var"] = np.random.rand(16).astype('float32') - _construct_model(func, params) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, params=params, dpu_target='DPUCADX8G') + verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + + +def test_add(): + """Test add operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + shape = (10, 10) + x = relay.var('x', shape=shape) + y = x + x + func = relay.Function([x], y) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + def test_global_avg_pool2d(): - shape = (10, 10, 10, 10) + """Test global_avg_pool2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + shape = (10, 10, 7, 7) x = relay.var('x', shape=shape) y = relay.nn.global_avg_pool2d(x) func = relay.Function([x], y) - _construct_model(func) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + def test_avg_pool2d(): + """Test avg_pool2d for operator Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + shape = (10, 10, 10, 10) x = relay.var('x', shape=shape) y = relay.nn.avg_pool2d(x, pool_size=(3, 3)) func = relay.Function([x], y) - _construct_model(func) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + + +def test_max_pool2d(): + """Test max_pool2d for operator Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + shape = (64, 512, 10, 10) + x = relay.var('x', shape=shape) + y = relay.nn.max_pool2d(x, pool_size=(3, 3)) + func = relay.Function([x], y) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + + +def test_global_max_pool2d(): + """Test global_maxpool2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + shape = (1, 512, 7, 7) + x = relay.var('x', shape=shape) + y = relay.nn.global_max_pool2d(x) + func = relay.Function([x], y) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + + +def test_upsampling(): + """Test upsampling operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + shape = (64, 512, 10, 10) + x = relay.var('x', shape=shape) + y = relay.nn.upsampling(x, scale_h=2, scale_w=2) + func = relay.Function([x], y) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, dpu_target='DPUCADX8G') + verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + + +def test_conv2d_transpose(): + """Test conv2d_transpose operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + dshape = (1, 3, 18, 18) + kshape = (3, 10, 3, 3) + x = relay.var("x", shape=dshape) + w = relay.const(np.zeros(kshape, dtype='float32')) + y = relay.nn.conv2d_transpose(x, w, + channels=10, kernel_size=(3, 3), strides=(1, 1), + padding=(1, 1)) + func = relay.Function([x], y) + params = {} + dtype = "float32" + params["x"] = np.random.uniform(size=dshape).astype(dtype) + params["w"] = np.random.uniform(size=kshape).astype(dtype) + mod = tvm.IRModule() + mod["main"] = func + verify_codegen(mod, params=params, dpu_target='DPUCADX8G') + verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + def test_annotate(): - """Test annotation with Vitis-AI DP (DPUCADX8G)""" - def partition(): + """Test annotation operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" + + def partition(dpu_target): data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) @@ -158,7 +263,7 @@ def partition(): params["bn_beta"] = np.random.rand(16).astype('float32') params["bn_mean"] = np.random.rand(16).astype('float32') params["bn_var"] = np.random.rand(16).astype('float32') - mod = annotation(mod, params, "DPUCADX8G") + mod = annotation(mod, params, dpu_target) opt_pass = tvm.transform.Sequential([ transform.MergeCompilerRegions(), @@ -206,10 +311,13 @@ def expected(): bn_mvar], call0) return mod - partitioned = partition() + partitioned_dpuczdx8g_zcu104 = partition('DPUCZDX8G-zcu104') + partitioned_dpucadx8g = partition('DPUCADX8G') + ref_mod = expected() - assert tvm.ir.structural_equal(partitioned, ref_mod, map_free_vars=True) + assert tvm.ir.structural_equal(partitioned_dpuczdx8g_zcu104, ref_mod, map_free_vars=True) + assert tvm.ir.structural_equal(partitioned_dpucadx8g, ref_mod, map_free_vars=True) if __name__ == "__main__": @@ -217,10 +325,16 @@ def expected(): print("Skip test on Windows for now") sys.exit(0) - test_annotate() - test_add() - test_relu() test_conv2d() + test_depthwise_conv() + test_bias_add() + test_relu() + test_add() + test_max_pool2d() + test_global_max_pool2d() test_batchnorm() test_global_avg_pool2d() test_avg_pool2d() + test_upsampling() + test_conv2d_transpose() + test_annotate() diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py deleted file mode 100644 index 45ce7aaf34b2..000000000000 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime.py +++ /dev/null @@ -1,109 +0,0 @@ -# Licensed to the Apache Software Foundation (ASF) under one -# or more contributor license agreements. See the NOTICE file -# distributed with this work for additional information -# regarding copyright ownership. The ASF licenses this file -# to you under the Apache License, Version 2.0 (the -# "License"); you may not use this file except in compliance -# with the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, -# software distributed under the License is distributed on an -# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -# KIND, either express or implied. See the License for the -# specific language governing permissions and limitations -# under the License. -# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 - -""" Vitis-AI runtime test """ - -import sys -import numpy as np - -import pytest -pytest.importorskip('pyxir') -import pyxir.contrib.target.DPUCADX8G - -import tvm -import tvm.relay.testing -from tvm import relay -from tvm import runtime -from tvm.relay import transform -from tvm.contrib import util -from tvm.relay.backend import compile_engine -from tvm.relay.build_module import bind_params_by_name -from tvm.relay.op.contrib.vitis_ai import annotation -from tvm.contrib.target import vitis_ai - - - -def check_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", - ctx=tvm.cpu(), params=None): - """ To check the result between reference and byoc vitis-ai flow""" - - def update_lib(lib): - tmp_path = util.tempdir() - lib_name = 'lib.so' - lib_path = tmp_path.relpath(lib_name) - lib.export_library(lib_path) - lib = runtime.load_module(lib_path) - return lib - - def check_graph_runtime_result(): - compile_engine.get().clear() - with tvm.transform.PassContext(opt_level=3, - config={'relay.ext.vitis_ai.options.target' : 'DPUCADX8G'}): - json, lib, param = relay.build(mod, target=target, params=params) - lib = update_lib(lib) - rt_mod = tvm.contrib.graph_runtime.create(json, lib, ctx) - - for name, data in map_inputs.items(): - rt_mod.set_input(name, data) - rt_mod.set_input(**param) - rt_mod.run() - - out_shapes = out_shape if isinstance(out_shape, list) else [out_shape] - results = result if isinstance(result, list) else [result] - - for idx, shape in enumerate(out_shapes): - out = tvm.nd.empty(shape, ctx=ctx) - out = rt_mod.get_output(idx, out) - - tvm.testing.assert_allclose(out.asnumpy(), results[idx], rtol=tol, atol=tol) - - check_graph_runtime_result() - - -def test_extern_vai_resnet18(): - """Test resnet18 model using Vitis-AI byoc flow""" - if sys.platform == "win32": - print("Skip test on Windows for now") - return - - if not tvm.get_global_func("relay.ext.vitis_ai", True): - print("skip because VITIS-AI codegen is not available") - return - - dtype = 'float32' - ishape = (1, 3, 224, 224) - - mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) - mod["main"] = bind_params_by_name(mod["main"], params) - mod = annotation(mod, params, "DPUCADX8G") - mod = transform.MergeCompilerRegions()(mod) - mod = transform.PartitionGraph()(mod) - - ref_mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) - ref_ex = relay.create_executor("graph", mod=ref_mod, ctx=tvm.cpu(0)) - i_data = np.random.uniform(0, 1, ishape).astype(dtype) - - ref_res = ref_ex.evaluate()(i_data, **params) - - check_result(mod, {"data": i_data}, - (1, 1000), ref_res.asnumpy(), tol=1e-5, params=params) -if __name__ == "__main__": - if sys.platform == "win32": - print("Skip test on Windows for now") - sys.exit(0) - test_extern_vai_resnet18() diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py new file mode 100644 index 000000000000..dd364fb8d0a0 --- /dev/null +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py @@ -0,0 +1,74 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +# pylint: disable=no-else-return, unidiomatic-typecheck, invalid-name, W0611, C0413 + +"""Vitis-AI runtime test for CPU only part + +This test verifies as much as possible whether the a model can be correctly offloaded +and executed for Vitis-AI acceleration. This entails: + - Annotating and partitioning model for Vitis-AI acceleration + - Building a Vitis-AI PyXIR runtime module with on-the-fly quantization enabled + - Run first iteration of on-the-fly quantization flow. This will always be run + on CPU as the first N (parameter) will be used for collecting calibration data + for quantization. + +NOTE This is not a full end-to-end test as we need the full Vitis-AI docker environment +and access to an FPGA instance for that. This test verifies the Vitis-AI flow as much as +possible without requiring access to dedicated docker environment and/or hardware setup. +NOTE Quantization is not being tested (we need to be inside Vitis-AI docker environment +for that) buth the internal representation used for quantization is being generated and +functionally tested (CPU). +""" + +import sys +import numpy as np + +import pytest +pytest.importorskip('pyxir') +import pyxir.contrib.target.DPUCADX8G + +import tvm +import tvm.relay.testing +from tvm import relay + +from .infrastructure import skip_test, verify_result + +def test_extern_vitis_ai_resnet18(): + """Test first part of Vitis-AI on-the-fly quantization runtime with ResNet 18 model""" + if skip_test(): + return + + dtype = 'float32' + ishape = (1, 3, 224, 224) + mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) + ref_mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) + + ref_ex = relay.create_executor("graph", mod=ref_mod, ctx=tvm.cpu(0)) + i_data = np.random.uniform(0, 1, ishape).astype(dtype) + + ref_res = ref_ex.evaluate()(i_data, **params) + verify_result(mod, {"data": i_data}, + (1, 1000), ref_res.asnumpy(), + tol=1e-5, params=params, + dpu_target='DPUCADX8G', tvm_ops=4) + + +if __name__ == "__main__": + if sys.platform == "win32": + print("Skip test on Windows for now") + sys.exit(0) + test_extern_vitis_ai_resnet18() From f022b41cca1b837db9c5439b339b8e3cf669a2b0 Mon Sep 17 00:00:00 2001 From: Jorn Tuyls Date: Fri, 25 Sep 2020 12:55:45 -0700 Subject: [PATCH 15/22] Fix linter errors --- .../contrib/vitis_ai/vitis_ai_runtime.cc | 92 +++++++------------ .../contrib/vitis_ai/vitis_ai_runtime.h | 34 +++---- 2 files changed, 49 insertions(+), 77 deletions(-) diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index 1ef041bc2e7c..6544b8bae57d 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -21,13 +21,13 @@ * \file vitis_ai_runtime.cc */ -#include -#include +#include "vitis_ai_runtime.h" -#include #include +#include -#include "vitis_ai_runtime.h" +#include +#include using namespace pyxir::runtime; @@ -45,14 +45,12 @@ TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.export_runtime_modul /*! \brief (Optional config) Load PyXIR runtime module from disk */ TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.load_runtime_module", String); - -VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, - const Array const_names, +VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, const Array const_names, const std::string& serialized_rt_mod, const std::string& export_rt_mod_path) - : symbol_name_(symbol_name), const_names_(const_names), + : symbol_name_(symbol_name), + const_names_(const_names), export_rt_mod_path_(export_rt_mod_path) { - std::istringstream sstream(serialized_rt_mod); rt_mod_.reset(new RuntimeModule()); rt_mod_->deserialize(sstream); @@ -60,17 +58,13 @@ VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, out_tensor_names_ = rt_mod_->get_out_tensor_names(); } - -VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, - const std::string& xgraph_str, - const Array const_names, - const std::string& target, - const std::string& build_dir, - const std::string& work_dir, +VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, const std::string& xgraph_str, + const Array const_names, const std::string& target, + const std::string& build_dir, const std::string& work_dir, const std::string& export_rt_mod_path) - : symbol_name_(symbol_name), const_names_(const_names), + : symbol_name_(symbol_name), + const_names_(const_names), export_rt_mod_path_(export_rt_mod_path) { - std::istringstream xgraph_sstream(xgraph_str); pyxir::XGraphHolder xgraph = std::make_shared(""); pyxir::read(xgraph, xgraph_sstream); @@ -82,38 +76,28 @@ VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, pyxir::RunOptionsHolder run_options(new pyxir::runtime::RunOptions()); run_options->on_the_fly_quantization = true; run_options->build_dir = build_dir; - if (!work_dir.empty()) - run_options->work_dir = work_dir; - rt_mod_ = pyxir::build_rt(xgraph, target, in_tensor_names_, out_tensor_names_, - "vai", run_options); + if (!work_dir.empty()) run_options->work_dir = work_dir; + rt_mod_ = + pyxir::build_rt(xgraph, target, in_tensor_names_, out_tensor_names_, "vai", run_options); } - -Module VitisAIRuntimeCreate(const std::string& name, - const std::string& xgraph_str, - const std::string& target, - const std::string& build_dir, - const std::string& work_dir, - const std::string& export_rt_mod_path) { +Module VitisAIRuntimeCreate(const std::string& name, const std::string& xgraph_str, + const std::string& target, const std::string& build_dir, + const std::string& work_dir, const std::string& export_rt_mod_path) { Array const_vars; - auto exec = make_object(name, xgraph_str, const_vars, target, - build_dir, work_dir, + auto exec = make_object(name, xgraph_str, const_vars, target, build_dir, work_dir, export_rt_mod_path); return Module(exec); } TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.from_xgraph").set_body([](TVMArgs args, TVMRetValue* rv) { - *rv = VitisAIRuntimeCreate(args[0], args[1], args[2], - args[3], args[4], args[5]); + *rv = VitisAIRuntimeCreate(args[0], args[1], args[2], args[3], args[4], args[5]); }); - -Module VitisAIRuntimeCreate(const std::string& name, - const std::string& serialized_rt_mod, +Module VitisAIRuntimeCreate(const std::string& name, const std::string& serialized_rt_mod, const std::string& export_rt_mod_path) { Array const_vars; - auto exec = make_object(name, const_vars, serialized_rt_mod, - export_rt_mod_path); + auto exec = make_object(name, const_vars, serialized_rt_mod, export_rt_mod_path); return Module(exec); } @@ -128,7 +112,6 @@ TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.from_rt_mod").set_body([](TVMArgs args *rv = VitisAIRuntimeCreate(args[0], serialized_rt_mod, args[2]); }); - Module VitisAIRuntimeLoadFromBinary(void* strm ) { dmlc::Stream* stream = static_cast(strm); std::string symbol_name; @@ -143,8 +126,8 @@ Module VitisAIRuntimeLoadFromBinary(void* strm ) { for (const auto& it : const_vars) { const_names.push_back(it); } - auto exec = make_object(symbol_name, const_names, serialized_rt_mod, - export_rt_mod_path); + auto exec = + make_object(symbol_name, const_names, serialized_rt_mod, export_rt_mod_path); return Module(exec); } @@ -172,7 +155,6 @@ void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { } } - PackedFunc VitisAIRuntime::GetFunction(const std::string& name, const ObjectPtr& sptr_to_self) { if (name == "get_symbol") { @@ -189,34 +171,30 @@ PackedFunc VitisAIRuntime::GetFunction(const std::string& name, *rv = 0; }); } else if (this->symbol_name_ == name) { - return PackedFunc( - [sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { + return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) { // Initialize input tensors DLTensor* inputs = args[0]; std::vector in_tensors; std::vector in_shape; - for (int i = 0; i < inputs->ndim; ++i) - in_shape.push_back(inputs->shape[i]); - in_tensors.push_back( - std::shared_ptr(new pyxir::XBuffer(reinterpret_cast(static_cast(inputs->data)), - 4, "f", in_shape.size(), in_shape, false, false))); + for (int i = 0; i < inputs->ndim; ++i) in_shape.push_back(inputs->shape[i]); + in_tensors.push_back(std::shared_ptr( + new pyxir::XBuffer(reinterpret_cast(static_cast(inputs->data)), 4, "f", + in_shape.size(), in_shape, false, false))); // Initialize output tensors std::vector out_tensors; for (unsigned i = 0; i < out_tensor_names_.size(); ++i) { - DLTensor* output_tensor = args[args.size() - out_tensor_names_.size()+i]; + DLTensor* output_tensor = args[args.size() - out_tensor_names_.size() + i]; std::vector out_shape; - for (int i = 0; i < output_tensor->ndim; ++i) - out_shape.push_back(output_tensor->shape[i]); - void* output_data = reinterpret_cast (static_cast(output_tensor->data)); - out_tensors.push_back( - std::shared_ptr(new pyxir::XBuffer(output_data, 4, "f", out_shape.size(), - out_shape, false, false))); + for (int i = 0; i < output_tensor->ndim; ++i) out_shape.push_back(output_tensor->shape[i]); + void* output_data = reinterpret_cast(static_cast(output_tensor->data)); + out_tensors.push_back(std::shared_ptr( + new pyxir::XBuffer(output_data, 4, "f", out_shape.size(), out_shape, false, false))); } // Execute the subgraph. rt_mod_->execute(in_tensors, out_tensors); - }); + }); } else { return PackedFunc(); } diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h index 06086083d0e0..05021832e8e2 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h @@ -22,20 +22,20 @@ * containing only tvm PackedFunc. * \file vitis_ai_runtime.h */ -#ifndef TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ -#define TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ +#ifndef TVM_RUNTIME_CONTRIB_VITIS_AI_VITIS_AI_RUNTIME_H_ +#define TVM_RUNTIME_CONTRIB_VITIS_AI_VITIS_AI_RUNTIME_H_ + +#include +#include +#include #include #include #include -#include -#include -#include #include #include - namespace tvm { namespace runtime { @@ -47,7 +47,7 @@ namespace runtime { */ class VitisAIRuntime : public ModuleNode { public: - /*! + /*! * \brief Create VitisAI runtime from serialized XGraph * \param symbol_name The name of the function. * \param const_names The names of each constant in the sub-graph. @@ -55,10 +55,8 @@ class VitisAIRuntime : public ModuleNode { * \param export_rt_mod_path The path to the file to be used for exporting the * PyXIR runtime module. */ - VitisAIRuntime(const std::string& symbol_name, - const Array const_names, - const std::string& serialized_rt_mod, - const std::string& export_rt_mod); + VitisAIRuntime(const std::string& symbol_name, const Array const_names, + const std::string& serialized_rt_mod, const std::string& export_rt_mod); /*! * \brief Create VitisAI runtime from serialized XGraph @@ -71,12 +69,9 @@ class VitisAIRuntime : public ModuleNode { * \param export_rt_mod_path The path to the file to be used for exporting the * PyXIR runtime module. */ - VitisAIRuntime(const std::string& symbol_name, - const std::string& xgraph_str, - const Array const_names, - const std::string& target, - const std::string& build_dir, - const std::string& work_dir, + VitisAIRuntime(const std::string& symbol_name, const std::string& xgraph_str, + const Array const_names, const std::string& target, + const std::string& build_dir, const std::string& work_dir, const std::string& export_runtime_module_path); /*! @@ -85,8 +80,7 @@ class VitisAIRuntime : public ModuleNode { * \param sptr_to_self The pointer to the module node. * \return The corresponding member function. */ - virtual PackedFunc GetFunction(const std::string& name, - const ObjectPtr& sptr_to_self); + virtual PackedFunc GetFunction(const std::string& name, const ObjectPtr& sptr_to_self); /*! * \return The type key of the executor. @@ -119,4 +113,4 @@ class VitisAIRuntime : public ModuleNode { } // namespace runtime } // namespace tvm -#endif // TVM_RUNTIME_CONTRIB_VITIS_AI_RUNTIME_H_ +#endif // TVM_RUNTIME_CONTRIB_VITIS_AI_VITIS_AI_RUNTIME_H_ From 287f285facf58edbce2ec27a0b350c28ca55c35d Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Mon, 28 Sep 2020 06:34:45 -0600 Subject: [PATCH 16/22] Fix linter errors --- python/tvm/contrib/target/vitis_ai.py | 90 ++++---- python/tvm/relay/op/contrib/vitis_ai.py | 18 +- .../contrib/vitis_ai/vitis_ai_runtime.cc | 12 +- .../contrib/vitis_ai/vitis_ai_runtime.h | 5 +- .../contrib/test_vitis_ai/infrastructure.py | 78 ++++--- .../test_vitis_ai/test_vitis_ai_codegen.py | 200 +++++++++--------- .../test_vitis_ai_runtime_cpu_part.py | 20 +- 7 files changed, 236 insertions(+), 187 deletions(-) diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index 4ae35853b2f3..2fa95b31b7ce 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -38,8 +38,9 @@ def __init__(self, model_name, function): def convert_pyxir(self, target): """Convert Relay expression to PyXIR XGraph""" - xgraph = pyxir.frontend.tvm.from_relay(self.function, - params=self.params, postprocessing=None) + xgraph = pyxir.frontend.tvm.from_relay( + self.function, params=self.params, postprocessing=None + ) xgraph = pyxir.partition(xgraph, targets=[target]) return xgraph @@ -71,40 +72,56 @@ def vitis_ai_compiler(ref): pass_context = tvm.get_global_func("transform.GetCurrentPassContext")() # The target Vitis-AI accelerator device - target = str(pass_context.config['relay.ext.vitis_ai.options.target']) \ - if 'relay.ext.vitis_ai.options.target' in pass_context.config else None + target = ( + str(pass_context.config["relay.ext.vitis_ai.options.target"]) + if "relay.ext.vitis_ai.options.target" in pass_context.config + else None + ) # (Optional configs) The build and work directories to be used by Vitis-AI - vai_build_dir = str(pass_context.config['relay.ext.vitis_ai.options.build_dir']) \ - if 'relay.ext.vitis_ai.options.build_dir' in pass_context.config else \ - tvm.contrib.util.tempdir().relpath("") - vai_work_dir = str(pass_context.config['relay.ext.vitis_ai.options.work_dir']) \ - if 'relay.ext.vitis_ai.options.work_dir' in pass_context.config else \ - tvm.contrib.util.tempdir().relpath("") + vai_build_dir = ( + str(pass_context.config["relay.ext.vitis_ai.options.build_dir"]) + if "relay.ext.vitis_ai.options.build_dir" in pass_context.config + else tvm.contrib.util.tempdir().relpath("") + ) + vai_work_dir = ( + str(pass_context.config["relay.ext.vitis_ai.options.work_dir"]) + if "relay.ext.vitis_ai.options.work_dir" in pass_context.config + else tvm.contrib.util.tempdir().relpath("") + ) # (Optional configs) Export and load PyXIR runtime module to file if provided. This is used to # compile and quantize a model on the host and deploy it at the edge - export_runtime_module = \ - str(pass_context.config['relay.ext.vitis_ai.options.export_runtime_module']) \ - if 'relay.ext.vitis_ai.options.export_runtime_module' in pass_context.config else "" - load_runtime_module = \ - str(pass_context.config['relay.ext.vitis_ai.options.load_runtime_module']) \ - if 'relay.ext.vitis_ai.options.load_runtime_module' in pass_context.config else "" + export_runtime_module = ( + str(pass_context.config["relay.ext.vitis_ai.options.export_runtime_module"]) + if "relay.ext.vitis_ai.options.export_runtime_module" in pass_context.config + else "" + ) + load_runtime_module = ( + str(pass_context.config["relay.ext.vitis_ai.options.load_runtime_module"]) + if "relay.ext.vitis_ai.options.load_runtime_module" in pass_context.config + else "" + ) # Config checks - if load_runtime_module != "" and target is not None: - warnings.warn("Both `load_runtime_module` and `target` configs were specified." - " The `load_runtime_module` points to a prebuilt runtime module with" - " an internal target so the `target` config will be ignored") - if load_runtime_module != "" and 'relay.ext.vitis_ai.options.build_dir' in pass_context.config: - warnings.warn("Both `load_runtime_module` and `build_dir` configs were specified." - " The `load_runtime_module` points to a prebuilt runtime module with" - " an internal build directory so the `build_dir` config will be ignored") - if load_runtime_module != "" and 'relay.ext.vitis_ai.options.work_dir' in pass_context.config: - warnings.warn("Both `load_runtime_module` and `work_dir` configs were specified." - " The `load_runtime_module` points to a prebuilt runtime module with" - " an internal work directory so the `work_dir` config will be ignored") - + if load_runtime_module and target is not None: + warnings.warn( + "Both `load_runtime_module` and `target` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal target so the `target` config will be ignored" + ) + if load_runtime_module and "relay.ext.vitis_ai.options.build_dir" in pass_context.config: + warnings.warn( + "Both `load_runtime_module` and `build_dir` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal build directory so the `build_dir` config will be ignored" + ) + if load_runtime_module and "relay.ext.vitis_ai.options.work_dir" in pass_context.config: + warnings.warn( + "Both `load_runtime_module` and `work_dir` configs were specified." + " The `load_runtime_module` points to a prebuilt runtime module with" + " an internal work directory so the `work_dir` config will be ignored" + ) # If load_runtime_module is not set, we will build the PyXIR runtime module from scratch if load_runtime_module == "": @@ -118,20 +135,21 @@ def vitis_ai_compiler(ref): out_tensor_names = [] for layer in layers: if not layer.internal: - for relay_id in layer.attrs['relay_id']: + for relay_id in layer.attrs["relay_id"]: if relay_id in output_relay_ids: out_tensor_names.append(layer.name) break - if len(out_tensor_names) == 0: - raise ValueError("During codegeneration the loading of subexpression \ - failed due to output tensor name mismatch in Relay PyXIR interface.") - xgraph.meta_attrs['tvm_out_tensors'] = out_tensor_names + if not out_tensor_names: + raise ValueError( + "During codegeneration the loading of subexpression \ + failed due to output tensor name mismatch in Relay PyXIR interface." + ) + xgraph.meta_attrs["tvm_out_tensors"] = out_tensor_names xgraph_str = pyxir.get_xgraph_str(xgraph) runtime_func = "tvm.vitis_ai_runtime.from_xgraph" fcreate = tvm._ffi.get_global_func(runtime_func) - return fcreate(name, xgraph_str, target, vai_build_dir, vai_work_dir, - export_runtime_module) + return fcreate(name, xgraph_str, target, vai_build_dir, vai_work_dir, export_runtime_module) runtime_func = "tvm.vitis_ai_runtime.from_rt_mod" fcreate = tvm._ffi.get_global_func(runtime_func) diff --git a/python/tvm/relay/op/contrib/vitis_ai.py b/python/tvm/relay/op/contrib/vitis_ai.py index f9ba4e1a6916..fa17c63fc00a 100644 --- a/python/tvm/relay/op/contrib/vitis_ai.py +++ b/python/tvm/relay/op/contrib/vitis_ai.py @@ -60,9 +60,8 @@ def visit_tuple(self, tup): def visit_tuple_getitem(self, op): """Add compiler_begin and compiler_end annotations to TupleGetItem""" - if int(hash(op.tuple_value)) in annotator.relay_ids: - tuple_value = compiler_begin(super().visit(op.tuple_value), - annotator.compiler) + if int(hash(op.tuple_value)) in annotator.relay_ids: + tuple_value = compiler_begin(super().visit(op.tuple_value), annotator.compiler) return compiler_end(TupleGetItem(tuple_value, op.index), annotator.compiler) else: tuple_value = super().visit(op.tuple_value) @@ -73,11 +72,9 @@ def visit_call(self, call): if int(hash(call)) in annotator.relay_ids: new_args = [] for arg in call.args: - ann = compiler_begin(super().visit(arg), - annotator.compiler) + ann = compiler_begin(super().visit(arg), annotator.compiler) new_args.append(ann) - new_call = relay.Call(call.op, new_args, call.attrs, - call.type_args) + new_call = relay.Call(call.op, new_args, call.attrs, call.type_args) return compiler_end(new_call, annotator.compiler) else: @@ -92,8 +89,11 @@ def annotation(mod, params, target): xgraph = pyxir.partition(xgraph, targets=[target]) layers = xgraph.get_layers() - relay_ids = [list(np.array(layer.attrs['relay_id']).flatten()) - for layer in layers if layer.target == target] + relay_ids = [ + list(np.array(layer.attrs["relay_id"]).flatten()) + for layer in layers + if layer.target == target + ] relay_ids_flatten = [item for sublist in relay_ids for item in sublist] mod = VitisAIAnnotationPass("vitis_ai", relay_ids_flatten)(mod) diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index 6544b8bae57d..35075aba4a93 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -28,6 +28,8 @@ #include #include +#include +#include using namespace pyxir::runtime; @@ -112,7 +114,7 @@ TVM_REGISTER_GLOBAL("tvm.vitis_ai_runtime.from_rt_mod").set_body([](TVMArgs args *rv = VitisAIRuntimeCreate(args[0], serialized_rt_mod, args[2]); }); -Module VitisAIRuntimeLoadFromBinary(void* strm ) { +Module VitisAIRuntimeLoadFromBinary(void* strm) { dmlc::Stream* stream = static_cast(strm); std::string symbol_name; std::vector const_vars; @@ -132,9 +134,9 @@ Module VitisAIRuntimeLoadFromBinary(void* strm ) { } TVM_REGISTER_GLOBAL("runtime.module.loadbinary_VitisAIRuntime") - .set_body_typed(VitisAIRuntimeLoadFromBinary); + .set_body_typed(VitisAIRuntimeLoadFromBinary); -void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { +void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { std::ostringstream sstream; rt_mod_->serialize(sstream); stream->Write(sstream.str()); @@ -143,10 +145,10 @@ void VitisAIRuntime::SaveToBinary(dmlc::Stream* stream) { std::vector consts; for (const auto& it : const_names_) { consts.push_back(it); - } + } stream->Write(consts); - // If export_runtime_module_ member variable is set, we will additionally export the PyXIR + // If export_rt_mod_path_ member variable is set, we will additionally export the PyXIR // runtime_module to the specified file if (!export_rt_mod_path_.empty()) { std::ofstream out_file(export_rt_mod_path_); diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h index 05021832e8e2..1092bc0ba27b 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.h @@ -24,15 +24,14 @@ */ #ifndef TVM_RUNTIME_CONTRIB_VITIS_AI_VITIS_AI_RUNTIME_H_ #define TVM_RUNTIME_CONTRIB_VITIS_AI_VITIS_AI_RUNTIME_H_ - #include #include #include - +// clang-format off #include #include #include - +// clang-format on #include #include diff --git a/tests/python/contrib/test_vitis_ai/infrastructure.py b/tests/python/contrib/test_vitis_ai/infrastructure.py index 1799f42cdd59..ff7330602776 100644 --- a/tests/python/contrib/test_vitis_ai/infrastructure.py +++ b/tests/python/contrib/test_vitis_ai/infrastructure.py @@ -22,7 +22,8 @@ import numpy as np import pytest -pytest.importorskip('pyxir') + +pytest.importorskip("pyxir") import pyxir.contrib.target.DPUCADX8G import pyxir.contrib.target.DPUCZDX8G @@ -39,6 +40,7 @@ def get_cpu_op_count(mod): """Traverse graph counting ops offloaded to TVM.""" + class Counter(tvm.relay.ExprVisitor): def __init__(self): super().__init__() @@ -54,6 +56,7 @@ def visit_call(self, call): c.visit(mod["main"]) return c.count + def skip_test(): """Skip test if it requires the Vitis-AI codegen and it's not present.""" if not tvm.get_global_func("relay.ext.vitis_ai", True): @@ -61,39 +64,51 @@ def skip_test(): return True return False -def build_module(mod, target, dpu_target='DPUCADX8G', params=None, - enable_vitis_ai=True, tvm_ops=0, - vitis_ai_partitions=1): + +def build_module( + mod, + target, + dpu_target="DPUCADX8G", + params=None, + enable_vitis_ai=True, + tvm_ops=0, + vitis_ai_partitions=1, +): """Build module for Vitis-AI codegen.""" if isinstance(mod, tvm.relay.expr.Call): mod = tvm.IRModule.from_expr(mod) if params is None: params = {} - with tvm.transform.PassContext(opt_level=3, - config={'relay.ext.vitis_ai.options.target': dpu_target}): + with tvm.transform.PassContext( + opt_level=3, config={"relay.ext.vitis_ai.options.target": dpu_target} + ): if enable_vitis_ai: mod["main"] = bind_params_by_name(mod["main"], params) mod = annotation(mod, params, dpu_target) mod = transform.MergeCompilerRegions()(mod) mod = transform.PartitionGraph()(mod) tvm_op_count = get_cpu_op_count(mod) - assert tvm_op_count == tvm_ops, \ - "Got {} TVM operators, expected {}".format(tvm_op_count, tvm_ops) + assert tvm_op_count == tvm_ops, "Got {} TVM operators, expected {}".format( + tvm_op_count, tvm_ops + ) partition_count = 0 for global_var in mod.get_global_vars(): if "vitis_ai" in global_var.name_hint: partition_count += 1 - assert vitis_ai_partitions == partition_count, \ - "Got {} Vitis-AI partitions, expected {}".format( - partition_count, vitis_ai_partitions) + assert ( + vitis_ai_partitions == partition_count + ), "Got {} Vitis-AI partitions, expected {}".format( + partition_count, vitis_ai_partitions + ) relay.backend.compile_engine.get().clear() return relay.build(mod, target, params=params) + def update_lib(lib, cross_compile=None): tmp_path = util.tempdir() - lib_name = 'lib.so' + lib_name = "lib.so" lib_path = tmp_path.relpath(lib_name) if cross_compile: lib.export_library(lib_path, cc=cross_compile) @@ -102,29 +117,42 @@ def update_lib(lib, cross_compile=None): lib = runtime.load_module(lib_path) return lib + def extract_vitis_ai_modules(module): """Get the Vits-AI runtime module from llvm module.""" - return list(filter(lambda mod: mod.type_key == "VitisAIRuntime", - module.get_lib().imported_modules)) + return list( + filter(lambda mod: mod.type_key == "VitisAIRuntime", module.get_lib().imported_modules) + ) + -def verify_codegen(module, num_vitis_ai_modules=1, params=None, - target='llvm', dpu_target='DPUCADX8G'): +def verify_codegen( + module, num_vitis_ai_modules=1, params=None, target="llvm", dpu_target="DPUCADX8G" +): """Check Vitis-AI codegen against a known good output.""" module = build_module(module, target, params=params, dpu_target=dpu_target) vitis_ai_modules = extract_vitis_ai_modules(module) - assert len(vitis_ai_modules) == num_vitis_ai_modules, \ - f"The number of Vitis-AI modules produced ({len(vitis_ai_modules)}) does not " \ + assert len(vitis_ai_modules) == num_vitis_ai_modules, ( + f"The number of Vitis-AI modules produced ({len(vitis_ai_modules)}) does not " f"match the expected value ({num_vitis_ai_modules})." - - -def verify_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", - ctx=tvm.cpu(), params=None, - dpu_target='DPUCADX8G', tvm_ops=0): + ) + + +def verify_result( + mod, + map_inputs, + out_shape, + result, + tol=1e-5, + target="llvm", + ctx=tvm.cpu(), + params=None, + dpu_target="DPUCADX8G", + tvm_ops=0, +): """To check the result between reference and byoc vitis-ai flow""" - lib = build_module(mod, target, params=params, - dpu_target=dpu_target, tvm_ops=tvm_ops) + lib = build_module(mod, target, params=params, dpu_target=dpu_target, tvm_ops=tvm_ops) lib = update_lib(lib) ctx = tvm.cpu() rt_mod = graph_runtime.GraphModule(lib["default"](tvm.cpu())) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index 8528bdd0d7fa..4310ab6d9ced 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -22,7 +22,8 @@ import numpy as np import pytest -pytest.importorskip('pyxir') + +pytest.importorskip("pyxir") import pyxir.contrib.target.DPUCADX8G import pyxir.contrib.target.DPUCZDX8G @@ -35,6 +36,7 @@ from .infrastructure import skip_test, verify_codegen + def set_func_attr(func, compile_name, symbol_name): func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) func = func.with_attr("Inline", tvm.tir.IntImm("int32", 1)) @@ -42,50 +44,48 @@ def set_func_attr(func, compile_name, symbol_name): func = func.with_attr("global_symbol", symbol_name) return func + def test_conv2d(): """Test conv2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" - x = relay.var('x', shape=(1, 3, 224, 224)) - w = relay.const(np.zeros((16, 3, 3, 3), dtype='float32')) + x = relay.var("x", shape=(1, 3, 224, 224)) + w = relay.const(np.zeros((16, 3, 3, 3), dtype="float32")) y = relay.nn.conv2d(x, w, strides=[2, 2], padding=[1, 1, 1, 1], kernel_size=[3, 3]) func = relay.Function([x], y) params = {} - params["x"] = np.zeros((1, 3, 224, 224), dtype='float32') - params["w"] = np.random.rand(16, 3, 3, 3).astype('float32') + params["x"] = np.zeros((1, 3, 224, 224), dtype="float32") + params["w"] = np.random.rand(16, 3, 3, 3).astype("float32") mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, params=params, dpu_target='DPUCADX8G') - verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, params=params, dpu_target="DPUCADX8G") + verify_codegen(mod, params=params, dpu_target="DPUCZDX8G-zcu104") + def test_depthwise_conv(): """Test depthwise_conv operator for Vitis-AI DPUCZDX8G-zcu104 target""" - dtype = 'float32' + dtype = "float32" ishape = (1, 32, 14, 14) wshape = (32, 1, 3, 3) data = relay.var("data", shape=(ishape), dtype=dtype) weights = relay.var("weights", shape=(wshape), dtype=dtype) - depthwise_conv2d = relay.nn.conv2d(data, - weights, - kernel_size=(3, 3), - padding=(1, 1), - groups=32) + depthwise_conv2d = relay.nn.conv2d(data, weights, kernel_size=(3, 3), padding=(1, 1), groups=32) func = relay.Function([data, weights], depthwise_conv2d) params = {} params["weights"] = np.random.randn(32, 1, 3, 3).astype(dtype) params["data"] = np.random.randn(1, 32, 14, 14).astype(dtype) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, params=params, dpu_target="DPUCZDX8G-zcu104") def test_bias_add(): """Test bias_add operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" - dtype = 'float32' + dtype = "float32" ishape = (1, 32, 14, 14) data = relay.var("data", shape=(ishape), dtype=dtype) - bias = relay.var("bias", relay.TensorType((32, ), dtype)) + bias = relay.var("bias", relay.TensorType((32,), dtype)) out = relay.nn.bias_add(data, bias) func = relay.Function([data, bias], out) params = {} @@ -93,123 +93,121 @@ def test_bias_add(): params["data"] = np.random.randn(1, 32, 14, 14).astype(dtype) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, params=params, dpu_target='DPUCADX8G') - verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, params=params, dpu_target="DPUCADX8G") + verify_codegen(mod, params=params, dpu_target="DPUCZDX8G-zcu104") def test_relu(): """Test relu operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (10, 10) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.relu(x) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_batchnorm(): """Test batchnorm operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" - data = relay.var('data', shape=(1, 16, 112, 112)) - bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) - bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) - bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) - bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) - bn_output = relay.nn.batch_norm(data, bn_gamma, bn_beta, bn_mmean, - bn_mvar) - func = relay.Function([data, bn_gamma, bn_beta, bn_mmean, - bn_mvar], bn_output[0]) + data = relay.var("data", shape=(1, 16, 112, 112)) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16,), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16,), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16,), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16,), "float32")) + bn_output = relay.nn.batch_norm(data, bn_gamma, bn_beta, bn_mmean, bn_mvar) + func = relay.Function([data, bn_gamma, bn_beta, bn_mmean, bn_mvar], bn_output[0]) params = {} - params["data"] = np.zeros((1, 16, 112, 112), dtype='float32') - params["bn_gamma"] = np.random.rand(16).astype('float32') - params["bn_beta"] = np.random.rand(16).astype('float32') - params["bn_mean"] = np.random.rand(16).astype('float32') - params["bn_var"] = np.random.rand(16).astype('float32') + params["data"] = np.zeros((1, 16, 112, 112), dtype="float32") + params["bn_gamma"] = np.random.rand(16).astype("float32") + params["bn_beta"] = np.random.rand(16).astype("float32") + params["bn_mean"] = np.random.rand(16).astype("float32") + params["bn_var"] = np.random.rand(16).astype("float32") mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, params=params, dpu_target='DPUCADX8G') - verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, params=params, dpu_target="DPUCADX8G") + verify_codegen(mod, params=params, dpu_target="DPUCZDX8G-zcu104") def test_add(): """Test add operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (10, 10) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = x + x func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_global_avg_pool2d(): """Test global_avg_pool2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (10, 10, 7, 7) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.global_avg_pool2d(x) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_avg_pool2d(): """Test avg_pool2d for operator Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (10, 10, 10, 10) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.avg_pool2d(x, pool_size=(3, 3)) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_max_pool2d(): """Test max_pool2d for operator Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (64, 512, 10, 10) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.max_pool2d(x, pool_size=(3, 3)) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_global_max_pool2d(): """Test global_maxpool2d operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (1, 512, 7, 7) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.global_max_pool2d(x) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_upsampling(): """Test upsampling operator for Vitis-AI DPUCADX8G and DPUCZDX8G-zcu104 targets""" shape = (64, 512, 10, 10) - x = relay.var('x', shape=shape) + x = relay.var("x", shape=shape) y = relay.nn.upsampling(x, scale_h=2, scale_w=2) func = relay.Function([x], y) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, dpu_target='DPUCADX8G') - verify_codegen(mod, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, dpu_target="DPUCADX8G") + verify_codegen(mod, dpu_target="DPUCZDX8G-zcu104") def test_conv2d_transpose(): @@ -218,10 +216,10 @@ def test_conv2d_transpose(): dshape = (1, 3, 18, 18) kshape = (3, 10, 3, 3) x = relay.var("x", shape=dshape) - w = relay.const(np.zeros(kshape, dtype='float32')) - y = relay.nn.conv2d_transpose(x, w, - channels=10, kernel_size=(3, 3), strides=(1, 1), - padding=(1, 1)) + w = relay.const(np.zeros(kshape, dtype="float32")) + y = relay.nn.conv2d_transpose( + x, w, channels=10, kernel_size=(3, 3), strides=(1, 1), padding=(1, 1) + ) func = relay.Function([x], y) params = {} dtype = "float32" @@ -229,8 +227,8 @@ def test_conv2d_transpose(): params["w"] = np.random.uniform(size=kshape).astype(dtype) mod = tvm.IRModule() mod["main"] = func - verify_codegen(mod, params=params, dpu_target='DPUCADX8G') - verify_codegen(mod, params=params, dpu_target='DPUCZDX8G-zcu104') + verify_codegen(mod, params=params, dpu_target="DPUCADX8G") + verify_codegen(mod, params=params, dpu_target="DPUCZDX8G-zcu104") def test_annotate(): @@ -239,36 +237,35 @@ def test_annotate(): def partition(dpu_target): data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) - bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) - bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) - bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) - bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16,), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16,), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16,), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16,), "float32")) conv = relay.nn.conv2d( - data=data, - weight=weight, - kernel_size=(3, 3), - channels=16, - padding=(1, 1)) - bn_output = relay.nn.batch_norm(conv, bn_gamma, bn_beta, bn_mmean, - bn_mvar) - - func = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, - bn_mvar], bn_output.astuple()) + data=data, weight=weight, kernel_size=(3, 3), channels=16, padding=(1, 1) + ) + bn_output = relay.nn.batch_norm(conv, bn_gamma, bn_beta, bn_mmean, bn_mvar) + + func = relay.Function( + [data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar], bn_output.astuple() + ) mod = tvm.IRModule() mod["main"] = func params = {} - params["weight"] = np.random.rand(16, 3, 3, 3).astype('float32') - params["bn_gamma"] = np.random.rand(16).astype('float32') - params["bn_beta"] = np.random.rand(16).astype('float32') - params["bn_mean"] = np.random.rand(16).astype('float32') - params["bn_var"] = np.random.rand(16).astype('float32') + params["weight"] = np.random.rand(16, 3, 3, 3).astype("float32") + params["bn_gamma"] = np.random.rand(16).astype("float32") + params["bn_beta"] = np.random.rand(16).astype("float32") + params["bn_mean"] = np.random.rand(16).astype("float32") + params["bn_var"] = np.random.rand(16).astype("float32") mod = annotation(mod, params, dpu_target) - opt_pass = tvm.transform.Sequential([ - transform.MergeCompilerRegions(), - transform.PartitionGraph(), - ]) + opt_pass = tvm.transform.Sequential( + [ + transform.MergeCompilerRegions(), + transform.PartitionGraph(), + ] + ) with tvm.transform.PassContext(opt_level=3): mod = opt_pass(mod) @@ -280,20 +277,18 @@ def expected(): data0 = relay.var("data0", relay.TensorType((1, 3, 224, 224), "float32")) weight0 = relay.var("weight0", relay.TensorType((16, 3, 3, 3), "float32")) conv = relay.nn.conv2d( - data=data0, - weight=weight0, - kernel_size=(3, 3), - channels=16, - padding=(1, 1)) + data=data0, weight=weight0, kernel_size=(3, 3), channels=16, padding=(1, 1) + ) # function variables for batch_norm - bn_gamma0 = relay.var("bn_gamma0", relay.TensorType((16, ), "float32")) - bn_beta0 = relay.var("bn_beta0", relay.TensorType((16, ), "float32")) - bn_mmean0 = relay.var("bn_mean0", relay.TensorType((16, ), "float32")) - bn_mvar0 = relay.var("bn_var0", relay.TensorType((16, ), "float32")) + bn_gamma0 = relay.var("bn_gamma0", relay.TensorType((16,), "float32")) + bn_beta0 = relay.var("bn_beta0", relay.TensorType((16,), "float32")) + bn_mmean0 = relay.var("bn_mean0", relay.TensorType((16,), "float32")) + bn_mvar0 = relay.var("bn_var0", relay.TensorType((16,), "float32")) bn = relay.nn.batch_norm(conv, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0) - func0 = relay.Function([data0, weight0, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0], - bn.astuple()) + func0 = relay.Function( + [data0, weight0, bn_gamma0, bn_beta0, bn_mmean0, bn_mvar0], bn.astuple() + ) func0 = set_func_attr(func0, "vitis_ai", "vitis_ai_0") gv0 = relay.GlobalVar("vitis_ai_0") mod = tvm.IRModule() @@ -302,17 +297,16 @@ def expected(): # main function data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) weight = relay.var("weight", relay.TensorType((16, 3, 3, 3), "float32")) - bn_gamma = relay.var("bn_gamma", relay.TensorType((16, ), "float32")) - bn_beta = relay.var("bn_beta", relay.TensorType((16, ), "float32")) - bn_mmean = relay.var("bn_mean", relay.TensorType((16, ), "float32")) - bn_mvar = relay.var("bn_var", relay.TensorType((16, ), "float32")) + bn_gamma = relay.var("bn_gamma", relay.TensorType((16,), "float32")) + bn_beta = relay.var("bn_beta", relay.TensorType((16,), "float32")) + bn_mmean = relay.var("bn_mean", relay.TensorType((16,), "float32")) + bn_mvar = relay.var("bn_var", relay.TensorType((16,), "float32")) call0 = gv0(data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar) - mod["main"] = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, - bn_mvar], call0) + mod["main"] = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar], call0) return mod - partitioned_dpuczdx8g_zcu104 = partition('DPUCZDX8G-zcu104') - partitioned_dpucadx8g = partition('DPUCADX8G') + partitioned_dpuczdx8g_zcu104 = partition("DPUCZDX8G-zcu104") + partitioned_dpucadx8g = partition("DPUCADX8G") ref_mod = expected() diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py index dd364fb8d0a0..030dda372cfe 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_runtime_cpu_part.py @@ -38,7 +38,8 @@ import numpy as np import pytest -pytest.importorskip('pyxir') + +pytest.importorskip("pyxir") import pyxir.contrib.target.DPUCADX8G import tvm @@ -47,12 +48,13 @@ from .infrastructure import skip_test, verify_result + def test_extern_vitis_ai_resnet18(): """Test first part of Vitis-AI on-the-fly quantization runtime with ResNet 18 model""" if skip_test(): return - dtype = 'float32' + dtype = "float32" ishape = (1, 3, 224, 224) mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) ref_mod, params = relay.testing.resnet.get_workload(num_layers=18, batch_size=1) @@ -61,10 +63,16 @@ def test_extern_vitis_ai_resnet18(): i_data = np.random.uniform(0, 1, ishape).astype(dtype) ref_res = ref_ex.evaluate()(i_data, **params) - verify_result(mod, {"data": i_data}, - (1, 1000), ref_res.asnumpy(), - tol=1e-5, params=params, - dpu_target='DPUCADX8G', tvm_ops=4) + verify_result( + mod, + {"data": i_data}, + (1, 1000), + ref_res.asnumpy(), + tol=1e-5, + params=params, + dpu_target="DPUCADX8G", + tvm_ops=4, + ) if __name__ == "__main__": From dcdcf5c157d61a155b254a590dd0582159de2248 Mon Sep 17 00:00:00 2001 From: Jorn Tuyls Date: Mon, 28 Sep 2020 08:40:37 -0700 Subject: [PATCH 17/22] Address sphinx warnings --- docs/deploy/index.rst | 1 + docs/deploy/vitis_ai.rst | 21 +++++++++++---------- 2 files changed, 12 insertions(+), 10 deletions(-) diff --git a/docs/deploy/index.rst b/docs/deploy/index.rst index b38a7f561ab3..1770b5674d16 100644 --- a/docs/deploy/index.rst +++ b/docs/deploy/index.rst @@ -69,3 +69,4 @@ target device without relying on RPC. see the following resources on how to do s integrate hls arm_compute_lib + vitis_ai diff --git a/docs/deploy/vitis_ai.rst b/docs/deploy/vitis_ai.rst index f5bda10cddea..f0bd3edcd6e2 100755 --- a/docs/deploy/vitis_ai.rst +++ b/docs/deploy/vitis_ai.rst @@ -100,7 +100,7 @@ Hardware setup and docker build .. code:: bash - git clone --recurse-submodules https://github.com/Xilinx/Vitis-AI + git clone --recurse-submodules https://github.com/Xilinx/Vitis-AI 2. Install Docker, and add the user to the docker group. Link the user to docker installation instructions from the following docker's @@ -139,11 +139,12 @@ Hardware setup and docker build them at once. To do so: - Run the following commands: - .. code:: bash + + .. code:: bash - cd Vitis-AI/alveo/packages - sudo su - ./install.sh + cd Vitis-AI/alveo/packages + sudo su + ./install.sh - Power cycle the system. @@ -339,14 +340,14 @@ interface between TVM and Vitis-AI tools. 1. First install the PyXIR h5py and pydot dependencies: - .. code:: bash: + .. code:: bash apt-get install libhdf5-dev pip3 install pydot h5py 2. Install PyXIR - .. code:: bash: + .. code:: bash git clone --recursive https://github.com/Xilinx/pyxir.git cd pyxir @@ -354,7 +355,7 @@ interface between TVM and Vitis-AI tools. 3. Build TVM with Vitis-AI - .. code:: bash: + .. code:: bash git clone --recursive https://github.com/apache/incubator-tvm cd incubator-tvm @@ -367,14 +368,14 @@ interface between TVM and Vitis-AI tools. 4. Install TVM - .. code:: bash: + .. code:: bash cd incubator-tvm/python pip3 install -e . --user 5. Check whether the setup was successful in the Python shell: - .. code:: bash: + .. code:: bash python3 -c 'import pyxir; import tvm' From 128115ebcc7238e2cb2ad796b6d62f3eb5ae5c56 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 20 Oct 2020 08:53:18 -0600 Subject: [PATCH 18/22] Add infertype to fix Vitis-AI annotation test --- tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py index 4310ab6d9ced..4d5d5dc92c41 100644 --- a/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py +++ b/tests/python/contrib/test_vitis_ai/test_vitis_ai_codegen.py @@ -293,6 +293,7 @@ def expected(): gv0 = relay.GlobalVar("vitis_ai_0") mod = tvm.IRModule() mod[gv0] = func0 + mod = relay.transform.InferType()(mod) # main function data = relay.var("data", relay.TensorType((1, 3, 224, 224), "float32")) @@ -303,6 +304,7 @@ def expected(): bn_mvar = relay.var("bn_var", relay.TensorType((16,), "float32")) call0 = gv0(data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar) mod["main"] = relay.Function([data, weight, bn_gamma, bn_beta, bn_mmean, bn_mvar], call0) + mod = relay.transform.InferType()(mod) return mod partitioned_dpuczdx8g_zcu104 = partition("DPUCZDX8G-zcu104") From 07c84d91a4912a3886a575f8061fa6fbcc4afc0e Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Mon, 2 Nov 2020 10:08:27 -0700 Subject: [PATCH 19/22] Renaming util to utils --- python/tvm/contrib/target/vitis_ai.py | 4 ++-- tests/python/contrib/test_vitis_ai/infrastructure.py | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/python/tvm/contrib/target/vitis_ai.py b/python/tvm/contrib/target/vitis_ai.py index 2fa95b31b7ce..d4931d9e3f48 100644 --- a/python/tvm/contrib/target/vitis_ai.py +++ b/python/tvm/contrib/target/vitis_ai.py @@ -82,12 +82,12 @@ def vitis_ai_compiler(ref): vai_build_dir = ( str(pass_context.config["relay.ext.vitis_ai.options.build_dir"]) if "relay.ext.vitis_ai.options.build_dir" in pass_context.config - else tvm.contrib.util.tempdir().relpath("") + else tvm.contrib.utils.tempdir().relpath("") ) vai_work_dir = ( str(pass_context.config["relay.ext.vitis_ai.options.work_dir"]) if "relay.ext.vitis_ai.options.work_dir" in pass_context.config - else tvm.contrib.util.tempdir().relpath("") + else tvm.contrib.utils.tempdir().relpath("") ) # (Optional configs) Export and load PyXIR runtime module to file if provided. This is used to diff --git a/tests/python/contrib/test_vitis_ai/infrastructure.py b/tests/python/contrib/test_vitis_ai/infrastructure.py index ff7330602776..df7836a37647 100644 --- a/tests/python/contrib/test_vitis_ai/infrastructure.py +++ b/tests/python/contrib/test_vitis_ai/infrastructure.py @@ -35,7 +35,7 @@ from tvm.relay.build_module import bind_params_by_name from tvm.contrib.target import vitis_ai from tvm.contrib import graph_runtime -from tvm.contrib import util +from tvm.contrib import utils def get_cpu_op_count(mod): @@ -107,7 +107,7 @@ def build_module( def update_lib(lib, cross_compile=None): - tmp_path = util.tempdir() + tmp_path = utils.tempdir() lib_name = "lib.so" lib_path = tmp_path.relpath(lib_name) if cross_compile: From 753322e83080ca57dcc57cfb6fc8d603e0cffb1e Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 3 Nov 2020 03:17:51 -0700 Subject: [PATCH 20/22] Add Vitis-AI flag to config.cmake file --- tests/scripts/task_config_build_cpu.sh | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/scripts/task_config_build_cpu.sh b/tests/scripts/task_config_build_cpu.sh index 521ab9b8ccdc..4fd8bba8b29b 100755 --- a/tests/scripts/task_config_build_cpu.sh +++ b/tests/scripts/task_config_build_cpu.sh @@ -45,3 +45,4 @@ echo set\(USE_TENSORFLOW_PATH \"/tensorflow\"\) >> config.cmake echo set\(USE_FLATBUFFERS_PATH \"/flatbuffers\"\) >> config.cmake echo set\(USE_ETHOSN /opt/arm/ethosn-driver\) >> config.cmake echo set\(USE_ETHOSN_HW OFF\) >> config.cmake +echo set\(USE_VITIS_AI ON\) >> config.cmake From bdf2bbd1b1ce2c153049e0e00dd1d79e9b4b2a87 Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 3 Nov 2020 09:59:36 -0700 Subject: [PATCH 21/22] Move vitis-ai config options to compiler sources instead of runtime sources --- cmake/modules/contrib/VitisAI.cmake | 3 ++ .../contrib/vitis_ai/config_vitis_ai.cc | 47 +++++++++++++++++++ .../contrib/vitis_ai/vitis_ai_runtime.cc | 12 ----- 3 files changed, 50 insertions(+), 12 deletions(-) create mode 100644 src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc diff --git a/cmake/modules/contrib/VitisAI.cmake b/cmake/modules/contrib/VitisAI.cmake index 896a3c8d1513..083bd6d7adc8 100644 --- a/cmake/modules/contrib/VitisAI.cmake +++ b/cmake/modules/contrib/VitisAI.cmake @@ -38,6 +38,9 @@ if(USE_VITIS_AI) message(STATUS "Build with contrib.vitisai") include_directories(${PYXIR_INCLUDE_DIR}) file(GLOB VAI_CONTRIB_SRC src/runtime/contrib/vitis_ai/*.cc) + file(GLOB COMPILER_VITIS_AI_SRCS + CONFIGURE_DEPENDS src/relay/backend/contrib/vitis_ai/*) + list(APPEND COMPILER_SRCS ${COMPILER_VITIS_AI_SRCS}) link_directories(${PYXIR_LIB_DIR}) list(APPEND TVM_RUNTIME_LINKER_LIBS "pyxir") list(APPEND RUNTIME_SRCS ${VAI_CONTRIB_SRC}) diff --git a/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc b/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc new file mode 100644 index 000000000000..0598ff4eedcc --- /dev/null +++ b/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc @@ -0,0 +1,47 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +/*! + * \file src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc + * \brief Register Vitis-AI codegen options. Main codegen is implemented in python. + */ + +#include + + +namespace tvm { +namespace relay { +namespace contrib { +namespace vitis_ai { + +/*! \brief The target Vitis-AI accelerator device */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.target", String); +/*! \brief (Optional config) The build directory to be used by Vitis-AI */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.build_dir", String); +/*! \brief (Optional config) The work directory to be used by Vitis-AI */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.work_dir", String); +/*! \brief (Optional config) Export PyXIR runtime module to disk during serialization if provided */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.export_runtime_module", String); +/*! \brief (Optional config) Load PyXIR runtime module from disk */ +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.load_runtime_module", String); + +} // namespace vitis_ai +} // namespace contrib +} // namespace relay +} // namespace tvm diff --git a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc index 35075aba4a93..37dc767d31af 100755 --- a/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc +++ b/src/runtime/contrib/vitis_ai/vitis_ai_runtime.cc @@ -23,7 +23,6 @@ #include "vitis_ai_runtime.h" -#include #include #include @@ -36,17 +35,6 @@ using namespace pyxir::runtime; namespace tvm { namespace runtime { -/*! \brief The target Vitis-AI accelerator device */ -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.target", String); -/*! \brief (Optional config) The build directory to be used by Vitis-AI */ -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.build_dir", String); -/*! \brief (Optional config) The work directory to be used by Vitis-AI */ -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.work_dir", String); -/*! \brief (Optional config) Export PyXIR runtime module to disk during serialization if provided */ -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.export_runtime_module", String); -/*! \brief (Optional config) Load PyXIR runtime module from disk */ -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.vitis_ai.options.load_runtime_module", String); - VitisAIRuntime::VitisAIRuntime(const std::string& symbol_name, const Array const_names, const std::string& serialized_rt_mod, const std::string& export_rt_mod_path) From 4cf883819fa5f4fbb4750f5c8d6e8f3d10de770d Mon Sep 17 00:00:00 2001 From: Anil Martha Date: Tue, 3 Nov 2020 10:35:11 -0700 Subject: [PATCH 22/22] Fix clang-format errors --- src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc b/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc index 0598ff4eedcc..f74b5306c5f4 100644 --- a/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc +++ b/src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc @@ -19,12 +19,11 @@ /*! * \file src/relay/backend/contrib/vitis_ai/config_vitis_ai.cc - * \brief Register Vitis-AI codegen options. Main codegen is implemented in python. + * \brief Register Vitis-AI codegen options. Main codegen is implemented in python. */ #include - namespace tvm { namespace relay { namespace contrib {