基于GPU编程的三维重建系统

 一、介绍  

 在机器时间领域我们做算法开发一般都会先在CPU上执行。CPU的全称是Central Processing Unit，而GPU的全称是Graphics Processing Unit。在命名上。这两种器件相同点是它们都是Processing Unit——处理单元；不同点是CPU是“核心的”，而GPU是用于“图像”处理的。在我们一般理解里，这些名称的确非常符合大众印象中它们的用途——一个是电脑的“大脑核心”，一个是图像方面的“处理器件”。但是聪明的人类并不会被简单的名称所束缚，他们发现GPU在一些场景下可以提供优于CPU的计算能力。

   

 二、GPU线程与架构

在CUDA架构下, 显示芯片执行时的最小单位是thread. 数个thread可以组成一个block. 一个block中的thread能存取同一块共享的内存, 而且可以快速进行同步的动作. 不同block中的thread无法存取同一个共享的内存, 因此无法直接互通或进行同步. 因此, 不同block中的thread能合作的程度是比较低的.

   

 然后依据thread, block和grid, 有着不同的存储. 核心就是thread. 可以结合下图进行理解:

    

三、基于GPU的双目三维重建

      之前所做的多频外差法三维重建，单目三重建都是基于CPU架构使用c++编程实现的，目前工业上已经大范围需要使用三维重建技术来完成一些测量需求，因此三维重建的质量与效率都显的十分重要，哪个环境满足不了要求都无法实现工业化量产的要求。基于此背景我将之前开发的双目结构光和单目结构光均实现了GPU化，提高了三维重建效率。

   实验环境，vs2013+CUDA9.0

  硬件：结构光双目三维测量系统，相机分辨率130w和500w  

void decMulPha5_GPU(unsigned char* GPUImageLeft, unsigned char* GPUImageRight,vector<unsigned char*> leftImage, vector<unsigned char*> rightImage, float* leftPhaseImg, float* rightPhaseImg, int w, int h)
{
	//InitCUDA();
	//使用event计算时间
	float time_elapsed = 0;
	cudaEvent_t start, stop;
	cudaEventCreate(&start);    //创建Event
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);    //记录数据准备时间

	// compute the size of data
	const size_t imageDataSize = sizeof(unsigned char) * w * h * PATTERNSNUM;
	// part1a. prepare the result data ,在设备端为结果开辟空间
	const size_t phaseDataSize = sizeof(float) * w * h;
	float* leftPhaseImage;
	float* rightPhaseImage;
	cudaMalloc(&leftPhaseImage, phaseDataSize);
	cudaMalloc(&rightPhaseImage, phaseDataSize);
	unsigned char* leftImageTexure; //left Image
	unsigned char* rightImageTexure; //right Image
	cudaMalloc(&leftImageTexure, imageDataSize);
	cudaMalloc(&rightImageTexure, imageDataSize);
	cudaDeviceSynchronize();	
	cudaMemcpy(leftImageTexure, GPUImageLeft, imageDataSize, cudaMemcpyHostToDevice);
	cudaMemcpy(rightImageTexure, GPUImageRight, imageDataSize, cudaMemcpyHostToDevice);
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time_elapsed, start, stop);

	printf("A: Time of prepare data: %f (ms)\n", time_elapsed);
	cudaEventRecord(start, 0);    //记录零均值化时间
	// part2. run kernel
	dim3 BlockDim(TX, TY);
	unsigned int bx = (w + BlockDim.x - 1) / BlockDim.x;
	unsigned int by = (h + BlockDim.y - 1) / BlockDim.y;
	dim3 GridDim(bx, by);
	DecMulPha5_Texture << < GridDim, BlockDim >> > (leftImageTexure, rightImageTexure, leftPhaseImage, rightPhaseImage, w, h);
	cudaDeviceSynchronize();
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time_elapsed, start, stop);
	printf("B: Time of DecMulPha5_Texture : %f (ms)\n", time_elapsed);
	// part3. copy the data from device

	cudaEventRecord(start, 0);    //记录查找相关点时间
	cudaMemcpy(leftPhaseImg, leftPhaseImage, sizeof(float) * w * h, cudaMemcpyDeviceToHost);
	cudaMemcpy(rightPhaseImg, rightPhaseImage, sizeof(float) * w * h, cudaMemcpyDeviceToHost);
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time_elapsed, start, stop);
	printf("C: Time of copy memory from Device: %f (ms)\n", time_elapsed);

	// par4. release data
	cudaEventRecord(start, 0);    //记录释放内存时间
	cudaFree(leftImageTexure);
	cudaFree(rightImageTexure);
	cudaFree(leftPhaseImg);
	cudaFree(rightPhaseImg);
	free(GPUImageLeft);
	free(GPUImageRight);
	//free(leftPhaseImage);
	//free(rightPhaseImage);
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time_elapsed, start, stop);
	printf("D: Time of release data: %f (ms)\n", time_elapsed);
}

四、结论

  采用的硬件与软件系统如图所示

 

 最终的结果比对是用现在的GPU方式与之前的CPU方式进行重建时间的对比，如下表所示

相机分辨率	CPU处理时间/ms	GPU处理时间/ms
1280*1024	1200ms~1300ms	230ms~280ms
2448*2048	3500ms~3800ms	900ms~1100ms

 交流微信