From afa87f9c36c39bc087465f293a7c08ca23481b8b Mon Sep 17 00:00:00 2001 From: liuxinwei Date: Sun, 29 Sep 2024 09:39:29 +0800 Subject: [PATCH] =?UTF-8?q?=E6=9B=B4=E6=96=B0=E6=96=87=E6=A1=A3?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- doc/read/cpp-runtime/c-runtime-api.ipynb | 163 +++++++++++++++ doc/read/cpp-runtime/index.md | 1 + doc/read/relay-frontend/common/set_env.py | 6 - doc/read/relay-frontend/common/tag-span.ipynb | 93 ++++++--- .../common/utils/assert_diagnostic.py | 66 ------ .../common/utils/external_codegen.py | 125 ------------ .../relay-frontend/common/utils/ref_funcs.py | 48 ----- .../relay-frontend/common/utils/tag_span.py | 107 ---------- doc/read/source-map/index.md | 2 +- doc/read/source-map/span.ipynb | 2 +- doc/topics/frontend/caffe/network.ipynb | 4 +- tests/a.toml | 2 - tests/test.ipynb | 191 +++++++++--------- 13 files changed, 333 insertions(+), 477 deletions(-) create mode 100644 doc/read/cpp-runtime/c-runtime-api.ipynb delete mode 100755 doc/read/relay-frontend/common/set_env.py delete mode 100644 doc/read/relay-frontend/common/utils/assert_diagnostic.py delete mode 100644 doc/read/relay-frontend/common/utils/external_codegen.py delete mode 100644 doc/read/relay-frontend/common/utils/ref_funcs.py delete mode 100644 doc/read/relay-frontend/common/utils/tag_span.py delete mode 100644 tests/a.toml diff --git a/doc/read/cpp-runtime/c-runtime-api.ipynb b/doc/read/cpp-runtime/c-runtime-api.ipynb new file mode 100644 index 00000000..07a57e1a --- /dev/null +++ b/doc/read/cpp-runtime/c-runtime-api.ipynb @@ -0,0 +1,163 @@ +{ + "cells": [ + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "# TVM Runtime api\n", + "\n", + "源码见：`tvm/src/runtime/c_runtime_api.h`" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "`tvm/runtime/c_runtime_api.h` 包含了 TVM 运行时库的一些函数和宏定义。\n", + "\n", + "TVM 项目的理念是定制编译阶段以生成可以透明地被其他项目使用的代码。因此，这个头文件中包含了一些最小的运行时代码粘合以及有限的内存管理代码，以便进行快速测试。\n", + "\n", + "运行时 API 独立于 TVM 编译堆栈，可以通过链接 `libtvm_runtime` 来使用。\n", + "\n", + "常见的流程如下：\n", + "\n", + "- 使用 `TVMFuncListGlobalNames` 获取全局函数名称\n", + "- 使用 `TVMFuncCall` 调用这些函数\n", + "\n", + "API函数的可能返回值有：\n", + "- `0`：成功\n", + "- `-1`：错误可以通过 `TVMGetLastError` 检索\n", + "- `-2`：前端发生错误并记录在前端" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "```cpp\n", + "// Macros to do weak linking\n", + "#ifdef _MSC_VER\n", + "#define TVM_WEAK __declspec(selectany)\n", + "#else\n", + "#define TVM_WEAK __attribute__((weak))\n", + "#endif\n", + "```\n", + "\n", + "这段代码是 C++ 头文件中的宏定义，用于控制 TVM 库在不同编译器环境下进行弱链接的方式。\n", + "\n", + "首先，`#ifdef _MSC_VER` 检查是否使用的是 Microsoft Visual C++ 编译器（即 Windows 平台）。如果是，那么使用 `__declspec(selectany)` 来声明弱符号。`__declspec(selectany)` 是 Visual C++ 特有的关键字，用于指示编译器在多个目标文件中选择一个符号的定义，而不是报错。\n", + "\n", + "如果不是 Visual C++ 编译器，那么使用 `__attribute__((weak))` 来声明弱符号。这是 GCC 和 Clang 编译器的一种方式，用于设置符号为弱符号。弱符号允许在链接时被其他同名强符号覆盖，如果没有找到强符号，则使用弱符号作为备选。\n", + "\n", + "通过这样的宏定义，TVM库可以在不同编译器环境下灵活地进行弱链接，以支持跨平台的编译和链接。" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "```cpp\n", + "#ifdef __EMSCRIPTEN__\n", + "#include \n", + "#define TVM_DLL EMSCRIPTEN_KEEPALIVE\n", + "#endif\n", + "```\n", + "\n", + "这段代码是 C++ 头文件中的宏定义，用于在 Emscripten 编译器环境下控制 TVM 库的导出方式。\n", + "\n", + "首先，`#ifdef __EMSCRIPTEN__` 检查是否使用的是 Emscripten 编译器。Emscripten 是一个将 C/C++ 代码编译成 JavaScript 的工具链，通常用于 Web 开发。\n", + "\n", + "如果当前是 Emscripten 编译器环境，那么包含 `` 头文件，这是 Emscripten 提供的一组 API 函数和宏定义。\n", + "\n", + "接下来，使用 `#define TVM_DLL EMSCRIPTEN_KEEPALIVE` 来声明 TVM 库中的符号为保持活跃状态。`EMSCRIPTEN_KEEPALIVE` 是 Emscripten 提供的一个宏，用于告诉编译器不要优化掉这个符号，确保它在运行时可用。这对于需要在 JavaScript 环境中调用的函数或变量非常有用。\n", + "\n", + "通过这样的宏定义，TVM 库可以在 Emscripten 编译器环境下正确地导出符号，以便在 Web 环境中使用。" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "```cpp\n", + "// helper macro to suppress unused warning\n", + "#if defined(__GNUC__)\n", + "#define TVM_ATTRIBUTE_UNUSED __attribute__((unused))\n", + "#else\n", + "#define TVM_ATTRIBUTE_UNUSED\n", + "#endif\n", + "```\n", + "\n", + "这段代码是 C 语言的宏定义，用于在编译时抑制未使用变量或函数的警告。它使用了条件编译来检查是否使用的是 GNU 编译器（如 GCC）。如果是，则定义名为 `TVM_ATTRIBUTE_UNUSED` 的宏，该宏带有 `__attribute__((unused))` 属性，告诉编译器这个变量或函数可能不会被使用，从而避免产生未使用的警告。如果不是 GNU 编译器，则简单地定义 `TVM_ATTRIBUTE_UNUSED` 为空。" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "```cpp\n", + "#ifndef TVM_DLL\n", + "#ifdef _WIN32\n", + "#ifdef TVM_EXPORTS\n", + "#define TVM_DLL __declspec(dllexport)\n", + "#else\n", + "#define TVM_DLL __declspec(dllimport)\n", + "#endif\n", + "#else\n", + "#define TVM_DLL __attribute__((visibility(\"default\")))\n", + "#endif\n", + "#endif\n", + "```" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "这段代码是 C++ 头文件中的宏定义，用于控制 TVM 库在不同平台和编译环境下的动态链接库（DLL）导出方式。\n", + "\n", + "首先，`#ifndef TVM_DLL` 检查是否已经定义了 `TVM_DLL` 这个宏。如果没有定义，那么继续执行下面的代码块。\n", + "\n", + "接下来，`#ifdef _WIN32` 检查当前是否是 Windows 平台。如果是 Windows 平台，再进一步检查 `TVM_EXPORTS` 宏是否被定义。\n", + "\n", + "- 如果 `TVM_EXPORTS` 被定义，说明我们正在构建 TVM 库本身，需要将函数或变量导出为 DLL 接口，因此使用 `__declspec(dllexport)` 来声明。\n", + "- 如果 `TVM_EXPORTS` 没有被定义，说明我们正在使用 TVM 库，需要导入库中的函数或变量，因此使用 `__declspec(dllimport)` 来声明。\n", + "\n", + "如果不是 Windows 平台，那么使用 `__attribute__((visibility(\"default\")))` 来声明，这是 GCC 编译器的一种方式，用于设置符号的可见性。在这种情况下，默认情况下所有符号都是可见的。\n", + "\n", + "最后，如果 `TVM_DLL` 已经被定义，那么这个宏定义块不会再次执行，避免了重复定义的问题。" + ] + }, + { + "cell_type": "markdown", + "metadata": {}, + "source": [ + "```cpp\n", + "#ifdef __cplusplus\n", + "extern \"C\" {\n", + "#endif\n", + "// 函数原型等内容\n", + "#ifdef __cplusplus\n", + "} // TVM_EXTERN_C\n", + "#endif\n", + "```\n", + "\n", + "这段代码是 C++ 头文件中的宏定义，用于控制 TVM 库在不同编译器环境下进行 C 语言链接的方式。\n", + "\n", + "首先，`#ifdef __cplusplus` 检查是否使用的是 C++ 编译器。如果是，那么使用 `extern \"C\"` 来声明一个 C 语言链接区块。这是为了确保在这个区块中的函数和变量遵循 C 语言的命名规则和调用约定，而不是 C++ 的规则。\n", + "\n", + "接下来，可以在这个区块中添加函数原型、类型定义等 C 语言相关的代码。这些代码将被编译器视为 C 语言代码，并按照 C 语言的方式进行编译和链接。\n", + "\n", + "最后，再次使用 `#ifdef __cplusplus` 来结束 C 语言链接区块，并在结束时加上 `} // TVM_EXTERN_C` 注释，以便于阅读和维护代码。\n", + "\n", + "通过这样的宏定义，TVM 库可以在不同编译器环境下灵活地进行 C 语言链接，以支持跨平台的编译和链接。" + ] + } + ], + "metadata": { + "language_info": { + "name": "python" + } + }, + "nbformat": 4, + "nbformat_minor": 2 +} diff --git a/doc/read/cpp-runtime/index.md b/doc/read/cpp-runtime/index.md index a86fbe77..c52429ad 100644 --- a/doc/read/cpp-runtime/index.md +++ b/doc/read/cpp-runtime/index.md @@ -1,6 +1,7 @@ # 运行时(C++) ```{toctree} +c-runtime-api runtime-base Object ObjectPtr diff --git a/doc/read/relay-frontend/common/set_env.py b/doc/read/relay-frontend/common/set_env.py deleted file mode 100755 index 48ffaa93..00000000 --- a/doc/read/relay-frontend/common/set_env.py +++ /dev/null @@ -1,6 +0,0 @@ -# from pathlib import Path -from tvm_book.config.env import set_tvm -# TVM_ROOT = Path(__file__).resolve().parents[5] -# print(TVM_ROOT) -TVM_ROOT = "/media/pc/data/lxw/ai/tvm/" -set_tvm(TVM_ROOT) \ No newline at end of file diff --git a/doc/read/relay-frontend/common/tag-span.ipynb b/doc/read/relay-frontend/common/tag-span.ipynb index a7aae0af..f767bd22 100644 --- a/doc/read/relay-frontend/common/tag-span.ipynb +++ b/doc/read/relay-frontend/common/tag-span.ipynb @@ -11,9 +11,19 @@ "cell_type": "code", "execution_count": 1, "metadata": {}, - "outputs": [], + "outputs": [ + { + "name": "stdout", + "output_type": "stream", + "text": [ + "/media/pc/data/lxw/ai/tvm-book/doc/read\n" + ] + } + ], "source": [ - "import set_env" + "%cd ../..\n", + "import set_env\n", + "from tools.tag_span import _set_span, _create_span, _verify_structural_equal_with_span" ] }, { @@ -25,8 +35,7 @@ "import numpy as np\n", "\n", "from tvm import relay, testing\n", - "from tvm.relay.frontend.common import StrAttrsDict, set_span\n", - "from utils.tag_span import _set_span, _create_span, _verify_structural_equal_with_span" + "from tvm.relay.frontend.common import StrAttrsDict, set_span" ] }, { @@ -57,7 +66,7 @@ }, { "cell_type": "code", - "execution_count": 7, + "execution_count": 4, "metadata": {}, "outputs": [], "source": [ @@ -72,7 +81,7 @@ }, { "cell_type": "code", - "execution_count": 8, + "execution_count": 5, "metadata": {}, "outputs": [], "source": [ @@ -85,7 +94,7 @@ }, { "cell_type": "code", - "execution_count": 9, + "execution_count": 6, "metadata": {}, "outputs": [ { @@ -117,17 +126,9 @@ }, { "cell_type": "code", - "execution_count": 10, + "execution_count": 7, "metadata": {}, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "(Constant([[[1]]]), Constant([[[0]]]))\n" - ] - } - ], + "outputs": [], "source": [ "def _res():\n", " a = relay.const(np.ones([1, 1, 1]), dtype=\"int64\", span=_create_span(\"a\"))\n", @@ -147,7 +148,7 @@ }, { "cell_type": "code", - "execution_count": 12, + "execution_count": 8, "metadata": {}, "outputs": [ { @@ -172,7 +173,7 @@ }, { "cell_type": "code", - "execution_count": 7, + "execution_count": 9, "metadata": {}, "outputs": [], "source": [ @@ -207,7 +208,7 @@ }, { "cell_type": "code", - "execution_count": 8, + "execution_count": 10, "metadata": {}, "outputs": [], "source": [ @@ -225,7 +226,7 @@ }, { "cell_type": "code", - "execution_count": 9, + "execution_count": 11, "metadata": {}, "outputs": [], "source": [ @@ -245,7 +246,7 @@ }, { "cell_type": "code", - "execution_count": 13, + "execution_count": 12, "metadata": {}, "outputs": [], "source": [ @@ -270,7 +271,7 @@ }, { "cell_type": "code", - "execution_count": 14, + "execution_count": 13, "metadata": {}, "outputs": [ { @@ -297,7 +298,7 @@ }, { "cell_type": "code", - "execution_count": 11, + "execution_count": 14, "metadata": {}, "outputs": [], "source": [ @@ -325,7 +326,7 @@ }, { "cell_type": "code", - "execution_count": 12, + "execution_count": 15, "metadata": {}, "outputs": [], "source": [ @@ -355,7 +356,7 @@ }, { "cell_type": "code", - "execution_count": 13, + "execution_count": 16, "metadata": {}, "outputs": [], "source": [ @@ -391,7 +392,7 @@ }, { "cell_type": "code", - "execution_count": 14, + "execution_count": 17, "metadata": {}, "outputs": [], "source": [ @@ -427,7 +428,7 @@ }, { "cell_type": "code", - "execution_count": 15, + "execution_count": 18, "metadata": {}, "outputs": [], "source": [ @@ -450,6 +451,40 @@ "_verify_structural_equal_with_span(_res(), _golden())" ] }, + { + "cell_type": "code", + "execution_count": 23, + "metadata": {}, + "outputs": [], + "source": [ + "x = relay.var(\"x\", shape=(1, 64, 56, 56), span=_create_span(\"x_var\"))\n", + "w = relay.const(np.ones([64, 64, 3, 3]), dtype=\"int64\", span=_create_span(\"func\"))\n", + "y = _set_span(\n", + " relay.nn.conv2d(x, w, channels=64, kernel_size=(3, 3), padding=(1, 1)), \"func\"\n", + ")\n", + "f = relay.Function([x], y, span=_create_span(\"func\"))" + ] + }, + { + "cell_type": "code", + "execution_count": 24, + "metadata": {}, + "outputs": [ + { + "name": "stdout", + "output_type": "stream", + "text": [ + "fn (%x: Tensor[(1, 64, 56, 56), float32] /* span=x_var:0:0 */) {\n", + " nn.conv2d(%x, meta[relay.Constant][0] /* span=func:0:0 */, padding=[1, 1, 1, 1], channels=64, kernel_size=[3, 3]) /* span=func:0:0 */\n", + "} /* span=func:0:0 */\n", + "\n" + ] + } + ], + "source": [ + "print(f)" + ] + }, { "cell_type": "code", "execution_count": null, @@ -474,7 +509,7 @@ "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", - "version": "3.12.2" + "version": "3.12.4" } }, "nbformat": 4, diff --git a/doc/read/relay-frontend/common/utils/assert_diagnostic.py b/doc/read/relay-frontend/common/utils/assert_diagnostic.py deleted file mode 100644 index 5fcd1c20..00000000 --- a/doc/read/relay-frontend/common/utils/assert_diagnostic.py +++ /dev/null @@ -1,66 +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. -import tvm -from tvm import IRModule, get_global_func, register_func, relay -from tvm.error import DiagnosticError -from tvm.ir.diagnostics import get_renderer, override_renderer -from tvm.relay import SpanCheck -from tvm.relay.transform import AnnotateSpans -from tvm.runtime import Object - -DEFAULT_RENDERER = get_renderer() - -__TESTING__ = None - - -def testing_renderer(diag_ctx): - global __TESTING__ - if __TESTING__ and __TESTING__.mirror: - DEFAULT_RENDERER.render(diag_ctx) - - if __TESTING__: - __TESTING__._render(diag_ctx) - - -class DiagnosticTesting: - def __init__(self, mirror=False): - self.mirror = mirror - self.messages = [] - - def __enter__(self): - global __TESTING__ - __TESTING__ = self - override_renderer(testing_renderer) - return self - - def __exit__(self, type, value, traceback): - global __TESTING__ - __TESTING__ = None - override_renderer(None) - if type is DiagnosticError and self.matches: - return True - - def assert_message(self, in_message): - self.messages.append(in_message) - - def _render(self, diag_ctx): - self.matches = False - for diagnostic in diag_ctx.diagnostics: - message = diagnostic.message - for partial_msg in self.messages: - if partial_msg in message: - self.matches = True diff --git a/doc/read/relay-frontend/common/utils/external_codegen.py b/doc/read/relay-frontend/common/utils/external_codegen.py deleted file mode 100644 index bb06d3bb..00000000 --- a/doc/read/relay-frontend/common/utils/external_codegen.py +++ /dev/null @@ -1,125 +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. -"""Utilities for testing external code generation""" - -import os -import sys - -import pytest - -import tvm -from tvm import relay, runtime, testing -from tvm.contrib import utils - - -skip_windows = pytest.mark.skipif(sys.platform == "win32", reason="Skip test on Windows for now") -skip_micro = pytest.mark.skipif( - tvm.support.libinfo().get("USE_MICRO", "OFF") != "ON", - reason="MicroTVM support not enabled. Set USE_MICRO=ON in config.cmake to enable.", -) - - -def parametrize_external_codegen_checks(test): - """Parametrize over the various check_result functions which are available""" - return pytest.mark.parametrize( - "check_result", - [ - pytest.param(check_aot_executor_result, marks=[skip_windows, skip_micro]), - pytest.param(check_graph_executor_result, marks=[skip_windows]), - pytest.param(check_vm_result, marks=[skip_windows]), - ], - )(test) - - -def parametrize_external_json_codegen_checks(test): - """Parametrize over the various check_result functions which are available for JSON""" - return pytest.mark.parametrize( - "check_result", - [ - pytest.param(check_graph_executor_result, marks=[skip_windows]), - pytest.param(check_vm_result, marks=[skip_windows]), - ], - )(test) - - -def update_lib(lib): - test_dir = os.path.dirname(os.path.realpath(os.path.expanduser(__file__))) - source_dir = os.path.join(test_dir, "..", "..", "..", "..") - contrib_path = os.path.join(source_dir, "src", "runtime", "contrib") - - kwargs = {} - kwargs["options"] = ["-O2", "-std=c++17", "-I" + contrib_path] - tmp_path = utils.tempdir() - lib_name = "lib.so" - lib_path = tmp_path.relpath(lib_name) - lib.export_library(lib_path, fcompile=False, **kwargs) - lib = tvm.runtime.load_module(lib_path) - - return lib - - -def check_vm_result(mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", device=tvm.cpu()): - with tvm.transform.PassContext(opt_level=3, disabled_pass=["AlterOpLayout"]): - exe = relay.vm.compile(mod, target=target) - code, lib = exe.save() - lib = update_lib(lib) - exe = runtime.vm.Executable.load_exec(code, lib) - vm = runtime.vm.VirtualMachine(exe, device) - out = vm.run(**map_inputs) - tvm.testing.assert_allclose(out.numpy(), result, rtol=tol, atol=tol) - - -def check_graph_executor_result( - mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", device=tvm.cpu() -): - with tvm.transform.PassContext(opt_level=3, disabled_pass=["AlterOpLayout"]): - executor_factory = relay.build(mod, target=target) - lib = update_lib(executor_factory.lib) - rt_mod = tvm.contrib.graph_executor.create(executor_factory.graph_json, lib, device) - - for name, data in map_inputs.items(): - rt_mod.set_input(name, data) - rt_mod.run() - out = tvm.nd.empty(out_shape, device=device) - out = rt_mod.get_output(0, out) - - tvm.testing.assert_allclose(out.numpy(), result, rtol=tol, atol=tol) - - -def check_aot_executor_result( - mod, map_inputs, out_shape, result, tol=1e-5, target="llvm", device=tvm.cpu() -): - # Late import to avoid breaking test with USE_MICRO=OFF. - from tvm.testing.aot import AOTTestModel, compile_and_run - from tvm.micro.testing.aot_test_utils import AOT_DEFAULT_RUNNER - - interface_api = "packed" - use_unpacked_api = False - test_runner = AOT_DEFAULT_RUNNER - compile_and_run( - AOTTestModel(module=mod, inputs=map_inputs, outputs={"output": result}), - test_runner, - interface_api, - use_unpacked_api, - ) - - -def set_external_func_attr(func, compiler, ext_symbol): - func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) - func = func.with_attr("Compiler", compiler) - func = func.with_attr("global_symbol", ext_symbol) - return func diff --git a/doc/read/relay-frontend/common/utils/ref_funcs.py b/doc/read/relay-frontend/common/utils/ref_funcs.py deleted file mode 100644 index 924805b2..00000000 --- a/doc/read/relay-frontend/common/utils/ref_funcs.py +++ /dev/null @@ -1,48 +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. -import numpy as np - - -def gather_nd(data_np, indices_np, batch_dims=0): - """gather_nd implemented using numpy""" - data_shape = data_np.shape - indices_shape = indices_np.shape - - def gather_nd_batch_dims_1_ref(data, indices): - res = [] - for i, row in enumerate(data): - indices_tuple = tuple(indices[:, i]) # the indices for the i-th batch - res.append(row[indices_tuple]) - # stack on the batch dim - return np.stack(res, 0) - - if batch_dims > 1: - data_np_reshape = np.reshape(data_np, (-1,) + data_shape[batch_dims:]) - indices_np_reshape = np.reshape( - indices_np, (indices_shape[0], -1) + indices_shape[(batch_dims + 1) :] - ) - - ref_res = gather_nd_batch_dims_1_ref(data_np_reshape, indices_np_reshape) - - out_shape = indices_shape[1 : (batch_dims + 1)] + ref_res.shape[1:] - ref_res = np.reshape(ref_res, out_shape) - elif batch_dims == 1: - ref_res = gather_nd_batch_dims_1_ref(data_np, indices_np) - else: - ref_res = data_np[tuple(indices_np)] - - return ref_res diff --git a/doc/read/relay-frontend/common/utils/tag_span.py b/doc/read/relay-frontend/common/utils/tag_span.py deleted file mode 100644 index ac79344b..00000000 --- a/doc/read/relay-frontend/common/utils/tag_span.py +++ /dev/null @@ -1,107 +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. -import tvm -from tvm import relay, tir -from tvm.relay import expr as _expr -# from tvm.relay.expr_functor import ExprVisitor - -def _create_span(src): - if isinstance(src, list): - tmp_list = [] - for s in src: - if isinstance(s, str): - tmp_list.append(_create_span(s)) - elif isinstance(s, relay.Span): - tmp_list.append(s) - elif isinstance(s, relay.SequentialSpan): - tmp_list.extend(s.spans) - elif s is None: - tmp_list.append(s) - else: - assert False, f"unsupported type {type(s)}" - return relay.SequentialSpan(tmp_list) - return relay.Span(relay.SourceName(src), 0, 0, 0, 0) - - -def _set_span(expr, src): - if isinstance(expr, _expr.Call): - return _expr.CallWithFields( - expr, expr.op, expr.args, expr.attrs, expr.type_args, None, _create_span(src) - ) - elif isinstance(expr, _expr.Var): - return _expr.VarWithFields(expr, expr.vid, expr.type_annotation, None, _create_span(src)) - elif isinstance(expr, _expr.TupleGetItem): - return _expr.TupleGetItemWithFields( - expr, expr.tuple_value, expr.index, None, _create_span(src) - ) - elif isinstance(expr, _expr.Constant): - return _expr.ConstantWithFields(expr, expr.data, None, _create_span(src)) - elif isinstance(expr, _expr.Tuple): - return _expr.TupleWithFields(expr, expr.fields, None, _create_span(src)) - elif isinstance(expr, _expr.TupleWrapper): - return _expr.TupleWrapper(_set_span(expr.tuple_value, src), expr.size) - - assert False, f"unsupported type {type(expr)}" - - -def _collect_spans(objref): - class Collector: - def __init__(self): - self._spans = [] - - def collect(self, objref): - if hasattr(objref, "span"): - self._spans.append(objref.span) - - @property - def get_spans(self): - return self._spans - - pov = None - if isinstance(objref, relay.Expr): - pov = relay.analysis.post_order_visit - elif isinstance(objref, (tir.Stmt, tir.expr.PrimExprWithOp)): - pov = tir.stmt_functor.post_order_visit - else: - assert False, f"unsupported type {type(objref)}" - - c = Collector() - pov(objref, c.collect) - return c.get_spans - - -def _verify_span(lhs, rhs): - lhs_spans, rhs_spans = _collect_spans(lhs), _collect_spans(rhs) - - assert len(lhs_spans) == len(rhs_spans) - - for i in range(len(lhs_spans)): - assert tvm.ir.structural_equal(lhs_spans[i], rhs_spans[i]) - - -def _verify_structural_equal_with_span(lhs, rhs, assert_mode=False, map_free_vars=False): - if isinstance(lhs, relay.Var) and isinstance(rhs, relay.Var): - # SEqualReduce compares the vid of Var type. Threrfore we only compare span here. - _verify_span(lhs, rhs) - return - - if assert_mode: - tvm.ir.assert_structural_equal(lhs, rhs, map_free_vars) - else: - assert tvm.ir.structural_equal(lhs, rhs, map_free_vars) - - _verify_span(lhs, rhs) diff --git a/doc/read/source-map/index.md b/doc/read/source-map/index.md index bf702cc3..6dea5657 100644 --- a/doc/read/source-map/index.md +++ b/doc/read/source-map/index.md @@ -1,4 +1,4 @@ -# Span +# 源码映射机制 ```{toctree} parse-source-map diff --git a/doc/read/source-map/span.ipynb b/doc/read/source-map/span.ipynb index 430ae74f..41e5e989 100644 --- a/doc/read/source-map/span.ipynb +++ b/doc/read/source-map/span.ipynb @@ -4,7 +4,7 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "# relay.Span" + "# `relay.Span` 示例" ] }, { diff --git a/doc/topics/frontend/caffe/network.ipynb b/doc/topics/frontend/caffe/network.ipynb index a329f05e..d871fd83 100644 --- a/doc/topics/frontend/caffe/network.ipynb +++ b/doc/topics/frontend/caffe/network.ipynb @@ -20,7 +20,7 @@ }, { "cell_type": "code", - "execution_count": 3, + "execution_count": 2, "metadata": {}, "outputs": [], "source": [ @@ -38,7 +38,7 @@ }, { "cell_type": "code", - "execution_count": 4, + "execution_count": 3, "metadata": {}, "outputs": [], "source": [ diff --git a/tests/a.toml b/tests/a.toml deleted file mode 100644 index bdfac22c..00000000 --- a/tests/a.toml +++ /dev/null @@ -1,2 +0,0 @@ -[a] -a.b = "d" \ No newline at end of file diff --git a/tests/test.ipynb b/tests/test.ipynb index 5d4a4f17..af3275b1 100644 --- a/tests/test.ipynb +++ b/tests/test.ipynb @@ -26,125 +26,136 @@ }, { "cell_type": "code", - "execution_count": 37, + "execution_count": 2, "metadata": {}, "outputs": [], "source": [ - "import torch\n", - "from torchvision.models import mobilenet_v2, MobileNet_V2_Weights\n", + "# import torch\n", + "# from torchvision.models import mobilenet_v2, MobileNet_V2_Weights\n", "\n", - "model = mobilenet_v2(weights=MobileNet_V2_Weights.DEFAULT).eval()\n", - "shape = 1, 3, 224, 224\n", - "trace = torch.jit.trace(model, torch.rand(shape).float()).eval()\n", - "torch.jit.save(trace, \".temp/mobilenet_v2.pt\")" + "# model = mobilenet_v2(weights=MobileNet_V2_Weights.DEFAULT).eval()\n", + "# shape = 1, 3, 224, 224\n", + "# trace = torch.jit.trace(model, torch.rand(shape).float()).eval()\n", + "# torch.jit.save(trace, \".temp/mobilenet_v2.pt\")" ] }, { "cell_type": "code", - "execution_count": 114, + "execution_count": 3, "metadata": {}, "outputs": [], "source": [ - "from pathlib import Path\n", - "from tomlkit import dumps, loads, load\n", - "\n", - "default_config = \"\"\"\n", - "[compile]\n", - "target = \"vta2.0\" # vta2.0, sim_vta2.0, vta1.0, vta2.0_lite, sim_vta2.0_lite\n", - "quantized_dtype = \"a8w8\" # 支持对称量化 a8w8,a8w4,a8w4_custom\n", - "quantized_method = \"channel\" # 量化方法 layer,channel,channel_custom\n", - "activation_mode = \"percentile\" # 针对激活的量化算法 max,percentile,kl_divergence,mse\n", - "weight_mode = \"max\" # 针对权重的量化算法 max,mse\n", - "channel_custom_layers = [] # 支持指定层(conv2d,conv2d_transpose,dense)的权重进行逐通道量化,在channel_custom\n", - "a8w4_custom_layers = [] # 支持指定层(conv2d,conv2d_transpose,dense)进行4bit量化,在a8w4_custom下有效\n", - "cle = false # cross-layer equalization 跨层均衡提高量化精度\n", - "output_format = \"float\" # 输出数据格式 float,fixpoint,fixpoint_6axis\n", - "weight_pad = true # 是否使用硬件对weight进行16对齐操作\n", - "fuse_conv = true # 是否进行卷积和卷积融合\n", - "\n", - "[inference]\n", - "priority = \"bandwidth_first\"\t\t# 优先级，可选参数perf_first_L0,perf_first_L1,perf_first_L2,bandwidth_first\n", - "device = \"XM_V500\"\t\t\t# 板端芯片型号, 在目标平台为vta2.0时有效, 可选设备XM_V500,XM680V200,FPGA_XM_V500,FPGA_XM680V200,SOC_XM_V500,SOC_XM680V200,XM210V200,FPGA_XM210V200,SOC_XM210V200\n", - "device_mount_dir = \"/tmp\"\t\t# 板端挂载目录，用于指明推理所需资源在板端的目录\n", - "runtime_ext = 0\t\t\t\t\t# 深度可分离卷积weight扩展功能，在推理时动态扩展或者预热阶段扩展，可选参数0,1。0:预热阶段扩展，占用内存多，性能高 1：推理时动态扩展，省内存，性能低\n", + "import torch\n", + "from torchvision.models import resnet18, ResNet18_Weights\n", + "from tvm import relay\n", "\n", - "[performance]\n", - "frequency = 37.125\t\t\t# NPU工作频率(MHz), SOC_XM_V500工作频率620MHz, SOC_XM210V200工作频率500MHz\n", - "ref_frequency = 37.125 # SOC 参考频率(MHz)\n", - "mac = 256\t\t\t\t # NPU MAC数\n", - "roofline.SOC_XM_V500 = [123.30401032, 0.96533333] # SOC_XM_V500 空载下Imax点，[Imax, MAC利用率]\n", - "roofline.SOC_XM210V200 = [122.13009695, 0.9465] # SOC_XM210V200 空载下Imax点，[Imax, MAC利用率]\n", - "bandwidth_percent = [100, 50, 20] # 空闲带宽，单位百分比\n", - "ratio_discount = 100 # 模型性能折扣，单位百分比\n", - "\"\"\"" - ] - }, - { - "cell_type": "code", - "execution_count": 115, - "metadata": {}, - "outputs": [], - "source": [ - "default_config = loads(default_config)\n", - "path = \"/media/pc/data/board/arria10/lxw/tasks/tools/npuusertools/models/pytorch/resnet18/config.toml\"\n", - "with open(path, \"rb\") as fp:\n", - " config = load(fp)\n", - "default_config.update(config)\n", - "config = default_config\n", - "# config[\"model_dir\"] = str(Path(path).resolve().parent)\n", - "# with open(\"test.toml\", \"w\") as fp:\n", - "# dump(config, fp)" - ] - }, - { - "cell_type": "code", - "execution_count": 127, - "metadata": {}, - "outputs": [], - "source": [ - "import tomlkit" - ] - }, - { - "cell_type": "code", - "execution_count": 128, - "metadata": {}, - "outputs": [], - "source": [ - "with open(\"a.toml\") as fp:\n", - " w = tomlkit.load(fp)" + "model = resnet18(weights=ResNet18_Weights.DEFAULT).eval()\n", + "shape = 1, 3, 224, 224\n", + "trace = torch.jit.trace(model, torch.rand(shape).float()).eval()\n", + "torch.jit.save(trace, \".temp/resnet18.pt\")" ] }, { "cell_type": "code", - "execution_count": 131, + "execution_count": 4, "metadata": {}, "outputs": [], "source": [ - "config = tomlkit.loads(\"\"\"\n", - "[a]\n", - "a.b = \"d\"\n", - "\"\"\")" + "mod, params = relay.frontend.pytorch.from_pytorch(trace, [(\"x\", shape)], use_parser_friendly_name=True)" ] }, { "cell_type": "code", - "execution_count": 132, + "execution_count": 5, "metadata": {}, "outputs": [ { - "data": { - "text/plain": [ - "{'a': {'a': {'b': 'd'}}}" - ] - }, - "execution_count": 132, - "metadata": {}, - "output_type": "execute_result" + "name": "stdout", + "output_type": "stream", + "text": [ + "fn (%x: Tensor[(1, 3, 224, 224), float32] /* ty=Tensor[(1, 3, 224, 224), float32] span=aten___convolution_0_x:0:0 */) -> Tensor[(1, 1000), float32] {\n", + " %0 = nn.conv2d(%x, meta[relay.Constant][0] /* ty=Tensor[(64, 3, 7, 7), float32] */, strides=[2, 2], padding=[3, 3, 3, 3], channels=64, kernel_size=[7, 7]) /* ty=Tensor[(1, 64, 112, 112), float32] */;\n", + " %1 = add(%0, meta[relay.Constant][1] /* ty=Tensor[(64, 1, 1), float32] */) /* ty=Tensor[(1, 64, 112, 112), float32] */;\n", + " %2 = nn.relu(%1) /* ty=Tensor[(1, 64, 112, 112), float32] span=aten__relu__0:0:0 */;\n", + " %3 = nn.max_pool2d(%2, pool_size=[3, 3], strides=[2, 2], padding=[1, 1, 1, 1]) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__max_pool2d_0:0:0 */;\n", + " %4 = nn.conv2d(%3, meta[relay.Constant][2] /* ty=Tensor[(64, 64, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=64, kernel_size=[3, 3]) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %5 = add(%4, meta[relay.Constant][3] /* ty=Tensor[(64, 1, 1), float32] */) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %6 = nn.relu(%5) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__relu__1:0:0 */;\n", + " %7 = nn.conv2d(%6, meta[relay.Constant][4] /* ty=Tensor[(64, 64, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=64, kernel_size=[3, 3]) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %8 = add(%7, meta[relay.Constant][5] /* ty=Tensor[(64, 1, 1), float32] */) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %9 = add(%8, %3) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__add__0:0:0 */;\n", + " %10 = nn.relu(%9) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__relu__2:0:0 */;\n", + " %11 = nn.conv2d(%10, meta[relay.Constant][6] /* ty=Tensor[(64, 64, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=64, kernel_size=[3, 3]) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %12 = add(%11, meta[relay.Constant][7] /* ty=Tensor[(64, 1, 1), float32] */) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %13 = nn.relu(%12) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__relu__3:0:0 */;\n", + " %14 = nn.conv2d(%13, meta[relay.Constant][8] /* ty=Tensor[(64, 64, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=64, kernel_size=[3, 3]) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %15 = add(%14, meta[relay.Constant][9] /* ty=Tensor[(64, 1, 1), float32] */) /* ty=Tensor[(1, 64, 56, 56), float32] */;\n", + " %16 = add(%15, %10) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__add__1:0:0 */;\n", + " %17 = nn.relu(%16) /* ty=Tensor[(1, 64, 56, 56), float32] span=aten__relu__4:0:0 */;\n", + " %18 = nn.conv2d(%17, meta[relay.Constant][10] /* ty=Tensor[(128, 64, 3, 3), float32] */, strides=[2, 2], padding=[1, 1, 1, 1], channels=128, kernel_size=[3, 3]) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %19 = add(%18, meta[relay.Constant][11] /* ty=Tensor[(128, 1, 1), float32] */) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %20 = nn.relu(%19) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__relu__5:0:0 */;\n", + " %21 = nn.conv2d(%20, meta[relay.Constant][12] /* ty=Tensor[(128, 128, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=128, kernel_size=[3, 3]) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %22 = nn.conv2d(%17, meta[relay.Constant][14] /* ty=Tensor[(128, 64, 1, 1), float32] */, strides=[2, 2], padding=[0, 0, 0, 0], channels=128, kernel_size=[1, 1]) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %23 = add(%21, meta[relay.Constant][13] /* ty=Tensor[(128, 1, 1), float32] */) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %24 = add(%22, meta[relay.Constant][15] /* ty=Tensor[(128, 1, 1), float32] */) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %25 = add(%23, %24) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__add__2:0:0 */;\n", + " %26 = nn.relu(%25) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__relu__6:0:0 */;\n", + " %27 = nn.conv2d(%26, meta[relay.Constant][16] /* ty=Tensor[(128, 128, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=128, kernel_size=[3, 3]) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %28 = add(%27, meta[relay.Constant][17] /* ty=Tensor[(128, 1, 1), float32] */) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %29 = nn.relu(%28) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__relu__7:0:0 */;\n", + " %30 = nn.conv2d(%29, meta[relay.Constant][18] /* ty=Tensor[(128, 128, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=128, kernel_size=[3, 3]) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %31 = add(%30, meta[relay.Constant][19] /* ty=Tensor[(128, 1, 1), float32] */) /* ty=Tensor[(1, 128, 28, 28), float32] */;\n", + " %32 = add(%31, %26) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__add__3:0:0 */;\n", + " %33 = nn.relu(%32) /* ty=Tensor[(1, 128, 28, 28), float32] span=aten__relu__8:0:0 */;\n", + " %34 = nn.conv2d(%33, meta[relay.Constant][20] /* ty=Tensor[(256, 128, 3, 3), float32] */, strides=[2, 2], padding=[1, 1, 1, 1], channels=256, kernel_size=[3, 3]) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %35 = add(%34, meta[relay.Constant][21] /* ty=Tensor[(256, 1, 1), float32] */) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %36 = nn.relu(%35) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__relu__9:0:0 */;\n", + " %37 = nn.conv2d(%36, meta[relay.Constant][22] /* ty=Tensor[(256, 256, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=256, kernel_size=[3, 3]) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %38 = nn.conv2d(%33, meta[relay.Constant][24] /* ty=Tensor[(256, 128, 1, 1), float32] */, strides=[2, 2], padding=[0, 0, 0, 0], channels=256, kernel_size=[1, 1]) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %39 = add(%37, meta[relay.Constant][23] /* ty=Tensor[(256, 1, 1), float32] */) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %40 = add(%38, meta[relay.Constant][25] /* ty=Tensor[(256, 1, 1), float32] */) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %41 = add(%39, %40) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__add__4:0:0 */;\n", + " %42 = nn.relu(%41) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__relu__10:0:0 */;\n", + " %43 = nn.conv2d(%42, meta[relay.Constant][26] /* ty=Tensor[(256, 256, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=256, kernel_size=[3, 3]) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %44 = add(%43, meta[relay.Constant][27] /* ty=Tensor[(256, 1, 1), float32] */) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %45 = nn.relu(%44) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__relu__11:0:0 */;\n", + " %46 = nn.conv2d(%45, meta[relay.Constant][28] /* ty=Tensor[(256, 256, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=256, kernel_size=[3, 3]) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %47 = add(%46, meta[relay.Constant][29] /* ty=Tensor[(256, 1, 1), float32] */) /* ty=Tensor[(1, 256, 14, 14), float32] */;\n", + " %48 = add(%47, %42) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__add__5:0:0 */;\n", + " %49 = nn.relu(%48) /* ty=Tensor[(1, 256, 14, 14), float32] span=aten__relu__12:0:0 */;\n", + " %50 = nn.conv2d(%49, meta[relay.Constant][30] /* ty=Tensor[(512, 256, 3, 3), float32] */, strides=[2, 2], padding=[1, 1, 1, 1], channels=512, kernel_size=[3, 3]) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %51 = add(%50, meta[relay.Constant][31] /* ty=Tensor[(512, 1, 1), float32] */) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %52 = nn.relu(%51) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__relu__13:0:0 */;\n", + " %53 = nn.conv2d(%52, meta[relay.Constant][32] /* ty=Tensor[(512, 512, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=512, kernel_size=[3, 3]) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %54 = nn.conv2d(%49, meta[relay.Constant][34] /* ty=Tensor[(512, 256, 1, 1), float32] */, strides=[2, 2], padding=[0, 0, 0, 0], channels=512, kernel_size=[1, 1]) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %55 = add(%53, meta[relay.Constant][33] /* ty=Tensor[(512, 1, 1), float32] */) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %56 = add(%54, meta[relay.Constant][35] /* ty=Tensor[(512, 1, 1), float32] */) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %57 = add(%55, %56) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__add__6:0:0 */;\n", + " %58 = nn.relu(%57) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__relu__14:0:0 */;\n", + " %59 = nn.conv2d(%58, meta[relay.Constant][36] /* ty=Tensor[(512, 512, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=512, kernel_size=[3, 3]) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %60 = add(%59, meta[relay.Constant][37] /* ty=Tensor[(512, 1, 1), float32] */) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %61 = nn.relu(%60) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__relu__15:0:0 */;\n", + " %62 = nn.conv2d(%61, meta[relay.Constant][38] /* ty=Tensor[(512, 512, 3, 3), float32] */, padding=[1, 1, 1, 1], channels=512, kernel_size=[3, 3]) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %63 = add(%62, meta[relay.Constant][39] /* ty=Tensor[(512, 1, 1), float32] */) /* ty=Tensor[(1, 512, 7, 7), float32] */;\n", + " %64 = add(%63, %58) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__add__7:0:0 */;\n", + " %65 = nn.relu(%64) /* ty=Tensor[(1, 512, 7, 7), float32] span=aten__relu__16:0:0 */;\n", + " %66 = nn.adaptive_avg_pool2d(%65, output_size=[1, 1]) /* ty=Tensor[(1, 512, 1, 1), float32] span=aten__adaptive_avg_pool2d_0:0:0 */;\n", + " %67 = reshape(%66, newshape=[0, -1, 1, 1]) /* ty=Tensor[(1, 512, 1, 1), float32] span=aten__flatten_0:0:0 */;\n", + " %68 = squeeze(%67, axis=[2, 3]) /* ty=Tensor[(1, 512), float32] span=aten__flatten_0:0:0 */;\n", + " %69 = nn.dense(%68, meta[relay.Constant][40] /* ty=Tensor[(1000, 512), float32] */, units=None) /* ty=Tensor[(1, 1000), float32] span=aten__linear_0:0:0 */;\n", + " add(%69, meta[relay.Constant][41] /* ty=Tensor[(1000), float32] */) /* ty=Tensor[(1, 1000), float32] */\n", + "} /* ty=fn (Tensor[(1, 3, 224, 224), float32]) -> Tensor[(1, 1000), float32] */\n", + "\n" + ] } ], "source": [ - "config" + "import tvm\n", + "\n", + "with tvm.transform.PassContext(opt_level=3):\n", + " mod = relay.quantize.prerequisite_optimize(mod, params)\n", + "print(mod[\"main\"])" ] }, {