CUDA实现JPEG编码

最新推荐文章于 2024-10-16 17:04:06 发布
最新推荐文章于 2024-10-16 17:04:06 发布
阅读量1w 收藏 13
点赞数 5
分类专栏： 图像处理
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：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，感谢作者的贡献。

MikeDai
关注
专栏目录
cuda 编码jpeg
littleleafboy的博客
03-26 1618
一、nvjpegImage_t赋值： 二、接口调用eg：
CUDA JPEG编码
08-08
基于英伟达的jpegNPP工程，分离实现独立的JPEG压缩。
彩色24位图像编码/压缩(基于nvjpeg.h库)
最新发布
weixin_53173024的博客
10-16 219
3、通过调节CHECK_NVJPEG(nvjpegEncoderParamsSetQuality(encoder_params, 80, NULL));输出图像：彩色24位JPG/JPEG图像（其他格式的图像如PNG应该也可以，不过我没测试，可以自己试试。1、运行环境：cuda11.6+opencv3.4.16+vs2022。输入图像：彩色24位RGB颜色空间的BMP图像。函数中的质量因子，可调节图像的压缩比。
CUDA JPEG编码
chenxiuli0810的专栏
06-04 1205
https://www.cnblogs.com/betterwgo/p/7300920.html CUDA JPEG编码 基于英伟达的jpegNPP工程，分离实现独立的JPEG压缩。 由于原工程是直接把解码时的jpeg图片的信息直接作为编码时的信息，所以在做独立的JPEG编码时，需要自己来填充各种信息。 1.JPEG编码流程 从网上一片文章中解出来的图，...
使用CUDA8.0封装的JPEG压缩接口
12-15
使用CUDA8.0封装的JPEG压缩接口，其中实现了JPEG的CPU和GPU两种方式压缩,其中GPU的压缩效率还需要优化，不过实现方法可以借鉴。
使用CUDA将NV12格式转换为JPEG的代码
12-22
使用CUDA，将NV12格式的缓冲区编码为JPEG。代码可以编译，并直接运行。代码是https://blog.csdn.net/mikedai/article/details/79084092，进行了接口、对齐、编译功能。
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 <<以获取更多信息。
基于CUDA的JPEG解码
tkp2014的专栏
03-06 2578
JPEG图片解码的思路： 1、将JPEG并行解码的算法在两种并行结构上实现：GPU - 多核系统，SESC - SMP系统. JPEG on CUDA and SESC.        fused multiply-add RGB转化为YUV的一种优化方法：
Cuda Jpeg Decoder-开源
05-09
它使用cuda解码gpu中的jpeg文件，并将其映射为direct3d纹理，我将此作为概念证明而编写，希望继续的人可以继续.....
jpg图片编码算法代码示例(bmp转jpg)
06-13
学习流媒体技术需要熟悉各种文件格式，了解一些编解码的算法，此程序演示了jpg图片的编码算法，算法流程相对简单，希望能对初学者一些启示和帮助。对各种图片和音视频文件格式感兴趣的朋友，请进群：417099529，一起探讨学习研究，进群时请介绍自己对流媒体哪些方面比较熟悉和擅长。
JPEG编码 解码代码 C++
12-01
JPEG图象编解码代码，代码没有注释，西北大学（美国那个西北大学）放在网上供下载的，没有别的乱七八糟的东西
nvcomp：用于在GPU上进行快速无损压缩解压缩的库
02-11
什么是nvcomp？ nvcomp是CUDA库，具有通用压缩接口，使开发人员能够在其应用程序中使用高性能GPU压缩器。 nvcomp 1.2包括级联和LZ4压缩方法。 级联压缩方法显示出高达500GB / s的吞吐量的高性能，并且对来自分析工作负载的数值数据的压缩率高达80倍。 LZ4方法具有高达60 GB / s的解压缩吞吐量和任意字节流的良好压缩率。 下面是针对TPC-H和Fanny Mae的抵押数据集的级联压缩（R1D1B1）的一些数字列以及针对LZ4压缩的TPC-H的字符串列的性能与压缩比散点图。 对于TPC-H数据集，我们使用SF10 lineitem表和以下列：列0（L_ORDERKEY）作为8B整数，而列8、9、13、14、15作为字节流。 在抵押数据集中，我们使用了2009年第2季度绩效表：第0列（LOAN_ID）为8B整数，第10列（CURRENT_LOAN_DEL
使用OpenCV操作CUDA 图像
jacke121的专栏
12-10 2159
void makeJPEG(PyObject* _py_list_src, int idx, int _target_size, bool _crop_to_square, PyObject* _py_list_tgt, bool withGpu) { cv::Mat _resized_mat_buffer; cv::cuda::GpuMat _resized_mat_buffer_gpu; std::vector<uchar> _output_jpeg_buf...
CUDA库之nvjpeg(一)：入门介绍
hjxu2016的博客
03-27 2003
nvJPEG 库提供了高性能，GPU加速的JPEG图像格式的解码函数，并且普遍应用在深度学习领域和超大规模的多媒体应用中。这个库提供单张图或者多张图同时解码的能力，可以充分利用GPU资源和优化效率，并且使用者也可以管理需要解码的内存，灵活性还是比较强的。使用JPEG图像数据流作为输入从数据流中获取图像的宽和高使用以上获取的信息来管理GPU内存并执行解码操作nvJPEG提供了专用的API，用于从原始JPEG图像数据流中检索图像信息。
CUDA：nvJPEG库的用法nvJPEG编码器支持单个和多个图像编码实例
希望我的博客，能帮上你解决学习中工作中所遇到的问题
12-22 841
CUDA：nvJPEG库的用法nvJPEG编码器支持单个和多个图像编码实例
CUDA处理jpeg缩略图
a5615024的博客
05-24 286
　　缩略图的处理有很多种，比如之前写的go自带的image库，imagick库，libjpeg-turbo等等，今天来总结一下如何用CUDA处理jpeg的缩略图。 　　首先简单介绍一下CUDA，CUDA是Nvidia公司的GPU编程平台，通俗来说就是将CPU作为主机端，显卡作为设备端，将大批量的运算任务放在GPU上完成，这样做的好处就是能充分利用GPU的多核心运算能力优化程序，使程序效...
评论 42
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值