专利名称：一种用gpu实现快速小波变换的方法
技术领域：
本发明涉及图像处理技术领域，具体涉及一种使用计算机图形处理器GPU实现快 速小波变换的方法，以适用在通用计算机上实现高速的图像编码。
背景技术：
在数字图像处理领域，小波变换由于其在JPEG2000图像压缩标准中的优秀表现， 得到了广泛的应用。然而小波变换运算复杂，需要处理的数据量一般比较大，因而速度较 慢，严重限制了小波变换的应用。于是应当采取一定的优化措施。目前对小波变换的各种 优化多是基于专用集成电路ASIC或可编程逻辑器件FPGA的，具有开发周期长，难度大，成 本高等缺点。针对这一情况，有学者提出了在GPU设备上进行小波变换的方案。GPU设备凭借多 核共同驱动以及很高的内存带宽，获得了很高的运算能力。南京邮电大学陈大伟等人在其 文章《DaUbechieS9/7离散小波变换算法在GPU上的实现》(中国多媒体通信2007年12期 第45页到50页)中针对GPU设备特性，提出了一种快速小波变换方法。该方法用并行策 略实现小波变换，着重解决GPU设备共享内存限制的问题，将数据的每行或每列数据分段。 这种方法虽然具有较好的通用性，可以处理尺寸较大的图像，但是正如作者所言，不可避免 地引入了新的问题。其中最影响效率的就是，每行或每列数据各个段的段首和段尾都必须 要对称延拓。每次延拓都要一个线程单独地处理段首和段尾数据，而在此时间内，其他线程 皆是空闲状态，造成了资源的浪费。该方法的另一个缺点是，行变换之后要进行列变换，必 须将数据的每列提取出来，而该方法没有提出一个有效的重新排布数据的方法，导致列变 换效率不高。
发明内容
本发明针对上述问题，提出一种使用GPU实现的，不进行数据分段，且具有专门行 列转置功能的快速小波变换方法。本发明的步骤如下
第1步将需要小波变换的图像数据从计算机主机内存复制到GPU设备内存区域 A中；第2步使用水平滤波器，对第1步得到的GPU设备内存区域A中的数据进行行变 换，并将变换结果仍然保存于GPU设备的内存区域A中；第3步使用行列转置器，对第2步得到的GPU设备内存区域A中的数据进行行列 转置，将转置后的数据存储于GPU设备的内存区域B中；第4步使用水平滤波器，对第3步得到的GPU设备内存区域B中的数据进行再次 行变换，并将变换结果仍然保存于GPU设备的内存区域B中；第5步使用行列转置器，对第4步得到的GPU设备内存区域B中数据进行再次行 列转置，将转置后的数据存储于GPU设备的内存区域A中；
第6步将GPU设备内存区域A中的数据复制到计算机主机内存。本发明与现有技术相比，具有以下优点第一，本发明仅需要GPU设备共享内存缓存一行或一列的数据，而且随着GPU设备设计水平的提高，可供使用的共享内存也相应增加。因此，本发明不容易受GPU设备共享内 存不足的问题的限制。第二，本发明设计的水平滤波器充分使用了 GPU设备的多处理器和高速共享内存 资源，因而具有很高的处理速度。第三，本发明设计的行列转置器使用了 GPU设备的多处理器和高速共享内存资 源，因而具有很高的处理速度。通过使用此高速的行列转置器，可以避免进行列变换时由于 列数据存储不连续而导致的效率降低。同时，由于列变换通过行变换实现，本发明不再需要 单独的垂直滤波器来实现列变换，使得由本方法实现的系统更加简单。
图1是本发明的流程图。
具体实施例方式本发明采用CUDA语言，可以在NVIDIA的任何一款支持CUDA架构的GPU设备上实 现。在实施本发明所述的方法前，应先调用cudaMalloc函数在GPU设备上分配两个大小相 同的内存区域，分别记为内存区段A、B。在使用本发明所述的方法后，也应调用cudaFree函 数释放这两段内存区域。参照图1，本发明可通过以下步骤实现步骤1，将数据从计算机主机内存复制到GPU设备内存。数据从计算机主机内存到 GPU设备内存的复制可使用函数cudaMemcpy或cudaMemcpy2D来实现。步骤2，使用水平滤波器，对步骤1得到的GPU设备内存区域A中的数据进行行变 换，并将变换结果仍然保存于GPU设备的内存区域A中。 本发明设计了 一个行滤波器来完成对图像一行数据的行变换，这样水平滤波器就 可以通过使用多个行滤波器来实现整个图像数据的行变换。在对GPU设备资源的使用上， 每个行滤波器使用一个由256个线程组成的线程块，以及一行数据量大小的共享内存。通 过向行滤波器提供一行数据的内存首地址COef_rOW以及行宽w，行滤波器就可以完成对该 行数据的滤波处理。为完成上述功能，行滤波器的具体实施方法如下
2a)将一行的数据按奇数项与偶数项进行分离，即解交织，并将分离后的数据存储 到该线程块的共享内存里。本发明使用一种特殊的循环控制方法，使数据被多个线程并行 地处理。GPU设备每个线程块以及每个线程块中的所有线程都有一个索引值，本发明使用 此索引值完成循环的控制。因为行滤波器由一个线程块完成，所以设计循环的时候使用各 个线程在该线程块里的索引值threadldx. χ。设循环变量为i，i初值设置为threadldx. X，这样在该线程块里的256个线程里，i的初值分别为0至255。循环的控制条件设为i < width,这样，如果一行的宽度如果小于该块中线程的数量，则只有特定的线程会工作，如 果宽度大于该块中线程的数量，则会执行循环。每次执行循环体后i自增线程块内线程的 数量blockDim. X，这样，下次循环则刚好继续上个循环剩下的数据工作。通过这256个线程的协作，这一行的每个数据都将被依次处理，而需要的循环次数大大减少。对于在准备步骤 已经开辟的共享内存，将其等分为两段，前一段的首地址标记为S，后一段标记为P。按照上 述的循环控制方法，将COef_rOW地址偏移量为2i的数据放至s地址偏移量为i的位置；将 coef_row地址偏移量为2i+l的数据放至ρ地址偏移量为i的位置。经过至多w/2*256+l 次循环，即可完成数据的搬移。在结尾处应调用函数—syncthreads来实现一个块中所有 线程的同步，只有进行过同步，各个线程执行后的数据对其他线程才是可见的。2b)按照2a中所述的循环控制方法对解交织后共享内存中的数据分别进行四次 提升。第一次提升是将s[i]和s[i+l]两个数据相加，乘以提升系数后加到p[i]上，所有 数据处理完成后进行一次线程同步。第二次的提升与第一次过程类似，所不同的是将P[i] 和P[i+1]相加后乘以提升系数加到s [i+1]上。第三、四次的提升则与第一、二次的提升相 同。2c)按照2a所述的循环控制方法，将以上四步提升后存放于共享内存里s和ρ地 址开头的数据，即低频和高频数据分别乘以归一化系数送回共享内存并覆盖该行原来的数 据。 完成行滤波器的设计后，水平滤波器可通过使用512个行滤波器来完成所有行的 处理。水平滤波器的具体实现方法为按照行滤波器相同的循环控制方法，循环变量i的初 值为线程块的索引号blockldx.x，控制条件为i < h，每次执行后i自增512，在循环体内 部，向行滤波器提供一行数据的内存首地址coef_all+i*W及行宽w，使用行滤波器进行相 应行的滤波处理。水平滤波器使用512个行滤波器并行地工作，而每个行滤波器使用多个线程并行 地工作，实现了高度并行的数据处理。同时，行滤波器对数据的主要操作都在共享内存里完 成，因而数据具有很高的传输速度。高度并行的数据处理和高速的数据传输使得水平滤波 器具有极高的速度。步骤3，使用行列转置器，对步骤2得到的GPU设备内存区域A中的数据进行行列 转置，将转置后的数据存储于GPU设备的内存区域B中。在对GPU设备资源的使用上，行列转置器使用256个线程块，而每个线程块使用 256个线程以及一列数据量大小的共享内存。向行列转置器提供两个分别指示转置前和转 置后数据存放的位置的地址，以及图像的行宽与列高，行列转置器即可完成一幅图像的转置。行列转置器主要由一个二重循环来完成，外层循环调度整个GPU设备中所有的线 程块，内层循环调度一个线程块中的线程。外层循环的循环变量j初值设为线程块索引号， 控制条件为j小于行宽，每循环一次j自增256，循环体为内层循环。内层循环的循环变量 i的初值为该线程块中该线程的索引值，控制条件为i小于列高，每循环一次i自增256，循 环体完成数据搬移。数据搬移的过程是先将输入数据中地址偏移量为i*w+j的数据存放于 共享内存的第i个位置，再将共享内存中第i个位置的数据放至输出地址偏移量为i+j*h 的位置中。行列转置器通过二重循环，实现了以下两个功能第一，行列转置器通过循环控制 大量的线程，实现了高度并行的数据操作。第二，行列转置器不是直接将数据从输入地址的 相位位置搬移到输出地址的相应位置，而是先将数据缓存到共享内存中，再从共享内存中取出到输出位置，实现了高速的数据传输。高度并行的数据操作和高速的数据传输使得行 列转置器具有极高的速度。步骤4，使用水平滤波器，对步骤3得到的GPU设备内存区域B中的数据进行再次 行变换，并将变换结果仍然保存于GPU设备的内存区域B中。本步骤中的水平滤波器于步骤 2中水平滤波器完全相同，对转置后的数据进行水平滤波，相当于完成转置前数据的垂直方 向的滤波。步骤5，使用行列转置器，对步骤4得到的GPU设备内存区域B中数据进行再次行 列转置，将转置后的数据存储于GPU设备的内存区域A中。对步骤4提升结果重复步骤再 次行列转置操作，即完成了步骤3转置的恢复。步骤6，将GPU设备内存 区域A中的数据复制到计算机主机内存。将数据复制到计 算机主机内存后，原来的图像数据就被覆盖为小波变换后的数据。本发明通过以上步骤，实现了高速的小波变换，可用于在通用计算机上实现高速 的图像编码。
权利要求
一种用GPU实现快速小波变换的方法，包括第1步将需要小波变换的图像数据从计算机主机内存复制到GPU设备内存区域A中；第2步使用水平滤波器，对第1步得到的GPU设备内存区域A中的数据进行行变换，并将变换结果仍然保存于GPU设备的内存区域A中；第3步使用行列转置器，对第2步得到的GPU设备内存区域A中的数据进行行列转置，将转置后的数据存储于GPU设备的内存区域B中；第4步使用水平滤波器，对第3步得到的GPU设备内存区域B中的数据进行再次行变换，并将变换结果仍然保存于GPU设备的内存区域B中；第5步使用行列转置器，对第4步得到的GPU设备内存区域B中数据进行再次行列转置，将转置后的数据存储于GPU设备的内存区域A中；第6步将GPU设备内存区域A中的数据复制到计算机主机内存。
2.根据权利要求1所述的一种用GPU实现快速小波变换的方法，其中第2步与第4步 中所述的水平滤波器使用多线程进行并行的数据处理，并使用共享内存进行数据的传输。
3.根据权利要求1所述的一种用GPU实现快速小波变换的方法，其中第3步与第5步 中所述的行列转置器使用多线程进行并行的数据处理，并使用共享内存进行数据的传输。
全文摘要
一种用GPU实现快速小波变换的方法，包括(1)将数据从计算机主机内存复制到GPU设备内存区域A。(2)使用水平滤波器，对内存区域A中的数据进行行变换，并将变换结果保存于内存区域A。(3)使用行列转置器，对内存区域A中的数据进行行列转置，将转置后的数据存储于GPU设备内存区域B。(4)使用水平滤波器，对内存区域B中的数据进行再次行变换，并将变换结果保存于内存区域B。(5)使用行列转置器，对内存区域B中数据进行再次行列转置，将转置后的数据存储于内存区域A。(6)将GPU设备内存区域A中的数据复制到计算机主机内存。本发明使用了并行的数据处理，实现过程简单，提高了小波变换的处理速度。
文档编号G06T9/00GK101866493SQ20101020423
公开日2010年10月20日 申请日期2010年6月18日 优先权日2010年6月18日
发明者刘凯, 吴宪云, 宋长贺, 李云松, 王柯俨, 肖嵩 申请人:西安电子科技大学