基于GPGPU 的JPEG2000 图像压缩方法*

李玉峰 ，吴 蔚，王 恺，崔迎炜

(1.东南大学移动通信国家重点实验室，南京210096;2.沈阳航空航天大学电子信息工程学院，沈阳110136;3.北京方天长久科技有限公司，北京100085)

随着多媒体技术在计算机科学领域应用的广泛普及，图像压缩技术成为现代数字图像传输、处理、存储中的关键技术。不论是在网络传输等方面，还是在移动通信等领域都有着重要的意义和用途。JPEG2000 是在JPEG 基础上提出的一项新的静态图像压缩标准，同JPEG 压缩标准相比，不但在压缩性能上做了优化，能够以更高的压缩比率压缩图像数据。而且也具有同时支持有损压缩和无损压缩的优势。由于图像像素可以看成二维数组［1］，而计算二维数组则相当于计算大量不相关的数据，特别是质量较高的原始位图，利用传统CPU 串行的结构特点处理会消耗大量时间，无法满足图像压缩在现代多媒体技术应用方面的实时性要求。在硬件实现方面，传统的图像压缩一般采用DSP 和FPGA 等硬件平台来实现，但是DSP 和FPGA 等硬件的实现，要求研究人员需要对硬件内部结构有深入的研究，并且在移植性方面也存在着一定的困难。然而通用计算图形处理器(GPGPU)的推出，除了传统GPU 所拥有的图形处理架构以外，GPGPU 还增加了并行计算架构，该架构为计算密集型处理和高强度并行加速计算提供了可能。NVIDIA 公司为其GPGPU 提供了全新的软硬件开发平台CUDA(Compute Unified Device Architecture)，利用CUDA 技术在GPGPU 上优化JPEG2000 图像压缩核心算法同在CPU 上实现其算法相比较，在计算速度上有了显著提升，很大程度的提高了图像压缩的效率。

1 JPEG2000 图像压缩方法

JPEG2000 图像压缩标准同JPEG 相比，最大的不同是在算法方面做了进一步的改进［2］。首先，它摒弃了JPEG 采用的以离散余弦变换(DCT)为主的区块编码方式，而是选择了以离散小波变换(DWT)为主的全帧多解析编码方式，以此来减少图像中包含的数据冗余信息，避免了在低比特率的情况下JPEG 压缩标准会产生方块噪声的缺点。其次，在熵编码算法方面，JPEG2000 采用优化截断的嵌入式块编码(EBCOT)算法，取代了JPEG 的哈夫曼编码算法。

JPEG2000 核心编码系统主要包括七个模块，如图1 所示，首先对预处理后的原始图像进行正向小波变换得到小波系数，然后根据具体需要对变换后的小波系数进行量化;接着将量化后的小波系数划分成码块，对每一个码块进行独立的嵌入式编码;得到的所有码流按照率失真最优原则分层组织;最后按照一定的码流格式对这些不同质量的层打包后输出压缩码流，即完成了整个图像的压缩过程。

图1 JPEG2000 核心编码器原理

2 GPGPU 技术及CUDA 平台

2.1 GPGPU 技术

随着3D 时代的到来，庞大的3D 图形数据计算量已经远远超出了传统CPU 的计算能力范围。为此，图形处理器GPU 随之产生，GPU 是专门为图形计算而设计的，同CPU 相比，GPU 有高浮点运算性能，高带宽，高效并行计算的诸多优势。但如此强大的计算能力仅用于图形渲染，这对于计算资源来说无疑是一种浪费［3］。为了充分利用GPU 强大的计算能力，同时满足除图形计算以外的其他科学计算领域的需求，通用计算图形处理器GPGPU(General Purpose GPU)应运而生，并且已经取得了巨大的成果。GPGPU 的硬件是采用单指令多数据SIMD(Single Instruction Multiple Data)的结构，同时将图形处理架构和并行计算架构完美结合［4］，这使得GPGPU 不仅能够作为图形显卡用于图形渲染的本职工作，同时更大程度上涉足其他非图形的科学计算领域，充分发挥其强大的并行计算能力。如图2所示，一个硬件GPGPU 芯片是由多个流多处理器SM(Stream Multiprocessor)组成，每个流多处理器包含8 个流处理器SP(Stream Processor)和两个特殊功能单元SFU(Special Function Unit)以及一些片内存储器资源，例如共享存储器(shared memory)和寄存器(Register)等。

图2 GPGPU 硬件结构模型

如图3 所示。其中寄存器是高速存储器，与硬件结构对应，每个流处理器拥有一个私有的32 bit寄存器。共享存储器的访问速度几乎同寄存器一样快，但是存储空间较小，按照当前硬件支持，其默认大小为16 K，可被同一个线程块中的所有线程读写访问。在利用GPGPU 并行加速计算时需要根据不同存储器的访问速度和功能特点对存储器进行合理分配，这是提高GPGPU 计算性能的关键。

图3 GPGPU 硬件SM 结构

2.2 CUDA 平台

CUDA 是NVIDIA 公司在2007 年6 月推出的专门为其GPGPU 产品设计的异构开发平台，从此也彻底改变了GPGPU 并行计算的命运。CUDA 提供了硬件直接访问接口，无需像传统GPGPU 开发需要借助于Open GL 和Direct X 等图形学API 来实现。同时，CUDA 对广泛使用的C 语言进行了扩展，更进一步降低了GPGPU 并行加速计算的编程难度，使开发人员能够很容易地从C 语言的应用开发过渡到GPGPU 的应用开发。CUDA 提供了与GPGPU 的SIMD 结构相对应的单指令多线程SIMT(Single Instruction Multiple Thread)异构编程执行模型［5］，如图4 所示。CUDA的分层次线程结构包括线程(Thread)、线程块TB(Thread Block)和网格(Grid)。每个网格由一定数量的线程块组成，而每一个线程块最多包含512 个线程。一个CUDA 程序由运行在Host(CPU)端的程序和运行在Device(GPU)端的程序组成。Host 端执行串行指令用来调度任务分配，Device 作为Host 的协处理器执行并行计算部分，在Device 端执行的程序称为Kernel(内核)函数，Kernel 以一个Grid 的形式执行。一个简单的Device 程序需要完成以下两个过程:(1)在kernel 函数调用之前，需要将所需处理的数据从主机内存(Host Memory)复制到Device 端的全局存储器(Global Memory)中;(2)计算完成后，将计算结果从全局存储器返回到主机内存。CUDA 并行加速计算的实质是将一个任务分割成多个相互独立的任务块，利用成千上万的线程对任务块同时进行处理，从而提高整个任务的计算速度［7］。因此，将图像数据分割成数据块，利用CUDA 技术能够更进一步提高JPEG2000 图像压缩的速度。

图4 CUDA 异构编程执行模型

3 DWT 在GPGPU 上的实现

JPEG2000 标准中采用的离散小波变换算法有2 种:5/3 整数型小波提升算法和9/7 浮点型小波提升算法，其中5/3 小波提升算法适用于有损压缩和无损压缩，实现算法如下:

在较低比特率的情况下，9/7 浮点型小波提升算法能够发挥出最优越的性能，在有损压缩时推荐使用。同5/3 小波提升算法相比，9/7 小波提升算法比较复杂，如下:

其中:

3.1 CUDA 程序核心步骤实现

由JPEG2000 的小波变换算法可知，真正涉及到大量的数据运算是在一维行列变换的时候［7－8］。因此可以确定，将基本的行变换和列变换提升操作设计成kernel 函数，通过GPU 调用完成计算。其他部分工作任务均交给CPU 完成［9］。由此可以得出如图5 所示的小波变换CPU－GPU 异构并行实现任务分配原理图。

图5 DWT 的异构并行实现原理图

该过程的具体实现步骤如下:

(1)在CPU 端分配主机内存空间X 和Y 分别用于存放输入图像数据和压缩后的输出图像数据，将图像数据读取到CPU 内存，调用库函数中的cudaMalloc在设备端分配两个完全相同的全局存储器空间(显存)X1和X2;cudaMalloc((void＊＊)＆d_Data，sizeof(unsigned char)* dataSize);//分配输入显存空间X1cudaMalloc((void＊＊)＆d_new，sizeof(unsigned char)* dataSize);//分配输出显存空间X2;

(2)通过调用库函数中的cudaMemcpy 函数将CPU 内存中的图像数据拷贝至GPU 端已设置好大小的全局存储器空间X1，供其进行后续小波列变换计算;cudaMemcpy(d_Data，h_Data，sizeof(unsigned char)* dataSize，cudaMemcpyHostToDevice);

(3)初始化输入数据，设置执行参数及所需共享存储器大小，将分段的图像数据从全局存储器X1中拷贝至已经设置好的共享存储器中，对其中的数据进行小波提升变换，以9/7 小波变换为例，其流程如图6 所示。由于提升小波操作需要多个Thread同时进行处理来实现并行加速，由提升算法可知，每个Thread 一次只能处理2 个数据，然而在接下来的计算中需要用到前一步提升计算的结果。为了确保计算结果的准确性，在每一步提升计算操作完成之后都需要调用线程同步函数__syncthread()，用来等待所有的Thread 全部完成之前的提升操作以后，再进行后续的计算。然后将计算结果按顺序存入之前设置好的显存空间X2中;

__shared__unsigned char space_s1［16］［16］;

//声明未处理数据的Shared Memory 空间

__shared__unsigned char space_s2［16］［16］;

//声明处理后数据的Shared Memory 空间

图6 9/7 小波提升kernel 流程

(4)将上一步骤中的结果数据按照相同的方法操作，进行一维小波行变换，将计算结果存入已空置的全局存储器空间X1中;

(5)按照小波分解等级控制，重复小波提升变换过程，直至将所有等级的小波提升变换过程全部完成，最后将全局存储器的结果图像数据返回到CPU主机内存中，释放GPU 显存空间和CPU 内存空间。

cudaMemcpy(h_Image，d_new，sizeof(unsigned char)* dataSize，cudaMemcpyDeviceToHost);

cudaFree(d_new);//释放GPU 显存空间

Free(＆(img));//释放CPU 内存空间

3.2 共享存储器优化

由于全局存储器读取速度较慢［10］，若将待处理数据存储在全局存储器中进行反复存取，则会大幅度的降低程序执行的效率。然而共享存储器拥有较快的读取访问速度，因此选用共享存储器的优化策略来加速DWT 算法。但当多个线程对共享存储器的同一位置读取数据时，会产生Bank conflict，影响执行性能，所以在优化过程中必须避免Bank conflict。

首先，将全局存储器中的图像数据分割成(n+1)×(n+1)个数据块DB(Data Block)，同时创建相同数目的线程块TB(Thread Block)，将线程块一一映射到分割后的每一个数据块，如图7 所示。为每个TB 声明一块共享存储器空间，用于存放映射到线程块内的每个线程需要执行的数据块中的数据。

图7 共享存储器数据存取优化

利用共享存储器访问速度快这一优点，按照如上方法，将全局存储器中的源图像数据分块复制到共享存储器以后，线程块内的每一个线程直接对共享存储器内对应的数据进行计算，这样就避免了可能出现的Bank conflict。此外，在CUDA 程序优化时，要保证访问全局存储器时需要满足合并访问条件，该条件同样是优化CUDA 程序性能的重要因素之一。同时，还需考虑计算精度、访存延迟、计算数据量等多方面因素［11］。

4 实验平台及测试结果分析

测试所需的软硬件环境如下所示:

CPU:Intel E7400 酷睿双核2.80 GHz CPU，主频2 800 MHz，主机内存为2 GB;

GPU:采用NVIDIA GeForce GTX 560 Ti 设备，流多处理器数量8 个，CUDA 流处理器数量384 个，核心频率822 MHz，流处理器频率1 645 MHz，显存频率4 008 MHz，显存容量1 024 MB，显存带宽128 GB/s，显存位宽256 bit，计算能力2. 1，总线接口PCI－E 2.0x16;

编程环境:GPU 硬件驱动版本为301.42，使用CUDA4. 1 版本的编程环境，windows 7 操作系统，VS2010。

由于CPU 上和GPGPU 上的JPEG2000 图像压缩均是按照传统定义来实现的。因此，图像的压缩质量基本上是相同的，通过对CPU 和GPGPU 上测试时间进行比较分析，实验结果分析详见表1。

表1 DWT 在CPU 和GPGPU 上的时间对比

从实验数据可以看出，同未经过优化的CPU 测试结果相比，对于所需处理数据相对较少的像素为640×480 的图像来说，计算速度提高了9 倍多，相比之前的研究—基于CUDA 的小波Mallat 算法及提升方案的设计与研究［12］中对于像素为512×512 的图像其提升小波算法的计算时间为6.785 ms，可见优化后的DWT 提升算法的计算速度有所提升。而对于所需处理数据较多的图像来说，计算速度提高超过了50倍。显而易见，经过CUDA 优化的JPEG2000 静态图像压缩标准中的DWT 算法在GPGPU 上的计算时间要远远小于CPU 的计算时间，并且随着计算数据量的增加，CPU 的运算时间呈现大幅度增长趋势，而GPGPU 的计算时间增长较小。从加速比也可以看出，随着数据量计算的增加，GPGPU 表现出更优越的加速计算性能。

此外，由于基于GPU 的JPEG2000 图像压缩已经有了不少研究成果，为了证明本实验所选用的方法在原有的结果上做了相应的优化改进。选用之前普遍使用的NVIDIA GT 240 GPU 进行DWT 提升计算测试，并与本实验的测试结果进行对比，如表2 所示。

表2 DWT 在不同GPU 上的测试时间对比

由以上测试结果可以看出，本实验的优化方法在一定程度上取得了相应的效果。尽管如此，由于CUDA 程序的优化还需要结合指令流优化等多种因素。因此，在以后的程序优化工作中，还需要综合考虑各种优化原则进行反复试验，权衡各方面因素得出一个尽量接近最优的优化结果。

5 结束语

通过JPEG2000 图像压缩算法分别在CPU 和GPGPU 两种不同的处理器上的实现可以得出，在科学计算领域，特别是针对类似图像压缩这种数据之间相关性不大的大规模密集型的浮点型数据计算时，利用CUDA 平台在GPGPU 上并行加速实现能够表现出更大的性能优势。尽管当前GPGPU 并行加速计算还受到一些诸如兼容性、算法移植和优化等方面的技术难点限制，但是GPGPU 并行计算拥有的巨大潜力和无可比拟的性能优势已经受到了世界各行业领域的广泛关注，相信基于GPGPU 并行计算必将成为今后的发展方向。

［1］ 郭静，陈庆奎.基于CUDA 的快速图像压缩［J］.计算机工程与设计，2010，31(14):3302－3304，3308.

［2］ ISO/IEC 15444 － 1 Information Technology:JPEG2000 Image Coding System［S］.USA:ISO/IEC，2002.

［3］ 张舒，褚艳丽， ，等. GPU 高性能计算之CUDA［M］. 中国水利水电出版社，2009:12.

［4］ NVIDIA Corporation. NVIDIA CUDA Programming Guide，version 3.0，2010.

［5］ John D O，David L，Naga G，et al. A Survey of General－Purpose Computation on Graphics Hardware［J］. Computer Graphics Forum，2007(26):80－113.

［6］ 宋晓丽，王庆.基于GPGPU 的数字图像并行化预处理［J］. 计算机测量与控制，2009，17(6):1169－1171.

［7］ Ding W，Wu F，Li X，et al. Adaptive Directional Lifting－Based Wavelet Transform for Image Coding［j］. IEEE Trans Image Processing，2007，16(2):416－428.

［8］ 宋凯，臧晶. 图像处理中小波提升方案和Mallat 算法的比较［j］.微处理机，2004，10(5):39－41.

［9］ Joaquín Franco，Gregorio Bernabé，Juan Fernández，et al. A Parallel Implementation of the 2D Wavelet Transform Using CUDA［J］.IEEE Computer Society，2009，111－118.

［10］ Sunpyo Hong，Hyesoon Kim. Ananalytical Model for a GPU Architecture with Memory－Level and Thread－Level Parallelism Awareness［J］.ACM Sigarch Computer Architectu Renews，2009，37(3):152－163.

［11］ Han T D，Abdelrahman T S.High－Level GPGPU Programming［J］.Parallel and Distributed Systems，IEEE Transactions on，2011，2(1):78－90.

［12］ 孙自龙.基于CUDA 的小波Mallat 算法及提升方案的设计与研究［D］.华中科技大学，2011:37－39.