Ubuntu16.04+RTX2080TI+Cuda10.1调试ElasticFusion小记

Ubuntu16.04+RTX2080TI+Cuda10.1调试ElasticFusion小记

ElasticFusion是Thomas Whelan博士在Dyson Research做研究员时的经典工作。

环境是Ubuntu16.04+RTX2080TI+Cuda10.1

依赖安装

各种依赖安装参考项目的Github主页，Ubuntu16.04需要注意的是openjdk-7-jdk的安装，因为Ubuntu16.04的安装源已经默认没有openjdk7了，所以要自己手动添加仓库：

sudo add-apt-repository ppa:openjdk-r/ppa
sudo apt-get update
sudo apt-get install openjdk-7-jdk

代码编译

打开项目根目录的build.sh，根据里面的编译部分进行代码编译：

cd ../Core
mkdir build
cd build
cmake ../src
make -j8
cd ../../GPUTest
mkdir build
cd build
cmake ../src
make -j8
cd ../../GUI
mkdir build
cd build
cmake ../src
make -j8

可以看到代码目录的三个子项目都需要分别编译，中间也会有各种问题存在，下面展开叙述。

Core

首先是编译Core项目，这是核心代码，也是最常出问题的。

在编译前，需要根据自己的显卡配置和Cuda对/Core/src目录下的CMakeLists.txt进行修改，主要是第32行的CUDA_ARCH_BIN，这个算力是通过编译Cuda目录下的samples，然后调用./deviceQuery进行查询的。通常1080卡可以把这项改成61，2080是75。

set(CUDA_ARCH_BIN "75" CACHE STRING "Specify 'real' GPU arch to build binaries for, BIN(PTX) format is supported. Example: 1.3 2.1(1.3) or 13 21(13)")

修改完进行编译：

cd Core
mkdir build
cd build
cmake ../src
make -j8

在make时一直报错：/Core/src/Cuda/reduce.cu文件“__shfl_down” is undefined错误，查了很多资料也没有很好解决这个问题，有个资料提到cuda9.0之后弃用了该函数。

github issues里luruibo1995提到‘modify the function __shfl_down to __shfl_down_sync, the mask i choice is 0xFFFFFFFF’，但是文件里要改的地方太多，也没有尝试这个方法（后来有测试这个方法，确实是可行的，跟后面取巧的方法不论是重建过程还是重建结果都基本没有什么差别）。

为了尽快跑起来程序，这里采用了取巧的方式，查看这个reduce.cu文件，发现从第56行到第80行有对算力不够的device声明__shfl_down函数：

#if __CUDA_ARCH__ < 300
__inline__ __device__
float __shfl_down(float val, int offset, int width = 32)
{
    static __shared__ float shared[MAX_THREADS];
    int lane = threadIdx.x % 32;
    shared[threadIdx.x] = val;
    __syncthreads();
    val = (lane + offset < width) ? shared[threadIdx.x + offset] : 0;
    __syncthreads();
    return val;
}

__inline__ __device__
int __shfl_down(int val, int offset, int width = 32)
{
    static __shared__ int shared[MAX_THREADS];
    int lane = threadIdx.x % 32;
    shared[threadIdx.x] = val;
    __syncthreads();
    val = (lane + offset < width) ? shared[threadIdx.x + offset] : 0;
    __syncthreads();
    return val;
}
#endif

将#if __CUDA_ARCH__ < 300和#endif两行注释掉，编译就可以顺利通过了。

GPUTest

顺利编译：

cd GPUTest
mkdir build
cd build
cmake ../src
make -j8

然后运行程序：

./GPUTest ../

一阵眼花缭乱之后，输出对应显卡的4行配置，将其复制到/Core/src/Utils/GPUConfig.h：

icpStepMap["GeForce RTX 2080 Ti"] = std::pair<int, int>(128, 128);
rgbStepMap["GeForce RTX 2080 Ti"] = std::pair<int, int>(96, 128);
rgbResMap["GeForce RTX 2080 Ti"] = std::pair<int, int>(128, 304);
so3StepMap["GeForce RTX 2080 Ti"] = std::pair<int, int>(96, 112);

如果不修改的话，运行完ElasticFusion会报Your GPU ‘’ ‘’ isn't in the ICP Step performance database, please add it的警告。

改了Core的文件后，重新编译下Core。

GUI

顺利编译：

cd GUI
mkdir build
cd build
cmake ../src
make -j8

下载数据集，然后运行：

./ElasticFusion -l dyson_lab.klg

程序顺利运行。