TVM调试指南

1. TVM安装

这部分之前就写过，为了方便，这里再复制一遍。
首先下载代码
git clone —recursive https://github.com/dmlc/tvm
这个地方最好使用—recursive选项，不然会缺dlpack这些库，原因是
子模组 ‘HalideIR’ (https://github.com/dmlc/HalideIR) 未对路径 ‘3rdparty/HalideIR’ 注册
子模组 ‘dlpack’ (https://github.com/dmlc/dlpack) 未对路径 ‘3rdparty/dlpack’ 注册
子模组 ‘dmlc-core’ (https://github.com/dmlc/dmlc-core) 未对路径 ‘3rdparty/dmlc-core’ 注册
子模组 ‘3rdparty/rang’ (https://github.com/agauniyal/rang) 未对路径 ‘3rdparty/rang’ 注册
sudo apt-get update
sudo apt-get install -y python python-dev python-setuptools gcc libtinfo-dev zlib1g-dev
创建要编译生成so的文件夹（该文件夹位于tvm源码目录下，与src同级）
mkdir build
拷贝一份官方的cmake文件进行（测试时先使用官方的，之后这个config.camke文件我们要进行修改以支持更多的设备）
cp cmake/config.cmake build
修改该文件，这里我们的服务器上支持CUDA和LLVM的环境，因此将这两个配置打开
set(USE_CUDA OFF)
set(USE_LLVM OFF)
修改为：
set(USE_CUDA ON)
set(USE_LLVM ON)
修改好配置文件后，进行编译。因为修改了两个编译选项，因此首先需要cmake重新生成Makefile，以后每次新添加了文件和文件夹，一定要重新cmake，否则文件很可能没有编译。
cd build
好像最新版本编译出来的默认不是debug版本，为了保险，手动选择Debug选项
cmake -DCMAKE_BUILD_TYPE=Debug ..
make -j4
上边的三个步骤，非常关键，建议不要随便改变
在无错误编译完成后，build目录下形成了libtvm*.so之类的文件，我们因为要修改tvm，所以不建议移动这些so文件到python目录下，建议添加响应的配置。具体配置如下：
在.bashrc文件中，添加
export TVM_PATH=path/to/tvm
export PYTHONPATH=$TVM_PATH/python:$TVM_PATH/topi/python:$TVM_PATH/nnvm/python:${PYTHONPATH}

2. TVM测试代码

在安装完成后，进入python的命令行，使用
>>>import tvm
测试tvm使用可以使用
报错找不到module的原因可能是，配置不对，或者配置没有生效。简单的查看配置路径是否正确，可以按照以下命令进行
>>>import sys
>>>sys.path
查看python寻找module的文件夹，以排查错误。
下边列出生成CUDA代码的的一段python测试程序，程序来源于tvm.ai，略有改动
生成的结果如下：
import tvm import numpy as np import timeit import pdb import os

raw_input(os.getpid())

The size of the matrix

(M, K) x (K, N)

You are free to try out different shapes, sometimes TVM optimization outperforms numpy with MKL.

M = 1024 K = 1024 N = 1024

The default tensor type in tvm

dtype = “float32”

using Intel AVX2(Advanced Vector Extensions) ISA for SIMD

To get the best performance, please change the following line

to llvm -mcpu=core-avx2, or specific type of CPU you use

tgt = “cuda” tgt_host=”llvm”

n = tvm.var(“n”) A = tvm.placeholder((n,), name=’A’) B = tvm.placeholder((n,), name=’B’) C = tvm.compute(A.shape, lambda i: A[i] + B[i], name=”C”) s = tvm.create_schedule(C.op)

pdb.set_trace()

bx, tx = s[C].split(C.op.axis[0], factor=64) if tgt == “cuda”: s[C].bind(bx, tvm.thread_axis(“blockIdx.x”)) s[C].bind(tx, tvm.thread_axis(“threadIdx.x”))

fadd = tvm.build(s, [A, B, C], tgt, target_host=tgt_host, name=”myadd”) ctx = tvm.context(tgt, 0)

n = 1024 a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), ctx) b = tvm.nd.array(np.random.uniform(size=n).astype(B.dtype), ctx) c = tvm.nd.array(np.zeros(n, dtype=C.dtype), ctx)

fadd(a, b, c)

tvm.testing.assert_allclose(c.asnumpy(), a.asnumpy() + b.asnumpy())

if tgt == “cuda”: dev_module = fadd.imported_modules[0] print(“——-GPU code——-“) print(dev_module.get_source()) else: print(fadd.get_source())

3. gdb与pdb配置

Pdb一般是随着python安装包进行安装的，如果使用pdb命令失败，可以重新安装python。
上边的测试程序有可能会失败，报缺失decorator的错误，这是一个python的wheel，需要手动安装。
先说pdb的调试，pdb调试与gdb使用方式类似，都是使用pdb xxx.py进行。这个时候，程序会自动运行到程序的第一行。之后使用命令进行，网上相关的文档非常多，不再进行赘述。
这里介绍下python pdb特有的一种调试方法，在源码中可以使用
import pdb
pdb.set_trace()
然后使用python xxx.py运行程序，程序会自动断在pdb.set_trace()那一行，从该行起开始调试，这里仅作为介绍。
因为tvm是一个使用python接口，但是大部分实现是使用的C++的开源包。存在很多python和C++的交互，因此调试tvm的过程需要python和C++的联合调试工具。就是python代码需要pdb，C++代码需要gdb。因此后边介绍gdb对C++代码的调试。
由于进入的时候是python代码，因此想要使用gdb下断点非常困难。我们需要使用gdb附加进程的方式进行。这个过程需要root的支持。但是目前使用的Ubuntu系统中没有root用户。使用以下命令来添加：
echo “0” | sudo tee /proc/sys/kernel/yama/ptrace_scope
为了方便调试，我在刚才的测试代码中加入了一个
import os
raw_input(os.getpid())
来方便进行调试。

4. 调试过程

使用命令运行程序
jourluohua@jour:~/work/python/tvm$ python gen_cuda.py
19143
19143是运行该python程序时的pid
在另一个窗口中使用gdb attach该pid劫持python程序的线程
jourluohua@jour:~$ gdb attach 19143
然后对想要进行下断点的函数或者行进行下断点
(gdb) b tvm::codegen::CodeGenCUDA::PrintType
Breakpoint 1 at 0x7f22ae4e9df0 (2 locations)
在gdb所在的窗口使用c命令使程序执行起来
(gdb) c
Continuing.
然后再python对应的窗口输入回车，继续执行，就会断到断点所在的位置。