ubuntu 16.04下驱动TP-LINK TL-WDN6200H免驱版无线网卡

1 先编译构建驱动

1
2
3
4
git clone https://github.com/jurobystricky/Netgear-A6210
cd /usr/src/netgear-a6210-2.5.0/
make
sudo make install

DKMS Install

On Debian-based distros, you can add the module to DKMS so it will automatically
build and install on each successive kernel upgrade. To do this, issue the following
commands from within the repo’s folder:

1
2
3
4
5
6
7
    $ cd ..
$ sudo mv Netgear-A6210/ /usr/src/netgear-a6210-2.5.0
$ sudo dkms install netgear-a6210/2.5.0

To remove:

$ sudo dkms remove netgear-a6210/2.5.0 --all

2 将无线网卡储存区中的SetupInstall.exe拷贝出来(如果没有出现无线网卡的usb储存,需要重新插拔一下无线网卡)

注意一定要把这个文件拷贝出来,不能直接再储存区中运行。运行这个程序后,储存区可能会消失掉。

3 每次启动后

1
2
3
wine SetupInstall.exe,进行初始化
sudo modprobe mt7662u_sta 或者 sudo insmod /lib/modules/4.4.0-96-generic/kernel/drivers/net/wireless/mt7662u_sta.ko
sudo service network-manager restart

如果还不行,可能需要重新插拔无线网卡,再重新走一次上述流程。

4 ubuntu 18.04 补充

18.04上原来的代码无法编译了,可以从https://github.com/kaduke/Netgear-A6210/tree/port-to-4.15下载适配高版本内核的驱动。
18.04上使用wine SetupInstall.exe无法进行初始化了,原因未明。这样就是加载了网卡驱动也无法找到无线网卡。
无奈转而直接分析wine SetupInstall.exe到底做了什么事情。分析了dmesg,lsusb的信息后,注意到wine执行之后,无线网卡对应的设备发生了变化。
lsusb初始化前id是2870,初始化后是7612。dmesg显示wine初始化时,出现了一次旧设备的的disconnect和新设备的加入。
综合各种情况来看,wine初始化时,是将无线网卡中对应的usb储存器断开了,然后将网卡设备挂上来了(这正好能解释第二节的现象)。
剩下的事情,就是看看不通过wine SetupInstall.exe来复现这一系列动作。尝试了各种方案后,终于在一个老外的论坛上看到,只需要eject usb storage,新的网卡设备就会出现。
马上尝试了一下,果然可以了。完全没有想到就是这么简单。。。
理顺以后整个思路就非常简单了。先手工弹出网卡带的储存盘,获得网卡设备,然后装入网卡驱动就可以上网了。

tvm 构建结果调试

背景

使用tvm的irbuilder直接构建了一小段计算程序,运行时有段错误。
tvm对内存边界的检查在较为上层的流程中,直接构建ir无法自动做检查。

尝试的方法

添加调试信息

把target设置为llvm,在tvm build后使用print(build_f.get_source()),可以打印出程序的llvm-ir。
但是,由于tvm是直接发射的llvm-ir,没有上层的文本,所以也没有发射有意义的调试信息。

将LLVM-IR反向翻译回C/C++

由于tvm可以给出llvm-ir的文本形式,如果能将llvm-ir翻译会c/c++代码,那么调试错误(甚至使用sanitizer系列工具)会简单很多。

llvm曾经有一个c/c++ 后端,可以将llvm-ir翻译为c/c++代码。但是后面由于维护问题,主线已经将该功能移除。
经过搜索,发现有一个第三方的项目https://github.com/JuliaComputing/llvm-cbe,可以与llvm-8一起工作。

实际测试发现,使用clang 编译出的hello world llvm-ir能正常翻译会C,但是tvm给出的ir会导致cbe工具出现assert错误。
简单看了一下,可能是tvm用的llvm-ir特性较新,cbe的llvm8还不认识。

从搜索过程来看,cbe的支持和需求都不是很强烈,于是不再考虑投入时间来分析和解决其问题。

使用TVM 的C backend

TVM使用cuda时,其输出的就是文本形式c代码交给cuda编译器。
这样看来,其支持一个c代码的backend应该是顺理成章的事情。

浏览tvm代码,果然其已经存在c backend了。参考https://tvm.apache.org/docs/dev/relay_bring_your_own_codegen.html ,生成c 代码除了帮助调试外,还比较容易与定制的优化库交互,甚至也可以把它当成代码模板,再进行人工修改。

使用c backend也比较简单,只需在tvm.build是传入target=’c’即可。

1
2
build_f  = tvm.build(ls, [], target='c',  name='prune_conv', binds=None)
print(build_f.get_source())

实际使用时,发现其c backend还不支持分配作用域为local的storeage。

1
2
3
4
5
6
void CodeGenC::PrintStorageSync(const CallNode* op) {  // NOLINT(*)
}

void CodeGenC::PrintStorageScope(const std::string& scope, std::ostream& os) { // NOLINT(*)
CHECK_EQ(scope, "global");
}

临时将程序内的 irb.allocate 调用中的 scope=’local’ 改为’global’,可以正常的输出c代码了,如下所示。

1
2
3
4
5
6
7
#include "tvm/runtime/c_runtime_api.h"
#include "tvm/runtime/c_backend_api.h"
void* __tvm_module_ctx = NULL;
#ifdef __cplusplus
extern "C"
#endif
TVM_DLL int32_t main(void* args, void* arg_type_ids, int32_t num_args, void* out_ret_value, void* out_ret_tcode) {

作为调试,这里打印出来已经能满足要求了。
如果需要后续的定制开发,可以参考tvm自带的示例apps/howto_deploy/cpp_deploy.cc 。把tvm生成的代码作为与自己的程序进一步组合。

利用反向调试,直接分析生成的汇编

gdb支持反向调试,虽然其功能不太完善稳定,但是比较适合tvm生成的这类逻辑较为简单的场景。

1
2
3
4
5
6
7
8
b thread_pool.cc中的launch函数
c
停下后,set scheduler-locking on
si进入tvm生成的运算函数
record full
si走到故障处
然后就可以用reverse-stepi等进行反向调试
退出前可以record stop

注意record full是,一定要set scheduler-locking on。
因为full模式需要获取进程的所有信息,而当前gdb还没支持好多线程程序的序列执行功能。不锁住当前线程执行,会导致gdb报下面的assert。

1
2
../../gdb/nat/x86-linux-dregs.c:146: internal-error: void x86_linux_update_debug_registers(lwp_info*): Assertion `lwp_is_stopped (lwp)' failed.
A problem internal to GDB has been detected,

对于简单的程序,知道
可以在python端打印出tvm.nd.array的_tvm_handle。
在gdb中x _tvm_handle的数值,可以找到array对应的数据区地址。
结合反汇编,也可以进行简单的分析。

TVM C backend生成代码直接与python的runtime整合

前面提到可以使用tvm的c backend将tvm的结果输出为c文件,但是对于分析问题来说,还是能直接运行起来更为方便。
最为直接的两个需求就是:使用asan查找内存越界故障, 使用vtune/perf等查找性能瓶颈。
前期本来打算使用c++ runtime来调用c backend的代码,但是工作量稍大,和python这端的配合也比较麻烦。
经过调试分析,发现可以直接将c backend生成的代码放到python运行时中运行。
大致方法如下:

生成c 文件

使用前面介绍的方法,去除local指定后,可以在target=c的情况下,生成出c代码。如下实例命令。

1
2
3
4
build_f = tvm.build(ls, [Input, Offset, Reorder, Index, Stride, Weight, Output],  
target='c', name='prune_conv', binds=None)
with open( './out_csrc.c', 'w' ) as f:
f.write(build_f.get_source())

修改获得的c代码

tvm的runtime装载模块时,希望看到一个名为tvm_main的符号,并且该符号中应该以字符串形式存放模块的实际入口。
所以,需要稍微修改一下前面获得的c代码。将其头部稍微修改一下,将main改为main_t(clang将main视为特殊符号,参数如果和c标准不一致会拒绝编译),并且添加tvm_main符号。

1
2
const char __tvm_main__[] = "main_t";
TVM_DLL int32_t main_t(void* args, void* arg_type_ids, int32_t num_args, void* out_ret_value, void* out_ret_tcode) {

使用clang或者gcc将获得的c文件编译为so

然后就可以使用clang将文件编译为so文件了,此时可以加上调试信息和asan等功能

1
clang ./out_csrc.c  -I ../incubator-tvm/include/ -I ../incubator-tvm/3rdparty/dlpack/include/  -fsanitize=address -save-temps -fPIC -O0 -g3 -shared

在python中运行获得的so

参考面的示例代码,可以将前面获得的so装入python中运行。

1
2
3
4
mod_prune = '/mnt/d/opensource/tvm_files/out_csrc.so'
loaded_prune = tvm.runtime.load_module(mod_prune)
evaluator = loaded_prune.time_evaluator(loaded_prune.entry_name, ctx, number=10)
print('OPt_paper: %f' % evaluator(input_np, Offset_np, Reorder_np, Index_np, Stride_np, Weight_np, output_array).mean)

初步结论

可以考虑优先使用tvm 的c backend进行调试分析。
结合反向调试的直接汇编分析也有帮助。

arduino开发体验

背景

了解一下物联网的现状和开发过程。

甲醛浓度传感器制作

硬件准备

为了简单便宜,考虑使用成熟而廉价的arduino, uno版本30块左右。
但是考虑到后续接入网络的方便,选择商家的套装,包含了各种传感器和传输辅助部件。
甲醛浓度传感器为了简单,选择可直接从串口上报数据的商品。

硬件连接也比较简单,传感器商家会给接线图,照着接就可以了,如下图。

注意传感器的Tx接到arduno的Rx,Rx接到Tx就可以了。
另外,如果物联网板子只有一个串口并且串口已经被用来做与pc机器之间的连接(例如我的这个uno),不要将传感器的串口直接对接到板子的串口上(可以使用其他通用针脚,然后用软串口机制来读取数据)。
否则会因为相互干扰出现很多莫名其妙的问题。例如程序烧写失败,传感器读不出数据等等。

软件

arduino软件开发环境很成熟,可以直接选择最方便的Web方式进行开发。
https://create.arduino.cc/ 申请一个免费的账户,再到https://create.arduino.cc/getting-started/plugin/install 去下载连接uno的浏览器插件并安装,
然后启动插件(其实就是一个go语言编写的转发代理,安装后在桌面上有一个Arduino Create Agent 的图标可供快速启动)。
随后就可以进入https://create.arduino.cc/editor/ 的WebIde界面。
在这个主界面下,可以完成整套软件开发的过程,包括代码编写、编译、烧写,也可以查看串口的输出(和进行串口输入),如下图所示。

整个Ide界面比较直观,Example和Library中已经有大量代码,许多应用场景都有可以借鉴的示例。唯一个稍微有点隐晦的是,串口输入输出是在侧边菜单的Monitor中。

按照传感器厂商给出的说明,传感器的输出是8个byte的数据+1个byte的校验,可以使用如下的示例程序将传感器周期性上报的数据打印出来。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
#include <SoftwareSerial.h>

SoftwareSerial mySerial(2, 4); // RX, TX (Rx和Tx针脚注意与接线图中的连线对应)

void setup()
{
Serial.begin(9600);
while (!Serial) {
}
Serial.println("Goodnight moon!");

mySerial.begin(9600);
}

void loop() {
int len = 9;
String result;
while (len)
{
if(mySerial.available())
{
String stringOne = String(mySerial.read(), HEX);
Serial.println(stringOne);
result += stringOne;
--len;
}
delay(1);
}
Serial.println(result.length());

for (int i=0;i<9;i++)
Serial.print(result[i]);
Serial.println(" Hello Arduino ");
delay(1000);

}

剩下的事情就是将收集到的数据上报到云端了,可参考https://create.arduino.cc/projecthub/133030/iot-cloud-getting-started-c93255 中的说明。这个步骤各个云供应商不完全一样,但是总体流程都差不多。大致都是需要将单板联网,然后将数据通过指定格式上报到云端,云端通过定义好的数据格式解析并展示。

总体感受

硬件

IOT的硬件成本已经比较低了。50块左右就能买到直接通过2g/4g网络连接到云端的物联网模块。通常IOT应用的数据量不大,数据流量现在也很便宜,所以持续维护的成本也不高。

软件

整个软件栈已经比较完整成熟了。大部分复杂的功能逻辑已经被很好地封装到厂商提供的API内了,并且从单板到云端都有大量示例可控借鉴。
甚至连商业模式都比较成熟了,如下图所示。arduino的收费模式设计得非常平衡,确保个人实验时可以访问绝大部分功能,感受到方便好用而愿意付费。而商业使用时又不太可能免费占便宜。

个人历程感受

纯软件开发做IOT应用,最困难和耗时的还是和硬件相关的琐碎部分。
本次体验中,分析和解决硬件串口冲突导致的各种问题(烧写失败,传感器无法输出)耗费了最主要的时间。其余的所有步骤总共加起来也没有耗费多少时间。

phoronix 扫描2020-05

硬件

amd 桌面cpu性能重新占据优势。同级比较,优势明显。
https://www.phoronix.com/scan.php?page=article&item=amd-ryzen-313&num=8

amd核显吊打intel。。。Ryzen 7 4700U was coming out about 39% faster than the Core i7 1065G7 with this given set of tests
https://www.phoronix.com/scan.php?page=article&item=amd-renoir-icelake&num=5

亚马逊的graviton2 arm64 cpu 性能有提升。在云上用物理core对抗intel和amd的vcpu已经有性能优势。但是物理机的裸核心还是差了很多,比8核心16线程的EPYC 7F32还要略慢(graviton2有64个物理核心)。
https://www.phoronix.com/scan.php?page=article&item=amazon-graviton2-benchmarks&num=12

编译器、开发库

在intel corei7 5960X cpu 上,gcc5到gcc10的性能几乎没有变化(性能提升2%以下),
这说明编译器的通用优化技术最近几年几乎没有进步。
这对编译器从业者来说是一个很悲哀的结论。
An Intel Core i7 5960X Haswell-E system was used for testing rather than a newer CPU in order to rule out back-end/micro-architecture specific optimizations across the tested compilers. Intel Haswell has offered tuned GCC support since before the GCC 5 release. Ubuntu 19.10 was running on this Core i7 5960X system with the Linux 5.3 kernel.
https://www.phoronix.com/scan.php?page=article&item=gcc5-gcc10-benchmarks&num=4

clang9已经和gcc性能持平了,但构建速度反而是gcc更快了。以前的讽刺成真,clang编译器快是因为优化没有做够。。。
https://www.phoronix.com/scan.php?page=article&item=gcc-clang-3960x&num=7
https://www.phoronix.com/scan.php?page=article&item=gcc10-clang10-x86&num=5

GraalVM 和openjdk的性能差异看起来不大,这很不错,为多语言融合奠定了性能基础。
https://www.phoronix.com/scan.php?page=article&item=openjdk-corretto-graalvm&num=5

浏览器

chrome的性能还是明显比firefox更好,但是firefox最近性能在逐渐靠近chrome。
https://www.phoronix.com/scan.php?page=article&item=firefox-chrome-icelake&num=7
https://www.phoronix.com/scan.php?page=article&item=chrome-80-benchmarks&num=5

内核

x86允许直接操作FS、GS基址寄存器指令在linux内核得到了支持。
IO和redis的性能有明显提高。
https://www.phoronix.com/scan.php?page=article&item=linux-fsgsbase-2020&num=3

amd linux 开源显卡驱动似乎有明显进步,平均分数超过了厂商闭源驱动。
https://www.phoronix.com/scan.php?page=article&item=radeon-software-20&num=6

linux内核缓解幽灵等cpu漏洞带来的性能损失相当明显, E3-1280 v5 只有不缓解77%的性能,最新的Xeon Platinum 8280 Cascade Lake 可以到95%的性能。
https://www.phoronix.com/scan.php?page=article&item=spectre-meltdown-2&num=11

不编译apk直接在android上部署tvm编译结果

技术路线

理论上android就是一个linux内核加用户态库。因此tvm部署模型到android并不一定需要做一个apk,只需构建一个elf程序提供rpc的功能就可以了。
查看到https://github.com/apache/incubator-tvm/pull/4281,tvm主线已经添加了c++版本的rpc实现。
因此,可以使用android ndk中的工具链编译tvm rpc c++实现,绕开复杂且不必要的android apk构建(当然有一个可能的问题是,由于android的权限管控,编译好的程序在非root情况下可能无法启动。本次是在android7.1上/data/local/tmp目录可以用于执行)。

编译构建

根据手机型号(坚果pro)和android版本7.1.1,下载android-ndk-r21并选择aarch64-linux-android24-clang++作为交叉编译器。

tvm创建build目录并编辑config.make

1
2
3
4
在tvm目录下
mkdir build_arm64
cd build_arm64
cp ../config.make ./

为了支持rpc和gpu运算,编辑config.make确保下面两项正确。其中vulkan的设置目录dep_dirs会在后面的步骤中配置好。

1
2
3
# Whether enable RPC runtime
set(USE_RPC ON)
set(USE_VULKAN /home/majiang/hd/opensource/tvm/build_arm64/dep_dirs)

构建交叉版本的spirv-tools

1
2
3
4
5
6
7
8
9
10
11
git clone https://github.com/KhronosGroup/SPIRV-Tools.git
cd SPIRV-Tools/
git clone https://github.com/KhronosGroup/SPIRV-Headers.git external/spirv-headers
mkdir build
cd build
cmake .. -DCMAKE_CXX_COMPILER="/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++" -DCMAKE_C_COMPILER="/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang"
make -j 8
mkdir inst
make install DESTDIR=`pwd`/inst
#本来使用make install-headers 应该更为标准,但是spirv-headers的makefile没有写好,其忽略了DESTDIR变量,直接把头文件拷贝到了/usr/local下。规避方案直接copy
cp ../external/spirv-headers/include/* inst/usr/local/include/ -r

构建交叉版本的runtime

进入build_arm64目录
mkdir dep_dirs
cd dep_dirs/
mkdir include
cp ~/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/sysroot/usr/include/vulkan ./include -r
cp ~/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/sysroot/usr/lib/aarch64-linux-android/24 ./lib -r
cp ~/hd/opensource/android_sdk/spirv-tools/SPIRV-Tools/build/inst/usr/local/include/spirv-tools/ ./include/ -r
cp ~/hd/opensource/android_sdk/spirv-tools/SPIRV-Tools/build/inst/usr/local/lib/* ./lib/
#(可选的strip -g)
cp ~/hd/opensource/android_sdk/spirv-tools/SPIRV-Tools/external/spirv-headers/include/spirv/ ./include/ -r
cd ../
cmake .. -DCMAKE_CXX_COMPILER=”/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++” -DCMAKE_C_COMPILER=”/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang”
(修改config.cmake 将修改 USE_VULKAN指向 dep_dirs set(USE_VULKAN /.xxx…/tvm/build_arm64/dep_dirs))
make runtime -j8

构建cpp版本的rpc服务程序

在build_arm64目录下执行如下命令。

1
make -C ../apps/cpp_rpc CXX=/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++  TVM_RUNTIME_DIR=/home/majiang/hd/opensource/tvm/build_arm64/

tvm当前的makefile会把所有的cc都加进去编译(cmake文件不会),其中包括windows的win32_process.cc。为了阻止编译错误,手动将其改名为win32_process.cc-nouse。
有可能因为搜索路径的问题,找不到vulkan库,可以使用如下命令手动链接(添加-Wl,-rpath-link到对应api的lib目录;-Wl,-rpath-link=/home/majiang/hd/opensource/android_sdk/android-ndk-r21/platforms/android-24/arch-arm64/usr/lib/)。

1
/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++ -std=c++14 -O2 -fPIC -Wall -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/include -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/3rdparty/dmlc-core/include -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/3rdparty/dlpack/include -o tvm_rpc main.cc rpc_env.cc rpc_server.cc -L/home/majiang/hd/opensource/tvm/build_arm64_new/  -ltvm_runtime -ldl -Wl,-R/home/majiang/hd/opensource/tvm/build_arm64_new/ -Wl,-rpath-link=/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/lib/aarch64-linux-android/24/

更新本地的tvm使其支持opencl和vulkan代码生成

修改本地tvm build目录下的config.cmake,确保USE_OPENCL/VULKAN是ON状态。
并且确保装好了opencl和vulkan 的sdk(可参考https://www.codenong.com/cs105410317/,直接去https://vulkan.lunarg.com/sdk/home 下载vulkan sdk)。

配置手机

打开USB调试

在手机设置的“关于本机”页面中连续点击 “软件版本” 条目,可以打开开发者模式。然后在全局高级设置中会出现 “开发者选项”,进入其条目打开“USB调试”即可。

安装adb

apt install adb -y

上传文件并设置权限

android高版本在没有root的情况下,不能直接给sd卡中的程序加上可执行权限,参考https://my.oschina.net/jerikc/blog/497090 ,可以拷贝到/data/local/tmp 的特殊路径下,并添加执行权限。
在tvm的目录下将两个必须的文件上传。

1
2
adb push apps/cpp_rpc/tvm_rpc  /data/local/tmp
adb push build_arm64_new/libtvm_runtime.so /data/local/tmp

另外,由于使用了c++,还需上传 libc++_shared.so(参考https://developer.android.com/ndk/guides/cpp-support#libc)。

1
adb push /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/lib/aarch64-linux-android/libc++_shared.so  /data/local/tmp

然后进入手机,为程序设置可执行权限。

1
2
3
adb shell
cd /data/local/tmp
chmod 777 *

测试rpc程序

使用下面的命令,初步测试程序能否正常启动。

1
2
3
4
adb shell
cd /data/local/tmp
export LD_LIBRARY_PATH=`pwd`
./tvm_rpc

如果正常,应该能看到help信息,示例如下。

1
2
3
4
5
6
7
8
9
[10:17:53] main.cc:289: Command line usage
server - Start the server
--host - The hostname of the server, Default=0.0.0.0
--port - The port of the RPC, Default=9090
--port-end - The end search port of the RPC, Default=9199
--tracker - The RPC tracker address in host:port format e.g. 10.1.1.2:9190 Default=""
--key - The key used to identify the device type in tracker. Default=""
--custom-addr - Custom IP Address to Report to RPC Tracker. Default=""
--silent - Whether to run in silent mode. Default=False

启动rpc服务,进行测试

启动cpp版本的rpc后,测试其功能是否正常。
首先在host主机上启动rpc tracker。使用如下命令。
应该会看到”INFO:RPCTracker:bind to 0.0.0.0:9190”这样的提示。

1
2
3
export TVM_HOME=/home/majiang/hd/opensource/tvm
export PYTHONPATH=$TVM_HOME/python:$TVM_HOME/topi/python:${PYTHONPATH}
python3 -m tvm.exec.rpc_tracker

然后在手机上启动cpp 版本的rpc server。注意tracker选项中的ip地址是电脑主机的ip,不是手机的ip,9190是前面启动tracker给出的port。–key一定写成android,否则后面的android_rpc_test.py会找不到设备(它写死了设备的key为android)。

1
2
3
4
adb shell
cd /data/local/tmp
export LD_LIBRARY_PATH=`pwd`
./tvm_rpc server --tracker=192.168.3.4:9190 --key=android

此时,在电脑主机上可以查询到手机了。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
 python3 -m tvm.exec.query_rpc_tracker
Tracker address 192.168.3.4:9190

Server List
----------------------------
server-address key
----------------------------
192.168.3.33:38151 server:android
----------------------------

Queue Status
-------------------------------
key total free pending
-------------------------------
android 1 1 0
-------------------------------

最后,进入tvm/apps/android_rpc目录,启动android_rpc的测试。

1
2
3
4
5
6
export TVM_HOME=/home/majiang/hd/opensource/tvm
export PYTHONPATH=$TVM_HOME/python:$TVM_HOME/topi/python:${PYTHONPATH}
export TVM_TRACKER_HOST=192.168.3.4
export TVM_TRACKER_PORT=9190
export TVM_NDK_CC=/home/majiang//hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++
python3 tests/android_rpc_test.py

正常时可以看到如下的提示(打开GPU测试需要修改android_rpc_test.py,设置test_vulkan = True)。

1
2
3
4
5
Run CPU test ...
0.000340646 secs/op

Run GPU(Vulkan Flavor) test ...
4.40886e-05 secs/op

android端json解析报错问题分析

使用简单的rpc测试正常,但是使用apps/benchmark/mobile_gpu_imagenet_bench.py ( python3 ./mobile_gpu_imagenet_bench.py –model rk3399 –network mobilenet –rpc-key android)等复杂测试,会出现手机端报json格式错误。具体的表现是运行到runtime.create时,手机端的runtime解析json格式assert报错,形式不固定(JSONReader::BeginObject等期望的字符没有读到)。

使用gdb进行初步调试定位

在没有任何背景信息的情况下,可以先行利用gdb继续初步查看。
android-ndk提供了arm64版本的gdb-server和x86版本的gdb。

1
2
在ndk目录下
adb push ./prebuilt/android-arm64/gdbserver/gdbserver /data/local/tmp

然后。
android端
cd /data/local/tmp
./gdbserver 192.168.3.33:8888(手机ip和希望使用的端口) –attach 2126 (使用ps |grep tvm_rpc看到的rpc进程)

host端
./android-ndk-r21/prebuilt/linux-x86_64/bin/gdb ../tvm/apps/cpp_rpc/tvm_rpc

使用asan排查可疑内存问题

编译

需要将libtvm_runtime.so和tvm_rpc都加上asan重新编译(注意如果不重新编译tvm_rpc,asan的检查可能无法准确输出信息)。
对于前者,需要修改config.make,在其尾部加上如下语句,然后重新cmake一次(如果是干净的环境下cmake,可能会出现找不到pthread.h的错误。去掉下面语句成功cmake一次,再重新加上后cmake一次,就可以了,原因暂未调查。)。

1
2
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsanitize=address")

对于后者,可以在make的时候添加CXXFLAGS,也可以直接手动加编译参数(因为编译tvm_rpc只需要单条命令)。

1
/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-clang++ -std=c++14 -O2 -fPIC -Wall -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/include -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/3rdparty/dmlc-core/include -I/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/tvm/3rdparty/dlpack/include -o tvm_rpc main.cc rpc_env.cc rpc_server.cc -L/home/majiang/hd/opensource/tvm/build_arm64/  -ltvm_runtime -ldl -Wl,-R/home/majiang/hd/opensource/tvm/build_arm64/ -Wl,-rpath-link=/home/majiang/hd/opensource/android_sdk/android-ndk-r21/platforms/android-24/arch-arm64/usr/lib/ -fsanitize=address

运行

使用adb push将新的rumtime和tvm_rpc上传。
然后使用下面命令上传asan需要的动态库。

1
adb push /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/lib64/clang/9.0.8/lib/linux/libclang_rt.asan-aarch64-android.so /data/local/tmp

然后使用正常方式启动tvm_rpc即可。

再次运行触发json解析错误的测试用例,这次asan给出了准确的输出,确实有堆内存越界。在启动tvm_rpc的终端上可以看到如下输出。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
==5713==ERROR: AddressSanitizer: heap-buffer-overflow on address 0x0076a6ec7100 at pc 0x007fa7f66f50 bp 0x005fa41fc130 sp 0x005fa41fb8d8
WRITE of size 20942 at 0x0076a6ec7100 thread T1
#0 0x7fa7f66f4c (/data/local/tmp/libclang_rt.asan-aarch64-android.so+0x85f4c)
#1 0x7fa7f66c64 (/data/local/tmp/libclang_rt.asan-aarch64-android.so+0x85c64)
#2 0x7fa7757a68 (/data/local/tmp/libtvm_runtime.so+0x55ba68)
#3 0x7fa7757458 (/data/local/tmp/libtvm_runtime.so+0x55b458)
#4 0x7fa770a0cc (/data/local/tmp/libtvm_runtime.so+0x50e0cc)
#5 0x7fa76faf00 (/data/local/tmp/libtvm_runtime.so+0x4fef00)
#6 0x7fa76f9ce0 (/data/local/tmp/libtvm_runtime.so+0x4fdce0)
#7 0x7fa76fcb3c (/data/local/tmp/libtvm_runtime.so+0x500b3c)
#8 0x7fa774d3c8 (/data/local/tmp/libtvm_runtime.so+0x5513c8)
#9 0x555558b444 (/data/local/tmp/tvm_rpc+0x36444)
#10 0x5555583db0 (/data/local/tmp/tvm_rpc+0x2edb0)
#11 0x5555585180 (/data/local/tmp/tvm_rpc+0x30180)
#12 0x5555585354 (/data/local/tmp/tvm_rpc+0x30354)
#13 0x7fa7e6a41c (/system/lib64/libc.so+0x6841c)
#14 0x7fa7e1fe00 (/system/lib64/libc.so+0x1de00)

0x0076a6ec7100 is located 0 bytes to the right of 4096-byte region [0x0076a6ec6100,0x0076a6ec7100)
allocated by thread T1 here:
#0 0x7fa7f8b374 (/data/local/tmp/libclang_rt.asan-aarch64-android.so+0xaa374)
#1 0x7fa73a3460 (/data/local/tmp/libtvm_runtime.so+0x1a7460)
#2 0x7fa73a3438 (/data/local/tmp/libtvm_runtime.so+0x1a7438)
#3 0x7fa73a2bf0 (/data/local/tmp/libtvm_runtime.so+0x1a6bf0)
#4 0x7fa7683008 (/data/local/tmp/libtvm_runtime.so+0x487008)
#5 0x7fa7682240 (/data/local/tmp/libtvm_runtime.so+0x486240)
#6 0x7fa7681448 (/data/local/tmp/libtvm_runtime.so+0x485448)
#7 0x7fa76fab4c (/data/local/tmp/libtvm_runtime.so+0x4feb4c)
#8 0x7fa76f9ce0 (/data/local/tmp/libtvm_runtime.so+0x4fdce0)
#9 0x7fa76fcb3c (/data/local/tmp/libtvm_runtime.so+0x500b3c)
#10 0x7fa774d3c8 (/data/local/tmp/libtvm_runtime.so+0x5513c8)
#11 0x555558b444 (/data/local/tmp/tvm_rpc+0x36444)
#12 0x5555583db0 (/data/local/tmp/tvm_rpc+0x2edb0)
#13 0x5555585180 (/data/local/tmp/tvm_rpc+0x30180)
#14 0x5555585354 (/data/local/tmp/tvm_rpc+0x30354)
#15 0x7fa7e6a41c (/system/lib64/libc.so+0x6841c)
#16 0x7fa7e1fe00 (/system/lib64/libc.so+0x1de00)

Thread T1 created by T0 here:
#0 0x7fa7f725a0 (/data/local/tmp/libclang_rt.asan-aarch64-android.so+0x915a0)
#1 0x5555584ebc (/data/local/tmp/tvm_rpc+0x2febc)
#2 0x5555584c20 (/data/local/tmp/tvm_rpc+0x2fc20)
#3 0x5555582608 (/data/local/tmp/tvm_rpc+0x2d608)
#4 0x5555581224 (/data/local/tmp/tvm_rpc+0x2c224)
#5 0x555556cdd4 (/data/local/tmp/tvm_rpc+0x17dd4)
#6 0x555556d8ac (/data/local/tmp/tvm_rpc+0x188ac)
#7 0x7fa7e1c7d8 (/system/lib64/libc.so+0x1a7d8)
#8 0x5555566b10 (/data/local/tmp/tvm_rpc+0x11b10)
#9 0x7fa8460d54 (/system/bin/linker64+0x6d54)

SUMMARY: AddressSanitizer: heap-buffer-overflow (/data/local/tmp/libclang_rt.asan-aarch64-android.so+0x85f4c)
Shadow bytes around the buggy address:
0x001ed4dd8dd0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
0x001ed4dd8de0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
0x001ed4dd8df0: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
0x001ed4dd8e00: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
0x001ed4dd8e10: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
=>0x001ed4dd8e20:[fa]fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
0x001ed4dd8e30: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
0x001ed4dd8e40: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
0x001ed4dd8e50: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
0x001ed4dd8e60: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
0x001ed4dd8e70: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
Shadow byte legend (one shadow byte represents 8 application bytes):
Addressable: 00
Partially addressable: 01 02 03 04 05 06 07
Heap left redzone: fa
Freed heap region: fd
Stack left redzone: f1
Stack mid redzone: f2
Stack right redzone: f3
Stack after return: f5
Stack use after scope: f8
Global redzone: f9
Global init order: f6
Poisoned by user: f7
Container overflow: fc
Array cookie: ac
Intra object redzone: bb
ASan internal: fe
Left alloca redzone: ca
Right alloca redzone: cb
Shadow gap: cc
==5713==ABORTING

由于是交叉编译运行,asan无法直接给出了文件和行号信息。
为了便于理解日志,可以将上面的日志信息复制并保存到host机器的文件中,再使用asan提供的专用symbolize工具获得文件和行号。

日志解析

为了解析交叉编译的日志,asan提供了专门的symbolizer工具。
该工具的核心任务就是将 “binary文件+offset” 翻译为 “文件:函数:行号”。
llvm的asan提供的工具在, https://llvm.org/svn/llvm-project/compiler-rt/trunk/lib/asan/scripts/asan_symbolize.py。它对翻译工作做了封装,可以使用llvm-symbolizer/address2line等多种底层工具完成翻译。

其使用也非常简单,如下一条命令即可完成翻译。
error.log是含有错误信息的文件。
-c 是交叉编译的prefix
-s 是sysroot,该路径下需要含有带调试信息的binary文件

1
./asan_symbolize.py  -d -c "/home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android24-"  -s "/home/majiang/hd/opensource/tvm/debug_arm64/" < error.log

这个脚本错误信息不是非常友好,需要先行确认底层的symbolizer能正常工作,sysroot中含有正确的binary文件。
例如,如果系统PATH中没有llvm-symbolizer,只有llvm-symbolizer-9(没有安装默认的llvm版本,而是安装了新的9版本),需要先export ASAN_SYMBOLIZER_PATH=llvm-symbolizer-9,再运行脚本。
sysroot中的路径布置必须与log中的日志完全一致,否则脚本报错信息也比较难以理解。

故障修复

使用asan,可以获得如下的故障日志。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
==5713==ERROR: AddressSanitizer: heap-buffer-overflow on address 0x0076a6ec7100 at pc 0x007fa7f66f50 bp 0x005fa41fc130 sp 0x005fa41fb8d8
WRITE of size 20942 at 0x0076a6ec7100 thread T1
#0 0x7fa7f66f4c in recvfrom /toolchain/llvm-project/compiler-rt/lib/asan/../sanitizer_common/sanitizer_common_interceptors.inc:6404:5
#1 0x7fa7f66c64 in recv /toolchain/llvm-project/compiler-rt/lib/asan/../sanitizer_common/sanitizer_common_interceptors.inc:6385:17
#2 0x7fa7757a68 in tvm::support::TCPSocket::Recv(void*, unsigned long, int) /home/majiang/hd/opensource/tvm/src/runtime/rpc/../../support/socket.h:483:12
#3 0x7fa7757458 in tvm::runtime::SockChannel::Recv(void*, unsigned long) /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_socket_impl.cc:53:23
#4 0x7fa770a0cc in tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*)::$_1::operator()(void*, unsigned long) const /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_session.cc:880:28
#5 0x7fa76faf00 in unsigned long tvm::support::RingBuffer::WriteWithCallback<tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*)::$_1>(tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*)::$_1, unsigned long) /home/majiang/hd/opensource/tvm/src/runtime/rpc/../../support/ring_buffer.h:160:25
#6 0x7fa76f9ce0 in tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*) /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_session.cc:879:26
#7 0x7fa76fcb3c in tvm::runtime::RPCSession::ServerLoop() /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_session.cc:952:3
#8 0x7fa774d3c8 in tvm::runtime::RPCServerLoop(int) /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_socket_impl.cc:113:30
#9 0x555558b444 in tvm::runtime::RPCServer::ServerLoopProc(tvm::support::TCPSocket, tvm::support::SockAddr) ??:0:0
#10 0x5555583db0 in tvm::runtime::RPCServer::ListenLoopProc() ??:0:0
#11 0x5555585180 in std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::__execute() ??:0:0
#12 0x5555585354 in void* std::__ndk1::__thread_proxy<std::__ndk1::tuple<std::__ndk1::unique_ptr<std::__ndk1::__thread_struct, std::__ndk1::default_delete<std::__ndk1::__thread_struct> >, void (std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::*)(), std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >*> >(void*) ??:0:0
#13 0x7fa7e6a41c in __pthread_start(void*) ??:0:0
#14 0x7fa7e1fe00 in __start_thread ??:0:0

0x0076a6ec7100 is located 0 bytes to the right of 4096-byte region [0x0076a6ec6100,0x0076a6ec7100)
allocated by thread T1 here:
#0 0x7fa7f8b374 in operator new(unsigned long) _asan_rtl_:3
#1 0x7fa73a3460 in std::__ndk1::__libcpp_allocate(unsigned long, unsigned long) /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/include/c++/v1/new:253:10
#2 0x7fa73a3438 in std::__ndk1::allocator<char>::allocate(unsigned long, void const*) /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/include/c++/v1/memory:1813:37
#3 0x7fa73a2bf0 in std::__ndk1::allocator_traits<std::__ndk1::allocator<char> >::allocate(std::__ndk1::allocator<char>&, unsigned long) /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/include/c++/v1/memory:1546:21
#4 0x7fa7683008 in std::__ndk1::__split_buffer<char, std::__ndk1::allocator<char>&>::__split_buffer(unsigned long, unsigned long, std::__ndk1::allocator<char>&) /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/include/c++/v1/__split_buffer:318:29
#5 0x7fa7682240 in std::__ndk1::vector<char, std::__ndk1::allocator<char> >::shrink_to_fit() /home/majiang/hd/opensource/android_sdk/android-ndk-r21/toolchains/llvm/prebuilt/linux-x86_64/bin/../sysroot/usr/include/c++/v1/vector:1598:57
#6 0x7fa7681448 in tvm::support::RingBuffer::Reserve(unsigned long) /home/majiang/hd/opensource/tvm/src/runtime/rpc/../../support/ring_buffer.h:74:15
#7 0x7fa76fab4c in unsigned long tvm::support::RingBuffer::WriteWithCallback<tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*)::$_1>(tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*)::$_1, unsigned long) /home/majiang/hd/opensource/tvm/src/runtime/rpc/../../support/ring_buffer.h:148:11
#8 0x7fa76f9ce0 in tvm::runtime::RPCSession::HandleUntilReturnEvent(tvm::runtime::TVMRetValue*, bool, tvm::runtime::PackedFunc const*) /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_session.cc:879:26
#9 0x7fa76fcb3c in tvm::runtime::RPCSession::ServerLoop() /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_session.cc:952:3
#10 0x7fa774d3c8 in tvm::runtime::RPCServerLoop(int) /home/majiang/hd/opensource/tvm/src/runtime/rpc/rpc_socket_impl.cc:113:30
#11 0x555558b444 in tvm::runtime::RPCServer::ServerLoopProc(tvm::support::TCPSocket, tvm::support::SockAddr) ??:0:0
#12 0x5555583db0 in tvm::runtime::RPCServer::ListenLoopProc() ??:0:0
#13 0x5555585180 in std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::__execute() ??:0:0
#14 0x5555585354 in void* std::__ndk1::__thread_proxy<std::__ndk1::tuple<std::__ndk1::unique_ptr<std::__ndk1::__thread_struct, std::__ndk1::default_delete<std::__ndk1::__thread_struct> >, void (std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::*)(), std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >*> >(void*) ??:0:0
#15 0x7fa7e6a41c in __pthread_start(void*) ??:0:0
#16 0x7fa7e1fe00 in __start_thread ??:0:0

Thread T1 created by T0 here:
#0 0x7fa7f725a0 in pthread_create _asan_rtl_:3
#1 0x5555584ebc in std::__ndk1::thread::thread<void (std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::*)(), std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >*, void>(void (std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >::*&&)(), std::__ndk1::__async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >*&&) ??:0:0
#2 0x5555584c20 in std::__ndk1::future<void> std::__ndk1::__make_async_assoc_state<void, std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*> >(std::__ndk1::__async_func<void (tvm::runtime::RPCServer::*)(), tvm::runtime::RPCServer*>&&) ??:0:0
#3 0x5555582608 in tvm::runtime::RPCServer::Start() ??:0:0
#4 0x5555581224 in tvm::runtime::RPCServerCreate(std::__ndk1::basic_string<char, std::__ndk1::char_traits<char>, std::__ndk1::allocator<char> >, int, int, std::__ndk1::basic_string<char, std::__ndk1::char_traits<char>, std::__ndk1::allocator<char> >, std::__ndk1::basic_string<char, std::__ndk1::char_traits<char>, std::__ndk1::allocator<char> >, std::__ndk1::basic_string<char, std::__ndk1::char_traits<char>, std::__ndk1::allocator<char> >, bool) ??:0:0
#5 0x555556cdd4 in RpcServer(int, char**) ??:0:0
#6 0x555556d8ac in main ??:0:0
#7 0x7fa7e1c7d8 in __libc_init ??:0:0
#8 0x5555566b10 in _start_main ??:0:0

从这个日志可以很快定位到核心的错误逻辑。
ring_buffer.h 中的Reserve函数实现逻辑有问题,没有为数据预留足够的空间,后续直接recv到buffer中的数据会导致overflow。
如下逻辑所示,当请求reserver的size n小于当前buffer的size
时,reserve函数会减小buffer保有的内存,节约资源。
但是其减小内存后,只保留kInitCapacity个byte,忽略了输入请求n可能大于kInitCapacity的情况。
发生错误时,输入的n为25038 (graph json的string长度)kInitCapacity 只有 4096。recv到的json字符串破坏了buffer,导致后续逻辑混乱。
修复的逻辑也简单,只需要保证收缩后的尺寸不小于n就可以了。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
void Reserve(size_t n) {
if (ring_.size() < n) {
//扩大ring buffer的size
}
} else if (ring_.size() > n * 8 && ring_.size() > kInitCapacity && bytes_available_ > 0) {
// shrink too large temporary buffer to avoid out of memory on some embedded devices
size_t old_bytes = bytes_available_;

std::vector<char> tmp(old_bytes);

Read(&tmp[0], old_bytes);
//ring_.resize(kInitCapacity); this may cause overflow when n>kInitCapacity
ring_.resize(kInitCapacity > n? kInitCapacity : n);
ring_.shrink_to_fit();

memcpy(&ring_[0], &tmp[0], old_bytes);
head_ptr_ = 0;
bytes_available_ = old_bytes;
}
}

查看错误输出

rpc时看不到错误信息。查看apps/android_camera/app/src/main/jni/tvm_runtime.h 可以发现,原因是android上需要特殊的打印指令,但是编译时我们没有打开对应的宏,也没有添加对应的打印函数,如下代码所示。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
/* Enable custom logging - this will cause TVM to pass every log message
* through CustomLogMessage instead of LogMessage. By enabling this, we must
* implement dmlc::CustomLogMessage::Log. We use this to pass TVM log
* messages to Android logcat.
*/
#define DMLC_LOG_CUSTOMIZE 1

/* Ensure that fatal errors are passed to the logger before throwing
* in LogMessageFatal
*/
#define DMLC_LOG_BEFORE_THROW 1



#include <android/log.h>

void dmlc::CustomLogMessage::Log(const std::string& msg) {
// This is called for every message logged by TVM.
// We pass the message to logcat.
__android_log_write(ANDROID_LOG_DEBUG, "TVM_RUNTIME", msg.c_str());
}

添加adreon opencl支持

从手机中pull出libOpenCL_system.so libion.so放入dep_libs_from_phone
从高通网站下载opencl sdk
同样在build_arm64中
cmake ./. -DCMAKE_CXX_COMPILER=”/mnt/d/opensource/android_ndk/android-ndk-r21b/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang++” -DCMAKE_C_COMPILER=”/mnt/d/opensource/android_ndk/android-ndk-r21b/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang” -DOpenCL_INCLUDE_DIR=/mnt/d/opensource/opencl-sdk-1.2.2/inc -DOpenCL_LIBRARY=/mnt/d/opensource/opencl-sdk-1.2.2/dep_libs_from_phone

make -j 32 runtime

make -C ../apps/cpp_rpc CXX=/mnt/d/opensource/android_ndk/android-ndk-r21b/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang++ TVM_RUNTIME_DIR=/mnt/d/opensource/opensrc_tvm/tvm/build_arm64

链接cpp_rpc仍然有错,手动添加 -lOpenCL_system -Wl,-rpath-link=xx 通过(需要把android 的system/lib64下的库全拉过来)
/mnt/d/opensource/android_ndk/android-ndk-r21b/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android28-clang++ -std=c++14 -O2 -fPIC -Wall -I/mnt/d/opensource/opensrc_tvm/tvm/include -I/mnt/d/opensource/opensrc_tvm/tvm/3rdparty/dmlc-core/include -I/mnt/d/opensource/opensrc_tvm/tvm/3rdparty/dlpack/include -o tvm_rpc main.cc rpc_env.cc rpc_server.cc -L/mnt/d/opensource/opensrc_tvm/tvm/build_arm64 -ltvm_runtime -ldl -Wl,-R/mnt/d/opensource/opensrc_tvm/tvm/build_arm64 -L /mnt/d/opensource/opencl-sdk-1.2.2/dep_libs_from_phone -lOpenCL_system -Wl,-rpath-link=/mnt/d/opensource/opencl-sdk-1.2.2/dep_libs_from_phone/lib64 -Wl,-v -v -Wl,-t

启动 apps/android_rpc//tests/android_rpc_test.py 段错误

同时支持opencl和vulkan后,启动android_rpc_test.py 出现段错误。
使用pdb启动,可以看到打印出的错误信息。

1
2
3
4
5
  File "/usr/lib/python3.6/ctypes/__init__.py", line 348, in __init__
self._handle = _dlopen(self._name, mode)
OSError: /home/majiang/opensrc/tvm/build/libtvm.so: undefined symbol: spvContextDestroy
Uncaught exception. Entering post mortem debugging
Running 'cont' or 'step' will restart the program

google了这个错误,找到了https://github.com/google/shaderc/issues/470,原因应该是少链接了一个库。
删除libtvm.so后,make VERBOSE=1拷贝出链接命令,在尾部添加-lSPIRV-Tools后重新链接。
链接完成后,故障消失。
怀疑与cmake版本有关系,缺失了库的依赖,暂不进一步分析。

tvm graph_runtime 分析

runtime总体逻辑

runtime总体逻辑是:读出编译好的运算图(包含了二进制代码和描述信息);根据运算图信息为各个存储节点分配储存;构建可执行OP的函数体(实际是调用已经编译好的代码);逐个执行可执行的OP。

代码逻辑

用户编译和运行深度学习模型的典型python代码片段如下所示

1
2
3
4
5
6
with relay.build_config(opt_level=3):
graph, lib, params = relay.build(net, target=target_n, params=params)

ctx = [cpu_ctx, gpu_ctx]
module = graph_runtime.create(graph, lib, ctx)
module.run()

在python端调用graph_runtime.create,会走到GraphRuntimeCreate,然后再到
GraphRuntime::Init创建runtime结构,并返回Module结构给python。python端通过module.run()方法来运行模型。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
Module GraphRuntimeCreate(const std::string& sym_json,
const tvm::runtime::Module& m,
const std::vector<TVMContext>& ctxs) {
auto exec = make_object<GraphRuntime>();
exec->Init(sym_json, m, ctxs);
return Module(exec);
}

void GraphRuntime::Init(const std::string& graph_json,
tvm::runtime::Module module,
const std::vector<TVMContext>& ctxs) {
std::istringstream is(graph_json);
dmlc::JSONReader reader(&is);
this->Load(&reader);
module_ = module;
ctxs_ = ctxs;
this->SetupStorage();
this->SetupOpExecs();
for (size_t i = 0; i < input_nodes_.size(); i++) {
const uint32_t nid = input_nodes_[i];
std::string& name = nodes_[nid].name;
input_map_[name] = i;
}
}

GraphRuntime::Init所做的主要工作包括两个部分,第一个是从json格式的string中读取出编译好的运算图(this->Load(&reader)),第二个是初始化运行环境(SetupStorage和SetupOpExecs)。
Load比较简单不展开。
SetupStorage的核心逻辑是从json读出需要存空间的各个矩阵信息,然后为其在对应的计算设备上分配内存(通过调用NDArray::Empty(shape, DLDataType{kDLFloat, 32, 1}, ctx)))。需要注意的点是,每一个设备上实际上只进行一次分配(分配最大所需的储存)。
SetupOpExecs的核心逻辑是把构建 OP函数体(实际功能前面已经编译好了,这里的函数体实际上只是去调用)和其所需要的参数args结构。

不支持OP级别的并行

当前tvm的graph_runtime就是一个简单的静态执行器。
比较典型的示例点就是下面的run函数。它的逻辑只是串行地逐个运行OP的函数体。

1
2
3
4
5
6
void GraphRuntime::Run() {
// setup the array and requirements.
for (size_t i = 0; i < op_execs_.size(); ++i) {
if (op_execs_[i]) op_execs_[i]();
}
}

并且,由于每一个OP都是同步执行的(也就是必须等待执行结果出来后,OP函数体才返回),所以runtime的顶层是不具备并行能力的。
理论上,tvm runtime当前不能支持cpu和gpu同时执行计算。(除非在一个OP植入异构执行代码,但是当前又没有构造对应OP的方法?)

多线程运行是由生成的函数来调用的。

cuda运算和copy操作都是直接执行,没有调用多线程执行。
#0 TVMBackendParallelLaunch (flambda=0x7ff20bc98a20, cdata=0x7fffce203b60, num_task=0) at /home/majiang/hd/opensource/tvm/src/runtime/thread_pool.cc:398
#1 0x00007ff20bc98688 in ?? ()
#2 0x00007ff1dbf3ab4e in tvm::runtime::<lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue)>::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue *) const (
__closure=0x3664bf0, args=…, rv=0x7fffce203eb0) at /home/majiang/hd/opensource/tvm/src/runtime/library_module.cc:88
#3 0x00007ff1dbf3bfbc in std::_Function_handler<void(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue
), tvm::runtime::WrapPackedFunc(TVMBackendPackedCFunc, const tvm::runtime::ObjectPtrtvm::runtime::Object&)::<lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue)> >::Minvoke(const std::_Any_data &, tvm::runtime::TVMArgs &&, tvm::runtime::TVMRetValue *&&) (__functor=…, __args#0=…, __args#1=@0x7fffce203e10: 0x7fffce203eb0) at /usr/include/c++/7/bits/std_function.h:316
#4 0x00007ff1db3b52ec in std::function<void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue
)>::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const (
this=0x36e8f10, args#0=…, __args#1=0x7fffce203eb0) at /usr/include/c++/7/bits/std_function.h:706
#5 0x00007ff1db3b4e32 in tvm::runtime::PackedFunc::CallPacked (this=0x36e8f10, args=…, rv=0x7fffce203eb0)
at /home/majiang/hd/opensource/tvm/include/tvm/runtime/packed_func.h:1040
#6 0x00007ff1dbfa6c55 in tvm::runtime::GraphRuntime::<lambda()>::operator()(void) const (
closure=0x36e8f00)
at /home/majiang/hd/opensource/tvm/src/runtime/graph/graph_runtime.cc:402
#7 0x00007ff1dbfaa837 in std::_Function_handler<void(), tvm::runtime::GraphRuntime::CreateTVMOp(const tvm::runtime::TVMOpParam&, const std::vector&, size_t)::<lambda()> >::Minvoke(const std::_Any_data &) (__functor=…) at /usr/include/c++/7/bits/std_function.h:316
#8 0x00007ff1db439068 in std::function<void ()>::operator()() const (this=0x392ff70) at /usr/include/c++/7/bits/std_function.h:706
#9 0x00007ff1dbfa2fb9 in tvm::runtime::GraphRuntime::Run (this=0x4076790) at /home/majiang/hd/opensource/tvm/src/runtime/graph/graph_runtime.cc:56

vscode中使用clang-tidy

介绍

clang-tidy是一个开源的lint工具。
它的主要作用:
a) 自动化检查代码格式是否满足要求
b) 增强编译器的检查功能,提示可能出错或有性能问题的代码

背靠clang/llvm的强大能力,clang-tidy提供了极强的定制和扩展能力。
这使得很多新的大型C/C++项目从项目初始就启用clang-tidy。

vscode上使用clang-tidy

在vscode上使用clang-tidy很简单,只需要安装Clang-Tidy插件就可以了。
该插件的基本原理是调用clang-tidy –export-fixes=- 输出文本,然后解析文本后组装为vs能识别的告警信息。

准备环境

安装clang-tidy并配置好插件

首先需要安装clang-tidy,使用apt安装或者自行编译都可以。
然后安装Clang-Tidy插件,并确保插件配置能找到clang-tidy的程序(确保路径或者PATH正确)。

为工程中的代码生成compile_commands.json

clang-tidy和许多clang体系工具一样,知道源代码编译命令后可以工作得更好。
由于源代码文件众多,实际上可操作的方法只有使用编译系统自动生成的编译命令记录compile_commands.json。使用cmake的体系,添加-DCMAKE_EXPORT_COMPILE_COMMANDS=ON就能自动生成该文件。其他构建体系也有类似的解决方案,可参考https://sarcasm.github.io/notes/dev/compilation-database.html
生成该文件后,还需要注意把这个.json放置到源代码的父目录下,否则clang-tidy会找不到。如果在${top_dir}/build中构建工程并生成了compile_commands.json,但是代码在${top_dir}/src中,则clang-tidy无法自动找到compile_commands.json,需要把其拷贝到${top_dir}下。

修复Clang-Tidy不支持中文的bug

Clang-Tidy使用了clang-tidy文本输出YAML格式的部分(来自 –export-fixes部分)。
示例如下:

1
2
3
4
5
6
7
8
MainSourceFile:  '/media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/cpp_exercise/llvm_study_Kaleidoscope/./main.cpp'
Diagnostics:
- DiagnosticName: cppcoreguidelines-pro-type-vararg
DiagnosticMessage:
Message: do not call c-style vararg functions
FilePath: '/tmp/test_main.cpp'
FileOffset: 1266
Replacements: []

最核心的信息是FilePath和FileOffset,这两个信息给出了Vscode界面应该在哪里显示告警。
但不幸的是,FileOffset这个值是clang-tidy给其自动修复工具用的,所以其值是一个以byte计数的偏移。
而在vscode中,文件位置的offset不是以byte记的,而是以字符来计算的。如果混入了中文等多byte字符,则vscode中的offset数值将小于clang-tidy给出的FileOffset。

更加糟糕的是,vscode当前没有给出把一个FileOffset转换为行号和列号的接口。其只提供了TextDocument.positionAt(offset: number)。这里的offset是以字符记的。看起来vscode是把单个字符当做了最小单元(哪怕这个字符实际上对应多个byte,可能这样对上层抽象的处理更加容易)。
由于上面描述的问题,一旦代码中出现中文等多byte字符,Clang-Tidy插件给出的告警就会向下漂移(由于其调用了TextDocument.positionAt,并且传入的是以byte记的offset,所以计算出的lineno要更大)。

参考 https://github.com/notskm/vscode-clang-tidy/issues/13,已经有人提到了这个问题,并且作者也给出了与我同样的分析,但是没有提出解决方案。

但是,实际上clang-tidy在非YAML部分其实已经给出了正确的行号和列号,如下所示。

1
main.cpp:46:3: warning: do not call c-style vararg functions [cppcoreguidelines-pro-type-vararg]

很奇怪的是Clang-Tidy插件专门从这一行中提取了warning这个关键字用来计算提示信息的严重程度,但是没用这里的行号和列号。
一种快速的规避方案,可以就从这里提取行号和列号。参考如下补丁。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
--- /home/majiang/.vscode/extensions/notskm.clang-tidy-0.4.1/out/tidy.js
+++ /home/majiang/.vscode/extensions/notskm.clang-tidy-0.4.1/out/tidy-fix.js
@@ -97,6 +97,7 @@
"FilePath": diag.DiagnosticMessage.FilePath,
"FileOffset": diag.DiagnosticMessage.FileOffset,
"Replacements": diag.DiagnosticMessage.Replacements,
+ "Lineno": 0,
"Severity": vscode.DiagnosticSeverity.Warning
}
});
@@ -109,6 +110,7 @@
"FilePath": diag.FilePath,
"FileOffset": diag.FileOffset,
"Replacements": diag.Replacements ? diag.Replacements : [],
+ "Lineno": 0,
"Severity": vscode.DiagnosticSeverity.Warning
}
});
@@ -117,7 +119,8 @@
let diagnostics = structuredResults.Diagnostics;
const severities = collectDiagnosticSeverities(clangTidyOutput);
for (let i = 0; i < diagnostics.length || i < severities.length; i++) {
- diagnostics[i].DiagnosticMessage.Severity = severities[i];
+ diagnostics[i].DiagnosticMessage.Severity = severities[i].severity;
+ diagnostics[i].DiagnosticMessage.Lineno = severities[i].lineno;
}
return structuredResults;
}
@@ -129,10 +132,9 @@
if (diagnosticMessage.Replacements.length > 0) {
diagnosticMessage.Replacements
.forEach(replacement => {
- const beginPos = document.positionAt(replacement.Offset);
- const endPos = document.positionAt(replacement.Offset + replacement.Length);
+ const line = Number(diagnosticMessage.Lineno) - 1;
const diagnostic = {
- range: new vscode.Range(beginPos, endPos),
+ range: new vscode.Range(line, 0, line, Number.MAX_VALUE),
severity: diagnosticMessage.Severity,
message: diagnosticMessage.Message,
code: diag.DiagnosticName,
@@ -142,7 +144,7 @@
});
}
else {
- const line = document.positionAt(diagnosticMessage.FileOffset).line;
+ const line = Number(diagnosticMessage.Lineno) - 1;
results.push({
range: new vscode.Range(line, 0, line, Number.MAX_VALUE),
severity: diagnosticMessage.Severity,
@@ -157,28 +159,28 @@
exports.collectDiagnostics = collectDiagnostics;
function collectDiagnosticSeverities(clangTidyOutput) {
const data = clangTidyOutput.split('\n');
- const regex = /^.*:\d+:\d+:\s+(warning|error|info|hint):\s+.*$/;
+ const regex = /^.*:(\d{1,})+:(\d{1,})+:\s+(warning|error|info|hint):\s+.*$/;
let severities = [];
data.forEach(line => {
const matches = regex.exec(line);
if (matches === null) {
return;
}
- switch (matches[1]) {
+ switch (matches[3]) {
case 'error':
- severities.push(vscode.DiagnosticSeverity.Error);
+ severities.push({severity: vscode.DiagnosticSeverity.Error, lineno: matches[1]});
break;
case 'warning':
- severities.push(vscode.DiagnosticSeverity.Warning);
+ severities.push({severity: vscode.DiagnosticSeverity.Warning, lineno: matches[1]});
break;
case 'info':
- severities.push(vscode.DiagnosticSeverity.Information);
+ severities.push({severity: vscode.DiagnosticSeverity.Information, lineno: matches[1]});
break;
case 'hint':
- severities.push(vscode.DiagnosticSeverity.Hint);
+ severities.push({severity: vscode.DiagnosticSeverity.Hint, lineno: matches[1]} );
break;
default:
- severities.push(vscode.DiagnosticSeverity.Warning);
+ severities.push({severity: vscode.DiagnosticSeverity.Warning, lineno: matches[1]});
break;
}
});

nvidia 性能分析工具

总体结构

nsight system

是个系统级的工具,可以综合看系统的全貌。
特别适合于查看那些负载卸载到了GPU上,哪些操作各自耗时多少,CPU是否在等待数据等等。
如下图所示。

另外,只能启动程序并测量,不能attach进程,稍微有点不方便。
可以选择手动start测量。

nsight compute

介绍

可参考https://devblogs.nvidia.com/using-nsight-compute-to-inspect-your-kernels/

Important Features
Interactive kernel profiler and API debugger
Graphical profile report
Result comparison across one or multiple reports within the tool
Fast Data Collection
UI and Command Line interface
Fully customizable reports and analysis rules

试用问题

图形界面直接启动程序始终无法连接上

问题现象:
试图使用该工具观察jupyter-notebook时,图形界面上使用launch启动后,始终报错连接不上(把程序改为/usr/bin/python3也一样),而attach列表中也始终看不到任何进程(使用命令行工具/opt/nvidia/nsight-compute/2019.5.0/nv-nsight-cu-cli –mode=launch 启动也一样)。

问题解决:
nsight compute分析有几个先决条件:
1 必须由compute来启动程序
2 程序必须要走到调用cuda库的地方,compute才能看到并连接上
3 compute默认只监控其启动的主程序,如果是主程序的child启动cuda(jupyter-notebook就属于这类),并且希望使用Interactive Profile模式,需要调用命令行工具nv-nsight-cu-cli –mode=launch –target-processes all启动程序,然后再在图形界面上attach。
注:
Profile模式下有Target Process选项,选择all就无需使用cli了,如下图所示。

问题分析过程:
由于图形界面上看不到任何有效的提示信息,转而考虑使用命令工具,看看有没有有用的提示。
使用如下方式启动nv-nsight-cu-cli –mode=launch后再nv-nsight-cu-cli –mode=attach –hostname 127.0.0.1给出了一个有用的提示。

1
==WARNING== Profiling kernels launched by child processes requires the --target-processes all option

再看nv-nsight-cu-cli 的help信息,原来默认情况下compute只监控主进程不监控child。于是改为如下命令启动。

1
/opt/nvidia/nsight-compute/2019.5.0/nv-nsight-cu-cli  --mode=launch --target-processes all jupyter-notebook

在图形界面中attach仍然看不到任何进程。
考虑到compute需要连接cuda,可能是没有执行到cuda。
在浏览器中连接jupyter-notebook并启动一个工作脚本,
此时图形界面attach列表中就出现了进程。

采集性能时报没有权限,没有出现数据

参考https://developer.nvidia.com/nvidia-development-tools-solutions-ERR_NVGPUCTRPERM-permission-issue-performance-counters。
为了简单起见,在个人电脑上可以直接允许所有用户采集GPU 性能。
在/etc/modprobe.d新建一个文件,写入如下一行,并重启一下就可以了。

1
options nvidia "NVreg_RestrictProfilingToAdminUsers=0"

试用感受

compute是一个比较综合和强大的工具,与intel的vtune类似,它还提供了基本的优化建议。其示例界面如下图所示。

在nvidia gpu调优时应该会起到很好的辅助作用。

NVIDIA Visual Profiler

似乎无法处理较新的GPU,采集不到数据。
考虑到它似乎属于较为旧的工具,暂没有进一步分析不能使用的原因。

TVM

简介

TVM是一套端到端的深度学习编译系统。
它的主要特性如下图所示。

第一,它支持将多种前端模型( Keras, MXNet, PyTorch, Tensorflow, CoreML, DarkNet等)编译到多种后端硬件上(包括传统的CPU/GPU,还包括FPGA、TPU等专用加速硬件)。
第二,它提供了一整套自动优化基础设施,能够帮助用户快速在一个新的硬件体系下建立起较高的性能。

独到之处

TVM 相对它的前辈如(Halide),有两个创新的点值得关注。

其一:
section4部分引入了一个可扩展的tensor compute primitives (就是硬件支持的用于特定矩阵运算的原语)描述方法。这使得TVM能快速支持新的硬件加速指令。
如下图所示。

使用类似RISC的思路,只需要提供少量细粒度的基本步骤,就可以通过组合配置建模出复杂的硬件加速指令。

其二:
TVM提供的autotuner相当强大。

TVM支持常规的黑盒自动优化,也就是使用黑盒优化算法反复到硬件上运行程序获得性能。
也支持基于预测的自动优化,在这种模式下,TVM根据硬件上获得的性能测试结果,训练出了一个性能预测模型。使用这个预测模型,TVM能实现快速的调优空间探索(论文中的效果比黑盒算法好)。因为模型预测很快,耗时低于1ms,而真实运行测试可能需要多耗费几十倍的时间。并且模型能持续从硬件中学习,而黑盒优化算法每次都必须从头开始。

TVM和其他项目的关系

https://github.com/apache/incubator-tvm/blob/master/docs/faq.rst

和Halide关系

https://github.com/apache/incubator-tvm/issues/682
http://docs.tvmlang.org/faq.html#tvm-s-relation-to-other-ir-dsl-projects answers the difference from existing projects, including Halide. In short, we specifically focus on deep learning, and optimize for multiple hardware backends (GPUs and other accelerators).

The major challenge is to make the schedule space complete enough to cover the state of art kernels for hardware back-ends we want to support, specifically gpu and other hardwares. The second challenge is to build the dsl representation to cover things we care about in deep learning(e.g. recurrence). The other issues include the ease of deployment and interpolation.

These challenges are not well addressed by existing frameworks(including Halide) and requires rethink and design of the stack as opposed to simply reuse an existing one.

You can also find that the TVM’s IR itself is evolving, and we continuously learn new lessons from hand optimization and tuning for various backends.

和MLIR的关系

参考https://discuss.tvm.ai/t/google-lasted-work-mlir-primer/1721/15
中TVM作者的如下回复。

Interpretation of MLIR’s Vision

I think what you answered reflects MLIR’s vision. Make the abstract class of IR and derive dialects. But not necessarily provide specific pass for the dialect, so if X-IR is a dialect of MLIR, then there are dialect specific passes that is needed in the pass.

Polyhedral dialect is a dialect in MLIR. In the current case, the polyhedral IR is part of the mlir codebase, which gives the view of “native”, but non-the-less it is a dialect just like the other automatic optimization dialect. The fact that it is part of the native code base does give an opinionated view of what what automatic optimization should be like in MLIR ecosystem. I think it is still very much an open problem, TVM has done a lot in this direction, and we can collectively innovate on this area.
How TVM can work with MLIR

First of all, MLIR won’t make TVM obsolete. In the contrary, it can help TVM stack by providing insights in IR design and possibly some lowering infrastructure.The community will keep improving our current IR infrastructure toward a better unified TVM-IR infra. We will try to define TVM dialects in MLIR to see if it makes sense to allow bi-directional translation between MLIR and TVM-IR, this way we can take benefit of some of the infra provided by MLIR and make TVM work together with MLIR’s ecosystem.

一些初步的认知

从已有的信息看,针对新的硬件体系或者新的运算逻辑,TVM应该是一个不错的选择。
初步查看的结果,其文档比较丰富,如果使用python接口编程,易用性也不错。

调试

使用O0 -g构建

1
cmake ../ -DCMAKE_BUILD_TYPE=Debug

获取各个op的运行时间

只需使用debug_runtime替代普通runtime即可

1
from tvm.contrib.debugger import debug_runtime as graph_runtime

具体实现可参考src/runtime/graph/debug/graph_runtime_debug.cc中的RunIndividual函数。其实核心就是对每一个op运行计时。
输出的有价值信息主要包括两类:
a 是在后台输出按执行顺序排列的op运行时间(如果是notebook,这个信息不会出现在浏览器中,会出现在启动notebook的终端)
b 是在notebook中打印按耗时占比排序的op执行时间

获取各个pass执行后的IR

使用git bisect定位问题

在tvm的目录下:
使用git bisect start开始二分查找
然后使用git bisect good $commit_id和git bisect bad ${commit_id}指定搜索区间。
就可以反复使用 cd build;cmake ../ -DCMAKE_BUILD_TYPE=Debug ; make -j 6构建并运行tvm,观察行为是否正常。
如果正常就git bisect good,如果异常就git bisect bad;如果中途某个版本遇到其他问题(例如还有其他bug干扰),可以使用git bisect skip。
找到问题后,使用git bisect reset还原。

HALIDE

介绍

参考 https://halide-lang.org/http://stellar.mit.edu/S/course/6/sp15/6.815/courseMaterial/topics/topic2/lectureNotes/14_Halide_print/14_Halide_print.pdf。
Halide 的核心思想是把图像处理(可以理解为矩阵运算)的算法(需要计算什么内容)和调度(如何优化执行计算)分开。

WHY

大规模的矩阵运算性能优化空间很大,但是目前已有的人工优化和编译器优化都有一些问题。
人工优化,有两种范式,一种是针对具体的场景手工优化,一种是提供BLAS, IPP, MKL, OpenCV这类高度优化的库。前一种效率太低(场景众多,还要针对不同的后端硬件,优化工作量太大),后者则只能提供局部最优的模块,无法在全局进行调度和融化优化。
编译器优化,可以看见完整的运算pipeline,但是优化的效果相对手动优化差了很多(就是一个最简单的矩阵乘法,编译器的输出都可能比手工优化要慢数倍)。同时,编译器中的很多核心优化决策都没有开放外部控制(比较典型的是,连循环展开的次数,GCC等编译器都是最近几年才通过#pragma unroll等方式提供了支持,更不要说直接控制cache block的大小等等),这导致在编译器的基础上人工再调优(纠正编译器的错误优化决策)很困难。
综合以上信息,高效和高质量的全局优化,还是要靠编译器。Halide的创造者也是沿着这个思路解决问题。

WHAT

Halide 是一种DSL(领域语言),也是该DSL的编译器。
它的核心思想是把运算的逻辑和运算的过程分离。将运算过程剥离出后,再将各个典型优化决策的控制变量和控制逻辑暴露出来,以便人工或者黑盒优化算法(如遗传算法等)能持续调整优化决策,达到更好的性能。
它主页中的示例代码很好的说明了其思想,如下所示。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
Func blur_3x3(Func input) {
Func blur_x, blur_y;
Var x, y, xi, yi;

// The algorithm - no storage or order
blur_x(x, y) = (input(x-1, y) + input(x, y) + input(x+1, y))/3;
blur_y(x, y) = (blur_x(x, y-1) + blur_x(x, y) + blur_x(x, y+1))/3;

// The schedule - defines order, locality; implies storage
blur_y.tile(x, y, xi, yi, 256, 32)
.vectorize(xi, 8).parallel(y);
blur_x.compute_at(blur_y, x).vectorize(x, 8);

return blur_y;
}

上面代码将blur的运算逻辑,与具体运算的实施过程进行了分离。用很简单的几个方法就指定了tile/vectorize等重要优化,以及其对应的参数。直观看起来,代码很简洁,并且要修改优化的类型和参数,工作量也很小。
而要迫使编译器实现同样的优化,需要写下面一大段代码。
并且,这样的代码由于直接使用intel 的SIMD原语,可移植性大幅度下降。
更加严重的是,想要微调各项优化参数(为了针对不同的硬件做优化),都需要对代码做大幅度的修改,工作量很大。
而上面的Halide代码,只需修改几个入参就可以完成优化决策的调整。哪怕Halide不提供内置的autotuner,使用一个简单的python脚本接入opentuner等黑盒优化框架也都会非常简单(毕竟只是改几个参数而已)。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
void box_filter_3x3(const Image &in, Image &blury) 
{
__m128i one_third = _mm_set1_epi16(21846);
#pragma omp parallel for
for (int yTile = 0; yTile < in.height(); yTile += 32) {
__m128i a, b, c, sum, avg;
__m128iblurx[(256/8)*(32+2)]; // allocate tile blurx array
for (int xTile = 0; xTile < in.width(); xTile += 256){
__m128i *blurxPtr = blurx;
for (int y = -1; y < 32+1; y++) {
const uint16_t *inPtr = &(in[yTile+y][xTile]);
for (int x = 0; x < 256; x += 8){
a = _mm_loadu_si128((__m128i*)(inPtr-1));
b = _mm_loadu_si128((__m128i*)(inPtr+1));
c = _mm_load_si128((__m128i*)(inPtr));
sum = _mm_add_epi16(_mm_add_epi16(a, b), c);
avg = _mm_mulhi_epi16(sum, one_third);
_mm_store_si128(blurxPtr++, avg);
inPtr += 8;
}}
blurxPtr = blurx;
for (int y = 0; y < 32; y++) {
__m128i *outPtr = (__m128i *)(&(blury[yTile+y][xTile]));
for (int x = 0; x < 256; x += 8) {
a = _mm_load_si128(blurxPtr+(2*256)/8);
b = _mm_load_si128(blurxPtr+256/8);
c = _mm_load_si128(blurxPtr++);
sum = _mm_add_epi16(_mm_add_epi16(a, b), c);
avg = _mm_mulhi_epi16(sum, one_third);
_mm_store_si128(outPtr++, avg);
}}}}}

HOW

Halide介绍ppt中的一页示意图很好地展示了它的工作过程。

图中飘逸的寥寥几笔注释已经把核心的工作原理讲清楚了。
如果对编译技术比较熟悉,看完注释后最核心的几个疑问应该就豁然开朗了。
简单的说,Halide语言没有独立的语法定义,也就不需要独立的lexer和parser。
它利用了C++语言的元编程能力,直接构造出了Halide语言的中间表达IR。
具体情况,随后一节会有详细的分析。

语言实现分析

为了分析Halide编译器的具体实现,下载并编译了Halide的代码(使用Ubuntu18.04自带的clang/llvm8,按照官方命令编译,比较简单)。
然后编译、运行和调试其tutorial目录中的各个示例。
可以对其实现有一个大致的了解。

语言定义

从最简单的示例入手,理解整体概念更容易。
参考下面代码注释。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
//这个示例用Halide完成了一个矩阵
int main(int argc, char **argv) {
/*这里声明了三个核心概念
Func 是一系列运算(expr)的合集
Var 表达运算中涉及的变量
Expr 表达单个运算过程
*/
Halide::Func gradient;
Halide::Var x, y;
Halide::Expr e = x + y;
//这里才完成了函数定义 f(x,y) = x+ y
gradient(x, y) = e;
/*
上面的声明式定义,其实已经体现出Halide是一个新的语言了。
有几个比较细节的点:
1 注意到我们没有对Var x和y进行赋值,就直接在expr中使用它们了。
可以这样做的原因是,它们只是对应二维数组的两个轴向而已,并不代表具体的值。
2 Halide::Expr e = x + y; 中的'='和'+'显然都不是常规语义。这一句实际上构建了一个Halide expr IR节点,op为+,LHS是x,RHS是y;
3 gradient(x, y) = e; 把expr关联到函数上,同样对应了IR上的操作。
*/

//调用realize完成了编译和运行,并得到了结果
Halide::Buffer<int32_t> output = gradient.realize(800, 600);
//下面只是校验Halide和通常的运算结果一致
for (int j = 0; j < output.height(); j++) {
for (int i = 0; i < output.width(); i++) {
if (output(i, j) != i + j) {
printf("Something went wrong!\n"
"Pixel %d, %d was supposed to be %d, but instead it's %d\n",
i, j, i + j, output(i, j));
return -1;
}
}
}
printf("Success!\n");
return 0;
}

总结一下,Halide的核心概念就是Func、Var和Expr。它没有文本源代码的格式,直接是寄生在C++上。Func、Var和Expr都是C++的class。在完成声明和赋值的同时,利用对=、+和()的重载,完成了Halide的IR构建。

编译、调试和源码分析

编译准备

Halide的编译只需要依赖LLVM,在ubuntu18.04上安装llvm8就可以了。
clone下Halide代码后,进入目录后执行如下命令,可构建出带有调试信息的版本。

1
2
3
4
5
mkdir build
cd build
export CXXFLAGS="-O0 -g3"
export OPTIMIZE="-O0"
make -e -f ../Makefile -j 8

编译示例并调试

Halide 前端

在上一步构建完成的build目录中,继续执行如下命令。

1
2
3
cd distrib/tutorial/
g++ lesson_01*.cpp -g -I ../include -L ../bin -lHalide -lpthread -ldl -o lesson_01 -std=c++11 -g3
gdb ./lesson_01

这个lesson01就是前面语言定义一节中已经给出过的示例代码。
由于Function和Var的定义都没有传参,可以跳过,直接单步跟踪expr的赋值。这里给出的结果就非常典型了。
首先变量x和y被转换为了Expr。

1
2
3
4
15518	    /** A Var can be treated as an Expr of type Int(32) */
15519 operator const Expr &() const {
15520 return e;
15521 }

然后Expr的operator+,就调用了Internal::Add::make构建了Expr这个IR。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
1139	Expr operator+(Expr a, Expr b) {
1140 user_assert(a.defined() && b.defined()) << "operator+ of undefined Expr\n";
1141 Internal::match_types(a, b);
1142 return Internal::Add::make(std::move(a), std::move(b));
1143 }
--->
21 Expr Add::make(Expr a, Expr b) {
22 internal_assert(a.defined()) << "Add of undefined\n";
23 internal_assert(b.defined()) << "Add of undefined\n";
24 internal_assert(a.type() == b.type()) << "Add of mismatched types\n";
25
(gdb)
26 Add *node = new Add;
27 node->type = a.type();
28 node->a = std::move(a);
29 node->b = std::move(b);
30 return node;
31 }

调试到这里已经基本能确认Halide前端的工作原理了,通过operator重载,Halide直接构造了IR的,跳过了Lexer和Parser部分。
下一个问题是,Halide的IR如何,或者在什么时机进行Codegen。

Halide的代码生成流程

如前所述,Halide中通过realize方法完成了代码的编译和运行。接着调试上面程序的gradient.realize调用。
可以看到如下的调用链条。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
#0  Halide::Internal::lower (output_funcs=std::vector of length 1, capacity 1 = {...}, pipeline_name="f0", t=..., args=std::vector of length 1, capacity 1 = {...}, 
linkage_type=Halide::LinkageType::ExternalPlusMetadata, requirements=std::vector of length 0, capacity 0, trace_pipeline=false,
custom_passes=std::vector of length 0, capacity 0) at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Lower.cpp:87
#1 0x00007ffff3cb7b6b in Halide::Pipeline::compile_to_module (this=0x7fffffffd920, args=std::vector of length 1, capacity 1 = {...}, fn_name="f0", target=...,
linkage_type=Halide::LinkageType::ExternalPlusMetadata) at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:506
#2 0x00007ffff3cb819b in Halide::Pipeline::compile_jit (this=0x7fffffffd920, target_arg=...)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:573
#3 0x00007ffff3cbbda7 in Halide::Pipeline::realize (this=0x7fffffffd920, outputs=..., t=..., param_map=...)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:1099
#4 0x00007ffff3cb98b0 in Halide::Pipeline::realize (this=0x7fffffffd920, sizes=std::vector of length 2, capacity 2 = {...}, target=..., param_map=...)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:703
#5 0x00007ffff3ac078c in Halide::Func::realize (this=0x7fffffffdbf0, sizes=std::vector of length 0, capacity 0, target=..., param_map=...)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2922
#6 0x00007ffff3ac0a7d in Halide::Func::realize (this=0x7fffffffdbf0, x_size=800, y_size=600, target=..., param_map=...)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2937
#7 0x000055555555d56a in main (argc=1, argv=0x7fffffffdd98) at lesson_01_basics.cpp:78

看到lower,老司机应该已经心领神会找到门路了。一般lower意味着高层表达向硬件层级扩展,表达的内容将越发具体完整。

lower函数的过程,可以看到大致有两个主要的工作,第一个是补充最终程序需要的系列流程,如初始化环境,建立循环,已经插入一些等等;第二个是进行各项高层优化(优化越接近源码,执行起来越简单。)。但是lower部分看到结尾,仍然没有向另外一种IR或者机器指令转换。
从lower返回后,在compile_jit函数中继续向下调试,可以最终找到如下堆栈回溯中,Halide完成了IR到LLVM-IR的codegen过程(当然如果结合代码分析,查找LLVM的相关流程,找到这里会更快)。
#0 Halide::Internal::CodeGen_LLVM::compile (this=0x5555557b60e0, input=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/CodeGen_LLVM.cpp:637
#1 0x00007ffff399a42c in Halide::codegen_llvm (module=…, context=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/CodeGen_LLVM.cpp:46
#2 0x00007ffff3c326e1 in Halide::compile_module_to_llvm_module (module=…, context=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/LLVM_Output.cpp:381
#3 0x00007ffff3c13c91 in Halide::Internal::JITModule::JITModule (this=0x7fffffffbf90, m=…, fn=…, dependencies=std::vector of length 0, capacity 0)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/JITModule.cpp:251
#4 0x00007ffff3cb86dd in Halide::Pipeline::compile_jit (this=0x7fffffffd920, target_arg=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:607
#5 0x00007ffff3cbbda7 in Halide::Pipeline::realize (this=0x7fffffffd920, outputs=…, t=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:1099
#6 0x00007ffff3cb98b0 in Halide::Pipeline::realize (this=0x7fffffffd920, sizes=std::vector of length 2, capacity 2 = {…}, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:703
#7 0x00007ffff3ac078c in Halide::Func::realize (this=0x7fffffffdbf0, sizes=std::vector of length 0, capacity 0, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2922
#8 0x00007ffff3ac0a7d in Halide::Func::realize (this=0x7fffffffdbf0, x_size=800, y_size=600, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2937
#9 0x000055555555d56a in main (argc=1, argv=0x7fffffffdd98) at lesson_01_basics.cpp:78
后面的流程更加直接一些,CodeGen_LLVM.cpp包含了主要的转换内容,compile_func中的 f.body.accept(this); 发起了LLVM-IR的发射动作。
后面就是CodeGen_LLVM.cpp中的一堆visit函数完成了针对不同类型Halide IR的LLVMIR代码生成。

遗留的学习

Halide自带的autotuner如何工作?

有意思的一些编程技巧

1 把可变参数的输入转成vector处理

1
2
3
4
5
6
template<typename... Args>
HALIDE_NO_USER_CODE_INLINE typename std::enable_if<Internal::all_are_convertible<Var, Args...>::value, FuncRef>::type
operator()(Args &&... args) const {
std::vector<Var> collected_args{std::forward<Args>(args)...};
return this->operator()(collected_args);
}

2 在父类中访问子类成员
https://en.wikipedia.org/wiki/Curiously_recurring_template_pattern

Typically, the base class template will take advantage of the fact that member function bodies (definitions) are not instantiated until long after their declarations, and will use members of the derived class within its own member functions, via the use of a cast; e.g.:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
template <class T> 
struct Base
{
void interface()
{
// ...
static_cast<T*>(this)->implementation();
// ...
}

static void static_func()
{
// ...
T::static_sub_func();
// ...
}
};

struct Derived : Base<Derived>
{
void implementation();
static void static_sub_func();
};