前言

本篇部落格是在我學習別人的部落格內容上實驗出來的，所以大部分並非原創內容，參考地址：祕籍傳送。

在此篇部落格的基礎上，我將其中的程式實驗，並自己閱讀了一番，下面是具體內容。

主要目標

是對一張轉換為灰度影象進行腐蝕和膨脹，通過CUDA在GPU實現此過程，將其輸出在cpu上，儲存在本地檔案中（*.jpg）。

具體實現

影象腐蝕：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<opencv2/opencv.hpp>
#include <stdio.h>

using namespace cv;
using namespace std;

__global__ void test(unsigned char *d_in, unsigned char *d_out, int iwidth, int iheight, Size erodeElement){
	int xIndex = threadIdx.x + blockDim.x*blockIdx.x;
	int yIndex = threadIdx.y + blockDim.y*blockIdx.y;
	
	int ewidth = erodeElement.width/2;
	int ehight = erodeElement.height/2;
	d_out[xIndex + yIndex * iwidth] = d_in[xIndex + yIndex * iwidth];
	if ((xIndex > ewidth) && (yIndex > ehight) && (xIndex+ewidth < iwidth) && (yIndex+ehight < iheight)){
		for (int i = -ehight; i <= ehight; i++){
			for (int j = -ewidth; j <= ewidth; j++){
				if (d_in[xIndex + yIndex * iwidth] > d_in[(yIndex + j) * iwidth + xIndex + i]){
					d_out[xIndex + yIndex * iwidth] = d_in[(yIndex + j) * iwidth + xIndex + i];
				}
			}
		}
	}
}

int main(){

	Mat imgsrc = imread("D:\\che.jpg", 0);
	if (imgsrc.data == NULL){
		printf("圖片讀取錯誤！\n");
		return -1;
	}

	unsigned char *d_in;
	unsigned char *d_out;
	
	int width = imgsrc.cols;
	int height = imgsrc.rows;
	//unsigned char out[];	是否可以通過陣列和Mat轉換的方式來實現將資料儲存，雖然相比於直接放在Mat.data上要麻煩
	cudaError cudaStatus;
	cudaStatus = cudaMalloc((void**)&d_in, width*height*sizeof(unsigned char));
	if (cudaStatus != cudaSuccess){
		fprintf(stderr, "cudaMalloc failed!\n");
		cudaFree(d_in);
		return 1;
	}

	cudaStatus = cudaMalloc((void**)&d_out, width*height*sizeof(unsigned char));
	if (cudaStatus != cudaSuccess){
		fprintf(stderr, "cudaMalloc failed!\n");
		cudaFree(d_in);
		cudaFree(d_out);
		return 1;
	}

	cudaStatus = cudaMemcpy(d_in, imgsrc.data, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess){
		fprintf(stderr, "cudaMemcpy failed!\n");
		cudaFree(d_in);
		cudaFree(d_out);
		return 1;
	}

	dim3 threads(32, 32);
	dim3 blocks((width + threads.x - 1)/threads.x, (height + threads.y - 1)/threads.y);
	Size Element(3, 5);

	test << <blocks, threads >> >(d_in, d_out, width, height, Element);

	Mat m(height, width, CV_8UC1, cv::Scalar::all(0));

	cudaStatus = cudaMemcpy(m.data, d_out, width*height*sizeof(unsigned char), cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess){
		fprintf(stderr, "cudaMemcpy2 failed!\n");
		cudaFree(d_in);
		cudaFree(d_out);
		return 1;
	}

	imwrite("D:\\test.jpg", m);

	printf("%d,%d", width, height);

	getchar();
	return 0;
}

 

程式碼大致解釋：

程式碼主要分幾塊：首先，對存放陣列的分配地址並初始化（主要是對傳入GPU的陣列初始化）；其次，是呼叫核函式執行具體任務；最後，將執行後的結果陣列（d_out）複製給記憶體中負責顯示的Mat資料型別（m），並將其儲存到本地。

具體解釋：

•         Mat imgsrc = imread("D:\\che.jpg", 0);
  	if (imgsrc.data == NULL){
  		printf("圖片讀取錯誤！\n");
  		return -1;
  	}
  	..........
  
  	cudaStatus = cudaMemcpy(d_in, imgsrc.data, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
  	if (cudaStatus != cudaSuccess){
  		fprintf(stderr, "cudaMemcpy failed!\n");
  		cudaFree(d_in);
  		cudaFree(d_out);
  		return 1;
  	}
   
  

將圖片按灰度影象讀入，併為之後的GPU所需陣列分配空間，其中，cudaStatus是cudaError型別，來提示錯誤，當然也可以加一句：printf("%s\n", cudaGetErrorString(cudaStatus));來具體輸出什麼錯誤（雖然有時候提示的錯誤莫名其妙），網上有關於各種錯誤提示的解決方法。

•         dim3 threads(32, 32);
  	dim3 blocks((width + threads.x - 1)/threads.x, (height + threads.y - 1)/threads.y);
  	Size Element(3, 5);
  
  	test << <blocks, threads >> >(d_in, d_out, width, height, Element);
   
  

這裡是定義核函式的引數和執行核函式，這裡的引數在參考的部落格中有詳細的解釋。

特別的，對於定義的blocks，我的理解是：平常對於Grid含有的block數量是：總的行數/breads（即每個block所含有的thread，當然也有例外，比如參考部落格的上一篇：祕籍傳送2），而這裡將breads設定確定了，就是（32，32），則若是不同的圖片載入進來，需要對其有些適應性的改變，舉個特殊的例子，若是影象為1x1的，那麼總不能將blocks設定為1/32，而它實際上或許只需要一個執行緒來處理（我不太肯定，懇請各位大神給出確定答案，非常感謝！），則應該將一個block塊交給它，所以算下來的blocks應該是1塊（在上面的定義式算下來的結果是1塊）。

下面是核函式：

•         int ewidth = erodeElement.width/2;
  	int ehight = erodeElement.height/2;

在腐蝕上，可以上網查詢有關概念，這裡的程式的邏輯是：將一個畫素點和周圍比較，將最小值賦值給這個畫素點，至於“周圍”有多大，是在主程式中的Size Element(3, 5);定義的。寬3高5的矩形，通過此矩形來腐蝕影象。而具體實現起來，則是以現畫素點為原點，減加一半的寬度和高度來遍歷周圍。

• if ((xIndex > ewidth) && (yIndex > ehight) && (xIndex+ewidth < iwidth) && (yIndex+ehight < iheight))

 這一句是防止越界，因為要從中心點加減來遍歷，要是超出影象的邊界，則要不無效要不出錯，上面四個關係對應於四周邊界。

• if (d_in[xIndex + yIndex * iwidth] > d_in[(yIndex + j) * iwidth + xIndex + i]){
  	d_out[xIndex + yIndex * iwidth] = d_in[(yIndex + j) * iwidth + xIndex + i];
  }

 這個判斷則是取最小值存入d_out陣列中。

最後是儲存：

•         Mat m(height, width, CV_8UC1, cv::Scalar::all(0));
  
          cudaStatus = cudaMemcpy(m.data, d_out, width*height*sizeof(unsigned char), cudaMemcpyDeviceToHost);
  	if (cudaStatus != cudaSuccess){
  		fprintf(stderr, "cudaMemcpy2 failed!\n");
  		cudaFree(d_in);
  		cudaFree(d_out);
  		return 1;
  	}
  
  	imwrite("D:\\test.jpg", m);

將資料複製到記憶體的m變數上，將其儲存到本地。

總結

以上是我對程式的理解，得到的圖片結果和參考部落格相同，上述的理解供參考，若是有理解不對的地方，懇請不吝賜教，我在cuda上還是新手，謝謝各位！