comaniac commented on a change in pull request #6343: URL: https://github.com/apache/incubator-tvm/pull/6343#discussion_r479578714
########## File path: 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 <tvm/runtime/registry.h> +#include <tvm/ir/transform.h> + +#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); Review comment: ditto. ########## File path: 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.") Review comment: indent ########## File path: 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 <tvm/runtime/registry.h> +#include <tvm/ir/transform.h> + +#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<pyxir::graph::XGraph> 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<std::string>{target}, ""); Review comment: Out of curiousity, what does this partition do? ########## File path: 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 Review comment: It might be a problem to directly import Vitis-AI dependencies, because most developers may not have this environment. Please make all Vitis-AI specific imports lazy, and safely skip all unit tests if the required environment is not ready. ########## File path: 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() Review comment: It seems not good to use the current directory as the model path. If it is necessary for XGraph to serialize the artifacts out for further processing, we should use `/tmp` or let user assign. ########## File path: docs/deploy/vitis_ai.rst ########## @@ -0,0 +1,617 @@ +Vitis-AI Integration +==================== + +`Vitis-AI <https://github.com/Xilinx/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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +respectively +`Alveo <https://www.xilinx.com/products/boards-and-kits/alveo.html>`__ +(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: + :: Review comment: Better to use `.. code:: bash` (or other languages) to be more specific. ########## File path: 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(): Review comment: It seems to me that you're expecting two subgreaphs (functions)? But your previous comment says the flow currently only supports one subgraph in a model. ########## File path: 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 """ Review comment: Put two test files under `tests/python/contrib/test_vitis_ai/`. ########## File path: 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): Review comment: 4 of test cases rely on this function, but this function does not check anything. It only makes sure the single op can go through the flow without crashing. I believe there are more to cover. ########## File path: docs/deploy/vitis_ai.rst ########## @@ -0,0 +1,617 @@ +Vitis-AI Integration +==================== + +`Vitis-AI <https://github.com/Xilinx/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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +respectively +`Alveo <https://www.xilinx.com/products/boards-and-kits/alveo.html>`__ +(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. Review comment: Didn't understand this statement and its functionality. ########## File path: docs/deploy/vitis_ai.rst ########## @@ -0,0 +1,617 @@ +Vitis-AI Integration Review comment: Missing license head. ########## File path: docs/deploy/vitis_ai.rst ########## @@ -0,0 +1,617 @@ +Vitis-AI Integration +==================== + +`Vitis-AI <https://github.com/Xilinx/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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +respectively +`Alveo <https://www.xilinx.com/products/boards-and-kits/alveo.html>`__ +(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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +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 <http://www.pynq.io/>`__ platform but Petalinux based flows are also supported. Review comment: Use `.. note::` to create a note block. ########## File path: docs/deploy/vitis_ai.rst ########## @@ -0,0 +1,617 @@ +Vitis-AI Integration +==================== + +`Vitis-AI <https://github.com/Xilinx/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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +respectively +`Alveo <https://www.xilinx.com/products/boards-and-kits/alveo.html>`__ +(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 <https://www.xilinx.com/products/silicon-devices/soc/zynq-ultrascale-mpsoc.html>`__ +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 <http://www.pynq.io/>`__ 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 <https://pynq.readthedocs.io/en/latest/getting_started.html>`__ +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 <https://github.com/Xilinx/DPU-PYNQ>`__ +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. Review comment: Ditto. ########## File path: 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""" Review comment: To me this is the BYOC codegen and should be under `src/relay/backend/contrib`. ########## File path: 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 Review comment: Put this check in `test_extern_vai_resnet18` directly so that we won't waste time running that test on Windows. ########## File path: 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): Review comment: Same question. ########## File path: 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 Review comment: ditto. ########## File path: 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 + + [email protected]_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) Review comment: It seems to me that it leverages its own annotation (and partition) flow. After that, the partition result is stored in `relay_ids_flatten`. The use of `VitisAIAnnotationPass` and the rest `MergeCompilerRegion`, `PartitionGraph` passes are more like wrappers to repeat the same logic in Relay. My concern to this flow is that it makes the Vitis-AI codegen less general compared other BYOC backends, but according to the above comment about replacing the current annotator with op-based one, will this problem be addressed as well? ########## File path: 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 <tvm/runtime/registry.h> +#include <tvm/ir/transform.h> + +#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<pyxir::graph::XGraph> 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); +} Review comment: Same opinion. You could write all of them together to a single binary stream in `SaveToBinary`. As long as you have a consistent format to decode it, you can avoid multiple files/paths issues. ---------------------------------------------------------------- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. For queries about this service, please contact Infrastructure at: [email protected]
