程序參考文章：http://blog.csdn.net/gamesdev/article/details/17535755? 程序優化2

簡介：CUDA ，MPI，Hadoop都是并行運算的工具。CUDA是基于NVIDIA GPU芯片計算。

闡述：GPU有很多個核(幾百個)，每個核可以跑一個線程，多個線程組成一個單位叫做塊。

舉個例子：
有三個向量 int a, b, c; 我們要計算a和b的向量之和存放到c中。

一般C語言：for(int i=0; i<10; i++)? c = a + b; 這個程序是順序執行的！

CUDA編程做法：

GPU中的每個線程(核)有一個獨立序號叫index，那么只要序號從0到9的線程執行c[index] = a[index] + b[index]; 就可以實現以上的for循環。

GPU的可貴之處就是，可以并發運行多個線程，相當于一個時間內賦值10次。

cuda.cu#include <stdio.h> #include <cuda_runtime.h>//運行在GPU __global__ void vectorADD(int* a, int* b, int* c) {int index = threadIdx.x; //獲得當前線程的序號if(index < blockDim.x)c = a + b; } int main () {//定義10個GPU運算線程int N = 10; // 本地開辟三個數組存放我們要計算的內容int* h_a = (int*) malloc (N * sizeof(int));int* h_b = (int*) malloc (N * sizeof(int));int* h_c = (int*) malloc (N * sizeof(int));// 初始化數組A, B和Cfor(int i=0; i<N; i++) {h_a = i;h_b = i;h_c = 0;}// 計算10個int型需要的空間int size = N * sizeof(int); // 在GPU上分配同樣大小的三個數組int* d_a;int* d_b;int* d_c;cudaMalloc((void**)&d_a, size);cudaMalloc((void**)&d_b, size);cudaMalloc((void**)&d_c, size);// 把本地的數組拷貝進GPU內存cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);cudaMemcpy(d_c, h_c, size, cudaMemcpyHostToDevice);// 定義一個GPU運算塊 由 10個運算線程組成dim3 DimBlock = N;// 通知GPU用10個線程執行函數vectorADDvectorADD<<<1, DimBlock>>>(d_a, d_b, d_c);// 將GPU運算完的結果復制回本地cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost); // 釋放GPU的內存cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);// 驗證計算結果for(int j=0; j<N; j++)printf("%d ", h_c);printf("\n"); }

警告！：這個例子是編譯不通過的；

首先：對 threadidx的使用，只能在CU文件里面；

其次：在cu文件里初始化數組是錯誤的： int * a ； a = new int [x]；是錯誤的；? 并且 malloc也是不可以的；

再者：文件路徑里面不能包含中文，否則會出現 MSB8791 這種錯誤！

2. 利用CUDA并行計算點云法線

兩個函數都存在于CU文件里! 通過外部CPP文件函數進行調用

void normalEstimate(pcl::PointCloud<pcl::PointXYZRGB> &input ,pcl::PointCloud<pcl::PointXYZRGB> &output,int k_,float search_parameter_,int THREAD_NUM)

//運行在GPU//cal the Normal

__global__ void normalEstimateSingle(pcl::PointCloud<pcl::PointXYZRGB> &input ,pcl::PointCloud<pcl::PointXYZRGB> &output, int* nn_indices ,int* nn_dists, int Gap, float search_parameter_) { const size_t computeSize =input.size() / Gap; const size_t tID = size_t(threadIdx.x ); int Mark;clock_t startTime; // 開始計時if ( tID == 0 ) startTime =clock( );// 選擇任意一個線程進行計時 //Thread loop!//循環發現鄰域！尋找法線！for ( size_t idx = tID *computeSize; idx < ( tID + 1 ) * computeSize && idx < input.size(); ++idx ) { // pOut[threadIdx.x] += pIn[i] * pIn[i]; Mark = pcl::searchForNeighbors (idx, search_parameter_, nn_indices, nn_dists);//對第IDX個建立索引!if (Mark == 0){output.points[idx].normal[0] = output.points[idx].normal[1] = output.points[idx].normal[2] = output.points[idx].curvature = std::numeric_limits<float>::quiet_NaN ();continue;}else {if (!isFinite (input[idx]) || Mark == 0){output.points[idx].normal[0] = output.points[idx].normal[1] = output.points[idx].normal[2] = output.points[idx].curvature = std::numeric_limits<float>::quiet_NaN ();continue;}}pcl::computePointNormal (input, nn_indices,output.points[idx].normal[0], output.points[idx].normal[1], output.points[idx].normal[2], output.points[idx].curvature);pcl::flipNormalTowardsViewpoint (input_->points[idx], vpx_, vpy_, vpz_,output.points[idx].normal[0], output.points[idx].normal[1], output.points[idx].normal[2]);} if ( tID == 0 ) *pElapsed =clock( ) - startTime;// 結束計時，返回至主程序 }

//運行在CPU端！

// as the input extern "C" void normalEstimate(pcl::PointCloud<pcl::PointXYZRGB> &input ,pcl::PointCloud<pcl::PointXYZRGB> &output,int k_,float search_parameter_,int THREAD_NUM) {// 在GPU上分配同樣大小的三個數組pcl::PointCloud<pcl::PointXYZRGB> &inputX ;pcl::PointCloud<pcl::PointXYZRGB> &outputX;int* nn_indices ;int* nn_dists;// 1、設置設備 cudaError_t cudaStatus = cudaSetDevice( 0 );// 只要機器安裝了英偉達顯卡，那么會調用成功 if ( cudaStatus != cudaSuccess ) { fprintf( stderr, "調用cudaSetDevice()函數失敗！" ); return ;//false; } // 使用CUDA內存分配器分配host端 //cudaError_t cudaStatus = cudaMallocHost( &inputX, input.size() * sizeof( pcl::pointXYZRGB ) ); //cudaError_t cudaStatus = cudaMallocHost( &outputX, output.size() * sizeof( pcl::Normal ) ); // 2、分配顯存空間 cudaError_t cudaStatus = cudaMalloc( &inputX, input.size() * sizeof( pcl::pointXYZRGB ) ); cudaError_t cudaStatusX = cudaMalloc( &outputX, output.size() * sizeof( pcl::Normal ) );// cudaStatus = cudaMalloc( (void**)&pData, DataSize * sizeof( int) ); if ( cudaStatus != cudaSuccess) { fprintf( stderr, "調用cudaMalloc()函數初始化顯卡中數組時失敗！" ); break; } // 3、將宿主程序數據復制到顯存中 cudaError_t cudaStatus2 = cudaMemcpy( inputX, input, input.size() * sizeof( pcl::pointXYZRGB ),cudaMemcpyHostToDevice ); cudaError_t cudaStatusX2 = cudaMemcpy(outputX,output,output.size() * sizeof( pcl::pointXYZRGB ),cudaMemcpyHostToDevice );if ( cudaStatus != cudaSuccess) { fprintf( stderr, "調用cudaMemcpy()函數初始化宿主程序數據數組到顯卡時失敗！" ); break; } //cudaMalloc( (void**)&nn_dists, k_ * sizeof( int) ); //cudaMalloc( (void**)&nn_indices, k_ * sizeof( int) ); //cudaMalloc( (void**)&Normal3f, 3 * sizeof( float) ); // 4、執行程序，宿主程序等待顯卡執行完畢 normalEstimateSingle<<<1, THREAD_NUM>>>( inputX,outputX, nn_indices, nn_dists, THREAD_NUM ,search_parameter_);//normalEstimateSingle(pcl::PointCloud<pcl::PointXYZRGB> &input ,pcl::PointCloud<pcl::PointXYZRGB> &output, int* nn_indices ,int* nn_dists, int Gap)// 5、查詢內核初始化的時候是否出錯 cudaStatus = cudaGetLastError( ); if ( cudaStatus != cudaSuccess) { fprintf( stderr, "顯卡執行程序時失敗！" ); break; }// 6、與內核同步等待執行完畢 cudaStatus = cudaDeviceSynchronize( ); if ( cudaStatus != cudaSuccess) { fprintf( stderr, "在與內核同步的過程中發生問題！" ); break; } // 7、獲取數據 //只復制出法線即可！cudaStatus = cudaMemcpy(output,outputX,output.size() * sizeof( pcl::pointXYZRGB ),cudaMemcpyHostToDevice ); if ( cudaStatus != cudaSuccess) { fprintf( stderr, "在將結果數據從顯卡復制到宿主程序中失敗！" ); break; } cudaFree( outputX );cudaFree( inputX ); }
注意事項：運行在GPU的函數，只能是原子函數，詳情請見 《高性能并行編程實踐》

總結

以上是生活随笔為你收集整理的CUDA 编程实例:计算点云法线的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。