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