CUDA JPEG编码
Posted 啊基米舍的博客
tags:
篇首语：本文由小常识网(cha138.com)小编为大家整理，主要介绍了CUDA JPEG编码相关的知识，希望对你有一定的参考价值。
基于英伟达的jpegNPP工程，分离实现独立的JPEG压缩。
由于原工程是直接把解码时的jpeg图片的信息直接作为编码时的信息，所以在做独立的JPEG编码时，需要自己来填充各种信息。
1.JPEG编码流程
从网上一片文章中解出来的图，红色框框中的流程图算是JPEG编码的一个流程图，对JPEG编码流程的了解有助于对代码的理解。
2.Huffman表和量化表
Huffman表和量化表采用标准Huffman表和标准量化表。
//标准亮度信号量化模板 static unsigned char std_Y_QT[64] = { 16, 11, 10, 16, 24, 40, 51, 61, 12, 12, 14, 19, 26, 58, 60, 55, 14, 13, 16, 24, 40, 57, 69, 56, 14, 17, 22, 29, 51, 87, 80, 62, 18, 22, 37, 56, 68, 109,103,77, 24, 35, 55, 64, 81, 104,113,92, 49, 64, 78, 87, 103,121,120,101, 72, 92, 95, 98, 112,100,103,99 }; //标准色差信号量化模板 static unsigned char std_UV_QT[64] = { 17, 18, 24, 47, 99, 99, 99, 99, 18, 21, 26, 66, 99, 99, 99, 99, 24, 26, 56, 99, 99, 99, 99, 99, 47, 66, 99 ,99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99, 99 }; // 标准Huffman表 (cf. JPEG standard section K.3) static unsigned char STD_DC_Y_NRCODES[17]={0, 0, 1, 5, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0}; static unsigned char STD_DC_Y_VALUES[12] ={0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; static unsigned char STD_DC_UV_NRCODES[17]={0, 0, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0}; static unsigned char STD_DC_UV_VALUES[12] ={0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; static unsigned char STD_AC_Y_NRCODES[17]={0,0,2,1,3,3,2,4,3,5,5,4,4,0,0,1,0X7D}; static unsigned char STD_AC_Y_VALUES[162]= { 0x01, 0x02, 0x03, 0x00, 0x04, 0x11, 0x05, 0x12, 0x21, 0x31, 0x41, 0x06, 0x13, 0x51, 0x61, 0x07, 0x22, 0x71, 0x14, 0x32, 0x81, 0x91, 0xa1, 0x08, 0x23, 0x42, 0xb1, 0xc1, 0x15, 0x52, 0xd1, 0xf0, 0x24, 0x33, 0x62, 0x72, 0x82, 0x09, 0x0a, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x25, 0x26, 0x27, 0x28, 0x29, 0x2a, 0x34, 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x43, 0x44, 0x45, 0x46, 0x47, 0x48, 0x49, 0x4a, 0x53, 0x54, 0x55, 0x56, 0x57, 0x58, 0x59, 0x5a, 0x63, 0x64, 0x65, 0x66, 0x67, 0x68, 0x69, 0x6a, 0x73, 0x74, 0x75, 0x76, 0x77, 0x78, 0x79, 0x7a, 0x83, 0x84, 0x85, 0x86, 0x87, 0x88, 0x89, 0x8a, 0x92, 0x93, 0x94, 0x95, 0x96, 0x97, 0x98, 0x99, 0x9a, 0xa2, 0xa3, 0xa4, 0xa5, 0xa6, 0xa7, 0xa8, 0xa9, 0xaa, 0xb2, 0xb3, 0xb4, 0xb5, 0xb6, 0xb7, 0xb8, 0xb9, 0xba, 0xc2, 0xc3, 0xc4, 0xc5, 0xc6, 0xc7, 0xc8, 0xc9, 0xca, 0xd2, 0xd3, 0xd4, 0xd5, 0xd6, 0xd7, 0xd8, 0xd9, 0xda, 0xe1, 0xe2, 0xe3, 0xe4, 0xe5, 0xe6, 0xe7, 0xe8, 0xe9, 0xea, 0xf1, 0xf2, 0xf3, 0xf4, 0xf5, 0xf6, 0xf7, 0xf8, 0xf9, 0xfa }; static unsigned char STD_AC_UV_NRCODES[17]={0,0,2,1,2,4,4,3,4,7,5,4,4,0,1,2,0X77}; static unsigned char STD_AC_UV_VALUES[162]= { 0x00, 0x01, 0x02, 0x03, 0x11, 0x04, 0x05, 0x21, 0x31, 0x06, 0x12, 0x41, 0x51, 0x07, 0x61, 0x71, 0x13, 0x22, 0x32, 0x81, 0x08, 0x14, 0x42, 0x91, 0xa1, 0xb1, 0xc1, 0x09, 0x23, 0x33, 0x52, 0xf0, 0x15, 0x62, 0x72, 0xd1, 0x0a, 0x16, 0x24, 0x34, 0xe1, 0x25, 0xf1, 0x17, 0x18, 0x19, 0x1a, 0x26, 0x27, 0x28, 0x29, 0x2a, 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x43, 0x44, 0x45, 0x46, 0x47, 0x48, 0x49, 0x4a, 0x53, 0x54, 0x55, 0x56, 0x57, 0x58, 0x59, 0x5a, 0x63, 0x64, 0x65, 0x66, 0x67, 0x68, 0x69, 0x6a, 0x73, 0x74, 0x75, 0x76, 0x77, 0x78, 0x79, 0x7a, 0x82, 0x83, 0x84, 0x85, 0x86, 0x87, 0x88, 0x89, 0x8a, 0x92, 0x93, 0x94, 0x95, 0x96, 0x97, 0x98, 0x99, 0x9a, 0xa2, 0xa3, 0xa4, 0xa5, 0xa6, 0xa7, 0xa8, 0xa9, 0xaa, 0xb2, 0xb3, 0xb4, 0xb5, 0xb6, 0xb7, 0xb8, 0xb9, 0xba, 0xc2, 0xc3, 0xc4, 0xc5, 0xc6, 0xc7, 0xc8, 0xc9, 0xca, 0xd2, 0xd3, 0xd4, 0xd5, 0xd6, 0xd7, 0xd8, 0xd9, 0xda, 0xe2, 0xe3, 0xe4, 0xe5, 0xe6, 0xe7, 0xe8, 0xe9, 0xea, 0xf2, 0xf3, 0xf4, 0xf5, 0xf6, 0xf7, 0xf8, 0xf9, 0xfa };
2.RGB转YUV
编码的输入时YUV数据，所以对于原始数据需要先把RGB转成YUV。要说的是，我这里输入的RGB数据的存储形式为BBBBBBB……….GGGGGGG………RRRRRRR…..，其他形式的RGB数据需要对kernel函数中的坐标做点调整。
RGB转YUV的公式网上有许多中，我这里采用：
Y = 0.299R + 0.587G + 0.114BU = -0.147R - 0.289G + 0.436BV = 0.615R - 0.515G - 0.100B
代码：
#include "cuda_kernels.h" typedef unsigned char uint8; typedef unsigned int uint32; typedef int int32; namespace cuda_common { __device__ unsigned char clip_value(unsigned char x, unsigned char min_val, unsigned char max_val){ if (x>max_val){ return max_val; } else if (x<min_val){ return min_val; } else{ return x; } } extern "C" __global__ void kernel_rgb2yuv(float *src_img, unsigned char* Y, unsigned char* u, unsigned char* v, int src_width, int src_height, size_t yPitch) { const int x = blockIdx.x * blockDim.x + threadIdx.x; const int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= src_width) return; //x = width - 1; if (y >= src_height) return; // y = height - 1; float B = src_img[y * src_width + x]; float G = src_img[src_width * src_height + y * src_width + x]; float R = src_img[src_width * src_height * 2 + y * src_width + x]; Y[y * yPitch + x] = clip_value((unsigned char)(0.299 * R + 0.587 * G + 0.114 * B), 0, 255); u[y * src_width + x] = clip_value((unsigned char)(-0.147 * R - 0.289 * G + 0.436 * B + 128), 0, 255); v[y * src_width + x] = clip_value((unsigned char)(0.615 * R - 0.515 * G - 0.100 * B + 128), 0, 255); //Y[y * yPitch + x] = clip_value((unsigned char)(0.257 * R + 0.504 * G + 0.098 * B + 16), 0, 255); //u[y * src_width + x] = clip_value((unsigned char)(-0.148 * R - 0.291 * G + 0.439 * B + 128), 0, 255); //v[y * src_width + x] = clip_value((unsigned char)(0.439 * R - 0.368 * G - 0.071 * B + 128), 0, 255); } extern "C" __global__ void kernel_resize_UV(unsigned char* src_img, unsigned char *dst_img, int src_width, int src_height, int dst_width, int dst_height, int nPitch) { const int x = blockIdx.x * blockDim.x + threadIdx.x; const int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= dst_width) return; //x = width - 1; if (y >= dst_height) return; // y = height - 1; float fx = (x + 0.5)*src_width / (float)dst_width - 0.5; float fy = (y + 0.5)*src_height / (float)dst_height - 0.5; int ax = floor(fx); int ay = floor(fy); if (ax < 0) { ax = 0; } else if (ax > src_width - 2) { ax = src_width - 2; } if (ay < 0){ ay = 0; } else if (ay > src_height - 2) { ay = src_height - 2; } int A = ax + ay*src_width; int B = ax + ay*src_width + 1; int C = ax + ay*src_width + src_width; int D = ax + ay*src_width + src_width + 1; float w1, w2, w3, w4; w1 = fx - ax; w2 = 1 - w1; w3 = fy - ay; w4 = 1 - w3; unsigned char val = src_img[A] * w2*w4 + src_img[B] * w1*w4 + src_img[C] * w2*w3 + src_img[D] * w1*w3; dst_img[y * nPitch + x] = clip_value(val,0,255); } cudaError_t RGB2YUV(float* d_srcRGB, int src_width, int src_height, unsigned char* Y, size_t yPitch, int yWidth, int yHeight, unsigned char* U, size_t uPitch, int uWidth, int uHeight, unsigned char* V, size_t vPitch, int vWidth, int vHeight) { unsigned char * u ; unsigned char * v ; cudaError_t cudaStatus; cudaStatus = cudaMalloc((void**)&u, src_width * src_height * sizeof(unsigned char)); cudaStatus = cudaMalloc((void**)&v, src_width * src_height * sizeof(unsigned char)); dim3 block(32, 16, 1); dim3 grid((src_width + (block.x - 1)) / block.x, (src_height + (block.y - 1)) / block.y, 1); dim3 grid1((uWidth + (block.x - 1)) / block.x, (uHeight + (block.y - 1)) / block.y, 1); dim3 grid2((vWidth + (block.x - 1)) / block.x, (vHeight + (block.y - 1)) / block.y, 1); kernel_rgb2yuv << < grid, block >> >(d_srcRGB, Y, u, v, src_width, src_height, yPitch); cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "kernel_rgb2yuv launch failed: %s\\n", cudaGetErrorString(cudaStatus)); goto Error; } cudaStatus = cudaDeviceSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching kernel_rgb2yuv!\\n", cudaStatus); goto Error; } kernel_resize_UV << < grid1, block >> >(u, U, src_width, src_height, uWidth, uHeight, uPitch); cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "kernel_resize_UV launch failed: %s\\n", cudaGetErrorString(cudaStatus)); goto Error; } cudaStatus = cudaDeviceSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching kernel_resize_UV!\\n", cudaStatus); goto Error; } kernel_resize_UV << < grid2, block >> >(v, V, src_width, src_height, vWidth, vHeight, vPitch); cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "kernel_resize_UV launch failed: %s\\n", cudaGetErrorString(cudaStatus)); goto Error; } cudaStatus = cudaDeviceSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching kernel_resize_UV!\\n", cudaStatus); goto Error; } Error : cudaFree(u); cudaFree(v); return cudaStatus; } }
RGB转码分了两步，第一步是按照RGB转YUV公式做转换，由kernel_rgb2yuv函数来完成；第二部把U、V的宽高缩小到原来的一半，由kernel_resize_UV函数来做，这个地方会影响到最终压缩后的图片的效果，我这里用的是一个双线性插值来缩小的，效果凑合，但和原图一比较还是能看出明显的色差，网上查到怎么缩小U、V的，这里的做法可能与标准做法不符，有待更好的方法。
3.JPEG压缩
jpegNPP工程包含了一个解码然后编码的过程，要分离出解码的话把编码部分去掉就行了，应该来说是比较容易的；要分离出编码部分就有点困难了，因为工程中编码所使用的参数是直接使用的解码读取到的图片的参数，所以要分离出编码，就需要自己来填充这些参数，这对于不懂JPEG压缩的来说就有点难度了，好在网上资料多。
int jpegNPP(const char *szOutputFile, float* d_srcRGB, int img_width, int img_height) { NppiDCTState *pDCTState; NPP_CHECK_NPP(nppiDCTInitAlloc(&pDCTState)); // Parsing and Huffman Decoding (on host) FrameHeader oFrameHeader; QuantizationTable aQuantizationTables[4]; Npp8u *pdQuantizationTables; cudaMalloc(&pdQuantizationTables, 64 * 4); HuffmanTable aHuffmanTables[4]; HuffmanTable *pHuffmanDCTables = aHuffmanTables; HuffmanTable *pHuffmanACTables = &aHuffmanTables[2]; ScanHeader oScanHeader; memset(&oFrameHeader,0,sizeof(FrameHeader)); memset(aQuantizationTables,0, 4 * sizeof(QuantizationTable)); memset(aHuffmanTables,0, 4 * sizeof(HuffmanTable)); int nMCUBlocksH = 0; int nMCUBlocksV = 0; int nRestartInterval = -1; NppiSize aSrcSize[3]; Npp16s *apdDCT[3] = {0,0,0}; Npp32s aDCTStep[3]; Npp8u *apSrcImage[3] = {0,0,0}; Npp32s aSrcImageStep[3]; size_t aSrcPitch[3]; /*************************** * * Output 编码部分 * ***************************/ unsigned char STD_DC_Y_NRCODES[16] = { 0, 1, 5, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0 }; unsigned char STD_DC_Y_VALUES[12] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11 }; unsigned char STD_DC_UV_NRCODES[16] = { 0, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0 }; unsigned char STD_DC_UV_VALUES[12] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11 }; unsigned char STD_AC_Y_NRCODES[16] = { 0,以上是关于CUDA JPEG编码的主要内容，如果未能解决你的问题，请参考以下文章
RuntimeError: ‘lengths’ argument should be a 1D CPU int64 tensor, but got 1D cuda:0 Long tensor(代码片段