CN114549277A - 一种基于编译的核函数自动多流调度方法 - Google Patents

一种基于编译的核函数自动多流调度方法 Download PDF

Info

Publication number
CN114549277A
CN114549277A CN202210172808.2A CN202210172808A CN114549277A CN 114549277 A CN114549277 A CN 114549277A CN 202210172808 A CN202210172808 A CN 202210172808A CN 114549277 A CN114549277 A CN 114549277A
Authority
CN
China
Prior art keywords
node
kernel function
kernel
flow
precursor
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.)
Pending
Application number
CN202210172808.2A
Other languages
English (en)
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.)
Sun Yat Sen University
Original Assignee
Sun Yat Sen 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 Sun Yat Sen University filed Critical Sun Yat Sen University
Priority to CN202210172808.2A priority Critical patent/CN114549277A/zh
Publication of CN114549277A publication Critical patent/CN114549277A/zh
Pending legal-status Critical Current

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/20—Processor architectures; Processor configuration, e.g. pipelining
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/44—Encoding
    • G06F8/447—Target code generation
    • 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
    • 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)
  • Theoretical Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Devices For Executing Special Programs (AREA)

Abstract

本发明公开了一种基于编译的核函数自动多流调度方法，包括：读取源代码，识别核函数的输入与输出；根据输入与输出构建DAG；对DAG进行层次化，并根据拓扑排序为核函数赋予适当的流；为跨流同步的核函数插入事件同步代码；生成可执行文件。本发明是首个通过LLVM编译器框架将串行的GPU核函数通过CUDA框架的多流机制并行化，从而提高程序性能与GPU硬件利用率。该发明不引入额外的编程模型与运行时开销，减小了开发者的学习成本与开发成本。同时该发明使用的流调度算法将存在数据依赖关系的核函数的运行时机临近安排，尽可能提高GPU对数据的访存效率，从而提高程序性能。

Description

一种基于编译的核函数自动多流调度方法
技术领域
本发明涉及计算机技术领域，尤其涉及一种基于编译的核函数自动多流调度方法。
背景技术
随着GPU算力的提升，计算量与数据吞吐量日益增长的AI应用在GPU上的应用越来越广泛。AI应用的基本运算单位为算子，多种算子以特定的网络结构组织起来形成了完整的AI模型。因此AI模型的性能取决于两个部分，一个是算子内的计算性能，另一个是算子之间的网络结构决定的资源占用率、访存效率等等。由于实际应用中单算子的计算与数据读写难以占用全部的GPU资源，因此如果能够让单一时刻GPU上有多个算子同时在运行，那么就可以使用更多的GPU资源，提高程序性能。开发者为了达到核函数(即算子)并行的效果，可以使用CUDA框架中的流与事件机制实现。
GPU支持用户同时使用多个流，不同流之间的任务则是异步执行的，这些任务在同一时刻可能会同时执行。因此为了确保不同流之间存在依赖的任务有序执行，可以使用事件进行跨流同步。
但是这需要用户深度学习CUDA框架中流、事件的相关知识以及一定的异构程序开发经验。同时也需要用户正确分析复杂程序中可能多达上百个核函数之间的数据依赖关系。核函数之间的数据依赖关系可以分为写后读、写后写、读后写和读后读。其中写后读指的是当前核函数的读取变量会被前面的核函数修改；写后写指的是当前核函数的修改变量会被前面的核函数修改；读后写指的是当前核函数的修改变量会被前面的核函数读取；读后读指的是当前核函数的读取变量会被前面的核函数读取。
同一时刻并行的核函数之间不能存在数据依赖关系，否则会出现最终结果不正确的错误。在复杂程序中将核函数剥离出来进行并行需要开发者付出较多的开发成本修改得到正确的核函数并行代码。这些门槛与问题都阻碍了GPU的普及，特别是原本在CPU上实现的传统HPC应用向GPU的移植。
为了简化GPU程序的开发并提高其性能，任务图计算系统成为了首选。其中包含的面向任务的编程框架将每个核函数抽象为一个任务，将任务之间的数据依赖关系抽象为一个有向无环图(DAG)。然后再对该图使用拓扑排序，将无依赖的核函数安排在不同的gpu流中达到并行效果。对于存在依赖的核函数，若这些核函数放在不同的流中，需要在存在依赖的函数之间加入事件同步，从而保证在不同流中的核函数有序执行。另一方面，其中包含的运行时软件根据实时GPU硬件资源使用情况调度核函数，决定其何时在哪个GPU上运行或哪个CPU上运行等等。
近年来有许多新的任务图计算系统诞生，但是现有技术有如下缺点：
1.开发者需要学习新的面向任务编程框架，增加了开发者的学习负担。
2.开发者在将串行核函数修改为多流并行核函数时需要大幅度调整现有代码框架，增加了开发成本。
3.现有技术引入了运行时软件实时调度核函数的运行，引入了额外的运行时开销，降低了程序性能。
发明内容
有鉴于现有技术的上述缺陷，本发明所要解决的技术问题是提供一种基于编译的核函数自动多流调度方法，用户只需要根据基本的功能逻辑使用C++编写串行的核函数，本发明即可借助编译器自动分析核函数之间的数据依赖，将核函数构建为DAG，并将核函数分配到不同的流中并行，从而提高性能。该方法不引入额外的编程框架和运行时开销，最大程度降低开发者开发使用高性能异构程序的成本。
为实现上述目的，本发明提供了一种基于编译的核函数自动多流调度方法，包括以下步骤：
步骤1、读取源代码，识别核函数的输入与输出；
步骤2、根据输入与输出构建DAG；
步骤3、对DAG进行层次化，并根据拓扑排序为核函数赋予适当的流；
步骤4、为跨流同步的核函数插入事件同步代码；
步骤5、生成可执行文件。
进一步的，所述读取源代码，识别核函数的输入与输出，具体为：将源代码通过LLVM编译框架转变为IR文件，然后使用opt工具对所述IR文件进行处理，分析每个核函数的只读变量与写变量。
进一步的，所述根据输入与输出构建DAG，具体为：假设当前核函数有N个传入参数，对每一个传入参数将DAG的结束节点作为起点，使用倒序BFS算法找到第一个拥有同个传入参数的核函数，将该核函数添加为当前核函数的前驱节点，将当前核函数代表的节点设为DAG中结束节点的前驱节点。
进一步的，所述对DAG按层次顺序将每一层次的每个核函数赋予适当的流，具体为：对于每一层的任务节点顺序赋予递增id，使得每个任务节点都有当前层级中唯一的id；先对那些没有前驱节点的任务节点计算分配的流；然后再根据前驱节点数量从小到大的顺序遍历当前层级的任务节点，并按照以下策略进行流分配：
若当前节点的某个前驱节点为其所在流的末尾节点，则将当前节点分配到该流中；
若有多个符合上述条件的前驱节点，则选择未完成的后驱节点最少的那一个前驱节点所在的流；
若不存在上述前驱节点，则用当前任务节点计算分配的流。
进一步的，所述为跨流同步的核函数插入事件同步代码，具体为：
对于前驱节点与自身不在同一个流的所有节点，在拥有其前驱节点的每个流中，选择该流种最后一个前驱节点，在其核函数调用代码之后加入以CUDA/HIP编程框架中的事件触发器和CUDA/HIP编程框架下的指定任务流为参数的代码；
然后对于与该前驱节点存在依赖关系且不处于同个流中的任务节点，在其调用核函数的语句之前加入以CUDA/HIP编程框架中的事件触发器和CUDA/HIP编程框架下的指定任务流为参数的代码，以等待不同流的前驱函数执行完毕再继续执行该函数。
进一步的，所述生成可执行文件具体为通过LLVM将中间语言表示文件转化为可执行文件。
本发明的有益效果是：
本发明是首个通过LLVM编译器框架将串行的GPU核函数通过CUDA框架的多流机制并行化，从而提高程序性能与GPU硬件利用率。该发明不引入额外的编程模型与运行时开销，减小了开发者的学习成本与开发成本。同时该发明使用的流调度算法将存在数据依赖关系的核函数的运行时机临近安排，尽可能提高GPU对数据的访存效率，从而提高程序性能。
以下将结合附图对本发明的构思、具体结构及产生的技术效果作进一步说明，以充分地了解本发明的目的、特征和效果。
附图说明
图1是本发明的整体优化过程流程图。
图2是本发明的使用定义链对核函数做深度优先搜索的代码截图。
图3是本发明的根据输入与输出构建DAG流程图。
图4是本发明的宽度优先的倒序遍历算法流程图。
图5是本发明的对DAG按层次顺序将每一层次的每个核函数赋予适当的流的算法代码截图。
具体实施方式
本发明的相关变量或者词句的定义如下：
相关变量：对于任一变量，将与之出现在同一代码语句中的其他变量。包括了由该变量作为操作数时的返回变量以及由该变量作为返回变量时的所有操作数变量。
直接相关变量：与目标变量出现在同一代码语句中的相关变量。
间接相关变量：不与目标变量出现在同一代码语句中但可通过多个直接相关变量形成的图联通的相关变量。
定义-使用链：对于目标变量，一个所有使用该定义值作为操作数的语句的链表。
使用-定义链：对于目标变量，一个定义该值的语句的链表。
有向无环图(DAG)：如果一个有向图无法从任意顶点出发经过若干条边回到该点，则这个图就是有向无环图。其中节点表示一个函数或者一个任务，节点与节点之间的边表示依赖关系。
宽度优先搜索(BFS)：对于一个有向无环图，每一步都从当前节点出发，得到所有只经过一条边即可达到的其他节点，并将这些节点作为下一步的出发点。
拓扑排序：在一个有向无环图中，对于当前无前驱节点的节点使用BFS得到可达节点，然后将该节点与可达节点之间的边删去，再重复对那些无前驱节点的节点重复上述过程。最终得到遵循任务依赖关系的任务执行顺序队列。
层次化：对于一个有向无环图做拓扑排序，每一步可达到的所有节点均为同一层次的节点，从而将DAG图划分为多个层次。
CUDA:Nvidia提出的异构编程框架，允许用户编写GPU计算函数，将计算任务卸载到GPU上以及CPU和GPU之间通信的代码。
HIP:是AMD提出的异构编程框架，与CUDA功能相似，可兼容CUDA。
如图1所示，本发明提供了一种基于编译的核函数自动多流调度方法，包括以下步骤：
步骤1、读取源代码，识别核函数的输入与输出；
步骤2、根据输入与输出构建DAG；
步骤3、对DAG进行层次化，并根据拓扑排序为核函数赋予适当的流；
步骤4、为跨流同步的核函数插入事件同步代码；
步骤5、生成可执行文件。
本实施例中，读取源代码，识别核函数的输入与输出，具体为：将源代码通过LLVM编译框架转变为IR文件，然后使用opt工具对IR文件进行处理，分析每个核函数的只读变量与写变量。其中LLVM是一组编译器和工具链技术，可用于开发任何编程语言的前端和任何指令集架构的后端。LLVM是围绕独立于语言的中间表示(IR)设计的，它作为一种可移植的高级汇编语言，可以通过多次通过各种转换进行优化。Opt是LLVM工具链中负责对中间表示(IR)做优化的工具，优化的具体逻辑由用户实现。
在每个核函数的定义部分中，对于函数签名中的每一个源变量，根据LLVM中的定义-使用链及使用-定义链对其做如图2所示的深度优先搜索，寻找其所有直接/间接相关变量。
该算法对于每个当前搜索目标变量的使用语句进行遍历，这些语句可以是普通的赋值语句，也可以是函数调用语句。对于赋值语句，将该语句定义的变量作为新的搜索目标变量进行深度优先搜索。对于函数调用语句，则在该调用函数内部对当前搜索目标变量对应的传入参数进行深度优先搜索。当源变量涉及的某个搜索目标遍变量作为某个代码语句(比如数据写存语句)的目的变量且该语句不为该相关变量的定义语句时，则将该源变量作为该核函数的写变量，反之则为读变量。
本实施例中，根据输入与输出构建DAG，如图3所示，具体为：假设当前核函数有N个传入参数，对每一个传入参数将DAG的结束节点作为起点，使用倒序BFS算法找到第一个拥有同个传入参数的核函数，将该核函数添加为当前核函数的前驱节点，将当前核函数代表的节点设为DAG中结束节点的前驱节点。宽度优先的倒序遍历算法如图4所示。
本实施例中，对DAG按层次顺序将每一层次的每个核函数赋予适当的流，具体算法如图5所示，具体为：对于每一层的任务节点顺序赋予递增id，使得每个任务节点都有当前层级中唯一的id；先对那些没有前驱节点的任务节点计算分配的流；然后再根据前驱节点数量从小到大的顺序遍历当前层级的任务节点，并按照以下策略进行流分配：
若当前节点的某个前驱节点为其所在流的末尾节点，则将当前节点分配到该流中；
若有多个符合上述条件的前驱节点，则选择未完成的后驱节点最少的那一个前驱节点所在的流；
若不存在上述前驱节点，则用当前任务节点计算分配的流。
本实施例中，为跨流同步的核函数插入事件同步代码，具体为：
对于前驱节点与自身不在同一个流的所有节点，在拥有其前驱节点的每个流中，选择该流种最后一个前驱节点，在其核函数调用代码hipLaunchKernel()之后加入hipEventRecord(event,stream)，其中event是CUDA/HIP编程框架中的事件触发器，stream是CUDA/HIP编程框架下的指定任务流。其中，hipLaunchKernel()将用户编写的、在GPU进行计算的函数在程序运行时部署到GPU上开始计算。hipEventRecord(event,stream)在指定的任务流的末尾中插入一个监控器，只有当该监控器之前的任务全部执行完，该监视器的状态才会改变，从而改变依赖于该监视器状态的监视器状态，使这些监视器之后的任务得以开始运行。前述核函数的代码这里不做具体罗列。
然后对于与该前驱节点存在依赖关系且不处于同个流中的任务节点，在其调用核函数的hipLaunchKernel()语句之前加入hipStreamWaitEvent(stream,event)，以等待不同流的前驱函数执行完毕再继续执行该函数。hipStreamWaitEvent(stream,event)在指定的任务流的末尾插入一个依赖于指定监视器的监视器，当被依赖的监视器状态改变时，该监视器状态也会改变，从而启动该任务流中后续任务的执行。前述核函数的代码这里不做具体罗列。
本实施例中，生成可执行文件具体为通过LLVM将中间语言表示文件转化为可执行文件。
通过以上步骤，用户仅需要编写串行的核函数，本工作就可以通过LLVM编译器自动地识别每个核函数的只读变量与修改变量，从而构建核函数之间的依赖关系图DAG，然后使用一种启发式的调度算法将可并行的核函数调度到不同的任务流中。不同任务流可以并行地执行各自流内的任务，从而实现最大化利用GPU资源并提升程序性能。
本发明是首个通过LLVM编译器框架将串行的GPU核函数通过CUDA框架的多流机制并行化，从而提高程序性能与GPU硬件利用率。该发明不引入额外的编程模型与运行时开销，减小了开发者的学习成本与开发成本。同时该发明使用的流调度算法将存在数据依赖关系的核函数的运行时机临近安排，尽可能提高GPU对数据的访存效率，从而提高程序性能。
以上详细描述了本发明的较佳具体实施例。应当理解，本领域的普通技术人员无需创造性劳动就可以根据本发明的构思做出诸多修改和变化。因此，凡本技术领域中技术人员依本发明的构思在现有技术的基础上通过逻辑分析、推理或者有限的实验可以得到的技术方案，皆应在由权利要求书所确定的保护范围内。

Claims (6)

1.一种基于编译的核函数自动多流调度方法，其特征在于，包括以下步骤：
步骤1、读取源代码，识别核函数的输入与输出；
步骤2、根据输入与输出构建DAG；
步骤3、对DAG进行层次化，并根据拓扑排序为核函数赋予适当的流；
步骤4、为跨流同步的核函数插入事件同步代码；
步骤5、生成可执行文件。
2.如权利要求1所述的一种基于编译的核函数自动多流调度方法，其特征在于，所述读取源代码，识别核函数的输入与输出，具体为：将源代码通过LLVM编译框架转变为IR文件，然后使用opt工具对所述IR文件进行处理，分析每个核函数的只读变量与写变量。
3.如权利要求1所述的一种基于编译的核函数自动多流调度方法，其特征在于：所述根据输入与输出构建DAG，具体为：假设当前核函数有N个传入参数，对每一个传入参数将DAG的结束节点作为起点，使用倒序BFS算法找到第一个拥有同个传入参数的核函数，将该核函数添加为当前核函数的前驱节点，将当前核函数代表的节点设为DAG中结束节点的前驱节点。
4.如权利要求1所述的一种基于编译的核函数自动多流调度方法，其特征在于，所述对DAG按层次顺序将每一层次的每个核函数赋予适当的流，具体为：对于每一层的任务节点顺序赋予递增id，使得每个任务节点都有当前层级中唯一的id；先对那些没有前驱节点的任务节点计算分配的流；然后再根据前驱节点数量从小到大的顺序遍历当前层级的任务节点，并按照以下策略进行流分配：
若当前节点的某个前驱节点为其所在流的末尾节点，则将当前节点分配到该流中；
若有多个符合上述条件的前驱节点，则选择未完成的后驱节点最少的那一个前驱节点所在的流；
若不存在上述前驱节点，则用当前任务节点计算分配的流。
5.如权利要求1所述的一种基于编译的核函数自动多流调度方法，其特征在于，所述为跨流同步的核函数插入事件同步代码，具体为：
对于前驱节点与自身不在同一个流的所有节点，在拥有其前驱节点的每个流中，选择该流种最后一个前驱节点，在其核函数调用代码之后加入以CUDA/HIP编程框架中的事件触发器和CUDA/HIP编程框架下的指定任务流为参数的代码；
然后对于与该前驱节点存在依赖关系且不处于同个流中的任务节点，在其调用核函数的语句之前加入以CUDA/HIP编程框架中的事件触发器和CUDA/HIP编程框架下的指定任务流为参数的代码，以等待不同流的前驱函数执行完毕再继续执行该函数。
6.如权利要求1所述的一种基于编译的核函数自动多流调度方法，其特征在于：所述生成可执行文件具体为通过LLVM将中间语言表示文件转化为可执行文件。
CN202210172808.2A 2022-02-24 2022-02-24 一种基于编译的核函数自动多流调度方法 Pending CN114549277A (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202210172808.2A CN114549277A (zh) 2022-02-24 2022-02-24 一种基于编译的核函数自动多流调度方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202210172808.2A CN114549277A (zh) 2022-02-24 2022-02-24 一种基于编译的核函数自动多流调度方法

Publications (1)

Publication Number Publication Date
CN114549277A true CN114549277A (zh) 2022-05-27

Family

ID=81678119

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202210172808.2A Pending CN114549277A (zh) 2022-02-24 2022-02-24 一种基于编译的核函数自动多流调度方法

Country Status (1)

Country Link
CN (1) CN114549277A (zh)

Similar Documents

Publication Publication Date Title
JP4042604B2 (ja) プログラム並列化装置，プログラム並列化方法およびプログラム並列化プログラム
JP6141365B2 (ja) 逐次コンピュータプログラムコードを並列処理する方法及びシステム
CN111738434A (zh) 在异构处理单元上执行深度神经网络的方法
Lin et al. Efficient GPU computation using task graph parallelism
Sioutas et al. Schedule synthesis for halide pipelines on gpus
Moren et al. Automatic mapping for OpenCL-programs on CPU/GPU heterogeneous platforms
de Andrade et al. Software deployment on heterogeneous platforms: A systematic mapping study
Doroshenko et al. A mixed method of parallel software auto-tuning using statistical modeling and machine learning
CN116861359A (zh) 面向深度学习推理任务编译器的算子融合方法和系统
CN114549277A (zh) 一种基于编译的核函数自动多流调度方法
US9003383B2 (en) Analytic engine to parallelize serial code
Delestrac et al. Demystifying the TensorFlow Eager Execution of Deep Learning Inference on a CPU-GPU Tandem
CN114127681A (zh) 用于实现数据流ai应用的自主加速的方法和装置
Li et al. An Adaptive Thread Partitioning Approach in Speculative Multithreading
Czejdo et al. Practical Approach to Introducing Parallelism in Sequential Programs
Han et al. Genetic algorithm based parallelization planning for legacy real-time embedded programs
Grelck et al. Engineering concurrent software guided by statistical performance analysis
Ferrandi et al. Automatic parallelization of sequential specifications for symmetric mpsocs
CN115469879A (zh) 一种基于多面体模型的自动调度生成方法
Himstedt et al. D5. 1 documentation of recommendations
Acosta et al. Xfeatur: Hardware feature extraction for dnn auto-tuning
Daszczuk et al. Adding parallelism to sequential programs–a combined method
Süß et al. And now for something completely different: running Lisp on GPUs
Nicolau A development environment for scientific parallel programs
Aiken et al. A development environment for horizontal microcode programs

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