CUDA实现JPEG编码

最新推荐文章于 2024-08-24 09:43:14 发布
最新推荐文章于 2024-08-24 09:43:14 发布
阅读量2.2k 收藏 3
点赞数
分类专栏： gpu

CUDA实现JPEG编码

https://blog.csdn.net/mikedai/article/details/79084092

2018年01月17日 13:11:41 MikeDai 阅读数 1963

 版权声明：本文为博主原创文章，未经博主允许不得转载。 https://blog.csdn.net/MikeDai/article/details/79084092

JPEG编码流程图： 
这里写图片描述 
详细可以参考：www.eecg.toronto.edu/~moshovos/CUDA08/arx/JPEG_report.pdf 
代码参考nvidia提供的例子，不过那个例子是先读取一张jpg图片，然后缩放（处理的是YV12格式的数据），最后编码。很多东西都混在一起，不便宜使用。现在就提取其中有用的，去掉缩放的功能。

jpeg数据结构定义：
//量化表数据
struct QuantizationTable
{
    unsigned char nPrecisionAndIdentifier;
    unsigned char aTable[64];
};

//图片信息
struct FrameHeader
{
    unsigned char nSamplePrecision;
    unsigned short nHeight;
    unsigned short nWidth;
    unsigned char nComponents;
    unsigned char aComponentIdentifier[3];
    unsigned char aSamplingFactors[3];
    unsigned char aQuantizationTableSelector[3];
};

//扫描头
struct ScanHeader
{
    unsigned char nComponents;
    unsigned char aComponentSelector[3];
    unsigned char aHuffmanTablesSelector[3];
    unsigned char nSs;
    unsigned char nSe;
    unsigned char nA;
};
//霍夫曼编码表数据，一般是固定的
struct HuffmanTable
{
    unsigned char nClassAndIdentifier;
    unsigned char aCodes[16];
    unsigned char aTable[256];
};
  •  

类封装，这里也只是处理yuv数据

class CudaJpegEncode
{
public:
    CudaJpegEncode();
    ~CudaJpegEncode();
public:
    void Init();
    void Uninit();
public:
    void EncodeJpeg(const char* output_file);
    void SetData(uint8_t* yuv_data, uint32_t data_size, int yuv_fmt, int w, int h);
public:
    NppiDCTState*        pDCTState;
    FrameHeader          oFrameHeader;
    QuantizationTable    aQuantizationTables[4];
    Npp8u*               pdQuantizationTables;

    HuffmanTable         aHuffmanTables[4];
    HuffmanTable*        pHuffmanDCTables = aHuffmanTables;
    HuffmanTable*        pHuffmanACTables = &aHuffmanTables[2];
    ScanHeader           oScanHeader;

    NppiSize aSrcSize[3];
    Npp16s *apdDCT[3];
    Npp32s aDCTStep[3];

    Npp8u *apSrcImage[3];
    Npp32s aSrcImageStep[3];
    size_t aSrcPitch[3];

    NppiEncodeHuffmanSpec *apHuffmanDCTable[3];
    NppiEncodeHuffmanSpec *apHuffmanACTable[3];

    int nMCUBlocksH = 0;
    int nMCUBlocksV = 0;

    Npp8u *pdScan;
    Npp32s nScanSize;

    Npp8u *pJpegEncoderTemp;
    size_t nTempSize;

    unsigned char *pDstJpeg;
public:
    uint8_t*    mY;
    uint8_t*    mU;
    uint8_t*    mV;
    int         mWidth;
    int         mHeight;
};
  •  

类实现：

支持的最大图片尺寸

#define  IMG_WIDTH_MAX  1920
#define  IMG_HEIGHT_MAX 1080

#define  IMG_WIDTH_MAX_HALF 960
#define  IMG_HEIGHT_MAX_HALF    540
#define  MEMORY_ALGN_DEVICE  511

#define  YUV_FMT_NV12           1
#define  YUV_FMT_NV21           2
#define  YUV_FMT_YUV420         3

#ifndef max 
#define max(X, Y) ((X) > (Y) ?(X):(Y))
#endif 
//霍夫曼编码表
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, 2, 1, 3, 3, 2, 4, 3, 5, 5, 4, 4, 0, 0, 1, 0X7D };
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
};

unsigned char STD_AC_UV_NRCODES[16] = { 0, 2, 1, 2, 4, 4, 3, 4, 7, 5, 4, 4, 0, 1, 2, 0X77 };
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
};
//量化表，可以根据图像质量要求更改，这里图像是90%
unsigned char std_Y_QT[64] =
{
    3, 2, 2, 3, 2, 2, 3, 3,
    3, 3, 4, 3, 3, 4, 5, 8,
    5, 5, 4, 4, 5, 10, 7, 7,
    6, 8, 12, 10, 12, 12, 11, 10,
    11, 11, 13, 14, 18, 16, 13, 14,
    17, 14, 11, 11, 16, 22, 16, 17,
    19, 20, 21, 21, 21, 12, 15, 23,
    24, 22, 20, 24, 18, 20, 21, 20
};

unsigned char std_UV_QT[64] =
{
    3, 4, 4, 5, 4, 5, 9, 5,
    5, 9, 20, 13, 11, 13, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20,
    20, 20, 20, 20, 20, 20, 20, 20
};

int DivUp(int x, int d)
{
    return (x + d - 1) / d;
}

void writeMarker(unsigned char nMarker, unsigned char *&pData)
{
    *pData++ = 0x0ff;
    *pData++ = nMarker;
}

template<class T>
void writeBigEndian(unsigned char *pData, T value)
{
    unsigned char *pValue = reinterpret_cast<unsigned char *>(&value);
    reverse_copy(pValue, pValue + sizeof(T), pData);
}


template<typename T>
void writeAndAdvance(unsigned char *&pData, T nElement)
{
    writeBigEndian<T>(pData, nElement);
    pData += sizeof(T);
}

void writeJFIFTag(unsigned char *&pData)
{
    const char JFIF_TAG[] =
    {
        0x4a, 0x46, 0x49, 0x46, 0x00,
        0x01, 0x02,
        0x00,
        0x00, 0x01, 0x00, 0x01,
        0x00, 0x00
    };

    writeMarker(0x0e0, pData);
    writeAndAdvance<unsigned short>(pData, sizeof(JFIF_TAG) + sizeof(unsigned short));
    memcpy(pData, JFIF_TAG, sizeof(JFIF_TAG));
    pData += sizeof(JFIF_TAG);
}

void writeQuantizationTable(const QuantizationTable &table, unsigned char *&pData)
{
    writeMarker(0x0DB, pData);
    writeAndAdvance<unsigned short>(pData, sizeof(QuantizationTable) + 2);
    memcpy(pData, &table, sizeof(QuantizationTable));
    pData += sizeof(QuantizationTable);
}


void writeHuffmanTable(const HuffmanTable &table, unsigned char *&pData)
{
    writeMarker(0x0C4, pData);

    // Number of Codes for Bit Lengths [1..16]
    int nCodeCount = 0;

    for (int i = 0; i < 16; ++i)
    {
        nCodeCount += table.aCodes[i];
    }

    writeAndAdvance<unsigned short>(pData, 17 + nCodeCount + 2);
    memcpy(pData, &table, 17 + nCodeCount);
    pData += 17 + nCodeCount;
}

void writeFrameHeader(const FrameHeader &header, unsigned char *&pData)
{
    unsigned char aTemp[128];
    unsigned char *pTemp = aTemp;

    writeAndAdvance<unsigned char>(pTemp, header.nSamplePrecision);
    writeAndAdvance<unsigned short>(pTemp, header.nHeight);
    writeAndAdvance<unsigned short>(pTemp, header.nWidth);
    writeAndAdvance<unsigned char>(pTemp, header.nComponents);

    for (int c = 0; c < header.nComponents; ++c)
    {
        writeAndAdvance<unsigned char>(pTemp, header.aComponentIdentifier[c]);
        writeAndAdvance<unsigned char>(pTemp, header.aSamplingFactors[c]);
        writeAndAdvance<unsigned char>(pTemp, header.aQuantizationTableSelector[c]);
    }

    unsigned short nLength = (unsigned short)(pTemp - aTemp);

    writeMarker(0x0C0, pData);
    writeAndAdvance<unsigned short>(pData, nLength + 2);
    memcpy(pData, aTemp, nLength);
    pData += nLength;
}

void writeScanHeader(const ScanHeader &header, unsigned char *&pData)
{
    unsigned char aTemp[128];
    unsigned char *pTemp = aTemp;

    writeAndAdvance<unsigned char>(pTemp, header.nComponents);

    for (int c = 0; c < header.nComponents; ++c)
    {
        writeAndAdvance<unsigned char>(pTemp, header.aComponentSelector[c]);
        writeAndAdvance<unsigned char>(pTemp, header.aHuffmanTablesSelector[c]);
    }

    writeAndAdvance<unsigned char>(pTemp, header.nSs);
    writeAndAdvance<unsigned char>(pTemp, header.nSe);
    writeAndAdvance<unsigned char>(pTemp, header.nA);

    unsigned short nLength = (unsigned short)(pTemp - aTemp);

    writeMarker(0x0DA, pData);
    writeAndAdvance<unsigned short>(pData, nLength + 2);
    memcpy(pData, aTemp, nLength);
    pData += nLength;
}


CudaJpegEncode::CudaJpegEncode()
{
    pDCTState = NULL;

    nMCUBlocksH = 0;
    nMCUBlocksV = 0;

    pdScan = NULL;
    nScanSize = 0;

    pdQuantizationTables = NULL;
    pJpegEncoderTemp = NULL;
    nTempSize = 0;

    pDstJpeg = NULL;
    mY = NULL;
    mU = NULL;
    mV = NULL;
}


CudaJpegEncode::~CudaJpegEncode()
{
    Uninit();
}

/*
brief: init the jpeg encoder
*/
void CudaJpegEncode::Init()
{
    NPP_CHECK_NPP(nppiDCTInitAlloc(&pDCTState));
    cudaMalloc(&pdQuantizationTables, 64 * 4);
    memset(&oFrameHeader, 0, sizeof(FrameHeader));
    memset(aQuantizationTables, 0, 4 * sizeof(QuantizationTable));
    memset(aHuffmanTables, 0, 4 * sizeof(HuffmanTable));

    //填充Huffman表
    aHuffmanTables[0].nClassAndIdentifier = 0;
    memcpy(aHuffmanTables[0].aCodes, STD_DC_Y_NRCODES, 16);
    memcpy(aHuffmanTables[0].aTable, STD_DC_Y_VALUES, 12);

    aHuffmanTables[1].nClassAndIdentifier = 1;
    memcpy(aHuffmanTables[1].aCodes, STD_DC_UV_NRCODES, 16);
    memcpy(aHuffmanTables[1].aTable, STD_DC_UV_VALUES, 12);

    aHuffmanTables[2].nClassAndIdentifier = 16;
    memcpy(aHuffmanTables[2].aCodes, STD_AC_Y_NRCODES, 16);
    memcpy(aHuffmanTables[2].aTable, STD_AC_Y_VALUES, 162);

    aHuffmanTables[3].nClassAndIdentifier = 17;
    memcpy(aHuffmanTables[3].aCodes, STD_AC_UV_NRCODES, 16);
    memcpy(aHuffmanTables[3].aTable, STD_AC_UV_VALUES, 162);

    aQuantizationTables[0].nPrecisionAndIdentifier = 0;
    memcpy(aQuantizationTables[0].aTable, std_Y_QT, 64);
    aQuantizationTables[1].nPrecisionAndIdentifier = 1;
    memcpy(aQuantizationTables[1].aTable, std_UV_QT, 64);

    // Copy DCT coefficients and Quantization Tables from host to device 
    //拷贝时之字型扫描不可少，否则会出现一下人为畸变
    Npp8u aZigzag[] = {
        0, 1, 5, 6, 14, 15, 27, 28,
        2, 4, 7, 13, 16, 26, 29, 42,
        3, 8, 12, 17, 25, 30, 41, 43,
        9, 11, 18, 24, 31, 40, 44, 53,
        10, 19, 23, 32, 39, 45, 52, 54,
        20, 22, 33, 38, 46, 51, 55, 60,
        21, 34, 37, 47, 50, 56, 59, 61,
        35, 36, 48, 49, 57, 58, 62, 63
    };

    for (int i = 0; i < 4; ++i)
    {
        Npp8u temp[64];

        for (int k = 0; k < 32; ++k)
        {
            temp[2 * k + 0] = aQuantizationTables[i].aTable[aZigzag[k + 0]];
            temp[2 * k + 1] = aQuantizationTables[i].aTable[aZigzag[k + 32]];
        }

        NPP_CHECK_CUDA(cudaMemcpyAsync((unsigned char *)pdQuantizationTables + i * 64, temp, 64, cudaMemcpyHostToDevice));
    }

    //NPP_CHECK_CUDA(cudaMemcpyAsync(pdQuantizationTables, aQuantizationTables[0].aTable, 64, cudaMemcpyHostToDevice));
    //NPP_CHECK_CUDA(cudaMemcpyAsync(pdQuantizationTables + 64, aQuantizationTables[1].aTable, 64, cudaMemcpyHostToDevice));

    oFrameHeader.nSamplePrecision = 8;
    oFrameHeader.nComponents = 3;
    oFrameHeader.aComponentIdentifier[0] = 1;
    oFrameHeader.aComponentIdentifier[1] = 2;
    oFrameHeader.aComponentIdentifier[2] = 3;
    oFrameHeader.aSamplingFactors[0] = 34;
    oFrameHeader.aSamplingFactors[1] = 17;
    oFrameHeader.aSamplingFactors[2] = 17;
    oFrameHeader.aQuantizationTableSelector[0] = 0;
    oFrameHeader.aQuantizationTableSelector[1] = 1;
    oFrameHeader.aQuantizationTableSelector[2] = 1;

    for (int i = 0; i < oFrameHeader.nComponents; ++i)
    {
        nMCUBlocksV = max(nMCUBlocksV, oFrameHeader.aSamplingFactors[i] & 0x0f);
        nMCUBlocksH = max(nMCUBlocksH, oFrameHeader.aSamplingFactors[i] >> 4);
    }

    oFrameHeader.nWidth = HD_IMG_WIDTH_MAX;
    oFrameHeader.nHeight = HD_IMG_HEIGHT_MAX;

    for (int i = 0; i < oFrameHeader.nComponents; ++i)
    {
        NppiSize oBlocks;
        NppiSize oBlocksPerMCU = { oFrameHeader.aSamplingFactors[i] >> 4, oFrameHeader.aSamplingFactors[i] & 0x0f };

        oBlocks.width = (int)ceil((oFrameHeader.nWidth + 7) / 8 *
            static_cast<float>(oBlocksPerMCU.width) / nMCUBlocksH);
        oBlocks.width = DivUp(oBlocks.width, oBlocksPerMCU.width) * oBlocksPerMCU.width;

        oBlocks.height = (int)ceil((oFrameHeader.nHeight + 7) / 8 *
            static_cast<float>(oBlocksPerMCU.height) / nMCUBlocksV);
        oBlocks.height = DivUp(oBlocks.height, oBlocksPerMCU.height) * oBlocksPerMCU.height;

        aSrcSize[i].width = oBlocks.width * 8;
        aSrcSize[i].height = oBlocks.height * 8;

        // Allocate Memory
        size_t nPitch;
        NPP_CHECK_CUDA(cudaMallocPitch(&apdDCT[i], &nPitch, oBlocks.width * 64 * sizeof(Npp16s), oBlocks.height));
        aDCTStep[i] = static_cast<Npp32s>(nPitch);

        NPP_CHECK_CUDA(cudaMallocPitch(&apSrcImage[i], &nPitch, aSrcSize[i].width, aSrcSize[i].height));
        aSrcPitch[i] = nPitch;
        aSrcImageStep[i] = static_cast<Npp32s>(nPitch);
    }

    /**
    * Allocates memory and creates a Huffman table in a format that is suitable for the encoder.
    */
    NppStatus t_status;
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanDCTables[0].aCodes, nppiDCTable, &apHuffmanDCTable[0]);
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanACTables[0].aCodes, nppiACTable, &apHuffmanACTable[0]);
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanDCTables[1].aCodes, nppiDCTable, &apHuffmanDCTable[1]);
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanACTables[1].aCodes, nppiACTable, &apHuffmanACTable[1]);
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanDCTables[1].aCodes, nppiDCTable, &apHuffmanDCTable[2]);
    t_status = nppiEncodeHuffmanSpecInitAlloc_JPEG(pHuffmanACTables[1].aCodes, nppiACTable, &apHuffmanACTable[2]);

    nScanSize = HD_IMG_WIDTH_MAX * HD_IMG_HEIGHT_MAX * 2;
    nScanSize = nScanSize > (4 << 20) ? nScanSize : (4 << 20);
    NPP_CHECK_CUDA(cudaMalloc(&pdScan, nScanSize));

    NPP_CHECK_NPP(nppiEncodeHuffmanGetSize(aSrcSize[0], 3, &nTempSize));
    NPP_CHECK_CUDA(cudaMalloc(&pJpegEncoderTemp, nTempSize));

    pDstJpeg = new unsigned char[nScanSize];

    uint32_t nPitch = (IMG_WIDTH_MAX + MEMORY_ALGN_DEVICE) & ~MEMORY_ALGN_DEVICE;
    mY = (uint8_t*)malloc(nPitch * IMG_HEIGHT_MAX);
    nPitch = (IMG_WIDTH_MAX_HALF + MEMORY_ALGN_DEVICE) & ~MEMORY_ALGN_DEVICE;
    mU = (uint8_t*)malloc(nPitch * IMG_HEIGHT_MAX_HALF);
    mV = (uint8_t*)malloc(nPitch * IMG_HEIGHT_MAX_HALF);

    uint32_t row_bytes = (IMG_WIDTH_MAX * 3 + 3) &  ~3;
    NPP_CHECK_CUDA(cudaMalloc(&mRGBData, row_bytes * HD_IMG_HEIGHT_MAX));
}

/*
brief:release host memory and device memory, no need always create the obj and release to induce the encoding performance
*/
void CudaJpegEncode::Uninit()
{
    cudaFree(pJpegEncoderTemp);
    cudaFree(pdQuantizationTables);
    cudaFree(pdScan);
    nppiDCTFree(pDCTState);
    if (pDstJpeg)
    {
        delete[] pDstJpeg;
        pDstJpeg = NULL;
    }
    if (mY)
    {
        free(mY);
        mY = NULL;
    }
    if (mU)
    {
        free(mU);
        mU = NULL;
    }
    if (mV)
    {
        free(mV);
        mV = NULL;
    }
    for (int i = 0; i < 3; ++i)
    {
        cudaFree(apdDCT[i]);
        cudaFree(apSrcImage[i]);
        nppiEncodeHuffmanSpecFree_JPEG(apHuffmanDCTable[i]);
        nppiEncodeHuffmanSpecFree_JPEG(apHuffmanACTable[i]);
    }
}


/*
brief: encode yuv data to jpeg format data, and write to file
*/
void CudaJpegEncode::EncodeJpeg(const char* output_file)
{
    oFrameHeader.nWidth = mWidth;
    oFrameHeader.nHeight = mHeight;
    for (int i = 0; i < oFrameHeader.nComponents; ++i)
    {
        NppiSize oBlocks;
        NppiSize oBlocksPerMCU = { oFrameHeader.aSamplingFactors[i] >> 4, oFrameHeader.aSamplingFactors[i] & 0x0f };

        oBlocks.width = (int)ceil((oFrameHeader.nWidth + 7) / 8 *
            static_cast<float>(oBlocksPerMCU.width) / nMCUBlocksH);
        oBlocks.width = DivUp(oBlocks.width, oBlocksPerMCU.width) * oBlocksPerMCU.width;

        oBlocks.height = (int)ceil((oFrameHeader.nHeight + 7) / 8 *
            static_cast<float>(oBlocksPerMCU.height) / nMCUBlocksV);
        oBlocks.height = DivUp(oBlocks.height, oBlocksPerMCU.height) * oBlocksPerMCU.height;

        aSrcSize[i].width = oBlocks.width * 8;
        aSrcSize[i].height = oBlocks.height * 8;
        // Allocate Memory
        size_t nPitch;
        nPitch = (oBlocks.width * 64 * sizeof(Npp16s) + HD_MEMORY_ALGN_DEVICE) & ~HD_MEMORY_ALGN_DEVICE;
        aDCTStep[i] = static_cast<Npp32s>(nPitch);

        nPitch = (aSrcSize[i].width + HD_MEMORY_ALGN_DEVICE) & ~HD_MEMORY_ALGN_DEVICE;
        aSrcPitch[i] = nPitch;
        aSrcImageStep[i] = static_cast<Npp32s>(nPitch);
    }

    NPP_CHECK_CUDA(cudaMemcpy(apSrcImage[0], mY, aSrcPitch[0] * mHeight, cudaMemcpyHostToDevice));
    NPP_CHECK_CUDA(cudaMemcpy(apSrcImage[1], mU, aSrcPitch[1] * mHeight / 2, cudaMemcpyHostToDevice));
    NPP_CHECK_CUDA(cudaMemcpy(apSrcImage[2], mV, aSrcPitch[2] * mHeight / 2, cudaMemcpyHostToDevice));
    /**
    * Forward DCT, quantization and level shift part of the JPEG encoding.
    * Input is expected in 8x8 macro blocks and output is expected to be in 64x1
    * macro blocks. The new version of the primitive takes the ROI in image pixel size and
    * works with DCT coefficients that are in zig-zag order.
    */
    int k = 0;
    NPP_CHECK_NPP(nppiDCTQuantFwd8x8LS_JPEG_8u16s_C1R_NEW(apSrcImage[0], aSrcImageStep[0],
        apdDCT[0], aDCTStep[0],
        pdQuantizationTables + k * 64,
        aSrcSize[0],
        pDCTState));
    k = 1;
    NPP_CHECK_NPP(nppiDCTQuantFwd8x8LS_JPEG_8u16s_C1R_NEW(apSrcImage[1], aSrcImageStep[1],
        apdDCT[1], aDCTStep[1],
        pdQuantizationTables + k * 64,
        aSrcSize[1],
        pDCTState));

    NPP_CHECK_NPP(nppiDCTQuantFwd8x8LS_JPEG_8u16s_C1R_NEW(apSrcImage[2], aSrcImageStep[2],
        apdDCT[2], aDCTStep[2],
        pdQuantizationTables + k * 64,
        aSrcSize[2],
        pDCTState));

    /**
    * Huffman Encoding of the JPEG Encoding.
    * Input is expected to be 64x1 macro blocks and output is expected as byte stuffed huffman encoded JPEG scan.
    */
    Npp32s nSs = 0;
    Npp32s nSe = 63;
    Npp32s nH = 0;
    Npp32s nL = 0;
    Npp32s nScanLength;
    NPP_CHECK_NPP(nppiEncodeHuffmanScan_JPEG_8u16s_P3R(apdDCT, aDCTStep,
        0, nSs, nSe, nH, nL,
        pdScan, &nScanLength,
        apHuffmanDCTable,
        apHuffmanACTable,
        aSrcSize,
        pJpegEncoderTemp));
    unsigned char *pDstOutput = pDstJpeg;

    writeMarker(0x0D8, pDstOutput);
    writeJFIFTag(pDstOutput);
    writeQuantizationTable(aQuantizationTables[0], pDstOutput);
    writeQuantizationTable(aQuantizationTables[1], pDstOutput);

    writeFrameHeader(oFrameHeader, pDstOutput);
    writeHuffmanTable(pHuffmanDCTables[0], pDstOutput);
    writeHuffmanTable(pHuffmanACTables[0], pDstOutput);
    writeHuffmanTable(pHuffmanDCTables[1], pDstOutput);
    writeHuffmanTable(pHuffmanACTables[1], pDstOutput);

    oScanHeader.nComponents = 3;
    oScanHeader.aComponentSelector[0] = 1;
    oScanHeader.aComponentSelector[1] = 2;
    oScanHeader.aComponentSelector[2] = 3;
    oScanHeader.aHuffmanTablesSelector[0] = 0;
    oScanHeader.aHuffmanTablesSelector[1] = 17;
    oScanHeader.aHuffmanTablesSelector[2] = 17;
    oScanHeader.nSs = 0;
    oScanHeader.nSe = 63;
    oScanHeader.nA = 0;

    writeScanHeader(oScanHeader, pDstOutput);
    NPP_CHECK_CUDA(cudaMemcpy(pDstOutput, pdScan, nScanLength, cudaMemcpyDeviceToHost));
    pDstOutput += nScanLength;
    writeMarker(0x0D9, pDstOutput);
    {
        // Write result to file.
        std::ofstream outputFile(output_file, ios::out | ios::binary);
        outputFile.write(reinterpret_cast<const char *>(pDstJpeg), static_cast<int>(pDstOutput - pDstJpeg));
    }
}

/*
brief: setting yuv image data, convert to yuv420
yuv_data:yuv data
yuv_fmt: yuv format, nv12 or nv21
w:       image width
h:       image height
size:    data size
*/

void CudaJpegEncode::SetData(uint8_t* yuv_data, uint32_t data_size, int yuv_fmt, int w, int h)
{
    if (!yuv_data)
    {
        return;
    }
    mWidth = w;
    mHeight = h;

    int i, j;
    uint32_t    nPitch;
    uint32_t    off;
    uint32_t    off_yuv;
    uint32_t    half_h;
    uint32_t    half_w;
    uint32_t    u_size;
    uint8_t*    yuv_ptr;
    uint8_t*    u_ptr;
    uint8_t*    v_ptr;

    nPitch = (w + HD_MEMORY_ALGN_DEVICE) & ~HD_MEMORY_ALGN_DEVICE;
    off = 0;
    off_yuv = 0;
    for (i = 0; i < h; i++)
    {
        memcpy(mY + off, yuv_data + off_yuv, w);
        off += nPitch;
        off_yuv += w;
    }

    half_w = w >> 1;
    half_h = h >> 1;
    u_size = half_w * half_h;
    nPitch = (half_w + HD_MEMORY_ALGN_DEVICE) & ~HD_MEMORY_ALGN_DEVICE;
    switch (yuv_fmt)
    {
    case YUV_FMT_NV12:
    {
        off_yuv = w * h;
        off = 0;
        for (i = 0; i < half_h; i++)
        {  
            yuv_ptr = yuv_data + off_yuv;
            u_ptr = mU + off;
            v_ptr = mV + off;
            for (j = 0; j < w; j += 2)
            {
                *u_ptr++ = *yuv_ptr++;
                *v_ptr++ = *yuv_ptr++;
            }
            off_yuv += w;
            off += nPitch;
        }
    }
    break;
    case YUV_FMT_NV21:
    {
        off_yuv = w * h;
        off = 0;
        for (i = 0; i < half_h; i++)
        {
            yuv_ptr = yuv_data + off_yuv;
            u_ptr = mU + off;
            v_ptr = mV + off;
            for (j = 0; j < w; j += 2)
            {
                *v_ptr++ = *yuv_ptr++;
                *u_ptr++ = *yuv_ptr++;
            }
            off_yuv += w;
            off += nPitch;
        }
    }
    break;
    case YUV_FMT_YUV420:
    {
        off_yuv = w * h;
        off = 0;
        for (i = 0; i < half_h; i++)
        {
            memcpy(mU + off, yuv_data + off_yuv, half_w);
            memcpy(mV + off, yuv_data + off_yuv + u_size, half_w);
            off_yuv += half_w;
            off += nPitch;
        }
    }
    break;
    default:
        break;
    }
}
  •  

使用时先Init，然后再setdata，最后encode。 
CUDA中不停的申请和拷贝数据会占用更多的时间，所以，需要的空间先分配，编码时直接使用，这样会提升很大效率。 
测试结果：采用cpu编码1920x1080的图片需要时间17ms，而采用上述代码，9ms，提升一倍以上。

以上部分代码参考：https://www.cnblogs.com/betterwgo/p/7300920.html，感谢作者的贡献。

yanyan_happy0506
关注
专栏目录
CUDA JPEG编码
08-08
基于英伟达的jpegNPP工程，分离实现独立的JPEG压缩。
CUDA JPEG编码
chenxiuli0810的专栏
06-04 1204
https://www.cnblogs.com/betterwgo/p/7300920.html CUDA JPEG编码 基于英伟达的jpegNPP工程，分离实现独立的JPEG压缩。 由于原工程是直接把解码时的jpeg图片的信息直接作为编码时的信息，所以在做独立的JPEG编码时，需要自己来填充各种信息。 1.JPEG编码流程 从网上一片文章中解出来的图，...
JPEG GPU处理库实战指南
最新发布
gitblog_00429的博客
08-24 834
JPEG GPU处理库实战指南 jpeg_gpuGPU accelerated JPEG decoder项目地址:https://gitcode.com/gh_mirrors/jp/jpeg_gpu 项目介绍 JPEG GPU 是一个致力于利用GPU加速JPEG图像编码与解码过程的开源项目。它通过高效的GPU算法优化，大大提升了在大规模图像处理场景下的性能。该项目托管于GitHub，地址为 h...
cuda 编码jpeg
littleleafboy的博客
03-26 1611
一、nvjpegImage_t赋值： 二、接口调用eg：
使用CUDA将NV12格式转换为JPEG的代码
12-22
使用CUDA，将NV12格式的缓冲区编码为JPEG。代码可以编译，并直接运行。代码是https://blog.csdn.net/mikedai/article/details/79084092，进行了接口、对齐、编译功能。
使用CUDA8.0封装的JPEG压缩接口
12-15
2. **JPEG压缩原理**：理解DCT（离散余弦变换）、量化和熵编码等JPEG压缩的核心步骤，这对实现GPU上的压缩至关重要。 3. **并行计算优化**：分析JPEG压缩算法的并行性，找出可以并行化的部分，如DCT计算，合理分配...
cuda编码成jpeg库
07-14
CUDA编码成JPEG库的主要目的是通过利用GPU的并行计算能力加快JPEG编码和解码的速度。JPEG是一种常用的图像压缩格式，其压缩算法复杂度较高，普通CPU处理速度较慢。而CUDA编码成JPEG库利用GPU中的数百个并行处理单元...
CUJ2K: Jpeg2000 encoder on Cuda-开源
04-19
在支持CUDA的GPU上运行的快速JPEG2000编码器。 支持无损和有损编码。 带有GUI和CUDA / C ++库的两个命令行工具。 请访问>> http://cuj2k.sourceforge.net <<以获取更多信息。
cudaEncode编码详细过程
尚钦
08-05 5806
cudaEncode编码详细过程如下： 一，ParseInputParams(argc, argv,&sEncoderParams) //配置参数：输入、输出文件，配置文件，基本的参数   二，pCudaEncoder = new VideoEncoder  (&sEnc
基于cuda平台的视频编码
06-17
基于cuda平台的视频编码实现 及各部分的详细介绍
CUDA H264编码源代码
09-24
演示使用CUAD SDK利用显卡硬件加速进行H264编码，充分利用GPU运算能力，CPU资源占用少。
通过cuda处理图像显示图片
07-14
非常简单的例子，把后缀改成.cu，安装好vs、opencv和cuda就能运行。cpu内存读取图片（单通道），拷贝到gpu纹理内存，再取出来放回cpu内存，主要体会关于内存的操作。非常适合初学者。
cuda编码入门学习笔记
xiaomu_347的博客
06-27 880
首先，确保你对CUDA编程的基本概念有所了解：- CUDA是NVIDIA提供的用于并行计算的平台和编程模型。- CUDA允许你利用NVIDIA GPU的并行处理能力来加速计算任务。
CUDA库之nvjpeg(一)：入门介绍
hjxu2016的博客
03-27 1928
nvJPEG 库提供了高性能，GPU加速的JPEG图像格式的解码函数，并且普遍应用在深度学习领域和超大规模的多媒体应用中。这个库提供单张图或者多张图同时解码的能力，可以充分利用GPU资源和优化效率，并且使用者也可以管理需要解码的内存，灵活性还是比较强的。使用JPEG图像数据流作为输入从数据流中获取图像的宽和高使用以上获取的信息来管理GPU内存并执行解码操作nvJPEG提供了专用的API，用于从原始JPEG图像数据流中检索图像信息。
使用CUDA的nvJPEG库进行图像编码
YOLO_CODE的博客
09-07 670
NVIDIA提供了一个高性能的图像编码库，即nvJPEG库，它利用CUDA技术来加速图像编码操作。这样，您就可以使用CUDA的nvJPEG库进行图像编码了。首先，确保您的系统中已经安装了CUDA，并正确设置了CUDA环境。从NVIDIA的开发者网站下载适用于您的CUDA版本的nvJPEG库，并按照提供的安装说明进行安装。在您的CUDA项目中，包含nvJPEG库所需的头文件和库文件。请注意，上述代码示例仅为演示目的，并可能需要根据您的具体应用进行适当的修改和调整。使用CUDA的nvJPEG库进行图像编码。
CUDA：nvJPEG库的用法nvJPEG编码器支持单个和多个图像编码实例
希望我的博客，能帮上你解决学习中工作中所遇到的问题
12-22 835
CUDA：nvJPEG库的用法nvJPEG编码器支持单个和多个图像编码实例
CUDA 实现JPEG图像解码为RGB数据
热门推荐
weixinhum
06-30 1万+
了解JPEG数据格式的人应该容易想到，其对图像以8*8像素块大小进行分割压缩的方法很好用并行处理的思想来实现。而事实上英伟达的CUDA自v5.5开始也提供了JPEG编解码的示例。该示例存储在CUDA的SDK中，即CUDA的默认安装路径“C:\ProgramData\NVDIA Corporation\CUDA Samples\v7.0\7_CUDALibraries\jpegNPP”(v后面的
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值