Method for parallel execution of kernel functions with many-core parallel processor

Abstract

The method for executing a kernel function in parallel and the many-core parallel processor provided by the embodiments of the present disclosure can reduce the amount of data stored in the registers of the many-core parallel processor, and meanwhile, reduce the number of bits required when passing control information between the modules of the many-core parallel processor, and thus, the reduction of the number of transmission bits can reduce the pressure on the control bus and improve the hardware execution efficiency.

Description

Technical Field

[0001] This disclosure relates to the field of artificial intelligence chips, and more specifically to a method for parallel execution of kernel functions and a many-core parallel processor. Background Technology

[0002] The rapid development of artificial intelligence (AI) has placed higher demands on computing power, leading to the development of AI chips specifically designed to accelerate AI workloads. These chips typically employ a many-core parallel processor architecture, leveraging its massive parallel computing capabilities to efficiently handle complex AI algorithms, such as neural network computations in deep learning. Processors employing a many-core parallel processor architecture will be referred to below as many-core parallel processors (e.g., GPUs).

[0003] A kernel function is the basic unit of parallel computing on a many-core parallel processor. Starting a kernel function on a many-core parallel processor initiates the entire parallel computing task corresponding to that kernel function. All the threads within this parallel computing task constitute a logical execution space called a grid. The grid is typically further organized into several thread blocks, each containing several threads. To improve scheduling flexibility and hardware resource utilization, mainstream many-core parallel processors also divide the grid into multiple sub-grids, each consisting of several thread blocks. The grid, sub-grids, and thread blocks are all organized using a three-dimensional index, containing X, Y, and Z dimensions respectively.

[0004] Currently, static parameter storage and passing mechanisms are commonly used to load the execution space configuration parameters of kernel functions (e.g., parameters for the execution configuration of the kernel function's grid, subgrids, and thread blocks) into hardware registers or constant memory, resulting in the consumption of significant register resources and transmission bandwidth. Therefore, it is necessary to improve the transmission and storage mechanisms of the kernel function's execution space configuration parameters to enhance the execution efficiency of kernel functions in many-core parallel processors. Summary of the Invention

[0005] This disclosure provides a method for parallel execution of kernel functions and a many-core parallel processor.

[0006] This disclosure provides a method for parallel execution of a kernel function. The method includes: receiving instructions for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor; the instructions including: a grid dimension parameter, a first grid dimension partitioning parameter, and a second grid dimension partitioning parameter, wherein the grid dimension parameter indicates the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function, and the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary positions of the grid of the kernel function in different dimensions; determining the grid configuration corresponding to the kernel function based on the grid dimension parameter, the first grid dimension partitioning parameter, and the second grid dimension partitioning parameter; configuring parallel threads corresponding to the kernel function based on the grid configuration; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0007] This disclosure provides a method for parallel execution of a kernel function. The method includes: receiving an instruction for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor via thread blocks, the instruction including: block dimension parameters, a first block dimension partitioning parameter, and a second block dimension partitioning parameter, wherein the block dimension parameters indicate the total number of threads contained in each dimension of the thread block corresponding to the kernel function, and the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function on different dimensions; determining a thread block configuration corresponding to the kernel function based on the block dimension parameters, the first block dimension partitioning parameter, and the second block dimension partitioning parameter; configuring parallel threads corresponding to the kernel function based on the thread block configuration; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0008] This disclosure provides a method for parallel execution of a kernel function. The method includes: receiving an instruction for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor; the instruction includes: dimension configuration parameters, which include grid dimension parameters, block dimension parameters, and block boundary parameters, wherein the grid dimension parameters indicate the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function, the block dimension parameters indicate the total number of threads contained in each dimension of each thread block corresponding to the kernel function, and the block boundary parameters indicate the partition boundary between the grid dimension and the block dimension parameters; determining the grid configuration and thread block configuration corresponding to the kernel function based on the grid dimension parameters, the block dimension parameters, and the block boundary parameters; configuring parallel threads corresponding to the kernel function based on the grid configuration and thread block configuration; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0009] This disclosure provides a many-core parallel processor configured to perform the above-described method.

[0010] This disclosure provides an electronic device, including: a processor; and a memory, wherein the memory stores a computer-executable program, and when the processor executes the computer-executable program, the above-described method is performed.

[0011] This disclosure provides a computer-readable storage medium storing computer instructions that, when executed by a processor, implement the above-described method.

[0012] According to another aspect of this disclosure, a computer program product or computer program is provided, comprising computer instructions stored in a computer-readable storage medium. A processor of a computer device reads the computer instructions from the computer-readable storage medium and executes the computer instructions, causing the computer device to perform the methods provided in the foregoing aspects or various alternative implementations of the foregoing aspects.

[0013] This disclosure can reduce the amount of data stored in the registers of a many-core parallel processor, and at the same time reduce the number of bits required to transfer control information between modules of the many-core parallel processor. In this way, reducing the number of transmission bits can reduce the pressure on the control bus and improve hardware execution efficiency. Attached Figure Description

[0014] To more clearly illustrate the technical solutions of the embodiments of this disclosure, the accompanying drawings used in the description of the embodiments will be briefly introduced below. The accompanying drawings in the following description are merely exemplary embodiments of this disclosure.

[0015] Figure 1 A schematic diagram of the hardware architecture of a many-core parallel processor is shown.

[0016] Figure 2 A schematic diagram showing the organization of grids, thread blocks, and threads is provided.

[0017] Figure 3 A flowchart of a method for parallel execution of kernel functions according to an embodiment of the present disclosure is shown.

[0018] Figure 4 A schematic diagram is shown of the grid dimension parameters, the first grid dimension partitioning parameters, the second grid dimension partitioning parameters, and the sub-grid partitioning parameters according to an embodiment of the present disclosure.

[0019] Figure 5 A schematic diagram of a method for parallel execution of kernel functions according to an embodiment of the present disclosure is shown.

[0020] Figure 6 A schematic diagram of block dimension parameters, a first block dimension partitioning parameter, and a second block dimension partitioning parameter according to an embodiment of the present disclosure is shown.

[0021] Figure 7 A schematic diagram of a method for parallel execution of kernel functions according to an embodiment of the present disclosure is shown.

[0022] Figure 8 A schematic diagram of grid dimension parameters, block dimension parameters, and block boundary parameters according to embodiments of the present disclosure is shown. Detailed Implementation

[0023] To make the objectives, technical solutions, and advantages of the embodiments of this disclosure clearer, the technical solutions of the embodiments of this disclosure 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 disclosure. All other embodiments obtained by those skilled in the art based on the described embodiments of this disclosure without creative effort are within the scope of protection of this disclosure.

[0024] Unless otherwise defined, the technical or scientific terms used in this disclosure shall have the ordinary meaning understood by one of ordinary skill in the art to which this disclosure pertains. The terms "first," "second," and similar terms used in this disclosure do not indicate any order, quantity, or importance, but are merely used to distinguish different components. Terms such as "comprising" or "including" mean that the element or object preceding the word encompasses the elements or objects listed following the word and their equivalents, without excluding other elements or objects. Terms such as "connected" or "linked" are not limited to physical or mechanical connections, but can include electrical connections, whether direct or indirect. Terms such as "upper," "lower," "left," and "right" are used only to indicate relative positional relationships, and these relative positional relationships may change accordingly when the absolute position of the described object changes. To keep the following description of the embodiments of this disclosure clear and concise, detailed descriptions of some known functions and components are omitted.

[0025] The rapid development of artificial intelligence (AI) has placed higher demands on computing power, giving rise to AI chips specifically designed to accelerate AI workloads. These chips typically employ a many-core parallel processor architecture, leveraging their massive parallel computing capabilities to efficiently handle complex AI algorithms, such as neural network computations in deep learning. Processors employing a many-core parallel processor architecture will be referred to below as many-core parallel processors (e.g., GPUs).

[0026] Figure 1 A schematic diagram of the hardware architecture of a many-core parallel processor is shown.

[0027] like Figure 1 As shown, a many-core parallel processor is essentially a highly programmable multiprocessor array, typically composed of multiple Streaming Processor Clusters (SPCs). These SPCs achieve efficient data sharing through global L2 cache or high-bandwidth memory (HBM). A complete computational task can be completed independently by a single SPC, or it can be accelerated on a larger scale through collaborative processing by multiple SPCs.

[0028] Once the host calls the kernel function and starts, the many-core parallel processor enters a parallel execution state. At this time, the hardware creates a complete grid at the global level, which represents the entire execution space of the kernel function and accommodates all the threads required for the computing task.

[0029] To match the computational load with physical hardware resources, the driver and hardware scheduler dynamically divide the entire grid into several subgrids, each corresponding to a streaming processor cluster, ensuring that each SPC receives a balanced and independent computational load. Within an SPC, subgrids are further subdivided into multiple thread blocks. The thread block scheduler is responsible for distributing these thread blocks to numerous Compute Units (CUs) within the SPC, and all threads within the same thread block reside and execute within the same CU, enabling them to efficiently access the CU's dedicated shared memory.

[0030] Upon arrival at the compute unit, each thread block is divided into several warps (typically containing 32 threads). The internal warp scheduler of the compute unit continuously monitors all ready warps and issues identical instructions simultaneously to multiple cores within the CU using SIMT (Single Instruction Multiple Threading) technology. This allows the 32 threads within the same warp to execute the same operation in parallel within the same clock cycle. Each core is equipped with complete integer and floating-point arithmetic logic units (ALUs), special function units, etc., capable of independently performing various computational instructions such as addition, subtraction, multiplication, division, logical operations, and matrix operations.

[0031] In terms of data storage and access, each thread has its own private register file; all threads within the same thread block share the high-speed shared memory of the computing unit for fast exchange of intermediate results, with performance close to that of registers; when a thread needs to access a larger range of data, the instructions will pass through the L1 cache and shared memory inside the computing unit, the buffer at the streaming processor cluster level, and the L2 cache partition in sequence, and finally access the on-chip high-bandwidth memory HBM when necessary.

[0032] When computing tasks require cross-SPC collaboration, data sharing between different streaming processor clusters is accomplished through a unified global L2 cache or direct read / write to HBM, thereby ensuring data consistency across the entire grid.

[0033] Figure 2 A schematic diagram showing the organization of grids, thread blocks, and threads is provided.

[0034] like Figure 2 As shown, when the kernel function starts, all threads participating in parallel execution will jointly form a logical grid, which is the complete execution space for this kernel function call. The grid, subgrids (not shown), thread blocks, and individual threads all support up to three dimensions, represented by the X, Y, and Z axes respectively, which can naturally adapt to the mapping requirements of various parallel computing scenarios such as one-dimensional vectors, two-dimensional images, and three-dimensional volume data.

[0035] The entire grid is first logically divided into several sub-grids, each corresponding one-to-one with a streaming processor cluster (SPC) in the hardware architecture. Each sub-grid is composed of several thread blocks arranged in a regular pattern, and each thread block contains a fixed number of threads, tightly arranged along the X, Y, and Z three-dimensional dimensions. Thus, the hierarchical relationship from macro to micro is: grid, sub-grid, thread block, thread.

[0036] The size of the grid is defined by the gridDim.x, gridDim.y, and gridDim.z parameters (abbreviated as GridDimX / Y / Z); the size of the subgrid is defined by the subGridDim.x, subGridDim.y, and subGridDim.z parameters (abbreviated as SubGridDimX / Y / Z); and the number of threads in each thread block is defined by the blockDim.x, blockDim.y, and blockDim.z parameters (abbreviated as BlockDimX / Y / Z).

[0037] When the kernel function starts, the above parameters are passed through<gridDim,subGridDim,blockDim,...> The form is explicitly loaded into hardware registers or constant memory through a static transfer mechanism, where the subGridDim parameter is usually automatically configured by the driver based on the hardware topology, or generated by the compiler during the compilation phase.

[0038] In traditional implementations, these dimensional parameters are typically encoded using a fixed-width direct encoding method. For example, each dimensional parameter is often stored with a fixed width of 16 or 32 bits. The actual value of the parameter is either exactly the same as the encoded value or the encoded value plus 1. While this static fixed-width encoding mechanism is simple to implement, it suffers from significant resource waste: regardless of whether the actual mesh size is in the tens or millions, it requires the entire fixed-width encoding, resulting in a large number of redundant bits needing to be transferred and stored to the hardware when starting the kernel function. This consumes valuable register file, constant cache space, and configuration channel bandwidth. In massively parallel programs, this inefficient configuration transfer mechanism restricts the kernel function startup speed and overall execution efficiency.

[0039] Therefore, it is urgent to optimize the encoding, transmission, and storage mechanisms of execution configuration parameters such as grids, subgrids, and thread blocks to reduce configuration overhead and improve the startup efficiency, running performance, and resource utilization of many-core parallel processor kernel functions.

[0040] Therefore, this disclosure is proposed to optimize the encoding, transmission, and storage mechanisms for execution configuration parameters of the grid and subgrids.

[0041] Specifically, this disclosure proposes a method for parallel execution of kernel functions. The method includes: receiving instructions for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor, the instructions including: a grid dimension parameter, a first grid dimension partitioning parameter, and a second grid dimension partitioning parameter, wherein the grid dimension parameter indicates the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function, and the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary positions of the grid of the kernel function in different dimensions; determining the grid configuration corresponding to the kernel function based on the grid dimension parameter, the first grid dimension partitioning parameter, and the second grid dimension partitioning parameter; configuring parallel threads corresponding to the kernel function based on the grid configuration corresponding to the kernel function; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0042] Next, combined Figure 3 and Figure 4 To further illustrate the method for parallel execution of kernel functions proposed in this disclosure.

[0043] Figure 3 A flowchart of a method 30 for parallel execution of kernel functions according to an embodiment of the present disclosure is shown. Figure 4 A schematic diagram illustrating the grid dimension parameters, first grid dimension partitioning parameters, second grid dimension partitioning parameters, and sub-grid partitioning parameters according to embodiments of the present disclosure is shown. As previously described, a kernel function is a basic parallel computing task unit on a many-core parallel processor. Of course, embodiments of the present disclosure may include more or fewer operations, and the present disclosure is not limited thereto.

[0044] In operation S301, an instruction is received for parallel execution of the computational task corresponding to the kernel function on a many-core parallel processor. The instruction includes: a grid dimension parameter, a first grid dimension partitioning parameter, and a second grid dimension partitioning parameter. The grid dimension parameter indicates the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function, and the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary position of the grid of the kernel function in different dimensions.

[0045] Optionally, the instructions used to execute the computational tasks corresponding to the kernel functions in parallel on a many-core parallel processor refer to the configuration instructions for starting the kernel functions. These instructions carry the necessary execution configuration information so that the many-core parallel processor can organize parallel threads and execute the computational tasks accordingly. For example, in practical applications, this instruction might be represented as the kernel startup syntax <<<grid,block> >> This includes relevant dimensional information about the grid and thread blocks, allowing many-core parallel processors to allocate resources and schedule execution. Of course, this disclosure is not limited thereto.

[0046] As an example, the grid dimension parameters include: a first dimension parameter, a second dimension parameter, and a third dimension parameter. The first dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the first dimension; the second dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the second dimension; and the third dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the third dimension. This disclosure is not limited thereto.

[0047] Optionally, the first, second, and third dimensions are orthogonal to each other. The product of the first, second, and third dimension parameters is less than the maximum number of thread blocks supported by the many-core parallel processor. The following explanation uses the first dimension as the X-axis, the second dimension as the Y-axis, and the third dimension as the Z-axis. However, this disclosure is not limited to this.

[0048] Optionally, the grid dimension parameter can be represented as Figure 4 In , by the first dimension parameter Second dimension parameters and third-dimensional parameters It is pieced together, that is, .in, , and These represent the number of thread blocks contained in the mesh along the X, Y, and Z axes, respectively. The constraints are as follows: ,in, This is the maximum number of thread blocks supported by a many-core parallel processor. Of course, this disclosure is not limited to this.

[0049] Optionally, the bit width of the grid dimension parameter is a fixed value, while the bit widths of the first dimension parameter, the second dimension parameter, and the third dimension parameter are all variable values. As an example, The bit width is a fixed value ,and , and The bit width is variable. As an example, satisfy For example, in a typical scenario, If the mesh has 1024 thread blocks in the X dimension, 1 thread block in the Y dimension, and 1 thread block in the Z dimension, then The value is 1024. and The value is 1, the whole Encoding using variable-width first, second, and third dimension parameters requires fewer bits than traditional schemes. Of course, this disclosure is not limited to this.

[0050] Optionally, such as Figure 4 As shown, the first grid dimension partitioning parameter can be expressed as: This parameter describes the boundary between the first and second dimensions. Specifically, for The starting position of the least significant bit (LSB). For example, if Setting it to 10 means The lower 10 bits are used for encoding. Bits above that position begin to be used .at this time, The maximum value is In this example, by using the "minus1" encoding method, there can be a maximum of 1023+1=1024 thread blocks in the first dimension (e.g., the X dimension).

[0051] Optionally, the second grid dimension partitioning parameter can be expressed as This parameter indicates the boundary between the second dimension (e.g., the dimension along the Y-axis) and the third dimension (e.g., the dimension along the Z-axis). Specifically, for The starting position of the least significant bit (LSB). For example, if Setting it to 20 means that the encoding will begin from the 20th bit. The preceding bits are used for encoding. and Of course, this disclosure is not limited to this.

[0052] Optionally, the instruction further includes sub-mesh partitioning parameters, which indicate the granularity at which the mesh is divided into sub-mesh sections in the first, second, and third dimensions. For example... Figure 4 As shown in the example, the submesh partitioning parameter is expressed as ,satisfy Of course, this disclosure is not limited to this.

[0053] In operation S302, the grid configuration corresponding to the kernel function is determined based on the grid dimension parameters, the first grid dimension partitioning parameters, and the second grid dimension partitioning parameters.

[0054] Optionally, the grid configuration corresponding to the kernel function refers to the configuration from the compressed encoding. Decoded , and The specific value of this configuration is used for subsequent thread organization. For example, in the sample scenario, if... The bitstream is a specific value, and the actual quantity of each dimension can be extracted by dividing the parameters. Of course, this disclosure is not limited to this.

[0055] Optionally, the mesh configuration corresponding to the kernel function can be obtained by the following formulas (1)-(3). This bit segment extraction method allows for the use of variable bit widths based on actual values, thereby avoiding redundancy of fixed widths.

[0056] Formula (1)

[0057] Formula (2)

[0058] Formula (3)

[0059] Optionally, the grid configuration and sub-grid configuration corresponding to the kernel function can be determined based on the grid dimension parameter, the first grid dimension partitioning parameter, the second grid dimension partitioning parameter, and the sub-grid partitioning parameter.

[0060] Optionally, the subgrid configuration corresponding to the kernel function can be obtained through the following rules, thereby determining the number of thread blocks in the three dimensions of the subgrid: , and .

[0061] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being less than the first mesh dimension partitioning parameter, determining that the partitioning granularity of the mesh in the first dimension being divided into sub-mesh is the number of thread blocks indicated by a base-two exponent equal to a power of the sub-mesh partitioning parameter, and determining that the partitioning granularity of the mesh in the second and third dimensions being divided into sub-mesh is one thread block. Of course, this disclosure is not limited thereto. For example, when... lie in Within the bit width range, that is hour, -1 (if) If so, it can be considered a special case, that is, the grid is not divided on the X-axis. , For example, if ,but .at this time, It also uses the "minus1" encoding method, with the value "15" representing a granularity of 16 thread blocks. That is, at this time, every 16 thread blocks along the X direction are divided into a sub-grid.

[0062] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being equal to the first mesh dimension partitioning parameter, determining that the granularity of the mesh being divided into sub-mesh in the first dimension is the number of thread blocks indicated by the first dimension parameter, and determining that the mesh is not divided into sub-mesh in the second dimension and the third dimension. Of course, this disclosure is not limited thereto.

[0063] For example, when hour, , , , indicating each Each dimension is treated as a separate subgrid, with both the Y and Z dimensions being at the granularity of a single thread block. In this case, the subgrid appears as a complete "line" along the X dimension.

[0064] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being greater than the first mesh dimension partitioning parameter and less than the second mesh dimension partitioning parameter, determining that the partitioning granularity of the mesh as a sub-mesh in the first dimension is the first dimension parameter, determining that the partitioning granularity of the mesh as a sub-mesh in the second dimension is the number of thread blocks indicated by a base-two exponent equal to the sub-mesh partitioning parameter minus a power of the first mesh dimension partitioning parameter, and determining that the partitioning granularity of the mesh in the third dimension is one thread block. Of course, this disclosure is not limited thereto.

[0065] For example, when lie in Within the bit width range, that is hour, , , For example, if and ,but In this example, by employing the "minus1" encoding method, where the encoded N represents the actual value N+1, the number of bits required to represent each dimension is reduced in most cases. Therefore, This actually means dividing the grid into sub-grids every 32 thread blocks along the Y direction. Therefore, in this example, the sub-grid shape should be (GridDimX + 1). 32 1.

[0066] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being equal to the second mesh dimension partitioning parameter, determining that the partitioning granularity of the mesh being divided into sub-mesh in the first dimension is the first dimension parameter, determining that the partitioning granularity of the mesh being divided into sub-mesh in the second dimension is the second dimension parameter, and determining that the partitioning granularity of the mesh being divided into sub-mesh in the third dimension is a thread block. Of course, this disclosure is not limited thereto.

[0067] For example, when hour, , , This indicates that each XY plane is a sub-grid, and there is only one thread block in the Z dimension.

[0068] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being greater than the second mesh dimension partitioning parameter and less than the bit width of the mesh dimension parameter minus one, determining that the partitioning granularity of the mesh being divided into sub-mesh in the first dimension is the first dimension parameter, determining that the partitioning granularity of the mesh being divided into sub-mesh in the second dimension is the number of thread blocks indicated by the second dimension parameter, and determining that the partitioning granularity of the mesh being divided into sub-mesh in the third dimension is the number of thread blocks indicated by a base-two exponent equal to the sub-mesh partitioning parameter minus the power of the second mesh dimension partitioning parameter. Of course, this disclosure is not limited thereto.

[0069] For example, when lie in Within the bit width range, that is hour, , , For example, if and ,but . In practice, this means dividing the grid along the Z direction into sub-grids every 32 thread blocks. The shape of each sub-grid is (GridDimX + 1). (GridDimY + 1) 32. At this point, each subgrid consists of several XY "planes" stacked together along the Z direction to form a block.

[0070] Optionally, operation S302 further includes: in response to the sub-mesh partitioning parameter being equal to the bit width of the mesh dimension parameter minus one, determining that the granularity of the mesh being divided into sub-mesh in the first dimension is the number of thread blocks indicated by the first dimension parameter, determining that the granularity of the mesh being divided into sub-mesh in the second dimension is the number of thread blocks indicated by the second dimension parameter, and determining that the granularity of the mesh being divided into sub-mesh in the third dimension is the number of thread blocks indicated by the third dimension parameter. Of course, this disclosure is not limited thereto.

[0071] For example, when hour, , , This indicates that the entire grid is treated as a subgrid, which is equivalent to not dividing the grid into subgrids.

[0072] In operation S303, the parallel thread corresponding to the kernel function is configured based on the grid configuration corresponding to the kernel function.

[0073] Optionally, the grid configuration determines how parallel threads are organized, as the grid dimensions specify the total number and distribution of thread blocks, which the hardware uses to allocate thread blocks to computational units. For example, in matrix computation scenarios, if... and Based on the row and column division of the corresponding matrix, parallel threads can cover the entire matrix data.

[0074] For example, suppose a mesh is configured as , , The total number of thread blocks is 4096. The hardware will schedule a corresponding number of thread blocks accordingly, with threads within each thread block processing local data in parallel. Of course, this disclosure is not limited to this.

[0075] Optionally, in operation S303, the parallel threads corresponding to the kernel function can be further configured based on both the subgrid configuration and the grid configuration corresponding to the kernel function. For example, thread blocks can be further grouped and mapped based on the subgrid configuration corresponding to the kernel function. Since the subgrid configuration determines the number of thread blocks contained in each dimension of the subgrid and divides the entire grid into several subgrids, the thread blocks are mapped to hardware computing units according to the subgrid division, so that thread blocks in the same subgrid are continuously scheduled on the hardware to improve data locality or achieve phased execution. Of course, this disclosure is not limited thereto.

[0076] In the S304 operation, the parallel threads corresponding to the kernel functions are used to execute the computational tasks corresponding to the kernel functions in parallel on many-core parallel processors.

[0077] Optionally, the execution of parallel threads depends on the grid configuration, as the indexes of thread blocks and threads are calculated based on the grid dimension (and / or sub-grid dimension), thus achieving data parallelism or task parallelism. For example, in convolution operations, different thread blocks process different output feature locations, achieving computational load balancing. For instance, assuming the kernel function implements vector addition, and the grid configuration and / or sub-grid configuration are linearly distributed, each thread block processes a portion of the vector, and the entire task is completed after all thread blocks are executed in parallel.

[0078] Therefore, compared to the traditional method of directly storing parameters of each dimension, method 30 can reduce the storage capacity required when performing parallel computing tasks, thereby saving register space. Simultaneously, this disclosure can also save bandwidth for hardware transmission between different modules because the number of bits transmitted is reduced. The bandwidth of the control bus transmitting control information is typically small, thus alleviating the bandwidth pressure on the control bus.

[0079] For example, assuming a grid limit of 1024 thread blocks, the traditional approach uses a fixed 10 bits per dimension to represent the maximum value of 1023, requiring a total of 30 bits. However, in method 30, for specific grid and thread block shapes... , , The actual value is much smaller than the maximum value, so variable bit-width encoding can represent each dimension using only 1 to 10 bits, thus significantly reducing the number of bits required for storage. For example, in the case of a one-dimensional grid, , , Only 12 bits are needed to represent the entire Traditional methods require a fixed 30 bits to represent three-dimensional dimensions.

[0080] In another aspect, method 30 unifies the encapsulation of grid, subgrid, and thread block attributes into a smaller bit width (e.g., Bit, , , Each register uses 4 bits (24 bits total), allowing for the storage of less data and reducing the number of bits required to transfer control information between modules. Since the bandwidth of the control bus is typically small, reducing the number of bits transmitted lowers the load on the control bus and improves hardware execution efficiency.

[0081] Next, refer to Figure 5 and Figure 6 To describe another embodiment of this disclosure, this embodiment uses the same principle as method 30 to configure thread blocks that execute kernel functions in parallel, in order to optimize the encoding, transmission and storage mechanism of the execution configuration parameters of the thread blocks.

[0082] Specifically, this disclosure provides a method for parallel execution of a kernel function, comprising: receiving an instruction for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor via thread blocks, the instruction comprising: a block dimension parameter, a first block dimension partitioning parameter, and a second block dimension partitioning parameter, wherein the block dimension parameter indicates the total number of threads contained in each dimension of the thread block corresponding to the kernel function, and the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function on different dimensions; determining a thread block configuration corresponding to the kernel function based on the block dimension parameter, the first block dimension partitioning parameter, and the second block dimension partitioning parameter; configuring parallel threads corresponding to the kernel function based on the thread block configuration; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0083] Figure 5 A schematic diagram of a method 50 for parallel execution of kernel functions according to an embodiment of the present disclosure is shown. Figure 6 A schematic diagram illustrating block dimension parameters, a first block dimension partitioning parameter, and a second block dimension partitioning parameter according to an embodiment of this disclosure is shown. As previously described, a kernel function is a basic parallel computing task unit on a many-core parallel processor. Of course, embodiments of this disclosure may include more or fewer operations, and this disclosure is not limited thereto.

[0084] In operation S501, an instruction is received for executing the computational task corresponding to the kernel function in parallel on a many-core parallel processor through thread blocks. The instruction includes: block dimension parameters, a first block dimension partitioning parameter, and a second block dimension partitioning parameter. The block dimension parameters indicate the total number of threads contained in each dimension of the thread block corresponding to the kernel function, and the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function in different dimensions.

[0085] Optionally, the instructions for executing the computational tasks corresponding to the kernel function in parallel on a many-core parallel processor refer to the configuration instructions for starting the kernel function. These instructions carry the necessary execution configuration information so that the many-core parallel processor can organize parallel threads and execute computational tasks accordingly. For example, in practical applications, these instructions may also be represented as kernel startup syntax, which includes relevant dimensional information about the grid and thread blocks, thereby allowing the many-core parallel processor to allocate resources and schedule execution. Of course, this disclosure is not limited thereto.

[0086] As an example, the block dimension parameters include: a first thread count parameter, a second thread count parameter, and a third thread count parameter. The first thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the first dimension; the second thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the second dimension; and the third thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the third dimension. Of course, this disclosure is not limited thereto.

[0087] Optionally, the first, second, and third dimensions are orthogonal to each other. The product of the first thread quantity parameter, the second thread quantity parameter, and the third thread quantity parameter is less than the maximum number of threads supported by the thread block. The following explanation uses the first dimension as the X-axis, the second dimension as the Y-axis, and the third dimension as the Z-axis. However, this disclosure is not limited to this.

[0088] Optionally, the block dimension parameter can be represented as Figure 6 In The number of threads is determined by the first thread count parameter. Second thread quantity parameter and the number of third threads parameter It is pieced together, that is, .in, , and These represent the number of threads contained in the thread block along the X, Y, and Z axes, respectively. The constraints are as follows: ,in, This represents the maximum number of threads supported by the thread block. Assume the thread block contains threads in three dimensions. , , If there are 100 threads, then the total number of threads in each thread block is 100. Less than the assumed Of course, this disclosure is not limited to this.

[0089] Optionally, the bit width of the block dimension parameter is a fixed value, while the bit widths of the first dimension parameter, the second dimension parameter, and the third dimension parameter are all variable values. As an example, The bit width is a fixed value ,and , and The bit width is variable. Assume... The bit allocation is as follows: A bit can represent 1 to 8 threads; A bit can represent 1 to 8 threads; Each bit can represent 1 to 64 threads; thus, three-dimensional thread block information can be fully represented with 12 bits, while saving storage and transmission overhead.

[0090] Optionally, such as Figure 6 As shown, the first block dimension partitioning parameter can be represented as This parameter describes the boundary between the first and second dimensions. Specifically, for The starting position of the least significant bit (LSB). For example, if Setting it to 10 means The lower 10 bits are used for encoding. Bits above that position begin to be used .

[0091] Optionally, the second grid dimension partitioning parameter can be expressed as This parameter indicates the boundary between the second dimension (e.g., the dimension along the Y-axis) and the third dimension (e.g., the dimension along the Z-axis). Specifically, for The starting position of the least significant bit (LSB). For example, if Setting it to 20 means that the encoding will begin from the 20th bit. The preceding bits are used for encoding. and Of course, this disclosure is not limited to this.

[0092] In operation S502, the thread block configuration corresponding to the kernel function is determined based on the block dimension parameters, the first block dimension partitioning parameters, and the second block dimension partitioning parameters.

[0093] Optionally, the thread block configuration corresponding to the kernel function refers to the size and arrangement of each thread block in three-dimensional space on a many-core parallel processor, that is, the number of threads contained in each thread block in the X, Y, and Z dimensions and their relative position in the entire thread block grid. For example, if BlockDimX=7, BlockDimY=3, and BlockDimZ=1, then each thread block contains 8×4×2=64 threads in three dimensions, and the entire thread block grid consists of multiple such thread blocks, covering the entire data space of the computation task. Of course, this disclosure is not limited to this.

[0094] Optionally, the thread block configuration corresponding to the kernel function can be obtained through the following formulas (4)-(6). This bit segment extraction method allows for the use of variable bit widths based on the actual values, thereby avoiding redundancy of fixed widths.

[0095] Formula (4)

[0096] Formula (5)

[0097] Formula (6)

[0098] In operation S503, the parallel thread corresponding to the kernel function is configured based on the thread block configuration corresponding to the kernel function.

[0099] Optionally, thread block configuration determines how parallel threads are organized, because the arrangement and size of thread blocks in various dimensions determine how thread indices are allocated and how threads are scheduled on hardware computing units. For example, threads within each thread block are arranged according to... , , The threads are numbered sequentially, and the hardware generates an index for each thread within the thread block, allowing each thread to access the data region corresponding to the computational task. This method precisely maps threads to data while ensuring the correctness and efficiency of parallel execution.

[0100] In the S504 operation, the parallel threads corresponding to the kernel functions are used to execute the computational tasks corresponding to the kernel functions in parallel on many-core parallel processors.

[0101] Optionally, each thread calculates its global index based on the thread block configuration and the thread's index within the block, accesses the corresponding data element, and executes the kernel function's computation logic. For example, in matrix computation, a thread can locate the row and column positions of the matrix using the global index and perform the corresponding addition or multiplication operations. Multiple threads simultaneously execute computations on different data elements, achieving task parallelization and thus fully utilizing the computing resources of the many-core parallel processor.

[0102] Therefore, compared to the traditional method of directly storing parameters of each dimension, Method 50 can reduce the storage capacity required when performing parallel computing tasks, thereby saving register space. Furthermore, if each dimension is directly represented by a fixed 10 bits, describing the thread block configuration would require 90 bits. This disclosure, through compression encoding, only requires 24 bits to completely describe the thread block configuration, thus saving register storage space and reducing the bandwidth required for transmitting control information between modules.

[0103] Next, refer to Figure 7 and Figure 8 To describe another embodiment of this disclosure, this embodiment employs the same principles as methods 30 and 50 to configure grids, subgrids, and thread blocks for parallel execution of kernel functions, in order to optimize the encoding, transmission, and storage mechanisms of execution configuration parameters for grids, subgrids, and thread blocks.

[0104] Specifically, this disclosure provides a method for parallel execution of a kernel function, comprising: receiving an instruction for parallel execution of a computational task corresponding to the kernel function on a many-core parallel processor, the instruction comprising: dimension configuration parameters, the dimension configuration parameters including grid dimension parameters, block dimension parameters, and block boundary parameters, wherein the grid dimension parameters indicate the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function, the block dimension parameters indicate the total number of threads contained in each dimension of each thread block corresponding to the kernel function, and the block boundary parameters indicate the partition boundary between the grid dimension and the block dimension parameters; determining the grid configuration and thread block configuration corresponding to the kernel function based on the grid dimension parameters, the block dimension parameters, and the block boundary parameters; configuring parallel threads corresponding to the kernel function based on the grid configuration and thread block configuration corresponding to the kernel function; and using the parallel threads corresponding to the kernel function to execute the computational task corresponding to the kernel function in parallel on a many-core parallel processor.

[0105] Figure 7 A schematic diagram of a method 70 for parallel execution of kernel functions according to an embodiment of the present disclosure is shown. Figure 8 A schematic diagram illustrating grid dimension parameters, block dimension parameters, and block boundary parameters according to embodiments of the present disclosure is shown. As previously described, a kernel function is a basic parallel computing task unit on a many-core parallel processor. Of course, embodiments of the present disclosure may include more or fewer operations, and the present disclosure is not limited thereto.

[0106] In operation S701, an instruction is received for parallel execution of the computational task corresponding to the kernel function on a many-core parallel processor. The instruction includes: dimension configuration parameters, which include grid dimension parameters, block dimension parameters, and block boundary parameters. The grid dimension parameters indicate the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function. The block dimension parameters indicate the total number of threads contained in each dimension of each thread block corresponding to the kernel function. The block boundary parameters indicate the partition boundary between the grid dimension and the block dimension parameters.

[0107] Optionally, the dimension configuration parameter can be represented as Figure 8 In , determined by the grid dimension parameter and block dimension parameters It is pieced together, that is, Block boundary parameters Indicates grid dimension parameters and block dimension parameters The dividing line between them.

[0108] Optionally, the bit width of the dimension configuration parameter is a fixed value, while the bit width of both the grid dimension parameter and the block dimension parameter are variable values. With a fixed bit width, the grid dimension parameter and block dimension parameters It can have a wider range of expression, thus describing a larger kernel function execution space with a smaller bit width.

[0109] As an example, the grid dimension parameters include: a first dimension parameter, a second dimension parameter, and a third dimension parameter. The first dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the first dimension. The second dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the second dimension. The third dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the third dimension.

[0110] Optionally, the block dimension parameters include: a first thread count parameter, a second thread count parameter, and a third thread count parameter. The first thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the first dimension; the second thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the second dimension; and the third thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the third dimension. Of course, this disclosure is not limited thereto.

[0111] Optionally, the first, second, and third dimensions are orthogonal to each other. The product of the first thread count parameter, the second thread count parameter, and the third thread count parameter is less than the maximum number of threads supported by the thread block. The following explanation uses the first dimension as the X-axis, the second dimension as the Y-axis, and the third dimension as the Z-axis.

[0112] Optionally, the instructions further include: a first grid dimension partitioning parameter and a second grid dimension partitioning parameter, wherein the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary positions of the kernel function's grid in different dimensions. Of course, this disclosure is not limited thereto.

[0113] Optionally, the instructions further include: sub-mesh partitioning parameters, which indicate the granularity at which the mesh is divided into sub-mesh sections in each dimension. Of course, this disclosure is not limited thereto.

[0114] Optionally, the instruction further includes: a first block dimension partitioning parameter and a second block dimension partitioning parameter, wherein the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function in different dimensions.

[0115] The parameters of grid dimension, block dimension, first grid dimension division, second grid dimension division, subgrid division, first block dimension division, and second block dimension division have been described in detail previously, and will not be repeated here.

[0116] For example, suppose we need to specify a support for many-core parallel processors. The thread's kernel function execution space. Dim in method 70 only needs to receive 37 bits of instructions. The grid shape can be... And the thread block shape is Or the grid shape is And the thread block shape is .at the same time, , , , , and Encoding with 6 bits each requires a total of 67 bits. In contrast, the traditional scheme... , , , , , , , and Each is represented using 32 bits, and a total of 288 bits are required when specifying the same kernel function execution space.

[0117] In operation S702, the mesh configuration and thread block configuration corresponding to the kernel function are determined based on the mesh dimension parameter, the block dimension parameter and the block boundary parameter.

[0118] For example, thread blocks are allocated to different computing units on the processor according to a defined grid configuration, and the arrangement and indexing of threads are organized within each thread block according to the defined thread block configuration. Optionally, the grid can be divided into multiple subgrids in conjunction with subgrid partitioning to facilitate phased scheduling and execution. Grid configuration and thread block configuration have already been described in detail previously and will not be repeated here.

[0119] In operation S703, the parallel thread corresponding to the kernel function is configured based on the grid configuration and thread block configuration corresponding to the kernel function.

[0120] Optionally, for example, a corresponding number of threads can be allocated to each thread block according to the thread block configuration, and a position index in the execution space can be allocated to each thread block according to the grid configuration. In this way, each parallel thread has a clear thread block index and thread index, thereby enabling it to access the corresponding data region during execution. The process of configuring the parallel threads corresponding to the kernel function has been described in detail previously, and will not be repeated here.

[0121] In operation S704, the computational task corresponding to the kernel function is executed in parallel on a many-core parallel processor using the parallel thread corresponding to the kernel function.

[0122] Therefore, compared to directly storing each dimension parameter separately, method 70 can reduce the required storage capacity when performing parallel computing tasks, thereby reducing the occupation of register space. Simultaneously, the compressed encoding format can reduce the number of bits required to transmit control information between different modules. For example, using a fixed 10 bits to represent each dimension requires 90 bits to fully describe the thread block configuration, while using the compressed encoding method in this embodiment, only 24 bits are needed to describe the corresponding thread block configuration, thereby reducing register storage overhead and bandwidth requirements when transmitting control information between modules.

[0123] As an example, with a mesh shape of 2×4×3, a thread block shape of 7×3×5, and a submesh size of 2×1×1, method 70 requires only 24 bits to describe the mesh, submesh, and thread block configuration, while the traditional approach requires 90 bits to achieve the same configuration. Therefore, in terms of register usage, control bus transmission, and instruction parsing, this method reduces storage and bandwidth consumption while maintaining accurate description of the mesh and thread block configuration.

[0124] In addition, this disclosure also provides a many-core parallel processor configured to perform the above-described methods.

[0125] Furthermore, this disclosure also provides an electronic device, including: a processor; and a memory, wherein the memory stores a computer-executable program, and when the processor executes the computer-executable program, the above-described method is performed.

[0126] The processor in this disclosure embodiment can be an integrated circuit chip with signal processing capabilities. The processor can be a general-purpose processor, a digital signal processor (DSP), an application-specific integrated circuit (ASIC), an off-the-shelf programmable gate array (FPGA), or other programmable logic devices, discrete gate or transistor logic devices, or discrete hardware components. It can implement or execute the methods, operations, and logic block diagrams disclosed in this disclosure embodiment. The general-purpose processor can be a microprocessor or any conventional processor, and can be based on an x86 architecture or an ARM architecture.

[0127] In general, the various exemplary embodiments of this disclosure can be implemented in hardware, dedicated circuitry, software, firmware, logic, or any combination thereof. Some aspects can be implemented in hardware, while others can be implemented in firmware or software that can be executed by a controller, microprocessor, or other computing device. When aspects of the embodiments of this disclosure are illustrated or described as block diagrams, flowcharts, or using some other graphical representation, it will be understood that the blocks, apparatuses, systems, techniques, or methods described herein can be implemented as non-limiting examples in hardware, software, firmware, dedicated circuitry or logic, general-purpose hardware or controllers or other computing devices, or some combination thereof.

[0128] This disclosure also provides a computer program product or computer program including computer instructions stored in a computer-readable storage medium. A processor of a computer device reads the computer instructions from the computer-readable storage medium and executes the computer instructions, causing the computer device to perform a method according to an embodiment of this disclosure.

[0129] It should be noted that the flowcharts and block diagrams in the accompanying drawings illustrate the architecture, functionality, and operation of possible implementations of systems, methods, and computer program products according to various embodiments of this disclosure. In this regard, each block in a flowchart or block diagram may represent a module, segment, or portion of code containing one or more executable instructions for implementing the specified logical function. It should also be noted that in some alternative implementations, the functions indicated in the blocks may occur in a different order than those indicated in the drawings. For example, two consecutively indicated blocks may actually be executed substantially in parallel, and they may sometimes be executed in reverse order, depending on the functions involved. It should also be noted that each block in the block diagrams and / or flowcharts, and combinations of blocks in the block diagrams and / or flowcharts, can be implemented using a dedicated hardware-based system that performs the specified function or operation, or using a combination of dedicated hardware and computer instructions.

[0130] In general, the various exemplary embodiments of this disclosure can be implemented in hardware or dedicated circuitry, software, firmware, logic, or any combination thereof. Some aspects can be implemented in hardware, while others can be implemented in firmware or software that can be executed by a controller, microprocessor, or other computing device. When aspects of the embodiments of this disclosure are illustrated or described as block diagrams, flowcharts, or using some other graphical representation, it will be understood that the blocks, apparatuses, systems, techniques, or methods described herein can be implemented as non-limiting examples in hardware, software, firmware, dedicated circuitry or logic, general-purpose hardware or controllers or other computing devices, or some combination thereof.

[0131] The exemplary embodiments of this disclosure described in detail above are merely illustrative and not restrictive. Those skilled in the art will understand that various modifications and combinations of these embodiments or their features can be made without departing from the principles and spirit of this disclosure, and such modifications should fall within the scope of this disclosure.

Claims

1. A method for parallel execution of kernel functions, characterized in that, The method includes: The system receives instructions for parallel execution of computational tasks corresponding to the kernel function on a many-core parallel processor. The instructions include: a grid dimension parameter, a first grid dimension partitioning parameter, and a second grid dimension partitioning parameter. The grid dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in each dimension, and the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary positions of the grid of the kernel function in different dimensions. Based on the grid dimension parameters, the first grid dimension partitioning parameters, and the second grid dimension partitioning parameters, determine the grid configuration corresponding to the kernel function; Based on the grid configuration corresponding to the kernel function, configure the parallel threads corresponding to the kernel function, and The computational task corresponding to the kernel function is executed in parallel on a many-core parallel processor using the parallel thread corresponding to the kernel function.

2. The method as described in claim 1, characterized in that, The grid dimension parameters include: a first dimension parameter, a second dimension parameter, and a third dimension parameter. The first dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the first dimension. The second dimension parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the second dimension. The third-dimensional parameter indicates the total number of thread blocks contained in the grid corresponding to the kernel function in the third dimension, and The first dimension, the second dimension, and the third dimension are orthogonal to each other.

3. The method as described in claim 2, characterized in that, The product of the first dimension parameter, the second dimension parameter, and the third dimension parameter is less than the maximum number of thread blocks supported by the many-core parallel processor.

4. The method as described in claim 2, characterized in that, The bit width of the grid dimension parameter is a fixed value, while the bit widths of the first dimension parameter, the second dimension parameter, and the third dimension parameter are all variable values.

5. The method as described in claim 2, characterized in that, The first grid dimension division parameter indicates the boundary between the first and second dimensions, and the second grid dimension division parameter indicates the boundary between the second and third dimensions.

6. The method as described in claim 2, characterized in that, The instruction further includes sub-mesh partitioning parameters, which indicate the granularity at which the mesh is divided into sub-mesh sections in the first, second, and third dimensions. Determining the mesh configuration corresponding to the kernel function further includes: Based on the grid dimension parameter, the first grid dimension partitioning parameter, the second grid dimension partitioning parameter, and the sub-grid partitioning parameter, the grid configuration and sub-grid configuration corresponding to the kernel function are determined.

7. The method as described in claim 6, characterized in that, Determining both the mesh configuration and the sub-mesh configuration corresponding to the kernel function includes: In response to the sub-mesh partitioning parameter being less than the first mesh dimension partitioning parameter, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the number of thread blocks indicated by a power of two (base 2) and an exponent equal to the power of the sub-mesh partitioning parameter; and the granularity of the mesh being divided into sub-mesh in the second and third dimensions is determined to be one thread block; and In response to the sub-mesh partitioning parameter being equal to the first mesh dimension partitioning parameter, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the number of thread blocks indicated by the first dimension parameter, and the granularity of the mesh being divided into sub-mesh in the second and third dimensions is determined to be one thread block.

8. The method as described in claim 6, characterized in that, Determining both the mesh configuration and the sub-mesh configuration corresponding to the kernel function includes: In response to the sub-mesh partitioning parameter being greater than the first mesh dimension partitioning parameter and less than the second mesh dimension partitioning parameter, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the first dimension parameter; the granularity of the mesh being divided into sub-mesh in the second dimension is determined to be the number of thread blocks indicated by a base-two exponent equal to the sub-mesh partitioning parameter minus the power of the first mesh dimension partitioning parameter; and the granularity of the mesh being divided into sub-mesh in the third dimension is determined to be one thread block; and In response to the sub-mesh partitioning parameter being equal to the second mesh dimension partitioning parameter, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the number of thread blocks indicated by the first dimension parameter; the granularity of the mesh being divided into sub-mesh in the second dimension is determined to be the number of thread blocks indicated by the second dimension parameter; and the granularity of the mesh being divided into sub-mesh in the third dimension is determined to be one thread block.

9. The method as described in claim 6, characterized in that, Determining both the mesh configuration and the sub-mesh configuration corresponding to the kernel function includes: In response to the sub-mesh partitioning parameter being greater than the second mesh dimension partitioning parameter and less than the bit width of the mesh dimension parameter minus one, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the number of thread blocks indicated by the first dimension parameter; the granularity of the mesh being divided into sub-mesh in the second dimension is determined to be the number of thread blocks indicated by the second dimension parameter; and the granularity of the mesh being divided into sub-mesh in the third dimension is determined to be the number of thread blocks base two, with an exponent equal to the sub-mesh partitioning parameter minus the power of the second mesh dimension partitioning parameter; and In response to the sub-mesh partitioning parameter being equal to the bit width of the mesh dimension parameter minus one, the granularity of the mesh being divided into sub-mesh in the first dimension is determined to be the number of thread blocks indicated by the first dimension parameter; the granularity of the mesh being divided into sub-mesh in the second dimension is determined to be the number of thread blocks indicated by the second dimension parameter; and the granularity of the mesh being divided into sub-mesh in the third dimension is determined to be the number of thread blocks indicated by the third dimension parameter.

10. The method as described in claim 6, characterized in that, The configuration of parallel threads corresponding to the kernel function based on the grid configuration includes: Based on both the grid configuration and the subgrid configuration corresponding to the kernel function, configure the parallel thread corresponding to the kernel function.

11. A method for parallel execution of kernel functions, characterized in that, The method includes: The system receives an instruction for executing the computational task corresponding to the kernel function in parallel on a many-core parallel processor through thread blocks. The instruction includes: a block dimension parameter, a first block dimension partitioning parameter, and a second block dimension partitioning parameter. The block dimension parameter indicates the total number of threads contained in each dimension of the thread block corresponding to the kernel function, and the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function in different dimensions. Based on the block dimension parameters, the first block dimension partitioning parameters, and the second block dimension partitioning parameters, the thread block configuration corresponding to the kernel function is determined. Based on the thread block configuration corresponding to the kernel function, configure the parallel thread corresponding to the kernel function, and The computational task corresponding to the kernel function is executed in parallel on a many-core parallel processor using the parallel thread corresponding to the kernel function.

12. The method as described in claim 11, characterized in that, The block dimension parameters include: a first thread count parameter, a second thread count parameter, and a third thread count parameter. The first thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the first dimension. The second thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the second dimension. The third thread count parameter indicates the total number of threads contained in the thread block corresponding to the kernel function in the third dimension, and The first dimension, the second dimension, and the third dimension are orthogonal to each other.

13. The method as described in claim 12, characterized in that, The product of the first thread count parameter, the second thread count parameter, and the third thread count parameter is less than the maximum number of threads supported by the thread block.

14. The method as described in claim 12, characterized in that, The bit width of the block dimension parameter is a fixed value, while the bit widths of the first thread quantity parameter, the second thread quantity parameter, and the third thread quantity parameter are all variable values.

15. The method as described in claim 12, characterized in that, The first block dimension partitioning parameter indicates the boundary position between the first and second dimensions, and the second block dimension partitioning parameter indicates the boundary position between the second and third dimensions.

16. A method for parallel execution of kernel functions, characterized in that, The method includes: The system receives instructions for executing the computational task corresponding to the kernel function in parallel on a many-core parallel processor. The instructions include: dimension configuration parameters, which include grid dimension parameters, block dimension parameters, and block boundary parameters. The grid dimension parameters indicate the total number of thread blocks contained in each dimension of the grid corresponding to the kernel function. The block dimension parameters indicate the total number of threads contained in each dimension of each thread block corresponding to the kernel function. The block boundary parameters indicate the partition boundary between the grid dimension parameters and the block dimension parameters. Based on the mesh dimension parameters, the block dimension parameters, and the block boundary parameters, determine the mesh configuration and thread block configuration corresponding to the kernel function; Based on the grid configuration and thread block configuration corresponding to the kernel function, configure the parallel threads corresponding to the kernel function, and The computational task corresponding to the kernel function is executed in parallel on a many-core parallel processor using the parallel thread corresponding to the kernel function.

17. The method as described in claim 16, characterized in that, The instruction further includes: a first grid dimension partitioning parameter and a second grid dimension partitioning parameter, wherein the first grid dimension partitioning parameter and the second grid dimension partitioning parameter respectively indicate the boundary positions of the kernel function's grid in different dimensions; or The instruction further includes: sub-mesh partitioning parameters, which indicate the granularity at which the mesh is divided into sub-mesh sections in each dimension; or The instruction further includes: a first block dimension partitioning parameter and a second block dimension partitioning parameter, wherein the first block dimension partitioning parameter and the second block dimension partitioning parameter respectively indicate the boundary positions of the thread blocks of the kernel function in different dimensions.

18. The method as described in claim 16, characterized in that, The bit width of the dimension configuration parameter is a fixed value, while the bit width of both the grid dimension parameter and the block dimension parameter are variable values.

19. A many-core parallel processor, characterized in that, The many-core parallel processor is configured to perform the method as described in any one of claims 1-18.

Images