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' 注冊

sudoapt-get update

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

總結

以上是生活随笔為你收集整理的TVM调试指南的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。