CN113722111A - 内存分配方法、系统、装置和计算机可读介质 - Google Patents

Abstract

提供内存分配方法、系统、装置和计算机可读介质。方法包括如下步骤：响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量；选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；为所述相应线程束分配内存空间；根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

Description

内存分配方法、系统、装置和计算机可读介质

技术领域

本申请涉及处理器领域，且更具体地涉及内存分配方法、系统、装置和计算机可读介质。

背景技术

现代处理器架构利用各种执行模型，例如（Single Instruction Multiple Data，SIMD）和单指令多线程（Single Instruction Multiple Threads，SIMT）。在图形处理单元（Graphic Processing Unit，GPU）中，为了有效地管理和执行多个单线程，多处理器通常采用单指令多线程SIMT架构。

在单指令多线程SIMT中，多个线程对不同的数据集执行相同的指令。每当图形处理单元GPU需要执行特定指令时，都会从内存中获取数据和指令，然后对其进行解码和执行。在这种情况下，需要使用同一条指令执行的所有数据集（达到一定限制）将使用处理器可用的各种线程同时预取并同时执行。相比于单指令多数据流SIMD，单指令多线程SIMT的好处是无需开发者费力把数据凑成合适的矢量长度，并且单指令多线程SIMT允许每个线程有不同的分支，且减少了指令预取带来的等待时间。

在线程执行的过程中，线程会请求处理器内核分配内存。当前，在处理器内核中存在许多不同的内存分配机制。例如每个线程采用malloc()函数来进行内存请求和分配，其参数包括请求分配的内存大小。但已有的内存分配方法仍然不够高性能和高速，存在改进的需求。

仍需要高性能、高速的内存分配机制。

发明内容

为解决现有技术中存在的一个或多个问题，提供本申请的各个方面。

根据本申请的一个方面，提供一种内存分配方法，包括如下步骤：响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量；选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；为所述相应线程束分配内存空间；根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

根据本申请的另一个方面，提供一种内存分配系统，包括如下装置：线程束分组装置，被配置为响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；线程束请求装置，被配置为计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量，并选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；内存分配装置，被配置为为所述相应线程束分配内存空间；线程束内存划分装置，被配置为根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

根据本申请的另一个方面，提供一种内存分配装置，包括：处理器；存储器，存储了计算器可执行指令，其被处理器运行时进行根据本申请的实施例的各个方法。

根据本申请的另一个方面，提供一种计算机可读介质，存储有计算机可读指令，其中当所述计算机可读指令由处理器运行时，进行根据本申请的实施例的各个方法。

本申请利用线程束的特点，多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，并且在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配。如此可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

附图说明

为了更清楚地说明本公开实施例或现有技术中的技术方案，下面将对实施例或现有技术描述中所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本公开的一些实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据这些附图获得其他的附图。

图1示出了为每个线程分配其自己的内存空间的malloc函数的操作示意图。

图2示出了根据本申请的实施例的内存分配方法的示意流程图。

图3示意性地示出了如图2所示的内存分配方法的操作过程。

图4示出了根据本申请的另一实施例的内存分配方法的示意流程图。

图5示出了图4示出的内存分配方法的各个步骤所产生的示例结果。

图6示出了根据本申请的实施例的内存分配系统的示意方框图。

图7示出了适于用来实现本申请实施方式的示例性计算机系统的框图。

图8示出了根据本公开的实施例的非暂时性计算机可读存储介质的示意图。

具体实施方式

现在将详细参照本申请的具体实施例，在附图中例示了本申请的例子。尽管将结合具体实施例描述本申请，但将理解，不是想要将本申请限于描述的实施例。相反，想要覆盖由所附权利要求限定的在本申请的精神和范围内包括的变更、修改和等价物。应注意，这里描述的方法步骤都可以由任何功能块或功能布置来实现，且任何功能块或功能布置可被实现为物理实体或逻辑实体、或者两者的组合。

在现有技术中，在多个处理器核（kernel）的场景下，处理器核需要动态大小的全局内存（该全局内存是被驱动器预留的预定空间），在一个处理器核中，存储不能被适应到共享内存中的大块数据。在多个处理器核之间可以交换数据。为线程分配内存空间例如采用如下软件代码来实现：

__global__ void example() {

cahr* v = malloc(threadIdx.x+1);

}

其中，malloc(threadIdx.x+1)指的是从全局内存中给整个线程块中的所有多个线程分配内存。线程块是OpenCL工作组汇总的概念，表示在相同执行单元中运行且共享相同的共享内存的一组线程。

然后，为整个线程块中的每个线程分配其自己的内存空间。图1示出了为每个线程分配其自己的内存空间的malloc函数的操作示意图。

如图1所示，假设在线程块中存在线程T0-T9。在为每个线程分配其自己的内存空间的过程中，10个线程分别发出10个各自的内存分配请求。

由于在内存分配过程中存在锁（lock）操作。锁是一种用于控制多个线程对共享的内存资源的访问的工具。通常，锁提供了对共享资源的独占访问，因此一次只能有一个线程可以获取该锁，并且对共享资源的所有访问都需要首先获取该锁。

每个线程首先要获取锁才能访问内存资源以为自己分配内存。例如，通过非阻塞尝试获取锁tryLock()函数来实现线程的获取锁操作。

因此所有线程都会尝试获取锁。假设线程T0获取了该锁，为该线程T0执行malloc内存分配操作。该malloc内存分配操作是一个原子操作，直到该原子操作完成之后，其他线程再会尝试获取锁。假设此时线程T3获取了该锁，为该线程T3执行malloc内存分配原子操作，直到该原子操作完成之后释放该锁，其他线程再会尝试获取锁。假设此时线程T2获取了该锁，为该线程T2执行malloc内存分配原子操作，直到该原子操作完成之后释放该锁，其他线程再会尝试获取锁。假设此时线程T9获取了该锁，为该线程T9执行malloc内存分配原子操作，直到该原子操作完成之后释放该锁，其他线程再会尝试获取锁。以此类推，直到所有线程都获取到锁且完成了内存分配原子操作。

这样操作的缺点是导致了处理器的低性能。首先，发出太多的原子操作，给例如内存总线和高速缓存器等硬件带来了压力。其次是存在非常大量的锁冲突，可能导致来回循环等待获取锁等待，因此执行了更多的指令，导致执行效率下降。

本申请考虑解决上述问题。在图形处理单元GPU硬件的线程调度中，线程调度的最小单位通常是一个线程束（warp），即在图形处理单元GPU中每次执行一个线程束。一般一个线程束可以包含32个线程或者其他数量个线程。本申请利用线程束的特点，多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，并且在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配。如此可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

接下来通过附图来描述本申请的各个实施例的细节。

图2示出了根据本申请的实施例的内存分配方法200的示意流程图。

如图2所示，内存分配方法200包括如下步骤：步骤202，响应于多个线程的各自的内存分配请求，将多个线程分组为多个线程束，其中，线程的内存分配请求包括线程的请求的内存量；步骤204，计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为相应线程束的请求的内存量；步骤206，选择相应线程束中的一个线程来请求分配相应线程束的内存空间，相应线程束的内存空间的量是相应线程束的请求的内存量；步骤208，为相应线程束分配内存空间；步骤210，根据相应线程束中的各个线程的请求的内存量，将被分配的相应线程束的内存空间划分给各个线程。

结合图3来详细描述上述步骤。图3示意性地示出了如图2所示的内存分配方法200的操作过程。

如图3所示，在步骤202中，响应于多个线程的各自的内存分配请求，将多个线程分组为多个线程束，例如线程束0、线程束1等等。每个线程的内存分配请求包括该线程的请求的内存量。假设线程束0中存在32个线程，线程束1中存在32个线程等等。

在步骤204中，计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为相应线程束的请求的内存量。即合并线程束中的各个线程的请求。例如，假设线程束0中的线程1请求1M内存，线程2请求2M内存，线程3请求1M内存……，线程束0中的32个线程各自请求的内存量的总和假设为50M内存，线程1中的32个线程各自请求的内存量的总和假设为30M内存，以此类推。

然后，在步骤206中，通过选择相应线程束中的一个线程作为请求方来请求分配相应线程束的内存空间。此时，相应线程束中的其他线程的内存分配请求可以被忽略，因为只要有一个线程作为请求方即可。

该线程请求分配的相应线程束的内存空间的量是相应线程束的请求的内存量，即先前计算的线程束0的50M内存、线程束1的30M内存等等。

在步骤208中，为相应线程束分配内存空间。例如，为线程束0分配50M内存块（地址范围为XXXXXXXXXXXX~XXXXXXXXXXXX，或起始地址为XXXXXXXXXXXX），为线程束1分配30M内存块（地址范围为XXXXXXXXXXXX~XXXXXXXXXXXX，或起始地址为XXXXXXXX）等。在这里为相应线程束分配内存空间同样也需要各个线程束进行获取锁和分配内存空间原子操作以及释放锁的操作，在此不赘述。

在步骤210中，根据相应线程束中的各个线程的请求的内存量，将被分配的相应线程束的内存空间划分给各个线程。例如，线程束0被分配了50M内存，然后将这50M内存中的1M内存分配给线程1，2M内存分配给线程2，1M内存分配给线程3，以此类推……。

本申请的实施例利用线程束的特点，由于线程束本身就是线程调度的最小单位，通过多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，能够利用现有的线程调度规则进行高效的内存分配，使得内存分配与线程调度的单位相同，从而也获得操作的一致性。

而且由于以线程束为单位进行内存分配，即获取锁和分配内存空间原子操作以及释放锁的操作，可以减少多个线程分别获取锁的次数和多个线程争抢锁带来的冲突，同时也执行了较少的原子操作。

并且本申请的实施例在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配，使得各个线程不需要争抢锁也能获得其各自所需的内存空间。后续还将讨论各个线程如何高效且不冲突地获得其各自所需的内存空间。

而且，由于以线程束为单位进行内存分配，因此在一个线程束的内存被分配好后，利用其它线程束继续争抢锁和分配内存的操作时间，可以在该操作时间中并行地进行将该线程束的内存划分给线程束中的各个线程的操作。相比于所有线程分别争抢锁和分配内存，本申请的实施例可以减少内存分配的操作时间，极大地提高内存分配效率。

如此，根据本申请的实施例，可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

在这里，本申请中的线程束的概念对应于warp，但是在其它处理器中可以对应于其它术语，本申请不限制具体的术语，而是只要将多个线程分组为适应于在处理器中共同处理的线程组（线程束）即可，而该线程组（线程束）如果正好与处理器中已有的线程处理过程中的规则相同，则可以更有效率地利用现有规则来开发本申请的各个实施例，而不用改变已有规则或增加硬件装置或增加软件指令。

图4示出了根据本申请的另一实施例的内存分配方法400的示意流程图。其中与图2所示的内存分配方法200中的步骤相同的步骤采用相同的编号。图5示出了图4示出的内存分配方法400的各个步骤所产生的示例结果。

下面结合图4和图5来介绍根据本申请的另一实施例的内存分配方法400的流程。

如图4所示，在步骤202中，响应于多个线程的各自的内存分配请求，将多个线程分组为多个线程束，其中，线程的内存分配请求包括线程的请求的内存量。

如图4所示，在一个实施例中，每个线程束的大小被配置为预定量个线程，响应于多个线程的各自的内存分配请求，将多个线程分组为多个线程束的步骤202包括：步骤2021，将所述多个线程分配到所述多个线程束中并添加或不添加不需要执行的附加线程使得每个线程束中包括的线程的数量等于所述预定量。

结合图5，假设线程块尺寸为64个线程，每个线程束的预定量为32，此时多个线程的数量能被预定量整除，则将64个线程划分为2个线程束，此时不需要添加不需要执行的附加线程。

当然，如果多个线程的数量不能被预定量整除，假如线程块尺寸为60个线程，每个线程束的预定量为32，则60个线程被划分为一个线程束包括32个线程，而另一个线程束包括28个线程，加上4个不需要执行的附加线程，该不需要执行的附加线程不活跃，且可以在实际分配地址的过程中不被分配地址且不被执行。

假设每个线程请求2字节（B）内存空间，如图5所示的“2”代表每个线程请求的2字节内存空间。当然这并非限制，实际上每个线程请求的内存空间可以并不完全相同，内存空间大小也可以是其他尺寸。

如图4所示，在一个实施例中，方法还可以包括：步骤203，如果线程的请求的内存量小于内存分配的预定最小量，将线程的请求的内存量扩大为内存分配的预定最小量。

结合图5，假设内存分配的预定最小量为4字节，则在该步骤203中，将线程请求的2字节内存量扩大为4字节，以满足内存分配的预定最小量。如果线程的请求的内存量大于或等于内存分配的预定最小量，则不改变线程的请求的内存量。

这样可以增加内存分配效率，也能符合内存分配的规则。

如图4所示，在步骤204中，计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为相应线程束的请求的内存量。

结合图5，在该例子中， 每个线程的请求的内存量被扩大为4字节，因此，第一个线程束中的32个线程的各个线程的请求的内存量的总和为128字节，作为第一个线程束的请求的内存量。第二个线程束中的32个线程的各个线程的请求的内存量的总和为128字节，作为第二个线程束的请求的内存量。

如图4所示，方法400还可以包括如下步骤205：根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的各自的前缀和（prefixSum），作为各个线程的内存地址偏移量。

如图4所示，在一个实施例中，根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的前缀和，作为各个线程的内存地址偏移量的步骤205包括如下步骤：步骤2051，将相应线程束中的第一线程的前缀和设置为0；步骤2052，计算相应线程束中的后续线程的前缀和为后续线程的前一线程的前缀和加上后续线程的请求的内存量。

图4描绘了步骤205在步骤204和步骤206之间，但这不是必须的，实际上步骤205可以任何时间进行，只要在步骤2101中能够得知各个线程的内存地址偏移量即可。

结合图5，在步骤205中，在第一个线程束中的线程T0的前缀和为0，在第一个线程束中的线程T1的前缀和为0+4=4字节，在第一个线程束中的线程T2的前缀和为4+4=8字节，在第一个线程束中的线程T3的前缀和为8+4=12字节，以此类推，在第一个线程束中的线程T31的前缀和为120+4=124字节。

同理，在第二个线程束中的线程T32的前缀和为0，在第二个线程束中的线程T33的前缀和为0+4=4字节，在第二个线程束中的线程T34的前缀和为4+4=8字节，在第二个线程束中的线程T3的前缀和为8+4=12字节，以此类推，在第二个线程束中的线程T63的前缀和为120+4=124字节。

由于计算前缀和可以用已有的函数prefixSum（），prefixSum（）的定义如下：

给定一个数组A[0..n-1]，对于i

的每个i（i是整数），前缀和被计算为： PrefixSum[i] = A[0]+A[1]+...+A[i-1]；

例如：A[5,8,9,2] --> PrefixSum[5,13,22,24]，即：

PrefixSum[0] =A[0] ;

PrefixSum[1] =A[0] + A[1] ;

PrefixSum[2] =A[0] + A[1] + A[2] ;

PrefixSum[3] =A[0] + A[1] + A[2] + A[3]。

上述例子中A[0]=5，而在本申请的实施例中，为了更好地适应于各个线程对内存的划分，设置A[0]=0。

即，根据本申请的实施例，巧妙地利用已有的前缀和计算方法，可以容易地将该前缀和与每个线程所需的内存地址的偏移量相关联，即用每个线程的所需内存空间的前缀和来得到每个线程所需的内存空间在地址上偏移量，以便后续在整个线程束被分配的内存地址空间中为各个线程进行划分。如此，并未增加后续划分的计算难度和硬件复杂性。

如图4所示，在步骤206中，选择相应线程束中的一个线程来请求分配相应线程束的内存空间，相应线程束的内存空间的量是相应线程束的请求的内存量。

结合图5，选择相应线程束中的一个线程来请求分配相应线程束的内存空间，即一个线程发起malloc（）。该线程可以是相应线程束中的第一个线程，例如第一个线程束中的线程T0，或者第二个线程束中的线程T32。当然，实际上相应线程束中的任一个线程都可以作为发起请求者。而其它线程不需要发起请求，因为每个线程束只需要请求分配内存空间一次。

如此，以每个线程束中的一个线程作为内存请求方，使得每个线程束只进行一次内存分配请求，减少了获取锁和分配内存空间原子操作以及释放锁的操作，可以减少多个线程分别获取锁的次数和多个线程争抢锁带来的冲突。

如图4所示，在步骤208中，为相应线程束分配内存空间。

结合图5，假设为第一个线程束分配了内存空间，起始地址为0x7babcd0000，长度为128B（即第一个线程束中的所有线程要求分配的内存空间总和），为第二个线程束分配了内存空间，起始地址为0x7bbcde0000，长度为128B（即第二个线程束中的所有线程要求分配的内存空间总和）。

如图4所示，在步骤210中，根据相应线程束中的各个线程的请求的内存量，将被分配的相应线程束的内存空间划分给各个线程。

如图4所示，根据相应线程束中的各个线程的请求的内存量，将被分配的相应线程束的内存空间划分给各个线程的步骤210可以包括：步骤2101，按照各个线程的内存地址偏移量，在被分配的相应线程束的内存空间的量中、将偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。

结合图5，第一个线程束的内存空间的起始地址为0x7babcd0000，根据前缀和计算的各个线程的内存地址偏移量，偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。例如线程T0的内存空间的起始地址为0x7babcd0000，其空间大小为4B，而线程T1的内存空间的起始地址为0x7babcd0004，其空间大小为4B，线程T2的内存空间的起始地址为0x7babcd0008，其空间大小为4B，线程T3的内存空间的起始地址为0x7babcd000C，其空间大小为4B，以此类推。

第二个线程束的内存空间的起始地址为0x7bbcde0000，根据前缀和计算的各个线程的内存地址偏移量，偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。例如线程T32的内存空间的起始地址为0x7bbcde 0000，其空间大小为4B，而线程T1的内存空间的起始地址为0x7bbcde 0004，其空间大小为4B，线程T2的内存空间的起始地址为0x7bbcde 0008，其空间大小为4B，线程T3的内存空间的起始地址为0x7bbcde 000C，其空间大小为4B，以此类推。

如此，由于通过计算各个线程的请求的内存量在相应线程束中的各自的前缀和，得到了各个线程的内存地址偏移量，可以使得在以线程束为单位获得内存分配的地址范围之后，根据内存地址偏移量将分配的地址范围按照偏移量作为起始地址划分给各个线程。

根据本申请的各个实施例，可以容易且高效地进行多个线程分组为线程束、为线程束分配内存空间、将内存空间划分给线程束中的各个线程的整个过程。

在一个实施例中，多个线程可以是在单指令多线程SIMT执行系统中用于执行单个指令的多个线程。如此，可以对执行单个指令的一批线程进行上述内存分配方法，提高了该单个指令的执行效率。

本申请的实施例利用线程束的特点，由于线程束本身就是线程调度的最小单位，通过多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，能够利用现有的线程调度规则进行高效的内存分配，使得内存分配与线程调度的单位相同，从而也获得操作的一致性，也节省了执行单独的线程分配内存操作的每个线程的循环逻辑。

而且由于以线程束为单位进行内存分配，其原子操作的数量是以线程为单位进行内存分配的例如32分之一，即执行了较少的原子操作。以每个线程束中的一个线程作为内存请求方，使得每个线程束只进行一次内存分配请求，减少了获取锁和分配内存空间原子操作以及释放锁的操作，可以减少多个线程分别获取锁的次数和多个线程争抢锁带来的冲突。

并且本申请的实施例在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配，使得各个线程不需要争抢锁也能获得其各自所需的内存空间。后续还将讨论各个线程如何高效且不冲突地获得其各自所需的内存空间。

而且，由于以线程束为单位进行内存分配，因此在一个线程束的内存被分配好后，利用其它线程束继续争抢锁和分配内存的操作时间，可以在该操作时间中并行地进行将该线程束的内存划分给线程束中的各个线程的操作。相比于所有线程分别争抢锁和分配内存，本申请的实施例可以减少内存分配的操作时间，极大地提高内存分配效率。

如此，根据本申请的实施例，每个线程束仅增加了少量线程束分组操作和一些指令，但可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

图6示出了根据本申请的实施例的内存分配系统的示意方框图。

如图6所示，内存分配系统600包括如下装置：线程束分组装置601，被配置为响应于来自系统的内存分配（Malloc）接口的多个线程的各自的内存分配请求，将多个线程分组为多个线程束，其中，线程的内存分配请求包括线程的请求的内存量；线程束请求装置602，被配置为计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为相应线程束的请求的内存量，并选择相应线程束中的一个线程来请求分配相应线程束的内存空间，相应线程束的内存空间的量是相应线程束的请求的内存量；内存分配装置603，被配置为从系统的全局内存中为相应线程束分配内存空间；线程束内存划分装置604，被配置为根据相应线程束中的各个线程的请求的内存量，将被分配的相应线程束的内存空间划分给各个线程。

在一个实施例中，线程束请求装置602还被配置为：根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的各自的前缀和，作为各个线程的内存地址偏移量。线程束内存划分装置604被配置为：按照各个线程的内存地址偏移量，在被分配的相应线程束的内存空间的量中、将偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。

在一个实施例中，线程束请求装置602还被配置为通过如下步骤来根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的前缀和，作为各个线程的内存地址偏移量：将相应线程束中的第一线程的前缀和设置为0；计算相应线程束中的后续线程的前缀和为后续线程的前一线程的前缀和加上后续线程的请求的内存量。

在一个实施例中，每个线程束的大小被配置为预定量个线程，线程束分组装置601被配置为：将所述多个线程分配到所述多个线程束中并添加或不添加不需要执行的附加线程使得每个线程束中包括的线程的数量等于所述预定量。

在一个实施例中，线程束请求装置602还被配置为：如果线程的请求的内存量小于内存分配的预定最小量，将线程的请求的内存量扩大为内存分配的预定最小量。

在一个实施例中，多个线程是在单指令多线程SIMT执行系统中用于执行单个指令的多个线程。

本申请的实施例利用线程束的特点，由于线程束本身就是线程调度的最小单位，通过多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，能够利用现有的线程调度规则进行高效的内存分配，使得内存分配与线程调度的单位相同，从而也获得操作的一致性。

而且由于以线程束为单位进行内存分配，即获取锁和分配内存空间原子操作以及释放锁的操作，可以减少多个线程分别获取锁的次数和多个线程争抢锁带来的冲突，同时也执行了较少的原子操作。

并且本申请的实施例在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配，使得各个线程不需要争抢锁也能获得其各自所需的内存空间。后续还将讨论各个线程如何高效且不冲突地获得其各自所需的内存空间。

而且，由于以线程束为单位进行内存分配，因此在一个线程束的内存被分配好后，利用其它线程束继续争抢锁和分配内存的操作时间，可以在该操作时间中并行地进行将该线程束的内存划分给线程束中的各个线程的操作。相比于所有线程分别争抢锁和分配内存，本申请的实施例可以减少内存分配的操作时间，极大地提高内存分配效率。

如此，根据本申请的实施例，可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

图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的情况下，可以进行如上描述的各个方法。

本申请的实施例利用线程束的特点，由于线程束本身就是线程调度的最小单位，通过多个线程以线程束为单位来进行内存分配而不是以一个线程为单位来进行内存分配，能够利用现有的线程调度规则进行高效的内存分配，使得内存分配与线程调度的单位相同，从而也获得操作的一致性。

而且由于以线程束为单位进行内存分配，即获取锁和分配内存空间原子操作以及释放锁的操作，可以减少多个线程分别获取锁的次数和多个线程争抢锁带来的冲突，同时也执行了较少的原子操作。

并且本申请的实施例在以线程束为单位来进行内存分配之后再根据线程束中的各个线程的实际需要对各个线程进行各自的内存分配，使得各个线程不需要争抢锁也能获得其各自所需的内存空间。后续还将讨论各个线程如何高效且不冲突地获得其各自所需的内存空间。

而且，由于以线程束为单位进行内存分配，因此在一个线程束的内存被分配好后，利用其它线程束继续争抢锁和分配内存的操作时间，可以在该操作时间中并行地进行将该线程束的内存划分给线程束中的各个线程的操作。相比于所有线程分别争抢锁和分配内存，本申请的实施例可以减少内存分配的操作时间，极大地提高内存分配效率。

如此，根据本申请的实施例，可以获得更好的操作性能和内存分配效率，例如可以执行少得多的原子操作，且有更少的原子操作冲突的概率，并且只需要执行很少的指令即可完成。

本申请提供如下方面：

方面1． 一种内存分配方法，包括如下步骤：

响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；

计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量；

选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；

为所述相应线程束分配内存空间；

根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

方面2． 根据方面1所述的方法，其中，所述方法还包括如下步骤：

根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的各自的前缀和，作为各个线程的内存地址偏移量；

其中，所述根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程的步骤包括：

按照所述各个线程的内存地址偏移量，在被分配的所述相应线程束的内存空间的量中、将偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。

方面3． 根据方面2所述的方法，其中，所述根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的前缀和，作为各个线程的内存地址偏移量的步骤包括如下步骤：

将所述相应线程束中的第一线程的前缀和设置为0；

计算所述相应线程束中的后续线程的前缀和为所述后续线程的前一线程的前缀和加上所述后续线程的请求的内存量。

方面4． 根据方面1所述的方法，其中所述每个线程束的大小被配置为预定量个线程，所述响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束的步骤包括：

将所述多个线程分配到所述多个线程束中并添加或不添加不需要执行的附加线程使得每个线程束中包括的线程的数量等于所述预定量。

方面5． 根据方面1所述的方法，其中，所述方法还包括：

如果所述线程的请求的内存量小于内存分配的预定最小量，将所述线程的请求的内存量扩大为所述内存分配的预定最小量。

方面6． 根据方面1所述的方法，其中，所述多个线程是在单指令多线程SIMT执行系统中用于执行单个指令的多个线程。

方面7． 一种内存分配系统，包括如下装置：

线程束分组装置，被配置为响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；

线程束请求装置，被配置为计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量，并选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；

内存分配装置，被配置为为所述相应线程束分配内存空间；

线程束内存划分装置，被配置为根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

方面8． 根据方面7所述的系统，其中，所述线程束请求装置还被配置为：

根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的各自的前缀和，作为各个线程的内存地址偏移量；

其中，所述线程束内存划分装置被配置为：

按照所述各个线程的内存地址偏移量，在被分配的所述相应线程束的内存空间的量中、将偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。

方面9． 根据方面8所述的系统，其中，所述线程束请求装置还被配置为通过如下步骤来根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的前缀和，作为各个线程的内存地址偏移量：

将所述相应线程束中的第一线程的前缀和设置为0；

计算所述相应线程束中的后续线程的前缀和为所述后续线程的前一线程的前缀和加上所述后续线程的请求的内存量。

方面10． 根据方面7所述的系统，其中所述每个线程束的大小被配置为预定量个线程，所述线程束分组装置被配置为：

将所述多个线程分配到所述多个线程束中并添加或不添加不需要执行的附加线程使得每个线程束中包括的线程的数量等于所述预定量。

方面11． 根据方面7所述的系统，其中，所述线程束请求装置还被配置为：

如果所述线程的请求的内存量小于内存分配的预定最小量，将所述线程的请求的内存量扩大为所述内存分配的预定最小量。

方面12． 根据方面7所述的系统，其中，所述多个线程是在单指令多线程SIMT执行系统中用于执行单个指令的多个线程。

方面13． 一种内存分配装置，包括：

处理器；

存储器，存储了计算器可执行指令，其被处理器运行时进行如方面1-6中任一所述的方法。

方面14． 一种计算机可读介质，存储有计算机可读指令，其中当所述计算机可读指令由处理器运行时，进行如方面1-6中任一所述的方法。

当然，上述的具体实施例仅是例子而非限制，且本领域技术人员可以根据本申请的构思从上述分开描述的各个实施例中合并和组合一些步骤和装置来实现本申请的效果，这种合并和组合而成的实施例也被包括在本申请中，在此不一一描述这种合并和组合。

注意，在本公开中提及的优点、优势、效果等仅是示例而非限制，不能认为这些优点、优势、效果等是本申请的各个实施例必须具备的。另外，上述公开的具体细节仅是为了示例的作用和便于理解的作用，而非限制，上述细节并不限制本申请为必须采用上述具体的细节来实现。

本公开中涉及的器件、装置、设备、系统的方框图仅作为例示性的例子并且不意图要求或暗示必须按照方框图示出的方式进行连接、布置、配置。如本领域技术人员将认识到的，可以按任意方式连接、布置、配置这些器件、装置、设备、系统。诸如“包括”、“包含”、“具有”等等的词语是开放性词汇，指“包括但不限于”，且可与其互换使用。这里所使用的词汇“或”和“和”指词汇“和/或”，且可与其互换使用，除非上下文明确指示不是如此。这里所使用的词汇“诸如”指词组“诸如但不限于”，且可与其互换使用。

本公开中的步骤示意流程图以及以上方法描述仅作为例示性的例子并且不意图要求或暗示必须按照给出的顺序进行各个实施例的步骤。如本领域技术人员将认识到的，可以按任意顺序进行以上实施例中的步骤的顺序。诸如“其后”、“然后”、“接下来”等等的词语不意图限制步骤的顺序；这些词语仅用于引导读者通读这些方法的描述。此外，例如使用冠词“一个”、“一”或者“该”对于单数的要素的任何引用不被解释为将该要素限制为单数。

另外，本文中的各个实施例中的步骤和装置并非仅限定于某个实施例中实行，事实上，可以根据本申请的概念来结合本文中的各个实施例中相关的部分步骤和部分装置以构思新的实施例，而这些新的实施例也包括在本申请的范围内。

以上描述的方法的各个操作可以通过能够进行相应的功能的任何适当的手段而进行。该手段可以包括各种硬件和/或软件组件和/或模块，包括但不限于硬件的电路、专用集成电路（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）。此外，措辞“示例的”不意味着描述的例子是优选的或者比其他例子更好。

可以不脱离由所附权利要求定义的教导的技术而进行对在此描述的技术的各种改变、替换和更改。此外，本公开的权利要求的范围不限于以上描述的处理、机器、制造、事件的组成、手段、方法和动作的具体方面。可以利用与在此描述的相应方面进行基本相同的功能或者实现基本相同的结果的当前存在的或者稍后要开发的处理、机器、制造、事件的组成、手段、方法或动作。因而，所附权利要求包括在其范围内的这样的处理、机器、制造、事件的组成、手段、方法或动作。

提供所公开的方面的以上描述以使本领域的任何技术人员能够做出或者使用本申请。对这些方面的各种修改对于本领域技术人员而言是非常显而易见的，并且在此定义的一般原理可以应用于其他方面而不脱离本申请的范围。因此，本申请不意图被限制到在此示出的方面，而是按照与在此公开的原理和新颖的特征一致的最宽范围。

为了例示和描述的目的已经给出了以上描述。此外，此描述不意图将本申请的实施例限制到在此公开的形式。尽管以上已经讨论了多个示例方面和实施例，但是本领域技术人员将认识到其某些变型、修改、改变、添加和子组合。

Claims (9)

1.一种内存分配方法，包括如下步骤：

响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；

计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量；

选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；

为所述相应线程束分配内存空间；

根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

2.根据权利要求1所述的方法，其中，所述方法还包括如下步骤：

根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的各自的前缀和，作为各个线程的内存地址偏移量；

其中，所述根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程的步骤包括：

按照所述各个线程的内存地址偏移量，在被分配的所述相应线程束的内存空间的量中、将偏移了相应内存地址偏移量的内存地址作为起始地址、按相应线程的请求的内存量划分给相应线程。

3.根据权利要求2所述的方法，其中，所述根据相应线程束中的各个线程的请求的内存量，计算各个线程的请求的内存量在相应线程束中的前缀和，作为各个线程的内存地址偏移量的步骤包括如下步骤：

将所述相应线程束中的第一线程的前缀和设置为0；

计算所述相应线程束中的后续线程的前缀和为所述后续线程的前一线程的前缀和加上所述后续线程的请求的内存量。

4.根据权利要求1所述的方法，其中所述每个线程束的大小被配置为预定量个线程，所述响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束的步骤包括：

将所述多个线程分配到所述多个线程束中并添加或不添加不需要执行的附加线程使得每个线程束中包括的线程的数量等于所述预定量。

5.根据权利要求1所述的方法，其中，所述方法还包括：

如果所述线程的请求的内存量小于内存分配的预定最小量，将所述线程的请求的内存量扩大为所述内存分配的预定最小量。

6.根据权利要求1所述的方法，其中，所述多个线程是在单指令多线程SIMT执行系统中用于执行单个指令的多个线程。

7.一种内存分配系统，包括如下装置：

线程束分组装置，被配置为响应于多个线程的各自的内存分配请求，将所述多个线程分组为多个线程束，其中，所述线程的内存分配请求包括所述线程的请求的内存量；

线程束请求装置，被配置为计算多个线程束中的相应线程束中的各个线程的请求的内存量的总和，作为所述相应线程束的请求的内存量，并选择所述相应线程束中的一个线程来请求分配所述相应线程束的内存空间，所述相应线程束的内存空间的量是所述相应线程束的请求的内存量；

内存分配装置，被配置为为所述相应线程束分配内存空间；

线程束内存划分装置，被配置为根据所述相应线程束中的各个线程的请求的内存量，将被分配的所述相应线程束的内存空间划分给各个线程。

8.一种内存分配装置，包括：

处理器；

存储器，存储了计算器可执行指令，其被处理器运行时进行如权利要求1-6中任一所述的方法。

9.一种计算机可读介质，存储有计算机可读指令，其中当所述计算机可读指令由处理器运行时，进行如权利要求1-6中任一所述的方法。

Images