PCL+CUDA編程（一）

點雲的操作對運算資源的消耗是十分高的。但利用GPU並行運算的優點可以解決這個問題。下面我將跟大家分享關於利用CUDA處理PCL點雲數據的一些經驗。

首先舉一個簡單的例子說明CUDA程序是如何運作的。

我們先寫一個簡單的C++程序helloworld.cpp

/*
 * helloworld.cpp
 *
 *  Created on: Nov 25, 2016
 *      Author: lzp
 */

#include <iostream>

#include <addition.h>


int main(int argc, char** argv)
{
    int a=1,b=2,c;

    if(addition(a,b,&c))
        std::cout<<"c="<<c<<std::endl;
    else
        std::cout<<"Addition failed!"<<std::endl;

    return 0;
}

我們將利用addition()函數將a和b相加，然后由c儲存它們的和。
addition()函數在頭文件聲明：

/*
 * addition.h
 *
 *  Created on: Nov 25, 2016
 *      Author: lzp
 */

#ifndef INCLUDES_ADDITION_H_
#define INCLUDES_ADDITION_H_

/*check if the compiler is of C++*/
#ifdef __cplusplus
extern "C" bool addition(int a, int b, int *c);

#endif



#endif /* INCLUDES_ADDITION_H_ */

 修飾符extern "C"是CUDA和C++混合編程時必須的。然后我們來看addition()的在CUDA上的實現：

#include <addition.h>
__global__ void add(int *a, int *b, int *c)
{
    *c=*a+*b;
}

extern "C" bool addition(int a, int b, int *c)
{
    int *d_a, *d_b, *d_c;
    int size=sizeof(int);
    
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);
    
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
    
    add<<<1,1>>>(d_a, d_b, d_c);
    
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
    
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return true;
}

其中，帶有__global__修飾符的函數稱為”核函數“，它負責處理GPU內存里的數據，是並行計算發生的地方。而bool addition(int a, int b, int *c)充當了CPU和GPU之間數據傳輸的角色。也就是Host和Device之間的數據傳輸。

最后，編寫CMake文件編譯。

cmake_minimum_required(VERSION 2.6 FATAL_ERROR)

project(helloworld)

find_package(CUDA REQUIRED)

include_directories(../../includes)

cuda_add_executable (helloworld helloworld.cpp addition.cu)

其中include_directories的參數為.h文件所在的目錄。

 

下面我們用相同的程序結構，寫一個最簡單的例子，用CUDA對PCL點雲中的一個點進行操作。

/*
 * pcl_points_gpu.cpp
 *
 *  Created on: Nov 24, 2016
 *      Author: lzp
 */



#include <gpu_draw_cloud.h>
#include <pcl/io/pcd_io.h>

int main(int argc, char** argv)
{
    pcl::PointCloud<pcl::PointXYZRGB> cloud;
    pcl::gpu::DeviceArray<pcl::PointXYZRGB> cloud_device;


    cloud.width = 1;
    cloud.height =1;
    cloud.is_dense=false;
    cloud.points.resize(cloud.width*cloud.height);

    std::vector<float> point_val;

    for(size_t i=0; i<3*cloud.points.size(); ++i)
    {
        point_val.push_back(1024*rand()/(RAND_MAX+1.0f));
    }

    for (size_t i = 0; i < cloud.points.size(); ++i) {
        cloud.points[i].x = point_val[3 * i];
        cloud.points[i].y = point_val[3 * i + 1];
        cloud.points[i].z = point_val[3 * i + 2];
    }

    std::cout<<"cloud.points="<<cloud.points[0]<<std::endl;

    cloud_device.upload(cloud.points);

    cloud2GPU(cloud_device);

    cloud_device.download(cloud.points);

    std::cout<<"cloud.points="<<cloud.points[0]<<std::endl;
    return (0);
}

 

這段代碼模仿了PCL中寫點雲的一個例子，生成了一個點，坐標是隨機生成的。關鍵點是pcl::gpu::DeviceArray<pcl::PointXYZRGB>，這是一個可以將點雲傳輸到GPU上的橋梁。它的upload() 和download()方法相當於前面例子中的cudaMemcpy()。詳情可參考PCL的源碼倉庫中/gpu/examples/和/gpu/octree/這兩個目錄的源碼。

接下來是頭文件：

/*
 * gpu_draw_cloud.h
 *
 *  Created on: Nov 25, 2016
 *      Author: lzp
 */

#ifndef INCLUDES_GPU_DRAW_CLOUD_H_
#define INCLUDES_GPU_DRAW_CLOUD_H_


#include <iostream>
#include <pcl/point_types.h>
#include <pcl/gpu/containers/device_array.h>

/*check if the compiler is of C++*/
#ifdef __cplusplus


/*
 * Try accessing GPU with pointcloud
 * */
extern "C" bool cloud2GPU(pcl::gpu::DeviceArray<pcl::PointXYZRGB>& cloud_device);


#endif


#endif /* INCLUDES_GPU_DRAW_CLOUD_H_ */

然后是函數實現體：

#include <gpu_draw_cloud.h>




__global__ void change_points(pcl::gpu::PtrSz<pcl::PointXYZRGB> cloud_device)
{
    cloud_device[0].x+=1;
    pcl::PointXYZRGB q=cloud_device.data[0];
    printf("x=%f, y=%f, z=%f, r=%d, g=%d, b=%d \n", q.x, q.y, q.z, q.r, q.g, q.b);
}



extern "C" bool 
cloud2GPU(pcl::gpu::DeviceArray<pcl::PointXYZRGB>& cloud_device)
{
    change_points<<<1,1>>>(cloud_device);
    return true;
}

在這個例子中，我將CPU和GPU的數據交互放到主函數中了，因此cloud2GPU函數只充當了一個調用核函數的接口。值得注意的是，在核函數的參數中，傳入的pcl::gpu::DeviceArray<pcl::PointXYZRGB>隱式轉換成pcl::gpu::PtrSz<pcl::PointXYZRGB>了。這兩個數據類型是實現C++和CUDA混合編程的關鍵。

最后附上CMakeLists。

project(pcl_points_gpu)

find_package(PCL 1.8 REQUIRED)
find_package(CUDA REQUIRED)
INCLUDE(FindCUDA)

include_directories(../../includes)

include_directories(${PCL_INCLUDE_DIRS})
link_directories(${PCL_LIBRARY_DIRS})
add_definitions(${PCL_DEFINITIONS})

get_directory_property(dir_defs DIRECTORY ${CMAKE_SOURCE_DIR} COMPILE_DEFINITIONS)
set(vtk_flags)

foreach(it ${dir_defs})
    if(it MATCHES "vtk*")
    list(APPEND vtk_flags ${it})
    endif()
endforeach()

foreach(d ${vtk_flags})
    remove_definitions(-D${d})
endforeach()

cuda_add_executable (pcl_points_gpu pcl_points_gpu.cpp gpu_draw_cloud.cu)
target_link_libraries (pcl_points_gpu ${PCL_LIBRARIES})

留意13-24行，如果沒有這幾行，nvcc編譯時會報出類似這樣的錯誤：

nvcc fatal   : A single input file is required for a non-link phase when an outputfile is specified
CMake Error at pcl_points_gpu_generated_gpu_draw_cloud.cu.o.cmake:209 (message):
  Error generating
 XXXXXXXXXXXXXXXX./pcl_points_gpu_generated_gpu_draw_cloud.cu.o

根據https://github.com/PointCloudLibrary/pcl/issues/776的描述，這是VTK的一個bug所致，因此在CMake中添加了這幾行腳本。

希望這些例子對剛接觸PCL和CUDA的人有幫助。本人也是新手，對很多概念仍然模糊不清，望體諒。

 

推薦閱讀：

https://devblogs.nvidia.com/parallelforall/easy-introduction-cuda-c-and-c/

http://www.nvidia.com/docs/IO/116711/sc11-cuda-c-basics.pdf

http://www.nvidia.com/object/SC09_Tutorial.html