减少GPU和CPU交互开销的方法及系统技术方案

技术编号:37711530 阅读:9 留言:0更新日期:2023-06-02 00:04
本发明专利技术公开了一种减少GPU和CPU交互开销的方法及系统,该方法包括:根据DNN模型中经过拓扑排序后的所有算子的类型及其参数信息,从最后一个算子开始逐个生成kernel算法代码,同时在每个已经生成的kernel算法代码的末尾,根据下一个要运行的算子的kernel算法代码的程序信息,插入串联函数;调用运行时库的在线编译器编译刚生成的算子的kernel算法代码,并把编译生成的二进制程序加载到GPU中作为输入数据准备运行;根据输入数据,在GPU上通过运行函数从第一个算子开始迭代运行各个算子,直至监听到运行结束事件发生时结束运行并将输出数据拷贝至CPU。本发明专利技术的技术方案能够消除每个算子的执行都需要GPU与CPU进行交互产生的交互开销,节约了计算资源,提升了模型运行的速度。度。度。

【技术实现步骤摘要】
减少GPU和CPU交互开销的方法及系统


[0001]本专利技术涉及神经网络
,尤其涉及一种减少GPU和CPU交互开销的方法及系统。

技术介绍

[0002]GPU(Graphics Processing Unit,图形处理器)当今已经成为DNN(Deep Neural Network,深度神经网络)加速运行的主流设备,无论是用于推理还是训练。DNN模型是一个由许多算子(operator,以下简称op)组成的DAG(Directed Acyclic Graph,有向无环图),对于GPU设备而言,每个op实现为一个在GPU上执行的kernel(GPU中的内核函数),在DNN完整一趟运行过程中,CPU按照拓扑排序好的op顺序逐个传入kernel参数并调度其在GPU上运行。该运行方式存在的一个最大的问题就是每次op的运行都需要CPU传递kernel参数并启动其在GPU执行并等待其运行结束,产生了大量的性能开销,同时在该段时间间隙内GPU的可执行单元(ExecutionUnit)处于空转状态,带来了不小的计算资源浪费。
[0003]如图1所示,当前在GPU设备上运行的DNN网络方式存在以下问题:
[0004]其一,在DNN网络中每个算子的Kernel算法在运行前都要经历以下步骤:

在线构建Kernel算法代码(即build kernel),该步骤可通过cache(缓存)结果进行优化;

通过参数设置函数clSetKernelArg设置Kernel的运行参数;

最后调用运行函数clEnqueueNDRangeKernel启动运行Kernel算法。
[0005]其二,这些重复的Kernel算法运行前的准备步骤给CPU带来了不少的性能开销以及CPU与GPU之间产生的交互开销。

技术实现思路

[0006]本专利技术的目的之一是为了克服现有技术中的不足,针对现有技术中存在的CPU在传递、调度kernel算法的参数时GPU空转及GPU执行kernel参数时CPU空转会导致CPU和GPU的性能开销和计算资源浪费的问题,提供一种减少GPU和CPU交互开销的方法及系统。
[0007]为实现以上目的,本专利技术通过以下技术方案实现:
[0008]第一方面,本专利技术提供了一种减少GPU和CPU交互开销的方法,所述方法包括:
[0009]S100:根据DNN模型中经过拓扑排序后的所有算子的类型及其参数信息,从最后一个算子开始,根据算子类型对应的代码模板和算子的参数信息按照倒序顺序,逐个生成各个算子对应的kernel算法代码,同时在每个已经生成的kernel算法代码的末尾,根据下一个要运行的算子的kernel算法代码的程序信息,插入用于在当前算子运行完毕后启动下一个要运行的算子的串联函数;
[0010]S200:调用运行时库的在线编译器编译刚生成的算子的kernel算法代码,并把编译生成的二进制程序加载到GPU中作为输入数据准备运行;
[0011]S300:根据所述输入数据,在GPU上通过运行函数从DNN模型中的第一个算子开始迭代运行各个算子,直至监听到运行结束事件发生时结束运行并将输出数据拷贝至CPU。
[0012]在本申请的一个优选实施例中,S100具体包括:
[0013]S101:对DNN模型中的各个算子进行拓扑排序,获取各个算子的类型及其参数信息;
[0014]S102:将kernel算法代码的生成顺序设置为倒序,将准备生成各个算子对应的kernel算法代码的起始位置i设置为最后一个算子的位置n;
[0015]S103:根据第i个算子的类型对应的代码模板和算子的参数信息,构建第i个算子对应的算法模型;
[0016]S104:根据所述第i个算子对应的算法模型和参数信息,生成所述第i个算子对应的kernel算法代码;
[0017]S105:在所述第i个算子对应的kernel算法代码的末尾,根据第i+1个算子的kernel算法代码的程序信息,插入用于在第i个算子运行完毕后启动要运行第i+1个算子的串联函数;
[0018]S106:判断i是否为1;若i=1,则执行S200;若i≠1,则执行S107;
[0019]S107:令i=i

1,返回执行S103。
[0020]在本申请的一个优选实施例中,所述算子信息包括各个算子的属性信息、输入输出张量形状及分配给张量用于存储数据的内存地址以及后继算子对应的kernel算法代码的入口地址信息;
[0021]所述输出张量的形状根据所述属性信息和输入张量的形状推算得出;
[0022]若待生成kernel算法代码的算子不是最后一个算子,则获取位于该算子之后的后继算子对应的kernel算法代码编译后得到的可执行程序的入口地址信息。
[0023]在本申请的一个优选实施例中,各个算子的属性信息、输入输出张量形状及分配给张量用于存储数据的内存地址以及后继算子对应的kernel算法代码的入口地址信息通过代码生成函数生成;各个算子的参数信息和算法模型通过编译函数构建对应的kernel算法代码。
[0024]在本申请的一个优选实施例中,S300具体包括:
[0025]S301:从CPU拷贝输入数据;
[0026]S302:通过运行函数在GPU上启动运行第一个算子;
[0027]S303:在运行过程中通过CPU异步监听从GPU发出的运行结束事件;
[0028]S304:在CPU监听到运行结束事件后,GPU将当前计算结果作为输出数据拷贝到CPU的内存并做后续处理;
[0029]重复执行S301至S304。
[0030]第二方面,本专利技术提供了一种减少GPU和CPU交互开销的系统,所述系统包括串联模块、编译模块和运行模块;
[0031]所述串联模块用于根据DNN模型中经过拓扑排序后的所有算子的类型及其参数信息,从最后一个算子开始,根据算子类型对应的代码模板和算子的参数信息按照倒序顺序,逐个生成各个算子对应的kernel算法代码,同时在每个已经生成的kernel算法代码的末尾,根据下一个要运行的算子的kernel算法代码的程序信息,插入用于在当前算子运行完毕后启动下一个要运行的算子的串联函数;
[0032]所述编译模块用于调用运行时库的在线编译器编译刚生成的算子的kernel算法
代码,并把编译生成的二进制程序加载到GPU中作为输入数据准备运行;
[0033]所述运行模块用于根据所述输入数据,在GPU上通过运行函数从DNN模型中的第一个算子开始重复迭代运行各个算子,直至监听到运行结束事件发生时结束运行并将输出数据拷贝至CPU。
[0034]第三方面,本专利技术提供了一种计算机可读存储介质,所述计算机可读存储介质中存储有计算机程序,当其在计算机上运行时,使得计算机执行如第一方面所述的减少GPU和CPU交互开销的方法。
[0035]第四方面本文档来自技高网
...

【技术保护点】

【技术特征摘要】
1.一种减少GPU和CPU交互开销的方法,其特征在于,所述方法包括:S100:根据DNN模型中经过拓扑排序后的所有算子的类型及其参数信息,从最后一个算子开始,根据算子类型对应的代码模板和算子的参数信息按照倒序顺序,逐个生成各个算子对应的kernel算法代码,同时在每个已经生成的kernel算法代码的末尾,根据下一个要运行的算子的kernel算法代码的程序信息,插入用于在当前算子运行完毕后启动下一个要运行的算子的串联函数;S200:调用运行时库的在线编译器编译刚生成的算子的kernel算法代码,并把编译生成的二进制程序加载到GPU中作为输入数据准备运行;S300:根据所述输入数据,在GPU上通过运行函数从DNN模型中的第一个算子开始迭代运行各个算子,直至监听到运行结束事件发生时结束运行并将输出数据拷贝至CPU。2.根据权利要求1所述的减少GPU和CPU交互开销的方法,其特征在于,S100具体包括:S101:对DNN模型中的各个算子进行拓扑排序,获取各个算子的类型及其参数信息;S102:将kernel算法代码的生成顺序设置为倒序,将准备生成各个算子对应的kernel算法代码的起始位置i设置为最后一个算子的位置n;S103:根据第i个算子的类型对应的代码模板和算子的参数信息,构建第i个算子对应的算法模型;S104:根据所述第i个算子对应的算法模型和参数信息,生成所述第i个算子对应的kernel算法代码;S105:在所述第i个算子对应的kernel算法代码的末尾,根据第i+1个算子的kernel算法代码的程序信息,插入用于在第i个算子运行完毕后启动要运行第i+1个算子的串联函数;S106:判断i是否为1;若i=1,则执行S200;若i≠1,则执行S107;S107:令i=i

1,返回执行S103。3.根据权利要求1所述的减少GPU和CPU交互开销的方法,其特征在于,所述算子信息包括各个算子的属性信息、输入输出张量形状及分配给张量用于存储数据的内存地址以及后继算子对应的kernel算法代码的入口地址信息;所述输出张量的形状根据所述属性信息和输入张量的形状推算得出;若待生成kernel算法代码的算子不是最后一个算子,则获取位于该算子之...

【专利技术属性】
技术研发人员:易祚超
申请(专利权)人:格兰菲智能科技有限公司
类型:发明
国别省市:

网友询问留言 已有0条评论
  • 还没有人留言评论。发表了对其他浏览者有用的留言会获得科技券。

1