编译器深度学习模型优化【免费下载链接】tvmOpen deep learning compiler stack for cpu, gpu and specialized accelerators项目地址https://gitcode.com/gh_mirrors/tvm7/tvm点击查看免费下载导读VTAVersatile Tensor Accelerator是 TVM 内置的可配置深度学习加速器涵盖 FPGA 硬件设计与配套软件工具链。本文以 VTA Python API 参考文档 为核心主线结合仓库源码逐层拆解vta包的 7 个核心 API硬件信息Environment/get_env、RPC 工具reconfig_runtime/program_fpga以及编译器 APIbuild_config/lower/build。读完本文你将掌握 VTA 的硬件配置机制、远程 FPGA 编程与运行时重建流程以及如何用 TVM 完整编译一个面向 VTA 后端的算子内核。一、VTA Python 包编译器工具链与硬件访问的统一入口vta是一个 Python 包位于 vta/python/vta/其__init__.py中说明VTA Package is a TVM backend extension to support VTA hardware。除了编译器工具链它还包含配置硬件环境、通过 RPC 访问远程设备的工具函数。包的顶层导出如下硬件环境Environment、get_env来自 environment.pyRPC 工具reconfig_runtime、program_fpga来自 rpc_client.py位流管理get_bitstream_path、download_bitstream来自 bitstream.py编译器 APIbuild_config、lower、build来自 build_module.py值得注意的是__init__.py中有一段条件导入逻辑if not tvm._ffi.base._RUNTIME_ONLY: from . import top from .build_module import build_config, lower, build即在 TVM 运行时runtime-only模式下vta.exec.rpc_server只加载 libvta 运行时库不会导入编译器相关模块。包内部还包括 top/TVM TOPI 连接器与 VTA 专用算子、autotvm.pyAutoTVM 集成、exec/rpc_server.pyVTA 定制 RPC 服务器等子模块。从源码结构看vta包的设计遵循文档中的定位We program VTA using TVM, so the compiler API in vta package is only a thin wrapper to provide VTA specific extensions——编译器主体完全复用 TVMvta只负责提供 VTA 特有的扩展。二、硬件信息vta.Environment与vta.get_env2.1 什么是 EnvironmentEnvironment是硬件配置对象包含了为特定 VTA 后端编译所需的全部信息。它的构造函数接受一个cfg字典key 为字符串、值为对应配置值并通过pkg_config()加载3rdparty/vta-hw/config/pkg_config.py中的PkgConfig来派生参数。关键派生逻辑见 environment.pydef __init__(self, cfg): self.pkg pkg_config(cfg) self.__dict__.update(self.pkg.cfg_dict) # data type width self.INP_WIDTH 1 self.LOG_INP_WIDTH self.WGT_WIDTH 1 self.LOG_WGT_WIDTH self.ACC_WIDTH 1 self.LOG_ACC_WIDTH self.OUT_WIDTH 1 self.LOG_OUT_WIDTH # tensor intrinsic shape self.BATCH 1 self.LOG_BATCH self.BLOCK_IN 1 self.LOG_BLOCK_IN self.BLOCK_OUT 1 self.LOG_BLOCK_OUT # buffer size self.UOP_BUFF_SIZE 1 self.LOG_UOP_BUFF_SIZE self.INP_BUFF_SIZE 1 self.LOG_INP_BUFF_SIZE self.WGT_BUFF_SIZE 1 self.LOG_WGT_BUFF_SIZE self.ACC_BUFF_SIZE 1 self.LOG_ACC_BUFF_SIZE self.OUT_BUFF_SIZE 1 self.LOG_OUT_BUFF_SIZE所有带LOG_前缀的配置项都以 log2 形式存储Environment会统一展开为实际的宽度、维度与缓冲大小并据此推导出字节数INP_ELEM_BYTES BATCH * BLOCK_IN * INP_WIDTH // 8等数据类型inp_dtype/wgt_dtype/acc_dtype/out_dtype例如默认配置下分别为int8、int8、int32、int8位流与模型标识BITSTREAM self.pkg.bitstreamMODEL TARGET _ BITSTREAM。2.2 vta_config.json架构规格的唯一来源VTA 是开箱即用的可配置加速器其架构规格由3rdparty/vta-hw/config/vta_config.json提供该文件同时参数化 TVM 编译器栈与 VTA 硬件栈。环境变量VTA_HW_PATH用于定位该目录默认即仓库内的3rdparty/vta-hw。参数说明表如下引自 docs/topic/vta/dev/config.rst属性格式说明TARGETStringTVM 设备目标可取pynq、ultra96、sim快速模拟、tsimverilator 周期精确模拟等HW_VERStringVTA 硬件版本号用于唯一标识硬件位流LOG_INP_WIDTHInt (log2)输入数据类型的有符号整型位宽LOG_WGT_WIDTHInt (log2)权重数据类型的有符号整型位宽LOG_ACC_WIDTHInt (log2)累加器数据类型的有符号整型位宽LOG_BATCHInt (log2)VTA 矩阵乘内禀的输入/输出第 0 维即 (A,B)×(B,C) 中的 ALOG_BLOCKInt (log2)VTA 矩阵乘内维度即 (A,B)×(B,C) 中的 B 与 CLOG_UOP_BUFF_SIZEInt (log2)微指令micro-op片上缓冲单位字节LOG_INP_BUFF_SIZEInt (log2)输入片上缓冲单位字节LOG_WGT_BUFF_SIZEInt (log2)权重片上缓冲单位字节LOG_ACC_BUFF_SIZEInt (log2)累加器片上缓冲单位字节由于这些参数只能取 2 的幂故以 log2 值描述例如 8 位输入整型对应LOG_INP_WIDTH364KB 微指令缓冲对应LOG_UOP_BUFF_SIZE16。修改该 JSON 后执行make即可按新参数重建 VTA见 VTA 安装指南。Environment会在初始化时自动读取该文件_init_env()见 environment.py并创建全局默认环境Environment.current。2.3 上下文管理器与vta.get_env()Environment实现了 Python 上下文管理器协议其内部维护了一个环境栈def __enter__(self): self._last_env Environment.current Environment.current self return self def __exit__(self, ptype, value, trace): Environment.current self._last_env因此可以在with块内临时切换到新配置例如 environment.py 文档中的用法new_cfg json.load(json.load(open(new_cfg.json))) with vta.Environment(new_cfg): env vta.get_env() # 获取当前生效的环境vta.get_env()直接返回Environment.current是获取当前 VTA 环境的标准入口测试脚本中随处可见例如 test_benchmark_topi_conv2d.py 中的env vta.get_env()。2.4 target 与 target_hostEnvironment根据TARGET自动推导 TVM 编译目标environment.pyproperty def target(self): return tvm.target.vta(modelself.MODEL) property def target_host(self): if self.TARGET in [pynq, de10nano]: return llvm -mtriplearmv7-none-linux-gnueabihf if self.TARGET ultra96: return llvm -mtripleaarch64-linux-gnu if self.TARGET in [sim, tsim, intelfocl]: return llvm raise ValueError(Unknown target %s % self.TARGET)target为 VTA 外设目标模型名形如pynq_f1_...编译时通过tvm.target.Target(ext_dev, hostenv.target_host)组合使用见 test_benchmark_gemm.pytarget_host是主机侧 LLVM 目标Pynq/DE10 为 ARMv7 交叉编译三元组Ultra96 为 AArch64模拟器场景直接使用llvm。2.5 开发者上下文 DevContextEnvironment.dev返回内部开发者上下文DevContext集中存放面向编译器的非用户可见常量environment.pyDMA 内存 IDMEM_ID_UOP0、MEM_ID_WGT1、MEM_ID_INP2、MEM_ID_ACC3、MEM_ID_OUT4、MEM_ID_ACC_8BIT5ALU 操作码ALU_OPCODE_MIN0、MAX1、ADD2、SHR3、MUL4任务队列 ID流水线阶段QID_LOAD_INP1、QID_LOAD_WGT1、QID_LOAD_OUT2、QID_STORE_OUT3、QID_COMPUTE2VTA 线程轴te.thread_axis(vta)、命令句柄tir.tvm_thread_context及 GEMM 内禀intrin.gemm。Environment还提供mock 模式env.mock将dma_copy与alupragma 分别替换为skip_dma_copy/skip_alu用于在不涉及真实硬件的场景下跳过相应内禀。三、RPC 工具reconfig_runtime与program_fpgaVTA 支持把编译产物发送到远程 FPGA 板卡执行。这两个函数rpc_client.py面向tvm.rpc.RPCSession是连接编译器与远端硬件的桥梁。3.1vta.reconfig_runtime(remote)根据当前硬件规格重建远端运行时def reconfig_runtime(remote): env get_env() freconfig remote.get_function(tvm.contrib.vta.reconfig_runtime) freconfig(env.pkg.cfg_json)它把当前Environment的配置 JSON 推送到远端触发远端 VTA 运行时重新编译。在 RPC 服务器端exec/rpc_server.py该回调会把收到的cfg_json与本地vta_config.json对比若配置不同则用cc.create_shared以-O2 -stdc17重新编译libvta共享库并写回配置文件。因此每当修改vta_config.json都必须重新执行reconfig_runtime确保 JIT 生成的加速器可执行代码与 FPGA 上的 VTA 设计规格一致Pynq 上约需 30 秒。3.2vta.program_fpga(remote, bitstreamNone)上传并编程位流def program_fpga(remote, bitstreamNone): env get_env() if bitstream: assert os.path.isfile(bitstream) else: bitstream get_bitstream_path() if not os.path.isfile(bitstream): if env.TARGET de10nano: return download_bitstream() if isinstance(remote, rpc.LocalSession): program_bitstream.bitstream_program(env.TARGET, bitstream) else: fprogram remote.get_function(tvm.contrib.vta.init) remote.upload(bitstream) fprogram(os.path.basename(bitstream))逻辑要点位流参数可选不传时通过get_bitstream_path()定位缓存位流默认缓存于$VTA_CACHE_PATH即~/.vta_cache/TARGET/HW_VER/BITSTREAM.bitbitstream.py不存在时自动调用download_bitstream()从 VTA 位流仓库下载DE10-Nano 特殊处理为直接返回。本地会话直接调用program_bitstream.bitstream_program按目标分发——Pynq/Ultra96 走pynq.BitstreamDE10-Nano 走vta.de10nano.programIntel FPGA OpenCL 走vta.oclfpga.programsim/tsim则为 no-opprogram_bitstream.py。远程会话先remote.upload(bitstream)上传再调用远端tvm.contrib.vta.init完成 FPGA 编程。远端回调会按TARGET重置 Xilinx 驱动xlnk.Xlnk().xlnk_reset()或加载 DE10-Nano 库exec/rpc_server.py。典型用法自定义位流场景见 install.rstvta.program_fpga(remote, bitstreamtvm root/3rdparty/vta-hw/build/hardware/xilinx/vivado/configuration/export/vta.bit)3.3 测试基础设施如何选择运行目标VTA 的测试工具vta.testing.run()testing/utils.py按TARGET分派sim/tsim/intelfocl优先连接本地 RPC环境变量VTA_LOCAL_SIM_RPC否则使用rpc.LocalSession()直跑模拟器pynq/ultra96/de10nano优先通过 trackerTVM_TRACKER_HOST/TVM_TRACKER_PORT获取远端否则要求设置VTA_RPC_HOST/VTA_RPC_PORT直连板卡。这解释了为何真实硬件测试需要先导出环境变量如 Pynq 的export VTA_RPC_HOST192.168.2.99、export VTA_RPC_PORT9091。模拟器场景下还可以用 simulator.py 的clear_stats()/stats()采集执行统计见 test_benchmark_gemm.py。四、编译器 APIbuild_config、lower与buildVTA 用 TVM 编程因此vta的编译器 API 只是薄封装自动注入 VTA 特有的扩展 pass。4.1vta.build_config(debug_flag0, **kwargs)返回一个配置好 VTA 降级 pass 序列的tvm.transform.PassContextbuild_module.pypass_list [ (0, transform.InjectConv2DTransposeSkip()), (1, transform.InjectDMAIntrin()), (1, transform.InjectSkipCopy()), (1, transform.AnnotateALUCoProcScope()), (1, tvm.tir.transform.LiftAttrScope(coproc_uop_scope)), (1, transform.LiftAllocToScopeBegin()), (1, tvm.tir.transform.LiftAttrScope(coproc_scope)), (1, transform.InjectCoProcSync()), (1, EarlyRewrite()), ] if debug_flag: pass_list.append((1, add_debug)) pass_list.append((2, transform.InjectALUIntrin())) pass_list.append((3, tvm.tir.transform.LowerDeviceStorageAccessInfo())) pass_list.append((3, transform.FoldUopLoop())) pass_list.append((3, transform.CPUAccessRewrite())) config {tir.add_lower_pass: pass_list} return tvm.transform.PassContext(configconfig, **kwargs)各 pass 的职责实现见 transform.pyPass作用InjectConv2DTransposeSkip跳过转置卷积中 stride1 的零权重生成tir.vta.uop_push微指令InjectDMAIntrin把dma_copypragma 注入为VTALoadBuffer2D/VTAStoreBuffer2D外部调用并完成 2D 模式的步长/填充分析支持 5 维内的 padding 加载InjectSkipCopy/AnnotateALUCoProcScope处理skip_dma_copy/skip_alu调试 pragma为alu标注coproc_uop_scopeLiftAllocToScopeBegin将分配提升到当前作用域起始处InjectCoProcSync注入coproc_sync同步并调用 TVM 的CoProcSyncEarlyRewrite提前做存储重写失败时回退原模块InjectALUIntrin将加/减/乘/取min/取max/移位等表达式翻译为 ALU 微指令操作码FoldUopLoop检测 GEMM 循环结构折叠为VTAUopLoopBegin/VTAUopLoopEndCPUAccessRewrite检测 CPU 对 VTA 缓冲的访问改写为VTABufferCPUPtr返回的指针debug_flag非 0 时会在函数体开头注入VTASetDebugMode调用。Environment定义了若干调试位如DEBUG_DUMP_INSN 1 1、DEBUG_DUMP_UOP 1 2、DEBUG_SKIP_READ_BARRIER 1 3、DEBUG_SKIP_WRITE_BARRIER 1 4。kwargs直接透传给PassContext因此可以传入disabled_pass等标准 TVM 配置如测试中的with vta.build_config(disabled_pass{tir.CommonSubexprElimTIR}):见 test_benchmark_topi_conv2d.py。4.2vta.lower与vta.build两者分别是tvm.lower/tvm.build的薄封装若当前PassContext中已存在tir.add_lower_pass则直接透传否则自动套用vta.build_config()build_module.py。这意味着with vta.build_config(): vta_module tvm.build(s, ...)与直接调用vta.build(s, ...)等价——后者会隐式创建 VTA 的 PassContext。测试中典型的完整调用为mod vta.build( s, [data, weight, res], tvm.target.Target(ext_dev, hostenv.target_host), namegemm, ) mod.save(temp.relpath(gemm.o)) remote.upload(temp.relpath(gemm.o)) f remote.load_module(gemm.o)见 test_benchmark_gemm.py。构建产物是.o目标文件上传到远端后由 RPC 服务器加载执行。4.3 算子与内存信息的底层注册build_module.py底部还完成了关键的全局注册build_module.py为tir.vta.coproc_sync、coproc_dep_push、coproc_dep_pop、uop_push、command_handle注册TCallEffectKindOpaque及全局符号如VTAUopPush、VTATLSCommandHandle通过tvm.register_func(tvm.info.mem.%s % scope)为三个片上缓冲作用域注册内存信息local.inp_buffer、local.wgt_buffer、local.acc_buffer单位位宽、最大 SIMD 位宽与缓冲大小均取自Environment注册内禀降级coproc_sync → VTASynchronize、coproc_dep_push → VTADepPush、coproc_dep_pop → VTADepPop把 TVM IR 中的抽象同步原语翻译为 VTA 运行时调用。五、上层扩展VTA 算子、AutoTVM 与图打包除文档列出的 7 个核心 API 外vta包还提供面向应用的上层能力VTA 专用算子top/conv2d_packed、conv2d_transpose_packed、group_conv2d_packed、dense_packed及对应 schedule以及bitpack数据位打包与graph_pack图打包把标准 Relay 图转换为 VTA 的 packed 表示。AutoTVM 集成autotvm.pymodule_loader(bitstreamNone)构造 VTA 专属的 ModuleLoader在每次测量前自动调用program_fpgareconfig_runtime重编程 FPGA保证搜索过程与目标硬件规格一致。自定义 RPC 服务器exec/rpc_server.pymain()支持--host、--port默认 9091、--port-end默认 9199、--key、--tracker参数也可配合apps/vta_rpc/start_rpc_server.sh启动。六、环境准备与验证对应 API 的落地运行前提要实际运行上述 API需要按 VTA 安装指南 完成环境搭建export TVM_PATHpath to TVM root export VTA_HW_PATH$TVM_PATH/3rdparty/vta-hw # 编译 TVM 时启用 VTA 功能模拟库 cd tvm-root mkdir build cp cmake/config.cmake build/. echo set(USE_VTA_FSIM ON) build/config.cmake cd build cmake .. make -j4 # 将 VTA Python 库加入路径 export PYTHONPATH/path/to/vta/python:${PYTHONPATH}随后可用 2D 卷积基准验证整套 API该脚本依次使用vta.get_env()、vta.lower()、vta.build_config()、vta.build()python tvm root/vta/tests/python/integration/test_benchmark_topi_conv2d.py自定义 VTA 参数化时编辑3rdparty/vta-hw/config/vta_config.json后执行make重建TARGET设为sim时所有 TVM 工作负载都在 VTA 模拟器上执行无需额外硬件即可验证本文所述的全部 Python API。总结vta包的 Python API 以Environment为单一事实来源一份vta_config.json同时参数化硬件、编译器 pass 与运行目标get_env()在任意位置获取当前环境program_fpga与reconfig_runtime通过 RPC 完成编程 FPGA 重建运行时的硬件部署闭环build_config/lower/build则以最小侵入的方式把 VTA 的 DMA、ALU、GEMM 内禀与同步原语注入 TVM 的标准编译流水线。无论是跑模拟器、Pynq 板卡还是 DE10-Nano这套 API 都是使用 VTA 的必经入口。赞分享编译器深度学习模型优化【免费下载链接】tvmOpen deep learning compiler stack for cpu, gpu and specialized accelerators项目地址https://gitcode.com/gh_mirrors/tvm7/tvm点击查看免费下载相关推荐TVM TIRx 编译 API 深度解析tvm.tirx.build 与 compilation_pipeline 编译流水线实战指南TVM TIRx 编译 API 深度解析tvm.tirx.build 与 compilation_pipeline 编译流水线实战指南 TVM 的 TIRx模型编译深度学习推理引擎PaddleOCR 通用 OCR 流水线 C 本地部署实战Windows 环境编译与推理指南PaddleOCR 通用 OCR 流水线 C 本地部署实战Windows 环境编译与推理指南 通用 OCR 流水线General OCR Pipelin人工智能计算机视觉OCR深度学习大模型RAGTVM VTA 安装指南从功能仿真到 Pynq / DE10-Nano FPGA 硬件部署与自定义比特流生成TVM VTA 安装指南从功能仿真到 Pynq / DE10 Nano FPGA 硬件部署与自定义比特流生成 导读 本指南以 TVM 开源深度学习编译器栈中的编译器深度学习模型优化创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考