Automatic operator optimization method, electronic equipment and storage medium

Abstract

The invention provides an operator automatic optimization method, electronic equipment and a storage medium, and relates to the technical field of artificial intelligence. The method comprises the following steps: acquiring a to-be-optimized parameter set of a to-be-optimized operator; generating a plurality of parameter value combinations based on the candidate values of the parameters in the to-be-optimized parameter set; aiming at each parameter value combination, loading the parameter value combination into a parameterization template of the operator to be optimized, and generating an executable code on the target hardware platform; the execution performance of each executable code is evaluated, and a target parameter combination is determined from the parameter value combinations according to the execution performance; and generating a target code executed by the operator to be optimized on the target hardware platform based on the target parameter combination. In the implementation mode, high-performance operator codes fusing inverse quantization and matrix vector multiplication calculation are generated for different GPU platforms without manual intervention, and operator development efficiency and cross-platform execution performance are remarkably improved.

Description

Technical Field

[0001] This application relates to the field of artificial intelligence technology, and in particular to an automatic operator optimization method, electronic device, and storage medium. Background Technology

[0002] With the rapid development of artificial intelligence technology, large-scale language models (LLMs) have been widely used in the field of natural language processing, such as intelligent dialogue, text generation, and code assistance. The main bottleneck in the inference performance of this large-scale language model lies in the core computational operators in the decoding stage. The execution efficiency of these operators directly affects the response speed and user experience of the large-scale language model.

[0003] To improve the execution efficiency of the aforementioned operators, related technologies mainly employ two optimization methods: manual optimization and automatic optimization. However, both have significant drawbacks. For example, manual optimization is heavily tied to the graphics processing unit (GPU) hardware platform, heavily reliant on engineers' experience, resulting in long development cycles, high costs, and poor portability. Automatic optimization, on the other hand, has a coarse-grained design space and fails to fully consider the microarchitectural differences between different GPUs, thus limiting its optimization effectiveness. Therefore, a new operator optimization method is urgently needed to address these issues. Summary of the Invention

[0004] This application provides an automatic operator optimization method, electronic device, and storage medium. The method constructs a set of parameters to be tuned containing three types of parameters and uses an automated traversal and evaluation method to generate high-performance operator code that integrates inverse quantization and matrix-vector multiplication calculations for different GPU platforms without manual intervention, which significantly improves operator development efficiency and cross-platform execution performance.

[0005] To achieve the above objectives, this application adopts the following technical solution: In a first aspect, embodiments of this application provide an automatic operator optimization method, the method comprising: obtaining a set of parameters to be tuned for an operator to be optimized; generating multiple parameter value combinations for the operator to be optimized based on candidate values ​​of each parameter in the set of parameters to be tuned; loading each parameter value combination in the multiple parameter value combinations into a parameterized template of the operator to be optimized, thereby generating executable code for the operator to be optimized on a target hardware platform; evaluating the execution performance of each executable code, and determining a target parameter combination from the multiple parameter value combinations based on the execution performance; and generating target code for the operator to be optimized to be executed on the target hardware platform based on the target parameter combination.

[0006] Optionally, the parameter set to be tuned is a collection of all parameters that can be adjusted during the automatic optimization process and can directly change the operator's implementation logic and runtime performance. The parameter set to be tuned includes a first type of parameters related to computational parallel scheduling, a second type of parameters related to the execution unit of the target hardware platform, and a third type of parameters related to the quantization and dequantization data formats of the operator to be optimized. The quantization and dequantization data formats are used to characterize the storage format of the quantization weight matrix of the matrix-vector multiplication operator.

[0007] Optionally, the first type of parameters includes at least one or more of the following: recursive axis parallelism parameter, non-recursive axis parallelism parameter, loop unrolling dimension parameter, memory cache enable parameter, vector computation width parameter, thread block size, data layout format, and global memory access alignment mark.

[0008] Optionally, the second type of parameters includes at least one or more of the following: thread bundle or subgroup or wavefront size parameter, data aggregation loading width parameter, single-cycle SIMD instruction width, maximum available shared memory capacity threshold, thread scheduling priority configuration, and global memory bandwidth matching coefficient. The data aggregation loading width parameter is determined based on the product of the thread bundle or subgroup size parameter and the quantization data packing bit width parameter, such as data aggregation loading width = thread bundle / subgroup size × quantization data packing bit width.

[0009] Optionally, the third type of parameters includes at least one or more of the following: quantization weight data type parameter, quantization weight data bit width parameter, quantization data packing bit width parameter, quantization group size, zero-point offset storage format, and dequantization intermediate data precision.

[0010] Optionally, the quantization weight data type parameter may include signed integer and unsigned integer, the quantization weight data bit width parameter may include 2 bits, 3 bits or 4 bits, and the quantization data packing bit width parameter may include 8 bits, 16 bits, 32 bits or 64 bits, etc.

[0011] Optionally, the target hardware platform refers to the hardware environment in which the operator ultimately runs, which may include a GPU hardware platform. The value of the thread bundle or subgroup size parameter is determined according to the GPU architecture type; specifically, the value of the thread bundle or subgroup size parameter is 32 for the first type of GPU architecture (such as Nvidia GPU), 8, 16, or 32 for the second type of GPU architecture (such as Intel GPU), and 64 for the third type of GPU architecture (such as AMD GPU).

[0012] Optionally, the operators to be optimized may include GEMV operators, GEMM operators, batch normalization operators, layer normalization operators, activation function operators, pooling operators, etc.

[0013] In this implementation, a set of parameters to be tuned is constructed, which includes parameters related to computational parallel scheduling, parameters related to hardware platform execution units, and parameters related to quantization and dequantization data formats. An automated search method is used to generate executable code that is compatible with the target hardware platform. This improves the efficiency of operator development and achieves high-performance cross-platform execution.

[0014] Optionally, in one possible implementation, after generating multiple combinations of parameter values ​​for the operator to be optimized based on the candidate values ​​of each parameter in the parameter set to be tuned, the method further includes: eliminating parameter value combinations that do not meet the hardware constraints from the multiple parameter value combinations according to the hardware constraints of the target hardware platform.

[0015] Optionally, the hardware constraints include at least one or more of the following: the limit on the execution unit size of the target hardware platform, the upper limit on the total computing resources of the target hardware platform, and the relationship between the parallelism of the operator to be optimized and the thread bundle or subgroup size parameter. The parallelism of the operator to be optimized is an integer multiple of the thread bundle or subgroup size parameter, and the parallelism of the operator to be optimized is determined by the product of the recursive axis parallelism parameter and the non-recursive axis parallelism parameter in the first type of parameters.

[0016] This implementation method eliminates invalid parameter value combinations in advance, effectively reducing the number of parameter value combinations that need to be evaluated subsequently, and significantly improving the overall optimization efficiency; it ensures that all parameter value combinations evaluated subsequently are within the range allowed by hardware constraints, avoiding compilation errors or runtime anomalies caused by invalid parameter value combinations; at the same time, it integrates the understanding of hardware characteristics into the automated process, realizing hardware-aware optimization preprocessing, and creating good input conditions for the automated search of subsequent modules.

[0017] Optionally, in one possible implementation, the parameter value combination is loaded into the parameterized template of the operator to be optimized, generating executable code of the operator to be optimized on the target hardware platform, including: assigning each parameter value in the parameter value combination to the corresponding variable in the parameterized template to obtain an instantiated operator template; and compiling the instantiated operator template into executable code adapted to the target hardware platform according to a compilation backend that matches the target hardware platform.

[0018] Optionally, when the target hardware platform is a first type of GPU architecture, the compilation backend is CUDA; when the target hardware platform is a second type of GPU architecture, the compilation backend is OpenCL; and when the target hardware platform is a third type of GPU architecture, the compilation backend is Rocm.

[0019] Optionally, CUDA code (.ptx or cubin files) is generated on Nvidia GPU platforms, OpenCL code (.cl files) is generated on Intel GPU platforms, and Rocm code (or HIP executable code) is generated on AMD GPU platforms.

[0020] This implementation method achieves one-click cross-platform generation of operator code by identifying the target GPU architecture type and automatically matching the corresponding compilation backend. Developers do not need to manually write low-level code for different GPUs to obtain executable files that are deeply adapted to the hardware, which significantly reduces the technical threshold and development cost of cross-platform deployment, while ensuring the execution efficiency of the generated code on their respective hardware platforms.

[0021] Optionally, in one possible implementation, each parameter value in the parameter value combination is assigned to the corresponding variable of the parameterized template one by one to obtain the instantiated operator template, including: loading the first type of parameters and the second type of parameters in the parameter value combination into the corresponding variable of the parameterized template; selecting the target dequantization logic branch from the predefined dequantization logic branches of the parameterized template according to the third type of parameters in the parameter value combination; wherein the predefined dequantization logic branches include signed type dequantization logic branches and unsigned type dequantization logic branches; and obtaining the instantiated operator template based on the loaded first type of parameters and the second type of parameters, and the target dequantization logic branch.

[0022] In this implementation, the dequantization operation is built into the parameterized template as a configurable logical branch and dynamically selected by a third type of parameter. This allows for deep integration of the dequantization process and matrix-vector multiplication at the code generation level, avoiding the data round-trip overhead of traditional solutions and significantly improving memory access efficiency and computational pipeline utilization. Simultaneously, this approach enables operators to flexibly adapt to different quantization algorithms, enhancing the model's deployment flexibility.

[0023] Optionally, in one possible implementation, the execution performance of each executable code is evaluated, and a target parameter combination is determined from multiple parameter value combinations based on the execution performance. This includes: sequentially evaluating the performance of the executable code corresponding to each parameter value combination to obtain the corresponding execution performance; comparing the execution performance of the current parameter value combination with the execution performance of the retained parameter value combinations; if the execution performance of the current parameter value combination is better, then updating the retained parameter value combinations with the current parameter value combination; and after traversing all parameter value combinations, determining the final retained parameter value combination as the target parameter combination. Optionally, the execution performance includes at least one or more of the following: inference latency, computational utilization, and data transmission efficiency of the operator to be optimized. In this implementation, a pairwise comparison and online update screening strategy is adopted, requiring only one optimal combination and its execution performance index to be stored in memory, without needing to store the evaluation results of all candidate combinations, thus greatly reducing memory overhead.

[0024] Optionally, in one possible implementation, based on the target parameter combination, target code for the operator to be optimized to be executed on the target hardware platform is generated, including: loading the target parameter combination into the parameterized template of the operator to be optimized to obtain the target operator template; compiling the target operator template based on a compilation backend that matches the target hardware platform to generate target code; wherein the target code integrates the weight dequantization operation and the computational logic of the operator to be optimized.

[0025] In this implementation, the optimal parameter combination obtained through automated search is transformed into the final deployed target code, and the dequantization operation is deeply integrated with matrix-vector multiplication calculations in the target code. Each thread directly reads the quantized data according to the bit width of the quantized data during memory access, and immediately participates in multiplication and accumulation calculations after unpacking and dequantization, completely eliminating redundant data movement in traditional solutions. The resulting target code not only has excellent performance but can also be directly deployed across platforms, truly achieving a single search that benefits multiple platforms, significantly improving the efficiency of the entire process of operator optimization from search to deployment.

[0026] Secondly, embodiments of this application provide an operator automatic optimization apparatus, which includes units for performing the method described in the first aspect above. This apparatus can correspond to performing the method described in the first aspect above. For a detailed description of the units within this apparatus, please refer to the description in the first aspect above; for brevity, it will not be repeated here.

[0027] The method described in the first aspect above can be implemented in hardware or by executing corresponding software within the hardware. The hardware or software includes one or more modules or units corresponding to the aforementioned functions. For example, an acquisition module or unit, a generation module or unit, etc.

[0028] Thirdly, embodiments of this application provide an electronic device, which includes a plurality of processors and a memory; the memory is coupled to the plurality of processors, and the memory is used to store computer program code, the computer program code including computer instructions, and the plurality of processors call the computer instructions to cause the electronic device to perform the method in the first aspect above.

[0029] Fourthly, embodiments of this application provide a computer-readable storage medium having a computer program (also referred to as instructions or code) stored thereon for implementing the method in the first aspect. For example, when the computer program is executed by a computer, it enables the computer to perform the method in the first aspect.

[0030] Fifthly, embodiments of this application provide a chip including a processor. The processor is used to read and execute a computer program stored in a memory to perform the methods in the first aspect and any possible implementation thereof. Optionally, the chip further includes a memory connected to the processor via a circuit or wire.

[0031] In a sixth aspect, embodiments of this application provide a computer program product, the computer program product including a computer program (also referred to as instructions or code), which, when executed by an electronic device, causes the electronic device to implement the method in the first aspect.

[0032] It is understood that the beneficial effects of the second to sixth aspects mentioned above can be found in the relevant descriptions in the first aspect mentioned above, and will not be repeated here. Attached Figure Description

[0033] Figure 1 This is a schematic diagram of an algorithm for a matrix-vector multiplication operator, as illustrated in an exemplary embodiment of this application. Figure 2 This is a system architecture diagram illustrating an exemplary embodiment of this application; Figure 3A and Figure 3B This is a schematic diagram illustrating the differences in thread bundle and subgroup sizes under different GPU architectures, as shown in an exemplary embodiment of this application; Figure 4A and Figure 4B This is a schematic diagram illustrating the difference in data aggregation loading width according to an exemplary embodiment of this application; Figure 5A and Figure 5B This is a schematic diagram illustrating the inverse quantization process of quantized data in an exemplary embodiment of this application; Figure 6 This is a schematic diagram illustrating the relationship between the quantization weight data bit width and the quantization data packing bit width, as shown in an exemplary embodiment of this application; Figure 7This is an overall flowchart of the operator automatic optimization method shown in an exemplary embodiment of this application; Figure 8 This is a detailed flowchart illustrating an exemplary embodiment of the operator automatic optimization method of this application; Figure 9 This is a schematic diagram of an inverse quantization and matrix-vector multiplication fusion operator template shown in an exemplary embodiment of this application; Figure 10 This is a flowchart illustrating an operator automatic optimization method according to an exemplary embodiment of this application; Figure 11 This is a schematic diagram illustrating the performance distribution of parameter combinations during the automatic optimization process, as shown in an exemplary embodiment of this application. Figure 12 This is a schematic diagram of the structure of an electronic device illustrated in an exemplary embodiment of this application; Figure 13 This is a software structure block diagram of an electronic device illustrated in an exemplary embodiment of this application; Figure 14 This is a schematic diagram of an operator automatic optimization device illustrated in an exemplary embodiment of this application; Figure 15 This is a schematic diagram of the structure of a chip shown in an exemplary embodiment of this application. Detailed Implementation

[0034] The technical solutions of the embodiments of this application will be clearly described below with reference to the accompanying drawings. Obviously, the described embodiments are only some, not all, of the embodiments of this application. Other embodiments obtained by those skilled in the art based on the embodiments of this application without creative effort are all within the protection scope of this application.

[0035] In the description of this application, unless otherwise stated, " / " means "or," for example, A / B can mean A or B. The "and / or" in this document is merely a description of the relationship between related objects, indicating that three relationships can exist. For example, A and / or B can represent: A alone, A and B simultaneously, and B alone. Furthermore, "at least one" means one or more, and "multiple" means two or more. The terms "first," "second," etc., do not limit the quantity or order of execution, and "first," "second," etc., do not necessarily imply differences.

[0036] It should be noted that, in this application, the terms "exemplary" or "for example" are used to indicate that something is being described as an example, illustration, or illustration. Any embodiment or design described as "exemplary" or "for example" in this application should not be construed as being more preferred or advantageous than other embodiments or design solutions. Specifically, the use of terms such as "exemplary" or "for example" is intended to present the relevant concepts in a concrete manner.

[0037] To facilitate understanding of the technical solutions in the embodiments of this application, some terms used in the embodiments of this application will be explained before introducing the technical solutions in the examples of this application.

[0038] 1. General Matrix-Vector Multiplication (GEMV) Matrix-vector multiplication is a special case of matrix multiplication, taking a row vector and a weight matrix as input and outputting a row vector. In the decoding phase of large-scale language models, GEMV is a core computational operation, and its execution efficiency directly affects the response speed of the large-scale language model. In this embodiment, the GEMV operator, as a type of operator to be optimized, is the object of application of the method in this application. The weight matrix of this GEMV operator is usually stored in quantized form and needs to be dequantized before computation.

[0039] 2. Operator In the field of deep learning, operators are the basic computational units for building neural network models, such as convolution, matrix multiplication, and activation. It should be understood that a deep learning model can be viewed as a computational graph composed of a series of interconnected operators.

[0040] 3. Graphics Processing Unit (GPU) A graphics processing unit (GPU) is a processor specifically designed for parallel computing and can be widely used in the training and inference of deep learning models. In the embodiments of this application, the target hardware platform mainly refers to different types of GPU architectures.

[0041] 4. Quantization and Dequantization Quantization is the process of mapping high-precision data to low-precision data, aiming to reduce model storage and memory bandwidth. Dequantization is the reverse process of quantization, which restores low-precision data to high-precision data for computation.

[0042] The above is a brief introduction to some of the terms involved in the embodiments of this application, and will not be repeated below.

[0043] With the rapid development of artificial intelligence technology, Large Language Models (LLMs) are increasingly widely used in natural language processing fields such as intelligent dialogue, text generation, code assistance, and intelligent question answering. Their performance directly affects the user experience when using LLMs for tasks such as dialogue, question answering, and programming. However, these applications have extremely high real-time requirements, and users expect to obtain a near-instantaneous response experience. Any perceptible latency will seriously affect user satisfaction and product competitiveness.

[0044] Specifically, the inference process of LLM is mainly divided into two stages: prefill and decoding. The prefill stage handles the parallel processing of input prompts, which has high computational complexity but is easily parallelized. The decoding stage, however, requires word-by-word autoregression to generate the output sequence, producing only one token at each time step. Furthermore, the computation at the current time step depends on the output of the previous time step, making it difficult to improve the efficiency of this stage through simple parallelization. This serial generation characteristic makes the decoding stage a key bottleneck restricting the performance of LLM inference, especially when generating long output sequences, where the time consumed by the decoding stage increases dramatically.

[0045] It should be understood that during the decoding phase, the core computational load is highly concentrated on a specific type of core computational operator. The execution efficiency of this operator directly determines the LLM's response speed and user experience. Taking a certain application scenario as an example, this core computational operator can be the General Matrix-Vector Multiplication (GEMV) operator. The GEMV operator is a special case of matrix multiplication; its input is a row vector and a weight matrix, and its output is a row vector. Although this operator has a relatively simple computational pattern, it is frequently called during the LLM decoding phase, meaning that a GEMV operation needs to be performed at every moment.

[0046] During the decoding phase, at each moment, a GEMV operation needs to be performed between the vector corresponding to the currently generated word (i.e., the hidden state vector) and the weight matrix stored in the LLM to calculate the probability distribution of the next word. This calculation process continues throughout every moment of the decoding phase, from the generation of the first word to the generation of the last word. More importantly, the weight matrix in the LLM is extremely large. Taking the current mainstream models with hundreds of billions of parameters as an example, the storage of its weight matrix often requires tens of gigabytes or even hundreds of gigabytes of GPU memory. This means that the GEMV operation at each moment needs to read massive amounts of weight data from GPU memory and perform corresponding multiplication and accumulation calculations.

[0047] For the reasons mentioned above, the execution time of the GEMV operator accounts for the majority of the total decoding time, becoming the core bottleneck restricting LLM inference performance. Given the critical bottleneck role of the GEMV operator in the decoding stage, how to efficiently optimize it has become a key technical issue determining the user experience of LLM applications. Currently, optimization of the GEMV operator is mainly achieved through two types of optimization methods: manual optimization and automatic optimization. However, both of these methods have inherent limitations and cannot meet the comprehensive requirements of LLM inference for high performance, cross-platform compatibility, and low development costs.

[0048] Specifically, manual optimization involves engineers manually writing assembly code or optimizing high-level language code to design customized operator implementations tailored to the architectural characteristics of a specific GPU hardware platform. GPUs, as processors specifically designed for parallel computing, contain a large number of parallel computing units. The architectures of GPUs from different manufacturers (such as Nvidia, Intel, and AMD) vary significantly, placing extremely high demands on engineers. They must possess in-depth knowledge of the target GPU's microarchitectural details, instruction set characteristics, memory hierarchy, and other specialized expertise. In other words, manual optimization methods are highly dependent on engineer experience. Furthermore, because it is manual optimization, the entire design cycle is long, development costs are high, and optimization results vary from person to person.

[0049] Furthermore, while GPU vendors typically provide GEMV operator optimizations for specific architectures, this further increases design costs, and optimization methods designed for one architecture are difficult to generalize and migrate to other architectures, severely limiting flexibility. In other words, manually optimized operators are strongly tied to the target GPU hardware, and their code logic is entirely dependent on the hardware characteristics of a specific architecture, making it impossible to directly migrate to GPUs of other architectures and failing to meet the requirements for cross-platform deployment.

[0050] Taking the industry-standard Intel OpenVINO framework as an example, its manually optimized GEMV operator, tailored for its own GPU architecture, exemplifies the aforementioned shortcomings. The following section uses a snippet of manually optimized operator code as an example to detail the deficiencies of manual optimization.

[0051] For example, the GEMV operator optimization code in the Intel OpenVINO framework, which is deeply customized for the Intel GPU architecture, is as follows: bool FullyConnected_GEMV_Validate(const Params¶ms) const { / / Data type re-check:only support f16:int4:f16 if (scale_group_size == 0 || (scale_group_size % 16 != 0)) { DO_NOT_USE_THIS_KERNEL(params.layerID); } The `bool FullyConnected_GEMV_Validate(const Params¶ms)const` function in this code defines the parameter validation function for the GEMV operator. This function checks whether the input parameters meet the hard constraints of the kernel before executing a specific optimized kernel for that GEMV operator. Next, it constrains the quantization group size, checking if the quantization group size (scale_group_size) is 0 or not a multiple of 16. If `scale_group_size` is found to be 0 or not a multiple of 16, the macro `DO_NOT_USE_THIS_KERNEL(params.layerID)` is called, indicating that the currently manually optimized kernel is not suitable for this configuration, and thus the kernel is rejected.

[0052] / / Data type re-check:only support f16:int4:f16 if (input.GetDType() != Datatype::F16 || (output.GetDType() !=Datatype::F16&&output.GetDType() !=Datatype::F32 || (weights.GetDType() != WeightsType::INT4&&weights.GetDType() !=WeightsType::UINT4)){ DO_NOT_USE_THIS_KERNEL(params.layerID); } This code continues to impose hard constraints on data types within the validation function. Specifically, the code logic checks the data types of the input, output, and weights. For example, it requires that the input data type must be half-precision floating-point (F16), the output data type can be F16 or single-precision floating-point (F32), and the weight data type must be a 4-bit integer (INT4) or an unsigned 4-bit integer (UINT4) quantized type.

[0053] If any data type in the input, output, or weight data does not meet the above conditions, the DO_NOT_USE_THIS_KERNEL(params.layerID) macro is also called, indicating that the current manually optimized kernel cannot adapt to this combination of data types and is rejected. This limitation makes the kernel only applicable to models with specific data types, and completely ineffective for other data types (such as input data type F32, weight data type INT3, etc.).

[0054] / / Only support vector data as input, the data size should be aligned by 16 elements auto input_size=get_input_bf_size(fc_params); if (input_size.first>1 || input_size.second== 0 || input_size.second% 16 != 0|| weights.IFM().v % 16 != 0) { DO_NOT_USE_THIS_KERNEL(params.layerID); } } This code further imposes hard constraints on the shape of the input data. The code stipulates that the input must be in vector form, the length of the input vector must be a multiple of 16, and the dimension of the input feature matrix of the weight matrix must be a multiple of 16.

[0055] If any of the above conditions are not met, the DO_NOT_USE_THIS_KERNEL macro can be invoked, indicating that the manually optimized kernel cannot handle the current input shape. This limitation makes the kernel only applicable to scenarios with dimensions that are multiples of 16, and it cannot be used for more general input sizes (such as vectors of arbitrary length), further highlighting the fundamental lack of flexibility in manual optimization.

[0056] In summary, the verification function, through the three checks mentioned above, constitutes a hard constraint on manually optimized kernels, making the implementation code of this operator deeply customized for the Intel GPU architecture. However, it also brings a series of strict limitations: First, there are strict restrictions on the quantization format. When the model's quantization format does not meet this alignment condition, the GEMV operator cannot be called directly, and developers need to implement an adapted version separately.

[0057] Second, there are strict restrictions on data types. This means that the operator cannot be used at all for models that use other data types, limiting its applicability in diverse model deployment scenarios.

[0058] Third, there are strict restrictions on the shape of the input data. When the dimension is not a multiple of 16, the existing loop boundaries and memory access patterns will fail, leading to incorrect calculation results or a severe performance degradation.

[0059] For application scenarios that do not meet any of the above requirements, developers cannot directly use the manually optimized operator and must reinvest significant manpower in adaptation development. They must either modify the original operator's code logic to adapt to the new parameter configuration or write a low-performance fallback implementation. This undoubtedly relies heavily on engineers' experience and leads to higher design costs and longer development cycles, severely limiting the operator's flexibility and reusability.

[0060] The shortcomings of manual optimization methods have been described above. The shortcomings of automatic optimization methods will be described below with reference to the code.

[0061] Automatic optimization methods (such as compiler-based automatic optimization methods) automatically explore parameter combinations by defining a design space composed of general optimization parameters, and have good portability. However, the design space of related technologies is relatively coarse-grained, and is usually independent of hardware platform. It cannot capture key architectural micro-features, that is, it cannot perceive the impact of the micro-feature differences of different GPU architectures on operator performance, which can easily lead to limited optimization results.

[0062] For example, in GPU architecture, warp or subgroup is the basic unit of thread organization. Nvidia GPUs have a fixed warp size of 32, while Intel GPUs can configure subgroup sizes to 8, 16, or 32. This difference has a decisive impact on thread scheduling and data aggregation memory access efficiency. However, existing automatic optimization methods do not include architecture-related parameters such as warp / subgroup size and data aggregation loading width in their design space, resulting in generated code that cannot fully utilize the hardware characteristics of a specific architecture, thus limiting performance optimization.

[0063] Furthermore, existing automatic optimization methods do not fully consider the impact of the dequantization process on the performance of GEMV operators, limiting the actual execution efficiency of automatic optimization operators. In LLM inference, to reduce memory usage and memory bandwidth pressure, the weight matrix is ​​usually stored using low-bit-width quantization. Before computation, a dequantization operation must be performed to restore the low-bit data to floating-point numbers. The dequantization process involves multiple parameters, such as the quantization weight data type, the quantization weight data bit width, and the quantization data packing bit width. The selection of these parameters directly affects memory access efficiency and computational overhead.

[0064] However, existing automatic optimization methods typically treat dequantization as a pre-process independent of GEMV computation, failing to optimize it jointly with GEMV memory access and computation processes. This results in dequantization and GEMV computation being isolated in the generated operators, increasing data round-trip overhead and limiting actual execution efficiency.

[0065] For example, the code snippet defining the GEMV operator design space in the Tensor Virtual Machine (TVM) compiler is as follows: class GEMV(GPUScheduleRule): "A rule for GEMV and DecodeGEMV." if target.kind.name == "cuda": VEC_C = 4 LOAD_V_SHARED = True LOAD_V_VEC = 8 VEC_LOAD= 4 UNROLL = 256 SUPPORT_WARP_SHUFFLE = True if isinstance(len.S.int): TS, TR = 16.32 else: TS, TR = 16, 64 elif target.kind.name == "opencl" and ( ( "android" in str(target.host) or "(adreno" in str(target.attrs) ): TAG_S, TAG_R = "threadldx.x", "threadldx.x" VEC_C = 8 LOAD_V_SHARED = False LOAD_V_VEC = -1 UNROLL = 8 TS, TR = 2,32 This code snippet shows that while the design space predefines different parameter values ​​for different backends, its core design still only includes general optimization parameters. Although these general optimization parameters can be tuned to some extent on different hardware platforms, the design space has two fundamental flaws: First, architecture-related parameters are not considered. Existing automatic optimization methods do not take into account the impact of GPU architecture-related parameters on operator performance. In fact, the differences in these parameters have a decisive impact on operator performance. However, due to the lack of such architecture-specific parameters in the design space, the generated code cannot be finely tuned for the hardware characteristics of a specific architecture, thus failing to fully unleash the hardware potential and severely limiting the optimization effect.

[0066] Second, the relevant parameters of the dequantization process are not considered. The design space does not include parameters related to the dequantization operation, which are crucial to the inference efficiency of the quantization model, directly determining the implementation method and memory access overhead of the dequantization operation. Because the dequantization parameters are not incorporated into the joint optimization, and the coupling between the weight dequantization process and the vector-matrix multiplication process is not considered, the GEMV operator generated by existing automatic optimization methods separates dequantization from GEMV computation, failing to achieve deep integration between the two. This increases data round-trip overhead and limits the actual execution efficiency of the operator.

[0067] To address the aforementioned issues, there is an urgent need for an operator optimization method that can perceive the differences between different GPU architectures and integrate the dequantization process, in order to improve the cross-platform performance and development efficiency of operators. This application is proposed based on this need. The core concept of this application is to adopt an automatic operator optimization method based on parameterized templates, and to achieve perception and adaptation to different GPU architectures by constructing a fine-grained joint design space.

[0068] It should be noted that, in this application, the design space refers to the multidimensional space consisting of all possible design decisions (i.e., adjustable parameters) and their values ​​when designing a system or algorithm. The joint design space here refers to the parametric design space formed by jointly modeling parameters from different optimization dimensions. The meaning of "joint" is reflected in integrating the three types of parameters—originally separate computational scheduling, hardware architecture, and quantization data format—into the same design space for collaborative optimization.

[0069] Specifically, this joint design space can include three types of parameters. For example, the first type consists of general parameters related to computational parallel scheduling; the second type consists of characteristic parameters related to GPU hardware architecture, including thread bundle / subgroup size and data aggregation loading width, used to enable the optimization process to perceive and adapt to the micro-characteristics of different GPU architectures; the third type consists of parameters related to the dequantization process, including quantization weight data type parameters, quantization weight data bit width parameters, and quantization data packing bit width parameters, used to deeply integrate dequantization operations with GEMV computation. Based on this, an automated search method is used to traverse the parameter combinations in this joint design space to generate optimal operator implementation code for different hardware platforms.

[0070] By employing the aforementioned automatic operator optimization method, this application effectively addresses the shortcomings of related technologies. Specifically, the problems of poor flexibility, high development costs, and long design cycles associated with manual optimization stem from the significant differences in GPU hardware architectures and the numerous types of operator optimization parameters, making comprehensive consideration difficult. This application replaces the rigid constraints of manual coding by constructing a fine-grained joint design space and an automated optimization process. This enables operators to automatically perceive and adapt to different GPU hardware architectures, achieving high-performance cross-platform deployment without manual intervention, fundamentally solving the problems of poor flexibility and high costs.

[0071] To address the issues of insufficient hardware awareness and inadequate dequantization fusion in automatic optimization, this application comprehensively incorporates GPU architecture characteristic parameters (such as thread bundle / subgroup size and data aggregation loading width) and dequantization process parameters (such as quantization weight data type parameters, quantization weight data bit width parameters, and quantization data packing bit width parameters) into the design space, forming a fine-grained fusion operator design space. This achieves the fusion optimization of hardware awareness, dequantization, and GEMV operators, thereby generating highly efficient operator code on different GPU architectures and solving the problem of limited effectiveness in traditional automatic optimization. Therefore, this application not only improves operator development efficiency but also achieves balanced high performance across different GPU architectures.

[0072] The automatic operator optimization method provided in the embodiments of this application will now be described with reference to the accompanying drawings.

[0073] Specifically, the automatic operator optimization method provided in this application embodiment may include obtaining a set of parameters to be tuned for the operator to be optimized; wherein, the set of parameters to be tuned includes a first type of parameters related to computational parallel scheduling, a second type of parameters related to the execution unit of the target hardware platform, and a third type of parameters related to the quantization and dequantization data formats of the operator to be optimized; generating multiple parameter value combinations for the operator to be optimized based on the candidate values ​​of each parameter in the set of parameters to be tuned; loading each parameter value combination in the multiple parameter value combinations into the parameterization template of the operator to be optimized to generate executable code of the operator to be optimized on the target hardware platform; evaluating the execution performance of each executable code, and determining a target parameter combination from the multiple parameter value combinations based on the execution performance; and generating target code for the operator to be optimized to be executed on the target hardware platform based on the target parameter combination.

[0074] In the embodiments of this application, the operator to be optimized refers to a deep learning computation unit that needs performance optimization. For example, the operator to be optimized may include the GEMV operator, the General Matrix Multiplication (GEMM) operator, the Batch Normalization operator, the Layer Normalization operator, the Activation operator, the Pooling operator, and so on.

[0075] It should be understood that this application uses the GEMV operator as an example for illustration, but the automatic optimization method of the operator provided in this application is also applicable to other types of operators, such as the various operators listed above, as long as their calculation mode can be described by a parameterized template and involves the dequantization operation of quantized data.

[0076] The quantized data may include weight parameters stored in a low-bit-width format, or it may include intermediate activation values ​​that have undergone quantization. The dequantization operation aims to restore this low-bit-width data to floating-point numbers so that high-precision calculations can be performed in the computing unit.

[0077] The set of parameters to be tuned refers to the collection of all parameters that can be adjusted to change the implementation and performance of the operator during the automatic optimization process. Each parameter in this set defines a set of possible candidate values, and these parameters and all combinations of their values ​​constitute a multidimensional search space, namely the parameterized design space described above in this application.

[0078] The parameterized design space is a concretization of the design space in the field of operator optimization. It specifically refers to the set defined for generating the target operator, consisting of multiple variable design parameters and all their possible discrete values. Each combination of parameters uniquely corresponds to a specific operator implementation.

[0079] The set of parameters to be tuned and the joint design space can be understood as two closely related, corresponding, but different concepts. The set of parameters to be tuned is the specific parameter content that constitutes the joint design space, while the joint design space is a multi-dimensional search space defined by these parameters and their candidate values. In this application, these two terms can be used interchangeably depending on the context. For example, when emphasizing the parameters themselves, "set of parameters to be tuned" is used; when emphasizing the possibility space constituted by these parameters, "joint design space" is used.

[0080] The operator automatic optimization method provided in this application constructs a joint design space that includes parameters related to computational parallel scheduling, parameters related to hardware platform execution units, and parameters related to quantization and dequantization data formats. It then uses an automated search method to generate executable code adapted to the target hardware platform, thereby improving operator development efficiency and achieving high-performance cross-platform execution.

[0081] Specifically, this method introduces a second type of parameter (such as thread bundle / subgroup size and data aggregation loading width) related to the execution unit of the target hardware platform, enabling it to perceive the micro-characteristic differences of different hardware platforms (such as different GPU architectures). This hardware-aware capability allows automated optimization to adjust parallel scheduling strategies for specific GPU architectures, generating executable code adapted to the characteristics of that hardware platform. This overcomes the performance bottleneck caused by traditional automated optimization ignoring GPU architecture differences, thus unlocking the performance potential of different GPU architectures.

[0082] By introducing a third type of parameter related to the quantization and dequantization data formats (such as quantization weight data type parameters, quantization weight data bit width parameters, and quantization data packing bit width parameters), the parameters involved in the dequantization process are incorporated into the joint design space. This enables deep integration of dequantization operations and matrix-vector multiplication calculations, achieving seamless connection at the code generation level. Specifically, each thread directly reads the quantized data according to the packing bit width during memory access, and immediately participates in multiplication and accumulation calculations after unpacking and dequantization. This avoids the data round-trip overhead of dequantization, storage, and then reading and calculating in traditional solutions, significantly improving memory access efficiency and computational pipeline utilization.

[0083] By incorporating three categories of features—parallel scheduling parameters, hardware platform execution unit parameters, and quantization and dequantization data format parameters—into the parameter set to be tuned, a richer and more granular joint design space is formed than that of related technologies. Through automated traversal and evaluation, the operator automatic optimization method provided in this application can accurately find the optimal configuration for the target hardware platform from a massive number of parameter combinations. This global optimization capability is impossible to achieve through manual design or single-dimensional automatic optimization, ensuring that the generated operator code achieves near-theoretical performance limits on the target platform.

[0084] Furthermore, the automatic operator optimization method provided in this application boasts high efficiency, requiring no manual intervention and automatically completing the optimization process through a preset search method. Simultaneously, by designing hardware-aware parameters, it overcomes the shortcomings of traditional automatic optimization methods, which are often "general-purpose but lack performance." Developers only need to specify the target hardware platform, and this method can automatically generate high-performance operator code adapted to that platform, eliminating the need for manual optimization for each hardware platform and significantly reducing the operator development cycle and technical barriers.

[0085] To more clearly illustrate the technical solution of this application, the GEMV operator is used as a specific example of the operator to be optimized, and the algorithm principle of the GEMV operator is explained in detail below.

[0086] Please refer to Figure 1 , Figure 1 This is a schematic diagram of an algorithm for a matrix-vector multiplication operator, as illustrated in an exemplary embodiment of this application.

[0087] like Figure 1 As shown, the GEMV operator receives two inputs. One input is a row vector of shape [1, K] (which can be simply referred to as the input vector), and the other input is a two-dimensional weight matrix of shape [K, M] (which can be simply referred to as the weight matrix). The GEMV operator performs a dot product operation between the input vector and each column of the weight matrix, generating an output vector of shape [1, M].

[0088] The code corresponding to the calculation process of the GEMV operator is as follows: for m in M; / / Iterate through each element of the output vector (non-recursive axis M) for k in K; / / Iterate through each element of the input vector (recursion axis K) Y[1, m] += X[1, k]* W[k, m]; / / Multiplication and summation calculation Here, the non-reduction axis M represents the dimension of the output vector. Computations along this axis are independent of each other and are suitable for parallel processing. In parallel computing architectures such as GPUs, the computational tasks along the M axis can be assigned to different threads or thread blocks for parallel execution.

[0089] The reduction axis K represents the dimension of the input vector, and is also the dimension that needs to be summed (reduced) in the computation of each output element. Computations on this axis are data-dependent and need to be completed through iterative summation. It should be understood that the size of the reduction axis K determines the computational cost of each output element, which can typically reach thousands or even tens of thousands.

[0090] The multiplication-accumulation calculation Y[1, m] += X[1, k]* W[k, m] is the basic computational unit of the operator, consisting of multiplication and addition operations.

[0091] It's worth noting that multiply-accumulate is the underlying operation for implementing the dot product in computers; that is, multiplication is performed first, and then the multiplied data is accumulated into the accumulator. Correspondingly, the dot product operation is implemented through K multiply-accumulate calculations. In the GEMV operator, performing a dot product between the input vector and each row of the weight matrix is ​​essentially a process of sequentially performing multiply-accumulate and summing K pairs of elements.

[0092] During the decoding phase of LLM, the GEMV operator is called word-by-word. That is, at each time step, the hidden state vector (i.e., the input vector) corresponding to the currently generated word needs to undergo a GEMV operation with the weight matrix in the model to calculate the probability distribution of the next word. This process continues throughout the decoding phase, and the weight matrix is ​​enormous. Therefore, the execution efficiency of the GEMV operator directly determines the response speed of LLM, becoming a core bottleneck restricting inference performance.

[0093] To reduce memory usage and bandwidth pressure, the weight matrix can be stored in a low-bit-width quantization format in LLM inference deployment. Quantization refers to the process of mapping high-precision floating-point numbers (such as FP32) to low-precision integers (such as INT8, INT4).

[0094] Specifically, taking INT4 quantization as an example, each weight value is represented by only 4 bits, which can be compressed to 1 / 8 of the storage space compared to FP32 (32 bits), greatly reducing memory access overhead.

[0095] However, since the GPU's computing units execute floating-point instructions, while the quantized data is of integer type, the quantized data cannot be directly used in GEMV multiplication and accumulation calculations. Therefore, before performing GEMV calculations, these low-bit-width quantization weights must first be dequantized back to floating-point numbers to restore their approximate original values.

[0096] For example, during calculation, INT4 data needs to be converted to half-precision floating-point (FP16) or single-precision floating-point (FP32) before participating in multiplication and accumulation calculations. The dequantization process involved in the embodiments of this application will be described in detail below with reference to the accompanying drawings.

[0097] It is worth noting that, in the actual execution of GEMV calculation, in order to further improve efficiency, the multiply-accumulate calculation process is mapped to the GPU's computing array, and multiple computing units complete the GEMV operation in parallel. This is also the key to breaking through the data dependency of the recursive axis K and improving the overall execution efficiency.

[0098] Specifically, GPUs employ a Single Instruction Multiple Thread (SIMT) execution model, achieving high throughput through the parallel execution of a large number of threads.

[0099] For example, parallel implementation methods may include: assigning the M elements of the output vector to different thread blocks for processing, with each thread block responsible for calculating a portion of the output elements; within each thread block, loading the input vector and corresponding weight data into shared memory to reduce the latency of accessing global memory; and having multiple threads collaboratively complete the multiply-accumulate calculation and perform the reduction operation through thread bundle shuffle or shared memory.

[0100] However, the parallel execution units of different GPU architectures differ significantly in size and organization. For example, Nvidia GPUs have a fixed thread bundle size of 32, Intel GPUs can configure subgroup sizes to 8, 16, or 32, and AMD GPUs typically have a wavefront size of 64. These differences directly determine thread scheduling strategies and data aggregation memory access efficiency; therefore, parallel optimization of the GEMV operator requires fine-tuning for specific GPU architectures. This application introduces a second type of parameter to enable automated optimization to perceive and adapt to these differences, which will be described in detail later.

[0101] It is understood that the automatic operator optimization method provided in this application is not limited to the GEMV operator exemplified above, but possesses good generalization and adaptation capabilities. In addition to the GEMV operator, various deep learning operators such as matrix multiplication, batch normalization, layer normalization, activation function, and pooling operators can all be applied to the automatic operator optimization scheme of this application. For example, by designing a corresponding parameterized template for each operator's specific computational mode and incorporating parameters related to quantization storage and dequantization processes into a unified joint design space, the adaptation and deployment of the optimization process can be completed, achieving automatic operator optimization.

[0102] The system architecture corresponding to the operator automatic optimization method provided in the embodiments of this application will be described below with reference to the accompanying drawings.

[0103] Please refer to Figure 2 , Figure 2 This is a system architecture diagram illustrating an exemplary embodiment of this application. (As shown...) Figure 2 As shown, the system architecture corresponding to the operator automatic optimization method provided in this application mainly includes five modules: a joint design space construction module, a design space exploration and control module, an operator code generation module, an operator code compilation and execution module, and an operator code output module. The functions, collaborative relationships, and hierarchical positioning of each module in the electronic device software architecture are described in detail below.

[0104] I. Joint Design Space Composition Modules The joint design space module is used to construct the core foundation of the method in this application, the "joint design space". The joint design space refers to the parametric design space formed by jointly modeling parameters from different optimization dimensions. The "joint" aspect is reflected in integrating three originally separate types of parameters into the same design space for collaborative optimization.

[0105] For example, the joint design space building module can be used to perform the following operations: (I) Parameter Definition and Classification The joint design space construction module first determines all parameters that need to be optimized and divides these parameters into three categories: the first category of parameters related to computational parallel scheduling, the second category of parameters related to the execution unit of the target hardware platform, and the third category of parameters related to the quantization and dequantization data format of the operator to be optimized.

[0106] 1. First type of parameters: Parameters related to parallel scheduling of computation.

[0107] Computational parallel scheduling refers to the process of decomposing computational tasks into parallelizable subtasks and allocating these subtasks to different computing resources (such as threads or thread blocks) in a parallel computing architecture (such as GPUs). The key aspects of computational parallel scheduling lie in how to divide computational tasks, how to organize threads, and how to manage data flow to maximize hardware resource utilization, reduce synchronization overhead, and hide memory access latency. In deep learning operator optimization scenarios, computational parallel scheduling strategies are typically described by a set of adjustable parameters that determine how computational tasks are mapped onto parallel hardware.

[0108] The first type of parameters in this application is precisely the set of parameters related to the aforementioned computational parallel scheduling strategy. In existing automatic optimization frameworks (such as TVM), these parameters can be called "general optimization parameters" or "scheduling parameters" because they are decoupled from specific hardware architectures and can be tuned to a certain extent on different hardware platforms. This application uses these general optimization parameters as the basic dimension of the joint design space, and introduces the second and third types of parameters on this basis to achieve finer-grained joint optimization.

[0109] Specifically, the first type of parameter may include, but is not limited to, one or more of the following parameters: The Reduction axis Tile size (TR) parameter controls the size of the parallel blocks along the reduction axis (i.e., the K-axis in GEMV), specifically the number of elements processed by each thread block or thread bundle on the reduction axis. In other words, this parameter describes the number of parallel threads along the reduction axis dimension; that is, on the K-axis, there are TR threads collaboratively completing the recursive computation. It should be understood that this parameter directly affects the granularity of the computational task partitioning.

[0110] The non-reduction axis tile size (TC) parameter controls the size of the parallel blocks on the non-reduction axis (i.e., the M-axis in GEMV), specifically the number of elements processed by each thread block or thread bundle on the non-reduction axis. In other words, this parameter describes the number of parallel thread bundles along the non-reduction axis dimension; specifically, on the M-axis, a total of TC thread bundles are allocated to compute the output results in parallel. This parameter determines the number of thread blocks or the number of output elements processed by each thread, directly affecting the degree of parallelism. In this embodiment, the non-reduction axis parallelism parameter can also be represented by TS.

[0111] The loop unroll factor (UNROLL) parameter specifies the number of times the inner loop is unrolled. In other words, it describes the explicit unrolled dimension of the loop body. Therefore, in this embodiment, the loop unroll factor parameter can also be called the loop unrolled dimension parameter. This parameter improves computational efficiency by increasing instruction-level parallelism, but may increase register pressure. A trade-off needs to be struck between computational density and resource consumption. The higher the number of unrolls, the larger the dimension of the unrolled loop body, and the higher the instruction-level parallelism.

[0112] The Shared Memory (USE_SMEM) flag indicates whether to use the GPU's shared memory to cache input vectors or weight data. In other words, it indicates whether to use a shared cache to temporarily cache data. It should be understood that shared memory is a high-speed programmable cache on the GPU, and its proper use can reduce latency in accessing global memory. Optionally, in some possible implementations, this shared memory (USE_SMEM) flag may also be named the "Use Shared Memory" parameter.

[0113] The Vectorization Width parameter (VEC_WIDTH) specifies the width of vectorized loading / computation, such as loading 2, 4, or 8 data elements at a time. In other words, it describes the number of computation operations performed at once using vectorized computation instructions. Vectorization can increase memory bandwidth utilization, but it requires data alignment. The larger the vectorization width, the more data is processed in a single instruction, and the higher the requirements for data alignment become.

[0114] The candidate values ​​for the above parameters are usually powers of 2. For example, the parallelism parameter of the recursive axis can be 8, 16, 32, 64, etc., which can be set according to hardware resources and model size. During the automated optimization process, these parameters will be searched together with the second and third types of parameters to find the target parameter combination (such as the optimal parameter combination).

[0115] 2. Second type of parameters: Parameters related to the execution unit of the target hardware platform. The second type of parameter is used to characterize the organization and capabilities of the underlying execution units of the target GPU hardware, enabling the optimization process to perceive the micro-features of different architectures. It should be understood that the second type of parameter can also be called architecture-related parameter, which is strongly related to the specific hardware architecture and may have different candidate values ​​or configurations on GPUs from different manufacturers and of different generations.

[0116] Specifically, the second type of parameter may include, but is not limited to, one or more of the following parameters: Thread bundle size, subgroup size, or wavefront size parameter: These parameters characterize the size of the basic parallel execution unit of the target GPU. It's worth noting that thread bundle, subgroup, and wavefront are essentially the same concept (the basic parallel execution unit of the GPU), differing only in naming by different manufacturers. Therefore, thread bundle size, subgroup size, and wavefront size parameters all characterize the size of the basic parallel execution unit of the target GPU.

[0117] In GPU architecture, a thread bundle / subgroup is the organizational unit of threads in the GPU, mapping to the SIMD execution unit within the GPU architecture. Threads are not scheduled independently, but rather scheduled and executed in fixed-size thread bundles, which are called basic parallel execution units. The size of this thread bundle / subgroup (i.e., the number of threads contained in each basic unit) varies across different architectures, and the SIMD execution units also differ between different GPU architectures. Therefore, defining a reasonable thread bundle / subgroup size directly impacts the performance of the GEMV operator.

[0118] For example, in Nvidia's GPU architecture, 32 threads form a thread bundle, which is the basic unit for executing the same instruction simultaneously. In Intel's GPU architecture, a similar concept is called a subgroup, which can be configured to be 8, 16, or 32 in size. In AMD's GPU architecture, the corresponding concept is called a wavefront, which is typically 64 in size.

[0119] Please refer to Figure 3A and Figure 3B , Figure 3A and Figure 3B This is a schematic diagram illustrating the differences in thread bundle and subgroup sizes under different GPU architectures, as shown in an exemplary embodiment of this application.

[0120] Figure 3A The corresponding architecture is the Nvidia GPU architecture, in which the thread bundle size is usually a fixed value of 32, which means that every 32 threads form a basic unit that executes the same instruction simultaneously.

[0121] Figure 3B The corresponding architecture is the Intel GPU architecture, in which the subgroup size can be adjusted from 8, 16, and 32 according to actual needs, providing greater flexibility for optimization.

[0122] Clearly, the size settings for thread bundles or subgroups differ significantly across GPU architectures. By treating thread bundle size, subgroup size, or wavefront size parameters as searchable variables, the operator's automatic optimization process can select the most suitable execution unit size for different GPU architectures, thereby affecting thread scheduling methods, data aggregation and memory access efficiency, and overall operator performance.

[0123] The Coalesced Load Width parameter describes the total data width when threads within a thread bundle / subgroup aggregate memory access. In other words, it represents the amount of data loaded at once using the Coalesced Load instruction.

[0124] To improve memory access efficiency, GPUs support merging independent memory access instructions from multiple threads within the same thread bundle (or subgroup) into a single, wider instruction (called a wide memory access instruction) for execution. The width of the data obtained from this one-time memory access needs to match the width of the storage units in the GPU architecture; therefore, the data aggregation loading width should be adjusted according to the GPU architecture. The total number of bits of data that can be read from this merged one-time memory access is the data aggregation loading width.

[0125] Please refer to Figure 4A and Figure 4B , Figure 4A and Figure 4B This is a schematic diagram illustrating the difference in data aggregation loading width as shown in an exemplary embodiment of this application.

[0126] in, Figure 4A This shows the data aggregation loading behavior when the memory access data width is 32 bits. Figure 4B This shows the data aggregation loading when the memory access data width is 64 bits.

[0127] Specifically, the data aggregation loading width is determined by two key parameters: the thread bundle / subgroup size and the data bit width of each thread's single memory access (which can be referred to as the single-thread memory access data width). In this embodiment, a single thread reads a complete quantized data packet unit at a time; therefore, the memory access data width of each thread is equal to the quantized data packet bit width in the dequantization-related parameters. In other words, the data aggregation loading width parameter is linked to the thread bundle / subgroup size parameter and the quantized data packet bit width parameter, and can be determined by their product, such as data aggregation loading width = thread bundle / subgroup size × quantized data packet bit width.

[0128] like Figure 4A As shown, when the memory access data width for a single thread is 32 bits, the corresponding quantized data packet width is also 32 bits. If the thread bundle / subgroup size is 32, then each of the 32 threads requests 32 bits of data. With hardware support, these independent memory access requests can be merged into a single 1024-bit wide memory access instruction, meaning the aggregated data loading width is 32 × 32 bits = 1024 bits. This implies that 32 threads within the same thread bundle can collaboratively complete a 1024-bit data loading operation.

[0129] like Figure 4B As shown, when the memory access data width of a single thread is 64 bits, the quantized data packet width is also 64 bits. Similarly, with a thread bundle / subgroup size of 32, if each of the 32 threads requests 64 bits of data, the aggregated data loading width after merging is 32 × 64 bits = 2048 bits.

[0130] By using the data aggregation loading width as a searchable variable and combining it with the linkage between the thread bundle / subgroup size parameter and the quantized data packing bit width parameter, the operator automatic optimization process can select the optimal memory access strategy for different GPU architectures and quantization configurations, thereby maximizing memory bandwidth utilization and improving overall operator performance.

[0131] 3. Third category of parameters: Parameters related to quantization and dequantization data formats. The third type of parameter describes the data representation and storage format of the weight matrix after quantization, enabling deep integration of dequantization operations with GEMV computation. In LLM inference, to reduce memory usage and bandwidth pressure, the weight matrix is ​​typically stored using a low-bit-width quantization format. However, quantized data cannot directly participate in floating-point operations; dequantization must be performed before computation to restore this low-bit-width data to floating-point numbers. The specific implementation of the dequantization operation depends on the configuration of the quantization parameters, the selection of which directly affects memory access efficiency and computational overhead. This application incorporates these parameters into the joint design space, enabling the automated optimization process to select the optimal dequantization implementation based on different quantization configurations and hardware platforms.

[0132] Specifically, the third type of parameter may include, but is not limited to, one or more of the following parameters: Quantized Weight Data Type (QDATA_TYPE): Defines the data type of the quantized weights, indicating whether the quantized data is signed or unsigned integers. Different quantization algorithms (such as symmetric quantization and asymmetric quantization) have different requirements for data types, which directly determines whether sign extension is needed during dequantization.

[0133] Please refer to Figure 5A and Figure 5B , Figure 5A and Figure 5B This is a schematic diagram illustrating the inverse quantization process of quantized data, as shown in an exemplary embodiment of this application. Figure 5A The dequantization process for unsigned quantized data is shown. Figure 5B The dequantization process for signed quantized data is shown.

[0134] like Figure 5A As shown, taking the case of quantization type uint4 and packing width of 16 bits as an example, the dequantization process of unsigned quantized data is as follows: First, data unpacking (i.e., extracting quantization weights) is performed. After reading the 16-bit packaged data from memory, four 4-bit unsigned integers need to be extracted. Specifically, for the nth quantization weight (n=1, 2, 3, 4), a right shift operation is performed, with a shift distance of (n-1) × 4 bits; the shifted data is then bitwise ANDed (&) with the mask to extract a single uint4 type quantization weight.

[0135] For example, the first quantization weight is shifted by 0 bits, the second by 4 bits, the third by 8 bits, and the fourth by 12 bits. The mask value depends on the quantization weight data bit width. For 4-bit quantization, the mask is 0xF (i.e., binary 1111).

[0136] Optionally, in the embodiments of this application, the quantization weight data bit width can be set and adjusted according to the actual situation. For example, when the quantization weight data bit width is 3 bits (uint3), it is shifted by (n-1)×3 bits, and the mask is 0x7.

[0137] Next, a data type conversion is performed (extending to a 32-bit integer). Specifically, the data type conversion function is called to convert the extracted uint4 type quantization weights to a 32-bit integer (int32). At this point, the data still retains its unsigned characteristic, with a value range of [0, 15].

[0138] Next, the value range is adjusted (i.e., the sign bit is restored). Specifically, to convert an unsigned integer to the range corresponding to a signed floating-point number, a constant value needs to be subtracted. This constant value can be the median offset of the value range corresponding to the quantization weight data type, or in other words, it can be half the maximum value that the quantization weight data type can represent. For example, for the uint4 quantization type, the constant value can be 0x8 (i.e., decimal 8). By subtracting 0x8, the value range is adjusted from [0, 15] to [-8, 7], thus restoring the positive and negative characteristics of the data.

[0139] Optionally, if the quantization weight data bit width is 3 bits (uint3), the constant value can be 0x4 (i.e., decimal 4), and the value range is adjusted from [0, 7] to [-4, 3].

[0140] Finally, the quantized value restoration and data type conversion are completed. Specifically, the data type conversion function is called to convert the adjusted int32 data into a half-precision floating-point type (float16) so that it can participate in subsequent multiplication and accumulation calculations.

[0141] For example, a numerical calculation function is called to multiply the adjusted int32 data with a value range of [-8, 7] by a preset scaling factor, thus mapping the integer value to the order of floating-point values. The scaling factor is determined by the calibration parameters of the quantization process and is used to restore the original numerical range of the weights. Then, a hardware-native data type conversion function is called to convert the scaled floating-point data to float16, resulting in floating-point weight data that can be directly used in GEMV calculations.

[0142] like Figure 5B As shown, taking the case of quantization type int4 and packing width of 16 bits as an example, the dequantization process of signed quantized data is as follows: First, data unpacking (quantization weight extraction) is performed. After reading the 16-bit packaged data from memory, four 4-bit signed integers are extracted. For details, please refer to [link to relevant documentation]. Figure 5A The descriptions in the corresponding embodiments will not be repeated here.

[0143] Next, a data type conversion is performed (extended to 32-bit integer). Specifically, the extracted 4-bit data is treated as an unsigned integer (uint4) and converted using a data type conversion function to a 32-bit integer (int32). It should be understood that at this point, the data value range is [0, 15], and the sign information has not yet been restored; that is, the data still retains its unsigned characteristics.

[0144] Next, the sign bit is restored and the value range is adjusted. Specifically, the above int32 data is XORed with the constant 0x8, and then the same constant 0x8 is subtracted. Through these two steps, the value range is adjusted from [0, 15] to [-8, 7], thus restoring the correct numerical representation of the signed integer.

[0145] Next, numerical scaling is performed. Specifically, the adjusted int32 data is multiplied by a preset scaling factor to map the integer value to the order of floating-point values.

[0146] Finally, data type conversion is performed. The hardware's native data type conversion function is called to convert the scaled floating-point data to float16, resulting in floating-point weighted data that can be directly used in GEMV calculations.

[0147] It should be understood that the difference between signed and unsigned dequantization lies in the different methods of sign recovery. Specifically, unsigned types directly adjust the value range by subtracting a constant, while signed types achieve sign bit flipping and value range adjustment through an XOR operation followed by subtracting the constant. This difference is reflected in the parameterized template as different predefined logical branches, controlled by the quantization weight data type parameter, which will be discussed in detail later.

[0148] In this embodiment, the quantization weight data type parameter is used as a searchable variable, enabling the operator automatic optimization process to select the correct dequantization code path according to the requirements of the quantization algorithm, thus ensuring the accuracy of the dequantization operation.

[0149] Quantized Weight Bit-width (QBITS): Defines the number of bits used for quantized weight data, such as 2 bits, 3 bits, or 4 bits.

[0150] The bit width parameter of the quantization weight data determines the compression ratio and accuracy of the model. The smaller the bit width, the less storage space is required for the model weights, and the memory bandwidth pressure is reduced accordingly. Taking int4 quantization as an example, each weight value is represented by only 4 bits, which can be compressed to 1 / 8 of the storage space compared to FP32 (32 bits); if int3 quantization is used, it can be further compressed to about 1 / 10 of the storage space.

[0151] However, the choice of quantization weight data bit width not only affects the model size but also significantly impacts the computational and memory access costs of the dequantization process. Specifically, the quantization weight data bit width parameter determines the number of weights contained in each packet unit. With the same packet bit width, a smaller quantization weight data bit width results in a larger number of weights per packet unit, and consequently, an increase in the number of shift and mask operations required for unpacking.

[0152] For example, when the packet width is 32 bits, for 4-bit quantization weights, each packet unit contains 8 weights, and unpacking requires 8 shift and mask operations; for 3-bit quantization weights, each packet unit can contain 10 weights, making the unpacking operation more complex.

[0153] Meanwhile, the quantization bit width parameter also determines the overall size of the weight matrix, thus affecting the time required to load data from global memory. For the GEMV operator with limited memory bandwidth, reducing the quantization bit width can significantly reduce memory access overhead, but may result in a loss of accuracy.

[0154] It's worth noting that the same quantization weight data bit width can lead to different performance characteristics of the GEMV operator on different GPU architectures. This is because different GPU architectures have differences in cache size, memory bandwidth, and computing power, making the matching relationship between quantization weight data bit width and hardware characteristics complex.

[0155] For example, on GPUs with large caches, slightly increasing the quantization weight data bit width may not significantly affect memory access efficiency, but it can improve model accuracy; while on GPUs with smaller caches, a smaller quantization weight data bit width may be a better choice.

[0156] Therefore, by incorporating the quantization weight data bit width parameter as a searchable variable into the joint design space, the operator automatic optimization process can find the optimal balance between model accuracy and inference performance among options such as 2-bit, 3-bit, and 4-bit, based on the characteristics of the target hardware platform.

[0157] The Quantized Data Packing Width (PACK_WIDTH) parameter specifies the granularity of quantized data packing in memory, or in other words, it describes the number of bits occupied when the quantized data is packaged and stored, such as 8 bits, 16 bits, 32 bits, or 64 bits. Since hardware typically cannot directly access data smaller than 8 bits, multiple low-width quantized weight data are packaged into a wider data type for storage and retrieval.

[0158] Please refer to Figure 6 , Figure 6 This is a schematic diagram illustrating the relationship between the quantization weight data bit width and the quantization data packing bit width, as shown in an exemplary embodiment of this application.

[0159] like Figure 6 As shown, assuming the quantization weight data type is a signed integer and the quantization weight data bit width is 4 bits (corresponding to the quantization weight data type Int4), when the quantization data packing bit width is defined as 16 bits, then 4 Int4 quantization weight data are compressed into 1 16-bit data for storage; when the quantization data packing bit width is defined as 32 bits, then 8 Int4 quantization weight data are compressed into 1 32-bit data for storage.

[0160] It is worth noting that the choice of quantized data packing bit width will affect the memory access granularity and unpacking computation, and a trade-off needs to be made based on the GPU architecture type.

[0161] Specifically, the larger the quantization data packet width, the more quantized data each memory access instruction can read, and the fewer memory accesses are required. For example, to read 32 4-bit quantization weight data, a 16-bit packet width requires 8 memory accesses (each read reads 16 bits, containing 4 data points); a 32-bit packet width requires only 4 memory accesses (each read reads 32 bits, containing 8 data points). Therefore, a larger quantization data packet width can improve memory access efficiency and reduce memory bandwidth pressure.

[0162] Meanwhile, the larger the bit width of the quantized data packet, the more weights are contained in each packet unit, and the computational complexity of the unpacking operation increases accordingly. The unpacking process requires shifting and masking operations based on the bit width of the quantized data packet to extract each quantization weight. The computational cost of the unpacking operation is directly proportional to the number of weights in the packet unit, and a balance needs to be struck between memory access efficiency and unpacking overhead.

[0163] This application incorporates the quantization data packing bit width parameter into the joint design space, enabling the automated optimization process to comprehensively consider memory access efficiency and unpacking overhead, and to find the packing granularity most suitable for the current hardware platform and quantization configuration among options such as 8 bits, 16 bits, 32 bits, and 64 bits.

[0164] It should be understood that the quantization data packing bit width parameter and the data aggregation loading width parameter in the second type of parameters are closely related. As mentioned earlier, the data aggregation loading width is determined by the product of the thread bundle / subgroup size and the quantization data packing bit width. This relationship allows the operator's automatic optimization process to simultaneously consider the matching degree between the hardware execution unit size and the data storage format.

[0165] During dequantization, each thread in the GPU reads the packaged data from memory using the quantized data packet width as the basic memory access unit. The memory access instructions within a thread bundle / subgroup then form an aggregated memory access instruction, the width of which is the aggregated data loading width. By using the quantized data packet width as a searchable variable and combining it with the relationship between the thread bundle / subgroup size parameter, the operator's automatic optimization process can select the optimal memory access configuration for different GPU architectures, thereby maximizing memory bandwidth utilization and improving overall operator performance.

[0166] As can be seen from the detailed explanation of the three types of parameters above, the quantization weight data type parameter, the quantization weight data bit width parameter, and the quantization data packing bit width parameter do not exist in isolation. Instead, they collectively describe, from different dimensions, how the quantized data is stored in memory and all the information required for the dequantization operation. These three types of parameters each play their respective roles in the dequantization process, while also cooperating with each other to jointly determine the generation logic of the dequantization code.

[0167] Specifically, the quantization weight data type parameter determines whether sign extension is required during dequantization, i.e., whether to choose a signed or unsigned dequantization logic branch, or in other words, whether to choose a signed or unsigned dequantization path. The quantization weight data bit width parameter determines the mask bit width and shift step size in the depackaging operation. The quantization data packing bit width parameter determines the granularity of memory reads (e.g., 16 bits, 32 bits, etc.) and the number of quantized data items contained in each packing unit.

[0168] By incorporating these three types of parameters into a joint design space, the automatic operator optimization process can generate the most suitable dequantization code for different quantization configurations and GPU hardware platforms, achieving a deep integration of dequantization operations and GEMV computation. It is precisely this design of jointly optimizing previously isolated parameters that enables this application to overcome the performance bottleneck of traditional automatic optimization methods, generating highly efficient operator code on different GPU architectures.

[0169] (ii) Definition of candidate values After clarifying the specific meaning of each parameter, all possible candidate values ​​are defined for each parameter to construct a complete joint design space. In this embodiment, the setting of these candidate values ​​comprehensively considers factors such as the general characteristics of the target hardware platform, the limitations of the hardware architecture, and engineering experience in the field.

[0170] For example, as shown in Table 1, this application provides candidate value ranges for each parameter in the joint design space.

[0171] Table 1

[0172] As shown in Table 1, the parameters in the joint design space are mainly divided into three categories: general GEMV optimization parameters, architecture-related parameters, and dequantization-related parameters. Among them, the general GEMV optimization parameters are the first category of parameters mentioned earlier in this application, which may include, but are not limited to, parameters whose names are recursive axis parallelism, non-recursive axis parallelism, loop unrolling dimension, whether to use shared cache, and vector computation width. It should be understood that the meanings of these parameters are as described above and will not be repeated here.

[0173] As shown in Table 1, the candidate value range for the parallelism of the recursive axis and the parallelism of the non-recursive axis can be powers of 2, such as 8, 16, 32, and 64; the candidate value range for the loop unrolling dimension can also be powers of 2, such as 16, 32, 64, and 128; the candidate value range for whether to use a shared cache can include: yes (or 1) or no (or 0); the candidate value range for the vector calculation width can include: 2 and 4. This is only an illustrative example and is not intended to be limiting.

[0174] Architecture-related parameters, namely the second type of parameters mentioned above in this application, may include, but are not limited to, parameters whose names are thread bundle / subgroup size and data aggregation loading width, respectively.

[0175] As shown in Table 1, the candidate value range for the thread beam / subgroup size can include 8, 16, and 32. It should be noted that this candidate value range is primarily applicable to GPU architectures that support configurable subgroup sizes (such as Intel GPUs); for Nvidia GPUs with fixed thread beam sizes, this parameter can be set to 32 in subsequent hardware constraint filtering; for AMD GPUs with a wavefront size of 64, it can be set to 64.

[0176] The candidate value range corresponding to the data aggregation loading width can be dynamically determined by the product of the thread bundle / subgroup size and the quantized data packing bit width, reflecting the linkage between the second type of parameter and the third type of parameter.

[0177] The dequantization-related parameters are the third type of parameters mentioned above in this application, which may include, but are not limited to, parameters whose names are quantization weight data type, quantization weight data bit width, and quantization data packing bit width. As shown in Table 1, the candidate value range corresponding to the quantization weight data type may include: signed int, unsigned int; the candidate value range corresponding to the quantization weight data bit width may include: 2, 3, 4, etc.; the candidate value range corresponding to the quantization data packing bit width may include: 8, 16, 32, 64, etc.

[0178] The candidate values ​​for the above parameters are for illustrative purposes only and can be adjusted appropriately according to the specific characteristics and optimization requirements of the target hardware platform in practical applications. After the candidate values ​​are defined, the joint design space composition module will perform subsequent parameter combination generation and hardware constraint screening operations based on these parameters and their candidate values.

[0179] (III) Generation of parameter value combinations The joint design space building block can generate a set containing all parameter combinations by performing a Cartesian product on all possible candidate values ​​of all parameters through enumeration or combination.

[0180] Specifically, for the first type of parameters (such as general GEMV optimization parameters), candidate values ​​of parameters such as TR, TC, loop unrolling factor, memory cache enabled, and vector calculation width are combined; for the second type of parameters (such as architecture-related parameters), candidate values ​​of parameters such as thread bundle / subgroup size and data aggregation loading width are included in the combination; for the third type of parameters (such as dequantization-related parameters), candidate values ​​of parameters such as quantization weight data type, quantization weight data bit width, and quantization data packing bit width are included in the combination.

[0181] For example, assuming there are 5 parameters of the first type (4 values ​​on average), 2 parameters of the second type (2 to 3 values ​​on average), and 3 parameters of the third type (2 to 4 values), the total number of combinations is approximately several thousand to tens of thousands, and each combination of parameter values ​​corresponds to a potential operator implementation scheme.

[0182] (iv) Hardware constraint screening The joint design space composition module can perform preliminary screening of the generated set of parameter value combinations based on the hardware constraints of the target hardware platform, and eliminate parameter value combinations that do not meet the hardware constraints.

[0183] Among them, hardware constraints may include, but are not limited to, the limit / constraint on the size of the execution unit of the target hardware platform, the upper limit constraint on the total amount of computing resources of the target hardware platform, the relationship constraint between the parallelism of the operator to be optimized and the thread bundle or subgroup size parameter, alignment constraints, etc.

[0184] For example, the limitation / constraint on the execution unit size of the target hardware platform can be understood as a thread bundle / subgroup size constraint. The thread bundle / subgroup size is the basic parallel execution unit size of GPU hardware, and different GPU architectures have different limitations on the value of this parameter.

[0185] For example, for the first type of GPU architecture (such as Nvidia GPUs), the thread bundle size is fixed at 32, so all parameter value combinations where the thread bundle / subgroup size parameter is not equal to 32 are eliminated. For the second type of GPU architecture (such as Intel GPUs), the subgroup size can be configured to 8, 16, or 32, so all parameter value combinations where the thread bundle / subgroup size parameter does not belong to {8, 16, 32} are eliminated. Optionally, for different generations of Intel GPUs, the subset of values ​​actually supported by the current hardware platform can be further filtered by querying hardware information or a preset compatibility list. For the third type of GPU architecture (such as AMD GPUs), the wavefront size is typically 64, so all parameter value combinations where the thread bundle / subgroup size parameter is not equal to 64 are eliminated.

[0186] The upper limit constraint on the total computing resources of the target hardware platform can be understood as a maximum parallelism constraint. GPU hardware typically has maximum parallelism limits, such as the maximum number of threads per thread block and the maximum number of parallel computing units supported by the device.

[0187] Taking the second type of GPU architecture as an example, its design document explicitly supports a maximum parallelism of 1024. Therefore, it is necessary to calculate the parallelism for each parameter value combination, which is determined by the product of the recursive axis parallelism and the non-recursive axis parallelism (e.g., Parallelism = TR × TC). If Parallelism is greater than the maximum parallelism limit of the target hardware platform (e.g., Parallelism > 1024), then the parameter value combination is discarded. Conversely, if Parallelism ≤ 1024, then the parameter value combination is retained.

[0188] The core constraint on the relationship between the parallelism of the operator to be optimized (i.e., the total number of threads in the thread block) and the thread bundle / subgroup size parameter is the matching constraint between the parallelism of a single thread block and the thread bundle / subgroup size. For example, to ensure performance, TRTC is usually set reasonably so that the product of the two is an integer multiple of the thread bundle / subgroup size (i.e., the total number of threads in the thread block is an integer multiple of the thread bundle or subgroup size parameter), thereby ensuring that the parallelism of the operator to be optimized (total number of threads in the thread block) is divisible by the thread bundle / subgroup size, thus improving computational efficiency.

[0189] For example, for the first type of GPU architecture (such as Nvidia GPUs), the constraint rule can be set as recursive axis parallelism = thread bundle / subgroup size, with a constraint value of 32. For the second type of GPU architecture, corresponding matching rules can be set according to the actual hardware characteristics. For example, the recursive axis parallelism can be set to an integer multiple of the subgroup size to fully utilize the collaborative computing capabilities of threads within the subgroup. For the third type of GPU architecture (such as AMD GPUs), the constraint rule can be set as recursive axis parallelism = thread bundle / subgroup size, with a constraint value of 64. If the parameter value combination does not satisfy this matching relationship, it can be eliminated or treated as a candidate combination with poor performance and given lower priority in subsequent searches.

[0190] Alignment constraints primarily concern the matching relationship between the quantized data packet width and the memory address alignment. For example, when the quantized data packet width is 16 bits, the data address must be aligned to 16 bits; when the quantized data packet width is 32 bits, the data address must be aligned to 32 bits. If parameters such as the vector calculation width in the parameter value combinations cause memory access to fail to meet the alignment requirements, these parameter value combinations can be eliminated to reduce the number of parameter combinations that need to be evaluated subsequently.

[0191] After the aforementioned hardware constraint screening, the joint design space composition module eliminates combinations that do not meet the hardware constraints from the original joint design space (i.e., the set of all possible parameter combinations), resulting in a valid set of parameter value combinations. This set can be stored in memory using data structures (such as lists or dictionaries) for use by the subsequent design space exploration control module. This screening process effectively reduces the number of parameter combinations that need to be evaluated subsequently, improves optimization efficiency, and ensures that all parameter combinations evaluated are within the limits allowed by the hardware constraints.

[0192] Through the above operations, the joint design space construction module integrates the previously fragmented general GEMV optimization parameters, architecture-related parameters, and dequantization-related parameters into a single design space, laying a solid foundation for subsequent collaborative optimization. Simultaneously, by using hardware constraint filtering to preemptively eliminate invalid parameter combinations, the search scope is effectively reduced, significantly improving overall optimization efficiency while ensuring optimization effectiveness. Compared to traditional automatic optimization methods that only consider general scheduling parameters, the design space constructed by this module includes architecture-related and dequantization-related parameters, providing richer search dimensions for discovering the global optimum and creating conditions for subsequent modules to generate cross-platform high-performance operator code.

[0193] Design Space Exploration Control Module The design space exploration control module is the central control hub of the entire operator automatic optimization process, responsible for coordinating other modules to complete the automated search process. This module obtains a set of valid parameter value combinations from the joint design space composition module, traverses each parameter value combination in a preset order, completes code generation and performance evaluation by calling the operator code generation module and the operator code compilation and execution module, and determines the target parameter combination (such as the optimal parameter combination) from all candidate parameter value combinations.

[0194] For example, the design space exploration control module can be used to perform the following operations: (a) Traversal initialization and parameter value combination retrieval This module first receives a set of valid parameter value combinations from the joint design space composition module. This set is stored in memory in the form of a data structure (such as a list or dictionary), which contains all feasible parameter value combinations after being filtered by hardware constraints.

[0195] This module numbers or indexes the elements in the set according to a preset traversal order (such as the index order generated by parameter value combinations), and then sequentially selects one parameter value combination as the current candidate solution to be evaluated. The selected parameter value combination can be a tuple or dictionary containing all the specific values ​​of the parameters, such as: { 'TR': 8, / / The parallelism of the recursive axis is 8 'TC': 8, / / The parallelism of non-recursive axes is 8 'UNROLL': 1, / / The value of the loop expansion factor is 1 'USE_SMEM': 0, / / Disable shared cache 'VEC_WIDTH': 1, / / The value of the vector's calculated width is 1. 'SUB_GROUP_SIZE': 8, / / The value for thread bundle / subgroup size is 8 'QDATA_TYPE': 'SIGNED', / / The quantization weight data type is a signed integer. 'QBITS': 4, / / The quantization weight data bit width is 4 bits. 'PACK_WIDTH': 16 / / The quantized data packet width is 16 bits. } This is merely an illustrative example and is not intended to be limiting.

[0196] (II) Code generation triggering and performance evaluation This module sends the retrieved parameter value combinations to the operator code generation module, triggering the code generation process. The operator code generation module loads the received parameter value combinations into a parameterized template, assigns values ​​to the variables in the parameterized template, and generates an instantiated operator template.

[0197] The operator code generation module sends the instantiated operator templates to the operator code compilation and execution module, which then performs compilation and performance evaluation. Specifically, the operator code compilation and execution module selects the corresponding compilation backend based on the type of the target hardware platform, such as the Open Computing Language (OpenCL) backend, the Compute Unified Device Architecture (CUDA) backend, or the Radeon Open Compute (Rocm) backend, and compiles the instantiated operator templates into executable code that can be directly executed on the target hardware platform.

[0198] Alternatively, during the search process, the instantiated operator template can be submitted to the virtual machine engine built into the TensorVirtual Machine (TVM) (such as TVM.relax.VirtualMachine) to run the generated executable code in the virtual machine environment and evaluate its execution performance, thereby quickly completing the initial screening without relying on actual hardware.

[0199] Subsequently, the executable code is run on the target hardware platform or virtual machine, its execution performance is measured, and the obtained execution performance data is returned to the design space exploration control module. The execution performance may include metrics such as inference latency, computing power utilization, data transmission efficiency, and throughput of the operator on the target hardware platform.

[0200] Optionally, in one possible implementation, to eliminate the impact of system noise and environmental fluctuations on performance measurement, the operator code compilation and execution module can run the generated executable code multiple times and take the average of the results from these multiple runs as the final execution performance. For example, it can run the code 10, 20, or 50 times consecutively and calculate the average inference latency as the execution performance to be evaluated. For scenarios with high real-time requirements, the minimum value (such as the optimal inference latency) from multiple runs can be taken, or the average can be taken after removing outliers through statistical analysis to obtain a more stable performance evaluation result.

[0201] Accordingly, during the search process, for each combination of parameter values, the design space exploration control module can record and compare these execution performances to determine a better parameter configuration. When a large number of combinations need to be quickly filtered, the search can be run only once to improve efficiency, and then more refined multi-round evaluations can be performed after candidate value combinations are selected. This implementation method can be flexibly chosen according to actual needs and is not limited thereto.

[0202] (III) Determination of target parameter combination The design space exploration control module evaluates the execution performance of each parameter value combination and then determines the target parameter combination (such as the optimal parameter combination).

[0203] In one possible implementation, a pairwise comparison and selection of the optimal parameter combination can be used to determine the target parameter combination. For example, for the first parameter value combination, it is directly used as the currently retained parameter value combination, and its corresponding execution performance is recorded. For each subsequent parameter value combination, its execution performance is compared with that of the currently retained parameter value combination. If the execution performance of the current parameter value combination is better (e.g., lower inference latency or higher throughput), then the retained parameter value combination is updated with the current parameter value combination, and the original retained parameter value combination is discarded. If the execution performance of the currently retained parameter value combination is better, then the currently retained parameter value combination is retained, and the current parameter value combination is discarded.

[0204] After traversing all parameter value combinations, the final retained parameter value combination is the target parameter combination. This method only needs to save one optimal parameter combination and its execution performance during the evaluation process, without storing the execution performance evaluation results of all combinations, thus significantly reducing memory overhead. It is suitable for scenarios with a large search space and a large number of parameter combinations.

[0205] In this embodiment, the search space refers to the set of all possible combinations of parameter values, that is, the range of all candidate solutions that need to be traversed and evaluated during the automated optimization process of the operator. Specifically, the joint design space construction module combines all possible candidate values ​​of the first type of parameter, the second type of parameter, and the third type of parameter by performing a Cartesian product to form a set containing all combinations of parameter values, which is the search space. After hardware constraint screening, combinations that do not meet the hardware constraints are eliminated, resulting in an effective search space. The design space exploration control module traverses and evaluates this effective search space to determine the optimal parameter combination.

[0206] Alternatively, another possible implementation approach is to use offline evaluation followed by unified selection. This involves first evaluating all parameter value combinations and collecting all execution performance data. Then, statistical analysis or machine learning methods are used to select the optimal parameter value combination as the target parameter combination. This approach provides a comprehensive understanding of the execution performance distribution of all parameter value combinations, facilitating subsequent performance analysis and debugging. It is suitable for scenarios with a small search space or where a full analysis of the evaluation results is required.

[0207] (iv) Traversal loop and search termination After evaluating and comparing the current parameter value combinations, the space exploration control module determines whether there are any unevaluated parameter value combinations. If there are still unprocessed combinations in the valid set of parameter value combinations, it returns to step (1) to continue traversing the next combination; if all combinations have been evaluated, the search process terminates.

[0208] After the search process terminates, the corresponding subsequent processing is performed according to the target parameter combination determination method adopted in step (III). For example, if a pairwise comparison and optimal selection method is used, the currently retained optimal parameter combination and its execution performance have been dynamically maintained during the traversal process. After the search terminates, the currently retained parameter value combination is the target parameter combination. If a method of completing all evaluations first and then making a unified selection is adopted, the execution performance of all parameter value combinations has been collected when the search terminates. At this time, the parameter value combination with the best execution performance is selected from the collected results as the target parameter combination.

[0209] By using the above-mentioned traversal loop and search termination methods, the space exploration control module achieves complete coverage of the effective search space, ensuring that no potential optimization opportunities are missed. At the same time, after the search is completed, it can correctly obtain the target parameter combination according to the selected determination method and terminate the loop in a timely manner to avoid invalid repeated calculations.

[0210] Through the above operations, the space exploration control module achieves fully automated parameter search and optimal solution discovery. This means that without manual intervention, it can automatically complete traversal and optimal solution discovery within massive parameter combinations, significantly improving operator development efficiency. Furthermore, the module adopts a modular design, allowing search algorithms (such as grid search, genetic algorithms, Bayesian optimization, etc.) to be flexibly replaced or extended according to actual needs. Only the implementation logic of this module needs to be adjusted to adapt to different search strategies, providing excellent scalability for diverse optimization scenarios.

[0211] Operator code generation module This operator code generation module generates instantiated operator templates based on parameter value combinations, providing input for subsequent compilation and execution. This module includes parameterized templates, which are operator code frameworks containing variable parameter placeholders and predefined dequantization logic branches.

[0212] For example, the operator code generation module can be used to perform the following operations: (a) Parametric Template Design The parameterized template in this application embodiment may include, but is not limited to, the following elements: Variable placeholders: The parameterized template defines variable placeholders corresponding to the first and second types of parameters. For example, the template of the GEMV operator contains variables such as TR, TC, UNROLL, USE_SMEM, VEC_WIDTH, and SUB_GROUP_SIZE, which are used to control the parallelism of the recursive axis, the parallelism of the non-recursive axis, the number of loop unrolls, the use of shared memory, the width of vectorized computation, and the size of the thread bundle / subgroup, respectively.

[0213] Predefined dequantization logic branches: The template predefines multiple sets of dequantization logic branches (or dequantization execution paths), each corresponding to different combinations of third-type parameters. For example, these branches may include signed type dequantization logic branches (or signed type dequantization paths) and unsigned type dequantization logic branches (or unsigned type dequantization paths).

[0214] Among them, the signed type dequantization logic branch is suitable for scenarios where the quantization weight data type is a signed integer, and it can include sign extension logic (such as...). Figure 5B As shown in the figure, the sign extension process is illustrated using int4 as an example. The unsigned type dequantization logic branch is suitable for scenarios where the quantization weight data type is an unsigned integer. It can include value range adjustment (subtracting constants) logic (such as...). Figure 5A As shown in the figure, the value range adjustment process is illustrated using uint4 as an example.

[0215] Unpacking logic for different packet widths: Based on the different packet widths of quantized data and quantized weight data, corresponding bitmasks and shift operations are generated. For example, when the packet width of quantized data is 16 and the packet width of quantized weight data is 4, the unpacking logic extracts four 4-bit data from the 16-bit data; when the packet width of quantized data is 32 and the packet width of quantized weight data is 4, it extracts eight 4-bit data.

[0216] Thread scheduling and resource binding instructions: The parameterized template can also include instructions related to thread scheduling and resource binding, used to map computational tasks to execution units on the target hardware platform. For example, for GPU architectures that support subgroup size configuration (such as Intel GPUs), the subgroup size can be set through the `intel_reqd_sub_group_size` attribute of OpenCL, allowing the compiler to generate SIMD instructions of the corresponding width; for GPU architectures with fixed thread bundle sizes (such as Nvidia GPUs), the thread index is bound to the computational task through CUDA's `threadIdx`, and the hardware automatically handles the thread bundle partitioning.

[0217] (ii) Instantiation of operator template generation After receiving the parameter value combination from the design space exploration control module, the operator code generation module can perform the following operations: Parameter assignment: Assign the values ​​of the first and second types of parameters to the corresponding variable placeholders in the parameterization template. Specifically, assign the values ​​of the first type of parameters, such as the parallelism of the recursive axis, the parallelism of the non-recursive axis, the loop unrolling factor, the memory cache enable parameter, and the vector calculation width, as well as the values ​​of the second type of parameters, such as the thread bundle / subgroup size, to the corresponding variable placeholders in the parameterization template.

[0218] For example, the value of TR is assigned to the variable BLOCK_SIZE_K, which controls the K-axis cyclic block partitioning; in the OpenCL template, the value of SUB_GROUP_SIZE is set to the subgroup size through the intel_reqd_sub_group_size attribute; in the CUDA template, the thread bundle partitioning is indirectly affected by configuring the thread block dimension.

[0219] Dequantization logic branch selection: Based on the value of the third type of parameter, the target branch is selected from multiple predefined dequantization logic branches in the template. Specifically, based on the quantization weight data type parameter, a signed or unsigned dequantization logic branch is selected; based on the quantization weight data bit width and quantization data packing bit width, the mask bit width and shift step size in the unpacking logic are determined; based on the assigned variables and the selected dequantization logic branch, this module can generate a complete, compilable, instantiated operator template (i.e., an operator implementation in source code form). In this instantiated operator template, all parameter placeholders have been replaced with specific values, the dequantization logic branch has been determined, and it can be directly sent to the compiler for compilation.

[0220] Through the above operations, the operator code generation module separates adjustable parameters from core computational logic, enabling the same set of parameterized templates to adapt to different parameter configurations and improving code reusability. By predefining dequantization logic branches in the template and dynamically selecting the target branch based on the third type of parameters, the dequantization operation is deeply integrated with matrix-vector multiplication calculations, laying the foundation for eliminating data round-trip overhead. At the same time, the template supports the characteristics of different hardware platforms (such as thread bundle size settings) through conditional compilation or template specialization, allowing the generated instantiated operator templates to naturally adapt to the target hardware characteristics.

[0221] Operator code compilation and execution module The operator code compilation and execution module is used to compile the instantiated operator template into binary code executable by the target hardware platform, run it on the target hardware platform to measure its execution performance, and return the measured execution performance to the design space exploration control module.

[0222] For example, the operator code compilation and execution module can be used to perform the following operations: (I) Backend Selection and Code Compilation After receiving the instantiated operator template from the operator code generation module, the operator code compilation and execution module selects the corresponding compilation backend based on the type of the target hardware platform.

[0223] For example, when the target hardware platform is an Intel GPU, the OpenCL backend is selected. This module can call an OpenCL compiler (such as clBuildProgram) to compile the instantiated operator template into executable binary code for the Intel GPU. During this compilation process, the intel_reqd_sub_group_size attribute set in the template instructs the compiler to generate SIMD instructions for the corresponding subgroup size.

[0224] For example, when the target hardware platform is an Nvidia GPU, the CUDA backend is selected. This module calls the Nvidia CUDA Compiler (NVCC) to compile the instantiated operator template into Parallel Thread Execution (PTX) code or a cubin file. Since the thread bundle size of Nvidia GPUs is fixed at 32, there is no need to explicitly set the thread bundle size in the template; the compiler will automatically divide the thread bundle according to the thread block dimension configuration.

[0225] For example, when the target hardware platform is an AMD GPU, the RoCM backend is selected. This module calls the Heterogeneous Interface Compiler (HIP) to compile the instantiated operator templates into executable code optimized for AMD GPUs. The wavefront size of AMD GPUs is typically 64, and this compiler can automatically handle wavefront partitioning based on the thread block dimension configuration.

[0226] Optionally, during compilation, the compiler can perform low-level optimizations based on the specific parameter values ​​in the instantiated operator template. For example, it can perform loop unrolling based on parameters such as UNROLL to reduce loop control overhead; automatically generate aggregate memory access instructions based on thread bundle / subgroup size and memory access continuity (i.e., whether multiple threads access contiguous memory addresses); and reorder instructions to hide memory access latency and improve pipeline efficiency.

[0227] (ii) Performance Evaluation For example, after compilation, the operator code compilation and execution module runs the generated executable code on the target hardware platform or virtual machine to measure its execution performance. The measured execution performance can include performance indicators such as inference latency, computing power utilization, data transmission efficiency, and throughput of the operator on the target hardware platform. For example, for LLM inference scenarios, inference latency is the most critical indicator, and the average time taken for a single operator execution can be measured as an evaluation value for inference latency.

[0228] Afterwards, the operator code compilation and execution module can return the measured execution performance to the design space exploration control module, which can then compare and filter parameter value combinations to determine the target parameter combination.

[0229] Through the above operations, this module runs code and measures performance on real target hardware, ensuring the accuracy and reliability of the evaluation results and avoiding errors that may be introduced by simulators or theoretical models. This module not only provides latency metrics but also multi-dimensional performance information such as throughput and computing power utilization, providing data support for developers to deeply analyze operator bottlenecks. Simultaneously, this module fully utilizes the optimization capabilities of the underlying compiler, transforming high-level parameter configurations into efficient low-level instruction sequences, ensuring accurate performance evaluation results for the design of the space exploration control module.

[0230] Operator code output module The operator code output module is used to generate and output the target code that the operator will ultimately execute on the target hardware platform based on the combination of target parameters determined by the design space exploration control module, so that users can call it in actual inference tasks.

[0231] Optionally, this module can also output the code corresponding to other parameter combinations generated during the search process for debugging, analysis, or subsequent reuse.

[0232] For example, the operator code output module can be used to perform the following operations: (a) Saving of target parameter combinations The operator code output module receives the target parameter combination sent from the design space exploration control module. This target parameter combination is the globally optimal solution determined through automated search, and it includes the specific values ​​of the first type of parameter, the second type of parameter, and the third type of parameter.

[0233] This module can save target parameter combinations to disk in a readable format for user reference and analysis. The saving format can be JavaScript Object Notation (JSON), Yet Another Markup Language (YAML), or a custom configuration file, and the saved content can include the specific values ​​of the optimal parameter combination and its corresponding execution performance.

[0234] Optionally, this module can also save some or all of the parameter combinations evaluated during the search process and their corresponding execution performance, making it easier for developers to perform performance analysis and debugging.

[0235] (II) Target Code Generation and Output This module loads the target parameter combination into the parameterized template and executes the same process as the operator code generation module to obtain the instantiated operator template (or target operator template) corresponding to the target parameter combination.

[0236] Specifically, the values ​​of the first and second types of parameters are assigned to the corresponding variable placeholders in the parameterized template; based on the value of the third type of parameter, the target dequantization logic branch is selected from multiple predefined dequantization logic branches in the template; based on the assigned variables and the selected target dequantization logic branch, the target operator template (i.e., the final implementation in source code form) is generated.

[0237] This module selects the corresponding compilation backend based on the type of the target hardware platform, compiles the target operator template into the final executable code (i.e., target code), and saves it to the disk in the backend format corresponding to the target hardware platform.

[0238] For example, for the second type of GPU architecture (such as Intel GPUs), the output is OpenCL kernel code (.cl file) or a pre-compiled binary file. For the first type of GPU architecture (such as Nvidia GPUs), the output is CUDA PTX code (.ptx file) or a cubin file. For the third type of GPU architecture (such as AMD GPUs), the output is a Rocm executable file.

[0239] Alternatively, in one possible implementation, the module can also directly output the target operator template (in source code form) instead of the compiled binary code, so that developers can make secondary modifications or integrate it into other compilation processes.

[0240] Optionally, in debug mode or scenarios requiring comprehensive performance analysis, this module can be configured to output executable code corresponding to some or all of the parameter combinations evaluated during the search process. Specifically, for each evaluated parameter value combination, its corresponding executable code is saved to disk and named according to the index of the parameter value combination or the parameter value for easy reference later. Simultaneously, the evaluation results corresponding to each parameter value combination are saved, forming a complete record of the optimization process.

[0241] This approach helps developers intuitively understand the impact of different parameter value combinations on performance, providing data support for subsequent manual optimization or algorithm improvement. It should be understood that this feature can be enabled selectively based on actual needs, and there are no restrictions on its implementation.

[0242] Through the above operations, the module saves the final result of automatic optimization as reusable code and configuration for users to call in actual inference tasks; it also supports outputting multiple backend formats, so that the same automatic optimization process can generate corresponding executable code for different hardware platforms; at the same time, it saves the configuration of the optimal parameter combination, which is convenient for developers to perform subsequent analysis and debugging.

[0243] The overall process of the automatic operator optimization method provided in the embodiments of this application will be briefly described below with reference to the accompanying drawings.

[0244] Please refer to Figure 7 , Figure 7 This is an exemplary embodiment of the automatic operator optimization method illustrated in this application. The flowchart corresponds to the collaborative relationship of the five modules in the aforementioned system architecture, demonstrating the complete process from the construction of the joint design space to the final code output.

[0245] S101, Joint Design Space Definition.

[0246] For example, this step can be performed by the joint design space construction module. First, all parameters requiring optimization are identified and categorized into three types: a first type of parameters related to computational parallel scheduling, a second type of parameters related to the execution units of the target hardware platform, and a third type of parameters related to the quantization and dequantization data formats of the operator to be optimized. For each parameter, all possible candidate values ​​are defined. By enumerating or combining all possible candidate values ​​of all parameters, a Cartesian product is performed to generate a set containing all combinations of parameter values, i.e., the original joint design space.

[0247] Subsequently, based on the hardware constraints of the target hardware platform (such as thread bundle / subgroup size limits, maximum parallelism limits, and the matching relationship between recursion axis parallelism and thread bundle / subgroup size), the generated parameter value combinations are initially screened, eliminating those that do not meet the hardware constraints, thus obtaining a set of valid parameter value combinations. This set of valid parameter value combinations constitutes the valid joint design space, which is stored in memory in the form of data structures (such as lists or dictionaries) for use by subsequent modules.

[0248] S102. Whether to iterate through all valid parameter value combinations.

[0249] For example, this step can be performed by the design space exploration control module. This module obtains a valid set of parameter value combinations from the joint design space composition module and determines whether all parameter value combinations in the set have been evaluated. If the traversal has not been completed (i.e., there are still unprocessed / unevaluated parameter value combinations), then step S103 is executed; if the traversal has been completed (i.e., all parameter value combinations have been evaluated), then step S107 is skipped.

[0250] S103, retrieve the next set of parameter value combinations.

[0251] For example, this step can be performed by the design space exploration control module. This module retrieves an unevaluated parameter value combination from the set of valid parameter value combinations according to a preset traversal order (such as the index order generated by parameter value combinations), as a candidate solution to be evaluated. It should be understood that the retrieved parameter value combination can be a tuple or dictionary containing all the specific values ​​of the parameters.

[0252] S104. Load into the parameterized template and compile to generate executable code.

[0253] For example, this step can be completed collaboratively by the operator code generation module and the operator code compilation and execution module. Specifically, the design space exploration control module sends the retrieved parameter value combinations to the operator code generation module. The operator code generation module loads the received parameter value combinations into a parameterized template, assigns the values ​​of the first and second types of parameters to the corresponding variable placeholders in the parameterized template, and selects the target dequantization logic branch from multiple predefined dequantization logic branches in the template based on the value of the third type of parameter, generating an instantiated operator template.

[0254] Subsequently, the operator code generation module sends the instantiated operator template to the operator code compilation and execution module. This module selects the appropriate compilation backend based on the target hardware platform type and calls the compiler corresponding to that backend to compile the instantiated operator template into executable code that can be directly executed on the target hardware platform.

[0255] The S105 and TVM virtual machines were run and their execution performance was extracted.

[0256] For example, this step can be performed by an operator code compilation and execution module. This module deploys the generated executable code to the target hardware platform and schedules its execution via a TVM virtual machine (such as the virtual machine engine built into the TVM virtual machine) to measure its execution performance. The measured execution performance can include performance metrics such as inference latency, computing power utilization, data transfer efficiency, and throughput of the operator on the target hardware platform.

[0257] Optionally, the operator code compilation and execution module returns the measured execution performance to the design space exploration control module.

[0258] S106, Update parameter value combinations and execution performance.

[0259] For example, this step can be performed by the design space exploration control module. This module receives the execution performance returned by the operator code compilation and execution module and compares it with the currently reserved best execution performance. The specific comparison method can be found in the previous description of the design space exploration control module, and will not be repeated here.

[0260] Optionally, after completing the comparison and update, return to step S102 to continue determining whether there are any unevaluated parameter value combinations.

[0261] S107, Output target parameter combination.

[0262] For example, this step can be completed collaboratively by the design space exploration control module and the operator code output module. After all parameter value combinations have been evaluated, the design space exploration control module determines the final optimal parameter value combination as the target parameter combination and sends it to the operator code output module.

[0263] The operator code output module saves the target parameter combination to disk in a readable format. The saved content may include the specific values ​​of the optimal parameter combination and the corresponding execution performance.

[0264] S108. Compile to generate target code.

[0265] For example, this step can be performed by the operator code output module. This module loads the target parameter combination into the parameterized template, performs the same process as the operator code generation module, and generates the target operator template (i.e., the final implementation in source code form).

[0266] Subsequently, the module selects the corresponding compilation backend based on the type of the target hardware platform, compiles the target operator template into the target code that the operator will ultimately execute on the target hardware platform, and saves it to the disk in the backend format corresponding to the target hardware platform.

[0267] Through the iterative and collaborative steps described above, this application achieves a fully automated operator optimization process from joint design space construction to target code generation. This process integrates general operator optimization parameters, architecture-related parameters, and dequantization-related parameters into the same design space, achieving collaborative optimization of hardware awareness and quantization fusion. Utilizing automated traversal and evaluation methods, it efficiently locates the optimal configuration from massive parameter combinations without manual intervention. The final executable code integrates weight dequantization and matrix-vector multiplication calculations, eliminating data round-trip overhead and improving operator development efficiency while achieving high-performance execution on different GPU architectures.

[0268] The detailed process of the automatic operator optimization method provided in the embodiments of this application will be described below with reference to the accompanying drawings. Please refer to... Figure 8 , Figure 8 This is a detailed flowchart illustrating an exemplary embodiment of the operator automatic optimization method of this application. Figure 7 Compared to the overall flowchart shown, Figure 8 The overall process is divided into three main stages: "joint design space definition", "template-based operator automatic optimization", and "target code generation". The operation steps involved in each stage are further subdivided to facilitate understanding of the specific implementation details of the technical solution in this application.

[0269] like Figure 8 As shown, the detailed process of the operator automatic optimization method provided in this application embodiment includes the following steps: I. Definition of Joint Design Space S201, Read the parameter list.

[0270] For example, this step can be performed by the joint design space construction module. When the operator automatic optimization process starts, this module first obtains the basic parameter information required to construct the joint design space. For example, it reads the names, meanings, and candidate value ranges of all parameters to be optimized from a preset parameter list.

[0271] The parameter list can be predefined by the user and stored and provided in various forms. This parameter list can include the names, meanings, and candidate value ranges of all parameters to be optimized. For example, the parameters to be optimized may include general GEMV optimization parameters, architecture-related parameters, dequantization-related parameters, etc. The names, meanings, and candidate value ranges of each parameter can be found in the previous descriptions and will not be repeated here.

[0272] For example, the parameter list can be pre-installed in the system code and stored in the local code or configuration file in the form of a data structure (such as a dictionary, mapping table, etc.). In this implementation, the dictionary keys are parameter names, and the values ​​are structured objects containing information such as the parameter meaning and a list of candidate value ranges. This parameter list can be pre-embedded in the system code, allowing the system to quickly load the default configuration each time it starts. At the same time, this parameter list also supports dynamic adjustment by users, such as by modifying the local configuration file or calling system interfaces to add, delete, modify, and query the candidate value range, thereby flexibly adapting to different hardware platforms and optimization needs.

[0273] It is worth noting that an exemplary parameter list may be shown in Table 1 above, but this application is not limited to this, and users may adjust it according to their actual needs.

[0274] Alternatively, the parameter list can be provided by the user through an external configuration file (such as JSON or YAML), or dynamically obtained from a cloud server at system startup. This flexible parameter acquisition method allows the method in this application to quickly adapt to different optimization needs and hardware platforms without modifying the core code.

[0275] For example, when the operator automatic optimization process starts, the joint design space construction module reads the parameter list and loads it into memory. The read operation can be performed only once, with subsequent steps directly using the data in memory. After the optimization process ends, the parameter list data in memory can be cleaned up and released, but the parameter list definition stored in the source code or configuration file is retained for use in the next optimization. This method of loading the parameter list into memory and cleaning it up after optimization effectively controls system resource consumption and avoids unnecessary memory overhead.

[0276] Through the above methods, this application achieves flexible definition and dynamic loading of the parameter list. The pre-set parameter list provides users with a default configuration out of the box, while also allowing users to adjust it according to actual needs to meet diverse optimization scenarios; and the method of reading the parameter list from local or cloud sources enables the method of this application to flexibly adapt to different hardware platforms and quantization schemes without modifying the core code.

[0277] S202, Generate all parameter value combinations.

[0278] For example, this step can be performed by the joint design space composition module. After reading the parameter list, the module enumerates and combines all possible candidate values ​​of all parameters based on the acquired basic parameter information to generate a set containing all parameter value combinations.

[0279] Specifically, this module combines all candidate values ​​of the first, second, and third types of parameters using a Cartesian product to generate a set containing all possible parameter value combinations. Each element in this set is a parameter value combination, corresponding to a potential operator implementation scheme.

[0280] For example, all possible values ​​of all parameters are traversed in sequence to generate all combinations of parameter values, and all combinations of parameter values ​​are put into a set (which can be the set of parameters to be tuned), or in other words, all combinations of parameter values ​​constitute the set of parameters to be tuned.

[0281] For example, in parameter value combination 1, the parameter "recursive axis parallelism" takes the first possible value "8", the parameter "non-recursive axis parallelism" takes the first possible value "8", and all other parameters take the first possible value, forming parameter combination 1: { Parallelism of recursive axes: 8 Non-recursive axis parallelism: 8, Loop unrolling dimension: 16, Use shared cache: Yes Vector calculation width: 2, Thread bundle / subgroup size: 8 Quantization weight data type: Signed integer. Quantization weight data bit width: 2, Quantized data packet width: 8 bits } For example, in parameter value combination 2, the parameter "recursive axis parallelism" takes the second possible value "16", while the other parameters still take the first possible value, forming parameter combination 2: { Parallelism of recursive axes: 16 Non-recursive axis parallelism: 8, Loop unrolling dimension: 16, Use shared cache: Yes Vector calculation width: 2, Thread bundle / subgroup size: 8 Quantization weight data type: Signed integer. Quantization weight data bit width: 2, Quantized data packet width: 8 bits } Following the above pattern, all possible values ​​of all parameters are iterated as needed to obtain a set containing all possible combinations of parameter values. The total number of parameter value combinations in this set is equal to the product of the number of candidate values ​​for each parameter. For example, assuming there are 5 parameters of the first type (an average of 4 values), 2 parameters of the second type (an average of 2 to 3 values), and 3 parameters of the third type (2 to 4 values), the total number of combinations is approximately several thousand to tens of thousands. This set is the original joint design space, containing all theoretically possible operator implementation schemes, providing a complete candidate range for subsequent automated search and optimization.

[0282] Alternatively, in one possible implementation, recursive generation or iterator patterns can be used to generate combinations of parameter values ​​to avoid consuming too much memory at once when the number of combinations is extremely large. This application does not limit this approach.

[0283] S203, Hardware Constraint Screening.

[0284] For example, this step can be performed by the joint design space composition module. After generating a set of all parameter value combinations, the module performs a preliminary screening of the parameter value combinations based on the hardware constraints of the target hardware platform, eliminating combinations that do not meet the hardware constraints, and obtaining a set of valid parameter value combinations.

[0285] Hardware constraints are a set of rules pre-determined based on the design specifications and architectural characteristics of the target hardware platform. These rules determine whether combinations of parameter values ​​fall within the range supported by the hardware platform. These constraints can include various types, such as execution unit size constraints, total computing resource limits, and parallelism matching constraints. Different constraints and constraint values ​​can be set for different hardware platforms.

[0286] For example, for the first type of GPU architecture (such as Nvidia GPUs), it can be seen from consulting the Nvidia GPU platform that its architecture typically only supports a fixed thread bundle size of 32. Therefore, the following constraint rule can be set: Execution unit size constraint: thread bundle / subgroup size = 32; Parallelism matching constraint: to ensure performance, the recursion axis parallelism = thread bundle / subgroup size, that is, the recursion axis parallelism must also be equal to 32. Or, the constraint value is "recursion axis parallelism = thread bundle / subgroup size = 32".

[0287] For example, for the second type of GPU architecture (such as Intel GPUs), by consulting its design manual, we know that this architecture supports subgroup sizes of configurable 8, 16, or 32, with different generations potentially supporting different subsets, and the maximum parallelism supported by this architecture is 1024. Therefore, the following constraint rules can be set: Execution unit size constraint: The thread bundle / subgroup size parameter must be {8, 16, 32}, and can be further filtered according to the actual hardware generation; Maximum constraint on total computing resources: Parallelism = Parallelism of recursive axis × Parallelism of non-recursive axis, with a constraint value of "parallelism ≤ 1024".

[0288] For example, for the third type of GPU architecture (such as AMD GPUs), consulting their design manuals reveals that the wavefront size is typically 64. Therefore, the following constraint rules can be set: Execution unit size constraint: thread bundle / subgroup size = 64; Parallelism matching constraint: to ensure performance, the recursion axis parallelism = thread bundle / subgroup size, that is, the recursion axis parallelism must also be equal to 64. Or, the constraint value is "recursion axis parallelism = thread bundle / subgroup size = 64".

[0289] Optionally, these constraints and constraint values ​​can be pre-defined by developers based on hardware design documents and embedded in the system's constraint configuration. Alternatively, a mapping database between hardware platforms and constraint rules can be established, and constraints can be dynamically retrieved by querying the database, thus adapting to new hardware platforms without modifying the code. Furthermore, constraints can also be dynamically distributed via cloud services, enabling flexible adaptation to different hardware platforms.

[0290] For example, the module iterates through the set of parameter value combinations generated in step S202 above, and applies the above hardware constraints to each parameter value combination in turn for judgment. For any parameter value combination, if it does not satisfy any of the hardware constraints, it is removed from the set.

[0291] For example, on the Intel GPU platform, the set of parameter value combinations is traversed, and the parallelism of each combination is calculated. For a certain parameter value combination, if the parallelism of the recursive axis is 32 and the parallelism of the non-recursive axis is 64, the calculated parallelism is 32 × 64 = 2048, which is greater than the maximum parallelism of 1024, so the combination is discarded.

[0292] Optionally, combinations of thread bundle / subgroup size parameters that do not belong to {8, 16, 32} are also excluded.

[0293] For example, for the Nvidia GPU platform, it iterates through the set of parameter value combinations, removing combinations where the thread bundle / subgroup size parameter value is not equal to 32, and also removing combinations where the recursion axis parallelism is not equal to 32.

[0294] For example, for the AMD GPU platform, it iterates through the set of parameter value combinations, removes combinations where the thread bundle / subgroup size parameter value is not equal to 64, and also removes combinations where the recursive axis parallelism is not equal to 64.

[0295] After filtering based on the aforementioned hardware constraints, this module eliminates combinations that do not meet the hardware constraints from the original joint design space (i.e., the set of all possible parameter value combinations), resulting in a valid set of parameter value combinations. This set can also be stored in memory as a data structure for use by the subsequent design space exploration and control module.

[0296] Through the aforementioned hardware constraint filtering operations, this application achieves preprocessing optimization of the search space. Eliminating invalid parameter combinations in advance effectively reduces the number of parameter combinations that need to be evaluated subsequently, significantly improving overall optimization efficiency. Simultaneously, it ensures that all parameter combinations evaluated later are within the limits allowed by hardware constraints, avoiding compilation errors or runtime anomalies caused by invalid parameter combinations. Furthermore, by presetting constraint rules and values, an understanding of hardware characteristics is integrated into the automated process, achieving hardware-aware optimization preprocessing. This design not only ensures the reliability of the optimization process but also creates favorable input conditions for the automated search of subsequent modules.

[0297] II. Automatic Optimization of Operators Based on Templates S204. Are there any unevaluated combinations of parameter values?

[0298] For example, this step can be performed by the design space exploration control module. After completing the screening of hardware constraints and obtaining a set of valid parameter value combinations, the module can number or index the elements in the set according to a preset order (such as the order of the index generated by the combination or the order in which they were added to the set) to complete the traversal initialization.

[0299] Subsequently, parameter value combinations are extracted sequentially from the set of valid parameter value combinations. For example, the first parameter value combination is extracted, and based on this parameter value combination, operations such as loading into the parameterized template, calling the compiler to generate executable code, and TVM virtual machine evaluating execution performance and updating the current optimal parameter combination are executed sequentially (i.e., steps S205 to S209).

[0300] After evaluating a parameter value combination (i.e., after executing steps S205 to S209), this module can determine whether there are any unevaluated parameter value combinations in the set of valid parameter value combinations. Alternatively, after evaluating a parameter value combination, this module returns to this step to determine whether there is another unevaluated parameter value combination in the set of valid parameter value combinations.

[0301] If there are unevaluated parameter value combinations, extract the next set of parameter value combinations and execute the following steps S205 to S209; if there are no unevaluated parameter value combinations (i.e. all valid parameter value combinations have been evaluated), proceed to step S210.

[0302] In this way, the module continuously maintains the current parameter value combinations to be evaluated during the traversal process, and decides whether to continue traversing or terminate the loop based on the evaluation completion status, ensuring complete coverage of the effective search space.

[0303] S205, Extract the next set of parameter value combinations.

[0304] For example, this step can be performed by the design space exploration control module. This step can be performed when step S204 determines that there are still unevaluated parameter value combinations in the set of valid parameter value combinations.

[0305] Specifically, this module retrieves the next unevaluated parameter value combination from the set of valid parameter value combinations in a preset order, as the current candidate solution to be evaluated. Then, based on the currently retrieved parameter value combination, it performs the operation of loading it into the parameterized template.

[0306] In this embodiment, one specific implementation of the parameterized template can be a fusion operator template of inverse quantization and matrix-vector multiplication. Please refer to... Figure 9 , Figure 9 This is a schematic diagram of an inverse quantization and matrix-vector multiplication fusion operator template shown in an exemplary embodiment of this application.

[0307] like Figure 9 As shown, this application introduces architecture-related parameters and dequantization-related parameters (marked in bold underline in the figure) on the basis of the vector matrix multiplication operator template provided by TVM (which only contains general GEMV optimization parameters), thus forming a dequantization and matrix vector multiplication fusion operator template. Its calculation process can include the following core steps: input loading, weight loading, weight unpacking, weight dequantization, vector multiplication (partial sum), partial sum accumulation, etc.

[0308] The input loading process reads the hidden state vector (of shape [1, K]) from global memory and determines whether to cache it in GPU shared memory. A related parameter is whether to use a shared cache (LOAD_V_SHARED), which controls whether the input vector is cached in shared memory to reduce the overhead of subsequent repeated accesses to global memory and lower global memory access latency.

[0309] For example, during the input loading phase, an input vector of shape [1, K] is loaded, LOAD_V_SHARED=1 is set, and the vector is cached in the shared memory of the Intel GPU.

[0310] Weight loading is used to read quantized weight data from global memory and control the total data width of aggregate memory access within a thread bundle / subgroup to improve memory access efficiency. Related parameters can include the aggregate memory loading width (a new variable, Coalescing Mem Req Width), a new architecture-related parameter (the second type of parameter). This parameter is linked to the thread bundle / subgroup size and the quantized data packing bit width, and is used to guide the compiler to generate aggregate memory access instructions. By merging independent memory access requests from multiple threads within the same thread bundle / subgroup into a single wide memory access instruction, memory bandwidth utilization is maximized.

[0311] For example, if the subgroup size is configured to be 16 and the quantized data packing width is 16 bits during the weight loading phase, then the data aggregation loading width = 16 × 16 bits = 256 bits; the packed weights are read from global memory in 16-bit basic memory access units (consistent with the quantized data packing width). The subgroup consisting of 16 threads aggregates their respective 16-bit memory access requests into a 256-bit wide memory access instruction, and reads 256 bits of data (containing 16 INT4 weights) at once.

[0312] Weight unpacking is used to split the packaged low-bit-width quantized data into individual quantized weight units. Associated parameters can include the quantized weight data bit width (Weight Width) and the quantized data packing bit width (a new variable, PackedWeight Width). Both of these are new inverse quantization-related parameters (a third type of parameter) used to determine the shift step size, mask bit width, and the number of weights contained in each packing unit during unpacking.

[0313] For example, if the quantization weight data bit width is configured to be 4 bits, the quantization data packing bit width is 16 bits, and the quantization weight data type is INT4 (signed) during the weight unpacking stage, then each packing unit contains 4 4-bit INT4 quantization weights; for the nth INT4 weight (n=1~4), a right shift of (n-1)×4 bits is performed, and a bitwise AND operation is performed with the mask 0xF to extract a single 4-bit INT4 quantization weight from the 16-bit packing data.

[0314] Weight dequantization is used to restore the unpacked low-bit-width integer weights to floating-point values, recovering their original numerical range. Related parameters can include the weight data type, a newly added dequantization-related parameter (the third type of parameter), used to select the signed or unsigned type dequantization logic branch to recover the data's sign information.

[0315] For example, in the weight dequantization stage, the quantization weight data type is configured as INT4 (signed). First, the extracted 4-bit data is temporarily treated as unsigned uint4 type and converted to a 32-bit integer (int32), at which point the value range is [0,15]. Then, a sign recovery operation is performed, that is, XORed with the constant 0x8 and then subtracted from the constant 0x8, adjusting the value range to [-8,7]. Finally, multiplied by the quantization calibration scaling factor, the integer value is mapped to the floating-point level and converted to half-precision floating-point type (FP16) to obtain the floating-point weight that can be directly used in the calculation.

[0316] Vector multiplication (partial sum) performs element-wise multiplication of the input vector elements with the dequantized weight elements, generating an intermediate partial sum. Related parameters may include recursive axis parallelism, non-recursive axis parallelism, loop unrolling dimension, vector computation width, and warp / subgroup size (a new variable, Warp / Subgroup Size). Specifically, recursive axis parallelism, non-recursive axis parallelism, loop unrolling dimension, and vector computation width are general GEMV optimization parameters (first type of parameters) used to control parallel block partitioning, loop unrolling, and vectorized computation; warp / subgroup size is a newly added architecture-related parameter (second type of parameter) used to control the granularity of thread organization, adapting to the hardware characteristics of different GPU architectures.

[0317] For example, in the vector multiplication stage, the thread bundle / subgroup size is configured to be 16, aligned with the recursive axis parallelism of 16. Each subgroup is responsible for the parallel multiplication of 16 recursive axis elements. The non-recursive axis parallelism is 16, splitting the output M axis into 16 parallel tasks, each task being processed by an independent subgroup. The loop unrolling dimension is 32, and the vector computation width is 4. Based on these parameters, the compiler automatically generates vectorized multiplication instructions, improving instruction-level parallelism and computational throughput.

[0318] Partial summation is used to reduce and sum the intermediate sums generated by vector multiplication along the recursion axis K, resulting in the final output vector (of shape [1, M]). Related parameters may include the recursion axis parallelism and the warp / subgroup size (a new variable, Warp / Subgroup Size). This warp / subgroup size is a new architecture-related parameter used to control the parallel granularity of the reduction operation, guiding the compiler to utilize shuffling instructions or shared memory within the warp / subgroup to perform efficient reduction and improve the efficiency of the summation calculation.

[0319] For example, in the partial summation stage, based on a subgroup size of 16, the Intel GPU's Subgroup Shuffle instruction is invoked to perform parallel reduction summation on the 16 intermediate parts within the subgroup; finally, the reduction results of each subgroup are summarized to obtain an output vector of shape [1, M], completing the complete computation process of one dequantization and vector-matrix multiplication fusion operator.

[0320] The above process integrates hardware architecture-related parameters, dequantization-related parameters, and general scheduling parameters into a joint design space, achieving collaborative optimization through hardware awareness and quantization fusion. Secondly, by using parameterized templates, the dequantization logic is directly embedded into the computation pipeline, avoiding the data round-trip overhead caused by the separation of dequantization and computation in traditional solutions. Furthermore, through automated traversal evaluation, the optimal configuration can be efficiently located from a massive number of parameter combinations without manual intervention, significantly improving operator development efficiency and cross-platform performance. Compared to manual optimization, this method does not require recoding for different hardware and quantization formats. Compared to traditional automatic optimization, this method, by introducing architecture awareness and dequantization-related parameters, constructs a finer-grained design space, fully exploring hardware potential and generating high-performance operator code on various GPU architectures.

[0321] S206, Load into parameterized template.

[0322] For example, this step can be performed by the operator code generation module. The design space exploration control module sends the retrieved parameter value combinations to the operator code generation module. The operator code generation module loads the received parameter value combinations into a parameterized template (such as a dequantization and matrix-vector multiplication fusion operator template) to generate an instantiated operator template.

[0323] Specifically, the operator code generation module extracts the values ​​of all parameters from a combination of parameter values ​​and assigns them to the corresponding variables in the parameterized template. For example, if the parameter "recursive axis parallelism" is 8 in a certain combination of parameter values, the corresponding variable TR in the operator template is assigned the value 8, and the number of recursive axis threads is set to TR=8 using the function "TVM.sch.bind(TR,threadIdx.x)" provided by TVM. If the parameter "non-recursive axis parallelism" is 8, the corresponding variable TC in the operator template is assigned the value 8, and the number of non-recursive axis threads is set to TC=8 using "TVM.sch.bind(TC, threadIdx.y)".

[0324] For other parameters, if their properties are similar to those of the recursive axis parallelism, they should be assigned values ​​using a similar method. For example, the value of the loop unrolling dimension is assigned to the variable in the template that controls the number of loop unrolls; the value of whether to use a shared cache is assigned to the flag variable in the template that controls whether the input vector is cached in shared memory; the value of the vector computation width is assigned to the variable in the template that controls the vectorization width; and the value of the thread bundle / subgroup size is assigned to the variable in the template that controls the granularity of thread organization. The subgroup size can be set in the template using appropriate compilation attributes based on the target hardware platform type. After the above assignment operations are completed, all variable placeholders in the template are replaced with specific numerical values.

[0325] For quantized weight data type parameters, the operator code generation module selects the corresponding branch from multiple predefined dequantization logic branches in the parameterization template based on their values. For example... Figure 5A and Figure 5B As shown, the template predefines two types of dequantization logic branches: unsigned and signed. When the quantization weight data type is unsigned integer, the following is selected: Figure 5A The unsigned type dequantization logic branch shown can adjust the value range from [0, 15] to [-8, 7] by subtracting a constant value; when the quantization weight data type is a signed integer, select... Figure 5B The signed type dequantization logic branch shown is used to recover the sign information by "XORing and then subtracting the constant".

[0326] For the quantization weight data bit width and quantization data packing bit width parameters, the operator code generation module determines the mask bit width and shift step size in the unpacking logic based on their values. For example... Figure 6As shown, when the quantization weight data bit width is 4 bits and the quantization data packing bit width is 16 bits, each packing unit contains 4 INT4 quantization data. During unpacking, the nth quantization data is right-shifted by (n-1)×4 bits and bitwise ANDed with the mask 0xF to extract a single 4-bit quantization weight. When the quantization data packing bit width is 32 bits, each packing unit contains 8 INT4 quantization data, the shift step size is 4, and the mask is still 0xF.

[0327] Based on the assigned variables and the selected target dequantization logic branch, the operator code generation module generates a complete, compilable instantiated operator template. In this instantiated operator template, all parameter placeholders have been replaced with specific values, and the dequantization logic branch has been determined, allowing it to be directly fed into the compiler for compilation.

[0328] Through the above operations, the operator code generation module separates adjustable parameters from core computational logic, enabling the same set of parameterized templates to adapt to different parameter configurations. At the same time, by predefining dequantization logic branches in the template and dynamically selecting the target branch based on the third type of parameters, the dequantization operation is deeply integrated with matrix-vector multiplication calculations, laying the foundation for eliminating data round-trip overhead. In addition, the template supports the characteristics of different hardware platforms (such as thread bundle size settings) through conditional compilation or template specialization, so that the generated instantiated operator templates can naturally adapt to the target hardware characteristics.

[0329] S207, Invoke the compiler to generate executable code.

[0330] For example, this step can be performed by the operator code compilation and execution module. The operator code generation module sends the instantiated operator template to the operator code compilation and execution module. This module selects the corresponding compilation backend based on the type of the target hardware platform and calls the compiler corresponding to the backend to compile the instantiated operator template into executable code that can be directly executed on the target hardware platform.

[0331] Specifically, the operator code compilation and execution module compiles the assigned instantiated operator templates into executable code using compilation functions provided by TVM (such as TVM.build()). When the target hardware platform is an Intel GPU architecture, the OpenCL backend is selected to generate OpenCL-style executable code.

[0332] Optionally, the executable code can be adjusted based on the second and third types of parameters during this process. For example, for the Intel GPU architecture, assuming the architecture-related parameter "thread bundle / subgroup size" corresponds to the variable SG_Size with a value of 8, the subgroup size can be set to 8 by calling the OpenCL function intel_reqd_sub_group_size(SG_Size), causing the compiler to generate SIMD instructions corresponding to the subgroup width.

[0333] Based on this, assuming the dequantization-related parameter "quantization data packing width" is 8, the data aggregation loading width can be determined to be the product of the thread bundle / subgroup size and the quantization data packing width, i.e., 8 × 8 = 64. By setting the data type accessed by each thread to 8-bit unsigned integer (unsigned char) in OpenCL, and setting the data addresses accessed by each thread to be adjacent to each other (e.g., setting thread 1 to access data at address 1, thread 2 to access data at address 2, thread 3 to access data at address 3, and so on), the OpenCL compiler can be instructed to automatically generate a data aggregation loading instruction with a width of 64, merging the independent memory access requests of multiple threads within the same subgroup into a single wide memory access instruction, thereby maximizing memory bandwidth utilization.

[0334] When the target hardware platform is an Nvidia GPU, the CUDA backend is selected, and the NVCC compiler is invoked to compile the instantiation operator template into PTX code or a cubin file. Since the thread bundle size of Nvidia GPUs is fixed at 32, there is no need to explicitly set the thread bundle size in the template; the compiler will automatically divide the thread bundle according to the thread block dimension configuration.

[0335] When the target platform is an AMD GPU, the Rocm backend is selected, and the heterogeneous HIP compiler is invoked to compile the instantiated operator template into executable code optimized for AMD GPUs. The wavefront size of AMD GPUs is typically 64, and the compiler can automatically handle wavefront partitioning based on the thread block dimension configuration.

[0336] Through the above operations, the compiler can perform low-level optimizations based on the specific parameter values ​​in the instantiated operator template. This includes automatically generating aggregate memory access instructions based on the thread bundle / subgroup size and the continuity of memory access, performing loop unrolling based on loop unrolling parameters, and reordering instructions to hide memory access latency. This transforms high-level parameter configurations into efficient low-level instruction sequences, providing accurate executable code for subsequent performance evaluation.

[0337] S208 and TVM virtual machines were used to evaluate execution performance.

[0338] For example, this step can be performed by an operator code compilation and execution module. This module deploys the generated executable code to the target hardware platform, schedules and executes it using a TVM virtual machine, and measures its execution performance. The measured execution performance can include performance metrics such as inference latency, compute utilization, data transfer efficiency, and throughput.

[0339] For example, this module can call the TVM virtual machine engine (such as TVM.relax.VirtualMachine) to load the compiled operator into the virtual machine engine, which will then schedule the operator to be executed on the target hardware platform and measure its execution performance.

[0340] Alternatively, in one possible implementation, to quickly filter a large number of parameter combinations, the executable code can be run only once, and its execution performance can be directly measured to improve search efficiency.

[0341] Alternatively, in another possible implementation, in order to eliminate the impact of system noise and environmental fluctuations on the measurement results of a single run, the operation can be repeated multiple times, and the average value of the results of multiple runs can be taken as the final performance index. For example, the average inference latency can be calculated as the performance to be evaluated.

[0342] This module returns the measured execution performance to the design space exploration control module, which then compares and filters parameter value combinations to determine the target parameter combination. Through these operations, the module accurately evaluates operator performance in a real target hardware or virtual machine environment, providing a reliable data foundation for subsequent determination of the optimal parameter combination.

[0343] S209. Update the current optimal parameter combination.

[0344] For example, this step can be performed by the design space exploration control module. This module receives the execution performance returned by the operator code compilation and execution module and compares it with the currently reserved best execution performance.

[0345] Specifically, for the first parameter value combination, it is directly used as the currently retained optimal parameter combination, and its corresponding execution performance is recorded. For each subsequent parameter value combination, its execution performance is compared with that of the currently retained optimal parameter combination. If the execution performance of the current parameter value combination is better, then the retained optimal parameter combination and its corresponding execution performance are updated with the current parameter value combination; if the execution performance of the current parameter value combination is worse than that of the currently retained optimal parameter combination, then the current parameter value combination is discarded, and the original optimal parameter combination remains unchanged.

[0346] After completing the comparison and update, the module returns to step S204 to continue determining whether there are any unevaluated parameter value combinations in the set of valid parameter value combinations. If there are unevaluated combinations, steps S205 to S209 are repeated to load, compile, evaluate, and compare the next parameter value combination; if all parameter value combinations have been evaluated, the loop exits and jumps to the subsequent step S210.

[0347] In this way, the module only needs to save one optimal parameter combination and its execution performance during the traversal process, without storing the evaluation results of all combinations, thus greatly reducing memory overhead; at the same time, it ensures complete coverage of all valid parameter combinations, and the optimal parameter combination that is finally retained is the global optimal solution determined by this automated search.

[0348] III. Target Code Generation S210, Save the target parameter combination to a file.

[0349] For example, this step can be completed collaboratively by the design space exploration control module and the operator code output module. After all parameter value combinations have been evaluated, the design space exploration control module determines the final optimal parameter combination as the target parameter combination and sends it to the operator code output module.

[0350] The operator code output module saves the target parameter combination to disk in a readable format for user reference and analysis. For example, the target parameter combination can be saved as a JSON file, with the filename containing a timestamp or target hardware platform information for easy identification and archiving. Optionally, the saved content should at least include the specific values ​​of each parameter in the optimal parameter combination and the corresponding execution performance (such as inference latency, throughput, and other performance metrics).

[0351] Optionally, this module can also save some or all of the parameter value combinations evaluated during the search process and their corresponding execution performance, forming a complete optimization process record to facilitate performance analysis and debugging by developers. Optionally, the saving format can also be YAML or a custom configuration file; this application does not limit this.

[0352] Through the above operations, the module will solidify the final results obtained from the automated search onto the disk, providing basic data support for users to subsequently call the optimal operator code, analyze the optimization effect, or carry out secondary development.

[0353] S211, Load into TVM scheduling template.

[0354] For example, this step can be performed by the operator code output module. Referring to the aforementioned step S206, the operator code output module loads the target parameter combination determined in step S210 into the parameterized template to obtain a compilable operator instance (i.e., an instantiated operator template generated for the target parameter combination).

[0355] Specifically, this module assigns the values ​​of the first and second types of parameters in the target parameter combination to the corresponding variable placeholders in the parameterized template, and selects the target dequantization logic branch from multiple predefined dequantization logic branches in the template based on the value of the third type of parameter, generating the target operator template. This target operator template can be directly sent to the compiler for compilation to generate the final executable target code.

[0356] In this way, this application transforms the optimal parameter combination obtained through automated search into a directly compileable operator instance, achieving a complete closed loop from parameter search to code generation. Since the target operator template reuses the validated template loading logic from step S206, it ensures that the quality of the final generated code remains consistent with the code evaluated during the search process, avoiding performance deviations caused by implementation differences. Simultaneously, saving the optimal parameter combination as source code as the target operator template provides a flexible interface for subsequent secondary development, customized modifications, or integration into other compilation processes, further enhancing the practicality and scalability of this method.

[0357] S212, Call the TVM compiler.

[0358] For example, this step can be performed by the operator code output module. This module selects the corresponding compilation backend according to the type of the target hardware platform and calls the corresponding compiler to compile the target operator template (i.e., the instantiated operator template for the target parameter combination) generated in step S211 into the final executable target code.

[0359] In this way, this application transforms the optimal parameter combination obtained through automated search into directly deployable target code, achieving a complete closed loop from parameter search to final deployment. This target code integrates the entire process logic of weight dequantization and matrix-vector multiplication, maintaining consistent optimization quality with the code used for performance evaluation during the search process, thus ensuring the effective implementation of the optimization results.

[0360] S213, Output the target code.

[0361] For example, this step can be performed by the operator code output module. This module saves the compiled target code to disk in the backend format corresponding to the target hardware platform, for users to call in actual inference tasks. For example, for Intel GPU platforms, it outputs OpenCL kernel code (.cl file) or pre-compiled binary files; for Nvidia GPU platforms, it outputs CUDA PTX code (.ptx file) or cubin files; for AMD GPU platforms, it outputs Rocm executable files. Optionally, the target operator template source code can also be directly output for secondary development.

[0362] Optionally, meaningful filenames (such as those containing information about the target hardware platform and quantization configuration) can be used when saving the file to facilitate user identification and management.

[0363] Through the detailed process described above, this application achieves fully automated operator optimization from joint design space construction to target code generation. This method integrates general computation scheduling parameters, hardware architecture-related parameters, and quantization data format parameters into the same design space, realizing collaborative optimization through hardware awareness and quantization fusion. Utilizing automated traversal and evaluation methods, it efficiently locates the optimal configuration from massive parameter combinations without manual intervention. The final executable code integrates weight dequantization and matrix-vector multiplication calculations, eliminating data round-trip overhead. While significantly improving operator development efficiency, it achieves high-performance execution on different GPU architectures, effectively solving the technical challenges of high cost and poor portability in existing manual optimization methods, as well as the lack of hardware awareness and insufficient quantization fusion in traditional automated optimization.

[0364] Please refer to Figure 10 , Figure 10 Here is a flowchart illustrating an automatic operator optimization method according to an exemplary embodiment of this application, as detailed below: S301. Obtain the set of parameters to be tuned for the operator to be optimized.

[0365] The parameter set to be tuned is a collection of all adjustable parameters that can directly alter the operator's implementation logic and performance during automatic operator optimization. It's worth noting that the parameter set to be tuned and the joint design space are two closely related but different concepts in automatic operator optimization; they can be understood as two aspects of the same optimization objective. The parameter set to be tuned comprises the specific parameter content that constitutes the joint design space, which is a multi-dimensional search space defined by each parameter in the parameter set and its preset candidate value range. Furthermore, the joint design space is a specific form of the parametric design space, composed of three types of parameters.

[0366] In this application embodiment, the operator to be optimized refers to a deep learning computation unit that requires performance optimization. This application mainly targets the inference scenario of LLM, where the core computational load in the decoding stage is concentrated in the GEMV operator. Therefore, this application uses the GEMV operator as the main example for illustration, but the automatic operator optimization method provided in this application is also applicable to other types of deep learning operators. For example, the operator to be optimized may include, but is not limited to: matrix-vector multiplication operator, matrix multiplication operator, convolution operator, batch normalization operator, layer normalization operator, activation function operator, pooling operator, etc.

[0367] Obtaining the set of parameters to be tuned for the operator to be optimized means that when the optimization process starts, the system collects all parameters that may affect the execution performance of the operator to be optimized by reading the preset configuration related to the operator to be optimized, querying hardware information, loading model quantization parameters, etc., and organizes them into a structured parameter set.

[0368] Specifically, at the beginning of the operator automatic optimization process, the acquisition of the parameter set to be tuned is performed. For example, after the optimization engine starts, it can read general scheduling parameters from the system's built-in deep learning operator scheduling parameter library, read the architectural feature information of the current hardware platform through the target GPU hardware driver interface, load quantization and dequantization related parameters from the model quantization configuration file, and combine them with user-specified custom parameters specified through configuration files or API interfaces, loading the parameter list read from the above multiple channels into the runtime memory. Based on this, combined with the currently specified target GPU hardware platform type, the parameters are adapted, filtered, and integrated to form a complete and unified parameter set to be tuned. This parameter set to be tuned serves as the basic input for subsequent generation of parameter value combinations and construction of the joint design space, providing complete parameter support for the fully automated optimization of the operator process.

[0369] Optionally, the set of parameters to be tuned is kept in the runtime memory as a structure until the entire automatic optimization process of this operator is completed.

[0370] For example, the set of parameters to be tuned obtained this time mainly includes three types of parameters. Among them, the first type of parameters are parameters related to computational parallel scheduling. These parameters are decoupled from the specific GPU hardware architecture and are used to control the granularity of operator computation tasks, thread mapping methods, loop optimization, and data caching strategies. They can be adapted to various general-purpose parallel computing devices and are the basic scheduling parameters in the operator automatic optimization framework.

[0371] The first type of parameters includes at least one or more of the following: recursive axis parallelism parameter, non-recursive axis parallelism parameter, loop unrolling dimension parameter, memory caching enable parameter, vector computation width parameter, thread block size, data layout format, and global memory access alignment flag.

[0372] The second type of parameter is related to the execution unit of the target hardware platform. This type of parameter is strongly bound to the underlying hardware execution unit of the GPU and is used to directly characterize the thread organization method, SIMD execution width and memory access controller, etc. It is the core parameter for realizing hardware-aware optimization.

[0373] The target hardware platform refers to the hardware environment in which the operator finally runs. In the embodiments of this application, the target hardware platform may include a GPU hardware platform. However, the automatic operator optimization method provided in this application is also applicable to other types of parallel computing hardware, such as neural processing units (NPU), tensor processing units (TPU), field programmable gate arrays (FPGA), etc. As long as they have programmable parallel computing capabilities and support the parameterized template method in this application, they can all be used as the target hardware platform.

[0374] Taking GPU hardware platforms as an example, GPUs from different manufacturers and generations exhibit significant differences in thread organization and SIMD execution width. These differences determine the parallel scheduling and memory access strategies of operators. For instance, the thread bundle / subgroup / wavefront size parameter characterizes the basic parallel execution unit size of the target GPU, and its value must be bound to the GPU architecture type. During automated optimization, this parameter is not searched independently. Instead, the optimization framework queries hardware information or a preset architecture mapping table to determine the architecture type of the current GPU and its supported value rules. This limits candidate values ​​to the range actually supported by the current GPU architecture when generating parameter value combinations and filtering hardware constraints. This design ensures that the generated parameter value combinations are truly effective on the hardware platform and integrates hardware characteristics into the optimization search, achieving hardware-aware automated optimization.

[0375] Specifically, the second type of parameters includes at least one or more of the following: thread bundle or subgroup or wavefront size parameter, data aggregation loading width parameter, single-cycle SIMD instruction width, maximum available shared memory capacity threshold, thread scheduling priority configuration, and global memory bandwidth matching coefficient. The data aggregation loading width parameter is determined based on the product of the thread bundle or subgroup size parameter and the quantization data packing bit width parameter, such as: data aggregation loading width = thread bundle / subgroup size × quantization data packing bit width.

[0376] GPU architectures can be categorized in an exemplary manner based on the size of their basic parallel execution units. For example, the first type of GPU architecture (such as Nvidia GPUs) typically has a fixed thread bundle size of 32; the second type of GPU architecture (such as Intel GPUs) has a subgroup size that can be configured to 8, 16, or 32; and the third type of GPU architecture (such as AMD GPUs) typically has a wavefront size of 64.

[0377] It should be understood that the above classifications and values ​​are for illustrative purposes only. In practical applications, the values ​​of the thread bundle / subgroup / wavefront size parameters must strictly match the hardware specifications of the target GPU. For architectures with a fixed thread bundle size, the value is that fixed value; for architectures with a configurable subgroup size, the value is a set of optional values ​​supported by the GPU hardware; for future new architectures, the value can be expanded accordingly with the hardware characteristics.

[0378] The third category of parameters relates to the quantization and dequantization data formats of the operator to be optimized. These quantization and dequantization data formats characterize the storage format of the quantization weight matrix of the matrix-vector multiplication operator, such as whether the quantization weights use signed or unsigned integers, the number of bits occupied by each weight, and the bit width when multiple quantized data are packaged and stored. This category of parameters manages the logic for low-bit-width quantization storage, packaging and unpacking, sign recovery, and numerical restoration of weights. It can deeply integrate the dequantization process with matrix-vector multiplication calculations, effectively eliminating the performance overhead caused by data read / write round trips.

[0379] The third category of parameters includes at least one or more of the following: quantization weight data type parameter, quantization weight data bit width parameter, quantization data packing bit width parameter, quantization group size, zero-point offset storage format, and dequantization intermediate data precision.

[0380] The quantization weight data type parameter can include signed integers and unsigned integers, the quantization weight data bit width parameter can be 2 bits, 3 bits or 4 bits, and the quantization data packing bit width parameter can be 8 bits, 16 bits, 32 bits or 64 bits, etc.

[0381] It should be understood that the above values ​​are merely illustrative examples. In practical applications, they can be extended or adjusted according to the specific quantization scheme and hardware platform requirements. For example, the quantization weight data bit width can be 1 bit or higher, and the quantization data packing bit width can also be selected according to the hardware bit width support. This application does not limit this.

[0382] S302. Based on the candidate values ​​of each parameter in the parameter set to be optimized, generate multiple parameter value combinations for the operator to be optimized.

[0383] In one example, after obtaining the set of parameters to be tuned, all possible candidate values ​​for each parameter are sequentially traversed to generate all possible combinations of parameter values ​​for the operator to be optimized. For example, all possible candidate values ​​for all parameters can be combined using enumeration or Cartesian product to generate a set containing all parameter value combinations. Subsequently, based on the hardware constraints of the target hardware platform, the generated set is filtered, and parameter value combinations that do not meet the hardware constraints are removed, resulting in a set of valid parameter value combinations. This method is simple to implement and suitable for scenarios where the number of parameter combinations is controllable.

[0384] In another example, candidate values ​​for each parameter can be pre-screened based on hardware constraints limiting the range of individual parameter values ​​(e.g., thread bundle / subgroup size can only take specific values ​​supported by the target GPU hardware), eliminating obviously invalid values. Then, the pre-screened candidate values ​​are combined using a Cartesian product to generate a set of parameter value combinations. This approach reduces the number of invalid combinations during the generation phase, lowering the overhead of subsequent screening. It should be noted that for constraints involving multiple parameter relationships (e.g., parallelism determined by the product of recursive axis parallelism and non-recursive axis parallelism), some combinations still require secondary screening based on the linked hardware constraints after generation.

[0385] In another example, a recursive or iterator pattern can be used to generate parameter value combinations step by step in the parameter order. Each time a parameter is added, a feasibility check is performed, retaining only the combinations that satisfy the currently known hardware constraints, ultimately generating a set that completely satisfies all hardware constraints. This approach is suitable for scenarios with a very large number of parameter combinations or where combinations need to be generated in real time, and it effectively controls memory usage.

[0386] It should be understood that regardless of which method is used, a valid set of parameter value combinations can ultimately be obtained.

[0387] The following describes the implementation method for obtaining a set of effective parameter value combinations by initially screening the parameter combinations according to the hardware constraints of the target hardware platform after generating multiple parameter value combinations for the operator to be optimized, eliminating combinations that do not meet the hardware constraints, based on the first example.

[0388] For example, hardware constraints may include at least one or more of the following: the limit on the size of the execution unit of the target hardware platform, the upper limit on the total amount of computing resources of the target hardware platform, the relationship between the parallelism of the operator to be optimized and the thread bundle or subgroup size parameter, and alignment constraints.

[0389] The upper limit of the total computing resources of the target hardware platform can include the parallelism of the operator to be optimized not exceeding the upper limit of the parallelism of the target hardware platform. The upper limit of the parallelism of the target hardware platform refers to the maximum parallelism supported by the hardware device (e.g., for Intel GPUs, this upper limit is 1024).

[0390] Optionally, the parallelism of the operator to be optimized is equal to an integer multiple of the thread bundle or subgroup size parameter, and the parallelism of the operator to be optimized is determined by the product of the recursive axis parallelism parameter and the non-recursive axis parallelism parameter.

[0391] Optionally, hardware constraints may also include register pressure constraints (such as estimating register usage based on loop unrolling factor and number of threads, not exceeding the hardware limit), thread block number constraints (the total number of thread blocks does not exceed the maximum number of thread blocks on the device), memory bandwidth constraints (estimate bandwidth usage based on data aggregation loading width and memory access mode), etc., which are not limited in this application.

[0392] By filtering based on the aforementioned hardware constraints, preprocessing optimization of the search space is achieved. Specifically, invalid parameter value combinations are eliminated in advance, effectively reducing the number of parameter value combinations that need to be evaluated subsequently, significantly improving overall optimization efficiency; it ensures that all parameter value combinations evaluated later are within the limits allowed by hardware constraints, avoiding compilation errors or runtime anomalies caused by invalid parameter value combinations; and it integrates an understanding of hardware characteristics into the automated process, achieving hardware-aware optimization preprocessing, creating favorable input conditions for the automated search of subsequent modules.

[0393] S303. For each parameter value combination in the multiple parameter value combinations, load the parameter value combination into the parameterized template of the operator to be optimized, and generate the executable code of the operator to be optimized on the target hardware platform.

[0394] In one example, multiple parameter value combinations can be processed in parallel using batch processing or pipelined methods. For instance, for multiple parameter value combinations, multiple threads or processes can be started simultaneously to perform parameter loading, template instantiation, and compilation operations respectively, thereby shortening the overall search time. This approach is suitable for scenarios with ample computing resources and a large number of parameter combinations, and can further improve operator optimization efficiency.

[0395] Alternatively, in another example, a multi-stage parallel processing pipeline based on a task dependency graph can be employed. The parameter value combination evaluation process is modeled as a directed acyclic graph, where nodes represent subtasks such as parameter loading, template instantiation, compilation, execution, and performance evaluation, and edges represent data dependencies. The asynchronous task scheduling mechanism of the TVM compiler is used to decouple and execute these subtasks in parallel. For example, evaluation can begin immediately after the compilation of a parameter value combination is complete, without waiting for the entire batch to finish; compilation itself can also be broken down into front-end, middle-end, and back-end subtasks for parallel processing. Furthermore, a priority queue can be introduced to prioritize exploring promising parameter combinations, significantly improving search efficiency.

[0396] S304. Evaluate the execution performance of each executable code and determine the target parameter combination from multiple parameter value combinations based on the execution performance.

[0397] For example, the target parameter combination refers to one or more parameter value combinations determined from all evaluated parameter value combinations according to preset selection rules (such as optimal performance, top N performance, or meeting specific constraints, etc.), which are used to generate target code subsequently.

[0398] In one possible implementation, the performance of each executable code is evaluated sequentially, its execution performance is recorded, and the optimal combination is determined as the target parameter combination by comparing each pair of code and retaining the best one.

[0399] Alternatively, in another possible implementation, incremental performance evaluation and early pruning strategies can be employed. For example, multiple checkpoints can be set (e.g., 10%, 30%, 50% of execution progress). If the performance of the current combination is significantly worse than the current optimal combination at an early checkpoint (e.g., latency exceeds the optimal value by 20%), the evaluation is terminated early to avoid unnecessary computation. Simultaneously, a lightweight regression model (e.g., random forest or Gaussian process) is trained using the evaluated data to predict the performance of unevaluated combinations. Only combinations predicted to have potential are evaluated in actual hardware, and the model is dynamically updated. This approach can significantly reduce the number of evaluations and improve operator optimization efficiency.

[0400] S305. Based on the target parameter combination, generate the target code for the operator to be optimized to be executed on the target hardware platform.

[0401] In one possible implementation, the target parameters are combined and loaded into a parameterized template, parameter assignment and dequantization logic branch selection are performed to generate a target operator template, and then the compilation backend corresponding to the target hardware platform is called to compile and generate the final target code.

[0402] Alternatively, in another possible implementation, when the target parameter combination contains multiple candidates (such as the top N optimal combinations), multiple target code versions are generated in parallel, and the most suitable version is dynamically selected for execution at runtime based on the actual hardware environment (such as temperature, power consumption, load, etc.) to improve the flexibility and robustness of deployment.

[0403] Alternatively, an incremental code generation and cache reuse strategy can be adopted, storing the target code generated each time and its corresponding target parameter combination hash value in a local cache. When the same or similar parameter combination is encountered again (such as different models sharing the same quantization configuration), the existing target code is directly read from the cache and reused, avoiding repeated compilation and significantly reducing the overhead of repeated optimization.

[0404] Optionally, target code with performance annotations can also be generated. For example, metadata, including parameter combination information, expected performance metrics, and hardware platform identifiers, can be embedded in the generated target code. The runtime system can then perform adaptive scheduling based on these annotations, such as allocating different operators to the optimal computing units or dynamically adjusting execution strategies according to the current load. This approach significantly improves the efficiency of the entire process from optimization to deployment while ensuring operator performance.

[0405] The aforementioned automated operator optimization process essentially automates the generation of operator kernel function code (i.e., parallel computation function code that executes directly on the GPU). Specifically, this application takes a dequantization and vector-matrix multiplication fusion operator template and a joint design space as input, and through a complete automated operator optimization process, generates executable code that can run directly on the target hardware platform. For example, users only need to provide a high-level description of the Python-style GEMV operator to obtain high-performance optimized operator code that can be directly executed on the target hardware platform through the automated optimization method of this application. This decouples the operator optimization process from the underlying hardware details, allowing users to focus only on the computational logic of the operator itself without needing to deeply understand the GPU architecture and underlying programming models, thus obtaining high-performance implementations for specific hardware platforms and quantization configurations, significantly reducing the barrier to operator development.

[0406] Optionally, in practical applications, users can specify the computation shape of the operator according to business needs (such as flexibly adjusting the size of the recursive axis K or the non-recursive axis M), and of course, they can also specify the target hardware platform type (including Intel, Nvidia and AMD GPU platforms, etc.), so as to customize the operator according to their own needs and achieve flexible deployment and cross-platform high-performance execution.

[0407] The aforementioned automatic operator optimization method effectively solves the technical problems of manual optimization relying on expert experience, high development costs, poor portability, and the limitations of traditional automatic optimization due to its coarse design space, inability to perceive hardware architecture differences, and lack of integration with the dequantization process. This application incorporates computational parallel scheduling parameters, hardware execution unit parameters, and dequantization-related parameters into a joint design space, achieving collaborative optimization through hardware awareness and quantization fusion. Through automated traversal and evaluation, it efficiently locates the optimal configuration for the target hardware platform from a massive number of parameter combinations without manual intervention. By embedding the dequantization logic directly into the computational pipeline using parameterized templates, it avoids data round-trip overhead. The final generated target code integrates weight dequantization and matrix-vector multiplication calculations, significantly improving operator development efficiency while achieving high-performance execution on different GPU architectures. It balances cross-platform versatility and extreme performance, overcoming the technical contradiction of the trade-off between performance and portability.

[0408] Optionally, in one possible implementation, the step S303 above, which loads each parameter value combination from multiple parameter value combinations into the parameterized template of the operator to be optimized, and generates executable code for the operator to be optimized on the target hardware platform, may further include: S3031. Assign each parameter value in the parameter value combination to the corresponding variable of the parameterized template one by one to obtain the instantiated operator template.

[0409] One possible implementation involves using template specialization or code generation techniques to directly compile the parameter value combinations into static, branchless instantiated code. For example, for the quantization weight data type in the third type of parameter, C++ template specialization or macro expansion can be used to directly generate dedicated code for the corresponding signed or unsigned dequantization path at compile time, thereby eliminating runtime branching and improving execution efficiency. This approach is suitable for scenarios with extremely high performance requirements, effectively reducing runtime overhead and improving operator execution efficiency.

[0410] Alternatively, in another possible implementation, the instantiated template can be generated through explicit assignment and conditional branch selection. For example, the first and second types of parameters in the parameter value combination are loaded into the corresponding variables of the parameterized template; based on the third type of parameter in the parameter value combination, the target dequantization logic branch is selected from the predefined dequantization logic branches of the parameterized template; wherein, the predefined dequantization logic branches include signed type dequantization logic branches and unsigned type dequantization logic branches; based on the loaded first and second types of parameters and the target dequantization logic branch, the instantiated operator template is obtained.

[0411] Specifically, taking the GEMV operator as an example, TR is assigned to the variable controlling the K-axis cyclic partitioning and bound to threadIdx.x; TC is assigned to the variable controlling the M-axis parallel partitioning and bound to threadIdx.y; the loop unrolling dimension, whether to use a shared cache, vector calculation width, thread bundle / subgroup size, etc., are assigned to the corresponding variables, and the subgroup size is set according to the target hardware platform through compilation attributes. The target dequantization logic branch is selected based on the quantization weight data type; for unsigned types, the value range is adjusted by subtracting a constant, and for signed types, the sign recovery logic is used. Simultaneously, the unpacking mask and shift step size are determined based on the quantization weight data bit width and the quantization data packing bit width. The assigned variables and the selected dequantization logic branch are integrated to generate a complete, compilable instantiated operator template.

[0412] S3032. Based on the compilation backend that matches the target hardware platform, compile the instantiated operator template into executable code adapted to the target hardware platform.

[0413] For example, the type of the current target hardware platform can be identified in advance by querying the underlying hardware driver information, reading the platform identifier in the system environment variables or configuration files, or by the user explicitly specifying it when the optimization process starts. This will determine whether the target hardware platform belongs to the first type of GPU architecture (such as NVIDIA GPU), the second type of GPU architecture (such as Intel GPU), the third type of GPU architecture (such as AMD GPU), or other possible GPU architectures.

[0414] The appropriate compilation backend is selected based on the target hardware platform type. Specifically, for Nvidia GPUs, the CUDA backend is selected, and the Nvidia NVCC compiler is invoked for compilation; for Intel GPUs, the OpenCL backend is selected, and the OpenCL compiler is invoked for compilation; for AMD GPUs, the Rocm backend is selected, and the HIP compiler is invoked for compilation. The compiler then performs low-level optimizations based on the specific parameter values ​​in the instantiated operator template, including aggregate memory access instruction generation, loop unrolling, and instruction reordering, ultimately generating executable code that can be directly executed by the target hardware platform.

[0415] Using the methods described above, this approach can automatically generate compatible executable code based on the selected hardware platform. For example, it generates CUDA code (.ptx or cubin files) on Nvidia GPU platforms, OpenCL code (.cl files or pre-compiled binaries) on Intel GPU platforms, and Rocm code (or HIP executable code) on AMD GPU platforms. This achieves full automation of the operator process from parameter configuration to hardware adaptation, significantly improving cross-platform deployment flexibility and development efficiency while ensuring execution efficiency.

[0416] Optionally, in one possible implementation, the process of evaluating the execution performance of each executable code and determining the target parameter combination from multiple parameter value combinations in S304 above may further include: S3041. The performance of the executable code corresponding to each combination of parameter values ​​is evaluated in turn to obtain the corresponding execution performance.

[0417] For example, the generated executable code is deployed to the target hardware platform and executed via TVM virtual machine scheduling to measure its execution performance. Execution performance includes at least one or more of the following: inference latency, computing power utilization, and data transfer efficiency of the operator to be optimized.

[0418] Alternatively, hardware performance counters can be used to collect more granular performance data, such as cache hit rate and instruction issuance efficiency, to provide richer decision-making basis for subsequent parameter combination selection.

[0419] S3042. Compare the execution performance of the current parameter value combination with the execution performance of the retained parameter value combination. If the execution performance of the current parameter value combination is better, then update the retained parameter value combination with the current parameter value combination.

[0420] For example, after receiving the execution performance data, the parameter value combinations can be updated using a "pairwise comparison, retain the best" approach. For instance, the first parameter combination is directly retained; subsequent combinations are compared with the currently retained best combination, and if the current performance is better, the retained combination and its execution performance are updated; otherwise, the current combination is discarded.

[0421] Optionally, to enhance the robustness of the search, the top N optimal combinations (N≥2) can be retained, and some unevaluated combinations can be randomly sampled in subsequent evaluations to balance the depth and breadth of the search and avoid getting trapped in local optima too early.

[0422] S3043. After traversing all parameter value combinations, the final retained parameter value combination is determined as the target parameter combination.

[0423] Once all parameter value combinations have been evaluated, the final retained parameter combination is determined as the target parameter combination. In the implementation of retaining the top N optimal combinations, the final target parameter combination can be further selected from the N candidate combinations based on auxiliary indicators (such as power consumption, code size, etc.) or user preferences. Optionally, the target parameter combination can be a single optimal combination, or multiple candidate combinations for the upper-layer scheduler to dynamically select.

[0424] In the above implementation, by sequentially evaluating the execution performance of each parameter value combination and employing a pairwise comparison and optimal dynamic update strategy, this application only needs to store one optimal combination and its performance index in memory, eliminating the need to store the evaluation results of all combinations, thus significantly reducing memory overhead. Simultaneously, this method can maintain the current optimal solution online in real time, ensuring that the globally optimal configuration is obtained after traversal, avoiding performance loss due to improper parameter selection, and providing a reliable data foundation for generating high-performance operator code. Furthermore, this evaluation and selection method requires no manual intervention and is fully automated, significantly improving operator optimization efficiency.

[0425] Optionally, in one possible implementation, the step S305 above, which generates the target code for the operator to be optimized to execute on the target hardware platform based on the combination of target parameters, may further include: S3051. Load the target parameter combination into the parameterized template of the operator to be optimized to obtain the target operator template.

[0426] Referring to step S3031 above, the first and second types of parameters in the target parameter combination are assigned to the template variable, and the dequantization logic branch is selected according to the third type of parameter to generate the target operator template (i.e., the instantiated operator template for the target parameter combination). This will not be elaborated further here.

[0427] S3052. The target operator template is compiled based on a compilation backend that matches the target hardware platform to generate target code; the target code integrates the weight dequantization operation and the calculation logic of the operator to be optimized.

[0428] For example, the corresponding compilation backend is selected according to the target hardware platform type, and the corresponding compiler is called to compile the target operator template into the final executable target code, which is then saved to disk in the backend format corresponding to the target hardware platform. This target code integrates the entire process logic of weight dequantization and matrix-vector multiplication calculation. Each thread directly reads the quantized data according to the bit width when accessing memory, and immediately participates in the multiplication and accumulation calculation after unpacking and dequantization, avoiding the data round-trip overhead of first dequantizing, then storing, and then reading and calculating in the traditional solution.

[0429] Through the above steps, this application realizes a fully automated operator optimization process from joint design space construction to target code generation. This method incorporates hardware architecture parameters and inverse quantization parameters into joint optimization, improving development efficiency while achieving high-performance execution across platforms. It effectively solves the technical problems of high cost and poor portability of manual optimization in related technologies, as well as the lack of hardware awareness and insufficient quantization fusion in traditional automatic optimization.

[0430] Optionally, to verify the effectiveness of the automatic operator optimization method provided in this application, tests were conducted on actual hardware platforms. The test model adopted the feed-forward network (FFN) layer of the Qwen3-4B large language model. This layer contains three matrix-vector multiplication operators, namely the Up operator, the Gate operator, and the Down operator. The test platforms were Intel Arrow Lake and Panther Lake GPU platforms.

[0431] The comparison objects include the default TVM operator without optimization using this method (denoted as TVM default), the operator optimized using this method with a weight quantization precision of INT4 (denoted as TVM-Int4), the operator optimized using this method with a weight quantization precision of INT3 (denoted as TVM-Int3), and the manually optimized INT4 operator in Intel's official inference engine Openvino (denoted as Openvino-Int4). Execution performance is the average inference latency of a single operator; lower latency indicates better performance.

[0432] The test results are shown in Table 2. The data shows that, using the same INT4 weighted quantization precision, the operator optimized using this method (TVM-Int4) achieves an average speedup of approximately 1.5 times compared to the default TVM operator (TVM default); and an average speedup of approximately 1.1 times compared to the Intel-optimized Openvino operator (Openvino-Int4). Furthermore, this method also supports INT3 weighted quantization precision (TVM-Int3), further improving overall performance. Some large-size operators may experience slight fluctuations due to computational characteristics, while Openvino currently does not support INT3 quantization.

[0433] Table 2

[0434] The test results above demonstrate that the automatic operator optimization method provided in this application can significantly improve the execution efficiency of matrix-vector multiplication operators. Compared to traditional automatic optimization methods, it achieves superior performance by constructing a fine-grained joint design space through hardware architecture-aware parameters and dequantization-related parameters. Compared to manual optimization methods, it not only maintains performance advantages but also supports new quantization formats such as INT3, providing a high-performance implementation path for model compression technology. This verifies the effectiveness of this method in improving development efficiency and cross-platform high-performance execution.

[0435] Please refer to Figure 11 , Figure 11This is a schematic diagram illustrating the performance distribution of parameter combinations during automatic optimization, as shown in an exemplary embodiment of this application. The diagram illustrates the delay distribution of all parameter value combinations and the iterative trend of optimal performance during the automatic optimization process.

[0436] like Figure 11 As shown, the horizontal axis represents the sampling period (i.e., the evaluation order / number of parameter value combinations), and the vertical axis represents the operator inference delay (the unit can be milliseconds or normalized values). The scatter plot (SampleTime) represents the actual inference delay of a single parameter value combination, and the curve (BestSoFar) represents the optimal delay value found up to the current sampling period.

[0437] As shown in the figure, the inference latency varies significantly with different parameter combinations, covering a wide range from approximately 0.42ms to 1.15ms. Some combinations exhibit poor performance and high latency, while others demonstrate excellent performance and low latency. This indicates that in the joint design space, different parameter configurations have a significant impact on operator performance, and it is difficult to directly locate the optimal combination relying solely on engineer experience or manual parameter selection.

[0438] This application employs an automated traversal evaluation method to systematically explore the entire joint design space, evaluating the performance of each parameter combination and recording the current optimal solution in real time. As the sampling period increases, the curve gradually decreases and converges after approximately 200 sampling periods, indicating that the automatic optimization process can continuously discover better parameter combinations and ultimately lock in the global optimum.

[0439] This figure fully demonstrates the effectiveness of the method in this application. By constructing a fine-grained joint design space and automating the evaluation, the optimal configuration can be efficiently located in a large number of parameter combinations without manual intervention. This avoids performance loss caused by improper parameter selection and ensures that the final operator code achieves near-theoretical performance on the target hardware platform.

[0440] It is worth noting that the electronic devices used in the technical solutions of this application are typically intelligent devices with one or more processors and supporting parallel computing capabilities. Specifically, the electronic devices in this application may include mobile phones, tablets, laptops, servers, in-vehicle devices, wearable devices, augmented reality (AR) / virtual reality (VR) devices, IoT terminals, etc. This application does not impose any limitations on the specific type and form of the electronic device; any device containing at least one GPU or general-purpose parallel computing unit is within the scope of this application.

[0441] Please refer to Figure 12 , Figure 12This is a schematic diagram of the structure of an electronic device illustrated in an exemplary embodiment of this application. The electronic device 100 may include a processor 110, an external memory interface 120, an internal memory 121, a universal serial bus (USB) interface 130, a charging management module 140, a power management module 141, a battery 142, an antenna 1, an antenna 2, a mobile communication module 150, a wireless communication module 160, an audio module 170, a speaker 170A, a receiver 170B, a microphone 170C, a headphone jack 170D, a sensor module 180, buttons 190, a motor 191, an indicator 192, a camera 193, a display screen 194, and a subscriber identification module (SIM) card interface 195, etc. The sensor module 180 may include a pressure sensor 180A, a gyroscope sensor 180B, a barometric pressure sensor 180C, a magnetic sensor 180D, an accelerometer sensor 180E, a distance sensor 180F, a proximity sensor 180G, a fingerprint sensor 180H, a temperature sensor 180J, a touch sensor 180K, an ambient light sensor 180L, a bone conduction sensor 180M, etc.

[0442] It should be noted that, Figure 12 The structure shown does not constitute a specific limitation on the electronic device 100. In other embodiments of this application, the electronic device 100 may include... Figure 12 The components shown may have more or fewer components, or some components may be combined, some components may be separated, or there may be different arrangements of components. The components shown may be implemented in hardware, software, or a combination of software and hardware.

[0443] In this embodiment, the processor 110 may include one or more processing units, such as a central processing unit (CPU), a GPU, a digital signal processor, a neural network processor, etc. The GPU serves as the target hardware platform for the method of this application, used to execute the generated target code and implement high-performance matrix-vector multiplication and other computational tasks. The processor 110 may also include other processing units, such as an application processor, a modem processor, an image signal processor, a controller, a video codec, a baseband processor, etc. Different processing units may be independent devices or integrated devices.

[0444] The processor 110 may also include a memory for storing instructions and data. In some embodiments, the memory in the processor 110 is a cache memory, which can store instructions or data that the processor 110 has just used or that are used repeatedly, thereby improving system efficiency.

[0445] The processor 110 can run the software code of the operator automatic optimization method provided in the embodiments of this application. For example, the processor 110 can execute the program instructions corresponding to the joint design space construction module, the design space exploration control module, the operator code generation module, the operator code compilation and execution module, and the operator code output module, thereby completing the automatic optimization process of the operator to be optimized and generating target code that can be executed efficiently on the GPU.

[0446] The electronic device 100 can implement display functions through a GPU, a display screen 194, and an application processor. The GPU is a microprocessor for image processing, connected to the display screen 194 and the application processor. The GPU can also be used to perform general-purpose parallel computing tasks; the high-performance implementation of operators such as matrix-vector multiplication in this embodiment runs on the GPU.

[0447] The wireless communication function of electronic device 100 can be implemented through antenna 1, antenna 2, mobile communication module 150, wireless communication module 160, etc., which will not be described in detail here. Sensor module 180 can provide environmental perception capability for electronic device 100.

[0448] In summary, Figure 12 The hardware architecture of the electronic device 100 shown provides a complete physical support for the automatic operator optimization method provided in this application embodiment. The computing power of the processor (such as a GPU), the data storage capacity of the memory, and the software scheduling of the operating system jointly support the specific implementation of the steps such as joint design space construction, parameter search, code generation and compilation in the method of this application, thereby achieving the effect of automatically generating high-performance operator code on different GPU architectures.

[0449] The software structure of the electronic device involved in the embodiments of this application will be briefly described below with reference to the accompanying drawings. Please refer to... Figure 13 , Figure 13 This is a software structure block diagram of an electronic device illustrated in an exemplary embodiment of this application.

[0450] The software system can adopt a layered architecture, event-driven architecture, microkernel architecture, microservice architecture or cloud architecture. This application embodiment takes a layered architecture as an example to exemplarily describe the software system of electronic device 100.

[0451] like Figure 13As shown, the layered architecture divides the software into several layers, each with a clear role and division of labor. Layers communicate with each other through software interfaces. In some embodiments, taking an operating system (such as Linux, Android, etc.) running on an electronic device as an example, the system can be divided into five layers, from top to bottom: application layer 210, application framework layer 220, system layer 230, hardware abstraction layer (HAL) 240, and kernel layer 250.

[0452] The application layer 210 may include a series of application packages, such as AI inference applications, model deployment applications / tools, compiler front-ends, etc. Optionally, users can trigger automatic operator optimization processes through this layer, or specify the parameter configuration of the operator to be optimized, the target hardware platform type, etc.

[0453] Application framework layer 220 provides an application programming interface (API) and programming framework for the applications of application layer 210. Application framework layer 220 includes some predefined functions.

[0454] In this embodiment, the application framework layer 220 may include a window manager, a content provider, a notification manager, an operator optimization configuration interface, and a parameter list management module. The operator optimization configuration interface receives user-specified operator types, target hardware platform information, quantization parameter configurations, etc., and passes this configuration information to lower-level modules. The parameter list management module manages the list of parameters to be tuned, including reading, storing, and dynamically adjusting the names, meanings, and candidate value ranges of the parameters, and provides the parameter list to lower-level modules (such as the joint design space composition module). It should be understood that the application framework layer 220 may also include... Figure 13 The resource manager, view system, etc., are not shown in the image.

[0455] System layer 230 may include the Android Runtime and system libraries. The Android Runtime includes core libraries and a virtual machine. The Android Runtime is responsible for the scheduling and management of the Android system, while the core libraries may consist of two parts: one part contains the functionalities that the Java language needs to call, and the other part contains the core Android libraries.

[0456] In this embodiment, system layer 230 is the core implementation layer of the operator automatic optimization method. System layer 230 may include a joint design space composition module, a design space exploration control module, an operator code generation module, an operator code compilation and execution module, and an operator code output module, etc.

[0457] The joint design space construction module is responsible for parameter definition and classification, candidate value definition, parameter combination generation, and hardware constraint screening. This module can reuse the scheduling and tensor expression functions of compilation optimization frameworks (such as TVM). The design space exploration control module, as the control center of the optimization process, is responsible for traversing parameter combinations, calling the code generation and compilation module, collecting execution performance data, and determining the target parameter combination. The operator code generation module generates instantiated operator templates based on the parameter value combinations. This module contains parameterized templates (variable parameter placeholders and predefined dequantization logic branches), and the template instantiation process is completed in memory. The operator code compilation and execution module is responsible for compiling the instantiated operator templates into code executable by the target hardware platform and running it on the target hardware platform to measure execution performance. The operator code output module generates the final target code based on the target parameter combination and saves it to disk in the backend format corresponding to the target hardware platform.

[0458] The above modules mainly run at system layer 230 and are closely integrated with the runtime environment of compilation optimization frameworks (such as TVM).

[0459] It should be understood that system layer 230 may also include Figure 13 The surface manager, media libraries, 3D image processing libraries (e.g., OpenGL ES), 2D graphics engines (e.g., SGL), system graphics services, etc., are not shown in the figure.

[0460] The hardware abstraction layer 240 encapsulates the underlying hardware drivers, providing a unified interface to shield hardware differences. In this embodiment, the hardware abstraction layer 240 may include a compilation backend adaptation interface (such as runtime libraries for OpenCL, CUDA, and Rocm) and a performance data acquisition interface, used to obtain hardware information and execution performance data of the target GPU.

[0461] Kernel layer 250 includes display drivers, GPU drivers, memory management drivers, etc. Among them, the GPU driver is used to directly manage and operate GPU hardware, execute compiled operator kernel functions, and return execution results and performance data.

[0462] Understandable, Figure 13 The software structure diagram shown does not constitute a specific limitation on the electronic device. In other embodiments of this application, the electronic device may include more or fewer components than shown, or combine some components, or split some components, or have different component arrangements.

[0463] The execution subject of the operator automatic optimization method provided in this application embodiment can be the above-mentioned electronic device, or a functional module and / or functional entity in the electronic device that can implement the method. The solution in this application can be implemented by hardware and / or software. The specific implementation can be determined according to actual usage requirements. This application embodiment does not limit the implementation.

[0464] Please refer to Figure 14 , Figure 14 This is a schematic diagram of an operator automatic optimization device illustrated in an exemplary embodiment of this application.

[0465] It should be understood that the operator automatic optimization device 300 can execute the operator automatic optimization method provided in this application; the operator automatic optimization device 300 may include: an acquisition unit 310, a first generation unit 320, a second generation unit 330, an evaluation unit 340, and a third generation unit 350.

[0466] The acquisition unit 310 is used to acquire a set of parameters to be tuned for the operator to be optimized; wherein, the set of parameters to be tuned includes a first type of parameters related to computational parallel scheduling, a second type of parameters related to the execution unit of the target hardware platform, and a third type of parameters related to the quantization and dequantization data formats of the operator to be optimized. The first generation unit 320 is used to generate multiple parameter value combinations of the operator to be optimized based on the candidate values ​​of each parameter in the parameter set to be tuned. The second generation unit 330 is used to load the parameter value combination into the parameterized template of the operator to be optimized for each parameter value combination in the multiple parameter value combinations, and generate the executable code of the operator to be optimized on the target hardware platform. Evaluation unit 340 is used to evaluate the execution performance of each executable code and determine the target parameter combination from multiple parameter value combinations based on the execution performance; The third generation unit 350 is used to generate target code for the operator to be optimized to be executed on the target hardware platform based on the target parameter combination.

[0467] It should be noted that the aforementioned operator automatic optimization device 300 is embodied in the form of a functional unit. The term "unit" here can be implemented in software and / or hardware, without specific limitations.

[0468] For example, a "unit" can be a software program, a hardware circuit, or a combination of both that implements the above functions. The hardware circuit may include an application-specific integrated circuit (ASIC), electronic circuitry, a processor (e.g., a shared processor, a proprietary processor, or a group processor) and memory for executing one or more software or firmware programs, integrated logic circuitry, and / or other suitable components that support the described functions.

[0469] Therefore, the units of the various examples described in the embodiments of this application can be implemented in electronic hardware, or a combination of computer software and electronic hardware. Whether these functions are implemented in hardware or software depends on the specific application and design constraints of the technical solution. Those skilled in the art can use different methods to implement the described functions for each specific application, but such implementation should not be considered beyond the scope of this application.

[0470] This application also provides a computer-readable storage medium storing computer instructions; when the computer-readable storage medium is used on an operator automatic optimization device, the operator automatic optimization device executes the operator automatic optimization method of any of the above embodiments.

[0471] The computer instructions may be stored in a computer-readable storage medium or transmitted from one computer-readable storage medium to another. For example, the computer instructions may be transmitted from one website, computer, server, or data center to another website, computer, server, or data center via wired (e.g., coaxial cable, fiber optic, digital subscriber line (DSL)) or wireless (e.g., infrared, wireless, microwave, etc.) means.

[0472] The computer-readable storage medium can be any available medium that a computer can access, or it can include one or more data storage devices such as servers and data centers that can be integrated with such media. The available medium can be magnetic media (e.g., floppy disks, hard disks, magnetic tapes), optical media, or semiconductor media (e.g., solid-state disks (SSDs)).

[0473] This application also provides a computer program product containing computer instructions, which, when run on an operator automatic optimization device, enables the operator automatic optimization device to execute the operator automatic optimization method of any of the above embodiments.

[0474] This application also provides a chip. Please refer to [link / reference]. Figure 15 , Figure 15This is a schematic diagram of the structure of a chip shown in an exemplary embodiment of this application. Figure 15 The chip shown can be a general-purpose processor or a special-purpose processor. The chip includes a processor 410. The processor 410 is used to execute the operator automatic optimization method of any of the above embodiments.

[0475] Optionally, the chip also includes a transceiver 420, which is used to receive control from the processor to support the operator automatic optimization device in executing the aforementioned technical solution.

[0476] Optional, Figure 15 The chip shown may also include: storage medium 430.

[0477] It should be noted that, Figure 15 The chip shown can be implemented using one or more field programmable gate arrays (FPGAs), programmable logic devices (PLDs), controllers, state machines, gate logic, discrete hardware components, any other suitable circuits, or any combination of circuits capable of performing the various functions described throughout this application.

[0478] In this embodiment, the electronic device, computer-readable storage medium, computer program product or chip are all used to execute the corresponding methods provided above. Therefore, the beneficial effects that can be achieved can be referred to the beneficial effects of the corresponding methods provided above, and will not be repeated here.

[0479] Through the above description of the embodiments, those skilled in the art will understand that, for the sake of convenience and brevity, only the division of the above functional modules is used as an example. In actual applications, the above functions can be assigned to different functional modules as needed, that is, the internal structure of the device can be divided into different functional modules to complete all or part of the functions described above.

[0480] In the several embodiments provided in this application, it should be understood that the disclosed apparatus and methods can be implemented in other ways. For example, the apparatus embodiments described above are merely illustrative; for instance, the division of modules or units is only a logical functional division, and in actual implementation, there may be other division methods. For example, multiple units or components may be combined or integrated into another apparatus, or some features may be ignored or not executed. Furthermore, the mutual coupling or direct coupling or communication connection shown or discussed may be through some interfaces; the indirect coupling or communication connection between apparatuses or units may be electrical, mechanical, or other forms.

[0481] The units described as separate components may or may not be physically separate. A component shown as a unit can be one or more physical units; that is, it can be located in one place or distributed in multiple different locations. Some or all of the units can be selected to achieve the purpose of this embodiment according to actual needs.

[0482] Furthermore, the functional units in the various embodiments of this application can be integrated into one processing unit, or each unit can exist physically separately, or two or more units can be integrated into one unit. The integrated unit can be implemented in hardware or as a software functional unit.

[0483] If the integrated unit is implemented as a software functional unit and sold or used as an independent product, it can be stored in a readable storage medium. Based on this understanding, the technical solutions of the embodiments of this application, in essence, or the parts that contribute to the prior art, or all or part of the technical solutions, can be embodied in the form of a software product. This software product is stored in a storage medium and includes several instructions to cause a device (which may be a microcontroller, chip, etc.) or processor to execute all or part of the steps of the methods of the various embodiments of this application. The aforementioned storage medium includes various media capable of storing program code, such as USB flash drives, portable hard drives, read-only memory (ROM), random access memory (RAM), magnetic disks, or optical disks.

[0484] The above description is merely a specific embodiment of this application, but the scope of protection of this application is not limited thereto. Any variations or substitutions that can be easily conceived by those skilled in the art within the scope of the technology disclosed in this application should be included within the scope of protection of this application. Therefore, the scope of protection of this application should be within the scope of the claims.

Claims

1. An automatic operator optimization method, characterized in that, include: Obtain a set of parameters to be tuned for the operator to be optimized; wherein, the set of parameters to be tuned includes a first type of parameters related to computational parallel scheduling, a second type of parameters related to the execution unit of the target hardware platform, and a third type of parameters related to the quantization and dequantization data formats of the operator to be optimized; Based on the candidate values ​​of each parameter in the set of parameters to be tuned, multiple combinations of parameter values ​​for the operator to be optimized are generated. For each parameter value combination among the multiple parameter value combinations, the parameter value combination is loaded into the parameterization template of the operator to be optimized, and the executable code of the operator to be optimized is generated on the target hardware platform; Evaluate the execution performance of each executable code, and determine a target parameter combination from the multiple parameter value combinations based on the execution performance; Based on the target parameter combination, target code for the operator to be optimized to be executed on the target hardware platform is generated.

2. The method according to claim 1, characterized in that, The third type of parameter includes at least one or more of the following: quantization weight data type parameter, quantization weight data bit width parameter, and quantization data packing bit width parameter.

3. The method according to claim 2, characterized in that, The second type of parameter includes at least one or more of the following: thread bundle or subgroup size parameter and data aggregation loading width parameter; wherein the data aggregation loading width parameter is determined based on the product of the thread bundle or subgroup size parameter and the quantized data packing bit width parameter.

4. The method according to claim 3, characterized in that, The target hardware platform includes a GPU hardware platform, and the value of the thread bundle or subgroup size parameter is determined according to the architecture type of the GPU; wherein, the value of the thread bundle or subgroup size parameter is 32 under the first type of GPU architecture, the value of the thread bundle or subgroup size parameter is 8, 16 or 32 under the second type of GPU architecture, and the value of the thread bundle or subgroup size parameter is 64 under the third type of GPU architecture.

5. The method according to claim 1, characterized in that, After generating multiple parameter value combinations for the operator to be optimized based on the candidate values ​​of each parameter in the parameter set to be tuned, the method further includes: Based on the hardware constraints of the target hardware platform, parameter value combinations that do not meet the hardware constraints are eliminated from the multiple parameter value combinations.

6. The method according to claim 5, characterized in that, The hardware constraints include at least one or more of the following: the limit on the size of the execution unit of the target hardware platform, the upper limit on the total computing resources of the target hardware platform, and the relationship between the parallelism of the operator to be optimized and the thread bundle or subgroup size parameter.

7. The method according to claim 1, characterized in that, The step of loading the parameter value combination into the parameterized template of the operator to be optimized, and generating executable code of the operator to be optimized on the target hardware platform, includes: Each parameter value in the parameter value combination is assigned to the corresponding variable in the parameterized template to obtain the instantiated operator template; The instantiated operator template is compiled into executable code adapted to the target hardware platform based on a compilation backend that matches the target hardware platform.

8. The method according to claim 7, characterized in that, The step of assigning each parameter value in the parameter value combination to the corresponding variable of the parameterized template to obtain the instantiated operator template includes: Load the first type of parameters and the second type of parameters in the parameter value combination into the corresponding variables of the parameterization template; Based on the third type of parameter in the parameter value combination, a target dequantization logic branch is selected from the predefined dequantization logic branches of the parameterization template; wherein, the predefined dequantization logic branches include signed type dequantization logic branches and unsigned type dequantization logic branches; Based on the loaded first type of parameters and second type of parameters, and the target dequantization logic branch, an instantiated operator template is obtained.

9. The method according to claim 7, characterized in that, When the target hardware platform is a first type of GPU architecture, the compilation backend is CUDA; when the target hardware platform is a second type of GPU architecture, the compilation backend is OpenCL; when the target hardware platform is a third type of GPU architecture, the compilation backend is Rocm.

10. The method according to claim 1, characterized in that, The process of evaluating the execution performance of each executable code and determining a target parameter combination from the multiple parameter value combinations based on the execution performance includes: The performance of the executable code corresponding to each combination of the parameter values ​​is evaluated sequentially to obtain the corresponding execution performance. Compare the execution performance of the current parameter value combination with the execution performance of the retained parameter value combination. If the execution performance of the current parameter value combination is better, then update the retained parameter value combination with the current parameter value combination. After traversing all parameter value combinations, the final retained parameter value combination is determined as the target parameter combination.

11. The method according to claim 10, characterized in that, The execution performance includes at least one or more of the following: inference latency, computing power utilization, and data transmission efficiency of the operator to be optimized.

12. The method according to claim 1, characterized in that, The step of generating target code for the operator to be optimized to execute on the target hardware platform based on the target parameter combination includes: The target parameter combination is loaded into the parameterized template of the operator to be optimized to obtain the target operator template; The target operator template is compiled based on a compilation backend that matches the target hardware platform to generate the target code; wherein, the target code integrates the weight dequantization operation and the calculation logic of the operator to be optimized.

13. The method according to claim 1, characterized in that, The operator to be optimized includes a matrix-vector multiplication operator; the quantization and dequantization data format is used to characterize the storage format of the quantization weight matrix of the matrix-vector multiplication operator.

14. The method according to claim 2, characterized in that, The quantization weight data type parameter includes signed integer and unsigned integer, the quantization weight data bit width parameter is 2, 3 or 4, and the quantization data packing bit width parameter is 8, 16, 32 or 64.

15. The method according to any one of claims 1 to 14, characterized in that, The first type of parameters includes at least one or more of the following: recursive axis parallelism parameter, non-recursive axis parallelism parameter, loop unrolling dimension parameter, memory cache enable parameter, and vector calculation width parameter.

16. The method according to claim 6, characterized in that, The correlation includes: the parallelism of the operator to be optimized is an integer multiple of the thread bundle or subgroup size parameter; wherein, the parallelism of the operator to be optimized is determined by the product of the recursive axis parallelism parameter and the non-recursive axis parallelism parameter in the first type of parameters.

17. An electronic device, characterized in that, The electronic device includes: a plurality of processors and a memory; the memory is coupled to the plurality of processors, the memory is used to store computer program code, the computer program code including computer instructions, and the plurality of processors call the computer instructions to cause the electronic device to perform the method as described in any one of claims 1 to 16.

18. A computer-readable storage medium, characterized in that, The computer-readable storage medium includes instructions that, when executed on an electronic device, cause the electronic device to perform the method as described in any one of claims 1 to 16.

19. A computer program product, characterized in that, The computer program product includes computer program code that, when run, causes the method as described in any one of claims 1 to 16 to be performed.

Images