基于CUDA的快速图像压缩

33022010,31(14)计算机工程与设计Computer Engineering and Design

0引言

图像压缩技术是计算机科学中数据压缩技术的一个分支。因为大多图像来自照相机或录像机，数据不规则，目前传统无损压缩算法做得不好，所以图像压缩技术一般是有损压缩。JPEG 图像压缩是有损压缩中的一项重要技术，除了在计算机、网络传输等方面有重要的意义，在扫描仪、打印机等图像处理领域也有广泛用途。图像是由像素矩阵构成的，图像处理实质上就是矩阵运算，而矩阵运算非常适合并行处理。现代推出的较新系列显卡上GPU 采用的都是第二代统一渲染架构，它包括是图形处理架构和并行计算架构。GPU 的并行计算架构上是为计算密集型处理、高强度并行计算而设计的，特别适合处理那些可以具有较高算法强度且可以表达为并行数据计算的问题。NVIDIA 专为GPU 通用并行计算提出

了全新的软硬件架构CUDA 。采用CUDA 技术在GPU 上实现JPEG 图像压缩比通过CPU 实现其算法速度有明显提升，效率有极大提高。

1

关键技术介绍

1.1

JPEG 原理简介

JPEG 主要采用了以DCT 为基础的有损压缩算法，它利用

了人的视角系统的特性，使用量化和无损压缩编码相结合来去掉视角的冗余信息和数据本身的冗余信息。

JPEG 算法分为基线JPEG (baseline system )和扩展JPEG (ex-tended system )，通过基本JPEG 得到的编码文件同样在扩展JPEG 中可以得到正确解码，所以基线JPEG 应用要广泛的多。本文主要讨论的就是基线JPEG 编码。JPEG 编码流程如图1所示。

收稿日期：2009-09-11；修订日期：2009-11-16。

开发与应用

郭静，陈庆奎：基于CUDA的快速图像压缩2010,31(14)3303

1.2GPU技术

在GPU概念提出的10年里，GPU发展极为迅速，甚至已

经超过CPU的发展速度。T&L技术的引入使得3D渲染大部

分计算交由GPU完成，解放了GPU。Shader技术引入又进一

步扩展了GPU功能，使原来的渲染管线和T&L带有了可编程

的能力，极大地增强了GPU绘制出来的3D图形的表现力。流

处理器技术的引入，即统一渲染架构，统一了顶点、像素、几何处理等各种Shader计算。

在此基础上，GPU计算能力得到了极大的提高，尤其是在并行计算方面。如此强大的GPU计算能力，除了用于图形渲染外，大部分都处在空闲状态。为了充分利用GPU计算资源，人们开始研究在GPU完成通常意义上数据计算GPGPU，并取得了很大成果。由此，近年来，为了适应未来并行运算的需要，GPU采用的是全新的第二代统一架构即图形处理架构和并行计算架构完美结合。GPU除了进行图形渲染的本职工作之外，开始越来越多地涉足非图形计算领域，GPU已成为一颗真正意义上的通用处理器。

NVIDIA GeForce9600GT就是采用了此架构，其除了在图形处理方面的升级外，同样是一款非常出色并行架构处理器。GeForce9600GT拥有64个流处理器。GeForce9600GT能够无缝将计算任务从CPU上释放出来，交由GPU进行处理。1.3CUDA技术

CUDA是一个基于GPU计算的全新软硬件架构，这种构架的灵活性很强，在流处理器加入了特定的缓存，既可以用来处理图形，也可以用来进行通用线程处理以及高性能计算。CUDA是NVIDA推出的一种CPU+GPU异构运算平台。在这个架构中，GPU可视为一个计算设备，是主机或者CPU处理器，用于处理高度并行的计算。GPU具备自己的存储器，可以并行地运行许多线程。CUDA提供了硬件的直接访问接口，不必像传统方式一样必须依赖图形API接中来实现GPU的访问。在架构上采用了一种全新的计算体系结构来使用GPU提供的硬件资源，从而给大规模的数据计算应用提供了一种比CPU更加强大的计算能力。CUDA采用C语言作为基础编程语言，提供大量的高性能计算指令开发能力，使开发者能够在GPU的强大计算能力的基础上建立起一种效率更高的密集数据计算解决方案。CUDA架构如图2所示：NVIDIA专为CUDA提供了些优化过Libraries，如math.h、FFT(快速傅里叶变换)等。CUDA架构的核心是NVIDIA C编译器NVCC。CUDA程序是GPU和CPU的混合代码，它经NVCC编译分离出GPU和CPU代码，CPU代码则由CPU编译执行，GPU代码被编译成GPU计算汇编码PTX且经CUDA Driver支持在GPU上运行。

Device上的代码被称为Kernel。Kernel不是一个完整的程序，而只是一个实现并行算法的关键函数。Kernel的编程模型是以CUDA线程层次结构为依据的。线程层次结构包括：线程、线程块和网格。一个Grid由若干个线程块以阵列形式组成，目前可以最多有65535×65535个线程块组成。一个Block 由若干个线程以阵列形式组成，目前可以一个Block最多可以有512个线程。Kernel是以一个网格的形式执行的。同一网格内的Block不可以相互通信，只能通过global memory共享数据。同一线程块内的Thread可以通过shared memory和同步实现通信。

2JPEG在GPU上实现

JPEG编码核心步骤是DCT变换和量化，同时DCT变换是多个矩阵相乘问题，是最耗时、最耗资源的。通过GPU快速实现DCT变换和量化，最能体现GPU优势所在。DCT公式如下所示

ｕ

i=0

7

c o s

u

16

c o s

v

16

cudaMemcpy2DToArray(Src,0,0,ImgSrcF,ImgSrcFStride* sizeof(float),Size.width,Size.height,cudaMemcpyHostToDevice) cudaBindTextureToArray(TexSrc,Src)

划分MCU：根据CUDA的线程组织结构特点，一批线程Threads可划分线程块Block，多个线程块Block组成格Grid，每个Thread，Block都有自己的ID。同一批线程执行GPU代码称为内核函数kernel。MCU是个8×8的最小编码单元，一幅图片像素阵可划分多个MCU。这样，一个MCU可交由一批线程处理，多个MCU就可由线程块Block及格Grid处理。在GPU的架构中，一个线程块最多可以包含512线程。从数据层面上看，一个线程块就可以同时处理8个MCU。所以，GPU 可以同时并行处理多个MCU。自定义线程块和网格的大小，每个线程有惟一标识，相关伪码如下：

dim3Block(BLOCK_SIZE,BLOCK_SIZE);

dim3Grid(Size.width/BLOCK_SIZE,Size.height/BLOCK_ SIZE);

int tx=blockIdx.x*BLOCK_SIZE+threadIdx.x;

intty=blockIdx.y*BLOCK_SIZE+threadIdx.y;

DCT变换是整个JPEG中最重要的也是最耗时的一步，由DCT定义知，二维的DCT变换可通过两个一维的DCT相乘得到。DCT=A T XA(A是余弦矩阵，A T是A和转置矩阵)。矩阵相乘是可并行计算中最典型的例子，其经典的做法是矩阵分块相乘。图像像素阵可以分成若干个块，每个块分别交由GPU 上不同的Block处理。GPU即可以并行处理多个像素分块相乘，运算速度大大提高。块并行运算核心代码如下：for(int k=0;k<BLOCK_SIZE;++k)

{Csub+=AS(ty,k)*BS(k,tx);}

Int c=wB*BLOCK_SIZE*blockIdx.y+BLOCK_SIZE* blockIdx.x;

C[c+wB*ty+tx]=Csub;

量化：量化过程是DCT系数除以量化表中相应的元素，是真正实现图像有损压缩的过程。通过Thread.ID来惟一标识数组和量化表中的元素，这样数据元素就可以映射到并行处理线程中。量化过程中的64个除操作也就可以通过线程并发的执行。系数和量化元素读取，量化过程伪码为：Coef=SrcDst[ty*Stride+tx]

Quant=Q[threadIdx.y*BLOCK_SIZE+threadIdx.x]

float quantized=round(Coef/Quant);

Zig-zag扫描：Zig-zag扫描是并行性差的锯齿排序过程。虽然CUDA架构是CPU+GPU的异构运算平台，可以将Zig-zag扫描放在CPU上执行，但是主机(CPU)和设备(GPU)间数传输开销很大，应尽力最小化主机和设备之间数传输。所以，应将Zig-zag扫描放在GPU中执行，来为下一步的huffman编码作准备。每个MCU的数据是8×8，相对应的Zig-zag下标是固定的。将下标数组Zz_coordinate声明为share型，所有Block 都可以使用，Block中的每个线程通过thread.ID访问已量化好的图像数据并按Zib-zag下标Zz_coordinate重新输出。相关伪码如下：

_share_float Zz_coordinate[64]={0,1,8,16,9,2…63}

HufDst[Zz_coordinate[threadIdx.y*BLOCK_SIZE+threadIdx.x]]=SrcDst[threadIdx.y*BLOCK_SIZE+threadIdx.x] Huffman编码：Huffman编码输入的是一个二元组的数组，同样可以用Thread.Id奇数项来标识0-length和偶数项标识value。惟一的DC系数Thread.id是为2的线程，它单独到DC表中查找编码。在AC编码中，剩余数组中的元素就通过Thread.id 并发地去查找给定的AC系数huffman表中相应的编码。

2.2优化

并行程序整体性能的优化主要有三方面策略：程序最大并行化执行，优化内存利用率获得最大内存带宽，优化指令利用率获得指令最大化吞吐量。上述JPEG编码实施中，已经基本实现了程序最大化并行执行，同时也伴随着对内存进行大量的并行读写操作，这样容易造成系统瓶颈。

在内存优化方面，充分利用常量内存和共享内存，少用设备内存。常量内存空间具有高速缓存。因此，将公用的常量，一维余弦阵及逆阵、量化表、huffman表放在常量内存，提高读取速度。共享内存是位于芯片上，访问共享内存与访问寄存器速度一样快，把MCU及DCT系数放在共享内存中，并发读写速度很快。

在指令优化方面，用移位方法实现线程ID操作，使用CUDA内部函数，如量化操作中的四舍五入函数。

3CPU平台算法

为了方便从直观上比较CPU和GPU在此问题上的计算能力，设计了DCT在CPU平台上的算法。计时器采用的是performance counter，它的计算CPU时钟数较精确，CPU算法是传统的通过多重循环实现的矩阵乘法。核心代码如下：QueryPerformanceCounter((LARGE_INTEGER*)&timer. performance_timer_start);

for(int i=0;i+BLOCK_SIZE-1<Size.height;i+=BLOCK_ SIZE)

for(int j=0;j+BLOCK_SIZE-1<Size.width;j+=BLOCK_ SIZE)

{mult8x8(DCTv8matrixT,BLOCK_SIZE,fSrc+i*Stride+j, Stride,tmpblock,BLOCK_SIZE);

}

QueryPerformanceCounter((LARGE_INTEGER*)&time);

return((float)(time-timer.performance_timer_start)*timer. resolution)*1000.0f;

4实验结果及分析

硬件平台：GPU为NVIDIA GeForce9600GT；CPU为AMD Athlon(tm)64x2Dual Core processor5000+；内存为2GB。

软件平台：操作系统：Windows XP；开发软件：VS2008；环境支持：CUDA驱动、工具包、开发包，CUDA_VS_Wizard_ W32.2.0。

选取256×256、512×512、1024×1024这3个BMP文件进行实验，因CPU和GPU上的算法是按照传统定义来实现的，图像的压缩质量基本上是相同的，CPU和GPU对比实验可以从时间上来比较分析。结果如表1所示。

(下转第3308页)