Memory allocation method, system, device and computer readable medium

Abstract

A memory allocation method, system, apparatus, and computer-readable medium are provided. The method includes the following steps: in response to respective memory allocation requests from multiple threads, grouping the multiple threads into multiple warps, wherein the memory allocation requests of the threads include memory amounts requested by the threads; calculating the sum of the memory amounts requested by the threads in a corresponding warp among the multiple warps as the memory amount requested by the corresponding warp; selecting a thread in the corresponding warp to request allocation of memory space for the corresponding warp, the memory amount of the corresponding warp being the memory amount requested by the corresponding warp; allocating memory space to the corresponding warp; and dividing the allocated memory space of the corresponding warp among the threads according to the memory amounts requested by the threads in the corresponding warp.

Description

Technical Field

[0001] The present application relates to the field of processors, and more particularly to a memory allocation method, system, device, and computer-readable medium. Background Art

[0002] Modern processor architectures utilize various execution models, such as Single Instruction Multiple Data (SIMD) and Single Instruction Multiple Threads (SIMT). In graphics processing units (GPUs), multiprocessors often employ a SIMT architecture to efficiently manage and execute multiple threads.

[0003] In SIMT, multiple threads execute the same instruction on different data sets. Whenever the graphics processing unit (GPU) needs to execute a specific instruction, it fetches the data and instruction from memory, decodes it, and executes it. In this case, all data sets that need to be executed using the same instruction (up to a certain limit) are prefetched and executed simultaneously using the various threads available to the processor. Compared to single instruction multiple data stream SIMD, the advantage of SIMT is that developers do not need to work hard to put the data into the appropriate vector length, and SIMT allows each thread to have different branches and reduces the waiting time caused by instruction prefetching.

[0004] During thread execution, the thread requests memory allocation from the processor core. Currently, many different memory allocation mechanisms exist within the processor core. For example, each thread uses the malloc() function to request and allocate memory, with its parameters including the requested memory size. However, existing memory allocation methods are still not performant and fast enough, and there is a need for improvement.

[0005] There is still a need for a high-performance, high-speed memory allocation mechanism. Summary of the Invention

[0006] Various aspects of the present application are provided to solve one or more problems existing in the prior art.

[0007] According to one aspect of the present application, a memory allocation method is provided, comprising the following steps: in response to respective memory allocation requests of a plurality of threads, grouping the plurality of threads into a plurality of thread bundles, wherein the memory allocation requests of the threads include the memory amounts requested by the threads; calculating the sum of the memory amounts requested by the threads in corresponding thread bundles among the plurality of thread bundles as the memory amount requested by the corresponding thread bundle; selecting one thread in the corresponding thread bundle to request allocation of memory space of the corresponding thread bundle, the amount of the memory space of the corresponding thread bundle being the memory amount requested by the corresponding thread bundle; allocating memory space to the corresponding thread bundle; and dividing the allocated memory space of the corresponding thread bundle to each thread according to the memory amounts requested by the threads in the corresponding thread bundle.

[0008] According to another aspect of the present application, a memory allocation system is provided, comprising the following apparatus: a warp grouping apparatus configured to group a plurality of threads into a plurality of warps in response to respective memory allocation requests of the plurality of threads, wherein the memory allocation requests of the threads include the memory amounts requested by the threads; a warp requesting apparatus configured to calculate the sum of the memory amounts requested by the respective threads in a corresponding warp among the plurality of warps as the memory amount requested by the corresponding warp, and select one thread in the corresponding warp to request allocation of memory space of the corresponding warp, the memory amount of the corresponding warp being the memory amount requested by the corresponding warp; a memory allocating apparatus configured to allocate memory space to the corresponding warp; and a warp memory partitioning apparatus configured to partition the allocated memory space of the corresponding warp to the respective threads according to the memory amounts requested by the respective threads in the corresponding warp.

[0009] According to another aspect of the present application, a memory allocation device is provided, comprising: a processor; and a memory storing computer executable instructions, which, when executed by the processor, perform various methods according to the embodiments of the present application.

[0010] According to another aspect of the present application, a computer-readable medium is provided, storing computer-readable instructions, wherein when the computer-readable instructions are executed by a processor, various methods according to the embodiments of the present application are performed.

[0011] This application leverages the characteristics of warps to allocate memory to multiple threads in warps rather than per thread. After allocating memory per warp, each thread in the warp is allocated its own memory based on its actual needs. This results in better operational performance and memory allocation efficiency, for example, requiring fewer atomic operations, a lower probability of atomic operation conflicts, and requiring only a few instructions to complete the operation. BRIEF DESCRIPTION OF THE DRAWINGS

[0012] In order to more clearly illustrate the embodiments of the present disclosure or the technical solutions in the prior art, the following briefly introduces the drawings required for use in the embodiments or the description of the prior art. Obviously, the drawings described below are only some embodiments of the present disclosure. For ordinary technicians in this field, other drawings can be obtained based on these drawings without any creative work.

[0013] Figure 1 A schematic diagram showing the operation of the malloc function that allocates its own memory space to each thread.

[0014] Figure 2 A schematic flow chart of a memory allocation method according to an embodiment of the present application is shown.

[0015] Figure 3 Schematically shows the Figure 2 The operation process of the memory allocation method is shown.

[0016] Figure 4 A schematic flowchart of a memory allocation method according to another embodiment of the present application is shown.

[0017] Figure 5 Shown Figure 4 Shows example results produced by the various steps of the memory allocation method.

[0018] Figure 6 A schematic block diagram of a memory allocation system according to an embodiment of the present application is shown.

[0019] Figure 7 A block diagram of an exemplary computer system suitable for implementing embodiments of the present application is shown.

[0020] Figure 8 A schematic diagram of a non-transitory computer-readable storage medium according to an embodiment of the present disclosure is shown. DETAILED DESCRIPTION

[0021] Reference will now be made in detail to specific embodiments of the present application, examples of which are illustrated in the accompanying drawings. Although the present application will be described in conjunction with specific embodiments, it will be understood that the present application is not intended to be limited to the described embodiments. On the contrary, it is intended to cover variations, modifications, and equivalents within the spirit and scope of the present application as defined by the appended claims. It should be noted that the method steps described herein can be implemented by any functional block or functional arrangement, and any functional block or functional arrangement can be implemented as a physical entity or a logical entity, or a combination of the two.

[0022] In the prior art, in scenarios with multiple processor cores, the processor cores require dynamically sized global memory (predetermined space reserved by the driver). Large blocks of data that cannot fit into shared memory are stored within a single processor core. Data can be exchanged between multiple processor cores. Allocating memory space to threads is achieved, for example, using the following software code:

[0023] __global__ void example() {

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

[0025] }

[0026] Here, malloc(threadIdx.x+1) allocates memory from global memory to all threads in the entire thread block. A thread block is a concept summarized by the OpenCL work group, representing a group of threads running in the same execution unit and sharing the same shared memory.

[0027] Each thread in the entire thread block is then allocated its own memory space. Figure 1 A schematic diagram showing the operation of the malloc function that allocates its own memory space to each thread.

[0028] like Figure 1 As shown, it is assumed that there are threads T0-T9 in the thread block. In the process of allocating its own memory space to each thread, the 10 threads respectively issue 10 respective memory allocation requests.

[0029] Because of the lock operation during the memory allocation process. A lock is a tool used to control access to shared memory resources by multiple threads. Typically, a lock provides exclusive access to a shared resource, so only one thread can acquire the lock at a time, and all access to the shared resource requires first acquiring the lock.

[0030] Each thread must first acquire a lock before it can access memory resources to allocate memory for itself. For example, the thread acquires the lock by using the non-blocking tryLock() function.

[0031] Therefore, all threads will try to acquire the lock. Suppose thread T0 acquires the lock and executes the malloc memory allocation operation for thread T0. The malloc memory allocation operation is an atomic operation. After the atomic operation is completed, other threads will try to acquire the lock again. Suppose thread T3 acquires the lock at this time and executes the malloc memory allocation atomic operation for thread T3. After the atomic operation is completed, the lock is released. Other threads will try to acquire the lock again. Suppose thread T2 acquires the lock at this time and executes the malloc memory allocation atomic operation for thread T2. After the atomic operation is completed, the lock is released. Other threads will try to acquire the lock again. Suppose thread T9 acquires the lock at this time and executes the malloc memory allocation atomic operation for thread T9. After the atomic operation is completed, the lock is released. Other threads will try to acquire the lock again. And so on, until all threads have acquired the lock and completed the memory allocation atomic operation.

[0032] The downside of this approach is that it leads to lower processor performance. First, issuing too many atomic operations puts pressure on hardware such as the memory bus and caches. Second, there is a high number of lock conflicts, which can cause loops to wait for locks, resulting in more instructions being executed and reduced execution efficiency.

[0033] The present application considers solving the above-mentioned problems. In the thread scheduling of the graphics processing unit (GPU) hardware, the minimum unit of thread scheduling is usually a warp, that is, one warp is executed at a time in the graphics processing unit (GPU). Generally, a warp can contain 32 threads or other numbers of threads. The present application utilizes the characteristics of warps, and multiple threads are allocated memory in warps instead of one thread, and after allocating memory in warps, each thread is allocated its own memory according to the actual needs of each thread in the warp. In this way, better operating performance and memory allocation efficiency can be obtained, for example, far fewer atomic operations can be executed, and there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete it.

[0034] The following describes the details of various embodiments of the present application through the accompanying drawings.

[0035] Figure 2 A schematic flow chart of a memory allocation method 200 according to an embodiment of the present application is shown.

[0036] like Figure 2As shown, the memory allocation method 200 includes the following steps: step 202, in response to respective memory allocation requests of the plurality of threads, grouping the plurality of threads into a plurality of warps, wherein the memory allocation requests of the threads include the memory amounts requested by the threads; step 204, calculating the sum of the memory amounts requested by the threads in a corresponding warp among the plurality of warps as the memory amount requested by the corresponding warp; step 206, selecting one thread in the corresponding warp to request allocation of memory space of the corresponding warp, the memory amount of the corresponding warp being the memory amount requested by the corresponding warp; step 208, allocating memory space to the corresponding warp; and step 210, dividing the allocated memory space of the corresponding warp to each thread according to the memory amounts requested by the threads in the corresponding warp.

[0037] Combine Figure 3 Let's describe the above steps in detail. Figure 3 Schematically shows the Figure 2 The operation process of the memory allocation method 200 is shown.

[0038] like Figure 3 As shown, in step 202, in response to respective memory allocation requests from the plurality of threads, the plurality of threads are grouped into a plurality of warps, such as warp 0, warp 1, and so on. The memory allocation request of each thread includes the requested amount of memory for the thread. Assume that there are 32 threads in warp 0, 32 threads in warp 1, and so on.

[0039] In step 204, the total memory requested by each thread in each of the multiple warps is calculated as the total memory requested by the corresponding warp. This is done by combining the requests of each thread in the warp. For example, suppose thread 1 in warp 0 requests 1MB of memory, thread 2 requests 2MB of memory, thread 3 requests 1MB of memory, and so on. The total memory requested by each of the 32 threads in warp 0 is assumed to be 50MB of memory, the total memory requested by each of the 32 threads in warp 1 is assumed to be 30MB of memory, and so on.

[0040] Then, in step 206, a thread in the corresponding warp is selected as the requester to request allocation of memory space of the corresponding warp. At this time, memory allocation requests of other threads in the corresponding warp can be ignored because only one thread is required as the requester.

[0041] The amount of memory space of the corresponding warp requested by the thread is the requested memory amount of the corresponding warp, that is, 50M memory for warp 0, 30M memory for warp 1, and so on, as calculated previously.

[0042] In step 208, memory space is allocated to the corresponding warps. For example, a 50MB memory block (with an address range of XXXXXXXXXXXX to XXXXXXXXXXXX, or a starting address of XXXXXXXXXXXX) is allocated to warp 0, and a 30MB memory block (with an address range of XXXXXXXXXXXX to XXXXXXXXXXXX, or a starting address of XXXXXXXX) is allocated to warp 1. Allocating memory space to the corresponding warps also requires each warp to perform atomic operations such as acquiring a lock, allocating memory space, and releasing a lock, which will not be detailed here.

[0043] In step 210, the allocated memory space of the corresponding warp is divided among the threads according to the memory requests of the threads in the corresponding warp. For example, if warp 0 is allocated 50MB of memory, 1MB of this 50MB is allocated to thread 1, 2MB to thread 2, 1MB to thread 3, and so on.

[0044] The embodiments of the present application utilize the characteristics of thread warps. Since the thread warp itself is the smallest unit of thread scheduling, memory allocation is performed by multiple threads in units of thread warps rather than in units of single threads. This allows for efficient memory allocation using existing thread scheduling rules, making the memory allocation and thread scheduling units the same, thereby achieving operational consistency.

[0045] Moreover, since memory allocation is performed in thread bundles, that is, the atomic operations of acquiring locks and allocating memory space and releasing locks, the number of times multiple threads acquire locks separately and the conflicts caused by multiple threads competing for locks can be reduced, and fewer atomic operations are also executed.

[0046] Furthermore, the embodiments of the present application allocate memory on a warp-by-warp basis, and then allocate memory to each thread based on its actual needs. This allows each thread to obtain the memory space it needs without having to contend for locks. We will discuss how each thread can efficiently and conflict-freely obtain the memory space it needs later.

[0047] Furthermore, because memory allocation is performed on a warp-by-warp basis, after a warp's memory is allocated, the warp's memory can be allocated to each thread in the warp in parallel, using the time that other warps continue to contend for locks and allocate memory. Compared to having all threads contend for locks and allocate memory separately, embodiments of the present application can reduce memory allocation operation time and significantly improve memory allocation efficiency.

[0048] In this way, according to the embodiments of the present application, better operating performance and memory allocation efficiency can be obtained, for example, far fewer atomic operations can be executed, there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete it.

[0049] Here, the concept of thread bundle in this application corresponds to warp, but it may correspond to other terms in other processors. This application does not limit the specific terminology, but only needs to group multiple threads into thread groups (thread bundles) that are suitable for common processing in the processor. If the thread group (thread bundle) happens to be the same as the rules in the existing thread processing process in the processor, the existing rules can be used more efficiently to develop various embodiments of this application without changing the existing rules or adding hardware devices or software instructions.

[0050] Figure 4 FIG. 4 is a schematic flow chart of a memory allocation method 400 according to another embodiment of the present application. Figure 2 The same steps in the illustrated memory allocation method 200 are numbered the same. Figure 5 Shown Figure 4 Example results generated by the various steps of the memory allocation method 400 are shown.

[0051] The following combination Figure 4 and Figure 5 The process of a memory allocation method 400 according to another embodiment of the present application is introduced.

[0052] like Figure 4 As shown, in step 202 , in response to respective memory allocation requests of the plurality of threads, the plurality of threads are grouped into a plurality of warps, wherein the memory allocation requests of the threads include the memory amounts requested by the threads.

[0053] like Figure 4 As shown, in one embodiment, the size of each thread warp is configured to be a predetermined number of threads, and in response to respective memory allocation requests of the plurality of threads, step 202 of grouping the plurality of threads into a plurality of thread warps includes: step 2021 of allocating the plurality of threads to the plurality of thread warps and adding or not adding additional threads that do not need to be executed so that the number of threads included in each thread warp is equal to the predetermined number.

[0054] Combine Figure 5 , assuming that the thread block size is 64 threads and the predetermined amount of each thread warp is 32, at this time the number of multiple threads can be divided by the predetermined amount, then the 64 threads are divided into 2 thread warps, and there is no need to add additional threads that do not need to be executed.

[0055] Of course, if the number of multiple threads cannot be divided evenly by the predetermined amount, if the thread block size is 60 threads and the predetermined amount of each thread warp is 32, then the 60 threads are divided into one thread warp including 32 threads, and another thread warp including 28 threads, plus 4 additional threads that do not need to be executed. The additional threads that do not need to be executed are inactive and may not be allocated addresses and executed during the actual address allocation process.

[0056] Assume that each thread requests 2 bytes (B) of memory space, such as Figure 5 The "2" shown here represents 2 bytes of memory space requested by each thread. Of course, this is not a limitation. In practice, the memory space requested by each thread may not be exactly the same, and the memory space size may also be other sizes.

[0057] like Figure 4 As shown, in one embodiment, the method may further include: step 203, if the memory amount requested by the thread is less than the predetermined minimum amount of memory allocation, expanding the memory amount requested by the thread to the predetermined minimum amount of memory allocation.

[0058] Combine Figure 5 Assuming that the predetermined minimum amount of memory allocation is 4 bytes, the 2-byte memory amount requested by the thread is expanded to 4 bytes in step 203 to meet the predetermined minimum amount of memory allocation. If the memory amount requested by the thread is greater than or equal to the predetermined minimum amount of memory allocation, the memory amount requested by the thread is not changed.

[0059] This can increase memory allocation efficiency and comply with memory allocation rules.

[0060] like Figure 4 As shown, in step 204 , the sum of the memory amounts requested by the threads in the corresponding warp among the multiple warps is calculated as the memory amount requested by the corresponding warp.

[0061] Combine Figure 5 In this example, the memory request amount for each thread is expanded to 4 bytes. Therefore, the total memory request amount for each of the 32 threads in the first warp is 128 bytes, which serves as the memory request amount for the first warp. The total memory request amount for each of the 32 threads in the second warp is also 128 bytes, which serves as the memory request amount for the second warp.

[0062] like Figure 4 As shown, the method 400 may further include the following step 205: according to the memory amount requested by each thread in the corresponding warp, calculating the prefix sum (prefixSum) of the memory amount requested by each thread in the corresponding warp as the memory address offset of each thread.

[0063] like Figure 4 As shown, in one embodiment, step 205 of calculating the prefix sum of the memory amount requested by each thread in the corresponding warp as the memory address offset of each thread includes the following steps: step 2051, setting the prefix sum of the first thread in the corresponding warp to 0; step 2052, calculating the prefix sum of subsequent threads in the corresponding warp as the prefix sum of the previous thread of the subsequent thread plus the memory amount requested by the subsequent thread.

[0064] Figure 4 Step 205 is depicted between step 204 and step 206 , but this is not necessary. In fact, step 205 can be performed at any time as long as the memory address offset of each thread can be known in step 2101 .

[0065] Combine Figure 5 In step 205, the prefix sum of thread T0 in the first warp is 0, the prefix sum of thread T1 in the first warp is 0+4=4 bytes, the prefix sum of thread T2 in the first warp is 4+4=8 bytes, the prefix sum of thread T3 in the first warp is 8+4=12 bytes, and so on. The prefix sum of thread T31 in the first warp is 120+4=124 bytes.

[0066] Similarly, the prefix sum of thread T32 in the second thread warp is 0, the prefix sum of thread T33 in the second thread warp is 0+4=4 bytes, the prefix sum of thread T34 in the second thread warp is 4+4=8 bytes, the prefix sum of thread T3 in the second thread warp is 8+4=12 bytes, and so on. The prefix sum of thread T63 in the second thread warp is 120+4=124 bytes.

[0067] Since the prefix sum can be calculated using the existing function prefixSum(), the definition of prefixSum() is as follows:

[0068] Given an array A[0..n-1], for i For each i (i is an integer), the prefix sum is calculated as: PrefixSum[i] = A[0]+A[1]+...+A[i-1];

[0069] For example: A[5,8,9,2] --> PrefixSum[5,13,22,24], that is:

[0070] PrefixSum[0] =A[0] ;

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

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

[0073] PrefixSum[3] =A[0] + A[1] + A[2] + A[3].

[0074] In the above example, A[0]=5, but in the embodiment of the present application, in order to better adapt to the division of memory by each thread, A[0]=0 is set.

[0075] Specifically, according to the embodiments of the present application, by cleverly utilizing existing prefix sum calculation methods, it is possible to easily associate the prefix sum with the memory address offset required by each thread. This means that the prefix sum of the memory space required by each thread can be used to determine the address offset of the memory space required by each thread, allowing for subsequent partitioning of the memory address space allocated to the entire warp for each thread. This approach does not increase the computational difficulty or hardware complexity of the subsequent partitioning.

[0076] like Figure 4 As shown, in step 206 , a thread in the corresponding warp is selected to request allocation of the memory space of the corresponding warp, and the amount of the memory space of the corresponding warp is the requested memory amount of the corresponding warp.

[0077] Combine Figure 5 , a thread in the corresponding warp is selected to request memory allocation for that warp. This thread initiates malloc(). This thread can be the first thread in the warp, for example, thread T0 in the first warp, or thread T32 in the second warp. Of course, any thread in the warp can be the requester. Other threads do not need to initiate requests, as each warp only needs to request memory allocation once.

[0078] In this way, one thread in each thread warp is used as the memory requester, so that each thread warp only makes one memory allocation request, reducing the atomic operations of acquiring locks and allocating memory space, as well as the operations of releasing locks. This can reduce the number of times multiple threads acquire locks separately and the conflicts caused by multiple threads competing for locks.

[0079] like Figure 4 As shown, in step 208 , memory space is allocated to the corresponding warp.

[0080] Combine Figure 5, assuming that memory space is allocated for the first thread warp, with a starting address of 0x7babcd0000 and a length of 128B (that is, the total memory space required to be allocated by all threads in the first thread warp), and memory space is allocated for the second thread warp, with a starting address of 0x7bbcde0000 and a length of 128B (that is, the total memory space required to be allocated by all threads in the second thread warp).

[0081] like Figure 4 As shown, in step 210 , the allocated memory space of the corresponding warp is divided among the threads according to the memory amount requested by the threads in the corresponding warp.

[0082] like Figure 4 As shown, step 210 of dividing the allocated memory space of the corresponding thread warp to each thread according to the memory amount requested by each thread in the corresponding thread warp may include: step 2101, according to the memory address offset of each thread, using a memory address offset by the corresponding memory address offset as a starting address in the amount of memory space allocated to the corresponding thread warp, and dividing the memory amount requested by the corresponding thread to the corresponding thread.

[0083] Combine Figure 5 The memory space of the first warp starts at 0x7babcd0000. The memory address offset of each thread is calculated based on the prefix sum. The memory address offset by the corresponding memory address offset is used as the starting address, and the memory is allocated to the corresponding thread according to the amount of memory requested by the thread. For example, the memory space of thread T0 starts at 0x7babcd0000 and has a size of 4 bytes. The memory space of thread T1 starts at 0x7babcd0004 and has a size of 4 bytes. The memory space of thread T2 starts at 0x7babcd0008 and has a size of 4 bytes. The memory space of thread T3 starts at 0x7babcd000C and has a size of 4 bytes, and so on.

[0084] The memory space of the second warp starts at 0x7bbcde0000. The memory address offset of each thread is calculated based on the prefix sum. The memory address offset by the corresponding memory address offset is used as the starting address, and the memory is allocated to each thread according to the amount of memory requested by the thread. For example, the memory space of thread T32 starts at 0x7bbcde 0000 and has a size of 4 bytes. The memory space of thread T1 starts at 0x7bbcde 0004 and has a size of 4 bytes. The memory space of thread T2 starts at 0x7bbcde 0008 and has a size of 4 bytes. The memory space of thread T3 starts at 0x7bbcde 000C and has a size of 4 bytes, and so on.

[0085] In this way, since the memory address offset of each thread is obtained by calculating the prefix sum of the memory amount requested by each thread in the corresponding thread warp, after the address range of memory allocation is obtained in units of thread warps, the allocated address range can be divided into each thread according to the memory address offset with the offset as the starting address.

[0086] According to various embodiments of the present application, the entire process of grouping multiple threads into warps, allocating memory space to the warps, and dividing the memory space among the threads in the warps can be performed easily and efficiently.

[0087] In one embodiment, the multiple threads may be multiple threads for executing a single instruction in a SIMT execution system. Thus, the memory allocation method may be performed on a group of threads executing a single instruction, thereby improving the execution efficiency of the single instruction.

[0088] The embodiments of the present application utilize the characteristics of thread warps. Since the thread warp itself is the smallest unit of thread scheduling, memory allocation is performed by multiple threads in units of thread warps rather than in units of single threads. This allows for efficient memory allocation using existing thread scheduling rules, making the memory allocation and thread scheduling units the same, thereby achieving operational consistency and saving the loop logic of each thread that executes separate thread memory allocation operations.

[0089] Furthermore, because memory allocation is performed on a warp-by-warp basis, the number of atomic operations is, for example, one-thirtieth of that performed on a thread-by-thread basis, resulting in fewer atomic operations. Using one thread in each warp as the memory requester ensures that each warp only makes one memory allocation request, reducing the number of atomic operations required to acquire locks, allocate memory space, and release locks. This reduces the number of lock acquisitions by multiple threads and the conflicts caused by lock contention.

[0090] Furthermore, the embodiments of the present application allocate memory on a warp-by-warp basis, and then allocate memory to each thread based on its actual needs. This allows each thread to obtain the memory space it needs without having to contend for locks. We will discuss how each thread can efficiently and conflict-freely obtain the memory space it needs later.

[0091] Furthermore, because memory allocation is performed on a warp-by-warp basis, after a warp's memory is allocated, the warp's memory can be allocated to each thread in the warp in parallel, using the time that other warps continue to contend for locks and allocate memory. Compared to having all threads contend for locks and allocate memory separately, embodiments of the present application can reduce memory allocation operation time and significantly improve memory allocation efficiency.

[0092] In this way, according to the embodiments of the present application, each thread warp only adds a small number of thread warp grouping operations and some instructions, but can obtain better operation performance and memory allocation efficiency. For example, far fewer atomic operations can be executed, and there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete.

[0093] Figure 6 A schematic block diagram of a memory allocation system according to an embodiment of the present application is shown.

[0094] like Figure 6 As shown, the memory allocation system 600 includes the following devices: a thread warp grouping device 601, configured to group multiple threads into multiple thread warps in response to respective memory allocation requests of multiple threads from a memory allocation (Malloc) interface of the system, wherein the memory allocation request of the thread includes the memory amount requested by the thread; a thread warp requesting device 602, configured to calculate the sum of the memory amounts requested by each thread in a corresponding thread warp among the multiple thread warps as the memory amount requested by the corresponding thread warp, and select one thread in the corresponding thread warp to request allocation of memory space of the corresponding thread warp, the amount of memory space of the corresponding thread warp being the memory amount requested by the corresponding thread warp; a memory allocating device 603, configured to allocate memory space for the corresponding thread warp from the global memory of the system; a thread warp memory partitioning device 604, configured to partition the allocated memory space of the corresponding thread warp to each thread according to the memory amount requested by each thread in the corresponding thread warp.

[0095] In one embodiment, the warp requesting means 602 is further configured to: calculate, based on the memory amounts requested by each thread in the corresponding warp, a prefix sum of the memory amounts requested by each thread in the corresponding warp, as the memory address offset of each thread. The warp memory partitioning means 604 is configured to: use the memory address offset of each thread as the starting address within the allocated memory space of the corresponding warp, and partition the memory amount requested by the thread to the thread in accordance with the memory address offset of the thread.

[0096] In one embodiment, the warp requesting device 602 is further configured to calculate the prefix sum of the memory amount requested by each thread in the corresponding warp according to the memory amount requested by each thread in the corresponding warp, as the memory address offset of each thread, by performing the following steps: setting the prefix sum of the first thread in the corresponding warp to 0; and calculating the prefix sum of subsequent threads in the corresponding warp as the prefix sum of the previous thread of the subsequent thread plus the memory amount requested by the subsequent thread.

[0097] In one embodiment, the size of each warp is configured to be a predetermined number of threads, and the warp grouping device 601 is configured to: allocate the multiple threads to the multiple warps and add or not add additional threads that do not need to be executed so that the number of threads included in each warp is equal to the predetermined number.

[0098] In one embodiment, the warp requesting device 602 is further configured to: if the memory amount requested by the thread is less than the predetermined minimum amount of memory allocation, increase the memory amount requested by the thread to the predetermined minimum amount of memory allocation.

[0099] In one embodiment, the multiple threads are multiple threads used to execute a single instruction in a single instruction multiple thread (SIMT) execution system.

[0100] The embodiments of the present application utilize the characteristics of thread warps. Since the thread warp itself is the smallest unit of thread scheduling, memory allocation is performed by multiple threads in units of thread warps rather than in units of single threads. This allows for efficient memory allocation using existing thread scheduling rules, making the memory allocation and thread scheduling units the same, thereby achieving operational consistency.

[0101] Moreover, since memory allocation is performed in thread bundles, that is, the atomic operations of acquiring locks and allocating memory space and releasing locks, the number of times multiple threads acquire locks separately and the conflicts caused by multiple threads competing for locks can be reduced, and fewer atomic operations are also executed.

[0102] Furthermore, the embodiments of the present application allocate memory on a warp-by-warp basis, and then allocate memory to each thread based on its actual needs. This allows each thread to obtain the memory space it needs without having to contend for locks. We will discuss how each thread can efficiently and conflict-freely obtain the memory space it needs later.

[0103] Furthermore, because memory allocation is performed on a warp-by-warp basis, after a warp's memory is allocated, the warp's memory can be allocated to each thread in the warp in parallel, using the time that other warps continue to contend for locks and allocate memory. Compared to having all threads contend for locks and allocate memory separately, embodiments of the present application can reduce memory allocation operation time and significantly improve memory allocation efficiency.

[0104] In this way, according to the embodiments of the present application, better operating performance and memory allocation efficiency can be obtained, for example, far fewer atomic operations can be executed, there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete it.

[0105] Figure 7 A block diagram of an exemplary computer system suitable for implementing embodiments of the present application is shown.

[0106] The computer system may include a processor (H1); a memory (H2) coupled to the processor (H1) and storing computer-executable instructions therein for performing the steps of the various methods of the embodiments of the present application when executed by the processor.

[0107] The processor (H1) may include, but is not limited to, one or more processors or microprocessors, etc.

[0108] The memory (H2) may include, but is not limited to, for example, random access memory (RAM), read-only memory (ROM), flash memory, EPROM memory, EEPROM memory, registers, computer storage media (such as hard disk, floppy disk, solid-state drive, removable disk, CD-ROM, DVD-ROM, Blu-ray disc, etc.).

[0109] In addition, the computer system may also include a data bus (H3), an input / output (I / O) bus (H4), a display (H5), and input / output devices (H6) (eg, keyboard, mouse, speakers, etc.).

[0110] The processor ( H1 ) may communicate with external devices ( H5 , H6 , etc.) via an I / O bus ( H4 ) via a wired or wireless network (not shown).

[0111] The memory (H2) may also store at least one computer-executable instruction for executing various functions and / or steps of methods in the embodiments described in this technology when the processor (H1) executes the instruction.

[0112] In one embodiment, the at least one computer executable instruction may also be compiled into or constitute a software product, wherein one or more computer executable instructions are executed by a processor to perform the various functions and / or method steps in the embodiments described in the present technology.

[0113] The embodiments of the present application utilize the characteristics of thread warps. Since the thread warp itself is the smallest unit of thread scheduling, memory allocation is performed by multiple threads in units of thread warps rather than in units of single threads. This allows for efficient memory allocation using existing thread scheduling rules, making the memory allocation and thread scheduling units the same, thereby achieving operational consistency.

[0114] Moreover, since memory allocation is performed in thread bundles, that is, the atomic operations of acquiring locks and allocating memory space and releasing locks, the number of times multiple threads acquire locks separately and the conflicts caused by multiple threads competing for locks can be reduced, and fewer atomic operations are also executed.

[0115] Furthermore, the embodiments of the present application allocate memory on a warp-by-warp basis, and then allocate memory to each thread based on its actual needs. This allows each thread to obtain the memory space it needs without having to contend for locks. We will discuss how each thread can efficiently and conflict-freely obtain the memory space it needs later.

[0116] Furthermore, because memory allocation is performed on a warp-by-warp basis, after a warp's memory is allocated, the warp's memory can be allocated to each thread in the warp in parallel, using the time that other warps continue to contend for locks and allocate memory. Compared to having all threads contend for locks and allocate memory separately, embodiments of the present application can reduce memory allocation operation time and significantly improve memory allocation efficiency.

[0117] In this way, according to the embodiments of the present application, better operating performance and memory allocation efficiency can be obtained, for example, far fewer atomic operations can be executed, there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete it.

[0118] Figure 8 A schematic diagram of a non-transitory computer-readable storage medium according to an embodiment of the present disclosure is shown.

[0119] like Figure 8As shown, computer-readable storage medium 820 stores instructions, such as computer-readable instructions 810. When computer-readable instructions 810 are executed by a processor, the various methods described above can be performed. Computer-readable storage media include, but are not limited to, volatile memory and / or non-volatile memory. Volatile memory may include, for example, random access memory (RAM) and / or cache memory. Non-volatile memory may include, for example, read-only memory (ROM), a hard disk, flash memory, etc. For example, computer-readable storage medium 820 may be connected to a computing device such as a computer. Then, when the computing device executes computer-readable instructions 810 stored on computer-readable storage medium 820, the various methods described above may be performed.

[0120] The embodiments of the present application utilize the characteristics of thread warps. Since the thread warp itself is the smallest unit of thread scheduling, memory allocation is performed by multiple threads in units of thread warps rather than in units of single threads. This allows for efficient memory allocation using existing thread scheduling rules, making the memory allocation and thread scheduling units the same, thereby achieving operational consistency.

[0121] Moreover, since memory allocation is performed in thread bundles, that is, the atomic operations of acquiring locks and allocating memory space and releasing locks, the number of times multiple threads acquire locks separately and the conflicts caused by multiple threads competing for locks can be reduced, and fewer atomic operations are also executed.

[0122] Furthermore, the embodiments of the present application allocate memory on a warp-by-warp basis, and then allocate memory to each thread based on its actual needs. This allows each thread to obtain the memory space it needs without having to contend for locks. We will discuss how each thread can efficiently and conflict-freely obtain the memory space it needs later.

[0123] Furthermore, because memory allocation is performed on a warp-by-warp basis, after a warp's memory is allocated, the warp's memory can be allocated to each thread in the warp in parallel, using the time that other warps continue to contend for locks and allocate memory. Compared to having all threads contend for locks and allocate memory separately, embodiments of the present application can reduce memory allocation operation time and significantly improve memory allocation efficiency.

[0124] In this way, according to the embodiments of the present application, better operating performance and memory allocation efficiency can be obtained, for example, far fewer atomic operations can be executed, there is less probability of atomic operation conflicts, and only a few instructions need to be executed to complete it.

[0125] This application provides the following aspects:

[0126] Aspect 1. A memory allocation method comprising the following steps:

[0127] grouping the plurality of threads into a plurality of warps in response to respective memory allocation requests of the plurality of threads, wherein the memory allocation requests of the threads include requested amounts of memory of the threads;

[0128] Calculating a sum of memory amounts requested by threads in a corresponding warp among the plurality of warps as the memory amount requested by the corresponding warp;

[0129] selecting a thread in the corresponding warp to request allocation of memory space of the corresponding warp, where the amount of the memory space of the corresponding warp is the requested memory amount of the corresponding warp;

[0130] Allocating memory space for the corresponding thread warp;

[0131] The allocated memory space of the corresponding warp is divided among the threads according to the memory amount requested by the threads in the corresponding warp.

[0132] Aspect 2. The method according to Aspect 1, further comprising the following steps:

[0133] According to the memory amount requested by each thread in the corresponding thread warp, a prefix sum of the memory amount requested by each thread in the corresponding thread warp is calculated as the memory address offset of each thread;

[0134] The step of dividing the allocated memory space of the corresponding warp to each thread according to the memory amount requested by each thread in the corresponding warp includes:

[0135] According to the memory address offset of each thread, in the amount of memory space allocated to the corresponding warp, a memory address offset by the corresponding memory address offset is used as a starting address, and the memory amount requested by the corresponding thread is divided to the corresponding thread.

[0136] Aspect 3. The method according to Aspect 2, wherein the step of calculating, based on the memory amounts requested by each thread in the corresponding warp, a prefix sum of the memory amounts requested by each thread in the corresponding warp and using the prefix sum as the memory address offset of each thread comprises the following steps:

[0137] Setting the prefix sum of the first thread in the corresponding warp to 0;

[0138] A prefix sum of a subsequent thread in the corresponding warp is calculated as the prefix sum of a previous thread of the subsequent thread plus the amount of memory requested by the subsequent thread.

[0139] Aspect 4. The method according to Aspect 1, wherein the size of each warp is configured to be a predetermined number of threads, and the step of grouping the plurality of threads into a plurality of warps in response to respective memory allocation requests of the plurality of threads comprises:

[0140] The plurality of threads are allocated to the plurality of warps and additional threads not required for execution are added or not added so that the number of threads included in each warp is equal to the predetermined amount.

[0141] Aspect 5. The method according to Aspect 1, further comprising:

[0142] If the memory amount requested by the thread is less than a predetermined minimum amount of memory allocation, the memory amount requested by the thread is increased to the predetermined minimum amount of memory allocation.

[0143] Aspect 6. The method according to aspect 1, wherein the plurality of threads are a plurality of threads for executing a single instruction in a single instruction multiple thread (SIMT) execution system.

[0144] Aspect 7. A memory allocation system comprising the following apparatus:

[0145] a warp grouping device configured to group the plurality of threads into a plurality of warps in response to respective memory allocation requests of the plurality of threads, wherein the memory allocation requests of the threads include requested memory amounts of the threads;

[0146] a warp requesting device configured to calculate a sum of memory amounts requested by threads in a corresponding warp among a plurality of warps as the requested memory amount of the corresponding warp, and select one thread in the corresponding warp to request allocation of memory space of the corresponding warp, the amount of the memory space of the corresponding warp being the requested memory amount of the corresponding warp;

[0147] A memory allocation device configured to allocate memory space to the corresponding thread warp;

[0148] The warp memory partitioning device is configured to partition the allocated memory space of the corresponding warp to each thread according to the memory amount requested by each thread in the corresponding warp.

[0149] Aspect 8. The system according to Aspect 7, wherein the warp requesting device is further configured to:

[0150] According to the memory amount requested by each thread in the corresponding thread warp, a prefix sum of the memory amount requested by each thread in the corresponding thread warp is calculated as the memory address offset of each thread;

[0151] The thread warp memory partitioning device is configured as follows:

[0152] According to the memory address offset of each thread, in the amount of memory space allocated to the corresponding warp, a memory address offset by the corresponding memory address offset is used as a starting address, and the memory amount requested by the corresponding thread is divided to the corresponding thread.

[0153] Aspect 9. The system according to Aspect 8, wherein the warp requesting device is further configured to calculate, based on the memory amounts requested by each thread in the corresponding warp, a prefix sum of the memory amounts requested by each thread in the corresponding warp as the memory address offset of each thread by performing the following steps:

[0154] Setting the prefix sum of the first thread in the corresponding warp to 0;

[0155] A prefix sum of a subsequent thread in the corresponding warp is calculated as the prefix sum of a previous thread of the subsequent thread plus the amount of memory requested by the subsequent thread.

[0156] Aspect 10. The system according to Aspect 7, wherein the size of each warp is configured to be a predetermined number of threads, and the warp grouping means is configured to:

[0157] The plurality of threads are allocated to the plurality of warps and additional threads not required for execution are added or not added so that the number of threads included in each warp is equal to the predetermined amount.

[0158] Aspect 11. The system according to Aspect 7, wherein the warp requesting device is further configured to:

[0159] If the memory amount requested by the thread is less than a predetermined minimum amount of memory allocation, the memory amount requested by the thread is increased to the predetermined minimum amount of memory allocation.

[0160] Aspect 12. The system according to aspect 7, wherein the plurality of threads are a plurality of threads for executing a single instruction in a single instruction multiple thread (SIMT) execution system.

[0161] Aspect 13. A memory allocation device comprising:

[0162] processor;

[0163] A memory stores computer executable instructions, which are executed by a processor to perform the method described in any one of aspects 1-6.

[0164] Aspect 14. A computer-readable medium storing computer-readable instructions, wherein when the computer-readable instructions are executed by a processor, the method according to any one of aspects 1 to 6 is performed.

[0165] Of course, the above-mentioned specific embodiments are merely examples and not limitations, and those skilled in the art can, according to the concept of the present application, merge and combine some steps and devices from the various embodiments described separately above to achieve the effects of the present application. Such merged and combined embodiments are also included in the present application, and such merges and combinations are not described one by one here.

[0166] Note that the advantages, strengths, and effects mentioned in this disclosure are merely examples and not limitations, and should not be construed as necessarily possessed by each embodiment of this application. Furthermore, the specific details disclosed above are provided for illustrative purposes and to facilitate understanding, not limitations, and the details do not necessarily limit this application to the use of the specific details.

[0167] The block diagrams of the devices, devices, equipment, and systems involved in this disclosure are merely illustrative examples and are not intended to require or imply that they must be connected, arranged, or configured in the manner shown in the block diagrams. As will be appreciated by those skilled in the art, these devices, devices, equipment, and systems can be connected, arranged, or configured in any manner. Words such as "include," "comprise," "have," and the like are open-ended words, meaning "including but not limited to," and can be used interchangeably therewith. The words "or" and "and" used herein refer to the words "and / or" and can be used interchangeably therewith, unless the context clearly indicates otherwise. The word "such as" used herein refers to the phrase "such as but not limited to," and can be used interchangeably therewith.

[0168] The schematic flow charts of steps in this disclosure and the above method descriptions are intended only as illustrative examples and are not intended to require or imply that the steps of the various embodiments must be performed in the order given. As will be appreciated by those skilled in the art, the order of the steps in the above embodiments can be performed in any order. Words such as "thereafter," "then," "next," and the like are not intended to limit the order of the steps; these words are only used to guide the reader through the description of these methods. In addition, any reference to an element in the singular, such as using the articles "a," "an," or "the," is not to be construed as limiting the element to the singular.

[0169] In addition, the steps and devices in the various embodiments of this document are not limited to being implemented in a certain embodiment. In fact, according to the concept of this application, relevant partial steps and partial devices in the various embodiments of this document can be combined to conceive new embodiments, and these new embodiments are also included in the scope of this application.

[0170] Each operation of the method described above may be performed by any suitable means capable of performing the corresponding functions, which may include various hardware and / or software components and / or modules, including but not limited to hardware circuits, application specific integrated circuits (ASICs), or processors.

[0171] The various illustrated logic blocks, modules, and circuits may be implemented or described using a general-purpose processor, a digital signal processor (DSP), an ASIC, a field programmable gate array signal (FPGA) or other programmable logic device (PLD), discrete gate or transistor logic, discrete hardware components, or any combination thereof designed to perform the functions described herein. A general-purpose processor may be a microprocessor, but as an alternative, the processor may be any commercially available processor, controller, microcontroller, or state machine. A processor may also be implemented as a combination of computing devices, such as a combination of a DSP and a microprocessor, multiple microprocessors, a microprocessor in conjunction with a DSP core, or any other such configuration.

[0172] The steps of the methods or algorithms described in conjunction with the present disclosure may be embedded directly in hardware, in a software module executed by a processor, or in a combination of the two. A software module may reside in any form of tangible storage medium. Some examples of storage media that may be used include random access memory (RAM), read-only memory (ROM), flash memory, EPROM memory, EEPROM memory, registers, hard disks, removable disks, CD-ROMs, and the like. A storage medium may be coupled to a processor so that the processor can read information from and write information to the storage medium. In an alternative embodiment, the storage medium may be integral to the processor. A software module may be a single instruction or many instructions and may be distributed across several different code segments, between different programs, and across multiple storage media.

[0173] The methods disclosed herein include actions for implementing the described methods. Methods and / or actions may be interchangeable with one another without departing from the scope of the claims. In other words, unless a specific order of actions is specified, the order and / or use of specific actions may be modified without departing from the scope of the claims.

[0174] The above functions can be implemented in hardware, software, firmware, or any combination thereof. If implemented in software, the functions can be stored as instructions on a tangible computer-readable medium. The storage medium can be any available tangible medium that can be accessed by a computer. By way of example and not limitation, such computer-readable media can include RAM, ROM, EEPROM, CD-ROM or other optical disk storage, magnetic disk storage or other magnetic storage devices, or any other tangible medium that can be used to carry or store desired program code in the form of instructions or data structures and can be accessed by a computer. As used herein, disk and disc include compact disc (CD), laser disc, optical disc, digital versatile disc (DVD), floppy disc, and Blu-ray disc, where disks typically reproduce data magnetically, while discs reproduce data optically using lasers.

[0175] Thus, a computer program product can perform the operations presented herein. For example, such a computer program product can be a computer-readable tangible medium having instructions tangibly stored (and / or encoded) thereon, the instructions being executable by a processor to perform the operations described herein. The computer program product can include packaging materials.

[0176] Software or instructions may also be transmitted over a transmission medium. For example, the software may be transmitted from a website, server, or other remote source using a transmission medium such as coaxial cable, fiber optic cable, twisted pair, digital subscriber line (DSL), or wireless technologies such as infrared, radio, or microwave.

[0177] In addition, the modules and / or other appropriate means for performing the methods and techniques described herein can be downloaded and / or otherwise obtained by the user terminal and / or base station, as appropriate. For example, such a device can be coupled to a server to facilitate the transmission of the means for performing the methods described herein. Alternatively, the various methods described herein can be provided via a storage component (e.g., RAM, ROM, a physical storage medium such as a CD or floppy disk) so that the user terminal and / or base station can obtain the various methods when coupled to the device or provided with the storage component. In addition, any other appropriate technology for providing the methods and techniques described herein to a device can be utilized.

[0178] Other examples and implementations are within the scope and spirit of the present disclosure and the appended claims. For example, due to the nature of software, the functions described above can be implemented using software executed by a processor, hardware, firmware, hardwiring, or any combination of these. Features that implement the functions can also be physically located in various locations, including being distributed so that parts of the functions are implemented at different physical locations. Moreover, as used herein, including in the claims, “or” used in a list of items beginning with “at least one of” indicates a separate list, so that, for example, a list of “at least one of A, B, or C” means A or B or C, or AB or AC or BC, or ABC (i.e., A and B and C). Furthermore, the word “exemplary” does not mean that the example described is preferred or better than other examples.

[0179] Various changes, substitutions, and modifications of the technology described herein may be made without departing from the teachings of the appended claims. Moreover, the scope of the claims of this disclosure is not limited to the specific aspects of the processes, machines, manufactures, compositions of things, means, methods, and actions described above. Currently existing or later developed processes, machines, manufactures, compositions of things, means, methods, or actions that perform substantially the same function or achieve substantially the same results as the corresponding aspects described herein may be utilized. Accordingly, the appended claims include within their scope such processes, machines, manufactures, compositions of things, means, methods, or actions.

[0180] The above description of the disclosed aspects is provided to enable any person skilled in the art to make or use the present application. Various modifications to these aspects will be readily apparent to those skilled in the art, and the general principles defined herein may be applied to other aspects without departing from the scope of the present application. Therefore, the present application is not intended to be limited to the aspects shown herein, but rather to be accorded the widest scope consistent with the principles and novel features disclosed herein.

[0181] The above description has been provided for the purpose of illustration and description. Furthermore, this description is not intended to limit the embodiments of the present application to the forms disclosed herein. Although a number of example aspects and embodiments have been discussed above, those skilled in the art will recognize certain variations, modifications, alterations, additions, and sub-combinations thereof.

Claims

1. A memory allocation method in a graphics processing unit (GPU), comprising the following steps: In response to respective memory allocation requests of a plurality of threads, the plurality of threads are grouped into a plurality of warps, wherein: The warp is the minimum unit of thread scheduling in a processing unit, the memory allocation request of the thread includes a requested memory amount of the thread, and the multiple threads are multiple threads for executing a single instruction in a single instruction multiple thread (SIMT) execution system; Calculating a sum of memory amounts requested by threads in a corresponding warp among the plurality of warps as the memory amount requested by the corresponding warp; Selecting one thread in the corresponding warp to request allocation of memory space of the corresponding warp by competing for a lock with threads in other warps, where the amount of the memory space of the corresponding warp is the requested memory amount of the corresponding warp; Allocate memory space for the corresponding thread warp in one atomic operation; During an operation in which other warps among the plurality of warps except the corresponding warp compete for the lock and allocate memory space, setting a prefix sum of a first thread in the corresponding warp to 0 according to a memory amount requested by each thread in the corresponding warp; A prefix sum of subsequent threads in the corresponding thread warp is calculated as the prefix sum of a thread preceding the subsequent thread plus the memory amount requested by the subsequent thread. The calculated prefix sum of the requests of each thread is used as the memory address offset of each thread. According to the memory address offset of each thread, a memory address that is offset by the corresponding memory address offset is used as a starting address in the amount of memory space allocated to the corresponding thread warp, and the memory amount requested by the corresponding thread is divided among the corresponding threads.

2. The method according to claim 1 , wherein the size of each warp is configured to be a predetermined number of threads, and the step of grouping the plurality of threads into a plurality of warps in response to respective memory allocation requests of the plurality of threads comprises: The plurality of threads are allocated to the plurality of warps and additional threads not required for execution are added or not added so that the number of threads included in each warp is equal to the predetermined amount.

3. The method according to claim 1, wherein The method further comprises: If the memory amount requested by the thread is less than a predetermined minimum amount of memory allocation, the memory amount requested by the thread is increased to the predetermined minimum amount of memory allocation.

4. A memory allocation system in a graphics processing unit (GPU), comprising the following apparatus: The thread warp grouping device is configured to group the plurality of threads into a plurality of thread warps in response to respective memory allocation requests of the plurality of threads, wherein The warp is the minimum unit of thread scheduling in a processing unit, the memory allocation request of the thread includes a requested memory amount of the thread, and the multiple threads are multiple threads for executing a single instruction in a single instruction multiple thread (SIMT) execution system; a warp requesting device configured to calculate a sum of memory amounts requested by threads in a corresponding warp among a plurality of warps as the requested memory amount of the corresponding warp, and select one thread in the corresponding warp to request allocation of memory space of the corresponding warp by competing for a lock with threads in other warps, the amount of the memory space of the corresponding warp being the requested memory amount of the corresponding warp; A memory allocation device configured to allocate memory space for the corresponding thread warp in one atomic operation; a warp memory partitioning device configured to set a prefix sum of a first thread in the corresponding warp to 0 according to a memory amount requested by each thread in the corresponding warp during an operation in which warps other than the corresponding warp among the plurality of warps compete for the lock and allocate memory space; A prefix sum of subsequent threads in the corresponding thread warp is calculated as the prefix sum of a thread preceding the subsequent thread plus the memory amount requested by the subsequent thread. The calculated prefix sum of the requests of each thread is used as the memory address offset of each thread. According to the memory address offset of each thread, a memory address that is offset by the corresponding memory address offset is used as a starting address in the amount of memory space allocated to the corresponding thread warp, and the memory amount requested by the corresponding thread is divided among the corresponding threads.

5. A memory allocation device comprising: processor; The memory stores computer executable instructions, which are executed by the processor to perform the method according to any one of claims 1 to 3.

6. A computer-readable medium storing computer-readable instructions, wherein when the computer-readable instructions are executed by a processor, the method according to any one of claims 1 to 3 is performed.

Images