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對應的窗口輸入回車，繼續執行，就會斷到斷點所在的位置。