Method, device, electronic equipment and storage medium for matrix loading

Abstract

Method, device, electronic equipment and storage medium for matrix loading. The method comprises: receiving logical layouts of input and output tensors and a physical layout type of an output register matrix; determining physical layout information of the output register matrix; performing a traversal process for the output register matrix based on the physical layout information; and loading the matrix from memory to registers based on the traversal process, wherein the traversal process comprises: calculating a current register address of a current register; determining a first element logical coordinate of a current output tensor corresponding to the current register; determining a first element logical coordinate of a current input tensor corresponding to the current output tensor; and calculating a current input memory address corresponding to the first element logical coordinate of the current input tensor. The method can directly support high-dimensional tensor and non-continuous memory / register data loading, while automatically and flexibly configuring instruction masks and supporting complex semantic scenarios such as broadcasting.

Description

Technical Field

[0001] Embodiments of this disclosure relate to methods, apparatus, electronic devices, and storage media for matrix loading. Background Technology

[0002] In the field of GPGPU (General-Purpose Graphics Processing Unit) operator development, intermediate representation (IR) techniques are commonly used to simplify the development process and improve code generation efficiency. Among these, matrix loading (ldmatrix) IR is a common intermediate representation used to load data from linear memory into registers across multiple threads. Current matrix loading methods for implementing IR typically require both input and output tensors to be two-dimensional, support only contiguous memory layouts (such as row-major or column-major), and the loading process is often limited to pointwise data mapping scenarios.

[0003] Matrix loading IR implemented using this matrix loading method has limitations: First, it has strict restrictions on tensor dimensions, making it unable to directly handle high-dimensional or non-contiguous memory / register data loading. Developers need to manually write loop code to implement related functions, increasing development complexity and the probability of errors. Second, in scenarios requiring instruction masks for data filtering, the masks must be manually calculated and configured by the developer, and all instruction masks generated in a single IR call must be consistent, resulting in insufficient flexibility. Furthermore, existing IRs lack support for complex data loading semantics such as broadcast, limiting their application scope. These problems collectively lead to low development efficiency and limited code performance. Summary of the Invention

[0004] At least one embodiment of this disclosure provides a method, apparatus, electronic device, and storage medium for matrix loading to address the above-mentioned problems.

[0005] According to at least one embodiment of this disclosure, a method for loading a matrix is ​​provided. The method includes: receiving a logical layout of an input tensor, a logical layout of an output tensor, and a physical layout type of an output register matrix specified by a user; determining physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix; performing a traversal process for the output register matrix based on the physical layout information; and loading the matrix from memory into registers based on the traversal process. The traversal process includes iteratively performing the following operations for each register in the output register matrix until all registers in the output register matrix have been traversed: calculating the current register address of the current register based on the physical layout information; determining the logical coordinates of the first element of the current output tensor corresponding to the current register based on the current register address and the physical layout information; determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on the logical coordinates of the first element of the current output tensor; and calculating the current input memory address corresponding to the logical coordinates of the first element of the current input tensor based on the logical coordinates of the first element of the current input tensor.

[0006] For example, according to at least one embodiment of the present disclosure, the physical layout information of the output register matrix includes the physical layout of the output register matrix and a one-to-one first mapping relationship between the logical layout of the output tensor and the physical layout of the output register matrix.

[0007] For example, according to at least one embodiment of this disclosure, the physical layout of the output register matrix includes a register index layout and a thread layout, and the logical layout of the input tensor, the logical layout of the output tensor, the register index layout, and the thread layout each include a dimension vector and a step vector of the layout, wherein the dimension vector indicates the dimension information of the layout, and the step vector indicates the step size of each dimension.

[0008] For example, according to at least one embodiment of this disclosure, the physical layout type of the output register matrix corresponds to the hardware architecture or instruction mode.

[0009] For example, according to at least one embodiment of this disclosure, the method further includes: allocating a memory base address in the memory for the matrix and a register base address in the register for the matrix before performing the traversal process.

[0010] For example, according to at least one embodiment of the present disclosure, performing a traversal process for the output register matrix based on physical layout information includes: performing a traversal process for the output register matrix based on the dimension vector and step vector of the register index layout.

[0011] For example, according to at least one embodiment of this disclosure, performing a traversal process for the output register matrix based on the dimension vector and step vector of the register index layout includes: performing a traversal process on the registers in the output register matrix in ascending order of step size for each dimension based on the dimension coordinates of the registers in the register index layout.

[0012] For example, according to at least one embodiment of this disclosure, calculating the current register address of the current register based on physical layout information includes calculating based on the following formula: current register address = register base address + dimension coordinate of the current register · (step vector / total number of threads) * size of a single element / size of a single register, where the "·" symbol represents dot product.

[0013] For example, according to at least one embodiment of this disclosure, determining the first element logical coordinates of the current output tensor corresponding to the current register based on the current register address and physical layout information includes: determining the first element logical coordinates of the current output tensor corresponding to the current register based on the current register address, taking the first thread of the current register as the starting point, and according to a first mapping relationship.

[0014] For example, according to at least one embodiment of this disclosure, determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on the logical coordinates of the first element of the current output tensor includes: determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on a second mapping relationship between the input tensor and the output tensor.

[0015] For example, according to at least one embodiment of this disclosure, calculating the current input memory address corresponding to the logical coordinates of the first element of the current input tensor based on the logical coordinates of the first element of the current input tensor includes calculating based on the following formula: current input memory address = memory base address + logical coordinates of the first element of the current input tensor · step vector * size of a single element, where the "·" symbol represents dot product.

[0016] For example, according to at least one embodiment of this disclosure, the second mapping relationship is based on the semantic relationship between the input tensor and the output tensor, and the semantic relationship includes point-to-point data loading or broadcast data loading.

[0017] For example, according to at least one embodiment of this disclosure, the traversal process further includes: determining the logical coordinates of each element in the current output tensor corresponding to each thread based on the current register address and thread layout, according to a first mapping relationship, and determining the current thread mask by comparing the logical coordinates of each element corresponding to each thread with the logical layout of the input tensor, the current thread mask being used to prohibit the execution of a portion of the threads.

[0018] For example, according to at least one embodiment of this disclosure, the traversal process further includes: generating a single assembly instruction based on the current register address, the current input memory address, and the current thread mask.

[0019] For example, according to at least one embodiment of this disclosure, the traversal process further includes: iteratively performing a successive traversal to continuously traverse a plurality of registers in the output register matrix; and generating a single assembly instruction for the plurality of registers.

[0020] For example, according to at least one embodiment of this disclosure, the continuous traversal includes: after determining the current register address, the current input memory address, and the current thread mask, continuing to the next traversal to determine the next register address, the next input memory address, and the next thread mask in the next traversal; comparing the current register address, the current input memory address, and the current thread mask with the next register address, the next input memory address, and the next thread mask; in response to the current register address being consecutive with the next register address, the current input memory address being consecutive with the next input memory address, the current thread mask being the same as the next thread mask, the number of registers that have been continuously traversed not reaching the upper limit of the number of registers that can be continuously traversed supported by the hardware, and the traversal of all registers not yet being completed, continuing to continuously traverse the registers in the output register matrix; otherwise, generating a single assembly instruction based on the first register address, the first input memory address, the thread mask of the continuous traversal, the execution granularity, and the user configuration information in the first traversal of the continuous traversal, and continuing the traversal process after generating the single assembly instruction, wherein the execution granularity indicates the number of registers that have been continuously traversed.

[0021] For example, according to at least one embodiment of this disclosure, loading a matrix from memory into a register based on a traversal process includes: loading the matrix from memory into a register by executing one or more assembly instructions generated by the traversal process.

[0022] For example, according to at least one embodiment of this disclosure, the method further includes: encapsulating the traversal process as an intermediate representation, and generating one or more assembly instructions by calling the intermediate representation.

[0023] According to at least one embodiment of this disclosure, an apparatus for loading a matrix is ​​provided. The apparatus includes: a receiving module configured to receive a logical layout of an input tensor, a logical layout of an output tensor, and a physical layout type of an output register matrix specified by a user; a determining module configured to determine physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix; and an execution module configured to perform a traversal process for the output register matrix based on the physical layout information, and to load the matrix from memory into registers based on the traversal process. The traversal process includes iteratively performing the following operations for each register in the output register matrix until all registers in the output register matrix have been traversed: calculating the current register address of the current register based on the physical layout information; determining the logical coordinates of the first element of the current output tensor corresponding to the current register based on the current register address and the physical layout information; determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on the logical coordinates of the first element of the current output tensor; and calculating the current input memory address corresponding to the logical coordinates of the first element of the current input tensor based on the logical coordinates of the first element of the current input tensor.

[0024] According to at least one embodiment of the present disclosure, an electronic device for matrix loading is provided, the electronic device comprising: one or more processors; and a memory storing one or more computer programs; wherein, when executed by the one or more processors, the one or more computer programs cause the one or more processors to implement the method according to any one of the at least one embodiment of the present disclosure.

[0025] According to at least one embodiment of the present disclosure, one or more non-transitory computer-readable storage media are provided, having stored thereon computer-executable instructions that, when executed by a processor, cause at least one processor to implement the method according to any one of the at least one embodiments of the present disclosure.

[0026] The method, apparatus, electronic device, and storage medium for matrix loading provided in at least one embodiment of this disclosure achieve a more granular mapping relationship between register addresses and input memory addresses by determining the corresponding input memory address based on the address of each register. This supports data loading of high-dimensional tensors and non-contiguous memory / registers, while automatically and flexibly configuring instruction masks and supporting complex semantic scenarios such as broadcasting. Attached Figure Description

[0027] To more clearly illustrate the technical solutions of the embodiments of this disclosure, the accompanying drawings of the embodiments will be briefly described below. Obviously, the drawings described below only relate to some embodiments of this disclosure and are not intended to limit this disclosure.

[0028] Figure 1 A schematic structural diagram of a general-purpose graphics processor (GPGPU) according to at least one embodiment of the present disclosure is shown;

[0029] Figure 2 A flow diagram of a method for matrix loading according to at least one embodiment of the present disclosure is shown;

[0030] Figure 3A Physical layout information of the output register matrix according to at least one embodiment of the present disclosure is shown;

[0031] Figure 3B Physical layout information of the output register matrix according to another embodiment of this disclosure is shown;

[0032] Figure 4 A flow diagram of a method for matrix loading according to another embodiment of this disclosure is shown;

[0033] Figure 5 At least one embodiment according to the present disclosure is shown. Figure 4 The traversal process in the middle;

[0034] Figure 6 Another embodiment according to this disclosure is shown. Figure 4 The traversal process in the middle;

[0035] Figure 7 Another embodiment according to this disclosure is shown. Figure 4 The traversal process in the middle;

[0036] Figure 8 A flow diagram of a user invoking a method for loading an intermediate representation of a matrix is ​​shown according to at least one embodiment of the present disclosure;

[0037] Figure 9 An apparatus for matrix loading according to at least one embodiment of the present disclosure is shown;

[0038] Figure 10 An electronic device according to at least one embodiment of the present disclosure is shown;

[0039] Figure 11 A non-transitory computer-readable storage medium according to at least one embodiment of the present disclosure is shown. Detailed Implementation

[0040] The following description, with reference to the accompanying drawings, is provided to aid in a thorough understanding of the various embodiments of this disclosure as defined by the claims and their equivalents. This description includes various specific details to aid understanding but should be considered exemplary only. Therefore, those skilled in the art will recognize that various changes and modifications can be made to the various embodiments described herein without departing from the scope and spirit of this disclosure. Furthermore, for clarity and brevity, descriptions of well-known functions and structures may be omitted.

[0041] The terms and wording used in the following description and claims are not limited to their dictionary meanings, but are merely intended by the discloser to enable a clear and consistent understanding of this disclosure. Therefore, it will be apparent to those skilled in the art that the following description of various embodiments of this disclosure is for illustrative purposes only and not for limiting the purpose of this disclosure as defined in the appended claims and their equivalents.

[0042] It should be understood that the singular forms of “one,” “an,” and “the” include plural references unless the context clearly indicates otherwise. Thus, for example, the reference to “component surface” includes one or more such surfaces.

[0043] The terms “comprising” or “may include” refer to the presence of a corresponding disclosed function, operation, or component that may be used in the various embodiments of this disclosure, rather than limiting the presence of one or more additional functions, operations, or features. Furthermore, the terms “comprising” or “having” may be interpreted as indicating certain characteristics, numbers, steps, operations, constituent elements, components, or combinations thereof, but should not be construed as excluding the possibility of the presence of one or more other characteristics, numbers, steps, operations, constituent elements, components, or combinations thereof.

[0044] The term "or" as used in the various embodiments of this disclosure includes any of the listed terms and all combinations thereof. For example, "A or B" may include A, may include B, or may include both A and B.

[0045] Unless otherwise defined, all terms used in this disclosure (including technical or scientific terms) have the same meaning as understood by one of those skilled in the art as described herein. Common terms as defined in dictionaries are to be interpreted as having a meaning consistent with the context in the relevant technical field and should not be interpreted ideally or overly formally unless expressly defined in this disclosure.

[0046] Throughout the embodiments described herein, the same reference numerals consistently denote the same elements. In this respect, the embodiments may take different forms and should not be construed as limited to the description set forth herein. Therefore, the embodiments are described below only with reference to the accompanying drawings to explain various aspects. As used herein, the term "and / or" includes any and all combinations of one or more of the associated listed items. Expressions such as "at least one of..." modify the entire list of elements when preceding it, but do not modify individual elements in the list. In the drawings, the same reference numerals denote the same elements, and for clarity and ease of explanation, the dimensions of components in the drawings may be exaggerated.

[0047] While terms such as "first" and "second" can be used to describe various components, these terms are only used to distinguish one component from another. These terms are not intended to specify that the components are made of different materials or have different structures.

[0048] A singular expression may encompass a plural expression unless it has a distinctly different meaning in the context. When a part "includes" a component, it may further include another component, rather than excluding the other component, unless otherwise stated.

[0049] Furthermore, the terms "part," "module," etc., refer to a unit that processes at least one function or operation, and can be implemented by hardware, software, or a combination thereof.

[0050] In the development of GPGPU operators, to reduce development costs and efficiently generate operators, assembler code generators or compilers typically encapsulate commonly used instructions in operators into intermediate representations (IRs) and provide interfaces for developers to call. This allows developers to implement operators with different functions through combinations of IRs, without having to insert assembly instructions line by line. Matrix loading (ldmatrix) intermediate representations are one such example.

[0051] This disclosure provides a method, apparatus, electronic device, and storage medium for matrix loading. The matrix loading method according to at least one embodiment of this disclosure can directly support loading high-dimensional tensors and non-contiguous memory / register data, automatically and flexibly configure instruction masks, and support complex semantic scenarios such as broadcasting. The matrix loading method according to this disclosure can be encapsulated as a matrix loading IR for developers to call, providing assembly language operator generators (codegens) and their users with an efficient intermediate representation of matrix loading based on tensor object logical layout. This reduces development difficulty and error probability while providing performance-optimized local implementations for operators, helping to improve development efficiency and operator performance.

[0052] The following is for reference. Figures 1-10 Methods, apparatus, electronic devices, and storage media for matrix loading according to at least one embodiment of the present disclosure are described.

[0053] Figure 1 A schematic structural diagram of a general-purpose graphics processor (GPGPU) according to at least one embodiment of the present disclosure is shown.

[0054] like Figure 1 As shown, a general-purpose graphics processor is actually an array of programmable multiprocessors. For example, a programmable multiprocessor can be a streaming processor cluster (SPC), such as including... Figure 1 The diagram shows streaming processor clusters 1, ..., M, where M is a positive integer greater than 1. In a general-purpose graphics processor, one streaming processor cluster handles one computational task, or multiple streaming processor clusters handle one computational task. Multiple streaming processor clusters share data through a global cache or global memory.

[0055] like Figure 1 As shown, taking streaming processor cluster 1 as an example, one streaming processor cluster includes multiple computing units, such as... Figure 1 The system is structured as Computation Unit 1, Computation Unit 2, ..., Computation Unit N, where N is a positive integer. Each Computation Unit (CU) performs arithmetic and logical operations, such as accumulation, reduction, and standard addition, subtraction, multiplication, and division. A Computation Unit includes multiple cores (also called computational kernels), each of which includes an Arithmetic Logic Unit (ALU), a floating-point unit, etc. These cores are used to execute specific computational tasks. Furthermore, the Computation Unit also includes registers (e.g., ...). Figure 1 The register file and shared memory in a computing unit are used to store source and destination data related to computing tasks in a hierarchical manner. The shared memory in a computing unit is used to share data between the cores of that computing unit.

[0056] like Figure 1 As shown, each computing unit also provides a tensor core for performing tensor-related computations, such as tensor shrinking operations. Tensor cores can accelerate tensor operations such as matrix multiplication. Tensor cores in multiple computing units can be scheduled and controlled uniformly.

[0057] like Figure 1 As shown, each streaming processor cluster also provides a buffer for caching data across the N computing units within the streaming processor cluster.

[0058] In parallel computing, computational tasks are typically executed by multiple threads. These threads are divided into multiple thread blocks before execution in a general-purpose graphics processor (or parallel computing processor), and then dispatched via a thread block distribution module. Figure 1 (Not shown in the image) Multiple thread blocks are distributed to various computation units. All threads in a thread block must be assigned to the same computation unit for execution. Simultaneously, thread blocks are broken down into minimum execution thread bundles (or simply warps), each containing a fixed number (or less than this fixed number) of threads, for example, 32 threads. Multiple thread blocks can execute in the same computation unit or in different computation units.

[0059] In each computing unit, the thread beam scheduling / distribution module ( Figure 1 (Not shown in the diagram) Thread bundles are scheduled and allocated so that multiple computing cores within the computing unit can run thread bundles. Depending on the number of computing cores in the computing unit, multiple thread bundles within a thread block can execute concurrently or in a time-sharing manner. Multiple threads within each thread bundle execute the same instructions. Memory-executed instructions are issued to shared memory within the computing unit or further issued to intermediate-level caches, global caches, or global memory (e.g., [example cache]). Figure 1 High Bandwidth Memory (HBM) is used for read and write operations.

[0060] Figure 2 A flow 100 of a method for matrix loading according to at least one embodiment of the present disclosure is shown. Flow 100 includes steps 102-118.

[0061] like Figure 2 As shown, in step 102, the logical layout of the input tensor, the logical layout of the output tensor, and the physical layout type of the output register matrix are received.

[0062] For example, according to at least one embodiment of this disclosure, the logical layout of the input tensor, the logical layout of the output tensor, and the physical layout type of the output register matrix can be received by the user. For example, the user can define an input tensor object to declare the logical layout of the input tensor, and the user can define an output tensor object to declare the logical layout of the output tensor and the physical layout type of the output register matrix. The logical layout of the input tensor and the logical layout of the output tensor respectively indicate the logical layout of the input tensor and the output tensor. According to at least one embodiment of this disclosure, the logical layout of the input tensor and the logical layout of the output tensor can be provided in the form of (shape, stride) or shape : stride, where shape can refer to the dimension vector of the tensor's logical layout, indicating the dimensional information of the logical layout, and stride can refer to the step vector of the tensor's logical layout, indicating the step information of the logical layout.

[0063] For example, according to at least one embodiment of this disclosure, when a tensor has a three-dimensional layout, its logical layout can be represented as ((shape1, shape2, shape3), (stride1, stride2, stride3)) or (shape1, shape2, shape3):(stride1, stride2, stride3). In this example, the vector (shape1, shape2, shape3) is the dimension vector shape of the logical layout, and (stride1, stride2, stride3) is the step vector stride of the logical layout. Here, shape1, shape2, and shape3 respectively indicate the size of each dimension in the logical layout, and stride1, stride2, and stride3 respectively indicate the step size of each dimension in the logical layout. For example, a logical layout can be represented as ((2, 2, 4), (1, 2, 4)), where the shape of the dimension vector of the logical layout is (2, 2, 4), which indicates that the size of each dimension in the three-dimensional logical layout is 2, 2, 4 respectively; the stride of the logical layout is (1, 2, 4), which indicates that the stride size of each dimension in the three-dimensional logical layout is 1, 2, 4.

[0064] According to at least one embodiment of this disclosure, the stride vector of the logical layout indicates how data elements are extracted from linear storage units (e.g., memory / registers) to construct a tensor with a specific logical layout.

[0065] For example, in the example of a tensor with a three-dimensional logical layout of ((2, 2, 4), (1, 2, 4)), the dimension vector shape is (2, 2, 4), which indicates that the size of the first dimension (e.g., dimension 0) of the tensor's three-dimensional logical layout is 2, the size of the second dimension (e.g., dimension 1) is 2, and the size of the third dimension (e.g., dimension 2) is 4. The stride vector is (1, 2, 4), which indicates that in the first dimension (e.g., dimension 0) of the tensor, the index difference between two adjacent elements in the linear storage unit is 1, that is, one element is extracted from the linear storage unit every 1-1=0 elements; in the second dimension (e.g., dimension 1), the index difference between two adjacent elements in the linear storage unit is 2, that is, one element is extracted from the linear storage unit every 2-1=1 elements; and in the third dimension (e.g., dimension 2), the index difference between two adjacent elements in the linear storage unit is 4, that is, one element is extracted from the linear storage unit every 4-1=3 elements. Therefore, in the example where the tensor has a three-dimensional logical layout ((2,2,4), (1, 2, 4)), the three-dimensional logical layout ((2, 2, 4), (1, 2, 4)) indicates that data elements are extracted sequentially from linear storage cells to construct the tensor.

[0066] Alternatively, according to at least one embodiment of this disclosure, the logical layout (shape, stride) can indicate the non-contiguous extraction of data elements from linear storage units to construct a tensor. For example, the logical layout of a three-dimensional tensor could be ((2,2, 4), (1, 2, 8)), where the dimension vector shape is (2, 2, 4), indicating that the first dimension (e.g., dimension 0) of the logical layout of the three-dimensional tensor has a size of 2, the second dimension (e.g., dimension 1) has a size of 2, and the third dimension (e.g., dimension 2) has a size of 4. The stride is (1, 2, 8), indicating that in the first dimension (e.g., dimension 0) of the tensor, the index difference between two adjacent elements in the linear storage unit is 1, meaning that an element is extracted from every 1-1=0 elements in the linear storage unit; in the second dimension (e.g., dimension 1), the index difference between two adjacent elements in the linear storage unit is 2, meaning that an element is extracted from every 2-1=1 elements in the linear storage unit; and in the third dimension (e.g., dimension 2), the index difference between two adjacent elements in the linear storage unit is 8, meaning that an element is extracted from every 8-1=7 elements in the linear storage unit. Therefore, in the example where the tensor has a three-dimensional logical layout ((2, 2, 4), (1, 2, 8)), the three-dimensional logical layout ((2, 2, 4), (1, 2, 8)) indicates that data elements are extracted non-contiguously from the linear storage units to construct the tensor. The above embodiments are described using a tensor with a three-dimensional layout as an example, but this disclosure is not limited thereto. According to at least one embodiment of the disclosure, the tensor can have a layout of any dimension.

[0067] According to at least one embodiment of this disclosure, the physical layout type of the output register matrix can be different for different hardware or different modes of the same hardware ldmatrix instruction.

[0068] For example, according to at least one embodiment of this disclosure, the physical layout type of the output register matrix corresponds to the hardware architecture or instruction pattern. The physical layout type determines the physical organization of threads and registers during instruction execution, and also determines which register and which thread each data point in the output tensor will be loaded into. That is, there is a one-to-one mapping relationship between logical layout and physical layout in the registers; different register physical layout types result in different logical-to-physical mapping relationships.

[0069] According to at least one embodiment of this disclosure, for example, matrix operations and convolution operations can correspond to different hardware structures, the physical layout type of the output register matrix for subsequent matrix loading can correspond to the hardware structure of the matrix operation, and the physical layout type of the output register matrix for subsequent convolution operation can correspond to the hardware structure of the convolution operation.

[0070] For example, in the first instruction mode of matrix loading, a matrix with 512 elements can be arranged as a tensor of (64, 8). In this case, the physical layout type of the output register matrix can be the physical layout type of the first output register matrix corresponding to the first instruction mode. Alternatively, in the second instruction mode of matrix loading, a matrix with 512 elements can be arranged as a tensor of (16, 8, 4). In this case, the physical layout type of the output register matrix can be the physical layout type of the second output register matrix corresponding to the second instruction mode.

[0071] In step 104, the method includes determining the physical layout information of the output register matrix.

[0072] For example, according to at least one embodiment of this disclosure, the physical layout information of the output register matrix can be determined based on the logical layout of the output tensor and the physical layout type of the output register matrix.

[0073] For example, according to at least one embodiment of this disclosure, after a user-defined output tensor object declares the logical layout of the output tensor and the physical layout type of the output register matrix, the output tensor object not only stores these two pieces of information but also automatically derives the physical layout of the output register matrix based on these two pieces of information. The physical layout information of the output register matrix may include the physical layout of the output register matrix and a one-to-one first mapping relationship between the logical layout of the output tensor and the physical layout of the output register matrix. Furthermore, the physical layout of the output register matrix may include a register index layout and a thread layout.

[0074] According to at least one embodiment of this disclosure, the physical layout of the output register matrix, including the register index layout and thread layout, can have a form similar to the logical layout of the input tensor and the logical layout of the output tensor. For example, the register index layout and thread layout can each include a dimension vector and a step vector, respectively. For instance, the register index layout can be ((2, 2, 4), (1*128, 2*128, 4*128)), and the thread layout can be ((4, 8, 4), (1, 4, 32)). Similarly, the register index layout and thread layout can indicate whether data elements are extracted sequentially from linear memory units (registers) to construct the tensor or non-sequentially from linear memory units (registers) to construct the tensor.

[0075] The following will refer to Figure 3A and Figure 3B A detailed description of physical layout information according to at least one embodiment of this disclosure.

[0076] Back Figure 2 In step 106, the method may include performing traversal initialization.

[0077] According to at least one embodiment of the present disclosure, a traversal process of the output register matrix can be performed based on physical layout information. This traversal process iteratively performs multiple operations on each register in the output register matrix until all registers in the output register matrix have been traversed.

[0078] For example, in step 106, traversal initialization can be performed to execute the traversal process starting from the first register of the output register matrix. According to at least one embodiment of this disclosure, an iterator can be set, and traversal initialization may include setting the iterator to point to the first register of the output register matrix and setting the value i of the iterator to 0. According to at least one embodiment of this disclosure, the value i of the iterator can be incremented with each register traversal; for example, the value of the iterator can be incremented by 1 after each register traversal is completed. Optionally, according to at least one embodiment of this disclosure, before executing the traversal process, a memory base address in memory for the matrix and a register base address in the registers of the matrix can be allocated.

[0079] After the traversal initialization is completed, a traversal process according to at least one embodiment of the present disclosure can be executed, the traversal process including steps 108-116, and the traversal process according to at least one embodiment of the present disclosure iteratively executes steps 108-116.

[0080] In step 108, the current register address of the current register can be calculated.

[0081] For example, the current register address of the current register can be calculated based on the determined physical layout information. For example, the current register address can be calculated based on the register index layout in the physical layout information.

[0082] In step 110, the logical coordinates of the first element of the current output tensor corresponding to the current register can be determined.

[0083] For example, according to at least one embodiment of this disclosure, the first element logical coordinates of the current output tensor corresponding to the current register can be determined based on the current register address and physical layout information.

[0084] According to at least one embodiment of this disclosure, the physical layout information of the output register matrix may include a one-to-one first mapping relationship between the logical layout of the output tensor and the physical layout of the output register matrix, and the first element logical coordinates of the current output tensor corresponding to the current register may be determined based on the first mapping relationship.

[0085] For example, according to at least one embodiment of this disclosure, the logical coordinates of the first element of the current output tensor corresponding to the current register can be determined based on the calculated current register address, starting from the first thread of the current register, according to the first mapping relationship.

[0086] In step 112, the logical coordinates of the first element of the current input tensor corresponding to the current output tensor can be determined.

[0087] According to at least one embodiment of this disclosure, the logical coordinates of the first element of the current input tensor corresponding to the current output tensor can be determined based on the logical coordinates of the first element of the current output tensor. For example, the logical coordinates of the first element of the current input tensor corresponding to the current output tensor can be determined based on a second mapping relationship between the input tensor and the output tensor.

[0088] According to at least one embodiment of the present disclosure, the second mapping relationship may be based on the semantic relationship between the input tensor and the output tensor, and the semantic relationship may include, for example, point-to-point data loading or broadcast data loading, but the present disclosure is not limited thereto.

[0089] In step 114, the current input memory address corresponding to the logical coordinates of the first element of the current input tensor can be calculated.

[0090] For example, according to at least one embodiment of this disclosure, the current input memory address corresponding to the logical coordinates of the first element of the current input tensor can be calculated based on the logical coordinates of the first element of the current input tensor.

[0091] After steps 108-114 in the traversal process, the current register address and the corresponding input memory address can be obtained according to the method of at least one embodiment of the present disclosure. Then, based on the current register address and the input memory address, data elements can be loaded from the input memory address to the current register address.

[0092] According to at least one embodiment of this disclosure, a single assembly instruction can be generated based on the current register address and the corresponding input memory address during each traversal, and matrix loading can be performed based on one or more assembly instructions from the traversal process after the traversal is completed. Alternatively, according to at least one embodiment of this disclosure, the current register address and the corresponding input memory address obtained in each traversal can be stored, and one or more assembly instructions can be generated based on the stored multiple register address-input memory address pairs after the traversal is completed for matrix loading. However, this disclosure is not limited thereto.

[0093] In step 116, the traversal process checks whether all registers have been traversed.

[0094] For example, the value i of the iterator can be compared to see if it is less than the total number of registers N in the output register matrix. If the value i of the iterator is less than N, it is determined that the traversal process has not ended and the process returns to step 108 to continue the traversal process. Otherwise, the process proceeds to step 118.

[0095] In step 118, the matrix is ​​loaded from memory into a register.

[0096] According to at least one embodiment of this disclosure, a matrix can be loaded from memory into a register based on a traversal process. For example, a matrix can be loaded from memory into a register based on one or more assembly instructions generated by the traversal process. Alternatively, a matrix can be loaded from memory into a register based on a register address stored in the traversal process and the corresponding input memory address. However, this disclosure is not limited thereto.

[0097] exist Figure 2 In the illustrated process 100, step 118 is shown to be executed after the traversal of all registers is completed, but embodiments of this disclosure are not limited thereto. For example, step 118 can be a step in the traversal process, that is, after traversing each register, the data elements in the matrix can be loaded from the input memory address to the current register address based on the current register address and the corresponding input memory address.

[0098] In such Figure 2 In the method shown according to at least one embodiment of the present disclosure, by using a high-dimensional logical layout of tensors to construct the input tensor and output tensor and to traverse the output register matrix to realize the mapping relationship between register addresses and input memory addresses in a finer-grained manner, data loading of non-contiguous memory / registers can be realized, and high-dimensional tensors can be directly supported for processing.

[0099] Figure 3A Physical layout information 300-1 of the output register matrix according to at least one embodiment of the present disclosure is shown. Figure 3A The illustrated embodiment shows a case comprising 16 registers, Register0-Register15, with each register corresponding to 128 threads, T0-T127.

[0100] like Figure 3A As shown, the physical layout information 300-1 may include the physical layout 302 of the output register matrix and the one-to-one mapping relationship between the output tensor logic layout 301 and the physical layout 302. The output tensor logic layout 301 may be (64, 32):(1, 64).

[0101] like Figure 3A As shown, physical layout 302 may include register index layout 303 and thread layout 304. The register index layout may be (2, 2, 4):(1*128, 2*128, 2*2*128), and the thread layout 304 may be (4,8, 4):(1, 4, 32).

[0102] exist Figure 3AIn the illustrated embodiment, register index layout 303 indicates 16 registers arranged in a layout of (2, 2, 4):(1*128, 2*128, 2*2*128), and thread layout 304 indicates 128 threads in each register arranged in a layout of (4, 8, 4):(1, 4, 32). Since each register has 128 threads T0-T127, register index layout 303 (2, 2, 4):(1*128, 2*128, 2*2*128) indicates 16 consecutive registers Register0-Register15. Similarly, thread layout 304 (4, 8, 4):(1, 4, 32) indicates consecutive threads T0-T127.

[0103] like Figure 3A As shown, the physical layout information 300-1 also includes a one-to-one mapping relationship (e.g., a first mapping relationship) between the physical layout of the output tensor logic layout 301 and the physical layout of the output register matrix 302. Figure 3A In one embodiment, thread layout size * register index layout size = output tensor logic layout size, i.e., (2*2*4=16)*(4*4*8=128)=64*32=2048. According to at least one embodiment of this disclosure, each thread in each register of the physical layout 302 can have a one-to-one mapping relationship with an element in the output tensor logic layout 301. For example, the first element a1 of the output tensor logic layout 301 can have a one-to-one mapping relationship with thread T0 in the first register Register0, the second element a2 of the output tensor logic layout 301 can have a one-to-one mapping relationship with thread T1 in the first register Register0, and so on.

[0104] exist Figure 3A In the illustrated embodiment, physical layout 302 indicates 16 consecutive registers and 128 consecutive threads; however, according to the present disclosure, physical layout 302 can indicate 16 non-consecutive registers and 128 non-consecutive threads. The following is in conjunction with... Figure 3B This example describes a physical layout indicating 16 non-contiguous registers and 128 non-contiguous threads. For simplicity, details related to... Figure 3A The embodiments shown are described in the same or similar manner.

[0105] Figure 3B Physical layout information 300-2 of the output register matrix according to another embodiment of the present disclosure is shown.

[0106] exist Figure 3BIn the illustrated embodiment, the total number of registers is 32 (i.e., Register0-Register31) (not all shown) and the total number of threads is 256 (i.e., T0-T255) (not all shown), while the output register matrix includes 16 registers (i.e., Figure 3B The Register0, Register2...Register30 shown in the diagram) each have 128 threads (i.e., Figure 3B (as shown in T0, T2...T254).

[0107] like Figure 3B As shown, the physical layout information 300-2 may include the physical layout 302 of the output register matrix and the one-to-one mapping relationship between the output tensor logic layout 301 and the physical layout 302. The output tensor logic layout 301 may be (64, 32):(1, 64).

[0108] like Figure 3B As shown, physical layout 302 may include register index layout 303 and thread layout 304, wherein the register index layout may be (2, 2, 4):(2*256, 4*256, 2*4*256), and thread layout 304 may be (4, 8, 4):(2, 8, 64). Figure 3B In the illustrated embodiment, the step vector for the register index layout is (2*256, 4*256, 2*4*256) instead of Figure 3A The diagram shows (1*128, 2*128, 2*2*128), and the thread layout's step vector is (2, 8, 64) instead of... Figure 3A As shown in the diagram (1, 4, 32). That is to say, Figure 3B Register index layout 303 indicates non-contiguous registers, and thread layout 304 indicates non-contiguous threads.

[0109] exist Figure 3BIn the illustrated embodiment, register index layout 303 indicates 16 registers arranged in a layout of (2, 2, 4): (2*256, 4*256, 2*4*256), and thread layout 304 indicates 128 threads in each register arranged in a layout of (4, 8, 4): (2, 8, 64). Since each register has 256 threads, register index layout 303 (2, 2, 4): (2*256, 4*256, 2*4*256) indicates 16 non-contiguous registers: Register0, Register2, Register4…Register30. Similarly, thread layout 304 (4, 8, 4): (1, 4, 32) indicates 128 non-contiguous threads: T0, T2, T4…T254.

[0110] like Figure 3B As shown, the physical layout information 300-2 also includes a one-to-one mapping relationship (e.g., a first mapping relationship) between the physical layout of the output tensor logic layout 301 and the physical layout of the output register matrix 302. Figure 3B In one embodiment, thread layout size * register index layout size = output tensor logic layout size, i.e., (2*2*4=16)*(4*4*8=128)=64*32=2048. According to at least one embodiment of this disclosure, each thread in each register of the physical layout 302 can have a one-to-one mapping relationship with an element in the output tensor logic layout 301. For example, the first element a1 of the output tensor logic layout 301 can have a one-to-one mapping relationship with thread T0 in the first register Register0, the second element a2 of the output tensor logic layout 301 can have a one-to-one mapping relationship with thread T2 in the first register Register0, and so on.

[0111] Figure 4 A flow 200 of a method for matrix loading according to another embodiment of the present disclosure is shown. Flow 200 includes steps 202-212.

[0112] like Figure 4 As shown, steps 202, 204, and 212 of process 200 are similar to those in... Figure 2 Steps 102, 104, and 118 of the process 100 shown are similar. For the sake of brevity, the descriptions of steps 202, 204, and 212 are omitted.

[0113] like Figure 4 As shown, a method for matrix loading according to at least one embodiment of this disclosure may include the following steps:

[0114] Step 206: Allocate the memory base address and register base address.

[0115] For example, according to at least one embodiment of this disclosure, before performing the traversal process, a memory base address in memory for the matrix and a register base address in the registers for the matrix are allocated. For example, as... Figure 4 As shown, after receiving the logical layout of the input tensor, the logical layout of the output tensor, and the physical layout type of the output register matrix, and determining the physical layout information of the output register matrix, the memory base address of the input tensor in memory and the register base address of the output tensor in the register can be allocated before the initialization of the traversal process. For example, according to at least one embodiment of this disclosure, the memory base address and the register base address can be stored in user-defined input tensor objects and output tensor objects, respectively. For example, the memory base address and register base address can be allocated by allocating hardware resources through an assembler.

[0116] In step 208, the initialization of the traversal process can be performed.

[0117] For example, such as Figure 4 As shown, an iterator can be set to perform a traversal process on each register in the output register matrix. According to at least one embodiment of this disclosure, during traversal initialization, an iterator can be set to point to the first register in the output register matrix, and the value of the iterator i can be set to i=0. During the traversal, after each traversal, the value of the iterator i can be incremented by 1, and the value of the iterator i can be compared with the total number of registers N. If i is less than N, the traversal process continues; otherwise, the traversal process ends. Setting an iterator to perform a traversal process on each register in the output register matrix is ​​merely an example, and embodiments of this disclosure are not limited thereto.

[0118] According to at least one embodiment of this disclosure, the output register matrix can be traversed in various ways. For example, pointers to registers in the output register matrix can be set. During traversal initialization, the pointers are set to point to the first register in the output register matrix. After each traversal, it is determined whether the pointers have already pointed to the last register in the output register matrix. If the pointers are not pointing to the last register in the output register matrix, the pointers can be set to point to the next register and the traversal process can continue. Otherwise, the traversal process ends.

[0119] In step 210, the traversal process can be performed.

[0120] The following will refer to Figures 5-7 A detailed example of a traversal process according to at least one embodiment of the present disclosure is provided.

[0121] Figure 5 At least one embodiment according to the present disclosure is shown. Figure 4 The traversal process in 2100.

[0122] Figure 5 The traversal process 2100 can continue Figure 4 Step 208 is executed, that is... Figure 5 The traversal process 2100 shown can be Figure 4 The step 210 is shown. The traversal process 2100 includes steps 2102-2112.

[0123] In the method for loading a matrix according to at least one disclosed embodiment, a traversal process for the output register matrix can be performed based on physical layout information. For example, the traversal process for the output register matrix can be performed based on the dimension vector and step vector of the register index layout included in the physical layout.

[0124] According to at least one embodiment of the present disclosure, performing a traversal of an output register matrix based on the dimension vector and step vector of the register index layout may include performing a traversal of the registers in the output register matrix in ascending order of step size for each dimension based on the dimension coordinates of the registers in the register index layout.

[0125] For example, regarding such Figure 3A The physical layout information 300-1 shown can be used to perform a traversal of the output register matrix based on the dimension vector (2,2,4) and step vector (1*128, 2*128, 2*2*128) of the register index layout 303. For example, the registers in the output register matrix can be traversed sequentially from the first dimension to the third dimension based on the dimensional coordinates (three-dimensional coordinates) of the registers in the register index layout 303, in ascending order of the step size (1*128, 2*128, 2*2*128) for each dimension.

[0126] like Figure 5 As shown, in step 2102, the current register address of the current register can be calculated.

[0127] According to at least one embodiment of this disclosure, calculating the current register address of the current register based on physical layout information may include calculation based on the following formula:

[0128] Current register address = register base address + current coordinate * (indexStride / threadNumber) * demsize / regSize

[0129] Where “·” symbol represents dot product, current coordinate represents the dimension coordinate of the current register, indexStride represents the step vector, threadNumber represents the total number of executing threads, demsize represents the size of a single element (in bytes), and regSize represents the size of a single register (in bytes).

[0130] For example, assuming the register base address is 0, the current coordinate is (1, 1, 1), the indexStride is (1*128, 2*128, 2*2*128), the threadNumber is 128, and demsize=regSize=16 bytes, then the current register address = 0 + (1, 1, 1)·((1*128, 2*128, 2*2*128) / 128)*16 / 16=7.

[0131] Continue to refer to Figure 5 In step 2104, the logical coordinates of the first element of the current output tensor corresponding to the current register can be determined.

[0132] For example, according to at least one embodiment of this disclosure, the first element logical coordinates of the current output tensor corresponding to the current register can be determined based on the current register address and physical layout information calculated in step 2102.

[0133] For example, determining the logical coordinates of the first element of the current output tensor corresponding to the current register based on the current register address and physical layout information calculated in step 2102 may include determining the logical coordinates of the first element of the current output tensor corresponding to the current register based on the current register address, starting from the first thread of the current register, and according to a first mapping relationship. According to at least one embodiment of this disclosure, the first mapping relationship may be a one-to-one mapping relationship between the logical layout of the output tensor and the physical layout included in the physical layout information. For example, using... Figure 3A Taking the physical layout information 300-1 as an example, the current register can be the first register Register0. Based on the register address of the first register Register0, the first element a1 of the current output tensor corresponding to the first register Register0 can be determined starting from the first thread T0 of the first register Register0. Thus, the logical coordinates of the first element of the current output tensor corresponding to the first register Register0 can be determined to be (0, 0).

[0134] refer to Figure 5 The traversal process 2100 continues to step 2106.

[0135] After determining the logical coordinates of the first element of the current output tensor corresponding to the current register, in step 2106, the logical coordinates of the first element of the current input tensor corresponding to the current output tensor can be determined.

[0136] According to at least one embodiment of this disclosure, the logical coordinates of the first element of the current input tensor corresponding to the current output tensor can be determined based on the logical coordinates of the first element of the current output tensor. For example, determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on the logical coordinates of the first element of the current output tensor may include determining the logical coordinates of the first element of the current input tensor corresponding to the current output tensor based on a second mapping relationship between the input tensor and the output tensor.

[0137] According to at least one embodiment of this disclosure, the second mapping relationship can be based on the semantic relationship between the input tensor and the output tensor, and the semantic relationship can include point-to-point data loading or broadcast data loading. For example, when the logical layout of the input tensor and the logical layout of the output tensor have the same dimension vector, that is, when performing point-wise data loading, the logical coordinate of the first element of the current input tensor equals the logical coordinate of the first element of the current output tensor; when a certain dimension axis of the logical layout of the input tensor needs to be broadcast on a certain dimension axis of the logical layout of the output tensor, then the logical coordinate of the element of the input tensor is set to 1 on that dimension axis, and the other dimensions are set to be the same as the logical coordinate of the element of the output tensor. For example, if the logical layout of the input tensor is (64, 1) and the logical layout of the output tensor is (64, 32), then when the logical coordinate of the element of the output tensor is (x, y), the logical coordinate of the corresponding element of the input tensor is (x, 1). The above semantic relationship is merely an embodiment, and the method according to at least one embodiment of this disclosure can be applied to other possible data loading semantic scenarios.

[0138] In step 2108, the traversal process 2100 may include calculating the current input memory address corresponding to the logical coordinates of the first element of the current input tensor. According to at least one embodiment of this disclosure, the current input memory address corresponding to the logical coordinates of the first element of the current input tensor determined in step 2106 can be calculated. For example, calculating the current input memory address corresponding to the logical coordinates of the first element of the current input tensor determined in step 2106 may include calculation based on the following formula:

[0139] Current input memory address = base memory address + current coordinate · logicStride * demsize

[0140] Where “·” symbol represents dot product, current coordinate represents the logical coordinate of the first element of the current input tensor, logicStride represents the step vector of the logical layout of the input tensor, and demsize represents the size of a single element (in bytes).

[0141] According to at least one embodiment of this disclosure, such as Figure 5 As shown, in step 2110, a single assembly instruction can be generated.

[0142] For example, a single assembly instruction can be generated based on the determined current register address and current input memory address. For example, an assembler can generate a single assembly instruction based on the determined current register address and current input memory address. Optionally, according to at least one embodiment of this disclosure, a single assembly instruction can be generated based on the current register address, the current input memory address, the output tensor and its correspondence, the physical layout type of the output register matrix, user-configured synchronization information, and other user-configured information. According to at least one embodiment of this disclosure, the user-configured synchronization information may only exist in the first and last instructions, wherein the first instruction may need to wait for data in memory to be ready and the register to be free, and the last instruction may need to send a message indicating that register data is ready and memory is free.

[0143] refer to Figure 5 After generating a single assembly instruction, in step 2112, the value of the iterator i can be set to i = i + 1. And in step 2114, the value of the iterator i can be compared with the total number of registers N.

[0144] If i is less than N, the iterator points to the next register and the traversal continues; otherwise, the traversal ends. After the traversal is complete, the following can be executed: Figure 4 Step 212 is shown.

[0145] Figure 6 Another embodiment according to this disclosure is shown. Figure 4 The traversal process in the middle is 2200. Figure 6 The traversal process 2200 can continue Figure 4 Step 208 is executed, that is... Figure 6 The traversal process 2200 shown can be Figure 4 The step 210 is shown. The traversal process 2200 includes steps 2202-2216. Steps 2202-2204, 2208-2210, and 2214-2216 are respectively related to... Figure 5Steps 2102-2104, 2106-2108, and 2112-2114 shown are the same, and for the sake of brevity, the same descriptions are omitted.

[0146] and Figure 5 Compared to the traversal process shown in 2100, Figure 6 The traversal process 2200 shown further includes step 2206, in which a current thread mask can be determined according to at least one embodiment of the present disclosure. For example, according to at least one embodiment of the present disclosure, the logical coordinates of each element in the current output tensor corresponding to each thread can be determined based on the current register address and thread layout, according to a first mapping relationship. The current thread mask can be determined by comparing the logical coordinates of each element corresponding to each thread with the logical layout of the input tensor. The current thread mask is used to prohibit the execution of a portion of the threads. For example, based on the current register address and the thread layout included in the physical layout, the logical coordinates of the data elements in the output tensor corresponding to each thread can be obtained according to a one-to-one first mapping relationship between the physical layout and the logical layout of the output tensor. Whether an out-of-bounds error has occurred can be determined by comparing the logical coordinates of the data elements corresponding to each thread with the logical layout of the input tensor. When it is determined that the logical coordinates of the data elements in the output tensor corresponding to a certain thread exceed the logical layout of the input tensor, it is determined that the operation of that thread is prohibited. For example, if the logical layout of the input tensor is (64, 7) and the logical layout of the output tensor is (64, 8), and the logical coordinates of a data element in the output tensor corresponding to a certain thread are (50, 8), then since the logical coordinates (50, 8) exceed the logical layout of the input tensor, it can be determined that the data element is out of bounds. Therefore, the work of that thread can be disabled. The above method can be used to determine the current thread mask during the current traversal.

[0147] In step 2212, according to at least one embodiment of the present disclosure, a single assembly instruction can be generated. For example, a single assembly instruction can be generated based on the determined current register address and current input memory address. For example, an assembler can generate a single assembly instruction based on the determined current register address and current input memory address. Optionally, according to at least one embodiment of the present disclosure, a single assembly instruction can be generated based on the current register address, current input memory address, current thread mask, output tensor and output tensor correspondence, physical layout type of output register matrix, user-configured synchronization information, and other user-configured information. According to at least one embodiment of the present disclosure, the user-configured synchronization information may only exist in the first and last instructions, wherein the first instruction may need to wait for data in memory to be ready and the register to be free, and the last instruction may need to send a message indicating that register data is ready and memory is free.

[0148] Figure 7 Another embodiment according to this disclosure is shown. Figure 4 The traversal process in the middle is 2300.

[0149] Figure 7 The traversal process 2300 can continue. Figure 4 Step 208 is executed, that is... Figure 7 The traversal process 2300 shown can be Figure 4 The step 210 is shown. The traversal process 2300 includes steps 2302-2324. Steps 2302-2312 are respectively related to... Figure 6 Steps 2202-2210 and 2214 shown are the same, and for the sake of brevity, the same descriptions are omitted.

[0150] According to at least one embodiment of this disclosure, in many hardware architectures, a single instruction can be generated for multiple registers instead of just one. In this case, the method according to at least one embodiment of this disclosure can perform sequential traversal to sequentially traverse multiple registers in the output register matrix and generate a single assembly instruction for multiple registers. For example, after completing the traversal of the current register, a next traversal can be attempted to obtain the next register address, the next input memory address, and the next thread mask. In response to the current register address and the next register address being consecutive, the current input memory address and the next input memory address being consecutive, and the current thread mask being the same as the next thread mask, a next traversal can be attempted until the above conditions are not met, the upper limit of the number of registers that can be sequentially traversed supported by the hardware is reached, or all registers in the output register matrix have been traversed.

[0151] For example, a continuous traversal according to at least one embodiment of the present disclosure may include, after determining the current register address, the current input memory address, and the current thread mask, proceeding to the next traversal to determine the next register address, the next input memory address, and the next thread mask in the next traversal, comparing the current register address, the current input memory address, and the current thread mask with the next register address, the next input memory address, and the next thread mask, and in response to the current register address being consecutive with the next register address, the current input memory address being consecutive with the next input memory address, the current thread mask being the same as the next thread mask, the number of consecutively traversed registers not reaching the upper limit of the number of registers for continuous traversal supported by the hardware, and not having completed traversing all registers, continuing to continuously traverse the registers in the output register matrix; otherwise, generating a single assembly instruction based on the first register address of the first traversal in the continuous traversal, the first input memory address, the thread mask of the continuous traversal, the execution granularity, and the user configuration information, and after generating the single assembly instruction, continuing the traversal process, where the execution granularity indicates the number of consecutively traversed registers.

[0152] As Figure 7 shown, after completing the traversal of the current register, in step 2314, it is determined whether the value i of the iterator is less than N, where N represents the total number of registers in the output register matrix. If i is less than N, the traversal process 2300 proceeds to step 2316 to perform a continuous traversal of the next register.

[0153] In step 2316, it may be determined to perform a traversal of the next register to determine the next register address, the next input memory address, and the next thread mask in the next traversal. The traversal of the next register may be performed based on the method according to the corresponding embodiment of the present disclosure.

[0154] In step 2318, the current register address, the current input memory address, and the current thread mask may be compared with the next register address, the next input memory address, and the next thread mask.

[0155] In step 2320, based on the above comparison, it is determined whether the conditions for continuing to perform the continuous traversal are satisfied. For example, according to at least one embodiment of the present disclosure, the conditions for continuing to perform the continuous traversal may include: the current register address being consecutive with the next register address, the current input memory address being consecutive with the next input memory address, the current thread mask being the same as the next thread mask, n < T, and i = i + 1 < N. Where n represents the number of consecutively traversed registers, T represents the upper limit of the number of registers for continuous traversal supported by the hardware, i represents the value of the iterator, and N represents the total number of registers in the output register matrix.

[0156] In response to determining in step 2320 that the conditions for continuing the sequential traversal are met, for example, determining that the current register address is consecutive with the next register address, the current input memory address is consecutive with the next input memory address, the current thread mask is the same as the next thread mask, n < T, and i = i + 1 < N, the traversal process 2300 returns to step 2316 to continue the sequential traversal of the registers in the output register matrix; otherwise, the traversal process 2300 proceeds to step 2322.

[0157] According to at least one embodiment of the present disclosure, consecutive addresses may refer to logical consecutiveness rather than being limited to consecutive address values. For example, in an example where each register has 128 threads, 128 threads in one register execute simultaneously to process 128 data elements, and 128 data elements are loaded into one register. Therefore, the address values of two consecutive registers differ by 128, but the two registers can be considered logically consecutive. Thus, in an example where each register has 128 threads, the current register address being consecutive with the next register address means that the address value of the current register address differs from the address value of the next register address by 128. According to at least one embodiment of the present disclosure, step 2320 may include setting the value of the iterator i to i = i + 1 and determining whether i is less than N.

[0158] In step 2322, a single assembly may be generated for the multiple registers of the sequential traversal based on the sequential traversal.

[0159] For example, according to at least one embodiment of the present disclosure, a single assembly instruction may be generated based on the first register address of the first traversal during the sequential traversal, the first input memory address, the thread mask of the sequential traversal, and the execution granularity. The execution granularity indicates the number of registers that have been sequentially traversed. For example, when the traversal process sequentially traverses two registers and generates a single assembly instruction for the two sequentially traversed registers, a single assembly instruction may be generated based on the first register address in the two registers, the first input memory address corresponding to the first register address, the thread mask of the two sequentially traversed registers, and the execution granularity (i.e., 2). Optionally, according to at least one embodiment of the present disclosure, a single assembly instruction may be generated based on the first register address, the first input memory address, the thread mask of the sequential traversal, the execution granularity, the output tensor and output tensor correspondence, the physical layout type of the output register matrix, the synchronization information configured by the user, and other information configured by the user.

[0160] According to at least one embodiment of the present disclosure, the synchronization information configured by the user may only exist in the first and last instructions. Among them, the first instruction may need to wait for the data in the memory to be ready and the registers to be idle, and the last instruction may need to send the register data ready and the memory idle.

[0161] In step 2324, it can be determined whether the value i of the iterator is less than N. If i is less than N, the traversal process 2300 returns to step 2302 to continue the traversal process; if i is not less than N, the traversal process 2300 ends, and the method can continue to execute. Figure 4 Step 212 is shown in the figure.

[0162] The traversal process according to at least one embodiment of this disclosure can be encapsulated as an intermediate representation for matrix loading, and the user can generate one or more assembly instructions for matrix loading by calling the intermediate representation. By encapsulating the traversal process according to embodiments of this disclosure as an intermediate representation for user invocation, the intermediate representation can be directly invoked to handle high-dimensional tensor and non-contiguous memory / register data loading, automatically and flexibly configure instruction masks, and support complex semantic scenarios such as broadcasting. While reducing development difficulty and error probability, it provides performance-optimized local implementations for operators, helping to improve development efficiency and operator performance.

[0163] Figure 8 A flow 800 of a user invoking a method for loading an intermediate representation of a matrix is ​​shown according to at least one embodiment of the present disclosure.

[0164] In step 802, the user can declare input tensors and output tensors. For example, the user can declare input tensor objects and output tensor objects. According to at least one embodiment of this disclosure, the input tensor object can define the logical layout of the input tensor, and the output tensor object can define the logical layout of the output tensor and the physical layout type of the output register matrix. According to at least one embodiment of this disclosure, the user can declare input tensors and output tensors through a user application programming interface (API).

[0165] In step 804, matrix loading initialization can be performed. For example, initialization may include allocating memory base addresses and register base addresses, and determining the physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix. For example, according to at least one embodiment of this disclosure, after a user-defined output tensor object declares the logical layout of the output tensor and the physical layout type of the output register matrix, the physical layout of the output register matrix can be automatically derived based on these two pieces of information and stored in the output tensor object.

[0166] In step 806, an intermediate representation for matrix loading can be invoked. The intermediate representation for matrix loading can encapsulate a method for matrix loading according to at least one embodiment of this disclosure. For example, the intermediate representation for matrix loading can encapsulate a traversal process according to at least one embodiment of this disclosure for user invocation, thereby reducing code volume, lowering development difficulty and error probability, and improving development efficiency.

[0167] In step 808, the matrix can be loaded from memory into a register based on the assembly code generated from the intermediate representation.

[0168] exist Figure 8 In the illustrated embodiments, steps 804 and 808 are shown as separate steps, but this disclosure is not limited thereto. According to at least one embodiment of this disclosure, at least one of steps 804 and 808 may be encapsulated in an intermediate representation for matrix loading.

[0169] Figure 9 An apparatus 900 for matrix loading according to at least one embodiment of the present disclosure is shown.

[0170] Figure 9 The illustrated apparatus 900 can be configured to perform a method for matrix loading according to at least one embodiment of the present disclosure. For example... Figure 9 As shown, the device 900 includes a receiving module 902, a determining module 904, and an execution module 906. For example, these modules can be implemented using digital circuits, analog circuits, firmware (e.g., field-programmable gate arrays (FPGAs)), or any combination thereof, and the embodiments of this disclosure are not limited thereto.

[0171] According to at least one embodiment of this disclosure, for example, the receiving module 902 may be configured to receive the logical layout of the input tensor, the logical layout of the output tensor, and the physical layout type of the output register matrix specified by the user.

[0172] According to at least one embodiment of this disclosure, for example, the determining module 904 may be configured to determine the physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix.

[0173] According to at least one embodiment of this disclosure, for example, execution module 906 may be configured to perform a traversal process for the output register matrix based on physical layout information. The traversal process includes iteratively performing the following operations for each register in the output register matrix until all registers in the output register matrix have been traversed: calculating the current register address of the current register based on the physical layout information; determining the first element logical coordinate of the current output tensor corresponding to the current register based on the current register address and the physical layout information; determining the first element logical coordinate of the current input tensor corresponding to the current output tensor based on the first element logical coordinate of the current output tensor; and calculating the current input memory address corresponding to the first element logical coordinate of the current input tensor based on the first element logical coordinate of the current input tensor.

[0174] According to at least one embodiment of this disclosure, for example, the receiving module 902 may also be configured to load a matrix from memory into a register based on a traversal process.

[0175] Figure 10 An electronic device according to at least one embodiment of the present disclosure is shown.

[0176] The method and apparatus for matrix loading provided in at least one embodiment of this disclosure can be applied to different systems or devices, such as those used in... Figure 10 The electronic device shown is 1000.

[0177] The electronic device 1000 can be a terminal, such as a mobile phone, tablet computer, laptop computer, AR device, VR device, vehicle terminal, etc., or it can be a server. The matrix loading method provided in at least one embodiment of this disclosure can be applied to matrix loading scenarios in the electronic device 1000 involving high-performance computing (HPC) and artificial intelligence (AI), such as tensor computation units. Of course, this disclosure is not limited to this; any scenario, device, or apparatus involving matrix loading can employ the matrix loading method or apparatus provided in at least one embodiment of this disclosure.

[0178] In some embodiments, the apparatus for matrix loading provided in at least one embodiment of this disclosure can be a chip, such as a system-on-a-chip (SoC). The SoC includes a processor, which can be a single-core or multi-core processor, memory, and I / O interfaces, etc. The processor can load data and applications from memory and then process the data, such as performing matrix loading.

[0179] It should be noted that the input and output tensors of the matrix loading method and apparatus provided in at least one embodiment of this disclosure may have different physical meanings depending on the application scenario. For example, the matrix loading method provided in at least one embodiment of this disclosure can be applied in fields such as speech processing, image processing, text processing, and video processing.

[0180] For example, in the field of speech processing, input tensors and output tensors can be input and output parameters in tasks such as feature extraction, speech enhancement, and speech recognition.

[0181] For example, in the field of image processing, input tensors and output tensors can be the input and output parameters for tasks such as image recognition, feature extraction, image segmentation, object detection, image classification, and scene reconstruction.

[0182] For example, in the field of text processing, input tensors and output tensors can be the input and output parameters for tasks such as text classification, sentiment analysis, and text generation.

[0183] For example, in the field of video processing, input tensors and output tensors can be parameters related to image processing, or input and output parameters specific to the field of video processing, such as optical flow operators (used to estimate motion between video frames) and target tracking operators (used to track specific targets in video).

[0184] Of course, this disclosure is not limited to this. For other application scenarios or fields, as long as matrix loading is required, the matrix loading method described in at least one embodiment of this disclosure can be applied, and will not be elaborated here.

[0185] like Figure 10 As shown, the electronic device 1000 includes a processing unit 1001, which can perform various appropriate actions and processes according to non-transitory computer-readable instructions stored in a memory to achieve various functions.

[0186] For example, when the computer-readable instruction processing device 1001 is executed, it can perform one or more steps of the method for matrix loading according to any of the above embodiments. It should be noted that a detailed description of the process of the method for matrix loading can be found in the relevant descriptions in the embodiments of the matrix loading method described above.

[0187] For example, the memory may include any combination of one or more computer program products, which may include various forms of computer-readable storage media, such as volatile memory and / or non-volatile memory. Volatile memory may include, for example, random access memory (RAM) 1003 and / or cache memory, etc., whereby computer-readable instructions can be loaded from storage device 1008 into RAM 1003 to execute. Non-volatile memory may include, for example, read-only memory (ROM) 1002, hard disk, erasable programmable read-only memory (EPROM), portable compact disc read-only memory (CD-ROM), USB memory, flash memory, etc. Various applications and various data, such as style images, and various data used and / or generated by the applications, may also be stored in the computer-readable storage medium.

[0188] For example, the processing device 1001, ROM 1002, and RAM 1003 are interconnected via bus 1004. Input / output (I / O) interface 1005 is also connected to bus 1004.

[0189] Typically, the following devices can be connected to the I / O interface 1005: input devices 1006 including, for example, a touchscreen, touchpad, keyboard, mouse, camera, microphone, accelerometer, gyroscope, etc.; output devices 1007 including, for example, a liquid crystal display (LCD), speaker, vibrator, etc.; storage devices 1008 including, for example, magnetic tape, hard disk, flash memory, etc.; and communication devices 1009. Communication device 1009 allows electronic device 1000 to communicate wirelessly or wiredly with other electronic devices to exchange data. Although Figure 10 An electronic device 1000 with various devices is shown, but it should be understood that it is not required to implement or have all of the devices shown, and the electronic device 1000 may alternatively implement or have more or fewer devices. For example, the processing device 1001 may control other components in the electronic device 1000 to perform desired functions.

[0190] For example, the processing device 1001 can be a central processing unit (CPU), a tensor processor (TPU), or a graphics processing unit (GPU) with data processing and / or program execution capabilities. The CPU can be based on x86, ARM, RISC-V, or other architectures. The GPU can be directly integrated into the SoC, directly integrated onto the motherboard, or built into the northbridge chip of the motherboard.

[0191] Figure 11 A non-transitory computer-readable storage medium 1100 according to at least one embodiment of the present disclosure is shown.

[0192] For example, such as Figure 11 As shown, storage medium 1100 can be a non-transitory computer-readable storage medium on which one or more computer-readable instructions 1101 can be stored non-transitory. For example, when the computer-readable instructions 1101 are executed by a processor, one or more steps in the method for matrix loading described above can be performed.

[0193] For example, the storage medium 1100 can be used in an electronic device 1000, such as the storage medium 1100 may include the storage device 1008 in the electronic device 1000.

[0194] For example, a storage device may include any combination of one or more computer program products, which may include various forms of computer-readable storage media, such as volatile memory and / or non-volatile memory. Volatile memory may include, for example, random access memory (RAM) and / or cache memory. Non-volatile memory may include, for example, read-only memory (ROM), hard disk, erasable programmable read-only memory (EPROM), portable compact disc read-only memory (CD-ROM), USB storage, flash memory, etc. One or more computer-readable instructions may be stored on the computer-readable storage medium, and a processor may execute these instructions to perform various functions of the processor. Various application programs and various data may also be stored in the storage medium.

[0195] For example, the storage medium may include a memory card for a smartphone, a cache component for a tablet computer, a hard disk for a personal computer, random access memory (RAM), read-only memory (ROM), erasable programmable read-only memory (EPROM), portable compact disc read-only memory (CD-ROM), flash memory, or any combination of the above storage media, or other suitable storage media.

[0196] 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 a 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.

[0197] The units described in the embodiments of this disclosure can be implemented in software or hardware. The names of the units are not, in some cases, intended to limit the specific unit.

[0198] The functions described above in this document can be performed at least in part by one or more hardware logic components. For example, exemplary types of hardware logic components that can be used, without limitation, include: field-programmable gate arrays (FPGAs), application-specific integrated circuits (ASICs), application-specific standard products (ASSPs), system-on-a-chip (SoCs), complex programmable logic devices (CPLDs), and so on.

[0199] Those skilled in the art will understand that the illustrative embodiments described above are not intended to be limiting. It should be understood that any two or more of the embodiments disclosed herein can be combined in any combination. Furthermore, other embodiments may be utilized and other changes may be made without departing from the spirit and scope of the subject matter presented herein. It will be readily understood that aspects of this disclosure, as generally described herein and illustrated in the accompanying drawings, can be arranged, substituted, combined, separated, and designed in a variety of different configurations, all of which are contemplated herein.

[0200] Various embodiments of this disclosure can be implemented as computer-readable code embodied on a computer-readable recording medium from a particular perspective. A computer-readable recording medium is any data storage device capable of storing data readable by a computer system. Examples of computer-readable recording media may include read-only memory (ROM), random access memory (RAM), optical disc read-only memory (CD-ROM), magnetic tape, floppy disk, optical data storage device, carrier wave (e.g., data transmission via the Internet), and the like. Computer-readable recording media can be distributed via computer systems connected via a network, and thus computer-readable code can be stored and executed in a distributed manner. Furthermore, the functional programs, code, and code segments used to implement the various embodiments of this disclosure can be readily interpreted by those skilled in the art applying the embodiments of this disclosure.

[0201] The medium can be read by a computer, stored in a memory, and executed by a processor. Various embodiments can be implemented by a computer or a portable terminal including a controller and a memory, and the memory can be an example of a non-transitory computer-readable recording medium suitable for storing programs(s) having instructions for implementing embodiments of this disclosure. This disclosure can be implemented by a program having code for specifically implementing the apparatus and methods described in the claims, said program being stored in a machine (or computer) readable storage medium. The program can be carried electronically on any medium, such as communication signals transmitted via wired or wireless connections, and this disclosure suitably includes its equivalents.

[0202] The embodiments described herein are not intended to be limiting. The aspects of this disclosure, as generally described herein and shown in the accompanying drawings, can be arranged, substituted, combined, separated, and designed in a variety of different configurations, all of which are conceivable herein. Furthermore, unless the context otherwise requires, the features shown in each drawing can be used in combination with each other. Therefore, the accompanying drawings should be considered... Figure 1 These features are generally considered to be part of one or more overall embodiments, but it should be understood that not all of the illustrated features are necessary for each embodiment.

[0203] The embodiments set forth in the foregoing description do not represent all embodiments consistent with the subject matter described herein. Rather, they are merely some examples consistent with aspects related to the described subject matter. Although some variations have been described in detail above, other modifications or additions are possible. In particular, further features and / or changes may be provided in addition to those set forth herein. For example, the above embodiments may be applicable to various combinations and sub-combinations of the disclosed features and / or combinations and sub-combinations of several other features disclosed above. Furthermore, the logical flows depicted in the drawings and / or described herein do not necessarily require the specific order or sequence shown to achieve the desired results. Other implementations are within the scope of the appended claims.

[0204] The above description is merely a specific embodiment of this disclosure, but the scope of protection of this disclosure is not limited thereto. Any person skilled in the art can make various changes or substitutions within the technical scope disclosed in this disclosure, and such changes or substitutions should all be covered within the scope of protection of this disclosure. Therefore, the scope of protection of this disclosure should be determined by the scope of the claims.

Claims

1. A method for matrix loading, characterized by, The method comprises: receiving a logical layout of an input tensor, a logical layout of an output tensor, and a physical layout type of an output register matrix specified by a user; determining physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix; performing a traversal process for the output register matrix based on the physical layout information; and loading the matrix from memory to registers based on the traversal process, wherein the traversal process comprises iteratively performing the following operations for each register in the output register matrix until the traversal of all registers in the output register matrix is completed: calculating a current register address of a current register based on the physical layout information; determining a first element logical coordinate of a current output tensor corresponding to the current register based on the current register address and the physical layout information; determining a first element logical coordinate of a current input tensor corresponding to the current output tensor based on the first element logical coordinate of the current output tensor; calculating a current input memory address corresponding to the first element logical coordinate of the current input tensor based on the first element logical coordinate of the current input tensor.

2. The method for matrix loading of claim 1, wherein, The physical layout information of the output register matrix comprises a physical layout of the output register matrix and a one-to-one first mapping relationship between the logical layout of the output tensor and the physical layout of the output register matrix.

3. The method for matrix loading of claim 2, wherein, The physical layout of the output register matrix comprises a register index layout and a thread layout, and The logical layout of the input tensor, the logical layout of the output tensor, the register index layout, and the thread layout each comprise a dimension vector and a stride vector of the layout, wherein the dimension vector indicates dimension information of the layout, and the stride vector indicates a stride size of each dimension.

4. The method for matrix loading of claim 1, wherein, The physical layout type of the output register matrix corresponds to a hardware structure or an instruction mode.

5. The method for matrix loading of claim 3, wherein, The method further comprises: allocating a memory base address in memory for the matrix and a register base address in registers for the matrix before performing the traversal process.

6. The method for matrix loading of claim 5, wherein, Performing the traversal process for the output register matrix based on the physical layout information comprises: performing the traversal process for the output register matrix based on the dimension vector and the stride vector of the register index layout.

7. The method for matrix loading of claim 6, wherein, Performing the traversal process for the output register matrix based on the dimension vector and the stride vector of the register index layout comprises: performing the traversal process for the registers in the output register matrix in each dimension from small to large based on the dimension coordinates of the registers in the register index layout and the stride size of each dimension.

8. The method for matrix loading of claim 7, wherein, Calculating the current register address of the current register based on the physical layout information comprises calculating based on the following formula: current register address = register base address + dimension coordinates of the current register · (stride vector / thread total number) * single element size / single register size, wherein the "·" symbol represents dot multiplication.

9. The method for matrix loading of claim 8, wherein, determining, based on the current register address and the physical layout information, a first-element logical coordinate of a current output tensor corresponding to the current register, comprising: determining, based on the current register address, the first-element logical coordinate of the current output tensor corresponding to the current register according to the first mapping relationship, starting from a first thread of the current register.

10. The method for matrix loading of claim 9, wherein, determining, based on the first-element logical coordinate of the current output tensor, a first-element logical coordinate of the current input tensor corresponding to the current output tensor, comprising: determining, based on a second mapping relationship between the input tensor and the output tensor, the first-element logical coordinate of the current input tensor corresponding to the current output tensor.

11. The method for matrix loading of claim 10, wherein, calculating, based on the first-element logical coordinate of the current input tensor, a current input memory address corresponding to the first-element logical coordinate of the current input tensor, comprising calculating based on the following formula: current input memory address = memory base address + first-element logical coordinate of the current input tensor · step vector * single-element size, wherein the symbol "·" represents dot multiplication.

12. The method for matrix loading of claim 10, wherein, the second mapping relationship is based on a semantic relationship between the input tensor and the output tensor, and the semantic relationship comprises point-to-point data loading or broadcast data loading.

13. The method for matrix loading of claim 11, wherein, the traversal process further comprises: determining, based on the current register address and the thread layout, the logical coordinates of each element corresponding to each thread in the current output tensor according to the first mapping relationship, and determining a current thread mask by comparing the logical coordinates of each element corresponding to each thread with the logical layout of the input tensor, the current thread mask being used to disable the execution of a part of threads.

14. The method for matrix loading of claim 13, wherein, the traversal process further comprises: generating a single assembly instruction based on the current register address, the current input memory address and the current thread mask.

15. The method for matrix loading of claim 13, wherein, the traversal process further comprises: iteratively performing continuous traversal to continuously traverse a plurality of registers in the output register matrix; and generating a single assembly instruction for the plurality of registers.

16. The method for matrix loading of claim 15, wherein, the continuous traversal further comprises: after determining the current register address, the current input memory address and the current thread mask, continuing to the next iteration to determine a next register address, a next input memory address and a next thread mask in the next iteration; comparing the current register address, the current input memory address and the current thread mask with the next register address, the next input memory address and the next thread mask; in response to the current register address being continuous with the next register address, the current input memory address being continuous with the next input memory address, the current thread mask being the same as the next thread mask, the number of registers that have been continuously traversed not reaching an upper limit of the number of registers supported by hardware for continuous traversal, and the traversal of all registers not being completed, continuing to continuously traverse the registers in the output register matrix, Otherwise, generating a single assembly instruction based on the first register address, the first input memory address, the thread mask of the consecutive iteration, the execution granularity and the user configuration information in the first iteration of the consecutive iterations, and continuing the iteration process after the single assembly instruction is generated, wherein the execution granularity indicates a number of registers that have been consecutively iterated.

17. The method for matrix loading of claim 1, wherein, loading the matrix from the memory to the registers based on the iteration process, comprises: loading the matrix from the memory to the registers by executing one or more assembly instructions generated by the iteration process.

18. The method for matrix loading of claim 17, wherein, The method further comprises: encapsulating the iteration process as an intermediate representation, and generating the one or more assembly instructions by invoking the intermediate representation.

19. An apparatus for matrix loading, the apparatus comprising: The apparatus comprises: a receiving module configured to receive a logical layout of an input tensor, a logical layout of an output tensor and a physical layout type of an output register matrix specified by a user; a determining module configured to determine physical layout information of the output register matrix based on the logical layout of the output tensor and the physical layout type of the output register matrix; an executing module configured to execute an iteration process for the output register matrix based on the physical layout information, and load the matrix from the memory to the registers based on the iteration process; The iteration process comprises, for each register in the output register matrix, iteratively performing the following operations until the iteration of all registers in the output register matrix is completed: calculating a current register address of a current register based on the physical layout information; determining a first element logical coordinate of a current output tensor corresponding to the current register based on the current register address and the physical layout information; determining a first element logical coordinate of a current input tensor corresponding to the current output tensor based on the first element logical coordinate of the current output tensor; calculating a current input memory address corresponding to the first element logical coordinate of the current input tensor based on the first element logical coordinate of the current input tensor.

20. An electronic device for matrix loading, comprising: The electronic device comprises: one or more processors; a memory storing one or more computer programs; wherein the one or more computer programs, when executed by the one or more processors, cause the one or more processors to implement the method for matrix loading according to any one of claims 1-18.

21. One or more non-transitory computer-readable storage media storing instructions that, when executed by one or more processors of a computing device, cause the computing device to perform any of the methods of claims 1-20. The one or more non-transitory computer-readable storage media store computer-executable instructions thereon that, when executed by a processor, cause at least one processor to implement the method for matrix loading according to any one of claims 1-18.

Images