具体实施方式
现在将详细参照本申请的具体实施例,在附图中例示了本申请的例子。尽管将结合具体实施例描述本申请,但将理解,不是想要将本申请限于描述的实施例。相反,想要覆盖由所附权利要求限定的在本申请的精神和范围内包括的变更、修改和等价物。应注意,这里描述的方法步骤都可以由任何功能块或功能布置来实现,且任何功能块或功能布置可被实现为物理实体或逻辑实体、或者两者的组合。
CUDA是Nvidia为GPU、Intel x86设备异构计算编程而设计的一套API,它可以和C、C++和Fortran编程语言一起工作。Nvidia GPU硬件上的计算单元被分为以下三个层次:核(core):GPU硬件的最小计算单元,真正执行一个线程的设备;流多处理器(StreamingMultiprocessors,SMs):执行一个线程块的计算单元,里面有多个核;线程束(warp):硬件资源调度的单位。线程束(warp)是GPU硬件资源调度的最小单元,一个线程束对应32个线程,同一个线程束内的线程并行地执行同一条指令(SIMD模型),当SMs执行一个线程块时不会直接地一次性为线程块中的所有线程分配全部资源,取而代之的是它会将线程块中的线程先拆分成为线程束,然后以线程束为单位调度硬件资源。
CUDA执行线程间并发是通过流(Stream),其是一个任务队列,单个CUDA 流内部的任务按顺序保序执行,即根据发送过来的任务依次执行;不同CUDA 流的任务的执行顺序不保序,即,不要求按顺序执行,可以同时执行,也可以不同时执行。
通过以下指令来举例说明CUDA流内部保序的特性。
Kernel_0<<<100, 512, 0, stream>>>(args...)
Kernel_1<<<200, 1024, 0, stream>>>(args...)
这个例子中,线程Kernel_0和Kernel_1在主机端先后被启动,并且先后被装到了同一个CUDA流(都是stream)里面,这种情况下在设备计算这两个CUDA Kernel的时候也会按照主机端装入CUDA流的顺序先执行Kernel0,然后执行Kernel1,也就是CUDA流内部保序。
通过以下指令来举例说明CUDA流之间不保序的特性。
Kernel_0<<<100, 512, 0, stream1>>>(args...)
Kernel_1<<<200, 1024, 0, stream2>>>(args...)
这个例子与前一个例子的区别是线程Kernel0和Kernel_1被装进了不同的CUDA流(stream1和stream2),这种情况下虽然在主机端是先启动Kernel_0,然后启动Kernel_1,但是在设备端执行的时候却无法保证这两个CUDA Kernel的执行顺序,也就是CUDA流之间不保序,可能同时执行,也可能不同时执行。
然而,即使线程在CUDA流之间不保序的情况下,不保序执行的各个线程内部的代码的各个代码段仍然是顺序执行的,即从前到后串行执行所有代码段。在线程内顺序排列的代码段之间的读取数据和执行指令之间依然由于按顺序执行而存在等待。例如Kernel0中包括如下多个代码段:
#include <mma.h>
using namespace nvcuda;
__global__ void wmma_ker (half *a, half *b, float *c) {
// Declare the fragments
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float) c_frag;
// Initialize the output to zero
wmma::fill_fragment (c_frag, 0.0f);
// Load the inputs
wmma::load_matrix_sync (a_frag, a, 16);
wmma::load_matrix_sync (b_frag, b, 16);
// Perform the matrix multiplication
wmma::mma_sync (c_frag, a_trag, b_frag, c_frag);
// Store the output
wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}
仍然是从最开始的wmma::fragment代码段顺序地运行到wmma::fill_fragment代码段,再到wmma::load_matrix_sync代码段,以此类推,直到最后一个代码段。
然而,在线程内部的各个代码段之间并没有并行执行的机制,仍然存在代码段之间需要相互等待、以便顺序执行而产生的延时。本申请要解决在线程内部的各个代码段之间并行执行的问题。
具体地,本申请在源代码中确定能够并行执行的各个代码段,通过对编译器前端的解析器模块和代码生成器模块进行修改,将该各个代码段翻译为在GPU上并行执行的代码段结构,该结构包括3个不同的基本块,来帮助或控制翻译后的代码段之间的并行执行。
图1示出了根据本申请的实施例的一种代码编译器的示意方框图。
如图1所示,代码编译器100包括:解析器101,被配置为基于标记代码段的标记符识别多个代码段,多个代码段中的至少一部分并行执行;代码生成器102,被配置为将识别的多个代码段中的N个代码段翻译为第一中间表达式组,其中第一中间表达式组至少包括控制并行/串行执行或不执行N个代码段的第一控制块,其中N是1或大于1的正整数。
通常将编译器分为前端和后端。其中,前端编译器会对所输入的程序进行词法分析、语法分析、语义分析,然后生成中间表达形式,也就是中间表达式( intermediaterepresentation,IR)。后端编译器会对中间表达式进行优化,然后生成目标代码。这里,中间表达式是编译器对于源程序进行扫描后生成的内部表示,代表源程序的语义和语法结构。执行块的中间表达式会保留代码段的语义和语法结构,即执行块是用来执行(这些)代码段的真实动作的。执行块本身的语法可以与原始的代码段不同,只要该执行块的中间表达式是根据代码段的抽象语法树的复合表达式节点来生成的,且能够执行代码段的真实动作即可。在此,不能被直接执行的目标语言,都可以视为广义的中间表达式,不同编译器的中间表达式IR可以是不一样的,在此不做限制。
中间表达式抛弃了所有的变量名称和函数名称,使用标号以及变量以及临时变量(temp_newtemp)来代替。
以下介绍中间表达式的一个示例:
//
int main()
{
double a=128;
}
//
它生成的中间表达式是这样的:
//
; ModuleID = ‘try5.c’
target datalayout = “e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:64-f32:32:32-f64:32:64-v64:64:64-v128:128:128-a0:0:64-f80:32:32-n8:16:32-S128”
target triple = “i386-pc-linux-gnu”
define i32 @main() nounwind {
entry:
%retval = alloca i32, align 4
%a = alloca double, align 8
store i32 0, i32* %retval
store double 1.280000e+02, double* %a, align 8
%0 = load i32* %retval
ret i32 %0
}
//
%a = alloca double, align 8
当然上述中间表达式仅是示例,根据不同的代码可以生成不同的中间表达式。
如此,通过将N个代码段翻译为具有控制其执行的第一控制块的第一中间表达式组,即可以作为单个执行主体,与(也是翻译为具有控制块的中间表达式组的)其他执行主体没有顺序执行的关系,即可以与其他执行主体同时并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。
在此,结合图2来描述具体例子。图2示出了根据本申请的实施例的代码编译器的操作例子。
在GPU Kernel处,至少存在一部分能够并行执行的多个代码段,例如load_matrix_sync (a_frag, a, 16)、load_matrix_sync (b_frag, b, 16)、mma_matrix_sync(c_frag, a_frag, b_frag, c_frag)、store_matrix_sync (c, c_frag, 16, wmma::mem_row_major)等等。可以对这些能够并行执行的代码段进行标记,例如如图2所示,用coroutine标记符来标记这些可以并行执行的代码段。当然,coroutine标记符仅是示例,但本申请并不限于此,也可以是async等等,只要能够标记这些代码段可以被并行执行即可。判断代码段是否可以被并行执行可以人工判断,也可以自动判断。在判断时,可以考虑各个代码段之间是否有相互依赖的参数、是否有必须先后执行的运算、是否有因果关系等等。对代码段标上标记符可以用于后续识别多个代码段。这里的代码段是可以是一个或多个函数。或者,当一个函数的各个部分可以并行执行时,代码段也可以是该函数的一部分。
标记完的能并行执行的代码段例如是:coroutine{load_matrix_sync (a_frag,a, 16)}、coroutine{load_matrix_sync (b_frag, b, 16)}、coroutine{mma_matrix_sync(c_frag, a_frag, b_frag, c_frag)}、coroutine{store_matrix_sync (c, c_frag, 16,wmma::mem_row_major)}等等。标记符coroutine标记了能够被并行执行的多个代码段。当然,该标记的形式也仅是示例,实际上,其他标记形式也是可行的,例如利用映射表或修改函数名等,只要能让其他装置知道这些代码段能够并行执行即可。
解析器101可以基于标记代码段的标记符识别多个代码段,多个代码段中的至少一部分并行执行。代码生成器102,被配置为将识别的多个代码段中的N个代码段翻译为第一中间表达式组,其中第一中间表达式组至少包括控制并行/串行执行或不执行N个代码段的第一控制块,其中N是1或大于1的正整数。
具体地,解析器101可以在基于标记代码段的标记符识别多个代码段后,将识别的多个代码段翻译为抽象语法树(AST)中的至少一个复合表达式节点。
抽象语法树(abstract syntax tree或者缩写为AST),或者语法树(syntaxtree),是源代码的抽象语法结构的树状表现形式。抽象语法树是由名为节点(Node)的数据结构构成的,树上的每一个节点都表示表示源代码中的一种结构,例如语句、表达式以及变量等。抽象语法树并不会表示出真实语法出现的每个细节。生成语法树的过程就是生成代码段的语句、表达式以及变量等的复合表达式节点。
具体地,代码生成器102可以根据至少一个复合表达式节点将多个代码段中的N个代码段翻译为第一中间表达式组。N是1或大于1的正整数。这里,根据抽象语法树的复合表达式节点来翻译为第一中间表达组可以保留代码段的所有语法信息,从而完整地执行代码段的动作。
这里从多个代码段中选择N个代码段作为一个组来进行翻译。这里的选择可以是有目的的选择,例如选择为同一任务服务的N个代码段,以便在执行这N个代码段时,能够尽快执行完同一任务。这里的选择也可以是随机的,即从多个代码段中随机选其中N个代码段来作为一个组。
这里,N=1就表示对于1个代码段就翻译成第一中间表达式组,使得该代码段能够与与其他代码段并行执行。
而N大于1就表示对于N个代码段翻译成第一中间表达式组,使得该N个代码段能够与其他代码段并行执行,而该N个代码段内部可以是顺序(或串行)执行的,也可以是并行执行的,或者是部分顺序执行、部分并行执行的。
在一个实施例中,第一中间表达式组还包括基于N个代码段生成的第一执行块,其中第一执行块的数量为小于或等于N的正整数。在N=1的情况下,第一中间表达式组包括基于1个代码段生成的1个第一执行块。在N大于1的情况下,第一中间表达式组包括基于N个代码段生成的数量为小于或等于N的正整数个第一执行块。例如,当N=3时,第一中间表达式组包括基于3个代码段生成的1个或2个或3个第一执行块。
具体来讲,当N个代码段中不存在顺序执行或者串行的代码段时,可以基于N个代码段中的每一个代码段生成1个第一执行块,第一中间表达式组包括基于N个代码段生成的N个第一执行块;当N个代码段中的所有代码段均为顺序执行或者串行时,可以基于N个代码段中的全部N个代码段生成1个第一执行块,第一中间表达式组包括基于N个代码段生成的1个第一执行块;当N个代码段中存在部分顺序执行或者串行的代码段时,可以基于N个代码段中的部分顺序执行或串行的代码段生成相应数量个第一执行块,第一中间表达式组包括基于N个代码段生成的数量为大于1且小于N的正整数个第一执行块。
应当理解,对于基于N个代码段生成的第一执行块,在每个第一执行块内部为顺序执行或者串行,而多个第一执行块之间为并行执行。
在一个实施例中,第一中间表达式组还包括结束执行和跳转到其他控制块、或仅结束执行的第一结束块。
具体地,在第一控制块既需要控制并行/串行执行又需要控制不执行N个代码段的情况下,第一结束块可以设置为同时包含结束执行和跳转到其他控制块。例如,当第一控制块控制并行/串行执行N个代码段时,N个代码段被执行之后执行第一结束块的结束执行;在第一控制块控制不执行N个代码段的情况下,N个代码段不被执行而跳转到第一结束块执行跳转到其他控制块。
在第一控制块仅需要控制并行/串行执行N个代码段的情况下,例如最后一个或最后一批代码段,第一结束块可以设置为仅包含结束执行。例如,第一控制块控制并行/串行执行N个代码段,N个代码段被执行之后执行第一结束块的结束执行。
第一结束块可以是与N个第一执行块相对应的N个第一结束块,或1个第一结束块。
图3A示出了根据本申请的实施例的控制块301、执行块302和结束块303的数量的一个实施例的示意图。
具体地,在基于N个代码段生成的1个第一控制块301、N个第一执行块302(例如图3A所示的执行块load_matrix_sync、执行块mma_sync、执行块store_matrix_sync)、使得该并行执行的N个代码段能够与其他代码段并行执行的情况下,可以为N个第一执行块302的每个设置一个第一结束块303,即为N个第一执行块302设置N个第一结束块303,以便这些第一执行块302的执行和结束与该中间表达式组中的其他第一执行块302的执行和结束无关、也与其他中间表达式组(例如第二中间表达式组)中的其他执行块的执行和结束无关,从而可以独立地执行。也就是说,当第一控制块301控制并行执行N个代码段时,这N个第一执行块302之间可以并行执行,且这些第一执行块302与其他中间表达式组的执行块之间也是并行执行的,即N个第一执行块302中的任意一个第一执行块302执行完毕后进入到与其对应的第一结束块303,就可以结束执行,而与其他执行块的执行完成与否无关。当第一控制块301控制不执行N个代码段时,N个第一执行块302不被执行而跳转到与其对应的第一结束块303,以执行跳转到其他控制块。
图3B示出了根据本申请的实施例的控制块301、执行块302和结束块303的数量的另一个实施例的示意图。
具体地,在基于N个代码段生成的1个第一控制块301、N个第一执行块302(例如图3B所示的执行块load_matrix_sync、执行块mma_sync、执行块store_matrix_sync)、使得该并行执行的N个代码段能够与其他代码段并行执行的情况下,也可以为N个第一执行块302仅设置1个第一结束块303,以便这些第一执行块302可以共用1个第一结束块303。同样地,当第一控制块301控制并行执行N个代码段时,这N个第一执行块302之间可以并行执行,且这些第一执行块302与其他中间表达式组的执行块之间也是并行执行的,即N个第一执行块302中的任意一个第一执行块302执行完毕后均可以独立地进入第一结束块303以结束执行,而与其他执行块的执行完成与否无关。当第一控制块301控制不执行N个代码段时,N个第一执行块302不被执行而跳转到其共同的第一结束块303,以执行跳转到其他控制块。
在此,第一控制块301起到了至关重要的控制作用,其要控制执行N个第一执行块302或者控制不执行N个第一执行块302而直接跳转到第一结束块303。即,第一控制块301在控制并行执行N个代码段的情况下控制执行第一执行块302。第一控制块301在控制不执行N个代码段的情况下控制不执行第一执行块302而直接跳转到第一结束块303。同时,在N大于1的情况下,即存在多个第一执行块302时,第一控制块301还要在第一表达式组内部通过对N个第一执行块302的跳转控制来实现N个第一执行块302之间的并行执行。
进一步地,第一控制块301在控制并行执行N个代码段的情况下,控制并行执行N个第一执行块302后执行第一结束块303的结束执行;第一控制块301在控制不执行N个代码段的情况下,控制不执行N个第一执行块302而直接跳转到第一结束块303后执行跳转到其他控制块。
图3C和图3D示出了根据本申请的实施例的控制块301、执行块302和结束块303的数量的再一些实施例的示意图。
在一个实施例中,第一控制块301的数量为1个或N个。第一控制块301的数量为1个时的情况在图3A和图3B中已经示出。图3C和图3D主要示出控制块的数量为N个时的情况。在此,可以为一个中间表达式组中的每一个执行块设置一个控制块。例如有N个第一执行块302,可以设置N个第一控制块301,N个第一控制块301与N个第一执行块302一一对应,即每个第一控制块301可以独立地控制其对应的第一执行块302的执行或跳过。当然,如前所述,如果一个中间表达式组中的控制块的数量是1,那么不管该中间表达式组中有多少个执行块,该1个控制块都能够控制该中间表达式组中的所有执行块的执行或跳过。
综上,当N个代码段中不存在顺序执行或者串行的代码段时,以一个中间表达式组中的“控制块的数量-执行块的数量-结束块的数量”为表示形式,本申请的第一中间表达式组可以包括如下几种组合: 1-N-1、N-N-1、1-N-N、N-N-N等(以该中间表达式组中总共有N个代码段为例,N是大于或等于1的正整数)。例如,当N=1时,每个中间表达式组包括基于1个代码段生成的1个控制块、1个执行块和1个结束块。又例如,当N=3时,第一中间表达式组包括基于3个代码段生成的1个第一控制块301、3个第一执行块302和1个第一结束块303(例如图3B),或者3个第一控制块301、3个第一执行块302和1个第一结束块303(例如图3C),或者1个第一控制块301、3个第一执行块302和3个第一结束块303(例如图3A),或者3个第一控制块301、3个第一执行块302和3个第一结束块303(例如图3D)。
此外,上述控制块和结束块可以称为调度基本块,而执行块可以称为异步执行基本块,调度基本块负责具体的指令调度逻辑,异步执行基本块负责实际指令执行逻辑。这种分块的逻辑非常便于映射到具体的硬件指令,加速了代码的翻译和执行速度。
以上的控制块、执行块和结束块都是中间表达式( intermediaterepresentation, IR)形式的,只要该控制块的中间表达式是用来控制执行执行块还是跳转到结束块的即可,且该结束块的中间表达式是用来结束执行或是跳转到其他控制块的即可。
将代码段翻译为中间表达式可以向上能支持多语言的映射,向下能适应多平台转换且易于进行各种优化,而且在翻译代码段的过程中插入控制块和结束块,使得单个代码段组(1个或N个代码段)作为单个执行主体,与其他执行主体没有顺序执行的关系,即可以与其他执行主体同时并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。
图4A示出了根据本申请的实施例的第二中间表达式组的示例。
在一个示例中,代码生成器102被配置为将识别的多个代码段中的另外M个代码段翻译为第二中间表达式组,其中第二中间表达式组至少包括控制并行/串行执行或不执行M个代码段的第二控制块401、基于M个代码段生成的第二执行块402,以及结束执行和跳转到其他控制块、或仅结束执行的第二结束块403。其中,第一结束块303跳转到第二控制块401,其中M是1或大于1的正整数。
在此,与如前的N个代码段类似地,将另外M个代码段翻译为第二中间表达式组,并通过第二控制块401来控制并行/串行执行或不执行第二执行块402,以及通过第二结束块403来结束执行或跳转到其他控制块。基于N个代码段生成的第一结束块303用于跳转到基于该M个代码段生成的第二控制块401,以便第一中间表达式组的第一执行块302可以和第二中间表达式组的第二执行块402并行执行。在此,该M个代码段内部可以是顺序执行的,也可以是并行执行的,或者是部分顺序执行、部分并行执行的,此处不再赘述。
第二中间表达式组中的控制块、执行块和结束块的数量与图3A-图3C描述的第一中间表达式组中的控制块、执行块和结束块的数量类似。即,当M个代码段中不存在顺序执行或者串行的代码段时,本申请的第二中间表达式组可以包括“控制块的数量-执行块的数量-结束块的数量”的如下几种组合:1-M-1、M - M -1、1- M - M、M - M - M等(以该中间表达式组中总共有N个代码段为例,M是大于或等于1的正整数),在此不赘述。其中,N和M可以相同或者不同,本申请不做限制。
图4B示出了根据本申请的实施例的第一中间表达式组和第二中间表达式组中的执行块执行时对应的线程的示例。
在一个实施例中,一个执行块的执行对应一个或多个线程来完成,要使得多个执行块之间同时并行执行(或互相独立地执行),那么它们占用的线程不能互相重合。例如,在一个实施例中,第一控制块301控制第一数量个线程执行第一执行块302,与此同时,第二控制块401控制第二数量个线程执行第二执行块402,第一数量个线程和第二数量个线程分别为对应多个代码段的多个线程中的一部分线程,且第一数量个线程与第二数量个线程不同。其中第一数量可以等于或不等于第二数量。优选地,在本发明的实施例中,第一数量等于第二数量。
举例来讲,对应多个代码段的多个线程为100个线程。在该100个线程中有20个线程被配置为在第一执行块302上执行,而另外80个线程被配置为不在第一执行块302上执行。第一控制块301控制上述20个线程执行第一执行块302。同时,第一控制块301控制上述80个线程不执行第一执行块302而直接跳转到第一结束块303,并由第一结束块303继续跳转到第二控制块401。同样地,当跳转到第二控制块401的80个线程中有30个线程被配置为在第二执行块402上执行而另外50个线程被配置为不在第二执行块402上执行时,第二控制块401控制上述30个线程执行第二执行块402,同时第二控制块401控制上述50个线程不执行第二执行块402而直接跳转到第二结束块403,并由第二结束块403继续跳转到下一中间表达式组的控制块,以此类推。由此,通过设于控制块和结束块中的跳转机制,实现了多个中间表达式组(的控制块)之间的并行执行。
进一步地,在上述示例中,当第一中间表达式组中存在多个第一执行块302时,第一控制块301还要在第一表达式组内部通过对N个第一执行块302的跳转控制来实现多个线程在N个第一执行块302之间的并行执行。
举例来讲,当N=3,即第一中间表达式组中存在3个第一执行块302时,上述对应多个代码段的100个线程中有20个线程被配置为在第一执行块302上执行。此时,第一控制块301可以将上述20个线程分配给3个第一执行块302并行执行。例如,第一控制块301可以控制20个线程中的5个线程执行3个第一执行块302中的第一个第一执行块302,并控制另外15个线程跳转至3个第一执行块302中的第二个第一执行块302。第一控制块301还控制另外15个线程中的7个线程执行第二个第一执行块302,并控制剩余的8个线程跳转至3个第一执行块302中的第三个第一执行块302。第一控制块301控制上述8个线程执行第三个第一执行块302。
应当理解,线程是否在某一执行块(例如第一执行块302或第二执行块402)上执行可以取决于线程是否被配置为在该执行块对应的硬件模块上执行。
在本申请的以上实施例中,控制块和结束块中设置的跳转机制可以通过特定的跳转语句来实现。此外,跳转机制具有的瞬时性使得执行多个执行块的线程可以在硬件模块上实现并行执行。
需要说明的是,在实际应用中,多个中间表达式组之间并不限于本申请公开的并行执行这一种方式。例如,多个中间表达式组中的一些中间表达式组之间可以采用本申请的并行的执行方式,而另一些中间表达式组之间仍然可以为采用现有技术的顺序执行或串行的方式。举例来说,3个中间表达式组中的第一中间表达式组和第二中间表达式组之间为顺序执行,且它们与第三中间表达式组之间为并行执行。这样,一定数量个线程中的A数量个线程被分配去执行第一中间表达式组的执行块,另外B数量个线程跳转去并行执行第三中间表达式组的执行块。而A数量个线程在执行完第一中间表达式组的执行块后还可以去顺序执行第二中间表达式组的执行块。
综上,本申请能够将用标记符、例如coroutine{}标记的多个代码段翻译为在GPU上并行执行的代码段。通过对编译器前端的解析器模块和代码生成器模块进行修改,把源代码中的coroutine标记的模块翻译为包括3个不同的基本块的中间表达式组,来帮助或控制翻译后的代码段之间的并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。本申请的技术能够直接应用到使用诸如c++编程语言和LLVM(Low LevelVirtual Machine,低级虚拟机)编译器的所有编程和编译领域。
图5示出了根据本申请的实施例的代码编译方法500的示意流程图。
该代码编译方法500包括:步骤501,基于标记代码段的标记符识别多个代码段,多个代码段中的至少一部分并行执行;步骤502,将识别的多个代码段中的N个代码段翻译为第一中间表达式组,其中第一中间表达式组至少包括控制并行/串行执行或不执行N个代码段的第一控制块,其中N是1或大于1的正整数。
如此,通过将N个代码段翻译为具有控制其执行的第一控制块的第一中间表达式组,即可以作为单个执行主体,与(也是翻译为具有控制块的中间表达式组的)其他执行主体没有顺序执行的关系,即可以与其他执行主体同时并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。
这里从多个代码段中选择N个代码段作为一个组来进行翻译。这里的选择可以是有目的的选择,例如选择为同一任务服务的N个代码段,以便在执行这N个代码段时,能够尽快执行完同一任务。这里的选择也可以是随机的,即从多个代码段中随机选其中N个代码段来作为一个组。
这里,N=1就表示对于1个代码段就翻译成第一中间表达式组,使得该代码段能够与与其他代码段并行执行。
而N大于1就表示对于N个代码段翻译成第一中间表达式组,使得该N个代码段能够与其他代码段并行执行,而该N个代码段内部可以是顺序执行的,也可以是并行执行的,或者是部分顺序执行、部分并行执行的。
在一个实施例中,第一中间表达式组还包括基于N个代码段生成的第一执行块,其中第一执行块的数量为小于或等于N的正整数。在N=1的情况下,第一中间表达式组包括基于1个代码段生成的1个第一执行块。在N大于1的情况下,第一中间表达式组包括基于N个代码段生成的数量为小于或等于N的正整数个第一执行块。例如,当N=3时,第一中间表达式组包括基于3个代码段生成的1个或2个或3个第一执行块。
在一个实施例中,第一中间表达式组还包括结束执行和跳转到其他控制块、或仅结束执行的第一结束块,其中第一结束块是与N个第一执行块相对应的N个第一结束块,或1个第一结束块。
在一个实施例中,第一控制块在控制并行/串行执行N个代码段的情况下控制执行第一执行块,第一控制块在控制不执行N个代码段的情况下控制不执行第一执行块而直接跳转到第一结束块。
在一个实施例中,该代码编译方法还包括将识别的多个代码段中的另外M个代码段翻译为第二中间表达式组,其中第二中间表达式组至少包括控制并行/串行执行或不执行M个代码段的第二控制块、基于M个代码段生成的第二执行块,以及结束执行和跳转到其他控制块、或仅结束执行的第二结束块。其中,第一结束块跳转到第二控制块,其中M是1或大于1的正整数。
在一个实施例中,第二执行块与第一执行块并行执行。
在一个实施例中,第一控制块控制第一数量个线程执行第一执行块,第二控制块控制第二数量个线程执行第二执行块,第一数量个线程和第二数量个线程分别为对应多个代码段的多个线程中的一部分线程,且第一数量个线程与第二数量个线程不同。
在一个实施例中,第一控制块的数量为1个或N个,第二控制块的数量为1个或M个。
在一个实施例中,N和M相同或者不同。
在一个实施例中,至少一部分标记符标记了能够被并行执行的多个代码段。
在一个实施例中,将识别的多个代码段中的N个代码段翻译为第一中间表达式组的步骤包括:基于标记代码段的标记符识别多个代码段后,将识别的多个代码段翻译为抽象语法树中的至少一个复合表达式节点;根据至少一个复合表达式节点将多个代码段中的N个代码段翻译为第一中间表达式组。这里,根据抽象语法树的复合表达式节点来翻译为第一中间表达组可以保留代码段的所有语法信息,从而完整地执行代码段的动作。
综上,本申请能够将用标记符、例如coroutine{}标记的多个代码段翻译为在GPU上并行执行的代码段。通过对编译器前端的解析器模块和代码生成器模块进行修改,把源代码中的coroutine标记的模块翻译为包括3个不同的基本块的中间表达式组,来帮助或控制翻译后的代码段之间的并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。
图6示出了根据本申请的实施例的整体的代码编译方法的一个示例的示意数据流向图。
如图6所示,将标记符coroutine包含的代码段通过解析器处理后变成抽象语法树(AST),然后抽象语法树经过代码生成器之后,变成中间表达式IR。本申请在解析器模块中添加对标记符coroutine关键字的识别,并把标记符coroutine{}翻译为抽象语法树中的一个复合表达式节点。在对抽象语法树进行处理的生成IR(中间表达式)的时候,每一个coroutine 表达式都会被翻译出来包括如下3个基本块的中间表达式的中间表达式组:控制块、执行块和结束块。
控制块负责调度逻辑,决定是否执行或跳过执行块。执行块里面对应coroutine{代码}中的代码实现逻辑。结束块里面包含了跳转到下一个标记符coroutine的控制块的逻辑,来帮助或控制翻译后的代码段之间的并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,GeneralPurpose GPU)等的硬件架构来讲。
图7示出了适于用来实现本申请实施方式的示例性代码编译系统的框图。
计算机系统可以包括处理器(H1);存储器(H2),耦合于处理器(H1),且在其中存储计算机可读指令,用于在由处理器执行时进行本申请的实施例的各个代码编译方法的步骤。
处理器(H1)可以包括但不限于例如一个或者多个处理器或者或微处理器等。
存储器(H2)可以包括但不限于例如,随机存取存储器(RAM)、只读存储器(ROM)、快闪存储器、EPROM存储器、EEPROM存储器、寄存器、计算机存储介质(例如硬碟、软碟、固态硬盘、可移动碟、CD-ROM、DVD-ROM、蓝光盘等)。
除此之外,该计算机系统还可以包括数据总线(H3)、输入/输出(I/O)总线(H4),显示器(H5)以及输入/输出设备(H6)(例如,键盘、鼠标、扬声器等)等。
处理器(H1)可以通过I/O总线(H4)经由有线或无线网络(未示出)与外部设备(H5、H6等)通信。
存储器(H2)还可以存储至少一个计算机可读指令,用于在由处理器(H1)运行时执行本技术所描述的实施例中的各个功能和/或方法的步骤。
在一个实施例中,该至少一个计算机可读指令也可以被编译为或组成一种软件产品,其中一个或多个计算机可读指令被处理器运行时执行本技术所描述的实施例中的各个功能和/或方法的步骤。
图8示出了根据本公开的实施例的非暂时性计算机可读存储介质的示意图。
如图8所示,计算机可读存储介质820上存储有指令,指令例如是计算机可读指令810。当计算机可读指令810由处理器运行时,可以执行参照以上描述的各个方法。计算机可读存储介质包括但不限于例如易失性存储器和/或非易失性存储器。易失性存储器例如可以包括随机存取存储器(RAM)和/或高速缓冲存储器(cache)等。非易失性存储器例如可以包括只读存储器(ROM)、硬盘、闪存等。例如,计算机可读存储介质820可以连接于诸如计算机等的计算设备,接着,在计算设备运行计算机可读存储介质820上存储的计算机可读指令810的情况下,可以进行如上描述的各个方法。
如此,通过将N个代码段翻译为具有控制其执行的第一控制块的第一中间表达式组,即可以作为单个执行主体,与(也是翻译为具有控制块的中间表达式组的)其他执行主体没有顺序执行的关系,即可以与其他执行主体同时并行执行,从而减少代码段之间的相互等待和延时,加快整体代码执行的速度,尤其是对于例如通用图形处理器(GPGPU,General Purpose GPU)等的硬件架构来讲。
当然,上述的具体实施例仅是例子而非限制,且本领域技术人员可以根据本申请的构思从上述分开描述的各个实施例中合并和组合一些步骤和装置来实现本申请的效果,这种合并和组合而成的实施例也被包括在本申请中,在此不一一描述这种合并和组合。
注意,在本公开中提及的优点、优势、效果等仅是示例而非限制,不能认为这些优点、优势、效果等是本申请的各个实施例必须具备的。另外,上述公开的具体细节仅是为了示例的作用和便于理解的作用,而非限制,上述细节并不限制本申请为必须采用上述具体的细节来实现。
本公开中涉及的器件、装置、设备、系统的方框图仅作为例示性的例子并且不意图要求或暗示必须按照方框图示出的方式进行连接、布置、配置。如本领域技术人员将认识到的,可以按任意方式连接、布置、配置这些器件、装置、设备、系统。诸如“包括”、“包含”、“具有”等等的词语是开放性词汇,指“包括但不限于”,且可与其互换使用。这里所使用的词汇“或”和“和”指词汇“和/或”,且可与其互换使用,除非上下文明确指示不是如此。这里所使用的词汇“诸如”指词组“诸如但不限于”,且可与其互换使用。
本公开中的步骤流程图以及以上方法描述仅作为例示性的例子并且不意图要求或暗示必须按照给出的顺序进行各个实施例的步骤。如本领域技术人员将认识到的,可以按任意顺序进行以上实施例中的步骤的顺序。诸如“其后”、“然后”、“接下来”等等的词语不意图限制步骤的顺序;这些词语仅用于引导读者通读这些方法的描述。此外,例如使用冠词“一个”、“一”或者“该”对于单数的要素的任何引用不被解释为将该要素限制为单数。
另外,本文中的各个实施例中的步骤和装置并非仅限定于某个实施例中实行,事实上,可以根据本申请的概念来结合本文中的各个实施例中相关的部分步骤和部分装置以构思新的实施例,而这些新的实施例也包括在本申请的范围内。
以上描述的方法的各个操作可以通过能够进行相应的功能的任何适当的手段而进行。该手段可以包括各种硬件和/或软件组件和/或模块,包括但不限于硬件的电路、专用集成电路(ASIC)或处理器。
可以利用被设计用于进行在此描述的功能的通用处理器、数字信号处理器(DSP)、ASIC、场可编程门阵列信号(FPGA)或其他可编程逻辑器件(PLD)、离散门或晶体管逻辑、离散的硬件组件或者其任意组合而实现或进行描述的各个例示的逻辑块、模块和电路。通用处理器可以是微处理器,但是作为替换,该处理器可以是任何商业上可获得的处理器、控制器、微控制器或状态机。处理器还可以实现为计算设备的组合,例如DSP和微处理器的组合,多个微处理器、与DSP核协作的微处理器或任何其他这样的配置。
结合本公开描述的方法或算法的步骤可以直接嵌入在硬件中、处理器执行的软件模块中或者这两种的组合中。软件模块可以存在于任何形式的有形存储介质中。可以使用的存储介质的一些例子包括随机存取存储器(RAM)、只读存储器(ROM)、快闪存储器、EPROM存储器、EEPROM存储器、寄存器、硬碟、可移动碟、CD-ROM等。存储介质可以耦接到处理器以便该处理器可以从该存储介质读取信息以及向该存储介质写信息。在替换方式中,存储介质可以与处理器是整体的。软件模块可以是单个指令或者许多指令,并且可以分布在几个不同的代码段上、不同的程序之间以及跨过多个存储介质。
在此公开的方法包括用于实现描述的方法的动作。方法和/或动作可以彼此互换而不脱离权利要求的范围。换句话说,除非指定了动作的具体顺序,否则可以修改具体动作的顺序和/或使用而不脱离权利要求的范围。
上述功能可以按硬件、软件、固件或其任意组合而实现。如果以软件实现,功能可以作为指令存储在切实的计算机可读介质上。存储介质可以是可以由计算机访问的任何可用的切实介质。通过例子而不是限制,这样的计算机可读介质可以包括RAM、ROM、EEPROM、CD-ROM或其他光碟存储、磁碟存储或其他磁存储器件或者可以用于携带或存储指令或数据结构形式的期望的程序代码并且可以由计算机访问的任何其他切实介质。如在此使用的,碟(disk)和盘(disc)包括紧凑盘(CD)、激光盘、光盘、数字通用盘(DVD)、软碟和蓝光盘,其中碟通常磁地再现数据,而盘利用激光光学地再现数据。
因此,计算机程序产品可以进行在此给出的操作。例如,这样的计算机程序产品可以是具有有形存储(和/或编码)在其上的指令的计算机可读的有形介质,该指令可由处理器执行以进行在此描述的操作。计算机程序产品可以包括包装的材料。
软件或指令也可以通过传输介质而传输。例如,可以使用诸如同轴电缆、光纤光缆、双绞线、数字订户线(DSL)或诸如红外、无线电或微波的无线技术的传输介质从网站、服务器或者其他远程源传输软件。
此外,用于进行在此描述的方法和技术的模块和/或其他适当的手段可以在适当时由用户终端和/或基站下载和/或其他方式获得。例如,这样的设备可以耦接到服务器以促进用于进行在此描述的方法的手段的传送。或者,在此描述的各种方法可以经由存储部件(例如RAM、ROM、诸如CD或软碟等的物理存储介质)提供,以便用户终端和/或基站可以在耦接到该设备或者向该设备提供存储部件时获得各种方法。此外,可以利用用于将在此描述的方法和技术提供给设备的任何其他适当的技术。
其他例子和实现方式在本公开和所附权利要求的范围和精神内。例如,由于软件的本质,以上描述的功能可以使用由处理器、硬件、固件、硬连线或这些的任意的组合执行的软件实现。实现功能的特征也可以物理地位于各个位置,包括被分发以便功能的部分在不同的物理位置处实现。而且,如在此使用的,包括在权利要求中使用的,在以“至少一个”开始的项的列举中使用的“或”指示分离的列举,以便例如“A、B或C的至少一个”的列举意味着A或B或C,或AB或AC或BC,或ABC(即A和B和C)。此外,措辞“示例的”不意味着描述的例子是优选的或者比其他例子更好。
可以不脱离由所附权利要求定义的教导的技术而进行对在此描述的技术的各种改变、替换和更改。此外,本公开的权利要求的范围不限于以上描述的处理、机器、制造、事件的组成、手段、方法和动作的具体方面。可以利用与在此描述的相应方面进行基本相同的功能或者实现基本相同的结果的当前存在的或者稍后要开发的处理、机器、制造、事件的组成、手段、方法或动作。因而,所附权利要求包括在其范围内的这样的处理、机器、制造、事件的组成、手段、方法或动作。
提供所公开的方面的以上描述以使本领域的任何技术人员能够做出或者使用本申请。对这些方面的各种修改对于本领域技术人员而言是非常显而易见的,并且在此定义的一般原理可以应用于其他方面而不脱离本申请的范围。因此,本申请不意图被限制到在此示出的方面,而是按照与在此公开的原理和新颖的特征一致的最宽范围。
为了例示和描述的目的已经给出了以上描述。此外,此描述不意图将本申请的实施例限制到在此公开的形式。尽管以上已经讨论了多个示例方面和实施例,但是本领域技术人员将认识到其某些变型、修改、改变、添加和子组合。