CUDA 實現JPEG圖像解碼為RGB數據

    了解JPEG數據格式的人應該easy想到。其對圖像以8*8像素塊大小進行切割壓縮的方法非常好用並行處理的思想來實現。而其實英偉達的CUDA自v5.5開始也提供了JPEG編解碼的演示樣例。該演示樣例存儲在CUDA的SDK中，即CUDA的默認安裝路徑“C:\ProgramData\NVDIA Corporation\CUDA Samples\v7.0\7_CUDALibraries\jpegNPP”(v后面的數字依據版本號的不同會變更)中。

    該演示樣例將圖片數據進行了解碼和再編碼，因為解碼僅僅是將數據轉為YUV，我們假設要利用演示樣例來將圖像轉為RGB數據還需進行YUV->RGB的轉換工作。這也正是本篇文章要重點介紹的內容。

此外。因為演示樣例本身存在一個bug，所以無法直接利用其解碼壓縮寬高比不同的圖像，這個會在下文再次提到，並給出比較取巧的修復方法。這個bug已經上報給英偉達。英偉達回復將在下個版本號（也就是v7.0之后的版本號）修復這個bug。

    轉載請注明出處：http://blog.csdn.net/weixinhum/article/details/46683509

    OK。以下就開始吧

    因為我們須要改動DEMO的源代碼。還是先到上面的路徑下將jpegNPP目錄備份出一份來。然后我們直接打開目錄里面的vsproject。project的主要代碼在jpegNPP.cpp中，到了

// Inverse DCT
for (int i = 0; i < 3; ++i)
{
    NPP_CHECK_NPP(nppiDCTQuantInv8x8LS_JPEG_16s8u_C1R_NEW(apdDCT[i], aDCTStep[i],
                                                          apSrcImage[i], aSrcImageStep[i],
                                                          pdQuantizationTables + oFrameHeader.aQuantizationTableSelector[i] * 64,
                                                          aSrcSize[i],
                                                          pDCTState));
}

這段代碼便已經實現了將JPEG圖像進行解碼轉化為YUV數據的功能。YUV數據存儲在apSrcImage[0],apSrcImage[1],apSrcImage[2] 中，而其步長(通道寬度)分別存在aSrcImageStep[0],aSrcImageStep[1],aSrcImageStep[2] 中，已知條件已經足夠了。我們能夠直接刪掉上述所貼代碼后面的全部代碼(那部分代碼是關於圖像編碼的)，然后寫一個CUDA處理函數將YUV轉為RGB。

    大致流程例如以下：

    配置OpenCV環境並包括頭文件（這一步並非必要的。僅僅是為了方便查看我們轉出來的圖像是不是對的，假設認為不是必需能夠忽略掉。僅僅要知道輸出RGB的數據指針和數據長度大小便可）：

#include <opencv2/core/core.hpp>//OpenCV包括頭文件  
#include <opencv2/highgui/highgui.hpp> 
#include <opencv2/opencv.hpp> 
using namespace std;

    編寫代碼實現YUV轉RGB：

    在上述所貼DEMOproject代碼后面加上例如以下代碼：

int pwidth = aSrcSize[0].width;
int pheight = aSrcSize[0].height;

IplImage *drawimg;//數據輸出圖像
drawimg = cvCreateImage(cvSize(pwidth, pheight), 8, 3);

Npp8u *Host_img;//主機內存
Npp8u *Device_img;//顯卡內存
size_t mPitch;
NPP_CHECK_CUDA(cudaMallocPitch(&Device_img, &mPitch, pwidth * 3, pheight));//開辟顯存空間以存儲RGB數據

//unsigned char* imgdata = (unsigned char*)drawimg->imageData;
YCrCb2RGB(apSrcImage[0], apSrcImage[1], apSrcImage[2], pwidth, pheight, aSrcImageStep[0],
	aSrcImageStep[1], aSrcImageStep[2], drawimg->widthStep / sizeof(uchar), Device_img, nMCUBlocksV, nMCUBlocksH);

NPP_CHECK_CUDA(cudaHostAlloc(&Host_img, pwidth*pheight * 3, cudaHostAllocDefault));//分配主機鎖頁內存
NPP_CHECK_CUDA(cudaMemcpy(Host_img, Device_img, pwidth*pheight * 3, cudaMemcpyDeviceToHost));//拷貝顯卡處理完圖像到主機
drawimg->imageData = (char*)Host_img;

cvShowImage("", drawimg);
cvWaitKey(0);
getchar();
for (int i = 0; i < 3; ++i)//內存釋放
{
	cudaFree(apSrcImage[i]);
	cudaFree(apdDCT[i]);
	cudaFreeHost(aphDCT[i]);
}
cudaFree(Device_img);
cudaFreeHost(Host_img);
cudaDeviceReset();
return EXIT_SUCCESS;

    加入一個“CudaYCrCb.cu”文件來定義YCrCb2RGB函數的功能，至於怎樣去設置.cu文件假設有疑問的話請參照之前的這篇文章的相關內容，此外YCrCb2RGB函數須要在jpegNPP.cpp文件頭聲明下。文件的內容例如以下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "Endianess.h"

__device__ unsigned char judge(int value)
{
	if (value >= 0 && value <= 255)
	{
		return value;
	}
	else if (value>255)
	{
		return 255;
	}
	else
	{
		return 0;
	}
}

__global__ void YCrCb2RGBConver(unsigned char *Device_Y, unsigned char *Device_Cr, unsigned char *Device_Cb, unsigned char *Device_img, int width, int height, int YStep, int CrStep, int CbStep, int img_step, int nMCUBlocksV, int nMCUBlocksH)//處理核函數
{
	//int tid = blockIdx.x*blockDim.x + threadIdx.x;
	int row = blockIdx.y*blockDim.y + threadIdx.y;
	int cols = blockIdx.x*blockDim.x + threadIdx.x;

	if (row >= height)
	{
		return;
	}
	if (cols >= width)
	{
		return;
	}

	int Y = Device_Y[row*YStep + cols];
	int U = Device_Cr[row / nMCUBlocksH*CrStep + cols / nMCUBlocksV] - 128;
	int V = Device_Cb[row / nMCUBlocksH*CbStep + cols / nMCUBlocksV] - 128;

	Device_img[row*img_step + cols * 3 + 0] =
		judge(Y + U + ((U * 198) >> 8));
	Device_img[row*img_step + cols * 3 + 1] =
		judge(Y - (((U * 88) >> 8) + ((V * 183) >> 8)));
	Device_img[row*img_step + cols * 3 + 2] =
		judge(Y + V + ((V * 103) >> 8));
}

extern "C" int YCrCb2RGB(unsigned char *Device_Y, unsigned char *Device_Cr, unsigned char *Device_Cb, int width, int height, int YStep, int CrStep, int CbStep, int img_step, unsigned char *Device_data, int nMCUBlocksV, int nMCUBlocksH)//顯卡處理函數 
{
	cudaEvent_t start, stop;
	float time;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);
	//這個部分可調
	dim3 threads(16, 16);//線程塊中的線程數1*1
	//dim3 threads(256, 40);//線程塊中的線程數1*1
	dim3 blocks((width + threads.x - 1) / threads.x, (height + threads.y - 1) / threads.y);//線程塊大小
	YCrCb2RGBConver << <blocks, threads >> >(Device_Y, Device_Cr, Device_Cb, Device_data, width, height, YStep, CrStep, CbStep, img_step, nMCUBlocksV, nMCUBlocksH);//調用顯卡處理數據
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&time, start, stop);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	printf("核函數消耗時間:%f\n", time);
	return 0;
}

    聲明例如以下：

extern "C" int YCrCb2RGB(unsigned char *Device_Y, unsigned char *Device_Cr, unsigned char *Device_Cb,int width, int height, int YStep, int CrStep, int CbStep, int img_step, unsigned char *Device_data, int nMCUBlocksV, int nMCUBlocksH);//顯卡處理函數 

    到此，實現了文章題目的內容，對於前文提到的英偉達的DEMO本身存在的bug（ 解碼壓縮寬高比不同的圖像內存報錯），是因為壓縮的寬高比弄錯導致的，能夠通過例如以下截圖的方式進行改動。