CN111124691B - 多进程共享的gpu调度方法、系统及电子设备 - Google Patents

多进程共享的gpu调度方法、系统及电子设备 Download PDF

Info

Publication number
CN111124691B
CN111124691B CN202010000285.4A CN202010000285A CN111124691B CN 111124691 B CN111124691 B CN 111124691B CN 202010000285 A CN202010000285 A CN 202010000285A CN 111124691 B CN111124691 B CN 111124691B
Authority
CN
China
Prior art keywords
kernel
gpu
resource allocation
percentage
function
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
CN202010000285.4A
Other languages
English (en)
Other versions
CN111124691A (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 CN202010000285.4A priority Critical patent/CN111124691B/zh
Publication of CN111124691A publication Critical patent/CN111124691A/zh
Application granted granted Critical
Publication of CN111124691B publication Critical patent/CN111124691B/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
    • 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
    • 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/5011—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals
    • G06F9/5016—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals the resource being the memory
    • 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/54—Interprogram communication
    • G06F9/544—Buffers; Shared memory; Pipes
    • 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/54—Interprogram communication
    • G06F9/545—Interprogram communication where tasks reside in different layers, e.g. user- and kernel-space
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F2209/00—Indexing scheme relating to G06F9/00
    • G06F2209/50—Indexing scheme relating to G06F9/50
    • G06F2209/5011—Pool

Abstract

本发明提供一种多进程共享的GPU调度方法、系统及电子设备，所述多进程共享的GPU调度方法包括：从同时运行多个相同用户程序的进程池中捕获所述用户程序的某一个内核，并将该内核确定为待调度的内核；获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；根据确定的内核的资源分配百分比向GPU发射所述内核。实现一种基于Volta MPS多进程共享GPU调度策略，利用资源互补策略提升GPU整体吞吐率，通过线下的profile结果决定任务的计算资源分配方式，设计并实现了一种进程池机制，在Volta MPS的基础上可以对进程的GPU计算资源进行动态分配，实现最大化GPU系统的整体吞吐量。

Description

多进程共享的GPU调度方法、系统及电子设备
技术领域
本发明涉及处理器技术领域，特别是涉及GPU技术领域。
背景技术
最近机器学习已然成为了一个热门方向，不管是学术领域，还是在工业领域都有占有重要地位。其应用方向也很广泛，比如，医学方面可以用来诊断疾病，药物生产等；刑侦领域利用人脸识别技术追踪嫌疑犯；银行可以利用机器学习技术检测信用卡欺诈行为；语音文字识别。机器学习离不开大数据的支持，普通的GPU(Graphics Processing Unit，图形处理器)已经不能满足性能和TCO(Total Cost of Ownership，总拥有成本)的需求，机器学习往往会利用加速器来弥补这些不足，其中GPU是常用的加速器之一。之所以GPU可以做为机器学习的加速器，主要原因是GPU采用SIMD(Single Instruction Multiple Data，单指令多数据流)的架构，更擅长处理对大数据的简单重复操作。尽管如此，复杂的深度模型训练仍然耗时严重，因此关于GPU上机器学习训练任务的效率性能提升等问题的研究也越来越重要。
如果单独一个程序独占GPU，往往是不能占满GPU资源的，因为一般来说单个程序是占不满整个GPU资源的。为了提高GPU利用率可以同时跑多个GPU程序，在Nvidia(英伟达)提出Volta架构之前，多程序共享GPU只能时分的共享，通过先来先服务的方式对GPU上的kernel任务进行调度。在Volta架构出现后，利用Nvidia提供的MPS(Multi-ProcessService，多进程服务)可以进一步实现空间上共享GPU。Volta MPS旨在通过减少任务的上下文切换、解决Block(线程块)或Thread(线程)数量过少造成的资源浪费来提升GPU的整体吞吐率，但它不会考虑GPU上多任务的资源竞争问题，这些资源包括GPU计算资源，L2-cache(GPU的二级缓存)，访存带宽等。不同程序对GPU资源的需求不同，有计算密集型任务，有访存密集型任务，可以利用这些任务对资源需求的不同，合理重分配资源以达到资源互补从而提高整体吞吐率。
发明内容
鉴于以上所述现有技术的缺点，本发明的目的在于提供一种多进程共享的GPU调度方法、系统及电子设备，用于提高GPU资源的利用率和GPU的吞吐率。
为实现上述目的及其他相关目的，本发明提供一种多进程共享的GPU调度方法，包括：从同时运行多个相同用户程序的进程池中捕获所述用户程序的某一个内核，并将该内核确定为待调度的内核；获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；根据确定的内核的资源分配百分比向GPU发射所述内核。
于本发明的一实施例中，获取所述内核的特征的一种实现方式包括：在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置；通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
于本发明的一实施例中，所述确定所述内核的资源分配百分比的一种实现方式为：内核的资源分配百分比满足以下条件：p1+p2+...+pn＝100；
Figure BDA0002352835440000021
其中，p1,p2...pn分别表示为用户程序1的资源分配百分比，用户程序2的资源分配百分比，用户程序n的资源分配百分比，n表示为用户程序的数量，
Figure BDA0002352835440000022
表示第n个用户程序资源分配百分比为p时吞吐率SM的大小。
于本发明的一实施例中，每一个所述进程池对应于共享内存的一个固定位置；在所述进程池捕获到某一个内核后，将所述内核的信息写入所述共享内存中，并修改所述内核的状态为等待调度状态，将该内核确定为待调度的内核，所述进程池自身进入忙等状态；在确定所述内核的资源分配百分比之后，修改所述内核的状态为完成分配状态；在所述内核在GPU上执行完毕之后，更新所述共享内存中的已完成内核的个数。
于本发明的一实施例中，所述多进程共享的GPU调度方法还包括：捕获所述用户程序的发射所述内核的CUDA函数。
于本发明的一实施例中，所述捕获所述用户程序的发射所述内核的函数的一种实现方式包括：建立一动态库，并于所述动态库中预设若干重写函数；在CUDA运算程序运行时预加载该动态库，在CUDA函数发射前跳转到动态库的重写函数中；从所述重写函数中捕获所述用户程序的发射所述内核的CUDA函数。
于本发明的一实施例中，捕获的CUDA函数包括申请显存的函数、数据拷贝的函数、发射内核的函数、等待计算完成函数以及释放显存函数。
本发明的实施例还提供一种多进程共享的GPU调度系统，所述多进程共享的GPU调度系统包括：进程池模块，包括多个进程池，每一个所述进程池同时运行多个相同用户程序，所述进程池用于捕获所述用户程序的某一个内核，将该内核确定为待调度的内核，并根据确定的内核的资源分配百分比向GPU发射所述内核；调度器，获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；共享内存，建立所述进程池模块与所述调度器之间的通讯连接。
于本发明的一实施例中，所述调度器在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置，并通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
本发明的实施例还提供一种电子设备，包括处理器和存储器，所述存储器存储有程序指令，所述处理器运行程序指令实现如上所述的多进程共享的GPU调度方法。
如上所述，本发明的多进程共享的GPU调度方法、系统及电子设备具有以下有益效果：
本发明实现一种基于Volta MPS多进程共享GPU调度策略，利用资源互补策略提升GPU整体吞吐率，通过线下的profile结果决定任务的计算资源分配方式，设计并实现了一种进程池机制，在Volta MPS的基础上可以对进程的GPU计算资源进行动态分配，实现最大化GPU系统的整体吞吐量。
附图说明
为了更清楚地说明本发明实施例中的技术方案，下面将对实施例描述中所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本发明的一些实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据这些附图获得其他的附图。
图1显示为本申请一实施例中的多进程共享的GPU调度方法的整体流程示意图。
图2显示为本申请一实施例中的多进程共享的GPU调度方法中的交互架构图。
图3显示为本申请一实施例中的多进程共享的GPU调度方法的调度过程协同图。
图4显示为本申请一实施例中的计算机密集型内核的SM利用率随计算资源变化而变化的曲线图。
图5显示为本申请一实施例中的访存机密集型内核的SM利用率随计算资源变化而变化的曲线图。
图6显示为本申请一实施例中的多进程共享的GPU调度系统的原理框图。
图7显示为本申请一实施例中的电子设备的结构示意图。
元件标号说明
100 多进程共享的GPU调度系统
110 进程池模块
120 调度器
130 共享内存
1101 处理器
1102 存储器
S110～S130 步骤
具体实施方式
以下通过特定的具体实例说明本发明的实施方式，本领域技术人员可由本说明书所揭露的内容轻易地了解本发明的其他优点与功效。本发明还可以通过另外不同的具体实施方式加以实施或应用，本说明书中的各项细节也可以基于不同观点与应用，在没有背离本发明的精神下进行各种修饰或改变。需说明的是，在不冲突的情况下，以下实施例及实施例中的特征可以相互组合。
本实施例的目的在于提供一种多进程共享的GPU调度方法、系统及电子设备，用于提高GPU资源的利用率和GPU的吞吐率。
以下将详细阐述本发明的多进程共享的GPU调度方法、系统及电子设备的原理及实施方式，使本领域技术人员不需要创造性劳动即可理解本发明的多进程共享的GPU调度方法、系统及电子设备。
Volta是Nvidia(英伟达)公布的新一代GPU架构。Volta将具有1000GB/s带宽(GTX680的带宽是192GB/s)。Volta将硅基板置于DRAM和GPU之间，然后在基板上穿孔并且连接各层组件，最终就达成了「五十分之一秒内传完整张蓝光光盘数据」的惊人效果。在Volta架构出现后，利用Nvidia提供的MPS(Multi-Process Service，多进程服务)可以进一步实现空间上共享GPU。Volta MPS旨在通过减少任务的上下文切换、解决Block(线程块)或Thread(线程)数量过少造成的资源浪费来提升GPU的整体吞吐率。
本实施例的多进程共享的GPU调度方法、系统及电子设备基于Volta MPS多进程共享GPU调度策略，利用资源互补策略提升GPU整体吞吐率。本实施例通过线下的profile结果决定任务的计算资源分配方式，设计并实现了一种进程池机制，在Volta MPS的基础上可以对进程的GPU计算资源进行动态分配。其中，profile是在不同的环境下配置用不同的配置文件或者不同的配置。
具体地，如图1所示，本实施例提供一种多进程共享的GPU调度方法，所述多进程共享的GPU调度方法包括：
步骤S110，从同时运行多个相同用户程序的进程池中捕获所述用户程序的某一个内核，并将该内核确定为待调度的内核；
步骤S120，获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；
步骤S130，根据确定的内核的资源分配百分比向GPU发射所述内核。
以下将结合图2和图3对本实施例的多进程共享的GPU调度方法进行详细说明。
步骤S110，从同时运行多个相同用户程序的进程池中捕获所述用户程序的某一个内核，并将该内核确定为待调度的内核。
Volta MPS可以指定CUDA程序运行时占用的最大计算资源百分比，即可以限制程序最多使用多少SM。其中，CUDA(Compute Unified Device Architecture)，是显卡厂商NVIDIA推出的运算平台。CUDA是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题，CUDA包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。
具体地，于本实施例中，当CUDA程序第一次初始化CUDA时，CUDA驱动程序将尝试连接到MPS守护进程。如果连接尝试成功，则MPS守护进程的用户id与CUDA程序相同时，则CUDA程序继续连接到MPS server并以当前MPS server指定的计算资源百分比占用GPU。之后当改变MPS server的资源百分比后，之前已经连接的程序的资源百分比不会改变，即MPS只能静态的分配程序的计算资源百分比，运行过程中不可以更改百分比。仅依靠MPS不能完整的实现上述算法，所以本时时乐在MPS的基础上实现了进程池机制，可以实现资源百分比的动态调整，程序中的不同kernel可以以不同的计算资源量执行。
由于单个用户程序只能对应一个百分比，单独一个用户程序不能完成百分比的动态调整。假如要完成单个用户程序的百分比动态调整，可以同时运行多个相同的该用户程序，每个用户程序对应不同的百分比，这些用户程序可以看做是一个进程池。对于该用户程序中的某一个内核(kernel)，可以根据资源要求只选取进程池中相对应的一个用户程序执行，其余用户程序跳过，不发射内核kernel。这样在GPU上执行的工作量相当于来自一个程序，而不是把进程池中所有程序的kernel都执行一遍，没有额外增加GPU的工作量。
实现上述情况，需要使用Hook技术，将GPU相关动作hook到，根据情况作出相应处理。Hook技术又叫做钩子函数，在系统没有调用该函数之前，钩子程序就先捕获该消息，钩子函数先得到控制权，这时钩子函数既可以加工处理(改变)该函数的执行行为，还可以强制结束消息的传递。简单来说，就是把系统的程序拉出来变成独立执行代码片段。
具体地，如图2和图3所示，于本实施例中，每一个所述进程池对应于共享内存的一个固定位置；在所述进程池捕获到某一个内核后，将所述内核的信息写入所述共享内存中，并修改所述内核的状态为等待调度状态，将该内核确定为待调度的内核，所述进程池自身进入忙等状态；在确定所述内核的资源分配百分比之后，修改所述内核的状态为完成分配状态；在所述内核在GPU上执行完毕之后，更新所述共享内存中的已完成内核的个数。
步骤S120，获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比。
于本实施例中，获取所述内核的特征的一种实现方式包括：
在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置；通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
调度器做决策的依据是每个cuda程序的kernel特征，需要知道kernel是否是计算密集型或者访存密集型、在一定计算资源下的运行速度、任务完成所需时间。不同kernel由于自身代码的特性以及GPU机器的特性，运行时会表现出不同的特征。
本实施例采用线下采集数据的方式，在cuda任务跑之前，对其进行profile，拿到kernel的运行特性。Profile工作利用Nvidia提供的工具——NVIDIA Nsight Compute完成。NVIDIA Nsight Graphics是一个独立的开发工具，提供了一种数据驱动方法来查看应用程序如何真正使用系统资源，以及在何处集中分析工作。Nsight Compute可以拿到每个kernel的SM吞吐率、内存吞吐，SM(计算单元，Streaming Multiprocessor)吞吐率衡量了SM的利用率，内存吞吐为平均每秒访存量。当分配给当前程序的计算资源为10％、20％、30％，…，直到100％时，依次测出以上SM吞吐率、内存吞吐的值。
计算密集型往往SM吞吐率高，内存吞吐低，SM吞吐率和内存吞吐随着分配的计算资源百分比的增加而增加。图4是一个典型的计算机密集型kernel所表现出的特点，为Rodinia中lavaMD的一个kernel的SM利用率随计算资源变化而变化的曲线。
访存密集型任务访存吞吐高，SM吞吐低，当资源开始增加时，吞吐会随之增加，但是当增加到一定程度时，访存会成为瓶颈，达到机器的最大访存带宽，吞吐不会再随着计算资源的增加而增加。图5是一个典型的访存机密集型kernel所表现出的特点，为Rodinia中hotspot3D的一个kernel的SM利用率随计算资源变化而变化的曲线。
当计算密集型任务与访存密集型同时运行时，不加控制时，可能GPU的访存速度一段时间内会很高，当访存密集任务完成后，访存速度又变低，希望在程序完成之前，GPU整体访存速度保持平缓，减少访存竞争，使不同类型的任务尽可能实现资源互补。可以降低访存密集型任务的计算资源百分比以降低整体访存速度，其余计算资源给计算密集型任务，这样当所有任务同时完成时，可以完成资源互补的目标。对于单个程序而言，当工作量一定时，完成时间与SM吞吐率大致成反比，SM吞吐率可以看成是任务处理速度，这样就有公式：SM吞吐率*时间＝工作量。SM吞吐率是受计算资源百分比影响的，这样目标就转换成如何为每个程序选取合适的计算资源百分比，使所有程序同时完成。
将上述问题描述转换为数学形式：假设每个程序只有一种kernel，有n个程序，它们单独运行时耗时相同，同时运行时给他们分配的资源百分比分别为p1％，p2％，…，pn％。
于本实施例中，所述确定所述内核的资源分配百分比的一种实现方式为：内核的资源分配百分比满足以下条件：
p1+p2+...+pn＝100；
Figure BDA0002352835440000071
其中，p1,p2...pn分别表示为用户程序1的资源分配百分比，用户程序2的资源分配百分比，用户程序n的资源分配百分比，n表示为用户程序的数量，
Figure BDA0002352835440000072
表示第n个用户程序资源分配百分比为p时吞吐率SM的大小。
方程的求解方式：由于pi(i＝1，…，n)的取值范围是有限个点，所以不能得到刚好满足条件的解，只能选取近似解。首先选取p1的值，得到
Figure BDA0002352835440000073
的值，找到p2与p2’使得其满足以下条件：
1)
Figure BDA0002352835440000074
2)p2与p2’之间没有其他合法取值，在本文中可以转换成p2’-p2＝10，因为profile是每隔10％取一次值的。
同理找到其他pi的取值范围，如果p1+p2+...+pn<＝100<＝p1+p2’+...+pn’，则找到pi的解。
由于在理想条件下不能使所有任务同时完成，而且实际运行时，由于kernel之间的互相影响，实际完成时间与预计完成时间也会不同。这样当有任务完成时，应该重新按照上述方法分配其余任务的百分比。
步骤S130，根据确定的内核的资源分配百分比向GPU发射所述内核。
于本实施例中，通讯功能在本实施例中是利用共享内存实现的，kernel(内核)的调度需要共享内存协同完成，过程如图3所示。每个进程池对应于共享内存的一个固定位置。首先当进程池hook到一个kernel后，会将当前kernel id更新到共享内存中，并修改kernel的调度状态为等待调度，自身进入忙等状态，直到调度器分配好kernel百分比。调度器查询到有等待调度的kernel后，根据调度算法为其分配百分比，并将百分比写回，修改状态为已完成百分比分配。进程池查看到状态为分配完成，跳出轮询，按照分配的百分比发射kernel。进程池会为发射kernel的函数添加callback函数，当kernel执行完毕后，更新共享内存中的已完成kernel个数。
于本实施例中，所述多进程共享的GPU调度方法还包括：捕获所述用户程序的发射所述内核的CUDA函数。
具体地，于本实施例中，所述捕获所述用户程序的发射所述内核的函数的一种实现方式包括：
建立一动态库，并于所述动态库中预设若干重写函数；在CUDA运算程序运行时预加载该动态库，在CUDA函数发射前跳转到动态库的重写函数中；从所述重写函数中捕获所述用户程序的发射所述内核的CUDA函数。
具体地，于本实施例中，需要使用hook技术，将GPU相关动作hook到，根据情况作出相应处理。实施例的hook实现的方法为：实现一个动态库，在动态库中重写所有需要hook的函数，利用LD_PRELOAD在程序运行时预加载该动态库，在cuda函数发射前会跳转到动态库的重写函数中；当需要调用cuda实现的原函数时，利用dlsym()函数和RTLD_NEXT参数找到原函数。这样就可以在相关cuda函数的执行前后进行代码注入。
一般来说，cuda程序会执行的gpu动作大致包括：cuda初始化，选取GPU作为kernel发射的目的地，申请GPU显存，将CPU的数据拷贝到GPU上，发射kernel在GPU上完成计算工作，等待GPU完成计算工作，从GPU上把计算结果拷贝回CPU，释放GPU显存。对于进程池的实现来说，只有cuda初始化和GPU的选取不需要控制，可以不考虑，其余有关显存以及计算的cuda函数都需要hook。
其中，于本实施例中，捕获的CUDA函数包括但不限于申请显存的函数、数据拷贝的函数、发射内核的函数、等待计算完成函数以及释放显存函数。
1)申请显存的函数：cudaMalloc。
函数api为：cudaError_t cudaMalloc(void**devPtr,size_t size)。作用是申请一段大小为size的显存，devPtr是指向分配的显存的指针。进程池只需一个进程完成显存分配即可，可以选取资源百分比对应为100的进程来完成，其他进程不需申请但可以使用申请到的显存。由于不同进程的地址空间是不同的，不可以直接使用devPtr指向的地址，可以使用cudaIpcMemHandle_t进行地址的转化与取得。如果当前程序百分比为100：使用cudaMalloc申请内存，调用cudaIpcGetMemHandle获取devPtr指向的显存的句柄，将句柄放入共享内存中以供其他程序取用。如果当前程序百分比不是100：等待直到显存申请完毕，拿到句柄，调用cudaIpcOpenMemHandle函数拿到指针，将指针存到devPtr中。
2)数据拷贝的函数，cudaMemcpy。
函数api为：cudaError_t cudaMemcpy(void*dst,const void*src,size_tcount,cudaMemcpyKind kind)。作用是把src指向的数据拷贝count大小到dst，kind指明拷贝种类(HostToDevice、DeviceToHost等)。同理，数据拷贝也是由资源百分比对应为100的程序来完成。只有当数据从GPU端拷贝到CPU端时，其余程序也需要完成。
3)发射内核的函数，cudaLaunchKernel。
函数api为：cudaError_t cudaLaunchKernel(const void*func,dim3 gridDim,dim3 blockDim,void**args,size_t sharedMem,cudaStream_t stream)。作用是发射kernel函数到GPU。判断当前kernel需要的资源百分比是否与程序相符，如果相符则发射kernel，不符合直接返回。
4)等待计算完成函数，cudaDeviceSynchronize。
函数api为：cudaError_t cudaDeviceSynchronize()。作用是等待GPU完成计算工作。在进程池中就变成了等待所有进程完成计算工作，即所有进程都会执行cudaDeviceSynchronize函数，然后等所有其他进程也执行完毕。
5)释放显存函数，5.cudaFree。次函数与cudaMalloc函数相对应，是释放显存的，同样只执行一次。
由上可见，本实施例提供了一套基于Volta MPS的GPU资源调度方案，可以利用profile的到的数据计算出每个程序的合适资源百分比，利用进程池机制对程序的资源百分比进行动态调整，实现最大化GPU系统的整体吞吐量。
本实施例对一些benchmark效果较为明显。比如，rodinia中的lavaMD、srad_v2、hotspot3D、kmean，分别丛中选取一个kernel，迭代数次，使他们的完成时间相同。其中前两个为计算密集型，后两个为访存密集型，分别从中选取一个，两两组合成4组。与都为其分配100％资源相比，本实施例中的方法效果提升为12.1％到21.8％；与都为其分配50％资源相比，本实施例中的方法效果提升为13.6％到26.7％。
如图6所示，本实施例还提供一种多进程共享的GPU调度系统100，所述多进程共享的GPU调度系统100包括：进程池模块110，调度器120以及共享内存130。
于本实施例中，所述进程池模块110包括多个进程池，每一个所述进程池同时运行多个相同用户程序，所述进程池用于捕获所述用户程序的某一个内核，将该内核确定为待调度的内核，并根据确定的内核的资源分配百分比向GPU发射所述内核；所述调度器120获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；所述共享内存130建立所述进程池模块110与所述调度器120之间的通讯连接。
所述调度器120在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置，并通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
本实施例中，进程池模块110可以hook到用户程序的发射kernel的函数，并可以以指定好的GPU计算资源百分比完成kernel；调度器120接收所有进程池发来的信息，根据预先profile的结果决定当前百分比分配方式；共享内存130构成通讯模块负责进程池与调度器120的交互。
本实施例的多进程共享的GPU调度系统100具体实现的技术特征与前述实施例中的多进程共享的GPU调度方法基本相同，实施例间可以通用的技术内容不作重复赘述。
如图7所示，展示本申请一实施例中的电子设备的结构示意图，所述电子设备包括处理器1101和存储器1102；处理器1101为GPU(Graphics Processing Unit，图形处理器)。存储器1102通过系统总线与处理器1101连接并完成相互间的通信，存储器1102用于存储计算机程序，处理器1101用于运行计算机程序，以使所述处理器1101执行所述的多进程共享的GPU调度方法。上述已经对所述多进程共享的GPU调度方法进行了详细说明，在此不再赘述。
另需说明的是，上述提到的系统总线可以是外设部件互连标准(PeripheralComponent Interconnect，简称PCI)总线或扩展工业标准结构(Extended IndustryStandard Architecture，简称EISA)总线等。该系统总线可以分为地址总线、数据总线、控制总线等。为便于表示，图中仅用一条粗线表示，但并不表示仅有一根总线或一种类型的总线。通信接口用于实现数据库访问系统与其他设备(例如客户端、读写库和只读库)之间的通信。存储器可能包含随机存取存储器(Random Access Memory，简称RAM)，也可能还包括非易失性存储器(non-volatile memory)，例如至少一个磁盘存储器。
此外，本实施例还提供一种计算机可读存储介质，其上存储有计算机程序，所述计算机程序被处理器执行时实现所述的多进程共享的GPU调度方法。上述已经对所述多进程共享的GPU调度方法进行了详细说明，在此不再赘述。
本领域普通技术人员可以理解：实现上述各方法实施例的全部或部分步骤可以通过计算机程序相关的硬件来完成。前述的计算机程序可以存储于一计算机可读存储介质中。该程序在执行时，执行包括上述各方法实施例的步骤；而前述的存储介质包括：ROM、RAM、磁碟或者光盘等各种可以存储程序代码的介质。
综上所述，本发明实现一种基于Volta MPS多进程共享GPU调度策略，利用资源互补策略提升GPU整体吞吐率，通过线下的profile结果决定任务的计算资源分配方式，设计并实现了一种进程池机制，在Volta MPS的基础上可以对进程的GPU计算资源进行动态分配，实现最大化GPU系统的整体吞吐量。所以，本发明有效克服了现有技术中的种种缺点而具高度产业利用价值。
上述实施例仅例示性说明本发明的原理及其功效，而非用于限制本发明。任何熟悉此技术的人士皆可在不违背本发明的精神及范畴下，对上述实施例进行修饰或改变。因此，举凡所属技术领域中包括通常知识者在未脱离本发明所揭示的精神与技术思想下所完成的一切等效修饰或改变，仍应由本发明的权利要求所涵盖。

Claims (9)

1.一种多进程共享的GPU调度方法，其特征在于：包括：
从同时运行多个相同用户程序的进程池中捕获所述用户程序的某一个内核，并将该内核确定为待调度的内核；
获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；
根据确定的内核的资源分配百分比向GPU发射所述内核；
所述确定所述内核的资源分配百分比的一种实现方式为：
内核的资源分配百分比满足以下条件：
p1+p2+...+pn＝100；
Figure FDA0003856703210000011
其中，p1,p2...pn分别表示为用户程序1的资源分配百分比，用户程序2的资源分配百分比，用户程序n的资源分配百分比，n表示为用户程序的数量，
Figure FDA0003856703210000012
表示第n个用户程序资源分配百分比为p时其SM吞吐率的大小。
2.根据权利要求1所述的多进程共享的GPU调度方法，其特征在于：获取所述内核的特征的一种实现方式包括：
在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置；
通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
3.根据权利要求1所述的多进程共享的GPU调度方法，其特征在于：每一个所述进程池对应于共享内存的一个固定位置；在所述进程池捕获到某一个内核后，将所述内核的信息写入所述共享内存中，并修改所述内核的状态为等待调度状态，将该内核确定为待调度的内核，所述进程池自身进入忙等状态；在确定所述内核的资源分配百分比之后，修改所述内核的状态为完成分配状态；在发射所述内核之后，更新所述共享内存中的已发射内核的个数。
4.根据权利要求1所述的多进程共享的GPU调度方法，其特征在于：所述多进程共享的GPU调度方法还包括：
捕获所述用户程序的发射所述内核的CUDA函数。
5.根据权利要求4所述的多进程共享的GPU调度方法，其特征在于：所述捕获所述用户程序的发射所述内核的函数的一种实现方式包括：
建立一动态库，并于所述动态库中预设若干重写函数；
在CUDA运算程序运行时预加载该动态库，在CUDA函数发射前跳转到动态库的重写函数中；
从所述重写函数中捕获所述用户程序的发射所述内核的CUDA函数。
6.根据权利要求5所述的多进程共享的GPU调度方法，其特征在于：捕获的CUDA函数包括申请显存的函数、数据拷贝的函数、发射内核的函数、等待计算完成函数以及释放显存函数。
7.一种多进程共享的GPU调度系统，其特征在于：所述多进程共享的GPU调度系统包括：
进程池模块，包括多个进程池，每一个所述进程池同时运行多个相同用户程序，所述进程池用于捕获所述用户程序的某一个内核，将该内核确定为待调度的内核，并根据确定的内核的资源分配百分比向GPU发射所述内核；
调度器，获取待调度的内核，根据所述内核的特征确定所述内核的资源分配百分比；确定所述内核的资源分配百分比的一种实现方式为：
内核的资源分配百分比满足以下条件：
p1+p2+...+pn＝100；
Figure FDA0003856703210000021
其中，p1,p2...pn分别表示为用户程序1的资源分配百分比，用户程序2的资源分配百分比，用户程序n的资源分配百分比，n表示为用户程序的数量，
Figure FDA0003856703210000022
表示第n个用户程序资源分配百分比为p时其SM吞吐率的大小；
共享内存，建立所述进程池模块与所述调度器之间的通讯连接。
8.根据权利要求7所述的多进程共享的GPU调度系统，其特征在于：所述调度器在GPU的CUDA运算程序运行之前，对所述CUDA运算程序预先进行Profile配置，并通过CUDA内核分析工具NVIDIA Nsight Compute获取所述内核的特征。
9.一种电子设备，其特征在于：包括处理器和存储器，所述存储器存储有程序指令，所述处理器运行程序指令实现如权利要求1至权利要求6任一权利要求所述的多进程共享的GPU调度方法。
CN202010000285.4A 2020-01-02 2020-01-02 多进程共享的gpu调度方法、系统及电子设备 Active CN111124691B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202010000285.4A CN111124691B (zh) 2020-01-02 2020-01-02 多进程共享的gpu调度方法、系统及电子设备

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202010000285.4A CN111124691B (zh) 2020-01-02 2020-01-02 多进程共享的gpu调度方法、系统及电子设备

Publications (2)

Publication Number Publication Date
CN111124691A CN111124691A (zh) 2020-05-08
CN111124691B true CN111124691B (zh) 2022-11-25

Family

ID=70507308

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202010000285.4A Active CN111124691B (zh) 2020-01-02 2020-01-02 多进程共享的gpu调度方法、系统及电子设备

Country Status (1)

Country Link
CN (1) CN111124691B (zh)

Families Citing this family (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114217976B (zh) * 2021-12-23 2023-02-28 北京百度网讯科技有限公司 任务处理方法、装置、设备及存储介质

Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN1514354A (zh) * 2002-12-31 2004-07-21 ����̩ƽ 面向构件基于系统内核的进程池/线程池管理方法
CN102253919A (zh) * 2011-05-25 2011-11-23 中国石油集团川庆钻探工程有限公司 基于gpu和cpu协同运算的并行数值模拟方法和系统
CN107122245A (zh) * 2017-04-25 2017-09-01 上海交通大学 Gpu任务调度方法及系统
CN109857564A (zh) * 2019-03-05 2019-06-07 上海交通大学 基于细粒度的gpu的资源管理方法及其应用的gpu
CN110083488A (zh) * 2019-04-21 2019-08-02 哈尔滨工业大学 一种面向gpgpu的细粒度低开销的容错系统

Family Cites Families (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN106325996B (zh) * 2015-06-19 2019-11-19 华为技术有限公司 一种gpu资源的分配方法及系统

Patent Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN1514354A (zh) * 2002-12-31 2004-07-21 ����̩ƽ 面向构件基于系统内核的进程池/线程池管理方法
CN102253919A (zh) * 2011-05-25 2011-11-23 中国石油集团川庆钻探工程有限公司 基于gpu和cpu协同运算的并行数值模拟方法和系统
CN107122245A (zh) * 2017-04-25 2017-09-01 上海交通大学 Gpu任务调度方法及系统
CN109857564A (zh) * 2019-03-05 2019-06-07 上海交通大学 基于细粒度的gpu的资源管理方法及其应用的gpu
CN110083488A (zh) * 2019-04-21 2019-08-02 哈尔滨工业大学 一种面向gpgpu的细粒度低开销的容错系统

Also Published As

Publication number Publication date
CN111124691A (zh) 2020-05-08

Similar Documents

Publication Publication Date Title
Warneke et al. Exploiting dynamic resource allocation for efficient parallel data processing in the cloud
US8713574B2 (en) Soft co-processors to provide a software service function off-load architecture in a multi-core processing environment
US7975269B2 (en) Parallel processor methods and apparatus
US9891949B2 (en) System and method for runtime scheduling of GPU tasks
US8141076B2 (en) Cell processor methods and apparatus
Kang et al. Lalarand: Flexible layer-by-layer cpu/gpu scheduling for real-time dnn tasks
WO2011123991A1 (zh) 并行计算的内存访问方法
CN111966504B (zh) 图形处理器中的任务处理方法及相关设备
US11880715B2 (en) Method and system for opportunistic load balancing in neural networks using metadata
KR20130080663A (ko) 멀티-쓰레딩을 사용하는 그래픽 처리를 위한 방법 및 장치
CN114721818A (zh) 一种基于Kubernetes集群的GPU分时共享方法和系统
CN111124691B (zh) 多进程共享的gpu调度方法、系统及电子设备
CN114217930A (zh) 一种基于混合任务调度的加速器系统资源优化管理方法
CN103262039A (zh) 用于处理装置的同步操作的方法和系统
CN114637536A (zh) 任务处理方法、计算协处理器、芯片及计算机设备
US20150212859A1 (en) Graphics processing unit controller, host system, and methods
CN114816777A (zh) 命令处理装置、方法、电子设备以及计算机可读存储介质
Ino et al. Cooperative multitasking for GPU‐accelerated grid systems
US20230236878A1 (en) Efficiently launching tasks on a processor
CN112114967B (zh) 一种基于服务优先级的gpu资源预留方法
CN111930485B (zh) 一种基于性能表现的作业调度方法
JP2023544911A (ja) 並列量子コンピューティングのための方法及び装置
Plauth et al. CloudCL: distributed heterogeneous computing on cloud scale
Butler et al. Improving application concurrency on GPUs by managing implicit and explicit synchronizations
Simsek et al. Leveraging data-flow task parallelism for locality-aware dynamic scheduling on heterogeneous platforms

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