CN115934102A - 通用寄存器动态分配方法、装置、计算机设备和存储介质 - Google Patents
通用寄存器动态分配方法、装置、计算机设备和存储介质 Download PDFInfo
- Publication number
- CN115934102A CN115934102A CN202211705663.4A CN202211705663A CN115934102A CN 115934102 A CN115934102 A CN 115934102A CN 202211705663 A CN202211705663 A CN 202211705663A CN 115934102 A CN115934102 A CN 115934102A
- Authority
- CN
- China
- Prior art keywords
- registers
- general
- register
- machine instruction
- general registers
- 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.)
- Granted
Links
- 238000000034 method Methods 0.000 title claims abstract description 122
- 230000008569 process Effects 0.000 claims abstract description 74
- 230000015654 memory Effects 0.000 claims description 97
- 238000012545 processing Methods 0.000 claims description 50
- 230000001360 synchronised effect Effects 0.000 claims description 21
- 238000004590 computer program Methods 0.000 claims description 20
- 238000006243 chemical reaction Methods 0.000 claims description 3
- 230000008094 contradictory effect Effects 0.000 abstract description 3
- 238000010586 diagram Methods 0.000 description 11
- 238000004891 communication Methods 0.000 description 4
- 238000005516 engineering process Methods 0.000 description 3
- 238000000638 solvent extraction Methods 0.000 description 3
- 230000004888 barrier function Effects 0.000 description 2
- 230000005540 biological transmission Effects 0.000 description 2
- 238000004422 calculation algorithm Methods 0.000 description 2
- 230000008859 change Effects 0.000 description 2
- 101150037603 cst-1 gene Proteins 0.000 description 2
- 230000006870 function Effects 0.000 description 2
- 238000013507 mapping Methods 0.000 description 2
- 238000005192 partition Methods 0.000 description 2
- OKTJSMMVPCPJKN-UHFFFAOYSA-N Carbon Chemical compound [C] OKTJSMMVPCPJKN-UHFFFAOYSA-N 0.000 description 1
- 238000004458 analytical method Methods 0.000 description 1
- 238000004364 calculation method Methods 0.000 description 1
- 230000015556 catabolic process Effects 0.000 description 1
- 230000001413 cellular effect Effects 0.000 description 1
- 238000006731 degradation reaction Methods 0.000 description 1
- 238000013461 design Methods 0.000 description 1
- 238000007667 floating Methods 0.000 description 1
- 229910021389 graphene Inorganic materials 0.000 description 1
- 238000003384 imaging method Methods 0.000 description 1
- 239000004973 liquid crystal related substance Substances 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 230000004048 modification Effects 0.000 description 1
- 230000003287 optical effect Effects 0.000 description 1
- 230000003068 static effect Effects 0.000 description 1
- 230000000007 visual effect Effects 0.000 description 1
Images
Classifications
-
- 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
- Devices For Executing Special Programs (AREA)
Abstract
本申请涉及一种通用寄存器动态分配方法、装置、计算机设备和存储介质。所述方法包括：若机器指令程序在执行过程中满足寄存器重新划分条件，则确定机器指令程序在执行过程中通用寄存器的实际使用数量；若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。采用本方法能够解决GPU处理器启动线程数目和通用寄存器资源相互矛盾的问题，可以不经过重编译而直接对编译器的工作组大小进行重新划分，从而减少编译器的编译时间。
Description
技术领域
本申请涉及通用寄存器分配技术领域，特别是涉及一种通用寄存器动态分配方法、装置、计算机设备和存储介质。
背景技术
OpenCL(Open Computing Language)是一个开放，免版税的面向异构计算系统的并行编程标准。OpenCL程序主要包括主机端代码和设备端代码两个部分，其中主机端代码由C/C++编译器编译生成主机端的可执行程序，设备端代码是由OpenCL C编译器编译生成特定设备的可执行目标代码。通常，OpenCL编译器指的是用来编译设备端Kernel代码的编译器工具。
在OpenCL编译器中，传统寄存器分配方法是设置寄存器的最大可使用数量为一个固定值，并且只分配一次。针对复杂例程，传统寄存器分配方法会存在寄存器不够用的情况，通常做法是把寄存器中的变量先暂存至内存中，使用时再从内存中取出，然而从内存读写暂存数据需要消耗很长的时间，极大降低设备端的性能。而且，对于特定的OpenCLKernel程序，由于工作组大小设置的不尽合理，导致硬件资源不能充分的利用，进而导致设备端硬件平台的执行效率不能充分发挥。
发明内容
基于此，有必要针对上述技术问题，提供一种能够动态分配通用寄存器的最大可使用数量，使得单个线程中使用的寄存器数目和该线程对应的计算单元启动的线程数目更合理，提高设备端硬件在执行并行计算代码时的执行效率的通用寄存器动态分配方法、装置、计算机设备、计算机可读存储介质和计算机程序产品。
第一方面，本申请提供了一种通用寄存器动态分配方法，所述方法包括：
通过编译器将源代码转换成机器指令程序；
若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
在其中一个实施例中，所述方法还包括：
若机器指令程序在执行过程中不满足寄存器重新划分条件，则将通用寄存器的最大可分配数量设置为固定值，根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量。
在其中一个实施例中，寄存器重新划分条件包括：工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件。
在其中一个实施例中，调整通用寄存器的最大可分配数量，包括：
确定通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及机器指令程序执行完成时所需的时钟数量的第二理论值；
若第一理论值与第二理论值的比值大于预设值，则以固定的步进值增加通用寄存器的最大可分配数量，返回执行根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，直至第一理论值与第二理论值的比值小于预设值。
在其中一个实施例中，所述方法还包括：
若根据通用寄存器的最大可分配数量以及通用寄存器的实际使用数量，判定GPU处理器中处理单元不满足启动线程条件，则通过编译器重新设置通用寄存器的最大可分配数量，返回执行通过编译器将源代码转换成机器指令程序的步骤，直至处理单元满足启动线程条件。
在其中一个实施例中，启动线程条件包括通用寄存器的最大可分配数量与通用寄存器的实际使用数量的比值、大于工作组的大小与线程粒度之间的比值。
在其中一个实施例中，根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，重新设定工作组的大小，包括：
根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，确定GPU处理器中处理单元可启动的线程数量范围；
取线程数量范围中任一数量，作为处理单元的目标启动线程数量；
根据目标启动线程数量，确定多个第一基元，取多种第一基元中的任一种第一基元，作为第一目标基元；第一基元包括三个维度，三个维度的数值的乘积等于目标启动线程数量；
根据线程粒度，确定多个第二基元，取多种第二基元中的任一种第二基元，作为第二目标基元；第二基元包括三个维度，三个维度的数值的乘积等于线程粒度；
将第一目标基元与第二目标基元的对应维度上的数值乘积，确定为工作组在不同维度的目标大小。
第二方面，本申请还提供了一种通用寄存器动态分配装置。所述装置包括：
转换模块，用于通过编译器将源代码转换成机器指令程序；
分配模块，用于若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
调整模块，用于若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
工作组划分模块，用于根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
第三方面，本申请还提供了一种计算机设备。所述计算机设备包括存储器和处理器，所述存储器存储有计算机程序，所述处理器执行所述计算机程序时实现以下步骤：
通过编译器将源代码转换成机器指令程序；
若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
第四方面，本申请还提供了一种计算机可读存储介质。所述计算机可读存储介质，其上存储有计算机程序，所述计算机程序被处理器执行时实现以下步骤：
通过编译器将源代码转换成机器指令程序；
若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
第五方面，本申请还提供了一种计算机程序产品。所述计算机程序产品，包括计算机程序，该计算机程序被处理器执行时实现以下步骤：
通过编译器将源代码转换成机器指令程序；
若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
上述通用寄存器动态分配方法、装置、计算机设备和存储介质，针对满足工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件的机器指令程序，根据机器指令程序执行过程中是否使用寄存器溢出指令，动态分配通用寄存器的最大可分配数量，可以平衡一个处理单元中启动的硬件线程数目和一个硬件线程中使用通用寄存器数目之间的关系，以及解决GPU处理器启动线程数目和通用寄存器资源相互矛盾的问题，提高GPU处理器执行效率；同时，针对满足工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件的机器指令程序，可以不经过重编译而直接对编译器的工作组大小进行重新划分，从而有效减少编译器的编译时间。
附图说明
图1为一个实施例中通用寄存器动态分配方法的应用环境图；
图2为一个实施例中OpenCL内存模型与GPU处理器的存储关系对照关系示意图；
图3为一个实施例中通用寄存器动态分配方法的流程示意图；
图4为另一个实施例中编译器和通用寄存器动态分配流程示意图；
图5为一个实施例中调整通用寄存器的最大可分配数量的流程示意图；
图6为一个实施例中生成目标程序的流程示意图；
图7为一个实施例中重新设定工作组的大小的流程示意图；
图8为一个实施例中编译器执行流程图；
图9为一个实施例中通用寄存器动态分配装置的结构框图；
图10为一个实施例中计算机设备的内部结构图。
具体实施方式
为了使本申请的目的、技术方案及优点更加清楚明白，以下结合附图及实施例，对本申请进行进一步详细说明。应当理解，此处描述的具体实施例仅仅用以解释本申请，并不用于限定本申请。
本申请实施例提供的一种通用寄存器动态分配方法，可以应用于如图1所示的应用环境中。其中，终端102通过编译器104将源代码转换成机器指令程序；若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。编译器104将机器指令程序转换为目标程序，将目标程序发放给GPU处理器106，GPU处理器106执行目标程序对应的任务。其中，终端102可以但不限于是包括编译器104和GPU处理器106的各种个人计算机、笔记本电脑、智能手机、平板电脑、物联网设备和便携式可穿戴设备。
如图2所示，OpenCL编译器的内存模型包括全局内存、常量内存、局部内存和私有内存，这四种内存分别代表了不同的内存区域，内存空间都与OpenCL Kernel源代码有关。一个内核中，可能会有global(全局)、constant(常量)、local(局部)或private(私有)关键字，不同关键字用来指定变量使用特定的内存进行创建数据的存储空间。从OpenCL内存模型层面考虑，这些内存区域在逻辑上是不相交的，并且不同区域的数据要被其他区域使用或进行数据搬移是由Kernel源代码控制的，每个内存区域都有其各自的性能特性。这些不同层级的存储划分在读取时具有很大的性能差异。其中每个内存区域的描述如下文所示：
全局内存，对于执行内核中的所有工作项(work item)都是可见的。数据从主机端和设备端的传输需要存储在全局内存中才能进行，对应存储在本发明中GPU的显卡内存(Video Memory)中。其关键字为global或_global，Kernel源代码中global int*A代表A指针指向的数据放在全局内存中。其中，工作项(work item)指一个循环中最里面的一次运算，称为一个工作项，访问相同处理资源的工作项组成一个工作组。能够支持工作组的处理资源被称为处理单元。各个工作组都可以在单个处理单元上的执行。由于工作项之间需要进行同步化处理，因此需要对工作项进行工作组划分，因为只有工作项属于同一个工作组，才能够实现工作项同步化。
常量内存，并非为只读数据设计，但其能让计算空间中所有工作项同时对同一数据进行访问。常量内存存储的值通常不会变化。Kernel源代码中使用关键字constant或_constant将相应的数据映射到常量内存。一般情况，OpenCL编译器根据Kernel源代码中使用的常量内存的不同，把常量映射到GPU处理器的常量缓冲区(Constant Buffer)或显卡内存(Video Memory)中，当然GPU处理器对存储在Constant Buffer中的数据读取的更快。
局部内存，存储在局部内存中的数据，只有在同一工作组(work group)内的工作项可以共享。局部内存会映射到GPU处理器中的顶点缓冲区(VB)或共享内存(ShareMemory)，比起全局内存，局部内存具有更短的访问延迟，以及更高的传输带宽。调用clSetKernelArg()设置局部内存时，只需要传递大小，而无需传递对应的指针，相应的局部内存会在运行时在设备端进行开辟存储空间。OpenCL内核中，使用local或_local关键字来描述局部内存。
私有内存，只能由工作项自己进行访问。实践中，私有变量通常都与通用寄存器对应，当通用寄存器不够私有数据使用时，会产生数据溢出现象，溢出数据通常会存储到全局内存中。需要注意的是：本实施例中对于Kernel源代码中使用private或_private关键字的变量使用的是全局内存进行存储，对应的是图3右侧GPU处理器的Video Memory中。
其中，图3主要介绍OpenCL的内存模型以及Kernel源代码中不同层级的内存在GPU处理器中的映射关系。本实施例采用图3右侧的GPU内存架构，在该GPU架构中，Kernel中的一个工作组(Work Group)只能配置在同一个处理单元(PE)中，那么在该架构的GPU中一个工作组的最大值为处理单元(PE)中最多可执行线程数目N x SIMD32或者N x SIMD64。
在一个实施例中，如图3所示，提供了一种一种通用寄存器动态分配方法，以该方法应用于图1中的终端为例进行说明，包括以下步骤：
步骤302，通过编译器将源代码转换成机器指令程序。
其中，编译器指用来编译设备端Kernel代码的编译器工具。源代码可以是设备端Kernel代码。机器指令程序指CPU处理器能直接识别并执行的指令，机器指令程序的表现形式是二进制编码。
可选地，终端通过驱动程序启动一个进程调用OpenCL编译器，OpenCL编译器将源代码转换为中间表示程序，中间表示程序为编译器对于源代码进行扫描后生成的内部表示，代表源代码的语义和语法结构；对中间表示程序进行指令降级和指令选择处理，将中间表示程序转换为机器指令程序。
步骤304，若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量。
其中，通用寄存器(CRF)是GPGPU中有限存储容量的高速存储部件，是每个线程执行的临时存储器，用于记录指令执行的原始数据和结果。通常一个计算单元或计算核心中只有几十KByte(2x512bitxM)的存储容量或者表述为只有M个通过寄存器，而该计算单元能够启动的硬件线程(Wave)数量N是动态变化的为1-nmax(一般可启动1～32个线程)，那么每个线程可以使用的通用寄存器空间为M/N，由于M是固定值，那么当线程数目越多，每个线程可以使用的通用寄存器就越少。对于一段Kernel代码，OpenCL编译器如何分配通用寄存器，使得单个线程中实际使用的寄存器数目m和该计算单元启动的线程数目n更合理，是提升硬件执行效率的关键一环。
通用寄存器分配是通过将程序变量尽可能地分配到寄存器，从而提高程序执行速度的一种方法。通用寄存器分配完成后，每一个被分配或使用的通用寄存器都有一个唯一的编号，可以通过统计机器指令程序中通用寄存器的最大编号值，将最大编号值换算为通用寄存器的实际使用数量mtotal。
在一些实施例中，若机器指令程序在执行过程中不满足寄存器重新划分条件，则将通用寄存器的最大可分配数量mmax设置为固定值，根据通用寄存器的最大可分配数量mmax进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量。其中，寄存器重新划分条件包括：工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件。
可选地，如图4所示，终端从机器指令程序中分析Kernel源代码对应的机器指令程序在执行过程中满足寄存器重新划分条件，即未出现使用局部内存语句(_local关键字)、同步语句(barrier()或sync())和将工作组索引信息参与除全局工作组索引之外的其他运算中任一种情况，则此时对应通用寄存器分配前的初始化工作，首先第一次设定通用寄存器最大可分配的数目mmax，同时复制一份机器指令作为备份，根据初始化的通用寄存器的最大可分配数量mmax进行寄存器分配操作，并统计机器指令程序中通用寄存器的最大编号值，将最大编号值换算为通用寄存器的实际使用数量mtotal；
若机器指令程序在执行过程中不满足寄存器重新划分条件，即出现使用局部内存语句(_local关键字)、同步语句(barrier()或sync())和将工作组索引信息参与除全局工作组索引之外的其他运算中至少一种情况，则将通用寄存器的最大可分配数量设置为固定值，根据通用寄存器的最大可分配数量mmax进行一次寄存器分配操作，设定一个状态位，用来标记机器指令程序中是否使用局部内存语句或同步语句或将工作组索引信息参与除全局工作组索引之外的其他运算，根据状态位判定是否需要重新设置工作组，如果需要设置工作组，则设定为true；如果不需要设置工作组，则为false。同时确定机器指令程序在执行过程中通用寄存器的实际使用数量mtotal。
步骤306，若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令。
其中，OpenCL的Kernel源代码通过编译器转化为机器指令程序后，机器指令程序中的指令在GPU中被执行的基础是以线程为尺度来衡量的，因此需要特别关注编译器生成的机器指令程序中使用的通用寄存器的数量和GPU处理器设备中可启动的线程数量，如果一味追求GPU处理器中中能够启动最大nmax个线程，那么此时1个线程中能够使用的通用寄存器数量就最小为mmax/nmax，因此这些指标需要提前告知编译器设计人员，编译器设计人员在设计编译器时，在寄存器分配阶段需要固定可用物理通用寄存器数量的上限为mmax/nmax，如果此时Kernel源代码中的算法比较复杂，mmax/nmax个通用寄存器就不能满足可分配寄存器的需要，这时候需要使用寄存器溢出(Spill)技术，即先把一些放在通用寄存器中的还未使用的临时变量或者已经使用并需要留待使用的临时变量存储到全局内存中，等到使用时再从全局内存中读取并写回通用寄存器中。虽然这时候保证能够启动数量最多的线程，但是每个线程在执行时额外多出的全局内存的读写需要耗费大量时间，从而大大增加了启动线程后执行完所有指令所需要的时间。
对于不同Kernel源代码，编译器在通用寄存器分配时能够动态调整可使用通用寄存器的数量变得十分重要。同时，根据OpenCL的标准-局部存储(local memory)的数据只可以在同一工作组中的不同工作项之间共享，可以从理论上知道在未使用局部内存语句或同步语句的Kernel源代码中，可以对工作组的大小进行重新设定，这样可以更灵活的调整工作组大小和一个处理单元(PE)中启动的线程数目以及一个线程中使用通用寄存器的数量的关系。
可选地，如图4所示，若机器指令程序执行过程中使用寄存器溢出指令，则说明通用寄存器的最大可分配数量mmax不能满足可分配寄存器的需要，调整初始化的通用寄存器的最大可分配数mmax量，再根据调整后的通用寄存器的最大可分配数量mmax进行寄存器分配，再次判断机器指令程序执行过程中是否使用寄存器溢出指令，若机器指令程序执行过程中使用寄存器溢出指令，则需要在第一次调整的基础上再次调整通用寄存器的最大可分配数量mmax，返回执行根据调整后的通用寄存器的最大可分配数量进行寄存器分配，直至机器指令程序执行过程中未使用寄存器溢出指令，设定一个状态位，用来标记机器指令程序中是否使用局部内存语句或同步语句或将工作组索引信息参与除全局工作组索引之外的其他运算，根据状态位判定是否需要重新设置工作组，如果需要设置工作组，则设定为true；如果不需要设置工作组，则为false。同时确定机器指令程序在执行过程中通用寄存器的实际使用数量mtotal。
步骤308，根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
其中，通过分析程序中使用的内存模式，对未使用局部内存语句或同步语句或将工作组索引信息参与除全局工作组索引之外的其他运算的情况，可以重新分割工作组大小，从而达到每个硬件线程合理使用通用寄存器的目的，同时会减少重编译。
可选地，终端根据通用寄存器的实际使用数量mtotal和调整后通用寄存器的最大可分配数量mmax，计算出一个处理单元(PE)中此时可启动的最大线程数量，根据可启动的最大线程数量重新设定编译器中工作组的大小。
上述通用寄存器动态分配方法中，针对满足工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件的机器指令程序，根据机器指令程序执行过程中是否使用寄存器溢出指令，动态分配通用寄存器的最大可分配数量mmax，可以平衡一个处理单元中启动的硬件线程数目和一个硬件线程中使用通用寄存器数目之间的关系，以及解决GPU处理器启动线程数目和通用寄存器资源相互矛盾的问题，提高GPU处理器执行效率；同时，针对满足工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件的机器指令程序，可以不经过重编译而直接对编译器的工作组大小进行重新划分，从而有效减少编译器的编译时间。
在一个实施例中，通常情况下，通用寄存器的分配方式是一次编译流程的一次分配的模式，即在编译器中设置最大可分配的通用寄存器数量mmax，通用寄存器分配阶段最多能分配的通用寄存器的数量为mmax，如果Kernel源代码逻辑十分复杂，mmax个通用寄存器不足以满足可分配的需要，那么此时会对机器指令的基本块做拆分(Split)以减少通用寄存器的使用，如果拆分还不能满足需要，那么此时会使用通用寄存器Spill技术，即把暂时不用的变量先存入内存，等待使用时再从内存中读出并写回到通用寄存器中使用。因此，若机器指令程序执行过程中使用寄存器溢出指令，则说明需要动态调整通用寄存器的最大可分配数量mmax，以使得机器指令程序执行过程中未使用寄存器溢出指令。具体的，如图5所示，调整通用寄存器的最大可分配数量包括：
步骤502，确定通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及机器指令程序执行完成时所需的时钟数量的第二理论值。
可选地，终端根据GPU处理器运算速度和时钟频率，确定通用寄存器溢出指令执行时所需的时钟数量的第一理论值Cspill，以及机器指令程序执行完成时所需的时钟数量的第二理论值Ctotal。
步骤504，若第一理论值与第二理论值的比值大于预设值，则以固定的步进值增加通用寄存器的最大可分配数量，返回执行根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，直至第一理论值与第二理论值的比值小于预设值。
可选地，若第一理论值Cspill与第二理论值Ctotal的比值大于预设值，则说明需要增大最大可分配寄存器数量以减少寄存器溢出操作，终端将一个固定的步进值mstep与当前的最大可分配寄存器数量mmax之和，确定为重新设定的最大可分配寄存器数量，即mmax＝mmax+mstep；返回执行根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，第一理论值Cspill与第二理论值Ctotal的比值小于预设值，则无需增大最大可分配寄存器数量。
本实施例中，通过通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及机器指令程序执行完成时所需的时钟数量的第二理论值之间的比值，判定当前的最大可分配寄存器数量是否满足需求，并在需要增大最大可分配寄存器数量以减少寄存器溢出操作时，以固定的步进值增加通用寄存器的最大可分配数量mmax，返回编译器重新编译，直至第一理论值与第二理论值的比值小于预设值，即无需增大最大可分配寄存器数量，采用多次动态寄存器分配的方式，通过分析机器指令程序中使用的内存模式和生成的指令间的依赖关系以及每条指令执行需要的时间，从而合理设定寄存器分配时的最大的可用寄存器数量，从而实现减少或不使用暂存临时变量到内存中的方法，可以极大减少执行程序所花费的时间。
在一个实施例中，由OpenCL编程标准和GPU处理器特性可知，工作组越大，启动工作组需要用到的通用寄存器数目越多。对于0×400大小的工作组，通过计算可得到GPU处理器理论上需要32个线程才能启动一个工作组。由此可见，工作组的大小可以直接影响处理单元的可启动的最大线程数量，以及一个处理单元所需的通用寄存器的最大可分配数量mmax。如图6所示，根据上述实施例的状态位判定工作组是否需要重新划分，若需要重新划分，则根据通用寄存器的实际使用数量mtotal和调整后的通用寄存器的最大可分配数量mmax，重新设定工作组的大小，并通过指令发射流程生成目标程序，其中，目标程序为GPU处理器可以识别并处理的二进制文件，将目标程序和状态位为true返回编译器；若不需要重新划分，则通过指令发射流程生成目标程序，将目标程序和状态位为false返回编译器。
如图7所示，根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，重新设定工作组的大小，包括：
步骤702，根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，确定GPU处理器中处理单元可启动的线程数量范围。
可选地，调整后的通用寄存器的最大可分配数量mmax与通用寄存器的实际使用数量mtotal之间的比值，为处理单元可启动的最大线程数量nmax，因此，GPU处理器中处理单元可启动的线程数量范围为[1-nmax]。
步骤704，取线程数量范围中任一数量，作为处理单元的目标启动线程数量。
步骤706，根据目标启动线程数量，确定多个第一基元，取多种第一基元中的任一种第一基元，作为第一目标基元；第一基元包括三个维度，三个维度的数值的乘积等于目标启动线程数量。
其中，第一基于指用于体现工作组不同维度大小的组合，通常是指xyz方向上的大小。将目标启动线程数量表示为N，第一基元表示为(N1，N2，N3)，则N＝N1*N2*N3。
例如，线程数量范围为[1-12]，目标启动线程数量为12，则第一基元中N1、N2和N3可以12、1和1组成的6种三元组，也可以是2、2和3组成的6种三元组。
步骤708，根据线程粒度，确定多个第二基元，取多种第二基元中的任一种第二基元，作为第二目标基元；第二基元包括三个维度，三个维度的数值的乘积等于线程粒度。
其中，线程粒度用于反映一个线程的计算量大小，一般是常量32。将第二基元表示为(x_cst，y_cst，z_cst)，则32＝x_cst*y_cst*z_cst，第二基元中x_cst、y_cst和z_cst可以36、1和1组成的6种三元组，可以是1、4和8组成的6种三元组，可以是2、2和8组成的6种三元组，可以是1、2和16组成的6中三元组，还可以是2、4和4组成的6中三元组，在第二基元中任一种第二基元，作为第二目标基元。
步骤710，将第一目标基元与第二目标基元的对应维度上的数值乘积，确定为工作组在不同维度的目标大小。
其中，工作组重新划分策略为：重新划分的工作组大小＝{x_cst*N1，y_cst*N2，z_cst*N3}，其中，N＝N1*N2*N3，N为目标启动线程数量。本实施例以第一基元为(12，1，1)，第二基元为(32，1，1)为例，则重新划分的工作组大小{32*12，0×1，0×1}，重新划分的工作组的大小范围为{32×N，0×1，0×1}，N为[1-12]。需要注意的是：{32×N，0×1，0×1}中0×1不是数字0乘以1，而是16进制表示方式，0×1就是1。
本实施例中，根据通用寄存器的实际使用数量mtotal和调整后的通用寄存器的最大可分配数量mmax，确定GPU处理器中处理单元可启动的线程数量范围，根据线程数量范围确定多个第一基元，根据线程粒度确定多个第二基元，任取一个第一基元和一个第二基元作为第一目标基元和第二目标基元，根据第一目标基元和第二目标基元对应位置的乘积，确定重新划分的工作组大小。采用上述方法确定工作组的大小，可以平衡工作组大小与工作组所需的通用寄存器数量之间的关系，保证每个硬件线程合理使用通用寄存器的目的，同时会减少重编译。
在一个实施例中，机器指令程序经过上述实施例在满足或不满足寄存器重新划分条件时，都会设定一个状态位，用true表示需要设置工作组，用false表示不需要设置工作组。上述实施例中，将目标程序和状态位返回编译器，编译器根据状态位执行不同的处理流程。如图8所示，若编译器根据状态位确定需要重新设置工作组，则编译器将当前的工作组大小调整至上述实施例计算得到的重新划分后的工作组大小，并请求GPU处理器执行目标程序对应的任务。若编译器根据状态位确定不需要重新设置工作组，则编译器根据通用寄存器的最大可分配数量mmax以及通用寄存器的实际使用数量mtotal，判定GPU处理器中处理单元是否满足启动线程条件。若判定GPU处理器中处理单元不满足启动线程条件，则通过编译器重新设置通用寄存器的最大可分配数量，返回执行通过编译器将源代码转换成机器指令程序的步骤，直至处理单元满足启动线程条件。若判定GPU处理器中处理单元满足启动线程条件，则请求GPU处理器执行目标程序对应的任务。
其中，启动线程条件包括通用寄存器的最大可分配数量mmax与通用寄存器的实际使用数量mtotal的比值、大于工作组的大小与线程粒度之间的比值，即：
mmax/mtotal>(工作组/32)。
需要注意的是：工作组的大小分两种模式，一种是对齐模式，另一种是非对齐模式，工作组的模式与编译器设计人员的编写方向相关。将初始工作组的大小标记为：工作组＝{x_ori，y_ori，z_ori}。其中，对齐模式下的工作组大小为：
工作组＝(((x_ori+x_cst-1)/x_cst)*x_cst)*(((y_ori+y_cst-1)/y_cst)*y_cst)
*(((z_ori+z_cst-1)/z_cst)*z_cst)
非对齐模式下的工作组大小为：
工作组＝(x_ori*y_ori*z_ori+32-1)/32*32。
需要注意的是：在不需要重新划分工作组时，针对不满足寄存器重新划分条件，并且不满足启动线程条件的机器指令程序，会出现当前工作组大小不能启动线程的情况，因此，出现当前工作组大小不能启动线程的情况时，重新设定编译Kernel可使用的最大寄存器数量，进行重新编译或发送当前工作组超出最大工作组范围并退出。
本实施例中，在不需要重新划分工作组时，进一步判断GPU处理器中处理单元是否满足启动线程条件，并在GPU处理器中处理单元不满足启动线程条件时，重新设定编译Kernel可使用的最大寄存器数量，进行重新编译或发送当前工作组超出最大工作组范围并退出，保证GPU处理器的执行效率。
在一个实施例中，Kernel源代码包括主机端代码和设备端内核代码。主机端代码包括程序中设定的全局工作组和局部工作组大小，局部工作组大小由程序开发者设定，设备端在执行内核程序时，不同大小的局部工作组会影响内核程序执行的时间。其中，内核程序的工作是计算处理单元中的存储器buffer a中数据的反正切函数(atan)结果，并存储到处理单元中的存储器buffer b中。该内核程序没用使用到局部内存或同步语句的操作，那么在通用寄存器分配时，首先设置通用寄存区的最大可分配数量mmax的初始值为64，然后根据最大可分配数量mmax分配寄存器，但由于反正切函数采用泰勒展开的算法交复杂，同时double类型在GPU处理器的SIMD指令数据模式中占用的通用寄存器空间是单精度浮点数的2倍，因此在mmax为64时，通用寄存器分配发生了寄存器溢出指令。此时更新mmax＝mmax+mstep，其中本实施例以步进值mstep为21，更新后的最大可分配数量mmax为85，经过分析，在mmax为85时继续进行寄存器分配，此时得到通用寄存器的实际使用数量mtotal为79个通用寄存器，并没有产生寄存器溢出指令，此时，设定状态位为true，标记机器指令程序中未使用局部内存语句或同步语句或将工作组索引信息参与除全局工作组索引之外的其他运算。
通过状态位可以确定机器指令程序不含有共享内存或同步语句，可以判定需要对工作组进行重新划分，即根据通用寄存器的实际使用数量mtotal，计算出一个处理单元(PE)中此时可启动的最大线程数量为4×mmax/mtotal(4是一个固定常数)，即通用寄存器为79时的可启动最大线程数量nmax为4x256/79＝12。则此时工作组可重新划分的范围为{32×N,0×1,0×1}，N为{1-12}。然后通过指令发射流程将机器指令程序转换为目标程序，将目标程序以及设定工作组重新划分的状态为true和工作组重新划分时的线程范围[1-N]发送至编译器。
编译器中的驱动程序接收到编译后的内核二进制文件，同时收到工作组重新划分的状态为true以及编译器计算到的工作组重新划分的线程范围[-N]，驱动可以对工作组进行重新设定，然后调用clEnqueueNDRangeKernel API启动执行内核程序。如主机端代码的工作组大小初始设定为[0×400,0×1,0×1]，则此时的工作组有0×400个工作项，由OpenCL编程标准和GPU处理器特性可知，工作组越大，启动工作组需要用到的通用寄存器数目越多。对于0×400大小的工作组，通过计算可得到本发明中的GPU处理器理论上需要32个线程才能启动一个工作组，则此时一个线程只能分得32个通用寄存器资源，至少需要79个可分配的通用寄存器才能不产生寄存器溢出，进而才能提高程序执行性能。同时，根据编译器发送的设定工作组重新划分的状态为true和工作组重新划分时的线程范围，驱动选择重新划分工作组大小，然后请求GPU处理器执行目标程序对应的任务。
对于工作组不可重新划分的情况，则判断机器指令程序在执行过程中不满足寄存器重新划分条件，若机器指令程序在执行过程中满足寄存器重新划分条件，即此时有足够的硬件资源能够启动线程，则请求GPU处理器执行目标程序对应的任务。若机器指令程序在执行过程中不满足寄存器重新划分条件，即此时的工作组大小不能启动线程，则重新设定最大可分配数量mmax大小进行二次编译或发送当前局部工作组超出最大局部工作组范围并退出。
在一个实施例中，提供一种通用寄存器动态分配方法的详细步骤，包括：
步骤1，通过编译器将源代码转换成机器指令程序。
步骤2，判断机器指令程序在执行过程中是否满足寄存器重新划分条件，寄存器重新划分条件包括：工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件；若不满足，则执行步骤3；若满足，则执行步骤4。
步骤3，将通用寄存器的最大可分配数量设置为固定值，根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量，执行步骤7。
步骤4，根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量，执行步骤5。
步骤5，判断机器指令程序执行过程中是否使用寄存器溢出指令，若使用寄存器溢出指令，则执行步骤6；若未使用寄存器溢出指令，则执行步骤7。
步骤6，确定通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及机器指令程序执行完成时所需的时钟数量的第二理论值；若第一理论值与第二理论值的比值大于预设值，则以固定的步进值增加通用寄存器的最大可分配数量，返回执行根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，直至第一理论值与第二理论值的比值小于预设值，以使得机器指令程序执行过程中未使用寄存器溢出指令，执行步骤7。
步骤7，设定一个状态位，用来标记机器指令程序中是否使用局部内存语句或同步语句或将工作组索引信息参与除全局工作组索引之外的其他运算，执行步骤8；其中，机器指令程序在执行过程中不满足寄存器重新划分条件时，状态位为false；机器指令程序在执行过程中满足寄存器重新划分条件时，状态位为true。
步骤8，根据状态位判断是否需要重新划分工作组；若状态位为true，则执行步骤9；若状态位为false，则执行步骤10。
步骤9，根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，确定GPU处理器中处理单元可启动的线程数量范围；取线程数量范围中任一数量，作为处理单元的目标启动线程数量；根据目标启动线程数量，确定多个第一基元，取多种第一基元中的任一种第一基元，作为第一目标基元；第一基元包括三个维度，三个维度的数值的乘积等于目标启动线程数量；根据线程粒度，确定多个第二基元，取多种第二基元中的任一种第二基元，作为第二目标基元；第二基元包括三个维度，三个维度的数值的乘积等于线程粒度；将第一目标基元与第二目标基元的对应维度上的数值乘积，确定为工作组在不同维度的目标大小，执行步骤10。
步骤10，通过指令发射流程将机器指令程序转换为目标程序，将目标程序、状态位和处理单元可启动的线程数量范围发送给编译器的驱动程序中。
步骤11，驱动程序根据状态位判定是否需要重新划分工作组大小，若状态位为true，则执行步骤12；若状态位为false，则执行步骤13。
步骤12，编译器将当前的工作组大小调整至计算得到的重新划分后的工作组大小，并执行步骤15。
步骤13，编译器根据通用寄存器的最大可分配数量以及通用寄存器的实际使用数量，判定GPU处理器中处理单元是否满足启动线程条件，若不满足启动线程条件，则执行步骤14；若满足启动线程条件，则执行步15。
步骤14，通过编译器重新设置通用寄存器的最大可分配数量，返回执行通过编译器将源代码转换成机器指令程序的步骤，直至处理单元满足启动线程条件，执行步骤15。
步骤15，请求GPU处理器执行目标程序对应的任务。
本实施例中，根据特定GPU架构，对于工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部存储或同步语句的程序可以动态分配通用寄存器的数量，从而平衡GPU处理器启动线程数目和通用寄存器资源的矛盾，提高硬件的执行效率；根据特定GPU处理器架构，对于工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部存储或同步语句的程序可以对用户设置的工作组大小进行重新合理的设定，以更适合当前GPU处理器架构；对于未使用局部存储或同步语句的程序可以避免二次编译的情况发生。
应该理解的是，虽然如上所述的各实施例所涉及的流程图中的各个步骤按照箭头的指示依次显示，但是这些步骤并不是必然按照箭头指示的顺序依次执行。除非本文中有明确的说明，这些步骤的执行并没有严格的顺序限制，这些步骤可以以其它的顺序执行。而且，如上所述的各实施例所涉及的流程图中的至少一部分步骤可以包括多个步骤或者多个阶段，这些步骤或者阶段并不必然是在同一时刻执行完成，而是可以在不同的时刻执行，这些步骤或者阶段的执行顺序也不必然是依次进行，而是可以与其它步骤或者其它步骤中的步骤或者阶段的至少一部分轮流或者交替地执行。
基于同样的发明构思，本申请实施例还提供了一种用于实现上述所涉及的通用寄存器动态分配方法的通用寄存器动态分配装置。该装置所提供的解决问题的实现方案与上述方法中所记载的实现方案相似，故下面所提供的一个或多个通用寄存器动态分配装置实施例中的具体限定可以参见上文中对于通用寄存器动态分配方法的限定，在此不再赘述。
在一个实施例中，如图9所示，提供了一种通用寄存器动态分配装置，包括：
转换模块100，用于通过编译器将源代码转换成机器指令程序；
分配模块200，用于若机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量；
调整模块300，用于若机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得机器指令程序执行过程中未使用寄存器溢出指令；
工作组划分模块400，用于根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定编译器中工作组的大小。
在一个实施例中，分配模块200还用于若机器指令程序在执行过程中不满足寄存器重新划分条件，则将通用寄存器的最大可分配数量设置为固定值，根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量。
在一个实施例中，调整模块300还用于确定通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及机器指令程序执行完成时所需的时钟数量的第二理论值；
若第一理论值与第二理论值的比值大于预设值，则以固定的步进值增加通用寄存器的最大可分配数量，返回执行根据通用寄存器的最大可分配数量进行寄存器分配操作，确定机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，直至第一理论值与第二理论值的比值小于预设值。
在一个实施例中，工作组划分模块400还用于若根据通用寄存器的最大可分配数量以及通用寄存器的实际使用数量，判定GPU处理器中处理单元不满足启动线程条件，则通过编译器重新设置通用寄存器的最大可分配数量，返回执行通过编译器将源代码转换成机器指令程序的步骤，直至处理单元满足启动线程条件。
其中，启动线程条件包括通用寄存器的最大可分配数量与通用寄存器的实际使用数量的比值、大于工作组的大小与线程粒度之间的比值。
在一个实施例中，工作组划分模块400还用于根据通用寄存器的实际使用数量和调整后的通用寄存器的最大可分配数量，确定GPU处理器中处理单元可启动的线程数量范围；
取线程数量范围中任一数量，作为处理单元的目标启动线程数量；
根据目标启动线程数量，确定多个第一基元，取多种第一基元中的任一种第一基元，作为第一目标基元；第一基元包括三个维度，三个维度的数值的乘积等于目标启动线程数量；
根据线程粒度，确定多个第二基元，取多种第二基元中的任一种第二基元，作为第二目标基元；第二基元包括三个维度，三个维度的数值的乘积等于线程粒度；
将第一目标基元与第二目标基元的对应维度上的数值乘积，确定为工作组在不同维度的目标大小。
上述通用寄存器动态分配装置中的各个模块可全部或部分通过软件、硬件及其组合来实现。上述各模块可以硬件形式内嵌于或独立于计算机设备中的处理器中，也可以以软件形式存储于计算机设备中的存储器中，以便于处理器调用执行以上各个模块对应的操作。
在一个实施例中，提供了一种计算机设备，该计算机设备可以是终端，其内部结构图可以如图10所示。该计算机设备包括处理器、存储器、输入/输出接口、通信接口、显示单元和输入装置。其中，处理器、存储器和输入/输出接口通过系统总线连接，通信接口、显示单元和输入装置通过输入/输出接口连接到系统总线。其中，该计算机设备的处理器用于提供计算和控制能力。该计算机设备的存储器包括非易失性存储介质和内存储器。该非易失性存储介质存储有操作系统和计算机程序。该内存储器为非易失性存储介质中的操作系统和计算机程序的运行提供环境。该计算机设备的输入/输出接口用于处理器与外部设备之间交换信息。该计算机设备的通信接口用于与外部的终端进行有线或无线方式的通信，无线方式可通过WIFI、移动蜂窝网络、NFC(近场通信)或其他技术实现。该计算机程序被处理器执行时以实现一种通用寄存器动态分配方法。该计算机设备的显示单元用于形成视觉可见的画面，可以是显示屏、投影装置或虚拟现实成像装置。显示屏可以是液晶显示屏或者电子墨水显示屏，该计算机设备的输入装置可以是显示屏上覆盖的触摸层，也可以是计算机设备外壳上设置的按键、轨迹球或触控板，还可以是外接的键盘、触控板或鼠标等。
本领域技术人员可以理解，图10中示出的结构，仅仅是与本申请方案相关的部分结构的框图，并不构成对本申请方案所应用于其上的计算机设备的限定，具体的计算机设备可以包括比图中所示更多或更少的部件，或者组合某些部件，或者具有不同的部件布置。
在一个实施例中，提供了一种计算机设备，包括存储器和处理器，存储器中存储有计算机程序，该处理器执行计算机程序时实现上述各方法实施例中的步骤。
在一个实施例中，提供了一种计算机可读存储介质，其上存储有计算机程序，计算机程序被处理器执行时实现上述各方法实施例中的步骤。
在一个实施例中，提供了一种计算机程序产品，包括计算机程序，该计算机程序被处理器执行时实现上述各方法实施例中的步骤。
需要说明的是，本申请所涉及的用户信息(包括但不限于用户设备信息、用户个人信息等)和数据(包括但不限于用于分析的数据、存储的数据、展示的数据等)，均为经用户授权或者经过各方充分授权的信息和数据，且相关数据的收集、使用和处理需要遵守相关国家和地区的相关法律法规和标准。
本领域普通技术人员可以理解实现上述实施例方法中的全部或部分流程，是可以通过计算机程序来指令相关的硬件来完成，所述的计算机程序可存储于一非易失性计算机可读取存储介质中，该计算机程序在执行时，可包括如上述各方法的实施例的流程。其中，本申请所提供的各实施例中所使用的对存储器、数据库或其它介质的任何引用，均可包括非易失性和易失性存储器中的至少一种。非易失性存储器可包括只读存储器(Read-OnlyMemory，ROM)、磁带、软盘、闪存、光存储器、高密度嵌入式非易失性存储器、阻变存储器(ReRAM)、磁变存储器(Magnetoresistive Random Access Memory，MRAM)、铁电存储器(Ferroelectric Random Access Memory，FRAM)、相变存储器(Phase Change Memory，PCM)、石墨烯存储器等。易失性存储器可包括随机存取存储器(Random Access Memory，RAM)或外部高速缓冲存储器等。作为说明而非局限，RAM可以是多种形式，比如静态随机存取存储器(Static Random Access Memory，SRAM)或动态随机存取存储器(Dynamic RandomAccess Memory，DRAM)等。本申请所提供的各实施例中所涉及的数据库可包括关系型数据库和非关系型数据库中至少一种。非关系型数据库可包括基于区块链的分布式数据库等，不限于此。本申请所提供的各实施例中所涉及的处理器可为通用处理器、中央处理器、图形处理器、数字信号处理器、可编程逻辑器、基于量子计算的数据处理逻辑器等，不限于此。
以上实施例的各技术特征可以进行任意的组合，为使描述简洁，未对上述实施例中的各个技术特征所有可能的组合都进行描述，然而，只要这些技术特征的组合不存在矛盾，都应当认为是本说明书记载的范围。
以上所述实施例仅表达了本申请的几种实施方式，其描述较为具体和详细，但并不能因此而理解为对本申请专利范围的限制。应当指出的是，对于本领域的普通技术人员来说，在不脱离本申请构思的前提下，还可以做出若干变形和改进，这些都属于本申请的保护范围。因此，本申请的保护范围应以所附权利要求为准。
Claims (10)
1.一种通用寄存器动态分配方法，其特征在于，所述方法包括：
通过编译器将源代码转换成机器指令程序；
若所述机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定所述机器指令程序在执行过程中通用寄存器的实际使用数量；
若所述机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得所述机器指令程序执行过程中未使用寄存器溢出指令；
根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定所述编译器中工作组的大小。
2.根据权利要求1所述的方法，其特征在于，所述方法还包括：
若所述机器指令程序在执行过程中不满足寄存器重新划分条件，则将所述通用寄存器的最大可分配数量设置为固定值，根据所述通用寄存器的最大可分配数量进行寄存器分配操作，确定所述机器指令程序在执行过程中通用寄存器的实际使用数量。
3.根据权利要求1或2所述的方法，其特征在于，所述寄存器重新划分条件包括：工作组索引信息参与除全局工作组索引之外的其他运算、未使用局部内存语句以未使用同步语句中任一种条件。
4.根据权利要求1所述的方法，其特征在于，所述调整所述通用寄存器的最大可分配数量，包括：
确定所述通用寄存器溢出指令执行时所需的时钟数量的第一理论值，以及所述机器指令程序执行完成时所需的时钟数量的第二理论值；
若所述第一理论值与所述第二理论值的比值大于预设值，则以固定的步进值增加所述通用寄存器的最大可分配数量，返回执行所述根据通用寄存器的最大可分配数量进行寄存器分配操作，确定所述机器指令程序在执行过程中通用寄存器的实际使用数量的步骤，直至所述第一理论值与所述第二理论值的比值小于预设值。
5.根据权利要求2所述的方法，其特征在于，所述方法还包括：
若根据所述通用寄存器的最大可分配数量以及所述通用寄存器的实际使用数量，判定GPU处理器中处理单元不满足启动线程条件，则通过所述编译器重新设置所述通用寄存器的最大可分配数量，返回执行所述通过编译器将源代码转换成机器指令程序的步骤，直至所述处理单元满足启动线程条件。
6.根据权利要求5所述的方法，其特征在于，所述启动线程条件包括所述通用寄存器的最大可分配数量与所述通用寄存器的实际使用数量的比值、大于所述工作组的大小与线程粒度之间的比值。
7.根据权利要求1所述的方法，其特征在于，所述根据所述通用寄存器的实际使用数量和调整后的所述通用寄存器的最大可分配数量，重新设定所述工作组的大小，包括：
根据所述通用寄存器的实际使用数量和调整后的所述通用寄存器的最大可分配数量，确定GPU处理器中处理单元可启动的线程数量范围；
取所述线程数量范围中任一数量，作为所述处理单元的目标启动线程数量；
根据所述目标启动线程数量，确定多个第一基元，取多种第一基元中的任一种第一基元，作为第一目标基元；所述第一基元包括三个维度，三个维度的数值的乘积等于所述目标启动线程数量；
根据线程粒度，确定多个第二基元，取多种第二基元中的任一种第二基元，作为第二目标基元；所述第二基元包括三个维度，三个维度的数值的乘积等于所述线程粒度；
将所述第一目标基元与所述第二目标基元的对应维度上的数值乘积，确定为所述工作组在不同维度的目标大小。
8.一种通用寄存器动态分配装置，其特征在于，所述装置包括：
转换模块，用于通过编译器将源代码转换成机器指令程序；
分配模块，用于若所述机器指令程序在执行过程中满足寄存器重新划分条件，则根据通用寄存器的最大可分配数量进行寄存器分配操作，确定所述机器指令程序在执行过程中通用寄存器的实际使用数量；
调整模块，用于若所述机器指令程序执行过程中使用寄存器溢出指令，则调整通用寄存器的最大可分配数量，以使得所述机器指令程序执行过程中未使用寄存器溢出指令；
工作组划分模块，用于根据通用寄存器的实际使用数量和调整后通用寄存器的最大可分配数量，重新设定所述编译器中工作组的大小。
9.一种计算机设备，包括存储器和处理器，所述存储器存储有计算机程序，其特征在于，所述处理器执行所述计算机程序时实现权利要求1至7中任一项所述的方法的步骤。
10.一种计算机可读存储介质，其上存储有计算机程序，其特征在于，所述计算机程序被处理器执行时实现权利要求1至7中任一项所述的方法的步骤。
Priority Applications (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202211705663.4A CN115934102B (zh) | 2022-12-29 | 2022-12-29 | 通用寄存器动态分配方法、装置、计算机设备和存储介质 |
Applications Claiming Priority (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202211705663.4A CN115934102B (zh) | 2022-12-29 | 2022-12-29 | 通用寄存器动态分配方法、装置、计算机设备和存储介质 |
Publications (2)
| Publication Number | Publication Date |
|---|---|
| CN115934102A true CN115934102A (zh) | 2023-04-07 |
| CN115934102B CN115934102B (zh) | 2023-12-12 |
Family
ID=86699083
Family Applications (1)
| Application Number | Title | Priority Date | Filing Date |
|---|---|---|---|
| CN202211705663.4A Active CN115934102B (zh) | 2022-12-29 | 2022-12-29 | 通用寄存器动态分配方法、装置、计算机设备和存储介质 |
Country Status (1)
| Country | Link |
|---|---|
| CN (1) | CN115934102B (zh) |
Cited By (2)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN117971437A (zh) * | 2024-03-26 | 2024-05-03 | 摩尔线程智能科技(北京)有限责任公司 | 任务分配方法、电路、设备、介质及程序 |
| CN118113486A (zh) * | 2024-04-30 | 2024-05-31 | 北京麟卓信息科技有限公司 | 一种gpu线程最大可用寄存器数量的快速测算方法 |
Citations (9)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN1662904A (zh) * | 2002-06-26 | 2005-08-31 | 国际商业机器公司 | 具有级联simd结构的数字信号处理器 |
| CN101014933A (zh) * | 2004-07-13 | 2007-08-08 | 辉达公司 | 使用具有较低端口数的存储器的模拟多端口存储器 |
| CN102640132A (zh) * | 2009-09-28 | 2012-08-15 | 辉达公司 | 用于并行处理器的高效断言执行 |
| CN107851004A (zh) * | 2015-08-17 | 2018-03-27 | 高通股份有限公司 | 针对通用寄存器(gpr)的寄存器溢出管理 |
| CN110032395A (zh) * | 2017-11-14 | 2019-07-19 | 辉达公司 | 用于提高资源利用率的统一寄存器文件 |
| CN110187977A (zh) * | 2018-02-23 | 2019-08-30 | 英特尔公司 | 用于基于软件提示和硬件线程切换来降低寄存器区块冲突的系统和方法 |
| WO2021051022A1 (en) * | 2019-09-11 | 2021-03-18 | Redpine Signals Inc. | Multi-threaded processor with thread granularity |
| CN112947998A (zh) * | 2019-11-26 | 2021-06-11 | Arm有限公司 | 寄存器提供操作码指令 |
| CN114924748A (zh) * | 2022-05-31 | 2022-08-19 | 上海阵量智能科技有限公司 | 编译方法、装置及设备 |
-
2022
- 2022-12-29 CN CN202211705663.4A patent/CN115934102B/zh active Active
Patent Citations (9)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN1662904A (zh) * | 2002-06-26 | 2005-08-31 | 国际商业机器公司 | 具有级联simd结构的数字信号处理器 |
| CN101014933A (zh) * | 2004-07-13 | 2007-08-08 | 辉达公司 | 使用具有较低端口数的存储器的模拟多端口存储器 |
| CN102640132A (zh) * | 2009-09-28 | 2012-08-15 | 辉达公司 | 用于并行处理器的高效断言执行 |
| CN107851004A (zh) * | 2015-08-17 | 2018-03-27 | 高通股份有限公司 | 针对通用寄存器(gpr)的寄存器溢出管理 |
| CN110032395A (zh) * | 2017-11-14 | 2019-07-19 | 辉达公司 | 用于提高资源利用率的统一寄存器文件 |
| CN110187977A (zh) * | 2018-02-23 | 2019-08-30 | 英特尔公司 | 用于基于软件提示和硬件线程切换来降低寄存器区块冲突的系统和方法 |
| WO2021051022A1 (en) * | 2019-09-11 | 2021-03-18 | Redpine Signals Inc. | Multi-threaded processor with thread granularity |
| CN112947998A (zh) * | 2019-11-26 | 2021-06-11 | Arm有限公司 | 寄存器提供操作码指令 |
| CN114924748A (zh) * | 2022-05-31 | 2022-08-19 | 上海阵量智能科技有限公司 | 编译方法、装置及设备 |
Non-Patent Citations (2)
| Title |
|---|
| C. -C. HSIAO 等: "An Adaptive Thread Scheduling Mechanism With Low-Power Register File for Mobile GPUs", 《IEEE TRANSACTIONS ON MULTIMEDIA》, vol. 16, no. 01, pages 60 - 67, XP011533902, DOI: 10.1109/TMM.2013.2281584 * |
| 邱亚琼 等: "基于两类寄存器互为缓存方法的DSP寄存器分配溢出处理优化算法", 《计算机科学》, vol. 46, no. 06, pages 196 - 200 * |
Cited By (2)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN117971437A (zh) * | 2024-03-26 | 2024-05-03 | 摩尔线程智能科技(北京)有限责任公司 | 任务分配方法、电路、设备、介质及程序 |
| CN118113486A (zh) * | 2024-04-30 | 2024-05-31 | 北京麟卓信息科技有限公司 | 一种gpu线程最大可用寄存器数量的快速测算方法 |
Also Published As
| Publication number | Publication date |
|---|---|
| CN115934102B (zh) | 2023-12-12 |
Similar Documents
| Publication | Publication Date | Title |
|---|---|---|
| US9354944B2 (en) | Mapping processing logic having data-parallel threads across processors | |
| US8145625B2 (en) | Methods and systems for optimizing data accesses | |
| US9779469B2 (en) | Register spill management for general purpose registers (GPRs) | |
| US11243816B2 (en) | Program execution on heterogeneous platform | |
| CN110008009B (zh) | 在运行时绑定常量以提高资源利用率 | |
| US8806513B2 (en) | Application programming interfaces for data parallel computing on multiple processors | |
| US9720726B2 (en) | Multi-dimensional thread grouping for multiple processors | |
| CN103309786B (zh) | 用于在非可抢占式图形处理单元上交互调试的方法和装置 | |
| KR101558831B1 (ko) | 서브버퍼 오브젝트 | |
| US9411715B2 (en) | System, method, and computer program product for optimizing the management of thread stack memory | |
| CN109154886B (zh) | 用于处理数据的方法和设备 | |
| CN115934102B (zh) | 通用寄存器动态分配方法、装置、计算机设备和存储介质 | |
| US9513886B2 (en) | Heap data management for limited local memory(LLM) multi-core processors | |
| KR101609079B1 (ko) | 그래픽 프로세싱 유닛에서의 명령 선별 | |
| WO2014190315A1 (en) | Graphics processing using dynamic resources | |
| CN103996216A (zh) | 用于曲面细分和几何着色器的电力高效属性处置 | |
| US11500828B1 (en) | Method and device for constructing database model with ID-based data indexing-enabled data accessing | |
| US12020065B2 (en) | Hierarchical processor selection | |
| US10599638B2 (en) | System and method for identifying maximal independent sets in parallel | |
| Mantas et al. | An introduction to GPU computing for numerical simulation | |
| CN103870247A (zh) | 用于保存和恢复线程组操作状态的技术 | |
| CN111699506B (zh) | 指令处理 | |
| CN118550705A (zh) | 工作组调度方法、装置、计算机设备、存储介质和计算机程序产品 | |
| CN114995908A (zh) | 一种多核启动配置确定的方法、装置、计算机设备和存储介质 | |
| Lechner | Optimization and Extension of Minimal-Cronos |
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 | ||
| GR01 | Patent grant | ||
| GR01 | Patent grant | ||
| CP03 | Change of name, title or address | ||
| CP03 | Change of name, title or address |
Address after: 200135, 11th Floor, Building 3, No. 889 Bibo Road, China (Shanghai) Pilot Free Trade Zone, Pudong New Area, Shanghai Patentee after: Granfei Intelligent Technology Co.,Ltd. Country or region after: China Address before: 200135 Room 201, No. 2557, Jinke Road, China (Shanghai) pilot Free Trade Zone, Pudong New Area, Shanghai Patentee before: Gryfield Intelligent Technology Co.,Ltd. Country or region before: China |