中国论文发表网为您提供优质服务，专业成就品牌。
   摘 要：三维锥束CT 图像重建运算量大，重建时间较长。本文提出了一种运用GPU 的统一计算设备架构(CUDA)实现用于三维锥束CT 图像重建的FDK 算法的方法。与CPU 实现方法相比，本文的方法在滤波卷积阶段的加速比达到了10 倍，而在耗时最多的反投影阶段加速比高达77 倍，重建整体加速比达到30 倍。而与传统的图形API 加速方法相比，本文的方法更易于实现，效率更高。
　　关键词：GPU，CUDA，图像重建，FDK 算法．
   引言采用三维锥束重建技术，具有空间纵向分辨率高、扫描速度快以及射线利用率高等方面的优势，但是其运算量庞大，重建时间较长，一般情况下难以满足实际应用。重建速度逐渐成为制约Micro CT 技术发展的瓶颈。目前，国内外出现了很多三维锥束CT 重建的加速算法，也出现了一些利用图形处理器(Graphics Processing Unit，GPU)加速三维锥束图像重建的应用。
　　和Fang Xu 等人对利用GPU 加速锥束CT 图像重建做了一系列研究[1] [2]。
　　和Klaus Mueller 使用1-D 纹理映射及2-D 纹理映射实现了FDK 算法[1]，并提出使用定点数光栅化处理及浮点数叠加相结合的方法来解决浮点数管道的运算速度明显慢于定点数管道的问题，以兼顾图像质量和运算速度[1]；为了进一步提高GPU 的加速性能，和Fang Xu 利用RGBA 通道的并行机制提高了FDK 算法的重建速度[2]。在这些加速算法中, 投影和反投影由GPU 的纹理映射过程执行，距离加权运算通过将加权系数存为纹理来实现。戴智晟等人在GPU 上近似地实现了T- FDK 算法加速[3]。之后，马车平，曾理等人提出了一种利用GPU 多重纹理加速FDK 图像重建的方法[4]。
　　但是，以上这些研究利用的GPU 通用计算模型都是基于复杂的图形处理API 函数，如：
　　和Direct 3D 等，这种传统的GPU 加速方法必须首先将算法映射到图形API 函数，因此其实现比较复杂。本文利用NVIDIA 公司的GPU 通用计算模型——统一计算设备架构，CUDA)，提出了一种易于实现的基于CUDA 的三维锥束图像重建方法。该方法将FDK 重建过程分为滤波和反投影两步，分别利用实现，首先将投影图像数据传输到GPU 内存，然后对每幅图像做细粒度的并行滤波，之后将滤波后的图像数据绑定到GPU 纹理，并对其做并行反投影，最后将结果图像传输到端内存。本文的方法与仅利用CPU 实现的图像重建方法相比达到了较高的加速比。而由于CUDA 是C 语言的扩展，因此不用像传统的图形API 加速方法那样必须将算法映射到图形处理函数，因此与传统的GPU 图形API 方法相比，本文的方法更加易于实现，效率更高。
　　．CUDA 实现FDK 图像重建的基础三维锥束 CT 重建最常用的算法是FDK 算法，该算法易于理解，便于实现，而且很容易将该算法分解为滤波和反投影等几个相对独立的部分。三维锥束图像重建数据量大，运算时间较长，但是，其待处理数据和计算结果都是相对独立的图像像素，很适合于细粒度的并行计算，因此，很容易想到利用GPU 强大的并行计算能力加速FDK 图像重建，而使得我们可以在充分利用GPU 的计算能力的同时避开繁琐的图形API，因此成为GPU 通用计算的首选。下面将分别介绍CUDA 实现FDK 图像重建的两个理论基础。
　　．FDK 图像重建算法简介算法是L.A.Feldkamp，L.C.Davis 和J.W.Kress 等人于1984 年提出的一种锥束三维重建算法[5]，该算法由于易于理解，便于实现而得到广泛的应用。FDK 算法投影几何。
　　．GPU 和CUDA 简介图形处理单元(graphics processing unit，GPU)已经成为现代主流计算机系统不可或缺的一部分。过去七年间，随着对实时、高分辨率图形处理的需求，可编程图形处理器不断发展，成为拥有超强运算能力和高带宽的高度并行化，高度多线程的多核处理器[6]。由于GPU 在可编程能力和并行运算能力方面的飞速发展，目前，其应用已经逐渐超越计算机图形学，而在各种计算密集型研究领域得到广泛应用。GPU 在通用计算中的应用，GPGPU)使得GPU 极有可能在未来的高性能计算领域取代传统的。
　　然而，传统的GPU 通用计算模型基于复杂的图形处理API 函数，如：OpenGL 和等，要利用GPU 做通用计算的话，必须首先将算法映射到图形API 函数，这使得没有图形处理基础的研究人员对GPU 望而却步。
　　是NVIDIA 提出的一种GPU 上的并行编程模型和软件环境，它将GPU 视为数据并行计算设备，在其上进行计算的分配和管理，而无需将其映射到图形API。
　　编程模型将GPU 看作并行运行多个线程的计算设备(compute device)，是主或称为宿主，host)的协处理器。CUDA 程序经编译器NVCC 编译后，被分为两部分：
　　在主CPU 上运行的Host 部分和在GPU 上运行的Device 部分。Host 部分一般为控制流程的串行程序，Device 端则是计算密集的并行程序。
　　端运行的是被称为内核(kernel)的C 语言函数，调用此函数时，它将以N 个不同的线程被并行执行N 次。这N 个线程被组织成很多个线程块(Thread block)，所有的线程块组成一个线程网格(Thread grid)。一个块内的线程可彼此协作，通过共享存储器来共享数据，并同步其执行来协调存储器访问。更具体地说，可以通过调用 __syncthreads()内置函数在内核中指定同步点；__syncthreads()起到屏障的作用，块中的所有线程都必须在这里等待处理。这里的线程组层次结构、共享存储器、屏蔽同步(barrier synchronization)是CUDA 的三个核心抽象概念[6]，这些抽象提供了细粒度的数据并行化和线程并行化，嵌套于粗粒度的数据并行化和任务并行化之中，细粒度并行计算区别于传统的粗粒度或任务级并行化，它允许在粗粒度的并行化中将任务划分为更小的子任务，以便更高效利用GPU 的并行计算资源。
　　在这一模型中，问题分解为更小的片段，以便通过协作的方法并行解决。这样的分解保留了语言表达，允许线程在解决各子问题时协作，同时支持透明的可伸缩性，使得可以安排在任何可用处理器内核上处理各子问题[6]。因而，编译后的 CUDA 程序可以在任何数量的处理器内核上执行，只有运行时系统需要了解物理处理器数量。
　　卷积滤波的CUDA 实现卷积滤波只需对投影图做水平滤波。我们采用shepp-Logan 窗。因为GPU 中全局内存没有缓存，对全局内存的存取需要400-600 个指令周期，存取速度较慢，而且，只有当全局内存对齐访问时才能达到理想的访存速度，而对图像的每一行做卷积时，卷积核是相同的，因此我们把卷积核存入GPU 的常量内存(Constant Memory)，常量内存为只读存储器，但是拥有高速缓存，因此这样大大提高了访存效率。而对于每一个像素的计算，由于要用到整行的原始投影数据，因此对内存访问较为频繁，为了减少对全局内存的访问次数，达到提高计算效率的目的，我们把一行数据存入共享内存，线程组划分策略采取每个线程块计算结果图像中的一行。我们采用的投影图校正后的大小是840×600，这样每个线程块的纬度为×1)，则每个线程块中线程号为thread.x 的线程计算每行的第thread.x、第256+thread.x、和第768+thread.x 个像素。这样达到了很好的效果。
　　反投影的CUDA 实现反投影是 FDK 算法中计算量最大，最耗时的部分，包括投影坐标计算、距离加权和求和运算。CUDA 实现时，每个线程计算一幅投影图像对断层图中的一个像素的反投影，并将其累加到断层图像中的对应像素。每一个线程先计算投影坐标(Y(r), Z(r))，然后，在投影图中插值计算出(Y(r), Z(r))处的值即为其反投影。由于对投影图的插值坐标是随机的，随机访问全局内存的效率很低，因此，我们采用GPU 纹理操作提高内存访问效率，把滤波后的投影图像数据绑定到GPU 纹理，由于GPU 纹理拥有速度较快的纹理缓存，因此大大提高了反投影效率，实际操作的时候速度比直接对全局内存区的数据插值加快了5 倍左右。然后计算出距离加权系数，对结果图像做距离加权，即得到最终断层图像。
　　．实验结果本文基于CUDA的FDK图像重建硬件平台配置为：Intel Xeon 四核CPU，主频1.6GHz，为NVIDIA 公司的Tesla C870。Tesla C870 拥有128 个流处理器，浮点处理能力可以达到430 GFLOPs（Giga floating point operations per second），拥有1.5G 设备内存，内存最高频宽可达76.8GB/秒。
　　．结束语本文提出了一种运用GPU 的通用计算模型(CUDA)加速FDK 锥束CT 图像重建的方法。该方法首先将FDK 算法分解为卷积滤波和反投影两部分，然后分别对这两部分利用实现细粒度的并行计算。实验结果证明，本文提出的方法在不牺牲图像质量的情况下极大地提高了FDK 三维锥束CT 重建的效率，达到了很好的效果，对于400 幅投影图像，投影图像大小为840×800，重建100 幅大小为840×840 的断层图像时，滤波阶段的效率提高了10 倍左右，反投影阶段效率提高多达77 倍，重建整体加速比达到大约30 倍。与传统的基于Open GL 等图形处理语言的GPU 加速相比，我们的CUDA 实现方法更易于实现，效率更高。