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