本系列是北航计算机学院于 2024 年春季学期开设的一般专业课《FPGA多核并行计算》课程的实验报告记录，由于学习过程中掌握并不牢靠，如有错误请读者不吝赐教！

FPGA深度学习加速器：VTA

实验说明与实验相关内容上传在 Github 仓库中。实验过程中使用到的环境与程序：

  • 运行环境：Ubuntu 16.04，在线实验平台
  • apache/tvm

Apache TVM 是一个开源的机器学习编译框架，用于 CPU、GPU 和机器学习的加速。它的设计目标是令机器学习工程师可以在不同的硬件后端高效地优化和运行计算过程。

实验分析

  • 了解深度学习加速器 VTA 的基本原理
  • 掌握 VTA 环境的基本搭建方法

配置VTA模拟器运行环境

前排提醒，文内出现的<tvm root> 请替换为你的 tvm 项目根目录，llvm 同理

配置 TVM 环境前，需要预先准备好其编译链工具：

  • gcc 9.3.0
  • llvm 9.0.0
  • python 3.7

首先需要在宿主机上配置好模拟和仿真 VTA 环境。

由于 TVM 更新较快，所以本实验需要下载特定版本的 TVM，需要注意如果选择从 Github 直接下载源码，也要把链接到的外部仓库文件也一并下载。

版本hash：09c55fd1f3354d2280bb792a252590ac6bd68e58
  • 配置环境变量：TVM_PATH、GCC、LLVM 的 PATH
export TVM_PATH=<tvm root>

export PATH=/usr/local/gcc-9/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/gcc-9/lib64
export MANPATH=/usr/local/gcc-9/share/man:$MANPATH

export PATH=$PATH:<llvm root>/bin
  • 配置 VTA 仿真功能库：复制配置文件并进行修改
cd <tvm root>
mkdir build
cp cmake/config.cmake build/.
echo 'set(USE_VTA_FSIM ON)' >> build/config.cmake
echo 'set(USE_LLVM ON)' >> build/config.cmake
cd build && cmake .. && make -j4

image-20240417193308196

  • 将 VTA python 运行库添加到 python 库路径中
export PYTHONPATH=$TVM_PATH/python:$TVM_PATH/vta/python:$TVM_PATH/topi/python:${PYTHONPATH}

最终测试宿主机的 VTA 安装状态，测试 VTA 对卷积运算的支持

python37 <tvm root>/vta/tests/python/integration/test_benchmark_topi_conv2d.py

image-20240417193459760

配置 VTA FPGA 运行环境

  • 通过 jupyter notebook 将 tvm-09c 上传到开发板上，解压后构建运行环境
tar -zxvf tvm.tar.gz
cd <tvm root>
mkdir build
cp cmake/config.cmake build/.
echo 'set(USE_VTA_FPGA ON)' >> build/config.cmake
cp vta/config/pynq_sample.json vta/config/vta_config.json
cd build
cmake ..
make runtime vta -j2

编译时遇到如下问题：

image-20240417194805005

解决方法

但是因为是第一次编译，所以时间也没啥大问题，最后也是正常跑完了

image-20240417195535593

  • 修改 VTA 运行的配置
{
  "TARGET": "pynq",
  "HW_VER": "0.0.1",
  "LOG_INP_WIDTH": 3,
  "LOG_WGT_WIDTH": 3,
  "LOG_ACC_WIDTH": 5,
  "LOG_BATCH": 0,
  "LOG_BLOCK": 4,
  "LOG_UOP_BUFF_SIZE": 15,
  "LOG_INP_BUFF_SIZE": 15,
  "LOG_WGT_BUFF_SIZE": 18,
  "LOG_ACC_BUFF_SIZE": 17,
  "LOG_BLOCK_IN": 4,
  "LOG_BLOCK_OUT": 4,
  "LOG_OUT_WIDTH": 3,
  "LOG_OUT_BUFF_SIZE": 15
}
  • 配置开发板环境变量
export TVM_HOME=<tvm root>
  • 启动 RPC 服务器
cd ..
python3 -m vta.exec.rpc_server --port 9091

使用宿主机调用 FPGA VTA 环境

  • 配置 RPC 服务的地址，测试 VTA FPGA 运行环境是否正常运行

由于实验环境中宿主机和 FPGA 开发板并不处于同一环境，所以需要将 VTA 服务器运行在 9091 端口（上文），并通过申请开发板时提供的 VTA 服务端口访问 VTA 服务器

export VTA_PYNQ_RPC_HOST=开发板IP
export VTA_PYNQ_RPC_PORT=开发板VTA服务映射端口
  • 配置宿主机 VTA 运行环境参数
cd /path/to/tvm/vta/config
mv vta_config.json vta_config_sim.json
cp pynq_sample.json vta_config.json
  • 测试 VTA RPC 连通性
python37 vta/tests/python/pynq/test_program_rpc.py

然后失败了（）

image-20240417210736901

经过助教学长检查是平台服务器端口出现了问题，修复后再进行实验成功

image-20240417232547280

再次运行 2D 卷积测试用例

python37 vta/tests/python/integration/test_benchmark_topi_conv2d.py

image-20240417232641755

成功完成 2D 卷积测试

问题思考

自学 VTA 矩阵乘法案例并提交结果

VTA 矩阵乘法案例

首先要提的是实验文档里给出的连接已经宕掉了，这个页面是从 gitee 上托管的一个远古版本找到的

首先将源码复制到本地，然后按照代码中的注释进行分析：

  • Line 47：按照路径修改配置文件
# Load VTA parameters from the 3rdparty/vta-hw/config/vta_config.json file
env = vta.get_env()

配置内容和 FPGA 平台上的配置内容一致即可

  • Line 50：配置 RPC 远程的 IP 和端口，注意到这里用的环境变量字段和我们之前添加的并不完全一样，所以要进行修改
# We read the Pynq RPC host IP address and port number from the OS environment
host = os.environ.get("VTA_RPC_HOST", "192.168.2.99")
port = int(os.environ.get("VTA_RPC_PORT", "9091"))

# VTA_RPC_HOST -> VTA_PYNQ_RPC_HOST
# VTA_RPC_PORT -> VTA_PYNQ_RPC_PORT
  • Line 55：若运行 target 是 PYNQ 或 DE-10 的开发板，则利用 RPC 服务连接 RPC 服务器
if env.TARGET == "pynq" or env.TARGET == "de10nano":

    # Make sure that TVM was compiled with RPC=1
    assert tvm.runtime.enabled("rpc")
    remote = rpc.connect(host, port)
    vta.reconfig_runtime(remote)
    vta.program_fpga(remote, bitstream=None)
  • Line 165：创建矩阵的占位符 A、B 和中间张量 A_buf、B_buf。

两个 buf 变量会在运算过程中存储在 VTA 的片上缓存中，计算乘法结果后产生 C_buf，复制回结果张量 C 中，完成运算。过程如图所示：

gemm_dataflow

  • Line 202：定义了输入数据内部和外部两个规约轴 ko 和 ki，并执行了 VTA 的矩阵乘法操作。
ko = te.reduce_axis((0, n), name="ko")
ki = te.reduce_axis((0, env.BLOCK_IN), name="ki")
# Describe the in-VTA matrix multiplication
C_buf = te.compute(
    (o, m, env.BATCH, env.BLOCK_OUT),
    lambda bo, co, bi, ci: te.sum(
        A_buf[bo, ko, bi, ki].astype(env.acc_dtype) * B_buf[co, ko, ci, ki].astype(env.acc_dtype),
        axis=[ko, ki],
    ),
    name="C_buf",
)
  • Line 239：为了让数据能够在 VTA 的 DRAM 上存储，并且减少内存数据的传输量，需要执行类型转换
C = te.compute(
    (o, m, env.BATCH, env.BLOCK_OUT), lambda *i: C_buf(*i).astype(env.inp_dtype), name="C"
)

tensor_core

TVM 要求用户提供一种称为调度（schedule）的计算实现。调度是对原始计算的一组变形，它改变了计算的实现方式，而不影响正确性。我们实际上可以通过调整调度来实现对计算方式、数据存储、DMA 传输等方面的控制。

  • Line 269：创建一个调度 s，并设置中间运算变量在 VTA 片上缓冲区所放的区域
s = te.create_schedule(C.op)
# Set the intermediate tensor's scope to VTA's on-chip buffers
s[A_buf].set_scope(env.inp_scope)
s[B_buf].set_scope(env.wgt_scope)
s[C_buf].set_scope(env.acc_scope)
  • env.inp_scope：输入缓冲区，属于只读 SRAM 缓冲区，存储形状为 (env.BATCH, env.BLOCK_IN) 并且类型为 env.inp_dtype 的输入矩阵
  • env.wgt_scope：权重缓冲区，属于只读 SRAM 缓冲区，存储形状为 (env.BLOCK_OUT, env.BLOCK_IN) 并且类型为 env.wgt_dtype 的权重矩阵
  • env.acc_scope：累加器缓冲区，可读写 SRAM 缓冲区，它存储形状为 (env.BATCH, env.BLOCK_OUT) 类型为 env.acc_dtype 的累加器矩阵。累加器缓冲区是 VTA 的通用寄存器文件：保存矩阵乘法、池化、批量归一化和激活层等计算的中间结果
  • Line 339：DMA 传输。使用 compute at 将缓冲区的复制操作结合在矩阵乘法的循环中；再使用 pragma 原语控制计算 A_buf、B_buf 和 C 时应该直接通过 env.dma_copy 操作从内存中直接传输
# Move buffer copy into matrix multiply loop
s[A_buf].compute_at(s[C_buf], ko)
s[B_buf].compute_at(s[C_buf], ko)
# Tag the buffer copies with the DMA pragma to insert a DMA transfer
s[A_buf].pragma(s[A_buf].op.axis[0], env.dma_copy)
s[B_buf].pragma(s[B_buf].op.axis[0], env.dma_copy)
s[C].pragma(s[C].op.axis[0], env.dma_copy)
  • Line 371：张量化（Tensorization），这部分没有看明白，这一小块好像是用 reorder 重排了循环的顺序再进行了张量化（？）这也是算法中关于调度部分的结束
  • Line 389：TVM 编译，在创建一个 GEMM 内核后，我们可以将调度内容封装为一个 TVM 函数，保存在内核里，并将其上传、加载在 RPC 服务器中
my_gemm = vta.build(s, [A, B, C], "ext_dev", env.target_host, name="my_gemm")
temp = utils.tempdir()
my_gemm.save(temp.relpath("gemm.o"))
remote.upload(temp.relpath("gemm.o"))
f = remote.load_module("gemm.o")
  • Line 418：运行函数，先随机初始化两个矩阵 A_orig 和 B_orig 再打包、格式化到 DLPack 标准，最后调用封装好的模块 f 进行计算
  • Line 447：最后通过 numpy 验证 TVM 模块计算后结果的正确性，通过所有的 assert_equal 后表明矩阵乘法的操作是成功的
# Compute reference result with numpy
C_ref = np.dot(A_orig.astype(env.acc_dtype), B_orig.T.astype(env.acc_dtype)).astype(C.dtype)
C_ref = C_ref.reshape(o, env.BATCH, m, env.BLOCK_OUT).transpose((0, 2, 1, 3))
np.testing.assert_equal(C_ref, C_nd.numpy())
  • 基于线上实验平台的 TVM 版本，官方帮助文档的代码仍有一些部分不可用，需要作出修改：
  • Line 43：报错找不到模块 utils，查看项目内文件路径发现并不含有 tvm.contrib.utils，实际上应该是 util；同步修改 Line 392 的使用
from tvm.contrib import utils

temp = utils.tempdir()
  • Line 50：在上面已经提及，需要更改获取环境变量的名称
  • Line 449：C_nd 不包含方法 numpy()，使用 numpy 进行 TVM 效果验证时 C_nd 应该使用 asnumpy 函数进行类型的转换
np.testing.assert_equal(C_ref, C_nd.numpy())

最终结果：image-20240417235302186

完整的输出结果包含了一部分打印的调度信息，可以从这里里看具体的内容