A cross-platform SYCL migration and optimization method and system for high-performance matrix multiplication

Abstract

This application discloses a cross-platform SYCL migration and optimization method for high-performance matrix multiplication. The method includes the following steps: parsing and classifying the CUDA code to be migrated, converting the common code and the code specific to each target platform, and generating target SYCL files; performing vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL files; generating a set of valid candidate parameters by combining the shape characteristics of the input matrix and the execution characteristics of the target platform; performing automatic optimization on the set of valid candidate parameters, and outputting the optimal configuration corresponding to the current input shape and the optimized SYCL code. This invention can complete the cross-platform migration of high-performance matrix multiplication code under a unified programming model, and optimizes and automatically tunes the parameters of the migrated SYCL matrix multiplication kernel by combining the characteristics of different platforms and the input shape.

Description

Technical Field

[0001] This invention relates to the fields of heterogeneous computing and high-performance computing technology, specifically to a cross-platform SYCL migration and optimization method and system for high-performance matrix multiplication. Background Technology

[0002] In recent years, artificial intelligence applications, exemplified by large language models, have developed rapidly, with model size and training complexity continuously increasing, leading to a rapid rise in the demand for underlying computing power. Under these computationally intensive loads, GPUs and various dedicated accelerators have become the primary computing platforms. Currently, NVIDIA holds a significant advantage in large model training and high-performance computing applications thanks to its well-developed CUDA software ecosystem. However, facing the ever-expanding demand for computing power, the industry is no longer limited to hardware systems from a single vendor. In the international market, manufacturers such as AMD and Intel have launched dedicated accelerators for matrix computation; domestically, heterogeneous computing chips from manufacturers such as Huawei and Hygon are also developing rapidly. A hardware landscape with multiple vendors and architectures is taking shape. Matrix multiplication is a core computational load in deep learning training and inference, and the implementation of high-performance matrix multiplication often relies on dedicated instructions provided by the underlying hardware platform. Therefore, as the trend of hardware heterogeneity deepens, high-performance matrix multiplication code is becoming increasingly difficult to reuse directly across different platforms, and the cost of repeated development and adaptation is also increasing.

[0003] The current high-performance computing platform landscape has entered a phase of multi-vendor coexistence, but the accumulation of the software ecosystem is still primarily based on the CUDA roadmap. As CUDA becomes the de facto standard in related fields, how to enable CUDA programs to run efficiently on heterogeneous platforms has become a real problem. Once the same application needs to be deployed on multiple types of hardware, developers need to handle code migration, maintenance, and optimization on different platforms separately, which significantly increases development costs. Therefore, in the context of heterogeneous hardware, the core issue facing the software ecosystem is no longer just expanding support for different vendor interfaces, but how to form a unified and stable cross-platform development and maintenance method. SYCL is an open programming model that has gained attention under this need. SYCL can provide a unified programming interface for different types of hardware backends, thereby providing a foundation for cross-platform code migration, maintenance, and reuse.

[0004] Existing cross-platform code migration tools can be mainly divided into three categories. The first category is code migration tools based on static semantic transformation, represented by DPCT. These tools essentially rewrite from source to source, mainly solving syntax transformation problems. They can quickly provide a compileable and runnable basic version, but they still struggle to effectively migrate specialized instructions and platform-specific implementations. The second category is code migration tools based on runtime or binary compatibility, represented by ZLUDA. These tools typically try not to modify the original CUDA source code, but instead directly execute existing programs through a compatibility layer. Their advantage is lower integration cost; however, this method is highly black-box, limiting developers' control over the underlying execution behavior and making it difficult to attribute problems and further optimize in high-performance scenarios. The third category is code migration tools based on large models, represented by UniPar. These tools typically generate candidate code from a large model first, and then iteratively fix it by combining the results of compilation, testing, and environment execution. Their advantage is a high degree of automation; however, there are still significant problems in terms of accuracy, stability, and system complexity. Especially in scenarios such as high-performance matrix multiplication, which rely on low-level dedicated instructions and platform-specific data organization, it is still difficult to directly obtain a usable high-performance equivalent implementation.

[0005] The challenge of migrating high-performance matrix multiplication across platforms lies not in the general CUDA syntax, but in the platform-specific implementations that depend on the hardware. On NVIDIA platforms, developers often need to use low-level instructions such as `mma.sync`, `ldmatrix`, `wgmma`, and `cp.async` to efficiently organize matrix computation and data transfer. AMD platforms, however, do not support the PTX instructions under the CUDA architecture, and their high-performance matrix multiplication implementations rely more on MFMA instructions. The two platforms differ significantly in matrix computation units, thread organization, and shared memory access paths. Therefore, migrating high-performance matrix multiplication across platforms requires more than just syntax rewriting; it also necessitates addressing the platform-specific instructions and data organization issues.

[0006] Even after code migration is complete, the migrated SYCL program often still suffers from performance degradation. Automatic migration tools typically focus more on feature migration and struggle to maintain the high performance of the original implementation simultaneously. For high-performance operators like matrix multiplication, direct migration can easily lead to increased instruction count, higher synchronization overhead, and decreased resource utilization, resulting in a significant performance drop. Therefore, the cross-platform challenge of high-performance matrix multiplication lies not only in achieving code migration but also in maintaining high performance after migration.

[0007] To address the challenges of migrating high-performance matrix multiplication code across platforms under a unified programming model, the difficulty of existing migration tools in handling platform-specific instructions, and the significant performance loss of SYCL programs after migration, this invention proposes a CUDA to SYCL cross-platform migration and optimization method for high-performance matrix multiplication, taking into account the shortcomings of existing technologies. Summary of the Invention

[0008] To address the challenges of migrating high-performance matrix multiplication code across platforms under a unified programming model, the difficulty of existing migration tools in handling platform-specific code, and the significant performance loss after migration, this invention provides a cross-platform SYCL migration and optimization method and system for high-performance matrix multiplication.

[0009] In a first aspect, embodiments of this application provide a cross-platform SYCL migration and optimization method for high-performance matrix multiplication, the method comprising the following steps:

[0010] Code migration steps: Analyze and classify the CUDA code to be migrated, convert the common parts of the code and the code specific to each target platform, and generate the target SYCL file;

[0011] Performance optimization steps: Perform vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file.

[0012] Optionally, in this embodiment of the invention, the cross-platform SYCL migration and optimization method for high-performance matrix multiplication further includes the following steps:

[0013] Candidate parameter generation steps: Combine the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters;

[0014] Automatic tuning steps: Perform automatic tuning on the set of legal candidate parameters and output the optimal configuration and optimized SYCL code corresponding to the current input shape.

[0015] Optionally, in this embodiment of the invention, the code migration step includes:

[0016] Code analysis and classification steps: Analyze and classify the source code to identify general CUDA code and code segments containing hardware-specific instructions;

[0017] General code conversion steps: Perform a unified conversion on general CUDA code to generate corresponding SYCL code snippets;

[0018] Specialized instruction code conversion steps: For code segments containing hardware-specific instructions, different specialized conversion paths are executed according to the target platform; when the target platform is an NVIDIA platform, the relevant code segments are converted into the corresponding Builtin function interfaces; when the target platform is an AMD platform, SYCL code equivalent to the code segments containing hardware-specific instructions is regenerated; the general code conversion results and hardware-specific code conversion results are integrated to generate the target SYCL file.

[0019] Optionally, in this embodiment of the invention, the performance optimization step includes:

[0020] Vectorized write-back optimization steps: Optimize the write-back process of the accumulator segment in the migrated target SYCL file;

[0021] Thread block remapping optimization steps: Adjust the logical mapping relationship of thread blocks in the two-dimensional grid, and modify the global memory read and write address calculation logic accordingly to improve memory access locality.

[0022] Optionally, in this embodiment of the invention, the candidate parameter generation step includes:

[0023] Based on the input shape and target platform, determine the type and value range of the parameters to be tuned;

[0024] Construct the original candidate parameter space based on the candidate parameter type and value range;

[0025] We construct input shape consistency constraints, matrix instruction mapping constraints, and constraints for different target platforms, and apply these constraints to the original candidate parameter space to filter and obtain a set of legal candidate parameters.

[0026] Optionally, in this embodiment of the invention, the automatic optimization step further includes:

[0027] Initialize the number of evaluations, historical records, and the current optimal configuration;

[0028] Based on the current search status, select the configuration to be tested from the set of legal candidate parameters, and perform compilation, running, and performance evaluation on the configuration to be tested;

[0029] Record the configuration and running results based on the evaluation results, and update the current optimal configuration when the current performance is better than the existing best performance.

[0030] After the stopping condition is met, output the optimal configuration and optimized target SYCL code corresponding to the current input shape.

[0031] Optionally, in this embodiment of the invention, the general code conversion step includes:

[0032] Perform abstract syntax tree analysis on general CUDA code, extract key information from the CUDA code, remove CUDA-specific header file dependencies, and supplement the header files required by SYCL;

[0033] Adapt the preset data types and their related interfaces to form a unified type implementation; remove CUDA-specific syntax and add SYCL-related parameters to the kernel function signature; migrate the CUDA host-side interface to SYCL calls;

[0034] The GPU scheduling parameters are reconstructed into kernel function startup forms based on kernel commit primitives; the CUDA index semantics are mapped to the corresponding expressions of SYCL.

[0035] The declarations of dynamic and static shared memory are refactored to form a unified shared memory organization method; and synchronization statements and atomic operations are converted into corresponding SYCL implementations.

[0036] Optionally, in this embodiment of the invention, when the target platform is an NVIDIA platform, the step of converting the relevant code segment into the corresponding Builtin function interface includes:

[0037] The extracted PTX code segment is parsed to identify the code segment parameters; based on the parameters, a direct table lookup and replacement method or a placeholder mapping and parameter recombination method is used.

[0038] If it involves placeholder mapping and parameter reassembly, the instruction template in the inline assembly string is split to extract the inline assembly template information; placeholders, constraints, and corresponding parameter expressions are extracted from the inline assembly template, constraint string, and input / output parameter list; a correspondence between placeholders and actual parameter expressions is established based on the placeholder number, input / output order, and constraint position; the corresponding target Builtin interface is selected based on the type, shape, and parameter layout of the PTX instruction, and the parameter order is reassembled according to the call signature of that interface to obtain the first interface mapping result;

[0039] If the direct table lookup and replacement method is used, for PTX instructions that meet the fixed template replacement conditions, the corresponding Builtin interface is determined directly by looking up the table according to the pre-established mapping relationship, and the second interface mapping result is obtained.

[0040] Based on the mapping results of the first and second interfaces, the corresponding Builtin call code is generated, while maintaining the original input and output semantics of the PTX instruction.

[0041] Optionally, in this embodiment of the invention, when the target platform is an AMD platform, the step of regenerating SYCL code equivalent to the code segment containing hardware-specific instructions includes:

[0042] Extract the relevant AMD platform code segments to be regenerated, and determine the target interface and its basic shape based on the source instruction shape and data type.

[0043] While keeping the thread block-level tile partitioning unchanged, we derive the wavefront scheduling mechanism organization method and thread block partitioning configuration parameters on the AMD platform; and determine the output sub-blocks that each wavefront scheduling mechanism is responsible for and their calculation loop methods in each direction of the input matrix.

[0044] The converted pipeline depth is determined based on the source pipeline depth and AMD hardware limitations; the corresponding AMD SYCL code is generated based on the obtained configuration parameters.

[0045] Optionally, in this embodiment of the invention, the optimization step for the write-back process of the accumulator segment includes:

[0046] Read the accumulator segment to be written back and its corresponding output address to determine the position of the output element that the current thread is responsible for writing back;

[0047] Determine whether the current output fragment layout and data type meet the prerequisites for vectorized write-back. If the current fragment layout can support consecutive write-back of adjacent elements, then perform vectorized write-back.

[0048] Determine the position of the output element corresponding to the current thread in the result matrix, and identify the adjacent elements that can be continuously written back among the output elements handled by the current thread, for subsequent packaging and writing back;

[0049] Determine whether the corresponding write-back addresses are contiguous and meet the preset length alignment. If the conditions are met, perform a vectorized write-back operation and use SYCL to explicitly construct a preset length pack for write-back of adjacent elements that meet the conditions.

[0050] If the conditions are not met, the original scalar write-back method is maintained, and the scalar write-back path is maintained for elements that do not meet the conditions.

[0051] Optionally, in this embodiment of the invention, the thread block remapping step includes:

[0052] Input the physical 2D thread block coordinates, thread block mesh size, and remapping parameters;

[0053] Calculate the height of the complete remapped region based on the thread block grid size and remapped parameters, and determine whether the current thread block is located within the complete remapped region.

[0054] If it is located within the complete remapping region, calculate the vertical group number of the current thread block based on the physical 2D thread block coordinates; calculate the linear offset of the thread block within the group based on the physical 2D thread block coordinates and the group number; calculate the logical column coordinates of the thread block based on the linear offset within the group to determine the column position of the thread block in the remapped 2D grid; calculate the logical row coordinates of the thread block based on the vertical group number and the logical column coordinates to determine the row position of the thread block in the remapped 2D grid.

[0055] If it is not located within the complete remapping region, the original thread block mapping relationship remains unchanged. Based on the obtained logical two-dimensional thread block coordinates, the global memory read / write address calculation logic is modified to change the access order of the thread block to the input data block and the output data block.

[0056] Secondly, embodiments of this application provide a cross-platform SYCL migration and optimization system for high-performance matrix multiplication. Employing the aforementioned cross-platform SYCL migration and optimization method for high-performance matrix multiplication, the system includes the following steps:

[0057] Code migration module: parses and classifies the CUDA code to be migrated, converts the common parts of the code and the code specific to each target platform, and generates the target SYCL file;

[0058] Performance optimization module: Performs vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file;

[0059] Candidate parameter generation module: Combines the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters;

[0060] Automatic tuning module: Performs automatic tuning on the set of legal candidate parameters and outputs the optimal configuration and optimized SYCL code corresponding to the current input shape.

[0061] Thirdly, embodiments of this application provide a computer-readable storage medium storing a computer program that, when executed by a processor, implements the steps of a cross-platform SYCL migration and tuning method for high-performance matrix multiplication.

[0062] Fourthly, embodiments of this application provide an electronic device, including a memory, a processor, and a computer program stored in the memory and executable on the processor. When the processor executes the program, it implements the steps of a cross-platform SYCL migration and tuning method for high-performance matrix multiplication.

[0063] Compared with existing technologies, it has the following outstanding advantages:

[0064] The method and system of this invention have the following advantages: 1) This invention can complete the cross-platform migration of high-performance matrix multiplication code under a unified programming model, thereby improving code reuse capability and reducing repetitive development costs; 2) This invention can adapt to the key implementation paths on NVIDIA and AMD platforms respectively, and handle the issues of dedicated instruction expression and low-level data organization, thereby improving the adaptability of the method to different hardware platforms; 3) After completing the code migration, this invention can optimize and automatically tune the parameters of the migrated SYCL matrix multiplication kernel based on the characteristics of different platforms and input shapes, thereby improving program execution performance; 4) This invention can combine code migration, platform adaptation and performance optimization, providing a more complete technical path for the cross-platform implementation of high-performance matrix multiplication under a unified programming model. Attached Figure Description

[0065] The accompanying drawings, which are included to provide a further understanding of this application and form part of this application, illustrate exemplary embodiments and are used to explain this application, but do not constitute an undue limitation of this application. In the drawings:

[0066] Figure 1 This is a schematic diagram of the cross-platform SYCL migration and optimization method for high-performance matrix multiplication according to the present invention;

[0067] Figure 2 This is a flowchart illustrating the cross-platform SYCL migration and optimization method for high-performance matrix multiplication according to an embodiment of the present invention.

[0068] Figure 3 This is a flowchart illustrating the general code conversion module workflow of this invention.

[0069] Figure 4 This is a flowchart illustrating the PTX code conversion process for the NVIDIA platform according to an embodiment of the present invention.

[0070] Figure 5 This is a flowchart illustrating the code regeneration process for the AMD platform according to an embodiment of the present invention.

[0071] Figure 6 This is a flowchart illustrating the vectorized write-back optimization module workflow of an embodiment of the present invention.

[0072] Figure 7 This is a flowchart illustrating the workflow of the thread block remapping optimization module in an embodiment of the present invention.

[0073] Figure 8 This is a flowchart of the candidate parameter screening process according to an embodiment of the present invention;

[0074] Figure 9 This is a flowchart illustrating the automatic optimization module's workflow in an embodiment of the present invention.

[0075] Figure 10 This is a schematic diagram of the cross-platform SYCL migration and optimization system for high-performance matrix multiplication according to the present invention;

[0076] Figure 11 This is a system architecture diagram of a cross-platform SYCL migration and optimization system for high-performance matrix multiplication according to an embodiment of the present invention;

[0077] Figure 12 This is a schematic diagram of the computer hardware of the present invention. Detailed Implementation

[0078] In this invention, "at least one" means one or more, and "more than one" means two or more. "At least one of the following" or similar expressions refer to any combination of these items, including any combination of a single item or a plurality of items. For example, at least one of a, b, or c can represent: a, b, c, ab, ac, bc, or abc, where a, b, and c can be a single item or multiple items.

[0079] It should also be understood that the term "and / or" in this article 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 existing alone, A and B existing simultaneously, and B existing alone. A and B can be singular or plural. Additionally, the character " / " in this article generally indicates an "or" relationship between the preceding and following related objects, but it can also represent an "and / or" relationship. Please refer to the context for a more accurate understanding.

[0080] It should also be understood that, in various embodiments of the present invention, the order of the above-mentioned processes does not imply the order of execution. The execution order of each process should be determined by its function and internal logic, and should not constitute any limitation on the implementation process of the embodiments of the present invention.

[0081] In the several embodiments provided by this invention, it should be understood that the disclosed devices, apparatuses, and methods can be implemented in other ways. For example, the apparatus embodiments described above are merely illustrative; for instance, the division of 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 device, or some features may be ignored or not executed. Furthermore, the coupling or direct coupling or communication connection shown or discussed may be through some interfaces; the indirect coupling or communication connection between devices or units may be electrical, mechanical, or other forms.

[0082] The units described as separate components may or may not be physically separate. The components shown as units may or may not be physical units; that is, they may be located in one place or distributed across multiple network units. Some or all of the units can be selected to achieve the purpose of this embodiment according to actual needs.

[0083] In addition, the functional units in the various embodiments of the present invention 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.

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

[0085] To make the above-mentioned features and effects of the present invention clearer and easier to understand, specific embodiments are described below in conjunction with the accompanying drawings. This specification discloses one or more embodiments incorporating the features of the present invention. The disclosed embodiments are merely illustrative. The scope of protection of the present invention is not limited to the disclosed embodiments, but is defined by the appended claims.

[0086] This invention aims to provide a method and system for cross-platform migration and optimization of CUDA to SYCL for high-performance matrix multiplication. Its purpose is to address the difficulties in migrating high-performance matrix multiplication code across platforms under a unified programming model, the limitations of existing migration tools in handling platform-dependent code, and insufficient performance of the migrated program. This improves the reusability of high-performance matrix multiplication code across different hardware platforms and enhances the execution performance of the migrated SYCL code.

[0087] This invention provides a CUDA to SYCL cross-platform migration and optimization method for high-performance matrix multiplication, comprising:

[0088] 1) General CUDA code conversion method: The CUDA code to be migrated is parsed and classified, and the general code parts that are independent of the hardware platform are converted into a unified SYCL code representation, thereby providing a foundation for subsequent platform-related code conversion.

[0089] 2) PTX code conversion method for NVIDIA platform: For key code segments in high-performance matrix multiplication code that rely on low-level instructions such as mma.sync, ldmatrix, wgmma, and cp.async, a PTX code conversion path for NVIDIA platform is established so that the semantics of related special instructions can be expressed under the SYCL programming model.

[0090] 3) Code regeneration method for AMD platform: In view of the differences in matrix instructions, thread organization and data read and write paths of AMD platform, the code of key code segments is regenerated according to the target MFMA interface and related configuration parameters, so as to realize the migration of high-performance matrix multiplication code for AMD platform.

[0091] 4) Post-migration SYCL code performance optimization methods: After completing the general code conversion and platform-related code conversion, vectorized write-back optimization and thread block remapping optimization are performed on the target SYCL code after migration to reduce the memory access overhead caused by migration and improve L2 cache locality.

[0092] 5) Platform-related constraint construction and candidate parameter generation method: Construct the original candidate parameter space based on the input shape and target platform, and establish constraints for different target platforms respectively. Filter the original candidate parameter space to obtain a set of legal candidate parameters.

[0093] 6) Automatic tuning method: Perform configuration search, compilation and execution, and performance evaluation on the set of legal candidate parameters. Update the current optimal configuration based on the evaluation results and output the optimal configuration and optimized SYCL code corresponding to the current input shape.

[0094] The methods of this application will be described in detail below with reference to specific embodiments:

[0095] like Figure 1 and Figure 2 As shown, this invention provides a high-performance matrix multiplication cross-platform migration and optimization method based on the above features, including the following steps:

[0096] Code migration step 101: Analyze and classify the CUDA code to be migrated, convert the common parts of the code and the code for each target platform separately, and generate the target SYCL file;

[0097] Performance optimization step 102: Perform vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file.

[0098] Optionally, in this embodiment of the invention, the cross-platform SYCL migration and optimization method for high-performance matrix multiplication further includes the following steps:

[0099] Candidate parameter generation step 103: Combine the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters;

[0100] Automatic tuning step 104: Perform automatic tuning on the set of legal candidate parameters and output the optimal configuration and optimized SYCL code corresponding to the current input shape.

[0101] Optionally, in this embodiment of the invention, code migration step 101 includes:

[0102] Code analysis and classification steps: Analyze and classify the source code to identify general CUDA code and code segments containing hardware-specific instructions;

[0103] General code conversion steps: Perform a unified conversion on general CUDA code to generate corresponding SYCL code snippets;

[0104] Specialized instruction code conversion steps: For code segments containing hardware-specific instructions, different specialized conversion paths are executed according to the target platform; when the target platform is an NVIDIA platform, the relevant code segments are converted into the corresponding Builtin function interfaces; when the target platform is an AMD platform, SYCL code equivalent to the code segments containing hardware-specific instructions is regenerated; the general code conversion results and hardware-specific code conversion results are integrated to generate the target SYCL file.

[0105] Optionally, in this embodiment of the invention, the general code conversion step includes:

[0106] Perform abstract syntax tree analysis on general CUDA code, extract key information from the CUDA code, remove CUDA-specific header file dependencies, and supplement the header files required by SYCL;

[0107] Adapt the preset data types and their related interfaces to form a unified type implementation; remove CUDA-specific syntax and add SYCL-related parameters to the kernel function signature; migrate the CUDA host-side interface to SYCL calls;

[0108] The GPU scheduling parameters are reconstructed into kernel function startup forms based on kernel commit primitives; the CUDA index semantics are mapped to the corresponding expressions of SYCL.

[0109] The declarations of dynamic and static shared memory are refactored to form a unified shared memory organization method; and synchronization statements and atomic operations are converted into corresponding SYCL implementations.

[0110] Optionally, in this embodiment of the invention, when the target platform is an NVIDIA platform, the step of converting the relevant code segment into the corresponding Builtin function interface includes:

[0111] The extracted PTX code segment is parsed to identify the code segment parameters; based on the parameters, a direct table lookup and replacement method or a placeholder mapping and parameter recombination method is used.

[0112] If it involves placeholder mapping and parameter reassembly, the instruction template in the inline assembly string is split to extract the inline assembly template information; placeholders, constraints, and corresponding parameter expressions are extracted from the inline assembly template, constraint string, and input / output parameter list; a correspondence between placeholders and actual parameter expressions is established based on the placeholder number, input / output order, and constraint position; the corresponding target Builtin interface is selected based on the type, shape, and parameter layout of the PTX instruction, and the parameter order is reassembled according to the call signature of that interface to obtain the first interface mapping result;

[0113] If the direct table lookup and replacement method is used, for PTX instructions that meet the fixed template replacement conditions, the corresponding Builtin interface is determined directly by looking up the table according to the pre-established mapping relationship, and the second interface mapping result is obtained.

[0114] Based on the mapping results of the first and second interfaces, the corresponding Builtin call code is generated, while maintaining the original input and output semantics of the PTX instruction.

[0115] Optionally, in this embodiment of the invention, when the target platform is an AMD platform, the step of regenerating SYCL code equivalent to the code segment containing hardware-specific instructions includes:

[0116] Extract the relevant AMD platform code segments to be regenerated, and determine the target interface and its basic shape based on the source instruction shape and data type.

[0117] While keeping the thread block-level tile partitioning unchanged, we derive the wavefront scheduling mechanism organization method and thread block partitioning configuration parameters on the AMD platform; and determine the output sub-blocks responsible for each wavefront scheduling mechanism and their calculation loop methods in each direction of the input matrix.

[0118] The converted pipeline depth is determined based on the source pipeline depth and AMD hardware limitations; the corresponding AMD SYCL code is generated based on the obtained configuration parameters.

[0119] Specifically, such as Figure 2 As shown in the specific embodiment of the present invention, the above code migration steps include:

[0120] Step 1. Perform syntax analysis and structure recognition on the source code to distinguish between general CUDA code and platform-specific code containing hardware-specific instructions, providing a foundation for subsequent path migration.

[0121] Step 2. Rewrite the hardware platform-independent header files, types, function signatures, host-side runtime interfaces, kernel function startup syntax, and thread index semantics to generate code snippets that can be expressed under the SYCL programming model.

[0122] Step 3. Identify the code segment containing hardware-specific instructions. If the target hardware platform is an NVIDIA platform, proceed to step 4; otherwise, proceed to step 5.

[0123] Step 4. Parse the instruction template, operands, and constraint information, and map them to the corresponding Builtin function calls to preserve the original platform-related semantics.

[0124] Step 5. Derive the target interface and parameters based on the source matrix instruction shape, data type, and block-level configuration, and regenerate semantically equivalent SYCL code accordingly.

[0125] Step 6. Integrate the general SYCL code snippet obtained in Step 2 with the platform-related conversion results obtained in Step 4 or Step 5 to generate the target SYCL file required for subsequent performance optimization and automatic tuning.

[0126] The above code migration steps have the following three technical benefits:

[0127] 1) Separation of General-Purpose Code and Platform-Specific Code: This module first parses and categorizes the CUDA matrix multiplication code to be migrated, dividing the source code into general-purpose CUDA code and platform-specific code containing hardware-specific instructions. Then, it performs migration processing separately for each type of code. Compared to migrating all source code using a single path, this approach is more suitable for handling complex code in high-performance matrix multiplication that simultaneously contains both general-purpose syntax and underlying platform-specific implementations.

[0128] 2) PTX-related code on the NVIDIA platform is converted into Builtin interface: For PTX-related code segments on the NVIDIA platform, this module can convert special instruction calls that are difficult for existing migration tools to handle directly into corresponding Builtin function interfaces, thereby preserving platform-related semantics such as matrix calculation, data loading and synchronization control under the SYCL programming model, avoiding only general syntax-level migration.

[0129] 3) Regenerating equivalent SYCL code on the AMD platform: For AMD platform-related code, this module does not directly adopt the instruction mapping method under the NVIDIA platform. Instead, it derives the target MFMA interface and related configuration parameters based on the source matrix instruction shape, data type and thread block level tile partitioning, and regenerates semantically equivalent AMD SYCL code accordingly, thereby adapting to the underlying execution method of the AMD platform.

[0130] Regarding the general code conversion process in step 2, this invention proposes a unified CUDA code conversion process, the specific implementation details of which are as follows: Figure 3 As shown, the detailed steps are as follows:

[0131] Step 201. Perform abstract syntax tree analysis on the general CUDA code to extract key information such as header file dependencies, type information, function signatures, host-side runtime calls, kernel function startup parameters, thread index access, shared memory declarations, and device-side synchronization and atomic operations.

[0132] Step 202. Remove CUDA-specific header file dependencies and add the header files required by SYCL.

[0133] Step 203. Adapt the half data type and its related interfaces to form a unified type implementation.

[0134] Step 204. Remove CUDA-specific syntax such as __global__ and __device__, and add parameters such as nd_item to the kernel function signature.

[0135] Step 205. Migrate host-side interfaces such as cudaMalloc, cudaMemcpy, and cudaFree to SYCL calls based on USM and queue operations.

[0136] Step 206. [The text appears to be incomplete and contains several errors. A more accurate translation would require the full context.<grid, block, shared, stream> >> Refactored into a kernel function startup form based on q.submit and parallel_for.

[0137] Step 207. Map the CUDA index semantics such as threadIdx, blockIdx, blockDim, gridDim to their corresponding SYCL representations.

[0138] Step 208. Refactor the dynamic and static shared memory declarations to form a unified shared memory organization method.

[0139] Step 209. Convert the synchronization statements and atomic operations into their corresponding SYCL implementations.

[0140] For the NVIDIA platform code conversion process in step 4, this invention proposes a PTX code conversion process, the specific implementation details of which are as follows: Figure 4 As shown, the detailed steps are as follows:

[0141] Step 401. Locate the PTX instruction call code given in inline assembly form from the platform-related code to provide input for subsequent template parsing and interface mapping.

[0142] Step 402. Parse the extracted PTX code segment to identify the instruction name, shape information, operand list, and input / output parameter relationship.

[0143] Step 403. When the instruction template is stable, the number of operands is small, and the parameter positions are fixed, the direct table lookup and replacement method is used, and the process proceeds to step 408; otherwise, the subsequent placeholder mapping and parameter reorganization process is entered, and the process proceeds to step 404.

[0144] Step 404. Split the instruction template in the inline assembly string and extract the opcode, placeholder position, and descriptive information related to the instruction shape.

[0145] Step 405. Extract placeholders, constraints, and corresponding parameter expressions from the inline assembly template, constraint strings, and input / output parameter lists.

[0146] Step 406. Establish the correspondence between placeholders and actual parameter expressions based on placeholder numbers, input / output order, and constraint positions.

[0147] Step 407. Based on the type, shape, and parameter layout of the PTX instruction, select the corresponding target Builtin interface and reorder the parameters according to the call signature of that interface.

[0148] Step 408. For PTX instructions that meet the fixed template replacement conditions, directly look up the corresponding Builtin interface according to the pre-established mapping relationship.

[0149] Step 409. Based on the interface mapping result obtained in step 407 or step 408, generate the corresponding Builtin call code while maintaining the original input / output semantics of the PTX instruction.

[0150] Regarding the AMD platform code regeneration process in step 5, this invention proposes an AMD platform code regeneration workflow, the specific implementation details of which are as follows: Figure 5 As shown, the detailed steps are as follows:

[0151] Step 501. Extract the relevant code segments for the AMD platform to be regenerated.

[0152] Step 502. Determine the target MFMA interface and its basic shape based on the source MFMA instruction shape and data type.

[0153] Step 503. While keeping the thread block-level tile partition unchanged, derive the wavefront organization method and configuration parameters such as BM, BN, and BK on the AMD platform.

[0154] Step 504. Determine the output sub-blocks responsible for each wavefront and their calculation loops in the M, N, and K directions.

[0155] Step 505. Determine the converted pipeline depth based on the source pipeline depth and AMD hardware limitations.

[0156] Step 506. Generate the corresponding AMD SYCL code based on the obtained configuration parameters.

[0157] Optionally, in this embodiment of the invention, performance optimization step 102 includes:

[0158] Vectorized write-back optimization steps: Optimize the write-back process of the accumulator segment in the migrated target SYCL file;

[0159] Thread block remapping optimization steps: Adjust the logical mapping relationship of thread blocks in the two-dimensional grid, and modify the global memory read and write address calculation logic accordingly to improve memory access locality.

[0160] Optionally, in this embodiment of the invention, the optimization step for the write-back process of the accumulator segment includes:

[0161] Read the accumulator segment to be written back and its corresponding output address to determine the position of the output element that the current thread is responsible for writing back;

[0162] Determine whether the current output fragment layout and data type meet the prerequisites for vectorized write-back. If the current fragment layout can support consecutive write-back of adjacent elements, then perform vectorized write-back.

[0163] Determine the position of the output element corresponding to the current thread in the result matrix, and identify the adjacent elements that can be continuously written back among the output elements handled by the current thread, for subsequent packaging and writing back;

[0164] Determine whether the corresponding write-back addresses are contiguous and meet the preset length alignment. If the conditions are met, perform a vectorized write-back operation and use SYCL to explicitly construct a preset length pack for write-back of adjacent elements that meet the conditions.

[0165] If the conditions are not met, the original scalar write-back method is maintained, and the scalar write-back path is maintained for elements that do not meet the conditions.

[0166] Optionally, in this embodiment of the invention, the thread block remapping step includes:

[0167] Input the physical 2D thread block coordinates, thread block mesh size, and remapping parameters;

[0168] Calculate the height of the complete remapped region based on the thread block grid size and remapped parameters, and determine whether the current thread block is located within the complete remapped region.

[0169] If it is located within the complete remapping region, calculate the vertical group number of the current thread block based on the physical 2D thread block coordinates; calculate the linear offset of the thread block within the group based on the physical 2D thread block coordinates and the group number; calculate the logical column coordinates of the thread block based on the linear offset within the group to determine the column position of the thread block in the remapped 2D grid; calculate the logical row coordinates of the thread block based on the vertical group number and the logical column coordinates to determine the row position of the thread block in the remapped 2D grid.

[0170] If it is not located within the complete remapping region, the original thread block mapping relationship remains unchanged. Based on the obtained logical two-dimensional thread block coordinates, the global memory read / write address calculation logic is modified to change the access order of the thread block to the input data block and the output data block.

[0171] Specifically, such as Figure 2 As shown in the specific embodiment of the present invention, the above performance optimization steps specifically include:

[0172] Step 7. Perform accumulator fragment vectorization write-back optimization on the target SYCL code.

[0173] Step 8. Perform thread block remapping optimization on the target SYCL code.

[0174] The above performance optimization steps have the following two technical effects:

[0175] 1) Reduced memory access overhead: By performing accumulator fragment vectorized write-back optimization on the migrated target SYCL code, 64-bit pack write-back is performed on adjacent output elements that meet the conditions of output fragment layout, data type, address continuity and alignment, thereby reducing the memory access overhead caused by scalar successive write-back.

[0176] 2) Improved L2 Cache locality: By performing thread block remapping optimization on the migrated target SYCL code, the logical mapping relationship of thread blocks in the two-dimensional grid is reconstructed, and the global memory read and write address calculation logic is modified accordingly, thereby shortening the data reuse distance between blocks and improving L2 Cache locality and overall memory access efficiency.

[0177] Regarding the vectorized write-back optimization process in step 7, this invention proposes a vectorized write-back process for accumulator fragments, the specific implementation details of which are as follows: Figure 6 As shown, the detailed steps are as follows:

[0178] Step 701. Read the accumulator segment to be written back and its corresponding output address to determine the position of the output element that the current thread is responsible for writing back.

[0179] Step 702. Determine whether the current output fragment layout and data type meet the prerequisites for vectorized write-back. If the current fragment layout can support continuous write-back of adjacent elements, proceed to the subsequent vectorized write-back process.

[0180] Step 703. Determine the mapping relationship between threads and output elements based on the output fragment layout, and clarify the position of the output element corresponding to the current thread in the result matrix.

[0181] Step 704. Identify adjacent elements that can be written back consecutively among the output elements handled by the current thread, for subsequent packing and writing back.

[0182] Step 705. Determine whether the corresponding write-back addresses are contiguous and 64-bit aligned. If the conditions are met, perform vectorized write-back and proceed to step 706; otherwise, maintain the original scalar write-back method and proceed to step 707.

[0183] Step 706. Use sycl::vec on adjacent elements that meet the condition.<float, 2> Explicitly construct 64-bit pack writebacks to reduce the overhead of scalar writebacks one by one.

[0184] Step 707. Maintain the scalar write-back path for elements that do not meet the conditions to ensure the correctness and universality of the target SYCL code.

[0185] Regarding the thread block remapping optimization process in step 8, this invention proposes a thread block remapping procedure, the specific implementation details of which are as follows: Figure 7 As shown, the detailed steps are as follows:

[0186] Step 801. Input the physical 2D thread block coordinates, thread block mesh size, and remapping parameter S as the basis for subsequent remapping calculations.

[0187] Step 802. Calculate the height H of the complete remapping region based on the thread block mesh size and remapping parameter S, which is used to determine whether the current thread block is located in the region where remapping needs to be performed.

[0188] Step 803. Determine whether the current thread block is within the complete remapping region; if it is within the complete remapping region, continue to perform subsequent remapping calculations and proceed to step 804; if it is not within the complete remapping region, keep the original thread block mapping relationship unchanged and proceed to step 808.

[0189] Step 804. Calculate the vertical grouping number of the current thread block based on the physical two-dimensional thread block coordinates to determine the grouping position of the thread block in the remapping region.

[0190] Step 805. Calculate the linear offset of the thread block within the group based on the physical two-dimensional thread block coordinates and the group number, for use in subsequent logical coordinate calculations.

[0191] Step 806. Calculate the logical column coordinates of the thread block based on the linear offset within the group, and determine the column position of the thread block in the remapped two-dimensional grid.

[0192] Step 807. Calculate the logical row coordinates of the thread block based on the vertical grouping number and logical column coordinates, and determine the row position of the thread block in the remapped two-dimensional grid.

[0193] Step 808. Based on the obtained logical two-dimensional thread block coordinates, modify the global memory read / write address calculation logic to change the access order of the thread block to the input data block and the output data block.

[0194] Optionally, in this embodiment of the invention, the candidate parameter generation step 103 includes:

[0195] Based on the input shape and target platform, determine the type and value range of the parameters to be tuned;

[0196] Construct the original candidate parameter space based on the candidate parameter type and value range;

[0197] We construct input shape consistency constraints, matrix instruction mapping constraints, and constraints for different target platforms, and apply these constraints to the original candidate parameter space to filter and obtain a set of legal candidate parameters.

[0198] Specifically, such as Figure 2 As shown in the specific embodiment of the present invention, the above-mentioned candidate parameter generation step specifically includes:

[0199] Step 9. Determine the candidate parameter types and their value ranges based on the input shape and target platform.

[0200] Step 10. Based on the candidate parameter type and its value range, combine various parameters to construct the original candidate parameter space.

[0201] Step 11. Construct and apply platform-related constraints to filter the original candidate parameter space and obtain a set of legal candidate parameters.

[0202] The candidate parameter generation step described above has the following two technical effects:

[0203] 1) Controlled execution overhead on NVIDIA platforms: This module can filter out parameter combinations that would significantly increase execution overhead based on the impact of candidate parameter combinations on register usage, thread block concurrency, and fixed instruction overhead, thereby narrowing the search range for subsequent automatic tuning.

[0204] 2) Low risk of additional memory access on AMD platform: This module can filter out parameter combinations that may cause additional memory access overhead for scratch memory based on the impact of candidate parameter combinations on register pressure and pipelining configuration, thereby avoiding high-risk configurations from entering the subsequent automatic tuning process.

[0205] Regarding the platform-related constraint construction and screening process in step 11, this invention proposes a candidate parameter screening process, the specific implementation details of which are as follows: Figure 8 As shown, the detailed steps are as follows:

[0206] Step 1101. Based on the input matrix shapes M, N, K and the block-level partitioning parameters in the candidate parameters, construct input shape consistency constraints and check whether the partitioning results of each parameter combination in the three directions of M, N, and K meet the current matrix multiplication calculation requirements.

[0207] Step 1102. Based on the block-level partitioning parameters and intra-block calculation parameters in the candidate parameters, construct matrix instruction mapping constraints and check whether the thread block partitioning and thread organization method meet the mapping requirements of the target matrix instruction in the output sub-block partitioning and K-direction granularity.

[0208] Step 1103. Determine if the target platform is an NVIDIA platform; if yes, continue building NVIDIA platform-related constraints and proceed to step 1104; if no, continue building AMD platform-related constraints and proceed to step 1107.

[0209] Step 1104. Based on the thread block size and shared memory usage corresponding to the candidate parameters, construct thread block size and shared memory constraints, and check whether each parameter combination meets the NVIDIA platform's limitations on the number of threads and shared memory capacity.

[0210] Step 1105. Construct register constraints and concurrency protection constraints based on the register occupancy of the candidate parameters, and remove parameter combinations that are too high in register pressure or will cause a significant decrease in the concurrency of thread blocks.

[0211] Step 1106. Based on the pipeline depth and K-direction slice settings in the candidate parameters, construct pipeline depth constraints, check whether the pipeline depth matches the current block configuration and number of slices, and then proceed to step 1109.

[0212] Step 1107. Construct wavefront organization and MFMA instruction mapping constraints based on the thread organization method and target MFMA instruction specification in the candidate parameters, and check whether the parameter combination meets the requirements of the AMD platform in terms of execution granularity and matrix instruction mapping.

[0213] Step 1108. Construct platform resource constraints and register overflow risk constraints based on the shared memory usage, register usage, and pipeline depth corresponding to the candidate parameters, remove parameter combinations that may cause register overflow and additional scratch memory access overhead, and then proceed to step 1109.

[0214] Step 1109. Apply the above constraints to the original candidate parameter space, and gradually remove parameter combinations that do not meet the requirements of input shape, matrix instruction mapping and platform resources to obtain a set of legal candidate parameters.

[0215] Optionally, in this embodiment of the invention, the automatic tuning step 104 further includes:

[0216] Initialize the number of evaluations, historical records, and the current optimal configuration;

[0217] Based on the current search status, select the configuration to be tested from the set of legal candidate parameters, and perform compilation, running, and performance evaluation on the configuration to be tested;

[0218] Record the configuration and running results based on the evaluation results, and update the current optimal configuration when the current performance is better than the existing best performance.

[0219] After the stopping condition is met, output the optimal configuration and optimized target SYCL code corresponding to the current input shape.

[0220] Specifically, such as Figure 2As shown in the specific embodiment of the present invention, the above-mentioned automatic optimization step includes:

[0221] Step 12. Perform automatic tuning on the set of legal candidate parameters, and output the optimal configuration and optimized target SYCL code corresponding to the current input shape.

[0222] The above-mentioned automatic tuning steps include the following three technical effects:

[0223] 1) Configuration search based on the set of valid candidates: This module uses the set of valid candidate parameters output by the candidate parameter generation module as the search space, and performs the subsequent automatic tuning process while ensuring the executability of the configuration.

[0224] 2) Model-guided candidate configuration selection: After obtaining basic performance samples in the initial exploration phase, this module selects the next set of test configurations from the remaining legal candidates based on the GP-EI model, thereby improving the targeting of subsequent searches.

[0225] 3) Optimal configuration update based on real performance feedback: This module compiles, runs, and evaluates the performance of the configuration under test, records the configuration and running results, and updates the current optimal configuration when the current performance is better than the existing best performance. Finally, it outputs the optimal configuration corresponding to the current input shape and the optimized target SYCL code.

[0226] Regarding the automatic optimization process in step 12, this invention proposes an automatic optimization process, the specific implementation details of which are as follows: Figure 9 As shown, the detailed steps are as follows:

[0227] Step 1201. Set the evaluation budget, historical record table, and current optimal configuration for the current automatic tuning process, and establish the initial state for subsequent candidate configuration selection and performance comparison.

[0228] Step 1202. If the set of valid candidate parameters is empty, it means that there is no available configuration under the current input shape and target platform conditions, and the automatic tuning process ends directly, proceeding to step 1211; if it is not empty, proceed to the subsequent point selection and evaluation process, proceeding to step 1203.

[0229] Step 1203. If the current number of samples is insufficient to support a stable search model fit, proceed to the initial exploration phase and go to step 1204; if the initial exploration phase has ended, proceed to the candidate configuration selection phase based on the search model and go to step 1205.

[0230] Step 1204. In the initial exploration phase, randomly select the configuration to be tested from the remaining legitimate candidates to obtain basic real performance observation data, and then proceed to step 1206.

[0231] Step 1205. After the initial exploration phase ends, the expected improvement criterion based on Gaussian process is used to select the next set of configurations to be tested from the remaining unevaluated legal candidates in order to improve the efficiency of subsequent searches.

[0232] Step 1206. Apply the current configuration to be tested to the target SYCL code, complete the compilation and execution on the target hardware platform, and obtain the actual performance evaluation results corresponding to the configuration.

[0233] Step 1207. Write the current configuration to be tested and its corresponding performance evaluation results into the history for subsequent search model updates and optimal configuration comparisons.

[0234] Step 1208. If the measured performance of the current configuration is better than the existing best performance, proceed to step 1209 to update the current best configuration; if it is not better than the current best performance, retain the original best configuration and proceed to step 1210.

[0235] Step 1209. Update the current configuration to be tested to the new optimal configuration, and simultaneously update the corresponding best performance record as the basis for subsequent searches and final output.

[0236] Step 1210. If there are still valid candidate configurations that have not been evaluated and the evaluation budget has not been exhausted, return to step 1203 to continue the next round of search; if there are no unevaluated configurations or the budget has been exhausted, proceed to step 1211.

[0237] Step 1211. After the automatic tuning process is completed, output the optimal configuration obtained under the current input shape and target platform conditions, and output the optimized target SYCL code after adopting the configuration.

[0238] As described above, the method of the present invention can be implemented well.

[0239] Compared with the prior art, the present invention has the following outstanding advantages and beneficial effects:

[0240] Compared with the prior art, the method of the present invention has the following beneficial effects:

[0241] 1) Separate processing of general-purpose code and platform-specific code: The CUDA matrix multiplication code to be migrated is first parsed and classified, dividing the source code into general-purpose CUDA code and platform-specific code containing hardware-specific instructions. Then, migration processing is performed separately for each type of code. Compared with migrating all source code using a single path, this approach is more suitable for handling complex code in high-performance matrix multiplication that contains both general-purpose syntax and underlying platform-specific implementations.

[0242] 2) PTX-related code on the NVIDIA platform is converted into Builtin interface: For PTX-related code segments on the NVIDIA platform, this module can convert special instruction calls that are difficult for existing migration tools to handle directly into corresponding Builtin function interfaces, thereby preserving platform-related semantics such as matrix calculation, data loading and synchronization control under the SYCL programming model, avoiding only general syntax-level migration.

[0243] 3) Regenerating equivalent SYCL code on the AMD platform: For AMD platform-related code, this module does not directly adopt the instruction mapping method under the NVIDIA platform. Instead, it derives the target MFMA interface and related configuration parameters based on the source matrix instruction shape, data type and thread block level tile partitioning, and regenerates semantically equivalent AMD SYCL code accordingly, thereby adapting to the underlying execution method of the AMD platform.

[0244] 4) Reduced memory access overhead: By performing accumulator fragment vectorized write-back optimization on the migrated target SYCL code, 64-bit pack write-back is performed on adjacent output elements that meet the conditions of output fragment layout, data type, address continuity and alignment, thereby reducing the memory access overhead caused by scalar successive write-back.

[0245] 5) Improved L2 Cache locality: By performing thread block remapping optimization on the migrated target SYCL code, the logical mapping relationship of thread blocks in the two-dimensional grid is reconstructed, and the global memory read and write address calculation logic is modified accordingly, thereby shortening the data reuse distance between blocks and improving L2 Cache locality and overall memory access efficiency.

[0246] 6) Controlled execution overhead on NVIDIA platforms: It can filter out parameter combinations that would significantly increase execution overhead based on the impact of candidate parameter combinations on register usage, thread block concurrency, and fixed instruction overhead, thereby narrowing the search range for subsequent automatic tuning.

[0247] 7) AMD platform has low risk of additional memory access: It can filter out parameter combinations that may cause additional memory access overhead for scratch memory based on the impact of candidate parameter combinations on register pressure and pipelining configuration, thereby avoiding high-risk configurations from entering the subsequent automatic tuning process.

[0248] 8) Configuration search based on the set of valid candidates: The set of valid candidate parameters output by the candidate parameter generation step is used as the search space, and the subsequent automatic tuning process is executed while ensuring the executability of the configuration.

[0249] 9) Model-guided candidate configuration selection: After obtaining basic performance samples in the initial exploration phase, the next set of test configurations is selected from the remaining legal candidates according to the GP-EI model, thereby improving the targeting of subsequent searches.

[0250] 10) Optimal configuration update based on real performance feedback: By compiling, running and evaluating the performance of the configuration under test, the configuration and running results are recorded. When the current performance is better than the existing best performance, the current optimal configuration is updated. Finally, the optimal configuration corresponding to the current input shape and the optimized target SYCL code are output.

[0251] Secondly, such as Figure 10 and Figure 11 As shown, this application provides a cross-platform SYCL migration and optimization system for high-performance matrix multiplication, employing the aforementioned cross-platform SYCL migration and optimization method for high-performance matrix multiplication. Figure 10 As shown, the system in this invention mainly includes four modules: a code migration module, a performance optimization module, a candidate parameter generation module, and an automatic tuning module.

[0252] Code migration module 201: Analyzes and classifies the CUDA code to be migrated, converts the common parts of the code and the code for each target platform, and generates the target SYCL file;

[0253] Performance optimization module 202: Performs vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file;

[0254] Candidate parameter generation module 203: Combines the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters;

[0255] Automatic tuning module 204: Performs automatic tuning on the set of legal candidate parameters and outputs the optimal configuration and optimized SYCL code corresponding to the current input shape.

[0256] Specifically, such as Figure 11 As shown in the specific embodiment of the present invention, a system for specifically performing the above method is also provided, comprising:

[0257] This invention provides a high-performance matrix multiplication cross-platform migration and optimization method based on the above-mentioned features, and offers a CUDA to SYCL cross-platform migration and optimization system for high-performance matrix multiplication, such as... Figure 11 As shown, the core components include:

[0258] 1) The code migration module is responsible for parsing, classifying, and transforming the CUDA code to be migrated, and generating the target SYCL file. This module includes a code parsing and classification module, a general code conversion module, and a dedicated instruction code conversion module. The code parsing and classification module identifies general CUDA code and code segments containing hardware-specific instructions in the source code; the general code conversion module converts general CUDA code into corresponding SYCL code fragments; and the dedicated instruction code conversion module performs platform-specific conversions on code segments containing hardware-specific instructions based on the target platform.

[0259] The code migration module consists of three parts: a code parsing and classification module, a general code conversion module, and a dedicated instruction code conversion module. This module is responsible for parsing, classifying, and migrating the CUDA matrix multiplication code to be migrated, and generating the target SYCL file. Specifically, this module first parses and classifies the source code, identifying general CUDA code and code segments containing hardware-specific instructions. Then, it performs a unified conversion on the general CUDA code, generating corresponding SYCL code snippets. For code segments containing hardware-specific instructions, it executes different dedicated conversion paths depending on the target platform. When the target platform is an NVIDIA platform, the relevant code segments are converted into corresponding Builtin function interfaces; when the target platform is an AMD platform, SYCL code equivalent to the code segments containing hardware-specific instructions is regenerated. Finally, the general code conversion results and the hardware-specific code conversion results are integrated to generate the target SYCL file.

[0260] 2) The performance optimization module is responsible for further optimizing the target SYCL code generated by the code migration module to reduce execution overhead and improve program performance after migration. This module includes a vectorized write-back optimization module and a thread block remapping optimization module. The vectorized write-back optimization module optimizes the write-back process of the accumulator segment; the thread block remapping optimization module adjusts the logical mapping relationship of thread blocks in the two-dimensional grid and modifies the global memory read / write address calculation logic accordingly, thereby improving memory access locality.

[0261] The performance optimization module consists of two parts: a vectorized write-back optimization module and a thread block remapping optimization module. This module is responsible for performing platform-specific optimizations on the target SYCL code generated by the code migration module to reduce the execution overhead of the migrated code and improve program performance. Based on the target SYCL code, this module also performs accumulator fragment vectorized write-back optimization and thread block remapping optimization to obtain the optimized SYCL code.

[0262] 3) The candidate parameter generation module is responsible for combining the input shape and the target platform to generate a set of legal candidate parameters required for subsequent automatic optimization. This module determines the type and value range of candidate parameters based on the shape characteristics of the input matrix in the M, N, and K directions and the execution characteristics of the target platform, constructs the original candidate parameter space, and filters it through platform-related constraints to output a set of legal candidate parameters.

[0263] Candidate Parameter Generation Module. This module is responsible for determining the types and value ranges of candidate parameters based on the input matrix shapes M, N, and K, and the target platform. It constructs an initial candidate parameter space and then filters for a set of valid candidate parameters using platform-related constraints. Specifically, this module first determines the types and value ranges of the parameters to be tuned based on the input shape and target platform. Then, it constructs an initial candidate parameter space based on the candidate parameter types and value ranges. Finally, it constructs input shape consistency constraints, matrix instruction mapping constraints, and constraints specific to different target platforms, and applies these constraints to the initial candidate parameter space to obtain a set of valid candidate parameters.

[0264] 4) The automatic tuning module is responsible for automatically tuning the set of valid candidate parameters output by the candidate parameter generation module, and outputting the optimal configuration corresponding to the current input shape and the optimized SYCL code. During the tuning process, this module uses a Bayesian optimization model to select the configuration to be tested, updates the search process based on the evaluation records, and maintains the current optimal configuration through an optimal configuration table.

[0265] Automatic tuning module. This module is responsible for automatically tuning the set of valid candidate parameters output by the candidate parameter generation module, and outputting the optimal configuration and optimized target SYCL code for the current input shape. Specifically, this module first initializes the number of evaluations, historical records, and the current optimal configuration; then, it selects a configuration to be tested from the set of valid candidate parameters based on the current search state, and performs compilation, running, and performance evaluation on that configuration; next, it records the current configuration and running results based on the evaluation results, and updates the current optimal configuration when the current performance is better than the existing best performance; finally, after the stopping condition is met, it outputs the optimal configuration and optimized target SYCL code for the current input shape. The optimal configuration generated by this module can be used for performance optimization of the target SYCL code under the current input shape.

[0266] Through the above methods and systems, this invention can complete the cross-platform migration of high-performance matrix multiplication code under a unified programming model, and optimize and tune the migrated SYCL kernel by combining the characteristics of different platforms and input shapes, thereby improving code reusability and program execution performance.

[0267] Thirdly, embodiments of this application provide a computer-readable storage medium storing a computer program thereon, which, when executed by a processor, implements the steps of the cross-platform SYCL migration and tuning method for high-performance matrix multiplication.

[0268] Fourthly, embodiments of this application provide an electronic device, including a memory, a processor, and a computer program stored in the memory and executable on the processor. When the processor executes the program, it implements the steps of the cross-platform SYCL migration and tuning method for high-performance matrix multiplication as described above.

[0269] In addition, combined Figure 1 The cross-platform SYCL migration and optimization method for high-performance matrix multiplication described in this application embodiment can be implemented by electronic devices, such as computer devices. Figure 12 This is a schematic diagram of the hardware structure of a computer device according to an embodiment of this application.

[0270] In some embodiments, the computer device may further include a communication interface 83 and a bus 80. For example, Figure 12 As shown, the processor 81, memory 82, and communication interface 83 are connected through bus 80 and complete communication with each other.

[0271] Specifically, the processor 81 may include a central processing unit (CPU), an application-specific integrated circuit (ASIC), or one or more integrated circuits that can be configured to implement the embodiments of this application.

[0272] The memory 82 can be used to store or cache various data files that need to be processed and / or communicated, as well as possible computer program instructions executed by the processor 81.

[0273] The processor 81 reads and executes the computer program instructions stored in the memory 82 to implement any of the cross-platform SYCL migration and optimization methods for high-performance matrix multiplication in the above embodiments.

[0274] The technical features of the above embodiments can be combined in any way. For the sake of brevity, not all possible combinations of the technical features in the above embodiments are described. However, as long as there is no contradiction in the combination of these technical features, they should be considered to be within the scope of this specification.

[0275] The embodiments described above are merely illustrative of several implementation methods of this application, and while the descriptions are relatively specific and detailed, they should not be construed as limiting the scope of the invention patent. It should be noted that those skilled in the art can make various modifications and improvements without departing from the concept of this application, and these all fall within the protection scope of this application. Therefore, the protection scope of this patent application should be determined by the appended claims.

Claims

1. A cross-platform SYCL migration and optimization method for high-performance matrix multiplication, characterized in that, The method includes the following steps: Code migration steps: Analyze and classify the code to be migrated, convert the common parts of the code and the code specific to each target platform, and generate the target SYCL file; Performance optimization steps: Perform vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file.

2. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 1, characterized in that, The method further includes the following steps: Candidate parameter generation steps: Combine the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters; Automatic tuning steps: Perform automatic tuning on the set of legal candidate parameters, and output the optimal configuration and optimized SYCL code corresponding to the current input shape.

3. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 1 or 2, characterized in that, The code migration steps include: Code analysis and classification steps: Analyze and classify the source code to identify general CUDA code and code segments containing hardware-specific instructions; General code conversion steps: Perform a unified conversion on general CUDA code to generate corresponding SYCL code snippets; Specialized instruction code conversion steps: For code segments containing hardware-specific instructions, different specialized conversion paths are executed according to the target platform; when the target platform is an NVIDIA platform, the relevant code segments are converted into the corresponding Builtin function interfaces; when the target platform is an AMD platform, SYCL code equivalent to the code segments containing hardware-specific instructions is regenerated; the general code conversion results and hardware-specific code conversion results are integrated to generate the target SYCL file.

4. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication according to claim 1 or 2, characterized in that, The performance optimization steps include: Vectorized write-back optimization steps: Optimize the write-back process of the accumulator segment for the migrated target SYCL file; Thread block remapping optimization steps: Adjust the logical mapping relationship of thread blocks in the two-dimensional grid, and modify the global memory read and write address calculation logic accordingly to improve memory access locality.

5. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 2, characterized in that, The candidate parameter generation step includes: Based on the input shape and target platform, determine the type and value range of the parameters to be tuned; Construct the original candidate parameter space based on the candidate parameter type and value range; Construct input shape consistency constraints, matrix instruction mapping constraints, and constraints for different target platforms, and apply these constraints to the original candidate parameter space to filter and obtain a set of legal candidate parameters.

6. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 2, characterized in that, The automatic tuning step further includes: Initialize the number of evaluations, historical records, and the current optimal configuration; Based on the current search status, a configuration to be tested is selected from the set of legal candidate parameters, and the configuration to be tested is compiled, run, and its performance is evaluated. Record the configuration and running results based on the evaluation results, and update the current optimal configuration when the current performance is better than the existing best performance. After the stopping condition is met, output the optimal configuration and optimized target SYCL code corresponding to the current input shape.

7. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 3, characterized in that, The general code conversion steps include: Perform abstract syntax tree analysis on general CUDA code, extract key information from the CUDA code, remove CUDA-specific header file dependencies, and supplement the header files required by SYCL; Adapt the preset data types and their related interfaces to form a unified type implementation; remove CUDA-specific syntax and add SYCL-related parameters to the kernel function signature; migrate the CUDA host-side interface to SYCL calls; The GPU scheduling parameters are reconstructed into kernel function startup forms based on kernel commit primitives; the CUDA index semantics are mapped to the corresponding expressions of SYCL. The declarations of dynamic and static shared memory are refactored to form a unified shared memory organization method; and synchronization statements and atomic operations are converted into corresponding SYCL implementations.

8. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 3, characterized in that, When the target platform is an NVIDIA platform, the step of converting the relevant code segments into the corresponding Builtin function interface includes: The extracted PTX code segment is parsed to identify the code segment parameters; based on the parameters, a direct table lookup and replacement method or a placeholder mapping and parameter recombination method is used. If it involves placeholder mapping and parameter reassembly, the instruction template in the inline assembly string is split to extract the inline assembly template information; placeholders, constraints, and corresponding parameter expressions are extracted from the inline assembly template, constraint string, and input / output parameter list; a correspondence between placeholders and actual parameter expressions is established based on the placeholder number, input / output order, and constraint position; the corresponding target Builtin interface is selected based on the type, shape, and parameter layout of the PTX instruction, and the parameter order is reassembled according to the call signature of that interface to obtain the first interface mapping result; If the direct table lookup and replacement method is used, for PTX instructions that meet the fixed template replacement conditions, the corresponding Builtin interface is determined directly by looking up the table according to the pre-established mapping relationship, and the second interface mapping result is obtained. Based on the mapping results of the first and second interfaces, the corresponding Builtin call code is generated, while maintaining the original input and output semantics of the PTX instruction.

9. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in claim 3, characterized in that, When the target platform is an AMD platform, the step of regenerating SYCL code equivalent to the code segment containing hardware-specific instructions includes: Extract the relevant AMD platform code segments to be regenerated, and determine the target interface and its basic shape based on the source instruction shape and data type. While keeping the thread block-level tile partitioning unchanged, we derive the wavefront scheduling mechanism organization method and thread block partitioning configuration parameters on the AMD platform; and determine the output sub-blocks responsible for each wavefront scheduling mechanism and their calculation loop methods in each direction of the input matrix. The converted pipeline depth is determined based on the source pipeline depth and AMD hardware limitations; the corresponding AMD SYCL code is generated based on the obtained configuration parameters.

10. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication according to claim 4, characterized in that, The optimization steps for the write-back process of the accumulator segment include: Read the accumulator segment to be written back and its corresponding output address to determine the position of the output element that the current thread is responsible for writing back; Determine whether the current output fragment layout and data type meet the prerequisites for vectorized write-back. If the current fragment layout can support consecutive write-back of adjacent elements, then perform vectorized write-back. Determine the position of the output element corresponding to the current thread in the result matrix, and identify the adjacent elements that can be continuously written back among the output elements handled by the current thread, for subsequent packaging and writing back; Determine whether the corresponding write-back addresses are continuous and meet the preset length alignment. If the conditions are met, perform a vectorized write-back operation and use SYCL to explicitly construct the preset length for write-back of adjacent elements that meet the conditions. If the conditions are not met, the original scalar write-back method is maintained, and the scalar write-back path is maintained for elements that do not meet the conditions.

11. The cross-platform SYCL migration and optimization method for high-performance matrix multiplication according to claim 4, characterized in that, The thread block remapping step includes: Input the physical 2D thread block coordinates, thread block mesh size, and remapping parameters; Calculate the height of the complete remapping region based on the thread block mesh size and remapping parameters, and determine whether the current thread block is located within the complete remapping region; If it is located within the complete remapping region, the vertical group number of the current thread block is calculated based on the physical two-dimensional thread block coordinates; the linear offset of the thread block within the group is calculated based on the physical two-dimensional thread block coordinates and the group number; the logical column coordinates of the thread block are calculated based on the linear offset within the group to determine the column position of the thread block in the remapped two-dimensional grid; the logical row coordinates of the thread block are calculated based on the vertical group number and the logical column coordinates to determine the row position of the thread block in the remapped two-dimensional grid. If it is not located within the complete remapping region, the original thread block mapping relationship remains unchanged. Based on the obtained logical two-dimensional thread block coordinates, the global memory read / write address calculation logic is modified to change the access order of the thread block to the input data block and the output data block.

12. A cross-platform SYCL migration and optimization system for high-performance matrix multiplication, employing the cross-platform SYCL migration and optimization method for high-performance matrix multiplication as described in any one of claims 1-11, characterized in that, The system includes the following steps: Code migration module: parses and classifies the code to be migrated, converts the common parts of the code and the code specific to each target platform, and generates the target SYCL file; Performance optimization module: Performs vectorized write-back optimization and thread block remapping optimization on the migrated target SYCL file; Candidate parameter generation module: Combines the shape features of the input matrix and the execution characteristics of the target platform to generate a set of legal candidate parameters; Automatic tuning module: Performs automatic tuning on the set of legal candidate parameters and outputs the optimal configuration and optimized SYCL code corresponding to the current input shape.

13. A computer-readable storage medium having a computer program stored thereon, characterized in that, When executed by the processor, the program implements the steps of the cross-platform SYCL migration and tuning method for high-performance matrix multiplication as described in any one of claims 1-11.

14. An electronic device comprising a memory, a processor, and a computer program stored in the memory and executable on the processor, characterized in that, When the processor executes the program, it implements the steps of the cross-platform SYCL migration and tuning method for high-performance matrix multiplication as described in any one of claims 1-11.

Images