基于GPUJPEG2000高速图像解压缩研究

基于GPU的JPEG2000高速图像解压缩系统研究专业名称：通信与信息系统

姓名：杨国鹏

学号：1101120210

指导教师：李云松教授绪论JPEG2000静止图像压缩系统介绍GPU硬件架构及CUDA编程介绍解码软件系统的设计与实现解码软件系统的测试与分析总结与展望内容提要绪论选题背景近年来，图像压缩芯片发展迅速，目前应用于航天航空领域的图像压缩芯片每秒可以输出数十兆的压缩码流。为了能实时恢复图像，地面接收站通常是采用多台服务器同时解码来完成。但是，这种方式需要专门的机房，占用空间多且耗电量大。近年来，GPU在高性能并行计算领域发展迅速，因此考虑用GPU实现。绪论GPU与CPU的计算能力比较绪论GPU与CPU的存储带宽绪论选题意义因此，采用GPU解码，实现高速解压缩所需的硬件资源将大大减少，进而占用的空间也将大大减少，实现了轻便化。JPEG2000静止图像压缩系统JPEG2000编解码系统结构GPU硬件架构GPU的硬件架构是一种两层结构，通常一块显卡由多个流处理器SM组成，各SM相互独立。每个SM都由指令缓存、指令分发和调度单元、流处理器SP、特殊功能单元SFU、双精度处理单元DFU和共享存储器等组成。SM之间通过全局存储器通信。一个SM中包含多个SP,每个SP都有一个32bit的算术逻辑单元和一个乘加器。SP是SM中的主要运算单元。GPU硬件架构及CUDA编程Fermi架构SM结构图GPU硬件架构及CUDA编程CUDA编程CUDA编程也是两层结构。运行在GPU上的函数称为kernel。GPU以线程网格Grid的形式组织线程，Grid由多个线程块block组成，每个block由多个线程thread组成。每个线程都会执行kernel函数，并且每个线程都有自己的block和thread标示号。对于tid=(blockIdx.y*gridDim.x+blockIdx.x)*blockDim.x*blockDim.y+(threadIdx.y*blockDim.x+threadIdx.x)这个式子，每个线程会得到一个不同的tid值。对于A[tid]=A[tid]+8，可实现单指令多线程的处理方式。GPU硬件架构及CUDA编程Grid结构GPU硬件架构及CUDA编程GPU软硬件结合在执行时，block被平均分发到SM上。一个SM上可以执行多个block，各block是相互独立；当SM在执行其中一个block时，其它block处于准备状态，当正在执行的block执行高延迟的操作时，SM可以切换执行另外一个block，可以有效地隐藏流水线延迟。Block中的线程被分发给SM中的SP去执行。硬件结构和软件编程的配合，实现了GPU较高的并行计算能力。GPU硬件架构及CUDA编程解码系统框架解码软件系统的设计与实现

解码服务器上有6块高性能显卡，每个显卡都可以独立地完成解码，可以根据需要的解码速度，启动相应数量的显卡。码流发送客户端可以通过网络向解码服务器提供JPEG2000压缩码流，因此解码服务器可以解远程码流。图像显示客户端的功能是显示解码服务器恢复后的图像数据。JPEG2000压缩码流在单GPU上的实现反量化和后处理对每个像素点都是做同样的操作，只是操作数不同，可并行度特别高，所以grid结构设计为多个block，每个block处理一行，block里的各线程处理对应行中的像素点。在JPEG2000中，二维图像的离散小波变换采用Mallat分解结构，先按行进行水平方向的一维小波变换，再按列进行垂直方向的一维小波变换。解码是压缩的逆过程。由于按行进行变换时，各行可以同时进行变换，按列进行变换时，各列也可以同时进行变换。所以grid结构设计为多个block，每个block处理一行，block里的各线程处理对应行中的像素点。解码软件系统的设计与实现

在进行小波变换时，数据时按行存储在全局存储器中，因此在进行行处理时，block中的各线程访问的数据是连续的，通过合并访问，通过一次读取全局存储器就可以得到所需要的数据。在进行列处理时，各线程访问的数据不连续，无法使用合并访问，访问数据速度很慢。为了加快速度，可以借助共享存储器使用矩阵转置，把列转为行，这样就可以连续访问。解码软件系统的设计与实现

矩阵转置解码软件系统的设计与实现

JPEG2000解码在单个GPU上的并行架构解码软件系统的设计与实现

多GPU联合解码单个GPU的解码方式，解码速度已经提高了很多，但有时需要更高的解码速度，这时就需要多个GPU联合解码。在CPU上分配多个线程：分配一个CPU线程专门负责将n幅图像的压缩码流组织在一起，把该线程称为写码流线程；分配一个CPU线程专门负责将恢复后的图像数据存盘，把该线程称为写文件线程；为每个GPU分配一个CPU线程，该CPU线程使用cudaSetDevice()函数启动自己负责的显卡，把该线程称为解码线程。定义了几块内存空间，每块内存空间用于一组n幅图像的解码，并且定义三个队列——等待写码流队列、等待解码队列以及等待写文件队列。解码软件系统的设计与实现

内存块地址的流转解码软件系统的设计与实现

码流发送模块码流发送端通过网络给解码服务器发送码流，因此解码服务器可以解远程码流。给码流发送端做了一个简单的界面，界面采用wxWidgets语言开发。解码软件系统的设计与实现

图像显示模块显示模块提供显示功能，显示恢复后的图像数据。服务器跟显示客户端的通信通过网络实现。解码软件系统的设计与实现

在菜单栏包含“文件”、“工具”和“帮助”三个菜单命令。“文件”包含“打开”和“退出”，“打开”可以静态地显示本地的图像，“退出”用于退出界面。“帮助”显示帮助信息。“工具”包含“参数设置”、“开始”、“暂停”和“停止”。解码软件系统的设计与实现

解码服务器硬件平台：CPU:两个6核cpu，软硬件配置：解码软件系统的测试与分析

型号主频缓存内存硬盘操作系统Inter®Xeon®E56452400.144GHz12288KB40GB1.8TBCentOS64bitGPU:6块C2075显卡，软硬件配置：显卡名称C2075CUDA驱动版本/运行版本4.0/4.0计算能力2全局存储器5GB多处理器

xCUDACores/多处理器14*32CUDACoresGPU时钟频率1.15GHz共享内存/多处理器49152Bytes单GPU解码速度测试：28幅2048*1024图像，8bit、12bit、16bit，LeGall(5,3)无损压缩，Dauhechies(9,7)2倍和4倍压缩。解码软件系统的测试与分析

码流大小(MB)t2(ms)t1(ms)反量化+小波+电平位移(ms)Total(ms)解码速度

(M像素点/s)解码速度(MB/s)8bit_53无损27.246.02773.3873.45974.5457.4627.918bit_97二倍28.346.28844.18113.071070.1952.3326.448bit_97四倍14.243.88588.18113.40823.4568.0117.2412bit_53无损53.949.201138.1973.331323.6142.3140.7212bit_97二倍41.647.761032.35113.691260.2644.4433.0112bit_97四倍21.245.68769.86113.701001.7655.9021.1616bit_53无损59.651.161310.8573.191501.5537.2939.6916bit_97二倍41.147.841044.41113.291272.2844.0232.3016bit_97四倍28.246.39891.36113.601127.8549.6525.00单GPU解码速度测试结果分析：解码速度（M像素点/s）指的是用28图像总的像素点除以Total的时间，反映的是像素点的吞吐率；解码速度(MB/s)指的是用28幅图像码流的大小除以Total的时间，反映的是压缩码流的吞吐率。从每秒解多少像素点来反映解码速度，可以看出压缩倍数高，解码速度快；比特精度越小，解码速度越高。从每秒解多少字节的压缩码流来反映解码速度，可以看出对于同样宽度和高度的图像，压缩倍数小和比特精度高，码流吞吐率高。解码软件系统的测试与分析

多GPU解码速度测试：多幅28幅2048*1024图像，8bit、12bit、16bit，LeGall(5,3)无损压缩，Dauhechies(9,7)2倍和4倍压缩。不存储恢复图像数据。不同GPU的解码速度：解码软件系统的测试与分析

1个GPU2个GPU3个GPU4个GPU5个GPU6个GPU8bit_53无损58.17112.07166.79217.13275.64309.508bit_97二倍52.6599.54153.75197.43246.16281.348bit_97四倍69.27128.46191.69258.18316.40333.7912bit_53无损41.7483.08123.33162.83191.61228.2912bit_97二倍44.2586.44127.36168.19207.01241.6112bit_97四倍56.95107.23163.54209.98258.20296.0016bit_53无损36.8472.27104.75138.14170.88196.6416bit_97二倍43.9686.39126.92169.24204.01239.6216bit_97四倍49.6995.98142.34185.01231.31262.43解码软件系统的测试与分析

不同GPU相比于单GPU的解码倍数1个GPU2个GPU3个GPU4个GPU5个GPU6个GPU8bit_53无损1.001.932.873.734.745.328bit_97二倍1.001.892.923.754.685.348bit_97四倍1.001.852.773.734.574.8212bit_53无损1.001.992.953.904.595.4712bit_97二倍1.001.952.883.804.685.4612bit_97四倍1.001.882.873.694.535.2016bit_53无损1.001.962.843.754.645.3416bit_97二倍1.001.972.893.854.645.4516bit_97四倍1.001.932.863.724.665.28在GPU数量为2、3、4时，解码倍数是2、3、4倍；在GPU数量为5时，对比1个GPU，解码倍数在4.5到4.7倍之间；在GPU数量为6时，对比1个GPU，解码倍数在5.5倍之下。由于内存块的数目都是10，当GPU数目较多，对于内存块的竞争比较激烈，所以在4个GPU的基础上，再往上增加1个GPU，解码速度增加的幅度有所减少。解码软件系统的测试与分析

6个GPU，不存储恢复图像与存储恢复图像的速度比较不存储恢复图像数据存储恢复图像数据8bit_53无损309.50273.028bit_97二倍281.34252.418bit_97四倍333.79285.2316bit_53无损196.64145.2516bit_97二倍239.62135.8916bit_97四倍262.43148.93测试固态硬盘的写文件速度，得出结果：写文件速度在300MB/s到350MB/s之间，这个速度限制了6个GPU的解码速度。解码软件系统的测试与分析

CPU与GPU解码速度的比较CPU：Inter®Core™2DuoCPUE7500@2.93GHz，内存2.00GBGPU：C2075测试码流：2048*1024图像，8bit、16bit，LeGall(5,3)无损压缩，Dauhechies(9,7)2倍和4倍压缩。CPU单幅图像解码时间(ms)CPU解码速度1个GPU解码速度GPU/CPU8bit_53无损858.82.3357.4624.678bit_97二倍1927.21.0452.335