CN101976207A - 一种面向gpu的数据流处理方法 - Google Patents

Abstract

本发明公开了一种面向GPU的数据流处理方法，首先，在CPU和GPU组成的异构并行系统中运行待处理的数据密集型计算模块，得到拷入时间Ti、计算时间Tc、拷回时间To；然后，根据三部分运行时间计算如下比例关系：a＝Tc/Ti，b＝Tc/To；最后，根据所总结出的将初始数据在划分为n个流时各数据流之间大小的比例：1∶a∶a²∶…∶a^i-1∶a^i-1/b∶a^i-1/b²∶…∶a^i-1/b^n-i，结合初始数据的总量确定各数据流的大小，将它们分批交给GPU去执行。本发明在尽量避免系统资源浪费的基础上，更有效地重叠了通信与计算的时间，进而更大程度地缩短整个应用问题的处理时间。

Description

一种面向GPU的数据流处理方法

技术领域

本发明涉及一种数据处理方法，具体涉及一种面向GPU(图形处理器，下文简称GPU)的数据流处理方法。

背景技术

高性能计算发展水平是衡量一个国家综合国力与国际竞争力的重要指标，是一个国家科技实力的综合体现。在航空航天、核武研制、石油勘探、气象预报、生命科学、海啸与地震自然灾害预测等关键领域有着巨大应用需求。因此，世界各国都竞相将高性能计算作为技术与经济争夺的战略制高点。目前，异构多核架构为解决高性能计算提供了契机，它采用“主核心+协处理器”的分工机制，主核心负责操作系统以及一些逻辑性较强的事务处理，而协处理器则负责数据密集型计算任务。但随着主、协处理器架构的使用，两架构之间的通信也需要一定的开销，使得高性能计算问题再次遇到了瓶颈。

由CPU和GPU组成的异构并行系统中，CPU是主核心，GPU是协处理器。在单独使用CPU处理一个涉及高性能计算的应用问题时，只需直接处理准备好的数据，然而，CPU对数据的计算方式是串行的，从而导致了该应用问题需要很长的处理时间。为此，引进协处理器GPU，CPU负责把该应用问题的数据密集型计算任务分离出来移植给GPU进行，利用GPU对数据的并行计算方式缩短计算任务的处理时间，进而缩短了整个应用问题所需要的处理时间。但是，这个过程中需要CPU将初始数据传输到GPU，而GPU处理的结果也需要传输给CPU，这就是前面所提到的架构之间的通信开销。一种有效的解决方法重叠通信与计算，将初始数据划分为不同的数据流分批交给GPU进行处理，以此实现CPU与GPU之间的通信和GPU计算的重叠。

在当前技术中，可以使用CUDA(统一计算设备架构，下文简称CUDA)编程模型中的流处理技术来实现CPU与GPU通信和GPU计算的重叠。所谓的流处理技术是指将初始数据划分为若干个计算时没有相关性的数据块，将它们从属于不同的流进行分批执行，利用CUDA函数特有的异步并行特性，从属于一个流的通信与计算是严格串行的，而从属于不同的流的通信与计算是可以并行执行，如下述代码所示，从属于stream[1]的数据拷贝就可以和从属于stream[0]的kernel计算同时执行，从而实现应用通信与计算的重叠，进而提升应用整体性能。

cudaMemcpyAsync(destination0，source0，direction，stream[0])；

cudaMempcyAsync(destination1，source1，direction，stream[1])；

kernel<<<dimGrid，dimBlock>>>(stream[0])；

然而，该技术有着以下不足：

(1)很难有效地重叠通信与计算的时间

第一，将初始数据分批执行时必然会产生一定的资源浪费，包括CPU和GPU之间的通信带宽以及GPU计算能力两个方面；第二，调度流也需要一定的开销，一味地增加流的数目只会产生负面效果；第三，将初始数据采用不同的划分方式划分所能重叠的通信与计算的时间也是不尽相同的。

由于这些问题的存在，又没有一个明确的数据流处理方法作指导，单纯地基于流处理技术很难有效地重叠通信与计算的时间。

(2)由于没有明确的方向而导致大量的冗余工作

使用不同数目的流可以产生不同的结果，如何能够更快速地得到最好的结果是至关重要的，然而，CUDA编程模型中的流处理技术并没有提供一个明确的方向，导致了大量的冗余工作。

发明内容

针对CUDA编程模型中的流处理技术存在的缺陷，本发明的目的在于，提供一种面向GPU实现通信与计算重叠时方便、有效的数据流处理方法。

为达到以上目的，本发明是采用如下技术方案予以实现的：

一种面向GPU的数据流处理方法，其特征在于，包括下述步骤：

(1)使用CPU和GPU组成的异构并行系统执行如下操作：将待处理的初始数据从主机内存拷贝到GPU，GPU对数据进行计算，计算结果从GPU拷贝回主机内存，运行时间分别为拷入时间Ti、计算时间Tc、拷回时间To，其中：Tc＞Ti+To；

(2)根据步骤(1)中所得到的拷入时间Ti、计算时间Tc、拷回时间To计算如下比例：a＝Tc/Ti，b＝Tc/To；根据该比例将原初始数据向左右两侧划分，采用数据流的规模向两侧逐渐减小的方式，设一共划分为n个数据流，n≥3，第i个数据流的规模最大，n个数据流大小的比例为1∶a∶a²∶…∶a^i-1∶(a^i-1/b)∶(a^i-1/b²)∶…∶(a^i-1/b^n-i)，从而使得：前i-1个数据流中各数据流的计算时间刚好可以隐藏其后一个数据流的拷入时间，而后n-i个数据流中各数据流的计算时间刚好可以隐藏其前一个数据流的拷回时间；

(3)向左侧划分数据直至出现下面的情况之一为止：第1个数据流的传输不能充分利用系统的传输资源，或者第1个数据流的计算不能充分利用系统的计算资源；向右侧划分数据直至出现下面的情况之一为止：第n个数据流的传输不能充分利用系统的传输资源，或者第n个数据流的计算不能充分利用系统的计算资源；

其中：所述的不能充分利用系统的传输资源是指，CPU和GPU间数据的传输依赖于PCIE总线，数据的传输所利用的总线带宽如果没有达到PCIE总线可以实际提供的最高带宽，即为数据流的传输不能充分利用系统的传输资源；所述的不能充分利用系统的计算资源是指，GPU片内拥有着数以百计的计算核心，数据的计算如果没有把这些计算核心全部利用起来，即为数据流的计算不能充分利用系统的计算资源。

本发明针对使用通信与计算重叠的方法处理CPU和GPU之间通信开销的问题，提出了一种面向GPU的数据流处理方法，该方法基于CUDA编程模型中的流处理技术，通过将待处理的数据在CPU和GPU组成的异构并行系统中运行一次所收集的运行时间信息，确定出原初始数据在划分为n个数据流时各数据流之间大小的比例关系，根据该比例关系将数据划分，分批交给GPU执行，以更有效地重叠通信与计算的时间。使用该方法解决了CUDA编程模型中流处理技术的局限性，在考虑了尽量减小系统资源浪费的基础上，一方面可以有效地重叠通信与计算的时间，另一方面根据所总结出的各数据流之间的比例关系划分数据节省了大量的冗余工作，同时可以非常快速地得到一个好的重叠效果。

附图说明

图1为不使用流处理应用问题时的时空图

图2为本发明中使用三个流处理应用问题时的时空图

图3为本发明中使用n个流处理应用问题时的时空图

具体实施方式

以下结合附图及具体实施例对本发明作进一步的详细描述。

一种面向GPU的数据流处理方法，其特征在于，包括下述步骤：

(1)使用CPU和GPU组成的异构并行系统将应用问题(利用计算机处理一个实际问题)中数据密集型计算模块完整运行一次，即将待处理的初始数据从主机内存拷贝到GPU(copyin)，GPU对数据进行计算(kernel)，计算结果从GPU拷贝回主机内存(copyout)，运行时间分别为拷入时间Ti、计算时间Tc、拷回时间To，如图1所示。(该方法针对于计算机时间大于通信时间的问题，即Tc＞Ti+To)；

(2)根据步骤(1)中所得到的拷入时间Ti、计算时间Tc、拷回时间To计算如下比例：a＝Tc/Ti，b＝Tc/To；

如图2所示，可先将初始数据划分为三个数据流，则三个数据流大小的比例为：1∶a(a/b)，从而使得：第一个数据流的计算时间Tc1刚好可以隐藏第二个数据流的拷入时间Ti2，即Tc1＝Ti2；第三个数据流的计算时间Tc3刚好可以隐藏第二个数据流的拷回时间To2，即Tc3＝To2；

为更多地使用计算时间隐藏通信时间，数据的具体划分方法如下，如图3所示，设一共划分为n个数据流，n≥3，第i个数据流的规模最大，n个数据流大小的比例为1∶a∶a²∶…∶a^i-1∶(a^i-1/b)∶(a^i-1/b²)∶…∶(a^i-1/b^n-i)，从而使得：前i-1个数据流中各数据流的计算时间刚好可以隐藏其后一个数据流的拷入时间；而后n-i个数据流中各数据流的计算时间刚好可以隐藏其前一个数据流的拷回时间。

(3)向左侧划分数据直至出现下面的情况之一为止：第1个数据流的传输不能充分利用系统的传输资源，或者第1个数据流的计算不能充分利用系统的计算资源；向右侧划分数据直至出现下面的情况之一为止：第n个数据流的传输不能充分利用系统的传输资源，或者第n个数据流的计算不能充分利用系统的计算资源；

其中：所述的不能充分利用系统的传输资源是指，CPU和GPU间数据的传输依赖于PCIE总线，数据的传输所利用的总线带宽如果没有达到PCIE总线可以实际提供的最高带宽，即为数据流的传输不能充分利用系统的传输资源；所述的不能充分利用系统的计算资源是指，GPU片内拥有着数以百计的计算核心，数据的计算如果没有把这些计算核心全部利用起来，即为数据流的计算不能充分利用系统的计算资源。

本发明选取常用的矩阵相乘C＝A×B为例，选择矩阵A、B和C都是维度为4096的方阵进行测试，方阵的元素均为单精度浮点数。给出使用发明中所提出的面向GPU的数据流处理方法在CPU和GPU组成的异构并行系统中计算矩阵相乘的一个具体实施过程：

(1)将矩阵相乘运行一次，得到其各部分操作的运行时间，如下表所示，其中t_iA和t_iB分别代表矩阵A和B从CPU拷贝到GPU所需要的时间，t_c代表GPU计算矩阵A和B相乘所需要的时间，t_oC代表计算结果C从设备显存拷贝回主机内存的时间。

表1矩阵相乘运行时间

由于矩阵相乘的特性，矩阵B需要一次性全部拷贝到GPU中，然后将矩阵A分批交给GPU执行，因此，这里是用计算时间隐藏矩阵A和矩阵C的通信时间。

通信时间：Ti+To＝t_iA+t_oC＝22.49ms

计算时间：Tc＝t_c＝477.64ms

Tc＞Ti+To，矩阵相乘属于计算时间大于通信时间的应用。

(2)计算比例：a＝Tc/Ti≈42，b＝Tc/To≈42。

总的数据量大小为256×16×4096，而由于矩阵相乘应用的特性，16×4096是不可分割的，因此，我们针对256划分即可，划分为三个数据流，比例为：1∶42∶1，通过计算可得，第一个数据流的大小为6×16×4096，第二数据流的大小为244×16×4096，第三个流的大小为6×16×4096。

(3)将初始数据划分为三个数据流时(n＝3)，最小的数据流为6×16×4096个元素，每个元素为4个字节(4B)，则这些元素的量为384KB。通过对本实验所使用的环境中CPU和GPU总线通信的带宽进行测试，一次性拷贝的数据量在接近1MB时才能充分地利用带宽，因此，无需再将数据进一步划分。否则，就很有可能会因为系统资源的浪费产生负面效果。

使用本发明的方法得到的运行时间如下：

在不使用流时整个矩阵相乘的运行时间为：511.26ms，

使用此方法时整个矩阵相乘的运行时间为：491.85ms。

运行时间减少了19.41ms。相当于隐藏了矩阵A和C通信总时间的86.3％，由于部分额外开销的影响，实际的隐藏效果在86.3％以上。

根据本发明中提出的数据流处理方法，我们可以基于CUDA编程模型中的流技术重叠了矩阵A和C 86.3％以上的通信时间，同时这个过程相当快捷，很好地解决了CUDA编程模型中流处理技术的局限性。

Claims (1)

1.一种面向GPU的数据流处理方法，其特征在于，包括下述步骤：

(1)使用CPU和GPU组成的异构并行系统执行如下操作：将待处理的初始数据从主机内存拷贝到GPU，GPU对数据进行计算，计算结果从GPU拷贝回主机内存，运行时间分别为拷入时间Ti、计算时间Tc、拷回时间To，其中：Tc＞Ti+To；

(2)根据步骤(1)中所得到的拷入时间Ti、计算时间Tc、拷回时间To计算如下比例：a＝Tc/Ti，b＝Tc/To；根据该比例将原初始数据向左右两侧划分，采用数据流的规模向两侧逐渐减小的方式，设一共划分为n个数据流，n≥3，第i个数据流的规模最大，则n个数据流大小的比例为1∶a∶a²∶…∶a^i-1∶(a^i-1/b)∶(a^i-1/b²)∶…∶(a^i-1/b^n-i)，从而使得：前i-1个数据流中各数据流的计算时间刚好可以隐藏其后一个数据流的拷入时间，而后n-i个数据流中各数据流的计算时间刚好可以隐藏其前一个数据流的拷回时间；

(3)向左侧划分数据直至出现下面的情况之一为止：第1个数据流的传输不能充分利用系统的传输资源，或者第1个数据流的计算不能充分利用系统的计算资源；向右侧划分数据直至出现下面的情况之一为止：第n个数据流的传输不能充分利用系统的传输资源，或者第n个数据流的计算不能充分利用系统的计算资源；

其中：所述的不能充分利用系统的传输资源是指，CPU和GPU间数据的传输依赖于PCIE总线，数据的传输所利用的总线带宽如果没有达到PCIE总线可以实际提供的最高带宽，即为数据流的传输不能充分利用系统的传输资源；所述的不能充分利用系统的计算资源是指，GPU片内拥有着数以百计的计算核心，数据的计算如果没有把这些计算核心全部利用起来，即为数据流的计算不能充分利用系统的计算资源。

Images