CUDA 实现JPEG图像解码为RGB数据

    了解JPEG数据格式的人应该容易想到，其对图像以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文件夹备份出一份来。然后我们直接打开文件夹里面的vs工程。工程的主要代码在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：

    在上述所贴DEMO工程代码后面加上如下代码：

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（ 解码压缩宽高比不同的图像内存报错），是由于压缩的宽高比弄错导致的，可以通过如下截图的方式进行修改。