具体实施方式

以下通过特定的具体实例说明本发明的实施方式，本领域技术人员可由本说明书所揭露的内容轻易地了解本发明的其他优点与功效。本发明还可以通过另外不同的具体实施方式加以实施或应用，本说明书中的各项细节也可以基于不同观点与应用，在没有背离本发明的精神下进行各种修饰或改变。需说明的是，在不冲突的情况下，以下实施例及实施例中的特征可以相互组合。

本实施例的目的在于提供一种Warp级别调度的任务调度方法、系统、GPU及设备，用于实现高吞吐的Warp级别调度的任务调度。

本实施例旨在设计、实现一个运行于数据中心上，管理大量程序间的共享资源的系统和方法，在同时保证两个程序的正确性的情况下，最大化整体任务的吞吐量。

以下将详细阐述本发明的Warp级别调度的任务调度方法、系统、GPU及设备的原理及实施方式，使本领域技术人员不需要创造性劳动即可理解本发明的Warp级别调度的任务调度方法、系统、GPU及设备。

实施例1

具体地，如图1所示，本实施例提供一种Warp级别调度的任务调度方法，所述Warp级别调度的任务调度方法包括：

步骤S100，在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；

步骤S200，基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度；

步骤S300，在任务为非首次任务时，基于任务的硬件信息和收集好的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。

以下将结合图2对本实施例的Warp级别调度的任务调度方法的步骤S100至步骤S300进行详细说明。

步骤S100，在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息。

于本实施例中，所述在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息包括：收集用户提交任务对各个数据通路的使用情况，以及用户提交任务在不同线程块配置下的任务性能。

具体地，于本实施例中，首先，分析任务对GPU内存系统的需求，收集用户提交的任务对各个内存通路的使用情况。其次，分析任务对提交线程块(block)数的性能变化，收集任务在保证性能时的提交线程块(block)区间。

步骤S200，基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度。

于本实施例中，最大化两个打包任务的并行度，即最大化两种计算硬件的并行度。最大化两个任务的并行度需要管理两个任务的内存系统竞争。一个任务对内存系统的需求，主要取决于其在单个SM(计算单元)上同时计算的线程块(block)数。当一个SM(计算单元)上同时计算的线程块(block)数越多，其并发的内存的请求就越多，也就是对内存带宽需求越高。因此，本实施例通过源对源编译的支持，动态调整两个任务提交的线程块(block)数，进而最大化两个任务在GPU上的并行度。

通过相应的调研发现，GPU任务可以分为三类，内存密集型任务，均衡型任务和计算密集型任务。计算密集型任务对带宽的需求较小。当block数较多时，任务的性能会提升，但当block数提升到一定数量时，性能提升则趋于放缓。内存密集型任务对内存带宽的需求较大且计算需求较小。当block数增大时，block间的访存竞争导致了任务的性能下降。均衡型任务主要的特点是，block数的变化对其性能不造成任何的影响。

为了最大化两个任务的并行度，本实施例设计了任务对最大化并行度算法。本算法具体的思想是：第一步，先刻画任务随着提交线程块(block)数变化的性能曲线，该线程块数为GPU任务在单个流处理器上的驻留的block数；第二步，判断任务的任务类型，决定任务的起始block；第三步，遍历所有任务对的所有block配置对，找到并行度最大的并行配置。第四步，返回所有的配置对。

具体地，于本实施例中，所述基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度的实现方式为：

1)基于所述任务的配置信息，获取任务性能不变情况下的线程块配置区间；

2)基于所述任务的配置信息及硬件信息，获取任务的计算访存类型；

3)基于所述任务计算访存类型及任务线程块配置区间，遍历两个任务在所有的线程块配置下的并行性能，以此来获得最大的并行度。

于本实施例中，任务最大化并行度的算法示例如下表1所示。

表1

步骤S300，在任务为非首次任务时，基于任务的硬件信息和收集好的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。

在大量任务到来时，实时决定那两个任务能够一起获得两种计算硬件的并行优势。由于两个混跑的任务要共享GPU的内存架构，当一个任务没有使用完所有的GPU内存资源时，另一个任务才能被调度到GPU上。

当任务被并行放大分析后，直接采用并行放大后的任务配置进行任务打包。在任务到来且没有被并行放大分析时，根据获取的内存使用信息，进行打包任务对的在线决策。相应的在线打包决策算法：判断所有数据通路的使用率是否超过设定的阈值，这个阈值最初设置为但不限于100，根据用户和历史信息而调整。当任意一条数据通路的利用率之后大于阈值，两个任务就被判定不能成为一个打包对。

具体地，于本实施例中，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策包括：

若到来任务中有任务对经过并行放大调优处理并获得相应的最大并行度决策，则根据调优好的最大并行度决策挑选任务对；对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断。

其中，于本实施例中，所述对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断包括：

基于所述任务获取任务的硬件信息，将两个任务的硬件信息中的数据通路使用率两两相加；

判断各条数据通路相加之和是否存在一条数据通路的使用率超过设定的阈值：

若是，则确认两个任务不能打包成一个新任务；

若否，则确认两个任务能打包成一个新任务。

具体地，于本实施例中，实现对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断算法示例如下表2所示。

表2

于本实施例中，所述将选取的任务对打包成一个新任务包括：将两个任务的网格维数和线程块维数调整到一维；计算新任务的网格维数和线程块维数；调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；构建新的核函数的内容；调整新的设备函数的传递参数和同步函数。

由于NVIDIA现有只开放了MPS进行相关的调度，而MPS的调度粒度是kernel级别的，这导致了MPS无法同时利用到两种计算硬件。与此同时，当任务被调度到GPU后，SM上的调度粒度是warp级别的。为了能够享受到warp级别的调度，本实施例实现了两个任务的打包，以保证两个任务同时提交到GPU上，进而利用到两种计算单元的计算并行性。

具体地，将所述两个任务打包成一个新任务的一种具体实施过程如下。

第一，调整两个GPU任务的grid dimension和block dimension到只有一维，并将原有的grid dimension和block dimension作为新kernel function的传参。

第二，计算新任务的grid dimension，新的grid dimension是两个任务griddimension的最大值。

第三，计算新任务的block dimension，新的block dimension是两个任务blockdimension的和。

第四，调整两个任务从global function到device function，并根据两个任务原有的传参和原有任务的维度构建新的kernel function的传参。

第五，构建新的kernel function的内容。在相应的block内，前M个线程负责kernel A,后N个线程负责kernel B。

第六，调整新的device function的传参。新的device function除了原有的参数外，还需要添加原有kernel function的维度信息。通过传入的kernel function的维度信息，计算出原有的thread id和block id，以保证计算的正确性。

第七，为了保证线程内同步的正确性，调整__syncthreads()函数为bar.sync()。

本实施例中，warp级别的任务打包使用源对源编译，并解决了部分线程同步，以支持warp级别的任务调度，进而保证两个任务能够同时利用到GPU的两种计算硬件。

如图2和图3所示，为使本领域技术人员进一步理解本实施例的warp级别调度的任务调度方法，以下对本实施例的warp级别调度的任务调度的实施过程进行说明。

S1，用户提交任务：用户通过调用相应的API，按照需求编写自己的程序。

S2，分析是否为第一次任务，即任务是否第一次到来，若是，则进入S3：离线分析，若否则进入S5。

S3，离线分析：收集任务的相关信息。

S4，并行度放大：基于离线分析的相关信息，得知每个任务性能变化的block区间。然后，通过调整任务每次提交的block数，尝试发现两个任务在GPU上最大的并行度。

S5，在线决策打包任务对：基于离线分析和并行放大获得的任务信息进行在线决策。若任务未被并行放大分析，则通过对内存系统的使用进行在线打包决策。

S6，任务打包：属于warp级别的任务打包，基于在线决策的打包决策进行打包操作。将两个任务通过源对源编译的方法组成一个新任务，提交到GPU上，进而两个任务可以享受到warp级别的调度。

S7，任务运行：任务开始在相应节点上进行执行，执行完毕后执行S8：输出结果。

实施例2

如图4所示，本实施例提供一种Warp级别调度的任务调度系统100，所述Warp级别调度的任务调度系统100包括：离线分析模块110，并行度放大模块120，决策模块130以及任务打包模块140。

于本实施例中，所述离线分析模块110用于在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息。

于本实施例中，所述在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息包括：收集用户提交任务对各个数据通路的使用情况，以及用户提交任务在不同线程块配置下的任务性能。

具体地，于本实施例中，首先，分析任务对GPU内存系统的需求，收集用户提交的任务对各个内存通路的使用情况。其次，分析任务对提交线程块(block)数的性能变化，收集任务在保证性能时的提交线程块(block)区间。

于本实施例中，所述并行度放大模块120用于基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度。

所述并行度放大模块120的原理与实施例1中的步骤S200相同，原理间相似或相同的技术特征不再赘述。

于本实施例中，所述决策模块130用于在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策。

在大量任务到来时，所述决策模块130实时决定那两个任务能够一起获得两种计算硬件的并行优势。由于两个混跑的任务要共享GPU的内存架构，当一个任务没有使用完所有的GPU内存资源时，另一个任务才能被调度到GPU上。

当任务被并行并行度放大模块120分析后，直接采用并行度放大模块120的任务配置进行任务打包。在任务到来且没有被并行并行度放大模块120分析时，所述决策模块130根据离线分析模块110获取的内存使用信息，进行打包任务对的在线决策。

于本实施例中，所述任务打包模块140用于将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。

具体地，于本实施例中，所述任务打包模块140包括：维数控制单元，参数构建单元，内容构建单元以及调整单元。

于本实施例中，所述维数控制单元用于将两个任务的网格维数和线程块维数调整到一维，并计算新任务的网格维数和线程块维数；于本实施例中，所述参数构建单元用于调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；于本实施例中，所述内容构建单元用于构建新的核函数的内容；于本实施例中，所述调整单元用于调整新的设备函数的传递参数和同步函数。

于本实施例中，所述决策模块130和所述任务打包模块140的原理与实施例1中的步骤S300相同，原理间相似或相同的技术特征不再赘述。

实施例3

本发明的实施例还提供一种GPU，用于运行所述计算机程序以实现实施例1所述的Warp级别调度的任务调度方法。

如图5所示，本实施例中，GPU是一个基于配置张量计算核心(Tensor core)的新型GPU。图5所示是一个配置张量计算核心(Tensor core)的GPU芯片示意图。图5中，SM(计算单元)代表GPU的粗粒度计算单元。每个SM内部有两种细粒度计算单元，分别是流处理核心核心(CUDA core)和张量计算核心(Tensor core)。每个SM内共享寄存器和第一级缓存(L1)，所有的SM共享第二级缓存(L2 cache)和DRAM内存。

实施例4

本发明的实施例还提供一种电子设备，所述电子设备例如为固定终端，例如服务器、台式机等。

所述电子设备应用实施例3中所述的GPU，以使所述电子设备执行如实施例1中Warp级别调度的任务调度方法的步骤。由于Warp级别调度的任务调度方法的步骤的具体实施过程已经在实施例1中进行了详细说明，在此不再赘述。

综上所述，本发明在无需用户感知提前下，实现高吞吐的Warp级别调度的任务调度，可以间接为潜在的配置多种计算单元的GPU提供调度技术的支持，可以使配置多种计算单元的GPU芯片的动态任务调度系统可以面向用户提供任务调度服务。所以，本发明有效克服了现有技术中的种种缺点而具高度产业利用价值。

上述实施例仅例示性说明本发明的原理及其功效，而非用于限制本发明。任何熟悉此技术的人士皆可在不违背本发明的精神及范畴下，对上述实施例进行修饰或改变。因此，举凡所属技术领域中包括通常知识者在未脱离本发明所揭示的精神与技术思想下所完成的一切等效修饰或改变，仍应由本发明的权利要求所涵盖。