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