CN109901840A - 一种线程间冗余删除的异构编译优化方法 - Google Patents

Abstract

本发明通过自动选择优化策略以及分维度进行线程间冗余删除的收益代价分析，能自动根据代码是否含有同步操作和硬件平台的特点，针对性地选择一种优化策略，通过对Kernel代码的编译分析，得出合适的合并维度和合并因子，利于提高Kernel代码的运行性能。采用本发明方法能面向多种异构平台，进行自动分析和变换，得以优化设备端代码的线程间冗余和同步开销，提高设备端代码(即OpenCL的Kernel代码)的性能。

Description

一种线程间冗余删除的异构编译优化方法

技术领域

本发明涉及编译器的研发和优化技术，特别是涉及一种面向异构平台的线程间冗余删除编译优化技术。

背景技术

随着摩尔定律的放缓、GPU等加速设备获得越来越多的高性能计算应用的支持，现代计算系统正在向装备有CPU和加速设备(如GPU或Xeon Phi)的异构平台过渡，从全球TOP500超级计算机榜单上排名前三的皆是异构平台、异构平台数量超过100中可见一斑。异构系统通常由CPU与一个或多个加速设备在片上或主板上相互连接组成，其中CPU负责复杂的控制、调度等工作，而加速设备则负责大规模并行计算或专业领域的计算任务。当今市场上加速设备种类很多，每种类型的加速设备具有不同的特性，包括：CPU、GPU、协处理器、DSP(信号处理器)、ASIC(专用集成电路)、FPGA(现场可编程门阵列)等等。

与异构系统相对应，异构程序亦由两部分组成——主机端代码和设备端代码(即kernel)组成，主机端代码通过提交命令来驱动设备进行并行计算，而并行执行的kernel代码是编译优化的重点。面向异构系统的编程模型，有只适用于特定异构系统的编程模型，比如NVIDIA的CUDA，微软提出的C++AMP，IBM的LIME，Intel提出的Merge等等；也有通用的编程模型，由Khronos Group提出的OpenCL异构编程模型为不同的设备架构提供了统一接口，为异构计算提供了一个开放的、通用标准。OpenCL执行模型大致如下：当主机端代码提交kernel代码执行时，系统生成一个n维(n≤3)的索引空间NDRange，kernel代码被实例化成若干线程并且按照NDRange的形式组织线程并发执行。NDRange中的每个线程称为工作项(work-item)，work-item之间以SPMD(Same Program Multiple Data)模式并行执行，即每个work-item在不同的数据上执行相同的kernel程序。NDRange还可以根据粗粒度分解成若干工作组work-group：其空间维度与全局空间的维度相同，每个工作组work-group包含数目相同的相邻工作项。work-item之间可以进行局部(work-group内部)或全局(NDRange)的同步。如何将OpenCL程序映射到具体硬件单元完成相应操作，并管理和组织执行线程，是由具体硬件平台的OpenCL运行时支持系统完成，很多厂商通过提供OpenCL SDK使得OpenCL程序能直接运行于各自平台。

OpenCL编程语言具有良好的功能移植性，程序可面向多种不同架构的平台。OpenCL程序的线程(work-item)之间往往存在着完全相同的或者值相同的一些代码，以及开销大的同步操作。一方面这是与程序并行度相关的、自然形成的，提高程序并行度，相应地线程数量增多、线程粒度变小，因此需要引入更多的同步操作；同时一些计算变成每个线程内需计算一次，这样的操作于单个线程内部是不冗余的、于多个线程之间则存在冗余性，这里我们称为线程间的冗余。线程间的冗余操作的存在不一定益于性能，可以借助编译技术分析和优化。而另一方面这也不能忽略程序员编程习惯的人为影响。程序员倾向于从功能角度划分独立线程、充分发掘算法的并行度，而对于线程之间的优化则缺少关注、留待编译器完成。

现有的异构程序编译器通常对Kernel代码进行静态分析，重点关注于线程内的优化，它们不考虑线程之间的代码冗余，不提供自动的线程间冗余删除或线程合并优化。

发明内容

因此，本发明的目的在于克服上述现有技术的缺陷，提供一种新线程间冗余删除的异构编译优化方法。

面对线程间的冗余计算和可能冗余的同步操作，我们的方法是通过编译阶段的线程合并，使得原本线程间的冗余计算(即结果值或者功能完全相同的普通指令)变成新线程内的冗余计算得以删除，使得进行同步的线程数量减少和同步开销降低，达到提高设备端Kernel代码运行性能的效果。另一方面，线程间冗余删除也可能对性能产生副作用，于是，考虑到线程合并带来并行度的下降，以及改变原有的线程执行顺序对性能的影响，我们建立收益代价模型，分析线程合并在冗余删除和优化同步上的收益，分析它影响并行度和访存效率所带来的代价，基于收益代价分析的结果，实施代码变换，得到优化后的Kernel代码。

本发明的目的是通过以下技术方案实现的：

根据本发明的第一方面，本发明提供一种优化冗余策略，进行线程合并来达到删除冗余提高代码运行性能的目的，包括如下步骤：

S1、获取Kernel源代码及Kernel执行信息；

S2、对所述Kernel源代码及Kernel执行信息，以工作组为单位，分维度识别出候选冗余计算；

S3、基于所述候选冗余计算分维度进行收益代价分析,得到每一维度的最优合并因子及其优化冗余的纯收益；

S4、基于每一维度的最优合并因子及其优化冗余的纯收益确定总合并维度及总合并因子；

S5、基于总合并维度及总合并因子，对Kernel代码进行线程优化冗余操作，合并线程，进行代码变换和优化；

S6、输出优化后的源码形式的Kernel代码及中间文件形式的Kernel执行信息。

其中，上述步骤S2包括如下步骤：

S21、以工作组为单位，分维度对Kernel代码的中间语言进行指令之间的依赖分析，得出依赖于每一维度线程的普通指令以及不依赖于该维度线程的普通指令；普通指令为同步指令以外的其他指令；

S22、将不依赖于线程的普通指令视为该线程所在维度的候选冗余计算。

上述步骤S3包括如下步骤：

S31、在不同合并因子下，以工作组为单位，分维度计算优化冗余的收益；在不同合并因子下，以工作组为单位，分维度计算优化冗余的收益时，不同合并因子下每一维度的收益为不依赖于该维度的所有普通指令在合并因子下做优化冗余操作时为工作组带来的收益之和；不同合并因子下，每一不依赖于线程的普通指令的收益为该指令执行频率、指令延迟、合并因子减一之差这三者的乘积；

S32、结合硬件平台特征，以工作组为单位，在不同合并因子下，分维度计算优化冗余的代价；在不同合并因子下，以工作组为单位，分维度计算优化冗余的代价时，以线程合并导致的并行度下降和访存效率下降引发的性能损失之和作为代价，包括并行度影响代价和访存效率的代价；

S33、针对同一工作组同一维度线程，在不同合并因子下，用步骤S31中计算的收益与步骤S32中计算的同一维度同一合并因子下的代价相减，得到该维度线程在不同合并因子下的纯收益，取其中最大纯收益所对应的合并因子作为该维度线程的最优合并因子；

S34、重复步骤S31-S33，直至计算得到每一维度线程的最优合并因子及其对应的纯收益为止。

上述步骤S4包括如下步骤：

S41、获取每一维度线程的最优合并因子及其对应的纯收益；

S42、对所有维度的最优合并因子对应的纯收益进行对比，选出纯收益最大者所对应的维度作为总合并维度，并将该维度的最优合并因子作为总合并因子。

上述步骤S5包括如下步骤：

S51、按照选出的总合并维度和总合并因子，对Kernel代码内依赖于总合并维度的普通指令逐条复制，复制的遍数为总合并因子减一，对不依赖于总合并维度的普通指令不做复制；并转换代码中对线程号的使用；

S52、对复制后的指令中的相关变量做扩展或者更换变量名，避免变量冲突，得到变换后的新的Kernel代码；

S53、根据总合并维度和总合并因子，相应修改主机端相关代码，包括分维度的线程总数和工作组内每一维线程数。

需要说明的是，上述步骤S32还包括：

在不同合并因子下，以工作组为单位，分维度计算并行度影响代价；其中，并行度影响代价为0或者IMAX，IAMX为目标平台所能表示的最大正整数；以及，当平台上合并导致目标平台不能充分发挥并行处理能力时，并行度影响代价为IMAX，否则为0；

在不同合并因子下，以工作组为单位，分维度计算访存效率的代价；其中，访存效率影响代价以一个工作组内的访存开销在做优化冗余合并操作前后的变化值来衡量；每个合并因子值下以合并后的工作组线程的访存开销减去合并前工作组线程的访存开销所得差值作为该合并因子下导致的访存效率的代价。

根据本发明的另一方面，本发明提供一种优化同步策略，删除多于的同步操作来达到降低同步开销，提高代码运行性能的目的。在步骤S1中，在获取Kernel源代码后，对Kernel代码做词法或语义分析，并根据分析结果选择优化策略；其中当Kernel代码含有同步指令且目标平台同步为CPU平台或其他以软件实现同步操作的平台时，选择优化同步策略，否则执行步骤S2。

优化同步策略包括如下步骤：

R1、对Kernel代码的所有维度进行合并，每一维度的合并因子为工作组内相应维度的线程总数，合并后，Kernel程序中的每一个工作组内的线程为1个线程；

R2、删除工作组内的所有同步操作，实现同步优化；

R3、根据步骤R2的同步优化，进行代码变换；

R4、输出步骤R3中代码变换后的结果。

其中，上述步骤R3,包括如下步骤：

R31、针对每一指令区域分别做代码变换，以同步语句为分界线，将Kernel代码划分成指令区域，指令区域的数量为同步语句数量加1；

R32、对每个指令区域的源代码进行变换，将其包装为一个含有至少一层循环的最内层循环体，循环体层数与Kernel代码维度的维数相同，每一层循环体对应一个维度，每一层循环体的内循环次数为其对应维度的线程总数；并转换代码中对线程号的使用；

R33、扩展相应变量或修改变量，避免变量冲突。

与现有技术相比，本发明的优点在于：采用本发明的方法能够减少线程间冗余计算、优化同步开销，达到提高加速设备端代码性能的效果。同时，本方法能面向多种异构平台，进行自动分析和变换，得以优化设备端代码的线程间冗余和同步开销，提高设备端代码(即OpenCL的Kernel代码)的性能。本方法面向OpenCL编程语言及其执行模型(不限于运行在哪种体系架构的硬件平台)，但同样适用于CUDA等其他SPMD(single program multipledata)的异构编程语言。

附图说明

以下参照附图对本发明实施例作进一步说明，其中：

图1为根据本发明实施例的优化冗余策略下Kernel代码优化前的代码编码示意图；

图2为根据本发明实施例的优化冗余策略下Kernel代码沿第0维以合并因子2合并优化后的代码编码示意图；

图3为根据本发明实施例的优化冗余策略下Kernel代码沿第1维以合并因子2合并优化后的代码编码示意图；

图4为根据本发明实施例的优化同步策略下Kernel代码优化前的代码编码示意图；

图5为根据本发明实施例的优化同步策略下Kernel代码优化后的代码编码示意图；

图6为根据本发明实施例的线程间冗余删除流程示意图。

具体实施方式

为了使本发明的目的，技术方案及优点更加清楚明白，以下结合附图通过具体实施例对本发明进一步详细说明。应当理解，此处所描述的具体实施例仅用以解释本发明，并不用于限定本发明。

本发明的目的在于提供一种面向多种异构平台的线程间冗余删除的编译优化技术，旨在优化同步、删除冗余计算，通过分析选择合并维度和合并因子进行线程合并，从而减少线程间冗余计算、优化同步开销，达到提高加速设备端代码性能的效果。

概括来说，面对线程间的冗余计算和可能冗余的同步操作，本发明的方法是通过编译阶段的线程合并，使得原本线程间的冗余计算即结果值或者功能完全相同的普通指令变成新线程内的冗余计算得以删除，使得进行同步的线程数量减少和同步开销降低，达到提高设备端Kernel代码运行性能的效果。另一方面，线程间冗余删除也可能对性能产生副作用，于是，考虑到线程合并带来并行度的下降，以及改变原有的线程执行顺序对性能的影响，本发明还建立了收益代价模型，分析线程合并在冗余删除和优化同步上的收益，分析它影响并行度和访存效率所带来的代价，基于收益代价分析的结果，实施代码变换，得到优化后的Kernel代码。本方法不仅面向OpenCL编程语言及其执行模型(不限于运行在哪种体系架构的硬件平台)，也同样适用于CUDA等其他SPMD(single programmultiple data)的异构编程语言。

根据本发明的一个实施例，如图6所示，本发明的一种线程间冗余删除的异构编译优化方法主要包括如下步骤：

步骤1、从平台获取Kernel源代码；

步骤2、对Kernel源代码进行词法或者语法分析；

步骤3、根据步骤2中的分析结果，结合平台信息选择优化策略；以一个工作组(OpenCL的一个work-group)作为考察对象，若Kernel含有同步指令且运行平台为CPU平台(或其他以软件实现的同步操作的平台)，则选择优化同步策略，转到步骤8；否则选择优化冗余策略，转到步骤4；

步骤4、从平台主机端代码获取Kernel执行信息，包括线程总数globalWorksize、各维度线程总数globalWorksize(i)、工作组内线程总数localWorksize、工作组内各维度线程总数localWorksize(i)，0≤i≤n-1，并以中间文件形式输入，转到步骤5；

步骤5、以一个工作组为考察对象，分维度识别候选冗余计算，转到步骤6；

步骤6、根据步骤5识别出的候选冗余计算，分维度进行收益代价分析，转到步骤7；

步骤7、根据步骤6中的收益代价分析，选出最大纯收益所对应的维度及合并因子作为总合并维度及总合并因子，对Kernel代码根据总合并维度及总合并因子进行合并，修改代码中对线程号的使用，扩展相应变量或修改变量，避免变量冲突，实现优化冗余操作，转到步骤9；

步骤8、采取对工作组完全合并策略对Kernel代码进行合并，合并维度为所有维度，每一维度合并因子为工作组内相应维度的线程总数，合并后使每个工作组变成一个线程，并对指令中出现的线程号的使用做相应转换，扩展相应变量或修改变量，避免变量冲突，然后删除同步操作，转到步骤9；

步骤9、对执行了优化同步策略或优化冗余策略的工作组，以中间文件形式输出优化后的Kernel执行相关信息并相应修改主机端代码信息，包括分维度的线程总数和工作组内每一维线程数，转到步骤10；

步骤10、输出优化后的源码形式的Kernel代码。

需要说明的是，有些步骤，比如步骤1、步骤2、步骤3、步骤4、步骤10等是本领域公知的，发明人不再展开描述，以下将分别针对步骤5、步骤6、步骤7、步骤8等进行详细描述。

在步骤5至步骤7中，执行优化冗余策略，通过线程合并使得一些线程之间的重复操作成为冗余，可做删除，提高Kernel代码运行性能。在本发明中通过分维度的收益代价分析，确定合适的合并维度和合并因子。首先，以一个工作组为考察对象，先对每一维度，分别计算不同合并因子带来的纯收益值，纯收益最大的合并因子被认为是该维度的最优合并因子；再从所有维度中选择纯收益最大者为最优合并维度和合并因子。

其中，在步骤5中，以一个工作组为考察对象，分维度识别候选冗余计算，通过对Kernel代码的中间语言进行指令之间的依赖分析实现。借助编译器指令调度优化中常用的依赖图DAG得出依赖于工作组某一个维度的指令或称操作，以及不依赖于这一维度的指令。这些不依赖于该维度的同步指令以外的指令即普通指令视为该维度的候选冗余计算，冗余计算是收益的来源。假设Kernel代码线程有n维，每一工作组内的线程号以id_i(i＝0,1,…,n-1)表示，根据OpenCL Specification，线程号最多为三维故而n的值可能为1、2或3，合并因子为cf，Kernel代码在编译器进行线程间冗余删除优化前的程序执行过程中，第i维的相邻cf个线程中的所有指令都执行一次即Kernel代码中的所有指令动态执行cf遍；在编译器进行优化冗余操作后，新的Kernel代码中依赖于idi的指令被复制了(cf-1)遍，不依赖于id_i的指令不做复制，那么，程序执行过程中，第i维的相邻cf个线程中的依赖于id_i的指令被执行cf次，不依赖id_i的指令执行一次，这些不依赖于id_i的指令通过一定的线程合并后成为冗余计算，是收益的来源。

根据本发明的一个实施例，如图1所示，图1是两维的Kernel代码，每一维以指令标号表示：对于第0维，依赖于id₀线程号的指令集合有DependOn(0)＝{insn₀,insn₁,insn₆,insn₈,insn₉}，不依赖id₀的指令集合有UnDependOn(0)＝{insn₂,insn₃,insn₄,insn₅,insn₇}；UnDependOn(0)是第0维收益的来源；对于第1维，依赖于id1线程号的指令集合有DependOn(1)＝{insn₀,insn₂,insn₇,insn₈,insn₉}，不依赖于id₁的指令集合有UnDependOn(1)＝{insn₁,insn₃,insn₄,insn₅,insn₆}。UnDependOn(1)是第1维收益的来源。

在步骤5识别出冗余计算后，根据步骤6要针对工作组的每一维进行收益代价分析，我们通过对一个工作组内的线程，通过收益代价分析，选择某一维度以合理的合并因子做线程合并，以获得的最大纯收益。

其中，在步骤6中分维度计算收益时，收益以拍数衡量(以正整数表示)，指合并带来的一个工作组内的总性能收益。根据本发明的一个实施例，计算第i维的收益为Benefit(i,cf)＝(∑(Freq_Thread(inst_j)*Latency(inst_j)*(cf-1)))*Wgsize(cf)，即合并后单线程内的收益乘以工作组内线程数的乘积，此处cf为合并因子，i表示工作组内第i维线程，inst_j为Kernel内任一不依赖于第i维线程的指令，Wgsize(cf)表示工作组在合并因子cf下合并后的线程数。单inst_j线程内的执行频率和指令延迟以及合并因子减一之差的乘积值作为该指令在单线程内的优化冗余收益；所有不依赖于第i维线程的指令的该乘积值求和得到维度i在合并因子cf下的单线程内收益；然后将该收益值乘以工作组内线程数所得乘积值即为维度i在合并因子cf下的总收益。具体地说，参考目标硬件平台上指令的拍数(即指令执行周期数或称指令延迟)，将指令大致分类成加减、乘除、访存(即load或store)、分支操作、复杂数学计算这几大类，其中按照OpenCL Specification，load/store进一步细分成global Memory load/store、constant Memory load/store、local Memory load/store三种，分别代表访问主存、访问常数存储器、访问局部存储器，平台因三种存储器的访存速度不同故而给予不同的拍数。对一个n维的工作组分维度计算收益,n≤3，对于第i(i＝{0,1，...n-1})维上的合并因子cf取值依次分析收益(cf＝{2,3,4,…,localWorkSize(i)})，包括如下步骤：

3.1.1)剔除不能整除localWorkSize(i)的合并因子。即若LocalWorkSize(i)/cf的余数不为0，则本cf值不合要求，转而考察cf的下一个值；否则，对本cf值继续执行步骤3.1.2；

3.1.2)计算本cf值下第i维的总收益。具体地：计算本cf值下的一个工作组内的总收益为Benefit(i,cf)＝(∑(Freq_Thread(inst_j)*Latency(inst_j)*(cf-1)))*Wgsize(cf)。先计算本cf值下的合并后的单线程内的收益，取Kernel内所有不依赖于id_i的普通指令在单线程内的收益之和，这里inst_j为Kernel内任一不依赖于第i维线程的普通指令，insn_j属于UnDependOn(i)集合，取其单线程内的执行频率Freq_Thread(inst_j)和指令延迟与合并因子减一之差的乘积值作为该条指令在单线程内的优化冗余收益。其次，计算本cf值下合并后的工作组内的总线程数(Wgsize(cf))，工作组内的总线程数变成原始工作组内线程数的1/cf，即Wgsize(cf)＝PRODUCT(localWorkSize(0),…,localWorkSize(i),…,localWorkSize(n-1))/cf。再次，将上述单线程内的收益乘以工作组内的线程总数得到本cf值下工作组的总收益。

其中，通过对第i维上的合并因子值的分析，得到候选的合并因子值的集合candidate_cf_set(i)＝{cf₁,cf₂,…,cf_t}，任一cf_p(1≤p≤t)均满足：2≤cf_p≤localWorkSize(i)，且localWorkSize(i)/cf_p的余数为0，第i维上的最优合并因子将从candidate_cf_set(i)中选取。同时，任一cf_p已计算得到相应的总收益值。

根据本发明的一个实施例，仍旧以图1中代码为例，其为2维Kernel代码，图2为图1中代码沿第0维以合并因子2进行合并后的代码，图3中代码为图1中代码沿第1维以合并因子2进行合并后的代码。设指令延迟为Latency(insn₂)＝Latency(insn₃)＝Latency(insn₄)＝Latency(insn₅)＝1,Latency(insn₁)＝Latency(insn₆)＝Latency(insn₇)＝Latency(insn₉)＝100，以k值指原始单个线程内指令执行次数(图1中即循环次数)，该值通过传统编译分析技术(比如常数传播)得到。以下为计算简便，假定此处k值为10，一个工作组合并前的线程数为(16，8)。

第0维候选的合并因子值的集合candidate_cf_set(0)＝{2,4,8,16}；第1维候选的合并因子值的集合candidate_cf_set(1)＝{2,4,8}；

对第0维计算合并收益：

Benefit(i,cf)＝(∑(Freq_Thread(inst_j)*Latency(inst_j)*(cf-1)))*Wgsize(cf)(i＝0，j＝2.3.4.5.7)

若cf＝2，则：工作组内线程数Wgsize(2)＝(16*8)/2＝64，Benefit(0,2)＝(1*(2-1)+3*k*(2-1)+100*k*(2-1))*64＝65984；；

若cf＝4，则：工作组内线程数Wgsize(4)＝(16*8)/4＝32，Benefit(0,4)＝(1*(4-1)+3*k*(4-1)+100*k*(4-1))*32＝98976；

若cf＝8，则：工作组内线程数Wgsize(8)＝(16*8)/8＝16，Benefit(0,8)＝(1*(8-1)+3*k*(8-1)+100*k*(8-1))*16＝115472；

若cf＝16，则：工作组内线程数Wgsize(16)＝(16*8)/16＝8，Benefit(0,16)＝(1*(16-1)+3*k*(16-1)+100*k*(16-1))*8＝123720；

对第1维计算合并收益，方法同第0维的计算方法。

Benefit(1,2)＝72320；

Benefit(1,4)＝108480；

Benefit(1,8)＝126560；

此外，在步骤6中分维度计算代价时，结合硬件平台特征，以线程合并导致的并行度下降和访存效率下降引发的性能损失之和作为代价，cost(i,cf,platform)＝cost_parallelism(i,cf,platform)+cost_efficiency(i,cf,platform)。

对于第i维和平台platform，考察cf取值为candidate_cf_set(i)中的每个候选值cf_p时，分别按照计算并行度影响代价和访存效率影响代价。

3.2.1)首先，计算并行度影响代价，结合以中间文件形式输入、从主机端代码中得的Kernel的执行信息和硬件平台特征分析，包括如下步骤：

A)计算合并后的线程总数和工作组内线程数。在第i维以合并因子值cf进行合并后，第i维的线程数变成原始线程数的1/cf，即globalWorkSize(i)变成globalWorkSize(i)/cf，localWorkSize(i)变成localWorkSize(i)/cf，即总合并维度所对应维度的线程数变成合并前的1/cf,非合并维度的线程数不变。

B)根据目标平台的OpenCL编程手册中描述的最小总线程数以及工作组内或某一维上的最小线程数设定并行度阈值，包括总线程数阈值Threshold_global，工作组内最小线程数阈值Threshold_local和第i维的最小线程数阈值Threshold_i。对于目标平台的OpenCL编程手册中未描述最小总线程数、工作组内或某一维上的最小线程数的，相应阈值则设为IMAX，IMAX为目标平台所能表示的最大正整数。

C)计算合并后的线程总数Sc、合并后的工作组内线程总数Lc、合并后第i维线程数Ic，

Sc＝PRODUCT(globalWorkSize(0),…,globalWorkSize(i)/cf,…,globalWorkSize(n-1))；

Lc＝PRODUCT(localWorkSize(0),…,localWorkSize(i)/cf,…,

localWorkSize(n-1))；

Ic＝localWorkSize(i)/cf。

D)计算并行度影响代价，若Sc<Threshold_global或者Lc<Threshold_local或者Ic<Threshold_i，则cost_parallelism(i,cf,platform)＝IMAX，否则cost_parallelism(i,cf,platform)＝0，即当某平台上在维度i和合并因子cf下的合并导致目标平台不能充分发挥并行处理能力，则相应合并代价为IMAX。

3.2.2)其次，计算访存效率影响代价，访存效率带来的代价以一个工作组内的访存开销的变化值即合并前后的变化值来衡量，以拍数计，该代价与具体硬件平台的线程执行和访存操作的特点以及Kernel代码的访存模式有关，具体包括如下步骤：

先计算合并前的一个工作组内的访存开销，然后计算在不同合并因子cf下的一个工作组内的访存开销，最后对每个合并因子值下以合并后的访存开销减去合并前的访存开销所得差值作为该合并因子下导致的访存效率的代价。对平台platform、在第i维以合并因子cf做考虑，访存效率的代价计算为cost_efficiency(i,cf,platform)＝cost_access(i,cf,platform)-cost_access(i,1,platform)，cf为1即合并前的情况。其中，先计算一个工作组内访存开销的总和即每条访存操作access_j在一个工作组内的执行次数Freq_Wg(access_j,cf)乘以其访存延迟所得乘积值再求和，即cost_access(i,cf,platform)＝∑(Freq_Wg(access_j,cf)*Latency(access_j))。

根据本发明的一个实施例，仍旧以图1中代码为例，其为2维Kernel代码，图2为图1中代码沿第0维以合并因子2进行合并后的代码，图3中代码为图1中代码沿第1维以合并因子2进行合并后的代码。假设Kernel代码的目标平台为Nvidia GPU Tesla C2050芯片(计算能力2.0，14个SMX，共448个核心)，至少可同时执行448个线程，以一个Warp(32个线程)为线程调度的基本单位，则Threshold_global＝448，Threshold_local＝32，Threshold₀＝IMAX，Threshold1＝IMAX，Threshold2＝IMAX，线程总数即globalWorkSize>20000。

对第0维计算并行度影响代价：

若cf＝2，合并后工作组内线程数(8，8)，线程总数Sc〉448，而工作组内线程数8*8>32，cost_parallelism(0,2,C2050)＝0；

若cf＝4，合并后工作组内线程数(4，8)，线程总数Sc〉448，而工作组内线程数4*8＝32，cost_parallelism(0,4,C2050)＝0；

若cf＝8，合并后工作组内线程数(2，8)，线程总数Sc〉448，而工作组内线程数2*8<32，cost_parallelism(0,8,C2050)＝IMAX；

若cf＝16，合并后工作组内线程数(1，8)。线程总数仍〉448，而工作组内线程数1*8<32，cost_parallelism(0,16,C2050)＝IMAX；

对第1维计算并行度影响代价，方法同第0维的计算方法一致。

cost_parallelism(1,2,C2050)＝0；

cost_parallelism(1,4,C2050)＝0；

cost_parallelism(1,8,C2050)＝IMAX。

对于图1代码而言，访存操作只有对global memory的读或写操作(或称指令)。而根据Nvidia-GPU的OpenCL编程手册，对一个warp内的访问globalmemory的指令(包括load/store)并不是来自不同线程的逐条指令依次去访存，而是根据指令中的访存地址自动转成若干条访存事务，故对该平台需要以访存事务的总次数来计算访存开销。比如，在NvidiaGPU Tesla C2050芯片上，访存指令会自动转成访存量为128Byte的若干条访存事务。

对第0维计算访存效率影响代价：

首先，原始代码(即合并前，或称cf＝1)的情况下，一个工作组内16*8个线程即4个warp，对数组A，B，C的访问涉及global memory访问。一个原始线程中，循环体内包含1条访问数组A的语句，其中，线程0访问地址为256(数组A的每个元素为4Byte)，线程1访问地址为256+1*4Byte，…，线程31访问地址为256+1*31*4Byte，这里一个warp内该语句的访问地址没有超过一条访存事务的访存量，对应1条访存事务；还包含一条访问B的语句对应1条访存事务；而循环外有一个访问C的语句即1条访存事务。各条访存事务乘以其执行次数乘以其访问开销得到的乘积值作为访存总开销，即Cost_access(0,1,C2050)＝(4*(1*k*Latency+1*k*Latency+1*Latency))＝8400；

若cf＝2，合并后工作组内线程数(8，8)即2个warp。如图2所示，合并后的新线程内有2条对数组A的访问、1条对数组B的访问、2条对数组C的访问，分别计算访存开销。对1个warp而言，访问数组A的第一条语句，对应2条访存事务(线程0访问地址为256+，线程1访问地址为256+2*4Byte，。。。，线程31访问地址为256+2*31*4Byte，即一个warp内该语句的访问地址超过一条访存事务的访存量，对应2条访存事务)，访问数组A的第二条语句也同样对应2条访存事务；数组B的访问和合并前一样，仍为1条访存事务；数组C的访问有2条语句，第一条语句对应2条访存事务(线程0访问地址为256+，线程31访问地址为256+2*31*4Byte，即一个warp内该语句的访问地址超过一条访存事务的访存量，对应2条访存事务)，第二条语句也同样对应2条访存事务，那么cf＝2的情况下，合并后的访存开销应为Cost_access(0，2,C2050)＝(2*(2*2*k*Latency+1*k*Latency+2*2*Latency))＝10800；因此，在合并因子cf＝2的情况下，其访存效率影响代价为Cost_efficiency(0,2,C2050)＝10800 8400＝2400；

若cf＝4，合并后工作组内线程数(4，8)即1个warp。合并后的新线程内有4条对数组A的访问、1条对数组B的访问、4条对数组C的访问，分别计算访存开销。对1个warp而言，访问数组A的第一条语句，对应4条访存事务(线程0访问地址为256，线程1访问地址为256+4*4Byte，线程31访问地址为256+4*31*4Byte，即一个warp内该语句的访问地址超过一条访存事务的访存量，对应4条访存事务)，访问数组A的第二条语句也同样对应4条访存事务，第三条访问A的语句和第四条语句均对应4条访存事务；数组B的访问和合并前一样，仍为1条访存事务；数组C的访问有4条语句，每条语句对应4条访存事务。那么cf＝4的情况下，合并后的访存开销应为Cost_access(0,4,C2050)＝(1*(4*k*Latency*4+1*k*Latency+4*Latency*4))＝18600；因此，在合并因子cf＝4的情况下，其访存效率影响代价为Cost_efficiency(0,4,C2050)＝18600 8400＝10200；

若cf＝8或者cf＝16，同样的方法计算访存效率的代价。

对第1维计算访存效率的代价，方法同第0维的计算。如图3所示为沿第1维以合并因子2进行合并后的代码。

Cost_access(1,1,C2050)＝8400；

Cost_access(1,2,C2050)＝10800；Cost_efficiency(1,2,C2050)＝2400；

Cost_access(1,4,C2050)＝18600；Cost_efficiency(1,4,C2050)＝10200；

若cf＝8，以同样的方法计算访存效率的代价。

计算出并行度影响代价和访存效率影响代价之后，计算每一维度在不同合并因子下的总代价，即并行度代价与访存效率代价之和。

计算出各个维度在不同合并因子下的总收益和总代价之后，计算得出各个维度的最优合并因子。具体地说，对第i维，在不同合并因子cf下，分别计算纯收益值，纯收益值等于其总收益减去其总代价得到的差值，并行度代价与访存效率代价之和作为总代价。取纯收益值最大的合并因子作为本维度的最优合并因子。

根据本发明的一个实施例，仍旧以图1中代码为例，根据前面的计算结果可知：

Benefit(0,2)＝65984；；

Benefit(0,4)＝98976；

Benefit(0,8)＝115472；

Benefit(0,16)＝1237200；

Benefit(1,2)＝72320；

Benefit(1,4)＝108480；

Benefit(1,8)＝126560；

cost_parallelism(0,2,C2050)＝0；

cost_parallelism(0,4,C2050)＝0；

cost_parallelism(0,8,C2050)＝IMAX；

cost_parallelism(0,16,C2050)＝IMAX；

cost_parallelism(1,2,C2050)＝0；

cost_parallelism(1,4,C2050)＝0；

cost_parallelism(1,8,C2050)＝IMAX。

Cost_efficiency(0,2,C2050)＝10800 8400＝2400；

Cost_efficiency(0,4,C2050)＝18600 8400＝10200；

Cost_efficiency(1,2,C2050)＝2400；

Cost_efficiency(1,4,C2050)＝10200；

每一维度在不同合并因子下的纯收益为相应同一合并因子下的总收益减去总代价之差。根据前面计算的结果可知：

PureB(0,2)＝65984 2400＝63584；

PureB(0,4)＝98976 10200＝88776；

因为cost_parallelism(0,8,C2050)＝IMAX，所以PureB(0,8)<0；

因为cost_parallelism(0,16,C2050)＝IMAX，所以PureB(0,16)<0；

所以，在第0维纯收益在cf＝4时最大，故第0维的最优合并因子为4，纯收益为88776。

根据同样计算第0维最优合并因子同样的方法计算第1维的最优合并因子和纯收益。

PureB(1,2)＝72320-2400＝69920；

PureB(1,4)＝108480 10200＝98280；

因为cost_parallelism(1,8,C2050)＝IMAX，PureB(1,8)<0；

所以，在第1维纯收益在cf＝4时最大，故第1维的最优合并因子为4，纯收益为98280。

在步骤6中计算得出每一维度的最优合并因子基础上，在步骤7中确定总合并维度和总合并因子。即每一个维度都选出最优合并因子后，从所有维度中选择纯收益最大者所对应的合并因子及其对应维度作为总合并维度和总合并因子。需要说明的是，若纯收益最大者为负数，则表明合并的代价大于收益，此时不应合并，则选定的合并维度不变，而合并因子改为1。

根据本发明的一个实施例，仍旧以图1中代码为例，根据前面的计算结果将第0维和第1维最优合并因子对应的纯收益做比较，选择纯收益最大的维度作为总合并维度，其最优合并因子作为总合并因子。此处PureB(1,4)>PureB(0,4),故而总合并维度为1，总合并因子为4。

在步骤7中选出总合并维度和总合并因子后，按照选定的总合并维度和合并因子cf，对Kernel代码内依赖于总合并维度的指令逐条复制(cf-1)遍，对指令中出现的本维度线程号的使用做转换以反映出合并进来的线程，其中，对全局线程号的使用get_global_id(i)转换成线性表达式cf*get_global_id(i)+c_i,c_i＝0,1,…(cf-1),对工作组内线程号的使用get_local_id(i)转换成表达式cf*get_local_id(i)+c_i。并对相关变量扩展成数组或者做变量换名；对不依赖于的指令不做复制，至此得到新的Kernel代码,将上述变换后的新的Kernel代码的中间语言，经代码输出阶段，以OpenCL源码形式输出优化后的Kernel代码，后续编译过程同一般的OpenCL编译过程，即可以调用本地OpenCL编译器，编译生成binary并运行。同时，根据总合并维度和总合并因子，主机端的Kernel执行信息须做相应修改，主要为线程总数(globalWorkSize)和工作组内线程数(localWorkSize)。假定沿第i维以合并因子cf做合并，则线程总数由(globalWorkSize(0),…,globalWorkSize(i),…,globalWorkSize(n-1))变成(globalWorkSize(0),…,globalWorkSize(i)/cf,…,globalWorkSize(n-1))，工作组内的线程数由(localWorkSize(0),…,localWorkSize(i),…,localWorkSize(n-1))变成(localWorkSize(0),…,localWorkSize(i)/cf,…,localWorkSize(n-1))，新的Kernel执行信息以中间文件形式进行记录和输出。

根据本发明的一个实施例，仍旧以图1中代码为例，根据选定的总合并维度1和总合并因子4，对代码进行变换，将依赖于第1维度的代码DependOn(1)＝{insn₀,insn₂,insn₇,insn₈,insn₉}分别复制3遍，对不依赖于第1维的代码UnDependOn(1)＝{insn₁,insn₃,insn₄,insn₅,insn₆}不做复制。其中，insn₂中出现对第1维线程号的使用需要做转换，复制后将有:insn2,insn21,insn22,insn23,将此四条指令中的出现的get_global_id(1)修改成4*get_global_id(1),4*get_global_id(1)+1,…,4*get_global_id(1)+3，以反映出合并进来的线程。Kernel执行信息在合并维度相应地改变、未发生合并的维度不变，即第1维的线程数由globalWorkSize(1)变成globalWorkSize(1)/4，工作组的第1维的线程数由localWorkSize(1)变成localWorkSize(1)/4，将这些Kernel执行信息以中间文件形式进行记录和输出。

需要说明的是，在步骤8中，执行优化同步策略，则是对一个工作组(如OpenCL的work-group)内的线程做完全合并，即将一个WorkGroup内的线程合并为一个线程(即一个Workitem)，同时删除work-group内的所有同步操作。即合并维度为该Kernel的所有维度，每一维度的合并因子为合并前工作组中与该维度对应的维度的线程总数,合并因子为(localWorkSize(0),localWorkSize(1),…,localWorkSize(n-1))即第i维对应的合并因子为localWorkSize(i)，i＝{0,1，…，n-1}，n≤3。。其中，为保证原始代码的指令执行顺序，需要按barrier语句(同步语句)划分指令区域并针对每个区域分别做代码变换，然后删除本Kernel代码中所有barrier语句，完成代码变换。根据本发明的一个实施例，包括如下步骤：

首先，以Kernel内的m条barrier语句作为分界线将Kernel代码划分成(m+1)个指令区域；

其次，对每个区域将该区域内的原始代码包装成一个n层循环的最内层循环体(该Kernel代码为n维则转成循环为n层循环)。该n层循环从最内层(第0层循环)到最外层(第(n-1)层循环)的索引值(index)范围为：第0层循环对应工作组的第0维，index₀取值范围是[0，…，(localWorkSize(0)-1)]，第1层循环对应工作组的第1维，index₁取值范围是[0，…，(localWorkSize(1)-1)]，。。。，第(n-1)层循环对应工作组的第(n-1)维，index_(n-1)取值范围是[0，…，(localWorkSize(n-1)-1)]。对指令中出现的线程号的使用相应做转换以反映出合并进来的线程，其中，get_local_id(i)替换成index_i,get_global_id(i)替换成由index_i计算得到，同时，将相关变量扩展成数组或者做变量换名，以避免变量同名的冲突；

再次，删除本Kernel的所有barrier语句，至此一个工作组内的所有线程合并成为一个线程(即一个work-item)，完成了代码变换。

根据本发明的一个实施例，如图4和图5所示，一个1维Knernel代码，代码含有2条barrier语句，因此将该Kernel代码划分成3个指令区域；将指令区域包装成1层的最内层循环体，每一层循环体的索引值为[0，…，localWorkSize(0)-1]。

本发明通过自动选择优化策略以及分维度进行线程间冗余删除的收益代价分析，能自动根据代码是否含有同步操作和硬件平台的特点，针对性地选择一种优化策略，通过对Kernel代码的编译分析，得出合适的合并维度和合并因子，利于提高Kernel代码的运行性能。采用本发明方法能面向多种异构平台，进行自动分析和变换，得以优化设备端代码的线程间冗余和同步开销，提高设备端代码(即OpenCL的Kernel代码)的性能。

需要说明的是，虽然上文按照特定顺序描述了各个步骤，但是并不意味着必须按照上述特定顺序来执行各个步骤，实际上，这些步骤中的一些可以并发执行，甚至改变顺序，只要能够实现所需要的功能即可。

本发明可以是系统、方法和/或计算机程序产品。计算机程序产品可以包括计算机可读存储介质，其上载有用于使处理器实现本发明的各个方面的计算机可读程序指令。

计算机可读存储介质可以是保持和存储由指令执行设备使用的指令的有形设备。计算机可读存储介质例如可以包括但不限于电存储设备、磁存储设备、光存储设备、电磁存储设备、半导体存储设备或者上述的任意合适的组合。计算机可读存储介质的更具体的例子(非穷举的列表)包括：便携式计算机盘、硬盘、随机存取存储器(RAM)、只读存储器(ROM)、可擦式可编程只读存储器(EPROM或闪存)、静态随机存取存储器(SRAM)、便携式压缩盘只读存储器(CD-ROM)、数字多功能盘(DVD)、记忆棒、软盘、机械编码设备、例如其上存储有指令的打孔卡或凹槽内凸起结构、以及上述的任意合适的组合。

以上已经描述了本发明的各实施例，上述说明是示例性的，并非穷尽性的，并且也不限于所披露的各实施例。在不偏离所说明的各实施例的范围和精神的情况下，对于本技术领域的普通技术人员来说许多修改和变更都是显而易见的。本文中所用术语的选择，旨在最好地解释各实施例的原理、实际应用或对市场中的技术改进，或者使本技术领域的其它普通技术人员能理解本文披露的各实施例。

Claims (14)

1.一种线程间冗余删除的异构编译优化方法，其特征在于，包括如下步骤：

S1、获取Kernel源代码及Kernel执行信息；

S2、对所述Kernel源代码及Kernel执行信息，以工作组为单位，分维度识别出候选冗余计算；

S3、基于所述候选冗余计算分维度进行收益代价分析,得到每一维度的最优合并因子及其优化冗余的纯收益；

S4、基于每一维度的最优合并因子及其优化冗余的纯收益确定总合并维度及总合并因子；

S5、基于总合并维度及总合并因子，对Kernel代码进行线程优化冗余操作，合并线程，进行代码变换和优化；

S6、输出优化后的源码形式的Kernel代码及中间文件形式的Kernel执行信息。

2.根据权利要求1所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

所述步骤S2包括如下步骤：

S21、以工作组为单位，分维度对Kernel代码的中间语言进行指令之间的依赖分析，得出依赖于每一维度线程的普通指令以及不依赖于该维度线程的普通指令；普通指令为同步指令以外的其他指令；

S22、将不依赖于线程的普通指令视为该线程所在维度的候选冗余计算。

3.根据权利要求2所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

所述步骤S3包括如下步骤：

S31、在不同合并因子下，以工作组为单位，分维度计算优化冗余的收益；

S32、结合硬件平台特征，以工作组为单位，在不同合并因子下，分维度计算优化冗余的代价；

S33、针对同一工作组同一维度线程，在不同合并因子下，用步骤S31中计算的收益与步骤S32中计算的同一维度同一合并因子下的代价相减，得到该维度线程在不同合并因子下的纯收益，取其中最大纯收益所对应的合并因子作为该维度线程的最优合并因子；

S34、重复步骤S31-S33，直至计算得到每一维度线程的最优合并因子及其对应的纯收益为止。

4.根据权利要求3所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

所述步骤S4包括如下步骤：

S41、获取每一维度线程的最优合并因子及其对应的纯收益；

S42、对所有维度的最优合并因子对应的纯收益进行对比，选出纯收益最大者所对应的维度作为总合并维度，并将该维度的最优合并因子作为总合并因子。

5.根据权利要求4所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

所述步骤S5包括如下步骤：

S51、按照选出的总合并维度和总合并因子，对Kernel代码内依赖于总合并维度的普通指令逐条复制，复制的遍数为总合并因子减一，对不依赖于总合并维度的普通指令不做复制；并转换代码中对线程号的使用；

S52、对复制后的指令中的相关变量做扩展或者更换变量名，避免变量冲突，得到变换后的新的Kernel代码；

S53、根据总合并维度和总合并因子，相应修改主机端相关代码，包括分维度的线程总数和工作组内每一维线程数。

6.根据权利要求3所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

在不同合并因子下，以工作组为单位，分维度计算优化冗余的收益时，不同合并因子下每一维度的收益为不依赖于该维度的所有普通指令在合并因子下做优化冗余操作时为工作组带来的收益之和；

不同合并因子下，每一不依赖于线程的普通指令的收益为该指令执行频率、指令延迟、合并因子减一之差这三者的乘积。

7.根据权利要求6所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

在不同合并因子下，以工作组为单位，分维度计算优化冗余的代价时，以线程合并导致的并行度下降和访存效率下降引发的性能损失之和作为代价，包括并行度影响代价和访存效率的代价。

8.根据权利要求7所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

在不同合并因子下，以工作组为单位，分维度计算并行度影响代价；

其中，并行度影响代价为0或者IMAX，IAMX为目标平台所能表示的最大正整数；以及

当平台上合并导致目标平台不能充分发挥并行处理能力时，并行度影响代价为IMAX，否则为0。

9.根据权利要求8所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

在不同合并因子下，以工作组为单位，分维度计算访存效率的代价；

其中，访存效率影响代价以一个工作组内的访存开销在做优化冗余合并操作前后的变化值来衡量；每个合并因子值下以合并后的工作组线程的访存开销减去合并前工作组线程的访存开销所得差值作为该合并因子下导致的访存效率的代价。

10.根据权利要求1所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

在步骤S1中，在获取Kernel源代码后，对Kernel代码做词法或语义分析，并根据分析结果选择优化策略；其中当Kernel代码含有同步指令且目标平台同步为CPU平台或其他以软件实现同步操作的平台时，选择优化同步策略，否则执行步骤S2。

11.根据权利要求10所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

优化同步策略包括如下步骤：

R1、对Kernel代码的所有维度进行合并，每一维度的合并因子为工作组内相应维度的线程总数，合并后，Kernel程序中的每一个工作组内的线程为1个线程；

R2、删除工作组内的所有同步操作，实现同步优化；

R3、根据步骤R2的同步优化，进行代码变换；

R4、输出步骤R3中代码变换后的结果。

12.根据权利要求11所述的一种线程间冗余删除的异构编译优化方法，其特征在于，

所述步骤R3,包括如下步骤：

R31、针对每一指令区域分别做代码变换，以同步语句为分界线，将Kernel代码划分成指令区域，指令区域的数量为同步语句数量加1；

R32、对每个指令区域的源代码进行变换，将其包装为一个含有至少一层循环的最内层循环体，循环体层数与Kernel代码维度的维数相同，每一层循环体对应一个维度，每一层循环体的内循环次数为其对应维度的线程总数；并转换代码中对线程号的使用；

R33、扩展相应变量或修改变量，避免变量冲突。

13.一种计算机可读存储介质，其上存储有计算机程序，其中，该程序被处理器执行时实现根据权利要求1至12中任一项所述方法的步骤。

14.一种计算机设备，包括存储器和处理器，在所述存储器上存储有能够在处理器上运行的计算机程序，其特征在于，所述处理器执行所述程序时实现权利要求1至12中任一项所述的方法的步骤。