CN103700123A - 基于cuda架构的gpu加速x光图像重建方法和装置 - Google Patents

Abstract

本发明公开了一种基于CUDA架构的GPU加速X光图像重建方法和装置，所述方法通过将获取的X光图像投影数据通过CPU传输至CPU内存中；通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；通过GPU并行核函数加速shiftandadd算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；将局部重建图像拼接合成为重建图像；通过CUDA架构下的GPU硬件加速实现并行shiftandadd算法，大幅度提高了运算速度并同时节省了成本，给用户带来了大大的方便。

Description

基于CUDA架构的GPU加速X光图像重建方法和装置

技术领域

 本发明涉及医疗图像处理领域，尤其涉及的是一种基于CUDA架构的GPU加速X光图像重建方法和装置。

背景技术

随着X光设备的不断发展，现有的X射线装置通过球管和平板，在固定SID（X线装置系统标识码系统标识）的情形下，通过线阵球管的一次扫描发射射线，获得多帧X光图像。然后针对多帧X光图像通过shift and add算法重建出断层图像。而现有采用的是传统串行shift and add算法或FPGA硬件加速shift and add算法来进行图像重建，但串行shift and add算法来进行图像重建的方法耗时过长，而FPGA硬件加速shift and add算法来进行图像重建的方法硬件成本过高。现有的X光图像重建技术中存在因数据量过大导致的算法耗时过长、以及硬件成本过高的问题，给用户带来了大大的不便。

因此，现有技术还有待于改进和发展。

发明内容

本发明要解决的技术问题在于，针对现有技术的上述缺陷，提供一种基于CUDA架构的GPU加速X光图像重建方法和装置，旨在解决现有的X光图像重建技术中因数据量过大导致的算法耗时过长、以及硬件成本过高的问题。

本发明解决技术问题所采用的技术方案如下：

一种基于CUDA架构的GPU加速X光图像重建方法，其中，包括以下步骤：

A、将获取的X光图像投影数据通过CPU传输至CPU内存中；

B、通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；

C、通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；

D、将局部重建图像拼接合成为重建图像。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，在所述步骤A之前还包括：

A10、通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，所述对获得的X光图像投影数据进行预处理具体包括：

A11、计算层析图像的像素之间的间距；

A12、计算两相邻源点对层析图像同一像素点所发的X射线被探测器平板所接收的位置之间的间距；

A13、计算同一源点对层析图像两相邻像素点所发X射线被探测器平板所接受的位置之间的间距。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，所述对线程进行分支预测具体包括：

A21、计算重建层析图像中心单个像素点所需源点的个数；

A22、根据层析图像不同位置处像素点所述源点的个数，通过CPU划分出层析图像的中心区域，并去除边缘区域。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，所述步骤B还包括：

B1、将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，相邻的投影数据区域设置为部分重叠。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，所述步骤B满足：

；

X表示总的数据量，表示分割后的数据量，K(x)表示核函数运行数据量为x时的运行时间，T(x)表示将数据量为x的数据从主机内存传输至GPU内存的时间。

所述的基于CUDA架构的GPU加速X光图像重建方法，其中，所述步骤C具体包括；

C1、分配CUDA核函数的线程块数以及每个线程块所包含的线程数，将每个线程与每个像素点一一对应，进行所有像素点的并行运算；

C2、对每一个线程进行平移、累加操作得到对应的局部重建图像；

C3、通过CPU控制多个CUDA流的运行顺序，同步运行当前CUDA流所进行的数据复制运算和相邻CUDA流所进行的核函数运算。

一种基于CUDA架构的GPU加速X光图像重建装置，其中，包括：

传输模块，用于将获取的X光图像投影数据通过CPU传输至CPU内存中；

分割模块，用于将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；

重建模块，用于通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；

拼接模块，用于将局部重建图像拼接合成为重建图像。

所述的基于CUDA架构的GPU加速X光图像重建装置，其中，还包括：

预处理模块，用于通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。

所述的基于CUDA架构的GPU加速X光图像重建装置，其中，所述分割模块，还用于将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，且相邻的投影数据区域有部分重叠。

本发明所提供的一种基于CUDA架构的GPU加速X光图像重建方法和装置，有效地解决了现有的X射线装置的图像重建过程采用传统串行shift and add算法或FPGA硬件加速shift and add算法导致的前者耗时过长而后者硬件成本过高的问题，其方法通过将获取的X光图像投影数据通过CPU传输至CPU内存中；通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；将所述多个局部重建图像拼接合成为重建图像；通过CUDA架构下的GPU硬件加速实现了并行shift and add算法，大幅度提高了运算速度并同时节省了成本，给用户带来了大大的方便。

附图说明

图1为本发明提供的基于CUDA架构的GPU加速X光图像重建方法较佳实施例的流程图。

图2a-2c为本发明提供的基于CUDA架构的GPU加速X光图像重建方法中CPU预处理shift and add数据的示意图。

图3为本发明提供的基于CUDA架构的GPU加速X光图像重建方法中分割数据重建区域图像的示意图。

图4为shift and add算法的原理示意图。

图5为本发明提供的基于CUDA架构的GPU加速X光图像重建方法中pipe-line模式的示意图。

图6为本发明提供的基于CUDA架构的GPU加速X光图像重建装置较佳实施例的结构框图。

具体实施方式

本发明提供一种基于CUDA架构的GPU加速X光图像重建方法和装置，为使本发明的目的、技术方案及优点更加清楚、明确，以下参照附图并举实施例对本发明进一步详细说明。应当理解，此处所描述的具体实施例仅仅用以解释本发明，并不用于限定本发明。

请参阅图1，图1为本发明提供的基于CUDA架构的GPU加速X光图像重建方法较佳实施例的流程图，所述方法包括以下步骤：

步骤S100、将X光图像投影数据通过CPU传输至CPU内存中；

步骤S200、通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；

步骤S300、通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；

步骤S400、将局部重建图像拼接合成为重建图像。

以下结合具体的实施例对上述步骤进行详细的描述。

在步骤S100中，将X光图像投影数据通过CPU传输至CPU内存中。具体来说，从X射线装置扫描的X射线强度数据获得X光图像投影数据，然后CPU将其发送至CPU内存中。也就是说，通过CPU读取该X光图像投影数据，并保存在主机（CPU对应的计算机）内存中。在实际应用时，通过DMA（直接内存存取）技术将X光图像投影数据从采集卡缓存空间地址直接复制到计算机内存空间地址。

进一步地，在所述步骤S100之前，还包括：步骤S90、通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。该步骤通过分析该X射线装置的SID、探测器分辨率、线阵球管扫描源间隔，计算出层析图像像素分辨率、层析图像像素相对探测器和球管的空间位置关系。具体来说，CPU还要对所述X光图像投影数据进行预处理，即预处理高精度浮点运算数据，并对线程进行分支预测。所述数据预处理具体包括：

计算层析图像的像素之间的间距；

计算两相邻源点对层析图像同一像素点所发的X射线被探测器平板所接收的位置之间的间距；

计算同一源点对层析图像两相邻像素点所发X射线被探测器平板所接受的位置之间的间距。

具体来说，请一并参阅图2a、图2b和图2c，图2a、图2b和图2c分别为CPU预处理shift and add数据的第一示意图、第二示意图和第三示意图。因侧重点不同，图2a、图2b和图2c的缩放比例略有差别。

为源面到探测器面的距离，

为层析平面距离线阵球管的距离，

表示缩放比例。

图2a中，

、

表示两个相邻的射线源，两条射向探测器中心的射线与层析平面相交于A、B两点，称其为主像素点，

、

称为A、B的主源点。A，B之间线性插值即可得到每一个像素点，假设插值数目N，则相邻两像素点之间的距离:

。

图2b中层析平面上两点

、表示两相邻的像素点，

表示两相邻源点对断层图像同一像素点所发X射线被探测器平板所接收的位置之间的间距，表示同一源点对断层图像两相邻像素点所发X射线被探测器平板所接受的位置之间的间距。；。

所述线程分支预测具体包括：

计算重建层析图像中心单个像素点所需源点的个数；

根据层析图像不同位置处像素点所述源点的个数，通过CPU划分出层析图像的中心区域，并去除边缘区域。

图2c中层析平面上点表示层析平面上中心的像素点，只有

到

之间的源是有意义的。对于图像边缘区域，显然过

点的源只有

一个，分支预测通过计算重建层析图像中心单个像素点所需源点的个数，由于重建层析图像边缘区域所需源点个数与中心区域不同，根据层析图像不同位置处像素点所述源点的个数不同，就可以计算排除掉与

源点数目不同的点。优选地，在CUDA架构下，每个线程代表一个像素点。为了确保每个线程对应源点个数相同,将由CPU划分出中心区域的大小，去除边缘区域。

在步骤S200中，通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存。具体来说，CPU合理分割主机内存中数据，并通过PCIE（一种总线和接口标准）总线将数据传输至GPU内存。在CUDA架构下的GPU加速运算模型中，GPU核函数的线程只能并行的访问GPU内存中的数据，因此需要将CPU内存中的数据传输至GPU内存中。传输速度受PCIE和显卡本身的传输速度所限，如表1所示，表1为在GTX460显卡和PCIE×162.0 的条件下，不同数据量的传输带宽。

数据量(MB)	传输时间(ms)	带宽(G/s)
			6.5	1.452064	4.47
13	2.805664	4.63
			22.75	4.854496	4.68
45.5	9.802944	4.64

表1

由表1可见，传输带宽基本是固定不变的，在PCIE 162.0的理论值范围内。

CPU分割数据指的是一次传输并不将所有X光图像投影数据传输至显卡，而只复制一部分。这样做一方面为了后续的pipeline运行方式(该方式将在后续步骤中详细阐述)；另一方面过大的数据量将消耗更多的线程数目从而增加运行时间。

受shift and add算法的限制，一块区域的源（即投影数据区域）只有中心一部分区域能重建出有意义的图像。因此两块相邻的重建图像所需源数据将会有重叠，需注意若分割过细会导致传输的总数据量增加。如图3所示，图3表征了分割总数据重建区域图像的过程。且图3中采用了将分为3块投影数据区域，相邻投影数据区域有部分重叠。

优选地，所述步骤S200满足以下条件：X表示总的数据量，表示分割后的数据量，K(x)表示核函数运行数据量为x时的运行时间，T(x)表示将数据量为x的数据从主机内存传输至GPU内存的时间，满足以下条件:

。

在步骤S300中，通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像。首先简单介绍一下shift and add算法，shift and add算法用于X光图像层析重建，该算法可以理解为由一系列投影数据平移累加获得层析图像，shift and add算法原理示意，请参阅图4，如图4所示，图中为射线源，为探测器。求解可得A点的像素值为：；B点为：

；通过shift and add算法可以求解各个层析平面上清楚的点，从而获得每一层的清晰图像。而在CUDA架构下GPU加速的原理为：在CUDA模型下，先由CPU将运算所需要的数据复制到GPU内存中，再由GPU内部核函数进行处理，最后将结果数据从GPU内存复制回CPU内存。运行在GPU上的程序称为Kernel(内核)函数，Kernel函数以Grid(线程网格)的形式组织，每个Grid由若干Block(线程块)组成，每个Block又包含若干Thread(线程)。在程序执行时，Grid被加载到SPA（流式处理器矩阵级别）上，Block则被分发到各个SM（流多处理器），Thread以Wrap(线程束，包含32个线程)为单位运行。Block可以是1维、2维或者3维。一个Block内的线程可彼此协作，通过共享存储器来共享数据，并同步其执行来协同存储器访问。通过Block的粗粒度分解和Thread的细粒度并行执行就可以完成任务的并行计算，达到加速的目的。

优选地，所述步骤S300具体包括：

分配CUDA核函数的线程块数以及每个线程块所包含的线程数，将每个线程与每个像素点一一对应，进行所有像素点的并行运算；

对每一个线程进行平移、累加操作得到对应的局部重建图像；

通过CPU控制多个CUDA流的运行顺序，同步运行当前CUDA流所进行的数据复制运算和相邻CUDA流所进行的核函数运算。

具体来说，传统的shift and add算法策略为：先计算当前源点对应断层平面像素点的像素值；再遍历所有源点并累加每个像素点的像素值。该算法顺序执行，每次耗时固定累加。

而本发明提供的基于CUDA架构下GPU加速的shift and add层析图像重建算法，从像素点出发，每次运算直接累加该像素点所需源点的探测器数据。CUDA架构下提供内建变量blockIdx.x和threadidx.x来进行寻址，使用二维的线程块直接索引像素点的坐标：

idx=blockIdx.x*blockDim.x+threadIdx.x；表示x方向像素点坐标；

idy=blockIdx.y*blockDim.y+threadIdx.y；表示y方向像素点坐标；

id=imagewidth*()+blockIdx.y*blockDim.y+threadIdx.y；表示一维数组中像素的索引。

这样就将每个线程与每个像素点进行了一一对应，为了最大限度的简化核函数避免浮点除法和多余的运算，核函数内策略为：

S310、令I(idx, idy)表示当前线程下所重建的像素点，分别计算idx和idy所属的主像素点坐标，因为线阵球管在y方向没有分部：

；

S320、计算I(idx, idy)偏移I(

,

)的像素个数：

；

S330、计算I(idx, idy)相对I(

，

)在探测器上的接收点偏移量。

尽管当前像素点和主像素点使用的源都是相同的，但是像素的偏移会导致在探测器上接收到源数据的点整体发生偏移，偏移量为：

；

S340、再循环遍历源点重建一个像素点I(idx, idy)。在遍历过程中因为主源点对主像素点发出的射线是经过探测器中心的，而I(idx, idy)相对I(

,

)在探测器上的接受点偏移量已经计算出，两相邻源点对断层图像同一像素点所发X射线被探测器平板所接收的位置之间的间距

已经在S90中详细给出。

表示当前源点在x方向距离主源点的源点偏移个数。则当前源点在探测器上接收点偏移探测器中心的距离为

；

则：当前源点对像素点idx, idy所发射线在探测器上接收点距离探测器中心的偏移量：

；

表示探测器上,

的读数，则idx, idy处的像素值为所有源点的累加：

。

经过S310到S340的步骤，一个像素点的重建就完成了，在该核函数中，通过相对位移偏移量累加的方式避免每次重复运算。CUDA体系下所有像素点同时并行运行，完成对一个区域的断层重建。本发明使用二维的线程块直接索引像素点的坐标。建立层析图像每个像素对线程的索引，多线程同时计算达到并行运算的效果。计算层析图像中一个像素与多帧探测器原始图像像素之间的对应关系，将多帧原始图像的像素点累加获得层析图像中的一个像素，该过程为多线程并行处理，大大提高了重建速度。请参阅表2，表2给出了不同图像大小下分割数据传输的速度和核函数运算的速度。

表2

由表2可知，若在本发明中重建图像大小为900×90，根据对X光图像投影数据的不同的分割方法，譬如1块，即不分割；2块，即分割成有交叠两部分；3块，即分割成3块，且相邻两块有部分数据交叠（如图3所示），那么根据表2的实际测试结果得知，当所述分为三个部分重建时，速度最快。此时每次重建大小为300×90。

然后，通过CPU控制多个CUDA流的运行顺序，同步运行当前CUDA流所进行的数据复制运算和相邻CUDA流所进行的核函数运算。具体来说，CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。在一个流中可以添加操作，例如核函数启动，内存复制，以及事件的启动和结束。GPU的硬件基础有两个基本的引擎，内存复制引擎和核函数执行引擎，这两个引擎是可以并行运作的，前提是核函数访问的数据必须是已经复制好的，图5给出了pipe-line的示意图：单流即表示不使用pipe-line,多流控制过程中；stream0的控制下第一步先进行数据的分割传输；接着第二步stream0控制核函数运行，stream1进行数据的分割传输，这两个行为是并行完成的；第三步stream1进行核函数运性，stream2进行数据的分割传输，最后stream2进行核函数的运行。显然，采用pipeline技术，即多流的方式，耗时大大减少。

令

表示数据传输耗时，

表示核函数运行耗时，N,表示分割的数据块数，则总耗时为：

。

优选地，为实现后续步骤S300中的算法的pipe-line运行，本发明将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，且相邻的投影数据区域设置为部分重叠。请继续参阅图3，优选分割为3块。请继续参阅图5，图5还表示了不采用pipeline和采用pipeline技术的时间对比，GPU并行运算的耗时主要分为两块，裸数据从内存传输至显卡的时间及并行运算的时间，步骤S400采用Pipeline的工作方式使用三条工作流，串行重建三部分区域层析图像，并最终合成为一幅，进一步减少了层析图像重建时间。

在步骤S400中，将所述局部重建图像拼接合成为重建图像。具体来说，在步骤S200中，CPU将所述X光图像投影数据分割为多块投影数据区域，经过步骤S300对每块投影数据区域进行层析重建之后，分别得到对应的部分重建图像，然后将所有的局部重建图像拼接合成为重建图像。请继续参阅图3，如图3所示，本发明将X光图像投影数据分为3个数据区域，且相邻两区域数据有部分重叠，每个数据区域经过上述运算对应生成一个局部重建图像，然后将所有的局部重建图像合成起来便是用户所需的完整的重建图像。请参阅表3，表3给出了在本发明提供的GPU加速与传统的串行shift and add耗时对比。

表3

由表3可知，相比传统的串行shift and add耗时，本发明提供的基于CUDA架构的GPU加速X光图像重建方法耗时大大减少；而在GPU加速中采用分割块数为3块时，CUDA加速耗时最短。因此本发明中优选分割为3块。

本发明采用的X射线装置，该X射线装置：线阵球管作为X射线源，该球管在水平方向含有多个X射线源，探测器接收平面获得不同位置下的物体X光投影图像，shift and add算法将这些投影图像平移、累加最终获得待测物体的层析图像。

本发明提供的基于CUDA架构的GPU加速X光图像重建方法，通过CUDA架构下的GPU加速，改变了传统shift and add串行运算的模式。具体为:在CUDA编程模型中，采用CPU完成X光投影图像数据的传输控制、待建层析图像边缘分支预测、重建过程子任务的调度分配以及高精度浮点运算的部分；GPU通过对数据的并行访问以及合理调用核函数完成shift and add层析图像重建；其中CPU调动GPU复制引擎和运算引擎实现了任务级的并行处理保证了算法的pipe-line运行模式，GPU依靠其多线程并行计算能力，从重建图像每个像素角度出发并行完成shift and add的图像重建算法。相比传统的shift and add算法本，发明通过GPU加速大大降低了算法的运行时间，同时节省了大量的硬件成本。

基于上述基于CUDA架构的GPU加速X光图像重建方法，本发明还提供了一种基于CUDA架构的GPU加速X光图像重建装置，请参阅图6，图6为本发明提供的基于CUDA架构的GPU加速X光图像重建装置较佳实施例的结构框图，如图6所示，所述装置包括：

传输模块10，用于将获取的X光图像投影数据通过CPU传输至CPU内存中；具体如步骤S100所述；

分割模块20，用于将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；具体如步骤S200所述；

重建模块30，用于通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；具体如步骤S300所述；

拼接模块40，用于将局部重建图像拼接合成为重建图像；具体如步骤S400所述。

进一步地，所述的基于CUDA架构的GPU加速X光图像重建装置，其中，还包括：

预处理模块，用于通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。

进一步地，所述分割模块20，还用于将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，且相邻的投影数据区域有部分重叠。

综上所述，本发明提供的一种基于CUDA架构的GPU加速X光图像重建方法和装置，通过将获取的X光图像投影数据通过CPU传输至CPU内存中；通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；将所有的局部重建图像拼接合成为重建图像；通过CUDA架构下的GPU硬件加速实现了并行shift and add算法，在大幅度提高了运算速度的同时节省了成本，大大减少了层析图像重建时间，给用户带来了大大的方便，其实现方法简单，成本较低。

应当理解的是，本发明的应用不限于上述的举例，对本领域普通技术人员来说，可以根据上述说明加以改进或变换，所有这些改进和变换都应属于本发明所附权利要求的保护范围。

Claims (10)

1. 一种基于CUDA架构的GPU加速X光图像重建方法，其特征在于，包括以下步骤：

A、将获取的X光图像投影数据通过CPU传输至CPU内存中；

B、通过CPU将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；

C、通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；

D、将局部重建图像拼接合成为重建图像。

2.根据权利要求1所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，在所述步骤A之前还包括：

A10、通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。

3.根据权利要求2所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，所述对获得的X光图像投影数据进行预处理具体包括：

A11、计算层析图像的像素之间的间距；

A12、计算两相邻源点对层析图像同一像素点所发的X射线被探测器平板所接收的位置之间的间距；

A13、计算同一源点对层析图像两相邻像素点所发X射线被探测器平板所接受的位置之间的间距。

4.根据权利要求2所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，所述对线程进行分支预测具体包括：

A21、计算重建层析图像中心单个像素点所需源点的个数；

A22、根据层析图像不同位置处像素点所述源点的个数，通过CPU划分出层析图像的中心区域，并去除边缘区域。

5.根据权利要求1所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，所述步骤B还包括：

B1、将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，相邻的投影数据区域设置为部分重叠。

6.根据权利要求5所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，所述步骤B满足：

；

X表示总的数据量，

表示分割后的数据量，K(x)表示核函数运行数据量为x时的运行时间，T(x)表示将数据量为x的数据从主机内存传输至GPU内存的时间。

7.根据权利要求5所述的基于CUDA架构的GPU加速X光图像重建方法，其特征在于，所述步骤C具体包括；

C1、分配CUDA核函数的线程块数以及每个线程块所包含的线程数，将每个线程与每个像素点一一对应，进行所有像素点的并行运算；

C2、对每一个线程进行平移、累加操作得到对应的局部重建图像；

C3、通过CPU控制多个CUDA流的运行顺序，同步运行当前CUDA流所进行的数据复制运算和相邻CUDA流所进行的核函数运算。

8.一种基于CUDA架构的GPU加速X光图像重建装置，其特征在于，包括：

传输模块，用于将获取的X光图像投影数据通过CPU传输至CPU内存中；

分割模块，用于将所述X光图像投影数据分割为多块投影数据区域，且使相邻的投影数据区域存在部分重叠，并将所述多块投影数据区域发送至GPU内存；

重建模块，用于通过GPU并行核函数加速shift and add算法对每块投影数据区域进行层析重建，得到相应的局部重建图像；

拼接模块，用于将局部重建图像拼接合成为重建图像。

9.根据权利要求8所述的基于CUDA架构的GPU加速X光图像重建装置，其特征在于，还包括：

预处理模块，用于通过CPU对获得的X光图像投影数据进行预处理，并对线程进行分支预测。

10.根据权利要求8所述的基于CUDA架构的GPU加速X光图像重建装置，其特征在于，所述分割模块，还用于将CPU内存中的X光图像投影数据分割为大小相等的多块投影数据区域，且相邻的投影数据区域有部分重叠。

Images