CN102708009A - 一种基于cuda实现多任务共享gpu的方法 - Google Patents

Abstract

本发明公布了一种基于CUDA实现多任务共享GPU的方法。包括：在Global Memory中建立映射表，确定合并后的Kernel中，每个Block执行的任务编号和任务分块编号；一次用一个Kernel启动N个Block，N等于所有任务的任务分块数目之和；用标记和阻塞等待的方法，满足原有任务之间的约束关系；对于Shared Memory，采用预申请和静态分配的方式进行多任务共享。通过本发明，可以简便地实现在现有GPU硬件架构上实现多任务共享，可以简化实际应用中的编程工作，并在一定情况下取得良好的性能。

Description

一种基于CUDA实现多任务共享GPU的方法

技术领域

本发明涉及一种多任务共享GPU的实现方法，具体涉及在NVIDA的CUDA架构中合并多个任务，实现任务并行的方法，属于GPGPU计算领域。

背景技术

GPGPU(General-purpose computing on graphics processing units)，是利用GPU来进行大规模计算的技术。CUDA是NVIDA公司提供的GPGPU架构。CUDA自从推出开始，就成为广泛应用的众核并行计算形式。

GPU具有远远高于CPU的浮点运算能力和内存带宽(附图1)，同时由于其高度的并行性，非常适合于大规模数据处理。

然而，由于GPU的硬件设计，GPU上的编程和CPU上的并行编程有所不同。一个显著的区别就是，GPU不支持多任务共享：每个任务在GPU上的运行都是对GPU的硬件资源独占的，不允许有其他的Kernel也在执行。例如，当多个任务都要使用GPU的时候，只能够一个一个的顺序地执行，而不能一起同时在GPU上运行。这一点和CPU上允许进程之间进行切换是有很大的不同。

目前，尚未发现有专利或者文献针对GPU上的多任务共享进行讨论。

发明内容

本发明所使用的一些术语定义如下：

Kernel：CUDA架构中，GPU一次运行所执行的代码。

Thread，Block，Grid：CUDA架构中，一个大的Kernel被划分为了很多小的基本单位，称为线程(Thread)。

所有Thread组织成为了两级结构。(附图2)

首先，这些Thread划分为了若干个线程块(Block)，每个Block包含相同数目的Thread。Thread是以Block为单位分发到硬件资源上来执行的。在Block内部，线程的ID编号可以采用一维、二维或者三维，这称为Block维数。在某一个确定的Kernel里面，每个Block都含有相同的维数，以及每一维上相同的大小(这称为维度)。维数和维度都是在GPU启动Kernel之前由程序员指定好的，执行过程中是不能改变的。

然后，所有的Block组成的整体(也就是所有的Thread)称为线程网格(Grid)。和Block的维数定义类似，Block的ID编号，可以组织成一维、二维或者三维，称为Grid的维数，每一维上的大小称为维度。

SM(Streaming Multiprocessor)：GPU上的硬件单元，包括运算单元、寄存器、片上存储器。每个Block都会被分配到SM上进行执行。一个Block只能在一个SM上执行，一个SM上可以同时执行一个或者多个Block(取决于单个Block消耗的寄存器、片上存储器资源的数目)。

Global Memory：显卡上的存储器，属于片外存储器，GPU可以从中读取，但是速度相对比较慢。

Shared Memory：GPU上的一种片上存储器，每个SM中拥有一定数量的Shared Memory，通常为KB级大小，可以视作一种可供程序员维护的Cache，在CUDA架构中有非常重要的意义。

任务分块：单个任务划分为多个子任务，每个子任务称为任务分块。一个任务分块将会对应到一个Block中完成计算。

本发明旨在提供一种方法使得在CUDA架构上实现多任务共享GPU，解决现有GT200架构不支持多任务共享GPU的问题。

本发明的原理包括三点：

1.对每个Block执行的任务分块，在GPU运行之初由程序员通过对一个映射表的赋值来确定。通常，每个Block执行的任务分块，在算法设计的时候就已经固定。本发明通过增加一个映射表，使得能够在GPU启动之前，程序员可以对每个Block执行的任务分块进行重排。

2.利用GPU高并行度的特点，把多个任务的Block合并到的一个Kernel中去。

CUDA架构中，可以同时启动成百上千的线程(Thread)，进行并行度非常高的计算任务。所有的Thread被按照一定数目组织成为若干的Block。本发明让一个GPU在一个Kernel中同时启动大量的Block，执行所有任务的任务分块。

3.考虑到任务之间可能具有约束性，利用GPU动态调度Block到SM上的特性，处理约束关系。

CPU上一种通常的做法是，如果有约束关系，那么在需要同步的地方，直接使用原子操作即可。虽然GPU也支持原子操作，但是原子操作的代价是非常大的，会带来性能上的严重损失，并且极易造成死锁。

本发明充分考虑到了GPU调度Block的规律性。

GPU通常拥有数十到上百个SM。GPU上的Block在启动的时候并不会拥有硬件资源，而是在执行的过程中，动态地把Block分发给SM。但SM的数量一般相对于Block少很多，所以每个时刻，在SM上执行的Block只是一部分，其他的Block只能等待。一旦有SM上Block的任务计算结束，GPU回收得到空闲资源的时候，就会从尚未执行的Block中选取一定的Block分发到有空闲资源的SM上去。

为Block分配资源的时候，是有一定的顺序性的，这种顺序性表现为：

1.总是优先分发ID编号较小的Block到SM上去。例如ID编号为0的Block分发到SM上去的时间总是不晚于ID编号为1的Block被分发的时间。

2.相对的，回收Block资源的时候，也是优先回收ID编号较小的Block的空闲资源。例如ID编号为0的Block尚未结束，ID编号为10的Block执行结束，但此时GPU不会回收编号为10的Block的资源，因为ID编号更小，ID编号为0的Block的资源并未被回收。

根据这种顺序性，本发明先把需要优先执行的任务，在映射表中适当排序：

1.被其他任务依赖的任务，使ID编号较小的Block来执行它的任务分块，这样它会优先获得资源，被调度到SM上去进行先执行；

2.依赖于其他任务的任务，使ID编号较大的Block来执行它的任务分块，同时，辅以适当的阻塞等待操作，保证其依赖的任务已经完全执行结束。

完整的技术方案如下(流程参见图4)：

一种基于CUDA实现多任务共享GPU的方法，包括如下步骤：

1)在Global Memory中建立映射表，确定合并后的Kernel中，每个Block执行的任务编号和任务分块编号；

2)一次用一个Kernel启动N个Block，N等于所有任务的任务分块数目之和；

3)用标记和阻塞等待的方法，满足原有任务之间的约束关系；

4)对于Shared Memory，采用预申请和静态分配的方式进行多任务共享。

其中，步骤1)的优选实现方法如下：

1.1)映射表要给出Block到任务和任务分块的映射关系，即确定某个Block执行哪个任务的哪个分块；

1.2)任务在映射表中的排布要求满足约束条件的拓扑顺序：如果任务A依赖于任务B，那么，执行任务A的所有Block的ID编号应该大于所有执行任务B的Block的ID编号；

1.3)在满足步骤1.2)所述的约束条件的情况下，其他无约束关系任务在映射表中以任意的方式进行排布。

步骤2)的优选实现方法如下：

2.1)将原有的任务的grid维数转换为一维：原有的任务可能是有不同的grid维数和维度，这里可以统一选取一维；如果原来的任务是多维的，只需要进行一维到多维的换算即可；

2.2)若Block的维数不一致，则将之统一转换为一维；若转成一维之后Block的维度不一致，则统一选取一个最大的维度，让其他较小的Block添加空线程补足，这样所有的任务都采用了相同的Block维数和维度；

2.3)所有的Block在开始执行计算之前，首先从映射表中获取该Block需要执行的任务编号和任务分块编号；

2.4)根据步骤2.3)中读取的任务编号，选择执行不同的任务代码；将步骤2.3)中读取的任务分块编号，使用到具体的任务计算中去。

步骤3)的优选实现方法如下：

3.1)给每个任务设置标志位mark，每个任务的每个任务分块设置标记数组tag[]，用以标记任务和任务分块的执行情况；

3.2)对于任务的每一个任务分块，在该任务分块完成之后，将对应的标记位tag置位，表明该任务分块已经执行结束；

3.3)用每个任务的最后一个Block，在计算返回前阻塞，循环检查同一任务的其他任务分块的对应tag标记位是否已经被置位，一旦全部被置位，则对标志位mark置位，表明该任务已经结束；

3.4)如果任务A需要依赖于任务B，那么在任务A计算开始之前阻塞，循环检测B的任务标志位mark，直到其置位。

步骤4)的优选实现方法如下：

4.1)预先在Kernel开始的时候，申请一个足够大的Shared Memory数组，其大小至少应该等于每个任务所需要的Shared Memory用量的最大值；

4.2)每个任务单独写成一个函数，将Shared Memory数组地址传给这个函数，函数中需要使用Shared Memory的时候，直接在这个数组中静态分配使用。

通过本发明，可以简便地实现在现有GPU硬件架构上实现多任务共享，可以简化实际应用中的编程工作，并在一定情况下取得良好的性能。

附图说明

图1GPU的浮点运算能力存储器带宽同CPU的比较(图片来源：NVIDIA CUDAProgramming Guide Version 2.3)

图2GPU中的Thread，Block，Grid结构(图片来源：NVIDIA CUDA Programming GuideVersion 2.3)

图3两种映射表的排布。图3(a)实施方案中的3个任务使用图形表示(约束关系用箭头表示，任务1依赖于任务0)；图3(b)一种合法的映射表排布；图3(c)一种不合法的映射表排布(任务1依赖于任务0，却有Block排布到了任务0前面)。

图4本发明所述方法的流程图。

具体实施方式

以下以一个具体的例子，对本发明做进一步的说明。但是需要注意的是，公布实施例的目的在于帮助进一步理解本发明，但是本领域的技术人员可以理解：在不脱离本发明及所附的权利要求的精神和范围内，各种替换和修改都是可能的。因此，本发明不应局限于实施例所公开的内容，本发明要求保护的范围以权利要求书界定的范围为准。

具体的例子是：3个计算任务(具体任务内容此处并无影响)。

任务存在以下的约束关系：任务1必须在任务0完成后才能进行，因为任务1需要使用任务0的结果，而任务2同任务0和任务1没有任何约束关系。(附图3(a)，圆圈代表任务，箭头代表依赖关系)

使用技术方案中所述的方法实现任务并行。

为了方便叙述，定义以下device函数分别完成3个计算任务，分别称为任务0，1和2。

任务0	_device_void compute0(...)；
		任务1	_device_void compute1(...)；
任务2	_device_void compute2(...)；

同时假设每个任务都进行了相同大小的任务分块，各自的计算任务划分为了N个任务分块。

实施过程分为以下步骤：

A.建立映射表

开辟两个一维的数组，长度等于所有任务的任务分块数之和，此例中为3*N。两个数组的具体含义如下

1.task_id[]，取值为0、1或者2。这个数组给出了Kernel中的Block需要执行的任务。例如，图3(b)，task_id[0]到task_id[N-1]的值都为0，代表了Block ID从0到N-1的Block都要执行任务0；其他值为1和2的元素，意义类似。

2.block_id[]，取值为0到N-1。这个数组给出了Kernel中Block需要执行的分块编号。例如，图3(b)，block_id[N]等于0，代表了编号为N的Block需要执行编号为0的任务分块(任务编号由1中的task_id指定)。

编号的顺序需要满足具体问题所要求的任务之间的约束关系。任务1依赖于任务0，所以执行任务1的分块的那些Block，ID编号应该大于执行任务0的Block ID编号。

例如简单地做法，把任务0和任务2的block排布到前端，任务1的block排布到后端，如图3(b)所示。

而图3(c)的排布是不满足约束关系的拓扑序的，任务1的分块被排布到了任务0前面，在GPU的调度过程中会首先执行，造成错误，所以是是不合法的。

确定好顺序之后，可以在Kernel启动之前直接对task_id[]和block_id[]进行赋值即可。

B.启动Kernel

定义一个合并后的Kernel，Block的数目等于原有所有任务的任务分块数之和，此例中为3*N。这个Kernel的参数列表应该传入三个任务需要的所有参数。

合并的Kernel应该先获取自身的Block ID号bidx

int bidx＝blockIdx.x+gridDim.x*blockIdx.y；

利用bidx从task_id和block_id数组中第bidx个位置读取任务和分块编号信息taskid和bid。

int bid＝block_id[bidx]，taskid＝task_id[bidx]；

然后根据taskid相应地选择执行不同的device函数，device函数需要做简单修改，把bid作为参数传给device函数使用。

如下面代码所示。

C.处理约束

Kernel启动前为每个任务在Global Memory中各自开辟一个标记数组tag0[]、tag1[]和tag2[]，每个元素取值0或者1，用以标记每个任务的Block是否执行结束。例如tag0[]的第10个元素为0，表明任务0的第10个Block尚未执行结束。

此外，为每个任务在Global Memory中各自设置一个标志变量mark0、mark1、mark2，取值0或者1，以标记任务的执行状态。例如，mark2＝0代表任务2已经执行结束。

tag0[]、tag1[]和tag2[]的所有元素，以及mark0、mark1和mark2，一开始都被设置为0。

某个任务的编号为bid任务分块执行结束，置对应的标记数组第bid个位置1。例如，任务0的第10个任务分块执行结束，就把tag0[]的第10个元素设置为1。

每个任务的最后一个Block(执行此任务的ID编号最大的Block)，在退出前，检查当前任务的其他任务分块的标记，查看是否都已经变成1，如果是，则对应任务的标志位为1，否则循环等待所有标记都变为1。

例如，对于任务0，任务0的最后一个Block，也就是bid＝N-1的Block在调用了device函数返回之后，检查任务0其他任务分块，循环阻塞，等待tag0[]全部变为1。

在需要满足约束的地方，对对应的mark进行阻塞检查。例如，任务1依赖于任务0，所以在任务1的device函数之前，需要满足约束。于是增加一个循环等待语句，阻塞等待任务0的标志位mark0变为1。

D.共享Shared Memory

预先在Kernel开始的地方，申请出一块足够大小的Shared Memory数组shared，大小应该是三个device函数使用的Shared Memory的最大值。

_shared_float shared[MAX]；

修改device函数，把shared传给device函数以作为Shared Memory使用。Device函数中所有用到Shared Memory的地方，都从传入的shared数组上，进行静态分配。

Claims (5)

1.一种基于CUDA实现多任务共享GPU的方法，包括如下步骤：

1)在Global Memory中建立映射表，确定合并后的Kernel中，每个Block执行的任务编号和任务分块编号；

2)一次用一个Kernel启动N个Block，N等于所有任务的任务分块数目之和；

3)用标记和阻塞等待的方法，满足原有任务之间的约束关系；

4)对于Shared Memory，采用预申请和静态分配的方式进行多任务共享。

2.如权利要求1所述的方法，其特征是，步骤1)的实现方法如下：

1.1)映射表要给出Block到任务和任务分块的映射关系，即确定某个Block执行哪个任务的哪个分块；

1.2)任务在映射表中的排布要求满足约束条件的拓扑顺序：如果任务A依赖于任务B，那么，执行任务A的所有Block的ID编号应该大于所有执行任务B的Block的ID编号；

1.3)在满足步骤1.2)所述的约束条件的情况下，其他无约束关系任务在映射表中以任意的方式进行排布。

3.如权利要求1所述的方法，其特征是，步骤2)的实现方法如下：

2.1)将原有的任务的grid维数转换为一维；

2.2)若Block的维数不一致，则将之统一转换为一维；若转成一维之后Block的维度不一致，则统一选取一个最大的维度，让其他较小的Block添加空线程补足，这样所有的任务都采用了相同的Block维数和维度；

2.3)所有的Block在开始执行计算之前，首先从映射表中获取该Block需要执行的任务编号和任务分块编号；

2.4)根据步骤2.3)中读取的任务编号，选择执行不同的任务代码；将步骤2.3)中读取的任务分块编号，使用到具体的任务计算中去。

4.如权利要求1所述的方法，其特征是，步骤3)的实现方法如下：

3.1)给每个任务设置标志位mark，每个任务的每个任务分块设置标记数组tag[]，用以标记任务和任务分块的执行情况；

3.2)对于任务的每一个任务分块，在该任务分块完成之后，将对应的标记位tag置位，表明该任务分块已经执行结束；

3.3)用每个任务的最后一个Block，在计算返回前阻塞，循环检查同一任务的其他任务分块的对应tag标记位是否已经被置位，一旦全部被置位，则对标志位mark置位，表明该任务已经结束；

3.4)如果任务A需要依赖于任务B，那么在任务A计算开始之前阻塞，循环检测B的任务标志位mark，直到其置位。

5.如权利要求1所述的方法，其特征是，步骤4)的实现方法如下：

4.1)预先在Kernel开始的时候，申请一个足够大的Shared Memory数组，其大小至少应该等于每个任务所需要的Shared Memory用量的最大值；

4.2)每个任务单独写成一个函数，将Shared Memory数组地址传给这个函数，函数中需要使用Shared Memory的时候，直接在这个数组中静态分配使用。

Images