CN113407333B - Warp级别调度的任务调度方法、系统、GPU及设备 - Google Patents

Warp级别调度的任务调度方法、系统、GPU及设备 Download PDF

Info

Publication number
CN113407333B
CN113407333B CN202011506442.5A CN202011506442A CN113407333B CN 113407333 B CN113407333 B CN 113407333B CN 202011506442 A CN202011506442 A CN 202011506442A CN 113407333 B CN113407333 B CN 113407333B
Authority
CN
China
Prior art keywords
task
tasks
scheduling
decision
gpu
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN202011506442.5A
Other languages
English (en)
Other versions
CN113407333A (zh
Inventor
陈�全
过敏意
赵涵
崔炜皞
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Shanghai Jiaotong University
Original Assignee
Shanghai Jiaotong University
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Shanghai Jiaotong University filed Critical Shanghai Jiaotong University
Priority to CN202011506442.5A priority Critical patent/CN113407333B/zh
Publication of CN113407333A publication Critical patent/CN113407333A/zh
Application granted granted Critical
Publication of CN113407333B publication Critical patent/CN113407333B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5027—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals
    • G06F9/5038—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals considering the execution order of a plurality of tasks, e.g. taking priority or time dependency constraints into consideration
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/48—Program initiating; Program switching, e.g. by interrupt
    • G06F9/4806—Task transfer initiation or dispatching
    • G06F9/4843—Task transfer initiation or dispatching by program, e.g. task dispatcher, supervisor, operating system
    • G06F9/4881—Scheduling strategies for dispatcher, e.g. round robin, multi-level priority queues
    • Y—GENERAL TAGGING OF NEW TECHNOLOGICAL DEVELOPMENTS; GENERAL TAGGING OF CROSS-SECTIONAL TECHNOLOGIES SPANNING OVER SEVERAL SECTIONS OF THE IPC; TECHNICAL SUBJECTS COVERED BY FORMER USPC CROSS-REFERENCE ART COLLECTIONS [XRACs] AND DIGESTS
    • Y02—TECHNOLOGIES OR APPLICATIONS FOR MITIGATION OR ADAPTATION AGAINST CLIMATE CHANGE
    • Y02D—CLIMATE CHANGE MITIGATION TECHNOLOGIES IN INFORMATION AND COMMUNICATION TECHNOLOGIES [ICT], I.E. INFORMATION AND COMMUNICATION TECHNOLOGIES AIMING AT THE REDUCTION OF THEIR OWN ENERGY USE
    • Y02D10/00—Energy efficient computing, e.g. low power processors, power management or thermal management

Landscapes

  • Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Theoretical Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • General Physics & Mathematics (AREA)
  • Management, Administration, Business Operations System, And Electronic Commerce (AREA)
  • Multi Processors (AREA)

Abstract

本发明提供一种Warp级别调度的任务调度方法、系统、GPU及设备，所述Warp级别调度的任务调度方法包括：在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度；在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。本发明在无需用户感知提前下，实现高吞吐的Warp级别调度的任务调度，可以间接为潜在的配置多种计算单元的GPU提供调度技术的支持。

Description

Warp级别调度的任务调度方法、系统、GPU及设备
技术领域
本发明涉及处理器技术领域，特别是涉及GPU设计控制技术领域。
背景技术
通用GPU的快速发展导致了其在云计算与数据中心的大量采用。一方面，传统科学比如数学，神经科学等学科采用GPU进行相应的科学计算。另一方面，计算机领域的多个方向都开始采用GPU加速相关计算任务。举例来说，机器学习和深度学习依赖GPU进行训练和服务，而数据流任务则依赖GPU加速数据处理。因此，为了支持越来越多样的GPU任务，大量的企业和学术机构开始构建基于GPU的私有云。
当应用逐渐多样化的同时，GPU架构也在不断发展。最初GPU的流多处理器SM只包含一类计算单元CUDA core，CUDA core进一步细分为INT32，FP32和FP64。随着深度学习的快速发展和大量应用，NVIDIA提出了专用的计算单元Tensor core。因为深度学习中最常用的计算是矩阵乘，所以Tensor core的唯一功能就是矩阵乘计算。容易看出，传统任务和数据流处理任务主要使用CUDA core，深度学习任务主要使用Tensor core。与此同时，CUDAcore和Tensor core是完全独立的两种计算硬件，这也意味着他们两个能够同时计算。所以，这里存在着以同时利用两种计算硬件的方式，提升GPU利用率的可能性，进而可以提高GPU整体的吞吐率。
现有GPU的调度依赖于GPU任务的提交配置。一个GPU任务内有一个或多个kernel函数。当GPU任务被提交到GPU上时，任务内的kernel函数顺序得到执行。当一个kernel函数被提交时，该kernel函数会以指定的配置启动指定数量的线程块Thread block。每个thread block内有指定的线程数。当一个thread block被提交到SM上时，每32个线程作为一个warp，同时在GPU上执行，当一个warp因为访存停止时，另一个warp被调度上GPU进行计算。本质上来说，GPU上的任务调度粒度是warp级别的。
而对于现有调度系统而言，这些系统都依赖于NVIDIA官方的提供的MPS进行调度。MPS支持GPU任务申请特定比例的SM资源，而任务只能使用这么多SM资源。当一个任务没有使用完当前SM上的计算资源时，另一个任务才能调度到SM上。这意味着，MPS的调度粒度是SM级别的，更准确的来说是任务级别的或者kernel级别的。所以kernel级别的调度粒度无法直接利用到两种计算单元。这意味着，并行使用两种硬件的第一个难题是粗粒度的任务调度。
除了调度粒度的问题之外，并行使用两种硬件的第二个难题是内存系统的竞争。由于两个任务共享一个GPU内存系统，内存竞争非常容易出现。现有的工具和之前的工作都没有管理内存系统的竞争。尽管两个计算硬件是理论上可并行的，当一个任务有对内存系统有特别高的要求时，两种计算单元仍然不能被并行使用。当任务大量到来时，如何决定哪些任务可以混跑是第三个难题。
NVIDIA MPS是NVIDIA官方提供的一种多任务并行的接口。MPS支持任务申请指定比例的SM资源，而任务只能使用这些SM资源。当一个任务没有使用完当前SM上的计算资源时，另一个任务才能调度到SM上。
当一个任务没有使用完当前SM上的计算资源时，另一个任务才能调度到SM上。这意味着，MPS的调度粒度是SM级别的，更准确的来说是任务级别的或者kernel级别的,所以Kernel级别的调度粒度无法直接利用到两种计算单元。这意味着，NVIDIA MPS的第一个缺点是粗粒度的任务调度。
除了调度粒度的问题之外，并行使用两种硬件的第二个缺点是无法感知内存系统的竞争。由于两个任务共享一个内存系统，内存竞争非常容易出现，进而影响到两个任务的并行性。最后，当大量任务到来的时候，MPS无法感知如何决定哪些任务可以混跑是其第三个缺点。
发明内容
鉴于以上所述现有技术的缺点，本发明的目的在于提供一种Warp级别调度的任务调度方法、系统、GPU及设备，用于实现高吞吐的Warp级别调度的任务调度。
为实现上述目的及其他相关目的，本发明提供一种Warp级别调度的任务调度方法，包括：在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度；在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。
于本发明的一实施例中，所述在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息包括：收集用户提交任务对各个数据通路的使用情况，以及用户提交任务在不同线程块配置下的任务性能。
于本发明的一实施例中，所述基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度的实现方式为：基于所述任务的配置信息，获取任务性能不变情况下的线程块配置区间；基于所述任务的配置信息及硬件信息，获取任务的计算访存类型；基于所述任务计算访存类型及任务线程块配置区间，遍历两个任务在所有的线程块配置下的并行性能，以此来获得最大的并行度。
于本发明的一实施例中，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策包括：若到来任务中有任务对经过并行放大调优处理并获得相应的最大并行度决策，则根据调优好的最大并行度决策挑选任务对；对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断。
于本发明的一实施例中，所述对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断包括：基于所述任务获取任务的硬件信息，将两个任务的硬件信息中的数据通路使用率两两相加；判断各条数据通路相加之和是否存在一条数据通路的使用率超过设定的阈值：若是，则确认两个任务不能打包成一个新任务；若否，则确认两个任务能打包成一个新任务。
于本发明的一实施例中，所述将选取的任务对打包成一个新任务包括：将两个任务的网格维数和线程块维数调整到一维；计算新任务的网格维数和线程块维数；调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；构建新的核函数的内容；调整新的设备函数的传递参数和同步函数。
于本发明的一实施例中，所述两个任务打包成一个新任务的一种实现方式包括：将两个任务的网格维数和线程块维数调整到一维；计算新任务的网格维数和线程块维数；调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；构建新的核函数的内容；调整新的设备函数的传递参数和同步函数。
本发明的实施例还提供一种Warp级别调度的任务调度系统，所述Warp级别调度的任务调度系统包括：离线分析模块，用于在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；并行度放大模块，基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度；决策模块，用于在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策；任务打包模块，用于将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。
于本发明的一实施例中，所述任务打包模块包括：维数控制单元，用于将两个任务的网格维数和线程块维数调整到一维，并计算新任务的网格维数和线程块维数；参数构建单元，用于调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；内容构建单元，用于构建新的核函数的内容；调整单元，用于调整新的设备函数的传递参数和同步函数。
本发明的实施例还提供一种GPU，用于运行所述计算机程序以实现如上所述的Warp级别调度的任务调度方法。
本发明的实施例还提供一种电子设备，应用如上所述的GPU。
如上所述，本发明的Warp级别调度的任务调度方法、系统、GPU及设备具有以下有益效果：
本发明在无需用户感知提前下，实现高吞吐的Warp级别调度的任务调度，可以间接为潜在的配置多种计算单元的GPU提供调度技术的支持，可以使配置多种计算单元的GPU芯片的动态任务调度系统可以面向用户提供任务调度服务。
附图说明
为了更清楚地说明本发明实施例中的技术方案，下面将对实施例描述中所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本发明的一些实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据这些附图获得其他的附图。
图1显示为本申请一实施例中的Warp级别调度的任务调度方法的整体流程示意图。
图2显示为本申请一实施例中的Warp级别调度的任务调度方法中的原理架构示意图。
图3显示为本申请一实施例中的Warp级别调度的任务调度方法中的具体执行过程示意图。
图4显示为本申请一实施例中的Warp级别调度的任务调度系统的原理框图。
图5显示为本申请一实施例中的GPU的原理框图。
元件标号说明
100 Warp级别调度的任务调度系统
110 离线分析模块
120 并行度放大模块
130 决策模块
140 任务打包模块
S100～S300 步骤
S1～S8 步骤
具体实施方式
以下通过特定的具体实例说明本发明的实施方式，本领域技术人员可由本说明书所揭露的内容轻易地了解本发明的其他优点与功效。本发明还可以通过另外不同的具体实施方式加以实施或应用，本说明书中的各项细节也可以基于不同观点与应用，在没有背离本发明的精神下进行各种修饰或改变。需说明的是，在不冲突的情况下，以下实施例及实施例中的特征可以相互组合。
本实施例的目的在于提供一种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
Figure SMS_1
/>
Figure SMS_2
步骤S300，在任务为非首次任务时，基于任务的硬件信息和收集好的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度。
在大量任务到来时，实时决定那两个任务能够一起获得两种计算硬件的并行优势。由于两个混跑的任务要共享GPU的内存架构，当一个任务没有使用完所有的GPU内存资源时，另一个任务才能被调度到GPU上。
当任务被并行放大分析后，直接采用并行放大后的任务配置进行任务打包。在任务到来且没有被并行放大分析时，根据获取的内存使用信息，进行打包任务对的在线决策。相应的在线打包决策算法：判断所有数据通路的使用率是否超过设定的阈值，这个阈值最初设置为但不限于100，根据用户和历史信息而调整。当任意一条数据通路的利用率之后大于阈值，两个任务就被判定不能成为一个打包对。
具体地，于本实施例中，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策包括：
若到来任务中有任务对经过并行放大调优处理并获得相应的最大并行度决策，则根据调优好的最大并行度决策挑选任务对；对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断。
其中，于本实施例中，所述对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断包括：
基于所述任务获取任务的硬件信息，将两个任务的硬件信息中的数据通路使用率两两相加；
判断各条数据通路相加之和是否存在一条数据通路的使用率超过设定的阈值：
若是，则确认两个任务不能打包成一个新任务；
若否，则确认两个任务能打包成一个新任务。
具体地，于本实施例中，实现对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断算法示例如下表2所示。
表2
Figure SMS_3
/>
Figure SMS_4
于本实施例中，所述将选取的任务对打包成一个新任务包括：将两个任务的网格维数和线程块维数调整到一维；计算新任务的网格维数和线程块维数；调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；构建新的核函数的内容；调整新的设备函数的传递参数和同步函数。
由于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芯片的动态任务调度系统可以面向用户提供任务调度服务。所以，本发明有效克服了现有技术中的种种缺点而具高度产业利用价值。
上述实施例仅例示性说明本发明的原理及其功效，而非用于限制本发明。任何熟悉此技术的人士皆可在不违背本发明的精神及范畴下，对上述实施例进行修饰或改变。因此，举凡所属技术领域中包括通常知识者在未脱离本发明所揭示的精神与技术思想下所完成的一切等效修饰或改变，仍应由本发明的权利要求所涵盖。

Claims (8)

1.一种Warp级别调度的任务调度方法，其特征在于：包括：
在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；
基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度，最大化两个打包任务的并行度；
通过任务对最大化并行度算法最大化两个打包任务的并行度，该算法思想是：第一步，先刻画任务随着提交线程块数变化的性能曲线，该线程块数为GPU任务在单个流处理器上的驻留的线程块数；第二步，判断任务的任务类型，决定任务的起始线程块；第三步，遍历所有任务对的所有线程块配置对，找到并行度最大的并行配置；第四步，返回所有的配置对；
在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策，将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度；
基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策包括：若到来任务中有任务对经过并行放大调优处理并获得相应的最大并行度决策，则根据调优好的最大并行度决策挑选任务对；对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断；
所述对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断包括：基于所述任务获取任务的硬件信息，将两个任务的硬件信息中的数据通路使用率两两相加；判断各条数据通路相加之和是否存在一条数据通路的使用率超过设定的阈值：若是，则确认两个任务不能打包成一个新任务；若否，则确认两个任务能打包成一个新任务。
2.根据权利要求1所述的Warp级别调度的任务调度方法，其特征在于：所述在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息包括：收集用户提交任务对各个数据通路的使用情况，以及用户提交任务在不同线程块配置下的任务性能。
3.根据权利要求1或2所述的Warp级别调度的任务调度方法，其特征在于：所述基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度的实现方式为：
基于所述任务的配置信息，获取任务性能不变情况下的线程块配置区间；
基于所述任务的配置信息及硬件信息，获取任务的计算访存类型；
基于所述任务计算访存类型及任务线程块配置区间，遍历两个任务在所有的线程块配置下的并行性能，以此来获得最大的并行度。
4.根据权利要求1所述的Warp级别调度的任务调度方法，其特征在于：所述将选取的任务对打包成一个新任务包括：
将两个任务的网格维数和线程块维数调整到一维；
计算新任务的网格维数和线程块维数；
调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；
构建新的核函数的内容；
调整新的设备函数的传递参数和同步函数。
5.一种Warp级别调度的任务调度系统，其特征在于：所述Warp级别调度的任务调度系统包括：
离线分析模块，用于在任务为首次任务时，离线分析用户提交任务的硬件信息及配置信息；
并行度放大模块，基于所述任务的硬件信息和配置信息，获取与主流任务并行时的最大并行度，最大化两个打包任务的并行度；
通过任务对最大化并行度算法最大化两个打包任务的并行度，该算法思想是：第一步，先刻画任务随着提交线程块数变化的性能曲线，该线程块数为GPU任务在单个流处理器上的驻留的线程块数；第二步，判断任务的任务类型，决定任务的起始线程块；第三步，遍历所有任务对的所有线程块配置对，找到并行度最大的并行配置；第四步，返回所有的配置对；
决策模块，用于在任务为非首次任务时，基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策；
任务打包模块，用于将选取的任务对打包成一个新任务，并将新任务提交到GPU上，以使得任务对中的原有两个GPU任务实现warp级别的任务调度；
基于任务的硬件信息和已收集的最大并行度决策，进行在线的任务对打包决策包括：若到来任务中有任务对经过并行放大调优处理并获得相应的最大并行度决策，则根据调优好的最大并行度决策挑选任务对；对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断；
所述对于剩余未经过并行放大处理的任务对，基于任务的硬件信息进行打包决策的判断包括：基于所述任务获取任务的硬件信息，将两个任务的硬件信息中的数据通路使用率两两相加；判断各条数据通路相加之和是否存在一条数据通路的使用率超过设定的阈值：若是，则确认两个任务不能打包成一个新任务；若否，则确认两个任务能打包成一个新任务。
6.根据权利要求5所述的Warp级别调度的任务调度系统，其特征在于：所述任务打包模块包括：
维数控制单元，用于将两个任务的网格维数和线程块维数调整到一维，并计算新任务的网格维数和线程块维数；
参数构建单元，用于调整两个任务从全局函数到设备函数，并基于两个任务原网格维数和原线程块维数构建新的核函数的传递参数；
内容构建单元，用于构建新的核函数的内容；
调整单元，用于调整新的设备函数的传递参数和同步函数。
7.一种GPU，其特征在于：用于运行计算机程序以实现如权利要求1至权利要求4任一权利要求所述的Warp级别调度的任务调度方法。
8.一种电子设备，其特征在于：应用如权利要求7所述的GPU。
CN202011506442.5A 2020-12-18 2020-12-18 Warp级别调度的任务调度方法、系统、GPU及设备 Active CN113407333B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202011506442.5A CN113407333B (zh) 2020-12-18 2020-12-18 Warp级别调度的任务调度方法、系统、GPU及设备

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202011506442.5A CN113407333B (zh) 2020-12-18 2020-12-18 Warp级别调度的任务调度方法、系统、GPU及设备

Publications (2)

Publication Number Publication Date
CN113407333A CN113407333A (zh) 2021-09-17
CN113407333B true CN113407333B (zh) 2023-05-26

Family

ID=77675644

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202011506442.5A Active CN113407333B (zh) 2020-12-18 2020-12-18 Warp级别调度的任务调度方法、系统、GPU及设备

Country Status (1)

Country Link
CN (1) CN113407333B (zh)

Families Citing this family (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114003359A (zh) * 2021-10-20 2022-02-01 上海交通大学 基于弹性持久的线程块的任务调度方法、系统及gpu

Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN107122245A (zh) * 2017-04-25 2017-09-01 上海交通大学 Gpu任务调度方法及系统
CN107357661A (zh) * 2017-07-12 2017-11-17 北京航空航天大学 一种针对混合负载的细粒度gpu资源管理方法
CN111736987A (zh) * 2020-05-29 2020-10-02 山东大学 一种基于gpu空间资源共享的任务调度方法

Family Cites Families (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102708009B (zh) * 2012-04-19 2014-04-02 华为技术有限公司 一种基于cuda实现多任务共享gpu的方法
EP2887219A1 (en) * 2013-12-23 2015-06-24 Deutsche Telekom AG System and method for mobile augmented reality task scheduling
CN106502771B (zh) * 2016-09-09 2019-08-02 中国农业大学 基于kernel函数的时间开销模型构建方法及系统
CN110088730B (zh) * 2017-06-30 2021-05-18 华为技术有限公司 任务处理方法、装置、介质及其设备
CN109522108B (zh) * 2018-10-30 2020-10-27 西安交通大学 一种基于Kernel合并的GPU任务调度系统及方法

Patent Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN107122245A (zh) * 2017-04-25 2017-09-01 上海交通大学 Gpu任务调度方法及系统
CN107357661A (zh) * 2017-07-12 2017-11-17 北京航空航天大学 一种针对混合负载的细粒度gpu资源管理方法
CN111736987A (zh) * 2020-05-29 2020-10-02 山东大学 一种基于gpu空间资源共享的任务调度方法

Also Published As

Publication number Publication date
CN113407333A (zh) 2021-09-17

Similar Documents

Publication Publication Date Title
CN104834561B (zh) 一种数据处理方法及装置
CN106598735B (zh) 一种分布式计算方法、主控节点和计算系统
US20070226735A1 (en) Virtual vector processing
KR101400577B1 (ko) Gpu를 이용한 희소행렬 곱셈 방법
US9665532B2 (en) Performing synchronized collective operations over multiple process groups
CN113407333B (zh) Warp级别调度的任务调度方法、系统、GPU及设备
Czarnul Investigation of parallel data processing using hybrid high performance CPU+ GPU systems and CUDA streams
WO2016041126A1 (zh) 基于gpu的数据流处理方法和装置
CN109410136A (zh) 基于最短传递路径的匀色方法及处理装置
KR101292670B1 (ko) 벡터 프로세싱 장치 및 방법
Prades et al. Multi-tenant virtual GPUs for optimising performance of a financial risk application
Andelfinger et al. Exploiting the parallelism of large-scale application-layer networks by adaptive GPU-based simulation
Haussmann et al. Cost-optimized parallel computations using volatile cloud resources
CN114003359A (zh) 基于弹性持久的线程块的任务调度方法、系统及gpu
Amestoy et al. Modeling 1D distributed-memory dense kernels for an asynchronous multifrontal sparse solver
CN113076190A (zh) 一种基于cpu与gpu协作的计算方法
Falcao et al. Heterogeneous implementation of a voronoi cell-based svp solver
Wang et al. GPARS: Graph predictive algorithm for efficient resource scheduling in heterogeneous GPU clusters
US20200202251A1 (en) Intelligent computation acceleration transform utility
CN111694672A (zh) 资源分配方法、任务提交方法、装置、电子设备和介质
Garg et al. FaaSter: Accelerated Functions-as-a-Service with Heterogeneous GPUs
Cao et al. SPMD performance analysis with parallel computing of MATLAB
CN117407643B (zh) 一种通用矩阵乘法的优化方法、系统、设备及介质
Yeh et al. CloudBook: Implementing an E-book reader on a cloud platform by an optimized vector graphic library
Chung et al. Proactive task offloading for load balancing in iterative applications

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant