Large model cross-platform system based on micro operator

Abstract

The invention relates to the field of electric digital data processing, in particular to a micro-operator-based large-model cross-platform system, which comprises a hardware execution layer, a hardware access layer, a pre-compiling layer, a micro-operator layer, a rear-end layer, a framework layer and a model layer. The embodiment of the invention can be suitable for different hardware platforms, can quickly realize the reasoning and fine tuning of a large model, effectively simplifies the operator development, debugging and deployment process, and improves the cross-platform model development efficiency. The invention provides a micro-operator development standard for a large model scene. The micro-operator development standard comprises a device access management micro-operator, a tensor calculation micro-operator and a large model typical operator realized by combination. In addition, the invention provides a software architecture scheme, which defines a large model development process based on the micro operator, covers a micro operator definition and calling mode, adaptation methods of different back ends, calculation core scheduling and a calculation framework and model building scheme based on the micro operator, so as to realize high efficiency and cross-platform adaptation capability of large model reasoning and fine tuning.

Description

Technical Field

[0001] The present invention relates to the field of electronic digital data processing, and in particular to a large-model cross-platform system based on micro-operators. Background Art

[0002] With the rapid development of large model technology, the computing power required for training and reasoning large models is increasing, which puts higher requirements on efficient and easy-to-use artificial intelligence computing frameworks. The current mainstream artificial intelligence computing frameworks, such as PyTorch, TensorFlow and Keras, provide the tools needed to build large models. However, these frameworks often rely too much on specific hardware platforms, resulting in many difficulties in their application on emerging or specific hardware architectures. Emerging or specific hardware architectures include domestic domain specific architectures (DSA), field programmable gate arrays (FPGA), general-purpose computing on graphics processing units (GPGPU), etc. There are currently two main solutions: one is to manually write a deep neural network (DNN) operator library that meets the definition of the mainstream artificial intelligence computing framework for a specific hardware platform; the other is to use multi-level compilation technology (MLIR) to automatically generate operator code through just-in-time (JIT). However, in the era of large models, both solutions face severe challenges.

[0003] For the development of traditional operator libraries, developers need to have an in-depth understanding of the underlying hardware architecture and algorithm optimization, and the development process is complex and time-consuming. In addition, with the continuous iteration of hardware and algorithms, the operator library needs to be continuously updated and maintained to ensure compatibility with new hardware and frameworks, resulting in high maintenance costs. In addition, traditional operator libraries also need to adapt to different versions of front-end deep learning frameworks, which involves complex interface adaptation and compatibility issues, further increasing the difficulty of development. In addition, the debugging of end-to-end models requires dealing with multiple levels of complexity superposition, making operator performance tuning and code debugging extremely difficult. In contrast, although the compilation scheme can automatically generate operator code for different platforms, it also has problems. For example, compilation takes a long time, cache management is difficult, and efficiency is easily reduced due to repeated compilation. In addition, the debugging of the compilation scheme is difficult, and the problem location and repair process is very complicated. Especially for dynamic input, the compilation scheme has poor support, and repeated compilation is required for variable-length input, which is difficult to meet the requirements of efficient reasoning and fine-tuning of large models. In summary, the prior art cannot provide a system that can quickly implement large model reasoning and fine-tuning on multiple hardware platforms. This makes the development, debugging, and deployment of operators for large models more difficult and challenging. Summary of the invention

[0004] The present invention provides a large-model cross-platform system based on micro kernels, aiming to solve the problem that the existing technology cannot realize fast large-model reasoning and fine-tuning across hardware platforms, thereby reducing the difficulty of operator development, debugging and deployment. The system architecture proposed by the present invention is simple and efficient, which can significantly reduce the complexity of large-model operator development on different hardware platforms and improve the efficiency of development and deployment.

[0005] The present invention provides a large model cross-platform system based on micro-operators, comprising: a hardware execution layer, a hardware access layer, a pre-compilation layer, a micro-operator layer, a back-end layer, a framework layer and a model layer; wherein the hardware execution layer is used to execute memory access instructions and / or various micro-operators compiled by the pre-compilation layer; the hardware access layer is used to perform high-level abstraction on various interfaces of various hardware devices accessing the large model cross-platform system, and form a consistent calling interface for the upper layer; the pre-compilation layer is used to call the platform compiler to compile the micro-operator code into intermediate instructions or executable instructions of the hardware device; the micro-operator layer is used to call various micro-operators compiled by the pre-compilation layer and provide a programming interface based on micro-operators to the upper layer, so that the upper layer can combine and call multiple micro-operators to realize various operators required by the mainstream large model; the back-end layer is connected to the framework layer upward and to the micro-operator layer downward, and is used to expose the functional interface of the micro-operator layer to the outside, and at the same time, the back-end layer can recognize the tensor data type, and the same tensor compiled by the pre-compilation layer is called by the micro-operator layer. Different data types of micro-operators are implemented; the framework layer contains various functional modules required for realizing large model reasoning and fine-tuning, including tensor modules, storage modules, device modules, weight loading modules, custom operator modules, neural network modules, weight quantization modules, word vector encoding modules, back propagation modules and output sampling modules. Each functional module realizes the corresponding large model reasoning function or large model fine-tuning function by calling the functional interface in the back-end layer; the model layer is used to provide users with model layer specifications for building large model instances based on the framework layer; wherein, providing users with model layer specifications for building large model instances based on the framework layer includes: using the tensor modules, neural network modules and custom operator modules provided by the framework layer to build a large model instance based on micro-operators according to the mainstream large model structure; using the weight loading module provided by the framework layer to load the large model weights and calling the weight quantization module provided by the framework layer according to the model requirements to quantize and dequantize the large model weights; calling the functional modules in the framework layer to perform reasoning and fine-tuning of the large model.

[0006] The present invention defines the hardware execution layer specification. Any hardware device that complies with this specification can be connected to the system of the present invention for large model reasoning and fine-tuning. Hardware devices include but are not limited to CPU, GPU, GPGPU, DSA, FPGA, etc.

[0007] Optionally, the hardware execution layer needs to have parallel computing capabilities and provide interfaces to support the launch of a specified number of threads to execute specified memory access instructions and / or various micro-operators compiled by the pre-compiled layer, and support event, asynchronous and streaming processing. The hardware execution layer supports half-precision (FP16, BF16) or integer (INT8) for reasoning, and supports higher precision (FP32, TF32, FP64) for fine-tuning. The hardware execution layer supports the launch of pre-compiled operator code and has a cache function, and supports parameter passing between the host and the device. The hardware execution layer (hardware device) needs to have a large capacity of internal high-speed storage (such as high-bandwidth memory HBM) and the ability to connect the host and the device or between devices at a high speed, and provide interfaces for memory access operations. Memory access instructions refer to instructions for data loading operations and data storage operations. The various types of micro-operators compiled are different types of operators for calculations or operations. Each type of micro-operator is an operator for specified calculations or specified operations.

[0008] The present invention uses a unified hardware access layer to be compatible with different hardware devices. The hardware access layer performs a higher level of abstraction on various interfaces in the hardware abstraction layer of different hardware manufacturers. Various interfaces are interfaces for performing various operations, including device query interfaces, computing core management / module interfaces, computing core interfaces, event interfaces, streaming interfaces, memory management interfaces, etc. Various interfaces are abstracted at a high level to form a calling interface for calling various hardware functions.

[0009] The present invention defines the function of the precompilation layer. When the large model cross-platform system is compiled into a binary program, the precompilation layer calls the platform compiler to compile the pre-written micro-operator code into an intermediate instruction or an executable instruction of the hardware device. The precompilation layer is also used to provide a cache function for precompiled instructions. When compiling and generating a binary program for the host-side reasoning and fine-tuning system, the precompilation layer calls the platform compiler to compile the micro-operator code into an intermediate instruction or an executable instruction of the hardware device, thereby obtaining the micro-operator. Specific error information is provided when the compilation fails. The micro-operator code of each type of micro-operator is a pre-written code that can obtain the type of micro-operator after compilation. The intermediate instructions or executable instructions of the hardware device obtained after the micro-operator code of each type of micro-operator is compiled are various types of micro-operators. Intermediate instructions refer to intermediate codes generated during the compilation process. The executable instructions of the hardware device refer to instructions that can be executed and acted on the hardware device connected to the large model cross-platform system. The platform compiler is a pre-set compiler for compiling micro-operator codes into intermediate instructions or executable instructions of hardware devices. The micro-operator codes of various types of micro-operators can be written by technicians and compiled and generated by the precompilation layer. The micro-operators compiled by calling the platform compiler must be compatible with the current hardware platform, have horizontal portability, and have backward compatibility, that is, the new hardware is compatible with the old instructions, or the old code can be recompiled to use the new hardware instructions. The precompilation layer provides a caching function for precompiled micro-operators, and there is no need to recompile the unchanged micro-operator source code. If the precompiled micro-operator library has been cached locally and the micro-operator source code has not been modified, this precompilation process can be skipped directly. The precompilation layer can also provide parallel compilation function. During the host program compilation process, multiple micro-operator codes can be precompiled at the same time to improve compilation efficiency.

[0010] The present invention also discloses a micro-operator layer for defining and implementing various operators required by mainstream large models. The micro-operator layer includes device memory access management micro-operators, tensor calculation micro-operators, and various operators required by mainstream large models realized by combining tensor calculation micro-operators. Device memory access management micro-operators include device memory allocation micro-operators, device memory release micro-operators, host-to-device communication micro-operators, device-to-host communication micro-operators, device-to-device communication micro-operators, device synchronization micro-operators, and device query micro-operators. Tensor calculation micro-operators include single-operand operators, double-operand operators, type conversion operators, index operators, general copy operators, conditional selection operators, aggregation operators, and matrix multiplication operators. Tensor calculation micro-operators are combined to realize various operators required by mainstream large models, and various operators required by mainstream large models include rotational position encoding operators, activation layer operators, normalization operators, standardization operators, and attention layer operators.

[0011] Optionally, the micro-operator layer includes device memory access management micro-operators. For example, device memory allocation micro-operator (alloc), device memory release micro-operator (free), host-to-device communication micro-operator (h2d), device-to-host communication micro-operator (d2h), device-to-device communication micro-operator (d2d), device synchronization micro-operator (sync), and device query micro-operator. The micro-operator layer also includes tensor calculation micro-operators. For example, single-operand operator (unary), double-operand operator (binary), type conversion operator (cast), indexing operator (indexing), general copy operator (general copy), conditional selection operator (select), aggregation operator (reduce), and matrix multiplication operator (matmul). Among them, single-operand operator and double-operand operator are general micro-operator names. Single-operand operator can further include multiple actual single-operand operators. For example, the sine operator (sin), the cosine operator (cos), the exponential operator (exp), the hyperbolic tangent operator (tanh), and the rectified linear unit operator (relu). The dual-operand operator may further include multiple actual dual-operand operators. For example, the addition operator (add), the subtraction operator (sub), the multiplication operator (mul), the division operator (div), and the differential operator (eq). Tensor calculation operators are combined to realize various operators required by mainstream large models. Various operators (fused kernels) required by mainstream large models include Rotary Position Embedding (RoPE), activation layer operator (activation), normalization operator (softmax), standardization operator (layernorm / rmsnorm), and attention layer operators (multiheaded attention, flash attention, paged attention), etc.

[0012] Optionally, the micro-operator layer exposes a micro-operator-based programming interface to the host code. The process of the host code calling the micro-operator to perform calculations is like a normal function call. The host code can implement various operators required by the mainstream large model by combining multiple micro-operators. The micro-operator-based programming interface is an interface for combining and calling various micro-operators. For the combined call of micro-operators, streaming computing can be used, that is, a group of micro-operators are executed on the same data stream / computation stream (stream) to improve the throughput of hardware execution. For the implementation of various micro-operators, generic programming can be used, that is, the same micro-operator supports the calculation of different data types. At the same time, the interface name exposed by the micro-operator layer is suffixed with the data type to distinguish it. For example, the micro-operator cast_f32_f16 converts single-precision floating point FP32 data to half-precision floating point FP16 data, and the micro-operator softmax_bf16 normalizes 16-bit floating point format BF16 data.

[0013] The present invention also discloses a backend layer (backend), which is used to connect to the framework layer upwards and the micro-operator layer downwards, and at the same time, the functional interface used to call various micro-operator layers is exposed to the outside. Corresponding to the tensor calculation micro-operator, the backend layer exposes the single operand operator interface (unary_impl), the double operand operator interface (binary_impl), the type conversion operator interface (binary_impl), the indexing operator interface (indexing_impl), the general copy operator interface (general_copy_impl), the conditional selection operator interface (select_impl), the aggregation operator interface (reduce_impl) and the matrix multiplication operator interface (matmul_impl). The backend layer can identify the data type of the tensor data to be processed. For the functional implementation of all micro-operators, the backend layer parses the interface call parameters of the framework layer and identifies the data type of the parameters, loads the corresponding micro-operator compiled by the pre-compiled layer in the form of a calculation kernel launch (kernel launch) according to the data type, starts execution through the hardware execution layer and encapsulates and returns the calculation results, and the encapsulated calculation results are returned to the backend layer. Corresponding to the device memory access management micro-operator, the backend layer exposes the device memory allocation micro-operator interface (alloc_impl), the device memory release micro-operator interface (free_impl), the host-to-device communication micro-operator interface, the device-to-host communication micro-operator interface, the device-to-device communication micro-operator interface, the device synchronization micro-operator interface (sync_impl), and the device query micro-operator interface.

[0014] Optionally, for different hardware platforms, corresponding functional interfaces (i.e., backend layers) and micro-operator codes of corresponding platforms can be implemented, and the framework layer can be accessed through the present invention to realize reasoning and fine-tuning of large models.

[0015] Optionally, the backend layer can convert non-contiguous memory blocks / non-contiguous tensors into contiguous tensors by calling a general copy micro-operator to give full play to the hardware's Single Instruction Multiple Data (SIMD) performance, while greatly reducing the difficulty of micro-operator development. That is, the backend layer uses a general copy micro-operator to ensure that other micro-operators only process contiguous tensors, without the need to develop computational kernel code for non-contiguous tensors. In order to convert all non-contiguous tensor calculations into contiguous tensor calculations, the backend layer also needs to track the layout change process of the tensor, which is used to record the original and target size (shape) and stride (stride) of the continuous tensor when it is converted to a non-contiguous tensor. Among them, the size and stride information are key information for calling the general copy micro-operator. The backend layer calls different implementations of micro-operators by identifying the tensor data type. For example, add_fp32, add_fp16, and add_bf16 called by the backend layer are hardware platform implementations of single precision (FP32), half precision (FP16), and half precision (BF16) of the addition micro-operator, respectively. Different data type implementations of the same micro-operator can refer to different operators generated on the basis of the same micro-operator for performing specified operations or specified calculations on data of different data types. Each operator performs a specified operation or specified calculation on data of one data type.

[0016] The present invention also discloses a framework layer. The framework layer includes various functional modules required for realizing large model reasoning and fine-tuning. Specifically, the framework layer includes a tensor module, a storage module, a device module, a weight loading module, a custom operator module, a neural network (layer) module, a weight quantization module (quantization), a tokenizer module (tokenizer), a backpropagation module, and an output sampling (sampling) module, etc.

[0017] Optionally, the tensor module is an abstract representation of tensors in a neural network. The tensor module contains tensor data and various common tensor operations. Various tensor operations are implemented by calling the functional interface of the corresponding micro-operator in the backend layer. The tensor module not only holds the corresponding tensor data, but also contains common operations applied to tensors. For example, reshape, transpose, squeeze, stack, and various micro-operator operations, such as single operand (unary), double operand (binary), broadcast, matrix multiplication (matmul), tensor conversion to continuous storage (to_contiguous), conditional selection (index_select), aggregation (reduce) and other micro-operator operations. Through the tensor module, you can query the data type, shape, layout, whether it is non-contiguous storage, the type of hardware where the tensor data is located, etc. In addition, the tensor module also overloads various operators (such as +, -, *, / , =, ==, etc.). Therefore, through the tensor module, you can easily create and initialize tensors and add various operations to tensors (or between tensors).

[0018] Optionally, the storage module is a low-level representation of tensors and tensor operations, which is used to identify the device class and data type of the tensor. The operation of identifying the device class and data type of the tensor is implemented by calling the functional interface of the corresponding micro-operator in the backend layer. For example, the first tensor is a tensor of data type single precision FP32 stored on the GPU device. Applying a linear rectification operation (ReLU) to the first tensor will call the relu_f32 micro-operator of the Compute Unified Device Architecture (CUDA) backend layer. It should be noted that only tensors of the same device and data type can be used for binary or multivariate operations. The second tensor is a tensor of data type half precision FP16 stored on the CPU. It is not feasible to directly add the first tensor to the second tensor. The general copy operator provided by the backend layer is required to move the second tensor from the CPU to the GPU, and then call the type conversion operator to convert the data type to single precision floating point FP32 before adding it to the first tensor. An exception is that for some special operators such as matrix multiplication operators, the data types of their binary operands can be different. For example, when the micro operator matmul_f32_f16 performs matrix multiplication, its left operand is single-precision FP32, and the right operand (usually weight) is half-precision FP16 data. In this case, the operator needs to support data type conversion or the hardware needs to support binary operations of different data types, but even so, the operands must be in the same device.

[0019] Optionally, the device module is a further abstraction of the device functions provided by the backend layer. The device module is used to provide the functional interfaces required for memory access operations and micro-operator execution operations on different hardware devices. Memory access operations include the operations required for allocating, recycling and migrating tensors between devices on different devices. For example, by calling the micro-operator alloc_uninit of the device module (while specifying the device type as GPU, the data type as single-precision FP32 and the shape as M×K×N), an uninitialized single-precision FP32 type area of ​​size 4×M×K×N is allocated on the GPU to represent the data held by the first tensor. By calling the data migration (to_device) functional interface of the first tensor to specify the target device to which the first tensor needs to be transferred, the functional interface copy_to of the device module is further called to realize the migration of the first tensor between different devices. The activation function is applied to the first tensor by calling the functional interface relu of the first tensor, and the functional interface launch of the device module is further called to determine the device and data type of the tensor, and to call the micro-operator of the corresponding backend layer, such as the micro-operator relu_f32 (note that loading and calling pre-compiled micro-operators such as micro-operator relu_f32 requires the runtime system runtime support of the hardware platform). At the same time, the calculation execution mode can be set through the device module, that is, synchronous execution or asynchronous execution. For synchronous execution (launch), the execution result will be returned immediately after calling the micro-operator relu_f32. For asynchronous execution, the calculation stream (stream) of the calculation operation needs to be specified, and the upper layer calls the synchronous interface of the device module at the appropriate time to wait for the result of the micro-operator or multiple micro-operators to return.

[0020] Optionally, the weight loading module is used to load the large model weights from the file. Among them, the weight loading process is triggered by the model layer. The weight loading module sets the target data type of the large model weight when creating the loader, converts the data type of the large model weight during the weight loading process, and converts the data type of the large model weight to the target data type. The data type conversion is performed dynamically during the weight loading process. For example, for a single-precision FP32 weight file, if the set data type is half-precision FP16, the function interface to_dtype operation will be called for each weight tensor during the weight loading process, which will further call the back-end layer micro operator cast_f32_f16 to convert the single-precision FP32 tensor to a half-precision FP16 type tensor. In addition, it should be pointed out that the weight loader is usually loaded in sequence according to the neural network layer, and the loading order and the weight name of each layer are generally specified by the upper layer (model layer).

[0021] Optionally, the custom operator module is used to provide a functional interface for extended micro-operators. The functional interface of the extended micro-operator includes an extended interface for a single-operand operator, a dual-operand operator, and a multi-operand operator. The functional interface of the extended micro-operator is used to implement a non-standard micro-operator or an aggregation operator. A non-standard micro-operator may be a new micro-operator that needs to be added. An aggregation operator may be an operator with new functions obtained by aggregating different micro-operators or a fusion operator directly written using a platform programming tool. For the implementation of non-standard micro-operators or when an aggregation operator needs to be implemented, the corresponding custom operator may be implemented in different backend layers. Only the type and parameters of the custom operator need to be specified, and the custom operator interface of the tensor may be called to perform the custom operator operation. For example, the custom operator RMSNorm is an extended single-operand operator, which is internally implemented as a combination of a group of micro-operators (composed of a square root operator, a broadcast operator, a division operator, a multiplication operator, and an addition operator, etc.) or a single micro-operator (written by the corresponding hardware platform code, such as a micro-operator CUDA kernel).

[0022] Optionally, the neural network module defines common large model operators, which are implemented by using micro operators or micro operator combinations, or by calling aggregation operators through custom operator modules. For example, the implementation forms of common large model operators such as linear layer operators, activation layer operators, normalization operators, rotational position encoding operators, and attention layer operators are usually divided into two types, namely, directly using micro operators or micro operator combinations to write. Linear layer operators can be composed of matrix multiplication operators, broadcast operators, and addition operators. Normalization operators can be composed of square root operators, broadcast operators, division operators, multiplication operators, and addition operators. Rotational position encoding operators and attention layer operators can be written directly by hardware platform code. The neural network module is also responsible for loading the corresponding weights. For example, the linear layer operator needs to load the weights required by the matrix multiplication operator.

[0023] Optionally, the weight quantization module is an extension of the tensor module. In addition to providing the standard operations defined by the tensor module, it also provides quantization and dequantization function interfaces to support low-precision large model reasoning and fine-tuning. The standard tensor module supports operations on common data types such as single-precision FP32, single-precision TF32, half-precision FP16, and half-precision BF16. However, for reasoning scenarios, lower precision such as 8-bit signed integer INT8, 4-bit signed integer INT4, and 6-bit floating point format FP6 are usually used to improve the operation speed. The weight quantization module provides additional quantization (quantize) and dequantization (dequantize) function interfaces in addition to the standard operations defined by the tensor module. The quantization and dequantization function interfaces can perform lossy compression (quantization) on the held tensor data and perform corresponding quantization operations (such as the micro operator qmatmul_f16_fp6, that is, matrix multiplication of half-precision FP16 type tensors and 6-bit precision floating point FP6 type tensors), and then perform dequantization on the output results to restore the accuracy. It should be pointed out that the quantization function requires additional micro-operator support. For example, micro-operator quantize_f16_fp6, micro-operator qmatmul_f16_fp6, micro-operator dequantize_fp6_f16, etc. Quantization, quantization matrix multiplication, and dequantization functions all require the writing of corresponding hardware platform codes. The quantization function access framework layer can use custom operator forms. Low-precision large model reasoning and fine-tuning refers to large model reasoning and fine-tuning based on quantized model weight data.

[0024] Optionally, the word vector encoding module is used to implement the word vector encoding operation required for large model reasoning or large model fine-tuning by calling the micro-operators in the micro-operator layer. The word vector encoding module is used to encode the input text sequence of the model into a word vector. The word vector is also called a token. There are many implementation forms of the word vector encoding module, usually using the BPE algorithm and the WordPiece algorithm, which will not be repeated here.

[0025] Optionally, the back-propagation module is used to fine-tune the large model. The back-propagation module relies on the tensor module to record the operation type and parameters during the tensor operation, and calculates the reverse gradient through the reverse calculation interface of the tensor module; the reverse gradient calculation operators corresponding to various tensor operations are defined in the back-propagation module, and the reverse gradient calculation operators corresponding to various tensor operations are implemented by combining various micro-operators or written and implemented by calling the aggregation operator through the custom operator module. The implementation of the back-propagation module requires the support of the tensor module (recording the operation type and parameters during the tensor operation for reverse gradient calculation), and provides the function interface backward upward. Among them, the function interface backward is used to calculate the gradient, and its process is to reversely traverse the tensor operations applied according to the difference (or loss) between the output tensor and the expected result, and calculate the gradient from back to front according to the loss. The back-propagation module defines the back-propagation operators corresponding to various tensor operations for calculating gradients. Its implementation is similar to that of custom operators, which are implemented by a combination of various basic micro-operators (such as addition operators, subtraction operators, multiplication operators, exponential operators, square root operators, type conversion operators, negative operators, etc.) or by hardware platform code. After the gradient calculation is completed, the upper layer fine-tunes the model weights according to the gradient and the corresponding optimization scheme (for example, the scheme using optimizer, learning rate and other hyperparameters for optimization).

[0026] Optionally, the output sampling module is used to convert the word vectors generated by the large model into character sequences. The output sampling module is mainly used to convert the word vectors (tokens) generated by model inference into corresponding character representations. Common sampling algorithms include the sampling algorithm top-k, the sampling algorithm argmax, and the sampling algorithm multinomial. Similar to the word vector encoding module, the output sampling module can be implemented using industry standards, so it will not be repeated here.

[0027] The present invention also discloses a model layer specification for providing users with building large model instances based on the framework layer. The model layer specification for providing users with building large model instances based on the framework layer includes: using the tensor module, neural network module and custom operator module provided by the framework layer to build a large model instance based on micro-operators according to the mainstream large model structure; using the weight loading module provided by the framework layer to load the large model weights and calling the weight quantization module provided by the framework layer according to the model requirements to quantize and dequantize the large model weights; calling the functional modules in the framework layer to perform reasoning and fine-tuning of the large model. Large model instances based on micro-operators refer to various types of large models written according to the mainstream large model structure according to the model writing specifications described in the present invention.

[0028] The embodiment of the present invention provides a large model cross-platform system based on micro-operators, which can be used to quickly perform reasoning and fine-tuning of large models on different hardware platforms, and is suitable for rapid development, debugging, reasoning and fine-tuning of large models on different hardware platforms, which simplifies the difficulty of large model operator development, has a simple system structure, and is conducive to model development, debugging and cross-platform deployment. It defines the hardware access method, micro-operator implementation form and calling method, and various functional modules required for large model reasoning and fine-tuning based on micro-operators. Its overall system structure design is simple, which can effectively reduce the difficulty of operator development and improve the efficiency of large model development and deployment.

[0029] It should be understood that the contents described in this section are not intended to identify the key or important features of the embodiments of the present invention, nor are they intended to limit the scope of the present invention. Other features of the present invention will become easily understood through the following description. BRIEF DESCRIPTION OF THE DRAWINGS

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

[0031] Figure 1 A schematic diagram of the structure of a large model cross-platform system based on micro-operators provided in an embodiment of the present invention.

[0032] Figure 2 A schematic diagram of an example directory structure of a precompiled layer provided in an embodiment of the present invention. DETAILED DESCRIPTION

[0033] In order to enable those skilled in the art to better understand the scheme of the present invention, the technical scheme in the embodiments of the present invention will be clearly and completely described below in conjunction with the drawings in the embodiments of the present invention. Obviously, the described embodiments are only part of the embodiments of the present invention, not all of the embodiments. Based on the embodiments of the present invention, all other embodiments obtained by ordinary technicians in this field without creative work should fall within the scope of protection of the present invention.

[0034] It should be noted that the terms "target", "first", "second", etc. in the specification and claims of the present invention and the above-mentioned drawings are used to distinguish similar objects, and are not necessarily used to describe a specific order or sequence. It should be understood that the data used in this way can be interchanged where appropriate, so that the embodiments of the present invention described herein can be implemented in an order other than those illustrated or described herein. In addition, the terms "comprise", "include" and "have" and any variations thereof are intended to cover non-exclusive inclusions, for example, a process, method, system, product or device that includes a series of steps or units is not necessarily limited to those steps or units that are clearly listed, but may include other steps or units that are not clearly listed or inherent to these processes, methods, products or devices.

[0035] The present invention proposes a large model reasoning and fine-tuning technical solution based on micro-operators. Compared with traditional operator library solutions and JIT real-time compilation solutions, the micro-operators defined in the present invention are simple to develop, and have high flexibility and scalability. At the same time, its pre-compilation capability solves the problems of long code generation time, low compatibility and poor dynamic support of just-in-time compilation solutions, thereby effectively reducing the difficulty of adapting mainstream large models on different hardware platforms.

[0036] Based on the micro-operator solution, the present invention proposes a complete end-to-end reasoning and fine-tuning solution, including hardware access forms, micro-operator definitions and calling methods, adaptation methods for different backends, and micro-operator-based computing frameworks and model building solutions. The overall system design structure is simple, the development difficulty is low, and the compatibility is high, which solves the problems of high operator development and maintenance costs and poor cross-platform compatibility in mainstream deep learning systems.

[0037] The experimental results of the present invention on the domestic DSA hardware platform show that the solution of the embodiment of the present invention can significantly reduce the difficulty of large model operator development. The large model reasoning system built based on this solution can be easily compatible with and adapt to a variety of mainstream large models and multimodal scenarios, and at the same time, the reasoning performance is comparable to the traditional operator library solution (developed by a professional operator development team, extreme optimization operators, and mainstream deep learning frameworks). Therefore, the present invention has significant technical advantages and can bring higher commercial value.

[0038] Figure 1A schematic diagram of the structure of a large model cross-platform system based on micro-operators provided in an embodiment of the present invention. This embodiment is applicable to the case of reasoning and fine-tuning large models used by enterprises. Mainstream large models include large language models (LLMs), large vision models (VLMs), vision-language models (VLMs) and large multimodal models (LargeMultimodalModels,LMMs). The large models used by enterprises are large language models for machine translation or intelligent question and answer, large visual models for face recognition, and large multimodal models for scene perception. Figure 1 As shown, the large model cross-platform system based on micro-operators specifically includes: a hardware execution layer 101, a hardware access layer 102, a pre-compilation layer 103, a micro-operator layer 104, a backend layer 105, a framework layer 106 and a model layer 107, and its structure and function are described below.

[0039] Among them, the hardware execution layer 101 is used to execute memory access instructions and / or various micro-operators compiled by the pre-compiled layer 103. The hardware access layer 102 is used to perform high-level abstraction on various interfaces of various hardware devices connected to the large model cross-platform system, and form a consistent calling interface for the upper layer. The pre-compiled layer 103 is used to call the platform compiler to compile the micro-operator code into intermediate instructions or executable instructions of the hardware device. The micro-operator layer 104 is used to call various micro-operators compiled by the pre-compiled layer 103 and provide a programming interface based on micro-operators to the upper layer, so that the upper layer can combine and call multiple micro-operators to realize various operators required by the mainstream large model. The back-end layer 105 is connected to the framework layer 106 upward and the micro-operator layer 104 downward, and is used to expose the functional interface of the micro-operator layer 104 to the outside. At the same time, the back-end layer 105 can recognize the tensor data type, and the different data types of the same micro-operator compiled by the pre-compiled layer 103 are implemented by the micro-operator layer 104. The framework layer 106 includes various functional modules required for realizing large model reasoning and fine-tuning, including tensor module, storage module, device module, weight loading module, custom operator module, neural network module, weight quantization module, word vector encoding module, back propagation module and output sampling module. Each functional module realizes the corresponding large model reasoning function or large model fine-tuning function by calling the functional interface in the backend layer 105. The model layer 107 is used to provide the user with the model layer specification for building a large model instance based on the framework layer 106; wherein, the model layer specification for building a large model instance based on the framework layer 106 is provided to the user, including: using the tensor module, neural network module and custom operator module provided by the framework layer 106 to build a large model instance based on micro-operators according to the mainstream large model structure; using the weight loading module provided by the framework layer 106 to load the large model weight and calling the weight quantization module provided by the framework layer 106 according to the model requirements to quantize and dequantize the large model weight; calling the functional modules in the framework layer 106 to reason and fine-tune the large model.

[0040] The following takes the domestic DSA hardware platform (Suiyuan GCU platform) access to the large-model cross-platform system based on micro-operators provided by the present invention as an example, and explains in detail various aspects involved in the present invention.

[0041] The precompilation layer is used to call the platform compiler to compile the micro-operator code into intermediate instructions or executable instructions for the hardware device. The function of the precompilation layer is to precompile the written micro-operator code into intermediate code or binary code (also called computing kernel code) that can be loaded and executed by the hardware platform before the program is executed. The hardware platform must provide the corresponding compiler. The implementation of the micro-operator is written in reference to the programming interface defined by the hardware platform. For example, the Suiyuan DSA hardware platform uses the extended C++ programming language to write operators, and provides a compiler topscc (Suiyuan platform compiler) to compile the C++ operator source code into assembly code (DSA hardware platform executable code).

[0042] The implementation of the precompiled layer can be written in the programming language Rust, and a new project can be created using the Rust manager Cargo. The written operator code files are placed in the project src directory, or in a separate directory (such as Figure 2 The kernels directory is shown in the figure), and the compilation of the operator code is processed in the project compilation file build.rs. The project compilation file build.rs is called during the project compilation process to process the pre-compilation related logic. Figure 2 A schematic diagram of an example directory structure of a precompiled layer provided in an embodiment of the present invention. The example directory structure of a precompiled layer is an exemplary directory structure for storing operator code files in a precompiled layer.

[0043] In the project compilation file build.rs, the modification status of each file in the operator directory can be detected and compared with the generated operator binary file. If the binary code of the operator is not generated, the operator file is directly compiled. Otherwise, it is determined whether the modification status of the operator file is later than the generation time of the corresponding binary file. If it is later than the binary file, it indicates that the operator file has been modified since the last precompilation and needs to be recompiled. Otherwise, the precompilation of the operator file is skipped. Before the operator code is precompiled, that is, before calling the platform compiler, it is also necessary to set the relevant environment variables and the header files that the operator code depends on, and pass the platform identifier to the compiler according to the hardware platform characteristics to correctly compile the operator binary file that conforms to the current hardware platform. To speed up precompilation, a thread pool can be used to start the precompilation of multiple operator codes at the same time and wait for the completion of all compilation tasks. In addition to the on-demand compilation of operator code, the precompilation layer needs to prompt the corresponding error information when the compilation fails, usually returning the compiler error information.

[0044] In addition to intensive computing operators such as matrix multiplication operators and two-dimensional convolution operators (conv2d), the micro-operator layer usually also includes single-operand operators, dual-operand operators, affine transformation operators (affine), type conversion operators, universal copy operators, index operators, aggregation operators, ternary operators, fill operators, etc. Operators required for large model reasoning and fine-tuning, such as rotational position encoding operators, attention operators, tensor mask filling operators (masked_fill), etc., can all be composed of the above micro-operators. Users can also use the custom operators mentioned in the present invention for expansion to achieve the development of fusion operators.

[0045] The micro-operator indexing (index_select) in the micro-operator layer can be written on the DSA hardware platform. As shown below, the operator code of the micro-operator indexing is declared as __device__ code, which uses the template template to support the type of type. The micro-operator indexing selects the corresponding data in the input data according to the given condition sequence and copies it to the output data.

[0046] template<typename ID_TYPE, typename T>

[0047] __device__ void index_select_kernel(size_t num_ids, ID_TYPE *ids,

[0048] T *in, T *out, size_t l_size, size_t dim_size, size_t r_size) {

[0049] int thread_id = GetThreadIdx();

[0050] __local__ ID_TYPE ids_buffer[MAX_IDS_SIZE];

[0051] tops_dte_ctx_t ctx;

[0052] dte_scope s(ctx);

[0053] memcpy(ctx, mdspan(Private, ids_buffer, num_ids), mdspan(Global, ids,num_ids));

[0054] int THREAD_STEP = 1;

[0055] int thread_step = 1;

[0056] / / determine thread steps

[0057] get_thread_steps(THREAD_STEP, thread_step, num_ids);

[0058] for (int i = 0; i < thread_step; i++) {

[0059] int idx = thread_id * THREAD_STEP + i;

[0060] for (int j = 0; j < l_size; ++j) {

[0061] int _idx = ids_buffer[idx];

[0062] mdspan hbm_in(Global, in + (j * dim_size + _idx) * r_size, r_size);

[0063] mdspan hbm_out(Global, out + (idx + j * num_ids) * r_size, r_size);

[0064] memcpy(ctx, hbm_out, hbm_in); / / copy selected buffer

[0065] }

[0066] }

[0067] }

[0068] To support multiple data types, the micro-operator layer can be extended using the C++ model and paradigm programming. Ten different types of micro-operators index_select can be defined, as shown below. For example, the micro-operator is_u32_bf16 represents an index_select micro-operator whose given condition data type is uint32 (u32) and whose input type is bfloat16 (bf16); the micro-operator is_u8_f16 represents an index_select micro-operator whose given condition data type is uint8 (u8) and whose input type is float16 (f16); the micro-operator is_u32_f32 represents an index_select micro-operator whose given condition data type is uint32 (u32) and whose input type is float32 (f32).

[0069] It should be pointed out that the external interfaces of the micro-operator layer are all defined by the implementation interface of the micro-operator. For example, the definition of the micro-operator index_select is:

[0070] #define IS_OP(TYPE, ID_TYPE, FN_NAME) \

[0071] extern "C" __global__ void FN_NAME( size_t id_num, \

[0072] ID_TYPE* ​​ids, TYPE *in, TYPE *out, \

[0073] size_t l_size, size_t dim_size, size_t r_size) \

[0074] { \

[0075] index_select_kernel<ID_TYPE, TYPE> \

[0076] (id_num, ids, in, out, l_size, dim_size, r_size); \

[0077] } \

[0078] IS_OP(__bf16, uint32_t, is_u32_bf16)

[0079] IS_OP(__fp16, uint32_t, is_u32_f16)

[0080] IS_OP(float, uint32_t, is_u32_f32)

[0081] IS_OP(uint8_t, uint32_t, is_u32_u8)

[0082] IS_OP(uint32_t, uint32_t, is_u32_u32)

[0083] IS_OP(__fp16, uint8_t, is_u8_f16)

[0084] IS_OP(__bf16, uint8_t, is_u8_bf16)

[0085] IS_OP(float, uint8_t, is_u8_f32)

[0086] IS_OP(uint8_t, uint8_t, is_u8_u8)

[0087] IS_OP(uint32_t, uint8_t, is_u8_u32)

[0088] Among them, the input and output (in, out) are of the same data type, the given input condition ids is of integer data type, and the other parameters are of size_t type. In a specific example, the input and output types are bf16, and the given input condition type is u32 type, the external interface of the micro operator is:

[0089] extern "C" __global__ void is_u32_bf16 (size_t id_num,

[0090] uint32_t * ids, __bf16 *in, __bf16 *out, size_t l_size,

[0091] size_t dim_size, size_t r_size)

[0092] Correspondingly, the backend layer provides an interface for the micro-operator index_select. If the types of input and output data and input conditions match it, the micro-operator is_u32_bf16 will be called in the form of kernel launch.

[0093] Multiple micro-operators can be combined arbitrarily to form the operators required for large model reasoning and fine-tuning. For example, the combination of micro-operators can be used to implement the commonly used normalization (layernorm / rmsnorm) operators in large models, as shown below:

[0094] impl Forward for LayerNorm {

[0095] fn forward(&self, x: &Tensor) -> Tensor {

[0096] let in_dtype = x.dtype();

[0097] let scale = 1 / x.dim(-1) as f64;

[0098] let x = x.to_dtype(Type::F32);

[0099] let x = if self.rms_norm{ x} else {

[0100] let mean = x.reduce_sum(-1) * scale;

[0101] x.sub(&mean)

[0102] };

[0103] let norm_x = x.sqr().reduce_sum(-1) * scale;

[0104] let x = x.div(&(norm_x + self.eps).sqrt());

[0105] let x = x.to_dtype(in_dtype).mul(&self.weight);

[0106] return x.add_if(&self.bias);

[0107] }

[0108] }

[0109] Among them, to ensure accuracy, the input data type is first converted to a high-precision representation (FP32), and then a combination of single-operand operators and double-operand operators (such as subtraction operators, square root operators, division operators, square root operators, addition operators, multiplication operators, etc.) are used for calculation according to the operator category (whether it is rmsnorm). The micro-operator Rmsnorm is a simplified normalization operator that removes the operation of subtracting the mean value. It should be pointed out that the type conversion to_dtype operation will call the type conversion operator. If the input type is FP16 and the output type is FP32, the type conversion operator cast_f16_f32 will be called. Similarly, other micro-operator calls, such as subtraction operators, addition operators, and division operators, will also call the micro-operator implementations of the relevant types.

[0110] Different hardware platforms need to implement their specific functions according to the interface defined by the backend layer. The backend layer loads and calls the pre-compiled computing kernel of the hardware platform according to the characteristics of the hardware platform through the interface call of the backend layer. Usually this call is launched in the form of computing kernel launch. For example, calling the micro operator index_select in the form of computing kernel launch is as follows:

[0111] impl IndexSelect {

[0112] fn call<T: DType> (&self, src: &GcuSlice <t>, src_l: &Layout,

[0113] dim: usize, dev: &GcuDevice) -> GcuSlice <t>{

[0114] let ids_l = &self.id_layout;

[0115] let ids_shape = ids_l.shape();

[0116] let ids_el = ids_shape.elem_count();

[0117] let src = match src_l.offsets() {

[0118] Some((o1, o2)) => src.slice(o1..o2),

[0119] _ => Err("index-select data must be contiguous!"),

[0120] };

[0121] let l_size: usize = src_l.dims()[..dim].iter().product();

[0122] let r_size: usize = src_l.dims()[dim + 1..].iter().product();

[0123] let dim_size = src_l.dims()[dim];

[0124] let out = dev.alloc:: <t>(ids_el * l_size * ri_size);

[0125] match &self.slice {

[0126] GcuStorageSlice::U32(slice) => {

[0127] let ptr = slice.slice(ids_l.offset()..);

[0128] let func = dev.load_func(&name:: <t>("is_u32"));

[0129] let params = (ids_el, ptr.device_ptr(),

[0130] src.device_ptr(), out.device_ptr(),

[0131] l_size, dim_size, r_size);

[0132] func.launch(&dev.launch_cfg, params);

[0133] }

[0134] GcuStorageSlice::U8(slice) => {

[0135] let ptr = slice.slice(ids_l.offset()..);

[0136] let func = dev.get_or_load_func(&name:: <t>("is_u8"));

[0137] let params = (ids_el, ptr.device_ptr(),

[0138] src.device_ptr(), out.device_ptr(),

[0139] l_size, dim_size, r_size);

[0140] func.launch(&dev.launch_cfg, params);

[0141] }

[0142] _ => Err("index_select ids should be u8 or u32"),

[0143] };

[0144] out

[0145] }

[0146] }

[0147] First, the calling parameters are organized according to the interface requirements of the micro-operator index_select, including the input and output data types, and then the output buffer is allocated using the device interface passed in by the backend layer (dev.alloc), and finally the computing kernel is loaded and the calculation is started (dev.load_func, dev.launch(&func)).

[0148] To simplify operator development, the backend layer converts all non-continuously stored tensors into continuously stored tensors before calling the micro-operator layer to perform calculations. This conversion process is implemented by a general copy operator and the backend layer. Specifically:

[0149] The micro-operator layer implements a universal copy operator, which can convert the input non-continuous storage tensor into a continuous tensor output according to the input and output target size (shape) and stride information. The specific conversion process is that for each output element, its position (index) in the input data is calculated according to the target size (shape) and stride information, and the data is copied to the corresponding output position according to the position. The core index calculation logic is as follows:

[0150] template<int RANK>

[0151] __device__ VType get_strided_index(VType &indexes, VType results,

[0152] VType dst_shape[], VType dst_strides[]) {

[0153] VType vec_rem[RANK];

[0154] vec_rem[0] = indexes;

[0155] for (int i = 0; i < RANK; i++) {

[0156] unsigned int dim_idx = RANK - 1 - i;

[0157] results = vadd(vmul(vrem(vec_rem[i], dst_shape[dim_idx]),

[0158] dst_strides[dim_idx]), results);

[0159] vec_rem[i + 1] = vdiv(vec_rem[i], dst_shape[dim_idx]);

[0160] }

[0161] return results;

[0162] }

[0163] As shown above, DSA vectorized instructions such as vadd, vmul, vdiv, etc. can be used to speed up the calculation. The use of these instructions allows the position calculation to be performed in batches. For example, each time 128 output positions are given (the length of indexes is 128, such as 0, 1, 2, 3...127), the function get_strided_index returns 128 calculation results, that is, the input element positions where the 128 elements are located. Then use the gather (vgather) instruction to batch obtain data from the input and write it to the corresponding output position at one time to complete the conversion of non-continuous tensors into continuously stored tensors.

[0164] Before calling each operator to implement the calculation, the backend layer first uses the universal copy operator to convert all non-contiguously stored tensors into continuously stored tensors, and then uses the continuously stored tensors as input data to call the corresponding operator implementation. For example, as shown in the following code, first determine whether the layout of the input data (self.slice) of the unary operation is non-contiguous. For the non-contiguously stored input data, first allocate temporary data, then call the universal copy operator to convert the input data into continuously stored data and place it in the temporary data dst, and at the same time modify the memory layout to be continuous. Finally, call the specific implementation of the unary operation with the temporary data dst as input.

[0165] fn unary_impl<T: UnaryOp> (&self, layout: &Layout) -> Self {

[0166] let device = self.device().clone();

[0167] if layout.is_contiguous() {

[0168] let slice = T::F.map(&self.slice, &device, layout);

[0169] Self { slice, device}

[0170] } else {

[0171] let mut dst= device.alloc(layout.shape(), self.dtype());

[0172] self.general_copy(&mut dst, 0, layout);

[0173] let slice = T::F.map(&dst.slice, &device, &layout.contiguous());

[0174] Self { slice, device}

[0175] }

[0176] }

[0177] As mentioned above, the framework layer connects to the backend layer downward and provides a functional interface for writing model instances upward. In addition to the word vector encoding and output sampling, which both use industry standards, the following introduces the implementation examples of other modules of the backend layer.

[0178] 1) Tensor module: holds tensor data and provides various computing interfaces for operating tensors. To represent different types of tensors, its structure is designed to store tensor data with storage members (managed by the storage module), the device member indicates the device where the tensor is stored, the dtype member indicates the data type of the tensor, the layout member indicates the memory layout of the tensor, and the backprop member records the forward operation applied to the tensor (for backpropagation and fine-tuning). To support multi-threaded access, the storage member of the tensor is protected by a read-write lock, allowing only one thread to write at a time, but allowing multiple threads to read at the same time. Tensors can be converted between different data types or different devices. For example, calling the to_dtype operation will call the to_dtype operation of the storage member, which will further call the type conversion (cast) micro-operator through the backend layer to convert the storage type of the tensor to the target type. At this time, other properties of the tensor will not change except the data type. Calling the to_device operation will match the target device based on the device where the current tensor is located. If they are in the same device, only the tensor is copied. If they are in different devices, the copy between the host and the device of the tensor storage member is triggered. In the case where the device to device (D2D) copy cannot be completed, the CPU memory is used as a relay. As mentioned earlier, the tensor module provides encapsulation of various interface calls, such as common reshape, transpose, squeeze, stack, and various micro-operator operations, such as unary, binary, broadcast, matmul, to_contiguous, index_select, reduce, etc. These operations will be called to different backend layers according to the storage type. For example, the GCU storage member will call the GCU backend layer to trigger the kernel launch on the Suiyuan DSA device. The layout attributes of the tensor, such as shape and stride, will change after operations such as reshape, broadcast, and reduce. For non-inplace operations, a copy after the operation is usually returned, and the original tensor remains unchanged. In order to implement the to_contiguous function to convert non-contiguously stored tensors into continuously stored tensors, the tensor module also records its layout change process, so that when calling to_contiguous, the original and target shape and stride information of the tensor can be passed to the generalcopy operator. In addition, in addition to querying the data type, shape, layout and device type, the tensor module can also query whether the tensor is non-contiguously stored (by judging the shape and stride information).The tensor module overloads various operators (such as +, -, *, / , =, ==, etc.), and its implementation is similar to ordinary binary operations and unary operations, so we will not go into details here.

[0179] 2) Storage module: It is a low-level representation of tensors and tensor operations. It directly connects to the backend layer. By identifying the device category and data type of the tensor, it calls different backend layers and different (data type) micro-operators to implement. Using the Rust programming language, the storage module can be defined as an enumeration type, as shown below, where CpuStorage, CudaStorage and GcuStorage represent the specific implementations of the CPU, GPU and GCU platforms respectively. Public interfaces can be defined in storage, and different platforms implement the corresponding interfaces to complete the call to the backend layer of the platform. Similarly, in order to distinguish the device type and data type of storage, storage also has device type and data type (dtype) attributes. Storage supports data copy between different devices to complete the to_device function, and its to_dtype function is responsible for calling the cast operation of different backends to complete the data type conversion. It should be pointed out that the calculation operation between tensors needs to be in the same device. This is because each tensor calculation is completed by its specific storage (specific backend) in the form of kernel launch, and different backends are isolated from each other. If binary or multi-element operands (tensors) are in different devices, you can use the to_device operation to pull all tensors to the same storage, such as GcuStorage, and then call the tensor calculation interface (that is, call the gcu backend and start the gcu kernel launch by calling the corresponding interface of GcuStorage). It should be noted that, except for to_device, the storage type returned by tensor calculation is always the same as the source input.

[0180] pub enum Storage {

[0181] Cpu(CpuStorage),

[0182] Cuda (CudaStorage),

[0183] Gcu(GcuStorage),

[0184] }

[0185] 3) Device module: It is a further abstraction of the device functions provided by the backend layer, which provides the functional interface required for device memory access and micro-operator execution (kernel launch). Generally speaking, the device module is a simple encapsulation of the runtime system (runtime, such as cudart) provided by the hardware platform. The basic runtime system runtime function should comply with the hardware execution layer specification of the present invention, and generally should provide functions such as creating devices, creating data streams / computation streams (streams), loading computing kernels (kernel code) and returning function addresses (kernel functions), executing computing kernel tasks (kernel launch), synchronization (sync), device memory allocation (alloc), device memory recovery (free), and communication between host and device memory (memcpy, d2d: device to device, d2h: device to host, h2d: host to device). It should be pointed out that the runtime system runtime needs to support both synchronous and asynchronous operations, that is, when the computing core task or memory access operation passes in data flow information or computing flow information, all operations should be executed asynchronously on the data flow or computing flow (return immediately after calling the computing core task or memory operation), and when there is no incoming data flow information or computing flow information, all operations are executed synchronously. When executing asynchronously, the upper layer controls the synchronization timing of the data flow information or computing flow information. Generally, for streaming computing, the runtime system runtime and the corresponding hardware platform need to provide a cache function, that is, during the execution of a group of operators, the intermediate results do not need to be frequently written back to the storage device (such as HBM), but are temporarily stored through a high-speed cache such as the L2 cache to accelerate streaming computing.

[0186] 4) Weight loading module: mainly used to load pre-trained network model weights from model weight files, such as Pytorch pth files, HuggingfaceSafetensors files or Numpy npy files. The model layer triggers the loading of weight files when building a large model. Usually, the loading process is carried out in sequence according to each layer of the network model, and the model layer provides a unique weight name when loading, such as the weight name layer1.weight1. Different loading methods can be used according to the type of weight file. Taking the standard model weight safetensors format of the open source community and Huggingface platform as an example, its loading process can be in the form of memory mapping, that is, mapping the model weight file to the physical host memory, and converting the mapped model weight file into a memory object (such as an instance safe tensor) in a deserialized manner. Instance safetensor deserialization is to first read the metadata (meta data) information according to the mapped memory address, and then read the corresponding tensor data according to the metadata information (such as tensor shape, tensor data type). After the tensor is loaded into the memory, it is decided whether to convert the weight tensor data type according to the target data type. It should be pointed out that the memory mapping method does not require all weights to be loaded into the host memory at once when loading the weight file. Each time a layer of weights is loaded in the model layer, the deserialization of the weight file corresponding to that layer, data type conversion (if necessary), and copying from the host to the device memory (if the weights are specified to be stored in the device memory) are triggered. When all model layers are loaded, the model's weight file is loaded into the specified device memory, and the model can be executed on the device.

[0187] 5) Custom operator module: provides a functional interface for extending micro-operators, that is, when the user considers connecting a new fusion operator or an operator that is not currently supported by the system of the present invention, the operator library can be extended through a custom operator (custom op). The steps for extending the operator library are: first, the user uses the programming interface provided by the hardware platform to write the corresponding operator code and connect to the pre-compiled layer to generate the corresponding operator binary file; second, define the new operator type according to the custom operator specification in the system of the present invention; finally, implement the forward calculation (forward) function of the custom operator, that is, load and call the newly written operator in the form of a computing kernel startup in the corresponding back-end layer. Taking the custom operator of binary operation as an example below, name() in the CustomOp interface indicates the operator name of the custom operator, and forward and backward are the forward calculation interface and reverse calculation interface of the custom operator, respectively.

[0188] pub trait CustomOp {

[0189] fn name(&self) -> String;

[0190] fn forward(&self, _: &GcuStorage, _: &Layout, _: &GcuStorage,

[0191] _: &Layout) -> (GcuStorage, Shape) {

[0192] todo!()

[0193] }

[0194] fn backward(&self, _arg1: &Tensor, _arg2: &Tensor, _res: &Tensor,

[0195] _grad_res: &Tensor) -> (Option <tensor>, Option <tensor>) {

[0196] todo!()

[0197] }

[0198] }

[0199] In addition, in order to connect the custom operator to the tensor module of the system of the present invention, the call to the custom operator needs to be completed in the implementation of the tensor, which is further called to the backend layer and the micro-operator layer through the storage module for implementation.

[0200] impl Tensor {

[0201] pub fn custom_forward<T: CustomOp> (&self,

[0202] rhs: &Tensor, c: &T) -> Tensor {

[0203] let (storage, shape) = self.storage()

[0204] .forward(self.layout(), &rhs.storage(), rhs.layout(), c);

[0205] from_storage(storage, shape, false)

[0206] }

[0207] pub fn custom_backward<T: CustomOp> (&self,

[0208] arg1: &Tensor, arg2: &Tensor,

[0209] res: &Tensor, grad: &Tensor, c: &T) -> Tensor {

[0210] let (storage, shape) = self.storage()

[0211] .backward(arg1, arg2, res, grad, c);

[0212] from_storage(storage, shape, false)

[0213] }

[0214] }

[0215] In a specific example, a new custom operator KVConcat is shown below, which mainly completes the concatenation (concat) operation of key and value in the forward calculation process of a large model. The dimension of the concatenation operation is specified by the concatenation dimension (concat_dim).

[0216] pub struct KVConcat {

[0217] pub concat_dim: i32,

[0218] }

[0219] pub fn kvconcat(lhs: &Tensor, rhs: &Tensor, dim: i32) -> Tensor {

[0220] let op = KVConcat { concat_dim: dim};

[0221] lhs.custom_forward(&rhs, op)

[0222] }

[0223] According to KVConcat, the kvconcat function can be further defined, which accepts two tensor inputs and a concatenation dimension (concat_dim). The function finally calls the custom_forward function interface of the tensor module, that is, the custom_forward function interface of the tensor module mentioned above. After the new operator completes the function and interface definition, the call to the new operator calculation kernel (kvconcat kernel launch) can be implemented in the corresponding backend layer, as shown below:

[0224] impl CustomOp for KVConcat {

[0225] fn name(&self) -> String { "kvconcat".to_string()}

[0226] fn forward(&self, l: &GcuStorage, l_l: &Layout,

[0227] r: &GcuStorage, r_l: &Layout) -> (GcuStorage, Shape) {

[0228] let dev = &l.device;

[0229] let cfg = &dev.launch_cfg;

[0230] let elem_count =

[0231] l_l.shape().elem_count() + r_l.shape().elem_count();

[0232] let dims = l_l.shape().dims().len();

[0233] let ds = dev

[0234] .htod_copy([l_l.shape().dims(), r_l.shape().dims()]

[0235] .concat());

[0236] let slice = match (&l.slice, &r.slice) {

[0237] (GcuStorageSlice::BF16(l_), GcuStorageSlice::BF16(r_)) => {

[0238] let out = dev.alloc:: <bf16>(elem_count);

[0239] let func = dev.load_func("kvconcat_bf16");

[0240] let params = ( l_.device_ptr(), r_.device_ptr(),

[0241] out.device_ptr(), ds.device_ptr(),

[0242] dims, self.concat_dim);

[0243] func.launch(cfg, params);

[0244] GcuStorageSlice::BF16(out)

[0245] }

[0246] (GcuStorageSlice::F16(l_), GcuStorageSlice::F16(r_)) => {

[0247] let out = dev.alloc:: <f16>(elem_count);

[0248] let func = dev.load_func("kvconcat_f16");

[0249] let params = (l_.device_ptr(), r_.device_ptr(),

[0250] out.device_ptr(), ds.device_ptr(),

[0251] dims, self.concat_dim);

[0252] func.launch(cfg, params);

[0253] GcuStorageSlice::F16(out)

[0254] }

[0255] _ => InternalError("dtype mismatch in kvconcat op"),

[0256] };

[0257] let dst_shape = concat_shape(l_l, r_l, self.concat_dim);

[0258] (GcuStorage { slice, device: dev.clone()}, dst_shape.into())

[0259] }

[0260] }

[0261] As shown above, in the implementation of the forward calculation (forward) function, the kernel call parameters are first organized, and then the kernel instance (kernel function) to be loaded is determined according to the data type of the tensor. For example, for the input tensor of FP16, the kernel kvconcat_f16 is loaded, and finally the kernel startup function of the device module is used to complete the call and return the calculation result. It should be pointed out that the input parameters (parameter type, order, and quantity) in the kernel need to be consistent with the written kernel interface. The main implementation logic of the kernel Kvconcat is to extract the corresponding slices of the two input tensors according to the specified splicing dimension and copy them to the specified area of ​​the output tensor through the device copy (d2d copy) operation.

[0262] 6) Neural network module: defines common large model operators. For example, linear layer operator, activation layer operator, normalization operator, rotational position encoding operator, attention operator, etc. There are two forms of operator implementation of neural network module. One is to write directly using micro-operators or micro-operator combinations, such as normalization operator or linear layer operator. The implementation of linear layer operator (as shown below) calls matrix multiplication operator, transposition operator and addition operator. The implementation of other neural network layers is similar to this, that is, they all need to implement the Forward interface. For the inplace operation, the InPlaceForward forward computing interface needs to be implemented, and its forward computing function directly modifies the input tensor without returning the modified copy. It should be pointed out that the operator or neural network layer defined by the neural network module can use the operator function provided in the tensor module, that is, there is no need to explicitly specify the backend, and the corresponding backend call is determined by the device where the tensor is located. At the same time, it can also directly load and execute the computing kernel of the specified hardware platform in the form of computing kernel startup, similar to the custom operator implementation.

[0263] pub struct Linear {

[0264] weight: Tensor,

[0265] bias: Option <tensor>,

[0266] }

[0267] impl Forward for Linear {

[0268] pub fn forward(&self, x: &Tensor) -> Tensor {

[0269] let x = x.matmul(&self.weight.transpose(1, 2));

[0270] match &self.bias {

[0271] Some(b) => x.add(&b),

[0272] None => x,

[0273] }

[0274] }

[0275] }

[0276] 7) Weight quantization module: It is an extension of the tensor module. The standard tensor module supports common data types such as FP16, BF16, FP32, etc. Lower precision data types such as INT8, INT4, FP6, etc. are often used to accelerate large model reasoning. In order to support the loading of low-precision model weights (or quantized weights) and low-precision reasoning of large models, the weight quantization module needs to implement the following functions: First, the weight loading module needs to support the loading of low-precision weights. Common large model low-precision weight file formats are GGML / GGUF and GPTQ. The GGML / GGUF file format is divided into Q4 (Q4_0, Q4_1, and Q4_K file formats), Q5 (Q5_0, Q5_1, and Q5_K file formats), Q8 (Q8_0, Q8_1, and Q8_K file formats) according to the quantization bit width. Therefore, the weight loading module needs to implement the dequantization of quantized weights. Secondly, the computing kernel (usually a matrix multiplication operator, that is, quantizing the weights of the matrix multiplication operator or the full linear layer operator) needs to support low-precision matrix operations or support real-time dequantization of low-precision weights (real-time dequantization of quantized weights to higher precision before high-precision matrix operations). Taking the file format Q8_0 in the common file format GGML / GGUF as an example, the Rust language definition of its data format is:

[0277] #[repr(C)]

[0278] pub struct Q8_0 {

[0279] pub scale: f16,

[0280] pub chunk: [i8; 32],

[0281] }

[0282] That is, a Q8_0 data contains a FP16 quantization scale factor member and 32 INT8 elements. The quantization process of Q8_0 is to quantize and scale a given high-precision tensor (32 elements need to be aligned) to the INT8 representation range (between -127 and +127) for each 32 elements as a group (chunk), where the scale factor is the scale member in Q8_0 (such as 8.0 if it is reduced by 8 times). In this way, only 32*1 bytes + 2 bytes = 34 bytes are needed to represent 32 high-precision values ​​(32*4 bytes = 128 bytes, taking FP32 as an example), which greatly saves the storage of weights and the occupation of device memory during inference.

[0283] Similarly, the weights can be dequantized during the loading process of quantized weights, or real-time dequantization can be performed during inference. The dequantization process is as follows: for a given low-precision tensor (such as Q8_0 format), every 34 bytes are parsed into a Q8_0 data, where the first two bytes represent the quantization scale, and the last 32 bytes are a group (chunk) of INT8 quantized elements. Each quantized element is first converted to a high-precision representation, and then multiplied by the scaling factor scale to restore the quantized element to a high-precision value (dequantization). The dequantized data can be used for normal matrix operations.

[0284] The inference process usually has high requirements for real-time dequantization speed. To speed up the dequantization speed, the hardware platform programming language can be used to write dequantization operators (micro operators, computing cores), making full use of the parallel execution capabilities of the hardware platform to improve the inference speed. The following is a pseudo-code implementation example of the Q8_0 format dequantization micro operator implemented on the Suiyuan GCU platform.

[0285] #define CHUNK_SIZE 32

[0286] typedef struct {

[0287] __fp16 scale;

[0288] int8_t chunk[CHUNK_SIZE];

[0289] } Q8_0;

[0290] template<typename T>

[0291] __device__ void dequantize_q8_0(const void *vx,

[0292] const int idx, T* out) {

[0293] const Q8_0* x = (const Q8_0*) vx;

[0294] const T scale = T(x[idx].scale);

[0295] auto chunk = vcast<T, uint8_t>(x[idx].chunk);

[0296] auto deq = vmul_scalar(chunk, scale);

[0297] vstore(deq, out + idx * CHUNK_SIZE);

[0298] }

[0299] #define DEQUANTIZE_OP(DST_TYPE, FN_NAME) \

[0300] extern "C" __global__ FN_NAME( \

[0301] const void* vx, DST_TYPE* out, \

[0302] const size_t elements) { \

[0303] size_t chunks = elements / CHUNK_SIZE; \

[0304] for (unsigned int i = \

[0305] blockIdx.x * blockDim.x + threadIdx.x; \

[0306] i < chunks; i += blockDim.x * gridDim.x) { \

[0307] dequantize_q8_0<DST_TYPE> (vx, i, out); \

[0308] } \

[0309] } \

[0310] DEQUANTIZE_OP(float, dequantize_q8_0_f32)

[0311] DEQUANTIZE_OP(__fp16, dequantize_q8_0_f16)

[0312] DEQUANTIZE_OP(__bf16, dequantize_q8_0_bf16)

[0313] Among them, first define the data structure Q8_0, and then define the dequantization function dequantize_q8_0, which is mainly used to dequantize the Q8_0 data of the specified group (chunk) index and store it in the specified position of the output data. The dequantization operators dequantize_q8_0_f32, dequantize_q8_0_fp16, and dequantize_q8_0_bf16 are three dequantization operators exposed to the upper layer to process inputs of different data types. The dequantization function dequantize_q8_0 is called according to the number of threads started, and the input data vx (every 34 bytes is a Q8_0 data) is dequantized in turn. Generics and template functions are used here to implement the dequantization of Q8_0 data into FP32, FP16, and BF16 data.

[0314] 8) Backward propagation module: mainly used for fine-tuning large models. The implementation of the backpropagation module requires the support of the tensor module, that is, the operation type and parameters are recorded during the tensor operation process, and the operations and corresponding parameters applied to the tensor are recorded in each tensor, so that the gradient of each step can be calculated backward from the output layer loss according to the chain rule during the reverse gradient calculation. The backpropagation module provides a backward calculation interface for the tensor. By calling the backward calculation interface of the tensor, the chain derivation is triggered, that is, first topological sorting is performed according to the operation sequence recorded by the tensor, and then the operation sequence is traversed in reverse order. The reverse calculation interface of each sequence element is called according to the output loss. After the gradient is calculated, it is passed down in sequence until the gradient of the specified level is calculated. Finally, the optimizer is used to adjust the weight parameters of the specified level to achieve model fine-tuning. The specific implementation of the backpropagation module is similar to the standard reverse calculation, and its implementation details are not repeated here. It should be pointed out that reverse calculation can also use a combination of micro-operators to reduce the difficulty of developing reverse operators. For example, the activation function SiLU, which is common in large models, can use the following combination of micro-operators for reverse calculation, which uses single-operand operators and double-operand operators, such as negative operators, exponential operators, reciprocal operators (recip), multiplication operators, addition operators, subtraction operators, etc.

[0315] fn silu_grad(&x: Tensor) -> Tensor {

[0316] let sigmoid_x = (x.neg().exp() + 1.).recip();

[0317] let grad = &sigmoid_x * (1 + x * (1 - &sigmoid_x));

[0318] return grad;

[0319] }

[0320] 9) Model layer: As mentioned above, the model layer uses various functional modules provided by the framework layer to build a large model instance based on micro-operators, where each network layer of the large model can be combined using the tensor module, neural network module (Layer module) and custom operator module of the framework layer, and the weights of each network layer can be loaded using the weight loading and quantization module of the framework layer.

[0321] Taking the mainstream large model network structure LLaMa as an example, its model contains multiple blocks, each of which consists of RmsNorm, Attention and MLP layers. The following is an example of writing a block, where the forward function is responsible for the forward calculation of the current block. Its writing method refers to the official example of LLaMa, and the load function is responsible for loading the current block weight (i.e., the weight of RmsNorm, Attention and MLP).

[0322] struct Block {

[0323] norm_1: RmsNorm,

[0324] norm_2: RmsNorm,

[0325] attention: Attention,

[0326] mlp: Mlp,

[0327] }

[0328] impl Block {

[0329] fn forward(&mut self, x: &Tensor, pos: usize,

[0330] layer_idx: usize,cache: &mut KVCache) -> Tensor {

[0331] let skip = x;

[0332] let x = self.norm_1.forward(x);

[0333] let x = (self.

[0334] attention.forward(&x, pos, layer_idx, cache) + skip);

[0335] let skip = &x;

[0336] self.mlp.forward(&self.norm_2.forward(&x)) + skip

[0337] }

[0338] fn load(ld: &Loader, w: &str, cfg: &Config) -> Block {

[0339] let ld = ld.sub(&w);

[0340] Block {

[0341] norm_1: RmsNorm::load(&ld, "input_layernorm", cfg),

[0342] norm_2: RmsNorm::load(&ld, "post_attention_layernorm", cfg),

[0343] attention: Attention::load(&ld, "self_attn", cfg),

[0344] mlp: Mlp::load(&ld, "mlp", cfg),

[0345] }

[0346] }

[0347] }

[0348] The RmsNorm module in Block is a simplified LayerNorm, which can be written in a custom operator or aggregated operator mode; and the Attention and MLP modules in Block can also be written in a similar way. For example, the structure of the attention layer is defined as follows, which includes the head, kv_heads and head_dim used to record the multi-head attention layer, and four Linear layers for projecting the input and output data:

[0349] struct Attention {

[0350] heads: usize,

[0351] kv_heads: usize,

[0352] head_dim: usize,

[0353] q_linear: Linear,

[0354] k_linear: Linear,

[0355] v_linear: Linear,

[0356] o_linear: Linear,

[0357] }

[0358] The forward calculation of the attention layer operator is as follows, which includes: first, the input tensor is projected with a linear layer of query vector (q), key vector (k), and value vector (v), and then the projected q, k, and v tensors are transformed in shape to adapt to the subsequent multi-headed attention calculation, followed by rotary embedding (rope) and key-value cache (kv cache) caching, and scaled dot-productattention (two matmul and softmax combinations) are used to calculate the attention value, and finally the output is transformed in shape and the output tensor is projected in a linear layer to obtain the final output of the attention layer operator.

[0359] impl Forward for Attention {

[0360] fn forward(&self, x: &Tensor, index: usize,

[0361] layer_idx: usize, cache: &mut Cache) -> Tensor {

[0362] let (batch, seq_len, hidden_size) = x.dims();

[0363] let q = self.q_linear.forward(x);

[0364] let k = self.k_linear.forward(x);

[0365] let v = self.v_linear.forward(x);

[0366] let q = q.reshape((batch, seq_len, self.heads, self.dim))

[0367] .transpose(1, 2);

[0368] let k = k.reshape((batch, seq_len, self.kv_heads, self.dim))

[0369] .transpose(1, 2);

[0370] let v = v.reshape((batch, seq_len, self.kv_heads, self.dim))

[0371] .transpose(1, 2);

[0372] let (q, mut k) = nn::rope(&q, &k, &cache, index);

[0373] if let Some((k_cache, v_cache)) = &cache.kvcache[layer_idx] {

[0374] k = nn::kvconcat(k_cache, &k, 2);

[0375] v = nn::kvconcat(v_cache, &v, 2);

[0376] }

[0377] cache.kvcache[layer_idx] = Some((k.clone(), v.clone()));

[0378] let k = nn::repeat(k);

[0379] let v = nn::repeat(v);

[0380] let attn = (q.matmul(&k.transpose()) / self.dim.sqrt());

[0381] let attn = match &cache.mask {

[0382] Some(mask) => {

[0383] nn::softmax(attn + mask)

[0384] }

[0385] _=> nn::softmax(attn)

[0386] };

[0387] let out = attn.matmul(&v).transpose(1, 2);

[0388] let out = out.reshape(&[batch, seq_len, hidden_size]);

[0389] self.o_linear.forward(&out)

[0390] }

[0391] }

[0392] The weight loading of the attention layer operator is similar to that of the module layer (Block). It uses the weight loading module to load the corresponding weights for the input and output projection layers, as shown below:

[0393] impl Attention {

[0394] fn load(ld: Loader, w: &str, c: &Config) -> Attention {

[0395] let ld = ld.sub(&w);

[0396] let h_sz = c.hidden_size;

[0397] let i_sz = c.inner_size;

[0398] let q_sz = (h_sz / c.heads) * c.heads;

[0399] let kv_sz = (h_sz / c.heads) * c.kv_heads;

[0400] Attention

[0401] heads: c.heads,

[0402] kv_heads: c.kv_heads,

[0403] head_dim: c.hidden_size / c.kv_heads,

[0404] q_linear: Linear::load(i_sz, q_sz, &ld, "q_proj"),

[0405] k_linear: Linear::load(i_sz, kv_sz, &ld, "k_proj"),

[0406] v_linear: Linear::load(i_sz, kv_sz, &ld, "v_proj"),

[0407] o_linear: Linear::load(q_sz, i_sz, &ld, "o_proj"),

[0408] }

[0409] }

[0410] }

[0411] After completing the writing of the attention layer operator, you can continue to write the multi-level perceptron layer (MLP). As shown below, the load function is responsible for loading the weights fc_1 and fc_2 required for the input linear projection, as well as the output projection weight fc_o. The forward calculation function first performs a linear projection of the input fc_1, then performs a nonlinear activation (silu activation) on the projected tensor, and multiplies it with the tensor result of the linear projection of fc_2, and finally performs an output projection to obtain the final output of the multi-level perceptron layer.

[0412] struct Mlp {

[0413] fc_1: Linear,

[0414] fc_2: Linear,

[0415] fc_o: Linear,

[0416] }

[0417] impl Mlp {

[0418] fn forward(&self, x: &Tensor) -> Tensor {

[0419] let x1 = self.fc_1.forward(x);

[0420] let x = nn::silu(&x1) * self.fc_2.forward(x);

[0421] self.fc_o.forward(&x)

[0422] }

[0423] fn load(ld: Loader, w: &str, cfg: &Config) -> Mlp {

[0424] let ld = ld.sub(&w);

[0425] let h_sz = cfg.hidden_size;

[0426] let i_sz = cfg.inner_size;

[0427] Mlp {

[0428] fc_1: Linear::load(h_sz, i_sz, &ld, "gate_proj"),

[0429] fc_2: Linear::load(h_sz, i_sz, &ld, "up_proj"),

[0430] fc_o: Linear::load(i_sz, h_sz, &ld, "down_proj"),

[0431] }

[0432] }

[0433] }

[0434] After writing the normalization operator, attention layer operator, multi-level perceptron layer and other modules, the model layer specification shown in the present invention can be used to write the mainstream large model LLaMa, as shown below. In the forward calculation function of LLaMa, the word vector encoding module is first used to encode the input tensor, and then the forward calculation function of each module is called in turn according to the number of model layers of LLaMa (different numbers of layers correspond to LLaMa models of different sizes, such as 7B, 13B, 70B) to perform forward calculations, and then the module output is normalized (RmsNorm normalization), and the slice operator is used to obtain the output result of the corresponding slice, and finally the output layer is linearly projected to obtain the final prediction result of LLaMa. LLaMa output also needs to undergo word vector encoding and output collection to obtain the text result, also known as token. Correspondingly, LLaMa's load function is responsible for loading the word vector encoding weights, the weights of each module layer of the model, the weights of the normalization operator, and the weights of the output projection layer.

[0435] pub struct LlamaModel {

[0436] emb: Embedding,

[0437] blocks: Vec <block>,

[0438] norm: RmsNorm,

[0439] head: Linear,

[0440] }

[0441] impl Forward for LlamaModel {

[0442] pub fn forward(&self, x: &Tensor, index: usize,

[0443] cache: &mut Cache) -> Tensor {

[0444] let (batch, seq_len) = x.dims();

[0445] let mut x = self.emb.forward(x);

[0446] for (idx, block) in self.blocks.iter().enumerate() {

[0447] x = block.forward(&x, index, idx, cache)?;

[0448] }

[0449] let x = self.norm.forward(&x);

[0450] let x = x.slice(seq_len - 1, 1);

[0451] self.head.forward(&x)

[0452] }

[0453] }

[0454] impl LlamaModel {

[0455] pub fn load(ld: Loader, w: &str, cfg: &Config) -> LlamaModel {

[0456] let ld = ld.sub(w);

[0457] let emb = Embedding::load(cfg.vocab_size, cfg.hidden_size,

[0458] &ld, "model.embed_tokens");

[0459] let blocks: Vec<_> = (0..cfg.num_layers)

[0460] .map(|i| Block::load(&ld, &format!("model.layers.{i}"), &cfg))

[0461] .collect();

[0462] let norm = RmsNorm::load(cfg.hidden_size, cfg.eps,

[0463] &ld, "model.norm");

[0464] let head = Linear::load(cfg.hidden_size, cfg.vocab_size,

[0465] &ld, "lm_head");

[0466] LlamaModel { emb, blocks, norm, head}

[0467] }

[0468] }

[0469] For fine-tuning of large models, it is generally necessary to add a small number of linear layers (low-rank weights) to the trained model. The reverse calculation micro-operators or micro-operator combinations of the corresponding linear layers can be written according to the back-propagation module in the framework layer, and the backward calculation function can be implemented for the large model (such as LLaMaModel). Among them, the backward calculation function calculates the gradients of the newly added linear layer according to the input loss and returns the gradient tensor. The returned gradient tensor is called by the back-propagation module in the framework layer to update the weights of the linear layer to achieve the purpose of fine-tuning the model weights. The backward calculation function of model fine-tuning is written in a similar way to the forward calculation function.

[0470] The embodiment of the present invention provides a large model cross-platform system based on micro-operators, which can be used to quickly perform reasoning and fine-tuning of large models on different hardware platforms, and is suitable for rapid development, debugging, reasoning and fine-tuning of large models on different hardware platforms, which simplifies the difficulty of large model operator development, has a simple system structure, and is conducive to model development, debugging and cross-platform deployment. It defines the hardware access method, micro-operator implementation form and calling method, and various functional modules required for large model reasoning and fine-tuning based on micro-operators. Its overall system structure design is simple, which can effectively reduce the difficulty of operator development and improve the efficiency of large model development and deployment.

[0471] It should be understood that the various forms of processes shown above can be used to reorder, add or delete steps. For example, the steps described in the present invention can be executed in parallel, sequentially or in different orders, as long as the desired results of the technical solution of the present invention can be achieved, and this document does not limit this.

[0472] The above specific implementations do not constitute a limitation on the protection scope of the present invention. It should be understood by those skilled in the art that various modifications, combinations, sub-combinations and substitutions can be made according to design requirements and other factors. Any modification, equivalent substitution and improvement made within the spirit and principle of the present invention should be included in the protection scope of the present invention.< / block> < / tensor> < / tensor> < / tensor> < / t> < / t> < / t> < / t> < / t>

Claims

1. A large model cross-platform system based on micro-operators, characterized in that: include: Hardware execution layer, hardware access layer, precompilation layer, micro-operator layer, backend layer, framework layer and model layer; The hardware execution layer is used to execute memory access instructions and / or various micro-operators compiled by the pre-compilation layer; The hardware access layer is used to perform high-level abstraction on various interfaces of various hardware devices connected to the large model cross-platform system, and form a consistent calling interface for the upper layer; The pre-compilation layer is used to call the platform compiler to compile the micro-operator code into intermediate instructions or executable instructions of the hardware device; The micro-operator layer is used to call various micro-operators compiled by the pre-compiled layer and provide a micro-operator-based programming interface to the upper layer, so that the upper layer can combine and call multiple micro-operators to realize various operators required by the mainstream large model; The backend layer is connected to the framework layer upwards and to the micro-operator layer downwards, so as to expose the functional interface of the micro-operator layer to the outside. Meanwhile, the backend layer can identify the tensor data type, and the different data types of the same micro-operator compiled by the pre-compiled layer are implemented by the micro-operator layer; The framework layer includes various functional modules required for realizing large model reasoning and fine-tuning, including tensor module, storage module, device module, weight loading module, custom operator module, neural network module, weight quantization module, word vector encoding module, back propagation module and output sampling module. Each functional module realizes the corresponding large model reasoning function or large model fine-tuning function by calling the functional interface in the backend layer; The model layer is used to provide users with model layer specifications for building large model instances based on the framework layer; wherein, providing users with model layer specifications for building large model instances based on the framework layer includes: using the tensor module, neural network module and custom operator module provided by the framework layer to build a large model instance based on micro-operators according to the mainstream large model structure; using the weight loading module provided by the framework layer to load the large model weights and calling the weight quantization module provided by the framework layer to quantize and dequantize the large model weights according to model requirements; calling the functional modules in the framework layer to perform reasoning and fine-tuning of the large model.

2. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: When the large model cross-platform system is compiled into a binary program, the pre-compiled layer calls the platform compiler to compile the pre-written micro-operator code into intermediate instructions or executable instructions of the hardware device; The pre-compilation layer is also used to provide a cache function for pre-compilation instructions.

3. The large-model cross-platform system based on micro-operators according to claim 1 is characterized in that: The micro-operator layer includes device memory access management micro-operators, tensor calculation micro-operators, and various operators required by mainstream large models realized by combining tensor calculation micro-operators; The device memory access management micro-operator includes a device memory allocation micro-operator, a device memory release micro-operator, a host-to-device communication micro-operator, a device-to-host communication micro-operator, a device-to-device communication micro-operator, a device synchronization micro-operator, and a device query micro-operator; The tensor computing micro-operators include single-operand operators, double-operand operators, type conversion operators, index operators, general copy operators, conditional selection operators, aggregation operators, and matrix multiplication operators; The tensor calculation micro-operators are combined to realize various operators required by the mainstream large model, including rotational position encoding operators, activation layer operators, normalization operators, standardization operators and attention layer operators.

4. The large-model cross-platform system based on micro-operators according to claim 1 is characterized in that: The backend layer parses the interface call parameters of the framework layer and identifies the parameter types, loads the micro-operators compiled by the pre-compiled layer in the form of computing core startup, starts execution through the hardware execution layer, and encapsulates and returns the calculation results.

5. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a tensor module; the tensor module is an abstract representation of a tensor in a neural network, and the tensor module includes tensor data and various common tensor operations, and various tensor operations are implemented by calling the functional interface of the corresponding micro-operator in the backend layer.

6. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a storage module; the storage module is a lower-level representation of tensors and tensor operations, and is used to identify the device category and data type of the tensor. The operation of identifying the device category and data type of the tensor is implemented by calling the functional interface of the corresponding micro-operator in the backend layer.

7. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a device module; the device module is a further abstraction of the device functions provided by the backend layer, and is used to provide the functional interface required for memory access operations and micro-operator execution operations on different hardware devices. The memory access operations include the operations required for allocating, recycling and migrating tensors between devices on different devices.

8. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a weight loading module; the weight loading module is used to load the large model weight from a file; wherein the weight loading process is triggered by the model layer, and the weight loading module sets the target data type of the large model weight when creating the loader, and converts the data type of the large model weight during the weight loading process, converting the data type of the large model weight to the target data type.

9. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a custom operator module; the custom operator module is used to provide a functional interface for an extended micro-operator; wherein the functional interface of the extended micro-operator includes an extended interface for a single-operand operator, a double-operand operator, and a multi-operand operator, and a non-standard micro-operator or an aggregation operator is implemented through the functional interface of the extended micro-operator.

10. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a neural network module; the neural network module defines common large model operators, and the implementation form of common large model operators is to use micro operators or micro operator combinations for writing, or to call aggregation operators for writing through the custom operator module.

11. The large-model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a weight quantization module; the weight quantization module is an extension of the tensor module. In addition to providing the standard operations defined by the tensor module, it also provides quantization and dequantization function interfaces to support low-precision large model reasoning and fine-tuning.

12. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a word vector encoding module; the word vector encoding module is used to implement the word vector encoding operation required for large model reasoning or large model fine-tuning by calling the micro-operator in the micro-operator layer.

13. The large model cross-platform system based on micro-operators according to claim 1 is characterized in that: The framework layer includes a back propagation module; the back propagation module is used to implement fine-tuning of large models, the back propagation module relies on the tensor module to record the operation type and parameters during the tensor operation, and calculates the reverse gradient through the reverse calculation interface of the tensor module; The reverse gradient calculation operators corresponding to various tensor operations are defined in the reverse propagation module. The reverse gradient calculation operators corresponding to various tensor operations are implemented by combining various micro-operators or by calling aggregation operators through the custom operator module.

Images