Inter-thread data sharing method, electronic equipment, storage medium and program product

Abstract

The invention relates to the technical field of high-performance computing, and provides an inter-thread data sharing method, electronic equipment, a storage medium and a program product.The method comprises the steps that a plurality of threads in a thread block are divided into at least one cooperative thread group, and the threads in each cooperative thread group are from at least two different thread bundle groups; when a write operation of a first thread in any cooperative thread group is received, writing target data into a shared register resource associated with any cooperative thread group; and when a read operation of a second thread in any cooperative thread group for the target data is received, executing a synchronous operation, and reading the target data from the shared register resource after the execution is completed. According to the method, the cooperative thread group spanning different thread beam groups is constructed, and the shared register resources are utilized for data exchange, so that efficient data sharing is directly carried out among the threads spanning the thread beam groups through the register, and the method has the advantages of low delay, high bandwidth and capability of remarkably improving parallel computing performance.

Description

Technical Field

[0001] This invention relates to the field of high-performance computing technology, and in particular to a method for inter-thread data sharing, an electronic device, a storage medium, and a program product. Background Technology

[0002] In modern parallel computing devices such as general-purpose computing on graphics processing units (GPGPUs), multiple threads often need to work together to complete complex computing tasks, which inevitably involves data exchange and sharing between threads.

[0003] Currently, there are two main ways to share data between threads. One way is to use shared memory, which is visible to all threads within a computing unit (such as a thread block), but its access latency and synchronization overhead are relatively high, affecting the overall performance of the program. The other way is to utilize specific hardware instructions, such as shuffle instructions. This method can directly exchange data through registers, which is fast and has low latency, but its application is limited to threads within a single warp. Therefore, when the threads that need to exchange data belong to different warps or even different warp groups, the efficient register-level data exchange method cannot be used, and the only option is to revert to using shared memory, which limits the performance of parallel computing. Summary of the Invention

[0004] This invention provides a method, electronic device, storage medium, and program product for inter-thread data sharing, which solves the problem of low performance caused by the reliance on shared memory for inter-thread data sharing across thread bundles or groups of thread bundles in related technologies.

[0005] This invention provides a method for inter-thread data sharing, comprising: Based on the thread index of each thread within the thread block, multiple threads within the thread block are divided into at least one cooperating thread group, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperating thread group come from at least two different thread bundles. When a write operation is received from the first thread in any cooperating thread group, the target data corresponding to the write operation is written into the shared register resource associated with the cooperating thread group. Upon receiving a read operation for the target data from the second thread in any of the cooperative thread groups, a synchronization operation is performed for any of the cooperative thread groups, and after the synchronization operation is completed, the target data is read from the shared register resource.

[0006] According to a method for inter-thread data sharing provided by the present invention, the step of dividing multiple threads within a thread block into at least one cooperating thread group based on the thread index of each thread within the thread block includes: Based on the thread index of each thread within the thread block and the preset number of threads, a modulo operation is performed to obtain the identifier of the cooperating thread group to which each thread belongs, wherein the preset number of threads is the number of threads included in each thread bundle group; Threads with the same collaborative thread group identifier are grouped into the same collaborative thread group.

[0007] According to a method for inter-thread data sharing provided by the present invention, the first thread and the second thread belong to different thread bundles.

[0008] According to a method for inter-thread data sharing provided by the present invention, before the first thread performs a write operation, the method further includes: In response to variable declarations of a preset memory type, a corresponding register resource is allocated to each of the cooperative thread groups as a shared register resource among the threads within each cooperative thread group. At least two threads within each cooperative thread group share data by accessing the same storage location in the shared register resource.

[0009] According to a method for inter-thread data sharing provided by the present invention, variables declared with the preset memory type are used as an array, and threads in each cooperating thread group access the shared register resources of the cooperating thread group through array indexing.

[0010] According to the inter-thread data sharing method provided by the present invention, different threads within the same cooperating thread group access the same physical storage unit using the same array index; Threads within different cooperating thread groups access different, isolated physical storage units using the same array index.

[0011] According to a method for inter-thread data sharing provided by the present invention, the synchronization operation is triggered by calling a preset synchronization interface function for the cooperating thread group.

[0012] The present invention also provides an inter-thread data sharing device, comprising: A thread grouping unit is used to divide multiple threads within a thread block into at least one cooperating thread group based on the thread index of each thread within the thread block, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperating thread group come from at least two different thread bundles. The data writing unit is used to write the target data corresponding to the write operation into the shared register resource associated with the first thread in any cooperative thread group when a write operation is received. The data reading unit is configured to perform a synchronization operation for any of the cooperative thread groups when it receives a read operation for the target data from the second thread in any of the cooperative thread groups, and read the target data from the shared register resource after the synchronization operation is completed.

[0013] The present invention also provides an electronic device, including a memory, a processor, and a computer program stored in the memory and running on the processor, wherein the processor executes the computer program to implement the inter-thread data sharing method as described above.

[0014] The present invention also provides a non-transitory computer-readable storage medium having a computer program stored thereon, which, when executed by a processor, implements the inter-thread data sharing method as described above.

[0015] The present invention also provides a computer program product, including a computer program that, when executed by a processor, implements the inter-thread data sharing method as described above.

[0016] The inter-thread data sharing method, electronic device, storage medium, and program product provided by this invention construct a cooperative thread group that spans the boundaries of traditional thread bundle groups through a novel thread grouping approach. It is no longer limited by traditional thread bundle or thread bundle group divisions, but rather dynamically organizes threads from at least two different thread bundle groups into the same cooperative thread group based on thread indexes. Since a thread bundle group is composed of multiple thread bundles, threads in the cooperative thread group spanning different thread bundle groups necessarily mean that these threads also span thread bundles. When these threads spanning thread bundle groups or thread bundles need to exchange data, this invention provides a high-speed register-based channel. The data to be exchanged is written to the shared register resource associated with the cooperative thread group, rather than traditional shared memory. This avoids the high latency and high power consumption problems caused by using shared memory, thereby improving the efficiency of data exchange. Furthermore, this invention employs a more refined synchronization mechanism. Synchronization operations are performed on the cooperative thread group, rather than globally synchronizing the entire thread block. This lightweight synchronization method reduces unnecessary thread waiting, further reducing synchronization overhead and making the entire data sharing process more efficient. Attached Figure Description

[0017] To more clearly illustrate the technical solutions in this invention or related technologies, the accompanying drawings used in the description of the embodiments or related technologies will be briefly introduced below. Obviously, the accompanying drawings described below are some embodiments of this invention. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0018] Figure 1 This is a flowchart illustrating the inter-thread data sharing method provided by the present invention; Figure 2 This is a schematic diagram of inter-thread data sharing across thread bundles provided by the present invention; Figure 3 This is a schematic diagram of the structure of the inter-thread data sharing device provided by the present invention; Figure 4 This is a schematic diagram of the structure of the electronic device provided by the present invention. Detailed Implementation

[0019] To make the objectives, technical solutions, and advantages of this invention clearer, the technical solutions of this invention will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some, not all, of the embodiments of this invention. All other embodiments obtained by those skilled in the art based on the embodiments of this invention without creative effort are within the scope of protection of this invention.

[0020] In modern parallel computing architectures (such as GPGPUs), a Single Instruction Multiple Threads (SIMT) execution model is typically used to improve computational efficiency. In this model, multiple threads are organized into a warp (usually containing 32 threads), and threads within the same warp can execute the same instructions in parallel. Multiple warps are further organized into a thread block, which is then scheduled for execution on a streaming multiprocessor.

[0021] To better organize and manage complex collaborative computing tasks, multiple thread bundles within the same thread block can be further divided into one or more warp groups. Each warp group typically contains multiple thread bundles and is assigned to execute specific subtasks within a larger computational task. For example, in complex operations such as matrix multiplication and accumulation, the thread bundles within a thread block can be divided into warp groups responsible for the core matrix computation and warp groups responsible for processing the computation results (such as data type conversion and writing back to memory). Each warp group can consist of four consecutive warps. This division allows computational units with different functions to collaborate efficiently in a pipelined manner.

[0022] To accomplish these collaborative computing tasks, efficient data exchange and synchronization are frequently required between different threads, especially those belonging to different thread bundles or even different thread bundle groups. Currently, achieving inter-thread data sharing mainly relies on the following methods: The first approach is to use shared memory. Shared memory is a read-write area on-chip that is visible to all threads within the same thread block (including all thread bundles and thread bundle groups). Threads can exchange data by reading and writing addresses in shared memory. However, although accessing shared memory is much faster than accessing global memory, it is still a form of memory access. Compared to directly accessing registers, it has higher access latency and requires explicit synchronization instructions to ensure the correct order of data reads and writes. This incurs some performance overhead, especially in scenarios where frequent data exchange is required.

[0023] The second approach utilizes instructions within a thread bundle, such as the shuffle instruction. This method allows threads within the same thread bundle to exchange data directly through registers, without needing shared memory. Because it directly manipulates registers, it is very fast and has low latency.

[0024] However, register-level data exchange based on instructions like shuffle is limited to threads within a single thread bundle. This approach becomes unsuitable when threads needing to exchange data belong to different thread bundles, particularly different thread bundle groups. In such scenarios, such as the aforementioned thread bundle group processing matrix calculations needing to pass calculation results to the thread bundle group processing results, this can only be accomplished through shared memory, which limits the overall performance of parallel computing.

[0025] Furthermore, even within a thread bundle, shuffle instructions have shortcomings. They often have high instruction overhead and long execution latency, and they limit the number of threads that can participate in data exchange (e.g., less than 32), which to some extent affects programming flexibility and program execution efficiency.

[0026] To address this issue, the present invention provides a method for inter-thread data sharing. First, a novel thread grouping approach is used to construct cooperative thread groups that span the boundaries of traditional thread bundles. Then, each cooperative thread group is provided with a register-based, mutually isolated shared data channel. Finally, a group-level lightweight synchronization mechanism is defined to ensure the correctness of data read and write operations. This allows threads belonging to different thread bundles or even different thread bundle groups that previously needed to communicate via slower shared memory to now directly share data via high-speed registers. This significantly reduces the latency of data exchange between threads across thread bundles or even thread bundle groups, improves the execution efficiency of parallel computing, and thus overcomes the aforementioned shortcomings.

[0027] It should be noted that the inter-thread data sharing method provided by this invention can be executed by a computing device with parallel processing capabilities. This computing device can be hardware with parallel computing capabilities such as a GPGPU, a graphics processing unit (GPU), a graphics processing unit integrated into a central processing unit (CPU), an application-specific integrated circuit (ASIC), or a field-programmable gate array (FPGA).

[0028] Taking the GPGPU architecture as an example, the execution entity can specifically be one or more streaming multiprocessors or computing units within the GPGPU. When a user-written parallel computing program (usually called a kernel) is scheduled to execute on this computing device, the instructions within the program are executed by the computing cores within these streaming multiprocessors or computing units. The program logic is decomposed into a large number of threads, which are organized and distributed to the streaming multiprocessors or computing units in the form of thread blocks or thread bundles. The method of this invention mainly focuses on the execution of multiple threads within a thread block within these streaming multiprocessors or computing units. By regrouping the threads and providing new data exchange and synchronization mechanisms, efficient inter-thread data sharing is achieved.

[0029] Figure 1 This is a flowchart illustrating the inter-thread data sharing method provided by the present invention, as shown below. Figure 1 As shown, the method includes: Step S10: Based on the thread index of each thread within the thread block, divide the multiple threads within the thread block into at least one cooperating thread group, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperating thread group come from at least two different thread bundles.

[0030] Specifically, the purpose of this step is to create a new, cross-thread bundle logical grouping of multiple threads within a thread block to form a cooperative thread group. Specifically, it is first necessary to determine the unique identifier of each thread within the thread block, i.e., the thread index. This index is a logical, linear number. For example, in a three-dimensionally organized thread block, the thread index of each thread can be calculated using the following formula: thread.index_rank=threadIdx.z blockDim.y blockDim.x+threadIdx.y blockDim.x+threadIdx.x Here, thread.index_rank represents the thread index of a certain thread within its own thread block, blockDim.x and blockDim.y represent the number of threads in the x and y dimensions of the thread block, respectively, and threadIdx.x, threadIdx.y, and threadIdx.z represent the index of the thread in the x, y, and z dimensions of its own thread block, respectively, which are the three-dimensional coordinates of the thread within its own thread block.

[0031] After obtaining the unique thread index for each thread, the threads can be assigned to different cooperative thread groups according to a preset partitioning rule. A key feature of this partitioning rule is that it can group threads belonging to different thread bundles into the same cooperative thread group. This breaks through the limitation of traditional technologies where cooperation is usually confined to a single thread bundle or thread bundle group.

[0032] For example, suppose a thread bundle consists of four consecutive thread bundles, each containing 32 threads. Threads numbered 0-127 within a thread block constitute thread bundle group one, threads numbered 128-255 constitute thread bundle group two, and so on. The grouping method in this embodiment can group the thread with index i and the thread with index 128+i (where 0 ≤ i < 128) into the same cooperative thread group. Clearly, thread i belongs to thread bundle group one, while thread 128+i belongs to thread bundle group two; they come from different thread bundle groups but logically belong to the same cooperative unit. This cross-thread bundle group division lays the foundation for more flexible direct data interaction between threads.

[0033] Step S20: When a write operation is received from the first thread in any collaborative thread group, the target data corresponding to the write operation is written into the shared register resource associated with the collaborative thread group.

[0034] Specifically, after dividing threads into cooperating thread groups, this step aims to provide an efficient data sharing channel for threads within the group by utilizing shared register resources.

[0035] In the specific implementation, for each cooperating thread group, when the first thread in the cooperating thread group (i.e., the thread performing the write operation, which can be any thread in the cooperating thread group) needs to share the target data it has calculated or obtained from other storage levels with other threads in its cooperating thread group, it will perform a write operation. This write operation is not written to regular global memory or shared memory, but to a shared register resource that is uniquely associated with the cooperating thread group.

[0036] The shared register resource here is a special type of storage resource. At the hardware level, it can be part of the physically existing register file within the streaming multiprocessor or computing unit of a GPGPU. It should be understood that in some GPGPU architectures, there may be some underutilized register resources. This embodiment of the invention can reallocate these resources as shared storage space between thread groups. Moreover, this resource is shared and accessible to all threads within the same cooperating thread group, but isolated from threads in different cooperating thread groups. In other words, data written by a thread in cooperating thread group A cannot be accessed by threads in cooperating thread group B, ensuring data isolation. This association between the shared register resource and the cooperating thread group ensures that the scope of data sharing is precisely limited to within the cooperating thread group, avoiding data conflicts and interference.

[0037] Step S30: When a read operation for the target data is received from the second thread in any of the cooperative thread groups, a synchronization operation is performed for any of the cooperative thread groups, and after the synchronization operation is completed, the target data is read from the shared register resource.

[0038] It should be noted that in parallel execution mode, due to the uncertainty of hardware scheduling, the execution progress of different threads cannot be guaranteed. A thread performing a read operation (referred to here as the second thread) may start before the first thread's write operation is completed. If the second thread attempts to read before the first thread has finished writing data, it will obtain invalid or stale data.

[0039] To address this issue, a synchronization operation is introduced in this step. When a second thread (which can be any thread in the cooperating thread group other than the first thread) needs to read target data written by other members of the group (such as the first thread), a synchronization operation for its cooperating thread group must be triggered first. Here, the synchronization operation is mainly used to suspend the execution of the thread requesting synchronization (such as the second thread) until all threads in the cooperating thread group have reached this synchronization point. This synchronization operation ensures that all write operations initiated by any thread in the group before the synchronization point (especially write operations to shared register resources) have been completed, and their results are visible to all threads in the group.

[0040] Once the synchronization operation is complete, data consistency is guaranteed. At this point, the second thread can safely perform a read operation, reading the target data written by the first thread from the same shared register resource associated with the cooperating thread group. Because the first and second threads belong to the same cooperating thread group, they access the same physical register space, thus ensuring correct data transfer.

[0041] It is understandable that the above synchronization operations are performed at a smaller granularity, namely the cooperating thread group, rather than the entire thread block. Compared to traditional thread block-level synchronization, this group-level synchronization has a smaller impact scope and involves fewer threads, thus achieving lower synchronization overhead and latency.

[0042] The method provided in this invention constructs a cooperative thread group that spans the boundaries of traditional thread bundle groups through a novel thread grouping approach. It is no longer limited by traditional thread bundle or thread bundle group divisions, but rather dynamically organizes threads from at least two different thread bundle groups into the same cooperative thread group based on thread indices. Since a thread bundle group is composed of multiple thread bundles, threads in the cooperative thread group spanning different thread bundle groups necessarily mean that these threads are also cross-thread bundles. When these threads spanning thread bundle groups or thread bundles need to exchange data, this invention provides a high-speed register-based channel. The data to be exchanged is written to the shared register resource associated with the cooperative thread group, rather than traditional shared memory. This avoids the high latency and high power consumption problems caused by using shared memory, thereby improving the efficiency of data exchange. Furthermore, this invention employs a more refined synchronization mechanism. Synchronization operations are performed on the cooperative thread group itself, rather than globally synchronizing the entire thread block. This lightweight synchronization method reduces unnecessary thread waiting, further reducing synchronization overhead and making the entire data sharing process more efficient.

[0043] Based on the above embodiments, step S10 specifically includes: Step S11: Based on the thread index of each thread within the thread block and the preset number of threads, perform a modulo operation to obtain the cooperative thread group identifier to which each thread belongs. The preset number of threads is the number of threads included in each thread bundle group. Step S12 assigns threads with the same collaborative thread group identifier to the same collaborative thread group.

[0044] Specifically, when grouping multiple threads within a thread block, a unique, linear thread index is first assigned to each thread within the block. Simultaneously, the system defines a preset thread count, which can be specifically defined as the number of threads contained in each thread bundle. For example, in a specific hardware architecture, if a thread bundle contains 32 threads, and a thread bundle group consists of 4 consecutive thread bundles, then the preset thread count is 32 × 4 = 128.

[0045] After obtaining the thread index and the preset number of threads, the system performs a modulo operation on each thread. Specifically, for a thread with a thread index `thread.index_rank`, its cooperating thread group identifier can be calculated using the following formula: co-operative_threads_group_ID=thread.index_rank%N Wherein, co-operative_threads_group_ID represents the cooperative thread group identifier, which is an integer used to uniquely identify a cooperative thread group; N represents the preset number of threads.

[0046] After the modulo operation described above, all threads that calculate the same cooperative thread group identifier will be logically grouped into the same cooperative thread group. This means that they will share the same associated shared register resource and participate in the same group-level synchronization operations.

[0047] For example, Figure 2 This is a schematic diagram of inter-thread data sharing across thread bundles provided by the present invention, as shown below. Figure 2The diagram illustrates a scenario where threads within a thread block are scheduled to execute on different execution units within a processing unit. Here, the processing unit can be a streaming multiprocessor or computing unit in a GPGPU architecture, and the execution unit can specifically be a computing core. Assume the thread block contains 256 threads, with thread indices ranging from 0 to 255. These threads are divided into eight thread bundles, warp0 to warp7, each containing 32 threads. Four consecutive thread bundles form a thread bundle group; warp0 to warp3 form one group, and warp4 to warp7 form another. The default number of threads (i.e., the size of a thread bundle group) is 128.

[0048] Threads within the same thread bundle can exchange and share data through registers, while threads belonging to different thread bundles or different thread bundle groups need to exchange data through shared memory. To improve the efficiency of data exchange between threads across thread bundle groups, this invention proposes a new grouping method: performing a modulo operation based on the thread index of each thread and a preset number of threads, thereby grouping threads from different thread bundle groups into the same cooperative thread group, enabling threads across thread bundle groups to quickly exchange data through shared register resources.

[0049] For example, in warp0, the thread indices range from 0 to 31, and in warp4, the thread indices range from 128 to 159. For the thread with index 0 in warp0, its cooperating thread group identifier is 0%128=0; for the thread with index 128 in warp4, its cooperating thread group identifier is 128%128=0. Therefore, the thread with index 0 in warp0 and the thread with index 128 in warp4 will be grouped into the same cooperating thread group, and these two threads can exchange data through the shared registers associated with this cooperating thread group. Similarly, the thread with index 1 in warp0 and the thread with index 129 in warp4 will also be grouped into the same cooperating thread group and exchange data through the corresponding register resources. Similarly, the threads corresponding to warp0 and warp4 can exchange data through their respective register resources; the threads corresponding to warp1 and warp5 can exchange data through their respective register resources; the threads corresponding to warp2 and warp6 can also exchange data through their respective register resources; and the threads corresponding to warp3 and warp7 can also exchange data through their respective register resources.

[0050] Since warp0 and warp4 belong to different thread bundles and different thread bundle groups, and warp1 and warp5 also belong to different thread bundles and different thread bundle groups, the grouping method based on modulo operation provided in this embodiment of the invention can simply and effectively divide threads from different thread bundle groups into the same cooperative thread group. This method is computationally simple and has extremely low overhead. It can establish a logical cooperative relationship across thread bundle groups in a deterministic and regular manner, so that the hardware or scheduler can easily identify the members of each cooperative thread group, providing a clear and consistent foundation for subsequent data sharing and synchronization operations.

[0051] Based on any of the above embodiments, the first thread and the second thread belong to different thread bundles.

[0052] Specifically, according to the grouping method in the above embodiments, the members of a collaborative thread group come from different thread bundle groups, and naturally from different thread bundles. Therefore, the first thread (write thread) and the second thread (read thread) belong to different thread bundles.

[0053] For example, such as Figure 2 As shown, assume the first thread is thread index 0 in warp0, and the second thread is thread index 128 in warp4, belonging to different thread bundles. When the first thread performs a computation task and needs to pass a result (i.e., the target data) to the second thread, the first thread performs a write operation for its cooperating thread group, writing the target data to the shared register resource associated with that cooperating thread group. Since the first and second threads are in different thread bundles, their execution progress may be completely out of sync. When the second thread needs to use the target data, it first performs a synchronization operation for the cooperating thread group. This synchronization operation ensures that the first thread's write operation has been completed and the data is visible to members of the group. After synchronization, the second thread can safely and correctly read the target data from the same shared register resource associated with the cooperating thread group.

[0054] The above process transforms the previously lengthy cross-thread bundle communication—which involved writing to shared memory, synchronizing thread blocks, and reading from shared memory—into a shorter path: writing to shared registers, synchronizing with the cooperative group, and reading from shared registers. Since register access is significantly faster than shared memory access, and the synchronization overhead at the cooperative thread group level is typically much smaller than that at the thread block level, this improves the efficiency of inter-thread data sharing across thread bundles.

[0055] Based on any of the above embodiments, before the first thread performs the write operation, the method further includes: In response to variable declarations of a preset memory type, a corresponding register resource is allocated to each of the cooperative thread groups as a shared register resource among the threads within each cooperative thread group. At least two threads within each cooperative thread group share data by accessing the same storage location in the shared register resource.

[0056] It should be noted that the embodiments of the present invention provide a specific description of the source and allocation method of the shared register resources associated with each cooperating thread group. The steps described in the embodiments of the present invention typically occur when the thread block is scheduled to a streaming multiprocessor or computing unit, and before any actual data read / write operation occurs.

[0057] Specifically, in actual parallel program development, developers need a way to inform the compiler and hardware that a certain variable needs to use the sharing mechanism proposed in this invention. To achieve this, embodiments of this invention introduce a preset memory type or type qualifier, which can be a new keyword, such as "__thread_shared_register__".

[0058] For example, when developers write their parallel computing kernel code, they can use the type qualifiers described above to declare a variable. When the compiler compiles this code, once it parses this predefined memory type declaration, it will recognize that the variable is neither a typical thread-private variable nor a traditional thread-block shared memory variable. The compiler will then generate specific instructions to trigger the subsequent resource allocation process.

[0059] Upon receiving this specific instruction, the hardware processing unit (such as a streaming multiprocessor or computing unit) or the runtime system performs resource allocation operations. Specifically, the system allocates corresponding register resources to each cooperating thread group. In the above embodiment, it has been described how to divide all threads within a thread block into M (M is an integer greater than or equal to 1) cooperating thread groups using modulo operations. Therefore, in this step, the system reserves M physical registers (or M blocks of register space, if the variable size exceeds one register) from the physical register file and binds these M physical registers to each of the M cooperating thread groups. This allocated register resource constitutes the shared register resource within each cooperating thread group.

[0060] It is understood that by allocating each cooperating thread group its own associated register resources, the physical basis for inter-thread data sharing is ensured. All threads within each cooperating thread group can share data by accessing the same storage location in the shared register resources associated with that cooperating thread group. In this embodiment of the invention, by providing a clear and convenient programming interface (i.e., a preset memory type), users can quickly utilize the efficient data sharing mechanism provided by the hardware.

[0061] Based on any of the above embodiments, the variables declared by the preset memory type are used as an array, and the threads in each of the cooperative thread groups access the shared register resources of the cooperative thread group through array indexing.

[0062] Specifically, in this embodiment of the invention, variables declared using a preset memory type can be arrays. This design enhances the flexibility and practicality of the invention because in many complex parallel computing algorithms, such as parallel reduction and histogram calculation, threads need to exchange a set of structured data, rather than a single piece of data. Designing shared variables as arrays satisfies this requirement.

[0063] For example, developers can declare an array tsr of size n using a preset memory type qualifier, such as "__thread_shared_register__ int tsr[n]". When this declaration is compiled and executed, the system will allocate a block of contiguous or logically contiguous register resources that can hold n integers for each cooperating thread group.

[0064] Accordingly, threads within a cooperative thread group can access the shared register resources of that group using array indexing. For example, a thread within the cooperative thread group can read or write a specific location in this shared register group using tsr[j] (where 0 ≤ j < n).

[0065] The method provided in this invention, by declaring variables as arrays and accessing them using array indices, enhances data processing capabilities and application breadth. It is no longer limited to the exchange of single data items but supports efficient, register-level collaborative operations on structured datasets among members within a thread group. This allows many complex parallel computing patterns (such as reduction) to be implemented at the extremely low-latency register level, enhancing algorithm execution performance and resource utilization efficiency. Furthermore, users can directly access shared register resources through array indices without needing to concern themselves with the complex underlying register index mapping, simplifying programming complexity.

[0066] Based on any of the above embodiments, different threads within the same collaborative thread group access the same physical storage unit using the same array index. Threads within different cooperating thread groups access different, isolated physical storage units using the same array index.

[0067] Specifically, different threads within the same cooperating thread group have access to the same shared register resources. Therefore, they access the same physical memory unit using the same array index, thus enabling them to access and read the same data. Specifically, when two or more threads in the same cooperating thread group attempt to access the same index location in a shared array, the hardware address translation mechanism directs all these access requests to a unique, identical physical memory unit. This ensures that when one thread writes data to that location, other threads in the group can read that new data from the exact same location, thereby achieving data sharing.

[0068] For threads within different cooperative thread groups, since they have access to different shared register resources, they will not access or write the same array even if they use the same array index, because they are accessing different physical storage units that are isolated from each other. This avoids data races in parallel computing. For example, suppose thread A belongs to cooperative thread group one and thread B belongs to cooperative thread group two. When thread A accesses tsr[0] and thread B also accesses tsr[0], although they use the exact same variable name and index at the code level, the hardware will map them to two completely different physical storage units. Thread A's access will be directed to the 0th shared register allocated to cooperative thread group one, while thread B's access will be directed to the 0th shared register allocated to cooperative thread group two. These two physical registers are independent of each other.

[0069] Therefore, thread A's write operation to tsr[0] is completely invisible to thread B, and vice versa. This design ensures that data operations between different cooperating thread groups do not interfere with each other, and each group has its own private shared space, thereby guaranteeing the correctness and stability of the program.

[0070] Based on any of the above embodiments, the synchronization operation is triggered by calling a preset synchronization interface function for the cooperative thread group.

[0071] Specifically, in parallel computing, ensuring the correct order of data dependencies is crucial. The synchronization operation in this embodiment of the invention is precisely to ensure that write operations are completed and visible to all members of the group before shared data is read.

[0072] Specifically, synchronization operations can be triggered by calling a pre-defined synchronization interface function for cooperative thread groups. This means that this embodiment of the invention provides developers with a clear Application Programming Interface (API) for inserting synchronization points into code. This API can be a built-in function or a library function, for example, it can be named "__sync_cooperative_threads()" or a similar form.

[0073] When a thread within the cooperative thread group (such as the second thread) reaches the line of code that calls the synchronization interface function, the execution flow of that thread is paused. Simultaneously, the processor's synchronization control unit begins checking the execution status of all other member threads within the cooperative thread group to which that thread belongs. Only when all threads within the cooperative thread group have reached this synchronization point (i.e., all have called the synchronization interface function), and the results of all memory write operations prior to this synchronization point (especially write operations to shared register resources) are globally visible, will the synchronization control unit release all threads in the cooperative thread group, allowing them to continue executing instructions after the synchronization point.

[0074] It is understandable that this synchronization interface function is designed for the granularity of cooperating thread groups. Traditional thread block-level synchronization functions suspend the execution of all threads within the entire thread block until all threads within the block reach the synchronization point. However, the synchronization operation in this embodiment only affects the cooperating thread group to which the thread that called it belongs; threads in other cooperating thread groups can continue execution unaffected. Because the synchronization scope is smaller and the number of threads participating in synchronization is fewer, cooperating thread group-level synchronization typically has lower latency and system overhead.

[0075] This invention provides a pre-defined synchronization interface function, offering users a simple, intuitive, and efficient means to control the execution flow within a collaborative thread group and ensure data consistency. This fine-grained synchronization mechanism, compared to traditional coarse-grained block-level synchronization, significantly reduces unnecessary thread waiting time, improves processor utilization, and increases the overall throughput of parallel programs.

[0076] The inter-thread data sharing device provided by the present invention is described below. The inter-thread data sharing device described below can be referred to in correspondence with the inter-thread data sharing method described above.

[0077] Based on any of the above embodiments Figure 3 This is a schematic diagram of the inter-thread data sharing device provided by the present invention, as shown below. Figure 3 As shown, the device includes: The thread grouping unit 310 is used to divide multiple threads in the thread block into at least one cooperative thread group based on the thread index of each thread in the thread block, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperative thread group come from at least two different thread bundles. The data writing unit 320 is used to write the target data corresponding to the write operation into the shared register resource associated with the cooperating thread group when a write operation is received from the first thread in any cooperating thread group. The data reading unit 330 is configured to perform a synchronization operation for any of the cooperative thread groups when it receives a read operation for the target data from the second thread in any of the cooperative thread groups, and read the target data from the shared register resource after the synchronization operation is completed.

[0078] The apparatus provided in this invention constructs a cooperative thread group that spans the boundaries of traditional thread bundle groups through a novel thread grouping method. It is no longer limited by the traditional division of thread bundles or thread bundle groups, but dynamically organizes threads from at least two different thread bundle groups into the same cooperative thread group based on thread indexes. Since a thread bundle group is composed of multiple thread bundles, threads in the cooperative thread group spanning different thread bundle groups necessarily mean that these threads also span thread bundles. When these threads spanning thread bundle groups or thread bundles need to exchange data, this invention provides a high-speed register-based channel. The data to be exchanged is written to the shared register resource associated with the cooperative thread group, rather than traditional shared memory. This avoids the high latency and high power consumption problems caused by using shared memory, thereby improving the efficiency of data exchange. Furthermore, this invention employs a more refined synchronization mechanism. Synchronization operations are performed on the cooperative thread group, rather than globally synchronizing the entire thread block. This lightweight synchronization method reduces unnecessary thread waiting, further reducing synchronization overhead and making the entire data sharing process more efficient.

[0079] Based on any of the above embodiments, the thread grouping unit 310 is specifically used for: Based on the thread index of each thread within the thread block and the preset number of threads, a modulo operation is performed to obtain the identifier of the cooperating thread group to which each thread belongs, wherein the preset number of threads is the number of threads included in each thread bundle group; Threads with the same collaborative thread group identifier are grouped into the same collaborative thread group.

[0080] Based on any of the above embodiments, the first thread and the second thread belong to different thread bundles.

[0081] Based on any of the above embodiments, the device further includes a resource configuration unit, the resource configuration unit being used for: In response to variable declarations of a preset memory type, a corresponding register resource is allocated to each of the cooperative thread groups as a shared register resource among the threads within each cooperative thread group. At least two threads within each cooperative thread group share data by accessing the same storage location in the shared register resource.

[0082] Based on any of the above embodiments, the variables declared by the preset memory type are used as an array, and the threads in each of the cooperative thread groups access the shared register resources of the cooperative thread group through array indexing.

[0083] Based on any of the above embodiments, different threads within the same collaborative thread group access the same physical storage unit using the same array index. Threads within different cooperating thread groups access different, isolated physical storage units using the same array index.

[0084] Based on any of the above embodiments, the synchronization operation is triggered by calling a preset synchronization interface function for the cooperative thread group.

[0085] Figure 4 An example is a schematic diagram of the physical structure of an electronic device, such as... Figure 4 As shown, the electronic device may include a processor 410, a communication interface 420, a memory 430, and a communication bus 440, wherein the processor 410, the communication interface 420, and the memory 430 communicate with each other through the communication bus 440. The processor 410 can call logical instructions in the memory 430 to execute an inter-thread data sharing method. This method includes: dividing multiple threads within a thread block into at least one cooperative thread group based on the thread index of each thread within the thread block, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperative thread group come from at least two different thread bundles; upon receiving a write operation from a first thread in any cooperative thread group, writing the target data corresponding to the write operation into a shared register resource associated with the cooperative thread group; upon receiving a read operation from a second thread in any cooperative thread group for the target data, performing a synchronization operation for the cooperative thread group, and reading the target data from the shared register resource after the synchronization operation is completed.

[0086] Furthermore, the logical instructions in the aforementioned memory 430 can be implemented as software functional units and, when sold or used as independent products, can be stored in a computer-readable storage medium. Based on this understanding, the technical solution of the present invention, or the part that contributes to related technologies, or a portion of the technical solution, can be embodied in the form of a software product. This computer software product is stored in a storage medium and includes several instructions to cause a computer device (which may be a personal computer, server, or network device, etc.) to execute all or part of the steps of the methods described in the various embodiments of the present invention. The aforementioned storage medium includes various media capable of storing program code, such as USB flash drives, portable hard drives, read-only memory (ROM), random access memory (RAM), magnetic disks, or optical disks.

[0087] On the other hand, the present invention also provides a computer program product, which includes a computer program that can be stored on a non-transitory computer-readable storage medium. When the computer program is executed by a processor, the computer is able to execute the inter-thread data sharing method provided by the above methods. The method includes: dividing multiple threads in a thread block into at least one cooperative thread group based on the thread index of each thread in the thread block, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperative thread group come from at least two different thread bundles; upon receiving a write operation from a first thread in any cooperative thread group, writing the target data corresponding to the write operation into a shared register resource associated with the cooperative thread group; upon receiving a read operation from a second thread in any cooperative thread group for the target data, performing a synchronization operation for the cooperative thread group, and reading the target data from the shared register resource after the synchronization operation is completed.

[0088] In another aspect, the present invention also provides a non-transitory computer-readable storage medium storing a computer program thereon, which, when executed by a processor, implements the inter-thread data sharing method provided by the above methods. The method includes: dividing multiple threads within a thread block into at least one cooperative thread group based on the thread index of each thread within the thread block, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperative thread group come from at least two different thread bundles; upon receiving a write operation from a first thread in any cooperative thread group, writing the target data corresponding to the write operation into a shared register resource associated with the cooperative thread group; upon receiving a read operation from a second thread in any cooperative thread group for the target data, performing a synchronization operation for the cooperative thread group, and after the synchronization operation is completed, reading the target data from the shared register resource.

[0089] The device embodiments described above are merely illustrative. The units described as separate components may or may not be physically separate. The components shown as units may or may not be physical units; that is, they may be located in one place or distributed across multiple network units. Some or all of the modules can be selected to achieve the purpose of this embodiment according to actual needs. Those skilled in the art can understand and implement this without any creative effort.

[0090] Through the above description of the embodiments, those skilled in the art can clearly understand that each embodiment can be implemented by means of software plus necessary general-purpose hardware platforms, and of course, it can also be implemented by hardware. Based on this understanding, the above technical solutions, in essence or the parts that contribute to the related technology, can be embodied in the form of software products. This computer software product can be stored in a computer-readable storage medium, such as ROM / RAM, magnetic disk, optical disk, etc., and includes several instructions to cause a computer device (which may be a personal computer, server, or network device, etc.) to execute the methods described in the various embodiments or some parts of the embodiments.

[0091] Finally, it should be noted that the above embodiments are only used to illustrate the technical solutions of the present invention, and not to limit them; although the present invention has been described in detail with reference to the foregoing embodiments, those skilled in the art should understand that modifications can still be made to the technical solutions described in the foregoing embodiments, or equivalent substitutions can be made to some of the technical features; and these modifications or substitutions do not cause the essence of the corresponding technical solutions to deviate from the spirit and scope of the technical solutions of the embodiments of the present invention.

Claims

1. A method for inter-thread data sharing, characterized in that, include: Based on the thread index of each thread within the thread block, multiple threads within the thread block are divided into at least one cooperating thread group, wherein the thread block is divided into multiple thread bundles, and the threads in each cooperating thread group come from at least two different thread bundles. When a write operation is received from the first thread in any cooperating thread group, the target data corresponding to the write operation is written into the shared register resource associated with the cooperating thread group. Upon receiving a read operation for the target data from the second thread in any of the cooperative thread groups, a synchronization operation is performed for any of the cooperative thread groups, and after the synchronization operation is completed, the target data is read from the shared register resource.

2. The inter-thread data sharing method according to claim 1, characterized in that, The step of dividing multiple threads within a thread block into at least one cooperating thread group based on the thread index of each thread within the thread block includes: Based on the thread index of each thread within the thread block and the preset number of threads, a modulo operation is performed to obtain the cooperative thread group identifier to which each thread belongs. The preset number of threads is the number of threads included in each thread bundle group. Threads with the same collaborative thread group identifier are grouped into the same collaborative thread group.

3. The inter-thread data sharing method according to claim 1, characterized in that, The first thread and the second thread belong to different thread bundles.

4. The inter-thread data sharing method according to claim 1, characterized in that, Before the first thread performs the write operation, the method further includes: In response to variable declarations of a preset memory type, a corresponding register resource is allocated to each of the cooperative thread groups as a shared register resource among the threads within each cooperative thread group. At least two threads within each cooperative thread group share data by accessing the same storage location in the shared register resource.

5. The inter-thread data sharing method according to claim 4, characterized in that, Using the variables declared with the preset memory type as an array, each thread in the cooperative thread group accesses the shared register resources of the cooperative thread group through array indexing.

6. The inter-thread data sharing method according to claim 5, characterized in that, Different threads within the same collaborative thread group access the same physical storage unit using the same array index; Threads within different cooperating thread groups access different, isolated physical storage units using the same array index.

7. The inter-thread data sharing method according to any one of claims 1 to 6, characterized in that, The synchronization operation is triggered by calling a preset synchronization interface function for the collaborative thread group.

8. An electronic device comprising a memory, a processor, and a computer program stored in the memory and running on the processor, characterized in that, When the processor executes the computer program, it implements the inter-thread data sharing method as described in any one of claims 1 to 7.

9. A non-transitory computer-readable storage medium having a computer program stored thereon, characterized in that, When the computer program is executed by the processor, it implements the inter-thread data sharing method as described in any one of claims 1 to 7.

10. A computer program product, comprising a computer program, characterized in that, When the computer program is executed by the processor, it implements the inter-thread data sharing method as described in any one of claims 1 to 7.

Images