CUDA學習筆記一：CUDA+OpenCV的影象轉置，採用Shared Memory進行CUDA程式優化
阿新 • • 發佈：2019-01-01
原創文章，轉載請註明出處......
一、問題背景
最近要做一個關於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 Memory
tile[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數為65535
int 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;
}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數為65535
dim3 dimBlock_Shared(W, W); //每個Block允許的最大執行緒數為512
cudaEvent_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); //時間單位ms
cvNamedWindow("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的時間都不同，這個不知道是什麼原因。。。