CUDA學習筆記一：CUDA+OpenCV的映像轉置，採用Shared Memory進行CUDA程式最佳化

原創文章，轉載請註明出處...... 一、問題背景        

最近要做一個關於CUDA的學習分享報告，想在報告中舉一個利用CUDA進行影像處理的例子，並使用Shared Memory避免Global Memory不合并訪存情況，提高影像處理效能。但是對於CUDA程式如何讀取映像有點困惑，網上找到了一篇“第二個cuda程式——映像展開”的博文點擊開啟連結，所示代碼涉及了映像互動部分，但是需要包含“cutil_inline.h”標頭檔（據說是開發人員編寫常式時用的標頭檔），悲催的是自從CUDA5.0之後“cutil.h”跟“cutil_inline.h”等標頭檔就被移除了，而我安裝的是CUDA6.5，所以用不了博文中讀取映像的方法。

後來在CUDA的Sample中看到了影像處理的樣本程式，但是說實話，對我這種剛入門的人來說，程式有點複雜，所以放棄了研究它的念頭。另外，有人跟我說CUDA有一個NPP庫可以支援映像的互動，但是不知道這個庫該怎麼調用。

於是，我想了個方法，用OpenCV函數來讀取、顯示映像，映像的處理則交由CUDA核函數完成，因為Windows平台上OpenCV與CUDA編程都是在Visual Studio上完成的，因此，此法是可行的。 二、實驗過程

1. 實驗平台：Visual Studio 2010，CUDA 6.5，OpenCV 2.4.9

2. OpenCV開發環境配置

要在VC上調用OpenCV函數庫函數，需要先進行OpenCV開發環境的配置，OpenCV的安裝與環境配置參考博文“【OpenCV入門教程之一】 安裝OpenCV：OpenCV 3.0、OpenCV 2.4.8、OpenCV 2.4.9 ”點擊開啟連結

3. 代碼

本文採用CUDA+OpenCV的環境進行映像轉置處理，分別採用CPU與GPU對讀入的映像進行轉置，其中GPU的實現分為Global Memory與Shared Memory兩個版本，實驗結果表明採用Shared Memory進行映像轉置，可以避免不合并訪存的情況，從而提高程式運行速度。

（1）CPU、GPU Global Memory與GPU Shared Memory進行映像轉置的函數定義均在標頭檔”imageTranspose.h“中：

#ifndef _IMAGETRANSPOSE_CU#define _IMAGETRANSPOSE_CU#include <stdio.h>#include <stdlib.h>#include <cuda_runtime.h>#include <device_launch_parameters.h>#include <Windows.h>//用於計時#include <time.h>#define W 16  //Block的尺寸#define N 1024  //Grid的尺寸//映像資料放在GlobalMemory上進行處理__global__ static void GPUImageTranspose_Global(unsigned char *imageDataSrc, unsigned char *imageDataDst, int Width, int Height){int tid = threadIdx.x + blockDim.x * blockIdx.x;  //得到線程id//越界判斷，線程可能沒有與之對應的像素if(tid >= Width * Height)return;int i, j;  i = tid / Width;j = tid % Width;//轉置imageDataDst[j * Height + i] = imageDataSrc[tid];               return;}//映像資料放在SharedMemory上進行處理__global__ static void GPUImageTranspose_Shared(unsigned char *imageDataSrc, unsigned char *imageDataDst, int Width, int Height){__shared__ unsigned char tile[W][W];  //聲明儲存映像資料的Shared Memory//計算當前線程處理的像素在輸入矩陣中的索引int x = threadIdx.x + blockIdx.x * W;int y = threadIdx.y + blockIdx.y * W;int index_in = x + y * Width; //這個越界判斷很關鍵，不然輸出結果錯誤if(index_in >= Width * Height)return;//將當前線程處理的像素值從Global Memory複製到Shared Memorytile[threadIdx.y][threadIdx.x] = imageDataSrc[index_in];     __syncthreads();  //線程同步語句//計算當前線程處理的像素在輸出矩陣中的索引x = threadIdx.x + blockIdx.y * W;y = threadIdx.y + blockIdx.x * W;int index_out = x + y * Height; //將當前線程處理的像素值從Shared Memory複製到Global Memory，通過座標變換完成轉置imageDataDst[index_out] = tile[threadIdx.x][threadIdx.y];return;}//CPU完成映像轉置void CPUImageTranspose(unsigned char *imageDataSrc, unsigned char *imageDataDst, int Width, int Height){      int i, j;       if(imageDataSrc == NULL || imageDataDst == NULL || Width <= 0 || Height <= 0)         return;    //遍曆映像資料完成映像轉置for(i=0; i<Height; i++) {          for(j=0; j<Width; j++) {             imageDataDst[j * Height + i] = imageDataSrc[i * Width + j];         }     }       }#endif

（2）主函數在”imageTranspose.cu“檔案中定義，主函數調用映像轉置函數進行影像處理，以下展示為使用CPU及GPU Global Memory進行映像轉置：

#include <cv.h>  //use OpenCV#include <highgui.h>#include <stdio.h>#include <time.h>  //clock_t clock()#include "imageTranspose_cu.h"int main(){//通過OpenCV函數讀取映像IplImage *ImgSrc = cvLoadImage("<span style="text-align: justify;">Lena.jpg</span>", CV_LOAD_IMAGE_GRAYSCALE);int Width = ImgSrc->width;int Height = ImgSrc->height;//輸出映像的寬高尺寸互換IplImage *ImgDst_GPU_Global = cvCreateImage(cvSize(Height, Width), IPL_DEPTH_8U, 1);IplImage *ImgDst_CPU = cvCreateImage(cvSize(Height, Width), IPL_DEPTH_8U, 1);//定義指向映像資料的指標，作為函數調用的參數unsigned char *pSrcData = (unsigned char*)(ImgSrc->imageData);unsigned char *pDstData_Global = (unsigned char*)(ImgDst_GPU_Global->imageData);unsigned char *cDstData = (unsigned char*)(ImgDst_CPU->imageData);//分配顯存用於儲存原映像數組和靶心圖表像數組unsigned char *device_ImgDataSrc = NULL;unsigned char *device_ImgDataDst_Global = NULL;cudaMalloc((void**)&device_ImgDataSrc, sizeof(unsigned char) * Width * Height);cudaMalloc((void**)&device_ImgDataDst_Global, sizeof(unsigned char) * Height * Width);//將原映像數組傳遞到顯存中cudaMemcpy(device_ImgDataSrc, pSrcData, sizeof(unsigned char) * Width * Height, cudaMemcpyHostToDevice);//GlobalMemory版本的參數設定int dimGrid_Global = 6000;  //每個Grid允許的最大Block數為65535int dimBlock_Global = 512;  //每個Block允許的最大線程數為512//建立事件，啟動定時cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);//啟動CUDA核函數，GPU進行映像轉置GPUImageTranspose_Global<<<dimGrid_Global, dimBlock_Global>>>(device_ImgDataSrc, device_ImgDataDst_Global,  ImgSrc->width, ImgSrc->height);cudaEventRecord(stop, 0);cudaEventSynchronize(start);  //事件同步語句cudaEventSynchronize(stop);  //事件同步語句//計算CUDA核函數進行映像轉置耗時，並顯示時間float GPUTime_Global = 0;cudaEventElapsedTime(&GPUTime_Global, start, stop);printf("GPU_Time_Global = %f\n", GPUTime_Global);//將結果傳遞至記憶體cudaMemcpy(pDstData_Global, device_ImgDataDst_Global, sizeof(unsigned char) * Width * Height, cudaMemcpyDeviceToHost);//計算CPU進行映像轉置耗時，並顯示時間clock_t t1 = clock();CPUImageTranspose(pSrcData, cDstData, ImgSrc->width, ImgSrc->height);  //CPU進行映像轉置clock_t t2 = clock();float time_cpu = 0;time_cpu = t2 - t1;printf("CPU_Time = %f\n", time_cpu*1000/CLOCKS_PER_SEC);  //時間單位ms//釋放資源cvNamedWindow("Src");cvShowImage("Src", ImgSrc);cvNamedWindow("Dst_CPU");cvShowImage("Dst_CPU", ImgDst_CPU);cvNamedWindow("Dst_GPU_Global");cvShowImage("Dst_GPU_Global", ImgDst_GPU_Global);cvWaitKey();cudaFree(device_ImgDataSrc);cudaFree(device_ImgDataDst_Global);cvDestroyAllWindows();cvReleaseImage(&ImgSrc);cvReleaseImage(&ImgDst_CPU);cvReleaseImage(&ImgDst_GPU_Global);return 0;}

（3）當對CUDA程式進行最佳化，使用GPU Shared Memory進行映像轉置時，”imageTranspose.cu“檔案要進行相應的修改：

int main(){IplImage *ImgSrc = cvLoadImage("<span style="text-align: justify;">Lena.jpg</span>", CV_LOAD_IMAGE_GRAYSCALE);int Width = ImgSrc->width;int Height = ImgSrc->height;//輸出映像的寬高尺寸互換IplImage *ImgDst_GPU_Shared = cvCreateImage(cvSize(Height, Width), IPL_DEPTH_8U, 1);IplImage *ImgDst_CPU = cvCreateImage(cvSize(Height, Width), IPL_DEPTH_8U, 1);unsigned char *pSrcData = (unsigned char*)(ImgSrc->imageData);unsigned char *pDstData_Shared = (unsigned char*)(ImgDst_GPU_Shared->imageData);unsigned char *cDstData = (unsigned char*)(ImgDst_CPU->imageData);//分配顯存用於儲存原映像數組和靶心圖表像數組unsigned char *device_ImgDataSrc = NULL;unsigned char *device_ImgDataDst_Shared = NULL;cudaMalloc((void**)&device_ImgDataSrc, sizeof(unsigned char) * Width * Height);cudaMalloc((void**)&device_ImgDataDst_Shared, sizeof(unsigned char) * Height * Width);//將原映像數組傳遞到顯存中cudaMemcpy(device_ImgDataSrc, pSrcData, sizeof(unsigned char) * Width * Height, cudaMemcpyHostToDevice);//SharedMemory版本的參數設定dim3 dimGrid_Shared(N/W, N/W);  //每個Grid允許的最大Block數為65535dim3 dimBlock_Shared(W, W);  //每個Block允許的最大線程數為512cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);GPUImageTranspose_Shared<<<dimGrid_Shared, dimBlock_Shared>>>(device_ImgDataSrc, device_ImgDataDst_Shared,  ImgSrc->width, ImgSrc->height);cudaEventRecord(stop, 0);cudaEventSynchronize(start);cudaEventSynchronize(stop);float GPUTime_Shared = 0;cudaEventElapsedTime(&GPUTime_Shared, start, stop);printf("GPU_Time_Shared = %f\n",  GPUTime_Shared);//將結果傳遞至記憶體cudaMemcpy(pDstData_Shared, device_ImgDataDst_Shared, sizeof(unsigned char) * Width * Height, cudaMemcpyDeviceToHost);clock_t t1 = clock();CPUImageTranspose(pSrcData, cDstData, ImgSrc->width, ImgSrc->height);  //CPU處理的映像clock_t t2 = clock();float time_cpu = 0;time_cpu = t2 - t1;printf("CPU_Time = %f\n", time_cpu*1000/CLOCKS_PER_SEC);  //時間單位mscvNamedWindow("Src");cvShowImage("Src", ImgSrc);cvNamedWindow("Dst_CPU");cvShowImage("Dst_CPU", ImgDst_CPU);cvNamedWindow("Dst_GPU_Shared");cvShowImage("Dst_GPU_Shared", ImgDst_GPU_Shared);cvWaitKey();cudaFree(device_ImgDataSrc);cudaFree(device_ImgDataDst_Shared);cvDestroyAllWindows();cvReleaseImage(&ImgSrc);cvReleaseImage(&ImgDst_CPU);cvReleaseImage(&ImgDst_GPU_Shared);return 0;}

4. 實驗結果

（1）採用經典測試映像”Lena.jpg“作為輸入映像：

（2）CPU、GPU Global Memory運行結果及時間比較

（3）GPU Shared Memory運行結果及時間

5. 實驗結果分析

經過測試，由輸出映像可以判斷，映像轉置的結果是正確的，耗時GPU Shared Memory < GPU Global Memory < CPU，但是比較不解的是兩次運行CPU的時間都不同，這個不知道是什麼原因。。。